NVIDIA / NVIDIA/cccl

Provide Run-Length Decode API

Open
#599 1 comment 3 reactions 0 assignees View on GitHub
Dominant language
C++
Stars
2.5k
Forks
486
Avg merge
2d 6h
Merged PRs (30d)
295

Description

With the recent addition of batched copy, we should be able to provide RLD API. This API is needed for tuning infrastructure but should be useful elsewhere. Here's a simplified example from batched copy test to illustrate the idea:

```cpp
const int num_segments = 4;
const int num_items = 12;
thrust::device_vector segment_offsets = {0, 2, 8, num_items};
thrust::device_vector out(num_items);
int *d_out = thrust::raw_pointer_cast(out.data());
int *d_offsets = thrust::raw_pointer_cast(segment_offsets.data());

thrust::counting_iterator iota(0);
auto d_range_srcs = //
thrust::make_transform_iterator(iota, [](int i) {
return thrust::constant_iterator(i);
});
auto d_range_dsts = //
thrust::make_transform_iterator(d_offsets, [d_out](int offset) {
return d_out + offset;
});
auto d_range_sizes = //
thrust::make_transform_iterator(iota, [d_offsets](int i) {
return d_offsets[i + 1] - d_offsets[i];
});

std::uint8_t *d_temp_storage = nullptr;
std::size_t temp_storage_bytes = 0;
cub::DeviceCopy::Batched(d_temp_storage,
temp_storage_bytes,
d_range_srcs,
d_range_dsts,
d_range_sizes,
num_segments);
```

We'll need two overloads:

### Overloads
- [ ] Implement segment-size-based RLD
- [ ] Implement segment-offset-based RLD

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.