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

Question

Я работаю на куске 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; 
     } 

Cheers!

Answer 1

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

Мой первый пропуск на этом получил достойное улучшение в производительности. Я заметил, что каждый вектор [k] читается в общей сложности D раз, поэтому я скопировал его в __local. Это возможно только потому, что D достаточно мал, чтобы это можно было сделать. Ядро, как у вас, выше, страдает от ужасного соотношения ALU: fetch 0,08 как на 5870, так и на 6970 gpus. Даже медленный 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; 
    } 

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

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

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

результаты: 5870 ALU: fetch = 0.59: 1, avg = 708. 6970 ALU: выборка = 0,72, avg = 590. Согласно профилировщику APP, это примерно в два раза быстрее, чем ваш первоначальный листинг.

#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); 
    } 
} 

EDIT: APP профайлер = AMD APP KernelAnalyzer