NVIDIA / NVIDIA/cccl

Implementing reduction for non-commutative operations

Open
#897 3 comments 0 reactions 0 assignees View on GitHub
cub
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

Open the contributing guide

Assessment

This issue has not been assessed yet.

Get new issues in your inbox

A short digest of beginner-friendly GitHub issues.