CURAND 和内核，在哪里生成？

Question

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

背景:我需要考虑随机噪声，它在我打算研究的复杂系统的演化过程中起着至关重要的作用。据我了解，与 CPU 上的类似操作相比，CUDA 中的随机数生成可能相当麻烦。在文档中，我看到必须存储 RNG 的状态，并将其继续提供给需要(生成和)使用随机数的内核(全局函数)。我找到了 these examples非常有启发性，也许您还推荐我阅读有关此内容的其他内容？

问题:生成 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

Answer 1

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

What is the advantage of generating n seed values, store them in an array on the device global memory, and then feeding them to the kernel which in turn generates a couple of random numbers to use, opposed to generating 2n random numbers, storing them in the device global memory, and then feeding them directly to the kernel which needs to use them?

两者都是有效的方法并且可以利用 GPU 加速:要么预先生成数字并存储它们，要么即时生成它们。

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

1. 仅当您知道需要多少(或多少的上限)时，预先生成数字才有用。如果您的算法变化很大(可能存在不同的数据集)，这可能很难确定。
2. 存储生成的数字可能是个问题。对于某些类型的算法(例如蒙特卡洛模拟)，可能需要生成如此多的随机数，以至于预先进行并存储它们可能会让人望而却步。在这些情况下，即时生成它们可能会让您绕过对大量随机数存储的需求。
3. 通过动态生成数字可能会稍微提高机器利用率，从而避免在使用数字之前调用额外的内核来生成数字的开销。

同样，CURAND 的主要优势在于性能。如果随机数生成只占应用程序总计算成本的一小部分，那么使用哪种方法可能无关紧要，甚至根本不使用 CURAND(例如，代替普通的基于 CPU 的 RNG 方法)。