Skip to content

Commit 4b47f7f

Browse files
haonanya1vmaksimo
haonanya1
authored and
vmaksimo
committed
Add support for cl_ext_float_atomics in SPIRVWriter
See details for KhronosGroup/OpenCL-Docs#552 Signed-off-by: Haonan Yang <haonan.yang@intel.com> Original commit: KhronosGroup/SPIRV-LLVM-Translator@89ecd25
1 parent 4e560b7 commit 4b47f7f

File tree

7 files changed

+318
-24
lines changed

7 files changed

+318
-24
lines changed

‎llvm-spirv/lib/SPIRV/OCLToSPIRV.cpp

Lines changed: 23 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -834,7 +834,7 @@ void OCLToSPIRVBase::transAtomicBuiltin(CallInst *CI,
834834
AttributeList Attrs = CI->getCalledFunction()->getAttributes();
835835
mutateCallInstSPIRV(
836836
M, CI,
837-
[=](CallInst *CI, std::vector<Value *> &Args) {
837+
[=](CallInst *CI, std::vector<Value *> &Args) -> std::string {
838838
Info.PostProc(Args);
839839
// Order of args in OCL20:
840840
// object, 0-2 other args, 1-2 order, scope
@@ -863,7 +863,28 @@ void OCLToSPIRVBase::transAtomicBuiltin(CallInst *CI,
863863
std::rotate(Args.begin() + 2, Args.begin() + OrderIdx,
864864
Args.end() - Offset);
865865
}
866-
return getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName));
866+
llvm::Type *AtomicBuiltinsReturnType =
867+
CI->getCalledFunction()->getReturnType();
868+
auto IsFPType = [](llvm::Type *ReturnType) {
869+
return ReturnType->isHalfTy() || ReturnType->isFloatTy() ||
870+
ReturnType->isDoubleTy();
871+
};
872+
auto SPIRVFunctionName =
873+
getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName));
874+
if (!IsFPType(AtomicBuiltinsReturnType))
875+
return SPIRVFunctionName;
876+
// Translate FP-typed atomic builtins. Currently we only need to
877+
// translate atomic_fetch_[add, max, min] and atomic_fetch_[add, max,
878+
// min]_explicit to related float instructions
879+
auto SPIRFunctionNameForFloatAtomics =
880+
llvm::StringSwitch<std::string>(SPIRVFunctionName)
881+
.Case("__spirv_AtomicIAdd", "__spirv_AtomicFAddEXT")
882+
.Case("__spirv_AtomicSMax", "__spirv_AtomicFMaxEXT")
883+
.Case("__spirv_AtomicSMin", "__spirv_AtomicFMinEXT")
884+
.Default("others");
885+
return SPIRFunctionNameForFloatAtomics == "others"
886+
? SPIRVFunctionName
887+
: SPIRFunctionNameForFloatAtomics;
867888
},
868889
&Attrs);
869890
}

‎llvm-spirv/lib/SPIRV/OCLUtil.cpp

Lines changed: 9 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -662,29 +662,32 @@ size_t getSPIRVAtomicBuiltinNumMemoryOrderArgs(Op OC) {
662662
return 1;
663663
}
664664

665+
// atomic_fetch_[add, min, max] and atomic_fetch_[add, min, max]_explicit
666+
// functions declared in clang headers should be translated to corresponding
667+
// FP-typed Atomic Instructions
665668
bool isComputeAtomicOCLBuiltin(StringRef DemangledName) {
666669
if (!DemangledName.startswith(kOCLBuiltinName::AtomicPrefix) &&
667670
!DemangledName.startswith(kOCLBuiltinName::AtomPrefix))
668671
return false;
669672

670673
return llvm::StringSwitch<bool>(DemangledName)
671-
.EndsWith("add", true)
672674
.EndsWith("sub", true)
675+
.EndsWith("atomic_add", true)
676+
.EndsWith("atomic_min", true)
677+
.EndsWith("atomic_max", true)
678+
.EndsWith("atom_add", true)
679+
.EndsWith("atom_min", true)
680+
.EndsWith("atom_max", true)
673681
.EndsWith("inc", true)
674682
.EndsWith("dec", true)
675683
.EndsWith("cmpxchg", true)
676-
.EndsWith("min", true)
677-
.EndsWith("max", true)
678684
.EndsWith("and", true)
679685
.EndsWith("or", true)
680686
.EndsWith("xor", true)
681-
.EndsWith("add_explicit", true)
682687
.EndsWith("sub_explicit", true)
683688
.EndsWith("or_explicit", true)
684689
.EndsWith("xor_explicit", true)
685690
.EndsWith("and_explicit", true)
686-
.EndsWith("min_explicit", true)
687-
.EndsWith("max_explicit", true)
688691
.Default(false);
689692
}
690693

