Shared Memoryのバンク衝突についての実験
作成日 2008/7/18
Shared Memoryのバンク衝突を回避するためにどのようなアドレッシングにすればよいのか、
また、バンク衝突が起きた場合、どの程度のパフォーマンスの低下が見られるのか、実験してみました。
まずは、ソースコードを。カーネル部分だけ示します。
もとは、SDKのtemplateサンプルとほぼ同じです。
| #ifndef _SharedTest_KERNEL_H_ #define _SharedTest_KERNEL_H_ #include <stdio.h> #define SDATA( index) CUT_BANK_CHECKER(sdata, index) //////////////////////////////////////////////////////////////////////////////// //! Simple test kernel for device functionality //! @param g_idata input data in global memory //! @param g_odata output data in global memory //////////////////////////////////////////////////////////////////////////////// //Culumn Major Access __global__ void SharedTestKernel1( float* g_idata, float* g_odata) { // shared memory // the size is determined by the host application extern __shared__ float sdata[]; // access thread id // access number of threads in this block const unsigned int base = blockIdx.x * 256; const unsigned int num_threads = blockDim.x; const unsigned int tid = threadIdx.x + threadIdx.y * num_threads; // read in input data from global memory // use the bank checker macro to check for bank conflicts during host // emulation SDATA(tid) = g_idata[tid+base]; __syncthreads(); // perform some computations SDATA(tid) = 2.0f * SDATA( tid); __syncthreads(); // write data to global memory g_odata[tid+base] = SDATA(tid); } //Row Major Access __global__ void SharedTestKernel2( float* g_idata, float* g_odata) { // shared memory // the size is determined by the host application extern __shared__ float sdata[]; // access thread id // access number of threads in this block const unsigned int base = blockIdx.x * 256; const unsigned int num_threads = blockDim.y; const unsigned int tid = threadIdx.y + threadIdx.x * num_threads; // read in input data from global memory // use the bank checker macro to check for bank conflicts during host // emulation SDATA(tid) = g_idata[tid+base]; __syncthreads(); // perform some computations SDATA(tid) = 2.0f * SDATA( tid); __syncthreads(); // write data to global memory g_odata[tid+base] = SDATA(tid); } #endif // #ifndef _SharedTest_KERNEL_H_ |
SharedTest1
は、スレッド番号と共有メモリのアドレッシングがリニアになっているもの。
SharedTest2は、スレッド番号と共有メモリのアドレッシングがリニアでなく、縦になっているものです。
Test1では、同じワープに所属するスレッド0,1,2,3...が、メモリの0-3,4-7,8-11,12-15...をアクセスします。
Test2では、同じワープに所属するスレッド0,1,2,3...が、メモリの0-3,64-67,128-131,192-195をアクセスします。
このため、Test1のアクセスはすべて異なるバンクをアクセスしますが、Test2では全てが同じバンクをアクセスします。
これを呼び出す側も一部抜粋しておきます。
これは、SDKのtemplateサンプルをベースにしています。
フルソースはそちらを参照して作ってください。
| void runTest( int argc, char** argv) { CUT_DEVICE_INIT(); unsigned int timer = 0; unsigned int num_threads = 16*16; unsigned int smem_size = sizeof( float) * num_threads; unsigned int num_blocks = 256; unsigned int mem_size = sizeof( float) * num_threads * num_blocks; // allocate host memory float* h_idata = (float*) malloc( mem_size); // initalize the memory for( unsigned int i = 0; i < num_threads * num_blocks; ++i) { h_idata[i] = (float) i; } // allocate device memory float* d_idata; CUDA_SAFE_CALL( cudaMalloc( (void**) &d_idata, mem_size)); // copy host memory to device CUDA_SAFE_CALL( cudaMemcpy( d_idata, h_idata, mem_size, cudaMemcpyHostToDevice) ); CUDA_SAFE_CALL( cudaThreadSynchronize() ); // allocate device memory for result float* d_odata; CUDA_SAFE_CALL( cudaMalloc( (void**) &d_odata, mem_size)); // allocate mem for check float* reference = (float*) malloc( mem_size); // allocate mem for the result on host side float* h_odata = (float*) malloc( mem_size); // setup execution parameters dim3 grid( num_blocks, 1, 1); dim3 threads( 16, 16, 1); printf("Execute Culumn Major Test\n"); CUT_SAFE_CALL( cutCreateTimer( &timer)); // execute the kernel for(int i=0; i<100; i++) { CUT_SAFE_CALL( cutStartTimer( timer)); SharedTestKernel1<<< grid, threads, smem_size >>>( d_idata, d_odata); CUDA_SAFE_CALL( cudaThreadSynchronize() ); CUT_SAFE_CALL( cutStopTimer( timer)); } // check if kernel execution generated and error CUT_CHECK_ERROR("Kernel execution failed"); // copy result from device to host CUDA_SAFE_CALL( cudaMemcpy( h_odata, d_odata, sizeof( float) * num_threads* num_blocks, cudaMemcpyDeviceToHost) ); CUDA_SAFE_CALL( cudaThreadSynchronize() ); printf( "Processing time: %f (ms)\n", cutGetAverageTimerValue( timer)); CUT_SAFE_CALL( cutDeleteTimer( timer)); // compute reference solution SharedTestGold( reference, h_idata, num_threads* num_blocks); // check result if( cutCheckCmdLineFlag( argc, (const char**) argv, "regression")) { // write file for regression test CUT_SAFE_CALL( cutWriteFilef( "./data/regression.dat", h_odata, num_threads, 0.0)); } else { // custom output handling when no regression test running // in this case check if the result is equivalent to the expected soluion CUTBoolean res = cutComparef( reference, h_odata, num_threads* num_blocks); printf( "Test %s\n", (1 == res) ? "PASSED" : "FAILED"); } //Row Major Test printf("Execute Row Major Test\n"); CUT_SAFE_CALL( cutCreateTimer( &timer)); // execute the kernel for( int i=0; i<100; i++) { CUT_SAFE_CALL( cutStartTimer( timer)); SharedTestKernel2<<< grid, threads, smem_size >>>( d_idata, d_odata); CUDA_SAFE_CALL( cudaThreadSynchronize() ); CUT_SAFE_CALL( cutStopTimer( timer)); } // check if kernel execution generated and error CUT_CHECK_ERROR("Kernel execution failed"); // copy result from device to host CUDA_SAFE_CALL( cudaMemcpy( h_odata, d_odata, sizeof( float) * num_threads* num_blocks, cudaMemcpyDeviceToHost) ); printf( "Processing time: %f (ms)\n", cutGetAverageTimerValue( timer)); CUT_SAFE_CALL( cutDeleteTimer( timer)); // compute reference solution SharedTestGold( reference, h_idata, num_threads* num_blocks); // check result if( cutCheckCmdLineFlag( argc, (const char**) argv, "regression")) { // write file for regression test CUT_SAFE_CALL( cutWriteFilef( "./data/regression.dat", h_odata, num_threads, 0.0)); } else { // custom output handling when no regression test running // in this case check if the result is equivalent to the expected soluion CUTBoolean res = cutComparef( reference, h_odata, num_threads* num_blocks); printf( "Test %s\n", (1 == res) ? "PASSED" : "FAILED"); } // cleanup memory free( h_idata); free( h_odata); free( reference); CUDA_SAFE_CALL(cudaFree(d_idata)); CUDA_SAFE_CALL(cudaFree(d_odata)); } |
速
度測定のために、配列を16*16*256のサイズにしています。
これが小さいと、実行のオーバーヘッドなどのためか、実行時間にばらつきが生じてしまいます。
また、呼び出し側でのスレッド同期の実行も必須です。
この結果、
Test1(衝突なし):0.0629[ms]
Test2(衝突あり):0.249[ms]
という結果になりました。実行速度に大きな差が生じていることが分かります。
ちなみに、Emulationモードで実行するとバンク衝突を検出できるようになっていますが、
これによってログを取ってみると、Test2においてのみバンク衝突が起きていることが確認できます。