CURAND和内核,在哪里生成?

我的动机:我使用算法来模拟种群动态,我希望使用CUDA,以便能够在数值模拟中考虑大量节点。 虽然这是我第一次在GPU上运行代码,但到目前为止结果看起来很有希望。

背景:我需要考虑随机噪声,它在我要研究的复杂系统的演化中起着至关重要的作用。 据我所知,与CPU上的类似操作相比,CUDA中的随机数生成可能非常麻烦。 在文档中,我看到必须存储RNG的状态并继续将其提供给需要(生成和)使用随机数的内核(全局函数)。 我发现这些例子很有启发性,也许还有其他什么你建议我阅读这个?

问题:生成n个种子值的优点是什么,将它们存储在设备全局内存中的数组中,然后将它们提供给内核,内核又生成一些随机数,而不是生成2n个随机数,存储它们在设备全局内存中,然后直接将它们提供给需要使用它们的内核? 我必须在这里遗漏一些真正重要的东西,因为它确实让我觉得在第二种情况下可以节省资源(在示例中从未使用过)。 似乎人们对生成的数字的分布更安全。

我的代码相当长,但我试着做一个我需要的简短例子。 这里是:

我的代码:

#include  #include  #include  #include  #include  __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 //ie n*(n-1) random numbers in total int idx_rand = (2*n-1-i)*i/2+j-1-i; if(idxi){ 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 <<>> (d_A, n*n, 0.0f); fill <<>> (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 <<>> (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 

通常,随机数生成是适合在GPU上并行化的过程,因此在许多情况下可以利用GPU来比在CPU上生成数字更快地生成数字。 这是使用像CURAND这样的API /库的主要动机。

生成n个种子值的优点是什么,将它们存储在设备全局内存中的数组中,然后将它们提供给内核,然后内核生成一些随机数,而不是生成2n个随机数,将它们存储在设备全局内存,然后直接将它们提供给需要使用它们的内核?

两者都是有效的方法,可以利用GPU加速:预先生成数字并存储它们,或者动态生成它们。

您可能想要考虑一种方法而不是另一种方法的一些原因是:

  1. 预先生成数字只有在您知道需要多少(或上限为多少)时才有用。 如果您的算法变化很大(可能存在不同的数据集),则可能难以确定。
  2. 存储生成的数字可能是个问题。 对于某些类型的算法(例如蒙特卡罗模拟),可能需要生成如此多的随机数,预先执行并将它们全部存储可能都是禁止的。 在这些情况下,动态生成它们可能会使您无需大量随机数存储。
  3. 通过动态生成数字可以获得略微更好的机器利用率,以避免额外内核调用的开销在它们被使用之前生成数字。

同样,CURAND的主要好处是性能。 如果随机数生成只是应用程序整体计算成本的一小部分,那么使用哪种方法或者根本不使用CURAND(例如代替基于普通CPU的RNG方法)可能无关紧要。

当您使用代码中所示的cuRand主机API时,您必须首先生成一些随机数,将它们存储在全局内存中,然后将它们加载到工作内核中。 存储和装载将花费一些时间,因此可能会损害性能。

但是,如果您使用链接中显示的设备API,则可以节省用于存储和加载这些随机数的带宽。

在这两种情况下,您都必须加载和存储相同数量的RNG状态数据。