Как использовать макросы или шаблоны C ++, чтобы разрешить повторное использование кода в коде хоста и устройства?

Question

Я пишу хобби raytracer на CUDA и C ++, и у меня возникла проблема, на которую я так и не смог получить ответ. Я написал код процессора и графического процессора таким образом, чтобы он мог выполняться на машинах с устройствами CUDA или без них. Однако это привело к некоторому дублированию кода в следующем смысле: Небольшой набор функций требует генерации случайных чисел, что достигается с помощью stdlib на хосте и curand на устройстве. Я хотел бы иметь __host__ __device__ функции, которые принимают структуру Sampler, которая либо вызывает rand() на хосте, либо curand_uniform() на устройстве. Я пробовал кое-что, но не могу заставить программу скомпилироваться - компилятор жалуется на то, что он не вызывает __device__ функций из __host__ кода и наоборот.

В идеале я бы хотел, чтобы мои функции рендеринга взяли Sampler *, который выглядит примерно так, как показано ниже.

Спасибо!

struct Sampler {
    __host__ virtual float getNextFloat() { return rand() / (RAND_MAX + 1.f); }
};

struct CudaSampler : Sampler { 
    curandState* p_curandState;
    __device__ float getNextFloat() { return curand_uniform(p_curandState); }
};

Answer 1

То, что вы просите, должно быть возможным. Мы не хотим пытаться перегрузить функцию отдельно через __host__ и __device__ (не разрешено), и мы не хотим пытаться сделать это с помощью наследования и виртуальных функций (таблица виртуальных функций не будет использоваться в такой объект передается от хоста к устройству).

Но если мы избежим этих проблем, основная идея состоит в том, чтобы использовать макрос __CUDA_ARCH__ nvcc для разграничения пути хоста и устройства для компилятора, как правило, в соответствии с тем, что предлагается здесь .

Вот один из возможных методов, примерно следуя вашей схеме:

$ cat t34.cu
#include <cstdio>
#include <cstdlib>
#include <iostream>
#include <curand_kernel.h>


struct Sampler {
  __host__ __device__ float operator()(curandState *s){
#ifdef __CUDA_ARCH__
    return curand_uniform(s);
#else
    return rand()/(float)RAND_MAX;
#endif
  }
};

__global__ void init_rng(curandState *state, size_t n){
        size_t idx = threadIdx.x+blockDim.x*blockIdx.x;
        if (idx < n)
          curand_init(1234, idx, 0, state+idx);
}

__global__ void gpu_sample(Sampler s, curandState *state, size_t n){
        size_t idx = threadIdx.x+blockDim.x*blockIdx.x;
        if (idx < n)
                printf("gpu id: %lu, val: %f\n", idx, s(state+idx));
}

__host__  void cpu_sample(Sampler s){
        curandState dummy;
    std::cout << "cpu: " << s(&dummy) << std::endl;
}

int main(){
        int n = 1;
        int nTPB = 256;
        curandState *s;
        Sampler my_op;
        cudaMalloc(&s, n*sizeof(curandState));
        init_rng<<<(n+nTPB-1)/nTPB, nTPB>>>(s,n);
        gpu_sample<<<(n+nTPB-1)/nTPB, nTPB>>>(my_op, s, n);
        cudaDeviceSynchronize();
        cpu_sample(my_op);
}



$ nvcc -o t34 t34.cu
$ cuda-memcheck ./t34
========= CUDA-MEMCHECK
gpu id: 0, val: 0.145468
cpu: 0.840188
========= ERROR SUMMARY: 0 errors
$