[BUG] HashCSR increases retained hash_join memory by 82% for large unique builds

[thirtiseven]

Describe the bug

An exact parent/child comparison confirms that #23640 increases the memory required for large unique-key cudf::hash_join builds. The same pure-libcudf workload that succeeds before the change fails with OOM after it, before allocating probe output.

Revisions:

• Parent: 884d2351bb37ee36dc8f2148615b6b1baadd3d4b
• Child: ae2fcb3a0b6f6d253347bb1095a794bf9a178d6a

Each input column occupies an additional 256,000,000 bytes. The table below excludes input storage from join memory and constructor peak.

Metric, 32M unique build rows, LF=0.5	Parent	Child
Persistent join bytes	512,001,919	933,306,447
Constructor peak bytes	512,001,919	1,189,370,958
Warm constructor median	10.264 ms	35.723 ms
Warm inner_join, 1M matching probe rows	0.513 ms	0.958 ms
Retain 30 objects	Passes, 3/3 runs	OOM constructing object 28, 3/3 runs

Persistent join memory increases 82.3%, and constructor peak increases 132.3%. When continuing until allocation fails, the parent retains 43 completed objects and the child retains 27, consistently across three runs. The parent's next failure is allocation of input column 44; the child's failure is a 128,000,000-byte allocation in constructor 28.

After all join objects and input columns are destroyed, synchronized RMM current bytes return to zero on both revisions, including expected-OOM cases. This appears to be increased representation/build workspace cost rather than a leak.

Steps/Code to reproduce bug

The GTest source hash_join_memory_tests.cpp uses only libcudf public APIs and RMM statistics. It can be added to cpp/tests/join and registered in the existing JOIN_TEST source list on either revision.

Each build input contains 32,000,000 unique int64 keys generated with cudf::sequence. The test retains the input columns and join objects, using disjoint key ranges across chunks. It sets an RMM statistics_resource_adaptor as the current resource and passes it explicitly to the constructor, with the same CUDA allocation backend on both revisions. It records synchronized input, persistent join, constructor peak, and post-destruction allocation counts.

Single-object measurement:

HASHCSR_ROWS=32000000 HASHCSR_CHUNKS=1 HASHCSR_LF=0.5 \
  ./cpp/build/gtests/JOIN_TEST \
    --gtest_filter=HashJoinMemory.RetainedGeneral --gtest_repeat=3

Retained-object comparison:

# Parent: all 30 objects succeed.
HASHCSR_ROWS=32000000 HASHCSR_CHUNKS=30 HASHCSR_LF=0.5 HASHCSR_EXPECT_OOM=0 \
  ./cpp/build/gtests/JOIN_TEST \
    --gtest_filter=HashJoinMemory.RetainedGeneral --gtest_repeat=3

# Child: expects the observed constructor OOM and verifies cleanup afterward.
HASHCSR_ROWS=32000000 HASHCSR_CHUNKS=30 HASHCSR_LF=0.5 HASHCSR_EXPECT_OOM=1 \
  ./cpp/build/gtests/JOIN_TEST \
    --gtest_filter=HashJoinMemory.RetainedGeneral --gtest_repeat=3

No query engine, file I/O, synthetic memory reservation, managed memory, or allocator limit is involved.

Complete reproducer: hash_join_memory_tests.cpp

/*
 * SPDX-FileCopyrightText: Copyright (c) 2026, NVIDIA CORPORATION.
 * SPDX-License-Identifier: Apache-2.0
 */
#include <cudf/column/column.hpp>
#include <cudf/filling.hpp>
#include <cudf/join/hash_join.hpp>
#include <cudf/join/distinct_hash_join.hpp>
#include <cudf/scalar/scalar.hpp>
#include <cudf/table/table_view.hpp>
#include <cudf/table/table.hpp>
#include <cudf/utilities/memory_resource.hpp>
#include <rmm/mr/statistics_resource_adaptor.hpp>
#include <gtest/gtest.h>
#include <cuda_runtime_api.h>
#include <algorithm>
#include <chrono>
#include <cstdlib>
#include <iostream>
#include <memory>
#include <string>
#include <type_traits>
#include <vector>