‎llvm-spirv/test/AtomicBuiltinsFloat.ll

Lines changed: 94 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,94 @@
1+
; Check that translator generate atomic instructions for atomic builtins
2+
; FP-typed atomic_fetch_sub and atomic_fetch_sub_explicit should be translated
3+
; to FunctionCall
4+
; RUN: llvm-as %s -o %t.bc
5+
; RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s
6+
; RUN: llvm-spirv %t.bc -o %t.spv
7+
; RUN: spirv-val %t.spv
8+
9+
; CHECK-LABEL: Label
10+
; CHECK: Store
11+
; CHECK-COUNT-3: AtomicStore
12+
; CHECK-COUNT-3: AtomicLoad
13+
; CHECK-COUNT-3: AtomicExchange
14+
; CHECK-COUNT-3: FunctionCall
15+
16+
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"
17+
target triple = "spir-unknown-unknown"
18+
19+
; Function Attrs: convergent norecurse nounwind
20+
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 {
21+
entry:
22+
%0 = addrspacecast float addrspace(3)* %ff to float addrspace(4)*
23+
tail call spir_func void @_Z11atomic_initPU3AS4VU7_Atomicff(float addrspace(4)* %0, float 1.000000e+00) #2
24+
tail call spir_func void @_Z12atomic_storePU3AS4VU7_Atomicff(float addrspace(4)* %0, float 1.000000e+00) #2
25+
tail call spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicff12memory_order(float addrspace(4)* %0, float 1.000000e+00, i32 0) #2
26+
tail call spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicff12memory_order12memory_scope(float addrspace(4)* %0, float 1.000000e+00, i32 0, i32 1) #2
27+
%call = tail call spir_func float @_Z11atomic_loadPU3AS4VU7_Atomicf(float addrspace(4)* %0) #2
28+
%call1 = tail call spir_func float @_Z20atomic_load_explicitPU3AS4VU7_Atomicf12memory_order(float addrspace(4)* %0, i32 0) #2
29+
%call2 = tail call spir_func float @_Z20atomic_load_explicitPU3AS4VU7_Atomicf12memory_order12memory_scope(float addrspace(4)* %0, i32 0, i32 1) #2
30+
%call3 = tail call spir_func float @_Z15atomic_exchangePU3AS4VU7_Atomicff(float addrspace(4)* %0, float 1.000000e+00) #2
31+
%call4 = tail call spir_func float @_Z24atomic_exchange_explicitPU3AS4VU7_Atomicff12memory_order(float addrspace(4)* %0, float 1.000000e+00, i32 0) #2
32+
%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
33+
%call6 = tail call spir_func float @_Z16atomic_fetch_subPU3AS3VU7_Atomicff(float addrspace(3)* %ff, float 1.000000e+00) #2
34+
%call7 = tail call spir_func float @_Z25atomic_fetch_sub_explicitPU3AS3VU7_Atomicff12memory_order(float addrspace(3)* %ff, float 1.000000e+00, i32 0) #2
35+
%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
36+
ret void
37+
}
38+
39+
; Function Attrs: convergent
40+
declare spir_func void @_Z11atomic_initPU3AS4VU7_Atomicff(float addrspace(4)*, float) local_unnamed_addr #1
41+
42+
; Function Attrs: convergent
43+
declare spir_func void @_Z12atomic_storePU3AS4VU7_Atomicff(float addrspace(4)*, float) local_unnamed_addr #1
44+
45+
; Function Attrs: convergent
46+
declare spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicff12memory_order(float addrspace(4)*, float, i32) local_unnamed_addr #1
47+
48+
; Function Attrs: convergent
49+
declare spir_func void @_Z21atomic_store_explicitPU3AS4VU7_Atomicff12memory_order12memory_scope(float addrspace(4)*, float, i32, i32) local_unnamed_addr #1
50+
51+
; Function Attrs: convergent
52+
declare spir_func float @_Z11atomic_loadPU3AS4VU7_Atomicf(float addrspace(4)*) local_unnamed_addr #1
53+
54+
; Function Attrs: convergent
55+
declare spir_func float @_Z20atomic_load_explicitPU3AS4VU7_Atomicf12memory_order(float addrspace(4)*, i32) local_unnamed_addr #1
56+
57+
; Function Attrs: convergent
58+
declare spir_func float @_Z20atomic_load_explicitPU3AS4VU7_Atomicf12memory_order12memory_scope(float addrspace(4)*, i32, i32) local_unnamed_addr #1
59+
60+
; Function Attrs: convergent
61+
declare spir_func float @_Z15atomic_exchangePU3AS4VU7_Atomicff(float addrspace(4)*, float) local_unnamed_addr #1
62+
63+
; Function Attrs: convergent
64+
declare spir_func float @_Z24atomic_exchange_explicitPU3AS4VU7_Atomicff12memory_order(float addrspace(4)*, float, i32) local_unnamed_addr #1
65+
66+
; Function Attrs: convergent
67+
declare spir_func float @_Z24atomic_exchange_explicitPU3AS4VU7_Atomicff12memory_order12memory_scope(float addrspace(4)*, float, i32, i32) local_unnamed_addr #1
68+
69+
; Function Attrs: convergent
70+
declare spir_func float @_Z16atomic_fetch_subPU3AS3VU7_Atomicff(float addrspace(3)*, float) local_unnamed_addr #1
71+
72+
; Function Attrs: convergent
73+
declare spir_func float @_Z25atomic_fetch_sub_explicitPU3AS3VU7_Atomicff12memory_order(float addrspace(3)*, float, i32) local_unnamed_addr #1
74+
75+
; Function Attrs: convergent
76+
declare spir_func float @_Z25atomic_fetch_sub_explicitPU3AS3VU7_Atomicff12memory_order12memory_scope(float addrspace(3)*, float, i32, i32) local_unnamed_addr #1
77+
78+
attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="false" }
79+
attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
80+
attributes #2 = { convergent nounwind }
81+
82+
!llvm.module.flags = !{!0}
83+
!opencl.ocl.version = !{!1}
84+
!opencl.spir.version = !{!1}
85+
!llvm.ident = !{!2}
86+
87+
!0 = !{i32 1, !"wchar_size", i32 4}
88+
!1 = !{i32 2, i32 0}
89+
!2 = !{!"clang version 14.0.0 (https://github.com/llvm/llvm-project.git 28c4f97a1dc8608cdd4db452b73d7d4afc89acc9)"}
90+
!3 = !{i32 3}
91+
!4 = !{!"none"}
92+
!5 = !{!"atomic_float*"}
93+
!6 = !{!"_Atomic(float)*"}
94+
!7 = !{!"volatile"}

