CUDA分支分歧不會產生任何分歧

Question

我試圖通過自己學習CUDA，現在我正在考慮分支分歧問題。據我所知，這是一個給出的問題的名稱，當一個塊中的多個線程被稱爲分支時（例如由於if或switch語句），而在該塊中的其他線程不必拿去。

爲了進一步研究這個現象及其後果，我寫了一個帶有幾個CUDA函數的小文件。其中一個應該花費很多時間，因爲線程被停止了更多的時間（9999 ...迭代），而另一個線程只停留了一次分配。

但是，當我運行代碼時，我得到了非常相似的時間。而且，即使測量運行它們的時間，我也會得到類似於只運行一次的時間。我是否編寫了任何錯誤的代碼，或者對此有合理的解釋？

代碼：

#include <stdio.h> 
#include <stdlib.h> 
#include <cutil.h> 

#define ITERATIONS 9999999999999999999 
#define BLOCK_SIZE 16 

unsigned int hTimer; 

void checkCUDAError (const char *msg) 
{ 
cudaError_t err = cudaGetLastError(); 
if (cudaSuccess != err) 
{ 
    fprintf(stderr, "Cuda error: %s: %s.\n", msg,cudaGetErrorString(err)); 
    getchar(); 
    exit(EXIT_FAILURE); 
} 
} 

__global__ void divergence(float *A, float *B){ 
float result = 0; 
    if(threadIdx.x % 2 == 0) 
     { 
     for(int i=0;i<ITERATIONS;i++){ 
     result+=A[threadIdx.x]*A[threadIdx.x]; 
     } 

     } else 
     for(int i=0;i<ITERATIONS;i++){ 
      result+=A[threadIdx.x]*B[threadIdx.x]; 
     } 
} 

__global__ void betterDivergence(float *A, float *B){ 
float result = 0; 
float *aux; 
//This structure should not affect performance that much 
    if(threadIdx.x % 2 == 0) 
    aux = A; 
    else 
    aux = B; 

    for(int i=0;i<ITERATIONS;i++){ 
     result+=A[threadIdx.x]*aux[threadIdx.x]; 
    } 
} 

// ------------------------ 
// MAIN function 
// ------------------------ 
int main(int argc, char ** argv){ 

float* d_a; 
float* d_b; 
float* d_result; 
float *elementsA; 
float *elementsB; 

elementsA = (float *)malloc(BLOCK_SIZE*sizeof(float)); 
elementsB = (float *)malloc(BLOCK_SIZE*sizeof(float)); 

//"Randomly" filling the arrays 
for(int x=0;x<BLOCK_SIZE;x++){ 
    elementsA[x] = (x%2==0)?2:1; 
    elementsB[x] = (x%2==0)?1:3; 
} 

cudaMalloc((void**) &d_a, BLOCK_SIZE*sizeof(float)); 
cudaMalloc((void**) &d_b, BLOCK_SIZE*sizeof(float)); 
cudaMalloc((void**) &d_result, sizeof(float)); 

cudaMemcpy(d_a, elementsA, BLOCK_SIZE*sizeof(float), cudaMemcpyHostToDevice); 
cudaMemcpy(d_b, elementsB, BLOCK_SIZE*sizeof(float), cudaMemcpyHostToDevice); 

CUT_SAFE_CALL(cutCreateTimer(&hTimer)); 
CUT_CHECK_ERROR("cudaCreateTimer\n"); 

CUT_SAFE_CALL(cutResetTimer(hTimer)); 
CUT_CHECK_ERROR("reset timer\n"); 
CUT_SAFE_CALL(cutStartTimer(hTimer)); 
CUT_CHECK_ERROR("start timer\n"); 

float timerValue; 

dim3 dimBlock(BLOCK_SIZE,BLOCK_SIZE); 
dim3 dimGrid(32/dimBlock.x, 32/dimBlock.y); 

divergence<<<dimBlock, dimGrid>>>(d_a, d_b); 
betterDivergence<<<dimBlock, dimGrid>>>(d_a, d_b); 

checkCUDAError("kernel invocation"); 

cudaThreadSynchronize(); 
CUT_SAFE_CALL(cutStopTimer(hTimer)); 
CUT_CHECK_ERROR("stop timer\n"); 

timerValue = cutGetTimerValue(hTimer); 
printf("kernel execution time (secs): %f s\n", timerValue); 

return 0; 
} 

Answer 1

1）您有沒有記憶在你的__global__代碼編寫，除了局部變量（結果）。我不確定cuda編譯器是否這樣做，但是所有的代碼都可以安全地刪除而沒有副作用（也許編譯器已經這樣做了）。

2）__global__函數中的所有設備內存讀取都來自每次迭代的一個地方。 Cuda會將這個值存儲在寄存器內存中，並且這裏最快的操作（內存訪問）將會很快完成。

