我的动机:我正在使用一种算法来模拟人口动态,并且我希望使用 CUDA 以便能够在数值模拟中考虑大量节点。虽然这是我第一次在 GPU 上运行代码,但到目前为止结果看起来很有希望。
背景:我需要考虑随机噪声,它在我打算研究的复杂系统的演变中起着至关重要的作用。据我了解,与 CPU 上的类似操作相比,CUDA 中的随机数生成可能相当麻烦。在文档中,我看到必须存储 RNG 的状态并继续将其提供给需要(生成和)使用随机数的内核(全局函数)。我发现这些示例很有启发性,也许您还推荐我阅读有关此内容的其他内容?
问题:生成 n 个种子值,将它们存储在设备全局内存上的数组中,然后将它们提供给内核,内核依次生成几个随机数供使用,而不是生成 2n 个随机数,存储它们在设备全局内存中,然后将它们直接提供给需要使用它们的内核?我必须在这里遗漏一些真正重要的东西,因为在我看来,在第二种情况下肯定会节省资源(示例中从未使用过)。对于生成的数字的分布,似乎也更安全。
我的代码相当长,但我试图做一个简短的例子来说明我需要什么。这里是:
我的代码:
#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