Skip to content

Commit

Permalink
Rename CUB uninitialized_copy (#1913)
Browse files Browse the repository at this point in the history 
CUB's uninitialized_copy is not related to std::uninitialized_copy since it only operates on a single value. Therefore, it should have a different name.
  • Loading branch information
@bernhardmgruber
bernhardmgruber committed Jun 25, 2024
1 parent df8109c commit 76288d5
Show file tree
Hide file tree
Showing 8 changed files with 33 additions and 33 deletions.
6 changes: 3 additions & 3 deletions cub/cub/agent/single_pass_scan_operators.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -1158,7 +1158,7 @@ struct TilePrefixCallbackOp
// Update our status with our tile-aggregate
if (threadIdx.x == 0)
{
detail::uninitialized_copy(&temp_storage.block_aggregate, block_aggregate);
detail::uninitialized_copy_single(&temp_storage.block_aggregate, block_aggregate);

tile_status.SetPartial(tile_idx, block_aggregate);
}
Expand Down Expand Up @@ -1190,9 +1190,9 @@ struct TilePrefixCallbackOp
inclusive_prefix = scan_op(exclusive_prefix, block_aggregate);
tile_status.SetInclusive(tile_idx, inclusive_prefix);

detail::uninitialized_copy(&temp_storage.exclusive_prefix, exclusive_prefix);
detail::uninitialized_copy_single(&temp_storage.exclusive_prefix, exclusive_prefix);

detail::uninitialized_copy(&temp_storage.inclusive_prefix, inclusive_prefix);
detail::uninitialized_copy_single(&temp_storage.inclusive_prefix, inclusive_prefix);
}

// Return exclusive_prefix
Expand Down
28 changes: 14 additions & 14 deletions cub/cub/block/block_exchange.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -225,7 +225,7 @@ private:
{
item_offset += item_offset >> LOG_SMEM_BANKS;
}
detail::uninitialized_copy(temp_storage.buff + item_offset, input_items[ITEM]);
detail::uninitialized_copy_single(temp_storage.buff + item_offset, input_items[ITEM]);
}

CTA_SYNC();
Expand Down Expand Up @@ -274,7 +274,7 @@ private:
{
item_offset += item_offset >> LOG_SMEM_BANKS;
}
detail::uninitialized_copy(temp_storage.buff + item_offset, input_items[ITEM]);
detail::uninitialized_copy_single(temp_storage.buff + item_offset, input_items[ITEM]);
}
}

Expand Down Expand Up @@ -330,7 +330,7 @@ private:
{
item_offset += item_offset >> LOG_SMEM_BANKS;
}
detail::uninitialized_copy(temp_storage.buff + item_offset, input_items[ITEM]);
detail::uninitialized_copy_single(temp_storage.buff + item_offset, input_items[ITEM]);
}

WARP_SYNC(0xffffffff);
Expand Down Expand Up @@ -369,7 +369,7 @@ private:
{
item_offset += item_offset >> LOG_SMEM_BANKS;
}
detail::uninitialized_copy(temp_storage.buff + item_offset, input_items[ITEM]);
detail::uninitialized_copy_single(temp_storage.buff + item_offset, input_items[ITEM]);
}

WARP_SYNC(0xffffffff);
Expand Down Expand Up @@ -401,7 +401,7 @@ private:
{
item_offset += item_offset >> LOG_SMEM_BANKS;
}
detail::uninitialized_copy(temp_storage.buff + item_offset, input_items[ITEM]);
detail::uninitialized_copy_single(temp_storage.buff + item_offset, input_items[ITEM]);
}

WARP_SYNC(0xffffffff);
Expand Down Expand Up @@ -440,7 +440,7 @@ private:
{
item_offset += item_offset >> LOG_SMEM_BANKS;
}
detail::uninitialized_copy(temp_storage.buff + item_offset, input_items[ITEM]);
detail::uninitialized_copy_single(temp_storage.buff + item_offset, input_items[ITEM]);
}

CTA_SYNC();
Expand Down Expand Up @@ -497,7 +497,7 @@ private:
{
item_offset += item_offset >> LOG_SMEM_BANKS;
}
detail::uninitialized_copy(temp_storage.buff + item_offset, input_items[ITEM]);
detail::uninitialized_copy_single(temp_storage.buff + item_offset, input_items[ITEM]);
}
}
}
Expand Down Expand Up @@ -547,7 +547,7 @@ private:
{
item_offset += item_offset >> LOG_SMEM_BANKS;
}
detail::uninitialized_copy(temp_storage.buff + item_offset, input_items[ITEM]);
detail::uninitialized_copy_single(temp_storage.buff + item_offset, input_items[ITEM]);
}

