3

私は自分でCUDAを学ぼうとしていますが、現在、ブランチの分岐の問題に取り組んでいます。私が理解している限りでは、これは、ブロック内の複数のスレッドが (たとえば、if または switch ステートメントによって) 分岐すると言われているが、そのブロック内の他のスレッドは分岐する必要がない場合に発生する問題に付けられた名前です。それを取る。

この現象とその結果をもう少し詳しく調査するために、いくつかの CUDA 関数を含む小さなファイルを作成しました。そのうちの 1 つは、他のスレッド (割り当てのためにのみ停止される) よりもはるかに多くの時間 (9999... 反復) スレッドが停止されるため、多くの時間がかかると想定されています。

ただし、コードを実行すると、非常によく似た時間になります。さらに、両方を実行するのにかかる時間を測定しても、1 つだけを実行するのと同じような時間が得られます。何か間違ったコードを書いたのですか、それとも論理的な説明はありますか?

コード：

#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

0

以下は、分岐分岐がある場合とない場合の 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;
}
于 2013-05-21T19:02:19.680 に答える