cuda::experimental::launch

Overloads

launch(__submitter, __conf, __kernel, __args)

template<typename ..._Args, typename ..._Config, typename _Submitter, typename _Dimensions, typename _Kernel>
auto cuda::experimental::launch(

_Submitter &&__submitter,

const kernel_config<_Dimensions, _Config...> &__conf,

const _Kernel &__kernel,

_Args&&... __args,

)

Launch a kernel functor with specified configuration and arguments.

Launches a kernel functor object on the specified stream and with specified configuration. Kernel functor object is a type with device operator(). Functor might or might not accept the configuration as its first argument.

Snippet

#include <cstdio>
#include <cuda/experimental/launch.cuh>

struct kernel {
    template <typename Configuration>
    __device__ void operator()(Configuration conf, unsigned int thread_to_print) {
        if (conf.dims.rank(cudax::thread, cudax::grid) == thread_to_print) {
            printf("Hello from the GPU\n");
        }
    }
};

void launch_kernel(cuda::stream_ref stream) {
    auto dims    = cudax::make_hierarchy(cudax::block_dims<128>(), cudax::grid_dims(4));
    auto config = cudax::make_config(dims, cudax::launch_cooperative());

    cudax::launch(stream, config, kernel(), 42);
}

Parameters:

• stream – cuda::stream_ref to launch the kernel into

• conf – configuration for this launch

• kernel – kernel functor to be launched

• args – arguments to be passed into the kernel functor

launch(__submitter, __conf, _ExpArgs...), __args)

template<typename ..._ExpArgs, typename ..._ActArgs, typename _Submitter, typename ..._Config, typename _Dimensions>
auto cuda::experimental::launch(

_Submitter &&__submitter,

const kernel_config<_Dimensions, _Config...> &__conf,

void (*__kernel)(kernel_config<_Dimensions, _Config...>, _ExpArgs...),

_ActArgs&&... __args,

)

Launch a kernel function with specified configuration and arguments.

Launches a kernel function on the specified stream and with specified configuration. Kernel function is a function with global annotation. Function might or might not accept the configuration as its first argument.

Snippet

#include <cstdio>
#include <cuda/experimental/launch.cuh>

template <typename Configuration>
__global__ void kernel(Configuration conf, unsigned int thread_to_print) {
    if (conf.dims.rank(cudax::thread, cudax::grid) == thread_to_print) {
        printf("Hello from the GPU\n");
    }
}

void launch_kernel(cuda::stream_ref stream) {
    auto dims    = cudax::make_hierarchy(cudax::block_dims<128>(), cudax::grid_dims(4));
    auto config = cudax::make_config(dims, cudax::launch_cooperative());

    cudax::launch(stream, config, kernel<decltype(config)>, 42);
}

Parameters:

• stream – cuda::stream_ref to launch the kernel into

• conf – configuration for this launch

• kernel – kernel function to be launched

• args – arguments to be passed into the kernel function

launch(__submitter, __conf, void(*__kernel)(_ExpArgs...), __args)

template<typename ..._ExpArgs, typename ..._ActArgs, typename _Submitter, typename ..._Config, typename _Dimensions>
auto cuda::experimental::launch(

_Submitter &&__submitter,

const kernel_config<_Dimensions, _Config...> &__conf,

void (*__kernel)(_ExpArgs...),

_ActArgs&&... __args,

)

Launch a kernel function with specified configuration and arguments.

Launches a kernel function on the specified stream and with specified configuration. Kernel function is a function with global annotation. Function might or might not accept the configuration as its first argument.

Snippet

#include <cstdio>
#include <cuda/experimental/launch.cuh>

template <typename Configuration>
__global__ void kernel(Configuration conf, unsigned int thread_to_print) {
    if (conf.dims.rank(cudax::thread, cudax::grid) == thread_to_print) {
        printf("Hello from the GPU\n");
    }
}

void launch_kernel(cuda::stream_ref stream) {
    auto dims    = cudax::make_hierarchy(cudax::block_dims<128>(), cudax::grid_dims(4));
    auto config = cudax::make_config(dims, cudax::launch_cooperative());

    cudax::launch(stream, config, kernel<decltype(config)>, 42);
}

Parameters:

• __stream – cuda::stream_ref to launch the kernel into

• __conf – configuration for this launch

• __kernel – kernel function to be launched

• __args – arguments to be passed into the kernel function

launch(>, _Fn, >)

template<class _Dimensions, class ..._Config, class _Fn, class ..._Args>
inline constexpr auto cuda::experimental::launch(

kernel_config<_Dimensions, _Config...> __config,

_Fn __fn,

_Args... __args,

) -> __kernel_t::__sndr_t<kernel_config<_Dimensions, _Config...>, _Fn, _Args...>