WARP_SYNC(0xffffffff);
Expand All @@ -560,7 +560,7 @@ private:
{
item_offset += item_offset >> LOG_SMEM_BANKS;
}
detail::uninitialized_copy(output_items + ITEM, temp_storage.buff[item_offset]);
detail::uninitialized_copy_single(output_items + ITEM, temp_storage.buff[item_offset]);
}
}

Expand Down Expand Up @@ -591,7 +591,7 @@ private:
{
item_offset += item_offset >> LOG_SMEM_BANKS;
}
detail::uninitialized_copy(temp_storage.buff + item_offset, input_items[ITEM]);
detail::uninitialized_copy_single(temp_storage.buff + item_offset, input_items[ITEM]);
}

WARP_SYNC(0xffffffff);
Expand Down Expand Up @@ -636,7 +636,7 @@ private:
{
item_offset = SHR_ADD(item_offset, LOG_SMEM_BANKS, item_offset);
}
detail::uninitialized_copy(temp_storage.buff + item_offset, input_items[ITEM]);
detail::uninitialized_copy_single(temp_storage.buff + item_offset, input_items[ITEM]);
}

CTA_SYNC();
Expand Down Expand Up @@ -690,7 +690,7 @@ private:
{
item_offset = SHR_ADD(item_offset, LOG_SMEM_BANKS, item_offset);
}
detail::uninitialized_copy(temp_storage.buff + item_offset, input_items[ITEM]);
detail::uninitialized_copy_single(temp_storage.buff + item_offset, input_items[ITEM]);
}
}

Expand Down Expand Up @@ -745,7 +745,7 @@ private:
{
item_offset = SHR_ADD(item_offset, LOG_SMEM_BANKS, item_offset);
}
detail::uninitialized_copy(temp_storage.buff + item_offset, input_items[ITEM]);
detail::uninitialized_copy_single(temp_storage.buff + item_offset, input_items[ITEM]);
}

CTA_SYNC();
Expand Down Expand Up @@ -800,7 +800,7 @@ private:
{
item_offset = SHR_ADD(item_offset, LOG_SMEM_BANKS, item_offset);
}
detail::uninitialized_copy(temp_storage.buff + item_offset, input_items[ITEM]);
detail::uninitialized_copy_single(temp_storage.buff + item_offset, input_items[ITEM]);
}
}

Expand Down
2 changes: 1 addition & 1 deletion cub/cub/block/specializations/block_reduce_warp_reductions.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -181,7 +181,7 @@ struct BlockReduceWarpReductions
// Share lane aggregates
if (lane_id == 0)
{
detail::uninitialized_copy(temp_storage.warp_aggregates + warp_id, warp_aggregate);
detail::uninitialized_copy_single(temp_storage.warp_aggregates + warp_id, warp_aggregate);
}

