功能指针(到其他内核）作为CUDA中的内核Arg

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

本文关键字：内核 CUDA Arg 作为 其他 功能 指针      更新时间：2023-10-16

使用CUDA中的动态并行性，您可以从特定版本开始在GPU侧启动内核。我有一个包装器功能，该功能将指针用于我要使用的内核，并且它可以在CPU上用于旧设备的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 GPUn");
        Test << <1, 1 >> >(kernel, count, args...);
    }
    else
    {
        printf("Iterate on CPUn");
        Test << <1, 1 >> >(kernel, count, args...);
    }
}

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

原始答案：

似乎不时出现这个问题，尽管我可以想到的以前的示例与调用__device__函数而不是__global__函数有关。

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

围绕此操作的一种可能的方法（尽管我的实用程序对我来说尚不清楚；似乎有更简单的调度机制）是提取"设备代码中"所需的设备地址，并将该值返回到主机，用于调度使用。在这种情况下，我正在创建一个简单的示例，该示例将所需的设备地址提取到__device__变量中，但是您还可以编写一个内核来进行此设置（即，在您的文字中，"给我一个在GPU上正确的指针地址"））。

这是一个粗略的示例，在您显示的代码上构建：

$ cat t746.cu
#include <stdio.h>
__global__ void ckernel1(){
  printf("hello1n");
}
__global__ void ckernel2(){
  printf("hello2n");
}
__global__ void ckernel3(){
  printf("hello3n");
}
__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 GPUn");
        Test << <1, 1 >> >(kernel, count, args...);
    }
    else
    {
        printf("Iterate on CPUn");
        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__变量）方法可能无法与模板的子核一起使用，但是可能可以创建一个模板的"提取器"内核，该内核返回（实例化的）模板子内核的地址。我链接的上一个答案中给出了"提取器" setup_kernel方法的大概。这是模板子内核/提取器内核方法的粗略示例：

$ 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 GPUn");
        Test << <1, 1 >> >(kernel, count, args...);
    }
    else
    {
        printf("Iterate on CPUn");
        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
$
相关文章：