cuda::experimental::launch

Defined in /home/runner/work/cccl/cccl/cudax/include/cuda/experimental/__launch/launch.cuh

template<typename ...ExpArgs, typename ...ActArgs, typename ...Config, typename Dimensions>
void cuda::experimental::launch(::cuda::stream_ref stream, 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 Congifuration>
__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 confing = 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