Skip to content
Merged
Show file tree
Hide file tree
Changes from all commits
Commits
File filter

Filter by extension

Filter by extension

Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
5 changes: 4 additions & 1 deletion cmake.spec
Original file line number Diff line number Diff line change
@@ -1,11 +1,14 @@
### RPM external cmake 3.31.7
### RPM external cmake 3.31.12
%define downloaddir %(echo %realversion | cut -d. -f1,2)
Source: http://www.cmake.org/files/v%{downloaddir}/%n-%realversion.tar.gz
# Note: this patch should be removed after updating to CMake 4.4+
Patch0: patches/cmake_cuda_std_23
Requires: bz2lib curl expat zlib
BuildRequires: gmake

%prep
%setup -n cmake-%realversion
%patch0 -p1

%build
cat > build-flags.cmake <<- EOF
Expand Down
4 changes: 2 additions & 2 deletions cuda-flags.file
Original file line number Diff line number Diff line change
Expand Up @@ -2,8 +2,8 @@
## INCLUDE cpp-standard
# define the CUDA compilation flags in a way that can be shared by SCRAM-based and regular tools

# build support for Volta (7.0), Turing (7.5), Ampere (8.x), Lovelace (8.9), Hopper (9.0) and Blackwell (10.0, 12.0)
%define cuda_arch 70 75 80 89 90 100 120
# build support for Turing (7.5), Ampere (8.x), Lovelace (8.9), Hopper (9.0) and Blackwell (10.0, 12.0)
%define cuda_arch 75 80 89 90 100 120

# LIBOMPTARGET_NVPTX_COMPUTE_CAPABILITIES style for listing the supported CUDA compute architectures
%define omptarget_cuda_archs %(echo $(for ARCH in %cuda_arch; do echo "$ARCH"; done) | sed -e"s/ /,/g")
Expand Down
13 changes: 8 additions & 5 deletions cuda.spec
Original file line number Diff line number Diff line change
@@ -1,15 +1,16 @@
### RPM external cuda 12.9.1
### RPM external cuda 13.3.0
## INITENV +PATH LD_LIBRARY_PATH %i/lib64

%define runpath_opts -m compute-sanitizer -m drivers -m nvvm
%define driversversion 575.57.08
%define driversversion 610.43.02

%ifarch x86_64
Source0: https://developer.download.nvidia.com/compute/cuda/%{realversion}/local_installers/%{n}_%{realversion}_%{driversversion}_linux.run
%endif
%ifarch aarch64
Source0: https://developer.download.nvidia.com/compute/cuda/%{realversion}/local_installers/%{n}_%{realversion}_%{driversversion}_linux_sbsa.run
%endif
Patch0: patches/cuda_cccl_8771
Requires: python3
AutoReq: no

