OpenCL代码是否可以优化？

我工作一段的OpenCL代码为一个专门的矩阵功能：为Dx1矢量v，二DxD矩阵A和B和恒定c，返回1xD矢量r其中r[i] = c * sum_over_j (v[j] * A[i][j] * B[i][j])OpenCL代码是否可以优化？

下面是我有这么远，但它运行速度非常缓慢。没有求和的版本返回DxD矩阵的速度大约快10倍。如果这有什么不同，它会从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; 
     }

干杯！

来源

2012-02-23 trolle3000

你肯定，你不是，那y的计算* d是由编译器抬出k个循环？而且常见的子表达式（y * D）+ k只在每次迭代中计算一次？ – 2012-02-23 13:20:32

你是否在NVIDIA GPU上运行这个机会？ – talonmies 2012-02-23 13:40:58

@talonmies，我不能确定。计算不是在我的电脑本地完成的;基本上它只需要成为OpenCL。 – trolle3000 2012-02-23 14:06:58

优化＃1：使向量__local。

我第一次在这方面的表现得到了不错的提升。我注意到每个矢量[k]总共读取了D次，所以我将它复制到__local。这是唯一可能的，因为D足够小以允许这样做。上面的内核受到5870和6970 gpus上可怕的ALU：取比为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：取0.20比率为5870和6970的GPU。平均时间从相同卡上的1513 - > 1034和1261 - > 861改变。低端gpus现在被ALU绑定而不是获取。（大于4：1的比例）

Opimization＃2：使用整个工作组计算每个结果[y]。

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

结果：5870 ALU：取= 0.59：1，平均= 708。 6970 ALU：取= 0.72，平均= 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探查= AMD APP KernelAnalyzer

来源

2012-02-23 16:24:42 mfa

OpenCL代码是否可以优化？

回答

相关问题