Skip to content
This repository was archived by the owner on Mar 21, 2024. It is now read-only.

Add libcu++ dependency; initial round of NV_IF_TARGET ports. #448

Merged
alliepiper merged 7 commits into from
May 17, 2022
Merged

Add libcu++ dependency; initial round of NV_IF_TARGET ports. #448

Changes from 1 commit
Commits
Show all changes
7 commits
Select commit Hold shift + click to select a range
f42070e
Add libcu++ dependency.
alliepiper Mar 23, 2022
f9beaa5
Remove checks for obsolete architectures.
alliepiper Mar 26, 2021
5799e96
Remove PTX arguments from compile-time contexts.
alliepiper Mar 26, 2021
c4299c4
Use NV_IF_TARGET to select between host/device/sm implementations.
alliepiper Nov 30, 2021
476c1b8
Don't use host-only functions in host-device contexts.
alliepiper May 11, 2021
f4d61fb
Use Thrust's kernel launch helper in DispatchRadixSort.
alliepiper Mar 23, 2022
4de961a
Skip large allocation tests that exceed device memory.
alliepiper Mar 24, 2022
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
Next Next commit
Use Thrust's kernel launch helper in DispatchRadixSort.
  • Loading branch information
@alliepiper
alliepiper committed May 16, 2022
commit f4d61fbb26c0c3b2ffd392bfbe6edbc60c0d5be3
14 changes: 10 additions & 4 deletions cub/device/dispatch/dispatch_radix_sort.cuh
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1344,9 +1344,14 @@ struct DispatchRadixSort :

// exclusive sums to determine starts
const int SCAN_BLOCK_THREADS = ActivePolicyT::ExclusiveSumPolicy::BLOCK_THREADS;
DeviceRadixSortExclusiveSumKernel<MaxPolicyT, OffsetT>
<<<num_passes, SCAN_BLOCK_THREADS, 0, stream>>>(d_bins);
if (CubDebug(error = cudaPeekAtLastError())) break;
error = THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron(
num_passes, SCAN_BLOCK_THREADS, 0, stream
).doit(DeviceRadixSortExclusiveSumKernel<MaxPolicyT, OffsetT>,
d_bins);
if (CubDebug(error))
{
break;
}

// use the other buffer if no overwrite is allowed
KeyT* d_keys_tmp = d_keys.Alternate();
Expand Down Expand Up @@ -1374,7 +1379,8 @@ struct DispatchRadixSort :
stream))) break;
auto onesweep_kernel = DeviceRadixSortOnesweepKernel<
MaxPolicyT, IS_DESCENDING, KeyT, ValueT, OffsetT, PortionOffsetT>;
errror = THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron(

error = THRUST_NS_QUALIFIER::cuda_cub::launcher::triple_chevron(
num_blocks, ONESWEEP_BLOCK_THREADS, 0, stream
).doit(onesweep_kernel,
d_lookback, d_ctrs + portion * num_passes + pass,
Expand Down