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;
}

- Jorge Antonio Díaz-Benito

---

检查您的代码后，我发现所有线程都在执行99999次迭代。哪些线程应该更快？ - Evans

编译选项是什么？ - Mikhail

你在内核调用中颠倒了dimBlock和dimGrid变量。dimGrid应该放在前面。我同意答案，编译器可能会优化掉代码。 - Robert Crovella

Answer 1

1) 在您的__global__代码中，除了本地变量（result）之外，您没有进行任何内存写操作。我不确定Cuda编译器是否会这样做，但是您的整个代码可以被安全地删除而不会产生任何副作用（也许编译器已经这样做了）。

2) 在__global__函数中，您从设备内存读取的所有数据都来自每次迭代中的同一位置。Cuda将在寄存器存储器中存储该值，并且最长操作（内存访问）将在此处非常快速地完成。

3) 可能编译器已将循环替换为单个乘法，例如`result=ITERATIONS*A[threadIdx.x]*B[threadIdx.x]`。

4) 如果您的函数中的所有代码将按照编写的方式执行，则您的betterDivergence大约比另一个函数快两倍，因为您在较慢的函数中具有包含循环的if分支，而在较快的函数中，分支中没有循环。但是，在执行相同循环的线程之间不会有任何空闲时间，因为所有线程都将在每个迭代中执行循环体。

我建议您编写另一个示例，将结果存储在某些设备内存中，然后将该内存复制回主机并进行更多的不可预测计算，以防止可能的优化。

- JustAnotherCurious

---

最终，在您的帮助下，我成功编写了一个适当的示例，并使用了此链接：http://courses.engr.illinois.edu/ece408/lectures/ece408-lecture4-CUDA%20parallelism-model-2012.pdf - Jorge Antonio Díaz-Benito

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;
}

- Jorge Antonio Díaz-Benito