cuda::experimental::launch

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

template<typename ...Args, typename ...Levels, typename Kernel>
void cuda::experimental::launch(::cuda::stream_ref stream, const hierarchy_dimensions<Levels...> &dims, const Kernel &kernel, Args... args)

Launch a kernel functor with specified thread hierarchy and arguments.

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

Snippet

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

struct kernel {
    template <typename Dimensions>
    __device__ void operator()(Dimensions dims, unsigned int thread_to_print) {
        if (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));

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

Parameters

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

• dims – thread hierarchy dimensions for this launch

• kernel – kernel functor to be launched

• args – arguments to be passed into the kernel functor