Implementing reduction for non-commutative operations
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 487
- Avg merge
- 2d 7h
- Merged PRs (30d)
- 296
Description
I'd like to investigate implementing a reduction for associative, but non-commutative operations. Related to https://github.com/NVIDIA/cccl/issues/774
The this kind of algorithm comes in handy when establishing the global context in parsing a regular language by chunks in parallel, like the token languages for JSON, CSV etc. in cuDF. On a smaller scale, it can also be used to parse digits into integers, if you think of a segmented version of the algorithm.
Interface-wise, I don't see a reason do deviate from what `DeviceReduce::Reduce` does, so I imagined something like
```cpp
template
static CUB_RUNTIME_FUNCTION cudaError_t cub::DeviceReduce::ReduceOrdered(
void * d_temp_storage,
size_t & temp_storage_bytes,
InputIteratorT d_in,
OutputIteratorT d_out,
int num_items,
ReductionOpT reduction_op,
T init,
cudaStream_t stream = 0,
bool debug_synchronous = false
)
```
For the implementation, I would use the following hierarchy:
* Host-side: Call the reduction kernel once to compute block-local reductions, and a second time in block partials to compute the final result.
* Kernel: Pretty much identical to `DeviceReduceKernel`
* Agent: Similar to `AgentReduce`, only using `LoadDirectBlocked` instead of `LoadDirectStriped`, and doing a local sequential reduction before a block reduction.
* Block reduction: I haven't dug deep enough into the different specializations to see where/if they reorder inputs, but at least BlockReduceRaking supports it according to the documentation.
Contributor guide
Assessment
This issue has not been assessed yet.