임의의 수의 스레드에 대한 CUDA의 Mersenne Twister

Question

Mersenne Twister (MT) 난수 생성기의 CUDA 구현은 256 및 200 블록/격자의 최대 스레드 수/블록으로 제한됩니다. 즉 최대 스레드 수는 51200입니다 .

따라서

kernel<<<blocksPerGrid, threadsPerBlock>>>(devMTGPStates, ...) 

int blocksPerGrid = (n+threadsPerBlock-1)/threadsPerBlock; 

및과 함께 MT를 사용하는 커널을 실행 할 수 없습니다은 총 스레드 수입니다.

스레드> 51200에 MT를 사용하는 가장 좋은 방법은 무엇입니까?

내 접근 방식은 blocksPerGrid 및 threadsPerBlock에 대한 상수 값을 사용하면됩니다. <<<128,128>>> 및 커널 코드에 다음 사용

__global__ void kernel(curandStateMtgp32 *state, int n, ...) { 

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

    while (id < n) { 

     float x = curand_normal(&state[blockIdx.x]); 
     /* some more calls to curand_normal() followed 
      by the algorithm that works with the data */ 

     id += blockDim.x*gridDim.x; 
    } 
} 

나는 이것이 올바른 방법인지 확실하지 않다 또는 그것을 원하지 않는 방법으로 MT 상태에 영향을 미칠 수 있다면?

감사합니다.

Answer 1

나는 CURAND documentation을 신중하고 철저히 읽는 것이 좋습니다.

MT API는 블록 당 256 스레드를 사용하여 최대 64 블록을 사용하여 숫자를 생성 할 때 가장 효율적입니다. - (예 : 64 개 블록, 256 개 스레드) 세트를,와 사이에이 숫자를 배포

1. 는 단순히 기존의 상태에서 더 많은 번호를 생성 : 당신이 그 이상이 필요한 경우

   , 당신은 다양한 옵션을 가지고 스레드가 필요합니다.

2. 블록 당 하나 이상의 상태를 사용합니다 (그러나 상태 집합 내에서 전체 제한을 초과 할 수는 없으며 단일 블록의 필요성 만 해결합니다).
3. 독립 시드로 여러 MT 생성기 만들기 (따라서 독립적 인 상태 집합).

일반적으로 내가 설명한 커널에는 문제가 나타나지 않으며 위의 선택 1과 대략 일치합니다. 그러나 51200 개의 스레드를 초과 할 수는 없습니다. (귀하의 예는 <<<128, 128>>>이므로 16384 스레드)