Workaround nvcc compiler hangs in libcudf debug build

[davidwendt]

Description

Adds noinline declaration to select functor operators that normally inline a significant amount code. Otherwise the compiler will appear to hang (run for many hours) trying to process and generate the ptx.
Recent changes appear to have pushed the size of the inlined code beyond some internal boundary.
Adding the noinline option may reduce the overall runtime for functions that use these utilities but only for a debug build and not for a release build where performance is required.

Checklist

• I am familiar with the Contributing Guidelines.
• New or existing tests cover these changes.
• The documentation is up to date with these changes.

[coderabbitai]

No actionable comments were generated in the recent review. 🎉

ℹ️ Recent review info

⚙️ Run configuration

Configuration used: Path: .coderabbit.yaml

Review profile: CHILL

Plan: Enterprise

Run ID: 619f3d55-92bf-44a7-a14d-28376ac5fd7f

📥 Commits

Reviewing files that changed from the base of the PR and between 3d52baa and 57ed4dc.

📒 Files selected for processing (2)

• cpp/src/groupby/streaming_groupby/common.cuh
• cpp/src/join/mixed_join_common_utils.cuh

---

📝 Walkthrough

Summary by CodeRabbit

• Chores
  • Internal compilation optimizations have been refined to improve performance debugging in development builds and maintain optimal runtime efficiency in production releases.

Walkthrough

This PR applies conditional inlining control to two device operator functions. The n_table_comparator::operator() in groupby and the pair_expression_equality<has_nulls>::operator() in join are each wrapped with #ifndef NDEBUG to force noinline behavior in debug builds while preserving release-build inlining strategy.

Changes

Device operator inlining control for debug vs release builds

Layer / File(s)	Summary
Debug-conditional inlining for device operators cpp/src/groupby/streaming_groupby/common.cuh, cpp/src/join/mixed_join_common_utils.cuh	n_table_comparator::operator() and pair_expression_equality<has_nulls>::operator() are each updated to apply __attribute__((noinline)) in debug builds via #ifndef NDEBUG guards, preserving release-build inlining behavior without changing operator logic or signatures.

---

Estimated code review effort

🎯 1 (Trivial) | ⏱️ ~3 minutes

🚥 Pre-merge checks | ✅ 5

✅ Passed checks (5 passed)

Check name	Status	Explanation
Title check	✅ Passed	The title clearly and accurately summarizes the main change: a workaround for nvcc compiler hangs in debug builds.
Description check	✅ Passed	The description is clearly related to the changeset, explaining the motivation, solution, and trade-offs of adding noinline declarations to functor operators.
Docstring Coverage	✅ Passed	No functions found in the changed files to evaluate docstring coverage. Skipping docstring coverage check.
Linked Issues check	✅ Passed	Check skipped because no linked issues were found for this pull request.
Out of Scope Changes check	✅ Passed	Check skipped because no linked issues were found for this pull request.

_{✏️ Tip: You can configure your own custom pre-merge checks in the settings.}

✨ Finishing Touches

🧪 Generate unit tests (beta)

• Create PR with unit tests

Warning

Review ran into problems

🔥 Problems

Stopped waiting for pipeline failures after 30000ms. One of your pipelines takes longer than our 30000ms fetch window to run, so review may not consider pipeline-failure results for inline comments if any failures occurred after the fetch window. Increase the timeout if you want to wait longer or run a @coderabbit review after the pipeline has finished.

---

_{Comment @coderabbitai help to get the list of available commands and usage tips.}

[davidwendt]

Here is a snippet of Claude's analysis of the build hanging src/groupby/streaming_groupby/insert_subsequent_nested.cu:

