NVIDIA / NVIDIA/cudf

[FEA] Improve cudf::apply_boolean_mask scalability and reduce kernel launch overhead

Open
#21,155 0 comments 0 reactions 0 assignees View on GitHub
feature request libcudf
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

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.