CURAND - 设备 API 和种子

Question

上下文:我目前正在学习如何正确使用 CUDA，特别是如何使用 CURAND 生成随机数。我学会了here当我需要随机数时，在我的代码中执行核心计算的内核中直接生成随机数可能是明智的。

在 documentation 之后，我决定尝试一下，尝试编写一段简单的运行代码，稍后我可以根据自己的需要进行调整。

我排除了 MTGP32，因为一个 block 中有 256 个并发线程的限制(并且只有 200 个预生成的参数集)。此外，我不想使用 double ，所以我决定坚持使用默认生成器 (XORWOW)。

问题:我很难理解为什么我的代码中的相同种子值会为大于 128 的每个 block 的多个线程生成不同的数字序列(当 blockSize<129 时，所有内容按我的预期运行)。做适当的事后 CUDA error checking ，正如罗​​伯特在他的评论中所建议的那样，很明显硬件限制起了作用。此外，在编译时不使用“-G -g”标志会将“阈值问题”从 128 提高到 384。

问题:到底是什么原因造成的？罗伯特在他的评论中写到“这可能是每个线程的寄存器问题”。这是什么意思？有没有一种简单的方法可以查看硬件规范并说明此限制在哪里？我可以在不必为每个线程生成更多随机数的情况下解决这个问题吗？

似乎已经讨论了一个相关问题here但我认为它不适用于我的情况。

我的代码(见下文)主要受 these examples 启发.

代码:

    #include <stdio.h>
    #include <stdlib.h>
    #include <cuda.h>
    #include <curand_kernel.h>

    #define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }
    inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true){
        if (code != cudaSuccess){ 
           fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line);
           if (abort) exit(code);
        }
    }

    __global__ void setup_kernel(curandState *state, int seed, int n){

        int id = threadIdx.x + blockIdx.x*blockDim.x;

        if(id<n){
            curand_init(seed, id, 0, &state[id]);
        }
    }

    __global__ void generate_uniform_kernel(curandState *state, float *result, int n){

        int id = threadIdx.x + blockIdx.x*blockDim.x;
        float x;

        if(id<n){
            curandState localState = state[id];
            x = curand_uniform(&localState);
            state[id] = localState;
            result[id] = x; 
        }
    }

    int main(int argc, char *argv[]){

        curandState *devStates;
        float *devResults, *hostResults;

        int n = atoi(argv[1]);
        int s = atoi(argv[2]);
        int blockSize = atoi(argv[3]);

        int nBlocks = n/blockSize + (n%blockSize == 0?0:1);

        printf("\nn: %d, blockSize: %d, nBlocks: %d, seed: %d\n", n, blockSize, nBlocks, s);

        hostResults = (float *)calloc(n, sizeof(float));
        cudaMalloc((void **)&devResults, n*sizeof(float));

        cudaMalloc((void **)&devStates, n*sizeof(curandState));
        setup_kernel<<<nBlocks, blockSize>>>(devStates, s, n);
        gpuErrchk( cudaPeekAtLastError() );
        gpuErrchk( cudaDeviceSynchronize() );

        generate_uniform_kernel<<<nBlocks, blockSize>>>(devStates, devResults, n);
        gpuErrchk( cudaPeekAtLastError() );
        gpuErrchk( cudaDeviceSynchronize() );

        cudaMemcpy(hostResults, devResults, n*sizeof(float), cudaMemcpyDeviceToHost);

        for(int i=0; i<n; i++) {
            printf("\n%10.13f", hostResults[i]);
        }

        cudaFree(devStates);
        cudaFree(devResults);
        free(hostResults);

        return 0;
    }

我编译了两个二进制文件，一个使用“-G -g”调试标志，另一个没有。我分别将它们命名为 rng_gen_d 和 rng_gen:

     $ nvcc -lcuda -lcurand -O3 -G -g --ptxas-options=-v rng_gen.cu -o rng_gen_d
    ptxas /tmp/tmpxft_00002257_00000000-5_rng_gen.ptx, line 2143; warning : Double is not supported. Demoting to float
    ptxas info    : 77696 bytes gmem, 72 bytes cmem[0], 32 bytes cmem[14]
    ptxas info    : Compiling entry function '_Z12setup_kernelP17curandStateXORWOWii' for 'sm_10'
    ptxas info    : Used 43 registers, 32 bytes smem, 72 bytes cmem[1], 6480 bytes lmem
    ptxas info    : Compiling entry function '_Z23generate_uniform_kernelP17curandStateXORWOWPfi' for 'sm_10'
    ptxas info    : Used 10 registers, 36 bytes smem, 40 bytes cmem[1], 48 bytes lmem

     $ nvcc -lcuda -lcurand -O3 --ptxas-options=-v rng_gen.cu -o rng_gen
    ptxas /tmp/tmpxft_00002b73_00000000-5_rng_gen.ptx, line 533; warning : Double is not supported. Demoting to float
    ptxas info    : 77696 bytes gmem, 72 bytes cmem[0], 32 bytes cmem[14]
    ptxas info    : Compiling entry function '_Z12setup_kernelP17curandStateXORWOWii' for 'sm_10'
    ptxas info    : Used 20 registers, 32 bytes smem, 48 bytes cmem[1], 6440 bytes lmem
    ptxas info    : Compiling entry function '_Z23generate_uniform_kernelP17curandStateXORWOWPfi' for 'sm_10'
    ptxas info    : Used 19 registers, 36 bytes smem, 4 bytes cmem[1]

首先，在编译时有一条奇怪的警告消息(见上文):

    ptxas /tmp/tmpxft_00002b31_00000000-5_rng_gen.ptx, line 2143; warning : Double is not supported. Demoting to float

