0

我的动机:我正在使用一种算法来模拟人口动态,并且我希望使用 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
4

2 回答 2

5

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

生成 n 个种子值有什么好处,将它们存储在设备全局内存的数组中,然后将它们馈送到内核,内核依次生成几个随机数以供使用,而不是生成 2n 个随机数,将它们存储在设备全局内存,然后将它们直接提供给需要使用它们的内核?

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

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

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

同样,CURAND 的主要优势在于性能。如果随机数生成只是应用程序总体计算成本的一小部分,那么您使用哪种方法可能并不重要,或者即使您完全使用 CURAND(例如,代替基于 CPU 的普通 RNG 方法)。

于 2013-10-30T20:05:34.147 回答
3

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

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

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

于 2013-10-30T20:05:58.663 回答