cuda - CUDA 共有メモリの書き込みにより、説明のつかない長いレイテンシが発生する

Question

これは私を夢中にさせています。1D ブロックの 3D グリッドがあります。各ブロックには 272 のスレッドが含まれます。各スレッドは、2 つのベクトルのドット積を計算し、その結果を、同じスレッド数である [272] のサイズの double の共有メモリ配列内の対応する場所に格納します。メインスレッドは複数のカーネルを呼び出しており、それらすべてを実行するのにかかった時間を合計しています。共有メモリに書き込む行をコメントアウトすると、約 2,401 ミリ秒の実行時間が得られます。共有メモリ書き込み行のコメントを外すと、450,309 ミリ秒のように非常に長い時間がかかります。double の代わりに int 値を使用してみました。また、可能性のあるバンクの競合を回避するために、1 つのスレッドだけが書き込みを実行できるようにする if(threadIdx.x ==0) ステートメントを用意してみました。何も機能していないようです。呼び出しスレッドのコードは次のとおりです。

  double theta=0;
int count=0;
cudaEventRecord(start,0);
while(theta <180)
{
theta+=0.18;

calc_LF<<<gridDim, blockDim>>>(ori_dev, X_dev, Y_dev, Z_dev, F_dev, F_grad_dev, g_oriD, r_vD, LF);
calc_S<<<gridDim, 272>>>(g_traD, LF, Ci, C);
count++;
}
cudaEventRecord( stop, 0 );
cudaEventSynchronize( stop );
cudaEventElapsedTime( &elapsedTime, start, stop );
err = cudaGetLastError();
if ( cudaSuccess != err )
    {
    fprintf( stderr, "Cuda error in file '%s' in line %i : %s.\n",
         __FILE__, __LINE__, cudaGetErrorString( err) );
    }
else
    {
    fprintf( stderr, "\n \n Cuda NO error in file '%s' in line %i : %s.\n",
         __FILE__, __LINE__, cudaGetErrorString( err) );
    printf("\n %d orientation updates: Total Time = %3.10f ms\n", count, elapsedTime);
    }

問題のカーネルは calc_S カーネルで、そのコードは次のとおりです。

__global__ void calc_S(double* g_traD, double* LF, double* Ci, double* C)
{


__shared__ double G[H];
int myTRA[W];
int tx= threadIdx.x;

for(int j=0; j<W; j++)
{
    myTRA[j]= getElement(g_traD, tx, j, W);
}
double sum;
for(int j=0; j<W; j++)
{
     sum += myTRA[j] * LF[j];
}       

// Write your sum to shared memory
    G[threadIdx.x]=sum;
     //__syncthreads();
}

MS Visual Studio 2008 を CUDA 4.2 とコンピューティング機能 2.0 の GPU (つまり、GeForce GTX 580) で使用しています。注: ブロックあたり 272 スレッド。H/W スレッド制限: 1,536 / 272 = 最大 5 ブロック共有メモリ制限: double の G[272] = 2,176 バイトが必要。48K / 2176 = 最大 22 ブロック (発生することはありませんが、共有メモリに制限がないことはわかっています) レジスタはまったく問題になりません。したがって、5 つのブロックを同時に実行できるはずです。

助けてくれてありがとう。

マイ

編集：

これは、コード全体の短縮版です。コード全体は、MatrixMul Nvidia SDK サンプルで実行できます。

