cuda::experimental::dynamic_shared_memory_option

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

template<typename Content, std::size_t Extent = 1, bool NonPortableSize = false>
struct dynamic_shared_memory_option : public cuda::experimental::detail::launch_option

Launch option specyfying dynamic shared memory configuration.

This launch option causes the launch to allocate amount of shared memory sufficient to store the specified number of object of the specified type. This type can be constructed directly or with dynamic_shared_memory helper function.

When launch configuration contains this option, that configuration can be then passed to dynamic_smem_span or dynamic_smem_ref function to get a span/reference to that shared memory allocation that is approprietly typed. It is also possible to obtain that memory through the original extern shared variable[] declaration.

CUDA guarantees that each device has at least 48kB of shared memory per block, but most devices have more than that. In order to allocate more dynamic shared memory than the portable limit, opt-in NonPortableSize template argument should be set to true, otherwise kernel launch will fail.

Snippet

#include <cudax/launch.cuh>

template <typename Configuration>
__global__ void kernel(Configuration conf)
{
    auto dynamic_shared = cudax::dynamic_smem_span(conf);
    dynamic_shared[0] = 1;
}

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

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

Template Parameters

• Content – Type intended to be stored in dynamic shared memory

• Extent – Statically specified number of Content objects in dynamic shared memory, or cuda::std::dynamic_extent, if its dynamic

• NonPortableSize – Needs to be enabled to exceed the portable limit of 48kB of shared memory per block

Public Types

using content_type = Content

Public Functions

inline constexpr dynamic_shared_memory_option(std::size_t set_size) noexcept

Public Members

const std::size_t size

Public Static Attributes

static constexpr std::size_t extent = Extent

static constexpr bool is_relevant_on_device = true

static constexpr detail::launch_option_kind kind = detail::launch_option_kind::dynamic_shared_memory