Передайте лямбду __device__ в качестве аргумента функции __global__

Question

Определение __device__ лямбда-символов весьма полезно.

Я хотел сделать то же самое, что и код ниже, но с лямбда-выражением, определенным в разных файлах из ядра, которое будет его использовать.

// Sample code that works
template<typename Func>
__global__ void kernel(Func f){
    f(threadIdx.x);
}

int main(){
    auto f = [] __device__ (int i){ printf("Thread n°%i\n",i); };
    kernel<<<1,16>>>(f);
}

Я пробовал эту (не работающую) реализацию.

main.cu

#include "kernelFile.h"

int main(){
    auto f = [] __device__ (int i){ printf("Thread n°%i\n",i); };
    kernelCaller(f);
}

kernelFile.cu

template<typename Func>
__global__ void kernel(Func f){
    f(threadIdx.x);
}

template<typename Func>
__host__ void kernelCaller(Func f){
    kernelCaller(f);
}

Но компилятор жалуется, потому что kernelCaller никогда не создается. Я не знаю, можно ли это создать или нет, или то, что я пытаюсь сделать, должно быть реализовано иначе. Любой намек на то, что я должен делать?

Answer 1

Невозможно обеспечить надлежащее создание экземпляра функции для шаблонной функции, если только требуемый тип не известен в модуле компиляции, где создается экземпляр функции. Это не указывается c для CUDA.

Следовательно, любой метод потребует некоторого знания требуемого типа в модуле компиляции, где функция ядра компилируется / создается. При условии, что один из возможных подходов охватывается здесь . Мы можем избежать неопределенности типа, связанной с лямбдой, обернув ее в nvstd::function объект. Затем пусть ваше ядро ​​примет объект nvstd::function (который может фактически являться упаковщиком типов для лямбды), и ваш вызывающий хост вставит нужную лямбду в объект nvstd::function.

Вот пример :

$ cat k.cu
#include <nvfunctional>
#include <cstdio>

typedef nvstd::function<int(unsigned)> v;
__global__ void kernel(v *f){

  printf("%d, %d\n", threadIdx.x, (*f)(threadIdx.x));
}

__host__ void kernelCaller(v *f){
  kernel<<<1,2>>>(f);
}

$ cat m.cu
#include <nvfunctional>
// prototype would normally be in a header file
void kernelCaller(nvstd::function<int(unsigned)> *);


template <typename T1, typename T2>
__global__ void inserter(T1 *f, T2 l){
  *f = l;
}


int main(){

  nvstd::function<int(unsigned)> *d_f;
  cudaMalloc(&d_f, sizeof(nvstd::function<int(unsigned)>));
  auto lam1 = [] __device__ (unsigned i) { return i+1;};
  inserter<<<1,1>>>(d_f, lam1);
  kernelCaller(d_f);
  auto lam2 = [] __device__ (unsigned i) { return (i+1)*2;};
  inserter<<<1,1>>>(d_f, lam2);
  kernelCaller(d_f);
  cudaDeviceSynchronize();
}
$ nvcc -o test k.cu m.cu -std=c++11 -expt-extended-lambda -rdc=true
$ cuda-memcheck ./test
========= CUDA-MEMCHECK
0, 1
1, 2
0, 2
1, 4
========= ERROR SUMMARY: 0 errors
$