我可以在不传递指针数组的情况下启动协作内核吗？
CUDA运行时API允许我们使用变量数参数三元雪佛龙语法启动内核：
my_kernel<<<grid_dims, block_dims, shared_mem_size>>>(
first_arg, second_arg, and_as_many, as_we, want_to, etc, etc);但是关于“协作”内核，数据自动化系统编程指南说(第C.3节)：
为了启用网格同步，在启动内核时，必须使用
cuLaunchCooperativeKernelCUDA运行时启动API，而不是使用cuLaunchCooperativeKernel执行配置语法： cudaLaunchCooperativeKernel( const *func，dim3 gridDim，dim3 blockDim，void **args，size_t sharedMem = 0，cudaStream_t stream =0)(或相当于CUDA驱动程序)。
我宁愿不必为构建指针数组编写自己的包装代码..。运行时API中真的没有避免这种情况的工具吗？
回答 3
Stack Overflow用户
发布于 2018-02-01 10:08:37
我们可以使用以下解决方法(需要--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);
}注意:这使用for_each_arg()一行程序.另见关于它的这篇文章 at FluentCPP。
Stack Overflow用户
发布于 2018-06-01 02:06:23
FWIW您可以通过void传递任意结构(在API文档中不是很明显)。在这种情况下，编译器从函数签名中计算出的大小并不明显，并且适当的大小被复制到内核中。API文档似乎没有详细说明这一点。
struct Param { int a, b; void* device_ptr; };
Param param{aa, bb, d_ptr};
void *kArgs = {¶m};
cudaLaunchCooperativeKernel(..., kArgs, ...);Stack Overflow用户
发布于 2018-02-01 10:14:57
答案是否定的。
在遮罩下，<<< >>>语法将按如下方式展开：
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中，您自己创建参数数组。同样的事情，只是不同的狗粮。
https://stackoverflow.com/questions/48547409
复制相似问题
腾讯云开发者