ファイル「MatrixMul.cu」内

    int main(int argc, char** argv)
    {
    // reading data from Matlab into double arrays
    //CUDA begins here:
if(shrCheckCmdLineFlag(argc, (const char**)argv, "device"))
    {
      cutilDeviceInit(argc, argv);
     }
   else
   {
     cutilSafeCall( cudaSetDevice(cutGetMaxGflopsDeviceId()) );
   }
  int devID;
  cudaDeviceProp props;

// get GPU props
cutilSafeCall(cudaGetDevice(&devID));
cutilSafeCall(cudaGetDeviceProperties(&props, devID));

printf("Device %d: \"%s\" with Compute %d.%d capability\n", devID, props.name, props.major, props.minor);

//Declare Device memory for matrices read from Matlab
 double *X_dev;    // size 19 x 1
 double *Y_dev;    // size 19 x 1
 double *Z_dev;    // size 17 x 1
 double *r_vD;     // size 544 x 3
 double *g_oriD;   // size 544 x 3
 double *g_traD;   // size 272 x 544
 double *cov_D;    // size 272 x 272
 double *cov_i_D;  // size 272 x 272

err= cudaMalloc((void**)&X_dev, sizeX*sizeof(double));
errorCheck(err);

err= cudaMalloc((void**)&Y_dev, sizeY*sizeof(double));
errorCheck(err);

err= cudaMalloc((void**)&Z_dev, sizeZ*sizeof(double));
errorCheck(err);

err= cudaMalloc((void**)&r_vD, sizeR_V*sizeof(double));
errorCheck(err);

err= cudaMalloc((void**)&g_oriD, sizeG_ori*sizeof(double));
errorCheck(err);

err= cudaMalloc((void**)&g_traD, sizeG_tra*sizeof(double));
errorCheck(err);

err= cudaMalloc((void**)&cov_D, sizeCov*sizeof(double));
errorCheck(err); 

err= cudaMalloc((void**)&cov_i_D, sizeCov_i*sizeof(double));
errorCheck(err); 

   //Transfer Xs, Ys, and Zs to GPU Global memory
cudaMemcpy(X_dev,dipole_x_coords, sizeX*sizeof(double), cudaMemcpyHostToDevice);
errorCheck(err);

cudaMemcpy(Y_dev,dipole_y_coords, sizeY*sizeof(double), cudaMemcpyHostToDevice);
errorCheck(err);

cudaMemcpy(Z_dev,dipole_z_coords, sizeZ*sizeof(double), cudaMemcpyHostToDevice);
errorCheck(err);

    // Transfer r_v, g_ori, and g_tra to GPU memory
cudaMemcpy(r_vD, r_v, sizeR_V*sizeof(double), cudaMemcpyHostToDevice);
errorCheck(err);

cudaMemcpy(g_oriD,g_ori, sizeG_ori*sizeof(double), cudaMemcpyHostToDevice);
errorCheck(err);

cudaMemcpy(g_traD,g_tra, sizeG_tra*sizeof(double), cudaMemcpyHostToDevice);
    errorCheck(err);

    // Transfer cov, and cov_i to GPU memory
cudaMemcpy(cov_D, cov_post, sizeCov*sizeof(double), cudaMemcpyHostToDevice);
errorCheck(err);

cudaMemcpy(cov_i_D,cov_post_i, sizeCov_i*sizeof(double), cudaMemcpyHostToDevice);
  //Specify dimensions of block and grid
dim3 gridDim(sizeX, sizeY, sizeZ);   // 19 x 19 x 17
int numThreads=(int) sizeR_V/3;      // numThreads = 544
dim3 blockDim(numThreads,1,1);       // 544 x 1 x 1 

//call Cuda wrapper
float cf = runB(X_dev, Y_dev, Z_dev, r_vD, g_oriD, g_traD, cov_i_D, cov_D, blockDim, gridDim, sizeG_tra, tra_W, tra_H);

int c=0;
scanf("%d", c);
return 0;

}


    float runB(double* X_dev, double* Y_dev, double* Z_dev, 
double* r_vD, double* g_oriD, double* g_traD, double* Ci, double* C,
dim3 blockDim, dim3 gridDim, int sizeG_tra, int tra_W, int tra_H)
    {  
   cudaError err;

   // Calculate the size of thread output in global memory
   size_t size_F = gridDim.x * gridDim.y * gridDim.z * blockDim.x;
   size_t size_F_grad = gridDim.x * gridDim.y * gridDim.z * blockDim.x * 3;

   // Make global memory space for F and F_grad 
double* F_dev;
double* F_grad_dev;
err= cudaMalloc((void**)&F_dev, size_F*sizeof(double));
errorCheck(err); 
err= cudaMalloc((void**)&F_grad_dev, size_F_grad*sizeof(double));
errorCheck(err); 

    //Allocate Device memory for LF 
double *LF;
err= cudaMalloc((void**)&LF, 544*sizeof(double));
errorCheck(err); 

    cudaEvent_t start, stop;
    float elapsedTime;

    cudaEventCreate(&start);
    cudaEventCreate(&stop);

double theta=0;
cudaEventRecord(start,0);
while(theta <180)
{
theta+=0.18;
calc_LF<<<gridDim, blockDim>>>(ori_dev, X_dev, Y_dev, Z_dev, F_dev, F_grad_dev, g_oriD, r_vD, LF);
   calc_S<<<gridDim, 272>>>(g_traD, LF, Ci, C);
   count++;
   }
   cudaEventRecord( stop, 0 );
   cudaEventSynchronize( stop );
cudaEventElapsedTime( &elapsedTime, start, stop );
err = cudaGetLastError();
if ( cudaSuccess != err )
    {
    fprintf( stderr, "Cuda error in file '%s' in line %i : %s.\n",
         __FILE__, __LINE__, cudaGetErrorString( err) );
    }
else
    {
    fprintf( stderr, "\n \n Cuda NO error in file '%s' in line %i : %s.\n",
         __FILE__, __LINE__, cudaGetErrorString( err) );
    printf("\n 180 orientation updates: Total Time = %3.10f ms\n",elapsedTime);
    }
     return 0;

   }

