这个 OpenCL 代码可以优化吗？答案

【问题标题】：Can this OpenCL code be optimized?这个 OpenCL 代码可以优化吗？
【发布时间】：2012-03-13 21:40:49
【问题描述】：

我正在为一个专门的矩阵函数编写一段 OpencL 代码：对于一个 Dx1 向量 v，两个 DxD 矩阵 A 和 B 和一个常量 c，返回 @987654327 @矢量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;
      }

干杯！

【问题讨论】：

您确定 yD 的计算是由您的编译器从 k 循环中提取出来的吗？并且公共子表达式 (yD)+k 在每次迭代中只计算一次？
您是否在 NVIDIA GPU 上运行此程序？
@talonmies，我不能确定。计算不是在我的计算机上本地完成的；基本上它只需要是 OpenCL。
@trolle3000：您的代码使用的内存访问模式对于 NVIDIA GPU 上的性能将非常不利，如果这恰好是目标的话。这个计算实际上是两个矩阵与向量的 Hadamard 积的点积？
@talonmies，没错。您能否详细说明内存访问模式？

标签： opencl gpgpu pyopencl

【解决方案1】：

优化 #1：将向量设为 __local。

我在这方面的第一次通过在性能上得到了不错的改进。我注意到每个 vector[k] 总共被读取了 D 次，所以我将它复制到了一个 __local。这仅是可能的，因为 D 小到足以允许这样做。上面的内核在 5870 和 6970 gpus 上的 ALU:fetch 比率都为 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 分析器显示 5870 和 6970 gpus 的新 ALU:fetch 比率为 0.20。在同一张卡片上，平均时间从 1513-->1034 和 1261-->861 变化。低端 GPU 现在由 ALU 绑定而不是 fetch。（大于 4:1 的比例）

优化#2：使用整个工作组计算每个结果[y]。

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

结果：5870 ALU:fetch=0.59:1，avg=708。 6970 ALU：获取=0.72，平均=590。根据 APP profiler，这大约是您原始列表的两倍。

#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 分析器 = AMD APP KernelAnalyzer

【讨论】：