NVIDIA / NVIDIA/cccl

[FEA]: POC Hierarchical Parallelism

Open
#6,410 0 comments 0 reactions 1 assignee Claimed by @NaderAlAwar View on GitHub
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

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.