namespace {
std::string setting(char const* name, char const* fallback)
{
  auto p = std::getenv(name);
  return p ? p : fallback;
}
struct measured_resource {
  rmm::mr::statistics_resource_adaptor stats{cudf::get_current_device_resource_ref()};
  measured_resource() { cudf::set_current_device_resource(stats); }
  ~measured_resource() { cudf::set_current_device_resource(stats.get_upstream_resource()); }
};

template <typename Join>
void measure()
{
  ASSERT_EQ(cudaSetDevice(0), cudaSuccess);
  auto const rows = std::stoi(setting("HASHCSR_ROWS", "10000"));
  auto const chunks = std::stoi(setting("HASHCSR_CHUNKS", "1"));
  auto const lf = std::stod(setting("HASHCSR_LF", "0.5"));
  auto const multiplicity = std::stoi(setting("HASHCSR_MULTIPLICITY", "1"));
  auto const probe_rows = std::stoi(setting("HASHCSR_PROBE_ROWS", "256"));
  ASSERT_GT(rows, 0);
  ASSERT_GT(chunks, 0);
  ASSERT_GT(multiplicity, 0);
  ASSERT_EQ(rows % multiplicity, 0);
  ASSERT_GT(probe_rows, 0);
  if constexpr (std::is_same_v<Join, cudf::distinct_hash_join>) {
    ASSERT_EQ(multiplicity, 1);
  }
  measured_resource memory;
  auto& stats = memory.stats;
  auto stream = cudf::get_default_stream();
  auto snapshot = [&](char const* stage, int chunk) {
    stream.sync();
    std::size_t free{}, total{};
    ASSERT_EQ(cudaMemGetInfo(&free, &total), cudaSuccess);
    auto c = stats.get_bytes_counter();
    std::cout << "MEM stage=" << stage << " chunk=" << chunk
              << " current=" << c.value << " peak=" << c.peak
              << " cuda_free=" << free << " cuda_total=" << total << std::endl;
  };
  std::cout << "CONFIG rows=" << rows << " chunks=" << chunks << " lf=" << lf
            << " multiplicity=" << multiplicity << " probe_rows=" << probe_rows << std::endl;
  snapshot("baseline", 0);
  bool oom = false;
  {
    std::vector<std::unique_ptr<cudf::column>> inputs;
    std::vector<std::unique_ptr<Join>> joins;
    try {
      for (int i = 0; i < chunks; ++i) {
        {
          cudf::numeric_scalar<int64_t> start{int64_t{i} * rows};
          auto keys = cudf::sequence(rows / multiplicity, start);
          if (multiplicity == 1) {
            inputs.push_back(std::move(keys));
          } else {
            auto repeated = cudf::repeat(cudf::table_view{{keys->view()}}, multiplicity);
            auto columns = repeated->release();
            inputs.push_back(std::move(columns.front()));
          }
        }
        snapshot("input", i);
        stats.push_counters();
        auto begin = std::chrono::steady_clock::now();
        try {
          if constexpr (std::is_same_v<Join, cudf::hash_join>) {
            joins.push_back(std::make_unique<Join>(cudf::table_view{{inputs.back()->view()}},
              cudf::nullable_join::NO, cudf::null_equality::EQUAL, lf, stream, stats));
          } else {
            joins.push_back(std::make_unique<Join>(cudf::table_view{{inputs.back()->view()}},
              cudf::null_equality::EQUAL, lf, stream, stats));
          }
          stream.sync();
        } catch (...) {
          stats.pop_counters();
          throw;
        }
        auto elapsed = std::chrono::duration<double, std::milli>(
          std::chrono::steady_clock::now() - begin).count();
        auto c = stats.pop_counters().first;
        std::cout << "BUILD chunk=" << i << " persistent=" << c.value
                  << " peak=" << c.peak << " ms=" << elapsed << std::endl;
        snapshot("built", i);
      }
      cudf::numeric_scalar<int64_t> start{0};
      auto probe = cudf::sequence(probe_rows, start);
      stream.sync();
      for (int iteration = 0; iteration < 8; ++iteration) {
        stats.push_counters();
        auto begin = std::chrono::steady_clock::now();
        auto result = joins.front()->inner_join(cudf::table_view{{probe->view()}});
        stream.sync();
        auto elapsed = std::chrono::duration<double, std::milli>(
          std::chrono::steady_clock::now() - begin).count();
        auto c = stats.pop_counters().first;
        EXPECT_EQ(result.first->size(), std::min(rows / multiplicity, probe_rows) * multiplicity);
        EXPECT_EQ(result.first->size(), result.second->size());
        std::cout << "PROBE iteration=" << iteration << " count=" << result.first->size()
                  << " output=" << c.value << " peak=" << c.peak << " ms=" << elapsed << std::endl;
        if (iteration == 7) {
          std::vector<cudf::size_type> left(result.first->size()), right(result.second->size());
          ASSERT_EQ(cudaMemcpy(left.data(), result.first->data(), left.size() * sizeof(left[0]),
                               cudaMemcpyDeviceToHost), cudaSuccess);
          ASSERT_EQ(cudaMemcpy(right.data(), result.second->data(), right.size() * sizeof(right[0]),
                               cudaMemcpyDeviceToHost), cudaSuccess);
          std::vector<int> counts(std::min(rows / multiplicity, probe_rows));
          std::vector<bool> seen(std::size_t{counts.size()} * multiplicity);
          for (std::size_t j = 0; j < left.size(); ++j) {
            ASSERT_GE(left[j], 0);
            ASSERT_LT(left[j], counts.size());
            ASSERT_GE(right[j], 0);
            ASSERT_LT(right[j], seen.size());
            ASSERT_EQ(right[j] / multiplicity, left[j]);
            ASSERT_FALSE(seen[right[j]]);
            seen[right[j]] = true;
            ++counts[left[j]];
          }
          EXPECT_TRUE(std::all_of(counts.begin(), counts.end(),
                                 [&](auto n) { return n == multiplicity; }));
        }
      }
      snapshot("probed", chunks);
    } catch (rmm::out_of_memory const& e) {
      oom = true;
      std::cout << "OOM built_objects=" << joins.size() << " what=" << e.what() << std::endl;
    }
  }
  snapshot("destroyed", 0);
  EXPECT_EQ(stats.get_bytes_counter().value, 0);
  EXPECT_EQ(oom, setting("HASHCSR_EXPECT_OOM", "0") == "1");
}
}  // namespace

