pyCUDA 커널에서 난수를 생성하는 방법은 무엇입니까?

Question

저는 CUDA 프로그래밍을 위해 pyCUDA를 사용하고 있습니다. 커널 함수 내에서 난수를 사용해야합니다. CURAND 라이브러리가 내부에서 작동하지 않습니다 (pyCUDA). GPU에서 수행해야 할 작업이 많으므로 CPU에서 임의의 번호를 생성 한 다음 GPU로 전송하면 GPU를 사용하는 동기가 사라지게됩니다.

보충 질문 :

1. 1 개 블록 1 개 실을 사용하여 GPU에 메모리를 할당하는 방법이 있나요.
2. 하나 이상의 커널을 사용하고 있습니다. 여러 개의 SourceModule 블록을 사용해야합니까?

Answer 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 인터페이스를 통해 조작 할 수 있습니다.