cuda::experimental::cooperative_launch

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

struct cooperative_launch : public cuda::experimental::detail::launch_option

Launch option enabling cooperative launch.

This launch option causes the launched grid to be restricted to a number of blocks that can simultaneously execute on the device. It means that every thread in the launched grid can eventualy observe execution of each other thread in the grid. It also enables usage of cooperative_groups::grid_group::sync() function, that synchronizes all threads in the grid.

Snippet

#include <cudax/launch.cuh>
#include <cooperative_groups.h>

template <typename Configuration>
__global__ void kernel(Configuration conf)
{
    auto grid = cooperative_groups::this_grid();
    grid.sync();
}

void kernel_launch(cuda::stream_ref stream) {
    auto dims = cudax::make_hierarchy(cudax::block<128>(), cudax::grid(4));
    auto conf = cudax::make_configuration(dims, cooperative_launch());

    cudax::launch(stream, conf, kernel);
}

Public Functions

constexpr cooperative_launch() = default

Public Static Attributes

static constexpr bool needs_attribute_space = true

static constexpr bool is_relevant_on_device = true

static constexpr detail::launch_option_kind kind = detail::launch_option_kind::cooperative_launch