具有函数指针和可变参数模板的 CUDA 内核

CUDA kernel with function pointer and variadic templates(具有函数指针和可变参数模板的 CUDA 内核)

问题描述

我正在尝试设计一个 cuda 框架，该框架将接受用户函数并通过设备函数指针将它们转发到内核.CUDA 可以与可变参数模板 (-stc=c++11) 一起使用，到目前为止一切正常.

但是，当内核调用设备函数指针时，我遇到了问题.显然内核运行没有问题，但 GPU 使用率为 0%.如果我只是用实际函数替换回调指针，那么 GPU 使用率为 99%.这里的代码非常简单，大循环范围只是为了让事情变得可衡量.我用以下方法测量了 gpu 状态:

nvidia-smi --query-gpu=utilization.gpu,utilization.mory,memory.used --format=csv -lms 100 -f out.txt

IIRC，用户函数需要与内核在同一个文件单元中(可能是#included)才能使 nvcc 成功.func_d 就在源代码中，它编译和运行良好，除了不使用函数指针(这是本设计的重点).

我的问题是:为什么带有回调设备函数指针的内核不起作用?

请注意，当我 printf noth 回调和 func_d 地址时，它们是相同的，如此示例输出中所示:

Args 的大小 = 1回调()地址 = 4024b0func_d() 地址 = 4024b0

另一件奇怪的事是，如果取消注释 kernel() 中的 callback() 调用，那么 GPU 使用率会回到 0%，即使使用 func_d() 调用仍在其中... func_d 版本大约需要 4 秒才能运行，而回调版本不需要任何时间(嗯，~0.1 秒).

系统规格和编译命令在下面代码的头部.

谢谢！

//编译://nvcc -g -G -O0 -std=c++11 -arch=sm_20 -x cu sample.cpp////Nvidia Quadro 6000(计算能力 2.0)//CUDA 6.5 (V6.5.12),//Arch Linux，Nvidia 驱动程序 343.22-4，gcc 4.9.1//2014 年 11 月#include <stdio.h>__设备__无效 func_d(双 * 卷){*卷 += 5.4321f;}//CUDA核函数模板<typename...类型>__global__ void kernel( void (*callback)(Types*...) ){双 val0 = 1.2345f；////不使用 gpu (0% gpu 利用率)//for (int i = 0; i <1000000; i++) {//回调( &val0 );/