CUDA 分支分歧没有任何区别

Question

我正在尝试自学 CUDA，现在遇到分支发散问题。据我所知，这是当一个 block 中的多个线程被称为采用分支(例如，由于 if 或 switch 语句)时出现的问题的名称，但该 block 中的其他线程不必拿去吧。

为了进一步研究这种现象及其后果，我编写了一个包含几个 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) 可能是编译器用像`result=ITERATIONS*A[threadIdx.x]*B[threadIdx.x] 这样的单次乘法替换了你的循环

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

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