TEST(HashJoinMemory, RetainedGeneral) { measure<cudf::hash_join>(); }
TEST(HashJoinMemory, RetainedDistinct) { measure<cudf::distinct_hash_join>(); }

Expected behavior

Preserve a comparable retained-memory footprint for large unique/low-duplicate general join builds while retaining the high-multiplicity improvements. Callers using the general API may not have an a priori uniqueness guarantee.

Environment overview (please complete the following information)

• Environment location: Docker on a local bare-metal development host; one GPU exposed to the container.
• Method of cuDF install: built from source, separately for the exact parent and child commits.

Both libraries were built in full from clean source trees with the same dependencies, GCC 14, CUDA 12.9.86, Release configuration, and CUDA architecture 120. The test binaries were verified to load their respective build-tree libraries. Tests ran sequentially on the same RTX PRO 4500 Blackwell Server Edition GPU (32,623 MiB reported by nvidia-smi), driver 595.45.04.

The toolchain image was built locally and is not published; there is no public docker pull command for it. cuDF itself was built from the two source revisions above, rather than taken from the image's installed cuDF. The build invocation was as follows (host mount paths represented by variables; run once per revision):

docker run --rm --gpus device=0 --entrypoint bash \
  -v "$CUDF_SOURCE:/src:ro" \
  -v "$AB_DIR:/ab" \
  -v "$RAPIDS_CMAKE_SOURCE:/rapids-cmake:ro" \
  hl-dev-pr179-cudf456-deps:20260917 -lc '
    variant=$1
    cmake -S /src/cpp -B /ab/$variant -G Ninja \
      -DFETCHCONTENT_SOURCE_DIR_RAPIDS-CMAKE=/rapids-cmake \
      -DCMAKE_CUDA_ARCHITECTURES=120 -DCMAKE_BUILD_TYPE=Release \
      -DBUILD_TESTS=OFF -DBUILD_BENCHMARKS=OFF \
      -DCUDF_BUILD_TESTUTIL=OFF -DCUDF_BUILD_STREAMS_TEST_UTIL=OFF \
      -DCUDF_BUILD_STATIC_DEPS=OFF -DCPM_USE_LOCAL_PACKAGES=OFF \
      -DCPM_DOWNLOAD_rtcx=ON -DCPM_DOWNLOAD_nanoarrow=ON \
      -DCPM_SOURCE_CACHE=/ab/cache \
      -DCCCL_ROOT=/usr/local/lib64/rapids/cmake/cccl \
      -DCMAKE_INSTALL_PREFIX=/ab/install-$variant &&
    cmake --build /ab/$variant --target cudf -j 24
  ' bash "$VARIANT"

