Skip to content

Commit cf446a4

Browse files
fwyzardfwyzard
committed
ONNXRuntime: CUDA 13 fixes
Backport microsoft/onnxruntime#28736 Patch Abseil to workaround NVIDIA bug #6302392
1 parent 9b5172e commit cf446a4

3 files changed

Lines changed: 97 additions & 0 deletions

File tree

‎onnxruntime/absl_nvidia_bug_6302392.patch‎

Lines changed: 61 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,61 @@
1+
diff --git a/absl/container/internal/common.h b/absl/container/internal/common.h
2+
index 5ef6c569..86de28c6 100644
3+
--- a/absl/container/internal/common.h
4+
+++ b/absl/container/internal/common.h
5+
@@ -68,6 +68,13 @@ struct IfRRef<T&&> {
6+
using AddPtr = Other*;
7+
};
8+
9+
+// Workaround for NVCC (cudafe++) which fails to parse
10+
+// `IfRRef<concrete_type>::AddPtr<Dependent>` when used as a template argument
11+
+// inside a heavily macro-expanded template parameter list. Top-level alias
12+
+// defers the member-template lookup outside the surrounding template-id.
13+
+template <class T, class Other>
14+
+using IfRRefAddPtr = typename IfRRef<T>::template AddPtr<Other>;
15+
+
16+
template <class, class = void>
17+
struct IsTransparent : std::false_type {};
18+
template <class T>
19+
diff --git a/absl/container/internal/raw_hash_map.h b/absl/container/internal/raw_hash_map.h
20+
index b42a4f22..540788b7 100644
21+
--- a/absl/container/internal/raw_hash_map.h
22+
+++ b/absl/container/internal/raw_hash_map.h
23+
@@ -106,8 +106,8 @@ class raw_hash_map : public raw_hash_set<Policy, Hash, Eq, Alloc> {
24+
typename K = key_type, class V = mapped_type, \
25+
ABSL_INTERNAL_IF_##KValue##_NOR_##VValue( \
26+
int = (EnableIf<LifetimeBoundKV<K, KValue, V, VValue, \
27+
- IfRRef<int KQual>::AddPtr<K>, \
28+
- IfRRef<int VQual>::AddPtr<V>>>()), \
29+
+ IfRRefAddPtr<int KQual, K>, \
30+
+ IfRRefAddPtr<int VQual, V>>>()), \
31+
ABSL_INTERNAL_SINGLE_ARG( \
32+
int &..., \
33+
decltype(EnableIf<LifetimeBoundKV<K, KValue, V, VValue>>()) = \
34+
diff --git a/absl/container/internal/btree_container.h b/absl/container/internal/btree_container.h
35+
index 21f00ae4..0a017f9f 100644
36+
--- a/absl/container/internal/btree_container.h
37+
+++ b/absl/container/internal/btree_container.h
38+
@@ -497,8 +497,8 @@ class btree_map_container : public btree_set_container<Tree> {
39+
typename K = key_type, class M, \
40+
ABSL_INTERNAL_IF_##KValue##_NOR_##MValue( \
41+
int = (EnableIf<LifetimeBoundKV<K, KValue, M, MValue, \
42+
- IfRRef<int KQual>::AddPtr<K>, \
43+
- IfRRef<int MQual>::AddPtr<M>>>()), \
44+
+ IfRRefAddPtr<int KQual, K>, \
45+
+ IfRRefAddPtr<int MQual, M>>>()), \
46+
ABSL_INTERNAL_SINGLE_ARG( \
47+
int &..., \
48+
decltype(EnableIf<LifetimeBoundKV<K, KValue, M, MValue>>()) = \
49+
@@ -598,10 +598,10 @@ class btree_map_container : public btree_set_container<Tree> {
50+
ABSL_INTERNAL_IF_##KValue( \
51+
class... Args, \
52+
int = (EnableIf< \
53+
- LifetimeBoundK<K, KValue, IfRRef<int KQual>::AddPtr<K>>>())), \
54+
+ LifetimeBoundK<K, KValue, IfRRefAddPtr<int KQual, K>>>())), \
55+
ABSL_INTERNAL_IF_##KValue( \
56+
decltype(EnableIf<LifetimeBoundK< \
57+
- K, KValue, IfRRef<int KQual>::AddPtr<K>>>()) = 0, \
58+
+ K, KValue, IfRRefAddPtr<int KQual, K>>>()) = 0, \
59+
class... Args), \
60+
std::enable_if_t<!std::is_convertible<K, const_iterator>::value, int> = \
61+
0> \

