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代碼是否可以優化？

回答

相關問題