VARIANT is parent or child; CUDF_SOURCE points to its clean checkout. AB_DIR and RAPIDS_CMAKE_SOURCE are shared between the two builds. The standalone GTest translation unit was then linked against each build-tree cudf::cudf target, and its ldd output was checked. Alternatively, register the provided source in JOIN_TEST as described above.

Environment details

Both libraries were built from source in the same Docker toolchain on this local development host. Both revisions used the same frozen dependencies (not independently resolved historical environments): RMM 26.10, CCCL 3.5.0.0, cuCollections 4b26118c99866221f99f35f4e3bc74afdbe063bc, nanoarrow 0.8.0, and rtcx a9f63f8cdd4b0b41a2d88a9f705576a61b4222ec. The libraries were built using their native CMake targets with CMAKE_BUILD_TYPE=Release, CMAKE_CUDA_ARCHITECTURES=120, and CUDF_BUILD_STATIC_DEPS=OFF. The attached test translation unit was linked separately with GTest; the full cuDF unit-test suite was not run.

For the reported timing medians, the first of three process-local repetitions is excluded. Each repetition probes eight times; probe iterations 0–2 are excluded. The 1M-probe measurements use HASHCSR_PROBE_ROWS=1000000.

The following is actual print_env.sh output collected after the experiments in the same toolchain image with GPU 0 exposed. It reports "Not inside a git repository" because the checkout was mounted without its external worktree Git metadata; the two commit IDs above were verified on the host. The CUDA version in nvidia-smi is driver capability; the compiler/toolkit used was 12.9.86, shown under nvcc.