‎onnxruntime/backport28736.patch‎

Lines changed: 29 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,29 @@
1+
diff --git a/onnxruntime/contrib_ops/cuda/quantization/matmul_4bits.cu b/onnxruntime/contrib_ops/cuda/quantization/matmul_4bits.cu
2+
index cb9eab579a..ac93777a58 100644
3+
--- a/onnxruntime/contrib_ops/cuda/quantization/matmul_4bits.cu
4+
+++ b/onnxruntime/contrib_ops/cuda/quantization/matmul_4bits.cu
5+
@@ -281,8 +281,9 @@ __global__ void __launch_bounds__(kWarpSize* kColsPerThreadBlock) MatMulFloatInt
6+
#define UnRollReduction(unroll_size) \
7+
do { \
8+
constexpr int kUnroll = unroll_size; \
9+
- constexpr int kUnrollMask = 0xffffffff & (~(kUnroll * k_per_iter - 1)); \
10+
- for (; k_id < (k & kUnrollMask); k_id += kUnroll * k_per_iter) { \
11+
+ constexpr int kUnrollStep = kUnroll * k_per_iter; \
12+
+ const int k_unroll_bound = k - k % kUnrollStep; \
13+
+ for (; k_id < k_unroll_bound; k_id += kUnrollStep) { \
14+
_Pragma("unroll") for (int i = 0; i < kUnroll; i++) { \
15+
uint32_t value = *(reinterpret_cast<const uint32_t*>(b_data_quant + k_per_iter / 2 * i)); \
16+
T scale = b_scale_vec[t_meta_k + k_per_iter / block_size * i]; \
17+
diff --git a/onnxruntime/core/providers/cuda/tensor/concat_impl.cu b/onnxruntime/core/providers/cuda/tensor/concat_impl.cu
18+
index f369ead0f9..8b27b0dd93 100644
19+
--- a/onnxruntime/core/providers/cuda/tensor/concat_impl.cu
20+
+++ b/onnxruntime/core/providers/cuda/tensor/concat_impl.cu
21+
@@ -80,7 +80,7 @@ Status ConcatSameConcatDimImpl(cudaStream_t stream, const size_t element_bytes,
22+
template Status ConcatSameConcatDimImpl<const void**>(cudaStream_t stream, const size_t element_bytes,
23+
const int block_size_including_axis_dim,
24+
const int block_size_inside_axis_dim, const int64_t concat_size,
25+
- void* output_data, const void** input_data,
26+
+ void* output_data, const void** const input_data,
27+
const size_t output_size);
28+
29+
// input tensor addresses passed by value

‎onnxruntime/spec‎

Lines changed: 7 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -6,6 +6,9 @@
66
%define tag v%{realversion}
77
Source: git+https://github.com/%{github_user}/%{n}.git?obj=%{branch}/%{tag}&export=%{n}-%{realversion}&submodules=1&output=/%{n}-%{realversion}.tgz
88
Patch0: onnxruntime/cms-changes
9+
# Note: these two patches should likely be removed for onnxruntime 1.27+
10+
Patch1: onnxruntime/backport28736
11+
Patch2: onnxruntime/absl_nvidia_bug_6302392
912

1013
BuildRequires: cmake ninja
1114
Requires: protobuf py3-numpy py3-onnx zlib libpng py3-pybind11 re2 eigen
@@ -14,6 +17,7 @@ Requires: protobuf py3-numpy py3-onnx zlib libpng py3-pybind11 re2 eigen
1417
%prep
1518
%setup -q -n %{n}-%{realversion}
1619
%patch0 -p1
20+
%patch1 -p1
1721

1822
%build
1923
rm -rf ../build; mkdir ../build; cd ../build
@@ -59,6 +63,9 @@ cmake ../%{n}-%{realversion}/cmake \
5963
# -Wno-error=overloaded-virtual
6064
# See https://github.com/intel/neural-speed/issues/188
6165

66+
# patch Abseil bundled with ONNX to workaround NVIDIA bug #6302392
67+
patch -d %{_builddir}/build/_deps/abseil_cpp-src -p1 < %{PATCH2}
68+
6269
ninja -v %{makeprocesses}
6370
python3 ../%{n}-%{realversion}/setup.py build
6471

0 commit comments

Comments
 (0)