Repository navigation
[EPIC] Use programmatic dependent launch in all CUB algorithms #3115
Description
Programmatic dependent launch (PDL) allows subsequent kernels to skip a device-wide synchronization (executed on the host) and instead synchronize on the device at a source location specified by the programmer. This can eliminate bubbles inside a stream when a kernel is finishing until the next kernel can ramp up.
We should evaluate where we can use this feature in CUB and then add it wherever it shows a benefit.
- DeviceFind (implemented from the start)
- DeviceHistogram: Use PDL in DeviceHistogram #6367
- DeviceMergeSort: Use programmatic dependent launch in CUB merge sort #3114
- DeviceScan: Use PDL in cub::DeviceScan #6639
- Use PDL in
cub::DeviceTransform#5250 - Use PDL in
DeviceRadixSort#8765 - DeviceAdjacentDifference
- DeviceCopy
- DeviceFor
- DeviceMemcpy
- DeviceMerge
- DevicePartition
- DeviceReduce
- DeviceRunLengthEncode
- DeviceSegmentedRadixSort
- DeviceSegmentedReduce
- DeviceSegmentedSort
- DeviceSelect
- DeviceTopK
Reacted by icyveins7
Metadata
Metadata
Assignees
Type
Projects
- Status
Activity
[-]Use programmatic dependent launch in all CUB algorithms[/-][+][EPIC] Use programmatic dependent launch in all CUB algorithms[/+]gevtushenko commentedon Dec 11, 2024
DeviceRadixSort#8765cub::DeviceTransform#5250charan-003 commentedon Oct 6, 2026
Hi @bernhardmgruber , I was reading about PDL and it seems really interesting. I’d like to try implementing it for DeviceCopy if it’s still available
bernhardmgruber commentedon Oct 6, 2026
Sure, give it a try!