The hang was introduced by commit 962d15b (PR #21924, "Add streaming_groupby for stateful streaming aggregation"). It affects both insert_subsequent.cu and insert_subsequent_nested.cu.

The trigger is n_table_comparator::operator() in cpp/src/groupby/streaming_groupby/common.cuh:89-108, which has three call sites to RowEqT (device_row_comparator<has_nested, DYNAMIC, nan_equal>):

• common.cuh:95 batch_self_eq(...)
• common.cuh:99 cross_eqs[loc.first](lhs - max_distinct_keys, loc.second)
• common.cuh:103 cross_eqs[loc.first](rhs - max_distinct_keys, loc.second)

That body gets inlined into cuco::static_set::ref<insert_and_find> (rebinded via .rebind_key_eq(comparator) at
insert_subsequent.cuh:48), which is itself inlined into the predicate of thrust::copy_if. At -O0 no dead-code
elimination occurs, so every cuco probe step keeps all three comparator call sites live; with has_nested=true the
comparator carries the full type_dispatcher tree per call site. cicc's IR blows up before it ever reaches code
generation.

[davidwendt]

In the same session, I also asked Claude to analyze why the mixed_join_kernel.cu now also appears to hang:

• streaming_groupby hangs in the C++ frontend — template instantiation explosion before any IR is even generated.
• mixed_join hangs in the PTX assembler — cicc emits a huge but valid PTX, and ptxas can't get through register allocation /scheduling on it.

Why mixed_join_kernel.cu is huge:
The kernel mixed_join<has_nulls> (mixed_join_kernel.cuh:81-143) body contains:

  if (is_outer_join) {
    retrieve_matches<true>(...);   // __forceinline__
  } else {
    retrieve_matches<false>(...);  // __forceinline__
  }

is_outer_join is a runtime bool, so both instantiations get fully inlined into the kernel. Each retrieve_matches has a probe loop that calls pair_expression_equality::operator() (also __forceinline__) twice per bucket (slot 0 + slot 1), and that calls the AST expression_evaluator::evaluate (expression_evaluator.cuh, 791 lines; ~50 AST operator dispatch cases). With -O0 and -g -lineinfo, ptxas tracks line info through every inlined copy.

It's the same root cause family as streaming_groupby (oversized kernel from __forceinline__ template fan-out), but the explosion lands at ptxas instead of cicc because cicc can handle the IR, just not at the size it produces.

[PointKernel]

@davidwendt Thanks for the fix. Good news is that we will replace mixed join internals with normal hash join and filter_join_indicies in this release so the mixed join kernels will be gone (see #22124).

[davidwendt]

@PointKernel I would be curious if the noinline change in cpp/src/groupby/streaming_groupby/common.cuh would effect performance (in a release build). Seems the change may reduce register pressure significantly.

[PointKernel]

👍 We can check the nightly benchmark once it's merged.

[davidwendt]

👍 We can check the nightly benchmark once it's merged.

The nightly benchmarks are run on release code and not debug built code. So this change would have to be measured locally or by the nightly run on a draft PR that removes the #ifndef NDEBUG wrapper. What is the name of the benchmark that could verify this?

[davidwendt]

/merge

[PointKernel]

What is the name of the benchmark that could verify this?

./GROUPBY_NVBENCH -b 5 -a api=streaming

voila

[davidwendt]

@PointKernel Here is results from that benchmark on my local machine using noinline on the n_table_comparator::operator() in a Release build:

# groupby_max_cardinality

## [0] NVIDIA RTX A6000

|  num_aggregations  |  cardinality  |   Ref Time |   Cmp Time |         Diff |   %Diff |
|--------------------|---------------|------------|------------|--------------|---------|
|         1          |      20       |  11.572 ms |  11.566 ms |    -5.770 us |  -0.05% |
|         2          |      20       |  20.778 ms |  20.561 ms |  -216.966 us |  -1.04% |
|         3          |      20       |  29.776 ms |  29.408 ms |  -367.660 us |  -1.23% |
|         4          |      20       |  38.591 ms |  38.433 ms |  -157.464 us |  -0.41% |
|         5          |      20       |  47.709 ms |  47.931 ms |   222.771 us |   0.47% |
|         6          |      20       |  57.048 ms |  56.741 ms |  -307.554 us |  -0.54% |
|         7          |      20       |  66.286 ms |  65.885 ms |  -400.842 us |  -0.60% |
|         8          |      20       |  75.908 ms |  75.817 ms |   -90.658 us |  -0.12% |
|         1          |      50       |  14.900 ms |  14.879 ms |   -21.199 us |  -0.14% |
|         2          |      50       |  27.196 ms |  27.369 ms |   172.918 us |   0.64% |
|         3          |      50       |  39.764 ms |  39.792 ms |    27.681 us |   0.07% |
|         4          |      50       |  52.335 ms |  52.279 ms |   -55.994 us |  -0.11% |
|         5          |      50       |  64.493 ms |  64.526 ms |    32.748 us |   0.05% |
|         6          |      50       |  76.919 ms |  76.889 ms |   -29.153 us |  -0.04% |
|         7          |      50       |  89.473 ms |  89.633 ms |   160.351 us |   0.18% |
|         8          |      50       | 101.539 ms | 101.834 ms |   294.929 us |   0.29% |
|         1          |      100      |  12.136 ms |  12.110 ms |   -26.572 us |  -0.22% |
|         2          |      100      |  21.256 ms |  21.446 ms |   189.980 us |   0.89% |
|         3          |      100      |  30.662 ms |  30.563 ms |   -98.973 us |  -0.32% |
|         4          |      100      |  39.747 ms |  39.705 ms |   -41.677 us |  -0.10% |
|         5          |      100      |  48.788 ms |  48.889 ms |   101.093 us |   0.21% |
|         6          |      100      |  57.825 ms |  58.251 ms |   426.049 us |   0.74% |
|         7          |      100      |  67.186 ms |  67.521 ms |   334.501 us |   0.50% |
|         8          |      100      |  76.977 ms |  76.955 ms |   -21.540 us |  -0.03% |
|         1          |     1000      |   5.658 ms |   5.648 ms |    -9.629 us |  -0.17% |
|         2          |     1000      |   7.696 ms |   7.714 ms |    18.395 us |   0.24% |
|         3          |     1000      |   9.794 ms |   9.771 ms |   -23.726 us |  -0.24% |
|         4          |     1000      |  13.419 ms |  11.878 ms | -1541.330 us | -11.49% |
|         5          |     1000      |  14.654 ms |  13.954 ms |  -699.957 us |  -4.78% |
|         6          |     1000      |  18.359 ms |  16.132 ms | -2227.036 us | -12.13% |
|         7          |     1000      |  19.741 ms |  18.164 ms | -1577.090 us |  -7.99% |
|         8          |     1000      |  21.966 ms |  20.370 ms | -1596.146 us |  -7.27% |
|         1          |     10000     |   5.316 ms |   5.317 ms |     0.145 us |   0.00% |
|         2          |     10000     |   6.241 ms |   6.236 ms |    -5.247 us |  -0.08% |
|         3          |     10000     |   7.149 ms |   7.153 ms |     4.161 us |   0.06% |
|         4          |     10000     |   8.100 ms |   8.056 ms |   -43.801 us |  -0.54% |
|         5          |     10000     |   9.038 ms |   9.049 ms |    11.261 us |   0.12% |
|         6          |     10000     |   9.923 ms |   9.994 ms |    70.206 us |   0.71% |
|         7          |     10000     |  10.901 ms |  10.926 ms |    25.083 us |   0.23% |
|         8          |     10000     |  11.842 ms |  11.837 ms |    -4.250 us |  -0.04% |
|         1          |    100000     |   9.847 ms |   9.968 ms |   121.809 us |   1.24% |
|         2          |    100000     |  10.850 ms |  10.841 ms |    -9.100 us |  -0.08% |
|         3          |    100000     |  11.748 ms |  11.729 ms |   -19.117 us |  -0.16% |
|         4          |    100000     |  12.621 ms |  12.640 ms |    18.991 us |   0.15% |
|         5          |    100000     |  13.535 ms |  13.575 ms |    39.998 us |   0.30% |
|         6          |    100000     |  14.376 ms |  14.438 ms |    61.369 us |   0.43% |
|         7          |    100000     |  15.411 ms |  15.412 ms |     1.339 us |   0.01% |
|         8          |    100000     |  16.334 ms |  16.320 ms |   -13.540 us |  -0.08% |
|         1          |    1000000    |  11.713 ms |  11.868 ms |   154.464 us |   1.32% |
|         2          |    1000000    |  12.900 ms |  12.927 ms |    26.940 us |   0.21% |
|         3          |    1000000    |  13.844 ms |  13.970 ms |   125.494 us |   0.91% |
|         4          |    1000000    |  15.000 ms |  15.066 ms |    65.585 us |   0.44% |
|         5          |    1000000    |  16.037 ms |  16.012 ms |   -25.173 us |  -0.16% |
|         6          |    1000000    |  16.993 ms |  17.024 ms |    30.910 us |   0.18% |
|         7          |    1000000    |  18.103 ms |  18.133 ms |    30.247 us |   0.17% |
|         8          |    1000000    |  19.068 ms |  19.233 ms |   164.815 us |   0.86% |

(removed the columns which did not change value: I32, 20000000, streaming)

There appears to be no regression using noinline IMO.
So I think it may be worth setting this to noinline without the DEBUG caveat.
What do you think?

[PointKernel]

Thank you for double checking the performance impact. That aligns with my expectations, since the overall workflow is fairly complex and whether the row comparator is inlined or not likely won’t have a noticeable effect on performance. Consistently marking it as noinline makes total sense.