Integrate CUDASTF -> CudaX

.github/ISSUE_TEMPLATE/feature_request.yml

@@ -21,6 +21,9 @@ body:
         - Thrust
         - CUB
         - libcu++
+        - CUDA Experimental (cudax)
+        - cuda.cooperative (Python)
+        - cuda.parallel (Python)
         - General CCCL
         - Infrastructure
         - Not sure

CMakePresets.json

@@ -62,18 +62,26 @@
         "THRUST_MULTICONFIG_ENABLE_SYSTEM_TBB": true,
         "cudax_ENABLE_HEADER_TESTING": true,
         "cudax_ENABLE_TESTING": true,
+        "cudax_ENABLE_EXAMPLES": true,
+        "cudax_ENABLE_CUDASTF_BOUNDSCHECK": false,
+        "cudax_ENABLE_CUDASTF_DEBUG": false,
+        "cudax_ENABLE_CUDASTF_MATHLIBS": false,
         "cudax_ENABLE_DIALECT_CPP17": true,
-        "cudax_ENABLE_DIALECT_CPP20": true
+        "cudax_ENABLE_DIALECT_CPP20": true,
+        "CCCL_C_Parallel_ENABLE_TESTING": true,
+        "CCCL_C_Parallel_ENABLE_HEADER_TESTING": true
       }
     },
     {
       "name": "all-dev-debug",
       "displayName": "all-dev debug",
       "inherits": "all-dev",
       "cacheVariables": {
-        "CCCL_ENABLE_BENCHMARKS": false,
         "CMAKE_BUILD_TYPE": "Debug",
-        "CMAKE_CUDA_FLAGS": "-G"
+        "CMAKE_CUDA_FLAGS": "-G",
+        "CCCL_ENABLE_BENCHMARKS": false,
+        "cudax_ENABLE_CUDASTF_BOUNDSCHECK": true,
+        "cudax_ENABLE_CUDASTF_DEBUG": true
       }
     },
     {
@@ -295,6 +303,9 @@
         "CCCL_ENABLE_CUDAX": true,
         "cudax_ENABLE_HEADER_TESTING": true,
         "cudax_ENABLE_TESTING": true,
+        "cudax_ENABLE_EXAMPLES": true,
+        "cudax_ENABLE_CUDASTF_BOUNDSCHECK": false,
+        "cudax_ENABLE_CUDASTF_MATHLIBS": false,
         "cudax_ENABLE_DIALECT_CPP17": false,
         "cudax_ENABLE_DIALECT_CPP20": false
       }

README.md

