带有模板的 CUDA 内核启动宏

Question

我做了一个宏来简化 CUDA 内核调用：

#define LAUNCH LAUNCH_ASYNC

#define LAUNCH_ASYNC(kernel_name, gridsize, blocksize, ...) \
    LOG("Async kernel launch: " #kernel_name);              \
    kernel_name <<< (gridsize), (blocksize) >>> (__VA_ARGS__);

#define LAUNCH_SYNC(kernel_name, gridsize, blocksize, ...)     \
    LOG("Sync kernel launch: " #kernel_name);                  \
    kernel_name <<< (gridsize), (blocksize) >>> (__VA_ARGS__); \
    cudaDeviceSynchronize();                                   \
    // error check, etc...

用法：

LAUNCH(my_kernel, 32, 32, param1, param2)

LAUNCH(my_kernel<int>, 32, 32, param1, param2)

这很好用；通过第一个定义，我可以启用同步调用和错误检查以进行调试。

但是它不适用于多个模板参数，如下所示：

LAUNCH(my_kernel<int,float>, 32, 32, param1, param3)

我在调用宏的行中收到的错误消息：

error : expected a ">"

是否可以使这个宏与多个模板参数一起工作？

Answer 1

问题是预处理器对尖括号嵌套一无所知，因此它将它们之间的逗号解释为宏参数分隔符。

如果内核启动语法支持内核名称周围的括号（我现在无法检查，而不是在 CUDA 机器上），您可以这样做：

LAUNCH((my_kernel<int, float>), 32, 32, param1, param3)