Fix check

Author: Hugh Delaney

libclc/amdgcn-amdhsa/libspirv/atomic/atomic_add.cl

@@ -32,11 +32,15 @@ AMDGPU_ATOMIC(_Z18__spirv_AtomicIAdd, unsigned long, m, __hip_atomic_fetch_add)
     return *(float *)&ret;                                                                                 \
   }
 
-AMDGPU_ATOMIC_FP32_ADD_IMPL(global, U3AS1, 1, __oclc_ISA_version >= 9010,
+AMDGPU_ATOMIC_FP32_ADD_IMPL(global, U3AS1, 1,
+                            __oclc_ISA_version >= 9010 &&
+                                __oclc_ISA_version < 10000,
                             __builtin_amdgcn_global_atomic_fadd_f32)
 AMDGPU_ATOMIC_FP32_ADD_IMPL(local, U3AS3, 1, __oclc_ISA_version >= 8000,
                             __builtin_amdgcn_ds_atomic_fadd_f32)
-AMDGPU_ATOMIC_FP32_ADD_IMPL(, , 0, __oclc_ISA_version >= 9400,
+AMDGPU_ATOMIC_FP32_ADD_IMPL(, , 0,
+                            __oclc_ISA_version >= 9400 &&
+                                __oclc_ISA_version < 10000,
                             __builtin_amdgcn_flat_atomic_fadd_f32)
 
 #define AMDGPU_ATOMIC_FP64_ADD_IMPL(AS, AS_MANGLED, SUB1, SUB2, CHECK,                                                          \
@@ -74,11 +78,17 @@ AMDGPU_ATOMIC_FP32_ADD_IMPL(, , 0, __oclc_ISA_version >= 9400,
   }
 
 #ifdef cl_khr_int64_base_atomics
-AMDGPU_ATOMIC_FP64_ADD_IMPL(global, U3AS1, 1, 5, __oclc_ISA_version >= 9010,
+AMDGPU_ATOMIC_FP64_ADD_IMPL(global, U3AS1, 1, 5,
+                            __oclc_ISA_version >= 9010 &&
+                                __oclc_ISA_version < 10000,
                             __builtin_amdgcn_global_atomic_fadd_f64)
-AMDGPU_ATOMIC_FP64_ADD_IMPL(local, U3AS3, 1, 5, __oclc_ISA_version >= 9010,
+AMDGPU_ATOMIC_FP64_ADD_IMPL(local, U3AS3, 1, 5,
+                            __oclc_ISA_version >= 9010 &&
+                                __oclc_ISA_version < 10000,
                             __builtin_amdgcn_ds_atomic_fadd_f64)
-AMDGPU_ATOMIC_FP64_ADD_IMPL(, , 0, 4, __oclc_ISA_version >= 9400,
+AMDGPU_ATOMIC_FP64_ADD_IMPL(, , 0, 4,
+                            __oclc_ISA_version >= 9400 &&
+                                __oclc_ISA_version < 10000,
                             __builtin_amdgcn_flat_atomic_fadd_f64)
 #endif
 

libclc/amdgcn-amdhsa/libspirv/atomic/atomic_max.cl

@@ -26,12 +26,14 @@ AMDGPU_ATOMIC_FP_MINMAX_IMPL(Max, >, float, f, int, i, , , 0, 4_ii, false, )
 
 #ifdef cl_khr_int64_base_atomics
 AMDGPU_ATOMIC_FP_MINMAX_IMPL(Max, >, double, d, long, l, global, U3AS1, 1, 5_ll,
-                             __oclc_ISA_version >= 9010,
+                             __oclc_ISA_version >= 9010 &&
+                                 __oclc_ISA_version < 10000,
                              __builtin_amdgcn_global_atomic_fmax_f64)
 AMDGPU_ATOMIC_FP_MINMAX_IMPL(Max, >, double, d, long, l, local, U3AS3, 1, 5_ll,
                              false, )
 AMDGPU_ATOMIC_FP_MINMAX_IMPL(Max, >, double, d, long, l, , , 0, 4_ll,
-                             __oclc_ISA_version >= 9010,
+                             __oclc_ISA_version >= 9010 &&
+                                 __oclc_ISA_version < 10000,
                              __builtin_amdgcn_flat_atomic_fmax_f64)
 #endif
 

libclc/amdgcn-amdhsa/libspirv/atomic/atomic_min.cl

@@ -26,12 +26,14 @@ AMDGPU_ATOMIC_FP_MINMAX_IMPL(Min, <, float, f, int, i, , , 0, 4_ii, false, )
 
 #ifdef cl_khr_int64_base_atomics
 AMDGPU_ATOMIC_FP_MINMAX_IMPL(Min, <, double, d, long, l, global, U3AS1, 1, 5_ll,
-                             __oclc_ISA_version >= 9010,
+                             __oclc_ISA_version >= 9010 &&
+                                 __oclc_ISA_version < 10000,
                              __builtin_amdgcn_global_atomic_fmin_f64)
 AMDGPU_ATOMIC_FP_MINMAX_IMPL(Min, <, double, d, long, l, local, U3AS3, 1, 5_ll,
                              false, )
 AMDGPU_ATOMIC_FP_MINMAX_IMPL(Min, <, double, d, long, l, , , 0, 4_ll,
-                             __oclc_ISA_version >= 9010,
+                             __oclc_ISA_version >= 9010 &&
+                                 __oclc_ISA_version < 10000,
                              __builtin_amdgcn_flat_atomic_fmin_f64)
 #endif