@@ -177,7 +177,7 @@ conda install -c conda-forge cccl
 CCCL uses [CMake](https://cmake.org/) for all build and installation infrastructure, including tests as well as targets to link against in other CMake projects.
 Therefore, CMake is the recommended way to integrate CCCL into another project.
 
-For a complete example of how to do this using CMake Package Manager see [our example project](examples/example_project).
+For a complete example of how to do this using CMake Package Manager see [our basic example project](examples/basic).
 
 Other build systems should work, but only CMake is tested.
 Contributions to simplify integrating CCCL into other build systems are welcome.

c/parallel/src/reduce.cu

@@ -394,7 +394,7 @@ extern "C" CCCL_C_API CUresult cccl_device_reduce_build(
         : std::format(
             "extern \"C\" __device__ {3} {4}(const void *self_ptr);\n"
             "extern \"C\" __device__ void {5}(void *self_ptr, {0} offset);\n"
-            "struct __align__({2}) input_iterator_state_t {{\n;"
+            "struct __align__({2}) input_iterator_state_t {{\n"
             "  using iterator_category = cuda::std::random_access_iterator_tag;\n"
             "  using value_type = {3};\n"
             "  using difference_type = {0};\n"

c/parallel/test/CMakeLists.txt

@@ -20,7 +20,7 @@ function(cccl_c_parallel_add_test target_name_var source)
 
   target_compile_definitions(${target_name} PRIVATE
     TEST_CUB_PATH="-I${CCCL_SOURCE_DIR}/cub"
-    TEST_THRUST_PATH="-I${CCCL_SOURCE_DIR}/cub"
+    TEST_THRUST_PATH="-I${CCCL_SOURCE_DIR}/thrust"
     TEST_LIBCUDACXX_PATH="-I${CCCL_SOURCE_DIR}/libcudacxx/include"
     TEST_CTK_PATH="-I${CUDAToolkit_INCLUDE_DIRS}"
   )

c/parallel/test/c2h.h

@@ -101,7 +101,7 @@ static std::string compile(const std::string& source)
   REQUIRE(NVRTC_SUCCESS == nvrtcGetLTOIR(prog, ltoir.get()));
   REQUIRE(NVRTC_SUCCESS == nvrtcDestroyProgram(&prog));
 
-  return std::string(ltoir.release(), ltoir_size);
+  return std::string(ltoir.get(), ltoir_size);
 }
 
 template <class T>

ci/build_common.sh

@@ -34,7 +34,7 @@ function usage {
     echo "  $ PARALLEL_LEVEL=8 $0"
     echo "  $ PARALLEL_LEVEL=8 $0 -cxx g++-9"
     echo "  $ $0 -cxx clang++-8"
-    echo "  $ $0 -configure -arch=80"
+    echo "  $ $0 -configure -arch 80"
     echo "  $ $0 -cxx g++-8 -std 14 -arch 80-real -v -cuda /usr/local/bin/nvcc"
     echo "  $ $0 -cmake-options \"-DCMAKE_BUILD_TYPE=Debug -DCMAKE_CXX_FLAGS=-Wfatal-errors\""
     exit 1

ci/matrix.yaml

@@ -8,6 +8,18 @@ workflows:
   #   - {jobs: ['test'], project: 'thrust', std: 17, ctk: 'curr', cxx: ['gcc12', 'llvm16']}
   #
   override:
+    - {jobs: ['build'], project: 'cudax', ctk: ['12.0', 'curr'], std: 'all', cxx: ['gcc9', 'gcc10', 'gcc11']}
+    - {jobs: ['build'], project: 'cudax', ctk: ['12.0', 'curr'], std: 'all', cxx: ['clang9', 'clang10', 'clang11', 'clang12', 'clang13']}
+    - {jobs: ['build'], project: 'cudax', ctk: [        'curr'], std: 'all', cxx: ['clang14', 'clang15', 'clang16', 'clang17']}
+    - {jobs: ['build'], project: 'cudax', ctk: ['12.0',       ], std: 20,    cxx: ['msvc14.36']}
+    - {jobs: ['build'], project: 'cudax', ctk: [        'curr'], std: 20,    cxx: ['msvc2022']}
+    - {jobs: ['build'], project: 'cudax', ctk: ['12.0'        ], std: 17,    cxx: ['gcc12'], sm: "90"}
+    - {jobs: ['build'], project: 'cudax', ctk: [        'curr'], std: 17,    cxx: ['gcc13'], sm: "90a"}
+    - {jobs: ['build'], project: 'cudax', ctk: [        'curr'], std: 'all', cxx: ['gcc13', 'clang16'], cpu: 'arm64'}
+    - {jobs: ['test'],  project: 'cudax', ctk: ['12.0'        ], std: 'min', cxx: ['gcc12']}
+    - {jobs: ['test'],  project: 'cudax', ctk: [        'curr'], std: 'all', cxx: ['gcc12']}
+    - {jobs: ['test'],  project: 'cudax', ctk: ['12.0'        ], std: 'max', cxx: ['clang14']}
+    - {jobs: ['test'],  project: 'cudax', ctk: [        'curr'], std: 'max', cxx: ['clang18']}
 
   pull_request:
     # Old CTK

cmake/CCCLConfigureTarget.cmake

@@ -10,6 +10,13 @@ function(cccl_configure_target target_name)
 
   get_target_property(type ${target_name} TYPE)
 
+  set_target_properties(${target_name}
+    PROPERTIES
+      # Disable compiler extensions:
+      CXX_EXTENSIONS OFF
+      CUDA_EXTENSIONS OFF
+  )
+
   if (DEFINED CCT_DIALECT)
     set_target_properties(${target_name}
       PROPERTIES

cub/cub/device/dispatch/dispatch_batch_memcpy.cuh

@@ -455,7 +455,7 @@ struct DispatchBatchMemcpy : SelectedPolicy
     // The number of thread blocks (or tiles) required to process all of the given buffers
     BlockOffsetT num_tiles = ::cuda::ceil_div(num_buffers, TILE_SIZE);
 
-    using BlevBufferSrcsOutT          = ::cuda::std::_If<IsMemcpy, void*, cub::detail::value_t<InputBufferIt>>;
+    using BlevBufferSrcsOutT          = ::cuda::std::_If<IsMemcpy, const void*, cub::detail::value_t<InputBufferIt>>;
     using BlevBufferDstOutT           = ::cuda::std::_If<IsMemcpy, void*, cub::detail::value_t<OutputBufferIt>>;
     using BlevBufferSrcsOutItT        = BlevBufferSrcsOutT*;
     using BlevBufferDstsOutItT        = BlevBufferDstOutT*;

cub/cub/device/dispatch/dispatch_merge.cuh

@@ -25,8 +25,8 @@ namespace detail
 {
 namespace merge
 {
-_LIBCUDACXX_INLINE_VAR constexpr int fallback_BLOCK_THREADS    = 64;
-_LIBCUDACXX_INLINE_VAR constexpr int fallback_ITEMS_PER_THREAD = 1;
+_CCCL_INLINE_VAR constexpr int fallback_BLOCK_THREADS    = 64;
+_CCCL_INLINE_VAR constexpr int fallback_ITEMS_PER_THREAD = 1;
 
 template <typename DefaultPolicy, class... Args>
 class choose_merge_agent

cub/cub/device/dispatch/dispatch_select_if.cuh

@@ -237,13 +237,6 @@ struct agent_select_if_wrapper_t
 /******************************************************************************
  * Kernel entry points
  *****************************************************************************/
-// TODO (elstehle) gird-private constants were introduced in CTK 11.7. The macro is temporarily placed here, do we want
-// to make this a CCCL macro?
-#if defined(_CCCL_CUDACC_BELOW_11_7) || (CUB_PTX_ARCH < 700)
-#  define _CUB_GRID_CONSTANT
-#else
-#  define _CUB_GRID_CONSTANT __grid_constant__
-#endif
 
 /**
  * Select kernel entry point (multi-block)
@@ -356,7 +349,7 @@ __launch_bounds__(int(
     EqualityOpT equality_op,
     OffsetT num_items,
     int num_tiles,
-    _CUB_GRID_CONSTANT const StreamingContextT streaming_context,
+    _CCCL_GRID_CONSTANT const StreamingContextT streaming_context,
     cub::detail::vsmem_t vsmem)
 {
   using VsmemHelperT = cub::detail::vsmem_helper_default_fallback_policy_t<

cub/cub/iterator/arg_index_input_iterator.cuh

@@ -64,8 +64,8 @@ CUB_NAMESPACE_BEGIN
  *        indices (forming \p KeyValuePair tuples).
  *
  * @par Overview
- * - ArgIndexInputIteratorTwraps a random access input iterator @p itr of type @p InputIteratorT.
- *   Dereferencing an ArgIndexInputIteratorTat offset @p i produces a @p KeyValuePair value whose
+ * - ArgIndexInputIterator wraps a random access input iterator @p itr of type @p InputIteratorT.
+ *   Dereferencing an ArgIndexInputIterator at offset @p i produces a @p KeyValuePair value whose
  *   @p key field is @p i and whose @p value field is <tt>itr[i]</tt>.
  * - Can be used with any data type.
  * - Can be constructed, manipulated, and exchanged within and between host and device
@@ -74,7 +74,7 @@ CUB_NAMESPACE_BEGIN
  * - Compatible with Thrust API v1.7 or newer.
  *
  * @par Snippet
- * The code snippet below illustrates the use of @p ArgIndexInputIteratorTto
+ * The code snippet below illustrates the use of @p ArgIndexInputIterator to
  * dereference an array of doubles
  * @par
  * @code
@@ -87,17 +87,16 @@ CUB_NAMESPACE_BEGIN
  * cub::ArgIndexInputIterator<double*> itr(d_in);
  *
  * // Within device code:
- * using Tuple = typename cub::ArgIndexInputIterator<double*>::value_type;
- * Tuple item_offset_pair.key = *itr;
- * printf("%f @ %d\n",
- *   item_offset_pair.value,
- *   item_offset_pair.key);   // 8.0 @ 0
+ * cub::ArgIndexInputIterator<double*>::value_type tup = *itr;
+ * printf("%f @ %ld\n",
+ *   tup.value,
+ *   tup.key);   // 8.0 @ 0
  *
  * itr = itr + 6;
- * item_offset_pair.key = *itr;
- * printf("%f @ %d\n",
- *   item_offset_pair.value,
- *   item_offset_pair.key);   // 9.0 @ 6
+ * tup = *itr;
+ * printf("%f @ %ld\n",
+ *   tup.value,
+ *   tup.key);   // 9.0 @ 6
  *
  * @endcode
  *

cub/cub/iterator/constant_input_iterator.cuh

@@ -61,15 +61,15 @@ CUB_NAMESPACE_BEGIN
  * @brief A random-access input generator for dereferencing a sequence of homogeneous values
  *
  * @par Overview
- * - Read references to a ConstantInputIteratorTiterator always return the supplied constant
+ * - Read references to a ConstantInputIterator always return the supplied constant
  *   of type @p ValueType.
  * - Can be used with any data type.
  * - Can be constructed, manipulated, dereferenced, and exchanged within and between host and device
  *   functions.
  * - Compatible with Thrust API v1.7 or newer.
  *
  * @par Snippet
- * The code snippet below illustrates the use of @p ConstantInputIteratorTto
+ * The code snippet below illustrates the use of @p ConstantInputIterator to
  * dereference a sequence of homogeneous doubles.
  * @par
  * @code

cub/cub/iterator/counting_input_iterator.cuh

@@ -67,14 +67,14 @@ CUB_NAMESPACE_BEGIN
  * @brief A random-access input generator for dereferencing a sequence of incrementing integer values.
  *
  * @par Overview
- * - After initializing a CountingInputIteratorTto a certain integer @p base, read references
+ * - After initializing a CountingInputIterator to a certain integer @p base, read references
  *   at @p offset will return the value @p base + @p offset.
  * - Can be constructed, manipulated, dereferenced, and exchanged within and between host and device
  *   functions.
  * - Compatible with Thrust API v1.7 or newer.
  *
  * @par Snippet
- * The code snippet below illustrates the use of @p CountingInputIteratorTto
+ * The code snippet below illustrates the use of @p CountingInputIterator to
  * dereference a sequence of incrementing integers.
  * @par
  * @code

cub/cub/iterator/transform_input_iterator.cuh

@@ -61,7 +61,7 @@ CUB_NAMESPACE_BEGIN
  * @brief A random-access input wrapper for transforming dereferenced values.
  *
  * @par Overview
- * - TransformInputIteratorTwraps a unary conversion functor of type
+ * - TransformInputIterator wraps a unary conversion functor of type
  *   @p ConversionOp and a random-access input iterator of type <tt>InputIteratorT</tt>,
  *   using the former to produce references of type @p ValueType from the latter.
  * - Can be used with any data type.
@@ -71,7 +71,7 @@ CUB_NAMESPACE_BEGIN
  * - Compatible with Thrust API v1.7 or newer.
  *
  * @par Snippet
- * The code snippet below illustrates the use of @p TransformInputIteratorTto
+ * The code snippet below illustrates the use of @p TransformInputIterator to
  * dereference an array of integers, tripling the values and converting them to doubles.
  * @par
  * @code

cub/cub/util_type.cuh

@@ -458,7 +458,7 @@ struct CubVector
 };
 
 /// The maximum number of elements in CUDA vector types
-_LIBCUDACXX_INLINE_VAR constexpr int MAX_VEC_ELEMENTS = 4;
+_CCCL_INLINE_VAR constexpr int MAX_VEC_ELEMENTS = 4;
 
 /**
  * Generic vector-1 type

cub/test/catch2_test_launch_helper.h

@@ -91,7 +91,7 @@
 
 #define DECLARE_LAUNCH_WRAPPER(API, WRAPPED_API_NAME)                                                  \
   DECLARE_INVOCABLE(API, WRAPPED_API_NAME, , );                                                        \
-  _LIBCUDACXX_INLINE_VAR constexpr struct WRAPPED_API_NAME##_t                                         \
+  _CCCL_INLINE_VAR constexpr struct WRAPPED_API_NAME##_t                                               \
   {                                                                                                    \
     template <class... As>                                                                             \
     void operator()(As... args) const                                                                  \

cub/test/test_device_batch_memcpy.cu

@@ -216,7 +216,8 @@ void RunTest(BufferOffsetT num_buffers,
 {
   try
   {
-    using SrcPtrT = uint8_t*;
+    using SrcPtrT = const uint8_t*;
+    using DstPtrT = uint8_t*;
 
     // Buffer segment data (their offsets and sizes)
     c2h::host_vector<BufferSizeT> h_buffer_sizes(num_buffers);
@@ -286,8 +287,8 @@ void RunTest(BufferOffsetT num_buffers,
       thrust::raw_pointer_cast(d_buffer_src_offsets.data()), src_transform_op);
 
     // Prepare d_buffer_dsts
-    OffsetToPtrOp<SrcPtrT> dst_transform_op{static_cast<SrcPtrT>(thrust::raw_pointer_cast(d_out.data()))};
-    cub::TransformInputIterator<SrcPtrT, OffsetToPtrOp<SrcPtrT>, ByteOffsetT*> d_buffer_dsts(
+    OffsetToPtrOp<DstPtrT> dst_transform_op{static_cast<DstPtrT>(thrust::raw_pointer_cast(d_out.data()))};
+    cub::TransformInputIterator<DstPtrT, OffsetToPtrOp<DstPtrT>, ByteOffsetT*> d_buffer_dsts(
       thrust::raw_pointer_cast(d_buffer_dst_offsets.data()), dst_transform_op);
 
     // Get temporary storage requirements

cudax/CMakeLists.txt

@@ -16,9 +16,19 @@ endif()
 
 option(cudax_ENABLE_HEADER_TESTING "Test that CUDA Experimental's public headers compile." ON)
 option(cudax_ENABLE_TESTING "Build CUDA Experimental's tests." ON)
+option(cudax_ENABLE_EXAMPLES "Build CUDA Experimental's tests." ON)
+option(cudax_ENABLE_CUDASTF_BOUNDSCHECK "Enable bounds checks for STF targets. Requires debug build." OFF)
+option(cudax_ENABLE_CUDASTF_DEBUG "Enable additional debugging for STF targets. Requires debug build." OFF)
+option(cudax_ENABLE_CUDASTF_MATHLIBS "Enable STF tests/examples that use cublas/cusolver." OFF)
+
+if ((cudax_ENABLE_CUDASTF_BOUNDSCHECK OR cudax_ENABLE_CUDASTF_DEBUG) AND
+    NOT CMAKE_BUILD_TYPE MATCHES "Debug" AND NOT CMAKE_BUILD_TYPE MATCHES "RelWithDebInfo")
+  message(FATAL_ERROR "cudax_ENABLE_CUDASTF_BOUNDSCHECK and cudax_ENABLE_CUDASTF_DEBUG require a Debug build.")
+endif()
 
 include(cmake/cudaxBuildCompilerTargets.cmake)
 include(cmake/cudaxBuildTargetList.cmake)
+include(cmake/cudaxSTFConfigureTarget.cmake)
 
 cudax_build_compiler_targets()
 cudax_build_target_list()
@@ -28,7 +38,9 @@ if (cudax_ENABLE_HEADER_TESTING)
 endif()
 
 if (cudax_ENABLE_TESTING)
-  include(CTest)
-  enable_testing() # Must be in root directory
   add_subdirectory(test)
 endif()
+
+if (cudax_ENABLE_EXAMPLES)
+  add_subdirectory(examples)
+endif()

cudax/cmake/cudaxBuildCompilerTargets.cmake

@@ -23,6 +23,18 @@ function(cudax_build_compiler_targets)
     # cudax requires dim3 to be usable from a constexpr context, and the CUDART headers require
     # __cplusplus to be defined for this to work:
     append_option_if_available("/Zc:__cplusplus" cxx_compile_options)
+
+    # cudax requires __VA_OPT__ for its unit tests
+    append_option_if_available("/Zc:preprocessor" cxx_compile_options)
+
+    # stf used getenv which is potentially unsafe but not in our context
+    list(APPEND cxx_compile_definitions "_CRT_SECURE_NO_WARNINGS")
+  endif()
+
+  if("Clang" STREQUAL "${CMAKE_CXX_COMPILER_ID}")
+    # stf heavily uses host device lambdas which break on clang due to a warning about the implicitly
+    # deleted copy constructor
+    append_option_if_available("-Wno-deprecated-copy" cxx_compile_options)
   endif()
 
   cccl_build_compiler_interface(cudax.compiler_interface