pyCUDAカーネル内で乱数を生成する方法は？

Question

私はpyCUDAをCUDAプログラミングに使用しています。カーネル関数内で乱数を使う必要があります。 CURANDライブラリはその内部で動作しません（pyCUDA）。 GPUには多くの作業が必要なので、CPU内で乱数を生成してからGPUに転送するとGPUを使用する動機が失われます。

補足質問：

1. 1つのブロックと1つのスレッドを使用して、GPU上のメモリを割り当てる方法はあります。
2. 私は複数のカーネルを使用しています。複数のSourceModuleブロックを使用する必要がありますか？

Answer 1

あなたの疑問に主張しているものの、PyCUDAはCUrandをかなり包括的にサポートしています。 GPUArrayモジュールには、ホストサイドAPIを使用してデバイスメモリを埋めるための直接インタフェースがあります（この場合、ランダムジェネレータはGPU上で実行されます）。

PyCUDAカーネルコードにCUrandからデバイス側のAPIを使用することも完全に可能です。この使用例では、もっともトリッキーな部分はスレッドジェネレータの状態にメモリを割り当てています。コード内に静的に、ホストメモリ側の割り当てを動的に使用すること、およびデバイス側のメモリ割り当てを動的に使用するという3つの選択肢があります。 （非常に軽くテストした）次の例では、あなたの質問にそれについて尋ねたとして見て、後者を示しています

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)) 

ここで発電状態にメモリを割り当て、初期化するために1回実行する必要のある初期化カーネルがありますそれらに種子、そしてそれらの状態を使用するカーネルがあります。多くのスレッドを実行したい場合は、mallocヒープサイズの制限に注意する必要がありますが、それらはPyCUDAドライバAPIインタフェースを介して操作できます。