Repository navigation
Conversation
…r<T>` (NVIDIA#24776) Wherever `rmm::device_uvector<T>` is used as a scratch buffer, use a `cuda::device_buffer<T>` instead, in order to reduce our reliance on RMM types in favor of CCCL types. Issue: rapidsai/build-planning#321 Authors: - Kyle Edwards (https://github.com/KyleFromNVIDIA) Approvers: - Muhammad Haseeb (https://github.com/mhaseeb123) - Yunsong Wang (https://github.com/PointKernel) URL: NVIDIA#24776
…DIA#24890) rapidsai/rapidsmpf#1289 added disk-to-disk support to `buffer_copy`. As a result, `table_chunk::copy` no longer throws `std::invalid_argument` when both the chunk and the target reservation are on disk. - Remove the stale `@throws` note from the `table_chunk::copy` docs. - Replace the `DiskToDiskCopyUnsupported` test with `DiskToDiskCopy`, which copies disk to disk and verifies the round trip back to device. - Stop skipping the second copy for the `DISK` spill target in the parameterized `StreamingTableChunk` test. Tested locally with `libcudf_streaming_single_tests --gtest_filter='*TableChunk*'` (76 passed). 🤖 Generated with [Claude Code](https://claude.com/claude-code) Authors: - Niranda Perera (https://github.com/nirandaperera) Approvers: - Richard (Rick) Zamora (https://github.com/rjzamora) - Muhammad Haseeb (https://github.com/mhaseeb123) - Tom Augspurger (https://github.com/TomAugspurger) - Bradley Dice (https://github.com/bdice) URL: NVIDIA#24890
… to Int128 (NVIDIA#24861) Polars computes the supertype of `UInt64` and a signed integer as `Int128` and casts both operands to it. libcudf has no `INT128` column type, so cudf-polars rejected the cast and the query fell back to the CPU. TPC-DS q9 hits this because `pl.len()` is `UInt64` in `polars-runtime-64` and is compared against a signed threshold. ```python import polars as pl lf = pl.LazyFrame( {"u": [1, 2**63, 2**64 - 1, None], "i": [-1, 5, 7, 8]}, schema={"u": pl.UInt64, "i": pl.Int64}, ) q = lf.select((pl.col("u") > pl.col("i")).alias("r")) q.collect(engine=pl.GPUEngine(raise_on_fail=True)) # before: NotImplementedError, now: works ``` ## Changes - `dsl/translate.py`: when a comparison (`==`, `!=`, `eq_missing`, `ne_missing`, `<`, `<=`, `>`, `>=`) or `is_between` has `Int128` operands that are casts of narrow integers or integer literals, cast the operands to `Decimal(38, 0)` instead. `DECIMAL128` with scale 0 holds every `Int128` value exactly. - `Int128` stays unsupported everywhere else (arithmetic, joins, `concat`, outputs, group keys), so those still fall back. This follows the existing `UInt128` special case. - Tests in `tests/expressions/test_int128_comparison.py`; `test_numeric_binop` now expects the comparison operators to translate for `UInt64` with signed integers (arithmetic still raises). ## Limitations - If common-subexpression elimination hoists a shared `cast(Int128)` into its own projection, the query still falls back. - Tested with Polars 2.0 on a GPU: `test_int128_comparison.py` and `test_numeric_binops.py` pass. Authors: - Matthew Murray (https://github.com/Matt711) Approvers: - Tom Augspurger (https://github.com/TomAugspurger) URL: NVIDIA#24861
For dependencies that cuDF's DSOs dynamically link to + expect to find from system installs / other wheels, we use `auditwheel repair --exclude` to prevent `auditwheel` from vendoring copies of whatever happened to be there at build time. Looking at something else today, I noticed that for a few of those exclusions cuDF is hard-coding a specific SOVERSION, e.g. `libnvcomp.so.6` for nvCOMP 6. This proposes using wildcards, to prevent bugs like "updating nvCOMP and forgot to bump the `auditwheel` exclusion" in the future. We have binary-size checks to help detect unwanted files, but if the size difference is small enough those might not catch bugs like that in CI. This also reduces the effort to bump dependencies a little 😁 Authors: - James Lamb (https://github.com/jameslamb) Approvers: - Bradley Dice (https://github.com/bdice) URL: NVIDIA#24648
Follow up to NVIDIA#24805 This PR separates row-ordering metadata from partition-ordering metadata in cudf-polars: - Migrates cudf-polars from `Ordering.locally_ordered` to explicit `RowOrdering` metadata. - Consolidates ordering propagation and remapping so row ordering and partition boundaries are handled together. - Fixes existing metadata-propagation gaps and bugs exposed by the migration. - Removes the now-unused `locally_ordered` field from `Ordering` and its Python bindings (technically a "breaking" change, but cudf-polars is the only known consumer). Most of `actor_graph/partitioning.py` is existing partitioning logic moved from `actor_graph/utils.py`. The remaining changes consolidate ordering propagation and extend it to `RowOrdering`. Authors: - Richard (Rick) Zamora (https://github.com/rjzamora) Approvers: - Tom Augspurger (https://github.com/TomAugspurger) URL: NVIDIA#24844
Part 1/n. This trivial PR migrates the Parquet and hybrid scan readers from `rmm::device_async_resource_ref` to `cudf::memory_resources`. Returned allocations use the output resource and reader intermediates use the temporary resource. Note that this is not an exhaustive migration PR. There's several libcudf utils, free functions and functors being used in Parquet internals that need to fully migrate in subsequent PRs before we can add tests. Authors: - Muhammad Haseeb (https://github.com/mhaseeb123) - Bradley Dice (https://github.com/bdice) Approvers: - Vukasin Milovanovic (https://github.com/vuule) - Yunsong Wang (https://github.com/PointKernel) - Bradley Dice (https://github.com/bdice) URL: NVIDIA#24849
… and dictionary (NVIDIA#24794) Replaces SFINAE with C++20 concepts in the binaryop, unary, round, transform, AST, strings, dictionary and nvtext sources, and switches `requires` clauses written with type traits to the equivalent concepts. Negated fallback overloads become unconstrained where the overloads have identical signatures. Where the constrained overload contains an extended `__device__` lambda, the fallback keeps its negated constraint, because nvcc otherwise reports the two overloads as ambiguous. Part of NVIDIA#24341. Authors: - Vukasin Milovanovic (https://github.com/vuule) Approvers: - Muhammad Haseeb (https://github.com/mhaseeb123) - David Wendt (https://github.com/davidwendt) URL: NVIDIA#24794
PR CI runs the upstream Polars tests that have previously completed a GPU query. A versioned run list selects those tests before sharding. Tests can still use CPU fallback for other queries, but a passing selected test must record GPU execution. The telemetry follows each test through setup and teardown and binds query callbacks to their originating test. This keeps worker-thread observations with the correct test and prevents late callbacks from updating another test. Reports include the Polars version, checkout revision, collection inventory, and completion status. Nightly CI keeps full in-memory and small-blocksize SPMD runs and adds full default-blocksize SPMD runs to discover changes in the configuration PR CI uses. Only default-blocksize SPMD has a maintained run list. The audit checks complete telemetry from its producers and shards. New GPU-executing nodes, lost GPU execution, stale entries, and incomplete default-blocksize evidence fail the audit. The other full runs remain diagnostic. Candidate lists and exact drift are published before the audit gate runs, and list updates remain reviewed commits. The Polars 1.44.2 selection ran 2,634 tests from 36,187 collected upstream items in the latest PR run. Local cudf-polars tests are unchanged. Authors: - Vyas Ramasubramani (https://github.com/vyasr) - Tom Augspurger (https://github.com/TomAugspurger) Approvers: - Tom Augspurger (https://github.com/TomAugspurger) - Matthew Roeschke (https://github.com/mroeschke) - James Lamb (https://github.com/jameslamb) URL: NVIDIA#24568
This PR adds peak-memory tracking to the remaining join benchmarks, compiles the JIT filter benchmark as C++, and skips invalid multiplicity configurations. Contributes to NVIDIA#10528; depends on NVIDIA#24730. Authors: - Yunsong Wang (https://github.com/PointKernel) Approvers: - Muhammad Haseeb (https://github.com/mhaseeb123) - Bradley Dice (https://github.com/bdice) URL: NVIDIA#24731
- Allows `Time` columns and `dt.time()` expressions to be used as sort keys. - Adds coverage for the multi-column `Time` sort used by NYSE TAQ q22 and q23 (NVIDIA#24634). - Follow-up to NVIDIA#24511 Authors: - Richard (Rick) Zamora (https://github.com/rjzamora) Approvers: - Matthew Roeschke (https://github.com/mroeschke) - Tom Augspurger (https://github.com/TomAugspurger) URL: NVIDIA#24900
- Rewrites supported `group_by_dynamic` operations using existing groupby logic. - Supports grouped and ungrouped fixed-duration tumbling windows needed by the TAQ queries. - Preserves index ordering metadata for ungrouped output. - Rejects unsupported options during translation. This is a narrower alternative to NVIDIA#24601. Broader API support and execution-time planning can be added separately. Closes NVIDIA#24734 Authors: - Richard (Rick) Zamora (https://github.com/rjzamora) Approvers: - Tom Augspurger (https://github.com/TomAugspurger) URL: NVIDIA#24701
Supports `drop_nans()` before certain grouped aggregations in cudf-polars. When the aggregation already ignores nulls (`count`, `min`, `max`, `sum`, `mean`, `std`, and `var`) we can simply replace `drop_nans` with `mask_nans` to produce the correct result. The primary target of this PR is `drop_nans().mean()` for NVIDIA#24634 Authors: - Richard (Rick) Zamora (https://github.com/rjzamora) Approvers: - Matthew Roeschke (https://github.com/mroeschke) URL: NVIDIA#24892
…VIDIA#24908) Restores the decompression backend default from before NVIDIA#24699: when `LIBCUDF_HW_DECOMPRESSION` is unset, nvCOMP picks the backend itself (`NVCOMP_DECOMPRESS_BACKEND_DEFAULT`), so the hardware decompression engine is used where supported. Unlike the pre-NVIDIA#24699 behavior, the environment variable still applies to all codecs. Authors: - Vukasin Milovanovic (https://github.com/vuule) Approvers: - Bradley Dice (https://github.com/bdice) URL: NVIDIA#24908
Second batch of string-producing ops. Authors: - https://github.com/brandon-b-miller Approvers: - Matthew Roeschke (https://github.com/mroeschke) URL: NVIDIA#24876
…e xfail/skip test list (NVIDIA#24835) Working towards supporting a new pandas version in NVIDIA#24812, a significant task is updating our cudf.pandas testing of the pandas test suite. This PR adds an agent skill dedicated for this task, mirroring our `update-narwhals-testing` skill, and applying it to the pandas test suite for 3.0.6 that we currently test Some differences: * Merged the nearly identical `test_python_pandas_cudf` and `test_cudf_pandas_pandas_tests` keys in `dependencies.yml` that specify dependencies to run the pandas test suite. Dependencies are under the same key now * The agent skill has scripts to 1) collect the pandas public API according to their docs and 2) statically analyze the pandas unit tests to check if a test uses a private API. The agent is instructed add test to the skip lists if they use a private API. This lead to +7000 test additional skipped test (skimming them, seemed reasonable). This should hopefully speed up running this job. Authors: - Matthew Roeschke (https://github.com/mroeschke) Approvers: - Matthew Murray (https://github.com/Matt711) - Kyle Edwards (https://github.com/KyleFromNVIDIA) - James Lamb (https://github.com/jameslamb) URL: NVIDIA#24835
…VIDIA#24738) Fixes device-side asserts in `UNARY_TEST` and `BINARYOP_TEST` when libcudf is built in Debug mode. The null-aware path of `cudf::detail::transform_kernel` (`cpp/src/jit/transform_kernel.cuh`) dereferenced every output `cuda::std::optional` before writing it to the output column. When a checked arithmetic operator (added in NVIDIA#23981) nullifies a row, for example on overflow or divide by zero, that optional is empty. Dereferencing it triggers the libcudacxx assertion `optional operator* called on a disengaged value` in the precompiled checked-arithmetic kernels: - `UnaryOperatorParityTest.CheckedUnaryPropagates` - `BinaryOperatorParityTest.CheckedArithmeticPropagates` Release builds aren't affected; they just write an undefined value under a null bit. This PR skips the data write for null outputs. The validity bit is still set by `warp_compact_validity` as before. Verified on a Debug build: `UNARY_TEST`, `BINARYOP_TEST`, and `TRANSFORM_TEST` all pass. Authors: - David Wendt (https://github.com/davidwendt) Approvers: - Muhammad Haseeb (https://github.com/mhaseeb123) - Bradley Dice (https://github.com/bdice) URL: NVIDIA#24738
…ode (NVIDIA#24507) Reduces the temporary memory used by `nvtext::normalize_unicode` from about 12x to about 5x the input size and improves performance, especially for long rows and columns with many rows. The output is unchanged. - **Decomposition:** an `int64` output position per input byte is replaced with a `uint8` codepoint count per byte. A count kernel produces per-block totals and the fill kernel uses a block scan to place each byte's codepoints. Codepoints without a decomposition skip the expansion buffers. - **Output:** the UTF-8 width of each codepoint is stored in unused bits of its packed slot. Row sizes are a segmented sum of the widths and a single scan writes the output bytes, replacing the thread-per-row `make_strings_children` passes. - **Reordering:** runs of combining marks are marked in a bitmask and each run is sorted by its own thread instead of a thread per row. - **Composition:** each row is split at starters that cannot compose with any earlier codepoint. Each segment is composed by its own thread instead of a thread per row. The composition table is only searched for codepoints that can be a second operand. - The normalizer tables and temporary buffers now use `cuda::device_buffer`. - The temporary memory note is updated in the C++ header and added to the pylibcudf and cudf docstrings. This also includes the fix from NVIDIA#24504: the composition table copies only the compacted entries. Output was verified byte-identical to the previous implementation using the full `UnicodeData.txt` table. The randomized input covered all four forms, sliced input, nulls, empty rows, Hangul, combining-mark runs and multi-step decompositions. On an RTX A6000, the benchmark shows a 1.5-8.1x speedup and 0.36-0.42x peak memory compared to the previous implementation, with the largest gains for long rows and many rows. Detailed results are in [this comment](NVIDIA#24507 (comment)). Authors: - David Wendt (https://github.com/davidwendt) Approvers: - Matthew Roeschke (https://github.com/mroeschke) - Yunsong Wang (https://github.com/PointKernel) - Basit Ayantunde (https://github.com/lamarrr) URL: NVIDIA#24507
…nize (NVIDIA#24468) Reworks `nvtext::wordpiece_tokenize` to reduce intermediate memory, reduce warp divergence in the tokenizer, and fix incorrect results in the `max_words_per_row` path. **Memory** - The intermediate token array (4 bytes per input byte) and its fill, segmented-reduce, and `remove_copy` passes are replaced with a per-word layout: a token count and a single token per word, with tokens for multi-token words stored in an overflow buffer. - Word starts are located exactly using a row-start bitmask and `count_if`/`copy_if` instead of an array sized for the worst case (4x the input size) merged with the row offsets. - Row token counts are computed with `cub::DeviceSegmentedReduce` over the per-word counts. **Tokenizer divergence** Tokenizing is now done in two phases: 1. A thread per word resolves most words with a single vocabulary lookup. 2. The remaining words are deferred to a kernel where a tile of 8 threads cooperates on each word. The greedy longest-prefix search checks 8 prefix sizes per step, so each step produces a token instead of one lookup per character. **`max_words_per_row` fix** The previous implementation produced incorrect results for some rows (about 0.5-1% of rows in randomized testing), including rows with a single byte which produced no tokens. This path now locates the first words of each row with a warp-per-row kernel and shares the same tokenizer. New tests verify the results against a host implementation of the tokenizer for several `max_words_per_row` values on full and sliced input. **Vocabulary lookups** Profiling showed about half of the deferred kernel time was spent in vocabulary map probes, most of which are for prefixes not in the vocabulary. - The map keys now include the hash of each entry, so non-matching slots are rejected without loading and comparing the strings. - The maps use a lower load factor (capacity of 4x the number of entries) to reduce the number of slots checked for each lookup. **Results** (RTX A6000, realistic benchmark from NVIDIA#24467) | num_rows | num_words | max_words | before | after | speedup | peak memory | |---|---|---|---|---|---|---| | 262144 | 32 | 0 | 9.13 ms | 4.53 ms | 2.02x | 242 → 136 MiB | | 2097152 | 32 | 0 | 72.12 ms | 35.32 ms | 2.04x | 1.89 → 1.06 GiB | | 32768 | 2048 | 0 | 77.14 ms | 32.11 ms | 2.40x | 1.90 → 1.06 GiB | | 262144 | 256 | 20 | 12.98 ms | 3.31 ms | 3.92x | 1.46 GiB → 129 MiB | | 32768 | 2048 | 20 | 7.61 ms | 0.63 ms | 12.08x | 1.41 GiB → 55 MiB | | 2097152 | 32 | 40 | 88.09 ms | 36.05 ms | 2.44x | 2.34 → 1.08 GiB | | 32768 | 2048 | 40 | 8.55 ms | 0.94 ms | 9.07x | 1.42 GiB → 66 MiB | `num_words` is words per row (~5.6 bytes/word: 32 ≈ 180 B, 256 ≈ 1.4 KB, 2048 ≈ 11.5 KB rows). `max_words=0` means no limit (all words are tokenized). Every benchmark configuration is faster (1.92x-12.08x) and uses less memory. Authors: - David Wendt (https://github.com/davidwendt) Approvers: - Yunsong Wang (https://github.com/PointKernel) - Basit Ayantunde (https://github.com/lamarrr) URL: NVIDIA#24468
…VIDIA#24905) - Adds support for `pl.corr(..., method="pearson")` in grouped aggregations. - Supports pairwise null handling and mixed numeric input types. - Falls back to single-partition execution in the streaming engine. This enables the correlation queries in the NYSE TAQ benchmark (NVIDIA#24634). Streaming and distributed decomposition is left for future work. Authors: - Richard (Rick) Zamora (https://github.com/rjzamora) Approvers: - Tom Augspurger (https://github.com/TomAugspurger) URL: NVIDIA#24905
closes NVIDIA#24555 Adds a `--query-mode` to specify if benchmarks should use the LazyFrame API defined in `pdsh.py` or SQL files using `pl.SQLContext` along with a `--sql-directory` to specify a directory of SQL files named like `q<N>.sql` where `<N>` is the query number (I chose a directory specification to make it easier to generically benchmark any SQL files) SQL benchmark runtimes will exclude the work done by `pl.SqlContext.execute(...)` (Polars SQL translation) Authors: - Matthew Roeschke (https://github.com/mroeschke) Approvers: - Matthew Murray (https://github.com/Matt711) URL: NVIDIA#24746
…ice_buffer<T>` (NVIDIA#24925) Use CCCL buffers instead of RMM vectors where possible. Issue: rapidsai/build-planning#321 Authors: - Kyle Edwards (https://github.com/KyleFromNVIDIA) Approvers: - Yunsong Wang (https://github.com/PointKernel) - Bradley Dice (https://github.com/bdice) URL: NVIDIA#24925
NVIDIA#24704) (Agent implemented) In cudf-polars, we want to implement the `polars.Categorical` type using libcudf's dictionary type. Polars represents their categorical indices as unsigned ints (the Arrow spec merely [suggests](https://github.com/apache/arrow/blob/61f24ce27b0597258938566665d353ce54bccf60/format/Schema.fbs#L503) to use signed ints like libcudf does) This PR implements casting logic to convert dictionary (host or device) data using unsigned ints to signed ints for libcudf. cc @davidwendt for your recent improvements to the dictionary type Authors: - Matthew Roeschke (https://github.com/mroeschke) Approvers: - Vukasin Milovanovic (https://github.com/vuule) - Muhammad Haseeb (https://github.com/mhaseeb123) URL: NVIDIA#24704
…vector<T>` to `cuda::device_buffer<T>` (NVIDIA#24921) Use CCCL buffers instead of RMM vectors where possible. Issue: rapidsai/build-planning#321 Authors: - Kyle Edwards (https://github.com/KyleFromNVIDIA) Approvers: - Bradley Dice (https://github.com/bdice) - David Wendt (https://github.com/davidwendt) URL: NVIDIA#24921
The current `PANDAS_CURRENT_SUPPORTED_VERSION` is fairly stale as 3.0.6 is the latest patch version. I don't think we have a good way of keeping "current" here up-to-date. It's used only in 2 places so attempting to remove for now. This will make future pandas upgrades easier xref NVIDIA#24812 Authors: - Matthew Roeschke (https://github.com/mroeschke) Approvers: - Matthew Murray (https://github.com/Matt711) - Tom Augspurger (https://github.com/TomAugspurger) URL: NVIDIA#24813
…A#24596) Follow up NVIDIA#24567 This PR fixes ORC writer `sum` statistics for d32/64 cols by accumulating them in `__int128_t`. Previously they were accumulated in `int64_t` but encoded as `__int128_t` so e.g. `-5` was written as `18446744073709551611`. We also now detect sum overflow exactly (including d128) by reusing the `sum_overflow_result` and `overflow_sum_op` helpers instead of the conservative `max / n` guard, and invalidate merged sums when any input chunk has no valid sum, making the ORC decimal statistics encoder honor `has_sum`. Authors: - Muhammad Haseeb (https://github.com/mhaseeb123) Approvers: - Vukasin Milovanovic (https://github.com/vuule) - Matt Gara (https://github.com/mattgara) URL: NVIDIA#24596
…IDIA#24346) ## Description Expand cudf-query-engine-builder's evaluation dataset from two tasks to eight to address the coverage concern in [NVIDIA/skills#620](NVIDIA/skills#620). The added tasks exercise asynchronous ownership, worker readiness, engine-buffer inputs, conflicting query/reference semantics and native-result acceptance. Clarify exception cleanup and completion requirements in the skill and supporting reference. Add two executable C++ CPU fixtures and an independent grader that verifies protected test files before compiling submitted repairs. Standard evaluator scores and thresholds remain unchanged. The CI dataset uses CPU simulations and source review. These tests do not establish native CUDA execution or production readiness. Native GPU comparisons are separate supplementary checks. NVSkills CI must regenerate BENCHMARK.md, skill-card.md and skill.oms.sig for this change before merge. Local validation passed schema/configuration checks, correct and deliberately defective fixture controls, 52 grader integrity/output-path regression controls, and git diff --check. Sixty local trials compared runs with and without the skill on eight CPU/source tasks and two native tasks across three model configurations. The evaluations retained failures, including two skill-assisted native owning-result implementations that failed allocation cleanup; the local results do not establish an NVSkills CI PASS. The full local SkillSpector scan completed at CAUTION 29 with four medium and three low findings, none suppressed. Remaining findings concern tool scope, intentional compilation/execution of submitted native code and bounded image-build cleanup. Native-code execution requires runner isolation; the scripts' environment and resource limits are not a sandbox. Pipeline-generated reports and signatures are unchanged. ## Checklist - [ ] I am familiar with the [Contributing Guidelines](https://github.com/NVIDIA/cudf/blob/HEAD/CONTRIBUTING.md). - [ ] New or existing tests cover these changes. - [ ] The documentation is up to date with these changes. --------- Signed-off-by: Joe Sack <jsack@nvidia.com> Signed-off-by: nvskills-svc-account <svc-nvskills-signing@nvidia.com> Co-authored-by: nvskills-svc-account <svc-nvskills-signing@nvidia.com>
…4659) Trivial PR: adds missing public constructors and sparse payload page materialization APIs to the `hybrid_scan_reader`. The public interfaces simply call the corresponding `impl` backends. Details: * Constructors added: Accept a pre-existing footer metadata `FileMetaData&&`. * APIs added: `payload_pages_byte_ranges` and `setup_chunking_for_payload_pages` * Rename multifile reader's `setup_chunking_for_payload_columns` API to `setup_chunking_for_payload_pages` (breaking but experimental namespace and only one user - notified) * Python bindings for all these * Gtests and Pytests for these Authors: - Muhammad Haseeb (https://github.com/mhaseeb123) Approvers: - Vukasin Milovanovic (https://github.com/vuule) - Matt Gara (https://github.com/mattgara) - Bradley Dice (https://github.com/bdice) - Matthew Murray (https://github.com/Matt711) - GALI PREM SAGAR (https://github.com/galipremsagar) URL: NVIDIA#24659
To hopefully capture cases like NVIDIA#23789 via CodeRabbit (Note: I think we should start breaking this up into library specific guidelines) Authors: - Matthew Roeschke (https://github.com/mroeschke) Approvers: - Matthew Murray (https://github.com/Matt711) URL: NVIDIA#24932
…benchmarking (NVIDIA#24929) Currently running our cudf-polars benchmarks with a remote file path an `--iterations 1` would always raise an error due to `--io-mode` * The `cold` and `lukewarm` mode would eventually hit `drop_file_page_cache_recursively` which raises an error for remote data * `hot` mode raises because it requires `iterations>1` (pre-existing) So for the `cold` and `lukewarm` cases instead of raising a `ValueError`, it just no-ops instead Authors: - Matthew Roeschke (https://github.com/mroeschke) Approvers: - Tom Augspurger (https://github.com/TomAugspurger) URL: NVIDIA#24929
Now that the `sccache-dist` issues with the `libc` crate have been resolved, enable `sccache-dist`. Authors: - Kyle Edwards (https://github.com/KyleFromNVIDIA) Approvers: - James Lamb (https://github.com/jameslamb) - Tom Augspurger (https://github.com/TomAugspurger) URL: NVIDIA#25068
…og (NVIDIA#25065) `test_scan_delta_dv_single` asserts on the verbose log Polars writes while it reads a Delta table. cudf-polars reads the table since NVIDIA#24345, so the log is never written and the nightly job fails. Deselect it with the reason. closes NVIDIA#25063 ## Verification | Check | Result | |---|---| | `run_cudf_polars_polars_tests.sh --run-slow --engine in-memory -k test_delta_deletion_vector`, before | 6 failed (the CI assertion) | | Same, after | 42 passed, 5 skipped, 1 xfailed | | Same with `--engine spmd`, after | 42 passed, 5 skipped, 1 xfailed | Authors: - Matthew Murray (https://github.com/Matt711) Approvers: - Tom Augspurger (https://github.com/TomAugspurger) URL: NVIDIA#25065
…DIA#25003) Several places in libcudf called `release()` on a `std::unique_ptr` and then moved from the raw pointer it returned (`std::move(*ptr.release())`, `ptr.release()[0]`, or `auto p = contents.data.release();`). The device memory was moved out and freed normally, but nothing deleted the small host-side object (`rmm::device_buffer`, `cuda::device_buffer<bitmask_type>`, or `cudf::column`), so each call leaked it. This PR moves from `*ptr` instead, so the `unique_ptr` keeps ownership and deletes the moved-from object. It fixes all the sites listed in NVIDIA#24926, plus the same pattern found in: - `cpp/src/groupby/sort/group_argmin.cu`, `group_argmax.cu` - `cpp/src/io/utilities/column_buffer.cpp` - `java/src/main/native/src/ColumnViewJni.cu` - `cpp/benchmarks/common/generate_input.cu` (5 sites), `cpp/benchmarks/io/text/multibyte_split.cpp` It also adds a one-line note about this pattern to the Resource Management section of `cpp/REVIEW_GUIDELINES.md`, so future reviews catch it. The remaining `unique_ptr::release()` calls in the tree hand ownership off on purpose: the JNI `ContiguousTableJni.cpp`, DLPack/Arrow interop, and custom deleters. Closes NVIDIA#24926 Authors: - David Wendt (https://github.com/davidwendt) Approvers: - Bradley Dice (https://github.com/bdice) - Muhammad Haseeb (https://github.com/mhaseeb123) - Alessandro Bellina (https://github.com/abellina) URL: NVIDIA#25003
Use uv for wheel build isolation (`python -m build --installer uv`), wheel tests, upstream pandas and Polars tests, and Java's Python build tools. Use the image-provided uv and shared `rapids-init-uv` / `rapids-uv-pip` helpers to preserve proxy/nightly routing, generated constraints, interpreter selection, and a shared job-local package cache. Conda workflows retain their existing installers. Earlier paired CI runs observed: - **13.2% less summed worker time** for staged native wheel-build steps. - **Approximately 60–87% shorter dependency-installation phases** in upstream pandas and Polars jobs. These observations compare earlier measured commits. [Full benchmark data and limitations](NVIDIA#25011 (comment)). Authors: - Vyas Ramasubramani (https://github.com/vyasr) Approvers: - Bradley Dice (https://github.com/bdice) - https://github.com/paul-aiyedun - Nghia Truong (https://github.com/ttnghia) URL: NVIDIA#25011
…O-IR (NVIDIA#24476) Adds `cudf::ast::jit::call`, a JIT-only AST node that calls a consumer-supplied device function precompiled to LTO-IR. Row IR lowers the node to a single call inside the generated row function, so the function runs in the same kernel as the surrounding operators and no intermediate column is materialized. A kernel that contains calls is compiled to LTO-IR and linked by nvJitLink with the modules the caller passes. Consumers have functions with no cuDF operator, such as Presto's `date_trunc` in Velox, which today run as separate kernels or chains of libcudf calls. In a Velox prototype on cuDF 26.08 (GB10, 10M rows), fusing `date_trunc` with this node was 4.8–10.6x faster than Velox's current chain of libcudf calls; LTO-IR callees ran at the same speed as CUDA-source ones. - **API:** `call(tree, udf{symbol, is_pure, output_type, input_types, is_null_aware}, args, error_policy)` in `cudf/ast/jit/udf.hpp`. The new `cudf/udf.hpp` defines `lto_fragment{type, data}` and `lto_fragments`; `lto_binary_type` moves there from `cudf/transform.hpp`, which includes it. - **Callee signature:** `extern "C" __device__ cudf::errc symbol(R* out, A... in)`, with arguments passed by value. With `is_null_aware = YES`, `R` and each `A` are `cuda::std::optional`. STRING inputs and outputs are `cudf::string_view`; a returned view must stay valid until the output column is built. Calls of one symbol in a kernel must declare the same output type, input types and `is_null_aware` (decimal scales may differ), or `std::invalid_argument` is thrown. - **Nulls and errors:** for a callee that is not null-aware, a row with a null argument skips the call and produces null. `PROPAGATE` and `NULLIFY` behave as for other JIT operators. `cudf::errc` is now backed by `int32_t` (`cudf::error_int`); codes from `cudf::errc::USER_RESERVED` (0x10000) up are reserved for user-defined functions; libcudf gives them no meaning. - **CSE:** callees are assumed pure (`is_pure` defaults to `true`), so equal calls may be merged; the caller sets `is_pure = false` for a function that is not deterministic. A call is evaluated at most once per row however many expressions refer to it, and separate calls of an impure function are never merged. - **Entry points:** `compute_column_jit`, `compute_table_jit`, `transform_program` and the AST overload of `filter_join_indices_jit` take the modules as `lto_fragments fragments = {}`, just before `stream`, so callers that pass `stream` positionally now pass the fragments first (`{}` for none). A module may define any number of functions and reference functions in other modules; its bytes must stay valid until the evaluation finishes or the program is constructed. The AST interpreter, `compute_column`, rejects calls. - **Codegen:** the generated row function keeps its name and kernels call it by that name (`transform_args::udf_expression`), instead of the source parser renaming it, because callee declarations now precede it. Authors: - Pramod Satya (https://github.com/pramodsatya) - Basit Ayantunde (https://github.com/lamarrr) Approvers: - Basit Ayantunde (https://github.com/lamarrr) - Muhammad Haseeb (https://github.com/mhaseeb123) - Matthew Roeschke (https://github.com/mroeschke) - Nghia Truong (https://github.com/ttnghia) URL: NVIDIA#24476
…n the current device resource (NVIDIA#24953) Part of NVIDIA#24268. Adds a trailing `cudf::memory_resources mr` parameter to internal and detail functions that allocated on `cudf::get_current_device_resource_ref()` without taking an MR. Inside them, scratch allocations use `mr.get_temporary_mr()` and returned allocations use `mr.get_output_mr()`. Behavior is unchanged. Callers that take an `rmm::device_async_resource_ref`, or no MR, pass `cudf::get_current_device_resource_ref()`. Callers that already take `cudf::memory_resources` pass `mr.as_temporary()`. Functions already covered by open PRs (NVIDIA#24084, NVIDIA#24087, NVIDIA#24088) are not included. Neither are functions with no caller that holds an MR (mostly I/O writer internals). ### Review-required Most of the diff is mechanical. These are the places that need judgment: - **Returned allocations not previously flagged:** `scatter_to_gather`, `scatter_to_gather_complement` and `nulls_per_group` now allocate their returned uvector from `mr.get_output_mr()`. It previously came from rmm's implicit default resource, which is the same resource for all current callers. - **Single local resource split in two:** `process_token_stream` (nested_json_gpu.cu) and `generate_merged_centroids` (tdigest_aggregation.cu) each used one local resource variable for both scratch and returned buffers. It is now split into temporary and output resources. - **Output vs. temporary:** each helper that returns buffers its callers use only as scratch puts them on `mr.get_output_mr()`. The callers pass the current resource for now. This covers avro `decompress_data`, JSON `reduce_to_column_tree`, `get_values_column_indices`, `is_all_nulls_each_column`, `hash_node_type_with_field_name` and `build_command_buffer`. - **Unused parameter:** the column/column `compiled::binary_operation` overload takes an `mr` it does not use. All three overloads must match because a single templated call site in binaryop.cpp calls them. - **`groupby::detail::compute_row_bitmask`:** a local `mr` alias of the current resource is replaced by the parameter, which is forwarded to its callees. - **Lambda capture:** `copy_strings_to_host_sync` now captures `mr`. - **Not migrated:** `sort_groupby_helper::key_sort_order` and `unsorted_keys_labels` are left alone. They cache their results in members, so they need a constructor-level resource rather than a per-call `mr`. Authors: - Bradley Dice (https://github.com/bdice) Approvers: - David Wendt (https://github.com/davidwendt) - Muhammad Haseeb (https://github.com/mhaseeb123) - Vukasin Milovanovic (https://github.com/vuule) - Nghia Truong (https://github.com/ttnghia) URL: NVIDIA#24953
Add the pure-MLIR cooperative reduction backend used by the JIT GroupBy.apply engine, as a standalone building block (no GroupBy.apply wiring yet):
- cooperative_types: CooperativeArrayType {meminfo, data, size, index} + model.
- cooperative_typing: reduction attributes (min/max/sum/mean/var/std/idxmin/ idxmax/corr), unsupported-dtype columns modelled as Poison, and rejection of operators applied directly to a grouped column.
- cooperative_device_funcs: per-dtype MLIR reduction modules (shared-memory tree reductions) with NaN skipping (pandas skipna) for min/max/sum, int->int64 sum promotion, float32 mean/var/std, ddof=1 var/std, Pearson corr (int-only), and (value,index) idxmin/idxmax with smallest-index tie-break.
- cooperative_lowering: thin lowerings that call the outlined device functions, plus the linker hook that adds the needed PTX at complete().
Self-contained device tests (test_cooperative_reductions.py) launch one cooperative block per reduction and verify the scalar result, covering all four numeric dtypes, NaN skipping, idxmin/idxmax, corr (incl. the float->UDFError case), and operator rejection. compute-sanitizer memcheck clean.
Authors:
- https://github.com/brandon-b-miller
Approvers:
- Matthew Roeschke (https://github.com/mroeschke)
- Vyas Ramasubramani (https://github.com/vyasr)
URL: NVIDIA#25018
… mismatch (NVIDIA#25078) Authors: - Tom Augspurger (https://github.com/TomAugspurger) Approvers: - Matthew Murray (https://github.com/Matt711) URL: NVIDIA#25078
…IDIA#25054) The single-range `cudf::detail::slice(column_view, begin, end, stream)` computes the slice's null count over null-mask bits `[begin, end)`, ignoring `input.offset()`. When the input is itself a slice, the result's `null_count()` therefore described the wrong rows. The new view's offset (`input.offset() + begin`) and null mask were already correct, so only the stored null count was wrong. The multi-range `slice` overload in the same file already shifts by the input's offset before counting. This is reachable through public APIs. For example, copying a STRUCT column view into a `cudf::column` slices each child to the parent's rows (`column.cu`, `create_column_from_view` for `struct_view`). If a child view already has an offset (such as a STRUCT view built over sliced columns), the copied child gets the wrong null count. It was found through cudf-spark-jni, where a STRUCT view with offset 0 over sliced children came back from a host round trip with wrong child null counts. The fix counts nulls over `[input.offset() + begin, input.offset() + end)`. ### Tests New `StructsColumnCopyTest.ChildWithOffset` in `COLUMN_TEST`: an `INT32` column with nulls at rows 7, 9, 11, 13 and 15 is sliced to rows `[3, 17)`, so it has 5 nulls. It is wrapped as the only child of an offset-0 STRUCT view, which is then copied. Before the fix, the copied child reports 4 nulls, the count for rows `[0, 14)`. With the fix it reports 5 and matches the sliced input. `COLUMN_TEST` and `COPYING_TEST` pass. Authors: - MithunR (https://github.com/mythrocks) Approvers: - David Wendt (https://github.com/davidwendt) - Nghia Truong (https://github.com/ttnghia) URL: NVIDIA#25054
This PR adds a check if there are any definition and repetition levels to process in the subpass, before launching the preprocess_levels_kernel. If there are no definition and repetition levels present, then there is no work to do for the kernel, and we can avoid the launch entirely. Authors: - https://github.com/y2kiran Approvers: - Muhammad Haseeb (https://github.com/mhaseeb123) - Vukasin Milovanovic (https://github.com/vuule) URL: NVIDIA#25087
…olumn.is_sorted` (NVIDIA#25083) closes NVIDIA#21914 Authors: - Matthew Roeschke (https://github.com/mroeschke) Approvers: - Richard (Rick) Zamora (https://github.com/rjzamora) URL: NVIDIA#25083
…ond range (NVIDIA#24541) Contributes to NVIDIA#23422. ### First issue ORC timestamp statistics were left on the input instants while NVIDIA#23449 re-based the data stream onto the writer's wall clock, so in a non-UTC file the two disagreed by the timezone offset. This PR makes the statistics of a non-UTC file agree with its data stream. Each value is moved onto the writer's wall clock as it is reduced, per value rather than on the reduced bounds, because daylight saving makes the offset vary between the values of one column. The pre-ORC-135 `minimum`/`maximum` pair is no longer written in a non-UTC file, same as the Apache's writer. ### Second issue Overflow in ORC timestamp statistics for timestamps outside ~1677-2262. Statistics are now reduced as a 128-bit nanosecond count. Widening it also lets the shift happen in the statistics kernel, rather than in copies of the timestamp columns, so no scratch buffer is allocated per column. The shift costs a timezone lookup per timestamp value, which is measurable only when writing a non-UTC file; UTC is unaffected. Authors: - Vukasin Milovanovic (https://github.com/vuule) Approvers: - Muhammad Haseeb (https://github.com/mhaseeb123) - Bradley Dice (https://github.com/bdice) - MithunR (https://github.com/mythrocks) URL: NVIDIA#24541
…VIDIA#24964) Moves `cudf/utilities/execution_env.hpp` (a detail-only utility) to `cudf/detail/utilities/` and adds `cudf::detail::exec_policy_nosync(stream, mr)`, the Thrust counterpart to `make_env`. Both take `cudf::memory_resources` and use only the temporary resource, so callers don't need to extract memory resources themselves. Existing `make_env` call sites only have their include path updated. Adopting `exec_policy_nosync` at other call sites is left for follow-ups. Follow-up to NVIDIA#24489. Precursor to NVIDIA#21272. Authors: - Bradley Dice (https://github.com/bdice) Approvers: - Muhammad Haseeb (https://github.com/mhaseeb123) - Lawrence Mitchell (https://github.com/wence-) - Gera Shegalov (https://github.com/gerashegalov) - Yunsong Wang (https://github.com/PointKernel) URL: NVIDIA#24964
…IDIA#25064) Polars treats NaN as greater than every other value and equal to itself. libcudf follows IEEE 754, so float comparisons (`==`, `!=`, `<`, `<=`, `>`, `>=`, `eq_missing`, `ne_missing`) give the wrong answer for NaN rows. This is the NaN-ordering gap tracked in NVIDIA#25062 and, in libcudf, NVIDIA#18930. It made the nightly cudf-polars Polars-tests job fail on upstream Polars Iceberg tests. This PR does not change engine behavior. It marks the tests that depend on Polars' NaN ordering as expected failures, so nightly stops failing until the comparisons are fixed: - `inject_gpu_engine.py`: adds 17 `EXPECTED_FAILURES` entries for upstream Polars `test_iceberg.py` and `test_iceberg_plugin_suite.py` tests that compare or filter on NaN. - `test_numeric_binops.py`: adds tests that document the NaN behavior of float comparisons, a Parquet scan filter and a conditional join. The cases where the GPU result differs from Polars are marked `xfail` with `NAN_ORDERING_XFAIL`. Both carry a TODO to remove the xfails once libcudf's comparison operators order NaN the way Polars does (NVIDIA#18930). A workaround in cudf-polars that gives the correct result without waiting for libcudf is in NVIDIA#25102. Whichever of the two PRs lands second removes the xfails. Relates to NVIDIA#25062. Authors: - Matthew Murray (https://github.com/Matt711) - Tom Augspurger (https://github.com/TomAugspurger) Approvers: - Tom Augspurger (https://github.com/TomAugspurger) URL: NVIDIA#25064
Part of NVIDIA#24268 (factories cycle work item). Migrates `column_factories.{cpp,cu}`, `strings_column_factories.{cu,cuh}`, `strings_children.cuh`, `dictionary_factories.cu`, `cast_ops.cu`, `binaryop.cpp`, and `compiled/binary_ops.cu` to `cudf::memory_resources` in one PR, since these files call each other. Depends on NVIDIA#24536 (`fill_in_place` memory resources); once it lands, `make_column_from_scalar` passes its resources to `fill_in_place`, removing the TODO added there. Notes: - Also includes `transform/checked_arithmetic.cu`, which is part of the same cycle: `binaryop.cpp` calls it, and it calls `make_fixed_width_column` / `make_column_from_scalar`. - `make_lists_column_from_scalar` (migrated in NVIDIA#24535) receives `mr`. `make_structs_column` (not migrated yet) receives `mr.get_output_mr()`. - Fixes fixed-point to fixed-point `cast` with equal scales, which allocated the returned data from the current device resource instead of `mr`. - Temporary scalars passed to the scalar factories are allocated from `temp_mr`. This PR also includes NVIDIA#24084 and NVIDIA#24087, combined here to reduce merge conflicts. Both were approved by two C++ reviewers. Authors: - Bradley Dice (https://github.com/bdice) - Niranda Perera (https://github.com/nirandaperera) Approvers: - Muhammad Haseeb (https://github.com/mhaseeb123) - Niranda Perera (https://github.com/nirandaperera) - Nghia Truong (https://github.com/ttnghia) URL: NVIDIA#24531
…A#25070) A multi-rank `Union` now keeps the rows of every child on every rank when all children are duplicated. Duplicated children are still dropped on ranks other than 0 when the output is not duplicated, so mixed unions, where some children are distributed, behave as before. Before this change, a `Union` whose children were all duplicated (every rank holds a full copy, e.g. small aggregates) marked its output `duplicated=True` but still dropped the children on every rank other than 0. Those ranks sent empty chunks while claiming a full copy, so a join that broadcast that side skipped the gather and joined against nothing on ranks other than 0. On two ranks this silently dropped about half the rows of the join. Authors: - Mads R. B. Kristensen (https://github.com/madsbk) Approvers: - Vyas Ramasubramani (https://github.com/vyasr) URL: NVIDIA#25070
…5056) Moves sparse row-slice exchange out of multi-rank rolling into `SparseExchange`. It supports copying row slices, moving whole chunks, and storing received data in spillable memory. Rolling keeps its existing overlap planning and local execution. ## Next steps - Use `SparseExchange` for inter-rank movement in cursor-based ordered groupby. - Use it for inter-rank movement in cursor-based ordered join. - Remove `adjust_ordering` and its tests once those paths replace its remaining uses. ### Further Background Ordered groupby and join currently rely on the existence of global/flat `Ordering` boundaries. More specifically, we use `adjust_ordering` to "strictify" or align boundaries, and then perform chunkwise logic. This approach is neither robust nor is it performant. There are many cases where we don't have full boundary information, and therefore it is much more robust for us to take advantage of the required ordering in the channel by using a cursor-based approach (like we do for Rolling). This PR is the first step in moving toward that more-general/robust approach. If/when boundaries are available, we can still use them to avoid or optimize inter-rank data movement. Authors: - Richard (Rick) Zamora (https://github.com/rjzamora) Approvers: - Matthew Roeschke (https://github.com/mroeschke) URL: NVIDIA#25056
Discovered while working on the issues for NVIDIA#21910. When the key columns are not sorted and `maintain_order=False` from an expression, the hash-based `plc.stream_compaction.distinct` can change the row order, so the flag was wrong Authors: - Matthew Roeschke (https://github.com/mroeschke) Approvers: - Mads R. B. Kristensen (https://github.com/madsbk) - Tom Augspurger (https://github.com/TomAugspurger) URL: NVIDIA#25095
Runs the Narwhals tests in nightly CI only. The `narwhals-tests` job stays in `test.yaml`, so it still runs nightly on every branch. This removes it from `pr.yaml`: - the `narwhals-tests` job, - its entry in `pr-builder` `needs`, - the `test_python_narwhals` changed-files group, which only that job used. Narwhals tests are slow and break whenever Narwhals or Polars releases, unrelated to the PR being tested. The plugin files in `ci/` and `ci/test_narwhals.sh` are unchanged. This is a follow-up to the Narwhals upgrade in NVIDIA#24862. The two PRs are independent and can merge in either order. ## Verification | Check | Result | | --- | --- | | `pr.yaml` parses as YAML | yes | | No `narwhals` references left in `pr.yaml` | yes | | `narwhals-tests` job still in `test.yaml` | yes | Authors: - Matthew Murray (https://github.com/Matt711) Approvers: - Vyas Ramasubramani (https://github.com/vyasr) URL: NVIDIA#25101
We're seeing errors in CI related to compiling wheels in our CUDA 13 jobs. This is possibly related to the recent cuda-bindings==13.5 release. Apparently, there's a requirement that the version of `cuda-bindings` used for building projects match the version of CUDA we get from `ci-imgs`. For now, ci-imgs is on CUDA 13.4.x, so we'll pin our cuda-bindings to `<13.5.0`. Previously, pylibcudf didn't declare `cuda-bindings` as a build dependency. I think this was incorrect, since it `cimports` `cuda.bindings.cydriver` in a Cython file. I've updated our dependencies to declare that as a build dependency of pylibcudf. I've *not* added a *runtime* upper-bound on cuda-bindings. Authors: - Tom Augspurger (https://github.com/TomAugspurger) Approvers: - Robert Maynard (https://github.com/robertmaynard) - Bradley Dice (https://github.com/bdice) URL: NVIDIA#25107
Add TornadoVM to notable projects that use cuDF
…rs (NVIDIA#25079) closes NVIDIA#21911 Authors: - Matthew Roeschke (https://github.com/mroeschke) Approvers: - Tom Augspurger (https://github.com/TomAugspurger) URL: NVIDIA#25079
TPC-DS queries 47 and 57 rank rows with `over(..., order_by=["d_year", "d_moy"])`. Polars 2.1 plans a list of order keys with a `row_encode` expression that its plan visitor does not expose, so cudf-polars cannot run the query and it falls back to the CPU (NVIDIA#24936). With `raise_on_fail` on, as in the benchmarks, the query fails instead. This changes the two queries to pass one `pl.struct("d_year", "d_moy")` key. The ordering is the same, a struct sorts by its first field and then its second, and Polars plans no `row_encode`, so both queries run on the GPU on today's Polars. This is a workaround for the two benchmark queries only. A user who writes `over(order_by=[a, b])` still falls back until Polars exposes `row_encode` and NVIDIA#24943 lands. When that is released, this change is harmless and can stay or be reverted. This also removes the q47 and q57 entries from `EXPECTED_FAILURES_TPCDS`, which NVIDIA#24935 added. Related: NVIDIA#24936 (the issue), NVIDIA#24943 (the general fix, waiting on Polars 2.2). ## Verification GPU 0, Polars 2.1.0, `--raise-on-fail`, results checked against DuckDB. CI runs SF1 with the default partition size, like the first rows. | Check | Result | | --- | --- | | q47 and q57, decimal and float data, SF1, 2 iterations | pass | | q47 and q57 with the old `order_by` list | fail with `NotImplementedError: row_encode` | | q57 at SF10 | pass | | q47 at SF10, or q47 and q57 with a small `--target-partition-size` | fail, see below | | Full TPC-DS suite, decimal data, SF1, 2 iterations | 97 passed, 2 expected failures (q04, q11) | With more than one partition, the final sort of q47 is by an expression, which the multi-partition sort does not support. The engine runs it on one partition on the GPU, warns, and the test treats the warning as a failure. This is not a CPU fallback and is not caused by this change. q47 and q57 never reached it before because they failed earlier on `row_encode`. Authors: - Matthew Murray (https://github.com/Matt711) Approvers: - Matthew Roeschke (https://github.com/mroeschke) URL: NVIDIA#25053
Updates the Narwhals version used in CI from 2.24.0 to 2.28.0 (the latest release) and updates the test plugins, following the `update-narwhals-testing` skill. No cudf or cudf-polars code changes. ## Changes - `dependencies.yaml`: `narwhals==2.28.0`. - `ci/narwhals_cudf_test_plugin.py`, `ci/narwhals_cudf_polars_test_plugin.py`, `ci/narwhals_cudf_pandas_test_plugin.py`: - Removed entries whose tests now pass (dataframe interchange protocol, `CategoricalDtype` hashing, `test_concat_vertical`, and others). - Added tests that Narwhals marks `xfail(strict)` but that pass here to `TESTS_TO_SKIP`. - Added the other new failures to `EXPECTED_FAILURES`, with the observed error as the reason. | Plugin | Entries removed | Added to `TESTS_TO_SKIP` | Added to `EXPECTED_FAILURES` | | --- | --- | --- | --- | | cudf | 20 | 15 | 68 | | cudf-polars | 7 | 4 | 2 | | cudf.pandas | 6 | 6 | 18 | ## New failures worth following up Most of the new cudf failures are tests that did not fail at 2.24.0 (`is_close`, `is_finite` with nulls, floor division by zero, `std`/`var` over a window with `ddof`, rename with duplicate names). Two are cudf-polars bugs: - `test_expr_is_in_with_only_null_in_other`: `RuntimeError: CUDF failure ... Invalid type_id.` (NVIDIA#24863) - `test_fill_null_scalar_receiver`: `RuntimeError: CUDF failure ... Column size mismatch` (NVIDIA#24864) ## Testing Ran the three variants from `ci/test_narwhals.sh` against Narwhals v2.28.0 on a GPU: 0 failures in each of cudf, cudf-polars and cudf.pandas. As a baseline, the same commands against v2.24.0 had no failures other than `tests/version_test.py::test_package_version`, which only fails because Narwhals was used from a source checkout. - cudf and cudf.pandas ran in a development environment built from the `main` checkout. The cudf-polars run used the 26.12.0a530 nightly (polars 2.0.0), not this exact commit. - `test_expr_is_in_with_null_in_other[pandas]` fails in isolation but passed in one of two parallel runs, so it is in `TESTS_TO_SKIP` and not `EXPECTED_FAILURES` (a strict xfail would fail when it passes). Authors: - Matthew Murray (https://github.com/Matt711) Approvers: - Matthew Roeschke (https://github.com/mroeschke) - Marcus D. Hanwell (https://github.com/cryos) URL: NVIDIA#24862
…es (NVIDIA#25106) - `test_cast_errors[values4]`: the string-cast message now matches polars on the in-memory engine (NVIDIA#24959), so only xfail it for the streaming engine, where the error is still wrapped in an ExceptionGroup. - `test_sink_iceberg_overwrite_filter_polars_semantics`: NVIDIA#25021 is the root cause. On main, they're all xfailed but some are passing with the spmd and spmd-small engines, probably since they see different subsets of the data. Moving them all to a non-strict xfail. Authors: - Tom Augspurger (https://github.com/TomAugspurger) Approvers: - Vyas Ramasubramani (https://github.com/vyasr) URL: NVIDIA#25106
…t tests (NVIDIA#25111) `make_status_buffer` in `VARIANT_EXTRACT_TEST` zeroed the status column with a synchronous `cudaMemset` on the legacy default stream. Under the CI stream preload, `cudf::get_default_stream()` is a non-blocking stream, so nothing ordered the memset before the `cast_variant` kernel. The kernel could then read stale bytes from a reused pool allocation as incoming non-success status. This caused an intermittent failure in `CastVariantStatusTest.SqlNullInputProducesRowNullStatus` ([log](https://github.com/NVIDIA/cudf/actions/runs/42062854048/job/127381022050?pr=25107#step:12:3597)). This PR moves the memset to `cudaMemsetAsync` on `cudf::test::get_default_stream()`. Authors: - Bradley Dice (https://github.com/bdice) Approvers: - Yunsong Wang (https://github.com/PointKernel) - Vyas Ramasubramani (https://github.com/vyasr) URL: NVIDIA#25111
…VIDIA#24967) Follow up to NVIDIA#24932, IMO the current `python/REVIEW_GUIDELINES.md` is very general and has minimal specifics for pylibcudf/cudf_polars. Instead, this PR creates 2 dedicated review guidelines for pylibcudf and cudf_polars. I seeded these with an agent looking back over ~8 months of PRs identifying significant fixes, conventions, consistent review feedback that are not enforced by a lint check, CI check, API constraint, etc. and did some manual pruning. Authors: - Matthew Roeschke (https://github.com/mroeschke) Approvers: - James Lamb (https://github.com/jameslamb) - Vyas Ramasubramani (https://github.com/vyasr) - Tom Augspurger (https://github.com/TomAugspurger) URL: NVIDIA#24967
…IDIA#24766) # Motivation JIT transforms currently use separate column wrappers and access rules for operations already supported by column device views. Maintaining these parallel implementations duplicates logic and makes it easier for JIT and non-JIT behavior to diverge as column types evolve. Moving element accessors from JIT into the column view types gives both paths a shared implementation and reduces the work needed to maintain consistent behavior. `repr_element` makes access to the stored representation explicit and allows cuDF to have a common representation across JIT and non-JIT code. They let callers express the access they need through the same column view interface, without relying on JIT-specific wrappers. The JIT transform kernel and column element accessors are also now used in pre-compiled kernels NVIDIA#23981, making it imperative to factor this out of JIT-only machinery. This pull request: - Merges the custom mutable-string and raw-storage device-view wrappers' `element` operations into the core column device views. - Adds `repr_element` for treating types by their byte representations. This is necessary for `string_view` array output that is not Arrow-encoded. Authors: - Basit Ayantunde (https://github.com/lamarrr) Approvers: - Bradley Dice (https://github.com/bdice) - Muhammad Haseeb (https://github.com/mhaseeb123) URL: NVIDIA#24766
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
Description
JIT string outputs currently use a generic character span while other string operations use string views. Introduce cudf::mutable_string_view and use it for preallocated string outputs, giving JIT and non-JIT callers a common string representation through column device views. Update reflection, precompiled kernel signatures, examples, and existing string-output tests.
Stacked on the accessor consolidation in NVIDIA/cudf#24766; this PR targets transform-storage-access.
Validation: before rebasing onto the latest accessor branch, each stack layer built in the CUDA 13.3 devcontainer with sccache on an RTX A6000 (sm_86). COLUMN_TEST, TRANSFORM_TEST, TRANSFORM_LTO_TEST, and JOIN_TEST passed: 1,527 tests passed and two join tests skipped per branch. After rebasing, git range-diff confirms the topic patches are unchanged apart from resolving accessor formatting, and git diff --check passes. Builds and runtime tests have not been rerun on the rebased tips.
Checklist