Update CDP support macros for if-target compatibility

cub/detail/detect_cuda_runtime.cuh

@@ -0,0 +1,100 @@
+/******************************************************************************
+ * Copyright (c) 2011, Duane Merrill.  All rights reserved.
+ * Copyright (c) 2011-2021, NVIDIA CORPORATION.  All rights reserved.
+ *
+ * Redistribution and use in source and binary forms, with or without
+ * modification, are permitted provided that the following conditions are met:
+ *     * Redistributions of source code must retain the above copyright
+ *       notice, this list of conditions and the following disclaimer.
+ *     * Redistributions in binary form must reproduce the above copyright
+ *       notice, this list of conditions and the following disclaimer in the
+ *       documentation and/or other materials provided with the distribution.
+ *     * Neither the name of the NVIDIA CORPORATION nor the
+ *       names of its contributors may be used to endorse or promote products
+ *       derived from this software without specific prior written permission.
+ *
+ * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS "AS IS" AND
+ * ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED TO, THE IMPLIED
+ * WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR PURPOSE ARE
+ * DISCLAIMED. IN NO EVENT SHALL NVIDIA CORPORATION BE LIABLE FOR ANY
+ * DIRECT, INDIRECT, INCIDENTAL, SPECIAL, EXEMPLARY, OR CONSEQUENTIAL DAMAGES
+ * (INCLUDING, BUT NOT LIMITED TO, PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES;
+ * LOSS OF USE, DATA, OR PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND
+ * ON ANY THEORY OF LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT
+ * (INCLUDING NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS
+ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE.
+ *
+ ******************************************************************************/
+
+/**
+ * \file
+ * Utilities for CUDA dynamic parallelism.
+ */
+
+#pragma once
+
+#include <cub/util_namespace.cuh>
+
+#include <cuda_runtime_api.h>
+
+CUB_NAMESPACE_BEGIN
+namespace detail
+{
+
+#ifdef DOXYGEN_SHOULD_SKIP_THIS // Only parse this during doxygen passes:
+
+/**
+ * \def CUB_RDC_ENABLED
+ *
+ * Defined if RDC is enabled.
+ */
+#define CUB_RDC_ENABLED
+
+/**
+ * \def CUB_RUNTIME_FUNCTION
+ *
+ * Execution space for functions that can use the CUDA runtime API (`__host__`
+ * when RDC is off, `__host__ __device__` when RDC is on).
+ */
+#define CUB_RUNTIME_FUNCTION
+
+/**
+ * \def CUB_RUNTIME_ENABLED
+ *
+ * Whether or not the active compiler pass is allowed to invoke device kernels
+ * or methods from the CUDA runtime API.
+ *
+ * This macro should not be used in CUB, as it depends on `__CUDA_ARCH__`
+ * and is not compatible with `NV_IF_TARGET`. It is provided for legacy
+ * purposes only.
+ *
+ * Replace any usages with `CUB_RDC_ENABLED` and `NV_IF_TARGET`.
+ */
+#define CUB_RUNTIME_ENABLED
+
+#else // Non-doxygen pass:
+
+#ifndef CUB_RUNTIME_FUNCTION
+
+#if defined(__CUDACC_RDC__)
+
+#define CUB_RDC_ENABLED
+#define CUB_RUNTIME_FUNCTION __host__ __device__
+
+#else // RDC disabled:
+
+#define CUB_RUNTIME_FUNCTION __host__
+
+#endif // RDC enabled
+
+#if !defined(__CUDA_ARCH__) || defined(__CUDACC_RDC__)
+// Legacy only -- do not use in new code.
+#define CUB_RUNTIME_ENABLED
+#endif
+
+#endif // CUB_RUNTIME_FUNCTION predefined
+
+#endif // Do not document
+
+} // namespace detail
+CUB_NAMESPACE_END

cub/detail/device_synchronize.cuh

@@ -16,6 +16,7 @@
 
 #pragma once
 
+#include <cub/detail/detect_cuda_runtime.cuh>
 #include <cub/detail/exec_check_disable.cuh>
 #include <cub/util_arch.cuh>
 #include <cub/util_namespace.cuh>
