如何在 pyCUDA 内核中生成随机数？答案

【问题标题】：How to generate random number inside pyCUDA kernel?如何在 pyCUDA 内核中生成随机数？
【发布时间】：2021-09-08 09:57:08
【问题描述】：

我正在使用 pyCUDA 进行 CUDA 编程。我需要在内核函数中使用随机数。 CURAND 库在其中不起作用（pyCUDA）。由于GPU有很多工作要做，在CPU内部生成随机数然后将它们传输到GPU是行不通的，反而消解了使用GPU的动机。

补充问题：

有没有办法使用 1 个块和 1 个线程在 GPU 上分配内存。
我正在使用多个内核。我需要使用多个 SourceModule 块吗？

【问题讨论】：

我不明白这个问题。 PyCUDA 有一个与 curand 的接口，可以直接用随机值填充设备内存。并且设备端代码可以在内核中使用。
我知道。您正在谈论的界面是CUDA中#include的别名。但是我需要的随机数只有在#include有对应的情况下才能生成。我没有得到你的第二部分，“设备端代码可以在内核中使用，只需一点点努力。”你是在说主机吗？
不，我说的是设备端接口
你能告诉我你是如何在内核中生成随机数的吗？

标签： cuda pycuda

【解决方案1】：

尽管您在问题中断言，PyCUDA 对 CUrand 的支持非常全面。 GPUArray 模块具有使用主机端 API 填充设备内存的直接接口（请注意，在这种情况下，随机生成器在 GPU 上运行）。

也完全可以在 PyCUDA 内核代码中使用来自 CUrand 的设备端 API。在这个用例中，最棘手的部分是为线程生成器状态分配内存。有三种选择——在代码中静态、动态使用主机内存端分配和动态使用设备端内存分配。以下（非常轻微测试）示例说明了后者，正如您在问题中询问的那样：

import numpy as np
import pycuda.autoinit
from pycuda.compiler import SourceModule
from pycuda import gpuarray

code = """
    #include <curand_kernel.h>

    const int nstates = %(NGENERATORS)s;
    __device__ curandState_t* states[nstates];

    __global__ void initkernel(int seed)
    {
        int tidx = threadIdx.x + blockIdx.x * blockDim.x;

        if (tidx < nstates) {
            curandState_t* s = new curandState_t;
            if (s != 0) {
                curand_init(seed, tidx, 0, s);
            }

            states[tidx] = s;
        }
    }

    __global__ void randfillkernel(float *values, int N)
    {
        int tidx = threadIdx.x + blockIdx.x * blockDim.x;

        if (tidx < nstates) {
            curandState_t s = *states[tidx];
            for(int i=tidx; i < N; i += blockDim.x * gridDim.x) {
                values[i] = curand_uniform(&s);
            }
            *states[tidx] = s;
        }
    }
"""

N = 1024
mod = SourceModule(code % { "NGENERATORS" : N }, no_extern_c=True, arch="sm_52")
init_func = mod.get_function("_Z10initkerneli")
fill_func = mod.get_function("_Z14randfillkernelPfi")

seed = np.int32(123456789)
nvalues = 10 * N
init_func(seed, block=(N,1,1), grid=(1,1,1))
gdata = gpuarray.zeros(nvalues, dtype=np.float32)
fill_func(gdata, np.int32(nvalues), block=(N,1,1), grid=(1,1,1))

这里有一个初始化内核，它需要运行一次来为生成器状态分配内存并使用种子初始化它们，然后是使用这些状态的内核。如果要运行大量线程，则需要注意 malloc 堆大小限制，但可以通过 PyCUDA 驱动程序 API 接口进行操作。

【讨论】：

运行 LogicError 时出现此错误：cuModuleLoadDataEx failed: device kernel image is invalid -
@BhaskarDhariyal：显然您需要在SourceModule 实例中设置构建架构以匹配您的GPU 模型。
包含我需要的随机数的 values 数组；用于变量 r1 ， r2 这是语句velocity[i] = X*(velocity[i]+c1*r1*(pBestPos[i] - x[i]) + c2*r2*(lBestIdx[i%d] - x[i])) 的一部分。根据上面的程序，我不能直接访问它，因为给定的语句在不同的内核中。如何从语句的内核访问 values 数组？
也许你误解了 SO 的工作原理。你应该问一个问题。当这个问题得到回答时，继续前进。这个特定的问题/答案不是您自己的个人聊天室或帮助台。您关于随机数的问题已得到解答。你现在问的是一个完全不同的问题。有新问题吗？问一个新问题。
_Z10initkerneli 和 _Z14randfillkernelPfi 这两个名字是怎么回事。你从哪里得到它们，为什么它们与原来的名字不符？

【解决方案2】：

我对接受的答案有一个问题。我们在那里有一个名称混乱，有点讨厌（这些_Z10initkerneli 和_Z14randfillkernelPfi）。为了避免这种情况，我们可以手动将代码包装在extern "C" {...} 子句中。

code = """
    #include <curand_kernel.h>

    const int nstates = %(NGENERATORS)s;
    __device__ curandState_t* states[nstates];
    extern "C" {

    __global__ void initkernel(int seed)
    { .... }

    __global__ void randfillkernel(float *values, int N)
    { .... }
    }
"""

那么代码还是用no_extern_c=True编译：

mod = SourceModule(code % { "NGENERATORS" : N }, no_extern_c=True)

这应该适用于

init_func = mod.get_function("initkernel")
fill_func = mod.get_function("randfillkernel")

希望对您有所帮助。

【讨论】：