‎llvm-spirv/test/negative/InvalidAtomicBuiltins.cl

Lines changed: 0 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -34,20 +34,12 @@ double __attribute__((overloadable)) atom_and(volatile __global double *p, doubl
3434
double __attribute__((overloadable)) atom_or(volatile __global double *p, double val);
3535
double __attribute__((overloadable)) atom_xor(volatile __global double *p, double val);
3636

37-
float __attribute__((overloadable)) atomic_fetch_add(volatile generic atomic_float *object, float operand, memory_order order);
38-
float __attribute__((overloadable)) atomic_fetch_sub(volatile generic atomic_float *object, float operand, memory_order order);
3937
float __attribute__((overloadable)) atomic_fetch_or(volatile generic atomic_float *object, float operand, memory_order order);
4038
float __attribute__((overloadable)) atomic_fetch_xor(volatile generic atomic_float *object, float operand, memory_order order);
4139
double __attribute__((overloadable)) atomic_fetch_and(volatile generic atomic_double *object, double operand, memory_order order);
42-
double __attribute__((overloadable)) atomic_fetch_max(volatile generic atomic_double *object, double operand, memory_order order);
43-
double __attribute__((overloadable)) atomic_fetch_min(volatile generic atomic_double *object, double operand, memory_order order);
44-
float __attribute__((overloadable)) atomic_fetch_add_explicit(volatile generic atomic_float *object, float operand, memory_order order);
45-
float __attribute__((overloadable)) atomic_fetch_sub_explicit(volatile generic atomic_float *object, float operand, memory_order order);
4640
float __attribute__((overloadable)) atomic_fetch_or_explicit(volatile generic atomic_float *object, float operand, memory_order order);
4741
float __attribute__((overloadable)) atomic_fetch_xor_explicit(volatile generic atomic_float *object, float operand, memory_order order);
4842
double __attribute__((overloadable)) atomic_fetch_and_explicit(volatile generic atomic_double *object, double operand, memory_order order);
49-
double __attribute__((overloadable)) atomic_fetch_max_explicit(volatile generic atomic_double *object, double operand, memory_order order);
50-
double __attribute__((overloadable)) atomic_fetch_min_explicit(volatile generic atomic_double *object, double operand, memory_order order);
5143

5244
__kernel void test_atomic_fn(volatile __global float *p,
5345
volatile __global double *pp,
@@ -79,18 +71,10 @@ __kernel void test_atomic_fn(volatile __global float *p,
7971
d = atom_or(pp, val);
8072
d = atom_xor(pp, val);
8173

82-
f = atomic_fetch_add(p, val, order);
83-
f = atomic_fetch_sub(p, val, order);
8474
f = atomic_fetch_or(p, val, order);
8575
f = atomic_fetch_xor(p, val, order);
8676
d = atomic_fetch_and(pp, val, order);
87-
d = atomic_fetch_min(pp, val, order);
88-
d = atomic_fetch_max(pp, val, order);
89-
f = atomic_fetch_add_explicit(p, val, order);
90-
f = atomic_fetch_sub_explicit(p, val, order);
9177
f = atomic_fetch_or_explicit(p, val, order);
9278
f = atomic_fetch_xor_explicit(p, val, order);
9379
d = atomic_fetch_and_explicit(pp, val, order);
94-
d = atomic_fetch_min_explicit(pp, val, order);
95-
d = atomic_fetch_max_explicit(pp, val, order);
9680
}

‎llvm-spirv/test/transcoding/AtomicFAddEXTForOCL.ll

Lines changed: 64 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,64 @@
1+
; RUN: llvm-as %s -o %t.bc
2+
; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_add -o %t.spv
3+
; RUN: spirv-val %t.spv
4+
; RUN: llvm-spirv -to-text %t.spv -o %t.spt
5+
; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
6+
7+
; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc
8+
; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL20
9+
10+
; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc
11+
; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV
12+
13+
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"
14+
target triple = "spir-unknown-unknown"
15+
16+
; CHECK-SPIRV: Capability AtomicFloat32AddEXT
17+
; CHECK-SPIRV: Capability AtomicFloat64AddEXT
18+
; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_add"
19+
; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
20+
; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
21+
22+
23+
; Function Attrs: convergent norecurse nounwind
24+
define dso_local spir_func void @test_atomic_float(float addrspace(1)* %a) local_unnamed_addr #0 {
25+
entry:
26+
; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]]
27+
; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}})
28+
; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+fiif]]({{.*}})
29+
%call = tail call spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2
30+
ret void
31+
}
32+
33+
; Function Attrs: convergent
34+
declare spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1
35+
; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
36+
37+
; Function Attrs: convergent norecurse nounwind
38+
define dso_local spir_func void @test_atomic_double(double addrspace(1)* %a) local_unnamed_addr #0 {
39+
entry:
40+
; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]]
41+
; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}})
42+
; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+diid]]({{.*}})
43+
%call = tail call spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2
44+
ret void
45+
}
46+
; Function Attrs: convergent
47+
declare spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1
48+
; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
49+
50+
; CHECK-LLVM-CL20: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
51+
; CHECK-LLVM-CL20: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
52+
53+
attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
54+
attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
55+
attributes #2 = { convergent nounwind }
56+
57+
!llvm.module.flags = !{!0}
58+
!opencl.ocl.version = !{!1}
59+
!opencl.spir.version = !{!1}
60+
!llvm.ident = !{!2}
61+
62+
!0 = !{i32 1, !"wchar_size", i32 4}
63+
!1 = !{i32 2, i32 0}
64+
!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"}

