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 5e6e484567..c299cdcfe6 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;