CUDA random number geration inside kernel

CUDA kernel에서는 host function을 사용하지 못한다.

즉 C++에서 rand(), time()과 같은 함수들을 사용할 수 없다ㅠㅠ

그래서 CUDA에는 cuRAND라는 library를 제공한다.(http://docs.nvidia.com/cuda/curand)

원하는 Integer 범위에서 random값을 얻고 싶지만 이 문서는 너무 길고 어렵다. ㅠㅠ

그래서 이 포스팅 에서는 원하는 Integer 범위에서 값을 얻어오는 방법을 소개 하도록 하겠다.

말로 대충 설명해 보자면 curand_uniform 함수를 이용하여 Integer를 생성 하는데,

curand_uniform은 위의 reference를 참고 하자면 0.0~1.0까지의 float를 반환 하는데 0은 포함하지 않고 1은 포함한다고 문서에 나와있다.(http://docs.nvidia.com/cuda/curand/device-api-overview.html#bit-generation-1)

그리하여 우리는 다음과 같은 과정으로 random integer를 생성 할 수 있도록 설계 해 보자.

1. curand header파일 include
   #include <curand_kernel.h>

2. getRand 함수 작성(A부터 B-1까지 범위에서 한개의 정수를 반환시키는 함수)

__device__ int getRand(curandState *s, int A, int B) {                 float rand_int = curand_uniform(s);                 rand_int = rand_int * (B-A) + A;                return rand_int; }

3. curand_init을 통해 getRand를 사용할 준비를 해보자.
먼저 reference를 참고 해 보면 curand_init시 seed값을 바꿔줘야 하는데, 평소에는 time값으로 했지만 여기서는 사용할 수 없다. 떄문에 각각의 block, thread마다 다른 값의 seed를 주기 위해 다음과 같이 id를 설계 했다.

__global__ void simulationKernel() {               int id = threadIdx.x + blockDim.x * blockIdx.x;               unsigned int seed = id;               curandState s;              curand_init(seed, 0, 0, &s); // 나머지 argument는 reference를 참조하자             int randNum = getRand(&s, 0, 10); // 0~9 까지 난수 생성 }

4. 함정
seed를 threadIdx.x + blockDim.x * blockIdx.x 로 주었기 떄문에 어떠한 커널을 실행 해도 해당 thread는 항상 똑같은 패턴의 랜덤 값을 얻어 오게 된다는 함정이 있다.

간단한건데 한글 자료가 없어서 고생했다. ㅠㅠ

끗끗끗