コンテキスト:現在、CUDA を適切に使用する方法、特に CURAND を使用して乱数を生成する方法を学んでいます。ここで、コードでコア計算を実行するカーネル内で、乱数が必要なときに直接乱数を生成するのが賢明かもしれないことを学びました。
ドキュメントに従って、少し遊んで、後で自分のニーズに適応できる簡単な実行コードを考え出すことにしました。
ブロック内の 256 の同時スレッドの制限 (および 200 の事前生成されたパラメーター セットのみ) のため、MTGP32 を除外しました。その上、私は double を使いたくないので、デフォルトのジェネレーター (XORWOW) に固執することにしました。
問題:コード内の同じシード値が、128 を超えるブロックあたりのスレッド数に対して異なる数のシーケンスを生成する理由を理解するのに苦労しています (blockSize<129 の場合、すべてが期待どおりに実行されます)。適切なCUDA エラー チェックを行った後、Robert のコメントで示唆されているように、ハードウェアの制限が役割を果たすことはある程度明らかです。さらに、コンパイル時に「-G -g」フラグを使用しないと、「しきい値の問題」が 128 から 384 に上昇します。
質問:正確には何が原因ですか? ロバートはコメントで、「スレッドごとのレジスターの問題かもしれない」と述べています。これは何を意味するのでしょうか？ハードウェアの仕様を見て、この制限がどこにあるのかを簡単に判断する方法はありますか? スレッドごとに乱数を生成しなくても、この問題を回避できますか?
関連する問題がここで議論されているようですが、私のケースには当てはまらないと思います。
私のコード (以下を参照) は、主にこれらの例に触発されました。
コード：
#include <stdio.h>
#include <stdlib.h>
#include <cuda.h>
#include <curand_kernel.h>
#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true){
if (code != cudaSuccess){
fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
if (abort) exit(code);
}
}
__global__ void setup_kernel(curandState *state, int seed, int n){
int id = threadIdx.x + blockIdx.x*blockDim.x;
if(id<n){
curand_init(seed, id, 0, &state[id]);
}
}
__global__ void generate_uniform_kernel(curandState *state, float *result, int n){
int id = threadIdx.x + blockIdx.x*blockDim.x;
float x;
if(id<n){
curandState localState = state[id];
x = curand_uniform(&localState);
state[id] = localState;
result[id] = x;
}
}
int main(int argc, char *argv[]){
curandState *devStates;
float *devResults, *hostResults;
int n = atoi(argv[1]);
int s = atoi(argv[2]);
int blockSize = atoi(argv[3]);
int nBlocks = n/blockSize + (n%blockSize == 0?0:1);
printf("\nn: %d, blockSize: %d, nBlocks: %d, seed: %d\n", n, blockSize, nBlocks, s);
hostResults = (float *)calloc(n, sizeof(float));
cudaMalloc((void **)&devResults, n*sizeof(float));
cudaMalloc((void **)&devStates, n*sizeof(curandState));
setup_kernel<<<nBlocks, blockSize>>>(devStates, s, n);
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk( cudaDeviceSynchronize() );
generate_uniform_kernel<<<nBlocks, blockSize>>>(devStates, devResults, n);
gpuErrchk( cudaPeekAtLastError() );
gpuErrchk( cudaDeviceSynchronize() );
cudaMemcpy(hostResults, devResults, n*sizeof(float), cudaMemcpyDeviceToHost);
for(int i=0; i<n; i++) {
printf("\n%10.13f", hostResults[i]);
}
cudaFree(devStates);
cudaFree(devResults);
free(hostResults);
return 0;
}
2 つのバイナリをコンパイルしました。1 つは "-G -g" デバッグ フラグを使用し、もう 1 つは使用しませんでした。それぞれrng_gen_dとrng_genという名前を付けました。
$ nvcc -lcuda -lcurand -O3 -G -g --ptxas-options=-v rng_gen.cu -o rng_gen_d
ptxas /tmp/tmpxft_00002257_00000000-5_rng_gen.ptx, line 2143; warning : Double is not supported. Demoting to float
ptxas info : 77696 bytes gmem, 72 bytes cmem[0], 32 bytes cmem[14]
ptxas info : Compiling entry function '_Z12setup_kernelP17curandStateXORWOWii' for 'sm_10'
ptxas info : Used 43 registers, 32 bytes smem, 72 bytes cmem[1], 6480 bytes lmem
ptxas info : Compiling entry function '_Z23generate_uniform_kernelP17curandStateXORWOWPfi' for 'sm_10'
ptxas info : Used 10 registers, 36 bytes smem, 40 bytes cmem[1], 48 bytes lmem
$ nvcc -lcuda -lcurand -O3 --ptxas-options=-v rng_gen.cu -o rng_gen
ptxas /tmp/tmpxft_00002b73_00000000-5_rng_gen.ptx, line 533; warning : Double is not supported. Demoting to float
ptxas info : 77696 bytes gmem, 72 bytes cmem[0], 32 bytes cmem[14]
ptxas info : Compiling entry function '_Z12setup_kernelP17curandStateXORWOWii' for 'sm_10'
ptxas info : Used 20 registers, 32 bytes smem, 48 bytes cmem[1], 6440 bytes lmem
ptxas info : Compiling entry function '_Z23generate_uniform_kernelP17curandStateXORWOWPfi' for 'sm_10'
ptxas info : Used 19 registers, 36 bytes smem, 4 bytes cmem[1]
まず、コンパイル時に奇妙な警告メッセージが表示されます (上記を参照)。
ptxas /tmp/tmpxft_00002b31_00000000-5_rng_gen.ptx, line 2143; warning : Double is not supported. Demoting to float
一部のデバッグでは、この警告の原因となっている行は次のとおりであることが示されました。
curandState localState = state[id];
double が宣言されていないため、これを解決する方法が正確にはわかりません (または、これを解決する必要がある場合でも)。
さて、私が直面している（実際の）問題の例：
$ ./rng_gen_d 5 314 127
n: 5, blockSize: 127, nBlocks: 1, seed: 314
0.9151657223701
0.3925153017044
0.7007563710213
0.8806988000870
0.5301177501678
$ ./rng_gen_d 5 314 128
n: 5, blockSize: 128, nBlocks: 1, seed: 314
0.9151657223701
0.3925153017044
0.7007563710213
0.8806988000870
0.5301177501678
$ ./rng_gen_d 5 314 129
n: 5, blockSize: 129, nBlocks: 1, seed: 314
GPUassert: too many resources requested for launch rng_gen.cu 54
54行目はsetup_kernel()の直後のgpuErrchk()です。
他のバイナリ (コンパイル時に「-G -g」フラグなし) では、「トラブルのしきい値」が 384 に引き上げられます。
$ ./rng_gen 5 314 129
n: 5, blockSize: 129, nBlocks: 1, seed: 314
0.9151657223701
0.3925153017044
0.7007563710213
0.8806988000870
0.5301177501678
$ ./rng_gen 5 314 384
n: 5, blockSize: 384, nBlocks: 1, seed: 314
0.9151657223701
0.3925153017044
0.7007563710213
0.8806988000870
0.5301177501678
$ ./rng_gen 5 314 385
n: 5, blockSize: 385, nBlocks: 1, seed: 314
GPUassert: too many resources requested for launch rng_gen.cu 54
最後に、これがこの予備テストに使用しているハードウェアに何らかの関連がある場合 (プロジェクトは後ではるかに強力なマシンで開始されます)、使用しているカードの仕様は次のとおりです。
./deviceQuery Starting...
CUDA Device Query (Runtime API) version (CUDART static linking)
Detected 1 CUDA Capable device(s)
Device 0: "Quadro NVS 160M"
CUDA Driver Version / Runtime Version 5.5 / 5.5
CUDA Capability Major/Minor version number: 1.1
Total amount of global memory: 256 MBytes (268107776 bytes)
( 1) Multiprocessors, ( 8) CUDA Cores/MP: 8 CUDA Cores
GPU Clock rate: 1450 MHz (1.45 GHz)
Memory Clock rate: 702 Mhz
Memory Bus Width: 64-bit
Maximum Texture Dimension Size (x,y,z) 1D=(8192), 2D=(65536, 32768), 3D=(2048, 2048, 2048)
Maximum Layered 1D Texture Size, (num) layers 1D=(8192), 512 layers
Maximum Layered 2D Texture Size, (num) layers 2D=(8192, 8192), 512 layers
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
Max dimension size of a thread block (x,y,z): (512, 512, 64)
Max dimension size of a grid size (x,y,z): (65535, 65535, 1)
Maximum memory pitch: 2147483647 bytes
Texture alignment: 256 bytes
Concurrent copy and kernel execution: No with 0 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
Device supports Unified Addressing (UVA): No
Device PCI Bus ID / PCI location ID: 1 / 0
Compute Mode:
< Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >
deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 5.5, CUDA Runtime Version = 5.5, NumDevs = 1, Device0 = Quadro NVS 160M
Result = PASS
で、これです。この問題に関するガイダンスは大歓迎です。ありがとう！
編集：
1) Robert の提案に従って、適切なcuda エラー チェックを追加しました。
2)とにかく役に立たなかった cudaMemset 行を削除しました。
3)「-G -g」フラグなしでコードをコンパイルして実行しました。
4)それに応じて出力を更新しました。