函数指针（指向其他内核）作为 CUDA 中的内核 arg

Question

借助 CUDA 中的动态并行性，您可以从特定版本开始在 GPU 端启动内核。我有一个包装函数，它接受一个指向我想使用的内核的指针，它要么在旧设备的 CPU 上执行此操作，要么在新设备的 GPU 上执行此操作。对于回退路径，它很好，对于 GPU，它不是，并且说内存对齐不正确。

有没有办法在 CUDA (7) 中做到这一点？是否有一些较低级别的调用会给我一个在 GPU 上正确的指针地址？

代码如下，模板"TFunc"试图让编译器做一些不同的事情，但我也试过它是强类型的。

template <typename TFunc, typename... TArgs>
__global__ void Test(TFunc func, int count, TArgs... args)
{
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 320)
    (*func)<< <1, 1 >> >(args...);
#else
    printf("What are you doing here!?\n");
#endif
}

template <typename... TArgs>
__host__ void Iterate(void(*kernel)(TArgs...), const systemInfo *sysInfo, int count, TArgs... args)
{
    if(sysInfo->getCurrentDevice()->compareVersion("3.2") > 0)
    {
        printf("Iterate on GPU\n");
        Test << <1, 1 >> >(kernel, count, args...);
    }
    else
    {
        printf("Iterate on CPU\n");
        Test << <1, 1 >> >(kernel, count, args...);
    }
}

Answer 1

编辑： 在我最初写这个答案的时候，我相信这些陈述是正确的：不可能在主机代码中获取内核地址。但是我相信从那时起 CUDA 中发生了一些变化，所以现在（在 CUDA 8 中，也许之前）可以在主机代码中使用 kernel 地址（仍然不可能但是，在主机代码中获取 __device__ 函数的地址。）

原始答案：

虽然 previous examples I can think of 与调用 __device__ 函数而不是 __global__ 函数有关，但这个问题似乎时常出现。

通常在主机代码中获取设备实体（变量、函数）的地址是非法的。

解决这个问题的一种可能方法（虽然我不清楚它的用途；似乎会有更简单的调度机制）是提取所需的设备地址 "in device code" 和 return 将该值发送给主机，以供调度使用。在这种情况下，我正在创建一个简单的示例，将所需的设备地址提取到 __device__ 变量中，但您也可以编写一个内核来执行此设置（即 "give me a pointer address that's correct on the GPU" 用您的话来说）。

这是一个粗略的示例，基于您展示的代码：

$ cat t746.cu
#include <stdio.h>

__global__ void ckernel1(){

  printf("hello1\n");
}
__global__ void ckernel2(){

  printf("hello2\n");
}
__global__ void ckernel3(){

  printf("hello3\n");
}

__device__ void (*pck1)() = ckernel1;
__device__ void (*pck2)() = ckernel2;
__device__ void (*pck3)() = ckernel3;

template <typename TFunc, typename... TArgs>
__global__ void Test(TFunc func, int count, TArgs... args)
{
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 350)
    (*func)<< <1, 1 >> >(args...);
#else
    printf("What are you doing here!?\n");
#endif
}

template <typename... TArgs>
__host__ void Iterate(void(*kernel)(TArgs...), const int sysInfo, int count, TArgs... args)
{
    if(sysInfo >= 350)
    {
        printf("Iterate on GPU\n");
        Test << <1, 1 >> >(kernel, count, args...);
    }
    else
    {
        printf("Iterate on CPU\n");
        Test << <1, 1 >> >(kernel, count, args...);
    }
}


int main(){

  void (*h_ckernel1)();
  void (*h_ckernel2)();
  void (*h_ckernel3)();
  cudaMemcpyFromSymbol(&h_ckernel1, pck1, sizeof(void *));
  cudaMemcpyFromSymbol(&h_ckernel2, pck2, sizeof(void *));
  cudaMemcpyFromSymbol(&h_ckernel3, pck3, sizeof(void *));
  Iterate(h_ckernel1, 350, 1);
  Iterate(h_ckernel2, 350, 1);
  Iterate(h_ckernel3, 350, 1);
  cudaDeviceSynchronize();
  return 0;
}

$ nvcc -std=c++11 -arch=sm_35 -o t746 t746.cu -rdc=true -lcudadevrt
$ cuda-memcheck ./t746
========= CUDA-MEMCHECK
Iterate on GPU
Iterate on GPU
Iterate on GPU
hello1
hello2
hello3
========= ERROR SUMMARY: 0 errors
$

上述（__device__ 变量）方法可能无法与模板化的子内核一起使用，但可以创建一个模板化的 "extractor" 内核 return s 一个（实例化的）模板化子内核的地址。 "extractor" setup_kernel 方法的粗略概念在我链接的上一个答案中给出。这是模板化子 kernel/extractor 内核方法的粗略示例：

$ cat t746.cu
#include <stdio.h>

template <typename T>
__global__ void ckernel1(T *data){

  int my_val = (int)(*data+1);
  printf("hello: %d \n", my_val);
}
template <typename TFunc, typename... TArgs>
__global__ void Test(TFunc func, int count, TArgs... args)
{
#if defined(__CUDA_ARCH__) && (__CUDA_ARCH__ >= 350)
    (*func)<< <1, 1 >> >(args...);
#else
    printf("What are you doing here!?\n");
#endif
}

template <typename... TArgs>
__host__ void Iterate(void(*kernel)(TArgs...), const int sysInfo, int count, TArgs... args)
{
    if(sysInfo >= 350)
    {
        printf("Iterate on GPU\n");
        Test << <1, 1 >> >(kernel, count, args...);
    }
    else
    {
        printf("Iterate on CPU\n");
        Test << <1, 1 >> >(kernel, count, args...);
    }
}

template <typename T>
__global__ void extractor(void (**kernel)(T *)){

  *kernel = ckernel1<T>;
}

template <typename T>
void run_test(T init){

  void (*h_ckernel1)(T *);
  void (**d_ckernel1)(T *);
  T *d_data;
  cudaMalloc(&d_ckernel1, sizeof(void *));
  cudaMalloc(&d_data, sizeof(T));
  cudaMemcpy(d_data, &init, sizeof(T), cudaMemcpyHostToDevice);
  extractor<<<1,1>>>(d_ckernel1);
  cudaMemcpy((void *)&h_ckernel1, (void *)d_ckernel1, sizeof(void *), cudaMemcpyDeviceToHost);
  Iterate(h_ckernel1, 350, 1, d_data);
  cudaDeviceSynchronize();
  cudaFree(d_ckernel1);
  cudaFree(d_data);
  return;
}

int main(){

  run_test(1);
  run_test(2.0f);

  return 0;
}

$ nvcc -std=c++11 -arch=sm_35 -o t746 t746.cu -rdc=true -lcudadevrt
$ cuda-memcheck ./t746
========= CUDA-MEMCHECK
Iterate on GPU
hello: 2
Iterate on GPU
hello: 3
========= ERROR SUMMARY: 0 errors
$

函数指针（指向其他内核）作为 CUDA 中的内核 arg

Function pointer (to other kernel) as kernel arg in CUDA

c++

cuda

gpgpu

function-pointers