diff --git a/cub/CHANGELOG.md b/cub/CHANGELOG.md index 7d3cd839eb5..1f03565c118 100644 --- 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 index 5f3183ad8fc..584499e9367 100644 --- 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 _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( num_row_pixels, num_rows, row_stride_samples, tiles_per_row, tile_queue, Int2Type()); diff --git a/cub/cub/agent/agent_radix_sort_onesweep.cuh b/cub/cub/agent/agent_radix_sort_onesweep.cuh index cbe75b3edf6..45cb2052e52 100644 --- 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 index f6254050192..a69bc14510e 100644 --- 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 index 9dc11ea1aff..8df7066aa7a 100644 --- a/cub/cub/agent/single_pass_scan_operators.cuh +++ b/cub/cub/agent/single_pass_scan_operators.cuh @@ -532,7 +532,7 @@ struct ScanTileState /// Constructor _CCCL_HOST_DEVICE _CCCL_FORCEINLINE ScanTileState() - : d_tile_descriptors(NULL) + : d_tile_descriptors(nullptr) {} /** @@ -543,7 +543,7 @@ struct ScanTileState * * @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 /// 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 * * @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 // 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 /// Constructor _CCCL_HOST_DEVICE _CCCL_FORCEINLINE ReduceByKeyScanTileState() - : d_tile_descriptors(NULL) + : d_tile_descriptors(nullptr) {} /** @@ -937,7 +937,7 @@ struct ReduceByKeyScanTileState * 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 index 5917227bf55..1fd9497bf49 100644 --- 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 index 7162b1ff422..dc817e0fe27 100644 --- 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 index db63745d4de..3ef444f3af1 100644 --- 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 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 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 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 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 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 index dabe8f9d4e4..ee85fd959a1 100644 --- 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 index 96570e31f4b..f19102e0cc3 100644 --- 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, diff --git a/cub/cub/device/device_scan.cuh b/cub/cub/device/device_scan.cuh index ff7b6d25677..44e9c81e8c7 100644 --- a/cub/cub/device/device_scan.cuh +++ b/cub/cub/device/device_scan.cuh @@ -128,7 +128,7 @@ struct DeviceScan //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceScan::ExclusiveSum( //! d_temp_storage, temp_storage_bytes, @@ -238,7 +238,7 @@ struct DeviceScan //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceScan::ExclusiveSum( //! d_temp_storage, temp_storage_bytes, @@ -342,7 +342,7 @@ struct DeviceScan //! //! // Determine temporary device storage requirements for exclusive //! // prefix scan - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceScan::ExclusiveScan( //! d_temp_storage, temp_storage_bytes, @@ -486,7 +486,7 @@ struct DeviceScan //! //! // Determine temporary device storage requirements for exclusive //! // prefix scan - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceScan::ExclusiveScan( //! d_temp_storage, temp_storage_bytes, @@ -616,7 +616,7 @@ struct DeviceScan //! //! // Determine temporary device storage requirements for exclusive //! // prefix scan - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceScan::ExclusiveScan( //! d_temp_storage, temp_storage_bytes, @@ -771,7 +771,7 @@ struct DeviceScan //! //! // Determine temporary device storage requirements for exclusive //! // prefix scan - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceScan::ExclusiveScan( //! d_temp_storage, temp_storage_bytes, @@ -1510,7 +1510,7 @@ struct DeviceScan //! //! // Determine temporary device storage requirements for exclusive //! // prefix scan - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceScan::ExclusiveScanByKey( //! d_temp_storage, temp_storage_bytes, @@ -1700,7 +1700,7 @@ struct DeviceScan //! ... //! //! // Determine temporary device storage requirements for inclusive prefix sum - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceScan::InclusiveSumByKey( //! d_temp_storage, temp_storage_bytes, @@ -1875,7 +1875,7 @@ struct DeviceScan //! ... //! //! // Determine temporary device storage requirements for inclusive prefix scan - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceScan::InclusiveScanByKey( //! d_temp_storage, temp_storage_bytes, diff --git a/cub/cub/device/device_segmented_radix_sort.cuh b/cub/cub/device/device_segmented_radix_sort.cuh index 4bae106e1de..cff9c22bce4 100644 --- a/cub/cub/device/device_segmented_radix_sort.cuh +++ b/cub/cub/device/device_segmented_radix_sort.cuh @@ -139,7 +139,7 @@ public: //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSegmentedRadixSort::SortPairs( //! d_temp_storage, temp_storage_bytes, @@ -357,7 +357,7 @@ public: //! cub::DoubleBuffer 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::DeviceSegmentedRadixSort::SortPairs( //! d_temp_storage, temp_storage_bytes, d_keys, d_values, @@ -549,7 +549,7 @@ public: //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSegmentedRadixSort::SortPairsDescending( //! d_temp_storage, temp_storage_bytes, @@ -773,7 +773,7 @@ public: //! cub::DoubleBuffer 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::DeviceSegmentedRadixSort::SortPairsDescending( //! d_temp_storage, temp_storage_bytes, d_keys, d_values, @@ -969,7 +969,7 @@ public: //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSegmentedRadixSort::SortKeys( //! d_temp_storage, temp_storage_bytes, d_keys_in, d_keys_out, @@ -1167,7 +1167,7 @@ public: //! cub::DoubleBuffer 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::DeviceSegmentedRadixSort::SortKeys( //! d_temp_storage, temp_storage_bytes, d_keys, @@ -1355,7 +1355,7 @@ public: //! cub::DoubleBuffer 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::DeviceSegmentedRadixSort::SortKeysDescending( //! d_temp_storage, temp_storage_bytes, d_keys_in, d_keys_out, @@ -1550,7 +1550,7 @@ public: //! cub::DoubleBuffer 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::DeviceSegmentedRadixSort::SortKeysDescending( //! d_temp_storage, temp_storage_bytes, d_keys, diff --git a/cub/cub/device/device_segmented_sort.cuh b/cub/cub/device/device_segmented_sort.cuh index 1b56e7e7435..2aeb145c5d1 100644 --- a/cub/cub/device/device_segmented_sort.cuh +++ b/cub/cub/device/device_segmented_sort.cuh @@ -105,7 +105,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::DeviceSegmentedSort::SortPairs( //! d_temp_storage, temp_storage_bytes, @@ -213,7 +213,7 @@ public: //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSegmentedSort::SortKeys( //! d_temp_storage, temp_storage_bytes, d_keys_in, d_keys_out, @@ -411,7 +411,7 @@ public: //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSegmentedSort::SortKeysDescending( //! d_temp_storage, temp_storage_bytes, d_keys_in, d_keys_out, @@ -615,7 +615,7 @@ public: //! cub::DoubleBuffer 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::DeviceSegmentedSort::SortKeys( //! d_temp_storage, temp_storage_bytes, d_keys, @@ -804,7 +804,7 @@ public: //! cub::DoubleBuffer 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::DeviceSegmentedSort::SortKeysDescending( //! d_temp_storage, temp_storage_bytes, d_keys, @@ -948,7 +948,7 @@ public: //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSegmentedSort::StableSortKeys( //! d_temp_storage, temp_storage_bytes, d_keys_in, d_keys_out, @@ -1109,7 +1109,7 @@ public: //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSegmentedSort::StableSortKeysDescending( //! d_temp_storage, temp_storage_bytes, d_keys_in, d_keys_out, @@ -1285,7 +1285,7 @@ public: //! cub::DoubleBuffer 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::DeviceSegmentedSort::StableSortKeys( //! d_temp_storage, temp_storage_bytes, d_keys, @@ -1440,7 +1440,7 @@ public: //! cub::DoubleBuffer 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::DeviceSegmentedSort::StableSortKeysDescending( //! d_temp_storage, temp_storage_bytes, d_keys, @@ -1629,7 +1629,7 @@ public: //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSegmentedSort::SortPairs( //! d_temp_storage, temp_storage_bytes, @@ -1855,7 +1855,7 @@ public: //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSegmentedSort::SortPairsDescending( //! d_temp_storage, temp_storage_bytes, @@ -2091,7 +2091,7 @@ public: //! cub::DoubleBuffer 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::DeviceSegmentedSort::SortPairs( //! d_temp_storage, temp_storage_bytes, d_keys, d_values, @@ -2312,7 +2312,7 @@ public: //! cub::DoubleBuffer 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::DeviceSegmentedSort::SortPairsDescending( //! d_temp_storage, temp_storage_bytes, d_keys, d_values, @@ -2488,7 +2488,7 @@ public: //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSegmentedSort::StableSortPairs( //! d_temp_storage, temp_storage_bytes, @@ -2676,7 +2676,7 @@ public: //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSegmentedSort::StableSortPairsDescending( //! d_temp_storage, temp_storage_bytes, @@ -2881,7 +2881,7 @@ public: //! cub::DoubleBuffer 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::DeviceSegmentedSort::StableSortPairs( //! d_temp_storage, temp_storage_bytes, d_keys, d_values, @@ -3071,7 +3071,7 @@ public: //! cub::DoubleBuffer 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::DeviceSegmentedSort::StableSortPairsDescending( //! d_temp_storage, temp_storage_bytes, d_keys, d_values, diff --git a/cub/cub/device/device_select.cuh b/cub/cub/device/device_select.cuh index f5cf533f4b3..2afa303b7f1 100644 --- a/cub/cub/device/device_select.cuh +++ b/cub/cub/device/device_select.cuh @@ -109,7 +109,7 @@ struct DeviceSelect //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSelect::Flagged( //! d_temp_storage, temp_storage_bytes, @@ -249,7 +249,7 @@ struct DeviceSelect //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSelect::Flagged( //! d_temp_storage, temp_storage_bytes, @@ -399,7 +399,7 @@ struct DeviceSelect //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSelect::If( //! d_temp_storage, temp_storage_bytes, @@ -485,7 +485,7 @@ struct DeviceSelect false>::Dispatch(d_temp_storage, temp_storage_bytes, d_in, - NULL, + nullptr, d_out, d_num_selected_out, select_op, @@ -552,7 +552,7 @@ struct DeviceSelect //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSelect::If( //! d_temp_storage, temp_storage_bytes, @@ -633,7 +633,7 @@ struct DeviceSelect may_alias>::Dispatch(d_temp_storage, temp_storage_bytes, d_data, // in - NULL, + nullptr, d_data, // out d_num_selected_out, select_op, @@ -914,7 +914,7 @@ struct DeviceSelect //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSelect::Unique( //! d_temp_storage, temp_storage_bytes, @@ -994,7 +994,7 @@ struct DeviceSelect false>::Dispatch(d_temp_storage, temp_storage_bytes, d_in, - NULL, + nullptr, d_out, d_num_selected_out, SelectOp(), @@ -1060,7 +1060,7 @@ struct DeviceSelect //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSelect::UniqueByKey( //! d_temp_storage, temp_storage_bytes, @@ -1221,7 +1221,7 @@ struct DeviceSelect //! ... //! //! // Determine temporary device storage requirements - //! void *d_temp_storage = NULL; + //! void *d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSelect::UniqueByKey( //! d_temp_storage, temp_storage_bytes, diff --git a/cub/cub/device/device_spmv.cuh b/cub/cub/device/device_spmv.cuh index 30e8545da92..93e51f2293f 100644 --- a/cub/cub/device/device_spmv.cuh +++ b/cub/cub/device/device_spmv.cuh @@ -116,7 +116,7 @@ struct DeviceSpmv //! ... //! //! // Determine temporary device storage requirements - //! void* d_temp_storage = NULL; + //! void* d_temp_storage = nullptr; //! size_t temp_storage_bytes = 0; //! cub::DeviceSpmv::CsrMV(d_temp_storage, temp_storage_bytes, d_values, //! d_row_offsets, d_column_indices, d_vector_x, d_vector_y, @@ -139,7 +139,7 @@ struct DeviceSpmv //! //! @param[in] d_temp_storage //! Device-accessible allocation of temporary storage. - //! When NULL, the required allocation size is written to `temp_storage_bytes` and no work is done. + //! When nullptr, the required allocation size is written to `temp_storage_bytes` and no work is done. //! //! @param[in,out] temp_storage_bytes //! Reference to size in bytes of `d_temp_storage` allocation diff --git a/cub/cub/device/dispatch/dispatch_histogram.cuh b/cub/cub/device/dispatch/dispatch_histogram.cuh index 7259b1fb298..cdefca6df0b 100644 --- a/cub/cub/device/dispatch/dispatch_histogram.cuh +++ b/cub/cub/device/dispatch/dispatch_histogram.cuh @@ -865,7 +865,7 @@ public: * * @param d_temp_storage * Device-accessible allocation of temporary storage. - * When NULL, the required allocation size is written to `temp_storage_bytes` and + * When nullptr, the required allocation size is written to `temp_storage_bytes` and * no work is done. * * @param temp_storage_bytes @@ -1067,7 +1067,7 @@ public: * * @param d_temp_storage * Device-accessible allocation of temporary storage. - * When NULL, the required allocation size is written to `temp_storage_bytes` and + * When nullptr, the required allocation size is written to `temp_storage_bytes` and * no work is done. * * @param temp_storage_bytes @@ -1230,7 +1230,7 @@ public: * * @param d_temp_storage * Device-accessible allocation of temporary storage. - * When NULL, the required allocation size is written to + * When nullptr, the required allocation size is written to * `temp_storage_bytes` and no work is done. * * @param temp_storage_bytes @@ -1449,7 +1449,7 @@ public: * * @param d_temp_storage * Device-accessible allocation of temporary storage. - * When NULL, the required allocation size is written to `temp_storage_bytes` and + * When nullptr, the required allocation size is written to `temp_storage_bytes` and * no work is done. * * @param temp_storage_bytes diff --git a/cub/cub/device/dispatch/dispatch_merge_sort.cuh b/cub/cub/device/dispatch/dispatch_merge_sort.cuh index e687cf5b010..8bb025fb687 100644 --- a/cub/cub/device/dispatch/dispatch_merge_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_merge_sort.cuh @@ -388,7 +388,7 @@ struct DispatchMergeSort : SelectedPolicy // Problem state - /// Device-accessible allocation of temporary storage. When NULL, the required + /// Device-accessible allocation of temporary storage. When nullptr, the required /// allocation size is written to \p temp_storage_bytes and no work is done. void* d_temp_storage; diff --git a/cub/cub/device/dispatch/dispatch_radix_sort.cuh b/cub/cub/device/dispatch/dispatch_radix_sort.cuh index 0606485bb2c..eec687fed8e 100644 --- a/cub/cub/device/dispatch/dispatch_radix_sort.cuh +++ b/cub/cub/device/dispatch/dispatch_radix_sort.cuh @@ -1740,7 +1740,7 @@ struct DispatchRadixSort : SelectedPolicy //------------------------------------------------------------------------------ /// Device-accessible allocation of temporary storage. - // When NULL, the required allocation size is written to `temp_storage_bytes` and no work is + // When nullptr, the required allocation size is written to `temp_storage_bytes` and no work is // done. void* d_temp_storage; @@ -1856,7 +1856,7 @@ struct DispatchRadixSort : SelectedPolicy do { // Return if the caller is simply requesting the size of the storage allocation - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { temp_storage_bytes = 1; break; @@ -2151,7 +2151,7 @@ struct DispatchRadixSort : SelectedPolicy // just return if no temporary storage is provided cudaError_t error = cudaSuccess; - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { return error; } @@ -2310,7 +2310,7 @@ struct DispatchRadixSort : SelectedPolicy .doit(onesweep_kernel, d_lookback, d_ctrs + portion * num_passes + pass, - portion < num_portions - 1 ? d_bins + ((portion + 1) * num_passes + pass) * RADIX_DIGITS : NULL, + portion < num_portions - 1 ? d_bins + ((portion + 1) * num_passes + pass) * RADIX_DIGITS : nullptr, d_bins + (portion * num_passes + pass) * RADIX_DIGITS, d_keys.Alternate(), d_keys.Current() + portion * PORTION_SIZE, @@ -2457,7 +2457,7 @@ struct DispatchRadixSort : SelectedPolicy } // Return if the caller is simply requesting the size of the storage allocation - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { return cudaSuccess; } @@ -2669,7 +2669,7 @@ struct DispatchRadixSort : SelectedPolicy * @brief Internal dispatch routine * * @param[in] d_temp_storage - * Device-accessible allocation of temporary storage. When NULL, the required + * Device-accessible allocation of temporary storage. When nullptr, the required * allocation size is written to `temp_storage_bytes` and no work is done. * * @param[in,out] temp_storage_bytes @@ -2816,7 +2816,7 @@ struct DispatchSegmentedRadixSort : SelectedPolicy // Parameter members //------------------------------------------------------------------------------ - /// 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 `temp_storage_bytes` and no work is done. void* d_temp_storage; @@ -3077,7 +3077,7 @@ struct DispatchSegmentedRadixSort : SelectedPolicy } // Return if the caller is simply requesting the size of the storage allocation - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { if (temp_storage_bytes == 0) { @@ -3208,7 +3208,7 @@ struct DispatchSegmentedRadixSort : SelectedPolicy * @brief Internal dispatch routine * * @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 `temp_storage_bytes` and no work is done. * * @param[in,out] temp_storage_bytes diff --git a/cub/cub/device/dispatch/dispatch_reduce.cuh b/cub/cub/device/dispatch/dispatch_reduce.cuh index f39613adb80..aee3a296e10 100644 --- a/cub/cub/device/dispatch/dispatch_reduce.cuh +++ b/cub/cub/device/dispatch/dispatch_reduce.cuh @@ -660,7 +660,7 @@ struct DispatchReduce : SelectedPolicy { // Return if the caller is simply requesting the size of the storage // allocation - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { temp_storage_bytes = 1; break; @@ -773,7 +773,7 @@ struct DispatchReduce : SelectedPolicy break; } - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { // Return if the caller is simply requesting the size of the storage // allocation @@ -1207,7 +1207,7 @@ struct DispatchSegmentedReduce : SelectedPolicy { // Return if the caller is simply requesting the size of the storage // allocation - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { temp_storage_bytes = 1; return cudaSuccess; diff --git a/cub/cub/device/dispatch/dispatch_rle.cuh b/cub/cub/device/dispatch/dispatch_rle.cuh index 9a2bed3fe3b..9c6c32a95ce 100644 --- a/cub/cub/device/dispatch/dispatch_rle.cuh +++ b/cub/cub/device/dispatch/dispatch_rle.cuh @@ -255,7 +255,7 @@ struct DeviceRleDispatch * * @param d_temp_storage * Device-accessible allocation of temporary storage. - * When NULL, the required allocation size is written to + * When nullptr, the required allocation size is written to * `temp_storage_bytes` and no work is done. * * @param temp_storage_bytes @@ -468,7 +468,7 @@ struct DeviceRleDispatch * * @param d_temp_storage * Device-accessible allocation of temporary storage. - * When NULL, the required allocation size is written to + * When nullptr, the required allocation size is written to * `temp_storage_bytes` and no work is done. * * @param temp_storage_bytes diff --git a/cub/cub/device/dispatch/dispatch_scan.cuh b/cub/cub/device/dispatch/dispatch_scan.cuh index bc97e3d4bcc..bcdb13bc81f 100644 --- a/cub/cub/device/dispatch/dispatch_scan.cuh +++ b/cub/cub/device/dispatch/dispatch_scan.cuh @@ -244,7 +244,7 @@ struct DispatchScan : SelectedPolicy // The input value type using InputT = cub::detail::value_t; - /// Device-accessible allocation of temporary storage. When NULL, the + /// Device-accessible allocation of temporary storage. When nullptr, the /// required allocation size is written to \p temp_storage_bytes and no work /// is done. void* d_temp_storage; @@ -392,7 +392,7 @@ struct DispatchScan : SelectedPolicy break; } - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { // Return if the caller is simply requesting the size of the storage // allocation diff --git a/cub/cub/device/dispatch/dispatch_scan_by_key.cuh b/cub/cub/device/dispatch/dispatch_scan_by_key.cuh index a8cb886a2bd..9a1cdad9704 100644 --- a/cub/cub/device/dispatch/dispatch_scan_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_scan_by_key.cuh @@ -408,7 +408,7 @@ struct DispatchScanByKey : SelectedPolicy break; } - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { // Return if the caller is simply requesting the size of the storage // allocation diff --git a/cub/cub/device/dispatch/dispatch_spmv_orig.cuh b/cub/cub/device/dispatch/dispatch_spmv_orig.cuh index 92ba683df10..a7ce309a8df 100644 --- a/cub/cub/device/dispatch/dispatch_spmv_orig.cuh +++ b/cub/cub/device/dispatch/dispatch_spmv_orig.cuh @@ -545,7 +545,7 @@ struct DispatchSpmv * * @param[in] d_temp_storage * Device-accessible allocation of temporary storage. - * When NULL, the required allocation size is written to + * When nullptr, the required allocation size is written to * `temp_storage_bytes` and no work is done. * * @param[in,out] temp_storage_bytes @@ -606,7 +606,7 @@ struct DispatchSpmv if (spmv_params.num_rows == 0 || spmv_params.num_cols == 0) { // Empty problem, no-op. - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { temp_storage_bytes = 1; } @@ -616,7 +616,7 @@ struct DispatchSpmv if (spmv_params.num_nonzeros == 0) { - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { // Return if the caller is simply requesting the size of the storage allocation temp_storage_bytes = 1; @@ -652,7 +652,7 @@ struct DispatchSpmv if (spmv_params.num_cols == 1) { - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { // Return if the caller is simply requesting the size of the storage allocation temp_storage_bytes = 1; @@ -758,7 +758,7 @@ struct DispatchSpmv { break; } - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { // Return if the caller is simply requesting the size of the storage allocation break; @@ -783,7 +783,7 @@ struct DispatchSpmv // if (num_merge_tiles < spmv_sm_occupancy * sm_count) { // Not enough spmv tiles to saturate the device: have spmv blocks search their own staring coords - d_tile_coordinates = NULL; + d_tile_coordinates = nullptr; } else { @@ -933,7 +933,7 @@ struct DispatchSpmv * * @param[in] d_temp_storage * Device-accessible allocation of temporary storage. - * When NULL, the required allocation size is written to + * When nullptr, the required allocation size is written to * `temp_storage_bytes` and no work is done. * * @param[in,out] temp_storage_bytes diff --git a/cub/cub/device/dispatch/dispatch_unique_by_key.cuh b/cub/cub/device/dispatch/dispatch_unique_by_key.cuh index f62d6d27537..1d097a93a0a 100644 --- a/cub/cub/device/dispatch/dispatch_unique_by_key.cuh +++ b/cub/cub/device/dispatch/dispatch_unique_by_key.cuh @@ -232,7 +232,7 @@ struct DispatchUniqueByKey : SelectedPolicy // Tile status descriptor interface type using ScanTileStateT = ScanTileState; - /// 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 `temp_storage_bytes` and no work is done. void* d_temp_storage; @@ -267,7 +267,7 @@ struct DispatchUniqueByKey : SelectedPolicy /** * @param[in] d_temp_storage * Device-accessible allocation of temporary storage. - * When NULL, the required allocation size is written to + * When nullptr, the required allocation size is written to * `temp_storage_bytes` and no work is done. * * @tparam temp_storage_bytes @@ -397,7 +397,7 @@ struct DispatchUniqueByKey : SelectedPolicy } // Compute allocation pointers into the single storage blob (or compute the necessary size of the blob) - void* allocations[2] = {NULL, NULL}; + void* allocations[2] = {nullptr, nullptr}; error = CubDebug(AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes)); if (cudaSuccess != error) @@ -405,7 +405,7 @@ struct DispatchUniqueByKey : SelectedPolicy break; } - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { // Return if the caller is simply requesting the size of the storage allocation break; @@ -548,7 +548,7 @@ struct DispatchUniqueByKey : SelectedPolicy * * @param[in] d_temp_storage * Device-accessible allocation of temporary storage. - * When NULL, the required allocation size is written to + * When nullptr, the required allocation size is written to * `temp_storage_bytes` and no work is done. * * @param[in,out] &temp_storage_bytes diff --git a/cub/cub/grid/grid_barrier.cuh b/cub/cub/grid/grid_barrier.cuh index 79c3c67d924..88cb59eb56b 100644 --- a/cub/cub/grid/grid_barrier.cuh +++ b/cub/cub/grid/grid_barrier.cuh @@ -64,7 +64,7 @@ public: * Constructor */ GridBarrier() - : d_sync(NULL) + : d_sync(nullptr) {} /** @@ -156,7 +156,7 @@ public: if (d_sync) { retval = CubDebug(cudaFree(d_sync)); - d_sync = NULL; + d_sync = nullptr; } sync_bytes = 0; return retval; diff --git a/cub/cub/grid/grid_queue.cuh b/cub/cub/grid/grid_queue.cuh index a4621c03bb0..ec98c6445ba 100644 --- a/cub/cub/grid/grid_queue.cuh +++ b/cub/cub/grid/grid_queue.cuh @@ -101,7 +101,7 @@ public: /// Constructs an invalid GridQueue descriptor _CCCL_HOST_DEVICE _CCCL_FORCEINLINE GridQueue() - : d_counters(NULL) + : d_counters(nullptr) {} /** diff --git a/cub/cub/iterator/discard_output_iterator.cuh b/cub/cub/iterator/discard_output_iterator.cuh index 3e4aeebc426..acd68c44a28 100644 --- a/cub/cub/iterator/discard_output_iterator.cuh +++ b/cub/cub/iterator/discard_output_iterator.cuh @@ -191,7 +191,7 @@ public: /// Cast to void* operator _CCCL_HOST_DEVICE _CCCL_FORCEINLINE operator void*() const { - return NULL; + return nullptr; } /// Equal to diff --git a/cub/cub/iterator/tex_obj_input_iterator.cuh b/cub/cub/iterator/tex_obj_input_iterator.cuh index 358f78c10c5..48ebd6d3353 100644 --- a/cub/cub/iterator/tex_obj_input_iterator.cuh +++ b/cub/cub/iterator/tex_obj_input_iterator.cuh @@ -160,7 +160,7 @@ private: public: /// Constructor _CCCL_HOST_DEVICE _CCCL_FORCEINLINE TexObjInputIterator() - : ptr(NULL) + : ptr(nullptr) , tex_offset(0) , tex_obj(0) {} @@ -193,7 +193,7 @@ public: res_desc.res.linear.desc = channel_desc; res_desc.res.linear.sizeInBytes = bytes; tex_desc.readMode = cudaReadModeElementType; - return CubDebug(cudaCreateTextureObject(&tex_obj, &res_desc, &tex_desc, NULL)); + return CubDebug(cudaCreateTextureObject(&tex_obj, &res_desc, &tex_desc, nullptr)); } /// Unbind this iterator from its texture reference diff --git a/cub/cub/util_allocator.cuh b/cub/cub/util_allocator.cuh index 5af3bf9bb7c..776bf650d24 100644 --- a/cub/cub/util_allocator.cuh +++ b/cub/cub/util_allocator.cuh @@ -155,7 +155,7 @@ struct CachingDeviceAllocator // Constructor (suitable for searching maps for a range of suitable blocks, given a device) BlockDescriptor(int device) - : d_ptr(NULL) + : d_ptr(nullptr) , bytes(0) , bin(INVALID_BIN) , device(device) @@ -453,7 +453,7 @@ struct CachingDeviceAllocator */ cudaError_t DeviceAllocate(int device, void** d_ptr, size_t bytes, cudaStream_t active_stream = 0) { - *d_ptr = NULL; + *d_ptr = nullptr; int entrypoint_device = INVALID_DEVICE_ORDINAL; cudaError_t error = cudaSuccess; diff --git a/cub/cub/util_temporary_storage.cuh b/cub/cub/util_temporary_storage.cuh index c97c96a5f02..6563664484d 100644 --- a/cub/cub/util_temporary_storage.cuh +++ b/cub/cub/util_temporary_storage.cuh @@ -56,7 +56,7 @@ CUB_NAMESPACE_BEGIN * * @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,out] temp_storage_bytes diff --git a/cub/cub/util_type.cuh b/cub/cub/util_type.cuh index e9c30f7d4a6..c4145c4801a 100644 --- a/cub/cub/util_type.cuh +++ b/cub/cub/util_type.cuh @@ -262,7 +262,7 @@ struct CUB_DEPRECATED RemoveQualifiers #ifndef DOXYGEN_SHOULD_SKIP_THIS // Do not document /** - * \brief A simple "NULL" marker type + * \brief A simple "null" marker type */ struct NullType { @@ -907,8 +907,8 @@ struct DoubleBuffer _CCCL_HOST_DEVICE _CCCL_FORCEINLINE DoubleBuffer() { selector = 0; - d_buffers[0] = NULL; - d_buffers[1] = NULL; + d_buffers[0] = nullptr; + d_buffers[1] = nullptr; } /// \brief Constructor @@ -1023,7 +1023,7 @@ private: public: /// Whether the functor BinaryOp has a third unsigned int index param - static constexpr bool HAS_PARAM = sizeof(Test(NULL)) == sizeof(char); + static constexpr bool HAS_PARAM = sizeof(Test(nullptr)) == sizeof(char); }; /****************************************************************************** diff --git a/cub/examples/block/example_block_radix_sort.cu b/cub/examples/block/example_block_radix_sort.cu index 404e3790a7d..1070b54a3b2 100644 --- a/cub/examples/block/example_block_radix_sort.cu +++ b/cub/examples/block/example_block_radix_sort.cu @@ -174,9 +174,9 @@ void Test() Initialize(h_in, h_reference, TILE_SIZE * g_grid_size, TILE_SIZE); // Initialize device arrays - Key* d_in = NULL; - Key* d_out = NULL; - clock_t* d_elapsed = NULL; + Key* d_in = nullptr; + Key* d_out = nullptr; + clock_t* d_elapsed = nullptr; CubDebugExit(cudaMalloc((void**) &d_in, sizeof(Key) * TILE_SIZE * g_grid_size)); CubDebugExit(cudaMalloc((void**) &d_out, sizeof(Key) * TILE_SIZE * g_grid_size)); CubDebugExit(cudaMalloc((void**) &d_elapsed, sizeof(clock_t) * g_grid_size)); diff --git a/cub/examples/block/example_block_reduce.cu b/cub/examples/block/example_block_reduce.cu index 92c1ce1ec64..295a5d82c21 100644 --- a/cub/examples/block/example_block_reduce.cu +++ b/cub/examples/block/example_block_reduce.cu @@ -139,9 +139,9 @@ void Test() int h_aggregate = Initialize(h_in, TILE_SIZE); // Initialize device arrays - int* d_in = NULL; - int* d_out = NULL; - clock_t* d_elapsed = NULL; + int* d_in = nullptr; + int* d_out = nullptr; + clock_t* d_elapsed = nullptr; cudaMalloc((void**) &d_in, sizeof(int) * TILE_SIZE); cudaMalloc((void**) &d_out, sizeof(int) * 1); cudaMalloc((void**) &d_elapsed, sizeof(clock_t)); diff --git a/cub/examples/block/example_block_reduce_dyn_smem.cu b/cub/examples/block/example_block_reduce_dyn_smem.cu index 5f60a7dc1ef..057884cdc56 100644 --- a/cub/examples/block/example_block_reduce_dyn_smem.cu +++ b/cub/examples/block/example_block_reduce_dyn_smem.cu @@ -144,8 +144,8 @@ void Test() int h_aggregate = Initialize(h_in, BLOCK_THREADS); // Initialize device arrays - int* d_in = NULL; - int* d_out = NULL; + int* d_in = nullptr; + int* d_out = nullptr; cudaMalloc((void**) &d_in, sizeof(int) * BLOCK_THREADS); cudaMalloc((void**) &d_out, sizeof(int) * BLOCK_THREADS); @@ -170,7 +170,7 @@ void Test() auto smem_size = (std::max)(1 * sizeof(int), block_reduce_temp_bytes); // use default stream - cudaStream_t stream = NULL; + cudaStream_t stream = nullptr; // Run reduction kernel BlockReduceKernel<<>>(d_in, d_out); diff --git a/cub/examples/block/example_block_scan.cu b/cub/examples/block/example_block_scan.cu index 78b400f053d..f7a58b20de9 100644 --- a/cub/examples/block/example_block_scan.cu +++ b/cub/examples/block/example_block_scan.cu @@ -167,9 +167,9 @@ void Test() int h_aggregate = Initialize(h_in, h_reference, TILE_SIZE); // Initialize device arrays - int* d_in = NULL; - int* d_out = NULL; - clock_t* d_elapsed = NULL; + int* d_in = nullptr; + int* d_out = nullptr; + clock_t* d_elapsed = nullptr; cudaMalloc((void**) &d_in, sizeof(int) * TILE_SIZE); cudaMalloc((void**) &d_out, sizeof(int) * (TILE_SIZE + 1)); cudaMalloc((void**) &d_elapsed, sizeof(clock_t)); diff --git a/cub/examples/device/example_device_partition_flagged.cu b/cub/examples/device/example_device_partition_flagged.cu index 03402e72cb7..1a5acd38337 100644 --- a/cub/examples/device/example_device_partition_flagged.cu +++ b/cub/examples/device/example_device_partition_flagged.cu @@ -173,8 +173,8 @@ int main(int argc, char** argv) fflush(stdout); // Allocate problem device arrays - int* d_in = NULL; - unsigned char* d_flags = NULL; + int* d_in = nullptr; + unsigned char* d_flags = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_in, sizeof(int) * num_items)); CubDebugExit(g_allocator.DeviceAllocate((void**) &d_flags, sizeof(unsigned char) * num_items)); @@ -184,13 +184,13 @@ int main(int argc, char** argv) CubDebugExit(cudaMemcpy(d_flags, h_flags, sizeof(unsigned char) * num_items, cudaMemcpyHostToDevice)); // Allocate device output array and num selected - int* d_out = NULL; - int* d_num_selected_out = NULL; + int* d_out = nullptr; + int* d_num_selected_out = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_out, sizeof(int) * num_items)); CubDebugExit(g_allocator.DeviceAllocate((void**) &d_num_selected_out, sizeof(int))); // Allocate temporary storage - void* d_temp_storage = NULL; + void* d_temp_storage = nullptr; size_t temp_storage_bytes = 0; CubDebugExit( DevicePartition::Flagged(d_temp_storage, temp_storage_bytes, d_in, d_flags, d_out, d_num_selected_out, num_items)); diff --git a/cub/examples/device/example_device_partition_if.cu b/cub/examples/device/example_device_partition_if.cu index 43c4af3b48a..5b72f1c86c9 100644 --- a/cub/examples/device/example_device_partition_if.cu +++ b/cub/examples/device/example_device_partition_if.cu @@ -192,20 +192,20 @@ int main(int argc, char** argv) fflush(stdout); // Allocate problem device arrays - int* d_in = NULL; + int* d_in = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_in, sizeof(int) * num_items)); // Initialize device input CubDebugExit(cudaMemcpy(d_in, h_in, sizeof(int) * num_items, cudaMemcpyHostToDevice)); // Allocate device output array and num selected - int* d_out = NULL; - int* d_num_selected_out = NULL; + int* d_out = nullptr; + int* d_num_selected_out = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_out, sizeof(int) * num_items)); CubDebugExit(g_allocator.DeviceAllocate((void**) &d_num_selected_out, sizeof(int))); // Allocate temporary storage - void* d_temp_storage = NULL; + void* d_temp_storage = nullptr; size_t temp_storage_bytes = 0; CubDebugExit( DevicePartition::If(d_temp_storage, temp_storage_bytes, d_in, d_out, d_num_selected_out, num_items, select_op)); diff --git a/cub/examples/device/example_device_radix_sort.cu b/cub/examples/device/example_device_radix_sort.cu index 89fd96c4be0..ea4413b54cb 100644 --- a/cub/examples/device/example_device_radix_sort.cu +++ b/cub/examples/device/example_device_radix_sort.cu @@ -168,7 +168,7 @@ int main(int argc, char** argv) // Allocate temporary storage size_t temp_storage_bytes = 0; - void* d_temp_storage = NULL; + void* d_temp_storage = nullptr; CubDebugExit(DeviceRadixSort::SortPairs(d_temp_storage, temp_storage_bytes, d_keys, d_values, num_items)); CubDebugExit(g_allocator.DeviceAllocate(&d_temp_storage, temp_storage_bytes)); diff --git a/cub/examples/device/example_device_reduce.cu b/cub/examples/device/example_device_reduce.cu index 4dd3a6fcb0d..004d8e9bfdc 100644 --- a/cub/examples/device/example_device_reduce.cu +++ b/cub/examples/device/example_device_reduce.cu @@ -137,18 +137,18 @@ int main(int argc, char** argv) Solve(h_in, h_reference, num_items); // Allocate problem device arrays - int* d_in = NULL; + int* d_in = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_in, sizeof(int) * num_items)); // Initialize device input CubDebugExit(cudaMemcpy(d_in, h_in, sizeof(int) * num_items, cudaMemcpyHostToDevice)); // Allocate device output array - int* d_out = NULL; + int* d_out = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_out, sizeof(int) * 1)); // Request and allocate temporary storage - void* d_temp_storage = NULL; + void* d_temp_storage = nullptr; size_t temp_storage_bytes = 0; CubDebugExit(DeviceReduce::Sum(d_temp_storage, temp_storage_bytes, d_in, d_out, num_items)); CubDebugExit(g_allocator.DeviceAllocate(&d_temp_storage, temp_storage_bytes)); diff --git a/cub/examples/device/example_device_scan.cu b/cub/examples/device/example_device_scan.cu index cfce8c7bdb9..58d05cbbb1f 100644 --- a/cub/examples/device/example_device_scan.cu +++ b/cub/examples/device/example_device_scan.cu @@ -137,18 +137,18 @@ int main(int argc, char** argv) Solve(h_in, h_reference, num_items); // Allocate problem device arrays - int* d_in = NULL; + int* d_in = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_in, sizeof(int) * num_items)); // Initialize device input CubDebugExit(cudaMemcpy(d_in, h_in, sizeof(int) * num_items, cudaMemcpyHostToDevice)); // Allocate device output array - int* d_out = NULL; + int* d_out = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_out, sizeof(int) * num_items)); // Allocate temporary storage - void* d_temp_storage = NULL; + void* d_temp_storage = nullptr; size_t temp_storage_bytes = 0; CubDebugExit(DeviceScan::ExclusiveSum(d_temp_storage, temp_storage_bytes, d_in, d_out, num_items)); CubDebugExit(g_allocator.DeviceAllocate(&d_temp_storage, temp_storage_bytes)); diff --git a/cub/examples/device/example_device_select_flagged.cu b/cub/examples/device/example_device_select_flagged.cu index fba269129b5..c7d408f48bb 100644 --- a/cub/examples/device/example_device_select_flagged.cu +++ b/cub/examples/device/example_device_select_flagged.cu @@ -173,8 +173,8 @@ int main(int argc, char** argv) fflush(stdout); // Allocate problem device arrays - int* d_in = NULL; - unsigned char* d_flags = NULL; + int* d_in = nullptr; + unsigned char* d_flags = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_in, sizeof(int) * num_items)); CubDebugExit(g_allocator.DeviceAllocate((void**) &d_flags, sizeof(unsigned char) * num_items)); @@ -184,13 +184,13 @@ int main(int argc, char** argv) CubDebugExit(cudaMemcpy(d_flags, h_flags, sizeof(unsigned char) * num_items, cudaMemcpyHostToDevice)); // Allocate device output array and num selected - int* d_out = NULL; - int* d_num_selected_out = NULL; + int* d_out = nullptr; + int* d_num_selected_out = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_out, sizeof(int) * num_items)); CubDebugExit(g_allocator.DeviceAllocate((void**) &d_num_selected_out, sizeof(int))); // Allocate temporary storage - void* d_temp_storage = NULL; + void* d_temp_storage = nullptr; size_t temp_storage_bytes = 0; CubDebugExit( DeviceSelect::Flagged(d_temp_storage, temp_storage_bytes, d_in, d_flags, d_out, d_num_selected_out, num_items)); diff --git a/cub/examples/device/example_device_select_if.cu b/cub/examples/device/example_device_select_if.cu index 6a8728dbc10..7feefcfa290 100644 --- a/cub/examples/device/example_device_select_if.cu +++ b/cub/examples/device/example_device_select_if.cu @@ -192,20 +192,20 @@ int main(int argc, char** argv) fflush(stdout); // Allocate problem device arrays - int* d_in = NULL; + int* d_in = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_in, sizeof(int) * num_items)); // Initialize device input CubDebugExit(cudaMemcpy(d_in, h_in, sizeof(int) * num_items, cudaMemcpyHostToDevice)); // Allocate device output array and num selected - int* d_out = NULL; - int* d_num_selected_out = NULL; + int* d_out = nullptr; + int* d_num_selected_out = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_out, sizeof(int) * num_items)); CubDebugExit(g_allocator.DeviceAllocate((void**) &d_num_selected_out, sizeof(int))); // Allocate temporary storage - void* d_temp_storage = NULL; + void* d_temp_storage = nullptr; size_t temp_storage_bytes = 0; CubDebugExit( DeviceSelect::If(d_temp_storage, temp_storage_bytes, d_in, d_out, d_num_selected_out, num_items, select_op)); diff --git a/cub/examples/device/example_device_select_unique.cu b/cub/examples/device/example_device_select_unique.cu index 8fb6a5caf1a..5db10045bb4 100644 --- a/cub/examples/device/example_device_select_unique.cu +++ b/cub/examples/device/example_device_select_unique.cu @@ -168,20 +168,20 @@ int main(int argc, char** argv) fflush(stdout); // Allocate problem device arrays - int* d_in = NULL; + int* d_in = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_in, sizeof(int) * num_items)); // Initialize device input CubDebugExit(cudaMemcpy(d_in, h_in, sizeof(int) * num_items, cudaMemcpyHostToDevice)); // Allocate device output array and num selected - int* d_out = NULL; - int* d_num_selected_out = NULL; + int* d_out = nullptr; + int* d_num_selected_out = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_out, sizeof(int) * num_items)); CubDebugExit(g_allocator.DeviceAllocate((void**) &d_num_selected_out, sizeof(int))); // Allocate temporary storage - void* d_temp_storage = NULL; + void* d_temp_storage = nullptr; size_t temp_storage_bytes = 0; CubDebugExit(DeviceSelect::Unique(d_temp_storage, temp_storage_bytes, d_in, d_out, d_num_selected_out, num_items)); CubDebugExit(g_allocator.DeviceAllocate(&d_temp_storage, temp_storage_bytes)); diff --git a/cub/examples/device/example_device_sort_find_non_trivial_runs.cu b/cub/examples/device/example_device_sort_find_non_trivial_runs.cu index fc95c51a1c9..e00d45c2e6a 100644 --- a/cub/examples/device/example_device_sort_find_non_trivial_runs.cu +++ b/cub/examples/device/example_device_sort_find_non_trivial_runs.cu @@ -259,7 +259,7 @@ int main(int argc, char** argv) // Allocate temporary storage for sorting size_t temp_storage_bytes = 0; - void* d_temp_storage = NULL; + void* d_temp_storage = nullptr; CubDebugExit(DeviceRadixSort::SortPairs(d_temp_storage, temp_storage_bytes, d_keys, d_values, num_items)); CubDebugExit(g_allocator.DeviceAllocate(&d_temp_storage, temp_storage_bytes)); @@ -284,15 +284,15 @@ int main(int argc, char** argv) gpu_rle_timer.Start(); // Allocate device arrays for enumerating non-trivial runs - int* d_offests_out = NULL; - int* d_lengths_out = NULL; - int* d_num_runs = NULL; + int* d_offests_out = nullptr; + int* d_lengths_out = nullptr; + int* d_num_runs = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_offests_out, sizeof(int) * num_items)); CubDebugExit(g_allocator.DeviceAllocate((void**) &d_lengths_out, sizeof(int) * num_items)); CubDebugExit(g_allocator.DeviceAllocate((void**) &d_num_runs, sizeof(int) * 1)); // Allocate temporary storage for isolating non-trivial runs - d_temp_storage = NULL; + d_temp_storage = nullptr; CubDebugExit(DeviceRunLengthEncode::NonTrivialRuns( d_temp_storage, temp_storage_bytes, diff --git a/cub/test/catch2_test_vsmem.cu b/cub/test/catch2_test_vsmem.cu index b8d3f7e37be..a4ae61e08c3 100644 --- a/cub/test/catch2_test_vsmem.cu +++ b/cub/test/catch2_test_vsmem.cu @@ -223,7 +223,7 @@ struct dispatch_dummy_algorithm_t : SelectedPolicy { using item_t = cub::detail::value_t; - /// Device-accessible allocation of temporary storage. When NULL, the required + /// Device-accessible allocation of temporary storage. When nullptr, the required /// allocation size is written to \p temp_storage_bytes and no work is done. void* d_temp_storage; diff --git a/cub/test/link_a.cu b/cub/test/link_a.cu index e47a64693f5..65e20189a94 100644 --- a/cub/test/link_a.cu +++ b/cub/test/link_a.cu @@ -7,5 +7,5 @@ void a() cub::DoubleBuffer d_keys; cub::DoubleBuffer d_values; size_t temp_storage_bytes = 0; - cub::DeviceRadixSort::SortPairs(NULL, temp_storage_bytes, d_keys, d_values, 1024); + cub::DeviceRadixSort::SortPairs(nullptr, temp_storage_bytes, d_keys, d_values, 1024); } diff --git a/cub/test/link_b.cu b/cub/test/link_b.cu index 0ef00a171b5..be174fd244e 100644 --- a/cub/test/link_b.cu +++ b/cub/test/link_b.cu @@ -7,5 +7,5 @@ void b() cub::DoubleBuffer d_keys; cub::DoubleBuffer d_values; size_t temp_storage_bytes = 0; - cub::DeviceRadixSort::SortPairs(NULL, temp_storage_bytes, d_keys, d_values, 1024); + cub::DeviceRadixSort::SortPairs(nullptr, temp_storage_bytes, d_keys, d_values, 1024); } diff --git a/cub/test/test_allocator.cu b/cub/test/test_allocator.cu index af942a2c3b7..4b4723fe997 100644 --- a/cub/test/test_allocator.cu +++ b/cub/test/test_allocator.cu @@ -382,7 +382,7 @@ int main(int argc, char** argv) // CPU performance comparisons vs cached. Allocate and free a 1MB block 2000 times CpuTimer cpu_timer; - char* d_1024MB = NULL; + char* d_1024MB = nullptr; allocator.debug = false; // Prime the caching allocator and the kernel diff --git a/cub/test/test_iterator.cu b/cub/test/test_iterator.cu index 8782a615371..90ba4b4077f 100644 --- a/cub/test/test_iterator.cu +++ b/cub/test/test_iterator.cu @@ -127,8 +127,8 @@ template void Test(InputIteratorT d_in, T (&h_reference)[TEST_VALUES]) { // Allocate device arrays - T* d_out = NULL; - InputIteratorT* d_itrs = NULL; + T* d_out = nullptr; + InputIteratorT* d_itrs = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_out, sizeof(T) * TEST_VALUES)); CubDebugExit(g_allocator.DeviceAllocate((void**) &d_itrs, sizeof(InputIteratorT) * 2)); @@ -235,7 +235,7 @@ void TestModified() } // Allocate device arrays - T* d_data = NULL; + T* d_data = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_data, sizeof(T) * TEST_VALUES)); CubDebugExit(cudaMemcpy(d_data, h_data, sizeof(T) * TEST_VALUES, cudaMemcpyHostToDevice)); @@ -290,7 +290,7 @@ void TestTransform() } // Allocate device arrays - T* d_data = NULL; + T* d_data = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_data, sizeof(T) * TEST_VALUES)); CubDebugExit(cudaMemcpy(d_data, h_data, sizeof(T) * TEST_VALUES, cudaMemcpyHostToDevice)); @@ -344,8 +344,8 @@ void TestTexObj() } // Allocate device arrays - T* d_data = NULL; - T* d_dummy = NULL; + T* d_data = nullptr; + T* d_dummy = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_data, sizeof(T) * TEST_VALUES)); CubDebugExit(cudaMemcpy(d_data, h_data, sizeof(T) * TEST_VALUES, cudaMemcpyHostToDevice)); @@ -405,7 +405,7 @@ void TestTexTransform() } // Allocate device arrays - T* d_data = NULL; + T* d_data = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_data, sizeof(T) * TEST_VALUES)); CubDebugExit(cudaMemcpy(d_data, h_data, sizeof(T) * TEST_VALUES, cudaMemcpyHostToDevice)); diff --git a/cub/test/test_iterator_deprecated.cu b/cub/test/test_iterator_deprecated.cu index 72f073ad19d..f6a8e3814ad 100644 --- a/cub/test/test_iterator_deprecated.cu +++ b/cub/test/test_iterator_deprecated.cu @@ -96,8 +96,8 @@ template void Test(InputIteratorT d_in, T (&h_reference)[TEST_VALUES]) { // Allocate device arrays - T* d_out = NULL; - InputIteratorT* d_itrs = NULL; + T* d_out = nullptr; + InputIteratorT* d_itrs = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_out, sizeof(T) * TEST_VALUES)); CubDebugExit(g_allocator.DeviceAllocate((void**) &d_itrs, sizeof(InputIteratorT) * 2)); @@ -160,8 +160,8 @@ void TestTexRef() } // Allocate device arrays - T* d_data = NULL; - T* d_dummy = NULL; + T* d_data = nullptr; + T* d_dummy = nullptr; CubDebugExit(g_allocator.DeviceAllocate((void**) &d_data, sizeof(T) * TEST_VALUES)); CubDebugExit(cudaMemcpy(d_data, h_data, sizeof(T) * TEST_VALUES, cudaMemcpyHostToDevice)); diff --git a/libcudacxx/examples/rtc_example.cpp b/libcudacxx/examples/rtc_example.cpp index 513e5805848..bc5f51eb784 100644 --- a/libcudacxx/examples/rtc_example.cpp +++ b/libcudacxx/examples/rtc_example.cpp @@ -126,8 +126,8 @@ int main(int argc, char* argv[]) trie, // buffer "trie.cu", // name 0, // numHeaders - NULL, // headers - NULL)); // includeNames + nullptr, // headers + nullptr)); // includeNames const char* opts[] = { "-std=c++11", diff --git a/libcudacxx/include/cuda/std/detail/libcxx/include/__threading_support b/libcudacxx/include/cuda/std/detail/libcxx/include/__threading_support index 9f5fbe92555..a7c4de867fb 100644 --- a/libcudacxx/include/cuda/std/detail/libcxx/include/__threading_support +++ b/libcudacxx/include/cuda/std/detail/libcxx/include/__threading_support @@ -433,7 +433,7 @@ int __libcpp_condvar_destroy(__libcpp_condvar_t* __cv) bool __libcpp_semaphore_init(__libcpp_semaphore_t* __sem, int __init) { - return (*__sem = dispatch_semaphore_create(__init)) != NULL; + return (*__sem = dispatch_semaphore_create(__init)) != nullptr; } bool __libcpp_semaphore_destroy(__libcpp_semaphore_t* __sem) diff --git a/libcudacxx/test/libcudacxx/force_include.h b/libcudacxx/test/libcudacxx/force_include.h index 1c4ca961f67..93752019a17 100644 --- a/libcudacxx/test/libcudacxx/force_include.h +++ b/libcudacxx/test/libcudacxx/force_include.h @@ -53,7 +53,7 @@ int cuda_cluster_size = 1; __global__ void fake_main_kernel(int* ret) { - *ret = fake_main(0, NULL); + *ret = fake_main(0, nullptr); } #define CUDA_CALL(err, ...) \ diff --git a/libcudacxx/test/support/asan_testing.h b/libcudacxx/test/support/asan_testing.h index 043033a49cf..4f41b621536 100644 --- a/libcudacxx/test/support/asan_testing.h +++ b/libcudacxx/test/support/asan_testing.h @@ -17,7 +17,7 @@ extern "C" int __sanitizer_verify_contiguous_container(const void* beg, const vo template bool is_contiguous_container_asan_correct(const std::vector& c) { - if (std::is_same>::value && c.data() != NULL) + if (std::is_same>::value && c.data() != nullptr) { return __sanitizer_verify_contiguous_container(c.data(), c.data() + c.size(), c.data() + c.capacity()) != 0; } diff --git a/libcudacxx/test/support/rapid-cxx-test.h b/libcudacxx/test/support/rapid-cxx-test.h index 5b98987cf64..61f31785fa0 100644 --- a/libcudacxx/test/support/rapid-cxx-test.h +++ b/libcudacxx/test/support/rapid-cxx-test.h @@ -539,7 +539,7 @@ struct test_case : file("") , func("") , line(0) - , invoke(NULL) + , invoke(nullptr) {} test_case(const char* file1, const char* func1, std::size_t line1, invoker_t invoke1) diff --git a/libcudacxx/test/utils/nvidia/nvrtc/nvrtcc_build.h b/libcudacxx/test/utils/nvidia/nvrtc/nvrtcc_build.h index cd059bd3d6c..0d2d2a7bf0d 100644 --- a/libcudacxx/test/utils/nvidia/nvrtc/nvrtcc_build.h +++ b/libcudacxx/test/utils/nvidia/nvrtc/nvrtcc_build.h @@ -55,7 +55,7 @@ GpuProg nvrtc_build_prog(const std::string& testCu, const ArchConfig& config, co fprintf(stderr, "Compiling program...\r\n"); nvrtcProgram prog; - NVRTC_SAFE_CALL(nvrtcCreateProgram(&prog, testCu.c_str(), "test.cu", 0, NULL, NULL)); + NVRTC_SAFE_CALL(nvrtcCreateProgram(&prog, testCu.c_str(), "test.cu", 0, nullptr, nullptr)); nvrtcResult compile_result = nvrtcCompileProgram(prog, optList.size(), optList.data()); diff --git a/libcudacxx/test/utils/nvidia/nvrtc/nvrtcc_common.h b/libcudacxx/test/utils/nvidia/nvrtc/nvrtcc_common.h index b48f2f56578..f4df7674e25 100644 --- a/libcudacxx/test/utils/nvidia/nvrtc/nvrtcc_common.h +++ b/libcudacxx/test/utils/nvidia/nvrtc/nvrtcc_common.h @@ -119,6 +119,6 @@ __host__ __device__ int fake_main(int argc, char ** argv); // extern "C" to stop the name from being mangled extern "C" __global__ void main_kernel() { - fake_main(0, NULL); + fake_main(0, nullptr); } )program"; diff --git a/libcudacxx/test/utils/nvidia/nvrtc/nvrtcc_run.h b/libcudacxx/test/utils/nvidia/nvrtc/nvrtcc_run.h index 679974c5b16..c8e49296fd8 100644 --- a/libcudacxx/test/utils/nvidia/nvrtc/nvrtcc_run.h +++ b/libcudacxx/test/utils/nvidia/nvrtc/nvrtcc_run.h @@ -68,7 +68,7 @@ static void load_and_run_gpu_code(const std::string inputFile, const RunConfig& CUDA_SAFE_CALL(cuCtxCreate(&context, 0, cuDevice)); CUDA_SAFE_CALL(cuModuleLoadDataEx(&module, code.data(), 0, 0, 0)); CUDA_SAFE_CALL(cuModuleGetFunction(&kernel, module, "main_kernel")); - CUDA_SAFE_CALL(cuLaunchKernel(kernel, 1, 1, 1, rc.threadCount, 1, 1, rc.shmemSize, NULL, NULL, 0)); + CUDA_SAFE_CALL(cuLaunchKernel(kernel, 1, 1, 1, rc.threadCount, 1, 1, rc.shmemSize, nullptr, nullptr, 0)); CUDA_API_CALL(cudaGetLastError()); CUDA_API_CALL(cudaDeviceSynchronize()); diff --git a/libcudacxx/test/utils/nvidia/nvrtc/utils/platform.win.h b/libcudacxx/test/utils/nvidia/nvrtc/utils/platform.win.h index 8e4453c107d..744d6293152 100644 --- a/libcudacxx/test/utils/nvidia/nvrtc/utils/platform.win.h +++ b/libcudacxx/test/utils/nvidia/nvrtc/utils/platform.win.h @@ -24,7 +24,8 @@ static void platform_exec(char const* process, char** args, size_t nargs) printf("Running command: %s\r\n", cl.data()); - bool exec_result = CreateProcess(NULL, (LPSTR) cl.data(), NULL, NULL, false, false, NULL, NULL, &si, &pi); + bool exec_result = + CreateProcess(nullptr, (LPSTR) cl.data(), nullptr, nullptr, false, false, nullptr, nullptr, &si, &pi); if (!exec_result) { diff --git a/thrust/testing/mr_disjoint_pool.cu b/thrust/testing/mr_disjoint_pool.cu index 22fa835ea0b..2102038ed90 100644 --- a/thrust/testing/mr_disjoint_pool.cu +++ b/thrust/testing/mr_disjoint_pool.cu @@ -268,7 +268,7 @@ void TestDisjointGlobalPool() { typedef PoolTemplate Pool; - ASSERT_EQUAL(thrust::mr::get_global_resource() != NULL, true); + ASSERT_EQUAL(thrust::mr::get_global_resource() != nullptr, true); } void TestUnsynchronizedDisjointGlobalPool() diff --git a/thrust/testing/mr_pool.cu b/thrust/testing/mr_pool.cu index 8483323eda3..9f3008ddf79 100644 --- a/thrust/testing/mr_pool.cu +++ b/thrust/testing/mr_pool.cu @@ -38,7 +38,7 @@ struct tracked_pointer std::size_t offset; void* ptr; - _CCCL_HOST_DEVICE explicit tracked_pointer(T* ptr = NULL) + _CCCL_HOST_DEVICE explicit tracked_pointer(T* ptr = nullptr) : id() , size() , alignment() @@ -358,7 +358,7 @@ void TestGlobalPool() { typedef PoolTemplate Pool; - ASSERT_EQUAL(thrust::mr::get_global_resource() != NULL, true); + ASSERT_EQUAL(thrust::mr::get_global_resource() != nullptr, true); } void TestUnsynchronizedGlobalPool() diff --git a/thrust/testing/regression/nvbug_1990211__scan_requires_assignability_from_zero.cu b/thrust/testing/regression/nvbug_1990211__scan_requires_assignability_from_zero.cu index 01a5d0925b8..4d934912455 100644 --- a/thrust/testing/regression/nvbug_1990211__scan_requires_assignability_from_zero.cu +++ b/thrust/testing/regression/nvbug_1990211__scan_requires_assignability_from_zero.cu @@ -9,7 +9,7 @@ inline _CCCL_HOST_DEVICE uint2 operator+(uint2 a, uint2 b) int main() { int num_elements = 32; - uint2 *input = NULL, *output = NULL; + uint2 *input = nullptr, *output = nullptr; const uint2 zero = make_uint2(0, 0); thrust::exclusive_scan( diff --git a/thrust/testing/regression/nvbug_1990211__scan_requires_assignability_from_zero.fixed0.cu b/thrust/testing/regression/nvbug_1990211__scan_requires_assignability_from_zero.fixed0.cu index 246b60f337b..3f9cb10c447 100644 --- a/thrust/testing/regression/nvbug_1990211__scan_requires_assignability_from_zero.fixed0.cu +++ b/thrust/testing/regression/nvbug_1990211__scan_requires_assignability_from_zero.fixed0.cu @@ -12,7 +12,7 @@ struct uint2_adder int main() { int num_elements = 32; - uint2 *input = NULL, *output = NULL; + uint2 *input = nullptr, *output = nullptr; const uint2 zero = make_uint2(0, 0); thrust::exclusive_scan( diff --git a/thrust/testing/regression/nvbug_1990211__scan_requires_assignability_from_zero.fixed1.cu b/thrust/testing/regression/nvbug_1990211__scan_requires_assignability_from_zero.fixed1.cu index 5396e718dd0..35cdad51560 100644 --- a/thrust/testing/regression/nvbug_1990211__scan_requires_assignability_from_zero.fixed1.cu +++ b/thrust/testing/regression/nvbug_1990211__scan_requires_assignability_from_zero.fixed1.cu @@ -9,7 +9,7 @@ inline _CCCL_HOST_DEVICE uint2 operator+(uint2 a, uint2 b) int main() { int num_elements = 32; - uint2 *input = NULL, *output = NULL; + uint2 *input = nullptr, *output = nullptr; const uint2 zero = make_uint2(0, 0); thrust::exclusive_scan( diff --git a/thrust/testing/unittest/runtime_static_assert.h b/thrust/testing/unittest/runtime_static_assert.h index e519ecccc3f..33446462e6b 100644 --- a/thrust/testing/unittest/runtime_static_assert.h +++ b/thrust/testing/unittest/runtime_static_assert.h @@ -41,7 +41,7 @@ _CCCL_HOST_DEVICE void assert_static(bool condition, const char* filename, int l triggered = static_cast(*device_ptr).triggered; \ } \ thrust::device_free(device_ptr); \ - raw_ptr = NULL; \ + raw_ptr = nullptr; \ ::cudaMemcpyToSymbol(unittest::detail::device_exception, &raw_ptr, sizeof(ex_t*)); \ if (!triggered) \ { \ @@ -100,7 +100,7 @@ namespace detail #if defined(_CCCL_COMPILER_GCC) || defined(_CCCL_COMPILER_CLANG) __attribute__((used)) #endif -_CCCL_DEVICE static static_assert_exception* device_exception = NULL; +_CCCL_DEVICE static static_assert_exception* device_exception = nullptr; } // namespace detail _CCCL_HOST_DEVICE void assert_static(bool condition, const char* filename, int lineno) diff --git a/thrust/thrust/mr/disjoint_tls_pool.h b/thrust/thrust/mr/disjoint_tls_pool.h index 186959c82e4..4b5f5a34d1b 100644 --- a/thrust/thrust/mr/disjoint_tls_pool.h +++ b/thrust/thrust/mr/disjoint_tls_pool.h @@ -52,7 +52,7 @@ namespace mr */ template _CCCL_HOST thrust::mr::disjoint_unsynchronized_pool_resource& -tls_disjoint_pool(Upstream* upstream = NULL, Bookkeeper* bookkeeper = NULL) +tls_disjoint_pool(Upstream* upstream = nullptr, Bookkeeper* bookkeeper = nullptr) { static thread_local auto adaptor = [&] { assert(upstream && bookkeeper); diff --git a/thrust/thrust/mr/tls_pool.h b/thrust/thrust/mr/tls_pool.h index e87676e7068..44f00ff9925 100644 --- a/thrust/thrust/mr/tls_pool.h +++ b/thrust/thrust/mr/tls_pool.h @@ -49,7 +49,7 @@ namespace mr * \param upstream the argument to the constructor, if invoked */ template -_CCCL_HOST thrust::mr::unsynchronized_pool_resource& tls_pool(Upstream* upstream = NULL) +_CCCL_HOST thrust::mr::unsynchronized_pool_resource& tls_pool(Upstream* upstream = nullptr) { static thread_local auto adaptor = [&] { assert(upstream); diff --git a/thrust/thrust/system/cuda/detail/core/agent_launcher.h b/thrust/thrust/system/cuda/detail/core/agent_launcher.h index b24d1602b25..14aceec5ab7 100644 --- a/thrust/thrust/system/cuda/detail/core/agent_launcher.h +++ b/thrust/thrust/system/cuda/detail/core/agent_launcher.h @@ -259,7 +259,7 @@ THRUST_DETAIL_KERNEL_ATTRIBUTES void __launch_bounds__(Agent::ptx_plan::BLOCK_TH _kernel_agent_vshmem(char* vshmem, Args... args) { extern __shared__ char shmem[]; - vshmem = vshmem == NULL ? shmem : vshmem + blockIdx.x * temp_storage_size::value; + vshmem = vshmem == nullptr ? shmem : vshmem + blockIdx.x * temp_storage_size::value; Agent::entry(args..., vshmem); } # else @@ -268,7 +268,7 @@ THRUST_DETAIL_KERNEL_ATTRIBUTES void __launch_bounds__(Agent::ptx_plan::BLOCK_TH _kernel_agent_vshmem(char* vshmem, _0 x0) { extern __shared__ char shmem[]; - vshmem = vshmem == NULL ? shmem : vshmem + blockIdx.x * temp_storage_size::value; + vshmem = vshmem == nullptr ? shmem : vshmem + blockIdx.x * temp_storage_size::value; Agent::entry(x0, vshmem); } template @@ -276,7 +276,7 @@ THRUST_DETAIL_KERNEL_ATTRIBUTES void __launch_bounds__(Agent::ptx_plan::BLOCK_TH _kernel_agent_vshmem(char* vshmem, _0 x0, _1 x1) { extern __shared__ char shmem[]; - vshmem = vshmem == NULL ? shmem : vshmem + blockIdx.x * temp_storage_size::value; + vshmem = vshmem == nullptr ? shmem : vshmem + blockIdx.x * temp_storage_size::value; Agent::entry(x0, x1, vshmem); } template @@ -284,7 +284,7 @@ THRUST_DETAIL_KERNEL_ATTRIBUTES void __launch_bounds__(Agent::ptx_plan::BLOCK_TH _kernel_agent_vshmem(char* vshmem, _0 x0, _1 x1, _2 x2) { extern __shared__ char shmem[]; - vshmem = vshmem == NULL ? shmem : vshmem + blockIdx.x * temp_storage_size::value; + vshmem = vshmem == nullptr ? shmem : vshmem + blockIdx.x * temp_storage_size::value; Agent::entry(x0, x1, x2, vshmem); } template @@ -292,7 +292,7 @@ THRUST_DETAIL_KERNEL_ATTRIBUTES void __launch_bounds__(Agent::ptx_plan::BLOCK_TH _kernel_agent_vshmem(char* vshmem, _0 x0, _1 x1, _2 x2, _3 x3) { extern __shared__ char shmem[]; - vshmem = vshmem == NULL ? shmem : vshmem + blockIdx.x * temp_storage_size::value; + vshmem = vshmem == nullptr ? shmem : vshmem + blockIdx.x * temp_storage_size::value; Agent::entry(x0, x1, x2, x3, vshmem); } template @@ -300,7 +300,7 @@ THRUST_DETAIL_KERNEL_ATTRIBUTES void __launch_bounds__(Agent::ptx_plan::BLOCK_TH _kernel_agent_vshmem(char* vshmem, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4) { extern __shared__ char shmem[]; - vshmem = vshmem == NULL ? shmem : vshmem + blockIdx.x * temp_storage_size::value; + vshmem = vshmem == nullptr ? shmem : vshmem + blockIdx.x * temp_storage_size::value; Agent::entry(x0, x1, x2, x3, x4, vshmem); } template @@ -308,7 +308,7 @@ THRUST_DETAIL_KERNEL_ATTRIBUTES void __launch_bounds__(Agent::ptx_plan::BLOCK_TH _kernel_agent_vshmem(char* vshmem, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5) { extern __shared__ char shmem[]; - vshmem = vshmem == NULL ? shmem : vshmem + blockIdx.x * temp_storage_size::value; + vshmem = vshmem == nullptr ? shmem : vshmem + blockIdx.x * temp_storage_size::value; Agent::entry(x0, x1, x2, x3, x4, x5, vshmem); } template @@ -316,7 +316,7 @@ THRUST_DETAIL_KERNEL_ATTRIBUTES void __launch_bounds__(Agent::ptx_plan::BLOCK_TH _kernel_agent_vshmem(char* vshmem, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6) { extern __shared__ char shmem[]; - vshmem = vshmem == NULL ? shmem : vshmem + blockIdx.x * temp_storage_size::value; + vshmem = vshmem == nullptr ? shmem : vshmem + blockIdx.x * temp_storage_size::value; Agent::entry(x0, x1, x2, x3, x4, x5, x6, vshmem); } template @@ -324,7 +324,7 @@ THRUST_DETAIL_KERNEL_ATTRIBUTES void __launch_bounds__(Agent::ptx_plan::BLOCK_TH _kernel_agent_vshmem(char* vshmem, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7) { extern __shared__ char shmem[]; - vshmem = vshmem == NULL ? shmem : vshmem + blockIdx.x * temp_storage_size::value; + vshmem = vshmem == nullptr ? shmem : vshmem + blockIdx.x * temp_storage_size::value; Agent::entry(x0, x1, x2, x3, x4, x5, x6, x7, vshmem); } template @@ -332,7 +332,7 @@ THRUST_DETAIL_KERNEL_ATTRIBUTES void __launch_bounds__(Agent::ptx_plan::BLOCK_TH _kernel_agent_vshmem(char* vshmem, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8) { extern __shared__ char shmem[]; - vshmem = vshmem == NULL ? shmem : vshmem + blockIdx.x * temp_storage_size::value; + vshmem = vshmem == nullptr ? shmem : vshmem + blockIdx.x * temp_storage_size::value; Agent::entry(x0, x1, x2, x3, x4, x5, x6, x7, x8, vshmem); } template @@ -340,7 +340,7 @@ THRUST_DETAIL_KERNEL_ATTRIBUTES void __launch_bounds__(Agent::ptx_plan::BLOCK_TH _kernel_agent_vshmem(char* vshmem, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9) { extern __shared__ char shmem[]; - vshmem = vshmem == NULL ? shmem : vshmem + blockIdx.x * temp_storage_size::value; + vshmem = vshmem == nullptr ? shmem : vshmem + blockIdx.x * temp_storage_size::value; Agent::entry(x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, vshmem); } template ::value; + vshmem = vshmem == nullptr ? shmem : vshmem + blockIdx.x * temp_storage_size::value; Agent::entry(x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA, vshmem); } template ::value; + vshmem = vshmem == nullptr ? shmem : vshmem + blockIdx.x * temp_storage_size::value; Agent::entry(x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA, xB, vshmem); } template ::value; + vshmem = vshmem == nullptr ? shmem : vshmem + blockIdx.x * temp_storage_size::value; Agent::entry(x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA, xB, xC, vshmem); } template ::value; + vshmem = vshmem == nullptr ? shmem : vshmem + blockIdx.x * temp_storage_size::value; Agent::entry(x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA, xB, xC, xD, vshmem); } template ::value; + vshmem = vshmem == nullptr ? shmem : vshmem + blockIdx.x * temp_storage_size::value; Agent::entry(x0, x1, x2, x3, x4, x5, x6, x7, x8, x9, xA, xB, xC, xD, xE, vshmem); } # endif @@ -701,7 +701,7 @@ struct AgentLauncher : Agent , stream(stream_) , name(name_) , grid(static_cast((count + plan.items_per_tile - 1) / plan.items_per_tile)) - , vshmem(NULL) + , vshmem(nullptr) , has_shmem((size_t) core::get_max_shared_memory_per_block() >= (size_t) plan.shared_memory_size) , shmem_size(has_shmem ? plan.shared_memory_size : 0) { @@ -729,7 +729,7 @@ struct AgentLauncher : Agent , stream(stream_) , name(name_) , grid(plan.grid_size) - , vshmem(NULL) + , vshmem(nullptr) , has_shmem((size_t) core::get_max_shared_memory_per_block() >= (size_t) plan.shared_memory_size) , shmem_size(has_shmem ? plan.shared_memory_size : 0) { @@ -1012,7 +1012,7 @@ struct AgentLauncher : Agent void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::true_type, Args... args) const { - assert(has_shmem && vshmem == NULL); + assert(has_shmem && vshmem == nullptr); print_info(_kernel_agent); launcher::triple_chevron(grid, plan.block_threads, shmem_size, stream) .doit(_kernel_agent, args...); @@ -1030,7 +1030,7 @@ struct AgentLauncher : Agent void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::false_type, Args... args) const { - assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0)); + assert((has_shmem && vshmem == nullptr) || (!has_shmem && vshmem != nullptr && shmem_size == 0)); print_info(_kernel_agent_vshmem); launcher::triple_chevron(grid, plan.block_threads, shmem_size, stream) .doit(_kernel_agent_vshmem, vshmem, args...); @@ -1047,7 +1047,7 @@ struct AgentLauncher : Agent template void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::false_type, _0 x0) const { - assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0)); + assert((has_shmem && vshmem == nullptr) || (!has_shmem && vshmem != nullptr && shmem_size == 0)); void (*ptr)(char*, _0) = _kernel_agent_vshmem; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, shmem_size, stream).doit(ptr, vshmem, x0); @@ -1055,7 +1055,7 @@ struct AgentLauncher : Agent template void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::false_type, _0 x0, _1 x1) const { - assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0)); + assert((has_shmem && vshmem == nullptr) || (!has_shmem && vshmem != nullptr && shmem_size == 0)); void (*ptr)(char*, _0, _1) = _kernel_agent_vshmem; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, shmem_size, stream).doit(ptr, vshmem, x0, x1); @@ -1063,7 +1063,7 @@ struct AgentLauncher : Agent template void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2) const { - assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0)); + assert((has_shmem && vshmem == nullptr) || (!has_shmem && vshmem != nullptr && shmem_size == 0)); void (*ptr)(char*, _0, _1, _2) = _kernel_agent_vshmem; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, shmem_size, stream).doit(ptr, vshmem, x0, x1, x2); @@ -1071,7 +1071,7 @@ struct AgentLauncher : Agent template void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3) const { - assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0)); + assert((has_shmem && vshmem == nullptr) || (!has_shmem && vshmem != nullptr && shmem_size == 0)); void (*ptr)(char*, _0, _1, _2, _3) = _kernel_agent_vshmem; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, shmem_size, stream).doit(ptr, vshmem, x0, x1, x2, x3); @@ -1079,7 +1079,7 @@ struct AgentLauncher : Agent template void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4) const { - assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0)); + assert((has_shmem && vshmem == nullptr) || (!has_shmem && vshmem != nullptr && shmem_size == 0)); void (*ptr)(char*, _0, _1, _2, _3, _4) = _kernel_agent_vshmem; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, shmem_size, stream).doit(ptr, vshmem, x0, x1, x2, x3, x4); @@ -1087,7 +1087,7 @@ struct AgentLauncher : Agent template void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5) const { - assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0)); + assert((has_shmem && vshmem == nullptr) || (!has_shmem && vshmem != nullptr && shmem_size == 0)); void (*ptr)(char*, _0, _1, _2, _3, _4, _5) = _kernel_agent_vshmem; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, shmem_size, stream).doit(ptr, vshmem, x0, x1, x2, x3, x4, x5); @@ -1096,7 +1096,7 @@ struct AgentLauncher : Agent void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6) const { - assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0)); + assert((has_shmem && vshmem == nullptr) || (!has_shmem && vshmem != nullptr && shmem_size == 0)); void (*ptr)(char*, _0, _1, _2, _3, _4, _5, _6) = _kernel_agent_vshmem; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, shmem_size, stream).doit(ptr, vshmem, x0, x1, x2, x3, x4, x5, x6); @@ -1105,7 +1105,7 @@ struct AgentLauncher : Agent void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7) const { - assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0)); + assert((has_shmem && vshmem == nullptr) || (!has_shmem && vshmem != nullptr && shmem_size == 0)); void (*ptr)(char*, _0, _1, _2, _3, _4, _5, _6, _7) = _kernel_agent_vshmem; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, shmem_size, stream) @@ -1115,7 +1115,7 @@ struct AgentLauncher : Agent void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8) const { - assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0)); + assert((has_shmem && vshmem == nullptr) || (!has_shmem && vshmem != nullptr && shmem_size == 0)); void (*ptr)(char*, _0, _1, _2, _3, _4, _5, _6, _7, _8) = _kernel_agent_vshmem; launcher::triple_chevron(grid, plan.block_threads, shmem_size, stream) @@ -1125,7 +1125,7 @@ struct AgentLauncher : Agent void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9) const { - assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0)); + assert((has_shmem && vshmem == nullptr) || (!has_shmem && vshmem != nullptr && shmem_size == 0)); void (*ptr)(char*, _0, _1, _2, _3, _4, _5, _6, _7, _8, _9) = _kernel_agent_vshmem; print_info(ptr); @@ -1136,7 +1136,7 @@ struct AgentLauncher : Agent void THRUST_RUNTIME_FUNCTION launch_impl( thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9, _xA xA) const { - assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0)); + assert((has_shmem && vshmem == nullptr) || (!has_shmem && vshmem != nullptr && shmem_size == 0)); void (*ptr)(char*, _0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _xA) = _kernel_agent_vshmem; print_info(ptr); @@ -1159,7 +1159,7 @@ struct AgentLauncher : Agent thrust::detail::false_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9, _xA xA, _xB xB) const { - assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0)); + assert((has_shmem && vshmem == nullptr) || (!has_shmem && vshmem != nullptr && shmem_size == 0)); void (*ptr)(char*, _0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _xA, _xB) = _kernel_agent_vshmem; print_info(ptr); @@ -1195,7 +1195,7 @@ struct AgentLauncher : Agent _xB xB, _xC xC) const { - assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0)); + assert((has_shmem && vshmem == nullptr) || (!has_shmem && vshmem != nullptr && shmem_size == 0)); void (*ptr)(char*, _0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _xA, _xB, _xC) = _kernel_agent_vshmem; print_info(ptr); @@ -1233,7 +1233,7 @@ struct AgentLauncher : Agent _xC xC, _xD xD) const { - assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0)); + assert((has_shmem && vshmem == nullptr) || (!has_shmem && vshmem != nullptr && shmem_size == 0)); void (*ptr)(char*, _0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _xA, _xB, _xC, _xD) = _kernel_agent_vshmem; print_info(ptr); @@ -1273,7 +1273,7 @@ struct AgentLauncher : Agent _xD xD, _xE xE) const { - assert((has_shmem && vshmem == NULL) || (!has_shmem && vshmem != NULL && shmem_size == 0)); + assert((has_shmem && vshmem == nullptr) || (!has_shmem && vshmem != nullptr && shmem_size == 0)); void (*ptr)(char*, _0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _xA, _xB, _xC, _xD, _xE) = _kernel_agent_vshmem; print_info(ptr); @@ -1288,7 +1288,7 @@ struct AgentLauncher : Agent template void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::true_type, _0 x0) const { - assert(has_shmem && vshmem == NULL); + assert(has_shmem && vshmem == nullptr); void (*ptr)(_0) = _kernel_agent; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, plan.shared_memory_size, stream).doit(ptr, x0); @@ -1296,7 +1296,7 @@ struct AgentLauncher : Agent template void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::true_type, _0 x0, _1 x1) const { - assert(has_shmem && vshmem == NULL); + assert(has_shmem && vshmem == nullptr); void (*ptr)(_0, _1) = _kernel_agent; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, plan.shared_memory_size, stream).doit(ptr, x0, x1); @@ -1304,7 +1304,7 @@ struct AgentLauncher : Agent template void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2) const { - assert(has_shmem && vshmem == NULL); + assert(has_shmem && vshmem == nullptr); void (*ptr)(_0, _1, _2) = _kernel_agent; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, plan.shared_memory_size, stream).doit(ptr, x0, x1, x2); @@ -1312,7 +1312,7 @@ struct AgentLauncher : Agent template void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3) const { - assert(has_shmem && vshmem == NULL); + assert(has_shmem && vshmem == nullptr); void (*ptr)(_0, _1, _2, _3) = _kernel_agent; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, plan.shared_memory_size, stream).doit(ptr, x0, x1, x2, x3); @@ -1320,7 +1320,7 @@ struct AgentLauncher : Agent template void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4) const { - assert(has_shmem && vshmem == NULL); + assert(has_shmem && vshmem == nullptr); void (*ptr)(_0, _1, _2, _3, _4) = _kernel_agent; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, plan.shared_memory_size, stream).doit(ptr, x0, x1, x2, x3, x4); @@ -1328,7 +1328,7 @@ struct AgentLauncher : Agent template void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5) const { - assert(has_shmem && vshmem == NULL); + assert(has_shmem && vshmem == nullptr); void (*ptr)(_0, _1, _2, _3, _4, _5) = _kernel_agent; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, plan.shared_memory_size, stream) @@ -1338,7 +1338,7 @@ struct AgentLauncher : Agent void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6) const { - assert(has_shmem && vshmem == NULL); + assert(has_shmem && vshmem == nullptr); void (*ptr)(_0, _1, _2, _3, _4, _5, _6) = _kernel_agent; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, plan.shared_memory_size, stream) @@ -1348,7 +1348,7 @@ struct AgentLauncher : Agent void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7) const { - assert(has_shmem && vshmem == NULL); + assert(has_shmem && vshmem == nullptr); void (*ptr)(_0, _1, _2, _3, _4, _5, _6, _7) = _kernel_agent; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, plan.shared_memory_size, stream) @@ -1358,7 +1358,7 @@ struct AgentLauncher : Agent void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8) const { - assert(has_shmem && vshmem == NULL); + assert(has_shmem && vshmem == nullptr); void (*ptr)(_0, _1, _2, _3, _4, _5, _6, _7, _8) = _kernel_agent; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, plan.shared_memory_size, stream) @@ -1368,7 +1368,7 @@ struct AgentLauncher : Agent void THRUST_RUNTIME_FUNCTION launch_impl(thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9) const { - assert(has_shmem && vshmem == NULL); + assert(has_shmem && vshmem == nullptr); void (*ptr)(_0, _1, _2, _3, _4, _5, _6, _7, _8, _9) = _kernel_agent; print_info(ptr); launcher::triple_chevron(grid, plan.block_threads, plan.shared_memory_size, stream) @@ -1378,7 +1378,7 @@ struct AgentLauncher : Agent void THRUST_RUNTIME_FUNCTION launch_impl( thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9, _xA xA) const { - assert(has_shmem && vshmem == NULL); + assert(has_shmem && vshmem == nullptr); void (*ptr)(_0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _xA) = _kernel_agent; print_info(ptr); @@ -1401,7 +1401,7 @@ struct AgentLauncher : Agent thrust::detail::true_type, _0 x0, _1 x1, _2 x2, _3 x3, _4 x4, _5 x5, _6 x6, _7 x7, _8 x8, _9 x9, _xA xA, _xB xB) const { - assert(has_shmem && vshmem == NULL); + assert(has_shmem && vshmem == nullptr); void (*ptr)(_0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _xA, _xB) = _kernel_agent; print_info(ptr); @@ -1437,7 +1437,7 @@ struct AgentLauncher : Agent _xB xB, _xC xC) const { - assert(has_shmem && vshmem == NULL); + assert(has_shmem && vshmem == nullptr); void (*ptr)(_0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _xA, _xB, _xC) = _kernel_agent; print_info(ptr); @@ -1475,7 +1475,7 @@ struct AgentLauncher : Agent _xC xC, _xD xD) const { - assert(has_shmem && vshmem == NULL); + assert(has_shmem && vshmem == nullptr); void (*ptr)(_0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _xA, _xB, _xC, _xD) = _kernel_agent; print_info(ptr); @@ -1515,7 +1515,7 @@ struct AgentLauncher : Agent _xD xD, _xE xE) const { - assert(has_shmem && vshmem == NULL); + assert(has_shmem && vshmem == nullptr); void (*ptr)(_0, _1, _2, _3, _4, _5, _6, _7, _8, _9, _xA, _xB, _xC, _xD, _xE) = _kernel_agent; print_info(ptr); diff --git a/thrust/thrust/system/cuda/detail/extrema.h b/thrust/thrust/system/cuda/detail/extrema.h index 890f4f05f9b..dd22cb22439 100644 --- a/thrust/thrust/system/cuda/detail/extrema.h +++ b/thrust/thrust/system/cuda/detail/extrema.h @@ -204,12 +204,12 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step( size_t vshmem_size = core::vshmem_size(reduce_plan.shared_memory_size, 1); // small, single tile size - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { temp_storage_bytes = max(1, vshmem_size); return status; } - char* vshmem_ptr = vshmem_size > 0 ? (char*) d_temp_storage : NULL; + char* vshmem_ptr = vshmem_size > 0 ? (char*) d_temp_storage : nullptr; reduce_agent ra(reduce_plan, num_items, stream, vshmem_ptr, "reduce_agent: single_tile only"); ra.launch(input_it, output_it, num_items, reduction_op); @@ -245,7 +245,7 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step( size_t vshmem_size = core::vshmem_size(reduce_plan.shared_memory_size, max_blocks); // Temporary storage allocation requirements - void* allocations[3] = {NULL, NULL, NULL}; + void* allocations[3] = {nullptr, nullptr, nullptr}; size_t allocation_sizes[3] = { max_blocks * sizeof(T), // bytes needed for privatized block reductions cub::GridQueue::AllocationSize(), // bytes needed for grid queue descriptor0 @@ -253,14 +253,14 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step( }; status = cub::AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes); CUDA_CUB_RET_IF_FAIL(status); - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { return status; } T* d_block_reductions = (T*) allocations[0]; cub::GridQueue queue(allocations[1]); - char* vshmem_ptr = vshmem_size > 0 ? (char*) allocations[2] : NULL; + char* vshmem_ptr = vshmem_size > 0 ? (char*) allocations[2] : nullptr; // Get grid size for device_reduce_sweep_kernel int reduce_grid_size = 0; @@ -321,14 +321,14 @@ extrema(execution_policy& policy, InputIt first, Size num_items, Binary status, doit_step, num_items, - (NULL, temp_storage_bytes, first, num_items_fixed, binary_op, reinterpret_cast(NULL), stream)); + (nullptr, temp_storage_bytes, first, num_items_fixed, binary_op, reinterpret_cast(nullptr), stream)); cuda_cub::throw_on_error(status, "extrema failed on 1st step"); size_t allocation_sizes[2] = {sizeof(T*), temp_storage_bytes}; - void* allocations[2] = {NULL, NULL}; + void* allocations[2] = {nullptr, nullptr}; size_t storage_size = 0; - status = core::alias_storage(NULL, storage_size, allocations, allocation_sizes); + status = core::alias_storage(nullptr, storage_size, allocations, allocation_sizes); cuda_cub::throw_on_error(status, "extrema failed on 1st alias storage"); // Allocate temporary storage. @@ -379,7 +379,7 @@ element(execution_policy& policy, ItemsIt first, ItemsIt last, BinaryPr zip_iterator begin = make_zip_iterator(iter_tuple); - T result = extrema(policy, begin, num_items, arg_min_t(binary_pred), (T*) (NULL)); + T result = extrema(policy, begin, num_items, arg_min_t(binary_pred), (T*) (nullptr)); return first + thrust::get<1>(result); } @@ -454,7 +454,7 @@ minmax_element(execution_policy& policy, ItemsIt first, ItemsIt last, B zip_iterator begin = make_zip_iterator(iter_tuple); two_pairs_type result = __extrema::extrema( - policy, transform_t(begin, duplicate_t()), num_items, arg_minmax_t(binary_pred), (two_pairs_type*) (NULL)); + policy, transform_t(begin, duplicate_t()), num_items, arg_minmax_t(binary_pred), (two_pairs_type*) (nullptr)); ret = thrust::make_pair(first + get<1>(get<0>(result)), first + get<1>(get<1>(result)));), // CDP Sequential impl: (ret = thrust::minmax_element(cvt_to_seq(derived_cast(policy)), first, last, binary_pred);)); diff --git a/thrust/thrust/system/cuda/detail/merge.h b/thrust/thrust/system/cuda/detail/merge.h index c4fed7ee1c7..5cdfe440f81 100644 --- a/thrust/thrust/system/cuda/detail/merge.h +++ b/thrust/thrust/system/cuda/detail/merge.h @@ -644,20 +644,20 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step( size_t temp_storage1 = (1 + num_tiles) * sizeof(Size); size_t temp_storage2 = core::vshmem_size(merge_plan.shared_memory_size, num_tiles); - void* allocations[2] = {NULL, NULL}; + void* allocations[2] = {nullptr, nullptr}; size_t allocation_sizes[2] = {temp_storage1, temp_storage2}; status = core::alias_storage(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes); CUDA_CUB_RET_IF_FAIL(status); - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { return status; } // partition data into work balanced tiles Size* merge_partitions = (Size*) allocations[0]; - char* vshmem_ptr = temp_storage2 > 0 ? (char*) allocations[1] : NULL; + char* vshmem_ptr = temp_storage2 > 0 ? (char*) allocations[1] : nullptr; { Size num_partitions = num_tiles + 1; @@ -713,7 +713,7 @@ THRUST_RUNTIME_FUNCTION pair merge( cudaError_t status; status = doit_step( - NULL, + nullptr, storage_size, keys1_first, keys2_first, diff --git a/thrust/thrust/system/cuda/detail/reduce.h b/thrust/thrust/system/cuda/detail/reduce.h index 05aaff51e42..c539db4b539 100644 --- a/thrust/thrust/system/cuda/detail/reduce.h +++ b/thrust/thrust/system/cuda/detail/reduce.h @@ -644,12 +644,12 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step( size_t vshmem_size = core::vshmem_size(reduce_plan.shared_memory_size, 1); // small, single tile size - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { temp_storage_bytes = max(1, vshmem_size); return status; } - char* vshmem_ptr = vshmem_size > 0 ? (char*) d_temp_storage : NULL; + char* vshmem_ptr = vshmem_size > 0 ? (char*) d_temp_storage : nullptr; reduce_agent ra(reduce_plan, num_items, stream, vshmem_ptr, "reduce_agent: single_tile only"); ra.launch(input_it, output_it, num_items, reduction_op, init); @@ -685,7 +685,7 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step( size_t vshmem_size = core::vshmem_size(reduce_plan.shared_memory_size, max_blocks); // Temporary storage allocation requirements - void* allocations[3] = {NULL, NULL, NULL}; + void* allocations[3] = {nullptr, nullptr, nullptr}; size_t allocation_sizes[3] = { max_blocks * sizeof(T), // bytes needed for privatized block reductions cub::GridQueue::AllocationSize(), // bytes needed for grid queue descriptor0 @@ -693,14 +693,14 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step( }; status = cub::AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes); CUDA_CUB_RET_IF_FAIL(status); - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { return status; } T* d_block_reductions = (T*) allocations[0]; cub::GridQueue queue(allocations[1]); - char* vshmem_ptr = vshmem_size > 0 ? (char*) allocations[2] : NULL; + char* vshmem_ptr = vshmem_size > 0 ? (char*) allocations[2] : nullptr; // Get grid size for device_reduce_sweep_kernel int reduce_grid_size = 0; @@ -760,14 +760,15 @@ reduce(execution_policy& policy, InputIt first, Size num_items, T init, cudaStream_t stream = cuda_cub::stream(policy); cudaError_t status; - status = doit_step(NULL, temp_storage_bytes, first, num_items, init, binary_op, reinterpret_cast(NULL), stream); + status = + doit_step(nullptr, temp_storage_bytes, first, num_items, init, binary_op, reinterpret_cast(nullptr), stream); cuda_cub::throw_on_error(status, "reduce failed on 1st step"); size_t allocation_sizes[2] = {sizeof(T*), temp_storage_bytes}; - void* allocations[2] = {NULL, NULL}; + void* allocations[2] = {nullptr, nullptr}; size_t storage_size = 0; - status = core::alias_storage(NULL, storage_size, allocations, allocation_sizes); + status = core::alias_storage(nullptr, storage_size, allocations, allocation_sizes); cuda_cub::throw_on_error(status, "reduce failed on 1st alias_storage"); // Allocate temporary storage. @@ -809,7 +810,7 @@ reduce_n_impl(execution_policy& policy, InputIt first, Size num_items, status, cub::DeviceReduce::Reduce, num_items, - (NULL, tmp_size, first, reinterpret_cast(NULL), num_items_fixed, binary_op, init, stream)); + (nullptr, tmp_size, first, reinterpret_cast(nullptr), num_items_fixed, binary_op, init, stream)); cuda_cub::throw_on_error(status, "after reduction step 1"); // Allocate temporary storage. diff --git a/thrust/thrust/system/cuda/detail/reduce_by_key.h b/thrust/thrust/system/cuda/detail/reduce_by_key.h index 5a7eeabfb2b..0ce140fe4c9 100644 --- a/thrust/thrust/system/cuda/detail/reduce_by_key.h +++ b/thrust/thrust/system/cuda/detail/reduce_by_key.h @@ -815,11 +815,11 @@ THRUST_RUNTIME_FUNCTION cudaError_t doit_step( status = ScanTileState::AllocationSize(static_cast(num_tiles), allocation_sizes[0]); CUDA_CUB_RET_IF_FAIL(status); - void* allocations[2] = {NULL, NULL}; + void* allocations[2] = {nullptr, nullptr}; status = cub::AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes); CUDA_CUB_RET_IF_FAIL(status); - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { return status; } @@ -832,7 +832,7 @@ THRUST_RUNTIME_FUNCTION cudaError_t doit_step( ia.launch(tile_state, num_tiles, num_runs_output_it); CUDA_CUB_RET_IF_FAIL(cudaPeekAtLastError()); - char* vshmem_ptr = vshmem_size > 0 ? (char*) allocations[1] : NULL; + char* vshmem_ptr = vshmem_size > 0 ? (char*) allocations[1] : nullptr; reduce_by_key_agent rbka(reduce_by_key_plan, num_items, stream, vshmem_ptr, "reduce_by_keys::reduce_by_key_agent"); rbka.launch( @@ -878,13 +878,13 @@ THRUST_RUNTIME_FUNCTION pair reduce_by_key_dispatc cudaError_t status; status = doit_step( - NULL, + nullptr, temp_storage_bytes, keys_first, values_first, keys_output, values_output, - reinterpret_cast(NULL), + reinterpret_cast(nullptr), equality_op, reduction_op, num_items, @@ -892,10 +892,10 @@ THRUST_RUNTIME_FUNCTION pair reduce_by_key_dispatc cuda_cub::throw_on_error(status, "reduce_by_key failed on 1st step"); size_t allocation_sizes[2] = {sizeof(Size), temp_storage_bytes}; - void* allocations[2] = {NULL, NULL}; + void* allocations[2] = {nullptr, nullptr}; size_t storage_size = 0; - status = core::alias_storage(NULL, storage_size, allocations, allocation_sizes); + status = core::alias_storage(nullptr, storage_size, allocations, allocation_sizes); cuda_cub::throw_on_error(status, "reduce failed on 1st alias_storage"); // Allocate temporary storage. diff --git a/thrust/thrust/system/cuda/detail/set_operations.h b/thrust/thrust/system/cuda/detail/set_operations.h index 663bb975e89..fd36024d63a 100644 --- a/thrust/thrust/system/cuda/detail/set_operations.h +++ b/thrust/thrust/system/cuda/detail/set_operations.h @@ -1101,13 +1101,13 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step( size_t vshmem_storage = core::vshmem_size(set_op_plan.shared_memory_size, num_tiles); size_t partition_agent_storage = (num_tiles + 1) * sizeof(Size) * 2; - void* allocations[3] = {NULL, NULL, NULL}; + void* allocations[3] = {nullptr, nullptr, nullptr}; size_t allocation_sizes[3] = {tile_agent_storage, partition_agent_storage, vshmem_storage}; status = core::alias_storage(d_temp_storage, temp_storage_size, allocations, allocation_sizes); CUDA_CUB_RET_IF_FAIL(status); - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { return status; } @@ -1117,7 +1117,7 @@ cudaError_t THRUST_RUNTIME_FUNCTION doit_step( CUDA_CUB_RET_IF_FAIL(status); pair* partitions = (pair*) allocations[1]; - char* vshmem_ptr = vshmem_storage > 0 ? (char*) allocations[2] : NULL; + char* vshmem_ptr = vshmem_storage > 0 ? (char*) allocations[2] : nullptr; init_agent ia(init_plan, num_tiles, stream, "set_op::init_agent"); ia.launch(tile_state, num_tiles); @@ -1189,7 +1189,7 @@ THRUST_RUNTIME_FUNCTION pair set_operations( doit_step, num_keys1, num_keys2, - (NULL, + (nullptr, temp_storage_bytes, keys1_first, keys2_first, @@ -1199,18 +1199,18 @@ THRUST_RUNTIME_FUNCTION pair set_operations( num_keys2_fixed, keys_output, values_output, - reinterpret_cast(NULL), + reinterpret_cast(nullptr), compare_op, set_op, stream)); cuda_cub::throw_on_error(status, "set_operations failed on 1st step"); size_t allocation_sizes[2] = {sizeof(std::size_t), temp_storage_bytes}; - void* allocations[2] = {NULL, NULL}; + void* allocations[2] = {nullptr, nullptr}; size_t storage_size = 0; - status = core::alias_storage(NULL, storage_size, allocations, allocation_sizes); + status = core::alias_storage(nullptr, storage_size, allocations, allocation_sizes); cuda_cub::throw_on_error(status, "set_operations failed on 1st alias_storage"); // Allocate temporary storage. @@ -1315,7 +1315,7 @@ OutputIt _CCCL_HOST_DEVICE set_intersection( CompareOp compare) { THRUST_CDP_DISPATCH( - (using items1_t = thrust::iterator_value_t; items1_t* null_ = NULL; + (using items1_t = thrust::iterator_value_t; items1_t* null_ = nullptr; auto tmp = __set_operations::set_operations( policy, items1_first, diff --git a/thrust/thrust/system/cuda/detail/sort.h b/thrust/thrust/system/cuda/detail/sort.h index 5cb1ed09cfe..6b20a894a54 100644 --- a/thrust/thrust/system/cuda/detail/sort.h +++ b/thrust/thrust/system/cuda/detail/sort.h @@ -136,7 +136,7 @@ THRUST_RUNTIME_FUNCTION void merge_sort( cudaStream_t stream = cuda_cub::stream(policy); cudaError_t status; - status = doit_step(NULL, storage_size, keys_first, items_first, count, compare_op, stream); + status = doit_step(nullptr, storage_size, keys_first, items_first, count, compare_op, stream); cuda_cub::throw_on_error(status, "merge_sort: failed on 1st step"); // Allocate temporary storage. @@ -261,8 +261,8 @@ THRUST_RUNTIME_FUNCTION void radix_sort(execution_policy& policy, Key* size_t temp_storage_bytes = 0; cudaStream_t stream = cuda_cub::stream(policy); - cub::DoubleBuffer keys_buffer(keys, NULL); - cub::DoubleBuffer items_buffer(items, NULL); + cub::DoubleBuffer keys_buffer(keys, nullptr); + cub::DoubleBuffer items_buffer(items, nullptr); Size keys_count = count; Size items_count = SORT_ITEMS::value ? count : 0; @@ -270,7 +270,7 @@ THRUST_RUNTIME_FUNCTION void radix_sort(execution_policy& policy, Key* cudaError_t status; status = - dispatch::doit(NULL, temp_storage_bytes, keys_buffer, items_buffer, keys_count, stream); + dispatch::doit(nullptr, temp_storage_bytes, keys_buffer, items_buffer, keys_count, stream); cuda_cub::throw_on_error(status, "radix_sort: failed on 1st step"); size_t keys_temp_storage = core::align_to(sizeof(Key) * keys_count, 128); diff --git a/thrust/thrust/system/cuda/detail/transform_reduce.h b/thrust/thrust/system/cuda/detail/transform_reduce.h index 912b790c2ea..e7544403e1c 100644 --- a/thrust/thrust/system/cuda/detail/transform_reduce.h +++ b/thrust/thrust/system/cuda/detail/transform_reduce.h @@ -82,7 +82,7 @@ THRUST_RUNTIME_FUNCTION T transform_reduce_n_impl( status, cub::DeviceReduce::TransformReduce, num_items, - (NULL, tmp_size, first, reinterpret_cast(NULL), num_items_fixed, binary_op, unary_op, init, stream)); + (nullptr, tmp_size, first, reinterpret_cast(nullptr), num_items_fixed, binary_op, unary_op, init, stream)); cuda_cub::throw_on_error(status, "after reduction step 1"); // Allocate temporary storage. diff --git a/thrust/thrust/system/cuda/detail/unique.h b/thrust/thrust/system/cuda/detail/unique.h index 10b06cba709..d8afe32d574 100644 --- a/thrust/thrust/system/cuda/detail/unique.h +++ b/thrust/thrust/system/cuda/detail/unique.h @@ -517,12 +517,12 @@ static cudaError_t THRUST_RUNTIME_FUNCTION doit_step( status = ScanTileState::AllocationSize(static_cast(num_tiles), allocation_sizes[0]); CUDA_CUB_RET_IF_FAIL(status); - void* allocations[2] = {NULL, NULL}; + void* allocations[2] = {nullptr, nullptr}; // status = cub::AliasTemporaries(d_temp_storage, temp_storage_bytes, allocations, allocation_sizes); CUDA_CUB_RET_IF_FAIL(status); - if (d_temp_storage == NULL) + if (d_temp_storage == nullptr) { return status; } @@ -541,7 +541,7 @@ static cudaError_t THRUST_RUNTIME_FUNCTION doit_step( return status; } - char* vshmem_ptr = vshmem_size > 0 ? (char*) allocations[1] : NULL; + char* vshmem_ptr = vshmem_size > 0 ? (char*) allocations[1] : nullptr; unique_agent ua(unique_plan, num_items, stream, vshmem_ptr, "unique_by_key::unique_agent"); ua.launch(items_in, items_out, binary_pred, num_selected_out, num_items, tile_status, num_tiles); @@ -566,21 +566,21 @@ THRUST_RUNTIME_FUNCTION ItemsOutputIt unique( cudaError_t status; status = doit_step( - NULL, + nullptr, temp_storage_bytes, items_first, items_result, binary_pred, - reinterpret_cast(NULL), + reinterpret_cast(nullptr), num_items, stream); cuda_cub::throw_on_error(status, "unique: failed on 1st step"); size_t allocation_sizes[2] = {sizeof(size_type), temp_storage_bytes}; - void* allocations[2] = {NULL, NULL}; + void* allocations[2] = {nullptr, nullptr}; size_t storage_size = 0; - status = core::alias_storage(NULL, storage_size, allocations, allocation_sizes); + status = core::alias_storage(nullptr, storage_size, allocations, allocation_sizes); cuda_cub::throw_on_error(status, "unique: failed on 1st step"); // Allocate temporary storage. diff --git a/thrust/thrust/system/detail/sequential/trivial_copy.h b/thrust/thrust/system/detail/sequential/trivial_copy.h index 9dd70045f8b..30492b13ae9 100644 --- a/thrust/thrust/system/detail/sequential/trivial_copy.h +++ b/thrust/thrust/system/detail/sequential/trivial_copy.h @@ -53,7 +53,7 @@ _CCCL_HOST_DEVICE T* trivial_copy_n(const T* first, std::ptrdiff_t n, T* result) return result; } - T* return_value = NULL; + T* return_value = nullptr; NV_IF_TARGET(NV_IS_HOST, (std::memmove(result, first, n * sizeof(T)); return_value = result + n;),