NVIDIA / NVIDIA/cudf

[FEA] Implement `cudf::batch_concatenate` to concatenate tables more efficiently with less kernel calls

Open
#21,115 0 comments 0 reactions 1 assignee Claimed by @ttnghia 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::concatenate` function processes columns recursively, launching separate kernels for data copy and mask concatenation for each being concatenated columns. For deeply nested struct columns or when concatenating many columns, this results in significant kernel launch overhead. This becomes a bottleneck in workloads that frequently concatenate columns, such as:
- DataFrame operations that require vertical concatenation
- Aggregation operations that combine partial results
- ETL pipelines with many small column chunks

**Describe the solution you'd like**

Add a new `batch_concatenate` API that:

1. **Batches data copy operations**: Collects all data copy operations across all nesting levels and executes them using `cub::DeviceMemcpy::Batched` in a single kernel launch

2. **Batches mask concatenation**: Processes bitmask concatenation for all columns with reduced kernel launches

3. **Pre-computes null counts on host**: Since input columns already have their null counts computed, the output null count can be calculated as a simple sum on the host, avoiding device-side atomic operations

**Describe alternatives you've considered**
N/A.

**Additional context**
We can totally remove the internal implementatio of `cudf::concatenate` and redirect the its call to `batch_concatenate` so we will have just one single implementation that is efficient for all cases.

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.