1

考虑以下两个函数模板:

template <typename... Params>
void foo(Params... params)
{
    /* etc etc */
    my_kernel<<<grid_dims, block_dims, shmem_size, stream_id>>>(params...);
}

和:

template <typename... Params> 
void bar(Params... params)
{
    /* etc etc */
    void* arguments_ptrs[sizeof...(Params)];
    auto arg_index = 0;

    for_each_argument(
        [&](auto param) {arguments_ptrs[arg_index++] = &param;},
        params...);

    cudaLaunchKernel<decltype(my_kernel)>(
        &my_kernel, grid_dims, block_dims, argument_ptrs, shmem_size, stream_id);
}

由Sean Parent定义for_each_argument

问题:

  • foo和的语义是否bar完全相同?
  • 使用其中一个有什么好处吗?(例如,也许第一种形式在引擎盖下进行堆分配或其他东西......)
  • 在第二个函数中使用转发引用是个好主意吗?两个功能?
4

2 回答 2

1

foo 和 bar 的语义是否完全相同?

我还没有签入 CUDA 9,但在此之前,没有。语法被内联<<<>>>扩展为 API 调用和包装函数调用。有趣的是,所使用的内核启动 API早已被弃用。但是包装函数允许在编译时进行显式的参数类型安全检查,这很有帮助。

[编辑:我检查了 CUDA 9.1,它仍然cudaLaunch像所有以前版本的运行时 API 一样使用]

使用其中一个有什么好处吗?(例如,也许第一种形式在引擎盖下进行堆分配或其他东西......)

不是我知道的。

在第二个函数中使用转发引用是个好主意吗?两个功能?

如果内核在与调用代码相同的编译单元范围内编译,则不会。工具链自动发出内核的前向声明。

于 2018-02-01T07:44:20.243 回答
0

请记住,最终,运行时 API 需要进行驱动程序 API 调用(假设它不会进行我们不知道的秘密 APIcuLaunchKernel()调用),所以最终使用的是:

CUresult cuLaunchKernel ( 
    CUfunction f, 
    unsigned int  gridDimX, 
    unsigned int  gridDimY,
    unsigned int  gridDimZ, 
    unsigned int  blockDimX,
    unsigned int  blockDimY, 
    unsigned int  blockDimZ,
    unsigned int  sharedMemBytes,
    CUstream hStream, 
    void** kernelParams, 
    void** extra ) 

那是一个非模板化的接口,它不关心引用类型等。

当然,有两种方法可以指定启动参数 - usingkernelParams和 using extra。因此,如果您想调整启动内核的方式,您可能只想尝试一下。

于 2022-01-25T16:22:07.893 回答