[BUG] `compute_column_jit` regression on fp64 arithmetic since #21704
- Dominant language
- C++
- Stars
- 9.8k
- Forks
- 1.1k
- Avg merge
- 3d 6m
- Merged PRs (30d)
- 278
Description
**Describe the bug**
Since PR #21704 ("[FEA] Support Multi-Output JIT Transforms", merged 2026-04-20), `cudf::compute_column_jit` shows an 8–20% GPC-cycle regression on the two fp64 arithmetic expressions used in TPC-H Q1:
- `disc_price = extended_price * (1 - discount)` — 2 columns + 1 scalar
- `charge = extended_price * (1 - discount) * (1 + tax)` — 3 columns + 2 scalars
The slowdown comes from the new `cudf::jit::transform_kernel` template
(`cuda::std::apply` + `cuda::std::tuple` accessor expansion) increasing
per-thread register usage. For the 3-column expression this pushes registers
from 32 → 40, which drops the block size chosen by
`configure_1d_max_occupancy` from 1024 → 768 and theoretical occupancy from
100% → 75%.
**Steps/Code to reproduce bug**
Bisect setup — everything except the cudf commit is held identical between
the "pre" and "post" builds:
- Pre-#21704 cudf: `048787d70a` (parent of the #21704 merge)
- Post-#21704 cudf: `dc4e64bc02` (#21704 merge commit)
- `rapids-cmake-sha=4836a1a` (rapids-cmake `main` HEAD on 2026-04-16, four days before #21704 merged)
- CCCL 3.4.0 at `c5594eb1148649308b1996f75a63cc083d354cb2` (pinned by that rapids-cmake commit's `versions.json`)
- RMM `release/26.04`
- `-DCMAKE_CUDA_ARCHITECTURES=100-real -DCMAKE_BUILD_TYPE=Release`
- CUDA 12.9, GCC 14 (conda-forge), NVIDIA B200 (sm_100)
Minimal driver (no external dependency beyond libcudf):
```cpp
#include
#include
#include
#include
#include
#include
#include
#include
int main(int argc, char** argv)
{
auto const num_rows = static_cast(std::strtoll(argv[1], nullptr, 10));
std::string_view const ex = argv[2]; // "disc_price" or "charge"
cudf::numeric_scalar ep_v{100.0}, disc_v{0.05}, tax_v{0.08};
auto ep = cudf::sequence(num_rows, ep_v);
auto disc = cudf::sequence(num_rows, disc_v);
auto tax = cudf::sequence(num_rows, tax_v);
auto tbl = cudf::table_view{{ep->view(), disc->view(), tax->view()}};
namespace ast = cudf::ast;
auto ep_ref = ast::column_reference{0};
auto disc_ref = ast::column_reference{1};
auto tax_ref = ast::column_reference{2};
auto one = ast::literal{1.0};
auto one_sub_disc = ast::operation{ast::ast_operator::SUB, one, disc_ref};
auto one_add_tax = ast::operation{ast::ast_operator::ADD, one, tax_ref};
auto disc_price = ast::operation{ast::ast_operator::MUL, ep_ref, one_sub_disc};
auto charge = ast::operation{ast::ast_operator::MUL, disc_price, one_add_tax};
auto const& expr = (ex == "charge") ? charge : disc_price;
// Warmup + measured.
for (int i = 0; i < 2; ++i) {
auto result = cudf::compute_column_jit(tbl, expr);
cudaDeviceSynchronize();
}
return 0;
}
```
Profiling command. The kernel name changes across the PR, so the regex
differs between the two builds:
- pre-#21704 kernel: `void cudf::transformation::jit::kernel<...>` → `regex:.*transformation::jit::kernel.*`
- post-#21704 kernel: `void cudf::jit::transform_kernel<...>` → `regex:.*jit::transform_kernel.*`
`--launch-skip 1 --launch-count 1` skips the warmup call (which includes JIT
compilation) and captures the second (steady-state) launch.
```bash
METRICS=gpc__cycles_elapsed.avg,launch__registers_per_thread,launch__block_size
# --- point LD_LIBRARY_PATH at the pre-#21704 libcudf build tree ---
# 1. pre-#21704 + disc_price
ncu --target-processes all --kernel-name-base demangled \
--kernel-name 'regex:.*transformation::jit::kernel.*' \
--launch-skip 1 --launch-count 1 --metrics $METRICS \
-o report_disc_price_pre ./driver 600037902 disc_price
# 2. pre-#21704 + charge
ncu --target-processes all --kernel-name-base demangled \
--kernel-name 'regex:.*transformation::jit::kernel.*' \
--launch-skip 1 --launch-count 1 --metrics $METRICS \
-o report_charge_pre ./driver 600037902 charge
# --- swap LD_LIBRARY_PATH to the post-#21704 libcudf build tree ---
# 3. post-#21704 + disc_price
ncu --target-processes all --kernel-name-base demangled \
--kernel-name 'regex:.*jit::transform_kernel.*' \
--launch-skip 1 --launch-count 1 --metrics $METRICS \
-o report_disc_price_post ./driver 600037902 disc_price
# 4. post-#21704 + charge
ncu --target-processes all --kernel-name-base demangled \
--kernel-name 'regex:.*jit::transform_kernel.*' \
--launch-skip 1 --launch-count 1 --metrics $METRICS \
-o report_charge_post ./driver 600037902 charge
```
Read the numbers back with:
```bash
ncu --import report_disc_price_pre.ncu-repz --page details \
| grep -E "block_size|registers_per_thread|gpc__cycles"
```
**Expected behavior**
For simple fp64 arithmetic expressions with 2–5 accessor template parameters,
`compute_column_jit` should not regress meaningfully in GPC cycles across a
single PR. In particular, register-per-thread and launch block size should
stay stable for a workload as common as `extended_price * (1 - discount)`.
**Environment overview**
- Environment location: **Bare-metal**
- Method of cuDF install: **from source**
- Reproducing the bisect requires two separate builds at the commits above,
sharing the same conda env, `rapids-cmake` pin, RMM checkout, and CCCL
pin.
**Environment details**
Click here to see environment details
**git***
commit dc4e64bc02efb749d2789ce41eb3a504fdc7d9bd (HEAD)
Author: Basit Ayantunde
Date: Mon Apr 20 21:16:22 2026 +0100
[FEA] Support Multi-Output JIT Transforms (#21704)
This Pull-Request adds support for multi-output JIT Transforms (0-N).
It also improves the efficiency of string output creation by supporting pre-allocating string offsets, which helps to reduce memory usage and efficiency in the case where the size of the string columns can be cheaply pre-computed or known ahead of time.
## Summary
- [x] Support output of multiple columns and allow for batching. i.e. null-counting, stencil creation, and other pre/post-transform computations.
- [x] Remove intermediate-nullmask boolean creation in null-aware transforms
- [x] Make the Transform kernel streaming-oriented (allowing for non-coherent memory load/stores via `--restrict`)
- [x] Ban memory spaces in the PTX UDF, this has prevented us from passing stack variable pointers to UDFs as it attempts to load from an incorrect memory space. Numba UDFs do not produce memory-space specific UDFs when compiled without array arguments.
- [x] Made `column_device_view` types bytewise-compatible to allow aliasing in the kernel (via wrappers)
- [x] Merged all `transform` kernels into a single kernel
- [x] Refactored and simplified the Transforms code
- [x] Added more assertions and support checks to the Transforms code
- [x] Added and updated tests for these changes
- [x] Validated performance impact of these changes
Closes https://github.com/rapidsai/cudf/issues/20155
Authors:
- Basit Ayantunde (https://github.com/lamarrr)
Approvers:
- Muhammad Haseeb (https://github.com/mhaseeb123)
- Nghia Truong (https://github.com/ttnghia)
URL: https://github.com/rapidsai/cudf/pull/21704
**git submodules***
***OS Information***
NAME="main-2026-07-31-130241-25.05.5"
IMAGE=""
MIRROR_IMAGE=""
DGX_NAME="DGX Server"
DGX_PRETTY_NAME="NVIDIA DGX Server"
DGX_SWBUILD_DATE="2025-11-12-02-34-35"
DGX_SWBUILD_VERSION="7.3.1"
DGX_COMMIT_ID="01871cb"
DGX_PLATFORM="DGX Server for DGX B200"
DGX_SERIAL_NUMBER=""
DISTRIB_ID=Ubuntu
DISTRIB_RELEASE=24.04
DISTRIB_CODENAME=noble
DISTRIB_DESCRIPTION="Ubuntu 24.04.4 LTS"
PRETTY_NAME="Ubuntu 24.04.4 LTS"
NAME="Ubuntu"
VERSION_ID="24.04"
VERSION="24.04.4 LTS (Noble Numbat)"
VERSION_CODENAME=noble
ID=ubuntu
ID_LIKE=debian
HOME_URL="https://www.ubuntu.com/"
SUPPORT_URL="https://help.ubuntu.com/"
BUG_REPORT_URL="https://bugs.launchpad.net/ubuntu/"
PRIVACY_POLICY_URL="https://www.ubuntu.com/legal/terms-and-policies/privacy-policy"
UBUNTU_CODENAME=noble
LOGO=ubuntu-logo
Linux 6.8.0-110-generic #110-Ubuntu SMP PREEMPT_DYNAMIC Thu Mar 19 15:09:20 UTC 2026 x86_64 x86_64 x86_64 GNU/Linux
***GPU Information***
Sun Aug 2 23:35:31 2026
+-----------------------------------------------------------------------------------------+
| NVIDIA-SMI 595.58.03 Driver Version: 595.58.03 CUDA Version: 13.2 |
+-----------------------------------------+------------------------+----------------------+
| GPU Name Persistence-M | Bus-Id Disp.A | Volatile Uncorr. ECC |
| Fan Temp Perf Pwr:Usage/Cap | Memory-Usage | GPU-Util Compute M. |
| | | MIG M. |
|=========================================+========================+======================|
| 0 NVIDIA B200 On | 00000000:61:00.0 Off | 0 |
| N/A 42C P0 246W / 1000W | 0MiB / 183359MiB | 0% Default |
| | | Disabled |
+-----------------------------------------+------------------------+----------------------+
+-----------------------------------------------------------------------------------------+
| Processes: |
| GPU GI CI PID Type Process name GPU Memory |
| ID ID Usage |
|=========================================================================================|
| No running processes found |
+-----------------------------------------------------------------------------------------+
***CPU***
Architecture: x86_64
CPU op-mode(s): 32-bit, 64-bit
Address sizes: 52 bits physical, 57 bits virtual
Byte Order: Little Endian
CPU(s): 224
On-line CPU(s) list: 0-223
Vendor ID: GenuineIntel
Model name: INTEL(R) XEON(R) PLATINUM 8570
CPU family: 6
Model: 207
Thread(s) per core: 2
Core(s) per socket: 56
Socket(s): 2
Stepping: 2
CPU(s) scaling MHz: 25%
CPU max MHz: 4000.0000
CPU min MHz: 800.0000
BogoMIPS: 4200.00
Flags: fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush dts acpi mmx fxsr sse sse2 ss ht tm pbe syscall nx pdpe1gb rdtscp lm constant_tsc art arch_perfmon pebs bts rep_good nopl xtopology nonstop_tsc cpuid aperfmperf tsc_known_freq pni pclmulqdq dtes64 monitor ds_cpl vmx smx est tm2 ssse3 sdbg fma cx16 xtpr pdcm pcid dca sse4_1 sse4_2 x2apic movbe popcnt tsc_deadline_timer aes xsave avx f16c rdrand lahf_lm abm 3dnowprefetch cpuid_fault epb cat_l3 cat_l2 cdp_l3 intel_ppin cdp_l2 ssbd mba ibrs ibpb stibp ibrs_enhanced tpr_shadow flexpriority ept vpid ept_ad fsgsbase tsc_adjust bmi1 avx2 smep bmi2 erms invpcid cqm rdt_a avx512f avx512dq rdseed adx smap avx512ifma clflushopt clwb intel_pt avx512cd sha_ni avx512bw avx512vl xsaveopt xsavec xgetbv1 xsaves cqm_llc cqm_occup_llc cqm_mbm_total cqm_mbm_local split_lock_detect user_shstk avx_vnni avx512_bf16 wbnoinvd dtherm ida arat pln pts hwp hwp_act_window hwp_epp hwp_pkg_req vnmi avx512vbmi umip pku ospke waitpkg avx512_vbmi2 gfni vaes vpclmulqdq avx512_vnni avx512_bitalg tme avx512_vpopcntdq la57 rdpid bus_lock_detect cldemote movdiri movdir64b enqcmd fsrm md_clear serialize tsxldtrk pconfig arch_lbr ibt amx_bf16 avx512_fp16 amx_tile amx_int8 flush_l1d arch_capabilities ibpb_exit_to_user
Virtualization: VT-x
L1d cache: 5.3 MiB (112 instances)
L1i cache: 3.5 MiB (112 instances)
L2 cache: 224 MiB (112 instances)
L3 cache: 600 MiB (2 instances)
NUMA node(s): 2
NUMA node0 CPU(s): 0-55,112-167
NUMA node1 CPU(s): 56-111,168-223
Vulnerability Gather data sampling: Not affected
Vulnerability Indirect target selection: Not affected
Vulnerability Itlb multihit: Not affected
Vulnerability L1tf: Not affected
Vulnerability Mds: Not affected
Vulnerability Meltdown: Not affected
Vulnerability Mmio stale data: Not affected
Vulnerability Reg file data sampling: Not affected
Vulnerability Retbleed: Not affected
Vulnerability Spec rstack overflow: Not affected
Vulnerability Spec store bypass: Mitigation; Speculative Store Bypass disabled via prctl
Vulnerability Spectre v1: Mitigation; usercopy/swapgs barriers and __user pointer sanitization
Vulnerability Spectre v2: Mitigation; Enhanced / Automatic IBRS; IBPB conditional; PBRSB-eIBRS SW sequence; BHI BHI_DIS_S
Vulnerability Srbds: Not affected
Vulnerability Tsa: Not affected
Vulnerability Tsx async abort: Not affected
Vulnerability Vmscape: Mitigation; IBPB before exit to userspace
***CMake***
/usr/bin/cmake
cmake version 3.28.3
CMake suite maintained and supported by Kitware (kitware.com/cmake).
***g++***
/usr/bin/g++
g++ (Ubuntu 13.3.0-6ubuntu2~24.04.1) 13.3.0
Copyright (C) 2023 Free Software Foundation, Inc.
This is free software; see the source for copying conditions. There is NO
warranty; not even for MERCHANTABILITY or FITNESS FOR A PARTICULAR PURPOSE.
***nvcc***
/bin/nvcc
nvcc: NVIDIA (R) Cuda compiler driver
Copyright (c) 2005-2025 NVIDIA Corporation
Built on Tue_May_27_02:21:03_PDT_2025
Cuda compilation tools, release 12.9, V12.9.86
Build cuda_12.9.r12.9/compiler.36037853_0
***Python***
/bin/python
Python 3.12.13
***Environment Variables***
PATH :
LD_LIBRARY_PATH : .:/usr/lib
NUMBAPRO_NVVM :
NUMBAPRO_LIBDEVICE :
CONDA_PREFIX :
PYTHON_PATH :
***conda packages***
/conda
# packages in environment at :
#
# Name Version Build Channel
_openmp_mutex 4.5 20_gnu conda-forge
annotated-types 0.7.0 pypi_0 pypi
antlr4-python3-runtime 4.9.3 pypi_0 pypi
anyio 4.13.0 pypi_0 pypi
aws-c-auth 0.7.25 h15d0e8c_6 conda-forge
aws-c-cal 0.7.3 h8dac057_2 conda-forge
aws-c-common 0.9.27 h4bc722e_0 conda-forge
aws-c-compression 0.2.19 haa50ccc_0 conda-forge
aws-c-event-stream 0.4.3 h570d160_0 conda-forge
aws-c-http 0.8.7 h1c59cda_5 conda-forge
aws-c-io 0.14.18 hf5b9b93_6 conda-forge
aws-c-mqtt 0.10.4 hc14a930_17 conda-forge
aws-c-s3 0.6.4 h558cea2_8 conda-forge
aws-c-sdkutils 0.1.19 h038f3f9_2 conda-forge
aws-checksums 0.1.18 h038f3f9_10 conda-forge
aws-crt-cpp 0.27.5 h6e4e78f_8 conda-forge
aws-sdk-cpp 1.11.379 hce093eb_4 conda-forge
azure-core-cpp 1.13.0 h935415a_0 conda-forge
azure-identity-cpp 1.8.0 hd126650_2 conda-forge
azure-storage-blobs-cpp 12.12.0 hd2e3451_0 conda-forge
azure-storage-common-cpp 12.7.0 h10ac4d7_1 conda-forge
azure-storage-files-datalake-cpp 12.11.0 h325d260_1 conda-forge
binutils 2.45.1 default_h4852527_102 conda-forge
binutils_impl_linux-64 2.45.1 default_hfdba357_102 conda-forge
binutils_linux-64 2.45.1 default_h4852527_102 conda-forge
boost-cpp 1.85.0 h3c6214e_4 conda-forge
bzip2 1.0.8 hda65f42_9 conda-forge
c-ares 1.34.6 hb03c661_0 conda-forge
c-compiler 1.11.0 h4d9bdce_0 conda-forge
ca-certificates 2026.5.20 hbd8a1cb_0 conda-forge
certifi 2026.2.25 pypi_0 pypi
cffi 2.0.0 py312h460c074_1 conda-forge
cfgv 3.5.0 pyhd8ed1ab_0 conda-forge
charset-normalizer 3.4.7 pypi_0 pypi
clang-format 20.1.4 default_h1df26ce_0 conda-forge
clang-format-20 20.1.4 default_h1df26ce_0 conda-forge
clang-tools 20.1.4 default_h1df26ce_0 conda-forge
cmake 4.3.1 hc85cc9f_0 conda-forge
conda-gcc-specs 14.3.0 he8ccf15_18 conda-forge
conda-pack 0.9.1 pyhcf101f3_0 conda-forge
contourpy 1.3.3 pypi_0 pypi
cuda-cccl 12.9.27 0 nvidia
cuda-cccl_linux-64 12.9.27 0 nvidia
cuda-crt-dev_linux-64 12.9.86 ha770c72_2 conda-forge
cuda-crt-tools 12.9.86 ha770c72_2 conda-forge
cuda-cudart 12.9.79 0 nvidia
cuda-cudart-dev 12.9.79 0 nvidia
cuda-cudart-dev_linux-64 12.9.79 0 nvidia
cuda-cudart-static 12.9.79 0 nvidia
cuda-cudart-static_linux-64 12.9.79 0 nvidia
cuda-cudart_linux-64 12.9.79 0 nvidia
cuda-cupti 12.9.79 0 nvidia
cuda-cupti-dev 12.9.79 0 nvidia
cuda-driver-dev_linux-64 12.9.79 h3f2d84a_0 conda-forge
cuda-nvcc 12.9.86 hcdd1206_6 conda-forge
cuda-nvcc-dev_linux-64 12.9.86 he91c749_2 conda-forge
cuda-nvcc-impl 12.9.86 h85509e4_2 conda-forge
cuda-nvcc-tools 12.9.86 he02047a_2 conda-forge
cuda-nvcc_linux-64 12.9.86 he0b4e1d_6 conda-forge
cuda-nvml-dev 12.9.79 1 nvidia
cuda-nvrtc 12.9.86 0 nvidia
cuda-nvrtc-dev 12.9.86 0 nvidia
cuda-nvtx 12.9.79 0 nvidia
cuda-nvtx-dev 12.9.79 0 nvidia
cuda-nvvm-dev_linux-64 12.9.86 ha770c72_2 conda-forge
cuda-nvvm-impl 12.9.86 h4bc722e_2 conda-forge
cuda-nvvm-tools 12.9.86 h4bc722e_2 conda-forge
cuda-version 12.9 3 nvidia
cxx-compiler 1.11.0 hfcd1e18_0 conda-forge
cycler 0.12.1 pypi_0 pypi
database-benchmarking-tools 0.0.3 pypi_0 pypi
distlib 0.4.0 pyhd8ed1ab_0 conda-forge
distro 1.9.0 pypi_0 pypi
et-xmlfile 2.0.0 pypi_0 pypi
exp-oven 0.7.8 pypi_0 pypi
filelock 3.29.0 pyhd8ed1ab_0 conda-forge
fmt 11.1.4 h07f6e7f_1 conda-forge
fonttools 4.63.0 pypi_0 pypi
gcc 14.3.0 h0dff253_18 conda-forge
gcc_impl_linux-64 14.3.0 hbdf3cc3_18 conda-forge
gcc_linux-64 14.3.0 h298d278_23 conda-forge
gflags 2.2.2 h5888daf_1005 conda-forge
glog 0.7.1 hbabe93e_0 conda-forge
gqe-bench 0.0.1 pypi_0 pypi
gxx 14.3.0 h76987e4_18 conda-forge
gxx_impl_linux-64 14.3.0 h2185e75_18 conda-forge
gxx_linux-64 14.3.0 h91b0f8e_23 conda-forge
h11 0.16.0 pypi_0 pypi
httpcore 1.0.9 pypi_0 pypi
httpx 0.28.1 pypi_0 pypi
icu 75.1 he02047a_0 conda-forge
identify 2.6.19 pyhd8ed1ab_0 conda-forge
idna 3.11 pypi_0 pypi
importlib-metadata 8.8.0 pyhcf101f3_0 conda-forge
jiter 0.15.0 pypi_0 pypi
json5 0.14.0 pypi_0 pypi
kernel-headers_linux-64 4.18.0 he073ed8_9 conda-forge
keyutils 1.6.3 hb9d3cd8_0 conda-forge
kiwisolver 1.5.0 pypi_0 pypi
krb5 1.22.2 ha1258a1_0 conda-forge
ld_impl_linux-64 2.45.1 default_hbd61a6d_102 conda-forge
libabseil 20240116.2 cxx17_he02047a_1 conda-forge
libarrow 16.1.0 h03aeac6_19_cpu conda-forge
libarrow-acero 16.1.0 he02047a_19_cpu conda-forge
libarrow-dataset 16.1.0 he02047a_19_cpu conda-forge
libarrow-substrait 16.1.0 hc9a23c6_19_cpu conda-forge
libblas 3.11.0 6_h4a7cf45_openblas conda-forge
libboost 1.85.0 h0ccab89_4 conda-forge
libboost-devel 1.85.0 h00ab1b0_4 conda-forge
libboost-headers 1.85.0 ha770c72_4 conda-forge
libbrotlicommon 1.1.0 hb03c661_4 conda-forge
libbrotlidec 1.1.0 hb03c661_4 conda-forge
libbrotlienc 1.1.0 hb03c661_4 conda-forge
libcap 2.77 hd0affe5_1 conda-forge
libcblas 3.11.0 6_h0358290_openblas conda-forge
libclang-cpp20.1 20.1.8 default_h99862b1_14 conda-forge
libclang13 21.1.0 default_h746c552_1 conda-forge
libcrc32c 1.1.2 h9c3ff4c_0 conda-forge
libcurl 8.19.0 hcf29cc6_0 conda-forge
libedit 3.1.20250104 pl5321h7949ede_0 conda-forge
libev 4.33 hd590300_2 conda-forge
libevent 2.1.12 hf998b51_1 conda-forge
libexpat 2.7.5 hecca717_0 conda-forge
libfabric 2.5.1 ha770c72_0 conda-forge
libfabric1 2.5.1 hf621623_0 conda-forge
libffi 3.5.2 h3435931_0 conda-forge
libgcc 15.2.0 he0feb66_18 conda-forge
libgcc-devel_linux-64 14.3.0 hf649bbc_118 conda-forge
libgcc-ng 15.2.0 h69a702a_18 conda-forge
libgfortran 15.2.0 h69a702a_18 conda-forge
libgfortran5 15.2.0 h68bc16d_18 conda-forge
libgomp 15.2.0 he0feb66_18 conda-forge
libgoogle-cloud 2.28.0 h26d7fe4_0 conda-forge
libgoogle-cloud-storage 2.28.0 ha262f82_0 conda-forge
libgrpc 1.62.2 h15f2491_0 conda-forge
libhwloc 2.12.1 default_h3d81e11_1000 conda-forge
libiconv 1.18 h3b78370_2 conda-forge
liblapack 3.11.0 6_h47877c9_openblas conda-forge
libllvm20 20.1.8 hecd9e04_0 conda-forge
libllvm21 21.1.0 hecd9e04_0 conda-forge
liblzma 5.8.3 hb03c661_0 conda-forge
liblzma-devel 5.8.3 hb03c661_0 conda-forge
libnghttp2 1.68.1 h877daf1_0 conda-forge
libnl 3.11.0 hb9d3cd8_0 conda-forge
libnsl 2.0.1 hb9d3cd8_1 conda-forge
libnuma 2.0.18 hb03c661_3 conda-forge
libnvcomp 5.2.0.10 hb7e823c_0 conda-forge
libnvcomp-dev 5.2.0.10 hb7e823c_0 conda-forge
libnvjitlink 12.9.86 0 nvidia
libnvjitlink-dev 12.9.86 0 nvidia
libnvptxcompiler-dev 12.9.86 ha770c72_2 conda-forge
libnvptxcompiler-dev_linux-64 12.9.86 ha770c72_2 conda-forge
libnvshmem-dev 3.2.5 h2b6041c_0 nvidia
libnvshmem-static 3.2.5 hf45bede_0 nvidia
libnvshmem3 3.2.5 h0b1ad75_0 nvidia
libopenblas 0.3.32 pthreads_h94d23a6_0 conda-forge
libparquet 16.1.0 h9e5060d_19_cpu conda-forge
libpmix 5.0.8 h4bd6b51_2 conda-forge
libprotobuf 4.25.3 hd5b35b9_1 conda-forge
libre2-11 2023.09.01 h5a48ba9_2 conda-forge
libsanitizer 14.3.0 h8f1669f_18 conda-forge
libsqlite 3.53.0 h0c1763c_0 conda-forge
libssh2 1.11.1 hcf80075_0 conda-forge
libstdcxx 15.2.0 h934c35e_18 conda-forge
libstdcxx-devel_linux-64 14.3.0 h9f08a49_118 conda-forge
libstdcxx-ng 15.2.0 hdf11a46_18 conda-forge
libsystemd0 260.1 h6569c3e_0 conda-forge
libthrift 0.19.0 hb90f79a_1 conda-forge
libudev1 260.1 h6569c3e_0 conda-forge
liburing 2.14 hb700be7_0 conda-forge
libutf8proc 2.8.0 hf23e847_1 conda-forge
libuuid 2.42 h5347b49_0 conda-forge
libuv 1.51.0 hb03c661_1 conda-forge
libxcrypt 4.4.36 hd590300_1 conda-forge
libxml2 2.13.9 h04c0eec_0 conda-forge
libzlib 1.3.2 h25fd6f3_2 conda-forge
lz4-c 1.9.4 hcb278e6_0 conda-forge
matplotlib 3.10.9 pypi_0 pypi
mpi 1.0.1 openmpi conda-forge
mpi4py 4.1.1 py312hd140a38_103 conda-forge
nccl 2.29.7 h321a1a8_0 nvidia
ncurses 6.5 h2d0b736_3 conda-forge
nodeenv 1.10.0 pyhd8ed1ab_0 conda-forge
numactl 2.0.18 hb03c661_3 conda-forge
numpy 2.4.3 py312h33ff503_0 conda-forge
nvidia-ml-py 13.595.45 pypi_0 pypi
nvtx 0.2.15 pypi_0 pypi
omegaconf 2.3.0 pypi_0 pypi
openai 2.41.1 pypi_0 pypi
openmpi 5.0.8 h2fe1745_110 conda-forge
openpyxl 3.1.5 pypi_0 pypi
openssl 3.6.3 h35e630c_0 conda-forge
orc 2.0.1 h17fec99_1 conda-forge
packaging 26.1 pyhc364b38_0 conda-forge
pandas 3.0.2 py312h8ecdadd_0 conda-forge
pillow 12.2.0 pypi_0 pypi
pip 26.0.1 pyh8b19718_0 conda-forge
platformdirs 4.9.6 pyhcf101f3_0 conda-forge
pre-commit 4.5.1 pyha770c72_0 conda-forge
pyarrow 16.1.0 py312h9cebb41_6 conda-forge
pyarrow-core 16.1.0 py312h01725c0_6_cpu conda-forge
pybind11 2.13.6 pyhc790b64_3 conda-forge
pybind11-global 2.13.6 pyh217bc35_3 conda-forge
pycparser 2.22 pyh29332c3_1 conda-forge
pydantic 2.13.4 pypi_0 pypi
pydantic-core 2.46.4 pypi_0 pypi
pyparsing 3.3.2 pypi_0 pypi
python 3.12.13 hd63d673_0_cpython conda-forge
python-dateutil 2.9.0.post0 pyhe01879c_2 conda-forge
python-discovery 1.2.2 pyhcf101f3_0 conda-forge
python_abi 3.12 8_cp312 conda-forge
pyyaml 6.0.3 py312h8a5da7c_1 conda-forge
rdma-core 61.0 h192683f_0 conda-forge
re2 2023.09.01 h7f4b329_2 conda-forge
readline 8.3 h853b02a_0 conda-forge
requests 2.33.1 pypi_0 pypi
rhash 1.4.6 hb9d3cd8_1 conda-forge
rust 1.94.0 h53717f1_0 conda-forge
rust-std-x86_64-unknown-linux-gnu 1.94.0 h2c6d0dc_0 conda-forge
s2n 1.5.1 h3400bea_0 conda-forge
setuptools 82.0.1 pyh332efcf_0 conda-forge
six 1.17.0 pyhe01879c_1 conda-forge
snappy 1.2.2 h03e3b7b_1 conda-forge
sniffio 1.3.1 pypi_0 pypi
spdlog 1.15.3 h10b92b3_0 conda-forge
sqlglot 30.4.3 pypi_0 pypi
sqlite 3.53.0 hbc0de68_0 conda-forge
sysroot_linux-64 2.28 h4ee821c_9 conda-forge
thrift-compiler 0.19.0 h59595ed_1 conda-forge
tk 8.6.13 noxft_h366c992_103 conda-forge
tqdm 4.68.2 pypi_0 pypi
typing-inspection 0.4.2 pypi_0 pypi
typing_extensions 4.15.0 pyhcf101f3_0 conda-forge
tzdata 2025c hc9c84f9_1 conda-forge
ucc 1.6.0 hb729f83_1 conda-forge
ucx 1.19.1 h63b5c0b_0 conda-forge
ukkonen 1.1.0 py312hd9148b4_0 conda-forge
urllib3 2.6.3 pypi_0 pypi
uv 0.10.4 h6dd6661_0 conda-forge
virtualenv 21.2.4 pyhcf101f3_0 conda-forge
wheel 0.46.3 pyhd8ed1ab_0 conda-forge
xz 5.8.3 ha02ee65_0 conda-forge
xz-gpl-tools 5.8.3 ha02ee65_0 conda-forge
xz-tools 5.8.3 hb03c661_0 conda-forge
yaml 0.2.5 h280c20c_3 conda-forge
zipp 3.23.1 pyhcf101f3_0 conda-forge
zlib 1.3.2 h25fd6f3_2 conda-forge
zstd 1.5.7 hb78ec9c_6 conda-forge
**Additional context**
**Measured on 600M rows / B200, single NCU-captured launch:**
| Expression | Version | Block | Regs | GPC cycles (avg) | Δ |
|---|---|---|---|---|---|
| `disc_price` (2 col + 1 scalar) | pre-#21704 | 1024 | 32 | 4,296,667 | — |
| `disc_price` | post-#21704 | 1024 | 32 | 4,644,362 | **+8.1%** |
| `charge` (3 col + 2 scalar) | pre-#21704 | 1024 | 32 | 5,486,700 | — |
| `charge` | post-#21704 | **768** | **40** | 6,560,561 | **+19.6%** |
> Rebuilding `post-#21704` with `{"-restrict", "--dopt=on"}` removed from
> `cpp/src/transform/transform.cu:182` changes both numbers by less than
> 0.5% (`disc_price`: 4,634,006 cycles; `charge`: 6,558,672 cycles). So the
> regression comes from the new template structure, not from those NVRTC
> options.
**Where the register pressure comes from**
- pre-#21704 kernel signature: `void cudf::transformation::jit::kernel(...)` — direct parameter-pack expansion of `Out::element(outputs, i)` / `In::element(inputs, i)...`
- post-#21704 kernel signature: `void cudf::jit::transform_kernel(...)` — wraps accessors in `cuda::std::tuple` and dispatches through `cuda::std::apply` + `type_list::map` per row
The extra template layer keeps more intermediate objects live per row, which
pushes register usage past the threshold that would keep block size at 1024
for sm_100.
**Suggested next steps (not blocking)**
1. **CI coverage:** Consider adding `disc_price` / `charge`-shaped fp64
expressions to `TRANSFORM_BENCH`. The current fixture set doesn't appear
to exercise the common Q1 aggregation pattern, so a similar regression
would not be flagged at merge time.
2. **Template fast path:** In `cpp/src/transform/jit/kernel.cu`, a fast path
for the single-output / small-input-list case that skips the
`cuda::std::tuple` + `cuda::std::apply` layer would preserve multi-output
support while restoring pre-#21704 codegen for the common case.
Contributor guide
Assessment
This issue has not been assessed yet.