Use nullptr over NULL (#1805)

NVIDIA · Jun 5, 2024 · 0bfbe60 · 0bfbe60
1 parent 4808778
commit 0bfbe60
Show file tree

Hide file tree

Showing 77 changed files with 275 additions and 274 deletions.
diff --git a/cub/CHANGELOG.md b/cub/CHANGELOG.md
@@ -1269,7 +1269,7 @@ CUB 1.5.0 introduces segmented sort and reduction primitives.
     warp-reduction on non-primitive data types (e.g. user-defined structs).
 - Fix small radix sorting problems where 0 temporary bytes were required and
     users code was invoking `malloc(0)` on some systems where that returns
-    `NULL`.
+    `nullptr`.
   CUB assumed the user was asking for the size again and not running the sort.
 
 # CUB 1.4.1

diff --git a/cub/cub/agent/agent_histogram.cuh b/cub/cub/agent/agent_histogram.cuh
@@ -281,7 +281,7 @@ struct AgentHistogram
   /// Sample input iterator (with cache modifier applied, if possible)
   WrappedSampleIteratorT d_wrapped_samples;
 
-  /// Native pointer for input samples (possibly NULL if unavailable)
+  /// Native pointer for input samples (possibly nullptr if unavailable)
   SampleT* d_native_samples;
 
   /// The number of output bins for each channel
@@ -771,7 +771,7 @@ struct AgentHistogram
   template <typename IteratorT>
   _CCCL_DEVICE _CCCL_FORCEINLINE SampleT* NativePointer(IteratorT itr)
   {
-    return NULL;
+    return nullptr;
   }
 
   //---------------------------------------------------------------------
@@ -875,7 +875,7 @@ struct AgentHistogram
       ((row_bytes & pixel_mask) == 0); // number of row-samples is a multiple of the alignment of the pixel
 
     // Whether rows are aligned and can be vectorized
-    if ((d_native_samples != NULL) && (vec_aligned_rows || pixel_aligned_rows))
+    if ((d_native_samples != nullptr) && (vec_aligned_rows || pixel_aligned_rows))
     {
       ConsumeTiles<true>(
         num_row_pixels, num_rows, row_stride_samples, tiles_per_row, tile_queue, Int2Type<IS_WORK_STEALING>());

diff --git a/cub/cub/agent/agent_radix_sort_onesweep.cuh b/cub/cub/agent/agent_radix_sort_onesweep.cuh
@@ -459,7 +459,7 @@ struct AgentRadixSortOnesweep
   _CCCL_DEVICE _CCCL_FORCEINLINE void UpdateBinsGlobal(int (&bins)[BINS_PER_THREAD], int (&offsets)[BINS_PER_THREAD])
   {
     bool last_block = (block_idx + 1) * TILE_ITEMS >= num_items;
-    if (d_bins_out != NULL && last_block)
+    if (d_bins_out != nullptr && last_block)
     {
 #pragma unroll
       for (int u = 0; u < BINS_PER_THREAD; ++u)

diff --git a/cub/cub/agent/agent_spmv_orig.cuh b/cub/cub/agent/agent_spmv_orig.cuh
@@ -683,7 +683,7 @@ struct AgentSpmv
     // Read our starting coordinates
     if (threadIdx.x < 2)
     {
-      if (d_tile_coordinates == NULL)
+      if (d_tile_coordinates == nullptr)
       {
         // Search our starting coordinates
         OffsetT diagonal = (tile_idx + threadIdx.x) * TILE_ITEMS;

diff --git a/cub/cub/agent/single_pass_scan_operators.cuh b/cub/cub/agent/single_pass_scan_operators.cuh
@@ -532,7 +532,7 @@ struct ScanTileState<T, true>
 
   /// Constructor
   _CCCL_HOST_DEVICE _CCCL_FORCEINLINE ScanTileState()
-      : d_tile_descriptors(NULL)
+      : d_tile_descriptors(nullptr)
   {}
 
   /**
@@ -543,7 +543,7 @@ struct ScanTileState<T, true>
    *
    * @param[in] d_temp_storage
    *   Device-accessible allocation of temporary storage.
-   *   When NULL, the required allocation size is written to \p temp_storage_bytes and no work is
+   *   When nullptr, the required allocation size is written to \p temp_storage_bytes and no work is
    * done.
    *
    * @param[in] temp_storage_bytes
@@ -687,9 +687,9 @@ struct ScanTileState<T, false>
 
   /// Constructor
   _CCCL_HOST_DEVICE _CCCL_FORCEINLINE ScanTileState()
-      : d_tile_status(NULL)
-      , d_tile_partial(NULL)
-      , d_tile_inclusive(NULL)
+      : d_tile_status(nullptr)
+      , d_tile_partial(nullptr)
+      , d_tile_inclusive(nullptr)
   {}
 
   /**
@@ -700,7 +700,7 @@ struct ScanTileState<T, false>
    *
    * @param[in] d_temp_storage
    *   Device-accessible allocation of temporary storage.
-   *   When NULL, the required allocation size is written to \p temp_storage_bytes and no work is
+   *   When nullptr, the required allocation size is written to \p temp_storage_bytes and no work is
    *   done.
    *
    * @param[in] temp_storage_bytes
@@ -766,7 +766,7 @@ struct ScanTileState<T, false>
 
     // Set the necessary size of the blob
     void* allocations[3] = {};
-    return CubDebug(AliasTemporaries(NULL, temp_storage_bytes, allocations, allocation_sizes));
+    return CubDebug(AliasTemporaries(nullptr, temp_storage_bytes, allocations, allocation_sizes));
   }
 
   /**
@@ -927,7 +927,7 @@ struct ReduceByKeyScanTileState<ValueT, KeyT, true>
 
   /// Constructor
   _CCCL_HOST_DEVICE _CCCL_FORCEINLINE ReduceByKeyScanTileState()
-      : d_tile_descriptors(NULL)
+      : d_tile_descriptors(nullptr)
   {}
 
   /**
@@ -937,7 +937,7 @@ struct ReduceByKeyScanTileState<ValueT, KeyT, true>
    *   Number of tiles
    *
    * @param[in] d_temp_storage
-   *   Device-accessible allocation of temporary storage.  When NULL, the required allocation size
+   *   Device-accessible allocation of temporary storage.  When nullptr, the required allocation size
    *   is written to \p temp_storage_bytes and no work is done.
    *
    * @param[in] temp_storage_bytes

diff --git a/cub/cub/device/device_adjacent_difference.cuh b/cub/cub/device/device_adjacent_difference.cuh
@@ -93,7 +93,7 @@ CUB_NAMESPACE_BEGIN
 //!    //...
 //!
 //!    // Determine temporary device storage requirements
-//!    void     *d_temp_storage = NULL;
+//!    void     *d_temp_storage = nullptr;
 //!    size_t   temp_storage_bytes = 0;
 //!
 //!    cub::DeviceAdjacentDifference::SubtractLeft(
@@ -178,7 +178,7 @@ public:
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void     *d_temp_storage = NULL;
+  //!    void     *d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!
   //!    cub::DeviceAdjacentDifference::SubtractLeftCopy(
@@ -320,7 +320,7 @@ public:
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void     *d_temp_storage = NULL;
+  //!    void     *d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceAdjacentDifference::SubtractLeft(
   //!      d_temp_storage, temp_storage_bytes,
@@ -583,7 +583,7 @@ public:
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void *d_temp_storage = NULL;
+  //!    void *d_temp_storage = nullptr;
   //!    size_t temp_storage_bytes = 0;
   //!    cub::DeviceAdjacentDifference::SubtractRight(
   //!      d_temp_storage, temp_storage_bytes, d_data, num_items);

diff --git a/cub/cub/device/device_histogram.cuh b/cub/cub/device/device_histogram.cuh
@@ -113,7 +113,7 @@ struct DeviceHistogram
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void*    d_temp_storage = NULL;
+  //!    void*    d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceHistogram::HistogramEven(
   //!      d_temp_storage, temp_storage_bytes,
@@ -283,7 +283,7 @@ struct DeviceHistogram
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void*    d_temp_storage  = NULL;
+  //!    void*    d_temp_storage  = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceHistogram::HistogramEven(
   //!        d_temp_storage, temp_storage_bytes,
@@ -467,7 +467,7 @@ struct DeviceHistogram
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void*    d_temp_storage = NULL;
+  //!    void*    d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceHistogram::MultiHistogramEven<4, 3>(
   //!      d_temp_storage, temp_storage_bytes,
@@ -679,7 +679,7 @@ struct DeviceHistogram
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void*    d_temp_storage = NULL;
+  //!    void*    d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceHistogram::MultiHistogramEven<4, 3>(
   //!      d_temp_storage, temp_storage_bytes,
@@ -902,7 +902,7 @@ struct DeviceHistogram
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void*    d_temp_storage = NULL;
+  //!    void*    d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceHistogram::HistogramRange(
   //!      d_temp_storage, temp_storage_bytes,
@@ -1049,7 +1049,7 @@ struct DeviceHistogram
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void*    d_temp_storage = NULL;
+  //!    void*    d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceHistogram::HistogramRange(
   //!      d_temp_storage, temp_storage_bytes,
@@ -1220,7 +1220,7 @@ struct DeviceHistogram
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void*    d_temp_storage = NULL;
+  //!    void*    d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceHistogram::MultiHistogramRange<4, 3>(
   //!      d_temp_storage, temp_storage_bytes,
@@ -1412,7 +1412,7 @@ struct DeviceHistogram
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void*    d_temp_storage = NULL;
+  //!    void*    d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceHistogram::MultiHistogramRange<4, 3>(
   //!      d_temp_storage, temp_storage_bytes,

diff --git a/cub/cub/device/device_radix_sort.cuh b/cub/cub/device/device_radix_sort.cuh
@@ -263,7 +263,7 @@ public:
   //! ...
   //!
   //! // Determine temporary device storage requirements
-  //! void     *d_temp_storage = NULL;
+  //! void     *d_temp_storage = nullptr;
   //! size_t   temp_storage_bytes = 0;
   //! cub::DeviceRadixSort::SortPairs(d_temp_storage, temp_storage_bytes,
   //!     d_keys_in, d_keys_out, d_values_in, d_values_out, num_items);
@@ -739,7 +739,7 @@ public:
   //! cub::DoubleBuffer<int> d_values(d_value_buf, d_value_alt_buf);
   //!
   //! // Determine temporary device storage requirements
-  //! void     *d_temp_storage = NULL;
+  //! void     *d_temp_storage = nullptr;
   //! size_t   temp_storage_bytes = 0;
   //! cub::DeviceRadixSort::SortPairs(
   //!   d_temp_storage, temp_storage_bytes, d_keys, d_values, num_items);
@@ -1160,7 +1160,7 @@ public:
   //! ...
   //!
   //! // Determine temporary device storage requirements
-  //! void     *d_temp_storage = NULL;
+  //! void     *d_temp_storage = nullptr;
   //! size_t   temp_storage_bytes = 0;
   //! cub::DeviceRadixSort::SortPairsDescending(
   //!     d_temp_storage, temp_storage_bytes,
@@ -1627,7 +1627,7 @@ public:
   //! cub::DoubleBuffer<int> d_values(d_value_buf, d_value_alt_buf);
   //!
   //! // Determine temporary device storage requirements
-  //! void     *d_temp_storage = NULL;
+  //! void     *d_temp_storage = nullptr;
   //! size_t   temp_storage_bytes = 0;
   //! cub::DeviceRadixSort::SortPairsDescending(
   //!   d_temp_storage, temp_storage_bytes, d_keys, d_values, num_items);
@@ -2057,7 +2057,7 @@ public:
   //! ...
   //!
   //! // Determine temporary device storage requirements
-  //! void     *d_temp_storage = NULL;
+  //! void     *d_temp_storage = nullptr;
   //! size_t   temp_storage_bytes = 0;
   //! cub::DeviceRadixSort::SortKeys(
   //!   d_temp_storage, temp_storage_bytes, d_keys_in, d_keys_out, num_items);
@@ -2480,7 +2480,7 @@ public:
   //! cub::DoubleBuffer<int> d_keys(d_key_buf, d_key_alt_buf);
   //!
   //! // Determine temporary device storage requirements
-  //! void     *d_temp_storage = NULL;
+  //! void     *d_temp_storage = nullptr;
   //! size_t   temp_storage_bytes = 0;
   //! cub::DeviceRadixSort::SortKeys(
   //!   d_temp_storage, temp_storage_bytes, d_keys, num_items);
@@ -2867,7 +2867,7 @@ public:
   //! cub::DoubleBuffer<int> d_keys(d_key_buf, d_key_alt_buf);
   //!
   //! // Determine temporary device storage requirements
-  //! void     *d_temp_storage = NULL;
+  //! void     *d_temp_storage = nullptr;
   //! size_t   temp_storage_bytes = 0;
   //! cub::DeviceRadixSort::SortKeysDescending(
   //!   d_temp_storage, temp_storage_bytes, d_keys_in, d_keys_out, num_items);
@@ -3273,7 +3273,7 @@ public:
   //! cub::DoubleBuffer<int> d_keys(d_key_buf, d_key_alt_buf);
   //!
   //! // Determine temporary device storage requirements
-  //! void     *d_temp_storage = NULL;
+  //! void     *d_temp_storage = nullptr;
   //! size_t   temp_storage_bytes = 0;
   //! cub::DeviceRadixSort::SortKeysDescending(
   //!   d_temp_storage, temp_storage_bytes, d_keys, num_items);

diff --git a/cub/cub/device/device_reduce.cuh b/cub/cub/device/device_reduce.cuh
@@ -125,7 +125,7 @@ struct DeviceReduce
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void     *d_temp_storage = NULL;
+  //!    void     *d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceReduce::Reduce(
   //!      d_temp_storage, temp_storage_bytes,
@@ -254,7 +254,7 @@ struct DeviceReduce
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void     *d_temp_storage = NULL;
+  //!    void     *d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceReduce::Sum(
   //!      d_temp_storage, temp_storage_bytes, d_in, d_out, num_items);
@@ -374,7 +374,7 @@ struct DeviceReduce
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void     *d_temp_storage = NULL;
+  //!    void     *d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceReduce::Min(
   //!      d_temp_storage, temp_storage_bytes, d_in, d_out, num_items);
@@ -503,7 +503,7 @@ struct DeviceReduce
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void     *d_temp_storage = NULL;
+  //!    void     *d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceReduce::ArgMin(d_temp_storage, temp_storage_bytes, d_in, d_argmin, num_items);
   //!
@@ -630,7 +630,7 @@ struct DeviceReduce
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void     *d_temp_storage = NULL;
+  //!    void     *d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceReduce::Max(d_temp_storage, temp_storage_bytes, d_in, d_max, num_items);
   //!
@@ -761,7 +761,7 @@ struct DeviceReduce
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void     *d_temp_storage = NULL;
+  //!    void     *d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceReduce::ArgMax(
   //!      d_temp_storage, temp_storage_bytes, d_in, d_argmax, num_items);
@@ -1053,7 +1053,7 @@ struct DeviceReduce
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void     *d_temp_storage = NULL;
+  //!    void     *d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceReduce::ReduceByKey(
   //!      d_temp_storage, temp_storage_bytes,

diff --git a/cub/cub/device/device_run_length_encode.cuh b/cub/cub/device/device_run_length_encode.cuh
@@ -116,7 +116,7 @@ struct DeviceRunLengthEncode
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void     *d_temp_storage = NULL;
+  //!    void     *d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceRunLengthEncode::Encode(
   //!      d_temp_storage, temp_storage_bytes,
@@ -287,7 +287,7 @@ struct DeviceRunLengthEncode
   //!    ...
   //!
   //!    // Determine temporary device storage requirements
-  //!    void     *d_temp_storage = NULL;
+  //!    void     *d_temp_storage = nullptr;
   //!    size_t   temp_storage_bytes = 0;
   //!    cub::DeviceRunLengthEncode::NonTrivialRuns(
   //!      d_temp_storage, temp_storage_bytes,