6

我正在为一个专门的矩阵函数编写一段 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;
      }

干杯!

4

1 回答 1

6

优化 #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

于 2012-02-23T16:24:42.450 回答