Restore GCC 14 with CUDA 12.9 shuffle workaround (#4939)

Fixes #4937.

## Summary

- restore the JNI build toolchain to GCC 14
- replace the `shuffle_split` `std::pair` return type with an aggregate
result
- preserve structured-binding behavior and update the explicit test
container type

## Root cause

CUDA 12.9 `cudafe++` crashes while instantiating GCC 14's concepts-based
`std::pair`
constructor constraints for the shuffle result types. This previously
forced the JNI build back to
GCC 13. ARM64 nvCOMP 5.3 is built with GCC 14 and requires
`__cxa_call_terminate`, which the GCC 13
link does not provide.

Using an aggregate avoids the compiler crash without changing the
shuffle algorithm. GCC 14 can
then link nvCOMP 5.3 and supplies `__cxa_call_terminate` through its
nonshared compatibility
library, without adding a newer dynamic libstdc++ requirement.

This replaces the GCC 13 toolchain workaround introduced by #4468 while
retaining compatibility
with CUDA 12.9.

## Validation

- reproduced the ARM64 GCC 13 final-link failure: nvCOMP 5.3.0.16
contains 20 undefined
`__cxa_call_terminate` references, and GCC 13 does not provide the
symbol
- passed the matching ARM64 Rocky Linux 8 GCC 14 final-link probe; GCC
14 provides
  `__cxa_call_terminate` through `libstdc++_nonshared.a`
- reproduced the `shuffle_split.cu` crash with CUDA 12.9.86 and GCC
14.2.1 before the change
- compiled and linked the JNI native libraries with CUDA 12.9.86, GCC
14.2.1, and nvCOMP 5.3.0.16
  after the change
- verified the resulting `libcudf.so` has no unresolved relocations with
Rocky Linux 8
  `libstdc++.so.6.0.25`
- completed a clean x86_64 package build with CUDA 12.9.86 and GCC
14.2.1
- passed all 18 `SHUFFLE_SPLIT` native tests, including `MixedValidity`
- repository-pinned clang-format v20.1.8 passes

---------

Signed-off-by: Chong Gao <chongg@nvidia.com>
Co-authored-by: Chong Gao <res_life@163.com>

Author: res-life

build/env.sh

@@ -16,4 +16,4 @@
 #
 
 set -ex
-export sclCMD=${sclCMD:-"scl enable gcc-toolset-13"}
+export sclCMD=${sclCMD:-"scl enable gcc-toolset-14"}

ci/Dockerfile

@@ -42,7 +42,7 @@ ARG CMAKE_ARCH=x86_64
 
 ### Install basic requirements
 # pin urllib3<2.0 for https://github.com/psf/requests/issues/6432
-RUN dnf --enablerepo=powertools install -y scl-utils gcc-toolset-13 gcc-toolset-${TOOLSET_VERSION} python39 \
+RUN dnf --enablerepo=powertools install -y scl-utils gcc-toolset-14 gcc-toolset-${TOOLSET_VERSION} python39 \
     zlib-devel maven tar wget patch ninja-build git zip && \
   alternatives --set python /usr/bin/python3 && \
   python -m pip install requests 'urllib3<2.0'

src/main/cpp/src/shuffle_split.cu

@@ -794,11 +794,10 @@ shuffle_split_metadata compute_metadata(cudf::table_view const& input,
 /**
  * @copydoc spark_rapids_jni::shuffle_split
  */
-std::pair<shuffle_split_result, shuffle_split_metadata> shuffle_split(
-  cudf::table_view const& input,
-  std::vector<size_type> const& splits,
-  rmm::cuda_stream_view stream,
-  rmm::device_async_resource_ref mr)
+shuffle_split_output shuffle_split(cudf::table_view const& input,
+                                   std::vector<size_type> const& splits,
+                                   rmm::cuda_stream_view stream,
+                                   rmm::device_async_resource_ref mr)
 {
   SRJ_FUNC_RANGE();
 

src/main/cpp/src/shuffle_split.hpp

@@ -106,6 +106,13 @@ struct shuffle_split_result {
   rmm::device_uvector<size_t> offsets{0, cudf::get_default_stream()};
 };
 
+struct shuffle_split_output {
+  // Keep this as an aggregate instead of std::pair. CUDA 12.9 cudafe++ crashes when GCC 14's
+  // concepts-based std::pair constructor constraints are instantiated for these result types.
+  shuffle_split_result result;
+  shuffle_split_metadata metadata;
+};
+
 /**
  * @brief Performs a split operation on a cudf table, returning a buffer of data containing
  * all of the sub-tables as a contiguous buffer of anonymous bytes.
@@ -133,11 +140,10 @@ struct shuffle_split_result {
  * partition, and a shuffle_split_metadata struct which contains the metadata needed to reconstruct
  * a table using shuffle_assemble.
  */
-std::pair<shuffle_split_result, shuffle_split_metadata> shuffle_split(
-  cudf::table_view const& input,
-  std::vector<cudf::size_type> const& splits,
-  rmm::cuda_stream_view stream,
-  rmm::device_async_resource_ref mr);
+shuffle_split_output shuffle_split(cudf::table_view const& input,
+                                   std::vector<cudf::size_type> const& splits,
+                                   rmm::cuda_stream_view stream,
+                                   rmm::device_async_resource_ref mr);
 
 /**
  * @brief Buffer slice representing a portion of the shared allocated buffer

src/main/cpp/tests/shuffle_split.cu

@@ -1,5 +1,5 @@
 /*
- * Copyright (c) 2025, NVIDIA CORPORATION.
+ * Copyright (c) 2025-2026, NVIDIA CORPORATION.
  *
  * Licensed under the Apache License, Version 2.0 (the "License");
  * you may not use this file except in compliance with the License.
@@ -867,26 +867,24 @@ TEST_F(ShuffleSplitTests, MixedValidity)
       cudf::table_view expected_t{{static_cast<cudf::column_view>(*expected)}};
 
       // make the concatenated shuffle_split partitions
-      std::vector<
-        std::pair<spark_rapids_jni::shuffle_split_result, spark_rapids_jni::shuffle_split_metadata>>
-        shuf;
+      std::vector<spark_rapids_jni::shuffle_split_output> shuf;
       size_t total_size = 0;
       for (size_t idx = 0; idx < partition_views.size(); idx++) {
         shuf.push_back(spark_rapids_jni::shuffle_split(
           cudf::table_view{{partition_views[idx]}}, {}, stream, mr));
-        total_size += shuf.back().first.partitions->size();
+        total_size += shuf.back().result.partitions->size();
       }
       rmm::device_uvector<uint8_t> full{total_size, stream, mr};
       rmm::device_uvector<size_t> full_offsets{partition_views.size() + 1, stream, mr};
       std::vector<size_t> h_full_offsets(partition_views.size() + 1);
       size_t pos = 0;
       for (size_t idx = 0; idx < partition_views.size(); idx++) {
         cudaMemcpy(static_cast<uint8_t*>(full.data()) + pos,
-                   shuf[idx].first.partitions->data(),
-                   shuf[idx].first.partitions->size(),
+                   shuf[idx].result.partitions->data(),
+                   shuf[idx].result.partitions->size(),
                    cudaMemcpyDeviceToDevice);
         h_full_offsets[idx] = pos;
-        pos += shuf[idx].first.partitions->size();
+        pos += shuf[idx].result.partitions->size();
       }
       h_full_offsets[partition_views.size()] = pos;
       cudaMemcpy(full_offsets.data(),