Можно ли оптимизировать этот код OpenCL? - PullRequest
Новая
       15

Можно ли оптимизировать этот код OpenCL?

5 голосов
/ 23 февраля 2012

Я работаю над фрагментом кода OpencL для специализированной матричной функции: для Dx1 вектора v, двух DxD матриц A и B и константы c, возврат 1xD вектор r где r[i] = c * sum_over_j (v[j] * A[i][j] * B[i][j])

Ниже приведено то, что у меня есть, но оно работает очень медленно. Версия без суммирования, которая возвращает матрицу DxD, примерно в десять раз быстрее. Он вызывается из PyOpenCL, если это что-то меняет.

Что-то сделано не так? Может ли это быть оптимизировано? 

#define D 1000
...

   __kernel void element_mult(
      __global float *result,
      __global const float *vector,
      __global const float *matrix,
      __global const float *matrix2,
        const float factor)
      {
         int y = get_global_id(1);
         float sum = 0;
         for(int k = 0; k < D; k++)
         {
            sum += vector[k] * matrix[(y*D) + k]
            * matrix2[(y*D) + k ];
         }
         result[y] = sum * factor;
      }

Ура!

1 Ответ

6 голосов
/ 23 февраля 2012

Оптимизация # 1: сделать вектор __local.

Мой первый проход при этом получил приличное улучшение производительности. Я заметил, что каждый вектор [k] читается в общей сложности D раз, поэтому я скопировал его в __local. Это возможно только потому, что D достаточно мал, чтобы позволить это. Ядро, как оно есть у вас, страдает от ужасного соотношения ALU: выборка 0,08 на 5870 и 6970 gpus. Даже более медленный процессор все еще ожидает доступа к памяти. 

   #define D 1000
    __kernel void element_mult(
    __global float *result,
    __global const float *vector,
    __global const float *matrix,
    __global const float *matrix2,
    const float factor)
    {
        int y = get_global_id(0);
        float sum = 0;

        __local float vectCopy[D];
        int ls = get_local_size(0);
        int lid = get_local_id(0);
        for(int i=0;i<D;i+=ls){
            vectCopy[i+lid] = vector[i+lid];
        }
        mem_fence(CLK_LOCAL_MEM_FENCE);

        for(int k = 0; k < D; k++)
        {
            sum += vectCopy[k] * matrix[(y*D) + k] * matrix2[(y*D) + k ];
        }
        result[y] = sum * factor;
    }

С этим изменением профилировщик APP показывает новое соотношение ALU: выборка 0,20 для 5870 и 6970 графических процессоров. Среднее время изменилось с 1513 -> 1034 и 1261 -> 861 на одних и тех же картах. Младшие gpus теперь связаны ALU вместо выборки. (соотношение больше 4: 1)

Оптимизация # 2: вычислите каждый результат [y], используя всю рабочую группу.

Вы должны были бы сделать это id D были намного больше (100k +). Идея состоит в том, чтобы получить наилучший шаблон доступа к памяти, используя рабочую группу для вычисления одного элемента результата за раз. Я определил, что ls (локальный размер) здесь равен 64, потому что он работает на моем оборудовании, а также на большинстве поставщиков. Размер рабочей группы, который вы используете со стороны хоста, должен быть 64, если вы не измените это определение. Его нужно определить, чтобы создать хранилище sum [ls] как __local, и мне не нравится передавать __local переменные переменного размера в мои ядра.

Результаты: 5870 ALU: выборка = 0,59: 1, средняя = 708. 6970 АЛУ: выборка = 0,72, средняя = 590. Согласно профайлеру приложения, это примерно в два раза быстрее, чем ваш первоначальный список. 

#define D 1000
#define ls 64
__kernel void element_mult(
__global float *result,
__global const float *vector,
__global const float *matrix,
__global const float *matrix2,
const float factor)
{
    __local float vectCopy[D];
    int lid = get_local_id(0);
    for(int i=0;i<D;i+=ls){
        vectCopy[i+lid] = vector[i+lid];
    }
    mem_fence(CLK_LOCAL_MEM_FENCE);

    int ng = get_num_groups(0);
    int gid = get_group_id(0);
    int y, k;
    __local float sum[ls];
    for(y = gid; y < D; y+=ng){
        for(k = lid; k < D; k+=ls)
        {
            sum[lid] += vectCopy[k] * matrix[(y*D) + k] * matrix2[(y*D) + k ];
        }
        if(lid==0){
            result[y] = sum[0];
            for(k=1;k<ls;k++){
                result[y] += sum[k];
            }
            result[y] *= factor;
        }
        mem_fence(CLK_LOCAL_MEM_FENCE);
    }
}

РЕДАКТИРОВАТЬ: APP profiler = AMD APP KernelAnalyzer

...