CUDA問題-類似一維卷積函數的實作問題

最近在實作一些CUDA的kernel,將遇到的問題簡化成以下的例子,目前知道有兩種可能的平行化方法,想請問一下大家哪一種方法在實務上比較好,如果有其他更近一步的方法,也很歡迎可以放上來討論。
簡化後的數學問題如下:

A_{i} = A_{i} + \sum^{N-1}_{j=0}(i-j)*B_{j}\quad\forall i=0,\dots,(N-1)

這裡假設N不大,如:N=64。

// 非平行化版本
value_type A[N];
value_type B[N];
for(int i=0; i<N;++i){
  for(int j=0; j<N;++j){
    A[i] += (i - j) * B[j];
  }
}

平行化版本一,平行外迴圈,並使用share memory來減少對global memory(B[N])的access

// 平行化版本
value_type A[N];
value_type B[N];
tid = threadIdx.x;
if(tid < N){
  __shared__  value_type shared_B[N];
  shared_B[tid] = B[tid];
  for(int j=0; j<N;++j){
    A[tid] += (tid - j) * B[j];
  }
}

平行化版本二,平行內迴圈,但需要使用reuction

// 平行化版本二
value_type A[N];
value_type B[N];

for(int i=0; i<N; ++i){
  if(tid<N){
    value_type sum = (i - tid) * B[tid];
    // reuction by shuffle
    num_warp = blockDim.x/warpSize + 1;
    __shared__ value_type shared_sum[num_warp];
    int lane = tid % warpSize;
    int warpID = tid / warpSize;
    for (int offset = warpSize/2; offset > 0; offset /= 2){
      sum += __shfl_down(sum, offset);
    }
    if(lane == 0){
      shared_sum[warpID] = sum;
    }
    if(tid==0){
      for(int k=0; k<num_warp; ++k){
        A[i] += shared_sum[k];
      }
    }
  }
}

使用 __shared__ 要注意同步的問題
例如第一種平行化,將資料放進 shared memory 後,讀取前應該要用 __syncthreads() 確保資料都已經放好。

我個人會偏向類似於第二種,盡量的去使用 register
這裡又會分成你要怎麼去處理 A (這就要看你 A, B 大小)
是要每一個 block 處理一個 A_i ,還是每一個 warp 處理一個 A_i
當每一個 block 處理 A_i ,會類似於你第二種,
注意:在你第二種使用 A[tid] += shared_sum[k](應該是 tid 不是 i?) 會導致你讀取跟寫入 global memory num_warp 次,應該繼續使用 register 先把資料算好,最後在把資料加總進去

而另一種是每一個 warp (或者說 subwarp) 處理一個 A_i (如果沒有其他條件或資料種類的資訊,我會偏向這種)

const auto warp_id = tid / warp_size; // handle A[warp_id]
const auto id_in_warp = tid % warp_size;
value_type sum = 0;
for (auto i = id_in_warp; i < N; i += warp_size) {
  sum = (i - warp_id) * B[i];
} 
// perform the warp reduction on sum
auto total_sum = warp_reduction(sum);
if (id_in_warp == 0) {
  A[warp_id] += total_sum;
}

warp_reduction 通常會用 xor,但用你原先提的也可,效率上在這裡應該是沒啥差。
再往上,可以考慮一次將 block_size 讀到 shared memory,做完後再讀下一輪的 block_size 進 shared_memory
但這個我就不確定有沒有幫助,因為會多了一些 shared memory 跟同步的操作。
而且原先的不同 warp 間要的資料都一樣,所以當其他 warp 在處理時,可以從 Cache 拿而非 Global memory

1個讚