ファイル「MatrixMul_kernel.cu」内

     #define HDM_DIM 3 
      __global__ void calc_LF(double* ori_dev, double* X_dev, double* Y_dev, double* Z_dev, double* F_dev, double* F_grad_dev, 
                    double* g_oriD, double* r_vD, double* LF)
        { 
        // Get this block's global index 
     int blockId= blockIdx.x + gridDim.x*blockIdx.y + gridDim.x*gridDim.y*blockIdx.z;
     int tx= threadIdx.x;
    // This thread's global index
     int gtx= blockId*blockDim.x + threadIdx.x;
    double r_v[3];
    double g_ori[3];

// Each thread reads 1 row (3 values) of r_vD 
    r_v[0] = getElement(r_vD, tx, 0, HDM_DIM);
    r_v[1] = getElement(r_vD, tx, 1, HDM_DIM);
    r_v[2] = getElement(r_vD, tx, 2, HDM_DIM);
// Each thread reads 1 row (3 values) of g_oriD (which contains grad.ori data)
    g_ori[0] = getElement(g_oriD, tx, 0, HDM_DIM);
    g_ori[1] = getElement(g_oriD, tx, 1, HDM_DIM);
    g_ori[2] = getElement(g_oriD, tx, 2, HDM_DIM);

    //fetch d_ori from global memory
    double d_ori[3];
    for(int i=0; i< 3; i++){
       d_ori[i]= ori_dev[3*gtx+i];
    }   
   //read this block's X, Y, Z location
    double x= X_dev[blockIdx.x];
    double y= Y_dev[blockIdx.y];
    double z= Z_dev[blockIdx.z];     

    double c2[HDM_DIM];
    c2[0]= d_ori[1]*z - d_ori[2]*y;
    c2[1]= d_ori[2]*x - d_ori[0]*z;
    c2[2]= d_ori[0]*y - d_ori[1]*x;

    // Fetch F and F_grad from global memory
    double F = F_dev[gtx];
    double F_grad[3];
    for(int j=0; j<3; j++)
    {
        F_grad[j] = F_grad_dev[gtx*3+j];
    }


    double c1[HDM_DIM];
    c1[0]= F* c2[0];
    c1[1]= F* c2[1];
    c1[2]= F* c2[2];


    double d3= c2[0]*r_v[0] + c2[1]*r_v[1] + c2[2]*r_v[2];

    double s2[HDM_DIM];
    for(int j=0; j<HDM_DIM; j++)
    {
        s2[j] = d3*F_grad[j];
    }

    double s1[HDM_DIM];
    for(int j=0; j<HDM_DIM; j++)
    {
        s1[j] = c1[j] - s2[j];
    }

    double b_v[HDM_DIM];
    for(int j=0; j<HDM_DIM; j++)
    {
        b_v[j] = (10^-7)/(F*F) * s1[j]; 
    }   

    double sum=0;
    for(int j=0; j<HDM_DIM; j++)
    {
        sum += b_v[j]*g_ori[j];
    }   

// Write this thread's value to global memory
    LF[tx]= sum;

     }

言及する価値があるのは、この calc_LF カーネルが最終結果を共有メモリに書き込んでいたため、実行時間が約 500 ミリ秒以上から約 2,500 ミリ秒に増加したことです (つまり、共有メモリの書き込み行だけで、時間がおよそ 5 倍になりました)。

    __global__ void calc_S(double* g_traD, double* LF, double* Ci, double* C)
{
__shared__ double T[H];
__shared__ double G[H];

   // Get this block's global index 
  int blockId= blockIdx.x + gridDim.x*blockIdx.y + gridDim.x*gridDim.y*blockIdx.z;
  int tx= threadIdx.x;
// This thread's global index
   int gtx= blockId*blockDim.x + threadIdx.x;

int myTRA[W];
double my_LF[W];
for (int i=0; i<W; i++){
   my_LF[i]= LF[gtx];
}

for(int j=0; j<W; j++){
    myTRA[j]= getElement(g_traD, tx, j, W);
        }
    double sum;
    for(int j=0; j<W; j++)
    {
         sum += myTRA[j] * my_LF[j];
    }       

// Write your sum to shared memory
    G[tx]=sum;
    __syncthreads();
      }

