[FEA]: POC Hierarchical Parallelism
- Dominant language
- C++
- Stars
- 2.5k
- Forks
- 486
- Avg merge
- 2d 6h
- Merged PRs (30d)
- 295
Description
### Is this a duplicate?
- [x] I confirmed there appear to be no [duplicate issues](https://github.com/NVIDIA/cccl/issues) for this request and that I agree to the [Code of Conduct](CODE_OF_CONDUCT.md)
### Area
CUB
### Is your feature request related to a problem? Please describe.
Sequence algorithms are frequently insufficient for SOL implementation. Occasionally, user needs parallelism while processing a single input element. For instance, in the transformation `thrust::transform(begin, end, [](cuda::std::array element) { ... })` we might want to process a large element with multiple threads.
Another limitation is fusion of the algorithms with mismatching shapes. cudf has a use case for a custom kernel where a single thread of a thread block computes masks for the elements to be processed, and after synchronization, threads of a thread block apply transformation for unmasked elements cooperatively. This requires synchronization within a parallel algorithm and revealing threads-to-elements mapping, which is not possible with sequence algorithms, so user has to write a custom kernel or follow inefficient 2-kernel approach:
```c++
std::bitset<8> compute_group_mask(int group_index);
float compute_value(float value);
std::vector values(num_items);
std::vector> masks(num_items / 8);
auto iota = std::ranges::iota_view{0, num_items / 8};
std::transform(iota.begin(), iota.end(), masks.begin(), [](int group_index) {
return compute_group_mask(group_index);
});
iota = std::ranges::iota_view{0, num_items};
std::transform(iota.begin(), iota.end(), values.begin(), [&](int element_index) {
return masks[element_index / 8][element_index % 8]
? compute_value(values[element_index])
: values[element_index];
});
```
### Describe the solution you'd like
Common solution to described problems is called hierarchical parallelism in SYCL (`parallel_for_work_group`) and nested parallelism in TBB. We should investigate potential design for an interface like that.
### Describe alternatives you've considered
_No response_
### Additional context
_No response_
Contributor guide
Assessment
This issue has not been assessed yet.