Передача CUDA Random Generator State по ссылке

Question

Правильно ли следующий код при передаче состояния генератора случайных чисел (набор инструментов CUDA 3.2 curand.lib) по ссылке в функциях CalculateValue (curandState * localStat) и GetExponential (curandState * localState)?

Спасибо

__device__ double GetExponential(curandState *localState) { 
    double u1 = curand_uniform_double(localState); } 


__device__  double CalculateValue(curandState *localStat)  { 
  double x = GetExponential(localState);  
  return x; } 


__global__ void RunMonteCarloKernel(curandState *state, double *results) { 
    int i = threadIdx.x + blockIdx.x * blockDim.x; 

    /* Copy state to local memory for efficiency */ 
    curandState localState = state[threadIdx.x + blockIdx.x * blockDim.x];    

    results[i] = CalculateValue(&localState); 

    /* Copy state back to global memory */ 
    state[threadIdx.x + blockIdx.x * blockDim.x] = localState; }

__global__ void setup_kernel(curandState *state) { 
    int i = threadIdx.x + blockIdx.x * blockDim.x; 

    /* Each thread gets different seed, a different sequence number, no offset */ 
    curand_init(i, i, 0, &state[i]); } 

int main(void) { 
    double *devResults; 
    curandState *devStates; 

    /* Allocate space for prng states on device */ 
    CUDA_CALL(cudaMalloc((void **)&devStates, totalThreads * sizeof(curandState))); 

    /* Setup prng states */ 
   setup_kernel<<<totalBlocks, threadsPerBlock>>>(devStates); 

    for(int i=0; i< 1000; i++) 
    { 
            RunMonteCarloKernel(devStates, devResults); 
    } }

Answer 1

Есть ли проблема? Смотрится нормально.

Возможно, вы захотите проверить пример EstimatePiInlineP, который находится в каталоге MonteCarloCURAND в 3.2 SDK. Он использует стиль передачи C ++ по ссылке, чтобы избежать получения адреса локальной переменной. Вам нужно будет сохранить состояние обратно в память в конце ядра (как вы делаете в своем коде).

Передача по ссылке на C ++ может помочь компилятору, ясно продемонстрировав, что функция может работать с данными непосредственно в исходных регистрах. Получение адреса локального массива в графическом процессоре может отрицательно сказаться на производительности, если компилятор не может быть уверен, что все потоки обрабатывают указатель одинаково (т. Е. Идентичные операции над указателем), и в этом случае он будет проливать массив в локальную память. Это будет работать, но может быть медленнее.