‎llvm-spirv/test/transcoding/AtomicFMaxEXTForOCL.ll

Lines changed: 64 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,64 @@
1+
; RUN: llvm-as %s -o %t.bc
2+
; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_min_max -o %t.spv
3+
; RUN: spirv-val %t.spv
4+
; RUN: llvm-spirv -to-text %t.spv -o %t.spt
5+
; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV
6+
7+
; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc
8+
; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL20
9+
10+
; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc
11+
; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV
12+
13+
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"
14+
target triple = "spir-unknown-unknown"
15+
16+
; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT
17+
; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT
18+
; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_min_max"
19+
; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32
20+
; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64
21+
22+
; Function Attrs: convergent norecurse nounwind
23+
define dso_local spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 {
24+
entry:
25+
; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_32]]
26+
; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}})
27+
; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+fiif]]({{.*}})
28+
%call = tail call spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2
29+
ret void
30+
}
31+
32+
; Function Attrs: convergent
33+
declare spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1
34+
; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
35+
36+
; Function Attrs: convergent norecurse nounwind
37+
define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 {
38+
entry:
39+
; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_64]]
40+
; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}})
41+
; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+diid]]({{.*}})
42+
%call = tail call spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2
43+
ret void
44+
}
45+
46+
; Function Attrs: convergent
47+
declare spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1
48+
; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
49+
50+
; CHECK-LLVM-CL20: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float
51+
; CHECK-LLVM-CL20: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double
52+
53+
attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
54+
attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" }
55+
attributes #2 = { convergent nounwind }
56+
57+
!llvm.module.flags = !{!0}
58+
!opencl.ocl.version = !{!1}
59+
!opencl.spir.version = !{!1}
60+
!llvm.ident = !{!2}
61+
62+
!0 = !{i32 1, !"wchar_size", i32 4}
63+
!1 = !{i32 2, i32 0}
64+
!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"}

0 commit comments

Comments
 (0)