score 1 · Accepted Answer

表示されている効果は、コンパイラの最適化の結果です。基本的なカーネルコードのコンパイル可能なバージョンを取得します。

#define H (128)
#define W (128)

__device__
double getElement(const double *g, int t, int j, int w)
{
    return g[t + j*w];
}

__global__ 
void calc_S(double* g_traD, double* LF, double* Ci, double* C)
{
    __shared__ double G[H];

    // Get this block's global index 
    int blockId= blockIdx.x + gridDim.x*blockIdx.y + 
                   gridDim.x*gridDim.y*blockIdx.z;
    int tx= threadIdx.x;
    // This thread's global index
    int gtx= blockId*blockDim.x + threadIdx.x;

    int myTRA[W];
    double my_LF[W];
    for (int i=0; i<W; i++){
        my_LF[i]= LF[gtx];
    }

    for(int j=0; j<W; j++){
        myTRA[j]= getElement(g_traD, tx, j, W);
    }
    double sum;
    for(int j=0; j<W; j++)
    {
        sum += myTRA[j] * my_LF[j];
    }       

    // Write your sum to shared memory
    G[tx]=sum;
    __syncthreads();
}

CUDA 5でコンパイルすると、次のようになります。

$ nvcc -m64 -arch=sm_20 -cubin -Xptxas="-v"  dead_code.cu 
dead_code.cu(13): warning: variable "G" was set but never used

dead_code.cu(13): warning: variable "G" was set but never used

ptxas info    : 0 bytes gmem
ptxas info    : Compiling entry function '_Z6calc_SPdS_S_S_' for 'sm_20'
ptxas info    : Function properties for _Z6calc_SPdS_S_S_
    1536 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 23 registers, 1024 bytes smem, 64 bytes cmem[0]

共有メモリ変数が使用されていないという警告がありますGが、コンパイラはそれを尊重し、23 個のレジスタを消費するコードを出力します。G[tx]=sumしたがって、カーネルの最後にあるをコメントアウトすると、次のようにコンパイルされます。

$ nvcc -m64 -arch=sm_20 -cubin -Xptxas="-v"  dead_code.cu 
dead_code.cu(13): warning: variable "G" was declared but never referenced

dead_code.cu(13): warning: variable "G" was declared but never referenced

ptxas info    : 0 bytes gmem
ptxas info    : Compiling entry function '_Z6calc_SPdS_S_S_' for 'sm_20'
ptxas info    : Function properties for _Z6calc_SPdS_S_S_
    0 bytes stack frame, 0 bytes spill stores, 0 bytes spill loads
ptxas info    : Used 2 registers, 64 bytes cmem[0]

現在、使用されているレジスタは 2 つだけで、ツールチェーンはこれを発行しました。

$ cuobjdump -sass dead_code.cubin 

    code for sm_20
        Function : _Z6calc_SPdS_S_S_
    /*0000*/     /*0x00005de428004404*/     MOV R1, c [0x1] [0x100];
    /*0008*/     /*0xfc1fdc03207e0000*/     IMAD.U32.U32 RZ, R1, RZ, RZ;
    /*0010*/     /*0xffffdc0450ee0000*/     BAR.RED.POPC RZ, RZ;
    /*0018*/     /*0x00001de780000000*/     EXIT;

すなわち。4つの組み立て説明書。すべてのコードがなくなりました。

この影響の根本的な原因は、コンパイラのデッドコードの削除です。コンパイラは、グローバルまたは共有メモリの出力に影響を与えないコードが不要であり、削除できると判断するのに十分スマートです。この場合、書き込み先の 1 つGが削除され、カーネル全体が事実上無意味になり、コンパイラはすべてを最適化するだけです。デッドコードの削除とその効果の他の例を、こことここで見ることができます。後者は OpenCL にありますが、同じメカニズムが適用されます。

cuda - CUDA 共有メモリの書き込みにより、説明のつかない長いレイテンシが発生する

1 に答える 1

Related

Reference