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 /庫的主要動機。

生成n種子值的好處是，將它們存儲在設備全局內存的數組中，然後將它們送入內核，然後生成一對隨機數來使用，反對生成2n隨機數字，將它們存儲在設備的全局內存中，然後直接將它們提供給需要使用它們的內核？

這兩種方法都是有效的方法，可以利用GPU加速：預先生成數字並存儲它們，或者在飛行中生成它們。

你們當中有些人可能要考慮通過另一種方式的原因是：如果你知道

1. 生成的數字前面是唯一有用的多少（或上限多少），你會需要。如果你的算法變化很大（可能存在不同的數據集），這可能很難確定。
2. 存儲生成的數字可能是一個問題。對於某些類型的算法（例如蒙特卡羅模擬），可能需要生成如此多的隨機數，因此預先完成並將它們全部存儲起來可能會令人望而卻步。在這些情況下，動態生成它們可能會讓您繞過大量隨機數存儲的需要。
3. 通過動態生成數字可以獲得略高的機器利用率，從而避免額外的內核調用在使用之前生成數字的開銷。

同樣，CURAND的主要優點是性能。如果隨機數的產生只是應用程序整體計算成本的一小部分，那麼使用哪種方法或者甚至根本就不使用CURAND（例如，代替基於CPU的普通CPU）。

Answer 2

當您的代碼中顯示使用cuRand主機API時，您必須先生成一些隨機數，將它們存儲在全局mem中，然後將它們加載到工作內核中。存儲和加載會花費一些時間，因此可能會損害性能。

但是，如果您使用設備API（如鏈接所示），則可以節省用於存儲和加載這些隨機數的帶寬。

在這兩種情況下，您都必須加載和存儲相同數量的RNG狀態數據。