[FEA] Improve cudf::apply_boolean_mask scalability and reduce kernel launch overhead
- Dominant language
- C++
- Stars
- 9.8k
- Forks
- 1.1k
- Avg merge
- 3d 6m
- Merged PRs (30d)
- 278
Description
**Is your feature request related to a problem? Please describe.**
The current `cudf::apply_boolean_mask` implementation suffers from performance issues similar to those identified in #13509 for `cudf::gather`. When filtering tables with many columns or deeply nested types, the operation can take over 1 second even for moderately sized data.
The implementation in copy_if.cuh works as follows:
```c++
template
std::unique_ptr copy_if(table_view const& input,
Filter filter,
rmm::cuda_stream_view stream,
rmm::device_async_resource_ref mr)
{
// 1. Allocate indices buffer for all rows
auto indices = rmm::device_uvector(input.num_rows(), stream);
// 2. Stream compaction to build gather map
auto const indices_end =
thrust::copy_if(rmm::exec_policy_nosync(stream), begin, end, indices.begin(), filter);
// 3. Gather selected rows (processes columns one-by-one)
return cudf::detail::gather(input, map, ...);
}
```
_Bottlenecks:_
1. Gather scalability: The `cudf::detail::gather` call inherits the same column-by-column processing issue described in #13509. Kernel launch overhead grows linearly with the number of columns, and nested types (structs, lists) amplify this significantly.
2. Stream compaction overhead: `thrust::copy_if` performs a scan + scatter internally. For large row counts, this is a full pass over the data. When the mask has nulls, an additional validity bitmask read is required per element.
3. No batching across columns: Similar to the concatenate issue (#21115), each column is processed independently, missing opportunities to batch memory operations.
4. Memory allocation: The indices buffer is allocated for all `input.num_rows()`, even when selectivity is low. This can cause unnecessary memory pressure.
**Describe the solution you'd like**
1. Batch gather operations: Collect memory copy operations across all columns and execute them using `cub::DeviceMemcpy::Batched` or similar batched APIs to amortize kernel launch overhead (#13509) .
2. Consider allocating indices in chunks.
3. Optimize mask evaluation
**Describe alternatives you've considered**
- Redirect to improved gather: If #13509 is addressed with a batched gather implementation,` apply_boolean_mask` would automatically benefit since it delegates to gather. However, the stream compaction step would still benefit from dedicated optimizations.
**Additional context**
The `apply_boolean_mask` operation is frequently used in filtering pipelines. Performance improvements here would benefit workloads that:
- Filter wide tables with many columns
- Use deeply nested struct/list column types
- Perform repeated filtering operations in ETL pipelines
- Have high selectivity filters where gather cost dominates
Related Issues:
- #13509
- #21115
Contributor guide
Assessment
This issue has not been assessed yet.