Могу ли я запустить кооперативное ядро, не передавая массив указателей?

Question

Могу ли я запустить кооперативное ядро, не передавая массив указателей?

API среды выполнения CUDA позволяет нам запускать ядра с использованием синтаксиса тройного шеврона с переменным числом аргументов:

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

но что касается «совместных» ядер, в Руководстве по программированию CUDA сказано:раздел С.3):

Чтобы включить синхронизацию сетки, при запуске ядра это
необходимо использовать вместо <<<...>>> конфигурация исполнения
синтаксис cuLaunchCooperativeKernel API запуска CUDA:
cudaLaunchCooperativeKernel(
const T *func,
dim3 gridDim,
dim3 blockDim,
void **args,
size_t sharedMem = 0,
cudaStream_t stream = 0
)
(или эквивалент драйвера CUDA).

Я бы предпочел не писать свой собственный код-обертку для создания массива указателей … действительно ли в API времени выполнения нет средств, чтобы избежать этого?

1

c++cuda gpu-cooperative-groups variadic-templates

Решение

Другие решения

Мы можем использовать что-то вроде следующего обходного решения (требуется --std=c++11 или выше):

namespace detail {

template <typename F, typename... Args>
void for_each_argument_address(F f, Args&&... args) {
[](...){}((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);
}

0

FWIW вы можете передавать произвольные структуры (не сразу это видно из документации API), просто передавая их через void * args. Не очевидно, что sizeof вычисляется компилятором в этом случае из сигнатуры функции, и правильный размер копируется в ядро. Документы API, кажется, не детализируют это.

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

0

Источник

Accepted Answer

Ответ — нет.

Под капотом <<< >>> синтаксис расширяется так:

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 вы сами делаете массивы аргументов. То же самое, просто другой корм для собак.

2