From a9d9b083e425f102b1e1fbc9ed4d715afc392aab Mon Sep 17 00:00:00 2001 From: Ye Wang <52801275+wangyems@users.noreply.github.com> Date: Wed, 27 Mar 2024 15:59:35 -0700 Subject: [PATCH] Fix py package pipeline (#20065) ### Description ### Motivation and Context Fixes #20068 --- onnxruntime/contrib_ops/cuda/moe/ft_moe/moe_kernel.cu | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/onnxruntime/contrib_ops/cuda/moe/ft_moe/moe_kernel.cu b/onnxruntime/contrib_ops/cuda/moe/ft_moe/moe_kernel.cu index 5e6e484567988..c299cdcfe6a3d 100644 --- a/onnxruntime/contrib_ops/cuda/moe/ft_moe/moe_kernel.cu +++ b/onnxruntime/contrib_ops/cuda/moe/ft_moe/moe_kernel.cu @@ -656,7 +656,9 @@ inline __device__ float4 operator*(const float4 a, const float4 b) { return make_float4(a.x * b.x, a.y * b.y, a.z * b.z, a.w * b.w); } -#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 530 +// TODO(wy): use cuda common header and investigate pipeline build issue. +#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 530 && \ + ((__CUDACC_VER_MAJOR__ < 12) || ((__CUDACC_VER_MAJOR__ == 12) && (__CUDACC_VER_MINOR__ < 2))) inline __device__ half operator*(const half a, const half b) { return __float2half(__half2float(a) * __half2float(b)); } @@ -666,8 +668,10 @@ inline __device__ half2 operator*(const half2 a, const half2 b) { } #endif +// TODO(wy): use cuda common header and investigate pipeline build issue. inline __device__ Half4 operator*(const Half4 a, const Half4 b) { -#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 530 +#if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 530 && \ + ((__CUDACC_VER_MAJOR__ < 12) || ((__CUDACC_VER_MAJOR__ == 12) && (__CUDACC_VER_MINOR__ < 2))) Half4 result; result.x = a.x * b.x; result.y = a.y * b.y;