这个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;
      }

干杯!

Answer 1

优化#1:制作向量__local.

我的第一次传球在性能方面取得了不错的进步.我注意到每个向量[k]总共读了D次,所以我把它复制到__local.这是唯一可能的,因为D足够小以允许这样做.你上面的内核遭受了糟糕的ALU:5870和6970 gpus上的取值比为0.08.即使是较慢的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:5870和6970 gpus的取值比为0.20.平均时间从同一张卡上的1513 - > 1034和1261 - > 861变化.低端gpus现在由ALU绑定而不是fetch.(比例大于4:1)

Opimization#2:使用整个工作组计算每个结果[y].

你必须要做这个ID D更大(100k +).我们的想法是通过使用工作组一次计算结果的单个元素来获得最佳的内存访问模式.我在这里将ls(本地大小)定义为64,因为它适用于我的硬件以及大多数供应商.除非您更改该定义,否则从主机端使用的工作组大小必须为64.需要定义它来创建sum [ls]存储为__local,我不喜欢将变量大小的__local vars传递到我的内核中.

结果:5870 ALU:fetch = 0.59:1,avg = 708.6970 ALU:fetch = 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);
    }
}

编辑:APP profiler = AMD APP KernelAnalyzer