我可以在不传递指针数组的情况下启动协作内核吗？

Question

CUDA 运行时 API 允许我们使用变量数量的参数三重人字形语法启动内核：

my_kernel<<<grid_dims, block_dims, shared_mem_size>>>(
    first_arg, second_arg, and_as_many, as_we, want_to, etc, etc);

但是关于 "collaborative" 内核，CUDA 编程指南说 (section C.3):

To enable grid synchronization, when launching the kernel it is necessary to use, instead of the <<<...>>> execution configuration syntax, the cuLaunchCooperativeKernel CUDA runtime launch API:
cudaLaunchCooperativeKernel(
  const T *func,
  dim3 gridDim,
  dim3 blockDim,
  void **args,
  size_t sharedMem = 0,
  cudaStream_t stream = 0
)      
(or the CUDA driver equivalent).

我宁愿不必编写自己的包装器代码来构建指针数组...运行时中真的没有工具可以避免这种情况吗API？

Answer 1

我们可以使用类似以下的解决方法（需要 --std=c++11 或更新的 C++ 语言标准）：

namespace detail {

template <typename F, typename... Args>
void for_each_argument_address(F f, Args&&... args) {
    [](...){}((void)(f( (void*) &std::forward<Args>(args) ), 0)...);
}

} // namespace detail

template<typename KernelFunction, typename... KernelParameters>
inline void cooperative_launch(
    const KernelFunction&       kernel_function,
    stream::id_t                stream_id,
    launch_configuration_t      launch_configuration,
    KernelParameters...         parameters)
{
    void* arguments_ptrs[sizeof...(KernelParameters)];
    auto arg_index = 0;
    detail::for_each_argument_address(
        [&](void * x) {arguments_ptrs[arg_index++] = x;},
        parameters...);
    cudaLaunchCooperativeKernel<KernelFunction>(
        &kernel_function,
        launch_configuration.grid_dimensions,
        launch_configuration.block_dimensions,
        arguments_ptrs,
        launch_configuration.dynamic_shared_memory_size,
        stream_id);
}

注意：这使用 Sean Parent's classic for_each_arg() one-liner. See also this post about it at FluentCPP.

Answer 2

答案是否定的

在幕后，<<< >>> 语法被扩展如下：

deviceReduceBlockKernel0<<<nblocks, 256>>>(input, scratch, N);

变成：

(cudaConfigureCall(nblocks, 256)) ? (void)0 : deviceReduceBlockKernel0(input, scratch, N);

并发出样板包装函数：

void deviceReduceBlockKernel0(int *in, int2 *out, int N) ;

// ....

void deviceReduceBlockKernel0( int *__cuda_0,struct int2 *__cuda_1,int __cuda_2)
{
__device_stub__Z24deviceReduceBlockKernel0PiP4int2i(_cuda_0,__cuda_1,__cuda_2);
}

void __device_stub__Z24deviceReduceBlockKernel1P4int2Pii( struct int2 *__par0,  int *__par1,  int __par2) 
{  
    __cudaSetupArgSimple(__par0, 0UL); 
    __cudaSetupArgSimple(__par1, 8UL); 
    __cudaSetupArgSimple(__par2, 16UL); 
    __cudaLaunch(((char *)((void ( *)(struct int2 *, int *, int))deviceReduceBlockKernel1))); 
}

即。当您显式使用内核启动 API 时，无论是传统的单次启动 API 还是新的协作启动 API，工具链都会自动执行您必须在代码中手动（或通过花哨的生成器模板）执行的操作。在已弃用的 API 版本中，有一个内部堆栈为您完成肮脏的工作。在较新的 API 中，您可以自己创建参数数组。同样的事情，只是不同的狗粮。

Answer 3

FWIW 您可以通过 void* args 传递任意结构（在 API 文档中不是很明显）。在这种情况下，编译器根据函数签名计算 sizeof 并将正确的大小复制到内核并不明显。 API 文档似乎没有对此进行详细说明。

struct Param { int a, b; void* device_ptr; };
Param param{aa, bb, d_ptr};
void *kArgs = {&param};
cudaLaunchCooperativeKernel(..., kArgs, ...);

我可以在不传递指针数组的情况下启动协作内核吗？

Can I launch a cooperative kernel without passing an array of pointers?

c++

cuda

variadic-templates

gpu-cooperative-groups