一些调试显示导致此警告的行是:

    curandState localState = state[id];

没有声明 double ，所以我不知道如何解决这个问题(或者即使这需要解决)。

现在，我面临的(实际)问题的一个例子:

     $ ./rng_gen_d 5 314 127

    n: 5, blockSize: 127, nBlocks: 1, seed: 314

    0.9151657223701
    0.3925153017044
    0.7007563710213
    0.8806988000870
    0.5301177501678

     $ ./rng_gen_d 5 314 128

    n: 5, blockSize: 128, nBlocks: 1, seed: 314

    0.9151657223701
    0.3925153017044
    0.7007563710213
    0.8806988000870
    0.5301177501678

     $ ./rng_gen_d 5 314 129

    n: 5, blockSize: 129, nBlocks: 1, seed: 314
    GPUassert: too many resources requested for launch rng_gen.cu 54

第 54 行是在 setup_kernel() 之后的 gpuErrchk()。

对于其他二进制文件(编译时没有“-G -g”标志)，“故障阈值”提高到 384:

     $ ./rng_gen 5 314 129

    n: 5, blockSize: 129, nBlocks: 1, seed: 314

    0.9151657223701
    0.3925153017044
    0.7007563710213
    0.8806988000870
    0.5301177501678

     $ ./rng_gen 5 314 384 

    n: 5, blockSize: 384, nBlocks: 1, seed: 314

    0.9151657223701
    0.3925153017044
    0.7007563710213
    0.8806988000870
    0.5301177501678

     $ ./rng_gen 5 314 385

    n: 5, blockSize: 385, nBlocks: 1, seed: 314
    GPUassert: too many resources requested for launch rng_gen.cu 54

最后，如果这与我用于此初步测试的硬件有某种关系(该项目稍后将在功能更强大的机器上启动)，以下是我使用的卡的规范:

    ./deviceQuery Starting...

     CUDA Device Query (Runtime API) version (CUDART static linking)

    Detected 1 CUDA Capable device(s)

    Device 0: "Quadro NVS 160M"
      CUDA Driver Version / Runtime Version          5.5 / 5.5
      CUDA Capability Major/Minor version number:    1.1
      Total amount of global memory:                 256 MBytes (268107776 bytes)
      ( 1) Multiprocessors, (  8) CUDA Cores/MP:     8 CUDA Cores
      GPU Clock rate:                                1450 MHz (1.45 GHz)
      Memory Clock rate:                             702 Mhz
      Memory Bus Width:                              64-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:       Yes
      Alignment requirement for Surfaces:            Yes
      Device has ECC support:                        Disabled
      Device supports Unified Addressing (UVA):      No
      Device PCI Bus ID / PCI location ID:           1 / 0
      Compute Mode:
         < Default (multiple host threads can use ::cudaSetDevice() with device simultaneously) >

    deviceQuery, CUDA Driver = CUDART, CUDA Driver Version = 5.5, CUDA Runtime Version = 5.5, NumDevs = 1, Device0 = Quadro NVS 160M
    Result = PASS

就是这样。欢迎就此事提供任何指导。谢谢!

编辑:

1) 添加了正确的 cuda error checking ，正如罗​​伯特所建议的那样。

2) 删除了 cudaMemset 行，反正没用。

3) 在没有“-G -g”标志的情况下编译并运行代码。

4)相应地更新了输出。

Answer 1

首先，当您在使用 CUDA 代码时遇到问题时，始终建议您执行正确的操作 cuda error checking .它会消除一定程度的头痛，可能会节省您一些时间，并且肯定会提高人们在此类网站上为您提供帮助的能力。

现在您发现了每个线程的寄存器问题。编译器在生成代码时会出于各种目的使用寄存器。每个线程都需要这些寄存器的补充来运行它的线程代码。当您尝试启动内核时，必须满足的要求之一是每个线程所需的寄存器数量乘以启动中请求的线程数量必须小于每个 block 可用的寄存器总数。请注意，每个线程所需的寄存器数量可能必须四舍五入到某个粒度分配增量。另请注意，请求的线程数通常会四舍五入到下一个更高的增量 32(如果不能被 32 整除)，因为线程是在 32 的 warps 中启动的。还要注意最大寄存器每个 block 因计算能力而异，如您所示，可以通过 deviceQuery 示例检查此数量。正如您所发现的，某些命令行开关(如 -G)会影响 nvcc 使用寄存器的方式。

要提前通知这些类型的资源问题，您可以使用额外的命令行开关编译您的代码:

nvcc -arch=sm_11 -Xptxas=-v -o mycode mycode.cu

-Xptxas=-v 开关将生成 ptxas 汇编程序(将中间 ptx 代码转换为 sass 汇编代码，即机器代码)的资源使用输出，包括每个线程所需的寄存器。请注意，在这种情况下，输出将按内核传递，因为每个内核可能都有自己的要求。您可以在 documentation 中获得有关 nvcc 编译器的更多信息。 .

作为一种粗略的解决方法，您可以在编译时指定一个开关，以将所有内核编译限制为最大寄存器使用数:

nvcc -arch=sm_11 -Xptxas=-v -maxrregcount=16 -o mycode mycode.cu

这将限制每个内核每个线程使用不超过 16 个寄存器。当乘以 512(cc1.x 设备的每个 block 的线程硬件限制)时，得到的值为 8192，这是您设备的每个线程 block 的总寄存器的硬件限制。

然而，上述方法很粗糙，因为它对程序中的所有内核应用了相同的限制。如果您想针对每次内核启动对其进行定制(例如，如果您的程序中的不同内核启动不同数量的线程)，您可以使用启动边界方法，这在 here 中有描述。 .