Блог

Я пытаюсь подсчитать, сколько раз curand_uniform() возвращает 1.0. Однако, похоже, я не могу заставить следующий код работать для меня:

 #include <stdio.h>
#include <stdlib.h>  
#include <thrust/device_vector.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include <curand_kernel.h>
using namespace std;

__global__
void counts(int length, int *sum, curandStatePhilox4_32_10_t*  state) {
  int tempsum = int(0);
  int i = blockIdx.x * blockDim.x   threadIdx.x;
  curandStatePhilox4_32_10_t localState  =  state[i];
  for(; i < length; i  = blockDim.x * gridDim.x) {
    double thisnum = curand_uniform( amp;localState );
    if ( thisnum == 1.0 ){
      tempsum  = 1;
    }
  }
  atomicAdd(sum, tempsum);
}

__global__
void curand_setup(curandStatePhilox4_32_10_t *state, long seed) {
    int id = threadIdx.x   blockIdx.x * blockDim.x;
    curand_init(seed, id, 0, amp;state[id]);
}

int main(int argc, char *argv[]) {
  const int N = 1e5;

  int* count_h = 0;
  int* count_d;
  cudaMalloc(amp;count_d, sizeof(int) );
  cudaMemcpy(count_d, count_h, sizeof(int), cudaMemcpyHostToDevice);

  int threads_per_block = 64;
  int Nblocks = 32*6;

  thrust::device_vector<curandStatePhilox4_32_10_t> d_state(Nblocks*threads_per_block);
  curand_setup<<<Nblocks, threads_per_block>>>(d_state.data().get(), time(0));
  counts<<<Nblocks, threads_per_block>>>(N, count_d, d_state.data().get());

  cudaMemcpy(count_h, count_d, sizeof(int), cudaMemcpyDeviceToHost);

  cout << count_h << endl;

  cudaFree(count_d);
  free(count_h);
}
  

Я получаю ошибку терминала (на
linux):

 terminate called after throwing an instance of 'thrust::system::system_error'
  what():  parallel_for failed: cudaErrorInvalidValue: invalid argument
Aborted (core dumped)
  

И я компилирую следующим образом:

 nvcc -Xcompiler "-fopenmp" -o test uniform_one_hit_count.cu
  

Я не понимаю это сообщение об ошибке.

Эта строка:

 thrust::device_vector<curandStatePhilox4_32_10_t> d_state(Nblocks*threads_per_block);
  

инициализирует новый вектор на устройстве. Когда thrust делает это, он вызывает конструктор для используемого объекта, в данном случае curandStatePhilox4_32_10 , структуру, определение которой находится в /usr/local/cuda/include/curand_philox4x32_x.h (во всяком случае, в Linux). К сожалению, это определение структуры не предоставляет никаких конструкторов, украшенных __device__ , и это создает проблемы для thrust.

Простым обходным путем было бы собрать вектор на хосте и скопировать его на устройство:

 thrust::host_vector<curandStatePhilox4_32_10_t> h_state(Nblocks*threads_per_block);
thrust::device_vector<curandStatePhilox4_32_10_t> d_state = h_state;
  

В качестве альтернативы, просто используйте cudaMalloc для выделения пространства:

 curandStatePhilox4_32_10_t *d_state;
cudaMalloc(amp;d_state, (Nblocks*threads_per_block)*sizeof(d_state[0]));
  

У вас есть еще как минимум одна проблема. На самом деле это не обеспечивает надлежащего распределения памяти для того, на что должен указывать указатель:

 int* count_h = 0;
  

после этого вы должны сделать что-то вроде:

 count_h = (int *)malloc(sizeof(int));
memset(count_h, 0, sizeof(int));
  

и в вашей строке распечатки вы, скорее всего, захотите это сделать:

 cout << count_h[0] << endl;
  

Другим способом решения count_h проблемы было бы начать с:

 int count_h = 0;
  

и это потребовало бы другого набора изменений в вашем коде (в cudaMemcpy операциях).