带有模板的 CUDA 内核启动宏答案

【问题标题】：CUDA kernel launch macro with templates带有模板的 CUDA 内核启动宏
【发布时间】：2013-02-14 11:58:41
【问题描述】：

我做了一个宏来简化 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 ">"

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

【问题讨论】：

在你给出的情况下真的会发生吗，还是只有当模板有多个参数时才会发生？
是的，它只会因多个参数而失败。我要纠正这个问题。
问题是预处理器对尖括号嵌套一无所知，因此它将它们之间的逗号解释为宏参数分隔符。

标签： c++ templates macros cuda kernel

【解决方案1】：

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

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

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

【讨论】：

【解决方案2】：

您可以尝试的其他方法（基于您发布的宏）是将内核块大小和网格大小参数包装在它们自己的宏中：

#define KERNEL_ARGS2(grid, block) <<< grid, block >>>
#define KERNEL_ARGS3(grid, block, sh_mem) <<< grid, block, sh_mem >>>
#define KERNEL_ARGS4(grid, block, sh_mem, stream) <<< grid, block, sh_mem, stream >>>

现在你应该可以像这样使用你的宏了：

#define CUDA_LAUNCH(kernel_name, gridsize, blocksize, ...) \
kernel_name KERNEL_ARGS2(gridsize, blocksize)(__VA_ARGS__);

你可以像这样使用它：

CUDA_LAUNCH(my_kernel, grid_size, block_size, float* input, float* output, int size);

这将使用给定的网格和块大小以及输入参数启动名为“my_kernal”的内核。

【讨论】：

【解决方案3】：

考虑这个也会引发错误的解决方案

inline void echoError(cudaError_t e, const char *strs) {
    char a[255];
    if (e != cudaSuccess) {
        strncpy(a, strs, 255);
        fprintf(stderr, "Failed to %s,errorCode %s",
                a, cudaGetErrorString(e));
        exit(EXIT_FAILURE);
    }
}


#define CUDA_KERNEL_DYN(kernel, bpg, tpb, shd, ...){                     \
    kernel<<<bpg,tpb,shd>>>( __VA_ARGS__ );                              \
    cudaError_t err = cudaGetLastError();                                \
    echoError(err, #kernel);                                              \
}

【讨论】：