3

我正在尝试自己学习 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;
}
4

2 回答 2

4

__global__1）除了局部变量（结果）之外，您的代码中没有内存写入。我不确定 cuda 编译器是否会这样做，但是您的所有代码都可以安全地删除而没有副作用（也许编译器已经这样做了）。

2）您在函数中从设备内存中读取的所有内容__global__都来自每次迭代的一个地方。Cuda 会将值存储在寄存器内存中，最长的操作（内存访问）将在这里完成得非常快。

3）可能是编译器已经用单次乘法替换了你的循环，比如 `result=ITERATIONS*A[threadIdx.x]*B[threadIdx.x]

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

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

于 2013-05-20T13:27:24.187 回答
0

Below is shown the final, tested, right example of a code that allows to compare the performance between CUDA code with and without branch divergence:

#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;
}
于 2013-05-21T19:02:19.680 回答