CTA_SYNC();
Expand Down
16 changes: 8 additions & 8 deletions cub/cub/block/specializations/block_scan_raking.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -300,7 +300,7 @@ struct BlockScanRaking
{
// Place thread partial into shared memory raking grid
T* placement_ptr = BlockRakingLayout::PlacementPtr(temp_storage.raking_grid, linear_tid);
detail::uninitialized_copy(placement_ptr, input);
detail::uninitialized_copy_single(placement_ptr, input);

CTA_SYNC();

Expand Down Expand Up @@ -353,7 +353,7 @@ struct BlockScanRaking
{
// Place thread partial into shared memory raking grid
T* placement_ptr = BlockRakingLayout::PlacementPtr(temp_storage.raking_grid, linear_tid);
detail::uninitialized_copy(placement_ptr, input);
detail::uninitialized_copy_single(placement_ptr, input);

CTA_SYNC();

Expand Down Expand Up @@ -408,7 +408,7 @@ struct BlockScanRaking
{
// Place thread partial into shared memory raking grid
T* placement_ptr = BlockRakingLayout::PlacementPtr(temp_storage.raking_grid, linear_tid);
detail::uninitialized_copy(placement_ptr, input);
detail::uninitialized_copy_single(placement_ptr, input);

CTA_SYNC();

Expand Down Expand Up @@ -476,7 +476,7 @@ struct BlockScanRaking
{
// Place thread partial into shared memory raking grid
T* placement_ptr = BlockRakingLayout::PlacementPtr(temp_storage.raking_grid, linear_tid);
detail::uninitialized_copy(placement_ptr, input);
detail::uninitialized_copy_single(placement_ptr, input);

CTA_SYNC();

Expand Down Expand Up @@ -557,7 +557,7 @@ struct BlockScanRaking
{
// Place thread partial into shared memory raking grid
T* placement_ptr = BlockRakingLayout::PlacementPtr(temp_storage.raking_grid, linear_tid);
detail::uninitialized_copy(placement_ptr, input);
detail::uninitialized_copy_single(placement_ptr, input);

CTA_SYNC();

Expand Down Expand Up @@ -624,7 +624,7 @@ struct BlockScanRaking
{
// Place thread partial into shared memory raking grid
T* placement_ptr = BlockRakingLayout::PlacementPtr(temp_storage.raking_grid, linear_tid);
detail::uninitialized_copy(placement_ptr, input);
detail::uninitialized_copy_single(placement_ptr, input);

CTA_SYNC();

Expand Down Expand Up @@ -678,7 +678,7 @@ struct BlockScanRaking
{
// Place thread partial into shared memory raking grid
T* placement_ptr = BlockRakingLayout::PlacementPtr(temp_storage.raking_grid, linear_tid);
detail::uninitialized_copy(placement_ptr, input);
detail::uninitialized_copy_single(placement_ptr, input);

CTA_SYNC();

Expand Down Expand Up @@ -756,7 +756,7 @@ struct BlockScanRaking
{
// Place thread partial into shared memory raking grid
T* placement_ptr = BlockRakingLayout::PlacementPtr(temp_storage.raking_grid, linear_tid);
detail::uninitialized_copy(placement_ptr, input);
detail::uninitialized_copy_single(placement_ptr, input);

CTA_SYNC();

Expand Down
6 changes: 3 additions & 3 deletions cub/cub/block/specializations/block_scan_warp_scans.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -194,7 +194,7 @@ struct BlockScanWarpScans
// Last lane in each warp shares its warp-aggregate
if (lane_id == WARP_THREADS - 1)
{
detail::uninitialized_copy(temp_storage.warp_aggregates + warp_id, warp_aggregate);
detail::uninitialized_copy_single(temp_storage.warp_aggregates + warp_id, warp_aggregate);
}

CTA_SYNC();
Expand Down Expand Up @@ -417,7 +417,7 @@ struct BlockScanWarpScans
if (lane_id == 0)
{
// Share the prefix with all threads
detail::uninitialized_copy(&temp_storage.block_prefix, block_prefix);
detail::uninitialized_copy_single(&temp_storage.block_prefix, block_prefix);

exclusive_output = block_prefix; // The block prefix is the exclusive output for tid0
}
Expand Down Expand Up @@ -524,7 +524,7 @@ struct BlockScanWarpScans
if (lane_id == 0)
{
// Share the prefix with all threads
detail::uninitialized_copy(&temp_storage.block_prefix, block_prefix);
detail::uninitialized_copy_single(&temp_storage.block_prefix, block_prefix);
}
}

Expand Down
4 changes: 2 additions & 2 deletions cub/cub/detail/uninitialized_copy.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -56,15 +56,15 @@ _CCCL_HOST_DEVICE void uninitialized_copy(T* ptr, U&& val)
template <typename T,
typename U,
typename ::cuda::std::enable_if<::cuda::std::is_trivially_copyable<T>::value, int>::type = 0>
_CCCL_HOST_DEVICE void uninitialized_copy(T* ptr, U&& val)
_CCCL_HOST_DEVICE void uninitialized_copy_single(T* ptr, U&& val)
{
*ptr = ::cuda::std::forward<U>(val);
}

template <typename T,
typename U,
typename ::cuda::std::enable_if<!::cuda::std::is_trivially_copyable<T>::value, int>::type = 0>
_CCCL_HOST_DEVICE void uninitialized_copy(T* ptr, U&& val)
_CCCL_HOST_DEVICE void uninitialized_copy_single(T* ptr, U&& val)
{
new (ptr) T(::cuda::std::forward<U>(val));
}
Expand Down
2 changes: 1 addition & 1 deletion cub/cub/device/dispatch/dispatch_reduce.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -195,7 +195,7 @@ __launch_bounds__(int(ChainedPolicyT::ActivePolicy::ReducePolicy::BLOCK_THREADS)
// Output result
if (threadIdx.x == 0)
{
detail::uninitialized_copy(d_out + blockIdx.x, block_aggregate);
detail::uninitialized_copy_single(d_out + blockIdx.x, block_aggregate);
}
}

Expand Down
2 changes: 1 addition & 1 deletion cub/cub/util_type.cuh
View file
Original file line number Diff line number Diff line change
Expand Up @@ -303,7 +303,7 @@ struct InputValue
}
else
{
detail::uninitialized_copy(&m_immediate_value, other.m_immediate_value);
detail::uninitialized_copy_single(&m_immediate_value, other.m_immediate_value);
}
}

Expand Down

0 comments on commit 76288d5

Please sign in to comment.