diff --git a/hipamd/include/hip/amd_detail/amd_hip_bf16.h b/hipamd/include/hip/amd_detail/amd_hip_bf16.h index 0a5d1497fd..c1b2405cd8 100644 --- a/hipamd/include/hip/amd_detail/amd_hip_bf16.h +++ b/hipamd/include/hip/amd_detail/amd_hip_bf16.h @@ -1835,7 +1835,7 @@ __BF16_DEVICE_STATIC__ __hip_bfloat162 h2trunc(const __hip_bfloat162 h) { */ __BF16_DEVICE_STATIC__ __hip_bfloat162 unsafeAtomicAdd(__hip_bfloat162* address, __hip_bfloat162 value) { -#if __has_builtin(__builtin_amdgcn_flat_atomic_fadd_v2bf16) +#if __has_builtin(__builtin_amdgcn_flat_atomic_fadd_v2bf16) && !defined(__SPIRV__) typedef short __attribute__((ext_vector_type(2))) vec_short2; static_assert(sizeof(vec_short2) == sizeof(__hip_bfloat162_raw)); union { diff --git a/hipamd/include/hip/amd_detail/amd_hip_fp16.h b/hipamd/include/hip/amd_detail/amd_hip_fp16.h index 1a08bb8df1..b7b0d8a625 100644 --- a/hipamd/include/hip/amd_detail/amd_hip_fp16.h +++ b/hipamd/include/hip/amd_detail/amd_hip_fp16.h @@ -1530,7 +1530,7 @@ THE SOFTWARE. // Atomic #if defined(__clang__) && defined(__HIP__) inline __device__ __half2 unsafeAtomicAdd(__half2* address, __half2 value) { - #if __has_builtin(__builtin_amdgcn_flat_atomic_fadd_v2f16) + #if __has_builtin(__builtin_amdgcn_flat_atomic_fadd_v2f16) && !defined(__SPIRV__) // The api expects an ext_vector_type of half typedef _Float16 __attribute__((ext_vector_type(2))) vec_fp162; static_assert(sizeof(vec_fp162) == sizeof(__half2_raw));