add some simple utilities for manipulating lists of types

cub/test/CMakeLists.txt

@@ -21,11 +21,6 @@ endif()
 include("${CUB_SOURCE_DIR}/cmake/CPM.cmake")
 CPMAddPackage("gh:catchorg/[email protected]")
 
-option(METAL_BUILD_DOC OFF)
-option(METAL_BUILD_EXAMPLES OFF)
-option(METAL_BUILD_TESTS OFF)
-CPMAddPackage("gh:brunocodutra/[email protected]")
-
 CPMAddPackage(
   NAME NVTX
   GITHUB_REPOSITORY NVIDIA/NVTX
@@ -224,7 +219,6 @@ function(cub_add_test target_name_var test_name test_src cub_target launcher_id)
         target_link_libraries(${config_c2run_target} PRIVATE
           ${cub_target}
           ${config_c2h_target}
-          Metal
           Catch2::Catch2)
         cub_clone_target_properties(${config_c2run_target} ${cub_target})
         cub_configure_cuda_target(${config_c2run_target} RDC ${cdp_val})
@@ -264,7 +258,6 @@ function(cub_add_test target_name_var test_name test_src cub_target launcher_id)
     target_link_libraries(${test_target} PRIVATE
       ${cub_target}
       ${config_c2h_target}
-      Metal
       Catch2::Catch2
     )
     cub_clone_target_properties(${test_target} ${cub_target})

cub/test/catch2_test_device_histogram.cu

@@ -430,7 +430,7 @@ CUB_TEST("DeviceHistogram::Histogram* basic use", "[histogram][device]", types)
 
 // TODO(bgruber): float produces INFs in the HistogramRange test setup AND the HistogramEven implementation
 // This test covers int32 and int64 arithmetic for bin computation
-CUB_TEST("DeviceHistogram::Histogram* large levels", "[histogram][device]", metal::remove<types, float>)
+CUB_TEST("DeviceHistogram::Histogram* large levels", "[histogram][device]", c2h::remove<types, float>)
 {
   using sample_t             = c2h::get<0, TestType>;
   using level_t              = sample_t;

cub/test/catch2_test_helper.h

@@ -36,7 +36,6 @@
 
 #include "cuda/std/__cccl/diagnostic.h"
 #include "test_util_vec.h"
-#include <metal.hpp>
 
 #if __CUDACC_VER_MAJOR__ == 11
 _CCCL_NV_DIAG_SUPPRESS(177) // catch2 may contain unused variableds
@@ -60,31 +59,34 @@ namespace c2h
 {
 
 template <typename... Ts>
-using type_list = metal::list<Ts...>;
+using type_list = ::cuda::std::__type_list<Ts...>;
 
 template <typename TypeList>
-using size = metal::size<TypeList>;
+using size = ::cuda::std::__type_list_size<TypeList>;
 
 template <std::size_t Index, typename TypeList>
-using get = metal::at<TypeList, metal::number<Index>>;
+using get = ::cuda::std::__type_at_c<Index, TypeList>;
 
 template <class... TypeLists>
-using cartesian_product = metal::cartesian<TypeLists...>;
+using cartesian_product = ::cuda::std::__type_cartesian_product<TypeLists...>;
 
 template <typename T, T... Ts>
 using enum_type_list = c2h::type_list<std::integral_constant<T, Ts>...>;
 
 template <typename T0, typename T1>
-using pair = metal::pair<T0, T1>;
+using pair = ::cuda::std::__type_pair<T0, T1>;
 
 template <typename P>
-using first = metal::first<P>;
+using first = ::cuda::std::__type_pair_first<P>;
 
 template <typename P>
-using second = metal::second<P>;
+using second = ::cuda::std::__type_pair_second<P>;
 
 template <std::size_t start, std::size_t size, std::size_t stride = 1>
-using iota = metal::iota<metal::number<start>, metal::number<size>, metal::number<stride>>;
+using iota = ::cuda::std::__type_iota<start, size, stride>;
+
+template <typename TypeList, typename T>
+using remove = ::cuda::std::__type_remove<TypeList, T>;
 
 } // namespace c2h
 

cub/test/catch2_test_iterator.cu

@@ -43,7 +43,7 @@
 
 using scalar_types = c2h::type_list<std::int8_t, std::int16_t, std::int32_t, std::int64_t, float, double>;
 
-using types = metal::append<
+using types = ::cuda::std::__type_push_back<
   scalar_types,
   char2,
   short2,

cudax/samples/vector_add/vector.cuh

@@ -59,7 +59,7 @@ public:
   }
 
 private:
-  void sync_host_to_device(stream_ref __str, detail::__param_kind __p) const
+  void sync_host_to_device(::cuda::stream_ref __str, detail::__param_kind __p) const
   {
     if (__dirty_)
     {
@@ -78,7 +78,7 @@ private:
     }
   }
 
-  void sync_device_to_host(stream_ref __str, detail::__param_kind __p) const
+  void sync_device_to_host(::cuda::stream_ref __str, detail::__param_kind __p) const
   {
     if (__p != detail::__param_kind::_in)
     {
@@ -94,7 +94,7 @@ private:
     using __cv_vector = ::cuda::std::__maybe_const<_Kind == detail::__param_kind::_in, vector>;
 
   public:
-    explicit __action(stream_ref __str, __cv_vector& __v) noexcept
+    explicit __action(::cuda::stream_ref __str, __cv_vector& __v) noexcept
         : __str_(__str)
         , __v_(__v)
     {
@@ -116,25 +116,25 @@ private:
     }
 
   private:
-    stream_ref __str_;
+    ::cuda::stream_ref __str_;
     __cv_vector& __v_;
   };
 
   _CCCL_NODISCARD_FRIEND __action<detail::__param_kind::_inout>
-  __cudax_launch_transform(stream_ref __str, vector& __v) noexcept
+  __cudax_launch_transform(::cuda::stream_ref __str, vector& __v) noexcept
   {
     return __action<detail::__param_kind::_inout>{__str, __v};
   }
 
   _CCCL_NODISCARD_FRIEND __action<detail::__param_kind::_in>
-  __cudax_launch_transform(stream_ref __str, const vector& __v) noexcept
+  __cudax_launch_transform(::cuda::stream_ref __str, const vector& __v) noexcept
   {
     return __action<detail::__param_kind::_in>{__str, __v};
   }
 
   template <detail::__param_kind _Kind>
   _CCCL_NODISCARD_FRIEND __action<_Kind>
-  __cudax_launch_transform(stream_ref __str, detail::__box<vector, _Kind> __b) noexcept
+  __cudax_launch_transform(::cuda::stream_ref __str, detail::__box<vector, _Kind> __b) noexcept
   {
     return __action<_Kind>{__str, __b.__val};
   }

libcudacxx/include/cuda/std/__mdspan/static_array.h

@@ -59,7 +59,7 @@
 #include <cuda/std/__mdspan/macros.h>
 #include <cuda/std/__mdspan/maybe_static_value.h>
 #include <cuda/std/__mdspan/standard_layout_static_array.h>
-#include <cuda/std/__mdspan/type_list.h>
+#include <cuda/std/__type_traits/type_list.h>
 #include <cuda/std/__utility/integer_sequence.h>
 #include <cuda/std/array>
 #include <cuda/std/cstddef>
@@ -123,9 +123,8 @@ class __partially_static_array_impl<
 {
 private:
   template <size_t _Np>
-  using __base_n = typename __type_at<
-    _Np,
-    __type_list<__maybe_static_value<_Tp, _static_t, __values_or_sentinals, __sentinal, _Idxs>...>>::type;
+  using __base_n =
+    __type_index_c<_Np, __maybe_static_value<_Tp, _static_t, __values_or_sentinals, __sentinal, _Idxs>...>;
 
 public:
   static constexpr auto __size = sizeof...(_Idxs);