私は Cuda 開発の初心者であり、それがどのように機能するかを理解するために、小さな例のスクリプトを作成することにしました。私が作成し、2 つの等しいサイズの行列の対応する行間の 2 乗ユークリッド距離を計算するカーネル関数を共有することにしました。
__global__ void cudaEuclid( float* A, float* B, float* C, int rows, int cols )
{
int i, squareEuclDist = 0;
int r = blockDim.x * blockIdx.x + threadIdx.x; // rows
//int c = blockDim.y * blockIdx.y + threadIdx.y; // cols
if( r < rows ){ // take each row with var r (thread)
for ( i = 0; i < cols; i++ )//compute squared Euclid dist of each row
squareEuclDist += ( A[r + rows*i] - B[r + rows*i] ) * ( A[r + rows*i] - B[r + rows*i] );
C[r] = squareEuclDist;
squareEuclDist = 0;
}
}
カーネルの初期化は、
int threadsPerBlock = 256;
int blocksPerGrid = ceil( (double) numElements / threadsPerBlock);
// numElements = 1500x200 (matrix size) ==> 1172 blocks/grid
と呼ばれ、
cudaEuclid<<<blocksPerGrid, threadsPerBlock>>>( d_A, d_B, d_C, rows, cols );
d_A と d_B は挿入された行列で、この例ではサイズ 1500 x 200 です。
質問 1 : ブロックごとのスレッド数とグリッド数ごとのブロック数を選択する基本理論を読みましたが、まだ何かが欠けています。この単純なカーネルで最適なカーネル パラメーターの初期化を理解しようとしています。CUDA の方法で考え始めるために少し助けを求めています。
質問 2 : もう 1 つお聞きしたいのは、コードの効率を改善する方法について何か提案があれば教えてください。int c = blockDim.y * blockIdx.y + threadIdx.y物事をより並列にするために使用できますか?共有メモリはここで適用できますか?
以下に、私の GPU 情報が添付されています。
Device 0: "GeForce 9600 GT"
CUDA Driver Version / Runtime Version 5.5 / 5.0
CUDA Capability Major/Minor version number: 1.1
Total amount of global memory: 512 MBytes (536870912 bytes)
( 8) Multiprocessors x ( 8) CUDA Cores/MP: 64 CUDA Cores
GPU Clock rate: 1680 MHz (1.68 GHz)
Memory Clock rate: 700 Mhz
Memory Bus Width: 256-bit
Max Texture Dimension Size (x,y,z) 1D=(8192), 2D=(65536,32768), 3D=(2048,2048,2048)
Max Layered Texture Size (dim) x layers 1D=(8192) x 512, 2D=(8192,8192) x 512
Total amount of constant memory: 65536 bytes
Total amount of shared memory per block: 16384 bytes
Total number of registers available per block: 8192
Warp size: 32
Maximum number of threads per multiprocessor: 768
Maximum number of threads per block: 512
Maximum sizes of each dimension of a block: 512 x 512 x 64
Maximum sizes of each dimension of a grid: 65535 x 65535 x 1
Maximum memory pitch: 2147483647 bytes
Texture alignment: 256 bytes
Concurrent copy and kernel execution: Yes with 1 copy engine(s)
Run time limit on kernels: Yes
Integrated GPU sharing Host Memory: No
Support host page-locked memory mapping: Yes
Alignment requirement for Surfaces: Yes
Device has ECC support: Disabled
Concurrent kernel execution: No
Device supports Unified Addressing (UVA): No
Device PCI Bus ID / PCI location ID: 1 / 0
質問 3 : グローバル メモリの量を、GPU が持つ共有メモリやその他の種類のメモリの量で表現できますか? スレッドの数はそれと関係がありますか?
質問 4 : ブロックあたりのスレッドの最大数が 512 の場合、ブロックの各次元の最大サイズが 512x512x62 (= 16252628 スレッド) になる可能性はありますか? グリッドの各次元の最大サイズとの相関関係は?
質問 5 : メモリ クロック レートを使用して、1 秒あたりに処理されるスレッドの数を知ることができますか?
更新：
列スレッドに置き換えられた for ループ
__global__ void cudaEuclid( float* A, float* B, float* C, int rows, int cols ){
int r = blockDim.x * blockIdx.x + threadIdx.x; // rows
int c = blockDim.y * blockIdx.y + threadIdx.y; // cols
float x=0;
if(c < cols && r < rows){
x = ( A[c + r*cols] - B[c + r*cols] ) * ( A[c + r*cols] - B[c + r*cols] );
}
C[r] = x;
}
呼び出し:
int threadsPerBlock = 256;
int blocksPerGrid = ceil( (double) numElements / threadsPerBlock);
cudaEuclid<<<blocksPerGrid, threadsPerBlock>>>( d_A, d_B, d_C, rows, cols );