CUDA循環展開

Question

我在循環展開CUDA時遇到了一些問題。

在正常的串行代碼：

//serial basic: 
for(int i = 0; i < n; i++){ 
    c[i] = a[i] + b[i];} 

//serial loop unroll: 
for(int i = 0; i < n/4; i++){ 
    c[i] = a[i] + b[i]; 
    c[i+1] = a[i+1] + b[i+1]; 
    c[i+2] = a[i+2] + b[i+2]; 
    c[i+3] = a[i+3] + b[i+3];} 

所以我認爲CUDA循環展開如下：

int i = 2*(threadIdx.x + blockIdx.x * gridDim.x); 
a[i+0] = b[i+0] + c[i+0]; 
a[i+1] = b[i+1] + c[i+1]; 

但在CUDA手簿中展開的例子我能不明白

這是一個普通的GlobalWrite內核：

__global__ void GlobalWrites(T *out, T value, size_t N) 
{ 
for(size_t i = blockIdx.x*blockDim.x+threadIdx.x; 
    i < N; 
    i += blockDim.x*gridDim.x) { 
    out[i] = value; 
    } 
} 

展開內核：

template<class T, const int n> __global__ void Global_write(T* out, T value, size_t N){ 
size_t i; 
for(i = n*blockDim.x*blockIdx.x + threadIdx.x; 
    i < N - n*blockDim.x*blockIdx.x; 
    i += n*gridDim.x*blockDim.x;) 
    for(int j = 0; j < n; i++){ 
     size_t index = i + j * blockDim.x; 
     outp[index] = value; 
    } 
for (int j = 0; j < n; j++) { 
    size_t index = i+j*blockDim.x; 
    if (index<N) out[index] = value; 
}} 

我知道這個內核使用較少的街區，但有人可以解釋爲什麼它工作得更好（N = 4,10％的速度向上）。

Answer 1

如果它不明顯，因爲n是一個模板參數，它在編譯時是不變的。這意味着編譯器可以自由地通過展開來優化恆定跳閘計數循環。它是，因此，有教導以除去模板魔術和展開用手你提到的n = 4的情況下，循環：

template<class T> 
__global__ void Global_write(T* out, T value, size_t N) 
{ 
    size_t i; 
    for(i = 4*blockDim.x*blockIdx.x + threadIdx.x; 
     i < N - 4*blockDim.x*blockIdx.x; 
     i += 4*gridDim.x*blockDim.x;) { 
      out[i + 0 * blockDim.x] = value; 
      out[i + 1 * blockDim.x] = value; 
      out[i + 2 * blockDim.x] = value; 
      out[i + 3 * blockDim.x] = value; 
    } 
    if (i+0*blockDim.x < N) out[i+0*blockDim.x] = value; 
    if (i+1*blockDim.x < N) out[i+1*blockDim.x] = value; 
    if (i+2*blockDim.x < N) out[i+2*blockDim.x] = value; 
    if (i+3*blockDim.x < N) out[i+3*blockDim.x] = value; 
} 

展開的內環產率，其被合併4個完全獨立的寫入。正是這種指令級並行性給代碼提供了更高的指令吞吐量和更高的性能。如果你還沒有看到，我強烈推薦幾年前GTC會議上的瓦西里沃爾科夫的Unrolling Parallel Loops。他的演講闡述了爲什麼這種循環展開是CUDA優化的理論背景。

Answer 2

在模板化的內核中，編譯時已知const int n，允許編譯器實際展開for(int j = 0; j < n; i++)循環，從而刪除該循環上的條件檢查。如果在編譯時未知循環大小，編譯器將無法展開循環。就那麼簡單。