的OpenCL內核實現一個簡單的數學公式

error formula

誤差函數時要考慮的最佳實踐？

A，B和C是三維浮法陣列，\ delta是克羅內克三角洲。（N，M）=（2,7）或（N，M）=（2,23）的典型值。

天真的實現（下面給出）比CPU版本慢幾個數量級。

感謝，

__kernel void cl_bilinear_alg(
          __global float * A, 
          __global float * B, 
          __global float * C, 
          __global const int M, 
          __global const int N, 
          __global float * R) 
{ 
    int index = get_global_id(0); 
    int N2 = N * N; 
    int mat_offset = index * N2 * M; 
    float s1, s2, err = 0.0f; 

    for (int i = 0; i < N; ++i) 
    { 
     for (int j = 0; j < N; ++j) 
     { 
      for (int k = 0; k < N; ++k) 
      { 
       for (int l = 0; l < N; ++l) 
       { 
        for (int m = 0; m < N; ++m) 
        { 
         for (int n = 0; n < N; ++n) 
         { 
          s1 = (n == i) * (j == k) * (l == m); 
          s2 = 0; 

          for (int r = 0; r < M; ++r) 
          { 
           s2 += A[mat_offset + r * N2 + i * N + j] * 
             B[mat_offset + r * N2 + k * N + l] * 
             C[mat_offset + r * N2 + m * N + n]; 
          } 
          err += (s2 - s1) * (s2 - s1); 
         } 
        } 
       } 
      } 
     } 
    } 
    R[index] = err; 
}

UPDATE

主要的目標是一塊GeForce GTX 570，雖然這可能在未來改變。

UPDATE2

明確地矢量化的代碼，移動位本地內存，展開一些循環並通過預先計算克羅內克產品的內核後，代碼如下：

__kernel void cl_bilinear_alg(__global const float * A, 
           __global const float * B, 
           __global const float * C, 
           __global const int N, 
           __global const int M, 
           __global const float * kron, 
           __global float * R) 
{ 
    __private int index = get_global_id(0); 
    __private int cM = ceil(M/4.0f); 
    __private int N2 = N*N; 
    __private int N4 = N2*N2; 
    __private int mat_offset = index * N2 * M; 
    __private float s1, s2, err = 0; 
    __private float4 vzero = (float4) (0.0f, 0.0f, 0.0f, 0.0f); 
    __local float4 va[54], vb[54], vc[54]; 

for (int ij = 0, k = 0; ij < N2; ++ij) 
{ 
    int r = 0; 
    for (; r < M/4; r += 4, ++k) 
    { 
     int idx0 = mat_offset + N2 * r + ij; 
     int idx1 = mat_offset + N2 * (r + 1) + ij; 
     int idx2 = mat_offset + N2 * (r + 2) + ij; 
     int idx3 = mat_offset + N2 * (r + 3) + ij; 
     va[k] = (float4) (A[idx0], A[idx1], A[idx2], A[idx3]); 
     vb[k] = (float4) (B[idx0], B[idx1], B[idx2], B[idx3]); 
     vc[k] = (float4) (C[idx0], C[idx1], C[idx2], C[idx3]); 
    } 

    if (M % 4) 
    { 
     float buffa[4] = {0}, buffb[4] = {0}, buffc[4] = {0}; 
     for (; r < M; ++r) 
     { 
      int idx = mat_offset + N2 * r + ij; 
      buffa[r % 4] = A[idx]; 
      buffb[r % 4] = B[idx]; 
      buffc[r % 4] = C[idx]; 
     } 
     va[k] = vload4(0, buffa); 
     vb[k] = vload4(0, buffb); 
     vc[k++] = vload4(0, buffc); 
    } 
}  

for (int ij = 0; ij < N2; ++ij) 
{ 
    for (int kl = 0; kl < N2; ++kl) 
    { 
     for (int mn = 0; mn < N2; ++mn) 
     { 
      s1 = kron[ij * N4 + kl * N2 + mn]; 
      s2 = 0; 
      for (int r = 0; r < cM; ++r) 
       s2 += dot(va[cM * ij + r], mad(vb[cM * kl + r], vc[cM * mn + r], vzero)); 

      //the most expensive line 
      err += (s2 - s1) * (s2 - s1); 
     } 
    } 
} 

R[index] = err; 
}

通過將這些相比於天真的實施，觀察到4倍的速度增加。此外，顯示所有最昂貴的行是錯誤更新，即

err += (s2 - s1) * (s2 - s1);

有什麼建議嗎？

來源

2012-06-01 user92382

設備的架構是什麼？如果是帶矢量寄存器的intel，則可以使用SIMD寄存器獲得更高的速度，速度可提高5倍 – ardiyu07

「天真」版本作爲單個工作項目運行？如果是這樣，你正在浪費GPU的總運算峯值容量的99.8％，所以它應該不會太過出人意料，它的速度很慢。 – talonmies

@talonmies：不，每個工作項目都在工作一個單獨的問題。 – user92382

通常你會想打破一些這些循環起來的......很多...... - 外循環成爲多個劈了過來workgroup s，它自己的計算單元上運行（大約有16個計算單位每個GPU都不是很多） - 接下來的幾個循環將分割到每個工作組中的不同線程上

如果您嘗試同時運行所有計算，它們都會嘗試將數據加載到內存中同一時間，這將只是可怕地摔打。 GPU具有非常有限的內存。當然，全球內存聽起來足夠大，幾千兆字節，但全球GPU內存很慢。您希望將數據導入本地內存（每個計算單位），數據量爲32-64KB，並不比這更多。

通常你會希望以某種方式將你的任務成非常小的任務，並做到以下幾點，對每一個工作組：

負荷從全局內存的內存塊到本地內存
- 的線程的整個工作組經紗可以參與做拷貝，使用coallesced訪問
做這個工作記憶，就像做了一些資金，等等
結果寫回全局內存
然後，可以重複一下，或者乾脆退出，讓其他工作組來處理工作

在CPU的其他位，數學運算往往是這是一個主要的瓶頸，但在GPU上，一般來說，內核大部分是無用的，同時等待數據逐漸從全局內存中獲取。無論你如何優化這個過程，防止衝突的需求等等，都會使內核顯着更快。

來源

2015-02-19 04:16:49

的OpenCL內核實現一個簡單的數學公式

回答

相關問題