Expand Down Expand Up @@ -53,13 +54,11 @@ chmod a+x %_builddir/build/extras/CUPTI/lib64/*.so*
mv %_builddir/build/extras/CUPTI/lib64/*.so* %{i}/lib64/
mv %_builddir/build/extras/CUPTI/include/*.h %{i}/include/

# leave out the Nsight and NVVP graphical tools, and package the other binaries
rm -f %_builddir/build/bin/computeprof
# leave out the graphical tools, and package the other binaries
rm -f %_builddir/build/bin/cuda-uninstaller
rm -f %_builddir/build/bin/ncu*
rm -f %_builddir/build/bin/nsight*
rm -f %_builddir/build/bin/nsys*
rm -f %_builddir/build/bin/nvvp
mv %_builddir/build/bin %{i}/

# package the cuda-gdb support files, and rename the binary to use it via a wrapper
Expand Down Expand Up @@ -107,6 +106,10 @@ cp -p %_builddir/build/drivers/libnvidia-ml.so.%{driversversion}
ln -sf libnvidia-ml.so.%{driversversion} %{i}/lib64/stubs/libnvidia-ml.so.1
ln -sf libnvidia-ml.so.1 %{i}/lib64/stubs/libnvidia-ml.so

# fix invalid C++ syntax in cub/device/device_transform.cuh
# backport NVIDIA/cccl#8771 from the main branch
patch -d %{i}/include/cccl -p2 < %{PATCH0}

%post
# let nvcc find its components when invoked from the command line
sed \
Expand Down
4 changes: 2 additions & 2 deletions cudnn.spec
Original file line number Diff line number Diff line change
@@ -1,7 +1,7 @@
### RPM external cudnn 9.9.0.52
### RPM external cudnn 9.23.0.39
## INITENV +PATH LD_LIBRARY_PATH %i/lib64

%define cudaver 12
%define cudaver 13

# NVIDIA uses sbsa for aarch64, and the standard architecture name for ppc64le and x86_64
%ifarch aarch64
Expand Down
2 changes: 1 addition & 1 deletion gdrcopy.spec
Original file line number Diff line number Diff line change
@@ -1,4 +1,4 @@
### RPM external gdrcopy 2.5.1
### RPM external gdrcopy 2.6
## INITENV +PATH LD_LIBRARY_PATH %i/lib64
Source: https://github.com/NVIDIA/%{n}/archive/v%{realversion}.tar.gz
Requires: cuda
Expand Down
61 changes: 61 additions & 0 deletions onnxruntime/absl_nvidia_bug_6302392.patch
Original file line number Diff line number Diff line change
@@ -0,0 +1,61 @@
diff --git a/absl/container/internal/common.h b/absl/container/internal/common.h
index 5ef6c569..86de28c6 100644
--- a/absl/container/internal/common.h
+++ b/absl/container/internal/common.h
@@ -68,6 +68,13 @@ struct IfRRef<T&&> {
using AddPtr = Other*;
};

+// Workaround for NVCC (cudafe++) which fails to parse
+// `IfRRef<concrete_type>::AddPtr<Dependent>` when used as a template argument
+// inside a heavily macro-expanded template parameter list. Top-level alias
+// defers the member-template lookup outside the surrounding template-id.
+template <class T, class Other>
+using IfRRefAddPtr = typename IfRRef<T>::template AddPtr<Other>;
+
template <class, class = void>
struct IsTransparent : std::false_type {};
template <class T>
diff --git a/absl/container/internal/raw_hash_map.h b/absl/container/internal/raw_hash_map.h
index b42a4f22..540788b7 100644
--- a/absl/container/internal/raw_hash_map.h
+++ b/absl/container/internal/raw_hash_map.h
@@ -106,8 +106,8 @@ class raw_hash_map : public raw_hash_set<Policy, Hash, Eq, Alloc> {
typename K = key_type, class V = mapped_type, \
ABSL_INTERNAL_IF_##KValue##_NOR_##VValue( \
int = (EnableIf<LifetimeBoundKV<K, KValue, V, VValue, \
- IfRRef<int KQual>::AddPtr<K>, \
- IfRRef<int VQual>::AddPtr<V>>>()), \
+ IfRRefAddPtr<int KQual, K>, \
+ IfRRefAddPtr<int VQual, V>>>()), \
ABSL_INTERNAL_SINGLE_ARG( \
int &..., \
decltype(EnableIf<LifetimeBoundKV<K, KValue, V, VValue>>()) = \
diff --git a/absl/container/internal/btree_container.h b/absl/container/internal/btree_container.h
index 21f00ae4..0a017f9f 100644
--- a/absl/container/internal/btree_container.h
+++ b/absl/container/internal/btree_container.h
@@ -497,8 +497,8 @@ class btree_map_container : public btree_set_container<Tree> {
typename K = key_type, class M, \
ABSL_INTERNAL_IF_##KValue##_NOR_##MValue( \
int = (EnableIf<LifetimeBoundKV<K, KValue, M, MValue, \
- IfRRef<int KQual>::AddPtr<K>, \
- IfRRef<int MQual>::AddPtr<M>>>()), \
+ IfRRefAddPtr<int KQual, K>, \
+ IfRRefAddPtr<int MQual, M>>>()), \
ABSL_INTERNAL_SINGLE_ARG( \
int &..., \
decltype(EnableIf<LifetimeBoundKV<K, KValue, M, MValue>>()) = \
@@ -598,10 +598,10 @@ class btree_map_container : public btree_set_container<Tree> {
ABSL_INTERNAL_IF_##KValue( \
class... Args, \
int = (EnableIf< \
- LifetimeBoundK<K, KValue, IfRRef<int KQual>::AddPtr<K>>>())), \
+ LifetimeBoundK<K, KValue, IfRRefAddPtr<int KQual, K>>>())), \
ABSL_INTERNAL_IF_##KValue( \
decltype(EnableIf<LifetimeBoundK< \
- K, KValue, IfRRef<int KQual>::AddPtr<K>>>()) = 0, \
+ K, KValue, IfRRefAddPtr<int KQual, K>>>()) = 0, \
class... Args), \
std::enable_if_t<!std::is_convertible<K, const_iterator>::value, int> = \
0> \
29 changes: 29 additions & 0 deletions onnxruntime/backport28736.patch
Original file line number Diff line number Diff line change
@@ -0,0 +1,29 @@
diff --git a/onnxruntime/contrib_ops/cuda/quantization/matmul_4bits.cu b/onnxruntime/contrib_ops/cuda/quantization/matmul_4bits.cu
index cb9eab579a..ac93777a58 100644
--- a/onnxruntime/contrib_ops/cuda/quantization/matmul_4bits.cu
+++ b/onnxruntime/contrib_ops/cuda/quantization/matmul_4bits.cu
@@ -281,8 +281,9 @@ __global__ void __launch_bounds__(kWarpSize* kColsPerThreadBlock) MatMulFloatInt
#define UnRollReduction(unroll_size) \
do { \
constexpr int kUnroll = unroll_size; \
- constexpr int kUnrollMask = 0xffffffff & (~(kUnroll * k_per_iter - 1)); \
- for (; k_id < (k & kUnrollMask); k_id += kUnroll * k_per_iter) { \
+ constexpr int kUnrollStep = kUnroll * k_per_iter; \
+ const int k_unroll_bound = k - k % kUnrollStep; \
+ for (; k_id < k_unroll_bound; k_id += kUnrollStep) { \
_Pragma("unroll") for (int i = 0; i < kUnroll; i++) { \
uint32_t value = *(reinterpret_cast<const uint32_t*>(b_data_quant + k_per_iter / 2 * i)); \
T scale = b_scale_vec[t_meta_k + k_per_iter / block_size * i]; \
diff --git a/onnxruntime/core/providers/cuda/tensor/concat_impl.cu b/onnxruntime/core/providers/cuda/tensor/concat_impl.cu
index f369ead0f9..8b27b0dd93 100644
--- a/onnxruntime/core/providers/cuda/tensor/concat_impl.cu
+++ b/onnxruntime/core/providers/cuda/tensor/concat_impl.cu
@@ -80,7 +80,7 @@ Status ConcatSameConcatDimImpl(cudaStream_t stream, const size_t element_bytes,
template Status ConcatSameConcatDimImpl<const void**>(cudaStream_t stream, const size_t element_bytes,
const int block_size_including_axis_dim,
const int block_size_inside_axis_dim, const int64_t concat_size,
- void* output_data, const void** input_data,
+ void* output_data, const void** const input_data,
const size_t output_size);

// input tensor addresses passed by value
13 changes: 13 additions & 0 deletions onnxruntime/ciso646.patch
Original file line number Diff line number Diff line change
@@ -0,0 +1,13 @@
diff --git a/onnxruntime/core/providers/cuda/cuda_common.h b/onnxruntime/core/providers/cuda/cuda_common.h
index f90eb2813a..883491791f 100644
--- a/onnxruntime/core/providers/cuda/cuda_common.h
+++ b/onnxruntime/core/providers/cuda/cuda_common.h
@@ -7,7 +7,7 @@

// The following three lines were copied from ABSL
// cutlass needs them, because cutlass uses "and"/"or" keywords
-#ifdef __cplusplus
+#if defined(__cplusplus) && __cplusplus < 202002L
#include <ciso646>
#endif

9 changes: 9 additions & 0 deletions onnxruntime/spec
Original file line number Diff line number Diff line change
Expand Up @@ -6,6 +6,10 @@
%define tag v%{realversion}
Source: git+https://github.com/%{github_user}/%{n}.git?obj=%{branch}/%{tag}&export=%{n}-%{realversion}&submodules=1&output=/%{n}-%{realversion}.tgz
Patch0: onnxruntime/cms-changes
# Note: these two patches should likely be removed for onnxruntime 1.27+
Patch1: onnxruntime/backport28736
Patch2: onnxruntime/absl_nvidia_bug_6302392
Patch3: onnxruntime/ciso646

BuildRequires: cmake ninja
Requires: protobuf py3-numpy py3-onnx zlib libpng py3-pybind11 re2 eigen
Expand All @@ -14,6 +18,8 @@ Requires: protobuf py3-numpy py3-onnx zlib libpng py3-pybind11 re2 eigen
%prep
%setup -q -n %{n}-%{realversion}
%patch0 -p1
%patch1 -p1
%patch3 -p1

%build
rm -rf ../build; mkdir ../build; cd ../build
Expand Down Expand Up @@ -59,6 +65,9 @@ cmake ../%{n}-%{realversion}/cmake \
# -Wno-error=overloaded-virtual
# See https://github.com/intel/neural-speed/issues/188

# patch Abseil bundled with ONNX to workaround NVIDIA bug #6302392
patch -d %{_builddir}/build/_deps/abseil_cpp-src -p1 < %{PATCH2}

ninja -v %{makeprocesses}
python3 ../%{n}-%{realversion}/setup.py build

Expand Down
40 changes: 40 additions & 0 deletions patches/cmake_cuda_std_23.patch
Original file line number Diff line number Diff line change
@@ -0,0 +1,40 @@
commit d01c9a429b25fcea8f5b77547adc23f3fede0471
Author: Arha Gatram <agatram@nvidia.com>
Date: Wed Jun 10 15:53:55 2026 +0000

CUDA: Add support for cuda_std_23 for nvcc 13.3+

CUDA 13.3 has been released with support for `-std=c++23`.

diff --git a/Modules/Compiler/NVIDIA.cmake b/Modules/Compiler/NVIDIA.cmake
index 0d0f03ad00..5bf7044159 100644
--- a/Modules/Compiler/NVIDIA.cmake
+++ b/Modules/Compiler/NVIDIA.cmake
@@ -46,6 +46,14 @@ macro(__compiler_nvidia_cxx_standards lang)
set(CMAKE_${lang}_STANDARD_LATEST 20)
endif()
endif()
+
+ if (NOT CMAKE_${lang}_COMPILER_VERSION VERSION_LESS 13.3)
+ if(CMAKE_${lang}_SIMULATE_VERSION VERSION_GREATER_EQUAL 19.29.30129)
+ set(CMAKE_${lang}23_STANDARD_COMPILE_OPTION "-std=c++23")
+ set(CMAKE_${lang}23_EXTENSION_COMPILE_OPTION "-std=c++23")
+ set(CMAKE_${lang}_STANDARD_LATEST 23)
+ endif()
+ endif()
else()
set(CMAKE_${lang}03_STANDARD_COMPILE_OPTION "")
set(CMAKE_${lang}03_EXTENSION_COMPILE_OPTION "")
@@ -75,6 +83,12 @@ macro(__compiler_nvidia_cxx_standards lang)
set(CMAKE_${lang}20_EXTENSION_COMPILE_OPTION "-std=c++20")
set(CMAKE_${lang}_STANDARD_LATEST 20)
endif()
+
+ if (NOT CMAKE_${lang}_COMPILER_VERSION VERSION_LESS 13.3)
+ set(CMAKE_${lang}23_STANDARD_COMPILE_OPTION "-std=c++23")
+ set(CMAKE_${lang}23_EXTENSION_COMPILE_OPTION "-std=c++23")
+ set(CMAKE_${lang}_STANDARD_LATEST 23)
+ endif()
endif()

__compiler_check_default_language_standard(${lang} 6.0 03)
37 changes: 37 additions & 0 deletions patches/cuda_cccl_8771.patch
Original file line number Diff line number Diff line change
@@ -0,0 +1,37 @@
diff --git a/cub/cub/device/device_transform.cuh b/cub/cub/device/device_transform.cuh
index 1b21230619d..64d26c81be8 100644
--- a/cub/cub/device/device_transform.cuh
+++ b/cub/cub/device/device_transform.cuh
@@ -42,9 +42,12 @@ struct __return_constant
} // namespace detail
CUB_NAMESPACE_END

+namespace cuda
+{
template <typename T>
-struct ::cuda::proclaims_copyable_arguments<CUB_NS_QUALIFIER::detail::__return_constant<T>> : ::cuda::std::true_type
+struct proclaims_copyable_arguments<CUB_NS_QUALIFIER::detail::__return_constant<T>> : ::cuda::std::true_type
{};
+} // namespace cuda

CUB_NAMESPACE_BEGIN
//! DeviceTransform provides device-wide, parallel operations for transforming elements tuple-wise from multiple input
diff --git a/cub/cub/device/dispatch/tuning/tuning_transform.cuh b/cub/cub/device/dispatch/tuning/tuning_transform.cuh
index 499ac333072..97073b360fb 100644
--- a/cub/cub/device/dispatch/tuning/tuning_transform.cuh
+++ b/cub/cub/device/dispatch/tuning/tuning_transform.cuh
@@ -42,10 +42,12 @@ struct always_true_predicate
} // namespace detail::transform
CUB_NAMESPACE_END

+namespace cuda
+{
template <>
-struct ::cuda::proclaims_copyable_arguments<CUB_NS_QUALIFIER::detail::transform::always_true_predicate>
- : ::cuda::std::true_type
+struct proclaims_copyable_arguments<CUB_NS_QUALIFIER::detail::transform::always_true_predicate> : ::cuda::std::true_type
{};
+} // namespace cuda

CUB_NAMESPACE_BEGIN
namespace detail::transform
38 changes: 38 additions & 0 deletions patches/pytorch-typename-decltype.patch
Original file line number Diff line number Diff line change
@@ -0,0 +1,38 @@
commit c2671b853e1553b258a30c15d66f9647faae7aee
Author: Tom Fogal <60981+tfogal@users.noreply.github.com>
Date: Mon May 11 15:55:08 2026 +0000

Couple naming updates (#180996)

The template keyword in ivalue_inl.h is technically required for a correct parse; better to fix it now, before compilers start getting strict about it.

The 'List' commit is just a consistency with surrounding code change.
Pull Request resolved: https://github.com/pytorch/pytorch/pull/180996
Approved by: https://github.com/eqy

diff --git a/aten/src/ATen/core/List_inl.h b/aten/src/ATen/core/List_inl.h
index 96f78faea22..c32e77f3d8e 100644
--- a/aten/src/ATen/core/List_inl.h
+++ b/aten/src/ATen/core/List_inl.h
@@ -198,7 +198,7 @@ typename List<T>::internal_const_reference_type List<T>::operator[](size_type po
template<class T>
typename List<T>::internal_reference_type List<T>::operator[](size_type pos) {
static_cast<void>(impl_->list.at(pos)); // Throw the exception if it is out of range.
- return {impl_->list.begin() + static_cast<typename decltype(impl_->list)::difference_type>(pos)};
+ return {impl_->list.begin() + static_cast<typename c10::detail::ListImpl::list_type::difference_type>(pos)};
}

template<class T>
diff --git a/aten/src/ATen/core/ivalue_inl.h b/aten/src/ATen/core/ivalue_inl.h
index 1d17fe06906..0e09cd8bc08 100644
--- a/aten/src/ATen/core/ivalue_inl.h
+++ b/aten/src/ATen/core/ivalue_inl.h
@@ -1941,7 +1941,7 @@ Tuple generic_to_tuple_impl(
const ivalue::TupleElements& t,
std::index_sequence<INDEX...> /*unused*/) {
return std::make_tuple(
- t[INDEX].to<std::tuple_element_t<INDEX, Tuple>>()...);
+ t[INDEX].template to<std::tuple_element_t<INDEX, Tuple>>()...);
}
} // namespace detail

1 change: 1 addition & 0 deletions pip/torch.file
Original file line number Diff line number Diff line change
Expand Up @@ -7,6 +7,7 @@ Source97: pip/files/torch-build-env
Source98: FindEigen3.cmake
Source99: scram-tools.file/tools/eigen/env
Patch0: patches/pytorch-system-fmt
Patch1: patches/pytorch-typename-decltype

BuildRequires: cmake ninja python-python3 py3-packaging
Requires: fftw3 eigen fxdiv numactl openmpi protobuf psimd python3 py3-PyYAML
Expand Down