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 deleted file mode 100644 index 56fde0a2..00000000 --- a/patches/spirv/0001-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch +++ /dev/null @@ -1,1167 +0,0 @@ -From cc687f3c2399b83da0156ff24d09c24bab31e74e Mon Sep 17 00:00:00 2001 -From: haonanya -Date: Wed, 28 Jul 2021 14:24:23 +0800 -Subject: [PATCH] Add support for cl_ext_float_atomics in SPIRVWriter - -Signed-off-by: haonanya ---- - include/LLVMSPIRVExtensions.inc | 1 + - lib/SPIRV/OCLToSPIRV.cpp | 27 +++++- - lib/SPIRV/OCLUtil.cpp | 15 ++-- - lib/SPIRV/SPIRVToOCL.h | 3 + - lib/SPIRV/SPIRVToOCL12.cpp | 21 +++++ - lib/SPIRV/SPIRVToOCL20.cpp | 28 +++++- - lib/SPIRV/libSPIRV/SPIRVInstruction.h | 21 +++++ - lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h | 3 + - lib/SPIRV/libSPIRV/SPIRVOpCode.h | 8 +- - lib/SPIRV/libSPIRV/SPIRVOpCodeEnum.h | 2 + - lib/SPIRV/libSPIRV/spirv.hpp | 7 ++ - test/AtomicBuiltinsFloat.ll | 94 +++++++++++++++++++ - test/AtomicFAddEXT.ll | 72 +++++++++++++++ - test/AtomicFAddEXTForOCL.ll | 84 +++++++++++++++++ - test/AtomicFAddExt.ll | 119 ------------------------- - test/AtomicFMaxEXT.ll | 73 +++++++++++++++ - test/AtomicFMaxEXTForOCL.ll | 84 +++++++++++++++++ - test/AtomicFMinEXT.ll | 73 +++++++++++++++ - test/AtomicFMinEXTForOCL.ll | 81 +++++++++++++++++ - test/negative/InvalidAtomicBuiltins.cl | 18 +--- - 20 files changed, 688 insertions(+), 146 deletions(-) - create mode 100644 test/AtomicBuiltinsFloat.ll - create mode 100644 test/AtomicFAddEXT.ll - create mode 100644 test/AtomicFAddEXTForOCL.ll - delete mode 100644 test/AtomicFAddExt.ll - create mode 100644 test/AtomicFMaxEXT.ll - create mode 100644 test/AtomicFMaxEXTForOCL.ll - create mode 100644 test/AtomicFMinEXT.ll - create mode 100644 test/AtomicFMinEXTForOCL.ll - -diff --git a/include/LLVMSPIRVExtensions.inc b/include/LLVMSPIRVExtensions.inc -index e313a8e1..83469be1 100644 ---- a/include/LLVMSPIRVExtensions.inc -+++ b/include/LLVMSPIRVExtensions.inc -@@ -3,6 +3,7 @@ - #endif - - EXT(SPV_EXT_shader_atomic_float_add) -+EXT(SPV_EXT_shader_atomic_float_min_max) - EXT(SPV_KHR_no_integer_wrap_decoration) - EXT(SPV_KHR_float_controls) - EXT(SPV_INTEL_subgroups) -diff --git a/lib/SPIRV/OCLToSPIRV.cpp b/lib/SPIRV/OCLToSPIRV.cpp -index 7c65b9e8..7ea350ff 100644 ---- a/lib/SPIRV/OCLToSPIRV.cpp -+++ b/lib/SPIRV/OCLToSPIRV.cpp -@@ -386,7 +386,6 @@ void OCLToSPIRV::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)) -@@ -798,7 +797,7 @@ void OCLToSPIRV::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 -@@ -827,7 +826,29 @@ void OCLToSPIRV::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 2cc5d815..89ae7fe7 100644 ---- a/lib/SPIRV/OCLUtil.cpp -+++ b/lib/SPIRV/OCLUtil.cpp -@@ -655,29 +655,32 @@ 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 160fcf5c..31036bfa 100644 ---- a/lib/SPIRV/SPIRVToOCL.h -+++ b/lib/SPIRV/SPIRVToOCL.h -@@ -214,6 +214,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 75fce5a1..01b23e2c 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; - }; - -@@ -312,6 +315,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( -@@ -325,6 +343,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 b6838512..b0a54c3d 100644 ---- a/lib/SPIRV/SPIRVToOCL20.cpp -+++ b/lib/SPIRV/SPIRVToOCL20.cpp -@@ -83,6 +83,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; - }; - -@@ -161,11 +164,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); -@@ -232,7 +253,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/SPIRVInstruction.h b/lib/SPIRV/libSPIRV/SPIRVInstruction.h -index bd576d46..16d8383a 100644 ---- a/lib/SPIRV/libSPIRV/SPIRVInstruction.h -+++ b/lib/SPIRV/libSPIRV/SPIRVInstruction.h -@@ -2779,6 +2779,25 @@ public: - } - }; - -+class SPIRVAtomicFMinMaxEXTBase : public SPIRVAtomicInstBase { -+public: -+ llvm::Optional getRequiredExtension() const override { -+ return ExtensionID::SPV_EXT_shader_atomic_float_min_max; -+ } -+ -+ SPIRVCapVec getRequiredCapability() const override { -+ assert(hasType()); -+ if (getType()->isTypeFloat(16)) -+ return {CapabilityAtomicFloat16MinMaxEXT}; -+ if (getType()->isTypeFloat(32)) -+ return {CapabilityAtomicFloat32MinMaxEXT}; -+ if (getType()->isTypeFloat(64)) -+ return {CapabilityAtomicFloat64MinMaxEXT}; -+ llvm_unreachable( -+ "AtomicF(Min|Max)EXT can only be generated for f16, f32, f64 types"); -+ } -+}; -+ - #define _SPIRV_OP(x, ...) \ - typedef SPIRVInstTemplate SPIRV##x; - // Atomic builtins -@@ -2806,6 +2825,8 @@ _SPIRV_OP(MemoryBarrier, false, 3) - // Specialized atomic builtins - _SPIRV_OP(AtomicStore, AtomicStoreInst, false, 5) - _SPIRV_OP(AtomicFAddEXT, AtomicFAddEXTInst, true, 7) -+_SPIRV_OP(AtomicFMinEXT, AtomicFMinMaxEXTBase, true, 7) -+_SPIRV_OP(AtomicFMaxEXT, AtomicFMinMaxEXTBase, true, 7) - #undef _SPIRV_OP - - class SPIRVImageInstBase : public SPIRVInstTemplateBase { -diff --git a/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h b/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h -index 5d17e581..5519af83 100644 ---- a/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h -+++ b/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h -@@ -455,6 +455,9 @@ template <> inline void SPIRVMap::init() { - "PhysicalStorageBufferAddressesEXT"); - add(CapabilityAtomicFloat32AddEXT, "AtomicFloat32AddEXT"); - add(CapabilityAtomicFloat64AddEXT, "AtomicFloat64AddEXT"); -+ add(CapabilityAtomicFloat32MinMaxEXT, "AtomicFloat32MinMaxEXT"); -+ add(CapabilityAtomicFloat64MinMaxEXT, "AtomicFloat64MinMaxEXT"); -+ add(CapabilityAtomicFloat16MinMaxEXT, "AtomicFloat16MinMaxEXT"); - add(CapabilityComputeDerivativeGroupLinearNV, - "ComputeDerivativeGroupLinearNV"); - add(CapabilityCooperativeMatrixNV, "CooperativeMatrixNV"); -diff --git a/lib/SPIRV/libSPIRV/SPIRVOpCode.h b/lib/SPIRV/libSPIRV/SPIRVOpCode.h -index fda94d97..a4cfe2a8 100644 ---- a/lib/SPIRV/libSPIRV/SPIRVOpCode.h -+++ b/lib/SPIRV/libSPIRV/SPIRVOpCode.h -@@ -58,11 +58,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/lib/SPIRV/libSPIRV/SPIRVOpCodeEnum.h b/lib/SPIRV/libSPIRV/SPIRVOpCodeEnum.h -index bc70cdd6..ec6afc51 100644 ---- a/lib/SPIRV/libSPIRV/SPIRVOpCodeEnum.h -+++ b/lib/SPIRV/libSPIRV/SPIRVOpCodeEnum.h -@@ -342,6 +342,8 @@ _SPIRV_OP(FunctionPointerCallINTEL, 5601) - _SPIRV_OP(AsmTargetINTEL, 5609) - _SPIRV_OP(AsmINTEL, 5610) - _SPIRV_OP(AsmCallINTEL, 5611) -+_SPIRV_OP(AtomicFMinEXT, 5614) -+_SPIRV_OP(AtomicFMaxEXT, 5615) - _SPIRV_OP(VmeImageINTEL, 5699) - _SPIRV_OP(TypeVmeImageINTEL, 5700) - _SPIRV_OP(TypeAvcImePayloadINTEL, 5701) -diff --git a/lib/SPIRV/libSPIRV/spirv.hpp b/lib/SPIRV/libSPIRV/spirv.hpp -index f0e311c6..2a86f32e 100644 ---- a/lib/SPIRV/libSPIRV/spirv.hpp -+++ b/lib/SPIRV/libSPIRV/spirv.hpp -@@ -1020,6 +1020,9 @@ enum Capability { - CapabilityFunctionPointersINTEL = 5603, - CapabilityIndirectReferencesINTEL = 5604, - CapabilityAsmINTEL = 5606, -+ CapabilityAtomicFloat32MinMaxEXT = 5612, -+ CapabilityAtomicFloat64MinMaxEXT = 5613, -+ CapabilityAtomicFloat16MinMaxEXT = 5616, - CapabilityVectorComputeINTEL = 5617, - CapabilityVectorAnyINTEL = 5619, - CapabilitySubgroupAvcMotionEstimationINTEL = 5696, -@@ -1538,6 +1541,8 @@ enum Op { - OpAsmTargetINTEL = 5609, - OpAsmINTEL = 5610, - OpAsmCallINTEL = 5611, -+ OpAtomicFMinEXT = 5614, -+ OpAtomicFMaxEXT = 5615, - OpDecorateString = 5632, - OpDecorateStringGOOGLE = 5632, - OpMemberDecorateString = 5633, -@@ -2175,6 +2180,8 @@ inline void HasResultAndType(Op opcode, bool *hasResult, bool *hasResultType) { - case OpAsmTargetINTEL: *hasResult = true; *hasResultType = true; break; - case OpAsmINTEL: *hasResult = true; *hasResultType = true; break; - case OpAsmCallINTEL: *hasResult = true; *hasResultType = true; break; -+ case OpAtomicFMinEXT: *hasResult = true; *hasResultType = true; break; -+ case OpAtomicFMaxEXT: *hasResult = true; *hasResultType = true; break; - case OpDecorateString: *hasResult = false; *hasResultType = false; break; - case OpMemberDecorateString: *hasResult = false; *hasResultType = false; break; - case OpVmeImageINTEL: *hasResult = true; *hasResultType = true; break; -diff --git a/test/AtomicBuiltinsFloat.ll b/test/AtomicBuiltinsFloat.ll -new file mode 100644 -index 00000000..778c0cb0 ---- /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 norecurse nounwind -+define dso_local 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 norecurse nounwind "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"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="false" "unsafe-fp-math"="false" "use-soft-float"="false" } -+attributes #1 = { convergent "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"="true" "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 12.0.1 (https://github.com/llvm/llvm-project.git 23fe7b104a0adaaaecd52108105f49297c420c9b)"} -+!3 = !{i32 3} -+!4 = !{!"none"} -+!5 = !{!"atomic_float*"} -+!6 = !{!"_Atomic(float)*"} -+!7 = !{!"volatile"} -diff --git a/test/AtomicFAddEXT.ll b/test/AtomicFAddEXT.ll -new file mode 100644 -index 00000000..b012c904 ---- /dev/null -+++ b/test/AtomicFAddEXT.ll -@@ -0,0 +1,72 @@ -+; 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 -r %t.spv -o %t.rev.bc -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL12 -+ -+; 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-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 -+; 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 norecurse nounwind -+define dso_local spir_func float @_Z14AtomicFloatIncRf(float addrspace(4)* align 4 dereferenceable(4) %Arg) local_unnamed_addr #0 { -+entry: -+ %0 = addrspacecast float addrspace(4)* %Arg to float addrspace(1)* -+ ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]] -+ ; 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 -+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 nounwind -+define dso_local spir_func double @_Z15AtomicDoubleIncRd(double addrspace(4)* align 8 dereferenceable(8) %Arg) local_unnamed_addr #0 { -+entry: -+ %0 = addrspacecast double addrspace(4)* %Arg to double addrspace(1)* -+ ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]] -+ ; 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 -+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 -+ -+; 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} -+!opencl.spir.version = !{!1} -+!spirv.Source = !{!2} -+!llvm.ident = !{!3} -+ -+!0 = !{i32 1, !"wchar_size", i32 4} -+!1 = !{i32 1, i32 2} -+!2 = !{i32 4, i32 100000} -+!3 = !{!"clang version 13.0.0"} -diff --git a/test/AtomicFAddEXTForOCL.ll b/test/AtomicFAddEXTForOCL.ll -new file mode 100644 -index 00000000..4dc4564a ---- /dev/null -+++ b/test/AtomicFAddEXTForOCL.ll -@@ -0,0 +1,84 @@ -+; 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 norecurse nounwind -+define dso_local 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 norecurse nounwind -+define dso_local 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 norecurse nounwind "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"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } -+attributes #1 = { convergent "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"="true" "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 12.0.1 (https://github.com/llvm/llvm-project.git 23fe7b104a0adaaaecd52108105f49297c420c9b)"} -diff --git a/test/AtomicFAddExt.ll b/test/AtomicFAddExt.ll -deleted file mode 100644 -index 58e9f576..00000000 ---- a/test/AtomicFAddExt.ll -+++ /dev/null -@@ -1,119 +0,0 @@ --; RUN: llvm-as %s -o %t.bc --; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_add -o %t.spv --; RUN: llvm-spirv -to-text %t.spv -o %t.spt --; 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 -- --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 -- --$_ZTSZZ3addIdEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_37 = comdat any -- --@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 -- --; 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 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 { --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 -- ; 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 --} -- --; 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 -- --; 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 { --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 -- ; 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 --} -- --; 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 -- --attributes #0 = { convergent norecurse mustprogress "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" "sycl-module-id"="fadd.cpp" "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 } -- --!llvm.module.flags = !{!0} --!opencl.spir.version = !{!1} --!spirv.Source = !{!2} --!llvm.ident = !{!3} -- --!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} -diff --git a/test/AtomicFMaxEXT.ll b/test/AtomicFMaxEXT.ll -new file mode 100644 -index 00000000..67111c43 ---- /dev/null -+++ b/test/AtomicFMaxEXT.ll -@@ -0,0 +1,73 @@ -+; 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 -r %t.spv -o %t.rev.bc -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL12 -+ -+; 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-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 -+; 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 norecurse nounwind -+define dso_local spir_func float @_Z14AtomicFloatMaxRf(float addrspace(4)* align 4 dereferenceable(4) %Arg) local_unnamed_addr #0 { -+entry: -+ %0 = addrspacecast float addrspace(4)* %Arg to float addrspace(1)* -+ ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_32]] -+ ; 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 -+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 nounwind -+define dso_local spir_func double @_Z15AtomicDoubleMaxRd(double addrspace(4)* align 8 dereferenceable(8) %Arg) local_unnamed_addr #0 { -+entry: -+ %0 = addrspacecast double addrspace(4)* %Arg to double addrspace(1)* -+ ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_64]] -+ ; 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 -+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 -+ -+; 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 } -+ -+!llvm.module.flags = !{!0} -+!opencl.spir.version = !{!1} -+!spirv.Source = !{!2} -+!llvm.ident = !{!3} -+ -+!0 = !{i32 1, !"wchar_size", i32 4} -+!1 = !{i32 1, i32 2} -+!2 = !{i32 4, i32 100000} -+!3 = !{!"clang version 13.0.0"} -+ -diff --git a/test/AtomicFMaxEXTForOCL.ll b/test/AtomicFMaxEXTForOCL.ll -new file mode 100644 -index 00000000..2a5e947f ---- /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 norecurse nounwind -+define dso_local 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 norecurse nounwind -+define dso_local 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 norecurse nounwind "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"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } -+attributes #1 = { convergent "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"="true" "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 12.0.1 (https://github.com/llvm/llvm-project.git 23fe7b104a0adaaaecd52108105f49297c420c9b)"} -diff --git a/test/AtomicFMinEXT.ll b/test/AtomicFMinEXT.ll -new file mode 100644 -index 00000000..3322836d ---- /dev/null -+++ b/test/AtomicFMinEXT.ll -@@ -0,0 +1,73 @@ -+; 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 -r %t.spv -o %t.rev.bc -+; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL12 -+ -+; 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-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 -+; 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 norecurse nounwind -+define dso_local spir_func float @_Z14AtomicFloatMinRf(float addrspace(4)* align 4 dereferenceable(4) %Arg) local_unnamed_addr #0 { -+entry: -+ %0 = addrspacecast float addrspace(4)* %Arg to float addrspace(1)* -+ ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_32]] -+ ; 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 -+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 nounwind -+define dso_local spir_func double @_Z15AtomicDoubleMinRd(double addrspace(4)* align 8 dereferenceable(8) %Arg) local_unnamed_addr #0 { -+entry: -+ %0 = addrspacecast double addrspace(4)* %Arg to double addrspace(1)* -+ ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_64]] -+ ; 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 -+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 -+ -+; 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 } -+ -+!llvm.module.flags = !{!0} -+!opencl.spir.version = !{!1} -+!spirv.Source = !{!2} -+!llvm.ident = !{!3} -+ -+!0 = !{i32 1, !"wchar_size", i32 4} -+!1 = !{i32 1, i32 2} -+!2 = !{i32 4, i32 100000} -+!3 = !{!"clang version 13.0.0"} -+ -diff --git a/test/AtomicFMinEXTForOCL.ll b/test/AtomicFMinEXTForOCL.ll -new file mode 100644 -index 00000000..24a9ce5d ---- /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 norecurse nounwind -+define dso_local 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 norecurse nounwind -+define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 { -+entry: -+ %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]] -+ %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]] -+ ; 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]]({{.*}}) -+ %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 -+ ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_64]] -+ 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 "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"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } -+attributes #1 = { convergent "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"="true" "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 12.0.1 (https://github.com/llvm/llvm-project.git 23fe7b104a0adaaaecd52108105f49297c420c9b)"} -diff --git a/test/negative/InvalidAtomicBuiltins.cl b/test/negative/InvalidAtomicBuiltins.cl -index b8ec5b89..04d71665 100644 ---- a/test/negative/InvalidAtomicBuiltins.cl -+++ b/test/negative/InvalidAtomicBuiltins.cl -@@ -1,6 +1,8 @@ - // Check that translator doesn't generate atomic instructions for atomic builtins - // which are not defined in the spec. - -+// To drop `fdeclare-opencl-builtins` option, since FP-typed atomic function -+// TableGen definitions have not been introduced. - // RUN: %clang_cc1 -triple spir -O1 -cl-std=cl2.0 -fdeclare-opencl-builtins -finclude-default-header %s -emit-llvm-bc -o %t.bc - // RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s - // RUN: llvm-spirv %t.bc -o %t.spv -@@ -34,20 +36,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 +73,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 -