From 8e12d8fb3cdbdafca73fe8ed4f0cde773b1788b4 Mon Sep 17 00:00:00 2001 From: haonanya Date: Wed, 28 Jul 2021 11:43:20 +0800 Subject: [PATCH 2/2] Add support for cl_ext_float_atomics in SPIRVWriter Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-110/patches/spirv/0001-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch] Signed-off-by: haonanya Signed-off-by: Naveen Saini --- lib/SPIRV/OCLToSPIRV.cpp | 80 +++++++++++++++++++++++-- lib/SPIRV/OCLUtil.cpp | 26 -------- lib/SPIRV/OCLUtil.h | 4 -- test/negative/InvalidAtomicBuiltins.cl | 12 +--- test/transcoding/AtomicFAddEXTForOCL.ll | 64 ++++++++++++++++++++ test/transcoding/AtomicFMaxEXTForOCL.ll | 64 ++++++++++++++++++++ test/transcoding/AtomicFMinEXTForOCL.ll | 64 ++++++++++++++++++++ 7 files changed, 269 insertions(+), 45 deletions(-) create mode 100644 test/transcoding/AtomicFAddEXTForOCL.ll create mode 100644 test/transcoding/AtomicFMaxEXTForOCL.ll create mode 100644 test/transcoding/AtomicFMinEXTForOCL.ll diff --git a/lib/SPIRV/OCLToSPIRV.cpp b/lib/SPIRV/OCLToSPIRV.cpp index 04d51586..f00f5f7b 100644 --- a/lib/SPIRV/OCLToSPIRV.cpp +++ b/lib/SPIRV/OCLToSPIRV.cpp @@ -421,10 +421,63 @@ void OCLToSPIRVBase::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)) - return; + // Compute "atom" prefixed builtins do not support floating types. + if (CI.getType()->isFloatingPointTy()) { + if (DemangledName.find(kOCLBuiltinName::AtomPrefix) == 0) + return; + // handle functions which are "atomic_" prefixed. + StringRef Stem = DemangledName; + Stem = Stem.drop_front(strlen("atomic_")); + // FP-typed atomic_{add, sub, inc, dec, exchange, min, max, or, and, xor, + // fetch_or, fetch_xor, fetch_and, fetch_or_explicit, fetch_xor_explicit, + // fetch_and_explicit} should be identified as function call + bool IsFunctionCall = llvm::StringSwitch(Stem) + .Case("add", true) + .Case("sub", true) + .Case("inc", true) + .Case("dec", true) + .Case("cmpxchg", true) + .Case("min", true) + .Case("max", true) + .Case("or", true) + .Case("xor", true) + .Case("and", true) + .Case("fetch_or", true) + .Case("fetch_and", true) + .Case("fetch_xor", true) + .Case("fetch_or_explicit", true) + .Case("fetch_xor_explicit", true) + .Case("fetch_and_explicit", true) + .Default(false); + if (IsFunctionCall) + return; + if (F->arg_size() != 2) { + IsFunctionCall = llvm::StringSwitch(Stem) + .Case("exchange", true) + .Case("fetch_add", true) + .Case("fetch_sub", true) + .Case("fetch_min", true) + .Case("fetch_max", true) + .Case("load", true) + .Case("store", true) + .Default(false); + if (IsFunctionCall) + return; + } + if (F->arg_size() != 3 && F->arg_size() != 4) { + IsFunctionCall = llvm::StringSwitch(Stem) + .Case("exchange_explicit", true) + .Case("fetch_add_explicit", true) + .Case("fetch_sub_explicit", true) + .Case("fetch_min_explicit", true) + .Case("fetch_max_explicit", true) + .Case("load_explicit", true) + .Case("store_explicit", true) + .Default(false); + if (IsFunctionCall) + return; + } + } auto PCI = &CI; if (DemangledName == kOCLBuiltinName::AtomicInit) { @@ -839,7 +892,7 @@ void OCLToSPIRVBase::transAtomicBuiltin(CallInst *CI, 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 @@ -868,7 +921,22 @@ void OCLToSPIRVBase::transAtomicBuiltin(CallInst *CI, 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. + return llvm::StringSwitch(SPIRVFunctionName) + .Case("__spirv_AtomicIAdd", "__spirv_AtomicFAddEXT") + .Case("__spirv_AtomicSMax", "__spirv_AtomicFMaxEXT") + .Case("__spirv_AtomicSMin", "__spirv_AtomicFMinEXT"); }, &Attrs); } diff --git a/lib/SPIRV/OCLUtil.cpp b/lib/SPIRV/OCLUtil.cpp index 2de3f152..85155e39 100644 --- a/lib/SPIRV/OCLUtil.cpp +++ b/lib/SPIRV/OCLUtil.cpp @@ -662,32 +662,6 @@ size_t getSPIRVAtomicBuiltinNumMemoryOrderArgs(Op OC) { return 1; } -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("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); -} - BarrierLiterals getBarrierLiterals(CallInst *CI) { auto N = CI->getNumArgOperands(); assert(N == 1 || N == 2); diff --git a/lib/SPIRV/OCLUtil.h b/lib/SPIRV/OCLUtil.h index 4c05c672..c8577e9b 100644 --- a/lib/SPIRV/OCLUtil.h +++ b/lib/SPIRV/OCLUtil.h @@ -394,10 +394,6 @@ size_t getAtomicBuiltinNumMemoryOrderArgs(StringRef Name); /// Get number of memory order arguments for spirv atomic builtin function. size_t getSPIRVAtomicBuiltinNumMemoryOrderArgs(Op OC); -/// Return true for OpenCL builtins which do compute operations -/// (like add, sub, min, max, inc, dec, ...) atomically -bool isComputeAtomicOCLBuiltin(StringRef DemangledName); - /// Get OCL version from metadata opencl.ocl.version. /// \param AllowMulti Allows multiple operands if true. /// \return OCL version encoded as Major*10^5+Minor*10^3+Rev, diff --git a/test/negative/InvalidAtomicBuiltins.cl b/test/negative/InvalidAtomicBuiltins.cl index b8ec5b89..23dcc4e3 100644 --- a/test/negative/InvalidAtomicBuiltins.cl +++ b/test/negative/InvalidAtomicBuiltins.cl @@ -1,7 +1,9 @@ // Check that translator doesn't generate atomic instructions for atomic builtins // which are not defined in the spec. -// RUN: %clang_cc1 -triple spir -O1 -cl-std=cl2.0 -fdeclare-opencl-builtins -finclude-default-header %s -emit-llvm-bc -o %t.bc +// 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 -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 // RUN: spirv-val %t.spv @@ -41,13 +43,9 @@ float __attribute__((overloadable)) atomic_fetch_xor(volatile generic atomic_flo 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, @@ -86,11 +84,7 @@ __kernel void test_atomic_fn(volatile __global float *p, 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); } diff --git a/test/transcoding/AtomicFAddEXTForOCL.ll b/test/transcoding/AtomicFAddEXTForOCL.ll new file mode 100644 index 00000000..fb146fb9 --- /dev/null +++ b/test/transcoding/AtomicFAddEXTForOCL.ll @@ -0,0 +1,64 @@ +; 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_atomic_float(float addrspace(1)* %a) local_unnamed_addr #0 { +entry: + ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]] + ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}}) + ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+fiif]]({{.*}}) + %call = tail call spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2 + ret void +} + +; Function Attrs: convergent +declare spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1 +; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float + +; Function Attrs: convergent norecurse nounwind +define dso_local spir_func void @test_atomic_double(double addrspace(1)* %a) local_unnamed_addr #0 { +entry: + ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]] + ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}}) + ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+diid]]({{.*}}) + %call = tail call spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2 + ret void +} +; 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 + +; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float +; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double + +attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #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 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"} diff --git a/test/transcoding/AtomicFMaxEXTForOCL.ll b/test/transcoding/AtomicFMaxEXTForOCL.ll new file mode 100644 index 00000000..1f2530d9 --- /dev/null +++ b/test/transcoding/AtomicFMaxEXTForOCL.ll @@ -0,0 +1,64 @@ +; 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]] + ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}}) + ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+fiif]]({{.*}}) + %call = tail call spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2 + ret void +} + +; Function Attrs: convergent +declare spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1 +; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float + +; Function Attrs: convergent norecurse nounwind +define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 { +entry: + ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_64]] + ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}}) + ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+diid]]({{.*}}) + %call = tail call spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2 + ret void +} + +; 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 + +; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float +; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double + +attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #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 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"} diff --git a/test/transcoding/AtomicFMinEXTForOCL.ll b/test/transcoding/AtomicFMinEXTForOCL.ll new file mode 100644 index 00000000..6196b0f8 --- /dev/null +++ b/test/transcoding/AtomicFMinEXTForOCL.ll @@ -0,0 +1,64 @@ +; 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]] + ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}}) + ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+fiif]]({{.*}}) + %call = tail call spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2 + ret void +} + +; Function Attrs: convergent +declare spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1 +; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float + +; Function Attrs: convergent norecurse nounwind +define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 { +entry: + ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_64]] + ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}}) + ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+diid]]({{.*}}) + %call = tail call spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2 + ret void +} + +; 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 + +; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float +; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double + +attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } +attributes #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 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"} -- 2.17.1