diff --git a/patches/clang/0005-OpenCL-support-cl_ext_float_atomics.patch b/patches/clang/0005-OpenCL-support-cl_ext_float_atomics.patch index 0d358520..9957f162 100644 --- a/patches/clang/0005-OpenCL-support-cl_ext_float_atomics.patch +++ b/patches/clang/0005-OpenCL-support-cl_ext_float_atomics.patch @@ -1,15 +1,14 @@ -From 9b48f70bae77fdc752ee5e98949a7ed2c9373037 Mon Sep 17 00:00:00 2001 +From 6ef0d9afd03c80671393f4d749ddbeb08f7291fe Mon Sep 17 00:00:00 2001 From: haonanya Date: Fri, 13 Aug 2021 10:00:02 +0800 Subject: [PATCH] [OpenCL] support cl_ext_float_atomics Signed-off-by: haonanya -Signed-off-by: Haonan Yang --- clang/lib/Headers/opencl-c-base.h | 22 +++ - clang/lib/Headers/opencl-c.h | 208 ++++++++++++++++++++++++++ - clang/test/Headers/opencl-c-header.cl | 96 ++++++++++++ - 3 files changed, 326 insertions(+) + clang/lib/Headers/opencl-c.h | 232 ++++++++++++++++++++++++++ + clang/test/Headers/opencl-c-header.cl | 96 +++++++++++ + 3 files changed, 350 insertions(+) diff --git a/clang/lib/Headers/opencl-c-base.h b/clang/lib/Headers/opencl-c-base.h index 2cc688ccc3da..18d367de68ec 100644 @@ -45,10 +44,10 @@ index 2cc688ccc3da..18d367de68ec 100644 #ifndef __opencl_c_int64 #define __opencl_c_int64 1 diff --git a/clang/lib/Headers/opencl-c.h b/clang/lib/Headers/opencl-c.h -index d8173f0aa843..50515ac17a0c 100644 +index d8173f0aa843..454469991d59 100644 --- a/clang/lib/Headers/opencl-c.h +++ b/clang/lib/Headers/opencl-c.h -@@ -14354,6 +14354,214 @@ intptr_t __ovld atomic_fetch_max_explicit( +@@ -14354,6 +14354,238 @@ intptr_t __ovld atomic_fetch_max_explicit( // defined(cl_khr_int64_extended_atomics) #endif // (__OPENCL_C_VERSION__ >= CL_VERSION_3_0) @@ -102,10 +101,12 @@ index d8173f0aa843..50515ac17a0c 100644 +float __ovld atomic_fetch_max_explicit(volatile atomic_float *object, + float operand, memory_order order, + memory_scope scope); -+#endif // defined(__opencl_c_ext_fp32_global_atomic_min_max) && \ -+ defined(__opencl_c_ext_fp32_local_atomic_min_max) ++#endif // defined(__opencl_c_ext_fp32_global_atomic_min_max) && ++ // defined(__opencl_c_ext_fp32_local_atomic_min_max) + -+#if defined(__opencl_c_ext_fp64_global_atomic_min_max) ++#if defined(__opencl_c_ext_fp64_global_atomic_min_max) && \ ++ defined(cl_khr_int64_base_atomics) && \ ++ defined(cl_khr_int64_extended_atomics) +double __ovld atomic_fetch_min(volatile __global atomic_double *object, + double operand); +double __ovld atomic_fetch_max(volatile __global atomic_double *object, @@ -120,9 +121,13 @@ index d8173f0aa843..50515ac17a0c 100644 +double __ovld atomic_fetch_max_explicit(volatile __global atomic_double *object, + double operand, memory_order order, + memory_scope scope); -+#endif // defined(__opencl_c_ext_fp64_global_atomic_min_max) ++#endif // defined(__opencl_c_ext_fp64_global_atomic_min_max) && ++ // defined(cl_khr_int64_base_atomics) && ++ // defined(cl_khr_int64_extended_atomics) + -+#if defined(__opencl_c_ext_fp64_local_atomic_min_max) ++#if defined(__opencl_c_ext_fp64_local_atomic_min_max) && \ ++ defined(cl_khr_int64_base_atomics) && \ ++ defined(cl_khr_int64_extended_atomics) +double __ovld atomic_fetch_min(volatile __local atomic_double *object, + double operand); +double __ovld atomic_fetch_max(volatile __local atomic_double *object, @@ -137,10 +142,14 @@ index d8173f0aa843..50515ac17a0c 100644 +double __ovld atomic_fetch_max_explicit(volatile __local atomic_double *object, + double operand, memory_order order, + memory_scope scope); -+#endif // defined(__opencl_c_ext_fp64_local_atomic_min_max) ++#endif // defined(__opencl_c_ext_fp64_local_atomic_min_max) && ++ // defined(cl_khr_int64_base_atomics) && ++ // defined(cl_khr_int64_extended_atomics) + +#if defined(__opencl_c_ext_fp64_global_atomic_min_max) && \ -+ defined(__opencl_c_ext_fp64_local_atomic_min_max) ++ defined(__opencl_c_ext_fp64_local_atomic_min_max) && \ ++ defined(cl_khr_int64_base_atomics) && \ ++ defined(cl_khr_int64_extended_atomics) +double __ovld atomic_fetch_min(volatile atomic_double *object, double operand); +double __ovld atomic_fetch_max(volatile atomic_double *object, double operand); +double __ovld atomic_fetch_min_explicit(volatile atomic_double *object, @@ -153,8 +162,10 @@ index d8173f0aa843..50515ac17a0c 100644 +double __ovld atomic_fetch_max_explicit(volatile atomic_double *object, + double operand, memory_order order, + memory_scope scope); -+#endif // defined(__opencl_c_ext_fp64_global_atomic_min_max) && \ -+ defined(__opencl_c_ext_fp64_local_atomic_min_max) ++#endif // defined(__opencl_c_ext_fp64_global_atomic_min_max) && ++ // defined(__opencl_c_ext_fp64_local_atomic_min_max) && ++ // defined(cl_khr_int64_base_atomics) && ++ // defined(cl_khr_int64_extended_atomics) + +#if defined(__opencl_c_ext_fp32_global_atomic_add) +float __ovld atomic_fetch_add(volatile __global atomic_float *object, @@ -204,10 +215,12 @@ index d8173f0aa843..50515ac17a0c 100644 +float __ovld atomic_fetch_sub_explicit(volatile atomic_float *object, + float operand, memory_order order, + memory_scope scope); -+#endif // defined(__opencl_c_ext_fp32_global_atomic_add) && \ -+ defined(__opencl_c_ext_fp32_local_atomic_add) ++#endif // defined(__opencl_c_ext_fp32_global_atomic_add) && ++ // defined(__opencl_c_ext_fp32_local_atomic_add) + -+#if defined(__opencl_c_ext_fp64_global_atomic_add) ++#if defined(__opencl_c_ext_fp64_global_atomic_add) && \ ++ defined(cl_khr_int64_base_atomics) && \ ++ defined(cl_khr_int64_extended_atomics) +double __ovld atomic_fetch_add(volatile __global atomic_double *object, + double operand); +double __ovld atomic_fetch_sub(volatile __global atomic_double *object, @@ -222,9 +235,13 @@ index d8173f0aa843..50515ac17a0c 100644 +double __ovld atomic_fetch_sub_explicit(volatile __global atomic_double *object, + double operand, memory_order order, + memory_scope scope); -+#endif // defined(__opencl_c_ext_fp64_global_atomic_add) ++#endif // defined(__opencl_c_ext_fp64_global_atomic_add) && ++ // defined(cl_khr_int64_base_atomics) && ++ // defined(cl_khr_int64_extended_atomics) + -+#if defined(__opencl_c_ext_fp64_local_atomic_add) ++#if defined(__opencl_c_ext_fp64_local_atomic_add) && \ ++ defined(cl_khr_int64_base_atomics) && \ ++ defined(cl_khr_int64_extended_atomics) +double __ovld atomic_fetch_add(volatile __local atomic_double *object, + double operand); +double __ovld atomic_fetch_sub(volatile __local atomic_double *object, @@ -239,10 +256,14 @@ index d8173f0aa843..50515ac17a0c 100644 +double __ovld atomic_fetch_sub_explicit(volatile __local atomic_double *object, + double operand, memory_order order, + memory_scope scope); -+#endif // defined(__opencl_c_ext_fp64_local_atomic_add) ++#endif // defined(__opencl_c_ext_fp64_local_atomic_add) && ++ // defined(cl_khr_int64_base_atomics) && ++ // defined(cl_khr_int64_extended_atomics) + +#if defined(__opencl_c_ext_fp64_global_atomic_add) && \ -+ defined(__opencl_c_ext_fp64_local_atomic_add) ++ defined(__opencl_c_ext_fp64_local_atomic_add) && \ ++ defined(cl_khr_int64_base_atomics) && \ ++ defined(cl_khr_int64_extended_atomics) +double __ovld atomic_fetch_add(volatile atomic_double *object, double operand); +double __ovld atomic_fetch_sub(volatile atomic_double *object, double operand); +double __ovld atomic_fetch_add_explicit(volatile atomic_double *object, @@ -255,8 +276,10 @@ index d8173f0aa843..50515ac17a0c 100644 +double __ovld atomic_fetch_sub_explicit(volatile atomic_double *object, + double operand, memory_order order, + memory_scope scope); -+#endif // defined(__opencl_c_ext_fp64_global_atomic_add) && \ -+ defined(__opencl_c_ext_fp64_local_atomic_add) ++#endif // defined(__opencl_c_ext_fp64_global_atomic_add) && ++ // defined(__opencl_c_ext_fp64_local_atomic_add) && ++ // defined(cl_khr_int64_base_atomics) && ++ // defined(cl_khr_int64_extended_atomics) + +#endif // cl_ext_float_atomics + diff --git a/patches/spirv/0001-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch b/patches/spirv/0001-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch index 43d99be1..8e02f51f 100644 --- a/patches/spirv/0001-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch +++ b/patches/spirv/0001-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch @@ -1,31 +1,34 @@ -From 5b21454c542aea71a447afb5a652a713cf53b111 Mon Sep 17 00:00:00 2001 +From edc51a3da117861a7884162ac9646e4b1a9d6a0e Mon Sep 17 00:00:00 2001 From: haonanya Date: Mon, 19 Jul 2021 10:14:20 +0800 Subject: [PATCH] Add support for cl_ext_float_atomics in SPIRVWriter Signed-off-by: haonanya +Signed-off-by: Haonan Yang --- - lib/SPIRV/OCL20ToSPIRV.cpp | 25 +++++- - lib/SPIRV/OCLUtil.cpp | 4 - + lib/SPIRV/OCL20ToSPIRV.cpp | 26 +++++- + lib/SPIRV/OCLUtil.cpp | 19 +++-- lib/SPIRV/SPIRVToOCL.h | 3 + lib/SPIRV/SPIRVToOCL12.cpp | 21 +++++ lib/SPIRV/SPIRVToOCL20.cpp | 28 ++++++- lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h | 1 - lib/SPIRV/libSPIRV/SPIRVOpCode.h | 8 +- - test/AtomicFAddEXTForOCL.ll | 64 +++++++++++++++ + test/AtomicBuiltinsFloat.ll | 79 ++++++++++++++++++ + test/AtomicFAddEXTForOCL.ll | 88 ++++++++++++++++++++ test/AtomicFAddExt.ll | 111 ++++++++----------------- test/AtomicFMaxEXT.ll | 113 +++++++------------------- - test/AtomicFMaxEXTForOCL.ll | 64 +++++++++++++++ + test/AtomicFMaxEXTForOCL.ll | 88 ++++++++++++++++++++ test/AtomicFMinEXT.ll | 113 +++++++------------------- - test/AtomicFMinEXTForOCL.ll | 64 +++++++++++++++ - test/InvalidAtomicBuiltins.cl | 8 -- - 14 files changed, 366 insertions(+), 261 deletions(-) + test/AtomicFMinEXTForOCL.ll | 85 +++++++++++++++++++ + test/InvalidAtomicBuiltins.cl | 16 ---- + 15 files changed, 526 insertions(+), 273 deletions(-) + create mode 100644 test/AtomicBuiltinsFloat.ll create mode 100644 test/AtomicFAddEXTForOCL.ll create mode 100644 test/AtomicFMaxEXTForOCL.ll create mode 100644 test/AtomicFMinEXTForOCL.ll diff --git a/lib/SPIRV/OCL20ToSPIRV.cpp b/lib/SPIRV/OCL20ToSPIRV.cpp -index e30aa5be..faa5be2b 100644 +index e30aa5be..98d4289e 100644 --- a/lib/SPIRV/OCL20ToSPIRV.cpp +++ b/lib/SPIRV/OCL20ToSPIRV.cpp @@ -407,7 +407,6 @@ void OCL20ToSPIRV::visitCallInst(CallInst &CI) { @@ -45,7 +48,7 @@ index e30aa5be..faa5be2b 100644 Info.PostProc(Args); // Order of args in OCL20: // object, 0-2 other args, 1-2 order, scope -@@ -864,7 +863,27 @@ void OCL20ToSPIRV::transAtomicBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info) { +@@ -864,7 +863,28 @@ void OCL20ToSPIRV::transAtomicBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info) { std::rotate(Args.begin() + 2, Args.begin() + OrderIdx, Args.end() - Offset); } @@ -61,7 +64,8 @@ index e30aa5be..faa5be2b 100644 + if (!IsFPType(AtomicBuiltinsReturnType)) + return SPIRVFunctionName; + // Translate FP-typed atomic builtins. Currently we only need to -+ // translate atomic_fetch_[add, max, min]* to related float instructions ++ // translate atomic_fetch_[add, max, min] and atomic_fetch_[add, max, ++ // min]_explicit to related float instructions + auto SPIRFunctionNameForFloatAtomics = + llvm::StringSwitch(SPIRVFunctionName) + .Case("__spirv_AtomicIAdd", "__spirv_AtomicFAddEXT") @@ -75,10 +79,37 @@ index e30aa5be..faa5be2b 100644 &Attrs); } diff --git a/lib/SPIRV/OCLUtil.cpp b/lib/SPIRV/OCLUtil.cpp -index c7232623..9a4c8ab9 100644 +index c7232623..ecb97119 100644 --- a/lib/SPIRV/OCLUtil.cpp +++ b/lib/SPIRV/OCLUtil.cpp -@@ -136,13 +136,9 @@ bool isComputeAtomicOCLBuiltin(StringRef DemangledName) { +@@ -120,29 +120,32 @@ size_t getSPIRVAtomicBuiltinNumMemoryOrderArgs(Op OC) { + return 1; + } + ++// atomic_fetch_[add, sub, min, max] and atomic_fetch_[add, sub, min, ++// max]_explicit functions are defined on OpenCL headers, they are not ++// translated to function call + bool isComputeAtomicOCLBuiltin(StringRef DemangledName) { + if (!DemangledName.startswith(kOCLBuiltinName::AtomicPrefix) && + !DemangledName.startswith(kOCLBuiltinName::AtomPrefix)) + return false; + + return llvm::StringSwitch(DemangledName) +- .EndsWith("add", true) +- .EndsWith("sub", true) ++ .EndsWith("atomic_add", true) ++ .EndsWith("atomic_sub", true) ++ .EndsWith("atomic_min", true) ++ .EndsWith("atomic_max", true) ++ .EndsWith("atom_add", true) ++ .EndsWith("atom_sub", true) ++ .EndsWith("atom_min", true) ++ .EndsWith("atom_max", true) + .EndsWith("inc", true) + .EndsWith("dec", true) + .EndsWith("cmpxchg", true) +- .EndsWith("min", true) +- .EndsWith("max", true) .EndsWith("and", true) .EndsWith("or", true) .EndsWith("xor", true) @@ -245,12 +276,97 @@ index feec70f6..8e595e83 100644 } inline bool isBinaryOpCode(Op OpCode) { return ((unsigned)OpCode >= OpIAdd && (unsigned)OpCode <= OpFMod) || +diff --git a/test/AtomicBuiltinsFloat.ll b/test/AtomicBuiltinsFloat.ll +new file mode 100644 +index 00000000..18ee1c86 +--- /dev/null ++++ b/test/AtomicBuiltinsFloat.ll +@@ -0,0 +1,79 @@ ++; Check that translator generate atomic instructions for atomic builtins ++; RUN: llvm-as %s -o %t.bc ++; RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s ++; RUN: llvm-spirv %t.bc -o %t.spv ++; RUN: spirv-val %t.spv ++ ++; CHECK-LABEL: Label ++; CHECK: Store ++; CHECK-COUNT-3: AtomicStore ++; CHECK-COUNT-3: AtomicLoad ++; CHECK-COUNT-3: AtomicExchange ++ ++target datalayout = "e-p:32:32-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" ++target triple = "spir-unknown-unknown" ++ ++; Function Attrs: convergent nounwind ++define spir_kernel void @test_atomic_kernel(float addrspace(3)* %ff) local_unnamed_addr #0 !kernel_arg_addr_space !3 !kernel_arg_access_qual !4 !kernel_arg_type !5 !kernel_arg_base_type !6 !kernel_arg_type_qual !7 { ++entry: ++ %0 = addrspacecast float addrspace(3)* %ff to float addrspace(4)* ++ tail call spir_func void @_Z11atomic_initPU3AS4VU7_Atomicff(float addrspace(4)* %0, float 1.000000e+00) #2 ++ tail call spir_func void @_Z12atomic_storePU3AS4VU7_Atomicff(float addrspace(4)* %0, float 1.000000e+00) #2 ++ tail call spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicff12memory_order(float addrspace(4)* %0, float 1.000000e+00, i32 0) #2 ++ tail call spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicff12memory_order12memory_scope(float addrspace(4)* %0, float 1.000000e+00, i32 0, i32 1) #2 ++ %call = tail call spir_func float @_Z11atomic_loadPU3AS4VU7_Atomicf(float addrspace(4)* %0) #2 ++ %call1 = tail call spir_func float @_Z20atomic_load_explicitPU3AS4VU7_Atomicf12memory_order(float addrspace(4)* %0, i32 0) #2 ++ %call2 = tail call spir_func float @_Z20atomic_load_explicitPU3AS4VU7_Atomicf12memory_order12memory_scope(float addrspace(4)* %0, i32 0, i32 1) #2 ++ %call3 = tail call spir_func float @_Z15atomic_exchangePU3AS4VU7_Atomicff(float addrspace(4)* %0, float 1.000000e+00) #2 ++ %call4 = tail call spir_func float @_Z24atomic_exchange_explicitPU3AS4VU7_Atomicff12memory_order(float addrspace(4)* %0, float 1.000000e+00, i32 0) #2 ++ %call5 = tail call spir_func float @_Z24atomic_exchange_explicitPU3AS4VU7_Atomicff12memory_order12memory_scope(float addrspace(4)* %0, float 1.000000e+00, i32 0, i32 1) #2 ++ ret void ++} ++ ++; Function Attrs: convergent ++declare spir_func void @_Z11atomic_initPU3AS4VU7_Atomicff(float addrspace(4)*, float) local_unnamed_addr #1 ++ ++; Function Attrs: convergent ++declare spir_func void @_Z12atomic_storePU3AS4VU7_Atomicff(float addrspace(4)*, float) local_unnamed_addr #1 ++ ++; Function Attrs: convergent ++declare spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicff12memory_order(float addrspace(4)*, float, i32) local_unnamed_addr #1 ++ ++; Function Attrs: convergent ++declare spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicff12memory_order12memory_scope(float addrspace(4)*, float, i32, i32) local_unnamed_addr #1 ++ ++; Function Attrs: convergent ++declare spir_func float @_Z11atomic_loadPU3AS4VU7_Atomicf(float addrspace(4)*) local_unnamed_addr #1 ++ ++; Function Attrs: convergent ++declare spir_func float @_Z20atomic_load_explicitPU3AS4VU7_Atomicf12memory_order(float addrspace(4)*, i32) local_unnamed_addr #1 ++ ++; Function Attrs: convergent ++declare spir_func float @_Z20atomic_load_explicitPU3AS4VU7_Atomicf12memory_order12memory_scope(float addrspace(4)*, i32, i32) local_unnamed_addr #1 ++ ++; Function Attrs: convergent ++declare spir_func float @_Z15atomic_exchangePU3AS4VU7_Atomicff(float addrspace(4)*, float) local_unnamed_addr #1 ++ ++; Function Attrs: convergent ++declare spir_func float @_Z24atomic_exchange_explicitPU3AS4VU7_Atomicff12memory_order(float addrspace(4)*, float, i32) local_unnamed_addr #1 ++ ++; Function Attrs: convergent ++declare spir_func float @_Z24atomic_exchange_explicitPU3AS4VU7_Atomicff12memory_order12memory_scope(float addrspace(4)*, float, i32, i32) local_unnamed_addr #1 ++ ++attributes #0 = { convergent nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "uniform-work-group-size"="false" "unsafe-fp-math"="false" "use-soft-float"="false" } ++attributes #1 = { convergent "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } ++attributes #2 = { convergent nounwind } ++ ++!llvm.module.flags = !{!0} ++!opencl.ocl.version = !{!1} ++!opencl.spir.version = !{!1} ++!llvm.ident = !{!2} ++ ++!0 = !{i32 1, !"wchar_size", i32 4} ++!1 = !{i32 2, i32 0} ++!2 = !{!"clang version 10.0.1 (8560093eba963fba2edd47ca85404cdaff22f174)"} ++!3 = !{i32 3} ++!4 = !{!"none"} ++!5 = !{!"atomic_float*"} ++!6 = !{!"_Atomic(float)*"} ++!7 = !{!"volatile"} diff --git a/test/AtomicFAddEXTForOCL.ll b/test/AtomicFAddEXTForOCL.ll new file mode 100644 -index 00000000..fb146fb9 +index 00000000..cd6ea089 --- /dev/null +++ b/test/AtomicFAddEXTForOCL.ll -@@ -0,0 +1,64 @@ +@@ -0,0 +1,88 @@ +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_add -o %t.spv +; RUN: spirv-val %t.spv @@ -272,39 +388,63 @@ index 00000000..fb146fb9 +; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 +; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 + -+ -+; Function Attrs: convergent norecurse nounwind -+define dso_local spir_func void @test_atomic_float(float addrspace(1)* %a) local_unnamed_addr #0 { ++; Function Attrs: convergent nounwind ++define spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 { +entry: + ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]] ++ %call = tail call spir_func float @_Z16atomic_fetch_addPU3AS1VU7_Atomicff(float addrspace(1)* %a, float 0.000000e+00) #2 ++ ++ ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]] + ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}}) + ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+fiif]]({{.*}}) -+ %call = tail call spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2 ++ %call1 = tail call spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2 ++ ++ ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]] ++ %call2 = tail call spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order12memory_scope(float addrspace(1)* %a, float 0.000000e+00, i32 0, i32 1) #2 + ret void +} + +; Function Attrs: convergent ++declare spir_func float @_Z16atomic_fetch_addPU3AS1VU7_Atomicff(float addrspace(1)*, float) local_unnamed_addr #1 ++ ++; Function Attrs: convergent +declare spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1 +; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float + -+; Function Attrs: convergent norecurse nounwind -+define dso_local spir_func void @test_atomic_double(double addrspace(1)* %a) local_unnamed_addr #0 { ++; Function Attrs: convergent ++declare spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order12memory_scope(float addrspace(1)*, float, i32, i32) local_unnamed_addr #1 ++ ++; Function Attrs: convergent nounwind ++define spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 { +entry: + ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]] ++ %call = tail call spir_func double @_Z16atomic_fetch_addPU3AS1VU7_Atomicdd(double addrspace(1)* %a, double 0.000000e+00) #2 ++ ++ ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]] + ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}}) + ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+diid]]({{.*}}) -+ %call = tail call spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2 ++ %call1 = tail call spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2 ++ ++ ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]] ++ %call2 = tail call spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order12memory_scope(double addrspace(1)* %a, double 0.000000e+00, i32 0, i32 1) #2 + ret void +} ++ ++; Function Attrs: convergent ++declare spir_func double @_Z16atomic_fetch_addPU3AS1VU7_Atomicdd(double addrspace(1)*, double) local_unnamed_addr #1 ++ +; Function Attrs: convergent +declare spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1 +; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double + ++; Function Attrs: convergent ++declare spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order12memory_scope(double addrspace(1)*, double, i32, i32) local_unnamed_addr #1 ++ +; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float +; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double + -+attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -+attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } ++attributes #0 = { convergent nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } ++attributes #1 = { convergent "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #2 = { convergent nounwind } + +!llvm.module.flags = !{!0} @@ -314,7 +454,7 @@ index 00000000..fb146fb9 + +!0 = !{i32 1, !"wchar_size", i32 4} +!1 = !{i32 2, i32 0} -+!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"} ++!2 = !{!"clang version 10.0.1 (8560093eba963fba2edd47ca85404cdaff22f174)"} diff --git a/test/AtomicFAddExt.ll b/test/AtomicFAddExt.ll index 011dd8a7..42bdfeea 100644 --- a/test/AtomicFAddExt.ll @@ -619,10 +759,10 @@ index 1b81e53b..1c2eec93 100644 + diff --git a/test/AtomicFMaxEXTForOCL.ll b/test/AtomicFMaxEXTForOCL.ll new file mode 100644 -index 00000000..1f2530d9 +index 00000000..4097db02 --- /dev/null +++ b/test/AtomicFMaxEXTForOCL.ll -@@ -0,0 +1,64 @@ +@@ -0,0 +1,88 @@ +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_min_max -o %t.spv +; RUN: spirv-val %t.spv @@ -644,39 +784,63 @@ index 00000000..1f2530d9 +; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 +; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 + -+; Function Attrs: convergent norecurse nounwind -+define dso_local spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 { ++; Function Attrs: convergent nounwind ++define spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 { +entry: + ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_32]] ++ %call = tail call spir_func float @_Z16atomic_fetch_maxPU3AS1VU7_Atomicff(float addrspace(1)* %a, float 0.000000e+00) #2 ++ ++ ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_32]] + ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}}) + ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+fiif]]({{.*}}) -+ %call = tail call spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2 ++ %call1 = tail call spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2 ++ ++ ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_32]] ++ %call2 = tail call spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order12memory_scope(float addrspace(1)* %a, float 0.000000e+00, i32 0, i32 1) #2 + ret void +} + +; Function Attrs: convergent ++declare spir_func float @_Z16atomic_fetch_maxPU3AS1VU7_Atomicff(float addrspace(1)*, float) local_unnamed_addr #1 ++ ++; Function Attrs: convergent +declare spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1 +; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float + -+; Function Attrs: convergent norecurse nounwind -+define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 { ++; Function Attrs: convergent ++declare spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order12memory_scope(float addrspace(1)*, float, i32, i32) local_unnamed_addr #1 ++ ++; Function Attrs: convergent nounwind ++define spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 { +entry: + ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_64]] ++ %call = tail call spir_func double @_Z16atomic_fetch_maxPU3AS1VU7_Atomicdd(double addrspace(1)* %a, double 0.000000e+00) #2 ++ ++ ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_64]] + ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}}) + ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+diid]]({{.*}}) -+ %call = tail call spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2 ++ %call1 = tail call spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2 ++ ++ ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_64]] ++ %call2 = tail call spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order12memory_scope(double addrspace(1)* %a, double 0.000000e+00, i32 0, i32 1) #2 + ret void +} + +; Function Attrs: convergent ++declare spir_func double @_Z16atomic_fetch_maxPU3AS1VU7_Atomicdd(double addrspace(1)*, double) local_unnamed_addr #1 ++ ++; Function Attrs: convergent +declare spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1 +; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double + ++; Function Attrs: convergent ++declare spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order12memory_scope(double addrspace(1)*, double, i32, i32) local_unnamed_addr #1 ++ +; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float +; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double + -+attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -+attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } ++attributes #0 = { convergent nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } ++attributes #1 = { convergent "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #2 = { convergent nounwind } + +!llvm.module.flags = !{!0} @@ -686,7 +850,7 @@ index 00000000..1f2530d9 + +!0 = !{i32 1, !"wchar_size", i32 4} +!1 = !{i32 2, i32 0} -+!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"} ++!2 = !{!"clang version 10.0.1 (8560093eba963fba2edd47ca85404cdaff22f174)"} diff --git a/test/AtomicFMinEXT.ll b/test/AtomicFMinEXT.ll index 98c98b8e..9e40a669 100644 --- a/test/AtomicFMinEXT.ll @@ -841,10 +1005,10 @@ index 98c98b8e..9e40a669 100644 + diff --git a/test/AtomicFMinEXTForOCL.ll b/test/AtomicFMinEXTForOCL.ll new file mode 100644 -index 00000000..6196b0f8 +index 00000000..0ba86b97 --- /dev/null +++ b/test/AtomicFMinEXTForOCL.ll -@@ -0,0 +1,64 @@ +@@ -0,0 +1,85 @@ +; RUN: llvm-as %s -o %t.bc +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_min_max -o %t.spv +; RUN: spirv-val %t.spv @@ -866,39 +1030,60 @@ index 00000000..6196b0f8 +; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 +; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 + -+; Function Attrs: convergent norecurse nounwind -+define dso_local spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 { ++; Function Attrs: convergent nounwind ++define spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 { +entry: + ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_32]] ++ %call = tail call spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2 ++ ++ ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_32]] + ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}}) + ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+fiif]]({{.*}}) -+ %call = tail call spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2 ++ %call1 = tail call spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2 ++ ++ ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_32]] ++ %call2 = tail call spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order12memory_scope(float addrspace(1)* %a, float 0.000000e+00, i32 0, i32 1) #2 + ret void +} + +; Function Attrs: convergent +declare spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1 ++ ++; Function Attrs: convergent ++declare spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order12memory_scope(float addrspace(1)*, float, i32, i32) local_unnamed_addr #1 +; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float + -+; Function Attrs: convergent norecurse nounwind -+define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 { ++; Function Attrs: convergent nounwind ++define spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 { +entry: + ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_64]] ++ %call = tail call spir_func double @_Z16atomic_fetch_minPU3AS1VU7_Atomicdd(double addrspace(1)* %a, double 0.000000e+00) #2 ++ ++ ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_64]] + ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}}) + ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+diid]]({{.*}}) -+ %call = tail call spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2 ++ %call1 = tail call spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2 ++ ++ ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_64]] ++ %call2 = tail call spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order12memory_scope(double addrspace(1)* %a, double 0.000000e+00, i32 0, i32 1) #2 + ret void +} + +; Function Attrs: convergent ++declare spir_func double @_Z16atomic_fetch_minPU3AS1VU7_Atomicdd(double addrspace(1)*, double) local_unnamed_addr #1 ++ ++; Function Attrs: convergent +declare spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1 +; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double + ++; Function Attrs: convergent ++declare spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order12memory_scope(double addrspace(1)*, double, i32, i32) local_unnamed_addr #1 ++ +; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float +; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double + -+attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } -+attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } ++attributes #0 = { convergent nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } ++attributes #1 = { convergent "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="false" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } +attributes #2 = { convergent nounwind } + +!llvm.module.flags = !{!0} @@ -908,15 +1093,22 @@ index 00000000..6196b0f8 + +!0 = !{i32 1, !"wchar_size", i32 4} +!1 = !{i32 2, i32 0} -+!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"} ++!2 = !{!"clang version 10.0.1 (8560093eba963fba2edd47ca85404cdaff22f174)"} diff --git a/test/InvalidAtomicBuiltins.cl b/test/InvalidAtomicBuiltins.cl -index b8ec5b89..2182f070 100644 +index b8ec5b89..f8d159fe 100644 --- a/test/InvalidAtomicBuiltins.cl +++ b/test/InvalidAtomicBuiltins.cl -@@ -41,13 +41,9 @@ float __attribute__((overloadable)) atomic_fetch_xor(volatile generic atomic_flo +@@ -34,20 +34,12 @@ double __attribute__((overloadable)) atom_and(volatile __global double *p, doubl + double __attribute__((overloadable)) atom_or(volatile __global double *p, double val); + double __attribute__((overloadable)) atom_xor(volatile __global double *p, double val); + +-float __attribute__((overloadable)) atomic_fetch_add(volatile generic atomic_float *object, float operand, memory_order order); +-float __attribute__((overloadable)) atomic_fetch_sub(volatile generic atomic_float *object, float operand, memory_order order); + float __attribute__((overloadable)) atomic_fetch_or(volatile generic atomic_float *object, float operand, memory_order order); + float __attribute__((overloadable)) atomic_fetch_xor(volatile generic atomic_float *object, float operand, memory_order order); double __attribute__((overloadable)) atomic_fetch_and(volatile generic atomic_double *object, double operand, memory_order order); - double __attribute__((overloadable)) atomic_fetch_max(volatile generic atomic_double *object, double operand, memory_order order); - double __attribute__((overloadable)) atomic_fetch_min(volatile generic atomic_double *object, double operand, memory_order order); +-double __attribute__((overloadable)) atomic_fetch_max(volatile generic atomic_double *object, double operand, memory_order order); +-double __attribute__((overloadable)) atomic_fetch_min(volatile generic atomic_double *object, double operand, memory_order order); -float __attribute__((overloadable)) atomic_fetch_add_explicit(volatile generic atomic_float *object, float operand, memory_order order); -float __attribute__((overloadable)) atomic_fetch_sub_explicit(volatile generic atomic_float *object, float operand, memory_order order); float __attribute__((overloadable)) atomic_fetch_or_explicit(volatile generic atomic_float *object, float operand, memory_order order); @@ -927,10 +1119,17 @@ index b8ec5b89..2182f070 100644 __kernel void test_atomic_fn(volatile __global float *p, volatile __global double *pp, -@@ -86,11 +82,7 @@ __kernel void test_atomic_fn(volatile __global float *p, +@@ -79,18 +71,10 @@ __kernel void test_atomic_fn(volatile __global float *p, + d = atom_or(pp, val); + d = atom_xor(pp, val); + +- f = atomic_fetch_add(p, val, order); +- f = atomic_fetch_sub(p, val, order); + f = atomic_fetch_or(p, val, order); + f = atomic_fetch_xor(p, val, order); d = atomic_fetch_and(pp, val, order); - d = atomic_fetch_min(pp, val, order); - d = atomic_fetch_max(pp, val, order); +- d = atomic_fetch_min(pp, val, order); +- d = atomic_fetch_max(pp, val, order); - f = atomic_fetch_add_explicit(p, val, order); - f = atomic_fetch_sub_explicit(p, val, order); f = atomic_fetch_or_explicit(p, val, order); @@ -940,5 +1139,5 @@ index b8ec5b89..2182f070 100644 - d = atomic_fetch_max_explicit(pp, val, order); } -- -2.17.1 +2.18.1