From 8ecc944490bf3a5f324cf17e563bbfbdbae05468 Mon Sep 17 00:00:00 2001 From: pciolkosz Date: Tue, 16 Dec 2025 18:19:33 -0800 Subject: [PATCH 1/2] Remove _view from the shared memory getter --- .../include/cuda/__launch/configuration.h | 135 ++++++++++++------ .../ccclrt/launch/dynamic_shared_memory.cu | 6 +- .../cuda/ccclrt/launch/launch_smoke.cu | 4 +- 3 files changed, 96 insertions(+), 49 deletions(-) diff --git a/libcudacxx/include/cuda/__launch/configuration.h b/libcudacxx/include/cuda/__launch/configuration.h index 16e2511ddf04..d81a69ab4b57 100644 --- a/libcudacxx/include/cuda/__launch/configuration.h +++ b/libcudacxx/include/cuda/__launch/configuration.h @@ -196,7 +196,7 @@ inline constexpr ::cuda::std::size_t __max_portable_dyn_smem_size = 48 * 1024; * 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 + * then passed to dynamic_shared_memory 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. * @@ -213,14 +213,14 @@ inline constexpr ::cuda::std::size_t __max_portable_dyn_smem_size = 48 * 1024; * template * __global__ void kernel(Configuration conf) * { - * auto dynamic_shared = cuda::dynamic_shared_memory_view(conf); + * auto dynamic_shared = cuda::dynamic_shared_memory(conf); * dynamic_shared[0] = 1; * } * * void kernel_launch(cuda::stream_ref stream) { * auto dims = cuda::make_hierarchy(cuda::block<128>(), cuda::grid(4)); * auto conf = cuda::make_configuration(dims, - * dynamic_shared_memory()); + * cuda::dynamic_shared_memory()); * * cuda::launch(stream, conf, kernel); * } @@ -239,7 +239,7 @@ inline constexpr ::cuda::std::size_t __max_portable_dyn_smem_size = 48 * 1024; * per block */ template -class _CCCL_DECLSPEC_EMPTY_BASES dynamic_shared_memory +class _CCCL_DECLSPEC_EMPTY_BASES dynamic_shared_memory_option : __dyn_smem_option_base<_Tp> , public __detail::launch_option { @@ -258,48 +258,11 @@ class _CCCL_DECLSPEC_EMPTY_BASES dynamic_shared_memory using typename __base_type::value_type; //!< Value type of the dynamic //!< shared memory elements. using typename __base_type::view_type; //!< The view type returned by the - //!< cuda::dynamic_shared_memory_view(config). + //!< cuda::dynamic_shared_memory(config). static constexpr bool is_relevant_on_device = true; static constexpr __detail::launch_option_kind kind = __detail::launch_option_kind::dynamic_shared_memory; - _CCCL_HIDE_FROM_ABI constexpr dynamic_shared_memory() noexcept = default; - - _CCCL_HOST_API constexpr dynamic_shared_memory(non_portable_t) noexcept - : __non_portable_{true} - {} - - _CCCL_TEMPLATE(class _Tp2 = _Tp) - _CCCL_REQUIRES((!::cuda::std::is_unbounded_array_v<_Tp2>) ) - _CCCL_HOST_API constexpr dynamic_shared_memory() noexcept - { - static_assert(sizeof(_Tp2) <= __max_portable_dyn_smem_size, "portable dynamic shared memory limit exceeded"); - } - - _CCCL_TEMPLATE(class _Tp2 = _Tp) - _CCCL_REQUIRES((!::cuda::std::is_unbounded_array_v<_Tp2>) ) - _CCCL_HOST_API constexpr dynamic_shared_memory(non_portable_t) noexcept - : __non_portable_{true} - {} - - _CCCL_TEMPLATE(class _Tp2 = _Tp) - _CCCL_REQUIRES(::cuda::std::is_unbounded_array_v<_Tp2>) - _CCCL_HOST_API constexpr dynamic_shared_memory(::cuda::std::size_t __n) - : __base_type{__n} - { - if (__n * sizeof(value_type) > __max_portable_dyn_smem_size) - { - ::cuda::std::__throw_invalid_argument("portable dynamic shared memory limit exceeded"); - } - } - - _CCCL_TEMPLATE(class _Tp2 = _Tp) - _CCCL_REQUIRES(::cuda::std::is_unbounded_array_v<_Tp2>) - _CCCL_HOST_API constexpr dynamic_shared_memory(::cuda::std::size_t __n, non_portable_t) noexcept - : __base_type{__n} - , __non_portable_{true} - {} - //! @brief Gets the size of the dynamic shared memory in bytes. [[nodiscard]] _CCCL_API constexpr ::cuda::std::size_t size_bytes() const noexcept { @@ -328,11 +291,32 @@ class _CCCL_DECLSPEC_EMPTY_BASES dynamic_shared_memory return view_type{__ptr, __base_type::__n_}; } } + + // Helper function to access private constructors + static constexpr dynamic_shared_memory_option __create(bool __non_portable = false) noexcept + { + return dynamic_shared_memory_option{__non_portable}; + } + + static constexpr dynamic_shared_memory_option __create(::cuda::std::size_t __n, bool __non_portable = false) noexcept + { + return dynamic_shared_memory_option{__n, __non_portable}; + } + +private: + _CCCL_HOST_API constexpr dynamic_shared_memory_option(bool __non_portable = false) noexcept + : __non_portable_{__non_portable} + {} + + _CCCL_HOST_API constexpr dynamic_shared_memory_option(::cuda::std::size_t __n, bool __non_portable = false) noexcept + : __base_type{__n} + , __non_portable_{__non_portable} + {} }; template [[nodiscard]] ::cudaError_t __apply_launch_option( - const dynamic_shared_memory<_Tp>& __opt, ::CUlaunchConfig& __config, ::CUfunction __kernel) noexcept + const dynamic_shared_memory_option<_Tp>& __opt, ::CUlaunchConfig& __config, ::CUfunction __kernel) noexcept { ::cudaError_t __status = ::cudaSuccess; @@ -389,6 +373,69 @@ template return ::cudaSuccess; } +/** + * @brief Function that creates dynamic_shared_memory_option for non-unbounded array types + * + * @tparam _Tp Type intended to be stored in dynamic shared memory (must not be an unbounded array) + * @return dynamic_shared_memory_option<_Tp> instance + */ +_CCCL_TEMPLATE(class _Tp) +_CCCL_REQUIRES((!::cuda::std::is_unbounded_array_v<_Tp>) ) +[[nodiscard]] _CCCL_HOST_API constexpr dynamic_shared_memory_option<_Tp> dynamic_shared_memory() noexcept +{ + static_assert(sizeof(_Tp) <= __max_portable_dyn_smem_size, "portable dynamic shared memory limit exceeded"); + return dynamic_shared_memory_option<_Tp>::__create(false); +} + +/** + * @brief Function that creates dynamic_shared_memory_option for non-unbounded array types with non-portable flag + * + * @tparam _Tp Type intended to be stored in dynamic shared memory (must not be an unbounded array) + * @param __non_portable Flag indicating non-portable size + * @return dynamic_shared_memory_option<_Tp> instance + */ +_CCCL_TEMPLATE(class _Tp) +_CCCL_REQUIRES((!::cuda::std::is_unbounded_array_v<_Tp>) ) +[[nodiscard]] _CCCL_HOST_API constexpr dynamic_shared_memory_option<_Tp> dynamic_shared_memory(non_portable_t) noexcept +{ + return dynamic_shared_memory_option<_Tp>::__create(true); +} + +/** + * @brief Function that creates dynamic_shared_memory_option for unbounded array types + * + * @tparam _Tp Unbounded array type + * @param __n Number of elements in the dynamic shared memory + * @return dynamic_shared_memory_option<_Tp> instance + */ +_CCCL_TEMPLATE(class _Tp) +_CCCL_REQUIRES(::cuda::std::is_unbounded_array_v<_Tp>) +[[nodiscard]] _CCCL_HOST_API constexpr dynamic_shared_memory_option<_Tp> dynamic_shared_memory(::cuda::std::size_t __n) +{ + using value_type = typename dynamic_shared_memory_option<_Tp>::value_type; + if (__n * sizeof(value_type) > __max_portable_dyn_smem_size) + { + ::cuda::std::__throw_invalid_argument("portable dynamic shared memory limit exceeded"); + } + return dynamic_shared_memory_option<_Tp>::__create(__n, false); +} + +/** + * @brief Function that creates dynamic_shared_memory_option for unbounded array types with non-portable flag + * + * @tparam _Tp Unbounded array type + * @param __n Number of elements in the dynamic shared memory + * @param __non_portable Flag indicating non-portable size + * @return dynamic_shared_memory_option<_Tp> instance + */ +_CCCL_TEMPLATE(class _Tp) +_CCCL_REQUIRES(::cuda::std::is_unbounded_array_v<_Tp>) +[[nodiscard]] _CCCL_HOST_API constexpr dynamic_shared_memory_option<_Tp> +dynamic_shared_memory(::cuda::std::size_t __n, non_portable_t) noexcept +{ + return dynamic_shared_memory_option<_Tp>::__create(__n, true); +} + /** * @brief Launch option specifying launch priority * @@ -733,7 +780,7 @@ template # if _CCCL_CUDA_COMPILATION() template -_CCCL_DEVICE_API decltype(auto) dynamic_shared_memory_view(const kernel_config<_Dims, _Opts...>& __config) noexcept +_CCCL_DEVICE_API decltype(auto) dynamic_shared_memory(const kernel_config<_Dims, _Opts...>& __config) noexcept { auto& __opt = __detail::find_option_in_tuple<__detail::launch_option_kind::dynamic_shared_memory>(__config.options); using _Opt = ::cuda::std::remove_reference_t; diff --git a/libcudacxx/test/libcudacxx/cuda/ccclrt/launch/dynamic_shared_memory.cu b/libcudacxx/test/libcudacxx/cuda/ccclrt/launch/dynamic_shared_memory.cu index c0eaab93c05f..4b373a5b665d 100644 --- a/libcudacxx/test/libcudacxx/cuda/ccclrt/launch/dynamic_shared_memory.cu +++ b/libcudacxx/test/libcudacxx/cuda/ccclrt/launch/dynamic_shared_memory.cu @@ -25,10 +25,10 @@ struct TestKernel template __device__ void operator()(const Config& config) { - static_assert(cuda::std::is_same_v); - static_assert(noexcept(cuda::dynamic_shared_memory_view(config))); + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::dynamic_shared_memory(config))); - write_smem(cuda::dynamic_shared_memory_view(config)); + write_smem(cuda::dynamic_shared_memory(config)); } __device__ void write_smem(T& view) diff --git a/libcudacxx/test/libcudacxx/cuda/ccclrt/launch/launch_smoke.cu b/libcudacxx/test/libcudacxx/cuda/ccclrt/launch/launch_smoke.cu index 584aa4c2d1e6..7bcace845bd0 100644 --- a/libcudacxx/test/libcudacxx/cuda/ccclrt/launch/launch_smoke.cu +++ b/libcudacxx/test/libcudacxx/cuda/ccclrt/launch/launch_smoke.cu @@ -82,7 +82,7 @@ struct dynamic_smem_single template __device__ void operator()(Config config) { - decltype(auto) dynamic_smem = cuda::dynamic_shared_memory_view(config); + decltype(auto) dynamic_smem = cuda::dynamic_shared_memory(config); static_assert(::cuda::std::is_same_v); CCCLRT_REQUIRE_DEVICE(::cuda::device::is_object_from(dynamic_smem, ::cuda::device::address_space::shared)); kernel_run_proof = true; @@ -95,7 +95,7 @@ struct dynamic_smem_span template __device__ void operator()(Config config, int size) { - auto dynamic_smem = cuda::dynamic_shared_memory_view(config); + auto dynamic_smem = cuda::dynamic_shared_memory(config); static_assert(decltype(dynamic_smem)::extent == Extent); static_assert(::cuda::std::is_same_v); CCCLRT_REQUIRE_DEVICE(dynamic_smem.size() == size); From e067be9577aade432b94be26cdc91bb3525faf53 Mon Sep 17 00:00:00 2001 From: pciolkosz Date: Tue, 16 Dec 2025 23:24:53 -0800 Subject: [PATCH 2/2] Forgot about cudax --- cudax/test/launch/launch_smoke.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/cudax/test/launch/launch_smoke.cu b/cudax/test/launch/launch_smoke.cu index b6baab937bc5..130f8507c857 100644 --- a/cudax/test/launch/launch_smoke.cu +++ b/cudax/test/launch/launch_smoke.cu @@ -88,7 +88,7 @@ struct dynamic_smem_single template __device__ void operator()(Config config) { - decltype(auto) dynamic_smem = cuda::dynamic_shared_memory_view(config); + decltype(auto) dynamic_smem = cuda::dynamic_shared_memory(config); static_assert(::cuda::std::is_same_v); CUDAX_REQUIRE(::cuda::device::is_object_from(dynamic_smem, ::cuda::device::address_space::shared)); kernel_run_proof = true; @@ -101,7 +101,7 @@ struct dynamic_smem_span template __device__ void operator()(Config config, int size) { - auto dynamic_smem = cuda::dynamic_shared_memory_view(config); + auto dynamic_smem = cuda::dynamic_shared_memory(config); static_assert(decltype(dynamic_smem)::extent == Extent); static_assert(::cuda::std::is_same_v); CUDAX_REQUIRE(dynamic_smem.size() == size);