@@ -38,8 +39,6 @@ CUB_RUNTIME_FUNCTION inline cudaError_t device_synchronize()
 {
   cudaError_t result = cudaErrorUnknown;
 
-#ifdef CUB_RUNTIME_ENABLED
-
 #if defined(__CUDACC__) &&                                                     \
   ((__CUDACC_VER_MAJOR__ > 11) ||                                              \
    ((__CUDACC_VER_MAJOR__ == 11) && (__CUDACC_VER_MINOR__ >= 6)))
@@ -48,11 +47,6 @@ CUB_RUNTIME_FUNCTION inline cudaError_t device_synchronize()
   result = __cudaDeviceSynchronizeDeprecationAvoidance();
 #else // CUDA < 11.6
 #define CUB_TMP_DEVICE_SYNC_IMPL result = cudaDeviceSynchronize();
-#endif
-
-#else // Device code without the CUDA runtime.
-  // Device side CUDA API calls are not supported in this configuration.
-#define CUB_TMP_DEVICE_SYNC_IMPL result = cudaErrorInvalidConfiguration;
 #endif
 
   NV_IF_TARGET(NV_IS_HOST,

cub/detail/type_traits.cuh

@@ -32,8 +32,8 @@
 
 #pragma once
 
-#include "../util_cpp_dialect.cuh"
-#include "../util_namespace.cuh"
+#include <cub/util_cpp_dialect.cuh>
+#include <cub/util_namespace.cuh>
 
 #include <type_traits>
 

cub/device/dispatch/dispatch_adjacent_difference.cuh

@@ -259,12 +259,10 @@ struct DispatchAdjacentDifference : public SelectedPolicy
                 num_tiles,
                 tile_size);
 
-        if (debug_synchronous)
+        error = detail::DebugSyncStream(stream, debug_synchronous);
+        if (CubDebug(error))
         {
-          if (CubDebug(error = SyncStream(stream)))
-          {
-            break;
-          }
+          break;
         }
 
         // Check for failure to launch
@@ -302,12 +300,11 @@ struct DispatchAdjacentDifference : public SelectedPolicy
               difference_op,
               num_items);
 
-      if (debug_synchronous)
+
+      error = detail::DebugSyncStream(stream, debug_synchronous);
+      if (CubDebug(error))
       {
-        if (CubDebug(error = SyncStream(stream)))
-        {
-          break;
-        }
+        break;
       }
 
       // Check for failure to launch

cub/device/dispatch/dispatch_histogram.cuh

@@ -35,12 +35,12 @@
 #pragma once
 
 #include <cub/agent/agent_histogram.cuh>
+#include <cub/config.cuh>
+#include <cub/grid/grid_queue.cuh>
+#include <cub/thread/thread_search.cuh>
 #include <cub/util_debug.cuh>
 #include <cub/util_device.cuh>
 #include <cub/util_math.cuh>
-#include <cub/thread/thread_search.cuh>
-#include <cub/grid/grid_queue.cuh>
-#include <cub/config.cuh>
 
 #include <thrust/system/cuda/detail/core/triple_chevron_launch.h>
 
@@ -530,13 +530,6 @@ public:
         cudaStream_t                        stream,                                         ///< [in] CUDA stream to launch kernels within.  Default is stream<sub>0</sub>.
         bool                                debug_synchronous)                              ///< [in] Whether or not to synchronize the stream after every kernel launch to check for errors.  May cause significant slowdown.  Default is \p false.
     {
-    #ifndef CUB_RUNTIME_ENABLED
-
-        // Kernel launch not supported from this device
-        return CubDebug(cudaErrorNotSupported);
-
-    #else
-
         cudaError error = cudaSuccess;
         do
         {
@@ -674,17 +667,21 @@ public:
                 tile_queue);
 
             // Check for failure to launch
-            if (CubDebug(error = cudaPeekAtLastError())) break;
+            if (CubDebug(error = cudaPeekAtLastError()))
+            {
+              break;
+            }
 
             // Sync the stream if specified to flush runtime errors
-            if (debug_synchronous && (CubDebug(error = SyncStream(stream)))) break;
-
+            error = detail::DebugSyncStream(stream, debug_synchronous);
+            if (CubDebug(error))
+            {
+                break;
+            }
         }
         while (0);
 
         return error;
-
-    #endif // CUB_RUNTIME_ENABLED
     }
 
 

cub/device/dispatch/dispatch_merge_sort.cuh

@@ -704,12 +704,10 @@ struct DispatchMergeSort : SelectedPolicy
 
       block_sort_launcher.launch();
 
-      if (debug_synchronous)
+      error = detail::DebugSyncStream(stream, debug_synchronous);
+      if (CubDebug(error))
       {
-        if (CubDebug(error = SyncStream(stream)))
-        {
-          break;
-        }
+        break;
       }
 
       // Check for failure to launch
@@ -769,12 +767,10 @@ struct DispatchMergeSort : SelectedPolicy
                 target_merged_tiles_number,
                 tile_size);
 
-        if (debug_synchronous)
+        error = detail::DebugSyncStream(stream, debug_synchronous);
+        if (CubDebug(error))
         {
-          if (CubDebug(error = SyncStream(stream)))
-          {
-            break;
-          }
+          break;
         }
 
         // Check for failure to launch
@@ -786,12 +782,10 @@ struct DispatchMergeSort : SelectedPolicy
         // Merge
         merge_launcher.launch(ping, target_merged_tiles_number);
 
-        if (debug_synchronous)
+        error = detail::DebugSyncStream(stream, debug_synchronous);
+        if (CubDebug(error))
         {
-          if (CubDebug(error = SyncStream(stream)))
-          {
-            break;
-          }
+          break;
         }
 
         // Check for failure to launch