cuda::experimental::dynamic_shared_memory#

template<class _Tp>
class dynamic_shared_memory : private cuda::experimental::__dyn_smem_option_base<_Tp>, public cuda::experimental::__detail::launch_option#

Launch option specifying 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 with dynamic_shared_memory helper function.

When launch configuration contains this option, that configuration can be then passed to dynamic_shared_memory_view to get the view_type over the dynamic shared memory. 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_shared_memory_view(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 Functions

constexpr dynamic_shared_memory() noexcept = default#
inline constexpr dynamic_shared_memory(non_portable_t) noexcept#
template<class _Tp2 = _Tp, ::cuda::std::enable_if_t<(!::cuda::std::is_unbounded_array_v<_Tp2>), int> = 0>
inline constexpr dynamic_shared_memory(
) noexcept#
template<class _Tp2 = _Tp, ::cuda::std::enable_if_t<(!::cuda::std::is_unbounded_array_v<_Tp2>), int> = 0>
inline constexpr dynamic_shared_memory(
non_portable_t,
) noexcept#
template<class _Tp2 = _Tp, ::cuda::std::enable_if_t<::cuda::std::is_unbounded_array_v<_Tp2>, int> = 0>
inline constexpr dynamic_shared_memory(
::cuda::std::size_t __n,
)#
template<class _Tp2 = _Tp, ::cuda::std::enable_if_t<::cuda::std::is_unbounded_array_v<_Tp2>, int> = 0>
inline constexpr dynamic_shared_memory(
::cuda::std::size_t __n,
non_portable_t,
) noexcept#
inline constexpr ::cuda::std::size_t size_bytes() const noexcept#

Gets the size of the dynamic shared memory in bytes.

Public Static Attributes

static constexpr bool is_relevant_on_device = true#

< The view type returned by the cuda::device::dynamic_shared_memory_view(config).

static constexpr __detail::launch_option_kind kind = __detail::launch_option_kind::dynamic_shared_memory#