diff --git a/cudax/examples/simple_p2p.cu b/cudax/examples/simple_p2p.cu index 07a9b174f3f6..7b3b42b8a4b6 100644 --- a/cudax/examples/simple_p2p.cu +++ b/cudax/examples/simple_p2p.cu @@ -52,7 +52,7 @@ struct simple_kernel __device__ void operator()(Configuration config, ::cuda::std::span src, ::cuda::std::span dst) { // Just a dummy kernel, doing enough for us to verify that everything worked - const auto idx = config.dims.rank(cuda::thread); + const auto idx = config.dims.rank(cuda::gpu_thread); dst[idx] = src[idx] * 2.0f; } }; diff --git a/cudax/include/cuda/experimental/__execution/queries.cuh b/cudax/include/cuda/experimental/__execution/queries.cuh index 9a36b0217641..f89881c7d199 100644 --- a/cudax/include/cuda/experimental/__execution/queries.cuh +++ b/cudax/include/cuda/experimental/__execution/queries.cuh @@ -26,6 +26,7 @@ _CCCL_SUPPRESS_DEPRECATED_PUSH _CCCL_SUPPRESS_DEPRECATED_POP #include +#include #include #include #include diff --git a/cudax/include/cuda/experimental/__execution/stream/adaptor.cuh b/cudax/include/cuda/experimental/__execution/stream/adaptor.cuh index 5fca08ae21cf..66f8dcc402e2 100644 --- a/cudax/include/cuda/experimental/__execution/stream/adaptor.cuh +++ b/cudax/include/cuda/experimental/__execution/stream/adaptor.cuh @@ -23,6 +23,7 @@ #include #include +#include #include #include #include @@ -269,7 +270,7 @@ private: // the receiver tell us how to launch the kernel. auto const __launch_config = get_launch_config(execution::get_env(__state.__state_.__rcvr_)); using __launch_dims_t = decltype(__launch_config.dims); - constexpr int __block_threads = __launch_dims_t::static_count(thread, block); + constexpr int __block_threads = __launch_dims_t::static_count(gpu_thread, block); // Start the child operation state. This will launch kernels for all the predecessors // of this operation. @@ -295,7 +296,7 @@ private: _CCCL_DEVICE_API void __device_start() noexcept { using __launch_dims_t = __dims_of_t<__rcvr_config_t>; - constexpr int __block_threads = __launch_dims_t::static_count(thread, block); + constexpr int __block_threads = __launch_dims_t::static_count(gpu_thread, block); auto& __state = __get_state(); // without the following, the kernel in __host_start will fail to launch with diff --git a/cudax/include/cuda/experimental/__execution/stream/scheduler.cuh b/cudax/include/cuda/experimental/__execution/stream/scheduler.cuh index 09f866b2f4a6..7c83a6a25837 100644 --- a/cudax/include/cuda/experimental/__execution/stream/scheduler.cuh +++ b/cudax/include/cuda/experimental/__execution/stream/scheduler.cuh @@ -134,7 +134,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT stream_scheduler // the completion kernel, we will be completing the parent's receiver, so we must let // the receiver tell us how to launch the kernel. auto const __launch_dims = get_launch_config(execution::get_env(__rcvr_)).dims; - constexpr int __block_threads = decltype(__launch_dims)::static_count(cuda::thread, cuda::block); + constexpr int __block_threads = decltype(__launch_dims)::static_count(cuda::gpu_thread, cuda::block); int const __grid_blocks = __launch_dims.count(cuda::block, cuda::grid); static_assert(__block_threads != ::cuda::std::dynamic_extent); @@ -153,7 +153,7 @@ struct _CCCL_TYPE_VISIBILITY_DEFAULT stream_scheduler _CCCL_DEVICE_API void __device_start() noexcept { using __launch_dims_t = decltype(get_launch_config(execution::get_env(__rcvr_)).dims); - constexpr int __block_threads = __launch_dims_t::static_count(cuda::thread, cuda::block); + constexpr int __block_threads = __launch_dims_t::static_count(cuda::gpu_thread, cuda::block); // without the following, the kernel in __host_start will fail to launch with // cudaErrorInvalidDeviceFunction. diff --git a/cudax/include/cuda/experimental/__launch/launch.cuh b/cudax/include/cuda/experimental/__launch/launch.cuh index c3d042f49ebb..09519362f7b5 100644 --- a/cudax/include/cuda/experimental/__launch/launch.cuh +++ b/cudax/include/cuda/experimental/__launch/launch.cuh @@ -22,6 +22,7 @@ #endif // no system header #include +#include #include #include #include @@ -108,7 +109,7 @@ _CCCL_HOST_API auto __launch_impl(_Dst&& __dst, _Config __conf, ::CUfunction __k static_assert(!::cuda::std::is_same_v, "Can't launch a configuration without hierarchy dimensions"); ::CUlaunchConfig __config{}; - constexpr bool __has_cluster_level = has_level; + constexpr bool __has_cluster_level = has_level_v; constexpr unsigned int __num_attrs_needed = ::cuda::__detail::kernel_config_count_attr_space(__conf) + __has_cluster_level; ::CUlaunchAttribute __attrs[__num_attrs_needed == 0 ? 1 : __num_attrs_needed]; @@ -124,9 +125,9 @@ _CCCL_HOST_API auto __launch_impl(_Dst&& __dst, _Config __conf, ::CUfunction __k __config.gridDimX = static_cast(__conf.dims.extents(block, grid).x); __config.gridDimY = static_cast(__conf.dims.extents(block, grid).y); __config.gridDimZ = static_cast(__conf.dims.extents(block, grid).z); - __config.blockDimX = static_cast(__conf.dims.extents(thread, block).x); - __config.blockDimY = static_cast(__conf.dims.extents(thread, block).y); - __config.blockDimZ = static_cast(__conf.dims.extents(thread, block).z); + __config.blockDimX = static_cast(__conf.dims.extents(gpu_thread, block).x); + __config.blockDimY = static_cast(__conf.dims.extents(gpu_thread, block).y); + __config.blockDimZ = static_cast(__conf.dims.extents(gpu_thread, block).z); if constexpr (__has_cluster_level) { diff --git a/cudax/test/common/host_device.cuh b/cudax/test/common/host_device.cuh index ff4744d0ccfb..e74bc42ff1ad 100644 --- a/cudax/test/common/host_device.cuh +++ b/cudax/test/common/host_device.cuh @@ -67,10 +67,10 @@ void test_host_dev(const Dims& dims, const Lambda& lambda, const Filters&... fil cudaLaunchAttribute attrs[1]; config.attrs = &attrs[0]; - config.blockDim = dims.extents(cuda::thread, cuda::block); + config.blockDim = dims.extents(cuda::gpu_thread, cuda::block); config.gridDim = dims.extents(cuda::block, cuda::grid); - if constexpr (cuda::has_level) + if constexpr (cuda::has_level_v) { dim3 cluster_dims = dims.extents(cuda::block, cuda::cluster); config.attrs[config.numAttrs].id = cudaLaunchAttributeClusterDimension; diff --git a/cudax/test/launch/launch_smoke.cu b/cudax/test/launch/launch_smoke.cu index b6baab937bc5..5ade1c0cf03f 100644 --- a/cudax/test/launch/launch_smoke.cu +++ b/cudax/test/launch/launch_smoke.cu @@ -55,7 +55,7 @@ struct functor_taking_config template __device__ void operator()(Config config, int grid_size) { - static_assert(config.dims.static_count(cuda::thread, cuda::block) == BlockSize); + static_assert(config.dims.static_count(cuda::gpu_thread, cuda::block) == BlockSize); CUDAX_REQUIRE(config.dims.count(cuda::block, cuda::grid) == grid_size); kernel_run_proof = true; } @@ -248,7 +248,7 @@ void launch_smoke_test(StreamOrPathBuilder& dst) // Lambda { cudax::launch(dst, cuda::block_dims<256>() & cuda::grid_dims(1), [] __device__(auto config) { - if (config.dims.rank(cuda::thread, cuda::block) == 0) + if (config.dims.rank(cuda::gpu_thread, cuda::block) == 0) { printf("Hello from the GPU\n"); kernel_run_proof = true; @@ -354,7 +354,7 @@ void test_default_config() auto block = cuda::block_dims<256>; auto verify_lambda = [] __device__(auto config) { - static_assert(config.dims.count(cuda::thread, cuda::block) == 256); + static_assert(config.dims.count(cuda::gpu_thread, cuda::block) == 256); CUDAX_REQUIRE(config.dims.count(cuda::block) == 4); cooperative_groups::this_grid().sync(); }; diff --git a/libcudacxx/include/cuda/__fwd/hierarchy.h b/libcudacxx/include/cuda/__fwd/hierarchy.h new file mode 100644 index 000000000000..54cbeadeca15 --- /dev/null +++ b/libcudacxx/include/cuda/__fwd/hierarchy.h @@ -0,0 +1,69 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2024 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA___FWD_HIERARCHY_H +#define _CUDA___FWD_HIERARCHY_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include + +#include + +_CCCL_BEGIN_NAMESPACE_CUDA + +// hierarchy level + +template +struct hierarchy_level_base; + +template +struct __native_hierarchy_level_base; + +struct grid_level; +struct cluster_level; +struct block_level; +struct warp_level; +struct thread_level; + +template +inline constexpr bool __is_hierarchy_level_v = ::cuda::std::is_base_of_v, _Tp>; + +template +inline constexpr bool __is_native_hierarchy_level_v = + ::cuda::std::is_base_of_v<__native_hierarchy_level_base<_Tp>, _Tp>; + +// hierarchy + +template +struct hierarchy_dimensions; + +template +inline constexpr bool __is_hierarchy_v = false; +template +inline constexpr bool __is_hierarchy_v> = true; + +template +struct allowed_levels; + +_CCCL_END_NAMESPACE_CUDA + +#include + +#endif // _CUDA___FWD_HIERARCHY_H diff --git a/libcudacxx/include/cuda/__hierarchy/block_level.h b/libcudacxx/include/cuda/__hierarchy/block_level.h new file mode 100644 index 000000000000..cfa000df0d25 --- /dev/null +++ b/libcudacxx/include/cuda/__hierarchy/block_level.h @@ -0,0 +1,141 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA___HIERARCHY_BLOCK_LEVEL_H +#define _CUDA___HIERARCHY_BLOCK_LEVEL_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include +#include + +#include + +_CCCL_BEGIN_NAMESPACE_CUDA + +struct block_level : __native_hierarchy_level_base +{ + using product_type = unsigned; + using allowed_above = allowed_levels; + using allowed_below = allowed_levels; + + using __next_native_level = cluster_level; + + using __base_type = __native_hierarchy_level_base; + using __base_type::count_as; + using __base_type::extents_as; + using __base_type::index_as; + using __base_type::rank_as; + + // interactions with cluster level + + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp>) + [[nodiscard]] _CCCL_DEVICE_API static ::cuda::std::dims<3, _Tp> extents_as(const cluster_level&) noexcept + { + ::dim3 __dims{1u, 1u, 1u}; + NV_IF_TARGET(NV_PROVIDES_SM_90, (__dims = ::__clusterDim();)) + return ::cuda::std::dims<3, _Tp>{static_cast<_Tp>(__dims.x), static_cast<_Tp>(__dims.y), static_cast<_Tp>(__dims.z)}; + } + + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp>) + [[nodiscard]] _CCCL_DEVICE_API static _Tp count_as(const cluster_level&) noexcept + { + unsigned __count = 1; + NV_IF_TARGET(NV_PROVIDES_SM_90, (__count = ::__clusterSizeInBlocks();)) + return static_cast<_Tp>(__count); + } + + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp>) + [[nodiscard]] _CCCL_DEVICE_API static hierarchy_query_result<_Tp> index_as(const cluster_level&) noexcept + { + ::dim3 __idx{0u, 0u, 0u}; + NV_IF_TARGET(NV_PROVIDES_SM_90, (__idx = ::__clusterRelativeBlockIdx();)) + return {static_cast<_Tp>(__idx.x), static_cast<_Tp>(__idx.y), static_cast<_Tp>(__idx.z)}; + } + + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp>) + [[nodiscard]] _CCCL_DEVICE_API static _Tp rank_as(const cluster_level&) noexcept + { + unsigned __rank = 0; + NV_IF_TARGET(NV_PROVIDES_SM_90, (__rank = ::__clusterRelativeBlockRank();)) + return static_cast<_Tp>(__rank); + } + + // interactions with grid level + + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp>) + [[nodiscard]] _CCCL_DEVICE_API static ::cuda::std::dims<3, _Tp> extents_as(const grid_level&) noexcept + { + return ::cuda::std::dims<3, _Tp>{ + static_cast<_Tp>(gridDim.x), static_cast<_Tp>(gridDim.y), static_cast<_Tp>(gridDim.z)}; + } + + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp>) + [[nodiscard]] _CCCL_DEVICE_API static _Tp count_as(const grid_level&) noexcept + { + return static_cast<_Tp>(gridDim.x) * static_cast<_Tp>(gridDim.y) * static_cast<_Tp>(gridDim.z); + } + + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp>) + [[nodiscard]] _CCCL_DEVICE_API static hierarchy_query_result<_Tp> index_as(const grid_level&) noexcept + { + return {static_cast<_Tp>(blockIdx.x), static_cast<_Tp>(blockIdx.y), static_cast<_Tp>(blockIdx.z)}; + } + + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp>) + [[nodiscard]] _CCCL_DEVICE_API static _Tp rank_as(const grid_level& __level) noexcept + { + const auto __dims = dims_as<_Tp>(__level); + const auto __idx = index_as<_Tp>(__level); + return static_cast<_Tp>((__idx.z * __dims.y + __idx.y) * __dims.x + __idx.x); + } + + // interactions with grid level in hierarchy + + _CCCL_TEMPLATE(class _Tp, class _Hierarchy) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp> _CCCL_AND __is_hierarchy_v<_Hierarchy>) + [[nodiscard]] _CCCL_DEVICE_API static _Tp rank_as(const grid_level& __level, const _Hierarchy& __hier) noexcept + { + static_assert(has_unit_or_level_v, "_Hierarchy doesn't contain block level"); + static_assert(has_level_v, "_Hierarchy doesn't contain grid level"); + + const auto __dims = dims_as<_Tp>(__level, __hier); + const auto __idx = index_as<_Tp>(__level, __hier); + return static_cast<_Tp>((__idx.z * __dims.y + __idx.y) * __dims.x + __idx.x); + } +}; + +_CCCL_GLOBAL_CONSTANT block_level block; + +_CCCL_END_NAMESPACE_CUDA + +#include + +#endif // _CUDA___HIERARCHY_BLOCK_LEVEL_H diff --git a/libcudacxx/include/cuda/__hierarchy/cluster_level.h b/libcudacxx/include/cuda/__hierarchy/cluster_level.h new file mode 100644 index 000000000000..70ef8fdce89c --- /dev/null +++ b/libcudacxx/include/cuda/__hierarchy/cluster_level.h @@ -0,0 +1,74 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA___HIERARCHY_CLUSTER_LEVEL_H +#define _CUDA___HIERARCHY_CLUSTER_LEVEL_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include +#include + +#include + +_CCCL_BEGIN_NAMESPACE_CUDA + +struct cluster_level : __native_hierarchy_level_base +{ + using product_type = unsigned; + using allowed_above = allowed_levels; + using allowed_below = allowed_levels; + + using __next_native_level = grid_level; + + using __base_type = __native_hierarchy_level_base; + using __base_type::extents_as; + using __base_type::index_as; + + // interactions with grid level + + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp>) + [[nodiscard]] _CCCL_DEVICE_API static ::cuda::std::dims<3, _Tp> extents_as(const grid_level&) noexcept + { + ::dim3 __dims{gridDim}; + NV_IF_TARGET(NV_PROVIDES_SM_90, (__dims = ::__clusterGridDimInClusters();)) + return ::cuda::std::dims<3, _Tp>{static_cast<_Tp>(__dims.x), static_cast<_Tp>(__dims.y), static_cast<_Tp>(__dims.z)}; + } + + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp>) + [[nodiscard]] _CCCL_DEVICE_API static hierarchy_query_result<_Tp> index_as(const grid_level&) noexcept + { + ::dim3 __idx{blockIdx}; + NV_IF_TARGET(NV_PROVIDES_SM_90, (__idx = ::__clusterIdx();)) + return {static_cast<_Tp>(__idx.x), static_cast<_Tp>(__idx.y), static_cast<_Tp>(__idx.z)}; + } +}; + +_CCCL_GLOBAL_CONSTANT cluster_level cluster; + +_CCCL_END_NAMESPACE_CUDA + +#include + +#endif // _CUDA___HIERARCHY_CLUSTER_LEVEL_H diff --git a/libcudacxx/include/cuda/__hierarchy/dimensions.h b/libcudacxx/include/cuda/__hierarchy/dimensions.h index 2f52128b4fc5..395c4c6e996e 100644 --- a/libcudacxx/include/cuda/__hierarchy/dimensions.h +++ b/libcudacxx/include/cuda/__hierarchy/dimensions.h @@ -11,6 +11,16 @@ #ifndef _CUDA___HIERARCHY_DIMENSIONS_H #define _CUDA___HIERARCHY_DIMENSIONS_H +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + #include #include @@ -21,8 +31,7 @@ _CCCL_BEGIN_NAMESPACE_CUDA template using dimensions = ::cuda::std::extents<_Tp, _Extents...>; -// not unsigned because of a bug in ::cuda::std::extents -using dimensions_index_type = int; +using dimensions_index_type = unsigned; /** * @brief Type representing a result of a multi-dimensional hierarchy query. @@ -55,19 +64,19 @@ using dimensions_index_type = int; * Extents of the result */ template -struct hierarchy_query_result : public dimensions<_Tp, _Extents...> +struct hierarchy_query_result_org : public dimensions<_Tp, _Extents...> { using _Dims = dimensions<_Tp, _Extents...>; using _Dims::_Dims; - _CCCL_API constexpr hierarchy_query_result() + _CCCL_API constexpr hierarchy_query_result_org() : _Dims() , x(_Dims::extent(0)) , y(_Dims::rank() > 1 ? _Dims::extent(1) : 1) , z(_Dims::rank() > 2 ? _Dims::extent(2) : 1) {} - _CCCL_API explicit constexpr hierarchy_query_result(const _Dims& dims) + _CCCL_API explicit constexpr hierarchy_query_result_org(const _Dims& dims) : _Dims(dims) , x(_Dims::extent(0)) , y(_Dims::rank() > 1 ? _Dims::extent(1) : 1) @@ -135,7 +144,7 @@ __dims_sum(const dimensions<_T1, _E1...>& __h1, const dimensions<_T2, _E2...>& _ template [[nodiscard]] _CCCL_API constexpr auto __convert_to_query_result(const dimensions<_Tp, _Extents...>& __result) { - return hierarchy_query_result<_Tp, _Extents...>(__result); + return hierarchy_query_result_org<_Tp, _Extents...>(__result); } [[nodiscard]] _CCCL_API constexpr auto __dim3_to_dims(const ::dim3& dims) diff --git a/libcudacxx/include/cuda/__hierarchy/get_launch_dimensions.h b/libcudacxx/include/cuda/__hierarchy/get_launch_dimensions.h new file mode 100644 index 000000000000..ee3d07ee8ae7 --- /dev/null +++ b/libcudacxx/include/cuda/__hierarchy/get_launch_dimensions.h @@ -0,0 +1,90 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA___HIERARCHY_GET_LAUNCH_DIMENSIONS_H +#define _CUDA___HIERARCHY_GET_LAUNCH_DIMENSIONS_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include +#include + +#include + +_CCCL_BEGIN_NAMESPACE_CUDA + +/** + * @brief Returns a tuple of dim3 compatible objects that can be used to launch + * a kernel + * + * This function returns a tuple of hierarchy_query_result objects that contain + * dimensions from the supplied hierarchy, that can be used to launch that + * hierarchy. It is meant to allow for easy usage of hierarchy dimensions with + * the <<<>>> launch syntax or cudaLaunchKernelEx in case of a cluster launch. + * Contained hierarchy_query_result objects are results of extents() member + * function on the hierarchy passed in. The returned tuple has three elements if + * cluster_level is present in the hierarchy (extents(block, grid), + * extents(cluster, block), extents(thread, block)). Otherwise it contains only + * two elements, without the middle one related to the cluster. + * + * @par Snippet + * @code + * #include + * + * using namespace cuda; + * + * auto hierarchy = make_hierarchy(grid_dims(256), cluster_dims<4>(), + * block_dims<8, 8, 8>()); auto [grid_dimensions, cluster_dimensions, + * block_dimensions] = get_launch_dimensions(hierarchy); + * assert(grid_dimensions.x == 256); + * assert(cluster_dimensions.x == 4); + * assert(block_dimensions.x == 8); + * assert(block_dimensions.y == 8); + * assert(block_dimensions.z == 8); + * @endcode + * @par + * + * @param hierarchy + * Hierarchy that the launch dimensions are requested for + */ +template +constexpr auto _CCCL_HOST get_launch_dimensions(const hierarchy_dimensions<_Levels...>& __hierarchy) +{ + if constexpr (has_level_v>) + { + return ::cuda::std::make_tuple( + __hierarchy.extents(block_level{}, grid_level{}), + __hierarchy.extents(block_level{}, cluster_level{}), + __hierarchy.extents(thread_level{}, block_level{})); + } + else + { + return ::cuda::std::make_tuple( + __hierarchy.extents(block_level{}, grid_level{}), __hierarchy.extents(gpu_thread, block_level{})); + } +} + +_CCCL_END_NAMESPACE_CUDA + +#include + +#endif // _CUDA___HIERARCHY_GET_LAUNCH_DIMENSIONS_H diff --git a/libcudacxx/include/cuda/__hierarchy/grid_level.h b/libcudacxx/include/cuda/__hierarchy/grid_level.h new file mode 100644 index 000000000000..9cd6a111709d --- /dev/null +++ b/libcudacxx/include/cuda/__hierarchy/grid_level.h @@ -0,0 +1,44 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA___HIERARCHY_GRID_LEVEL_H +#define _CUDA___HIERARCHY_GRID_LEVEL_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include + +#include + +_CCCL_BEGIN_NAMESPACE_CUDA + +struct grid_level : __native_hierarchy_level_base +{ + using product_type = unsigned long long; + using allowed_above = allowed_levels<>; + using allowed_below = allowed_levels; +}; + +_CCCL_GLOBAL_CONSTANT grid_level grid; + +_CCCL_END_NAMESPACE_CUDA + +#include + +#endif // _CUDA___HIERARCHY_GRID_LEVEL_H diff --git a/libcudacxx/include/cuda/__hierarchy/hierarchy_dimensions.h b/libcudacxx/include/cuda/__hierarchy/hierarchy_dimensions.h index 0525c5671742..1d375843d833 100644 --- a/libcudacxx/include/cuda/__hierarchy/hierarchy_dimensions.h +++ b/libcudacxx/include/cuda/__hierarchy/hierarchy_dimensions.h @@ -11,7 +11,19 @@ #ifndef _CUDA___HIERARCHY_HIERARCHY_DIMENSIONS_H #define _CUDA___HIERARCHY_HIERARCHY_DIMENSIONS_H +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include #include +#include #include #include #include @@ -43,12 +55,6 @@ template } } // namespace __detail -template -using __level_type_of = typename _Level::level_type; - -template -struct hierarchy_dimensions; - namespace __detail { // Function to sometimes convince the compiler something is a constexpr and not @@ -67,22 +73,6 @@ template } } -template -struct __has_level_helper; - -template -struct __has_level_helper<_QueryLevel, hierarchy_dimensions<_Unit, _Levels...>> - : public ::cuda::std::__fold_or<::cuda::std::is_same_v<_QueryLevel, __level_type_of<_Levels>>...> -{}; - -template -struct __has_unit -{}; - -template -struct __has_unit<_QueryLevel, hierarchy_dimensions<_Unit, _Levels...>> : ::cuda::std::is_same<_QueryLevel, _Unit> -{}; - template struct __get_level_helper { @@ -101,15 +91,6 @@ struct __get_level_helper }; } // namespace __detail -template -inline constexpr bool has_level = - __detail::__has_level_helper<_QueryLevel, ::cuda::std::remove_cvref_t<_Hierarchy>>::value; - -template -inline constexpr bool has_level_or_unit = - __detail::__has_level_helper<_QueryLevel, ::cuda::std::remove_cvref_t<_Hierarchy>>::value - || __detail::__has_unit<_QueryLevel, ::cuda::std::remove_cvref_t<_Hierarchy>>::value; - namespace __detail { template @@ -377,7 +358,7 @@ struct __empty_hierarchy template struct hierarchy_dimensions { - static_assert(::cuda::std::is_base_of_v || ::cuda::std::is_same_v<_BottomUnit, void>); + static_assert(__is_hierarchy_level_v<_BottomUnit> || ::cuda::std::is_same_v<_BottomUnit, void>); ::cuda::std::tuple<_Levels...> levels; _CCCL_API constexpr hierarchy_dimensions(const _Levels&... __ls) noexcept @@ -419,8 +400,8 @@ struct hierarchy_dimensions [[nodiscard]] _CCCL_API static constexpr auto levels_range_static(const ::cuda::std::tuple<_Levels...>& __levels) noexcept { - static_assert(has_level<_Level, hierarchy_dimensions<_BottomUnit, _Levels...>>); - static_assert(has_level_or_unit<_Unit, hierarchy_dimensions<_BottomUnit, _Levels...>>); + static_assert(has_level_v<_Level, hierarchy_dimensions<_BottomUnit, _Levels...>>); + static_assert(has_unit_or_level_v<_Unit, hierarchy_dimensions<_BottomUnit, _Levels...>>); static_assert(__detail::__legal_unit_for_level<_Unit, _Level>); auto __fn = __detail::__get_levels_range<_Level, _Unit, _Levels...>; return ::cuda::std::apply(__fn, __levels); @@ -780,7 +761,7 @@ struct hierarchy_dimensions template _CCCL_API constexpr auto level(const _Level&) const noexcept { - static_assert(has_level<_Level, hierarchy_dimensions<_BottomUnit, _Levels...>>); + static_assert(has_level_v<_Level, hierarchy_dimensions<_BottomUnit, _Levels...>>); return ::cuda::std::apply(__detail::__get_level_helper<_Level>{}, levels); } @@ -811,8 +792,8 @@ struct hierarchy_dimensions // block) return ::cuda::std::apply(fragment_helper<_OtherUnit>(), ::cuda::std::tuple_cat(levels, __other.levels)); } - else if constexpr (has_level<__this_bottom_level, hierarchy_dimensions<_OtherUnit, _OtherLevels...>> - && (!has_level<__this_top_level, hierarchy_dimensions<_OtherUnit, _OtherLevels...>> + else if constexpr (has_level_v<__this_bottom_level, hierarchy_dimensions<_OtherUnit, _OtherLevels...>> + && (!has_level_v<__this_top_level, hierarchy_dimensions<_OtherUnit, _OtherLevels...>> || ::cuda::std::is_same_v<__this_top_level, __other_top_level>) ) { // Overlap with this on the top, e.g. this is (grid, cluster), other is @@ -837,8 +818,8 @@ struct hierarchy_dimensions { // Overlap with this on the bottom, e.g. this is (cluster, block), other // is (grid, cluster), can fully overlap - static_assert(has_level<__other_bottom_level, hierarchy_dimensions<_BottomUnit, _Levels...>> - && (!has_level<__this_bottom_level, hierarchy_dimensions<_OtherUnit, _OtherLevels...>> + static_assert(has_level_v<__other_bottom_level, hierarchy_dimensions<_BottomUnit, _Levels...>> + && (!has_level_v<__this_bottom_level, hierarchy_dimensions<_OtherUnit, _OtherLevels...>> || ::cuda::std::is_same_v<__this_bottom_level, __other_bottom_level>), "Can't combine the hierarchies"); @@ -856,54 +837,6 @@ struct hierarchy_dimensions #endif // _CCCL_DOXYGEN_INVOKED }; -/** - * @brief Returns a tuple of dim3 compatible objects that can be used to launch - * a kernel - * - * This function returns a tuple of hierarchy_query_result objects that contain - * dimensions from the supplied hierarchy, that can be used to launch that - * hierarchy. It is meant to allow for easy usage of hierarchy dimensions with - * the <<<>>> launch syntax or cudaLaunchKernelEx in case of a cluster launch. - * Contained hierarchy_query_result objects are results of extents() member - * function on the hierarchy passed in. The returned tuple has three elements if - * cluster_level is present in the hierarchy (extents(block, grid), - * extents(cluster, block), extents(thread, block)). Otherwise it contains only - * two elements, without the middle one related to the cluster. - * - * @par Snippet - * @code - * #include - * - * using namespace cuda; - * - * auto hierarchy = make_hierarchy(grid_dims(256), cluster_dims<4>(), - * block_dims<8, 8, 8>()); auto [grid_dimensions, cluster_dimensions, - * block_dimensions] = get_launch_dimensions(hierarchy); - * assert(grid_dimensions.x == 256); - * assert(cluster_dimensions.x == 4); - * assert(block_dimensions.x == 8); - * assert(block_dimensions.y == 8); - * assert(block_dimensions.z == 8); - * @endcode - * @par - * - * @param hierarchy - * Hierarchy that the launch dimensions are requested for - */ -template -constexpr auto _CCCL_HOST get_launch_dimensions(const hierarchy_dimensions<_Levels...>& __hierarchy) -{ - if constexpr (has_level>) - { - return ::cuda::std::make_tuple( - __hierarchy.extents(block, grid), __hierarchy.extents(block, cluster), __hierarchy.extents(thread, block)); - } - else - { - return ::cuda::std::make_tuple(__hierarchy.extents(block, grid), __hierarchy.extents(thread, block)); - } -} - // TODO consider having LUnit optional argument for template argument deduction /** * @brief Creates a hierarchy from passed in levels. diff --git a/libcudacxx/include/cuda/__hierarchy/hierarchy_level_base.h b/libcudacxx/include/cuda/__hierarchy/hierarchy_level_base.h new file mode 100644 index 000000000000..ccf8e319fb77 --- /dev/null +++ b/libcudacxx/include/cuda/__hierarchy/hierarchy_level_base.h @@ -0,0 +1,339 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA___HIERARCHY_HIERARCHY_LEVEL_BASE_H +#define _CUDA___HIERARCHY_HIERARCHY_LEVEL_BASE_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +_CCCL_BEGIN_NAMESPACE_CUDA + +[[nodiscard]] _CCCL_API _CCCL_CONSTEVAL ::cuda::std::size_t +__hierarchy_static_extents_mul_helper(::cuda::std::size_t __lhs, ::cuda::std::size_t __rhs) noexcept +{ + if (__lhs == ::cuda::std::dynamic_extent || __rhs == ::cuda::std::dynamic_extent) + { + return ::cuda::std::dynamic_extent; + } + else + { + return __lhs * __rhs; + } +} + +template +[[nodiscard]] _CCCL_API constexpr auto __hierarchy_static_extents_mul(::cuda::std::index_sequence<_Is...>) noexcept +{ + return ::cuda::std::extents< + _ResultIndex, + ::cuda::__hierarchy_static_extents_mul_helper((_Is < _LhsExts::rank()) ? _LhsExts::static_extent(_Is) : 1, + (_Is < _RhsExts::rank()) ? _RhsExts::static_extent(_Is) : 1)...>{}; +} + +//! @brief Multiplies 2 extents in column major order together, returning a new extents type. If the ranks don't match, +//! the extent with lower rank is padded with 1s on the right to match the rank of the other. +//! +//! @param __lhs The left hand side extents to multiply. +//! @param __rhs The right hand side extents to multiply. +//! +//! @return The result of multiplying the extents together. +template +[[nodiscard]] _CCCL_API constexpr auto +__hierarchy_extents_mul(const ::cuda::std::extents<_Index, _LhsExts...>& __lhs, + const ::cuda::std::extents<_Index, _RhsExts...>& __rhs) noexcept +{ + using _Lhs = ::cuda::std::extents<_Index, _LhsExts...>; + using _Rhs = ::cuda::std::extents<_Index, _RhsExts...>; + + constexpr auto __rank = ::cuda::std::max(_Lhs::rank(), _Rhs::rank()); + using _Ret = + decltype(::cuda::__hierarchy_static_extents_mul<_Index, _Lhs, _Rhs>(::cuda::std::make_index_sequence<__rank>{})); + + ::cuda::std::array<_Index, __rank> __ret{}; + for (::cuda::std::size_t __i = 0; __i < __rank; ++__i) + { + if (_Ret::static_extent(__i) == ::cuda::std::dynamic_extent) + { + __ret[__i] = static_cast<_Index>((__i < _Lhs::rank()) ? __lhs.extent(__i) : 1) + * static_cast<_Index>((__i < _Rhs::rank()) ? __rhs.extent(__i) : 1); + } + else + { + __ret[__i] = _Ret::static_extent(__i); + } + } + return _Ret{__ret}; +} + +template +[[nodiscard]] _CCCL_API constexpr ::cuda::std::extents<_Index, _StaticExts...> +__hierarchy_extents_cast(::cuda::std::extents<_OrgIndex, _StaticExts...> __org_exts) noexcept +{ + using _OrgExts = ::cuda::std::extents<_OrgIndex, _StaticExts...>; + ::cuda::std::array<_Index, _OrgExts::rank()> __ret{}; + for (::cuda::std::size_t __i = 0; __i < _OrgExts::rank(); ++__i) + { + if (_OrgExts::static_extent(__i) == ::cuda::std::dynamic_extent) + { + __ret[__i] = static_cast<_Index>(__org_exts.extent(__i)); + } + else + { + __ret[__i] = _OrgExts::static_extent(__i); + } + } + return ::cuda::std::extents<_Index, _StaticExts...>{__ret}; +} + +template +struct hierarchy_level_base +{ + using level_type = _Level; + + template + using __default_md_query_type = unsigned; + template + using __default_1d_query_type = ::cuda::std::size_t; + + _CCCL_TEMPLATE(class _InLevel, class _Hierarchy) + _CCCL_REQUIRES(__is_hierarchy_level_v<_InLevel> _CCCL_AND __is_hierarchy_v<_Hierarchy>) + [[nodiscard]] _CCCL_API static constexpr auto dims(const _InLevel& __level, const _Hierarchy& __hier) noexcept + { + return _Level::template dims_as<__default_md_query_type<_InLevel>>(__level, __hier); + } + + _CCCL_TEMPLATE(class _InLevel, class _Hierarchy) + _CCCL_REQUIRES(__is_hierarchy_level_v<_InLevel> _CCCL_AND __is_hierarchy_v<_Hierarchy>) + [[nodiscard]] _CCCL_API static constexpr auto static_dims(const _InLevel& __level, const _Hierarchy& __hier) noexcept + { + return __static_dims_impl(__level, __hier); + } + + _CCCL_TEMPLATE(class _InLevel, class _Hierarchy) + _CCCL_REQUIRES(__is_hierarchy_level_v<_InLevel> _CCCL_AND __is_hierarchy_v<_Hierarchy>) + [[nodiscard]] _CCCL_API static constexpr auto extents(const _InLevel& __level, const _Hierarchy& __hier) noexcept + { + return _Level::template extents_as<__default_md_query_type<_InLevel>>(__level, __hier); + } + + _CCCL_TEMPLATE(class _InLevel, class _Hierarchy) + _CCCL_REQUIRES(__is_hierarchy_level_v<_InLevel> _CCCL_AND __is_hierarchy_v<_Hierarchy>) + [[nodiscard]] _CCCL_API static constexpr ::cuda::std::size_t + count(const _InLevel& __level, const _Hierarchy& __hier) noexcept + { + return _Level::template count_as<__default_1d_query_type<_InLevel>>(__level, __hier); + } + +#if _CCCL_CUDA_COMPILATION() + _CCCL_TEMPLATE(class _InLevel, class _Hierarchy) + _CCCL_REQUIRES(__is_hierarchy_level_v<_InLevel> _CCCL_AND __is_hierarchy_v<_Hierarchy>) + [[nodiscard]] _CCCL_DEVICE_API static constexpr auto index(const _InLevel& __level, const _Hierarchy& __hier) noexcept + { + return _Level::template index_as<__default_md_query_type<_InLevel>>(__level, __hier); + } + + _CCCL_TEMPLATE(class _InLevel, class _Hierarchy) + _CCCL_REQUIRES(__is_hierarchy_level_v<_InLevel> _CCCL_AND __is_hierarchy_v<_Hierarchy>) + [[nodiscard]] _CCCL_DEVICE_API static constexpr ::cuda::std::size_t + rank(const _InLevel& __level, const _Hierarchy& __hier) noexcept + { + return _Level::template rank_as<__default_1d_query_type<_InLevel>>(__level, __hier); + } +#endif // _CCCL_CUDA_COMPILATION() + + _CCCL_TEMPLATE(class _Tp, class _InLevel, class _Hierarchy) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp> _CCCL_AND __is_hierarchy_level_v<_InLevel> _CCCL_AND + __is_hierarchy_v<_Hierarchy>) + [[nodiscard]] _CCCL_API static constexpr auto dims_as(const _InLevel& __level, const _Hierarchy& __hier) noexcept + { + return __dims_as_impl<_Tp>(__level, __hier); + } + + _CCCL_TEMPLATE(class _Tp, class _InLevel, class _Hierarchy) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp> _CCCL_AND __is_hierarchy_level_v<_InLevel> _CCCL_AND + __is_hierarchy_v<_Hierarchy>) + [[nodiscard]] _CCCL_API static constexpr auto extents_as(const _InLevel& __in_level, const _Hierarchy& __hier) noexcept + { + static_assert(has_unit_or_level_v<_Level, _Hierarchy>, "_Hierarchy doesn't contain _Level"); + static_assert(has_level_v<_InLevel, _Hierarchy>, "_Hierarchy doesn't contain _InLevel"); + + using _NextLevel = __next_hierarchy_level_t<_Level, _Hierarchy>; + using _CurrExts = decltype(::cuda::__hierarchy_extents_cast<_Tp>(__hier.level(_NextLevel{}).dims)); + + // Remove dependency on runtime storage. This makes the queries work for hierarchy levels with all static extents + // in constant evaluated context. + _CurrExts __curr_exts{}; + if constexpr (_CurrExts::rank_dynamic() > 0) + { + __curr_exts = ::cuda::__hierarchy_extents_cast<_Tp>(__hier.level(_NextLevel{}).dims); + } + + if constexpr (!::cuda::std::is_same_v<_NextLevel, _InLevel>) + { + const auto __next_exts = _NextLevel::template extents_as<_Tp>(__in_level, __hier); + return ::cuda::__hierarchy_extents_mul(__curr_exts, __next_exts); + } + else + { + return __curr_exts; + } + } + + _CCCL_TEMPLATE(class _Tp, class _InLevel, class _Hierarchy) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp> _CCCL_AND __is_hierarchy_level_v<_InLevel> _CCCL_AND + __is_hierarchy_v<_Hierarchy>) + [[nodiscard]] _CCCL_API static constexpr auto count_as(const _InLevel& __level, const _Hierarchy& __hier) noexcept + { + return __count_as_impl<_Tp>(__level, __hier); + } + +#if _CCCL_CUDA_COMPILATION() + _CCCL_TEMPLATE(class _Tp, class _InLevel, class _Hierarchy) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp> _CCCL_AND __is_hierarchy_level_v<_InLevel> _CCCL_AND + __is_hierarchy_v<_Hierarchy>) + [[nodiscard]] _CCCL_DEVICE_API static constexpr auto + index_as(const _InLevel& __level, const _Hierarchy& __hier) noexcept + { + static_assert(has_unit_or_level_v<_Level, _Hierarchy>, "_Hierarchy doesn't contain _Level"); + static_assert(has_level_v<_InLevel, _Hierarchy>, "_Hierarchy doesn't contain _InLevel"); + + using _NextLevel = __next_hierarchy_level_t<_Level, _Hierarchy>; + if constexpr (::cuda::std::is_same_v<_InLevel, _NextLevel>) + { + using _CurrExts = decltype(_Level::template extents_as<_Tp>(_NextLevel{}, __hier)); + auto __curr_idx = _Level::template index_as<_Tp>(_NextLevel{}); + for (::cuda::std::size_t __i = _CurrExts::rank(); __i < 3; ++__i) + { + __curr_idx[__i] = 0; + } + return __curr_idx; + } + else + { + const auto __curr_exts = _Level::template extents_as<_Tp>(_NextLevel{}, __hier); + const auto __next_idx = _NextLevel::template index_as<_Tp>(__level, __hier); + const auto __curr_idx = _Level::template index_as<_Tp>(_NextLevel{}, __hier); + + hierarchy_query_result<_Tp> __ret{}; + for (::cuda::std::size_t __i = 0; __i < 3; ++__i) + { + __ret[__i] = __curr_idx[__i] + ((__i < __curr_exts.rank()) ? __curr_exts.extent(__i) : 1) * __next_idx[__i]; + } + return __ret; + } + } + + _CCCL_TEMPLATE(class _Tp, class _InLevel, class _Hierarchy) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp> _CCCL_AND __is_hierarchy_level_v<_InLevel> _CCCL_AND + __is_hierarchy_v<_Hierarchy>) + [[nodiscard]] _CCCL_DEVICE_API static constexpr auto + rank_as(const _InLevel& __level, const _Hierarchy& __hier) noexcept + { + static_assert(has_unit_or_level_v<_Level, _Hierarchy>, "_Hierarchy doesn't contain _Level"); + static_assert(has_level_v<_InLevel, _Hierarchy>, "_Hierarchy doesn't contain _InLevel"); + + using _NextLevel = __next_hierarchy_level_t<_Level, _Hierarchy>; + + const auto __curr_exts = _Level::template extents_as<_Tp>(_NextLevel{}, __hier); + const auto __curr_idx = _Level::template index_as<_Tp>(_NextLevel{}, __hier); + + _Tp __ret = 0; + if constexpr (!::cuda::std::is_same_v<_InLevel, _NextLevel>) + { + __ret = _NextLevel::template rank_as<_Tp>(__level, __hier) * _Level::template count_as<_Tp>(_NextLevel{}, __hier); + } + + for (::cuda::std::size_t __i = __curr_exts.rank(); __i > 0; --__i) + { + _Tp __inc = __curr_idx[__i - 1]; + for (::cuda::std::size_t __j = __i - 1; __j > 0; --__j) + { + __inc *= __curr_exts.extent(__j - 1); + } + __ret += __inc; + } + return __ret; + } +#endif // _CCCL_CUDA_COMPILATION() + +private: + template + friend struct __native_hierarchy_level_base; + + _CCCL_EXEC_CHECK_DISABLE + template + [[nodiscard]] _CCCL_API static constexpr auto __dims_as_impl(const _Args&... __args) noexcept + { + auto __exts = _Level::template extents_as<_Tp>(__args...); + using _Exts = decltype(__exts); + + hierarchy_query_result<_Tp> __ret{1, 1, 1}; + for (::cuda::std::size_t __i = 0; __i < _Exts::rank(); ++__i) + { + __ret[__i] = __exts.extent(__i); + } + return __ret; + } + + template + [[nodiscard]] _CCCL_API static constexpr auto __static_dims_impl(const _Args&... __args) noexcept + { + using _Exts = decltype(_Level::extents(__args...)); + + hierarchy_query_result<::cuda::std::size_t> __ret{1, 1, 1}; + for (::cuda::std::size_t __i = 0; __i < _Exts::rank(); ++__i) + { + __ret[__i] = _Exts::static_extent(__i); + } + return __ret; + } + + _CCCL_EXEC_CHECK_DISABLE + template + [[nodiscard]] _CCCL_API static constexpr _Tp __count_as_impl(const _Args&... __args) noexcept + { + const auto __exts = _Level::template extents_as<_Tp>(__args...); + + _Tp __ret = 1; + for (::cuda::std::size_t __i = 0; __i < __exts.rank(); ++__i) + { + __ret *= __exts.extent(__i); + } + return __ret; + } +}; + +_CCCL_END_NAMESPACE_CUDA + +#include + +#endif // _CUDA___HIERARCHY_HIERARCHY_LEVEL_BASE_H diff --git a/libcudacxx/include/cuda/__hierarchy/hierarchy_levels.h b/libcudacxx/include/cuda/__hierarchy/hierarchy_levels.h index 6019b4a1ba59..a629cf938a9a 100644 --- a/libcudacxx/include/cuda/__hierarchy/hierarchy_levels.h +++ b/libcudacxx/include/cuda/__hierarchy/hierarchy_levels.h @@ -11,6 +11,17 @@ #ifndef _CUDA___HIERARCHY_HIERARCHY_LEVELS_H #define _CUDA___HIERARCHY_HIERARCHY_LEVELS_H +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include #include #include @@ -98,94 +109,6 @@ template inline constexpr bool __legal_unit_for_level<_Unit, void> = false; } // namespace __detail -// Base type for all hierarchy levels -struct hierarchy_level -{}; - -struct grid_level; -struct cluster_level; -struct block_level; -struct thread_level; - -/* - Types to represent CUDA threads hierarchy levels - All metadata about the hierarchy level goes here including certain forward - progress information or what adjacent levels are valid in the hierarchy for - validation. -*/ - -/** - * @brief Type representing the grid level in CUDA thread hierarchy - * - * This type can be used in hierarchy queries to refer to the - * grid level or to get that level from the hierarchy. - * There is a constexpr variable of this type available for convenience - * named grid. - */ -struct grid_level - : public hierarchy_level - , public __detail::__dimensions_query -{ - using product_type = unsigned long long; - using allowed_above = allowed_levels<>; - using allowed_below = allowed_levels; -}; -_CCCL_GLOBAL_CONSTANT grid_level grid; - -/** - * @brief Type representing the cluster level in CUDA thread hierarchy - * - * This type can be used in hierarchy queries to refer to the - * cluster level or to get that level from the hierarchy. - * There is a constexpr variable of this type available for convenience - * named cluster. - */ -struct cluster_level - : public hierarchy_level - , public __detail::__dimensions_query -{ - using product_type = unsigned int; - using allowed_above = allowed_levels; - using allowed_below = allowed_levels; -}; -_CCCL_GLOBAL_CONSTANT cluster_level cluster; - -/** - * @brief Type representing the block level in CUDA thread hierarchy - * - * This type can be used in hierarchy queries to refer to the - * block level or to get that level from the hierarchy. - * There is a constexpr variable of this type available for convenience - * named block. - */ -struct block_level - : public hierarchy_level - , public __detail::__dimensions_query -{ - using product_type = unsigned int; - using allowed_above = allowed_levels; - using allowed_below = allowed_levels; -}; -_CCCL_GLOBAL_CONSTANT block_level block; - -/** - * @brief Type representing the thread level in CUDA thread hierarchy - * - * This type can be used in hierarchy queries to specify threads as a - * unit of the query. - * There is a constexpr variable of this type available for convenience - * named thread. - */ -struct thread_level - : public hierarchy_level - , public __detail::__dimensions_query -{ - using product_type = unsigned int; - using allowed_above = allowed_levels; - using allowed_below = allowed_levels<>; -}; -_CCCL_GLOBAL_CONSTANT thread_level thread; - template constexpr bool is_core_cuda_hierarchy_level = ::cuda::std::is_same_v<_Level, grid_level> || ::cuda::std::is_same_v<_Level, cluster_level> diff --git a/libcudacxx/include/cuda/__hierarchy/hierarchy_query_result.h b/libcudacxx/include/cuda/__hierarchy/hierarchy_query_result.h new file mode 100644 index 000000000000..0192198ba036 --- /dev/null +++ b/libcudacxx/include/cuda/__hierarchy/hierarchy_query_result.h @@ -0,0 +1,149 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA___HIERARCHY_HIERARCHY_QUERY_RESULT_H +#define _CUDA___HIERARCHY_HIERARCHY_QUERY_RESULT_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include + +#include + +_CCCL_BEGIN_NAMESPACE_CUDA + +template +struct hierarchy_query_result +{ + using value_type = _Tp; + + _Tp x; + _Tp y; + _Tp z; + + [[nodiscard]] _CCCL_API constexpr _Tp& operator[](::cuda::std::size_t __i) noexcept + { + if (__i == 0) + { + return x; + } + else if (__i == 1) + { + return y; + } + else + { + return z; + } + } + [[nodiscard]] _CCCL_API constexpr const _Tp& operator[](::cuda::std::size_t __i) const noexcept + { + if (__i == 0) + { + return x; + } + else if (__i == 1) + { + return y; + } + else + { + return z; + } + } + + _CCCL_TEMPLATE(class _Tp2 = _Tp) + _CCCL_REQUIRES(::cuda::std::is_same_v<_Tp2, signed char>) + _CCCL_API constexpr operator char3() const noexcept + { + return {static_cast(x), static_cast(y), static_cast(z)}; + } + + _CCCL_TEMPLATE(class _Tp2 = _Tp) + _CCCL_REQUIRES(::cuda::std::is_same_v<_Tp2, short>) + _CCCL_API constexpr operator short3() const noexcept + { + return {static_cast(x), static_cast(y), static_cast(z)}; + } + + _CCCL_TEMPLATE(class _Tp2 = _Tp) + _CCCL_REQUIRES(::cuda::std::is_same_v<_Tp2, int>) + _CCCL_API constexpr operator int3() const noexcept + { + return {static_cast(x), static_cast(y), static_cast(z)}; + } + + _CCCL_TEMPLATE(class _Tp2 = _Tp) + _CCCL_REQUIRES(::cuda::std::is_same_v<_Tp2, long>) + _CCCL_API constexpr operator long3() const noexcept + { + return {static_cast(x), static_cast(y), static_cast(z)}; + } + + _CCCL_TEMPLATE(class _Tp2 = _Tp) + _CCCL_REQUIRES(::cuda::std::is_same_v<_Tp2, long long>) + _CCCL_API constexpr operator longlong3() const noexcept + { + return {static_cast(x), static_cast(y), static_cast(z)}; + } + + _CCCL_TEMPLATE(class _Tp2 = _Tp) + _CCCL_REQUIRES(::cuda::std::is_same_v<_Tp2, unsigned char>) + _CCCL_API constexpr operator uchar3() const noexcept + { + return {static_cast(x), static_cast(y), static_cast(z)}; + } + + _CCCL_TEMPLATE(class _Tp2 = _Tp) + _CCCL_REQUIRES(::cuda::std::is_same_v<_Tp2, unsigned short>) + _CCCL_API constexpr operator ushort3() const noexcept + { + return {static_cast(x), static_cast(y), static_cast(z)}; + } + + _CCCL_TEMPLATE(class _Tp2 = _Tp) + _CCCL_REQUIRES(::cuda::std::is_same_v<_Tp2, unsigned>) + _CCCL_API constexpr operator uint3() const noexcept + { + return {static_cast(x), static_cast(y), static_cast(z)}; + } + + _CCCL_TEMPLATE(class _Tp2 = _Tp) + _CCCL_REQUIRES(::cuda::std::is_same_v<_Tp2, unsigned long>) + _CCCL_API constexpr operator ulong3() const noexcept + { + return {static_cast(x), static_cast(y), static_cast(z)}; + } + + _CCCL_TEMPLATE(class _Tp2 = _Tp) + _CCCL_REQUIRES(::cuda::std::is_same_v<_Tp2, unsigned long long>) + _CCCL_API constexpr operator ulonglong3() const noexcept + { + return {static_cast(x), static_cast(y), static_cast(z)}; + } +}; + +_CCCL_END_NAMESPACE_CUDA + +#include + +#endif // _CUDA___HIERARCHY_HIERARCHY_QUERY_RESULT_H diff --git a/libcudacxx/include/cuda/__hierarchy/level_dimensions.h b/libcudacxx/include/cuda/__hierarchy/level_dimensions.h index 8dea5ca85d40..ec4fcd38d9fd 100644 --- a/libcudacxx/include/cuda/__hierarchy/level_dimensions.h +++ b/libcudacxx/include/cuda/__hierarchy/level_dimensions.h @@ -11,6 +11,20 @@ #ifndef _CUDA___HIERARCHY_LEVEL_DIMENSIONS_H #define _CUDA___HIERARCHY_LEVEL_DIMENSIONS_H +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include #include #include #include @@ -111,7 +125,7 @@ struct __dimensions_handler<::cuda::std::integral_constant<_Dims, _Val>> template struct level_dimensions { - static_assert(::cuda::std::is_base_of_v); + static_assert(__is_hierarchy_level_v<_Level>); using level_type = _Level; // Needs alignas to work around an issue with tuple diff --git a/libcudacxx/include/cuda/__hierarchy/native_hierarchy_level_base.h b/libcudacxx/include/cuda/__hierarchy/native_hierarchy_level_base.h new file mode 100644 index 000000000000..2d97796a7a41 --- /dev/null +++ b/libcudacxx/include/cuda/__hierarchy/native_hierarchy_level_base.h @@ -0,0 +1,203 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA___HIERARCHY_NATIVE_HIERARCHY_LEVEL_BASE_H +#define _CUDA___HIERARCHY_NATIVE_HIERARCHY_LEVEL_BASE_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +_CCCL_BEGIN_NAMESPACE_CUDA + +// cudafe++ makes the queries (that are device only) return void when compiling for host, which causes host compilers +// to warn about applying [[nodiscard]] to a function that returns void. +_CCCL_DIAG_PUSH +#if _CCCL_CUDA_COMPILER(NVCC) +_CCCL_DIAG_SUPPRESS_GCC("-Wattributes") +_CCCL_DIAG_SUPPRESS_CLANG("-Wignored-attributes") +_CCCL_DIAG_SUPPRESS_NVHPC(nodiscard_doesnt_apply) +#endif // _CCCL_CUDA_COMPILER(NVCC) + +template +struct __native_hierarchy_level_base : hierarchy_level_base<_Level> +{ + using __base_type = hierarchy_level_base<_Level>; + using __base_type::count; + using __base_type::count_as; + using __base_type::dims; + using __base_type::dims_as; + using __base_type::extents; + using __base_type::extents_as; + using __base_type::index; + using __base_type::index_as; + using __base_type::rank; + using __base_type::rank_as; + using __base_type::static_dims; + + template + using __default_md_query_type = unsigned; + template + using __default_1d_query_type = ::cuda::std::size_t; + +#if _CCCL_CUDA_COMPILATION() + + _CCCL_TEMPLATE(class _InLevel) + _CCCL_REQUIRES(__is_native_hierarchy_level_v<_InLevel>) + [[nodiscard]] _CCCL_DEVICE_API static auto dims(const _InLevel& __level) noexcept + { + return _Level::template dims_as<__default_md_query_type<_InLevel>>(__level); + } + + _CCCL_TEMPLATE(class _InLevel) + _CCCL_REQUIRES(__is_native_hierarchy_level_v<_InLevel>) + [[nodiscard]] _CCCL_DEVICE_API static constexpr auto static_dims(const _InLevel& __level) noexcept + { + return __base_type::__static_dims_impl(__level); + } + + _CCCL_TEMPLATE(class _InLevel) + _CCCL_REQUIRES(__is_native_hierarchy_level_v<_InLevel>) + [[nodiscard]] _CCCL_DEVICE_API static auto extents(const _InLevel& __level) noexcept + { + return _Level::template extents_as<__default_md_query_type<_InLevel>>(__level); + } + + _CCCL_TEMPLATE(class _InLevel) + _CCCL_REQUIRES(__is_native_hierarchy_level_v<_InLevel>) + [[nodiscard]] _CCCL_DEVICE_API static auto count(const _InLevel& __level) noexcept + { + return _Level::template count_as<__default_1d_query_type<_InLevel>>(__level); + } + + _CCCL_TEMPLATE(class _InLevel) + _CCCL_REQUIRES(__is_native_hierarchy_level_v<_InLevel>) + [[nodiscard]] _CCCL_DEVICE_API static auto index(const _InLevel& __level) noexcept + { + return _Level::template index_as<__default_md_query_type<_InLevel>>(__level); + } + + _CCCL_TEMPLATE(class _InLevel) + _CCCL_REQUIRES(__is_native_hierarchy_level_v<_InLevel>) + [[nodiscard]] _CCCL_DEVICE_API static auto rank(const _InLevel& __level) noexcept + { + return _Level::template rank_as<__default_1d_query_type<_InLevel>>(__level); + } + + _CCCL_TEMPLATE(class _Tp, class _InLevel) + _CCCL_REQUIRES(__is_native_hierarchy_level_v<_InLevel>) + [[nodiscard]] _CCCL_DEVICE_API static auto dims_as(const _InLevel& __level) noexcept + { + return __base_type::template __dims_as_impl<_Tp>(__level); + } + + _CCCL_TEMPLATE(class _Tp, class _InLevel) + _CCCL_REQUIRES(__is_native_hierarchy_level_v<_InLevel>) + [[nodiscard]] _CCCL_DEVICE_API static auto extents_as(const _InLevel& __level) noexcept + { + static_assert(__is_natively_reachable_hierarchy_level_v<_Level, _InLevel>, + "_InLevel must be reachable from _Level"); + + using _NextLevel = typename _Level::__next_native_level; + auto __next_exts = _NextLevel::template extents_as<_Tp>(__level); + auto __curr_exts = _Level::template extents_as<_Tp>(_NextLevel{}); + return ::cuda::__hierarchy_extents_mul(__curr_exts, __next_exts); + } + + _CCCL_TEMPLATE(class _Tp, class _InLevel) + _CCCL_REQUIRES(__is_native_hierarchy_level_v<_InLevel>) + [[nodiscard]] _CCCL_DEVICE_API static auto count_as(const _InLevel& __level) noexcept + { + return __base_type::template __count_as_impl<_Tp>(__level); + } + + _CCCL_TEMPLATE(class _Tp, class _InLevel) + _CCCL_REQUIRES(__is_native_hierarchy_level_v<_InLevel>) + [[nodiscard]] _CCCL_DEVICE_API static auto index_as(const _InLevel& __level) noexcept + { + static_assert(__is_natively_reachable_hierarchy_level_v<_Level, _InLevel>, + "_InLevel must be reachable from _Level"); + + using _NextLevel = typename _Level::__next_native_level; + const auto __curr_exts = _Level::template extents_as<_Tp>(_NextLevel{}); + const auto __next_idx = _NextLevel::template index_as<_Tp>(__level); + const auto __curr_idx = _Level::template index_as<_Tp>(_NextLevel{}); + + hierarchy_query_result<_Tp> __ret{}; + for (::cuda::std::size_t __i = 0; __i < 3; ++__i) + { + __ret[__i] = __curr_idx[__i] + ((__i < __curr_exts.rank()) ? __curr_exts.extent(__i) : 1) * __next_idx[__i]; + } + return __ret; + } + + _CCCL_TEMPLATE(class _Tp, class _InLevel) + _CCCL_REQUIRES(__is_native_hierarchy_level_v<_InLevel>) + [[nodiscard]] _CCCL_DEVICE_API static _Tp rank_as(const _InLevel& __level) noexcept + { + static_assert(__is_natively_reachable_hierarchy_level_v<_Level, _InLevel>, + "_InLevel must be reachable from _Level"); + + using _NextLevel = typename _Level::__next_native_level; + + const auto __curr_exts = _Level::template extents_as<_Tp>(_NextLevel{}); + const auto __curr_idx = _Level::template index_as<_Tp>(_NextLevel{}); + + _Tp __ret = 0; + if constexpr (!::cuda::std::is_same_v<_InLevel, _NextLevel>) + { + __ret = _NextLevel::template rank_as<_Tp>(__level) * _Level::template count_as<_Tp>(_NextLevel{}); + } + + for (::cuda::std::size_t __i = __curr_exts.rank(); __i > 0; --__i) + { + _Tp __inc = __curr_idx[__i - 1]; + for (::cuda::std::size_t __j = __i - 1; __j > 0; --__j) + { + __inc *= __curr_exts.extent(__j - 1); + } + __ret += __inc; + } + return __ret; + } + +#endif // _CCCL_CUDA_COMPILATION() +}; + +_CCCL_DIAG_POP + +template <> +struct __native_hierarchy_level_base : hierarchy_level_base +{}; + +_CCCL_END_NAMESPACE_CUDA + +#include + +#endif // _CUDA___HIERARCHY_NATIVE_HIERARCHY_LEVEL_BASE_H diff --git a/libcudacxx/include/cuda/__hierarchy/thread_level.h b/libcudacxx/include/cuda/__hierarchy/thread_level.h new file mode 100644 index 000000000000..f8beda576d1f --- /dev/null +++ b/libcudacxx/include/cuda/__hierarchy/thread_level.h @@ -0,0 +1,96 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA___HIERARCHY_THREAD_LEVEL_H +#define _CUDA___HIERARCHY_THREAD_LEVEL_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include +#include + +#include + +_CCCL_BEGIN_NAMESPACE_CUDA + +struct thread_level : __native_hierarchy_level_base +{ + using product_type = unsigned; + using allowed_above = allowed_levels; + using allowed_below = allowed_levels<>; + + using __next_native_level = block_level; + + using __base_type = __native_hierarchy_level_base; + using __base_type::extents_as; + using __base_type::index_as; + using __base_type::rank_as; + + // interactions with block level + + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp>) + [[nodiscard]] _CCCL_DEVICE_API static ::cuda::std::dims<3, _Tp> extents_as(const block_level&) noexcept + { + return ::cuda::std::dims<3, _Tp>{ + static_cast<_Tp>(blockDim.x), static_cast<_Tp>(blockDim.y), static_cast<_Tp>(blockDim.z)}; + } + + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp>) + [[nodiscard]] _CCCL_DEVICE_API static hierarchy_query_result<_Tp> index_as(const block_level&) noexcept + { + return {static_cast<_Tp>(threadIdx.x), static_cast<_Tp>(threadIdx.y), static_cast<_Tp>(threadIdx.z)}; + } + + // interactions with warp level + + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp>) + [[nodiscard]] + _CCCL_DEVICE_API static constexpr ::cuda::std::extents<_Tp, 32> extents_as(const warp_level&) noexcept + { + return {}; + } + + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp>) + [[nodiscard]] _CCCL_DEVICE_API static hierarchy_query_result<_Tp> index_as(const warp_level&) noexcept + { + return {static_cast<_Tp>(::cuda::ptx::get_sreg_laneid()), 0, 0}; + } + + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp>) + [[nodiscard]] _CCCL_DEVICE_API static _Tp rank_as(const warp_level&) noexcept + { + return static_cast<_Tp>(::cuda::ptx::get_sreg_laneid()); + } +}; + +_CCCL_GLOBAL_CONSTANT thread_level gpu_thread; + +_CCCL_END_NAMESPACE_CUDA + +#include + +#endif // _CUDA___HIERARCHY_THREAD_LEVEL_H diff --git a/libcudacxx/include/cuda/__hierarchy/traits.h b/libcudacxx/include/cuda/__hierarchy/traits.h new file mode 100644 index 000000000000..1e64e0aaa772 --- /dev/null +++ b/libcudacxx/include/cuda/__hierarchy/traits.h @@ -0,0 +1,119 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA___HIERARCHY_TRAITS_H +#define _CUDA___HIERARCHY_TRAITS_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include +#include +#include + +#include + +_CCCL_BEGIN_NAMESPACE_CUDA + +// __is_natively_reachable_hierarchy_level_v + +template +inline constexpr bool __is_natively_reachable_hierarchy_level_helper_v = false; +template +inline constexpr bool __is_natively_reachable_hierarchy_level_helper_v< + _FromLevel, + _CurrLevel, + _ToLevel, + ::cuda::std::void_t> = + __is_natively_reachable_hierarchy_level_helper_v<_FromLevel, typename _CurrLevel::__next_native_level, _ToLevel>; +template +inline constexpr bool __is_natively_reachable_hierarchy_level_helper_v<_Level, _Level, _ToLevel> = false; +template +inline constexpr bool __is_natively_reachable_hierarchy_level_helper_v<_FromLevel, _Level, _Level> = true; + +template +inline constexpr bool __is_natively_reachable_hierarchy_level_v = false; +template +inline constexpr bool __is_natively_reachable_hierarchy_level_v< + _FromLevel, + _ToLevel, + ::cuda::std::void_t> = + __is_native_hierarchy_level_v<_ToLevel> + && __is_natively_reachable_hierarchy_level_helper_v<_FromLevel, typename _FromLevel::__next_native_level, _ToLevel>; + +// __level_type_of + +template +using __level_type_of = typename _Level::level_type; + +// has_unit_v + +template +inline constexpr bool __has_unit_helper_v = false; +template +inline constexpr bool __has_unit_helper_v<_QueryLevel, hierarchy_dimensions<_QueryLevel, _Levels...>> = true; + +// has_level_v + +template +inline constexpr bool __has_level_helper_v = false; +template +inline constexpr bool __has_level_helper_v<_QueryLevel, hierarchy_dimensions<_Unit, _Levels...>> = + (::cuda::std::is_same_v<_QueryLevel, typename _Levels::level_type> || ...); + +template +inline constexpr bool has_level_v = __has_level_helper_v<_QueryLevel, ::cuda::std::remove_cvref_t<_Hierarchy>>; + +template +inline constexpr bool has_unit_v = __has_unit_helper_v<_QueryLevel, ::cuda::std::remove_cvref_t<_Hierarchy>>; + +// has_unit_or_level_v + +template +inline constexpr bool has_unit_or_level_v = has_unit_v<_QueryLevel, _Hierarchy> || has_level_v<_QueryLevel, _Hierarchy>; + +// __next_hierarchy_level + +template +struct __next_hierarchy_level; + +template +struct __next_hierarchy_level<_Level, hierarchy_dimensions<_BottomUnit, _Levels...>> +{ + static constexpr ::cuda::std::size_t __level_idx = + ::cuda::std::__find_exactly_one_t<_Level, typename _Levels::level_type...>::value; + using __type = ::cuda::std::__type_index_c<__level_idx - 1, typename _Levels::level_type...>; +}; + +template +struct __next_hierarchy_level<_Level, hierarchy_dimensions<_Level, _Levels...>> +{ + using __type = ::cuda::std::__type_index_c<(sizeof...(_Levels) - 1), typename _Levels::level_type...>; +}; + +template +using __next_hierarchy_level_t = typename __next_hierarchy_level<_Level, _Hierarchy>::__type; + +_CCCL_END_NAMESPACE_CUDA + +#include + +#endif // _CUDA___HIERARCHY_TRAITS_H diff --git a/libcudacxx/include/cuda/__hierarchy/warp_level.h b/libcudacxx/include/cuda/__hierarchy/warp_level.h new file mode 100644 index 000000000000..6e1d91703916 --- /dev/null +++ b/libcudacxx/include/cuda/__hierarchy/warp_level.h @@ -0,0 +1,67 @@ +//===----------------------------------------------------------------------===// +// +// Part of libcu++, the C++ Standard Library for your entire system, +// under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef _CUDA___HIERARCHY_WARP_LEVEL_H +#define _CUDA___HIERARCHY_WARP_LEVEL_H + +#include + +#if defined(_CCCL_IMPLICIT_SYSTEM_HEADER_GCC) +# pragma GCC system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_CLANG) +# pragma clang system_header +#elif defined(_CCCL_IMPLICIT_SYSTEM_HEADER_MSVC) +# pragma system_header +#endif // no system header + +#include +#include +#include +#include +#include +#include +#include +#include +#include + +#include + +_CCCL_BEGIN_NAMESPACE_CUDA + +struct warp_level : __native_hierarchy_level_base +{ + using __next_native_level = block_level; + + using __base_type = __native_hierarchy_level_base; + using __base_type::extents_as; + using __base_type::index_as; + + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp>) + [[nodiscard]] _CCCL_DEVICE_API static ::cuda::std::dims<1, _Tp> extents_as(const block_level&) noexcept + { + return ::cuda::std::dims<1, _Tp>{static_cast<_Tp>((gpu_thread.count(block) + 31) / 32)}; + } + + _CCCL_TEMPLATE(class _Tp) + _CCCL_REQUIRES(::cuda::std::__cccl_is_integer_v<_Tp>) + [[nodiscard]] _CCCL_DEVICE_API static hierarchy_query_result<_Tp> index_as(const block_level&) noexcept + { + return {static_cast<_Tp>(gpu_thread.rank(block) / 32), 0, 0}; + } +}; + +_CCCL_GLOBAL_CONSTANT warp_level warp; + +_CCCL_END_NAMESPACE_CUDA + +#include + +#endif // _CUDA___HIERARCHY_WARP_LEVEL_H diff --git a/libcudacxx/include/cuda/__launch/launch.h b/libcudacxx/include/cuda/__launch/launch.h index 8ed5350b3010..edf3e3e1b8a1 100644 --- a/libcudacxx/include/cuda/__launch/launch.h +++ b/libcudacxx/include/cuda/__launch/launch.h @@ -24,6 +24,11 @@ #if _CCCL_HAS_CTK() && !_CCCL_COMPILER(NVRTC) # include +# include +# include +# include +# include +# include # include # include # include @@ -82,7 +87,7 @@ _CCCL_HOST_API auto __launch_impl(_Dst&& __dst, _Config __conf, ::CUfunction __k static_assert(!::cuda::std::is_same_v, "Can't launch a configuration without hierarchy dimensions"); ::CUlaunchConfig __config{}; - constexpr bool __has_cluster_level = has_level; + constexpr bool __has_cluster_level = has_level_v; constexpr unsigned int __num_attrs_needed = __detail::kernel_config_count_attr_space(__conf) + __has_cluster_level; ::CUlaunchAttribute __attrs[__num_attrs_needed == 0 ? 1 : __num_attrs_needed]; __config.attrs = &__attrs[0]; @@ -97,9 +102,9 @@ _CCCL_HOST_API auto __launch_impl(_Dst&& __dst, _Config __conf, ::CUfunction __k __config.gridDimX = static_cast(__conf.dims.extents(block, grid).x); __config.gridDimY = static_cast(__conf.dims.extents(block, grid).y); __config.gridDimZ = static_cast(__conf.dims.extents(block, grid).z); - __config.blockDimX = static_cast(__conf.dims.extents(thread, block).x); - __config.blockDimY = static_cast(__conf.dims.extents(thread, block).y); - __config.blockDimZ = static_cast(__conf.dims.extents(thread, block).z); + __config.blockDimX = static_cast(__conf.dims.extents(gpu_thread, block).x); + __config.blockDimY = static_cast(__conf.dims.extents(gpu_thread, block).y); + __config.blockDimZ = static_cast(__conf.dims.extents(gpu_thread, block).z); if constexpr (__has_cluster_level) { diff --git a/libcudacxx/include/cuda/hierarchy b/libcudacxx/include/cuda/hierarchy index acabf9928d4d..dcb3ff104dad 100644 --- a/libcudacxx/include/cuda/hierarchy +++ b/libcudacxx/include/cuda/hierarchy @@ -8,8 +8,8 @@ // //===----------------------------------------------------------------------===// -#ifndef _CUDA_HIERARCHY_ -#define _CUDA_HIERARCHY_ +#ifndef _CUDA_HIERARCHY +#define _CUDA_HIERARCHY #include @@ -21,8 +21,18 @@ # pragma system_header #endif // no system header +#include +#include +#include +#include #include +#include #include +#include #include +#include +#include +#include +#include -#endif // _CUDA_HIERARCHY_ +#endif // _CUDA_HIERARCHY diff --git a/libcudacxx/test/libcudacxx/cuda/ccclrt/common/host_device.cuh b/libcudacxx/test/libcudacxx/cuda/ccclrt/common/host_device.cuh index 07202970ab82..424cd7fe0248 100644 --- a/libcudacxx/test/libcudacxx/cuda/ccclrt/common/host_device.cuh +++ b/libcudacxx/test/libcudacxx/cuda/ccclrt/common/host_device.cuh @@ -69,10 +69,10 @@ void test_host_dev(const Dims& dims, const Lambda& lambda, const Filters&... fil cudaLaunchAttribute attrs[1]; config.attrs = &attrs[0]; - config.blockDim = dims.extents(cuda::thread, cuda::block); + config.blockDim = dims.extents(cuda::gpu_thread, cuda::block); config.gridDim = dims.extents(cuda::block, cuda::grid); - if constexpr (cuda::has_level) + if constexpr (cuda::has_level_v) { dim3 cluster_dims = dims.extents(cuda::block, cuda::cluster); config.attrs[config.numAttrs].id = cudaLaunchAttributeClusterDimension; diff --git a/libcudacxx/test/libcudacxx/cuda/ccclrt/hierarchy/hierarchy_custom_types.cu b/libcudacxx/test/libcudacxx/cuda/ccclrt/hierarchy/hierarchy_custom_types.cu index 3235ac210287..d592700eefb5 100644 --- a/libcudacxx/test/libcudacxx/cuda/ccclrt/hierarchy/hierarchy_custom_types.cu +++ b/libcudacxx/test/libcudacxx/cuda/ccclrt/hierarchy/hierarchy_custom_types.cu @@ -13,7 +13,7 @@ #include #include -struct custom_level : public cuda::hierarchy_level +struct custom_level : public cuda::hierarchy_level_base { using product_type = unsigned int; using allowed_above = cuda::allowed_levels; @@ -51,7 +51,7 @@ struct custom_level_test auto custom_block_back = custom_dims.level(cuda::block); CCCLRT_REQUIRE(custom_block_back.dummy == 2); - auto custom_dims_fragment = custom_dims.fragment(cuda::thread, cuda::block); + auto custom_dims_fragment = custom_dims.fragment(cuda::gpu_thread, cuda::block); auto custom_block_back2 = custom_dims_fragment.level(cuda::block); CCCLRT_REQUIRE(custom_block_back2.dummy == 2); @@ -62,8 +62,8 @@ struct custom_level_test cuda::level_dimensions(custom_level_dims), cuda::block_dims<256>()); - static_assert(custom_hierarchy.extents(cuda::thread, custom_level()) == dim3(512, 2, 2)); - static_assert(custom_hierarchy.count(cuda::thread, custom_level()) == 2048); + static_assert(custom_hierarchy.extents(cuda::gpu_thread, custom_level()) == dim3(512, 2, 2)); + static_assert(custom_hierarchy.count(cuda::gpu_thread, custom_level()) == 2048); test_host_dev(custom_hierarchy, *this); } diff --git a/libcudacxx/test/libcudacxx/cuda/ccclrt/hierarchy/hierarchy_smoke.cu b/libcudacxx/test/libcudacxx/cuda/ccclrt/hierarchy/hierarchy_smoke.cu index e63e80a93fc0..2aa81ce57328 100644 --- a/libcudacxx/test/libcudacxx/cuda/ccclrt/hierarchy/hierarchy_smoke.cu +++ b/libcudacxx/test/libcudacxx/cuda/ccclrt/hierarchy/hierarchy_smoke.cu @@ -30,15 +30,15 @@ struct basic_test_single_dim // device-side require doesn't work with clang-cuda for now #if !_CCCL_CUDA_COMPILER(CLANG) CCCLRT_REQUIRE(dims.extents().x == grid_size * block_size); - CCCLRT_REQUIRE(dims.extents(cuda::thread).x == grid_size * block_size); - CCCLRT_REQUIRE(dims.extents(cuda::thread, cuda::grid).x == grid_size * block_size); + CCCLRT_REQUIRE(dims.extents(cuda::gpu_thread).x == grid_size * block_size); + CCCLRT_REQUIRE(dims.extents(cuda::gpu_thread, cuda::grid).x == grid_size * block_size); CCCLRT_REQUIRE(dims.count() == grid_size * block_size); - CCCLRT_REQUIRE(dims.count(cuda::thread) == grid_size * block_size); - CCCLRT_REQUIRE(dims.count(cuda::thread, cuda::grid) == grid_size * block_size); + CCCLRT_REQUIRE(dims.count(cuda::gpu_thread) == grid_size * block_size); + CCCLRT_REQUIRE(dims.count(cuda::gpu_thread, cuda::grid) == grid_size * block_size); - CCCLRT_REQUIRE(dims.extents(cuda::thread, cuda::block).x == block_size); + CCCLRT_REQUIRE(dims.extents(cuda::gpu_thread, cuda::block).x == block_size); CCCLRT_REQUIRE(dims.extents(cuda::block, cuda::grid).x == grid_size); - CCCLRT_REQUIRE(dims.count(cuda::thread, cuda::block) == block_size); + CCCLRT_REQUIRE(dims.count(cuda::gpu_thread, cuda::block) == block_size); CCCLRT_REQUIRE(dims.count(cuda::block, cuda::grid) == grid_size); #endif } @@ -47,34 +47,34 @@ struct basic_test_single_dim { auto dims = cuda::make_hierarchy(cuda::block_dims(), cuda::grid_dims()); static_assert(dims.extents().x == grid_size * block_size); - static_assert(dims.extents(cuda::thread).x == grid_size * block_size); - static_assert(dims.extents(cuda::thread, cuda::grid).x == grid_size * block_size); + static_assert(dims.extents(cuda::gpu_thread).x == grid_size * block_size); + static_assert(dims.extents(cuda::gpu_thread, cuda::grid).x == grid_size * block_size); static_assert(dims.count() == grid_size * block_size); - static_assert(dims.count(cuda::thread) == grid_size * block_size); - static_assert(dims.count(cuda::thread, cuda::grid) == grid_size * block_size); + static_assert(dims.count(cuda::gpu_thread) == grid_size * block_size); + static_assert(dims.count(cuda::gpu_thread, cuda::grid) == grid_size * block_size); static_assert(dims.static_count() == grid_size * block_size); - static_assert(dims.static_count(cuda::thread) == grid_size * block_size); - static_assert(dims.static_count(cuda::thread, cuda::grid) == grid_size * block_size); + static_assert(dims.static_count(cuda::gpu_thread) == grid_size * block_size); + static_assert(dims.static_count(cuda::gpu_thread, cuda::grid) == grid_size * block_size); static_assert(dims.static_extents()[0] == grid_size * block_size); - static_assert(dims.static_extents(cuda::thread)[0] == grid_size * block_size); - static_assert(dims.static_extents(cuda::thread, cuda::grid)[0] == grid_size * block_size); + static_assert(dims.static_extents(cuda::gpu_thread)[0] == grid_size * block_size); + static_assert(dims.static_extents(cuda::gpu_thread, cuda::grid)[0] == grid_size * block_size); - static_assert(dims.extents(cuda::thread, cuda::block).x == block_size); + static_assert(dims.extents(cuda::gpu_thread, cuda::block).x == block_size); static_assert(dims.extents(cuda::block, cuda::grid).x == grid_size); - static_assert(dims.count(cuda::thread, cuda::block) == block_size); + static_assert(dims.count(cuda::gpu_thread, cuda::block) == block_size); static_assert(dims.count(cuda::block, cuda::grid) == grid_size); - static_assert(dims.static_count(cuda::thread, cuda::block) == block_size); + static_assert(dims.static_count(cuda::gpu_thread, cuda::block) == block_size); static_assert(dims.static_count(cuda::block, cuda::grid) == grid_size); - static_assert(dims.static_extents(cuda::thread, cuda::block)[0] == block_size); + static_assert(dims.static_extents(cuda::gpu_thread, cuda::block)[0] == block_size); auto dims_dyn = cuda::make_hierarchy(cuda::block_dims(block_size), cuda::grid_dims(grid_size)); test_host_dev(dims_dyn, *this); - static_assert(dims_dyn.static_count(cuda::thread, cuda::block) == cuda::std::dynamic_extent); - static_assert(dims_dyn.static_count(cuda::thread, cuda::grid) == cuda::std::dynamic_extent); - static_assert(dims_dyn.static_extents(cuda::thread, cuda::block)[0] == cuda::std::dynamic_extent); - static_assert(dims_dyn.static_extents(cuda::thread, cuda::grid)[0] == cuda::std::dynamic_extent); + static_assert(dims_dyn.static_count(cuda::gpu_thread, cuda::block) == cuda::std::dynamic_extent); + static_assert(dims_dyn.static_count(cuda::gpu_thread, cuda::grid) == cuda::std::dynamic_extent); + static_assert(dims_dyn.static_extents(cuda::gpu_thread, cuda::block)[0] == cuda::std::dynamic_extent); + static_assert(dims_dyn.static_extents(cuda::gpu_thread, cuda::grid)[0] == cuda::std::dynamic_extent); // Test that we can also drop the empty parens in the level constructors: auto config = cuda::make_hierarchy(cuda::block_dims, cuda::grid_dims); @@ -92,18 +92,18 @@ struct basic_test_multi_dim // device-side require doesn't work with clang-cuda for now #if !_CCCL_CUDA_COMPILER(CLANG) CCCLRT_REQUIRE(dims.extents() == dim3(32, 12, 4)); - CCCLRT_REQUIRE(dims.extents(cuda::thread) == dim3(32, 12, 4)); - CCCLRT_REQUIRE(dims.extents(cuda::thread, cuda::grid) == dim3(32, 12, 4)); + CCCLRT_REQUIRE(dims.extents(cuda::gpu_thread) == dim3(32, 12, 4)); + CCCLRT_REQUIRE(dims.extents(cuda::gpu_thread, cuda::grid) == dim3(32, 12, 4)); CCCLRT_REQUIRE(dims.extents().extent(0) == 32); CCCLRT_REQUIRE(dims.extents().extent(1) == 12); CCCLRT_REQUIRE(dims.extents().extent(2) == 4); CCCLRT_REQUIRE(dims.count() == 512 * 3); - CCCLRT_REQUIRE(dims.count(cuda::thread) == 512 * 3); - CCCLRT_REQUIRE(dims.count(cuda::thread, cuda::grid) == 512 * 3); + CCCLRT_REQUIRE(dims.count(cuda::gpu_thread) == 512 * 3); + CCCLRT_REQUIRE(dims.count(cuda::gpu_thread, cuda::grid) == 512 * 3); - CCCLRT_REQUIRE(dims.extents(cuda::thread, cuda::block) == dim3(2, 3, 4)); + CCCLRT_REQUIRE(dims.extents(cuda::gpu_thread, cuda::block) == dim3(2, 3, 4)); CCCLRT_REQUIRE(dims.extents(cuda::block, cuda::grid) == dim3(16, 4, 1)); - CCCLRT_REQUIRE(dims.count(cuda::thread, cuda::block) == 24); + CCCLRT_REQUIRE(dims.count(cuda::gpu_thread, cuda::block) == 24); CCCLRT_REQUIRE(dims.count(cuda::block, cuda::grid) == 64); #endif } @@ -113,30 +113,30 @@ struct basic_test_multi_dim auto dims_multidim = cuda::make_hierarchy(cuda::block_dims<2, 3, 4>(), cuda::grid_dims<16, 4, 1>()); static_assert(dims_multidim.extents() == dim3(32, 12, 4)); - static_assert(dims_multidim.extents(cuda::thread) == dim3(32, 12, 4)); - static_assert(dims_multidim.extents(cuda::thread, cuda::grid) == dim3(32, 12, 4)); + static_assert(dims_multidim.extents(cuda::gpu_thread) == dim3(32, 12, 4)); + static_assert(dims_multidim.extents(cuda::gpu_thread, cuda::grid) == dim3(32, 12, 4)); static_assert(dims_multidim.extents().extent(0) == 32); static_assert(dims_multidim.extents().extent(1) == 12); static_assert(dims_multidim.extents().extent(2) == 4); static_assert(dims_multidim.count() == 512 * 3); - static_assert(dims_multidim.count(cuda::thread) == 512 * 3); - static_assert(dims_multidim.count(cuda::thread, cuda::grid) == 512 * 3); + static_assert(dims_multidim.count(cuda::gpu_thread) == 512 * 3); + static_assert(dims_multidim.count(cuda::gpu_thread, cuda::grid) == 512 * 3); static_assert(dims_multidim.static_count() == 512 * 3); - static_assert(dims_multidim.static_count(cuda::thread) == 512 * 3); - static_assert(dims_multidim.static_count(cuda::thread, cuda::grid) == 512 * 3); + static_assert(dims_multidim.static_count(cuda::gpu_thread) == 512 * 3); + static_assert(dims_multidim.static_count(cuda::gpu_thread, cuda::grid) == 512 * 3); static_assert(dims_multidim.static_extents() == cuda::std::array{32, 12, 4}); - static_assert(dims_multidim.static_extents(cuda::thread) == cuda::std::array{32, 12, 4}); + static_assert(dims_multidim.static_extents(cuda::gpu_thread) == cuda::std::array{32, 12, 4}); static_assert( - dims_multidim.static_extents(cuda::thread, cuda::grid) == cuda::std::array{32, 12, 4}); + dims_multidim.static_extents(cuda::gpu_thread, cuda::grid) == cuda::std::array{32, 12, 4}); - static_assert(dims_multidim.extents(cuda::thread, cuda::block) == dim3(2, 3, 4)); + static_assert(dims_multidim.extents(cuda::gpu_thread, cuda::block) == dim3(2, 3, 4)); static_assert(dims_multidim.extents(cuda::block, cuda::grid) == dim3(16, 4, 1)); - static_assert(dims_multidim.count(cuda::thread, cuda::block) == 24); + static_assert(dims_multidim.count(cuda::gpu_thread, cuda::block) == 24); static_assert(dims_multidim.count(cuda::block, cuda::grid) == 64); - static_assert(dims_multidim.static_count(cuda::thread, cuda::block) == 24); + static_assert(dims_multidim.static_count(cuda::gpu_thread, cuda::block) == 24); static_assert(dims_multidim.static_count(cuda::block, cuda::grid) == 64); static_assert( - dims_multidim.static_extents(cuda::thread, cuda::block) == cuda::std::array{2, 3, 4}); + dims_multidim.static_extents(cuda::gpu_thread, cuda::block) == cuda::std::array{2, 3, 4}); static_assert( dims_multidim.static_extents(cuda::block, cuda::grid) == cuda::std::array{16, 4, 1}); @@ -144,8 +144,8 @@ struct basic_test_multi_dim test_host_dev(dims_multidim_dyn, *this); - static_assert(dims_multidim_dyn.static_count(cuda::thread, cuda::block) == cuda::std::dynamic_extent); - static_assert(dims_multidim_dyn.static_count(cuda::thread, cuda::grid) == cuda::std::dynamic_extent); + static_assert(dims_multidim_dyn.static_count(cuda::gpu_thread, cuda::block) == cuda::std::dynamic_extent); + static_assert(dims_multidim_dyn.static_count(cuda::gpu_thread, cuda::grid) == cuda::std::dynamic_extent); } }; @@ -159,14 +159,14 @@ struct basic_test_mixed // device-side require doesn't work with clang-cuda for now #if !_CCCL_CUDA_COMPILER(CLANG) CCCLRT_REQUIRE(dims.extents() == dim3(2048, 4, 2)); - CCCLRT_REQUIRE(dims.extents(cuda::thread) == dim3(2048, 4, 2)); - CCCLRT_REQUIRE(dims.extents(cuda::thread, cuda::grid) == dim3(2048, 4, 2)); + CCCLRT_REQUIRE(dims.extents(cuda::gpu_thread) == dim3(2048, 4, 2)); + CCCLRT_REQUIRE(dims.extents(cuda::gpu_thread, cuda::grid) == dim3(2048, 4, 2)); CCCLRT_REQUIRE(dims.extents().extent(0) == 2048); CCCLRT_REQUIRE(dims.extents().extent(1) == 4); CCCLRT_REQUIRE(dims.extents().extent(2) == 2); CCCLRT_REQUIRE(dims.count() == 16 * 1024); - CCCLRT_REQUIRE(dims.count(cuda::thread) == 16 * 1024); - CCCLRT_REQUIRE(dims.count(cuda::thread, cuda::grid) == 16 * 1024); + CCCLRT_REQUIRE(dims.count(cuda::gpu_thread) == 16 * 1024); + CCCLRT_REQUIRE(dims.count(cuda::gpu_thread, cuda::grid) == 16 * 1024); CCCLRT_REQUIRE(dims.extents(cuda::block, cuda::grid) == dim3(8, 4, 2)); CCCLRT_REQUIRE(dims.count(cuda::block, cuda::grid) == 64); @@ -178,11 +178,11 @@ struct basic_test_mixed auto dims_mixed = cuda::make_hierarchy(cuda::block_dims(), cuda::grid_dims(dim3(8, 4, 2))); test_host_dev(dims_mixed, *this); - static_assert(dims_mixed.extents(cuda::thread, cuda::block) == block_size); - static_assert(dims_mixed.count(cuda::thread, cuda::block) == block_size); - static_assert(dims_mixed.static_count(cuda::thread, cuda::block) == block_size); + static_assert(dims_mixed.extents(cuda::gpu_thread, cuda::block) == block_size); + static_assert(dims_mixed.count(cuda::gpu_thread, cuda::block) == block_size); + static_assert(dims_mixed.static_count(cuda::gpu_thread, cuda::block) == block_size); static_assert(dims_mixed.static_count(cuda::block, cuda::grid) == cuda::std::dynamic_extent); - static_assert(dims_mixed.static_extents(cuda::thread, cuda::block)[0] == block_size); + static_assert(dims_mixed.static_extents(cuda::gpu_thread, cuda::block)[0] == block_size); // TODO include mixed static and dynamic info on a single level // Currently bugged in std::extents @@ -209,7 +209,7 @@ struct basic_test_cluster CCCLRT_REQUIRE(dims.extents(cuda::block, cuda::grid) == dim3(2, 6, 9)); CCCLRT_REQUIRE(dims.count(cuda::block, cuda::grid) == 108); CCCLRT_REQUIRE(dims.extents(cuda::cluster, cuda::grid) == dim3(1, 3, 9)); - CCCLRT_REQUIRE(dims.extents(cuda::thread, cuda::cluster) == dim3(512, 2, 1)); + CCCLRT_REQUIRE(dims.extents(cuda::gpu_thread, cuda::cluster) == dim3(512, 2, 1)); #endif } @@ -224,13 +224,13 @@ struct basic_test_cluster static_assert(dims.static_count() == 1024 * 1024); static_assert(dims.static_extents()[0] == 1024 * 1024); - static_assert(dims.extents(cuda::thread, cuda::block).x == 256); + static_assert(dims.extents(cuda::gpu_thread, cuda::block).x == 256); static_assert(dims.extents(cuda::block, cuda::grid).x == 4 * 1024); - static_assert(dims.count(cuda::thread, cuda::cluster) == 2 * 1024); + static_assert(dims.count(cuda::gpu_thread, cuda::cluster) == 2 * 1024); static_assert(dims.count(cuda::cluster) == 512); static_assert(dims.static_count(cuda::cluster) == 512); static_assert(dims.static_count(cuda::block, cuda::cluster) == 8); - static_assert(dims.static_extents(cuda::thread, cuda::block)[0] == 256); + static_assert(dims.static_extents(cuda::gpu_thread, cuda::block)[0] == 256); static_assert(dims.static_extents(cuda::block, cuda::grid)[0] == 4 * 1024); } SECTION("Mixed cluster dims") @@ -238,12 +238,12 @@ struct basic_test_cluster auto dims_mixed = cuda::make_hierarchy( cuda::block_dims<256>(), cuda::cluster_dims(dim3(2, 2, 1)), cuda::grid_dims(dim3(1, 3, 9))); test_host_dev(dims_mixed, *this, arch_filter, 90>); - static_assert(dims_mixed.extents(cuda::thread, cuda::block) == 256); - static_assert(dims_mixed.count(cuda::thread, cuda::block) == 256); - static_assert(dims_mixed.static_count(cuda::thread, cuda::block) == 256); + static_assert(dims_mixed.extents(cuda::gpu_thread, cuda::block) == 256); + static_assert(dims_mixed.count(cuda::gpu_thread, cuda::block) == 256); + static_assert(dims_mixed.static_count(cuda::gpu_thread, cuda::block) == 256); static_assert(dims_mixed.static_count(cuda::block, cuda::cluster) == cuda::std::dynamic_extent); static_assert(dims_mixed.static_count(cuda::block) == cuda::std::dynamic_extent); - static_assert(dims_mixed.static_extents(cuda::thread, cuda::block)[0] == 256); + static_assert(dims_mixed.static_extents(cuda::gpu_thread, cuda::block)[0] == 256); static_assert(dims_mixed.static_extents(cuda::block, cuda::cluster)[0] == cuda::std::dynamic_extent); static_assert(dims_mixed.static_extents(cuda::block)[0] == cuda::std::dynamic_extent); } @@ -294,16 +294,16 @@ C2H_TEST("Different constructions", "[hierarchy]") cuda::block_dims()); static_assert(std::is_same_v); - static_assert(config.dims.count(cuda::thread, cuda::block) == block_size); - static_assert(config.dims.count(cuda::thread, cuda::cluster) == cluster_cnt * + static_assert(config.dims.count(cuda::gpu_thread, cuda::block) == block_size); + static_assert(config.dims.count(cuda::gpu_thread, cuda::cluster) == cluster_cnt * block_size); static_assert(config.dims.count(cuda::block, cuda::cluster) == cluster_cnt); CCCLRT_REQUIRE(config.dims.count() == grid_size * cluster_cnt * block_size); - static_assert(cuda::has_level); - static_assert(cuda::has_level); - static_assert(cuda::has_level); - static_assert(!cuda::has_level); + static_assert(cuda::has_level_v); + static_assert(cuda::has_level_v); + static_assert(cuda::has_level_v); + static_assert(!cuda::has_level_v); */ } @@ -311,16 +311,16 @@ C2H_TEST("Replace level", "[hierarchy]") { const auto dimensions = cuda::make_hierarchy(cuda::block_dims<512>(), cuda::cluster_dims<8>(), cuda::grid_dims(256)); const auto fragment = dimensions.fragment(cuda::block, cuda::grid); - static_assert(!cuda::has_level); - static_assert(!cuda::has_level_or_unit); - static_assert(cuda::has_level); - static_assert(cuda::has_level); - static_assert(cuda::has_level_or_unit); + static_assert(!cuda::has_level_v); + static_assert(!cuda::has_unit_or_level_v); + static_assert(cuda::has_level_v); + static_assert(cuda::has_level_v); + static_assert(cuda::has_unit_or_level_v); const auto replaced = cuda::hierarchy_add_level(fragment, cuda::block_dims(256)); - static_assert(cuda::has_level); - static_assert(cuda::has_level_or_unit); - CCCLRT_REQUIRE(replaced.count(cuda::thread, cuda::block) == 256); + static_assert(cuda::has_level_v); + static_assert(cuda::has_unit_or_level_v); + CCCLRT_REQUIRE(replaced.count(cuda::gpu_thread, cuda::block) == 256); } template @@ -329,34 +329,34 @@ __global__ void kernel(Hierarchy hierarchy) auto grid = cg::this_grid(); auto block = cg::this_thread_block(); - CCCLRT_REQUIRE_DEVICE(grid.thread_rank() == (cuda::hierarchy::rank(cuda::thread, cuda::grid))); + CCCLRT_REQUIRE_DEVICE(grid.thread_rank() == (cuda::hierarchy::rank(cuda::gpu_thread, cuda::grid))); CCCLRT_REQUIRE_DEVICE(grid.block_rank() == (cuda::hierarchy::rank(cuda::block, cuda::grid))); - CCCLRT_REQUIRE_DEVICE(grid.thread_rank() == cuda::grid.rank(cuda::thread)); - CCCLRT_REQUIRE_DEVICE(grid.block_rank() == cuda::grid.rank(cuda::block)); + CCCLRT_REQUIRE_DEVICE(grid.thread_rank() == cuda::gpu_thread.rank(cuda::grid)); + CCCLRT_REQUIRE_DEVICE(grid.block_rank() == cuda::block.rank(cuda::grid)); CCCLRT_REQUIRE_DEVICE(grid.block_index() == (cuda::hierarchy::index(cuda::block, cuda::grid))); - CCCLRT_REQUIRE_DEVICE(grid.block_index() == cuda::grid.index(cuda::block)); + CCCLRT_REQUIRE_DEVICE(grid.block_index() == cuda::block.index(cuda::grid)); - CCCLRT_REQUIRE_DEVICE(grid.num_threads() == (cuda::hierarchy::count(cuda::thread, cuda::grid))); + CCCLRT_REQUIRE_DEVICE(grid.num_threads() == (cuda::hierarchy::count(cuda::gpu_thread, cuda::grid))); CCCLRT_REQUIRE_DEVICE(grid.num_blocks() == (cuda::hierarchy::count(cuda::block, cuda::grid))); - CCCLRT_REQUIRE_DEVICE(grid.num_threads() == (cuda::grid.count(cuda::thread))); - CCCLRT_REQUIRE_DEVICE(grid.num_blocks() == cuda::grid.count(cuda::block)); + CCCLRT_REQUIRE_DEVICE(grid.num_threads() == (cuda::gpu_thread.count(cuda::grid))); + CCCLRT_REQUIRE_DEVICE(grid.num_blocks() == cuda::block.count(cuda::grid)); CCCLRT_REQUIRE_DEVICE(grid.dim_blocks() == (cuda::hierarchy::extents())); - CCCLRT_REQUIRE_DEVICE(grid.dim_blocks() == cuda::grid.extents(cuda::block)); + CCCLRT_REQUIRE_DEVICE(grid.dim_blocks() == dim3{cuda::block.dims(cuda::grid)}); CCCLRT_REQUIRE_DEVICE(block.thread_rank() == (cuda::hierarchy::rank())); CCCLRT_REQUIRE_DEVICE(block.thread_index() == (cuda::hierarchy::index())); CCCLRT_REQUIRE_DEVICE(block.num_threads() == (cuda::hierarchy::count())); CCCLRT_REQUIRE_DEVICE(block.dim_threads() == (cuda::hierarchy::extents())); - CCCLRT_REQUIRE_DEVICE(block.thread_rank() == cuda::block.rank(cuda::thread)); - CCCLRT_REQUIRE_DEVICE(block.thread_index() == cuda::block.index(cuda::thread)); - CCCLRT_REQUIRE_DEVICE(block.num_threads() == cuda::block.count(cuda::thread)); - CCCLRT_REQUIRE_DEVICE(block.dim_threads() == cuda::block.extents(cuda::thread)); + CCCLRT_REQUIRE_DEVICE(block.thread_rank() == cuda::gpu_thread.rank(cuda::block)); + CCCLRT_REQUIRE_DEVICE(block.thread_index() == cuda::gpu_thread.index(cuda::block)); + CCCLRT_REQUIRE_DEVICE(block.num_threads() == cuda::gpu_thread.count(cuda::block)); + CCCLRT_REQUIRE_DEVICE(block.dim_threads() == dim3{cuda::gpu_thread.dims(cuda::block)}); - auto block_index = hierarchy.index(cuda::thread, cuda::block); + auto block_index = hierarchy.index(cuda::gpu_thread, cuda::block); CCCLRT_REQUIRE_DEVICE(block_index == block.thread_index()); auto grid_index = hierarchy.index(); CCCLRT_REQUIRE_DEVICE( @@ -370,7 +370,7 @@ __global__ void kernel(Hierarchy hierarchy) == static_cast(grid.block_index().z) * block.dim_threads().z + block.thread_index().z); CCCLRT_REQUIRE_DEVICE(hierarchy.rank(cuda::block) == grid.block_rank()); - CCCLRT_REQUIRE_DEVICE(hierarchy.rank(cuda::thread, cuda::block) == block.thread_rank()); + CCCLRT_REQUIRE_DEVICE(hierarchy.rank(cuda::gpu_thread, cuda::block) == block.thread_rank()); CCCLRT_REQUIRE_DEVICE(hierarchy.rank() == grid.thread_rank()); } @@ -400,14 +400,14 @@ C2H_TEST("Dims queries indexing and ambient hierarchy", "[hierarchy]") template __global__ void rank_kernel_optimized(Hierarchy hierarchy, unsigned int* out) { - auto thread_id = hierarchy.rank(cuda::thread, cuda::block); + auto thread_id = hierarchy.rank(cuda::gpu_thread, cuda::block); out[thread_id] = thread_id; } template __global__ void rank_kernel(Hierarchy hierarchy, unsigned int* out) { - auto thread_id = cuda::hierarchy::rank(cuda::thread, cuda::block); + auto thread_id = cuda::hierarchy::rank(cuda::gpu_thread, cuda::block); out[thread_id] = thread_id; } @@ -438,36 +438,36 @@ template __global__ void examples_kernel(Hierarchy hierarchy) { { - auto thread_index_in_block = hierarchy.index(cuda::thread, cuda::block); + auto thread_index_in_block = hierarchy.index(cuda::gpu_thread, cuda::block); CCCLRT_REQUIRE_DEVICE(thread_index_in_block == threadIdx); auto block_index_in_grid = hierarchy.index(cuda::block); CCCLRT_REQUIRE_DEVICE(block_index_in_grid == blockIdx); } { - int thread_rank_in_block = hierarchy.rank(cuda::thread, cuda::block); + int thread_rank_in_block = hierarchy.rank(cuda::gpu_thread, cuda::block); int block_rank_in_grid = hierarchy.rank(cuda::block); } { // Can be called with the instances of level types - int num_threads_in_block = cuda::hierarchy::count(cuda::thread, cuda::block); - int num_blocks_in_grid = cuda::grid.count(cuda::block); + int num_threads_in_block = cuda::hierarchy::count(cuda::gpu_thread, cuda::block); + int num_blocks_in_grid = cuda::block.count(cuda::grid); // Or using the level types as template arguments int num_threads_in_grid = cuda::hierarchy::count(); } { // Can be called with the instances of level types - int thread_rank_in_block = cuda::hierarchy::rank(cuda::thread, cuda::block); - int block_rank_in_grid = cuda::grid.rank(cuda::block); + int thread_rank_in_block = cuda::hierarchy::rank(cuda::gpu_thread, cuda::block); + int block_rank_in_grid = cuda::block.rank(cuda::grid); // Or using the level types as template arguments int thread_rank_in_grid = cuda::hierarchy::rank(); } { // Can be called with the instances of level types - auto block_dims = cuda::hierarchy::extents(cuda::thread, cuda::block); + auto block_dims = cuda::hierarchy::extents(cuda::gpu_thread, cuda::block); CCCLRT_REQUIRE_DEVICE(block_dims == blockDim); - auto grid_dims = cuda::grid.extents(cuda::block); + dim3 grid_dims{cuda::block.dims(cuda::grid)}; CCCLRT_REQUIRE_DEVICE(grid_dims == gridDim); // Or using the level types as template arguments @@ -475,10 +475,10 @@ __global__ void examples_kernel(Hierarchy hierarchy) } { // Can be called with the instances of level types - auto thread_index_in_block = cuda::hierarchy::index(cuda::thread, cuda::block); - CCCLRT_REQUIRE_DEVICE(thread_index_in_block == threadIdx); - auto block_index_in_grid = cuda::grid.index(cuda::block); - CCCLRT_REQUIRE_DEVICE(block_index_in_grid == blockIdx); + auto thread_index_in_block = cuda::gpu_thread.index(cuda::block); + CCCLRT_REQUIRE_DEVICE(static_cast(thread_index_in_block) == threadIdx); + auto block_index_in_grid = cuda::block.index(cuda::grid); + CCCLRT_REQUIRE_DEVICE(static_cast(block_index_in_grid) == blockIdx); // Or using the level types as template arguments auto thread_index_in_grid = cuda::hierarchy::index(); @@ -494,25 +494,25 @@ C2H_TEST("Examples", "[hierarchy]") auto hierarchy = cuda::make_hierarchy(cuda::grid_dims(256), cuda::cluster_dims<4>(), cuda::block_dims<8, 8, 8>()); auto fragment = hierarchy.fragment(cuda::block, cuda::grid); auto new_hierarchy = cuda::hierarchy_add_level(fragment, cuda::block_dims<128>()); - static_assert(new_hierarchy.count(cuda::thread, cuda::block) == 128); + static_assert(new_hierarchy.count(cuda::gpu_thread, cuda::block) == 128); } { auto hierarchy = cuda::make_hierarchy(cuda::grid_dims(256), cuda::cluster_dims<4>(), cuda::block_dims<8, 8, 8>()); - static_assert(hierarchy.count(cuda::thread, cuda::cluster) == 4 * 8 * 8 * 8); + static_assert(hierarchy.count(cuda::gpu_thread, cuda::cluster) == 4 * 8 * 8 * 8); CCCLRT_REQUIRE(hierarchy.count() == 256 * 4 * 8 * 8 * 8); CCCLRT_REQUIRE(hierarchy.count(cuda::cluster) == 256); } { [[maybe_unused]] auto hierarchy = cuda::make_hierarchy(cuda::grid_dims(256), cuda::cluster_dims<4>(), cuda::block_dims<8, 8, 8>()); - static_assert(hierarchy.static_count(cuda::thread, cuda::cluster) == 4 * 8 * 8 * 8); + static_assert(hierarchy.static_count(cuda::gpu_thread, cuda::cluster) == 4 * 8 * 8 * 8); CCCLRT_REQUIRE(hierarchy.static_count() == cuda::std::dynamic_extent); } { auto hierarchy = cuda::make_hierarchy(cuda::grid_dims(256), cuda::cluster_dims<4>(), cuda::block_dims<8, 8, 8>()); - static_assert(hierarchy.extents(cuda::thread, cuda::cluster).extent(0) == 4 * 8); - static_assert(hierarchy.extents(cuda::thread, cuda::cluster).extent(1) == 8); - static_assert(hierarchy.extents(cuda::thread, cuda::cluster).extent(2) == 8); + static_assert(hierarchy.extents(cuda::gpu_thread, cuda::cluster).extent(0) == 4 * 8); + static_assert(hierarchy.extents(cuda::gpu_thread, cuda::cluster).extent(1) == 8); + static_assert(hierarchy.extents(cuda::gpu_thread, cuda::cluster).extent(2) == 8); CCCLRT_REQUIRE(hierarchy.extents().extent(0) == 256 * 4 * 8); CCCLRT_REQUIRE(hierarchy.extents(cuda::cluster).extent(0) == 256); } @@ -568,11 +568,11 @@ C2H_TEST("Trivially constructable", "[hierarchy]") C2H_TEST("cuda::distribute", "[hierarchy]") { - int numElements = 50000; + unsigned numElements = 50000; constexpr int threadsPerBlock = 256; auto config = cuda::distribute(numElements); - CCCLRT_REQUIRE(config.dims.count(cuda::thread, cuda::block) == 256); + CCCLRT_REQUIRE(config.dims.count(cuda::gpu_thread, cuda::block) == 256); CCCLRT_REQUIRE(config.dims.count(cuda::block, cuda::grid) == (numElements + threadsPerBlock - 1) / threadsPerBlock); } @@ -583,45 +583,45 @@ C2H_TEST("hierarchy merge", "[hierarchy]") auto h1 = cuda::make_hierarchy(cuda::grid_dims<2>()); auto h2 = cuda::make_hierarchy(cuda::block_dims<3>()); auto combined = h1.combine(h2); - static_assert(combined.count(cuda::thread) == 6); - static_assert(combined.count(cuda::thread, cuda::block) == 3); + static_assert(combined.count(cuda::gpu_thread) == 6); + static_assert(combined.count(cuda::gpu_thread, cuda::block) == 3); static_assert(combined.count(cuda::block) == 2); auto combined_the_other_way = h2.combine(h1); static_assert(cuda::std::is_same_v); - static_assert(combined_the_other_way.count(cuda::thread) == 6); + static_assert(combined_the_other_way.count(cuda::gpu_thread) == 6); auto dynamic_values = cuda::make_hierarchy(cuda::cluster_dims(4), cuda::block_dims(5)); auto combined_dynamic = dynamic_values.combine(h1); - CCCLRT_REQUIRE(combined_dynamic.count(cuda::thread) == 40); + CCCLRT_REQUIRE(combined_dynamic.count(cuda::gpu_thread) == 40); } SECTION("Overlapping") { auto h1 = cuda::make_hierarchy(cuda::grid_dims<2>(), cuda::cluster_dims<3>()); auto h2 = cuda::make_hierarchy(cuda::block_dims<4>(), cuda::cluster_dims<5>()); auto combined = h1.combine(h2); - static_assert(combined.count(cuda::thread) == 24); - static_assert(combined.count(cuda::thread, cuda::block) == 4); + static_assert(combined.count(cuda::gpu_thread) == 24); + static_assert(combined.count(cuda::gpu_thread, cuda::block) == 4); static_assert(combined.count(cuda::block) == 6); auto combined_the_other_way = h2.combine(h1); static_assert(!cuda::std::is_same_v); - static_assert(combined_the_other_way.count(cuda::thread) == 40); - static_assert(combined_the_other_way.count(cuda::thread, cuda::block) == 4); + static_assert(combined_the_other_way.count(cuda::gpu_thread) == 40); + static_assert(combined_the_other_way.count(cuda::gpu_thread, cuda::block) == 4); static_assert(combined_the_other_way.count(cuda::block) == 10); auto ultimate_combination = combined.combine(combined_the_other_way); static_assert(cuda::std::is_same_v); - static_assert(ultimate_combination.count(cuda::thread) == 24); + static_assert(ultimate_combination.count(cuda::gpu_thread) == 24); auto block_level_replacement = cuda::make_hierarchy(cuda::block_dims<6>()); auto with_block_replaced = block_level_replacement.combine(combined); - static_assert(with_block_replaced.count(cuda::thread) == 36); - static_assert(with_block_replaced.count(cuda::thread, cuda::block) == 6); + static_assert(with_block_replaced.count(cuda::gpu_thread) == 36); + static_assert(with_block_replaced.count(cuda::gpu_thread, cuda::block) == 6); auto grid_cluster_level_replacement = cuda::make_hierarchy(cuda::grid_dims<7>(), cuda::cluster_dims<8>()); auto with_grid_cluster_replaced = grid_cluster_level_replacement.combine(combined); - static_assert(with_grid_cluster_replaced.count(cuda::thread) == 7 * 8 * 4); + static_assert(with_grid_cluster_replaced.count(cuda::gpu_thread) == 7 * 8 * 4); static_assert(with_grid_cluster_replaced.count(cuda::block, cuda::cluster) == 8); static_assert(with_grid_cluster_replaced.count(cuda::cluster) == 7); } diff --git a/libcudacxx/test/libcudacxx/cuda/ccclrt/launch/configuration.cu b/libcudacxx/test/libcudacxx/cuda/ccclrt/launch/configuration.cu index cf7e15183cda..2508df8b7c22 100644 --- a/libcudacxx/test/libcudacxx/cuda/ccclrt/launch/configuration.cu +++ b/libcudacxx/test/libcudacxx/cuda/ccclrt/launch/configuration.cu @@ -204,10 +204,10 @@ C2H_TEST("Hierarchy construction in config", "[launch]") static_assert(config.dims.count(cuda::block) == 2); auto config_larger = cuda::make_config(cuda::grid_dims<2>(), cuda::block_dims(256), cuda::cooperative_launch()); - CCCLRT_REQUIRE(config_larger.dims.count(cuda::thread) == 512); + CCCLRT_REQUIRE(config_larger.dims.count(cuda::gpu_thread) == 512); auto config_no_options = cuda::make_config(cuda::grid_dims(2), cuda::block_dims<128>()); - CCCLRT_REQUIRE(config_no_options.dims.count(cuda::thread) == 256); + CCCLRT_REQUIRE(config_no_options.dims.count(cuda::gpu_thread) == 256); [[maybe_unused]] auto config_no_dims = cuda::make_config(cuda::cooperative_launch()); static_assert(cuda::std::is_same_v); @@ -231,18 +231,18 @@ C2H_TEST("Configuration combine", "[launch]") static_assert(cuda::std::is_same_v); static_assert(cuda::std::is_same_v); static_assert(cuda::std::is_same_v); - CCCLRT_REQUIRE(combined.dims.count(cuda::thread) == 512); + CCCLRT_REQUIRE(combined.dims.count(cuda::gpu_thread) == 512); } SECTION("Combine with overlap") { auto config_part1 = make_config(grid, cluster, cuda::launch_priority(2)); auto config_part2 = make_config(cuda::cluster_dims<256>(), block, cuda::launch_priority(42)); auto combined = config_part1.combine(config_part2); - CCCLRT_REQUIRE(combined.dims.count(cuda::thread) == 2048); + CCCLRT_REQUIRE(combined.dims.count(cuda::gpu_thread) == 2048); CCCLRT_REQUIRE(cuda::std::get<0>(combined.options).priority == 2); auto replaced_one_option = cuda::make_config(cuda::launch_priority(3)).combine(combined); - CCCLRT_REQUIRE(replaced_one_option.dims.count(cuda::thread) == 2048); + CCCLRT_REQUIRE(replaced_one_option.dims.count(cuda::gpu_thread) == 2048); CCCLRT_REQUIRE(cuda::std::get<0>(replaced_one_option.options).priority == 3); [[maybe_unused]] auto combined_with_extra_option = combined.combine(cuda::make_config(cuda::cooperative_launch())); diff --git a/libcudacxx/test/libcudacxx/cuda/ccclrt/launch/launch_smoke.cu b/libcudacxx/test/libcudacxx/cuda/ccclrt/launch/launch_smoke.cu index 584aa4c2d1e6..232df2f0a96a 100644 --- a/libcudacxx/test/libcudacxx/cuda/ccclrt/launch/launch_smoke.cu +++ b/libcudacxx/test/libcudacxx/cuda/ccclrt/launch/launch_smoke.cu @@ -49,7 +49,7 @@ struct functor_taking_config template __device__ void operator()(Config config, int grid_size) { - static_assert(config.dims.static_count(cuda::thread, cuda::block) == BlockSize); + static_assert(config.dims.static_count(cuda::gpu_thread, cuda::block) == BlockSize); CCCLRT_REQUIRE_DEVICE(config.dims.count(cuda::block, cuda::grid) == grid_size); kernel_run_proof = true; } @@ -220,7 +220,7 @@ void launch_smoke_test(cudaStream_t dst) { cuda::launch(dst, cuda::block_dims<256>() & cuda::grid_dims(1), [] __device__(auto config) { - if (config.dims.rank(cuda::thread, cuda::block) == 0) { + if (config.dims.rank(cuda::gpu_thread, cuda::block) == 0) { printf("Hello from the GPU\n"); kernel_run_proof = true; } @@ -307,7 +307,7 @@ void test_default_config() { auto block = cuda::block_dims<256>; auto verify_lambda = [] __device__(auto config) { - static_assert(config.dims.count(cuda::thread, cuda::block) == 256); + static_assert(config.dims.count(cuda::gpu_thread, cuda::block) == 256); CCCLRT_REQUIRE(config.dims.count(cuda::block) == 4); cooperative_groups::this_grid().sync(); }; diff --git a/libcudacxx/test/libcudacxx/cuda/containers/buffer/transform.cu b/libcudacxx/test/libcudacxx/cuda/containers/buffer/transform.cu index 6fcdf4ea4d0e..d00c5f49c316 100644 --- a/libcudacxx/test/libcudacxx/cuda/containers/buffer/transform.cu +++ b/libcudacxx/test/libcudacxx/cuda/containers/buffer/transform.cu @@ -15,6 +15,7 @@ #include #include +#include #include #include #include @@ -68,8 +69,8 @@ struct add_kernel template __device__ void operator()(cuda::std::span a, cuda::std::span b) { - for (int i = cuda::hierarchy::rank(cuda::thread, cuda::grid); i < a.size(); - i += cuda::hierarchy::count(cuda::thread, cuda::grid)) + for (int i = cuda::hierarchy::rank(cuda::gpu_thread, cuda::grid); i < a.size(); + i += cuda::hierarchy::count(cuda::gpu_thread, cuda::grid)) { a[i] += b[i]; } diff --git a/libcudacxx/test/libcudacxx/cuda/hierarchy/block_level/hierarchy_queries.pass.cpp b/libcudacxx/test/libcudacxx/cuda/hierarchy/block_level/hierarchy_queries.pass.cpp new file mode 100644 index 000000000000..b0712e18cd80 --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/hierarchy/block_level/hierarchy_queries.pass.cpp @@ -0,0 +1,240 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// todo: enable with nvrtc +// UNSUPPORTED: nvrtc + +#include "hierarchy_queries.h" + +#include +#include +#include +#include +#include + +template +__device__ void test_block( + const Hierarchy& hier, const GridExts& grid_exts, const ClusterExts& cluster_exts, const BlockExts& block_exts) +{ + // 1. Test cuda::block.dims(x, hier) + if constexpr (cuda::has_level_v) + { + uint3 exp{1, 1, 1}; + NV_IF_TARGET(NV_PROVIDES_SM_90, (exp = __clusterDim();)) + test_dims(exp, cuda::block, cuda::cluster, hier); + } + test_dims(gridDim, cuda::block, cuda::grid, hier); + + // 2. Test cuda::block.static_dims(x, hier) + if constexpr (cuda::has_level_v) + { + const ulonglong3 exp{ + ClusterExts::static_extent(0), + ClusterExts::static_extent(1), + ClusterExts::static_extent(2), + }; + test_static_dims(exp, cuda::block, cuda::cluster, hier); + } + { + const ulonglong3 exp{ + mul_static_extents(GridExts::static_extent(0), ClusterExts::static_extent(0)), + mul_static_extents(GridExts::static_extent(1), ClusterExts::static_extent(1)), + mul_static_extents(GridExts::static_extent(2), ClusterExts::static_extent(2)), + }; + test_static_dims(exp, cuda::block, cuda::grid, hier); + } + + // 3. Test cuda::block.extents(x) + if constexpr (cuda::has_level_v) + { + uint3 dims{1, 1, 1}; + NV_IF_TARGET(NV_PROVIDES_SM_90, (dims = __clusterDim();)) + const cuda::std:: + extents + exp{dims.x, dims.y, dims.z}; + + test_extents(exp, cuda::block, cuda::cluster, hier); + } + { + const cuda::std::extents + exp{gridDim.x, gridDim.y, gridDim.z}; + test_extents(exp, cuda::block, cuda::grid, hier); + } + + // 4. Test cuda::block.count(x, hier) + if constexpr (cuda::has_level_v) + { + cuda::std::size_t exp = 1; + NV_IF_TARGET(NV_PROVIDES_SM_90, ({ + exp *= __clusterDim().x; + exp *= __clusterDim().y; + exp *= __clusterDim().z; + })) + test_count(exp, cuda::block, cuda::cluster, hier); + } + test_count(cuda::std::size_t{gridDim.z} * gridDim.y * gridDim.x, cuda::block, cuda::grid, hier); + + // 5. test cuda::block.index(x, hier) + if constexpr (cuda::has_level_v) + { + uint3 exp{0, 0, 0}; + NV_IF_TARGET(NV_PROVIDES_SM_90, (exp = __clusterRelativeBlockIdx();)) + test_index(exp, cuda::block, cuda::cluster, hier); + } + test_index(blockIdx, cuda::block, cuda::grid, hier); + + // 6. Test cuda::block.rank(x, hier) + if constexpr (cuda::has_level_v) + { + cuda::std::size_t exp = 0; + NV_IF_TARGET(NV_PROVIDES_SM_90, ({ + exp = ((__clusterRelativeBlockIdx().z * __clusterDim().y) + __clusterRelativeBlockIdx().y) + * __clusterDim().x + + __clusterRelativeBlockIdx().x; + })) + test_rank(exp, cuda::block, cuda::cluster, hier); + } + { + const cuda::std::size_t exp = (blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x; + test_rank(exp, cuda::block, cuda::grid, hier); + } +} + +__device__ void test_device() +{ + // todo: make hierarchy constructible on device + // test_thread(cuda::make_hierarchy(cuda::grid_dims(gridDim), cuda::block_dims(blockDim))); +} + +#if !_CCCL_COMPILER(NVRTC) +template +__global__ void test_kernel(Hierarchy hier, GridExts grid_exts, BlockExts block_exts) +{ + test_block(hier, grid_exts, cuda::std::extents{}, block_exts); +} + +template +__global__ void test_kernel(Hierarchy hier, GridExts grid_exts, ClusterExts cluster_exts, BlockExts block_exts) +{ + test_block(hier, grid_exts, cluster_exts, block_exts); +} + +template +void test_launch(GridExts grid_exts, BlockExts block_exts) +{ + const dim3 grid_dims{grid_exts.extent(0), grid_exts.extent(1), grid_exts.extent(2)}; + const dim3 block_dims{block_exts.extent(0), block_exts.extent(1), block_exts.extent(2)}; + + const cuda::std::dims<3, unsigned> grid_exts_dyn{grid_exts.extent(0), grid_exts.extent(1), grid_exts.extent(2)}; + const cuda::std::dims<3, unsigned> block_exts_dyn{block_exts.extent(0), block_exts.extent(1), block_exts.extent(2)}; + + // 1. Launch hierarchy with all static extents. + test_kernel<<>>( + cuda::make_hierarchy( + cuda::grid_dims(), + cuda::block_dims()), + grid_exts, + block_exts); + + // 2. Launch hierarchy with static grid extents and dynamic block extents. + test_kernel<<>>( + cuda::make_hierarchy( + cuda::grid_dims(), + cuda::block_dims(block_dims)), + grid_exts, + block_exts_dyn); + + // 3. Launch hierarchy with dynamic grid extents and static block extents. + test_kernel<<>>( + cuda::make_hierarchy( + cuda::grid_dims(grid_dims), + cuda::block_dims()), + grid_exts_dyn, + block_exts); + + // 4. Launch hierarchy with dynamic grid extents and dynamic block extents. + test_kernel<<>>( + cuda::make_hierarchy(cuda::grid_dims(grid_dims), cuda::block_dims(block_dims)), grid_exts_dyn, block_exts_dyn); +} + +template +void test_launch(GridExts grid_exts, ClusterExts cluster_exts, BlockExts block_exts) +{ + const dim3 grid_dims{grid_exts.extent(0), grid_exts.extent(1), grid_exts.extent(2)}; + const dim3 cluster_dims{cluster_exts.extent(0), cluster_exts.extent(1), cluster_exts.extent(2)}; + const dim3 block_dims{block_exts.extent(0), block_exts.extent(1), block_exts.extent(2)}; + + cuda::std::dims<3, unsigned> grid_exts_dyn{grid_exts.extent(0), grid_exts.extent(1), grid_exts.extent(2)}; + cuda::std::dims<3, unsigned> cluster_exts_dyn{cluster_exts.extent(0), cluster_exts.extent(1), cluster_exts.extent(2)}; + cuda::std::dims<3, unsigned> block_exts_dyn{block_exts.extent(0), block_exts.extent(1), block_exts.extent(2)}; + + cudaLaunchAttribute attribute[1]{}; + attribute[0].id = cudaLaunchAttributeClusterDimension; + attribute[0].val.clusterDim.x = cluster_dims.x; + attribute[0].val.clusterDim.y = cluster_dims.y; + attribute[0].val.clusterDim.z = cluster_dims.z; + + cudaLaunchConfig_t config{}; + config.gridDim = dim3{grid_dims.x * cluster_dims.x, grid_dims.y * cluster_dims.y, grid_dims.z * cluster_dims.z}; + config.blockDim = block_dims; + config.attrs = attribute; + config.numAttrs = 1; + + // 1. Launch hierarchy with all static extents. + { + auto hier = cuda::make_hierarchy( + cuda::grid_dims(), + cuda::cluster_dims(), + cuda::block_dims()); + auto kernel = test_kernel; + assert(cudaLaunchKernelEx(&config, kernel, hier, grid_exts, cluster_exts, block_exts) == cudaSuccess); + } + + // 2. Launch hierarchy with all dynamic extents. + { + auto hier = + cuda::make_hierarchy(cuda::grid_dims(grid_dims), cuda::cluster_dims(cluster_dims), cuda::block_dims(block_dims)); + auto kernel = + test_kernel; + assert(cudaLaunchKernelEx(&config, kernel, hier, grid_exts_dyn, cluster_exts_dyn, block_exts_dyn) == cudaSuccess); + } +} + +void test() +{ + int cc_major{}; + assert(cudaDeviceGetAttribute(&cc_major, cudaDevAttrComputeCapabilityMajor, 0) == cudaSuccess); + + // thread block clusters require compute capability at least 9.0 + const bool enable_clusters = cc_major >= 9; + + test_launch(cuda::std::extents{}, cuda::std::extents{}); + test_launch(cuda::std::extents{}, cuda::std::extents{}); + test_launch(cuda::std::extents{}, cuda::std::extents{}); + test_launch(cuda::std::extents{}, cuda::std::extents{}); + + if (enable_clusters) + { + test_launch(cuda::std::extents{}, + cuda::std::extents{}, + cuda::std::extents{}); + } + + assert(cudaDeviceSynchronize() == cudaSuccess); +} +#endif // !_CCCL_COMPILER(NVRTC) + +int main(int, char**) +{ + NV_IF_ELSE_TARGET(NV_IS_HOST, (test();), (test_device();)) + return 0; +} diff --git a/libcudacxx/test/libcudacxx/cuda/hierarchy/block_level/hierarchy_query_signatures.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/hierarchy/block_level/hierarchy_query_signatures.compile.pass.cpp new file mode 100644 index 000000000000..55f3c8b803e9 --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/hierarchy/block_level/hierarchy_query_signatures.compile.pass.cpp @@ -0,0 +1,128 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// todo: enable with nvrtc +// UNSUPPORTED: nvrtc + +#include +#include +#include +#include + +template +__device__ void test_query_signatures(const Level& level, const Hierarchy& hier) +{ + // 1. Test cuda::block_level::dims(x, hier) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::block_level::dims(level, hier))>); + static_assert(noexcept(cuda::block_level::dims(level, hier))); + + // 2. Test cuda::block_level::static_dims(x, hier) signature. + static_assert(cuda::std::is_same_v, + decltype(cuda::block_level::static_dims(level, hier))>); + static_assert(noexcept(cuda::block_level::static_dims(level, hier))); + + // 3. Test cuda::block_level::extents(x, hier) signature. + using ExtentsResult = decltype(cuda::block_level::extents(level, hier)); + static_assert(cuda::std::__is_cuda_std_extents_v); + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::block_level::extents(level, hier))); + + // 4. Test cuda::block_level::count(x, hier) signature. + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::block_level::count(level, hier))); + + // 5. Test cuda::block_level::index(x, hier) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::block_level::index(level, hier))>); + static_assert(noexcept(cuda::block_level::index(level, hier))); + + // 6. Test cuda::block_level::rank(x, hier) signature. + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::block_level::rank(level, hier))); +} + +template +__device__ void test_query_as_signatures(const Level& level, const Hierarchy& hier) +{ + // 1. Test cuda::block_level::dims_as(x, hier) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::block_level::dims_as(level, hier))>); + static_assert(noexcept(cuda::block_level::dims_as(level, hier))); + + // 2. Test cuda::block_level::extents_as(x, hier) signature. + using ExtentsResult = decltype(cuda::block_level::extents_as(level, hier)); + static_assert(cuda::std::__is_cuda_std_extents_v); + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::block_level::extents_as(level, hier))); + + // 3. Test cuda::block_level::count_as(x, hier) signature. + static_assert(cuda::std::is_same_v(level, hier))>); + static_assert(noexcept(cuda::block_level::count_as(level, hier))); + + // 4. Test cuda::block_level::index_as(x, hier) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::block_level::index_as(level, hier))>); + static_assert(noexcept(cuda::block_level::index_as(level, hier))); + + // 5. Test cuda::block_level::rank_as(x, hier) signature. + static_assert(cuda::std::is_same_v(level, hier))>); + static_assert(noexcept(cuda::block_level::rank_as(level, hier))); +} + +template +__device__ void test(const InLevel& in_level, const Hierarchy& hier) +{ + test_query_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); +} + +template +__device__ void test(const Hierarchy& hier) +{ + if constexpr (cuda::has_level_v) + { + test(cuda::cluster, hier); + } + test(cuda::grid, hier); +} + +template +__global__ void test_kernel(Hierarchy hier) +{ + test(hier); +} + +#define TEST_KERNEL_INSTANTIATE(...) \ + template __global__ void test_kernel( \ + decltype(cuda::make_hierarchy(__VA_ARGS__))) + +TEST_KERNEL_INSTANTIATE(cuda::grid_dims<1>(), cuda::block_dims<1>()); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims<1>(), cuda::block_dims(dim3{})); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims(dim3{}), cuda::block_dims<1>()); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims(dim3{}), cuda::block_dims(dim3{})); + +TEST_KERNEL_INSTANTIATE(cuda::grid_dims<1>(), cuda::cluster_dims<1>(), cuda::block_dims<1>()); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims<1>(), cuda::cluster_dims<1>(), cuda::block_dims(dim3{})); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims<1>(), cuda::cluster_dims(dim3{}), cuda::block_dims<1>()); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims<1>(), cuda::cluster_dims(dim3{}), cuda::block_dims(dim3{})); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims(dim3{}), cuda::cluster_dims<1>(), cuda::block_dims<1>()); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims(dim3{}), cuda::cluster_dims<1>(), cuda::block_dims(dim3{})); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims(dim3{}), cuda::cluster_dims(dim3{}), cuda::block_dims<1>()); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims(dim3{}), cuda::cluster_dims(dim3{}), cuda::block_dims(dim3{})); + +int main(int, char**) +{ + return 0; +} diff --git a/libcudacxx/test/libcudacxx/cuda/hierarchy/block_level/native_hierarchy_queries.pass.cpp b/libcudacxx/test/libcudacxx/cuda/hierarchy/block_level/native_hierarchy_queries.pass.cpp new file mode 100644 index 000000000000..270af7d4abad --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/hierarchy/block_level/native_hierarchy_queries.pass.cpp @@ -0,0 +1,125 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// todo: enable with nvrtc +// UNSUPPORTED: nvrtc + +#include +#include +#include +#include +#include + +#include "hierarchy_queries.h" + +__device__ void test_block() +{ + constexpr cuda::std::size_t dext = cuda::std::dynamic_extent; + + // 1. Test cuda::block.dims(x) + { + uint3 exp{1, 1, 1}; + NV_IF_TARGET(NV_PROVIDES_SM_90, (exp = __clusterDim();)) + test_dims(exp, cuda::block, cuda::cluster); + } + test_dims(gridDim, cuda::block, cuda::grid); + + // 2. Test cuda::block.static_dims(x) + test_static_dims(ulonglong3{dext, dext, dext}, cuda::block, cuda::cluster); + test_static_dims(ulonglong3{dext, dext, dext}, cuda::block, cuda::grid); + + // 3. Test cuda::block.extents(x) + { + uint3 exp{1, 1, 1}; + NV_IF_TARGET(NV_PROVIDES_SM_90, (exp = __clusterDim();)) + test_extents(cuda::std::dims<3, unsigned>{exp.x, exp.y, exp.z}, cuda::block, cuda::cluster); + } + test_extents(cuda::std::dims<3, unsigned>{gridDim.x, gridDim.y, gridDim.z}, cuda::block, cuda::grid); + + // 4. Test cuda::block.count(x) + { + cuda::std::size_t exp = 1; + NV_IF_TARGET(NV_PROVIDES_SM_90, ({ + exp *= __clusterDim().x; + exp *= __clusterDim().y; + exp *= __clusterDim().z; + })) + test_count(exp, cuda::block, cuda::cluster); + } + test_count(cuda::std::size_t{gridDim.z} * gridDim.y * gridDim.x, cuda::block, cuda::grid); + + // 5. test cuda::block.index(x) + { + uint3 exp{0, 0, 0}; + NV_IF_TARGET(NV_PROVIDES_SM_90, (exp = __clusterRelativeBlockIdx();)) + test_index(exp, cuda::block, cuda::cluster); + } + test_index(blockIdx, cuda::block, cuda::grid); + + // 6. Test cuda::block.rank(x) + { + cuda::std::size_t exp = 0; + NV_IF_TARGET(NV_PROVIDES_SM_90, ({ + exp = ((__clusterRelativeBlockIdx().z * __clusterDim().y) + __clusterRelativeBlockIdx().y) + * __clusterDim().x + + __clusterRelativeBlockIdx().x; + })) + test_rank(exp, cuda::block, cuda::cluster); + } + { + const cuda::std::size_t exp = (blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x + blockIdx.x; + test_rank(exp, cuda::block, cuda::grid); + } +} + +#if !_CCCL_COMPILER(NVRTC) +__global__ void test_kernel() +{ + test_block(); +} + +void test() +{ + int cc_major{}; + assert(cudaDeviceGetAttribute(&cc_major, cudaDevAttrComputeCapabilityMajor, 0) == cudaSuccess); + + // thread block clusters require compute capability at least 9.0 + const bool enable_clusters = cc_major >= 9; + + test_kernel<<<1, 128>>>(); + test_kernel<<<128, 1>>>(); + test_kernel<<>>(); + test_kernel<<>>(); + if (enable_clusters) + { + cudaLaunchAttribute attribute[1]{}; + attribute[0].id = cudaLaunchAttributeClusterDimension; + attribute[0].val.clusterDim.x = 4; + attribute[0].val.clusterDim.y = 2; + attribute[0].val.clusterDim.z = 1; + + cudaLaunchConfig_t config{}; + config.gridDim = {12, 10, 3}; + config.blockDim = {2, 8, 4}; + config.attrs = attribute; + config.numAttrs = 1; + + void* pargs[1]{}; + assert(cudaLaunchKernelExC(&config, (const void*) test_kernel, pargs) == cudaSuccess); + } + + assert(cudaDeviceSynchronize() == cudaSuccess); +} +#endif // !_CCCL_COMPILER(NVRTC) + +int main(int, char**) +{ + NV_IF_ELSE_TARGET(NV_IS_HOST, (test();), (test_block();)) + return 0; +} diff --git a/libcudacxx/test/libcudacxx/cuda/hierarchy/block_level/native_hierarchy_query_signatures.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/hierarchy/block_level/native_hierarchy_query_signatures.compile.pass.cpp new file mode 100644 index 000000000000..afd683f50d28 --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/hierarchy/block_level/native_hierarchy_query_signatures.compile.pass.cpp @@ -0,0 +1,93 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// todo: enable with nvrtc +// UNSUPPORTED: nvrtc + +#include +#include +#include +#include + +template +__device__ void test_query_signatures(const Level& level) +{ + // 1. Test cuda::block_level::dims(x) signature. + static_assert(cuda::std::is_same_v, decltype(cuda::block_level::dims(level))>); + static_assert(noexcept(cuda::block_level::dims(level))); + + // 2. Test cuda::block_level::static_dims(x) signature. + static_assert(cuda::std::is_same_v, + decltype(cuda::block_level::static_dims(level))>); + static_assert(noexcept(cuda::block_level::static_dims(level))); + + // 3. Test cuda::block_level::extents(x) signature. + static_assert(cuda::std::is_same_v, decltype(cuda::block_level::extents(level))>); + static_assert(noexcept(cuda::block_level::extents(level))); + + // 4. Test cuda::block_level::count(x) signature. + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::block_level::count(level))); + + // 5. Test cuda::block_level::index(x) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::block_level::index(level))>); + static_assert(noexcept(cuda::block_level::index(level))); + + // 6. Test cuda::block_level::rank(x) signature. + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::block_level::rank(level))); +} + +template +__device__ void test_query_as_signatures(const Level& level) +{ + // 1. Test cuda::block_level::dims(x) signature. + static_assert(cuda::std::is_same_v, decltype(cuda::block_level::dims_as(level))>); + static_assert(noexcept(cuda::block_level::dims_as(level))); + + // 2. Test cuda::block_level::extents(x) signature. + static_assert(cuda::std::is_same_v, decltype(cuda::block_level::extents_as(level))>); + static_assert(noexcept(cuda::block_level::extents_as(level))); + + // 3. Test cuda::block_level::count(x) signature. + static_assert(cuda::std::is_same_v(level))>); + static_assert(noexcept(cuda::block_level::count_as(level))); + + // 4. Test cuda::block_level::index(x) signature. + static_assert(cuda::std::is_same_v, decltype(cuda::block_level::index_as(level))>); + static_assert(noexcept(cuda::block_level::index_as(level))); + + // 5. Test cuda::block_level::rank(x) signature. + static_assert(cuda::std::is_same_v(level))>); + static_assert(noexcept(cuda::block_level::rank_as(level))); +} + +template +__device__ void test(const InLevel& in_level) +{ + test_query_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); +} + +__device__ void test() +{ + test(cuda::cluster); + test(cuda::grid); +} + +int main(int, char**) +{ + return 0; +} diff --git a/libcudacxx/test/libcudacxx/cuda/hierarchy/cluster_level/hierarchy_queries.pass.cpp b/libcudacxx/test/libcudacxx/cuda/hierarchy/cluster_level/hierarchy_queries.pass.cpp new file mode 100644 index 000000000000..ddb0d3695898 --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/hierarchy/cluster_level/hierarchy_queries.pass.cpp @@ -0,0 +1,185 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// todo: enable with nvrtc +// UNSUPPORTED: nvrtc + +#include "hierarchy_queries.h" + +#include +#include +#include +#include +#include + +template +__device__ void test_cluster(const Hierarchy& hier, const GridExts& grid_exts, const ClusterExts&, const BlockExts&) +{ + uint3 dims = gridDim; + NV_IF_TARGET(NV_PROVIDES_SM_90, (dims = __clusterGridDimInClusters();)) + + uint3 index = blockIdx; + NV_IF_TARGET(NV_PROVIDES_SM_90, (index = __clusterIdx();)) + + // 1. Test cuda::cluster.dims(x, hier) + test_dims(dims, cuda::cluster, cuda::grid, hier); + + // 2. Test cuda::cluster.static_dims(x, hier) + { + const ulonglong3 exp{ + GridExts::static_extent(0), + GridExts::static_extent(1), + GridExts::static_extent(2), + }; + test_static_dims(exp, cuda::cluster, cuda::grid, hier); + } + + // 3. Test cuda::cluster.extents(x) + { + const cuda::std::extents + exp{dims.x, dims.y, dims.z}; + test_extents(exp, cuda::cluster, cuda::grid, hier); + } + + // 4. Test cuda::cluster.count(x, hier) + test_count(cuda::std::size_t{dims.z} * dims.y * dims.x, cuda::cluster, cuda::grid, hier); + + // 5. test cuda::cluster.index(x, hier) + test_index(index, cuda::cluster, cuda::grid, hier); + + // 6. Test cuda::cluster.rank(x, hier) + { + const cuda::std::size_t exp = (index.z * dims.y + index.y) * dims.x + index.x; + test_rank(exp, cuda::cluster, cuda::grid, hier); + } +} + +__device__ void test_device() +{ + // todo: make hierarchy constructible on device + // test_thread(cuda::make_hierarchy(cuda::grid_dims(gridDim), cuda::cluster_dims(clusterDim))); +} + +#if !_CCCL_COMPILER(NVRTC) +template +__global__ void test_kernel(Hierarchy hier, GridExts grid_exts, BlockExts block_exts) +{ + test_cluster(hier, grid_exts, cuda::std::extents{}, block_exts); +} + +template +__global__ void test_kernel(Hierarchy hier, GridExts grid_exts, ClusterExts cluster_exts, BlockExts block_exts) +{ + test_cluster(hier, grid_exts, cluster_exts, block_exts); +} +template +void test_launch(GridExts grid_exts, BlockExts block_exts) +{ + const dim3 grid_dims{grid_exts.extent(0), grid_exts.extent(1), grid_exts.extent(2)}; + const dim3 block_dims{block_exts.extent(0), block_exts.extent(1), block_exts.extent(2)}; + + const cuda::std::dims<3, unsigned> grid_exts_dyn{grid_exts.extent(0), grid_exts.extent(1), grid_exts.extent(2)}; + const cuda::std::dims<3, unsigned> block_exts_dyn{block_exts.extent(0), block_exts.extent(1), block_exts.extent(2)}; + + // 1. Launch hierarchy with all static extents. + test_kernel<<>>( + cuda::make_hierarchy( + cuda::grid_dims(), + cuda::block_dims()), + grid_exts, + block_exts); + + // 2. Launch hierarchy with static grid extents and dynamic block extents. + test_kernel<<>>( + cuda::make_hierarchy( + cuda::grid_dims(), + cuda::block_dims(block_dims)), + grid_exts, + block_exts_dyn); + + // 3. Launch hierarchy with dynamic grid extents and static block extents. + test_kernel<<>>( + cuda::make_hierarchy( + cuda::grid_dims(grid_dims), + cuda::block_dims()), + grid_exts_dyn, + block_exts); + + // 4. Launch hierarchy with dynamic grid extents and dynamic block extents. + test_kernel<<>>( + cuda::make_hierarchy(cuda::grid_dims(grid_dims), cuda::block_dims(block_dims)), grid_exts_dyn, block_exts_dyn); +} + +template +void test_launch(GridExts grid_exts, ClusterExts cluster_exts, BlockExts block_exts) +{ + const dim3 grid_dims{grid_exts.extent(0), grid_exts.extent(1), grid_exts.extent(2)}; + const dim3 cluster_dims{cluster_exts.extent(0), cluster_exts.extent(1), cluster_exts.extent(2)}; + const dim3 block_dims{block_exts.extent(0), block_exts.extent(1), block_exts.extent(2)}; + + cuda::std::dims<3, unsigned> grid_exts_dyn{grid_exts.extent(0), grid_exts.extent(1), grid_exts.extent(2)}; + cuda::std::dims<3, unsigned> cluster_exts_dyn{cluster_exts.extent(0), cluster_exts.extent(1), cluster_exts.extent(2)}; + cuda::std::dims<3, unsigned> block_exts_dyn{block_exts.extent(0), block_exts.extent(1), block_exts.extent(2)}; + + cudaLaunchAttribute attribute[1]{}; + attribute[0].id = cudaLaunchAttributeClusterDimension; + attribute[0].val.clusterDim.x = cluster_dims.x; + attribute[0].val.clusterDim.y = cluster_dims.y; + attribute[0].val.clusterDim.z = cluster_dims.z; + + cudaLaunchConfig_t config{}; + config.gridDim = dim3{grid_dims.x * cluster_dims.x, grid_dims.y * cluster_dims.y, grid_dims.z * cluster_dims.z}; + config.blockDim = block_dims; + config.attrs = attribute; + config.numAttrs = 1; + + // 1. Launch hierarchy with all static extents. + { + auto hier = cuda::make_hierarchy( + cuda::grid_dims(), + cuda::cluster_dims(), + cuda::block_dims()); + auto kernel = test_kernel; + assert(cudaLaunchKernelEx(&config, kernel, hier, grid_exts, cluster_exts, block_exts) == cudaSuccess); + } + + // 2. Launch hierarchy with all dynamic extents. + { + auto hier = + cuda::make_hierarchy(cuda::grid_dims(grid_dims), cuda::cluster_dims(cluster_dims), cuda::block_dims(block_dims)); + auto kernel = + test_kernel; + assert(cudaLaunchKernelEx(&config, kernel, hier, grid_exts_dyn, cluster_exts_dyn, block_exts_dyn) == cudaSuccess); + } +} + +void test() +{ + int cc_major{}; + assert(cudaDeviceGetAttribute(&cc_major, cudaDevAttrComputeCapabilityMajor, 0) == cudaSuccess); + + // thread block clusters require compute capability at least 9.0 + const bool enable_clusters = cc_major >= 9; + + if (enable_clusters) + { + test_launch(cuda::std::extents{}, + cuda::std::extents{}, + cuda::std::extents{}); + } + + assert(cudaDeviceSynchronize() == cudaSuccess); +} +#endif // !_CCCL_COMPILER(NVRTC) + +int main(int, char**) +{ + NV_IF_ELSE_TARGET(NV_IS_HOST, (test();), (test_device();)) + return 0; +} diff --git a/libcudacxx/test/libcudacxx/cuda/hierarchy/cluster_level/hierarchy_query_signatures.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/hierarchy/cluster_level/hierarchy_query_signatures.compile.pass.cpp new file mode 100644 index 000000000000..4559783539fc --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/hierarchy/cluster_level/hierarchy_query_signatures.compile.pass.cpp @@ -0,0 +1,119 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// todo: enable with nvrtc +// UNSUPPORTED: nvrtc + +#include +#include +#include +#include + +template +__device__ void test_query_signatures(const Level& level, const Hierarchy& hier) +{ + // 1. Test cuda::cluster_level::dims(x, hier) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::cluster_level::dims(level, hier))>); + static_assert(noexcept(cuda::cluster_level::dims(level, hier))); + + // 2. Test cuda::cluster_level::static_dims(x, hier) signature. + static_assert(cuda::std::is_same_v, + decltype(cuda::cluster_level::static_dims(level, hier))>); + static_assert(noexcept(cuda::cluster_level::static_dims(level, hier))); + + // 3. Test cuda::cluster_level::extents(x, hier) signature. + using ExtentsResult = decltype(cuda::cluster_level::extents(level, hier)); + static_assert(cuda::std::__is_cuda_std_extents_v); + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::cluster_level::extents(level, hier))); + + // 4. Test cuda::cluster_level::count(x, hier) signature. + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::cluster_level::count(level, hier))); + + // 5. Test cuda::cluster_level::index(x, hier) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::cluster_level::index(level, hier))>); + static_assert(noexcept(cuda::cluster_level::index(level, hier))); + + // 6. Test cuda::cluster_level::rank(x, hier) signature. + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::cluster_level::rank(level, hier))); +} + +template +__device__ void test_query_as_signatures(const Level& level, const Hierarchy& hier) +{ + // 1. Test cuda::cluster_level::dims_as(x, hier) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::cluster_level::dims_as(level, hier))>); + static_assert(noexcept(cuda::cluster_level::dims_as(level, hier))); + + // 2. Test cuda::cluster_level::extents_as(x, hier) signature. + using ExtentsResult = decltype(cuda::cluster_level::extents_as(level, hier)); + static_assert(cuda::std::__is_cuda_std_extents_v); + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::cluster_level::extents_as(level, hier))); + + // 3. Test cuda::cluster_level::count_as(x, hier) signature. + static_assert(cuda::std::is_same_v(level, hier))>); + static_assert(noexcept(cuda::cluster_level::count_as(level, hier))); + + // 4. Test cuda::cluster_level::index_as(x, hier) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::cluster_level::index_as(level, hier))>); + static_assert(noexcept(cuda::cluster_level::index_as(level, hier))); + + // 5. Test cuda::cluster_level::rank_as(x, hier) signature. + static_assert(cuda::std::is_same_v(level, hier))>); + static_assert(noexcept(cuda::cluster_level::rank_as(level, hier))); +} + +template +__device__ void test(const InLevel& in_level, const Hierarchy& hier) +{ + test_query_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); +} + +template +__device__ void test(const Hierarchy& hier) +{ + test(cuda::grid, hier); +} + +template +__global__ void test_kernel(Hierarchy hier) +{ + test(hier); +} + +#define TEST_KERNEL_INSTANTIATE(...) \ + template __global__ void test_kernel( \ + decltype(cuda::make_hierarchy(__VA_ARGS__))) + +TEST_KERNEL_INSTANTIATE(cuda::grid_dims<1>(), cuda::cluster_dims<1>(), cuda::block_dims<1>()); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims<1>(), cuda::cluster_dims<1>(), cuda::block_dims(dim3{})); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims<1>(), cuda::cluster_dims(dim3{}), cuda::block_dims<1>()); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims<1>(), cuda::cluster_dims(dim3{}), cuda::block_dims(dim3{})); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims(dim3{}), cuda::cluster_dims<1>(), cuda::block_dims<1>()); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims(dim3{}), cuda::cluster_dims<1>(), cuda::block_dims(dim3{})); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims(dim3{}), cuda::cluster_dims(dim3{}), cuda::block_dims<1>()); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims(dim3{}), cuda::cluster_dims(dim3{}), cuda::block_dims(dim3{})); + +int main(int, char**) +{ + return 0; +} diff --git a/libcudacxx/test/libcudacxx/cuda/hierarchy/cluster_level/native_hierarchy_queries.pass.cpp b/libcudacxx/test/libcudacxx/cuda/hierarchy/cluster_level/native_hierarchy_queries.pass.cpp new file mode 100644 index 000000000000..4e58b131fa2c --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/hierarchy/cluster_level/native_hierarchy_queries.pass.cpp @@ -0,0 +1,97 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// todo: enable with nvrtc +// UNSUPPORTED: nvrtc + +#include +#include +#include +#include +#include + +#include "hierarchy_queries.h" + +__device__ void test_cluster() +{ + constexpr cuda::std::size_t dext = cuda::std::dynamic_extent; + + uint3 dims = gridDim; + NV_IF_TARGET(NV_PROVIDES_SM_90, (dims = __clusterGridDimInClusters();)) + + uint3 index = blockIdx; + NV_IF_TARGET(NV_PROVIDES_SM_90, (index = __clusterIdx();)) + + // 1. Test cuda::cluster.dims(x) + test_dims(dims, cuda::cluster, cuda::grid); + + // 2. Test cuda::cluster.static_dims(x) + test_static_dims(ulonglong3{dext, dext, dext}, cuda::cluster, cuda::grid); + + // 3. Test cuda::cluster.extents(x) + test_extents(cuda::std::dims<3, unsigned>{dims.x, dims.y, dims.z}, cuda::cluster, cuda::grid); + + // 4. Test cuda::cluster.count(x) + test_count(cuda::std::size_t{dims.z} * dims.y * dims.x, cuda::cluster, cuda::grid); + + // 5. test cuda::cluster.index(x) + test_index(index, cuda::cluster, cuda::grid); + + // 6. Test cuda::cluster.rank(x) + { + const cuda::std::size_t exp = (index.z * dims.y + index.y) * dims.x + index.x; + test_rank(exp, cuda::cluster, cuda::grid); + } +} + +#if !_CCCL_COMPILER(NVRTC) +__global__ void test_kernel() +{ + test_cluster(); +} + +void test() +{ + int cc_major{}; + assert(cudaDeviceGetAttribute(&cc_major, cudaDevAttrComputeCapabilityMajor, 0) == cudaSuccess); + + // thread block clusters require compute capability at least 9.0 + const bool enable_clusters = cc_major >= 9; + + test_kernel<<<1, 128>>>(); + test_kernel<<<128, 1>>>(); + test_kernel<<>>(); + test_kernel<<>>(); + if (enable_clusters) + { + cudaLaunchAttribute attribute[1]{}; + attribute[0].id = cudaLaunchAttributeClusterDimension; + attribute[0].val.clusterDim.x = 4; + attribute[0].val.clusterDim.y = 2; + attribute[0].val.clusterDim.z = 1; + + cudaLaunchConfig_t config{}; + config.gridDim = {12, 10, 3}; + config.blockDim = {2, 8, 4}; + config.attrs = attribute; + config.numAttrs = 1; + + void* pargs[1]{}; + assert(cudaLaunchKernelExC(&config, (const void*) test_kernel, pargs) == cudaSuccess); + } + + assert(cudaDeviceSynchronize() == cudaSuccess); +} +#endif // !_CCCL_COMPILER(NVRTC) + +int main(int, char**) +{ + NV_IF_ELSE_TARGET(NV_IS_HOST, (test();), (test_cluster();)) + return 0; +} diff --git a/libcudacxx/test/libcudacxx/cuda/hierarchy/cluster_level/native_hierarchy_query_signatures.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/hierarchy/cluster_level/native_hierarchy_query_signatures.compile.pass.cpp new file mode 100644 index 000000000000..64c0905c8d82 --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/hierarchy/cluster_level/native_hierarchy_query_signatures.compile.pass.cpp @@ -0,0 +1,94 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// todo: enable with nvrtc +// UNSUPPORTED: nvrtc + +#include +#include +#include + +template +__device__ void test_query_signatures(const Level& level) +{ + // 1. Test cuda::cluster_level::dims(x) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::cluster_level::dims(level))>); + static_assert(noexcept(cuda::cluster_level::dims(level))); + + // 2. Test cuda::cluster_level::static_dims(x) signature. + static_assert(cuda::std::is_same_v, + decltype(cuda::cluster_level::static_dims(level))>); + static_assert(noexcept(cuda::cluster_level::static_dims(level))); + + // 3. Test cuda::cluster_level::extents(x) signature. + static_assert(cuda::std::is_same_v, decltype(cuda::cluster_level::extents(level))>); + static_assert(noexcept(cuda::cluster_level::extents(level))); + + // 4. Test cuda::cluster_level::count(x) signature. + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::cluster_level::count(level))); + + // 5. Test cuda::cluster_level::index(x) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::cluster_level::index(level))>); + static_assert(noexcept(cuda::cluster_level::index(level))); + + // 6. Test cuda::cluster_level::rank(x) signature. + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::cluster_level::rank(level))); +} + +template +__device__ void test_query_as_signatures(const Level& level) +{ + // 1. Test cuda::cluster_level::dims(x) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::cluster_level::dims_as(level))>); + static_assert(noexcept(cuda::cluster_level::dims_as(level))); + + // 2. Test cuda::cluster_level::extents(x) signature. + static_assert(cuda::std::is_same_v, decltype(cuda::cluster_level::extents_as(level))>); + static_assert(noexcept(cuda::cluster_level::extents_as(level))); + + // 3. Test cuda::cluster_level::count(x) signature. + static_assert(cuda::std::is_same_v(level))>); + static_assert(noexcept(cuda::cluster_level::count_as(level))); + + // 4. Test cuda::cluster_level::index(x) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::cluster_level::index_as(level))>); + static_assert(noexcept(cuda::cluster_level::index_as(level))); + + // 5. Test cuda::cluster_level::rank(x) signature. + static_assert(cuda::std::is_same_v(level))>); + static_assert(noexcept(cuda::cluster_level::rank_as(level))); +} + +template +__device__ void test(const InLevel& in_level) +{ + test_query_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); +} + +__device__ void test() +{ + test(cuda::grid); +} + +int main(int, char**) +{ + return 0; +} diff --git a/libcudacxx/test/libcudacxx/cuda/hierarchy/hierarchy_objects.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/hierarchy/hierarchy_objects.compile.pass.cpp new file mode 100644 index 000000000000..c3d7fc4f917b --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/hierarchy/hierarchy_objects.compile.pass.cpp @@ -0,0 +1,24 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// todo: enable with nvrtc +// UNSUPPORTED: nvrtc + +#include +#include + +static_assert(cuda::std::is_same_v>); +static_assert(cuda::std::is_same_v>); +static_assert(cuda::std::is_same_v>); +static_assert(cuda::std::is_same_v>); + +int main(int, char**) +{ + return 0; +} diff --git a/libcudacxx/test/libcudacxx/cuda/hierarchy/hierarchy_query_result.pass.cpp b/libcudacxx/test/libcudacxx/cuda/hierarchy/hierarchy_query_result.pass.cpp new file mode 100644 index 000000000000..e9cc316462eb --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/hierarchy/hierarchy_query_result.pass.cpp @@ -0,0 +1,113 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// todo: enable with nvrtc +// UNSUPPORTED: nvrtc + +#include +#include +#include +#include + +template +__host__ __device__ constexpr void test() +{ + using HQR = cuda::hierarchy_query_result; + using Vec = cuda::__vector_type_t; + constexpr auto has_vec = !cuda::std::is_same_v; + + // 1. Test value_type + static_assert(cuda::std::is_same_v); + + // 2. Test constructors + static_assert(cuda::std::is_trivially_default_constructible_v); + static_assert(cuda::std::is_trivially_copyable_v); + + // 3. Test public members + { + HQR v{T{0}, T{1}, T{2}}; + assert(v.x == static_cast(0)); + assert(v.y == static_cast(1)); + assert(v.z == static_cast(2)); + } + + // 4. Test operator[] const + static_assert(cuda::std::is_same_v()[cuda::std::size_t{}])>); + static_assert(noexcept(cuda::std::declval()[cuda::std::size_t{}])); + { + const HQR v{T{0}, T{1}, T{2}}; + for (cuda::std::size_t i = 0; i < 3; ++i) + { + assert(v[i] == static_cast(i)); + } + } + + // 5. Test operator[] + static_assert(cuda::std::is_same_v()[cuda::std::size_t{}])>); + static_assert(noexcept(cuda::std::declval()[cuda::std::size_t{}])); + { + HQR v{T{0}, T{1}, T{2}}; + for (cuda::std::size_t i = 0; i < 3; ++i) + { + assert(v[i] == static_cast(i)); + } + } + + // 6. Test operator vector-type + static_assert(!has_vec || cuda::std::is_nothrow_convertible_v); + if constexpr (has_vec) + { + const HQR v{T{0}, T{1}, T{2}}; + Vec vec = v; + assert(vec.x == v.x); + assert(vec.y == v.y); + assert(vec.z == v.z); + } + + // 7. Test dim3 can be constructed from the query result + static_assert(!cuda::std::is_same_v || cuda::std::is_constructible_v); + if constexpr (cuda::std::is_same_v) + { + const HQR v{T{0}, T{1}, T{2}}; + dim3 vec{v}; + assert(vec.x == v.x); + assert(vec.y == v.y); + assert(vec.z == v.z); + } +} + +__host__ __device__ constexpr bool test() +{ + test(); + test(); + test(); + test(); + test(); +#if _CCCL_HAS_INT128() + test<__int128_t>(); +#endif // _CCCL_HAS_INT128(); + + test(); + test(); + test(); + test(); + test(); +#if _CCCL_HAS_INT128() + test<__uint128_t>(); +#endif // _CCCL_HAS_INT128(); + + return true; +} + +int main(int, char**) +{ + test(); + static_assert(test()); + return 0; +} diff --git a/libcudacxx/test/libcudacxx/cuda/hierarchy/thread_level/hierarchy_queries.pass.cpp b/libcudacxx/test/libcudacxx/cuda/hierarchy/thread_level/hierarchy_queries.pass.cpp new file mode 100644 index 000000000000..220d7107274d --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/hierarchy/thread_level/hierarchy_queries.pass.cpp @@ -0,0 +1,285 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// todo: enable with nvrtc +// UNSUPPORTED: nvrtc + +#include "hierarchy_queries.h" + +#include +#include +#include +#include +#include + +template +__device__ void test_thread( + const Hierarchy& hier, const GridExts& grid_exts, const ClusterExts& cluster_exts, const BlockExts& block_exts) +{ + // 1. Test cuda::gpu_thread.dims(x, hier) + test_dims(blockDim, cuda::gpu_thread, cuda::block, hier); + if constexpr (cuda::has_level_v) + { + uint3 exp = blockDim; + NV_IF_TARGET(NV_PROVIDES_SM_90, ({ + exp.x *= __clusterDim().x; + exp.y *= __clusterDim().y; + exp.z *= __clusterDim().z; + })) + test_dims(exp, cuda::gpu_thread, cuda::cluster, hier); + } + { + const uint3 exp{blockDim.x * gridDim.x, blockDim.y * gridDim.y, blockDim.z * gridDim.z}; + test_dims(exp, cuda::gpu_thread, cuda::grid, hier); + } + + // 2. Test cuda::gpu_thread.static_dims(x, hier) + test_static_dims(ulonglong3{BlockExts::static_extent(0), BlockExts::static_extent(1), BlockExts::static_extent(2)}, + cuda::gpu_thread, + cuda::block, + hier); + if constexpr (cuda::has_level_v) + { + const ulonglong3 exp{ + mul_static_extents(ClusterExts::static_extent(0), BlockExts::static_extent(0)), + mul_static_extents(ClusterExts::static_extent(1), BlockExts::static_extent(1)), + mul_static_extents(ClusterExts::static_extent(2), BlockExts::static_extent(2)), + }; + test_static_dims(exp, cuda::gpu_thread, cuda::cluster, hier); + } + { + const ulonglong3 exp{ + mul_static_extents(GridExts::static_extent(0), ClusterExts::static_extent(0), BlockExts::static_extent(0)), + mul_static_extents(GridExts::static_extent(1), ClusterExts::static_extent(1), BlockExts::static_extent(1)), + mul_static_extents(GridExts::static_extent(2), ClusterExts::static_extent(2), BlockExts::static_extent(2)), + }; + test_static_dims(exp, cuda::gpu_thread, cuda::grid, hier); + } + + // 3. Test cuda::gpu_thread.extents(x) + test_extents(block_exts, cuda::gpu_thread, cuda::block, hier); + if constexpr (cuda::has_level_v) + { + uint3 dims = blockDim; + NV_IF_TARGET(NV_PROVIDES_SM_90, ({ + dims.x *= __clusterDim().x; + dims.y *= __clusterDim().y; + dims.z *= __clusterDim().z; + })) + + const cuda::std::extents + exp{dims.x, dims.y, dims.z}; + + test_extents(exp, cuda::gpu_thread, cuda::cluster, hier); + } + { + const cuda::std::extents< + unsigned, + mul_static_extents(GridExts::static_extent(0), ClusterExts::static_extent(0), BlockExts::static_extent(0)), + mul_static_extents(GridExts::static_extent(1), ClusterExts::static_extent(1), BlockExts::static_extent(1)), + mul_static_extents(GridExts::static_extent(2), ClusterExts::static_extent(2), BlockExts::static_extent(2))> + exp{blockDim.x * gridDim.x, blockDim.y * gridDim.y, blockDim.z * gridDim.z}; + test_extents(exp, cuda::gpu_thread, cuda::grid, hier); + } + + // 4. Test cuda::gpu_thread.count(x, hier) + test_count(cuda::std::size_t{blockDim.z} * blockDim.y * blockDim.x, cuda::gpu_thread, cuda::block, hier); + if constexpr (cuda::has_level_v) + { + uint3 exp = blockDim; + NV_IF_TARGET(NV_PROVIDES_SM_90, ({ + exp.x *= __clusterDim().x; + exp.y *= __clusterDim().y; + exp.z *= __clusterDim().z; + })) + test_count(cuda::std::size_t{exp.z} * exp.y * exp.x, cuda::gpu_thread, cuda::cluster, hier); + } + { + const uint3 exp{blockDim.x * gridDim.x, blockDim.y * gridDim.y, blockDim.z * gridDim.z}; + test_count(cuda::std::size_t{exp.z} * exp.y * exp.x, cuda::gpu_thread, cuda::grid, hier); + } + + // 5. test cuda::gpu_thread.index(x, hier) + test_index(threadIdx, cuda::gpu_thread, cuda::block, hier); + if constexpr (cuda::has_level_v) + { + uint3 exp = threadIdx; + NV_IF_TARGET(NV_PROVIDES_SM_90, ({ + exp.x += blockDim.x * __clusterRelativeBlockIdx().x; + exp.y += blockDim.y * __clusterRelativeBlockIdx().y; + exp.z += blockDim.z * __clusterRelativeBlockIdx().z; + })) + test_index(exp, cuda::gpu_thread, cuda::cluster, hier); + } + { + const uint3 exp{ + threadIdx.x + blockDim.x * blockIdx.x, + threadIdx.y + blockDim.y * blockIdx.y, + threadIdx.z + blockDim.z * blockIdx.z, + }; + test_index(exp, cuda::gpu_thread, cuda::grid, hier); + } + + // 6. Test cuda::gpu_thread.rank(x, hier) + test_rank((threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x + threadIdx.x, cuda::gpu_thread, cuda::block, hier); + if constexpr (cuda::has_level_v) + { + cuda::std::size_t exp = 0; + NV_IF_ELSE_TARGET(NV_PROVIDES_SM_90, + ({ + exp = (((__clusterRelativeBlockIdx().z * __clusterDim().y * __clusterDim().x) + + __clusterRelativeBlockIdx().y * __clusterDim().x) + + __clusterRelativeBlockIdx().x) + * (blockDim.x * blockDim.y * blockDim.z) + + ((threadIdx.z * blockDim.y * blockDim.x) + threadIdx.y * blockDim.x) + threadIdx.x; + }), + ({ exp = ((threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x) + threadIdx.x; })) + test_rank(exp, cuda::gpu_thread, cuda::cluster, hier); + } + { + const cuda::std::size_t exp = + (blockIdx.z * gridDim.y * gridDim.x + blockIdx.y * gridDim.x + blockIdx.x) + * (blockDim.x * blockDim.y * blockDim.z) + + threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x; + test_rank(exp, cuda::gpu_thread, cuda::grid, hier); + } +} + +__device__ void test_device() +{ + // todo: make hierarchy constructible on device + // test_thread(cuda::make_hierarchy(cuda::grid_dims(gridDim), cuda::block_dims(blockDim))); +} + +#if !_CCCL_COMPILER(NVRTC) +template +__global__ void test_kernel(Hierarchy hier, GridExts grid_exts, BlockExts block_exts) +{ + test_thread(hier, grid_exts, cuda::std::extents{}, block_exts); +} + +template +__global__ void test_kernel(Hierarchy hier, GridExts grid_exts, ClusterExts cluster_exts, BlockExts block_exts) +{ + test_thread(hier, grid_exts, cluster_exts, block_exts); +} + +template +void test_launch(GridExts grid_exts, BlockExts block_exts) +{ + const dim3 grid_dims{grid_exts.extent(0), grid_exts.extent(1), grid_exts.extent(2)}; + const dim3 block_dims{block_exts.extent(0), block_exts.extent(1), block_exts.extent(2)}; + + const cuda::std::dims<3, unsigned> grid_exts_dyn{grid_exts.extent(0), grid_exts.extent(1), grid_exts.extent(2)}; + const cuda::std::dims<3, unsigned> block_exts_dyn{block_exts.extent(0), block_exts.extent(1), block_exts.extent(2)}; + + // 1. Launch hierarchy with all static extents. + test_kernel<<>>( + cuda::make_hierarchy( + cuda::grid_dims(), + cuda::block_dims()), + grid_exts, + block_exts); + + // 2. Launch hierarchy with static grid extents and dynamic block extents. + test_kernel<<>>( + cuda::make_hierarchy( + cuda::grid_dims(), + cuda::block_dims(block_dims)), + grid_exts, + block_exts_dyn); + + // 3. Launch hierarchy with dynamic grid extents and static block extents. + test_kernel<<>>( + cuda::make_hierarchy( + cuda::grid_dims(grid_dims), + cuda::block_dims()), + grid_exts_dyn, + block_exts); + + // 4. Launch hierarchy with dynamic grid extents and dynamic block extents. + test_kernel<<>>( + cuda::make_hierarchy(cuda::grid_dims(grid_dims), cuda::block_dims(block_dims)), grid_exts_dyn, block_exts_dyn); +} + +template +void test_launch(GridExts grid_exts, ClusterExts cluster_exts, BlockExts block_exts) +{ + const dim3 grid_dims{grid_exts.extent(0), grid_exts.extent(1), grid_exts.extent(2)}; + const dim3 cluster_dims{cluster_exts.extent(0), cluster_exts.extent(1), cluster_exts.extent(2)}; + const dim3 block_dims{block_exts.extent(0), block_exts.extent(1), block_exts.extent(2)}; + + cuda::std::dims<3, unsigned> grid_exts_dyn{grid_exts.extent(0), grid_exts.extent(1), grid_exts.extent(2)}; + cuda::std::dims<3, unsigned> cluster_exts_dyn{cluster_exts.extent(0), cluster_exts.extent(1), cluster_exts.extent(2)}; + cuda::std::dims<3, unsigned> block_exts_dyn{block_exts.extent(0), block_exts.extent(1), block_exts.extent(2)}; + + cudaLaunchAttribute attribute[1]{}; + attribute[0].id = cudaLaunchAttributeClusterDimension; + attribute[0].val.clusterDim.x = cluster_dims.x; + attribute[0].val.clusterDim.y = cluster_dims.y; + attribute[0].val.clusterDim.z = cluster_dims.z; + + cudaLaunchConfig_t config{}; + config.gridDim = dim3{grid_dims.x * cluster_dims.x, grid_dims.y * cluster_dims.y, grid_dims.z * cluster_dims.z}; + config.blockDim = block_dims; + config.attrs = attribute; + config.numAttrs = 1; + + // 1. Launch hierarchy with all static extents. + { + auto hier = cuda::make_hierarchy( + cuda::grid_dims(), + cuda::cluster_dims(), + cuda::block_dims()); + auto kernel = test_kernel; + assert(cudaLaunchKernelEx(&config, kernel, hier, grid_exts, cluster_exts, block_exts) == cudaSuccess); + } + + // 2. Launch hierarchy with all dynamic extents. + { + auto hier = + cuda::make_hierarchy(cuda::grid_dims(grid_dims), cuda::cluster_dims(cluster_dims), cuda::block_dims(block_dims)); + auto kernel = + test_kernel; + assert(cudaLaunchKernelEx(&config, kernel, hier, grid_exts_dyn, cluster_exts_dyn, block_exts_dyn) == cudaSuccess); + } +} + +void test() +{ + int cc_major{}; + assert(cudaDeviceGetAttribute(&cc_major, cudaDevAttrComputeCapabilityMajor, 0) == cudaSuccess); + + // thread block clusters require compute capability at least 9.0 + const bool enable_clusters = cc_major >= 9; + + test_launch(cuda::std::extents{}, cuda::std::extents{}); + test_launch(cuda::std::extents{}, cuda::std::extents{}); + test_launch(cuda::std::extents{}, cuda::std::extents{}); + test_launch(cuda::std::extents{}, cuda::std::extents{}); + + if (enable_clusters) + { + test_launch(cuda::std::extents{}, + cuda::std::extents{}, + cuda::std::extents{}); + } + + assert(cudaDeviceSynchronize() == cudaSuccess); +} +#endif // !_CCCL_COMPILER(NVRTC) + +int main(int, char**) +{ + NV_IF_ELSE_TARGET(NV_IS_HOST, (test();), (test_device();)) + return 0; +} diff --git a/libcudacxx/test/libcudacxx/cuda/hierarchy/thread_level/hierarchy_query_signatures.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/hierarchy/thread_level/hierarchy_query_signatures.compile.pass.cpp new file mode 100644 index 000000000000..8c693158df3c --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/hierarchy/thread_level/hierarchy_query_signatures.compile.pass.cpp @@ -0,0 +1,129 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// todo: enable with nvrtc +// UNSUPPORTED: nvrtc + +#include +#include +#include +#include + +template +__device__ void test_query_signatures(const Level& level, const Hierarchy& hier) +{ + // 1. Test cuda::thread_level::dims(x, hier) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::thread_level::dims(level, hier))>); + static_assert(noexcept(cuda::thread_level::dims(level, hier))); + + // 2. Test cuda::thread_level::static_dims(x, hier) signature. + static_assert(cuda::std::is_same_v, + decltype(cuda::thread_level::static_dims(level, hier))>); + static_assert(noexcept(cuda::thread_level::static_dims(level, hier))); + + // 3. Test cuda::thread_level::extents(x, hier) signature. + using ExtentsResult = decltype(cuda::thread_level::extents(level, hier)); + static_assert(cuda::std::__is_cuda_std_extents_v); + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::thread_level::extents(level, hier))); + + // 4. Test cuda::thread_level::count(x, hier) signature. + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::thread_level::count(level, hier))); + + // 5. Test cuda::thread_level::index(x, hier) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::thread_level::index(level, hier))>); + static_assert(noexcept(cuda::thread_level::index(level, hier))); + + // 6. Test cuda::thread_level::rank(x, hier) signature. + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::thread_level::rank(level, hier))); +} + +template +__device__ void test_query_as_signatures(const Level& level, const Hierarchy& hier) +{ + // 1. Test cuda::thread_level::dims_as(x, hier) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::thread_level::dims_as(level, hier))>); + static_assert(noexcept(cuda::thread_level::dims_as(level, hier))); + + // 2. Test cuda::thread_level::extents_as(x, hier) signature. + using ExtentsResult = decltype(cuda::thread_level::extents_as(level, hier)); + static_assert(cuda::std::__is_cuda_std_extents_v); + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::thread_level::extents_as(level, hier))); + + // 3. Test cuda::thread_level::count_as(x, hier) signature. + static_assert(cuda::std::is_same_v(level, hier))>); + static_assert(noexcept(cuda::thread_level::count_as(level, hier))); + + // 4. Test cuda::thread_level::index_as(x, hier) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::thread_level::index_as(level, hier))>); + static_assert(noexcept(cuda::thread_level::index_as(level, hier))); + + // 5. Test cuda::thread_level::rank_as(x, hier) signature. + static_assert(cuda::std::is_same_v(level, hier))>); + static_assert(noexcept(cuda::thread_level::rank_as(level, hier))); +} + +template +__device__ void test(const InLevel& in_level, const Hierarchy& hier) +{ + test_query_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); + test_query_as_signatures(in_level, hier); +} + +template +__device__ void test(const Hierarchy& hier) +{ + test(cuda::block, hier); + if constexpr (cuda::has_level_v) + { + test(cuda::cluster, hier); + } + test(cuda::grid, hier); +} + +template +__global__ void test_kernel(Hierarchy hier) +{ + test(hier); +} + +#define TEST_KERNEL_INSTANTIATE(...) \ + template __global__ void test_kernel( \ + decltype(cuda::make_hierarchy(__VA_ARGS__))) + +TEST_KERNEL_INSTANTIATE(cuda::grid_dims<1>(), cuda::block_dims<1>()); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims<1>(), cuda::block_dims(dim3{})); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims(dim3{}), cuda::block_dims<1>()); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims(dim3{}), cuda::block_dims(dim3{})); + +TEST_KERNEL_INSTANTIATE(cuda::grid_dims<1>(), cuda::cluster_dims<1>(), cuda::block_dims<1>()); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims<1>(), cuda::cluster_dims<1>(), cuda::block_dims(dim3{})); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims<1>(), cuda::cluster_dims(dim3{}), cuda::block_dims<1>()); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims<1>(), cuda::cluster_dims(dim3{}), cuda::block_dims(dim3{})); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims(dim3{}), cuda::cluster_dims<1>(), cuda::block_dims<1>()); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims(dim3{}), cuda::cluster_dims<1>(), cuda::block_dims(dim3{})); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims(dim3{}), cuda::cluster_dims(dim3{}), cuda::block_dims<1>()); +TEST_KERNEL_INSTANTIATE(cuda::grid_dims(dim3{}), cuda::cluster_dims(dim3{}), cuda::block_dims(dim3{})); + +int main(int, char**) +{ + return 0; +} diff --git a/libcudacxx/test/libcudacxx/cuda/hierarchy/thread_level/native_hierarchy_queries.pass.cpp b/libcudacxx/test/libcudacxx/cuda/hierarchy/thread_level/native_hierarchy_queries.pass.cpp new file mode 100644 index 000000000000..7ec30378b032 --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/hierarchy/thread_level/native_hierarchy_queries.pass.cpp @@ -0,0 +1,169 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// todo: enable with nvrtc +// UNSUPPORTED: nvrtc + +#include +#include +#include +#include +#include + +#include "hierarchy_queries.h" + +__device__ void test_thread() +{ + constexpr cuda::std::size_t dext = cuda::std::dynamic_extent; + + // 1. Test cuda::gpu_thread.dims(x) + test_dims(uint3{static_cast(warpSize), 1u, 1u}, cuda::gpu_thread, cuda::warp); + test_dims(blockDim, cuda::gpu_thread, cuda::block); + { + uint3 exp = blockDim; + NV_IF_TARGET(NV_PROVIDES_SM_90, ({ + exp.x *= __clusterDim().x; + exp.y *= __clusterDim().y; + exp.z *= __clusterDim().z; + })) + test_dims(exp, cuda::gpu_thread, cuda::cluster); + } + test_dims({blockDim.x * gridDim.x, blockDim.y * gridDim.y, blockDim.z * gridDim.z}, cuda::gpu_thread, cuda::grid); + + // 2. Test cuda::gpu_thread.static_dims(x) + test_static_dims(ulonglong3{cuda::std::size_t{32}, 1, 1}, cuda::gpu_thread, cuda::warp); + test_static_dims(ulonglong3{dext, dext, dext}, cuda::gpu_thread, cuda::block); + test_static_dims(ulonglong3{dext, dext, dext}, cuda::gpu_thread, cuda::cluster); + test_static_dims(ulonglong3{dext, dext, dext}, cuda::gpu_thread, cuda::grid); + + // 3. Test cuda::gpu_thread.extents(x) + test_extents(cuda::std::extents{}, cuda::gpu_thread, cuda::warp); + test_extents(cuda::std::dims<3, unsigned>{blockDim.x, blockDim.y, blockDim.z}, cuda::gpu_thread, cuda::block); + { + uint3 exp = blockDim; + NV_IF_TARGET(NV_PROVIDES_SM_90, ({ + exp.x *= __clusterDim().x; + exp.y *= __clusterDim().y; + exp.z *= __clusterDim().z; + })) + test_extents(cuda::std::dims<3, unsigned>{exp.x, exp.y, exp.z}, cuda::gpu_thread, cuda::cluster); + } + { + const uint3 exp{blockDim.x * gridDim.x, blockDim.y * gridDim.y, blockDim.z * gridDim.z}; + test_extents(cuda::std::dims<3, unsigned>{exp.x, exp.y, exp.z}, cuda::gpu_thread, cuda::grid); + } + + // 4. Test cuda::gpu_thread.count(x) + test_count(32, cuda::gpu_thread, cuda::warp); + test_count(cuda::std::size_t{blockDim.z} * blockDim.y * blockDim.x, cuda::gpu_thread, cuda::block); + { + uint3 exp = blockDim; + NV_IF_TARGET(NV_PROVIDES_SM_90, ({ + exp.x *= __clusterDim().x; + exp.y *= __clusterDim().y; + exp.z *= __clusterDim().z; + })) + test_count(cuda::std::size_t{exp.z} * exp.y * exp.x, cuda::gpu_thread, cuda::cluster); + } + { + const uint3 exp{blockDim.x * gridDim.x, blockDim.y * gridDim.y, blockDim.z * gridDim.z}; + test_count(cuda::std::size_t{exp.z} * exp.y * exp.x, cuda::gpu_thread, cuda::grid); + } + + // 5. test cuda::gpu_thread.index(x) + test_index(uint3{cuda::ptx::get_sreg_laneid(), 0, 0}, cuda::gpu_thread, cuda::warp); + test_index(threadIdx, cuda::gpu_thread, cuda::block); + { + uint3 exp = threadIdx; + NV_IF_TARGET(NV_PROVIDES_SM_90, ({ + exp.x += blockDim.x * __clusterRelativeBlockIdx().x; + exp.y += blockDim.y * __clusterRelativeBlockIdx().y; + exp.z += blockDim.z * __clusterRelativeBlockIdx().z; + })) + test_index(exp, cuda::gpu_thread, cuda::cluster); + } + { + const uint3 exp{ + threadIdx.x + blockDim.x * blockIdx.x, + threadIdx.y + blockDim.y * blockIdx.y, + threadIdx.z + blockDim.z * blockIdx.z, + }; + test_index(exp, cuda::gpu_thread, cuda::grid); + } + + // 6. Test cuda::gpu_thread.rank(x) + test_rank(cuda::ptx::get_sreg_laneid(), cuda::gpu_thread, cuda::warp); + test_rank((threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x + threadIdx.x, cuda::gpu_thread, cuda::block); + { + cuda::std::size_t exp = 0; + NV_IF_ELSE_TARGET(NV_PROVIDES_SM_90, + ({ + exp = (((__clusterRelativeBlockIdx().z * __clusterDim().y * __clusterDim().x) + + __clusterRelativeBlockIdx().y * __clusterDim().x) + + __clusterRelativeBlockIdx().x) + * (blockDim.x * blockDim.y * blockDim.z) + + ((threadIdx.z * blockDim.y * blockDim.x) + threadIdx.y * blockDim.x) + threadIdx.x; + }), + ({ exp = ((threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x) + threadIdx.x; })) + test_rank(exp, cuda::gpu_thread, cuda::cluster); + } + { + const cuda::std::size_t exp = + (blockIdx.z * gridDim.y * gridDim.x + blockIdx.y * gridDim.x + blockIdx.x) + * (blockDim.x * blockDim.y * blockDim.z) + + threadIdx.z * blockDim.y * blockDim.x + threadIdx.y * blockDim.x + threadIdx.x; + test_rank(exp, cuda::gpu_thread, cuda::grid); + } +} + +#if !_CCCL_COMPILER(NVRTC) +__global__ void test_kernel() +{ + test_thread(); +} + +void test() +{ + int cc_major{}; + assert(cudaDeviceGetAttribute(&cc_major, cudaDevAttrComputeCapabilityMajor, 0) == cudaSuccess); + + // thread block clusters require compute capability at least 9.0 + const bool enable_clusters = cc_major >= 9; + + test_kernel<<<1, 128>>>(); + test_kernel<<<128, 1>>>(); + test_kernel<<>>(); + test_kernel<<>>(); + if (enable_clusters) + { + cudaLaunchAttribute attribute[1]{}; + attribute[0].id = cudaLaunchAttributeClusterDimension; + attribute[0].val.clusterDim.x = 4; + attribute[0].val.clusterDim.y = 2; + attribute[0].val.clusterDim.z = 1; + + cudaLaunchConfig_t config{}; + config.gridDim = {12, 10, 3}; + config.blockDim = {2, 8, 4}; + config.attrs = attribute; + config.numAttrs = 1; + + void* pargs[1]{}; + assert(cudaLaunchKernelExC(&config, (const void*) test_kernel, pargs) == cudaSuccess); + } + + assert(cudaDeviceSynchronize() == cudaSuccess); +} +#endif // !_CCCL_COMPILER(NVRTC) + +int main(int, char**) +{ + NV_IF_ELSE_TARGET(NV_IS_HOST, (test();), (test_thread();)) + return 0; +} diff --git a/libcudacxx/test/libcudacxx/cuda/hierarchy/thread_level/native_hierarchy_query_signatures.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/hierarchy/thread_level/native_hierarchy_query_signatures.compile.pass.cpp new file mode 100644 index 000000000000..32250331ee17 --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/hierarchy/thread_level/native_hierarchy_query_signatures.compile.pass.cpp @@ -0,0 +1,102 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// todo: enable with nvrtc +// UNSUPPORTED: nvrtc + +#include +#include +#include +#include + +template +__device__ void test_query_signatures(const Level& level) +{ + // 1. Test cuda::thread_level::dims(x) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::thread_level::dims(level))>); + static_assert(noexcept(cuda::thread_level::dims(level))); + + // 2. Test cuda::thread_level::static_dims(x) signature. + static_assert(cuda::std::is_same_v, + decltype(cuda::thread_level::static_dims(level))>); + static_assert(noexcept(cuda::thread_level::static_dims(level))); + + // 3. Test cuda::thread_level::extents(x) signature. + using ExtentsRet = cuda::std::conditional_t, + cuda::std::extents, + cuda::std::dims<3, unsigned>>; + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::thread_level::extents(level))); + + // 4. Test cuda::thread_level::count(x) signature. + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::thread_level::count(level))); + + // 5. Test cuda::thread_level::index(x) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::thread_level::index(level))>); + static_assert(noexcept(cuda::thread_level::index(level))); + + // 6. Test cuda::thread_level::rank(x) signature. + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::thread_level::rank(level))); +} + +template +__device__ void test_query_as_signatures(const Level& level) +{ + // 1. Test cuda::thread_level::dims(x) signature. + static_assert(cuda::std::is_same_v, decltype(cuda::thread_level::dims_as(level))>); + static_assert(noexcept(cuda::thread_level::dims_as(level))); + + // 2. Test cuda::thread_level::extents(x) signature. + using ExtentsRet = cuda::std:: + conditional_t, cuda::std::extents, cuda::std::dims<3, T>>; + static_assert(cuda::std::is_same_v(level))>); + static_assert(noexcept(cuda::thread_level::extents_as(level))); + + // 3. Test cuda::thread_level::count(x) signature. + static_assert(cuda::std::is_same_v(level))>); + static_assert(noexcept(cuda::thread_level::count_as(level))); + + // 4. Test cuda::thread_level::index(x) signature. + static_assert( + cuda::std::is_same_v, decltype(cuda::thread_level::index_as(level))>); + static_assert(noexcept(cuda::thread_level::index_as(level))); + + // 5. Test cuda::thread_level::rank(x) signature. + static_assert(cuda::std::is_same_v(level))>); + static_assert(noexcept(cuda::thread_level::rank_as(level))); +} + +template +__device__ void test(const InLevel& in_level) +{ + test_query_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); +} + +__device__ void test() +{ + test(cuda::warp); + test(cuda::block); + test(cuda::cluster); + test(cuda::grid); +} + +int main(int, char**) +{ + return 0; +} diff --git a/libcudacxx/test/libcudacxx/cuda/hierarchy/traits/is_natively_reachable_hierarchy_level_v.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/hierarchy/traits/is_natively_reachable_hierarchy_level_v.compile.pass.cpp new file mode 100644 index 000000000000..cdfffc6d3725 --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/hierarchy/traits/is_natively_reachable_hierarchy_level_v.compile.pass.cpp @@ -0,0 +1,59 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// todo: enable with nvrtc +// UNSUPPORTED: nvrtc + +#include + +struct MyLevel : cuda::hierarchy_level_base +{}; + +template +inline constexpr bool trait_v = cuda::__is_natively_reachable_hierarchy_level_v; + +static_assert(!trait_v); +static_assert(!trait_v); +static_assert(trait_v); +static_assert(trait_v); +static_assert(trait_v); +static_assert(!trait_v); + +static_assert(!trait_v); +static_assert(!trait_v); +static_assert(trait_v); +static_assert(trait_v); +static_assert(trait_v); +static_assert(!trait_v); + +static_assert(!trait_v); +static_assert(!trait_v); +static_assert(!trait_v); +static_assert(trait_v); +static_assert(trait_v); +static_assert(!trait_v); + +static_assert(!trait_v); +static_assert(!trait_v); +static_assert(!trait_v); +static_assert(!trait_v); +static_assert(trait_v); +static_assert(!trait_v); + +static_assert(!trait_v); +static_assert(!trait_v); +static_assert(!trait_v); +static_assert(!trait_v); +static_assert(!trait_v); +static_assert(!trait_v); + +int main(int, char**) +{ + return 0; +} diff --git a/libcudacxx/test/libcudacxx/cuda/hierarchy/warp_level/native_hierarchy_queries.pass.cpp b/libcudacxx/test/libcudacxx/cuda/hierarchy/warp_level/native_hierarchy_queries.pass.cpp new file mode 100644 index 000000000000..52dbd70c3780 --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/hierarchy/warp_level/native_hierarchy_queries.pass.cpp @@ -0,0 +1,164 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// todo: enable with nvrtc +// UNSUPPORTED: nvrtc + +#include +#include +#include +#include +#include + +#include "hierarchy_queries.h" + +__device__ void test_warp() +{ + constexpr cuda::std::size_t dext = cuda::std::dynamic_extent; + + const unsigned count_in_block = (blockDim.x * blockDim.y * blockDim.z + warpSize - 1) / warpSize; + const unsigned rank_in_block = ((threadIdx.z * blockDim.y + threadIdx.y) * blockDim.x + threadIdx.x) / warpSize; + const uint3 dims_in_block{count_in_block, 1, 1}; + const uint3 index_in_block{rank_in_block, 0, 0}; + + // 1. Test cuda::warp.dims(x) + test_dims(dims_in_block, cuda::warp, cuda::block); + { + uint3 exp = dims_in_block; + NV_IF_TARGET(NV_PROVIDES_SM_90, ({ + exp.x *= __clusterDim().x; + exp.y *= __clusterDim().y; + exp.z *= __clusterDim().z; + })) + test_dims(exp, cuda::warp, cuda::cluster); + } + test_dims({count_in_block * gridDim.x, gridDim.y, gridDim.z}, cuda::warp, cuda::grid); + + // 2. Test cuda::warp.static_dims(x) + test_static_dims(ulonglong3{dext, 1, 1}, cuda::warp, cuda::block); + test_static_dims(ulonglong3{dext, dext, dext}, cuda::warp, cuda::cluster); + test_static_dims(ulonglong3{dext, dext, dext}, cuda::warp, cuda::grid); + + // 3. Test cuda::warp.extents(x) + test_extents(cuda::std::dims<1, unsigned>{count_in_block}, cuda::warp, cuda::block); + { + uint3 exp = dims_in_block; + NV_IF_TARGET(NV_PROVIDES_SM_90, ({ + exp.x *= __clusterDim().x; + exp.y *= __clusterDim().y; + exp.z *= __clusterDim().z; + })) + test_extents(cuda::std::dims<3, unsigned>{exp.x, exp.y, exp.z}, cuda::warp, cuda::cluster); + } + { + const uint3 exp{count_in_block * gridDim.x, gridDim.y, gridDim.z}; + test_extents(cuda::std::dims<3, unsigned>{exp.x, exp.y, exp.z}, cuda::warp, cuda::grid); + } + + // 4. Test cuda::warp.count(x) + test_count(count_in_block, cuda::warp, cuda::block); + { + uint3 exp = dims_in_block; + NV_IF_TARGET(NV_PROVIDES_SM_90, ({ + exp.x *= __clusterDim().x; + exp.y *= __clusterDim().y; + exp.z *= __clusterDim().z; + })) + test_count(cuda::std::size_t{exp.z} * exp.y * exp.x, cuda::warp, cuda::cluster); + } + { + const uint3 exp{count_in_block * gridDim.x, gridDim.y, gridDim.z}; + test_count(cuda::std::size_t{exp.z} * exp.y * exp.x, cuda::warp, cuda::grid); + } + + // 5. test cuda::warp.index(x) + test_index(index_in_block, cuda::warp, cuda::block); + { + uint3 exp = index_in_block; + NV_IF_TARGET(NV_PROVIDES_SM_90, ({ + exp.x += count_in_block * __clusterRelativeBlockIdx().x; + exp.y += __clusterRelativeBlockIdx().y; + exp.z += __clusterRelativeBlockIdx().z; + })) + test_index(exp, cuda::warp, cuda::cluster); + } + { + const uint3 exp{ + rank_in_block + count_in_block * blockIdx.x, + blockIdx.y, + blockIdx.z, + }; + test_index(exp, cuda::warp, cuda::grid); + } + + // 6. Test cuda::warp.rank(x) + test_rank(rank_in_block, cuda::warp, cuda::block); + { + cuda::std::size_t exp = 0; + NV_IF_ELSE_TARGET(NV_PROVIDES_SM_90, + ({ + exp = (__clusterRelativeBlockIdx().z * __clusterDim().y + __clusterRelativeBlockIdx().y) + * __clusterDim().x * count_in_block + + __clusterRelativeBlockIdx().x * count_in_block + rank_in_block; + }), + ({ exp = rank_in_block; })) + test_rank(exp, cuda::warp, cuda::cluster); + } + { + const cuda::std::size_t exp = + (blockIdx.z * gridDim.y + blockIdx.y) * gridDim.x * count_in_block + blockIdx.x * count_in_block + rank_in_block; + test_rank(exp, cuda::warp, cuda::grid); + } +} + +#if !_CCCL_COMPILER(NVRTC) +__global__ void test_kernel() +{ + test_warp(); +} + +void test() +{ + int cc_major{}; + assert(cudaDeviceGetAttribute(&cc_major, cudaDevAttrComputeCapabilityMajor, 0) == cudaSuccess); + + // thread block clusters require compute capability at least 9.0 + const bool enable_clusters = cc_major >= 9; + + test_kernel<<<1, 128>>>(); + test_kernel<<<128, 1>>>(); + test_kernel<<>>(); + test_kernel<<>>(); + if (enable_clusters) + { + cudaLaunchAttribute attribute[1]{}; + attribute[0].id = cudaLaunchAttributeClusterDimension; + attribute[0].val.clusterDim.x = 4; + attribute[0].val.clusterDim.y = 2; + attribute[0].val.clusterDim.z = 1; + + cudaLaunchConfig_t config{}; + config.gridDim = {12, 10, 3}; + config.blockDim = {2, 8, 4}; + config.attrs = attribute; + config.numAttrs = 1; + + void* pargs[1]{}; + assert(cudaLaunchKernelExC(&config, (const void*) test_kernel, pargs) == cudaSuccess); + } + + assert(cudaDeviceSynchronize() == cudaSuccess); +} +#endif // !_CCCL_COMPILER(NVRTC) + +int main(int, char**) +{ + NV_IF_ELSE_TARGET(NV_IS_HOST, (test();), (test_warp();)) + return 0; +} diff --git a/libcudacxx/test/libcudacxx/cuda/hierarchy/warp_level/native_hierarchy_query_signatures.compile.pass.cpp b/libcudacxx/test/libcudacxx/cuda/hierarchy/warp_level/native_hierarchy_query_signatures.compile.pass.cpp new file mode 100644 index 000000000000..175e73cb9153 --- /dev/null +++ b/libcudacxx/test/libcudacxx/cuda/hierarchy/warp_level/native_hierarchy_query_signatures.compile.pass.cpp @@ -0,0 +1,95 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +// todo: enable with nvrtc +// UNSUPPORTED: nvrtc + +#include +#include +#include +#include + +template +__device__ void test_query_signatures(const Level& level) +{ + // 1. Test cuda::warp_level::dims(x) signature. + static_assert(cuda::std::is_same_v, decltype(cuda::warp_level::dims(level))>); + static_assert(noexcept(cuda::warp_level::dims(level))); + + // 2. Test cuda::warp_level::static_dims(x) signature. + static_assert(cuda::std::is_same_v, + decltype(cuda::warp_level::static_dims(level))>); + static_assert(noexcept(cuda::warp_level::static_dims(level))); + + // 3. Test cuda::warp_level::extents(x) signature. + using ExtentsRet = cuda::std::dims<(cuda::std::is_same_v) ? 1 : 3, unsigned>; + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::warp_level::extents(level))); + + // 4. Test cuda::warp_level::count(x) signature. + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::warp_level::count(level))); + + // 5. Test cuda::warp_level::index(x) signature. + static_assert(cuda::std::is_same_v, decltype(cuda::warp_level::index(level))>); + static_assert(noexcept(cuda::warp_level::index(level))); + + // 6. Test cuda::warp_level::rank(x) signature. + static_assert(cuda::std::is_same_v); + static_assert(noexcept(cuda::warp_level::rank(level))); +} + +template +__device__ void test_query_as_signatures(const Level& level) +{ + // 1. Test cuda::warp_level::dims(x) signature. + static_assert(cuda::std::is_same_v, decltype(cuda::warp_level::dims_as(level))>); + static_assert(noexcept(cuda::warp_level::dims_as(level))); + + // 2. Test cuda::warp_level::extents(x) signature. + using ExtentsRet = cuda::std::dims<(cuda::std::is_same_v) ? 1 : 3, T>; + static_assert(cuda::std::is_same_v(level))>); + static_assert(noexcept(cuda::warp_level::extents_as(level))); + + // 3. Test cuda::warp_level::count(x) signature. + static_assert(cuda::std::is_same_v(level))>); + static_assert(noexcept(cuda::warp_level::count_as(level))); + + // 4. Test cuda::warp_level::index(x) signature. + static_assert(cuda::std::is_same_v, decltype(cuda::warp_level::index_as(level))>); + static_assert(noexcept(cuda::warp_level::index_as(level))); + + // 5. Test cuda::warp_level::rank(x) signature. + static_assert(cuda::std::is_same_v(level))>); + static_assert(noexcept(cuda::warp_level::rank_as(level))); +} + +template +__device__ void test(const InLevel& in_level) +{ + test_query_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); + test_query_as_signatures(in_level); +} + +__device__ void test() +{ + test(cuda::block); + test(cuda::cluster); + test(cuda::grid); +} + +int main(int, char**) +{ + return 0; +} diff --git a/libcudacxx/test/support/hierarchy_queries.h b/libcudacxx/test/support/hierarchy_queries.h new file mode 100644 index 000000000000..27ebd65e2c15 --- /dev/null +++ b/libcudacxx/test/support/hierarchy_queries.h @@ -0,0 +1,115 @@ +//===----------------------------------------------------------------------===// +// +// Part of the libcu++ Project, under the Apache License v2.0 with LLVM Exceptions. +// See https://llvm.org/LICENSE.txt for license information. +// SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception +// SPDX-FileCopyrightText: Copyright (c) 2025 NVIDIA CORPORATION & AFFILIATES. +// +//===----------------------------------------------------------------------===// + +#ifndef SUPPORT_HIERARCHY_QUERIES_H +#define SUPPORT_HIERARCHY_QUERIES_H + +#include +#include +#include +#include + +template +__device__ void test_result(cuda::hierarchy_query_result res, Vec exp) +{ + assert(res.x == static_cast(exp.x)); + assert(res.y == static_cast(exp.y)); + assert(res.z == static_cast(exp.z)); +} + +template +__device__ void test_result(cuda::std::extents res, cuda::std::extents exp) +{ + for (cuda::std::size_t i = 0; i < sizeof...(Exts); ++i) + { + assert(res.extent(i) == static_cast(exp.extent(i))); + } +} + +template +__device__ void test_dims(const uint3 exp, const Level& level, Args... args) +{ + test_result(level.dims(args...), exp); + test_result(level.template dims_as(args...), exp); + test_result(level.template dims_as(args...), exp); + test_result(level.template dims_as(args...), exp); + test_result(level.template dims_as(args...), exp); + test_result(level.template dims_as(args...), exp); + test_result(level.template dims_as(args...), exp); +} + +template +__device__ void test_static_dims(const ulonglong3 exp, Level level, Args... args) +{ + static_assert(level.static_dims(args...).x != 0); + test_result(level.static_dims(args...), exp); +} + +template +__device__ void test_extents(const Exp exp, const Level& level, Args... args) +{ + test_result(level.extents(args...), exp); + test_result(level.template extents_as(args...), exp); + test_result(level.template extents_as(args...), exp); + test_result(level.template extents_as(args...), exp); + test_result(level.template extents_as(args...), exp); + test_result(level.template extents_as(args...), exp); + test_result(level.template extents_as(args...), exp); +} + +template +__device__ void test_count(const cuda::std::size_t exp, const Level& level, Args... args) +{ + assert(level.count(args...) == exp); + assert(level.template count_as(args...) == static_cast(exp)); + assert(level.template count_as(args...) == static_cast(exp)); + assert(level.template count_as(args...) == static_cast(exp)); + assert(level.template count_as(args...) == static_cast(exp)); + assert(level.template count_as(args...) == static_cast(exp)); + assert(level.template count_as(args...) == static_cast(exp)); +} + +template +__device__ void test_index(const uint3 exp, const Level& level, Args... args) +{ + test_result(level.index(args...), exp); + test_result(level.template index_as(args...), exp); + test_result(level.template index_as(args...), exp); + test_result(level.template index_as(args...), exp); + test_result(level.template index_as(args...), exp); + test_result(level.template index_as(args...), exp); + test_result(level.template index_as(args...), exp); +} + +template +__device__ void test_rank(const cuda::std::size_t exp, const Level& level, Args... args) +{ + assert(level.rank(args...) == exp); + assert(level.template rank_as(args...) == static_cast(exp)); + assert(level.template rank_as(args...) == static_cast(exp)); + assert(level.template rank_as(args...) == static_cast(exp)); + assert(level.template rank_as(args...) == static_cast(exp)); + assert(level.template rank_as(args...) == static_cast(exp)); + assert(level.template rank_as(args...) == static_cast(exp)); +} + +template +__device__ constexpr cuda::std::size_t mul_static_extents(Args... args) +{ + if (((args == cuda::std::dynamic_extent) || ...)) + { + return cuda::std::dynamic_extent; + } + else + { + return (cuda::std::size_t{1} * ... * args); + } +} + +#endif // SUPPORT_HIERARCHY_QUERIES_H