diff --git a/patches/spirv/0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch b/patches/spirv/0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch deleted file mode 100644 index c601d113..00000000 --- a/patches/spirv/0002-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch +++ /dev/null @@ -1,1134 +0,0 @@ -From dba4306609dc3a9f8b01f6cdeace48b1a8974695 Mon Sep 17 00:00:00 2001 -From: haonanya -Date: Wed, 28 Jul 2021 18:27:54 +0800 -Subject: [PATCH] Add support for cl_ext_float_atomics in SPIRVWriter - -Signed-off-by: haonanya ---- - lib/SPIRV/OCL20ToSPIRV.cpp | 26 ++++++- - lib/SPIRV/OCLUtil.cpp | 16 +++-- - lib/SPIRV/SPIRVToOCL.h | 3 + - lib/SPIRV/SPIRVToOCL12.cpp | 21 ++++++ - lib/SPIRV/SPIRVToOCL20.cpp | 28 +++++++- - lib/SPIRV/libSPIRV/SPIRVOpCode.h | 8 ++- - test/AtomicBuiltinsFloat.ll | 94 +++++++++++++++++++++++++ - test/AtomicFAddEXTForOCL.ll | 88 ++++++++++++++++++++++++ - test/AtomicFAddExt.ll | 111 +++++++++--------------------- - test/AtomicFMaxEXT.ll | 113 +++++++++---------------------- - test/AtomicFMaxEXTForOCL.ll | 84 +++++++++++++++++++++++ - test/AtomicFMinEXT.ll | 113 +++++++++---------------------- - test/AtomicFMinEXTForOCL.ll | 81 ++++++++++++++++++++++ - test/InvalidAtomicBuiltins.cl | 16 ----- - 14 files changed, 531 insertions(+), 271 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 a742c8cf..a895307a 100644 ---- a/lib/SPIRV/OCL20ToSPIRV.cpp -+++ b/lib/SPIRV/OCL20ToSPIRV.cpp -@@ -407,7 +407,6 @@ void OCL20ToSPIRV::visitCallInst(CallInst &CI) { - } - if (DemangledName.find(kOCLBuiltinName::AtomicPrefix) == 0 || - DemangledName.find(kOCLBuiltinName::AtomPrefix) == 0) { -- - // Compute atomic builtins do not support floating types. - if (CI.getType()->isFloatingPointTy() && - isComputeAtomicOCLBuiltin(DemangledName)) -@@ -819,7 +818,7 @@ void OCL20ToSPIRV::transAtomicBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info) { - AttributeList Attrs = CI->getCalledFunction()->getAttributes(); - mutateCallInstSPIRV( - M, CI, -- [=](CallInst *CI, std::vector &Args) { -+ [=](CallInst *CI, std::vector &Args) -> std::string { - Info.PostProc(Args); - // Order of args in OCL20: - // object, 0-2 other args, 1-2 order, scope -@@ -864,7 +863,28 @@ void OCL20ToSPIRV::transAtomicBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info) { - std::rotate(Args.begin() + 2, Args.begin() + OrderIdx, - Args.end() - Offset); - } -- return getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName)); -+ llvm::Type *AtomicBuiltinsReturnType = -+ CI->getCalledFunction()->getReturnType(); -+ auto IsFPType = [](llvm::Type *ReturnType) { -+ return ReturnType->isHalfTy() || ReturnType->isFloatTy() || -+ ReturnType->isDoubleTy(); -+ }; -+ auto SPIRVFunctionName = -+ getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName)); -+ if (!IsFPType(AtomicBuiltinsReturnType)) -+ return SPIRVFunctionName; -+ // Translate FP-typed atomic builtins. Currently we only need to -+ // 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") -+ .Case("__spirv_AtomicSMax", "__spirv_AtomicFMaxEXT") -+ .Case("__spirv_AtomicSMin", "__spirv_AtomicFMinEXT") -+ .Default("others"); -+ return SPIRFunctionNameForFloatAtomics == "others" -+ ? SPIRVFunctionName -+ : SPIRFunctionNameForFloatAtomics; - }, - &Attrs); - } -diff --git a/lib/SPIRV/OCLUtil.cpp b/lib/SPIRV/OCLUtil.cpp -index 992f173f..f2626c04 100644 ---- a/lib/SPIRV/OCLUtil.cpp -+++ b/lib/SPIRV/OCLUtil.cpp -@@ -120,29 +120,31 @@ size_t getSPIRVAtomicBuiltinNumMemoryOrderArgs(Op OC) { - return 1; - } - -+// atomic_fetch_[add, min, max] and atomic_fetch_[add, min, max]_explicit -+// functions declared in clang headers should be translated to corresponding -+// FP-typed Atomic Instructions - 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_min", true) -+ .EndsWith("atomic_max", true) -+ .EndsWith("atom_add", 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) -- .EndsWith("add_explicit", true) - .EndsWith("sub_explicit", true) - .EndsWith("or_explicit", true) - .EndsWith("xor_explicit", true) - .EndsWith("and_explicit", true) -- .EndsWith("min_explicit", true) -- .EndsWith("max_explicit", true) - .Default(false); - } - -diff --git a/lib/SPIRV/SPIRVToOCL.h b/lib/SPIRV/SPIRVToOCL.h -index 746a7acf..af8dade9 100644 ---- a/lib/SPIRV/SPIRVToOCL.h -+++ b/lib/SPIRV/SPIRVToOCL.h -@@ -208,6 +208,9 @@ public: - - void translateOpaqueTypes(); - -+ // Transform FP atomic opcode to corresponding OpenCL function name -+ virtual std::string mapFPAtomicName(Op OC) = 0; -+ - private: - /// Transform uniform group opcode to corresponding OpenCL function name, - /// example: GroupIAdd(Reduce) => group_iadd => work_group_reduce_add | -diff --git a/lib/SPIRV/SPIRVToOCL12.cpp b/lib/SPIRV/SPIRVToOCL12.cpp -index 1a62c6b8..dc0ba9cc 100644 ---- a/lib/SPIRV/SPIRVToOCL12.cpp -+++ b/lib/SPIRV/SPIRVToOCL12.cpp -@@ -104,6 +104,9 @@ public: - /// cl_khr_int64_base_atomics and cl_khr_int64_extended_atomics extensions. - std::string mapAtomicName(Op OC, Type *Ty); - -+ // Transform FP atomic opcode to corresponding OpenCL function name -+ std::string mapFPAtomicName(Op OC) override; -+ - static char ID; - }; - -@@ -344,6 +347,21 @@ Instruction *SPIRVToOCL12::visitCallSPIRVAtomicBuiltin(CallInst *CI, Op OC) { - return NewCI; - } - -+std::string SPIRVToOCL12::mapFPAtomicName(Op OC) { -+ assert(isFPAtomicOpCode(OC) && "Not intended to handle other opcodes than " -+ "AtomicF{Add/Min/Max}EXT!"); -+ switch (OC) { -+ case OpAtomicFAddEXT: -+ return "atomic_add"; -+ case OpAtomicFMinEXT: -+ return "atomic_min"; -+ case OpAtomicFMaxEXT: -+ return "atomic_max"; -+ default: -+ llvm_unreachable("Unsupported opcode!"); -+ } -+} -+ - Instruction *SPIRVToOCL12::mutateAtomicName(CallInst *CI, Op OC) { - AttributeList Attrs = CI->getCalledFunction()->getAttributes(); - return mutateCallInstOCL( -@@ -357,6 +375,9 @@ Instruction *SPIRVToOCL12::mutateAtomicName(CallInst *CI, Op OC) { - std::string SPIRVToOCL12::mapAtomicName(Op OC, Type *Ty) { - std::string Prefix = Ty->isIntegerTy(64) ? kOCLBuiltinName::AtomPrefix - : kOCLBuiltinName::AtomicPrefix; -+ // Map fp atomic instructions to regular OpenCL built-ins. -+ if (isFPAtomicOpCode(OC)) -+ return mapFPAtomicName(OC); - return Prefix += OCL12SPIRVBuiltinMap::rmap(OC); - } - -diff --git a/lib/SPIRV/SPIRVToOCL20.cpp b/lib/SPIRV/SPIRVToOCL20.cpp -index 24f48956..c51c9189 100644 ---- a/lib/SPIRV/SPIRVToOCL20.cpp -+++ b/lib/SPIRV/SPIRVToOCL20.cpp -@@ -82,6 +82,9 @@ public: - /// compare_exchange_strong/weak_explicit - Instruction *visitCallSPIRVAtomicCmpExchg(CallInst *CI, Op OC) override; - -+ // Transform FP atomic opcode to corresponding OpenCL function name -+ std::string mapFPAtomicName(Op OC) override; -+ - static char ID; - }; - -@@ -150,11 +153,29 @@ void SPIRVToOCL20::visitCallSPIRVControlBarrier(CallInst *CI) { - &Attrs); - } - -+std::string SPIRVToOCL20::mapFPAtomicName(Op OC) { -+ assert(isFPAtomicOpCode(OC) && "Not intended to handle other opcodes than " -+ "AtomicF{Add/Min/Max}EXT!"); -+ switch (OC) { -+ case OpAtomicFAddEXT: -+ return "atomic_fetch_add_explicit"; -+ case OpAtomicFMinEXT: -+ return "atomic_fetch_min_explicit"; -+ case OpAtomicFMaxEXT: -+ return "atomic_fetch_max_explicit"; -+ default: -+ llvm_unreachable("Unsupported opcode!"); -+ } -+} -+ - Instruction *SPIRVToOCL20::mutateAtomicName(CallInst *CI, Op OC) { - AttributeList Attrs = CI->getCalledFunction()->getAttributes(); - return mutateCallInstOCL( - M, CI, - [=](CallInst *, std::vector &Args) { -+ // Map fp atomic instructions to regular OpenCL built-ins. -+ if (isFPAtomicOpCode(OC)) -+ return mapFPAtomicName(OC); - return OCLSPIRVBuiltinMap::rmap(OC); - }, - &Attrs); -@@ -221,7 +242,12 @@ CallInst *SPIRVToOCL20::mutateCommonAtomicArguments(CallInst *CI, Op OC) { - } - } - auto Ptr = findFirstPtr(Args); -- auto Name = OCLSPIRVBuiltinMap::rmap(OC); -+ std::string Name; -+ // Map fp atomic instructions to regular OpenCL built-ins. -+ if (isFPAtomicOpCode(OC)) -+ Name = mapFPAtomicName(OC); -+ else -+ Name = OCLSPIRVBuiltinMap::rmap(OC); - auto NumOrder = getSPIRVAtomicBuiltinNumMemoryOrderArgs(OC); - auto ScopeIdx = Ptr + 1; - auto OrderIdx = Ptr + 2; -diff --git a/lib/SPIRV/libSPIRV/SPIRVOpCode.h b/lib/SPIRV/libSPIRV/SPIRVOpCode.h -index feec70f6..8e595e83 100644 ---- a/lib/SPIRV/libSPIRV/SPIRVOpCode.h -+++ b/lib/SPIRV/libSPIRV/SPIRVOpCode.h -@@ -54,11 +54,17 @@ template <> inline void SPIRVMap::init() { - } - SPIRV_DEF_NAMEMAP(Op, OpCodeNameMap) - -+inline bool isFPAtomicOpCode(Op OpCode) { -+ return OpCode == OpAtomicFAddEXT || OpCode == OpAtomicFMinEXT || -+ OpCode == OpAtomicFMaxEXT; -+} -+ - inline bool isAtomicOpCode(Op OpCode) { - static_assert(OpAtomicLoad < OpAtomicXor, ""); - return ((unsigned)OpCode >= OpAtomicLoad && - (unsigned)OpCode <= OpAtomicXor) || -- OpCode == OpAtomicFlagTestAndSet || OpCode == OpAtomicFlagClear; -+ OpCode == OpAtomicFlagTestAndSet || OpCode == OpAtomicFlagClear || -+ isFPAtomicOpCode(OpCode); - } - 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..2714f066 ---- /dev/null -+++ b/test/AtomicBuiltinsFloat.ll -@@ -0,0 +1,94 @@ -+; Check that translator generate atomic instructions for atomic builtins -+; FP-typed atomic_fetch_sub and atomic_fetch_sub_explicit should be translated -+; to FunctionCall -+; 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 -+; CHECK-COUNT-3: FunctionCall -+ -+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 -+ %call6 = tail call spir_func float @_Z16atomic_fetch_subPU3AS3VU7_Atomicff(float addrspace(3)* %ff, float 1.000000e+00) #2 -+ %call7 = tail call spir_func float @_Z25atomic_fetch_sub_explicitPU3AS3VU7_Atomicff12memory_order(float addrspace(3)* %ff, float 1.000000e+00, i32 0) #2 -+ %call8 = tail call spir_func float @_Z25atomic_fetch_sub_explicitPU3AS3VU7_Atomicff12memory_order12memory_scope(float addrspace(3)* %ff, 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 -+ -+; Function Attrs: convergent -+declare spir_func float @_Z16atomic_fetch_subPU3AS3VU7_Atomicff(float addrspace(3)*, float) local_unnamed_addr #1 -+ -+; Function Attrs: convergent -+declare spir_func float @_Z25atomic_fetch_sub_explicitPU3AS3VU7_Atomicff12memory_order(float addrspace(3)*, float, i32) local_unnamed_addr #1 -+ -+; Function Attrs: convergent -+declare spir_func float @_Z25atomic_fetch_sub_explicitPU3AS3VU7_Atomicff12memory_order12memory_scope(float addrspace(3)*, 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" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-frame-pointer-elim"="false" "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" "less-precise-fpmad"="false" "no-frame-pointer-elim"="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 9.0.1 (338587c46f8c87845bd8b98e0338090655bf0313)"} -+!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..e8c2ab0b ---- /dev/null -+++ b/test/AtomicFAddEXTForOCL.ll -@@ -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 -+; RUN: llvm-spirv -to-text %t.spv -o %t.spt -+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV -+ -+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20 -+ -+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV -+ -+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" -+ -+; CHECK-SPIRV: Capability AtomicFloat32AddEXT -+; CHECK-SPIRV: Capability AtomicFloat64AddEXT -+; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_add" -+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 -+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 -+ -+; 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]]({{.*}}) -+ %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 -+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]]({{.*}}) -+ %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 nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-frame-pointer-elim"="false" "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" "less-precise-fpmad"="false" "no-frame-pointer-elim"="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 9.0.1 (7e3651de4027cff2d6afda804a914e56dc529487)"} -diff --git a/test/AtomicFAddExt.ll b/test/AtomicFAddExt.ll -index 011dd8a7..42bdfeea 100644 ---- a/test/AtomicFAddExt.ll -+++ b/test/AtomicFAddExt.ll -@@ -4,20 +4,16 @@ - ; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV - - ; RUN: llvm-spirv -r %t.spv -o %t.rev.bc --; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL12 - --target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" --target triple = "spir64-unknown-unknown-sycldevice" -- --%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } --%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] } --%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } -- --$_ZTSZZ3addIfEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_37 = comdat any -+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20 - --$_ZTSZZ3addIdEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_37 = comdat any -+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV - --@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 -+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" -+target triple = "spir64-unknown-unknown-sycldevice" - - ; CHECK-SPIRV: Capability AtomicFloat32AddEXT - ; CHECK-SPIRV: Capability AtomicFloat64AddEXT -@@ -25,62 +21,43 @@ $_ZTSZZ3addIdEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_3 - ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 - ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 - --; Function Attrs: convergent norecurse mustprogress --define weak_odr dso_local spir_kernel void @_ZTSZZ3addIfEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_37(float addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, float addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 { -+; Function Attrs: convergent norecurse nounwind -+define dso_local spir_func float @_Z14AtomicFloatIncRf(float addrspace(4)* align 4 dereferenceable(4) %Arg) local_unnamed_addr #0 { - entry: -- %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0 -- %1 = load i64, i64* %0, align 8 -- %add.ptr.i29 = getelementptr inbounds float, float addrspace(1)* %_arg_, i64 %1 -- %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0 -- %3 = load i64, i64* %2, align 8 -- %add.ptr.i = getelementptr inbounds float, float addrspace(1)* %_arg_4, i64 %3 -- %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !5 -- %5 = extractelement <3 x i64> %4, i64 0 -+ %0 = addrspacecast float addrspace(4)* %Arg to float addrspace(1)* - ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]] -- ; CHECK-LLVM: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+]]({{.*}}) -- %call3.i.i.i.i = tail call spir_func float @_Z21__spirv_AtomicFAddEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %add.ptr.i29, i32 1, i32 896, float 1.000000e+00) #2 -- %add.i.i = fadd float %call3.i.i.i.i, 1.000000e+00 -- %sext.i = shl i64 %5, 32 -- %conv5.i = ashr exact i64 %sext.i, 32 -- %ptridx.i.i = getelementptr inbounds float, float addrspace(1)* %add.ptr.i, i64 %conv5.i -- %ptridx.ascast.i.i = addrspacecast float addrspace(1)* %ptridx.i.i to float addrspace(4)* -- store float %add.i.i, float addrspace(4)* %ptridx.ascast.i.i, align 4, !tbaa !14 -- ret void -+ ; CHECK-LLVM-CL12: call spir_func float @[[FLOAT_FUNC_NAME:_Z10atomic_add[[:alnum:]]+ff]]({{.*}}) -+ ; 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]]({{.*}}) -+ %call3.i.i = tail call spir_func float @_Z21__spirv_AtomicFAddEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %0, i32 1, i32 896, float 1.000000e+00) #2 -+ ret float %call3.i.i - } - - ; Function Attrs: convergent --; CHECK-LLVM: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float addrspace(1)*, i32, i32, float) - declare dso_local spir_func float @_Z21__spirv_AtomicFAddEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)*, i32, i32, float) local_unnamed_addr #1 -+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float - --; Function Attrs: convergent norecurse mustprogress --define weak_odr dso_local spir_kernel void @_ZTSZZ3addIdEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_37(double addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, double addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 { -+; Function Attrs: convergent norecurse nounwind -+define dso_local spir_func double @_Z15AtomicDoubleIncRd(double addrspace(4)* align 8 dereferenceable(8) %Arg) local_unnamed_addr #0 { - entry: -- %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0 -- %1 = load i64, i64* %0, align 8 -- %add.ptr.i29 = getelementptr inbounds double, double addrspace(1)* %_arg_, i64 %1 -- %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0 -- %3 = load i64, i64* %2, align 8 -- %add.ptr.i = getelementptr inbounds double, double addrspace(1)* %_arg_4, i64 %3 -- %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !18 -- %5 = extractelement <3 x i64> %4, i64 0 -+ %0 = addrspacecast double addrspace(4)* %Arg to double addrspace(1)* - ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]] -- ; CHECK-LLVM: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+]]({{.*}}) -- %call3.i.i.i.i = tail call spir_func double @_Z21__spirv_AtomicFAddEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %add.ptr.i29, i32 1, i32 896, double 1.000000e+00) #2 -- %add.i.i = fadd double %call3.i.i.i.i, 1.000000e+00 -- %sext.i = shl i64 %5, 32 -- %conv5.i = ashr exact i64 %sext.i, 32 -- %ptridx.i.i = getelementptr inbounds double, double addrspace(1)* %add.ptr.i, i64 %conv5.i -- %ptridx.ascast.i.i = addrspacecast double addrspace(1)* %ptridx.i.i to double addrspace(4)* -- store double %add.i.i, double addrspace(4)* %ptridx.ascast.i.i, align 8, !tbaa !27 -- ret void -+ ; CHECK-LLVM-CL12: call spir_func double @[[DOUBLE_FUNC_NAME:_Z10atomic_add[[:alnum:]]+dd]]({{.*}}) -+ ; 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]]({{.*}}) -+ %call3.i.i = tail call spir_func double @_Z21__spirv_AtomicFAddEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %0, i32 1, i32 896, double 1.000000e+00) #2 -+ ret double %call3.i.i - } - - ; Function Attrs: convergent --; CHECK-LLVM: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double addrspace(1)*, i32, i32, double) - declare dso_local spir_func double @_Z21__spirv_AtomicFAddEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)*, i32, i32, double) local_unnamed_addr #1 -+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double - --attributes #0 = { convergent norecurse } --attributes #1 = { convergent } -+; 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 "disable-tail-calls"="false" "frame-pointer"="all" "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"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } -+attributes #1 = { convergent "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } - attributes #2 = { convergent nounwind } - - !llvm.module.flags = !{!0} -@@ -91,29 +68,5 @@ attributes #2 = { convergent nounwind } - !0 = !{i32 1, !"wchar_size", i32 4} - !1 = !{i32 1, i32 2} - !2 = !{i32 4, i32 100000} --!3 = !{!"clang version 12.0.0"} --!4 = !{i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1} --!5 = !{!6, !8, !10, !12} --!6 = distinct !{!6, !7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"} --!7 = distinct !{!7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"} --!8 = distinct !{!8, !9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"} --!9 = distinct !{!9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"} --!10 = distinct !{!10, !11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"} --!11 = distinct !{!11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"} --!12 = distinct !{!12, !13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"} --!13 = distinct !{!13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"} --!14 = !{!15, !15, i64 0} --!15 = !{!"float", !16, i64 0} --!16 = !{!"omnipotent char", !17, i64 0} --!17 = !{!"Simple C++ TBAA"} --!18 = !{!19, !21, !23, !25} --!19 = distinct !{!19, !20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"} --!20 = distinct !{!20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"} --!21 = distinct !{!21, !22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"} --!22 = distinct !{!22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"} --!23 = distinct !{!23, !24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"} --!24 = distinct !{!24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"} --!25 = distinct !{!25, !26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"} --!26 = distinct !{!26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"} --!27 = !{!28, !28, i64 0} --!28 = !{!"double", !16, i64 0} -+!3 = !{!"clang version 13.0.0"} -+ -diff --git a/test/AtomicFMaxEXT.ll b/test/AtomicFMaxEXT.ll -index 1b81e53b..1c2eec93 100644 ---- a/test/AtomicFMaxEXT.ll -+++ b/test/AtomicFMaxEXT.ll -@@ -4,20 +4,16 @@ - ; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV - - ; RUN: llvm-spirv -r %t.spv -o %t.rev.bc --; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL12 - --target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" --target triple = "spir64-unknown-unknown-sycldevice" -- --%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } --%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] } --%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } -- --$_ZTSZZ8max_testIfEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37 = comdat any -+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20 - --$_ZTSZZ8max_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37 = comdat any -+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV - --@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 -+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" -+target triple = "spir64-unknown-unknown-sycldevice" - - ; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT - ; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT -@@ -25,65 +21,42 @@ $_ZTSZZ8max_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4item - ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 - ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 - --; Function Attrs: convergent norecurse --define weak_odr dso_local spir_kernel void @_ZTSZZ8max_testIfEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37(float addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, float addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 { -+; Function Attrs: convergent norecurse nounwind -+define dso_local spir_func float @_Z14AtomicFloatMaxRf(float addrspace(4)* align 4 dereferenceable(4) %Arg) local_unnamed_addr #0 { - entry: -- %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0 -- %1 = load i64, i64* %0, align 8 -- %add.ptr.i29 = getelementptr inbounds float, float addrspace(1)* %_arg_, i64 %1 -- %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0 -- %3 = load i64, i64* %2, align 8 -- %add.ptr.i = getelementptr inbounds float, float addrspace(1)* %_arg_4, i64 %3 -- %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !5 -- %5 = extractelement <3 x i64> %4, i64 0 -- %conv.i = trunc i64 %5 to i32 -- %conv3.i = sitofp i32 %conv.i to float -- %add.i = fadd float %conv3.i, 1.000000e+00 -+ %0 = addrspacecast float addrspace(4)* %Arg to float addrspace(1)* - ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_32]] -- ; CHECK-LLVM: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+]]({{.*}}) -- %call3.i.i.i = tail call spir_func float @_Z21__spirv_AtomicFMaxEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %add.ptr.i29, i32 1, i32 896, float %add.i) #2 -- %sext.i = shl i64 %5, 32 -- %conv6.i = ashr exact i64 %sext.i, 32 -- %ptridx.i.i = getelementptr inbounds float, float addrspace(1)* %add.ptr.i, i64 %conv6.i -- %ptridx.ascast.i.i = addrspacecast float addrspace(1)* %ptridx.i.i to float addrspace(4)* -- store float %call3.i.i.i, float addrspace(4)* %ptridx.ascast.i.i, align 4, !tbaa !14 -- ret void -+ ; CHECK-LLVM-CL12: call spir_func float @[[FLOAT_FUNC_NAME:_Z10atomic_max[[:alnum:]]+ff]]({{.*}}) -+ ; 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.i.i.i = tail call spir_func float @_Z21__spirv_AtomicFMaxEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %0, i32 1, i32 896, float 1.000000e+00) #2 -+ ret float %call.i.i.i - } - - ; Function Attrs: convergent --; CHECK-LLVM: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float addrspace(1)*, i32, i32, float) - declare dso_local spir_func float @_Z21__spirv_AtomicFMaxEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)*, i32, i32, float) local_unnamed_addr #1 -+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float - --; Function Attrs: convergent norecurse --define weak_odr dso_local spir_kernel void @_ZTSZZ8max_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37(double addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, double addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 { -+; Function Attrs: convergent norecurse nounwind -+define dso_local spir_func double @_Z15AtomicDoubleMaxRd(double addrspace(4)* align 8 dereferenceable(8) %Arg) local_unnamed_addr #0 { - entry: -- %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0 -- %1 = load i64, i64* %0, align 8 -- %add.ptr.i29 = getelementptr inbounds double, double addrspace(1)* %_arg_, i64 %1 -- %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0 -- %3 = load i64, i64* %2, align 8 -- %add.ptr.i = getelementptr inbounds double, double addrspace(1)* %_arg_4, i64 %3 -- %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !18 -- %5 = extractelement <3 x i64> %4, i64 0 -- %conv.i = trunc i64 %5 to i32 -- %conv3.i = sitofp i32 %conv.i to double -- %add.i = fadd double %conv3.i, 1.000000e+00 -+ %0 = addrspacecast double addrspace(4)* %Arg to double addrspace(1)* - ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_64]] -- ; CHECK-LLVM: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+]]({{.*}}) -- %call3.i.i.i = tail call spir_func double @_Z21__spirv_AtomicFMaxEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %add.ptr.i29, i32 1, i32 896, double %add.i) #2 -- %sext.i = shl i64 %5, 32 -- %conv6.i = ashr exact i64 %sext.i, 32 -- %ptridx.i.i = getelementptr inbounds double, double addrspace(1)* %add.ptr.i, i64 %conv6.i -- %ptridx.ascast.i.i = addrspacecast double addrspace(1)* %ptridx.i.i to double addrspace(4)* -- store double %call3.i.i.i, double addrspace(4)* %ptridx.ascast.i.i, align 8, !tbaa !27 -- ret void -+ ; CHECK-LLVM-CL12: call spir_func double @[[DOUBLE_FUNC_NAME:_Z10atomic_max[[:alnum:]]+dd]]({{.*}}) -+ ; 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.i.i.i = tail call spir_func double @_Z21__spirv_AtomicFMaxEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %0, i32 1, i32 896, double 1.000000e+00) #2 -+ ret double %call.i.i.i - } - - ; Function Attrs: convergent --; CHECK-LLVM: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double addrspace(1)*, i32, i32, double) - declare dso_local spir_func double @_Z21__spirv_AtomicFMaxEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)*, i32, i32, double) local_unnamed_addr #1 -+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double - --attributes #0 = { convergent norecurse "disable-tail-calls"="false" "frame-pointer"="all" "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"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } -+; 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 "disable-tail-calls"="false" "frame-pointer"="all" "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"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } - attributes #1 = { convergent "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } - attributes #2 = { convergent nounwind } - -@@ -95,29 +68,5 @@ attributes #2 = { convergent nounwind } - !0 = !{i32 1, !"wchar_size", i32 4} - !1 = !{i32 1, i32 2} - !2 = !{i32 4, i32 100000} --!3 = !{!"clang version 12.0.0"} --!4 = !{i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1} --!5 = !{!6, !8, !10, !12} --!6 = distinct !{!6, !7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"} --!7 = distinct !{!7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"} --!8 = distinct !{!8, !9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"} --!9 = distinct !{!9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"} --!10 = distinct !{!10, !11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"} --!11 = distinct !{!11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"} --!12 = distinct !{!12, !13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"} --!13 = distinct !{!13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"} --!14 = !{!15, !15, i64 0} --!15 = !{!"float", !16, i64 0} --!16 = !{!"omnipotent char", !17, i64 0} --!17 = !{!"Simple C++ TBAA"} --!18 = !{!19, !21, !23, !25} --!19 = distinct !{!19, !20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"} --!20 = distinct !{!20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"} --!21 = distinct !{!21, !22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"} --!22 = distinct !{!22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"} --!23 = distinct !{!23, !24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"} --!24 = distinct !{!24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"} --!25 = distinct !{!25, !26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"} --!26 = distinct !{!26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"} --!27 = !{!28, !28, i64 0} --!28 = !{!"double", !16, i64 0} -+!3 = !{!"clang version 13.0.0"} -+ -diff --git a/test/AtomicFMaxEXTForOCL.ll b/test/AtomicFMaxEXTForOCL.ll -new file mode 100644 -index 00000000..a0bd8320 ---- /dev/null -+++ b/test/AtomicFMaxEXTForOCL.ll -@@ -0,0 +1,84 @@ -+; 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 -+; RUN: llvm-spirv -to-text %t.spv -o %t.spt -+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV -+ -+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20 -+ -+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV -+ -+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" -+ -+; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT -+; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT -+; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_min_max" -+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 -+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 -+ -+; 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]]({{.*}}) -+ %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 -+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]]({{.*}}) -+ %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 nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-frame-pointer-elim"="false" "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" "less-precise-fpmad"="false" "no-frame-pointer-elim"="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 9.0.1 (7e3651de4027cff2d6afda804a914e56dc529487)"} -diff --git a/test/AtomicFMinEXT.ll b/test/AtomicFMinEXT.ll -index 98c98b8e..9e40a669 100644 ---- a/test/AtomicFMinEXT.ll -+++ b/test/AtomicFMinEXT.ll -@@ -4,20 +4,16 @@ - ; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV - - ; RUN: llvm-spirv -r %t.spv -o %t.rev.bc --; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL12 - --target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" --target triple = "spir64-unknown-unknown-sycldevice" -- --%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } --%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] } --%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } -- --$_ZTSZZ8min_testIfEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37 = comdat any -+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20 - --$_ZTSZZ8min_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37 = comdat any -+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV - --@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 -+target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024-n8:16:32:64" -+target triple = "spir64-unknown-unknown-sycldevice" - - ; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT - ; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT -@@ -25,65 +21,42 @@ $_ZTSZZ8min_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4item - ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 - ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 - --; Function Attrs: convergent norecurse --define weak_odr dso_local spir_kernel void @_ZTSZZ8min_testIfEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37(float addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, float addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 { -+; Function Attrs: convergent norecurse nounwind -+define dso_local spir_func float @_Z14AtomicFloatMinRf(float addrspace(4)* align 4 dereferenceable(4) %Arg) local_unnamed_addr #0 { - entry: -- %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0 -- %1 = load i64, i64* %0, align 8 -- %add.ptr.i29 = getelementptr inbounds float, float addrspace(1)* %_arg_, i64 %1 -- %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0 -- %3 = load i64, i64* %2, align 8 -- %add.ptr.i = getelementptr inbounds float, float addrspace(1)* %_arg_4, i64 %3 -- %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !5 -- %5 = extractelement <3 x i64> %4, i64 0 -- %conv.i = trunc i64 %5 to i32 -- %conv3.i = sitofp i32 %conv.i to float -- %add.i = fadd float %conv3.i, 1.000000e+00 -+ %0 = addrspacecast float addrspace(4)* %Arg to float addrspace(1)* - ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_32]] -- ; CHECK-LLVM: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+]]({{.*}}) -- %call3.i.i.i = tail call spir_func float @_Z21__spirv_AtomicFMinEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %add.ptr.i29, i32 1, i32 896, float %add.i) #2 -- %sext.i = shl i64 %5, 32 -- %conv6.i = ashr exact i64 %sext.i, 32 -- %ptridx.i.i = getelementptr inbounds float, float addrspace(1)* %add.ptr.i, i64 %conv6.i -- %ptridx.ascast.i.i = addrspacecast float addrspace(1)* %ptridx.i.i to float addrspace(4)* -- store float %call3.i.i.i, float addrspace(4)* %ptridx.ascast.i.i, align 4, !tbaa !14 -- ret void -+ ; CHECK-LLVM-CL12: call spir_func float @[[FLOAT_FUNC_NAME:_Z10atomic_min[[:alnum:]]+ff]]({{.*}}) -+ ; 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.i.i.i = tail call spir_func float @_Z21__spirv_AtomicFMinEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %0, i32 1, i32 896, float 1.000000e+00) #2 -+ ret float %call.i.i.i - } - - ; Function Attrs: convergent --; CHECK-LLVM: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float addrspace(1)*, i32, i32, float) - declare dso_local spir_func float @_Z21__spirv_AtomicFMinEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)*, i32, i32, float) local_unnamed_addr #1 -+; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float - --; Function Attrs: convergent norecurse --define weak_odr dso_local spir_kernel void @_ZTSZZ8min_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37(double addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, double addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 { -+; Function Attrs: convergent norecurse nounwind -+define dso_local spir_func double @_Z15AtomicDoubleMinRd(double addrspace(4)* align 8 dereferenceable(8) %Arg) local_unnamed_addr #0 { - entry: -- %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0 -- %1 = load i64, i64* %0, align 8 -- %add.ptr.i29 = getelementptr inbounds double, double addrspace(1)* %_arg_, i64 %1 -- %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0 -- %3 = load i64, i64* %2, align 8 -- %add.ptr.i = getelementptr inbounds double, double addrspace(1)* %_arg_4, i64 %3 -- %4 = load <3 x i64>, <3 x i64> addrspace(4)* addrspacecast (<3 x i64> addrspace(1)* @__spirv_BuiltInGlobalInvocationId to <3 x i64> addrspace(4)*), align 32, !noalias !18 -- %5 = extractelement <3 x i64> %4, i64 0 -- %conv.i = trunc i64 %5 to i32 -- %conv3.i = sitofp i32 %conv.i to double -- %add.i = fadd double %conv3.i, 1.000000e+00 -+ %0 = addrspacecast double addrspace(4)* %Arg to double addrspace(1)* - ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_64]] -- ; CHECK-LLVM: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+]]({{.*}}) -- %call3.i.i.i = tail call spir_func double @_Z21__spirv_AtomicFMinEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %add.ptr.i29, i32 1, i32 896, double %add.i) #2 -- %sext.i = shl i64 %5, 32 -- %conv6.i = ashr exact i64 %sext.i, 32 -- %ptridx.i.i = getelementptr inbounds double, double addrspace(1)* %add.ptr.i, i64 %conv6.i -- %ptridx.ascast.i.i = addrspacecast double addrspace(1)* %ptridx.i.i to double addrspace(4)* -- store double %call3.i.i.i, double addrspace(4)* %ptridx.ascast.i.i, align 8, !tbaa !27 -- ret void -+ ; CHECK-LLVM-CL12: call spir_func double @[[DOUBLE_FUNC_NAME:_Z10atomic_min[[:alnum:]]+dd]]({{.*}}) -+ ; 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.i.i.i = tail call spir_func double @_Z21__spirv_AtomicFMinEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %0, i32 1, i32 896, double 1.000000e+00) #2 -+ ret double %call.i.i.i - } - - ; Function Attrs: convergent --; CHECK-LLVM: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double addrspace(1)*, i32, i32, double) - declare dso_local spir_func double @_Z21__spirv_AtomicFMinEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)*, i32, i32, double) local_unnamed_addr #1 -+; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double - --attributes #0 = { convergent norecurse "disable-tail-calls"="false" "frame-pointer"="all" "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"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } -+; 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 "disable-tail-calls"="false" "frame-pointer"="all" "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"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } - attributes #1 = { convergent "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } - attributes #2 = { convergent nounwind } - -@@ -95,29 +68,5 @@ attributes #2 = { convergent nounwind } - !0 = !{i32 1, !"wchar_size", i32 4} - !1 = !{i32 1, i32 2} - !2 = !{i32 4, i32 100000} --!3 = !{!"clang version 12.0.0 (https://github.com/otcshare/llvm.git 67add71766d55d6a8d8d894822f583d6365a3b7d)"} --!4 = !{i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1} --!5 = !{!6, !8, !10, !12} --!6 = distinct !{!6, !7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"} --!7 = distinct !{!7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"} --!8 = distinct !{!8, !9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"} --!9 = distinct !{!9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"} --!10 = distinct !{!10, !11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"} --!11 = distinct !{!11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"} --!12 = distinct !{!12, !13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"} --!13 = distinct !{!13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"} --!14 = !{!15, !15, i64 0} --!15 = !{!"float", !16, i64 0} --!16 = !{!"omnipotent char", !17, i64 0} --!17 = !{!"Simple C++ TBAA"} --!18 = !{!19, !21, !23, !25} --!19 = distinct !{!19, !20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"} --!20 = distinct !{!20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"} --!21 = distinct !{!21, !22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"} --!22 = distinct !{!22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"} --!23 = distinct !{!23, !24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"} --!24 = distinct !{!24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"} --!25 = distinct !{!25, !26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"} --!26 = distinct !{!26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"} --!27 = !{!28, !28, i64 0} --!28 = !{!"double", !16, i64 0} -+!3 = !{!"clang version 13.0.0"} -+ -diff --git a/test/AtomicFMinEXTForOCL.ll b/test/AtomicFMinEXTForOCL.ll -new file mode 100644 -index 00000000..db42af41 ---- /dev/null -+++ b/test/AtomicFMinEXTForOCL.ll -@@ -0,0 +1,81 @@ -+; 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 -+; RUN: llvm-spirv -to-text %t.spv -o %t.spt -+; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV -+ -+; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20 -+ -+; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV -+ -+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" -+ -+; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT -+; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT -+; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_min_max" -+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 -+; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 -+ -+; 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]]({{.*}}) -+ %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 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]]({{.*}}) -+ %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 nounwind "correctly-rounded-divide-sqrt-fp-math"="false" "denorms-are-zero"="false" "disable-tail-calls"="false" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-frame-pointer-elim"="false" "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" "less-precise-fpmad"="false" "no-frame-pointer-elim"="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 9.0.1 (7e3651de4027cff2d6afda804a914e56dc529487)"} -diff --git a/test/InvalidAtomicBuiltins.cl b/test/InvalidAtomicBuiltins.cl -index 111d54c5..8eca4b1d 100644 ---- a/test/InvalidAtomicBuiltins.cl -+++ b/test/InvalidAtomicBuiltins.cl -@@ -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); --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); - float __attribute__((overloadable)) atomic_fetch_xor_explicit(volatile generic atomic_float *object, float operand, memory_order order); - double __attribute__((overloadable)) atomic_fetch_and_explicit(volatile generic atomic_double *object, double operand, memory_order order); --double __attribute__((overloadable)) atomic_fetch_max_explicit(volatile generic atomic_double *object, double operand, memory_order order); --double __attribute__((overloadable)) atomic_fetch_min_explicit(volatile generic atomic_double *object, double operand, memory_order order); - - __kernel void test_atomic_fn(volatile __global float *p, - volatile __global double *pp, -@@ -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); -- f = atomic_fetch_add_explicit(p, val, order); -- f = atomic_fetch_sub_explicit(p, val, order); - f = atomic_fetch_or_explicit(p, val, order); - f = atomic_fetch_xor_explicit(p, val, order); - d = atomic_fetch_and_explicit(pp, val, order); -- d = atomic_fetch_min_explicit(pp, val, order); -- d = atomic_fetch_max_explicit(pp, val, order); - } --- -2.17.1 -