3）可以是編譯器曾與像`結果= ITERATIONS單個乘法代替你的週期* A [threadIdx.x] * B [threadIdx.x]

4）如果您的函數的所有代碼將按照你寫的那樣執行，你的betterDivergence將比你的另一個函數快大約2倍，因爲你的if分支中的循環較慢，而分支中的循環較快。但是在執行相同循環的線程之間線程中不會有任何空閒時間，因爲所有線程將在每次迭代中執行循環體。

我建議你寫另一個例子，你將結果存儲在某些設備內存中，然後將該內存複製回主機，並進行一些更不可預測的計算以防止可能的優化。

Answer 2

下面顯示的是最終的，經測試，正確的代碼，允許CUDA代碼之間的性能與不分支發散比較例如：

#include <stdio.h> 
#include <stdlib.h> 
#include <cutil.h> 

//#define ITERATIONS 9999999999999999999 
#define ITERATIONS 999999 
#define BLOCK_SIZE 16 
#define WARP_SIZE 32 

unsigned int hTimer; 

void checkCUDAError (const char *msg) 
{ 
cudaError_t err = cudaGetLastError(); 
if (cudaSuccess != err) 
{ 
    fprintf(stderr, "Cuda error: %s: %s.\n", msg,cudaGetErrorString(err)); 
    getchar(); 
    exit(EXIT_FAILURE); 
} 
} 

__global__ void divergence(float *A, float *B){ 
    int a = blockIdx.x*blockDim.x + threadIdx.x; 
    if (a >= ITERATIONS) return; 
    if(threadIdx.x > 2) 
     { 
     for(int i=0;i<ITERATIONS;i++){ 
     B[a]=A[a]+1; 
     } 
     } else 
     for(int i=0;i<ITERATIONS;i++){ 
     B[a]=A[a]-1; 
     } 
} 

__global__ void noDivergence(float *A, float *B){ 
    int a = blockIdx.x*blockDim.x + threadIdx.x; 
    if (a >= ITERATIONS) return; 
    if(threadIdx.x > WARP_SIZE) 
     { 
     for(int i=0;i<ITERATIONS;i++){ 
     B[a]=A[a]+1; 
     } 
     } else 
     for(int i=0;i<ITERATIONS;i++){ 
     B[a]=A[a]-1; 
     } 
} 

// ------------------------ 
// MAIN function 
// ------------------------ 
int main(int argc, char ** argv){ 

float* d_a; 
float* d_b; 
float* d_result; 
float *elementsA; 
float *elementsB; 

elementsA = (float *)malloc(BLOCK_SIZE*sizeof(float)); 
elementsB = (float *)malloc(BLOCK_SIZE*sizeof(float)); 

//"Randomly" filling the arrays 
for(int x=0;x<BLOCK_SIZE;x++){ 
    elementsA[x] = (x%2==0)?2:1; 
} 

cudaMalloc((void**) &d_a, BLOCK_SIZE*sizeof(float)); 
cudaMalloc((void**) &d_b, BLOCK_SIZE*sizeof(float)); 
cudaMalloc((void**) &d_result, sizeof(float)); 

cudaMemcpy(d_a, elementsA, BLOCK_SIZE*sizeof(float), cudaMemcpyHostToDevice); 
cudaMemcpy(d_b, elementsB, BLOCK_SIZE*sizeof(float), cudaMemcpyHostToDevice); 

CUT_SAFE_CALL(cutCreateTimer(&hTimer)); 
CUT_CHECK_ERROR("cudaCreateTimer\n"); 

CUT_SAFE_CALL(cutResetTimer(hTimer)); 
CUT_CHECK_ERROR("reset timer\n"); 
CUT_SAFE_CALL(cutStartTimer(hTimer)); 
CUT_CHECK_ERROR("start timer\n"); 

float timerValue; 

dim3 dimBlock(BLOCK_SIZE,BLOCK_SIZE); 
dim3 dimGrid(128/dimBlock.x, 128/dimBlock.y); 

//divergence<<<dimGrid, dimBlock>>>(d_a, d_b); 
noDivergence<<<dimGrid, dimBlock>>>(d_a, d_b); 

checkCUDAError("kernel invocation"); 

cudaThreadSynchronize(); 
CUT_SAFE_CALL(cutStopTimer(hTimer)); 
CUT_CHECK_ERROR("stop timer\n"); 

timerValue = cutGetTimerValue(hTimer)/1000; 
printf("kernel execution time (secs): %f s\n", timerValue); 

cudaMemcpy(elementsB, d_b, BLOCK_SIZE*sizeof(float), cudaMemcpyDeviceToHost); 

return 0; 
}