CUDA カーネルで 2D 配列をコピーする 2 つの異なるアプローチをテストしました。
最初のものは、TILE_DIM x TILE_DIM スレッドのブロックを起動します。各ブロックは、要素ごとに 1 つのスレッドを割り当てる配列のタイルをコピーします。
__global__ void simple_copy(float *outdata, const float *indata){
int x = blockIdx.x * TILE_DIM + threadIdx.x;
int y = blockIdx.y * TILE_DIM + threadIdx.y;
int width = gridDim.x * TILE_DIM;
outdata[y*width + x] = indata[y*width + x];
}
2 つ目は、NVIDIA ブログからの抜粋です。以前のカーネルに似ていますが、ブロックごとに TILE_DIM x BLOCK_ROWS スレッドを使用します。各スレッドは、行列の複数の要素をループします。
__global__ void fast_copy(float *outdata, const float *indata)
{
int x = blockIdx.x * TILE_DIM + threadIdx.x;
int y = blockIdx.y * TILE_DIM + threadIdx.y;
int width = gridDim.x * TILE_DIM;
for (int k = 0 ; k < TILE_DIM ; k += BLOCK_ROWS)
outdata[(y+k)*width + x] = indata[(y+k)*width + x];
}
これら 2 つのアプローチを比較するテストを実行します。両方のカーネルがグローバル メモリへの合体アクセスを実行しますが、2 番目のカーネルは著しく高速に見えます。
NVIDIA ビジュアル プロファイラーは、このテストを確認します。
では、2 番目のカーネルはどのようにしてより高速なコピーを実現するのでしょうか?
これは、カーネルのテストに使用した完全なコードです。
#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <conio.h>
#define TILE_DIM 32
#define BLOCK_ROWS 8
/* KERNELS */
__global__ void simple_copy(float *outdata, const float *indata){
int x = blockIdx.x * TILE_DIM + threadIdx.x;
int y = blockIdx.y * TILE_DIM + threadIdx.y;
int width = gridDim.x * TILE_DIM;
outdata[y*width + x] = indata[y*width + x];
}
//###########################################################################
__global__ void fast_copy(float *outdata, const float *indata)
{
int x = blockIdx.x * TILE_DIM + threadIdx.x;
int y = blockIdx.y * TILE_DIM + threadIdx.y;
int width = gridDim.x * TILE_DIM;
for (int k = 0 ; k < TILE_DIM ; k += BLOCK_ROWS)
outdata[(y+k)*width + x] = indata[(y+k)*width + x];
}
//###########################################################################
/* MAIN */
int main(){
float *indata,*dev_indata,*outdata1,*dev_outdata1,*outdata2,*dev_outdata2;
cudaEvent_t start, stop;
float time1,time2;
int i,j,k;
int n_iter = 100;
int N = 2048;
cudaEventCreate(&start);
cudaEventCreate(&stop);
dim3 grid(N/TILE_DIM, N/TILE_DIM);
dim3 threads1(TILE_DIM,TILE_DIM);
dim3 threads2(TILE_DIM,BLOCK_ROWS);
// Allocations
indata = (float *)malloc(N*N*sizeof(float));
outdata1 = (float *)malloc(N*N*sizeof(float));
outdata2 = (float *)malloc(N*N*sizeof(float));
cudaMalloc( (void**)&dev_indata,N*N*sizeof(float) );
cudaMalloc( (void**)&dev_outdata1,N*N*sizeof(float) );
cudaMalloc( (void**)&dev_outdata2,N*N*sizeof(float) );
// Initialisation
for(j=0 ; j<N ; j++){
for(i=0 ; i<N ; i++){
indata[i + N*j] = i + N*j;
}
}
// Transfer to Device
cudaMemcpy( dev_indata, indata, N*N*sizeof(float),cudaMemcpyHostToDevice );
// Simple copy
cudaEventRecord( start, 0 );
for(k=0 ; k<n_iter ; k++){
simple_copy<<<grid, threads1>>>(dev_outdata1,dev_indata);
}
cudaEventRecord( stop, 0 );
cudaEventSynchronize( stop );
cudaEventElapsedTime( &time1, start, stop );
printf("Elapsed time with simple copy: %f\n",time1);
// Fast copy
cudaEventRecord( start, 0 );
for(k=0 ; k<n_iter ; k++){
fast_copy<<<grid, threads2>>>(dev_outdata2,dev_indata);
}
cudaEventRecord( stop, 0 );
cudaEventSynchronize( stop );
cudaEventElapsedTime( &time2, start, stop );
printf("Elapsed time with fast copy: %f\n",time2);
// Transfer to Host
cudaMemcpy( outdata1, dev_outdata1, N*N*sizeof(float),cudaMemcpyDeviceToHost );
cudaMemcpy( outdata2, dev_outdata2, N*N*sizeof(float),cudaMemcpyDeviceToHost );
// Check for error
float error = 0;
for(j=0 ; j<N ; j++){
for(i=0 ; i<N ; i++){
error += outdata1[i + N*j] - outdata2[i + N*j];
}
}
printf("error: %f\n",error);
/*// Print the copied matrix
printf("Copy\n");
for(j=0 ; j<N ; j++){
for(i=0 ; i<N ; i++){
printf("%f\t",outdata1[i + N*j]);
}
printf("\n");
}*/
cudaEventDestroy( start );
cudaEventDestroy( stop );
free(indata);
free(outdata1);
free(outdata2);
cudaFree(dev_indata);
cudaFree(dev_outdata1);
cudaFree(dev_outdata2);
cudaDeviceReset();
getch();
return 0;
}
//###########################################################################