Click here to see environment details

 **git***
 Not inside a git repository
 
 ***OS Information***
 CentOS Stream release 9
 NAME="CentOS Stream"
 VERSION="9"
 ID="centos"
 ID_LIKE="rhel fedora"
 VERSION_ID="9"
 PLATFORM_ID="platform:el9"
 PRETTY_NAME="CentOS Stream 9"
 ANSI_COLOR="0;31"
 LOGO="fedora-logo-icon"
 CPE_NAME="cpe:/o:centos:centos:9"
 HOME_URL="https://centos.org/"
 BUG_REPORT_URL="https://issues.redhat.com/"
 REDHAT_SUPPORT_PRODUCT="Red Hat Enterprise Linux 9"
 REDHAT_SUPPORT_PRODUCT_VERSION="CentOS Stream"
 CentOS Stream release 9
 CentOS Stream release 9
 Linux 9a4862aa23e8 6.8.0-100-generic #100-Ubuntu SMP PREEMPT_DYNAMIC Tue Jan 13 16:40:06 UTC 2026 x86_64 x86_64 x86_64 GNU/Linux
 
 ***GPU Information***
 Fri Sep 18 05:50:58 2026
 +-----------------------------------------------------------------------------------------+
 | NVIDIA-SMI 595.45.04              Driver Version: 595.45.04      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 RTX PRO 4500 Blac...    On  |   00000000:10:00.0 Off |                    0 |
 | N/A   34C    P8             17W /  165W |       0MiB /  32623MiB |      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):                                  96
 On-line CPU(s) list:                     0-95
 Vendor ID:                               GenuineIntel
 Model name:                              Intel(R) Xeon(R) 6741P
 CPU family:                              6
 Model:                                   173
 Thread(s) per core:                      2
 Core(s) per socket:                      48
 Socket(s):                               1
 Stepping:                                1
 CPU(s) scaling MHz:                      99%
 CPU max MHz:                             3800.0000
 CPU min MHz:                             800.0000
 BogoMIPS:                                5000.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 hfi 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:                               2.3 MiB (48 instances)
 L1i cache:                               3 MiB (48 instances)
 L2 cache:                                96 MiB (48 instances)
 L3 cache:                                288 MiB (1 instance)
 NUMA node(s):                            1
 NUMA node0 CPU(s):                       0-95
 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 Not affected; 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/local/bin/cmake
 cmake version 4.3.2
 
 CMake suite maintained and supported by Kitware (kitware.com/cmake).
 
 ***g++***
 /opt/rh/gcc-toolset-14/root/usr/bin/g++
 g++ (GCC) 14.2.1 20250110 (Red Hat 14.2.1-13)
 Copyright (C) 2024 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***
 /usr/local/cuda/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***
 /usr/bin/python
 Python 3.9.25
 
 ***Environment Variables***
 PATH                            : /usr/share/Modules/bin:/opt/rh/gcc-toolset-14/root/usr/bin:/usr/lib/jvm/java-17-openjdk/bin:/usr/lib/jvm/java-1.8.0-openjdk/bin:/usr/local/cuda/bin:/usr/local/sbin:/usr/local/bin:/usr/sbin:/usr/bin:/sbin:/bin
 LD_LIBRARY_PATH                 :
 NUMBAPRO_NVVM                   :
 NUMBAPRO_LIBDEVICE              :
 CONDA_PREFIX                    :
 PYTHON_PATH                     :
 
 conda not found
 ***pip packages***
 /usr/bin/pip
 Package                      Version
 ---------------------------- -----------
 certifi                      2026.7.22
 charset-normalizer           3.5.1
 click                        8.1.8
 crc32c                       2.3
 dbus-python                  1.2.18
 distro                       1.5.0
 Flask                        2.2.4
 googleapis-common-protos     1.59.0
 googleapis-storage-testbench 0.33.0
 gpg                          1.15.1
 grpcio                       1.54.0
 gunicorn                     20.1.0
 idna                         3.19
 importlib_metadata           8.7.1
 itsdangerous                 2.2.0
 Jinja2                       3.1.6
 libcomps                     0.1.18
 MarkupSafe                   3.0.3
 pip                          21.3.1
 protobuf                     4.22.3
 python-dateutil              2.9.0.post0
 requests                     2.32.5
 requests-toolbelt            1.0.0
 rpm                          4.16.1.3
 scalpl                       0.4.2
 selinux                      3.6
 sepolicy                     3.6
 setools                      4.4.4
 setuptools                   82.0.1
 six                          1.15.0
 systemd-python               234
 urllib3                      2.6.3
 waitress                     2.1.2
 Werkzeug                     3.1.8
 zipp                         3.23.1

Additional context

Controls:

• distinct_hash_join at LF=0.5 retains exactly 512,000,087 bytes on both revisions and succeeds with 30 objects on both.
• With two build rows per key, general hash_join has the same memory footprints shown above. Thus the increase is not limited to strictly unique input.
• LF=0.75 reduces the parent's persistent state to 341,333,631 bytes, but the child remains at 933,306,447 bytes because its capacity rounds to the same power of two.
• Child LF=1.0 reduces persistent state to 530,653,263 bytes and allows 30 objects, but raises warm unique-build time to approximately 50.6 ms. It is not a performance-neutral workaround. Parent LF=1.0 has a different actual occupancy and very slow probing, so equal LF parameters do not imply equal occupancy across these representations.

HashCSR also delivers its intended benefit: for a separate 100K-row build with 1,000 rows per key, build time improves from 1.984 ms to 0.168 ms and inner_join from 3.646 ms to 0.0746 ms. That probe uses 256 keys, of which 100 match, producing 100K pairs. The concern is the tradeoff for large unique/low-duplicate build states, especially when callers retain multiple OO join objects.

The child retains _entries[capacity], _cumulative_ends[capacity], and _values[rows] even when all build keys are unique. Measured persistent bytes match 12 * capacity + 4 * rows, plus 79 bytes of tracked state. For 32M rows at LF=0.5, capacity is 67,108,864.

