Rework hierarchy levels

cudax/examples/simple_p2p.cu

@@ -52,7 +52,7 @@ struct simple_kernel
   __device__ void operator()(Configuration config, ::cuda::std::span<const float> src, ::cuda::std::span<float> 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;
   }
 };

cudax/include/cuda/experimental/__execution/queries.cuh

@@ -26,6 +26,7 @@ _CCCL_SUPPRESS_DEPRECATED_PUSH
 _CCCL_SUPPRESS_DEPRECATED_POP
 
 #include <cuda/__launch/configuration.h>
+#include <cuda/hierarchy>
 #include <cuda/std/__concepts/concept_macros.h>
 #include <cuda/std/__concepts/convertible_to.h>
 #include <cuda/std/__execution/env.h>

cudax/include/cuda/experimental/__execution/stream/adaptor.cuh

@@ -23,6 +23,7 @@
 
 #include <cuda/__launch/configuration.h>
 #include <cuda/__launch/launch.h>
+#include <cuda/hierarchy>
 #include <cuda/std/__concepts/concept_macros.h>
 #include <cuda/std/__memory/unique_ptr.h>
 #include <cuda/std/__type_traits/remove_cvref.h>
@@ -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

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.

cudax/include/cuda/experimental/__launch/launch.cuh

@@ -22,6 +22,7 @@
 #endif // no system header
 
 #include <cuda/__driver/driver_api.h>
+#include <cuda/__hierarchy/traits.h>
 #include <cuda/__launch/configuration.h>
 #include <cuda/__launch/launch.h>
 #include <cuda/__stream/launch_transform.h>
@@ -108,7 +109,7 @@ _CCCL_HOST_API auto __launch_impl(_Dst&& __dst, _Config __conf, ::CUfunction __k
   static_assert(!::cuda::std::is_same_v<decltype(__conf.dims), no_init_t>,
                 "Can't launch a configuration without hierarchy dimensions");
   ::CUlaunchConfig __config{};
-  constexpr bool __has_cluster_level = has_level<cluster_level, decltype(__conf.dims)>;
+  constexpr bool __has_cluster_level = has_level_v<cluster_level, decltype(__conf.dims)>;
   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<unsigned>(__conf.dims.extents(block, grid).x);
   __config.gridDimY  = static_cast<unsigned>(__conf.dims.extents(block, grid).y);
   __config.gridDimZ  = static_cast<unsigned>(__conf.dims.extents(block, grid).z);
-  __config.blockDimX = static_cast<unsigned>(__conf.dims.extents(thread, block).x);
-  __config.blockDimY = static_cast<unsigned>(__conf.dims.extents(thread, block).y);
-  __config.blockDimZ = static_cast<unsigned>(__conf.dims.extents(thread, block).z);
+  __config.blockDimX = static_cast<unsigned>(__conf.dims.extents(gpu_thread, block).x);
+  __config.blockDimY = static_cast<unsigned>(__conf.dims.extents(gpu_thread, block).y);
+  __config.blockDimZ = static_cast<unsigned>(__conf.dims.extents(gpu_thread, block).z);
 
   if constexpr (__has_cluster_level)
   {

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<cuda::cluster_level, decltype(dims)>)
+    if constexpr (cuda::has_level_v<cuda::cluster_level, decltype(dims)>)
     {
       dim3 cluster_dims                            = dims.extents(cuda::block, cuda::cluster);
       config.attrs[config.numAttrs].id             = cudaLaunchAttributeClusterDimension;

cudax/test/launch/launch_smoke.cu

@@ -55,7 +55,7 @@ struct functor_taking_config
   template <typename Config>
   __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();
   };

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 <cuda/std/detail/__config>
+
+#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 <cuda/std/__cstddef/types.h>
+#include <cuda/std/__type_traits/is_base_of.h>
+
+#include <cuda/std/__cccl/prologue.h>
+
+_CCCL_BEGIN_NAMESPACE_CUDA
+
+// hierarchy level
+
+template <class _Level>
+struct hierarchy_level_base;
+
+template <class _Level>
+struct __native_hierarchy_level_base;
+
+struct grid_level;
+struct cluster_level;
+struct block_level;
+struct warp_level;
+struct thread_level;
+
+template <class _Tp>
+inline constexpr bool __is_hierarchy_level_v = ::cuda::std::is_base_of_v<hierarchy_level_base<_Tp>, _Tp>;
+
+template <class _Tp>
+inline constexpr bool __is_native_hierarchy_level_v =
+  ::cuda::std::is_base_of_v<__native_hierarchy_level_base<_Tp>, _Tp>;
+
+// hierarchy
+
+template <class _BottomUnit, class... _Levels>
+struct hierarchy_dimensions;
+
+template <class _Tp>
+inline constexpr bool __is_hierarchy_v = false;
+template <class _BottomUnit, class... _Levels>
+inline constexpr bool __is_hierarchy_v<hierarchy_dimensions<_BottomUnit, _Levels...>> = true;
+
+template <typename... _Levels>
+struct allowed_levels;
+
+_CCCL_END_NAMESPACE_CUDA
+
+#include <cuda/std/__cccl/epilogue.h>
+
+#endif // _CUDA___FWD_HIERARCHY_H

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 <cuda/std/detail/__config>
+
+#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 <cuda/__fwd/hierarchy.h>
+#include <cuda/__hierarchy/hierarchy_query_result.h>
+#include <cuda/__hierarchy/native_hierarchy_level_base.h>
+#include <cuda/std/__concepts/concept_macros.h>
+#include <cuda/std/__mdspan/extents.h>
+#include <cuda/std/__type_traits/is_integer.h>
+
+#include <cuda/std/__cccl/prologue.h>
+
+_CCCL_BEGIN_NAMESPACE_CUDA
+
+struct block_level : __native_hierarchy_level_base<block_level>
+{
+  using product_type  = unsigned;
+  using allowed_above = allowed_levels<grid_level, cluster_level>;
+  using allowed_below = allowed_levels<thread_level>;
+
+  using __next_native_level = cluster_level;
+
+  using __base_type = __native_hierarchy_level_base<block_level>;
+  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<block_level, _Hierarchy>, "_Hierarchy doesn't contain block level");
+    static_assert(has_level_v<grid_level, _Hierarchy>, "_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 <cuda/std/__cccl/epilogue.h>
+
+#endif // _CUDA___HIERARCHY_BLOCK_LEVEL_H