saymrwulf/onnxruntime
1
0
Fork 0
mirror of https://github.com/saymrwulf/onnxruntime.git synced 2026-05-29 23:06:41 +00:00
Code Issues Projects Releases Packages Wiki Activity Actions

fix jetson build break caused by https://github.com/microsoft/onnxruntime/pull/12949 (#13310)

Browse source 
https://github.com/microsoft/onnxruntime/pull/12949 breaks Jetson/ARM
builds due to
invalid narrowing conversion from "char" to "signed char". (on ARM, char
is unsigned)
This was reported by
https://github.com/microsoft/onnxruntime/issues/13285#issuecomment-1276122505
This commit is contained in:
George Wu 2022-10-13 21:31:36 -07:00 committed by GitHub
parent dc324b1d90
commit 074e009e69
No known key found for this signature in database
GPG key ID: 4AEE18F83AFDEB23
2 changed files with 5 additions and 5 deletions
Show stats Download patch file Download diff file Expand all files Collapse all files

6
onnxruntime/contrib_ops/cuda/quantization/qordered_ops/qordered_common.cuh
View file

@ -74,7 +74,7 @@ __device__ inline char2 QuantizeHalf2Char2(const __half2 xy, const __half2 inver
s2xy.s2.x = __half2short_rn(scaled_xy.x);
s2xy.s2.y = __half2short_rn(scaled_xy.y);
s2xy.u1 = __vmaxs2(__vmins2(s2xy.u1, 0x007F007F), 0xFF80FF80);
return char2{(char)s2xy.s2.x, (char)s2xy.s2.y};
return char2{(signed char)s2xy.s2.x, (signed char)s2xy.s2.y};
}
__device__ inline char4 QuantizeHalf4Char4(const __half4 val4, const __half2 inverse_scale2) {
@ -87,7 +87,7 @@ __device__ inline char4 QuantizeHalf4Char4(const __half4 val4, const __half2 inv
shortzw.s2.y = __half2short_rn(__high2half(val4_zw));
shortxy.u1 = __vmaxs2(__vmins2(shortxy.u1, 0x007F007F), 0xFF80FF80);
shortzw.u1 = __vmaxs2(__vmins2(shortzw.u1, 0x007F007F), 0xFF80FF80);
return char4{(char)shortxy.s2.x, (char)shortxy.s2.y, (char)shortzw.s2.x, (char)shortzw.s2.y};
return char4{(signed char)shortxy.s2.x, (signed char)shortxy.s2.y, (signed char)shortzw.s2.x, (signed char)shortzw.s2.y};
}
__device__ inline char4 QuantizeHalf4Char4Strict(const __half4 val4, const float inverse_scale) {
@ -98,7 +98,7 @@ __device__ inline char4 QuantizeHalf4Char4Strict(const __half4 val4, const float
shortzw.s2.y = static_cast<short>(__float2int_rn(__half2float(val4.zw.y) * inverse_scale));
shortxy.u1 = __vmaxs2(__vmins2(shortxy.u1, 0x007F007F), 0xFF80FF80);
shortzw.u1 = __vmaxs2(__vmins2(shortzw.u1, 0x007F007F), 0xFF80FF80);
return char4{(char)shortxy.s2.x, (char)shortxy.s2.y, (char)shortzw.s2.x, (char)shortzw.s2.y};
return char4{(signed char)shortxy.s2.x, (signed char)shortxy.s2.y, (signed char)shortzw.s2.x, (signed char)shortzw.s2.y};
}
__device__ inline __half4 DeqantizeChar4Half4(const char4 ch4, const __half2 scale2) {

4
onnxruntime/contrib_ops/cuda/quantization/qordered_ops/qordered_qdq_impl.cu
View file

@ -42,7 +42,7 @@ struct DequantizeVec<float2> {
static __device__ inline QuantizedVecT Quantize(const float2 fpvals, const float inv_scale) {
float dqvalx = fmaxf(fminf(127.0f, fpvals.x * inv_scale), -128.0f);
float dqvaly = fmaxf(fminf(127.0f, fpvals.y * inv_scale), -128.0f);
return char2{static_cast<char>(__float2int_rn(dqvalx)), static_cast<char>(__float2int_rn(dqvaly))};
return char2{static_cast<signed char>(__float2int_rn(dqvalx)), static_cast<signed char>(__float2int_rn(dqvaly))};
}
static __device__ inline float2 Dequantize(const QuantizedVecT qvals, const float scale) {
@ -76,7 +76,7 @@ struct DequantizeVec<__half2> {
s2xy.s2.x = __half2short_rn(xy.x);
s2xy.s2.y = __half2short_rn(xy.y);
s2xy.u1 = __vmaxs2(__vmins2(s2xy.u1, 0x007F007F), 0xFF80FF80);
return char2{(char)s2xy.s2.x, (char)s2xy.s2.y};
return char2{(signed char)s2xy.s2.x, (signed char)s2xy.s2.y};
}
static __device__ inline __half2 Dequantize(const QuantizedVecT qvals, const __half scale) {