Could we reduce the unique/low-duplicate build footprint while preserving the high-multiplicity improvements? The unchanged distinct join provides a useful baseline, but general hash_join callers do not necessarily have a uniqueness guarantee.

The test validates output indices and cardinality for the exercised non-null int64 inner joins. Timing uses synchronized wall-clock measurements after warmup, with three process-local repetitions; clocks were not locked and this is not a full NVBench performance study. The byte counts and identical-workload pass/OOM comparison are the primary regression evidence.

The exact OOM boundary depends on available GPU memory and allocator configuration; the single-object byte comparison is the more portable reproducer.

Representative raw log excerpts (first repetition; allocator source paths omitted)

parent / unique-lf050

CONFIG rows=32000000 chunks=1 lf=0.5 multiplicity=1 probe_rows=256
MEM stage=baseline chunk=0 current=0 peak=0 cuda_free=33414905856 cuda_total=33689829376
BUILD chunk=0 persistent=512001919 peak=512001919 ms=13.6528
MEM stage=destroyed chunk=0 current=0 peak=768006111 cuda_free=33412808704 cuda_total=33689829376
[  PASSED  ] 1 test.

parent / retained30-lf050

CONFIG rows=32000000 chunks=30 lf=0.5 multiplicity=1 probe_rows=256
MEM stage=baseline chunk=0 current=0 peak=0 cuda_free=33414905856 cuda_total=33689829376
BUILD chunk=0 persistent=512001919 peak=512001919 ms=11.3258
MEM stage=destroyed chunk=0 current=0 peak=23040061762 cuda_free=33412808704 cuda_total=33689829376
[  PASSED  ] 1 test.

parent / boundary60-lf050

CONFIG rows=32000000 chunks=60 lf=0.5 multiplicity=1 probe_rows=256
MEM stage=baseline chunk=0 current=0 peak=0 cuda_free=33414905856 cuda_total=33689829376
BUILD chunk=0 persistent=512001919 peak=512001919 ms=13.6256
OOM built_objects=43 what=std::bad_alloc: out_of_memory: CUDA error (failed to allocate 256000000 bytes): cudaErrorMemoryAllocation out of memory
MEM stage=destroyed chunk=0 current=0 peak=33024082526 cuda_free=33412808704 cuda_total=33689829376
[  PASSED  ] 1 test.

child / unique-lf050

CONFIG rows=32000000 chunks=1 lf=0.5 multiplicity=1 probe_rows=256
MEM stage=baseline chunk=0 current=0 peak=0 cuda_free=33414905856 cuda_total=33689829376
BUILD chunk=0 persistent=933306447 peak=1189370958 ms=36.2147
MEM stage=destroyed chunk=0 current=0 peak=1445370958 cuda_free=33412808704 cuda_total=33689829376
[  PASSED  ] 1 test.

child / retained30-lf050

CONFIG rows=32000000 chunks=30 lf=0.5 multiplicity=1 probe_rows=256
MEM stage=baseline chunk=0 current=0 peak=0 cuda_free=33414905856 cuda_total=33689829376
BUILD chunk=0 persistent=933306447 peak=1189370958 ms=36.2269
OOM built_objects=27 what=std::bad_alloc: out_of_memory: CUDA error (failed to allocate 128000000 bytes): cudaErrorMemoryAllocation out of memory
MEM stage=destroyed chunk=0 current=0 peak=33172580516 cuda_free=33412808704 cuda_total=33689829376
[  PASSED  ] 1 test.

child / boundary60-lf050

CONFIG rows=32000000 chunks=60 lf=0.5 multiplicity=1 probe_rows=256
MEM stage=baseline chunk=0 current=0 peak=0 cuda_free=33414905856 cuda_total=33689829376
BUILD chunk=0 persistent=933306447 peak=1189370958 ms=36.2615
OOM built_objects=27 what=std::bad_alloc: out_of_memory: CUDA error (failed to allocate 128000000 bytes): cudaErrorMemoryAllocation out of memory
MEM stage=destroyed chunk=0 current=0 peak=33172580516 cuda_free=33412808704 cuda_total=33689829376
[  PASSED  ] 1 test.