私の動機:アルゴリズムを使用して人口動態をモデル化し、数値シミュレーションで多数のノードを考慮できるようにするために CUDA を使用したいと考えています。GPU でコードを実行するのはこれが初めてですが、これまでのところ結果は有望に見えます。
コンテキスト:研究対象の複雑なシステムの進化において重要な役割を果たす確率的ノイズを説明する必要があります。私が理解している限り、CUDA での乱数生成は、CPU での同様の操作に比べてかなり面倒です。ドキュメントでは、RNG の状態を保存し、乱数を (生成して) 使用する必要があるカーネル (グローバル関数) にこれを供給し続ける必要があることがわかります。これらの例は非常に啓発的であることがわかりました。おそらく、これについて読むことをお勧めする何かがありますか?
質問: n 個のシード値を生成し、それらをデバイスのグローバル メモリ上の配列に格納してからカーネルにフィードすると、2n 個の乱数を生成して格納するのではなく、使用する乱数を生成する利点は何ですか?それらをデバイスのグローバルメモリに入れ、それらを使用する必要があるカーネルに直接供給しますか? 2 番目のケース (例では使用されていません) でリソースを節約できるように見えるので、ここで本当に重要な何かが欠けているに違いありません。また、生成された数値の分布に関しては、より安全であるように思われます。
私のコードはかなり長いですが、必要なものの短い例を作成しようとしました。ここにあります：
私のコード:
#include <cstdlib>
#include <stdio.h>
#include <cuda.h>
#include <curand.h>
#include <math.h>
__global__ void update (int n, float *A, float *B, float p, float q, float *rand){
int idx = blockIdx.x*blockDim.x + threadIdx.x;
int n_max=n*n;
int i, j;
i=idx/n; //col
j=idx-i*n; //row
float status;
//A, B symmetric
//diagonal untouched, only need 2 random numbers per thread
//i.e. n*(n-1) random numbers in total
int idx_rand = (2*n-1-i)*i/2+j-1-i;
if(idx<n_max && j>i){
if(rand[idx_rand]<p){
status=A[idx];
if(status==1){
if(rand[idx_rand+n*(n-1)/2] < q){
B[idx]=-1.0f;
B[i+n*j]=-1.0f;
}
}
else if(status==0){
if(rand[idx_rand+n*(n-1)/2] < q){
B[idx]=1.0f;
B[i+n*j]=1.0f;
}
}
}
}
}
__global__ void fill(float *A, int n, float num){
int idx = blockIdx.x*blockDim.x + threadIdx.x;
if(idx<n){
A[idx]=num;
}
}
void swap(float** a, float** b) {
float* temp = *a;
*a = *b;
*b = temp;
}
int main(int argc, char* argv[]){
int t, n, t_max, seed;
seed = atoi(argv[1]);
n = atoi(argv[2]);
t_max = atoi(argv[3]);
int blockSize = 256;
int nBlocks = n*n/blockSize + ((n*n)%blockSize == 0?0:1);
curandGenerator_t prng;
curandCreateGenerator(&prng, CURAND_RNG_PSEUDO_DEFAULT);
curandSetPseudoRandomGeneratorSeed(prng, (unsigned long long) seed);
float *h_A = (float *)malloc(n * n * sizeof(float));
float *h_B = (float *)malloc(n * n * sizeof(float));
float *d_A, *d_B, *d_rand;
cudaMalloc(&d_A, n * n * sizeof(float));
cudaMalloc(&d_B, n * n * sizeof(float));
cudaMalloc(&d_rand, n * (n-1) * sizeof(float));
fill <<< nBlocks, blockSize >>> (d_A, n*n, 0.0f);
fill <<< nBlocks, blockSize >>> (d_B, n*n, 0.0f);
for(t=1; t<t_max+1; t++){
//generate random numbers
curandGenerateUniform(prng, d_rand, n*(n-1));
//update B
update <<< nBlocks, blockSize >>> (n, d_A, d_B, 0.5f, 0.5f, d_rand);
//do more stuff
swap(&d_A, &d_B);
}
cudaMemcpy(h_A, d_A, n*n*sizeof(float),cudaMemcpyDeviceToHost);
//print stuff
curandDestroyGenerator(prng);
cudaFree(d_A);
cudaFree(d_B);
cudaFree(d_rand);
free(h_A);
free(h_B);
return 0;
}
何が問題なのか (および修正方法のヒント) を教えてください。専門家が、考えられるすべてのパフォーマンス調整の後で、最善のシナリオで (実行時間で) どれだけ節約できるかを教えてくれれば、それは素晴らしいことです。・「勉強時間」の面でのメリットは非常に重要です。
これで終わりです。読んでくれてありがとう！
記録のために、私のハードウェア仕様は以下のとおりです。ただし、ある時点でこれに Amazon EC2 を使用する予定です。
私の（現在の）ハードウェア：
Device 0: "GeForce 8800 GTX"
CUDA Driver Version / Runtime Version 5.5 / 5.5
CUDA Capability Major/Minor version number: 1.0
Total amount of global memory: 768 MBytes (804978688 bytes)
(16) Multiprocessors, ( 8) CUDA Cores/MP: 128 CUDA Cores
GPU Clock rate: 1350 MHz (1.35 GHz)
Memory Clock rate: 900 Mhz
Memory Bus Width: 384-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: No
Alignment requirement for Surfaces: Yes
Device has ECC support: Disabled
Device supports Unified Addressing (UVA): No
Device PCI Bus ID / PCI location ID: 7 / 0