diff options
Diffstat (limited to 'dynamic-layers/clang-layer')
12 files changed, 1524 insertions, 144 deletions
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/0001-Fix-debug-info-of-work-item-builtin-translation-745.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/0001-Fix-debug-info-of-work-item-builtin-translation-745.patch deleted file mode 100644 index 923b871f..00000000 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/0001-Fix-debug-info-of-work-item-builtin-translation-745.patch +++ /dev/null | |||
@@ -1,119 +0,0 @@ | |||
1 | From 200c200eb19602ffd7c8f29d0b2df9df1fd311bf Mon Sep 17 00:00:00 2001 | ||
2 | From: Naveen Saini <naveen.kumar.saini@intel.com> | ||
3 | Date: Wed, 7 Apr 2021 17:44:20 +0800 | ||
4 | Subject: [PATCH] Fix debug info of work-item builtin translation (#745) | ||
5 | |||
6 | debug info of work-item builtins are lost in both llvm IR -> spirv and | ||
7 | spirv -> llvm IR translations. See #744 | ||
8 | |||
9 | Upstream-Status: Backport [https://github.com/KhronosGroup/SPIRV-LLVM-Translator/commit/c734c5c8bbd3012a09c610e4be68e90cc603c580] | ||
10 | Signed-off-by: Wenju He <wenju.he@intel.com> | ||
11 | Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com> | ||
12 | --- | ||
13 | lib/SPIRV/OCL20ToSPIRV.cpp | 5 ++- | ||
14 | lib/SPIRV/SPIRVReader.cpp | 1 + | ||
15 | test/DebugInfo/builtin-get-global-id.ll | 60 +++++++++++++++++++++++++ | ||
16 | 3 files changed, 65 insertions(+), 1 deletion(-) | ||
17 | create mode 100644 test/DebugInfo/builtin-get-global-id.ll | ||
18 | |||
19 | diff --git a/lib/SPIRV/OCL20ToSPIRV.cpp b/lib/SPIRV/OCL20ToSPIRV.cpp | ||
20 | index 1262c48c..a742c8cf 100644 | ||
21 | --- a/lib/SPIRV/OCL20ToSPIRV.cpp | ||
22 | +++ b/lib/SPIRV/OCL20ToSPIRV.cpp | ||
23 | @@ -1297,11 +1297,14 @@ void OCL20ToSPIRV::transWorkItemBuiltinsToVariables() { | ||
24 | for (auto UI = I.user_begin(), UE = I.user_end(); UI != UE; ++UI) { | ||
25 | auto CI = dyn_cast<CallInst>(*UI); | ||
26 | assert(CI && "invalid instruction"); | ||
27 | - Value *NewValue = new LoadInst(BV, "", CI); | ||
28 | + const DebugLoc &DLoc = CI->getDebugLoc(); | ||
29 | + Instruction *NewValue = new LoadInst(BV, "", CI); | ||
30 | + NewValue->setDebugLoc(DLoc); | ||
31 | LLVM_DEBUG(dbgs() << "Transform: " << *CI << " => " << *NewValue << '\n'); | ||
32 | if (IsVec) { | ||
33 | NewValue = | ||
34 | ExtractElementInst::Create(NewValue, CI->getArgOperand(0), "", CI); | ||
35 | + NewValue->setDebugLoc(DLoc); | ||
36 | LLVM_DEBUG(dbgs() << *NewValue << '\n'); | ||
37 | } | ||
38 | NewValue->takeName(CI); | ||
39 | diff --git a/lib/SPIRV/SPIRVReader.cpp b/lib/SPIRV/SPIRVReader.cpp | ||
40 | index 16a3dd38..528f6663 100644 | ||
41 | --- a/lib/SPIRV/SPIRVReader.cpp | ||
42 | +++ b/lib/SPIRV/SPIRVReader.cpp | ||
43 | @@ -307,6 +307,7 @@ bool SPIRVToLLVM::transOCLBuiltinFromVariable(GlobalVariable *GV, | ||
44 | auto Replace = [&](std::vector<Value *> Arg, Instruction *I) { | ||
45 | auto Call = CallInst::Create(Func, Arg, "", I); | ||
46 | Call->takeName(I); | ||
47 | + Call->setDebugLoc(I->getDebugLoc()); | ||
48 | setAttrByCalledFunc(Call); | ||
49 | SPIRVDBG(dbgs() << "[transOCLBuiltinFromVariable] " << *I << " -> " << *Call | ||
50 | << '\n';) | ||
51 | diff --git a/test/DebugInfo/builtin-get-global-id.ll b/test/DebugInfo/builtin-get-global-id.ll | ||
52 | new file mode 100644 | ||
53 | index 00000000..a4a00e63 | ||
54 | --- /dev/null | ||
55 | +++ b/test/DebugInfo/builtin-get-global-id.ll | ||
56 | @@ -0,0 +1,60 @@ | ||
57 | +; Check debug info of builtin get_global_id is preserved from LLVM IR to spirv | ||
58 | +; and spirv to LLVM IR translation. | ||
59 | + | ||
60 | +; Original .cl source: | ||
61 | +; kernel void test() { | ||
62 | +; size_t gid = get_global_id(0); | ||
63 | +; } | ||
64 | + | ||
65 | +; Command line: | ||
66 | +; ./clang -cc1 1.cl -triple spir64 -cl-std=cl2.0 -emit-llvm -finclude-default-header -debug-info-kind=line-tables-only -O0 | ||
67 | + | ||
68 | +; RUN: llvm-as %s -o %t.bc | ||
69 | +; RUN: llvm-spirv %t.bc -spirv-text -o - | FileCheck %s --check-prefix CHECK-SPIRV | ||
70 | +; RUN: llvm-spirv %t.bc -o %t.spv | ||
71 | +; RUN: llvm-spirv -r %t.spv -o - | llvm-dis -o - | FileCheck %s | ||
72 | + | ||
73 | +target datalayout = "e-i64:64-v16:16-v24:32-v32:32-v48:64-v96:128-v192:256-v256:256-v512:512-v1024:1024" | ||
74 | +target triple = "spir64" | ||
75 | + | ||
76 | +; CHECK-SPIRV: ExtInst {{.*}} DebugScope | ||
77 | +; CHECK-SPIRV-NEXT: Line {{[0-9]+}} 2 16 | ||
78 | +; CHECK-SPIRV-NEXT: Load {{[0-9]+}} [[LoadRes:[0-9]+]] | ||
79 | +; CHECK-SPIRV-NEXT: CompositeExtract {{[0-9]+}} {{[0-9]+}} [[LoadRes]] 0 | ||
80 | + | ||
81 | +; Function Attrs: convergent noinline norecurse nounwind optnone | ||
82 | +define spir_kernel void @test() #0 !dbg !7 !kernel_arg_addr_space !2 !kernel_arg_access_qual !2 !kernel_arg_type !2 !kernel_arg_base_type !2 !kernel_arg_type_qual !2 { | ||
83 | +entry: | ||
84 | + %gid = alloca i64, align 8 | ||
85 | + %call = call spir_func i64 @_Z13get_global_idj(i32 0) #2, !dbg !10 | ||
86 | +; CHECK: %call = call spir_func i64 @_Z13get_global_idj(i32 0) #1, !dbg [[DBG:![0-9]+]] | ||
87 | + store i64 %call, i64* %gid, align 8, !dbg !11 | ||
88 | + ret void, !dbg !12 | ||
89 | +} | ||
90 | + | ||
91 | +; Function Attrs: convergent nounwind readnone | ||
92 | +declare spir_func i64 @_Z13get_global_idj(i32) #1 | ||
93 | + | ||
94 | +attributes #0 = { convergent noinline norecurse nounwind optnone "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="false" "unsafe-fp-math"="false" "use-soft-float"="false" } | ||
95 | +attributes #1 = { convergent nounwind readnone "correctly-rounded-divide-sqrt-fp-math"="false" "disable-tail-calls"="false" "frame-pointer"="none" "less-precise-fpmad"="false" "no-infs-fp-math"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "unsafe-fp-math"="false" "use-soft-float"="false" } | ||
96 | +attributes #2 = { convergent nounwind readnone } | ||
97 | + | ||
98 | +!llvm.dbg.cu = !{!0} | ||
99 | +!llvm.module.flags = !{!3, !4} | ||
100 | +!opencl.ocl.version = !{!5} | ||
101 | +!opencl.spir.version = !{!5} | ||
102 | +!llvm.ident = !{!6} | ||
103 | + | ||
104 | +!0 = distinct !DICompileUnit(language: DW_LANG_C99, file: !1, producer: "clang version 12.0.0 (https://github.com/llvm/llvm-project.git b5bc56da8aa23dc57db9d286b0591dbcf9b1bdd3)", isOptimized: false, runtimeVersion: 0, emissionKind: LineTablesOnly, enums: !2, nameTableKind: None) | ||
105 | +!1 = !DIFile(filename: "<stdin>", directory: "") | ||
106 | +!2 = !{} | ||
107 | +!3 = !{i32 2, !"Debug Info Version", i32 3} | ||
108 | +!4 = !{i32 1, !"wchar_size", i32 4} | ||
109 | +!5 = !{i32 2, i32 0} | ||
110 | +!6 = !{!"clang version 12.0.0 (https://github.com/llvm/llvm-project.git b5bc56da8aa23dc57db9d286b0591dbcf9b1bdd3)"} | ||
111 | +!7 = distinct !DISubprogram(name: "test", scope: !8, file: !8, line: 1, type: !9, scopeLine: 1, flags: DIFlagPrototyped, spFlags: DISPFlagDefinition, unit: !0, retainedNodes: !2) | ||
112 | +!8 = !DIFile(filename: "1.cl", directory: "") | ||
113 | +!9 = !DISubroutineType(types: !2) | ||
114 | +!10 = !DILocation(line: 2, column: 16, scope: !7) | ||
115 | +!11 = !DILocation(line: 2, column: 10, scope: !7) | ||
116 | +!12 = !DILocation(line: 3, column: 1, scope: !7) | ||
117 | -- | ||
118 | 2.17.1 | ||
119 | |||
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-OpenCL-3.0-support.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0001-OpenCL-3.0-support.patch index 53395ea0..1ab00df0 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-OpenCL-3.0-support.patch +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0001-OpenCL-3.0-support.patch | |||
@@ -1,7 +1,7 @@ | |||
1 | From 31ec702cb365f4d02dd2146fb4329d642b8fc30b Mon Sep 17 00:00:00 2001 | 1 | From 8dbdb2f26674a938ff43b5bfe5b3bf3d1117f9e4 Mon Sep 17 00:00:00 2001 |
2 | From: Naveen Saini <naveen.kumar.saini@intel.com> | 2 | From: Naveen Saini <naveen.kumar.saini@intel.com> |
3 | Date: Wed, 7 Apr 2021 16:36:10 +0800 | 3 | Date: Wed, 7 Apr 2021 16:36:10 +0800 |
4 | Subject: [PATCH 1/2] OpenCL 3.0 support | 4 | Subject: [PATCH 1/7] OpenCL 3.0 support |
5 | 5 | ||
6 | Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-100/patches/clang/0001-OpenCL-3.0-support.patch] | 6 | Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-100/patches/clang/0001-OpenCL-3.0-support.patch] |
7 | Signed-off-by: Anton Zabaznov <anton.zabaznov@intel.com> | 7 | Signed-off-by: Anton Zabaznov <anton.zabaznov@intel.com> |
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-skip-building-tests.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0001-llvm-spirv-skip-building-tests.patch index 8e58ec25..84a4ba19 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-skip-building-tests.patch +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0001-llvm-spirv-skip-building-tests.patch | |||
@@ -1,7 +1,7 @@ | |||
1 | From 455ce9c25df5313f4a6649cc27075bdfbe25af18 Mon Sep 17 00:00:00 2001 | 1 | From 661021749a168c423d69d0ba7cdfa16fed860836 Mon Sep 17 00:00:00 2001 |
2 | From: Naveen Saini <naveen.kumar.saini@intel.com> | 2 | From: Naveen Saini <naveen.kumar.saini@intel.com> |
3 | Date: Wed, 21 Aug 2019 14:35:31 +0800 | 3 | Date: Wed, 21 Aug 2019 14:35:31 +0800 |
4 | Subject: [PATCH] llvm-spirv: skip building tests | 4 | Subject: [PATCH 1/3] llvm-spirv: skip building tests |
5 | 5 | ||
6 | Some of these need clang to be built and since we're building this in-tree, | 6 | Some of these need clang to be built and since we're building this in-tree, |
7 | that leads to problems when compiling libcxx, compiler-rt which aren't built | 7 | that leads to problems when compiling libcxx, compiler-rt which aren't built |
@@ -19,10 +19,10 @@ Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com> | |||
19 | 1 file changed, 10 deletions(-) | 19 | 1 file changed, 10 deletions(-) |
20 | 20 | ||
21 | diff --git a/CMakeLists.txt b/CMakeLists.txt | 21 | diff --git a/CMakeLists.txt b/CMakeLists.txt |
22 | index b718c00..9805140 100644 | 22 | index 92c50370..80999c98 100644 |
23 | --- a/CMakeLists.txt | 23 | --- a/CMakeLists.txt |
24 | +++ b/CMakeLists.txt | 24 | +++ b/CMakeLists.txt |
25 | @@ -24,13 +24,6 @@ if(LLVM_SPIRV_BUILD_EXTERNAL) | 25 | @@ -25,13 +25,6 @@ if(LLVM_SPIRV_BUILD_EXTERNAL) |
26 | set(CMAKE_CXX_STANDARD 14) | 26 | set(CMAKE_CXX_STANDARD 14) |
27 | set(CMAKE_CXX_STANDARD_REQUIRED ON) | 27 | set(CMAKE_CXX_STANDARD_REQUIRED ON) |
28 | 28 | ||
@@ -36,7 +36,7 @@ index b718c00..9805140 100644 | |||
36 | find_package(LLVM 10.0.0 REQUIRED | 36 | find_package(LLVM 10.0.0 REQUIRED |
37 | COMPONENTS | 37 | COMPONENTS |
38 | Analysis | 38 | Analysis |
39 | @@ -61,9 +54,6 @@ set(LLVM_SPIRV_INCLUDE_DIRS ${CMAKE_CURRENT_SOURCE_DIR}/include) | 39 | @@ -63,9 +56,6 @@ set(LLVM_SPIRV_INCLUDE_DIRS ${CMAKE_CURRENT_SOURCE_DIR}/include) |
40 | 40 | ||
41 | add_subdirectory(lib/SPIRV) | 41 | add_subdirectory(lib/SPIRV) |
42 | add_subdirectory(tools/llvm-spirv) | 42 | add_subdirectory(tools/llvm-spirv) |
@@ -47,5 +47,5 @@ index b718c00..9805140 100644 | |||
47 | install( | 47 | install( |
48 | FILES | 48 | FILES |
49 | -- | 49 | -- |
50 | 2.7.4 | 50 | 2.17.1 |
51 | 51 | ||
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/0002-Add-cl_khr_extended_subgroup-extensions.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0002-Add-cl_khr_extended_subgroup-extensions.patch index cbe492c4..3f1b24e7 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/0002-Add-cl_khr_extended_subgroup-extensions.patch +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0002-Add-cl_khr_extended_subgroup-extensions.patch | |||
@@ -1,7 +1,7 @@ | |||
1 | From 27d47f1a17c8921b07acc8cdc26e38cc609de4a9 Mon Sep 17 00:00:00 2001 | 1 | From 3f544cfe44ee5f113a3fb554aca2cf5d64996062 Mon Sep 17 00:00:00 2001 |
2 | From: Naveen Saini <naveen.kumar.saini@intel.com> | 2 | From: Naveen Saini <naveen.kumar.saini@intel.com> |
3 | Date: Wed, 7 Apr 2021 16:38:38 +0800 | 3 | Date: Wed, 7 Apr 2021 16:38:38 +0800 |
4 | Subject: [PATCH 2/2] Add cl_khr_extended_subgroup extensions. | 4 | Subject: [PATCH 2/7] Add cl_khr_extended_subgroup extensions. |
5 | 5 | ||
6 | Added extensions and their function declarations into | 6 | Added extensions and their function declarations into |
7 | the standard header. | 7 | the standard header. |
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/fix-shared-libs.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0002-Fix-building-in-tree-with-cmake-DLLVM_LINK_LLVM_DYLI.patch index d69d2a97..1aff65e7 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/fix-shared-libs.patch +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0002-Fix-building-in-tree-with-cmake-DLLVM_LINK_LLVM_DYLI.patch | |||
@@ -1,7 +1,7 @@ | |||
1 | From a6d4ccf082858e63e139ca06c02a071c343d2657 Mon Sep 17 00:00:00 2001 | 1 | From 331e323ae2633a8999a660314022491d670c442c Mon Sep 17 00:00:00 2001 |
2 | From: Andrea Bocci <andrea.bocci@cern.ch> | 2 | From: Andrea Bocci <andrea.bocci@cern.ch> |
3 | Date: Sun, 15 Mar 2020 17:35:44 +0100 | 3 | Date: Sun, 15 Mar 2020 17:35:44 +0100 |
4 | Subject: [PATCH] Fix building in-tree with cmake -DLLVM_LINK_LLVM_DYLIB=ON | 4 | Subject: [PATCH 2/3] Fix building in-tree with cmake -DLLVM_LINK_LLVM_DYLIB=ON |
5 | 5 | ||
6 | Building in-tree with LLVM 11.0 master with the LLVM_LINK_LLVM_DYLIB | 6 | Building in-tree with LLVM 11.0 master with the LLVM_LINK_LLVM_DYLIB |
7 | cmake flag fails to link with the LLVMSPIRVLib library. | 7 | cmake flag fails to link with the LLVMSPIRVLib library. |
@@ -28,3 +28,6 @@ index 9aa96d9c..501c0daf 100644 | |||
28 | target_link_libraries(llvm-spirv PRIVATE LLVMSPIRVLib) | 28 | target_link_libraries(llvm-spirv PRIVATE LLVMSPIRVLib) |
29 | endif() | 29 | endif() |
30 | 30 | ||
31 | -- | ||
32 | 2.17.1 | ||
33 | |||
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch new file mode 100644 index 00000000..49edd7e1 --- /dev/null +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch | |||
@@ -0,0 +1,982 @@ | |||
1 | From fbc9996d6490a5d4720b85b47f38335e7fdc99d9 Mon Sep 17 00:00:00 2001 | ||
2 | From: haonanya <haonan.yang@intel.com> | ||
3 | Date: Mon, 19 Jul 2021 10:14:20 +0800 | ||
4 | Subject: [PATCH 3/3] Add support for cl_ext_float_atomics in SPIRVWriter | ||
5 | |||
6 | Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-100/patches/spirv/0001-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch] | ||
7 | |||
8 | Signed-off-by: haonanya <haonan.yang@intel.com> | ||
9 | Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com> | ||
10 | --- | ||
11 | lib/SPIRV/OCL20ToSPIRV.cpp | 79 ++++++++++++++++-- | ||
12 | lib/SPIRV/SPIRVToOCL.h | 3 + | ||
13 | lib/SPIRV/SPIRVToOCL12.cpp | 21 +++++ | ||
14 | lib/SPIRV/SPIRVToOCL20.cpp | 28 ++++++- | ||
15 | lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h | 1 - | ||
16 | lib/SPIRV/libSPIRV/SPIRVOpCode.h | 8 +- | ||
17 | test/AtomicFAddEXTForOCL.ll | 64 +++++++++++++++ | ||
18 | test/AtomicFAddExt.ll | 111 ++++++++----------------- | ||
19 | test/AtomicFMaxEXT.ll | 113 +++++++------------------- | ||
20 | test/AtomicFMaxEXTForOCL.ll | 64 +++++++++++++++ | ||
21 | test/AtomicFMinEXT.ll | 113 +++++++------------------- | ||
22 | test/AtomicFMinEXTForOCL.ll | 64 +++++++++++++++ | ||
23 | test/InvalidAtomicBuiltins.cl | 8 -- | ||
24 | 13 files changed, 417 insertions(+), 260 deletions(-) | ||
25 | create mode 100644 test/AtomicFAddEXTForOCL.ll | ||
26 | create mode 100644 test/AtomicFMaxEXTForOCL.ll | ||
27 | create mode 100644 test/AtomicFMinEXTForOCL.ll | ||
28 | |||
29 | diff --git a/lib/SPIRV/OCL20ToSPIRV.cpp b/lib/SPIRV/OCL20ToSPIRV.cpp | ||
30 | index e30aa5be..b676a009 100644 | ||
31 | --- a/lib/SPIRV/OCL20ToSPIRV.cpp | ||
32 | +++ b/lib/SPIRV/OCL20ToSPIRV.cpp | ||
33 | @@ -408,10 +408,63 @@ void OCL20ToSPIRV::visitCallInst(CallInst &CI) { | ||
34 | if (DemangledName.find(kOCLBuiltinName::AtomicPrefix) == 0 || | ||
35 | DemangledName.find(kOCLBuiltinName::AtomPrefix) == 0) { | ||
36 | |||
37 | - // Compute atomic builtins do not support floating types. | ||
38 | - if (CI.getType()->isFloatingPointTy() && | ||
39 | - isComputeAtomicOCLBuiltin(DemangledName)) | ||
40 | - return; | ||
41 | + // Compute "atom" prefixed builtins do not support floating types. | ||
42 | + if (CI.getType()->isFloatingPointTy()) { | ||
43 | + if (DemangledName.find(kOCLBuiltinName::AtomPrefix) == 0) | ||
44 | + return; | ||
45 | + // handle functions which are "atomic_" prefixed. | ||
46 | + StringRef Stem = DemangledName; | ||
47 | + Stem = Stem.drop_front(strlen("atomic_")); | ||
48 | + // FP-typed atomic_{add, sub, inc, dec, exchange, min, max, or, and, xor, | ||
49 | + // fetch_or, fetch_xor, fetch_and, fetch_or_explicit, fetch_xor_explicit, | ||
50 | + // fetch_and_explicit} should be identified as function call | ||
51 | + bool IsFunctionCall = llvm::StringSwitch<bool>(Stem) | ||
52 | + .Case("add", true) | ||
53 | + .Case("sub", true) | ||
54 | + .Case("inc", true) | ||
55 | + .Case("dec", true) | ||
56 | + .Case("cmpxchg", true) | ||
57 | + .Case("min", true) | ||
58 | + .Case("max", true) | ||
59 | + .Case("or", true) | ||
60 | + .Case("xor", true) | ||
61 | + .Case("and", true) | ||
62 | + .Case("fetch_or", true) | ||
63 | + .Case("fetch_and", true) | ||
64 | + .Case("fetch_xor", true) | ||
65 | + .Case("fetch_or_explicit", true) | ||
66 | + .Case("fetch_xor_explicit", true) | ||
67 | + .Case("fetch_and_explicit", true) | ||
68 | + .Default(false); | ||
69 | + if (IsFunctionCall) | ||
70 | + return; | ||
71 | + if (F->arg_size() != 2) { | ||
72 | + IsFunctionCall = llvm::StringSwitch<bool>(Stem) | ||
73 | + .Case("exchange", true) | ||
74 | + .Case("fetch_add", true) | ||
75 | + .Case("fetch_sub", true) | ||
76 | + .Case("fetch_min", true) | ||
77 | + .Case("fetch_max", true) | ||
78 | + .Case("load", true) | ||
79 | + .Case("store", true) | ||
80 | + .Default(false); | ||
81 | + if (IsFunctionCall) | ||
82 | + return; | ||
83 | + } | ||
84 | + if (F->arg_size() != 3 && F->arg_size() != 4) { | ||
85 | + IsFunctionCall = llvm::StringSwitch<bool>(Stem) | ||
86 | + .Case("exchange_explicit", true) | ||
87 | + .Case("fetch_add_explicit", true) | ||
88 | + .Case("fetch_sub_explicit", true) | ||
89 | + .Case("fetch_min_explicit", true) | ||
90 | + .Case("fetch_max_explicit", true) | ||
91 | + .Case("load_explicit", true) | ||
92 | + .Case("store_explicit", true) | ||
93 | + .Default(false); | ||
94 | + if (IsFunctionCall) | ||
95 | + return; | ||
96 | + } | ||
97 | + } | ||
98 | |||
99 | auto PCI = &CI; | ||
100 | if (DemangledName == kOCLBuiltinName::AtomicInit) { | ||
101 | @@ -819,7 +872,7 @@ void OCL20ToSPIRV::transAtomicBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info) { | ||
102 | AttributeList Attrs = CI->getCalledFunction()->getAttributes(); | ||
103 | mutateCallInstSPIRV( | ||
104 | M, CI, | ||
105 | - [=](CallInst *CI, std::vector<Value *> &Args) { | ||
106 | + [=](CallInst *CI, std::vector<Value *> &Args) -> std::string { | ||
107 | Info.PostProc(Args); | ||
108 | // Order of args in OCL20: | ||
109 | // object, 0-2 other args, 1-2 order, scope | ||
110 | @@ -864,7 +917,21 @@ void OCL20ToSPIRV::transAtomicBuiltin(CallInst *CI, OCLBuiltinTransInfo &Info) { | ||
111 | std::rotate(Args.begin() + 2, Args.begin() + OrderIdx, | ||
112 | Args.end() - Offset); | ||
113 | } | ||
114 | - return getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName)); | ||
115 | + llvm::Type* AtomicBuiltinsReturnType = | ||
116 | + CI->getCalledFunction()->getReturnType(); | ||
117 | + auto IsFPType = [](llvm::Type *ReturnType) { | ||
118 | + return ReturnType->isHalfTy() || ReturnType->isFloatTy() || | ||
119 | + ReturnType->isDoubleTy(); | ||
120 | + }; | ||
121 | + auto SPIRVFunctionName = | ||
122 | + getSPIRVFuncName(OCLSPIRVBuiltinMap::map(Info.UniqName)); | ||
123 | + if (!IsFPType(AtomicBuiltinsReturnType)) | ||
124 | + return SPIRVFunctionName; | ||
125 | + // Translate FP-typed atomic builtins. | ||
126 | + return llvm::StringSwitch<std::string>(SPIRVFunctionName) | ||
127 | + .Case("__spirv_AtomicIAdd", "__spirv_AtomicFAddEXT") | ||
128 | + .Case("__spirv_AtomicSMax", "__spirv_AtomicFMaxEXT") | ||
129 | + .Case("__spirv_AtomicSMin", "__spirv_AtomicFMinEXT"); | ||
130 | }, | ||
131 | &Attrs); | ||
132 | } | ||
133 | diff --git a/lib/SPIRV/SPIRVToOCL.h b/lib/SPIRV/SPIRVToOCL.h | ||
134 | index ddeec0b6..006fb0b1 100644 | ||
135 | --- a/lib/SPIRV/SPIRVToOCL.h | ||
136 | +++ b/lib/SPIRV/SPIRVToOCL.h | ||
137 | @@ -178,6 +178,9 @@ public: | ||
138 | /// using separate maps for OpenCL 1.2 and OpenCL 2.0 | ||
139 | virtual Instruction *mutateAtomicName(CallInst *CI, Op OC) = 0; | ||
140 | |||
141 | + // Transform FP atomic opcode to corresponding OpenCL function name | ||
142 | + virtual std::string mapFPAtomicName(Op OC) = 0; | ||
143 | + | ||
144 | private: | ||
145 | /// Transform uniform group opcode to corresponding OpenCL function name, | ||
146 | /// example: GroupIAdd(Reduce) => group_iadd => work_group_reduce_add | | ||
147 | diff --git a/lib/SPIRV/SPIRVToOCL12.cpp b/lib/SPIRV/SPIRVToOCL12.cpp | ||
148 | index afddd596..d7f00de3 100644 | ||
149 | --- a/lib/SPIRV/SPIRVToOCL12.cpp | ||
150 | +++ b/lib/SPIRV/SPIRVToOCL12.cpp | ||
151 | @@ -104,6 +104,9 @@ public: | ||
152 | /// cl_khr_int64_base_atomics and cl_khr_int64_extended_atomics extensions. | ||
153 | std::string mapAtomicName(Op OC, Type *Ty); | ||
154 | |||
155 | + // Transform FP atomic opcode to corresponding OpenCL function name | ||
156 | + std::string mapFPAtomicName(Op OC) override; | ||
157 | + | ||
158 | static char ID; | ||
159 | }; | ||
160 | |||
161 | @@ -338,6 +341,21 @@ Instruction *SPIRVToOCL12::visitCallSPIRVAtomicBuiltin(CallInst *CI, Op OC) { | ||
162 | return NewCI; | ||
163 | } | ||
164 | |||
165 | +std::string SPIRVToOCL12::mapFPAtomicName(Op OC) { | ||
166 | + assert(isFPAtomicOpCode(OC) && "Not intended to handle other opcodes than " | ||
167 | + "AtomicF{Add/Min/Max}EXT!"); | ||
168 | + switch (OC) { | ||
169 | + case OpAtomicFAddEXT: | ||
170 | + return "atomic_add"; | ||
171 | + case OpAtomicFMinEXT: | ||
172 | + return "atomic_min"; | ||
173 | + case OpAtomicFMaxEXT: | ||
174 | + return "atomic_max"; | ||
175 | + default: | ||
176 | + llvm_unreachable("Unsupported opcode!"); | ||
177 | + } | ||
178 | +} | ||
179 | + | ||
180 | Instruction *SPIRVToOCL12::mutateAtomicName(CallInst *CI, Op OC) { | ||
181 | AttributeList Attrs = CI->getCalledFunction()->getAttributes(); | ||
182 | return mutateCallInstOCL( | ||
183 | @@ -351,6 +369,9 @@ Instruction *SPIRVToOCL12::mutateAtomicName(CallInst *CI, Op OC) { | ||
184 | std::string SPIRVToOCL12::mapAtomicName(Op OC, Type *Ty) { | ||
185 | std::string Prefix = Ty->isIntegerTy(64) ? kOCLBuiltinName::AtomPrefix | ||
186 | : kOCLBuiltinName::AtomicPrefix; | ||
187 | + // Map fp atomic instructions to regular OpenCL built-ins. | ||
188 | + if (isFPAtomicOpCode(OC)) | ||
189 | + return mapFPAtomicName(OC); | ||
190 | return Prefix += OCL12SPIRVBuiltinMap::rmap(OC); | ||
191 | } | ||
192 | |||
193 | diff --git a/lib/SPIRV/SPIRVToOCL20.cpp b/lib/SPIRV/SPIRVToOCL20.cpp | ||
194 | index d829ff42..01d088e9 100644 | ||
195 | --- a/lib/SPIRV/SPIRVToOCL20.cpp | ||
196 | +++ b/lib/SPIRV/SPIRVToOCL20.cpp | ||
197 | @@ -82,6 +82,9 @@ public: | ||
198 | /// compare_exchange_strong/weak_explicit | ||
199 | Instruction *visitCallSPIRVAtomicCmpExchg(CallInst *CI, Op OC) override; | ||
200 | |||
201 | + // Transform FP atomic opcode to corresponding OpenCL function name | ||
202 | + std::string mapFPAtomicName(Op OC) override; | ||
203 | + | ||
204 | static char ID; | ||
205 | }; | ||
206 | |||
207 | @@ -144,11 +147,29 @@ void SPIRVToOCL20::visitCallSPIRVControlBarrier(CallInst *CI) { | ||
208 | &Attrs); | ||
209 | } | ||
210 | |||
211 | +std::string SPIRVToOCL20::mapFPAtomicName(Op OC) { | ||
212 | + assert(isFPAtomicOpCode(OC) && "Not intended to handle other opcodes than " | ||
213 | + "AtomicF{Add/Min/Max}EXT!"); | ||
214 | + switch (OC) { | ||
215 | + case OpAtomicFAddEXT: | ||
216 | + return "atomic_fetch_add_explicit"; | ||
217 | + case OpAtomicFMinEXT: | ||
218 | + return "atomic_fetch_min_explicit"; | ||
219 | + case OpAtomicFMaxEXT: | ||
220 | + return "atomic_fetch_max_explicit"; | ||
221 | + default: | ||
222 | + llvm_unreachable("Unsupported opcode!"); | ||
223 | + } | ||
224 | +} | ||
225 | + | ||
226 | Instruction *SPIRVToOCL20::mutateAtomicName(CallInst *CI, Op OC) { | ||
227 | AttributeList Attrs = CI->getCalledFunction()->getAttributes(); | ||
228 | return mutateCallInstOCL( | ||
229 | M, CI, | ||
230 | [=](CallInst *, std::vector<Value *> &Args) { | ||
231 | + // Map fp atomic instructions to regular OpenCL built-ins. | ||
232 | + if (isFPAtomicOpCode(OC)) | ||
233 | + return mapFPAtomicName(OC); | ||
234 | return OCLSPIRVBuiltinMap::rmap(OC); | ||
235 | }, | ||
236 | &Attrs); | ||
237 | @@ -215,7 +236,12 @@ CallInst *SPIRVToOCL20::mutateCommonAtomicArguments(CallInst *CI, Op OC) { | ||
238 | } | ||
239 | } | ||
240 | auto Ptr = findFirstPtr(Args); | ||
241 | - auto Name = OCLSPIRVBuiltinMap::rmap(OC); | ||
242 | + std::string Name; | ||
243 | + // Map fp atomic instructions to regular OpenCL built-ins. | ||
244 | + if (isFPAtomicOpCode(OC)) | ||
245 | + Name = mapFPAtomicName(OC); | ||
246 | + else | ||
247 | + Name = OCLSPIRVBuiltinMap::rmap(OC); | ||
248 | auto NumOrder = getSPIRVAtomicBuiltinNumMemoryOrderArgs(OC); | ||
249 | auto ScopeIdx = Ptr + 1; | ||
250 | auto OrderIdx = Ptr + 2; | ||
251 | diff --git a/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h b/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h | ||
252 | index 13f93fbe..7b707993 100644 | ||
253 | --- a/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h | ||
254 | +++ b/lib/SPIRV/libSPIRV/SPIRVNameMapEnum.h | ||
255 | @@ -521,7 +521,6 @@ template <> inline void SPIRVMap<Capability, std::string>::init() { | ||
256 | add(CapabilityAtomicFloat64AddEXT, "AtomicFloat64AddEXT"); | ||
257 | add(CapabilityAtomicFloat32MinMaxEXT, "AtomicFloat32MinMaxEXT"); | ||
258 | add(CapabilityAtomicFloat64MinMaxEXT, "AtomicFloat64MinMaxEXT"); | ||
259 | - add(CapabilityAtomicFloat16MinMaxEXT, "AtomicFloat16MinMaxEXT"); | ||
260 | add(CapabilitySubgroupShuffleINTEL, "SubgroupShuffleINTEL"); | ||
261 | add(CapabilitySubgroupBufferBlockIOINTEL, "SubgroupBufferBlockIOINTEL"); | ||
262 | add(CapabilitySubgroupImageBlockIOINTEL, "SubgroupImageBlockIOINTEL"); | ||
263 | diff --git a/lib/SPIRV/libSPIRV/SPIRVOpCode.h b/lib/SPIRV/libSPIRV/SPIRVOpCode.h | ||
264 | index feec70f6..8e595e83 100644 | ||
265 | --- a/lib/SPIRV/libSPIRV/SPIRVOpCode.h | ||
266 | +++ b/lib/SPIRV/libSPIRV/SPIRVOpCode.h | ||
267 | @@ -54,11 +54,17 @@ template <> inline void SPIRVMap<Op, std::string>::init() { | ||
268 | } | ||
269 | SPIRV_DEF_NAMEMAP(Op, OpCodeNameMap) | ||
270 | |||
271 | +inline bool isFPAtomicOpCode(Op OpCode) { | ||
272 | + return OpCode == OpAtomicFAddEXT || OpCode == OpAtomicFMinEXT || | ||
273 | + OpCode == OpAtomicFMaxEXT; | ||
274 | +} | ||
275 | + | ||
276 | inline bool isAtomicOpCode(Op OpCode) { | ||
277 | static_assert(OpAtomicLoad < OpAtomicXor, ""); | ||
278 | return ((unsigned)OpCode >= OpAtomicLoad && | ||
279 | (unsigned)OpCode <= OpAtomicXor) || | ||
280 | - OpCode == OpAtomicFlagTestAndSet || OpCode == OpAtomicFlagClear; | ||
281 | + OpCode == OpAtomicFlagTestAndSet || OpCode == OpAtomicFlagClear || | ||
282 | + isFPAtomicOpCode(OpCode); | ||
283 | } | ||
284 | inline bool isBinaryOpCode(Op OpCode) { | ||
285 | return ((unsigned)OpCode >= OpIAdd && (unsigned)OpCode <= OpFMod) || | ||
286 | diff --git a/test/AtomicFAddEXTForOCL.ll b/test/AtomicFAddEXTForOCL.ll | ||
287 | new file mode 100644 | ||
288 | index 00000000..fb146fb9 | ||
289 | --- /dev/null | ||
290 | +++ b/test/AtomicFAddEXTForOCL.ll | ||
291 | @@ -0,0 +1,64 @@ | ||
292 | +; RUN: llvm-as %s -o %t.bc | ||
293 | +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_add -o %t.spv | ||
294 | +; RUN: spirv-val %t.spv | ||
295 | +; RUN: llvm-spirv -to-text %t.spv -o %t.spt | ||
296 | +; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV | ||
297 | + | ||
298 | +; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc | ||
299 | +; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20 | ||
300 | + | ||
301 | +; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc | ||
302 | +; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV | ||
303 | + | ||
304 | +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" | ||
305 | +target triple = "spir-unknown-unknown" | ||
306 | + | ||
307 | +; CHECK-SPIRV: Capability AtomicFloat32AddEXT | ||
308 | +; CHECK-SPIRV: Capability AtomicFloat64AddEXT | ||
309 | +; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_add" | ||
310 | +; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 | ||
311 | +; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 | ||
312 | + | ||
313 | + | ||
314 | +; Function Attrs: convergent norecurse nounwind | ||
315 | +define dso_local spir_func void @test_atomic_float(float addrspace(1)* %a) local_unnamed_addr #0 { | ||
316 | +entry: | ||
317 | + ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]] | ||
318 | + ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}}) | ||
319 | + ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+fiif]]({{.*}}) | ||
320 | + %call = tail call spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2 | ||
321 | + ret void | ||
322 | +} | ||
323 | + | ||
324 | +; Function Attrs: convergent | ||
325 | +declare spir_func float @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1 | ||
326 | +; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float | ||
327 | + | ||
328 | +; Function Attrs: convergent norecurse nounwind | ||
329 | +define dso_local spir_func void @test_atomic_double(double addrspace(1)* %a) local_unnamed_addr #0 { | ||
330 | +entry: | ||
331 | + ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]] | ||
332 | + ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}}) | ||
333 | + ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+diid]]({{.*}}) | ||
334 | + %call = tail call spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2 | ||
335 | + ret void | ||
336 | +} | ||
337 | +; Function Attrs: convergent | ||
338 | +declare spir_func double @_Z25atomic_fetch_add_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1 | ||
339 | +; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double | ||
340 | + | ||
341 | +; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float | ||
342 | +; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double | ||
343 | + | ||
344 | +attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } | ||
345 | +attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } | ||
346 | +attributes #2 = { convergent nounwind } | ||
347 | + | ||
348 | +!llvm.module.flags = !{!0} | ||
349 | +!opencl.ocl.version = !{!1} | ||
350 | +!opencl.spir.version = !{!1} | ||
351 | +!llvm.ident = !{!2} | ||
352 | + | ||
353 | +!0 = !{i32 1, !"wchar_size", i32 4} | ||
354 | +!1 = !{i32 2, i32 0} | ||
355 | +!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"} | ||
356 | diff --git a/test/AtomicFAddExt.ll b/test/AtomicFAddExt.ll | ||
357 | index 011dd8a7..42bdfeea 100644 | ||
358 | --- a/test/AtomicFAddExt.ll | ||
359 | +++ b/test/AtomicFAddExt.ll | ||
360 | @@ -4,20 +4,16 @@ | ||
361 | ; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV | ||
362 | |||
363 | ; RUN: llvm-spirv -r %t.spv -o %t.rev.bc | ||
364 | -; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM | ||
365 | +; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL12 | ||
366 | |||
367 | -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" | ||
368 | -target triple = "spir64-unknown-unknown-sycldevice" | ||
369 | - | ||
370 | -%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } | ||
371 | -%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] } | ||
372 | -%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } | ||
373 | - | ||
374 | -$_ZTSZZ3addIfEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_37 = comdat any | ||
375 | +; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc | ||
376 | +; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20 | ||
377 | |||
378 | -$_ZTSZZ3addIdEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_37 = comdat any | ||
379 | +; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc | ||
380 | +; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV | ||
381 | |||
382 | -@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 | ||
383 | +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" | ||
384 | +target triple = "spir64-unknown-unknown-sycldevice" | ||
385 | |||
386 | ; CHECK-SPIRV: Capability AtomicFloat32AddEXT | ||
387 | ; CHECK-SPIRV: Capability AtomicFloat64AddEXT | ||
388 | @@ -25,62 +21,43 @@ $_ZTSZZ3addIdEvvENKUlRN2cl4sycl7handlerEE19_14clES3_EUlNS1_4itemILi1ELb1EEEE23_3 | ||
389 | ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 | ||
390 | ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 | ||
391 | |||
392 | -; Function Attrs: convergent norecurse mustprogress | ||
393 | -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 { | ||
394 | +; Function Attrs: convergent norecurse nounwind | ||
395 | +define dso_local spir_func float @_Z14AtomicFloatIncRf(float addrspace(4)* align 4 dereferenceable(4) %Arg) local_unnamed_addr #0 { | ||
396 | entry: | ||
397 | - %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0 | ||
398 | - %1 = load i64, i64* %0, align 8 | ||
399 | - %add.ptr.i29 = getelementptr inbounds float, float addrspace(1)* %_arg_, i64 %1 | ||
400 | - %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0 | ||
401 | - %3 = load i64, i64* %2, align 8 | ||
402 | - %add.ptr.i = getelementptr inbounds float, float addrspace(1)* %_arg_4, i64 %3 | ||
403 | - %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 | ||
404 | - %5 = extractelement <3 x i64> %4, i64 0 | ||
405 | + %0 = addrspacecast float addrspace(4)* %Arg to float addrspace(1)* | ||
406 | ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_32]] | ||
407 | - ; CHECK-LLVM: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+]]({{.*}}) | ||
408 | - %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 | ||
409 | - %add.i.i = fadd float %call3.i.i.i.i, 1.000000e+00 | ||
410 | - %sext.i = shl i64 %5, 32 | ||
411 | - %conv5.i = ashr exact i64 %sext.i, 32 | ||
412 | - %ptridx.i.i = getelementptr inbounds float, float addrspace(1)* %add.ptr.i, i64 %conv5.i | ||
413 | - %ptridx.ascast.i.i = addrspacecast float addrspace(1)* %ptridx.i.i to float addrspace(4)* | ||
414 | - store float %add.i.i, float addrspace(4)* %ptridx.ascast.i.i, align 4, !tbaa !14 | ||
415 | - ret void | ||
416 | + ; CHECK-LLVM-CL12: call spir_func float @[[FLOAT_FUNC_NAME:_Z10atomic_add[[:alnum:]]+ff]]({{.*}}) | ||
417 | + ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}}) | ||
418 | + ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+fiif]]({{.*}}) | ||
419 | + %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 | ||
420 | + ret float %call3.i.i | ||
421 | } | ||
422 | |||
423 | ; Function Attrs: convergent | ||
424 | -; CHECK-LLVM: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float addrspace(1)*, i32, i32, float) | ||
425 | declare dso_local spir_func float @_Z21__spirv_AtomicFAddEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)*, i32, i32, float) local_unnamed_addr #1 | ||
426 | +; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float | ||
427 | |||
428 | -; Function Attrs: convergent norecurse mustprogress | ||
429 | -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 { | ||
430 | +; Function Attrs: convergent norecurse nounwind | ||
431 | +define dso_local spir_func double @_Z15AtomicDoubleIncRd(double addrspace(4)* align 8 dereferenceable(8) %Arg) local_unnamed_addr #0 { | ||
432 | entry: | ||
433 | - %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0 | ||
434 | - %1 = load i64, i64* %0, align 8 | ||
435 | - %add.ptr.i29 = getelementptr inbounds double, double addrspace(1)* %_arg_, i64 %1 | ||
436 | - %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0 | ||
437 | - %3 = load i64, i64* %2, align 8 | ||
438 | - %add.ptr.i = getelementptr inbounds double, double addrspace(1)* %_arg_4, i64 %3 | ||
439 | - %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 | ||
440 | - %5 = extractelement <3 x i64> %4, i64 0 | ||
441 | + %0 = addrspacecast double addrspace(4)* %Arg to double addrspace(1)* | ||
442 | ; CHECK-SPIRV: 7 AtomicFAddEXT [[TYPE_FLOAT_64]] | ||
443 | - ; CHECK-LLVM: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+]]({{.*}}) | ||
444 | - %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 | ||
445 | - %add.i.i = fadd double %call3.i.i.i.i, 1.000000e+00 | ||
446 | - %sext.i = shl i64 %5, 32 | ||
447 | - %conv5.i = ashr exact i64 %sext.i, 32 | ||
448 | - %ptridx.i.i = getelementptr inbounds double, double addrspace(1)* %add.ptr.i, i64 %conv5.i | ||
449 | - %ptridx.ascast.i.i = addrspacecast double addrspace(1)* %ptridx.i.i to double addrspace(4)* | ||
450 | - store double %add.i.i, double addrspace(4)* %ptridx.ascast.i.i, align 8, !tbaa !27 | ||
451 | - ret void | ||
452 | + ; CHECK-LLVM-CL12: call spir_func double @[[DOUBLE_FUNC_NAME:_Z10atomic_add[[:alnum:]]+dd]]({{.*}}) | ||
453 | + ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_add_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}}) | ||
454 | + ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFAddEXT[[:alnum:]]+diid]]({{.*}}) | ||
455 | + %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 | ||
456 | + ret double %call3.i.i | ||
457 | } | ||
458 | |||
459 | ; Function Attrs: convergent | ||
460 | -; CHECK-LLVM: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double addrspace(1)*, i32, i32, double) | ||
461 | declare dso_local spir_func double @_Z21__spirv_AtomicFAddEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)*, i32, i32, double) local_unnamed_addr #1 | ||
462 | +; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double | ||
463 | |||
464 | -attributes #0 = { convergent norecurse } | ||
465 | -attributes #1 = { convergent } | ||
466 | +; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float | ||
467 | +; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double | ||
468 | + | ||
469 | +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" } | ||
470 | +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" } | ||
471 | attributes #2 = { convergent nounwind } | ||
472 | |||
473 | !llvm.module.flags = !{!0} | ||
474 | @@ -91,29 +68,5 @@ attributes #2 = { convergent nounwind } | ||
475 | !0 = !{i32 1, !"wchar_size", i32 4} | ||
476 | !1 = !{i32 1, i32 2} | ||
477 | !2 = !{i32 4, i32 100000} | ||
478 | -!3 = !{!"clang version 12.0.0"} | ||
479 | -!4 = !{i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1} | ||
480 | -!5 = !{!6, !8, !10, !12} | ||
481 | -!6 = distinct !{!6, !7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"} | ||
482 | -!7 = distinct !{!7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"} | ||
483 | -!8 = distinct !{!8, !9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"} | ||
484 | -!9 = distinct !{!9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"} | ||
485 | -!10 = distinct !{!10, !11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"} | ||
486 | -!11 = distinct !{!11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"} | ||
487 | -!12 = distinct !{!12, !13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"} | ||
488 | -!13 = distinct !{!13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"} | ||
489 | -!14 = !{!15, !15, i64 0} | ||
490 | -!15 = !{!"float", !16, i64 0} | ||
491 | -!16 = !{!"omnipotent char", !17, i64 0} | ||
492 | -!17 = !{!"Simple C++ TBAA"} | ||
493 | -!18 = !{!19, !21, !23, !25} | ||
494 | -!19 = distinct !{!19, !20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"} | ||
495 | -!20 = distinct !{!20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"} | ||
496 | -!21 = distinct !{!21, !22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"} | ||
497 | -!22 = distinct !{!22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"} | ||
498 | -!23 = distinct !{!23, !24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"} | ||
499 | -!24 = distinct !{!24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"} | ||
500 | -!25 = distinct !{!25, !26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"} | ||
501 | -!26 = distinct !{!26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"} | ||
502 | -!27 = !{!28, !28, i64 0} | ||
503 | -!28 = !{!"double", !16, i64 0} | ||
504 | +!3 = !{!"clang version 13.0.0"} | ||
505 | + | ||
506 | diff --git a/test/AtomicFMaxEXT.ll b/test/AtomicFMaxEXT.ll | ||
507 | index 1b81e53b..1c2eec93 100644 | ||
508 | --- a/test/AtomicFMaxEXT.ll | ||
509 | +++ b/test/AtomicFMaxEXT.ll | ||
510 | @@ -4,20 +4,16 @@ | ||
511 | ; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV | ||
512 | |||
513 | ; RUN: llvm-spirv -r %t.spv -o %t.rev.bc | ||
514 | -; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM | ||
515 | +; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL12 | ||
516 | |||
517 | -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" | ||
518 | -target triple = "spir64-unknown-unknown-sycldevice" | ||
519 | - | ||
520 | -%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } | ||
521 | -%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] } | ||
522 | -%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } | ||
523 | - | ||
524 | -$_ZTSZZ8max_testIfEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37 = comdat any | ||
525 | +; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc | ||
526 | +; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20 | ||
527 | |||
528 | -$_ZTSZZ8max_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37 = comdat any | ||
529 | +; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc | ||
530 | +; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV | ||
531 | |||
532 | -@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 | ||
533 | +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" | ||
534 | +target triple = "spir64-unknown-unknown-sycldevice" | ||
535 | |||
536 | ; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT | ||
537 | ; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT | ||
538 | @@ -25,65 +21,42 @@ $_ZTSZZ8max_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4item | ||
539 | ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 | ||
540 | ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 | ||
541 | |||
542 | -; Function Attrs: convergent norecurse | ||
543 | -define weak_odr dso_local spir_kernel void @_ZTSZZ8max_testIfEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37(float addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, float addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 { | ||
544 | +; Function Attrs: convergent norecurse nounwind | ||
545 | +define dso_local spir_func float @_Z14AtomicFloatMaxRf(float addrspace(4)* align 4 dereferenceable(4) %Arg) local_unnamed_addr #0 { | ||
546 | entry: | ||
547 | - %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0 | ||
548 | - %1 = load i64, i64* %0, align 8 | ||
549 | - %add.ptr.i29 = getelementptr inbounds float, float addrspace(1)* %_arg_, i64 %1 | ||
550 | - %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0 | ||
551 | - %3 = load i64, i64* %2, align 8 | ||
552 | - %add.ptr.i = getelementptr inbounds float, float addrspace(1)* %_arg_4, i64 %3 | ||
553 | - %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 | ||
554 | - %5 = extractelement <3 x i64> %4, i64 0 | ||
555 | - %conv.i = trunc i64 %5 to i32 | ||
556 | - %conv3.i = sitofp i32 %conv.i to float | ||
557 | - %add.i = fadd float %conv3.i, 1.000000e+00 | ||
558 | + %0 = addrspacecast float addrspace(4)* %Arg to float addrspace(1)* | ||
559 | ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_32]] | ||
560 | - ; CHECK-LLVM: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+]]({{.*}}) | ||
561 | - %call3.i.i.i = tail call spir_func float @_Z21__spirv_AtomicFMaxEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %add.ptr.i29, i32 1, i32 896, float %add.i) #2 | ||
562 | - %sext.i = shl i64 %5, 32 | ||
563 | - %conv6.i = ashr exact i64 %sext.i, 32 | ||
564 | - %ptridx.i.i = getelementptr inbounds float, float addrspace(1)* %add.ptr.i, i64 %conv6.i | ||
565 | - %ptridx.ascast.i.i = addrspacecast float addrspace(1)* %ptridx.i.i to float addrspace(4)* | ||
566 | - store float %call3.i.i.i, float addrspace(4)* %ptridx.ascast.i.i, align 4, !tbaa !14 | ||
567 | - ret void | ||
568 | + ; CHECK-LLVM-CL12: call spir_func float @[[FLOAT_FUNC_NAME:_Z10atomic_max[[:alnum:]]+ff]]({{.*}}) | ||
569 | + ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}}) | ||
570 | + ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+fiif]]({{.*}}) | ||
571 | + %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 | ||
572 | + ret float %call.i.i.i | ||
573 | } | ||
574 | |||
575 | ; Function Attrs: convergent | ||
576 | -; CHECK-LLVM: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float addrspace(1)*, i32, i32, float) | ||
577 | declare dso_local spir_func float @_Z21__spirv_AtomicFMaxEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)*, i32, i32, float) local_unnamed_addr #1 | ||
578 | +; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float | ||
579 | |||
580 | -; Function Attrs: convergent norecurse | ||
581 | -define weak_odr dso_local spir_kernel void @_ZTSZZ8max_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37(double addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, double addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 { | ||
582 | +; Function Attrs: convergent norecurse nounwind | ||
583 | +define dso_local spir_func double @_Z15AtomicDoubleMaxRd(double addrspace(4)* align 8 dereferenceable(8) %Arg) local_unnamed_addr #0 { | ||
584 | entry: | ||
585 | - %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0 | ||
586 | - %1 = load i64, i64* %0, align 8 | ||
587 | - %add.ptr.i29 = getelementptr inbounds double, double addrspace(1)* %_arg_, i64 %1 | ||
588 | - %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0 | ||
589 | - %3 = load i64, i64* %2, align 8 | ||
590 | - %add.ptr.i = getelementptr inbounds double, double addrspace(1)* %_arg_4, i64 %3 | ||
591 | - %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 | ||
592 | - %5 = extractelement <3 x i64> %4, i64 0 | ||
593 | - %conv.i = trunc i64 %5 to i32 | ||
594 | - %conv3.i = sitofp i32 %conv.i to double | ||
595 | - %add.i = fadd double %conv3.i, 1.000000e+00 | ||
596 | + %0 = addrspacecast double addrspace(4)* %Arg to double addrspace(1)* | ||
597 | ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_64]] | ||
598 | - ; CHECK-LLVM: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+]]({{.*}}) | ||
599 | - %call3.i.i.i = tail call spir_func double @_Z21__spirv_AtomicFMaxEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %add.ptr.i29, i32 1, i32 896, double %add.i) #2 | ||
600 | - %sext.i = shl i64 %5, 32 | ||
601 | - %conv6.i = ashr exact i64 %sext.i, 32 | ||
602 | - %ptridx.i.i = getelementptr inbounds double, double addrspace(1)* %add.ptr.i, i64 %conv6.i | ||
603 | - %ptridx.ascast.i.i = addrspacecast double addrspace(1)* %ptridx.i.i to double addrspace(4)* | ||
604 | - store double %call3.i.i.i, double addrspace(4)* %ptridx.ascast.i.i, align 8, !tbaa !27 | ||
605 | - ret void | ||
606 | + ; CHECK-LLVM-CL12: call spir_func double @[[DOUBLE_FUNC_NAME:_Z10atomic_max[[:alnum:]]+dd]]({{.*}}) | ||
607 | + ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}}) | ||
608 | + ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+diid]]({{.*}}) | ||
609 | + %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 | ||
610 | + ret double %call.i.i.i | ||
611 | } | ||
612 | |||
613 | ; Function Attrs: convergent | ||
614 | -; CHECK-LLVM: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double addrspace(1)*, i32, i32, double) | ||
615 | declare dso_local spir_func double @_Z21__spirv_AtomicFMaxEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)*, i32, i32, double) local_unnamed_addr #1 | ||
616 | +; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double | ||
617 | |||
618 | -attributes #0 = { convergent norecurse "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } | ||
619 | +; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float | ||
620 | +; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double | ||
621 | + | ||
622 | +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" } | ||
623 | 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" } | ||
624 | attributes #2 = { convergent nounwind } | ||
625 | |||
626 | @@ -95,29 +68,5 @@ attributes #2 = { convergent nounwind } | ||
627 | !0 = !{i32 1, !"wchar_size", i32 4} | ||
628 | !1 = !{i32 1, i32 2} | ||
629 | !2 = !{i32 4, i32 100000} | ||
630 | -!3 = !{!"clang version 12.0.0"} | ||
631 | -!4 = !{i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1} | ||
632 | -!5 = !{!6, !8, !10, !12} | ||
633 | -!6 = distinct !{!6, !7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"} | ||
634 | -!7 = distinct !{!7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"} | ||
635 | -!8 = distinct !{!8, !9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"} | ||
636 | -!9 = distinct !{!9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"} | ||
637 | -!10 = distinct !{!10, !11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"} | ||
638 | -!11 = distinct !{!11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"} | ||
639 | -!12 = distinct !{!12, !13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"} | ||
640 | -!13 = distinct !{!13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"} | ||
641 | -!14 = !{!15, !15, i64 0} | ||
642 | -!15 = !{!"float", !16, i64 0} | ||
643 | -!16 = !{!"omnipotent char", !17, i64 0} | ||
644 | -!17 = !{!"Simple C++ TBAA"} | ||
645 | -!18 = !{!19, !21, !23, !25} | ||
646 | -!19 = distinct !{!19, !20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"} | ||
647 | -!20 = distinct !{!20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"} | ||
648 | -!21 = distinct !{!21, !22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"} | ||
649 | -!22 = distinct !{!22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"} | ||
650 | -!23 = distinct !{!23, !24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"} | ||
651 | -!24 = distinct !{!24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"} | ||
652 | -!25 = distinct !{!25, !26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"} | ||
653 | -!26 = distinct !{!26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"} | ||
654 | -!27 = !{!28, !28, i64 0} | ||
655 | -!28 = !{!"double", !16, i64 0} | ||
656 | +!3 = !{!"clang version 13.0.0"} | ||
657 | + | ||
658 | diff --git a/test/AtomicFMaxEXTForOCL.ll b/test/AtomicFMaxEXTForOCL.ll | ||
659 | new file mode 100644 | ||
660 | index 00000000..1f2530d9 | ||
661 | --- /dev/null | ||
662 | +++ b/test/AtomicFMaxEXTForOCL.ll | ||
663 | @@ -0,0 +1,64 @@ | ||
664 | +; RUN: llvm-as %s -o %t.bc | ||
665 | +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_min_max -o %t.spv | ||
666 | +; RUN: spirv-val %t.spv | ||
667 | +; RUN: llvm-spirv -to-text %t.spv -o %t.spt | ||
668 | +; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV | ||
669 | + | ||
670 | +; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc | ||
671 | +; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20 | ||
672 | + | ||
673 | +; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc | ||
674 | +; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV | ||
675 | + | ||
676 | +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" | ||
677 | +target triple = "spir-unknown-unknown" | ||
678 | + | ||
679 | +; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT | ||
680 | +; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT | ||
681 | +; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_min_max" | ||
682 | +; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 | ||
683 | +; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 | ||
684 | + | ||
685 | +; Function Attrs: convergent norecurse nounwind | ||
686 | +define dso_local spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 { | ||
687 | +entry: | ||
688 | + ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_32]] | ||
689 | + ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}}) | ||
690 | + ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+fiif]]({{.*}}) | ||
691 | + %call = tail call spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2 | ||
692 | + ret void | ||
693 | +} | ||
694 | + | ||
695 | +; Function Attrs: convergent | ||
696 | +declare spir_func float @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1 | ||
697 | +; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float | ||
698 | + | ||
699 | +; Function Attrs: convergent norecurse nounwind | ||
700 | +define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 { | ||
701 | +entry: | ||
702 | + ; CHECK-SPIRV: 7 AtomicFMaxEXT [[TYPE_FLOAT_64]] | ||
703 | + ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_max_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}}) | ||
704 | + ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMaxEXT[[:alnum:]]+diid]]({{.*}}) | ||
705 | + %call = tail call spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2 | ||
706 | + ret void | ||
707 | +} | ||
708 | + | ||
709 | +; Function Attrs: convergent | ||
710 | +declare spir_func double @_Z25atomic_fetch_max_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1 | ||
711 | +; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double | ||
712 | + | ||
713 | +; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float | ||
714 | +; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double | ||
715 | + | ||
716 | +attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } | ||
717 | +attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } | ||
718 | +attributes #2 = { convergent nounwind } | ||
719 | + | ||
720 | +!llvm.module.flags = !{!0} | ||
721 | +!opencl.ocl.version = !{!1} | ||
722 | +!opencl.spir.version = !{!1} | ||
723 | +!llvm.ident = !{!2} | ||
724 | + | ||
725 | +!0 = !{i32 1, !"wchar_size", i32 4} | ||
726 | +!1 = !{i32 2, i32 0} | ||
727 | +!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"} | ||
728 | diff --git a/test/AtomicFMinEXT.ll b/test/AtomicFMinEXT.ll | ||
729 | index 98c98b8e..9e40a669 100644 | ||
730 | --- a/test/AtomicFMinEXT.ll | ||
731 | +++ b/test/AtomicFMinEXT.ll | ||
732 | @@ -4,20 +4,16 @@ | ||
733 | ; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV | ||
734 | |||
735 | ; RUN: llvm-spirv -r %t.spv -o %t.rev.bc | ||
736 | -; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefix=CHECK-LLVM | ||
737 | +; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL12 | ||
738 | |||
739 | -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" | ||
740 | -target triple = "spir64-unknown-unknown-sycldevice" | ||
741 | - | ||
742 | -%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } | ||
743 | -%"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" = type { [1 x i64] } | ||
744 | -%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id" = type { %"class._ZTSN2cl4sycl6detail5arrayILi1EEE.cl::sycl::detail::array" } | ||
745 | - | ||
746 | -$_ZTSZZ8min_testIfEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37 = comdat any | ||
747 | +; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc | ||
748 | +; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20 | ||
749 | |||
750 | -$_ZTSZZ8min_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37 = comdat any | ||
751 | +; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc | ||
752 | +; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV | ||
753 | |||
754 | -@__spirv_BuiltInGlobalInvocationId = external dso_local local_unnamed_addr addrspace(1) constant <3 x i64>, align 32 | ||
755 | +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" | ||
756 | +target triple = "spir64-unknown-unknown-sycldevice" | ||
757 | |||
758 | ; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT | ||
759 | ; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT | ||
760 | @@ -25,65 +21,42 @@ $_ZTSZZ8min_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4item | ||
761 | ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 | ||
762 | ; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 | ||
763 | |||
764 | -; Function Attrs: convergent norecurse | ||
765 | -define weak_odr dso_local spir_kernel void @_ZTSZZ8min_testIfEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37(float addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, float addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 { | ||
766 | +; Function Attrs: convergent norecurse nounwind | ||
767 | +define dso_local spir_func float @_Z14AtomicFloatMinRf(float addrspace(4)* align 4 dereferenceable(4) %Arg) local_unnamed_addr #0 { | ||
768 | entry: | ||
769 | - %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0 | ||
770 | - %1 = load i64, i64* %0, align 8 | ||
771 | - %add.ptr.i29 = getelementptr inbounds float, float addrspace(1)* %_arg_, i64 %1 | ||
772 | - %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0 | ||
773 | - %3 = load i64, i64* %2, align 8 | ||
774 | - %add.ptr.i = getelementptr inbounds float, float addrspace(1)* %_arg_4, i64 %3 | ||
775 | - %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 | ||
776 | - %5 = extractelement <3 x i64> %4, i64 0 | ||
777 | - %conv.i = trunc i64 %5 to i32 | ||
778 | - %conv3.i = sitofp i32 %conv.i to float | ||
779 | - %add.i = fadd float %conv3.i, 1.000000e+00 | ||
780 | + %0 = addrspacecast float addrspace(4)* %Arg to float addrspace(1)* | ||
781 | ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_32]] | ||
782 | - ; CHECK-LLVM: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+]]({{.*}}) | ||
783 | - %call3.i.i.i = tail call spir_func float @_Z21__spirv_AtomicFMinEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)* %add.ptr.i29, i32 1, i32 896, float %add.i) #2 | ||
784 | - %sext.i = shl i64 %5, 32 | ||
785 | - %conv6.i = ashr exact i64 %sext.i, 32 | ||
786 | - %ptridx.i.i = getelementptr inbounds float, float addrspace(1)* %add.ptr.i, i64 %conv6.i | ||
787 | - %ptridx.ascast.i.i = addrspacecast float addrspace(1)* %ptridx.i.i to float addrspace(4)* | ||
788 | - store float %call3.i.i.i, float addrspace(4)* %ptridx.ascast.i.i, align 4, !tbaa !14 | ||
789 | - ret void | ||
790 | + ; CHECK-LLVM-CL12: call spir_func float @[[FLOAT_FUNC_NAME:_Z10atomic_min[[:alnum:]]+ff]]({{.*}}) | ||
791 | + ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}}) | ||
792 | + ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+fiif]]({{.*}}) | ||
793 | + %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 | ||
794 | + ret float %call.i.i.i | ||
795 | } | ||
796 | |||
797 | ; Function Attrs: convergent | ||
798 | -; CHECK-LLVM: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float addrspace(1)*, i32, i32, float) | ||
799 | declare dso_local spir_func float @_Z21__spirv_AtomicFMinEXTPU3AS1fN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEf(float addrspace(1)*, i32, i32, float) local_unnamed_addr #1 | ||
800 | +; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float | ||
801 | |||
802 | -; Function Attrs: convergent norecurse | ||
803 | -define weak_odr dso_local spir_kernel void @_ZTSZZ8min_testIdEvN2cl4sycl5queueEmENKUlRNS1_7handlerEE16_14clES4_EUlNS1_4itemILi1ELb1EEEE19_37(double addrspace(1)* %_arg_, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_1, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_2, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_3, double addrspace(1)* %_arg_4, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_6, %"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range"* byval(%"class._ZTSN2cl4sycl5rangeILi1EEE.cl::sycl::range") align 8 %_arg_7, %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* byval(%"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id") align 8 %_arg_8) local_unnamed_addr #0 comdat !kernel_arg_buffer_location !4 { | ||
804 | +; Function Attrs: convergent norecurse nounwind | ||
805 | +define dso_local spir_func double @_Z15AtomicDoubleMinRd(double addrspace(4)* align 8 dereferenceable(8) %Arg) local_unnamed_addr #0 { | ||
806 | entry: | ||
807 | - %0 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_3, i64 0, i32 0, i32 0, i64 0 | ||
808 | - %1 = load i64, i64* %0, align 8 | ||
809 | - %add.ptr.i29 = getelementptr inbounds double, double addrspace(1)* %_arg_, i64 %1 | ||
810 | - %2 = getelementptr inbounds %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id", %"class._ZTSN2cl4sycl2idILi1EEE.cl::sycl::id"* %_arg_8, i64 0, i32 0, i32 0, i64 0 | ||
811 | - %3 = load i64, i64* %2, align 8 | ||
812 | - %add.ptr.i = getelementptr inbounds double, double addrspace(1)* %_arg_4, i64 %3 | ||
813 | - %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 | ||
814 | - %5 = extractelement <3 x i64> %4, i64 0 | ||
815 | - %conv.i = trunc i64 %5 to i32 | ||
816 | - %conv3.i = sitofp i32 %conv.i to double | ||
817 | - %add.i = fadd double %conv3.i, 1.000000e+00 | ||
818 | + %0 = addrspacecast double addrspace(4)* %Arg to double addrspace(1)* | ||
819 | ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_64]] | ||
820 | - ; CHECK-LLVM: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+]]({{.*}}) | ||
821 | - %call3.i.i.i = tail call spir_func double @_Z21__spirv_AtomicFMinEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)* %add.ptr.i29, i32 1, i32 896, double %add.i) #2 | ||
822 | - %sext.i = shl i64 %5, 32 | ||
823 | - %conv6.i = ashr exact i64 %sext.i, 32 | ||
824 | - %ptridx.i.i = getelementptr inbounds double, double addrspace(1)* %add.ptr.i, i64 %conv6.i | ||
825 | - %ptridx.ascast.i.i = addrspacecast double addrspace(1)* %ptridx.i.i to double addrspace(4)* | ||
826 | - store double %call3.i.i.i, double addrspace(4)* %ptridx.ascast.i.i, align 8, !tbaa !27 | ||
827 | - ret void | ||
828 | + ; CHECK-LLVM-CL12: call spir_func double @[[DOUBLE_FUNC_NAME:_Z10atomic_min[[:alnum:]]+dd]]({{.*}}) | ||
829 | + ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}}) | ||
830 | + ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+diid]]({{.*}}) | ||
831 | + %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 | ||
832 | + ret double %call.i.i.i | ||
833 | } | ||
834 | |||
835 | ; Function Attrs: convergent | ||
836 | -; CHECK-LLVM: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double addrspace(1)*, i32, i32, double) | ||
837 | declare dso_local spir_func double @_Z21__spirv_AtomicFMinEXTPU3AS1dN5__spv5Scope4FlagENS1_19MemorySemanticsMask4FlagEd(double addrspace(1)*, i32, i32, double) local_unnamed_addr #1 | ||
838 | +; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double | ||
839 | |||
840 | -attributes #0 = { convergent norecurse "disable-tail-calls"="false" "frame-pointer"="all" "less-precise-fpmad"="false" "min-legal-vector-width"="0" "no-infs-fp-math"="false" "no-jump-tables"="false" "no-nans-fp-math"="false" "no-signed-zeros-fp-math"="false" "no-trapping-math"="true" "stack-protector-buffer-size"="8" "uniform-work-group-size"="true" "unsafe-fp-math"="false" "use-soft-float"="false" } | ||
841 | +; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float | ||
842 | +; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double | ||
843 | + | ||
844 | +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" } | ||
845 | 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" } | ||
846 | attributes #2 = { convergent nounwind } | ||
847 | |||
848 | @@ -95,29 +68,5 @@ attributes #2 = { convergent nounwind } | ||
849 | !0 = !{i32 1, !"wchar_size", i32 4} | ||
850 | !1 = !{i32 1, i32 2} | ||
851 | !2 = !{i32 4, i32 100000} | ||
852 | -!3 = !{!"clang version 12.0.0 (https://github.com/otcshare/llvm.git 67add71766d55d6a8d8d894822f583d6365a3b7d)"} | ||
853 | -!4 = !{i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1, i32 -1} | ||
854 | -!5 = !{!6, !8, !10, !12} | ||
855 | -!6 = distinct !{!6, !7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"} | ||
856 | -!7 = distinct !{!7, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"} | ||
857 | -!8 = distinct !{!8, !9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"} | ||
858 | -!9 = distinct !{!9, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"} | ||
859 | -!10 = distinct !{!10, !11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"} | ||
860 | -!11 = distinct !{!11, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"} | ||
861 | -!12 = distinct !{!12, !13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"} | ||
862 | -!13 = distinct !{!13, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"} | ||
863 | -!14 = !{!15, !15, i64 0} | ||
864 | -!15 = !{!"float", !16, i64 0} | ||
865 | -!16 = !{!"omnipotent char", !17, i64 0} | ||
866 | -!17 = !{!"Simple C++ TBAA"} | ||
867 | -!18 = !{!19, !21, !23, !25} | ||
868 | -!19 = distinct !{!19, !20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv: %agg.result"} | ||
869 | -!20 = distinct !{!20, !"_ZN7__spirv29InitSizesSTGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEE8initSizeEv"} | ||
870 | -!21 = distinct !{!21, !22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v: %agg.result"} | ||
871 | -!22 = distinct !{!22, !"_ZN7__spirvL22initGlobalInvocationIdILi1EN2cl4sycl2idILi1EEEEET0_v"} | ||
872 | -!23 = distinct !{!23, !24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv: %agg.result"} | ||
873 | -!24 = distinct !{!24, !"_ZN2cl4sycl6detail7Builder7getItemILi1ELb1EEENSt9enable_ifIXT0_EKNS0_4itemIXT_EXT0_EEEE4typeEv"} | ||
874 | -!25 = distinct !{!25, !26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE: %agg.result"} | ||
875 | -!26 = distinct !{!26, !"_ZN2cl4sycl6detail7Builder10getElementILi1ELb1EEEDTcl7getItemIXT_EXT0_EEEEPNS0_4itemIXT_EXT0_EEE"} | ||
876 | -!27 = !{!28, !28, i64 0} | ||
877 | -!28 = !{!"double", !16, i64 0} | ||
878 | +!3 = !{!"clang version 13.0.0"} | ||
879 | + | ||
880 | diff --git a/test/AtomicFMinEXTForOCL.ll b/test/AtomicFMinEXTForOCL.ll | ||
881 | new file mode 100644 | ||
882 | index 00000000..6196b0f8 | ||
883 | --- /dev/null | ||
884 | +++ b/test/AtomicFMinEXTForOCL.ll | ||
885 | @@ -0,0 +1,64 @@ | ||
886 | +; RUN: llvm-as %s -o %t.bc | ||
887 | +; RUN: llvm-spirv %t.bc --spirv-ext=+SPV_EXT_shader_atomic_float_min_max -o %t.spv | ||
888 | +; RUN: spirv-val %t.spv | ||
889 | +; RUN: llvm-spirv -to-text %t.spv -o %t.spt | ||
890 | +; RUN: FileCheck < %t.spt %s --check-prefix=CHECK-SPIRV | ||
891 | + | ||
892 | +; RUN: llvm-spirv --spirv-target-env=CL2.0 -r %t.spv -o %t.rev.bc | ||
893 | +; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-CL,CHECK-LLVM-CL20 | ||
894 | + | ||
895 | +; RUN: llvm-spirv --spirv-target-env=SPV-IR -r %t.spv -o %t.rev.bc | ||
896 | +; RUN: llvm-dis %t.rev.bc -o - | FileCheck %s --check-prefixes=CHECK-LLVM-SPV | ||
897 | + | ||
898 | +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" | ||
899 | +target triple = "spir-unknown-unknown" | ||
900 | + | ||
901 | +; CHECK-SPIRV: Capability AtomicFloat32MinMaxEXT | ||
902 | +; CHECK-SPIRV: Capability AtomicFloat64MinMaxEXT | ||
903 | +; CHECK-SPIRV: Extension "SPV_EXT_shader_atomic_float_min_max" | ||
904 | +; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_32:[0-9]+]] 32 | ||
905 | +; CHECK-SPIRV: TypeFloat [[TYPE_FLOAT_64:[0-9]+]] 64 | ||
906 | + | ||
907 | +; Function Attrs: convergent norecurse nounwind | ||
908 | +define dso_local spir_func void @test_float(float addrspace(1)* %a) local_unnamed_addr #0 { | ||
909 | +entry: | ||
910 | + ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_32]] | ||
911 | + ; CHECK-LLVM-CL20: call spir_func float @[[FLOAT_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicff[a-zA-Z0-9_]+]]({{.*}}) | ||
912 | + ; CHECK-LLVM-SPV: call spir_func float @[[FLOAT_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+fiif]]({{.*}}) | ||
913 | + %call = tail call spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)* %a, float 0.000000e+00, i32 0) #2 | ||
914 | + ret void | ||
915 | +} | ||
916 | + | ||
917 | +; Function Attrs: convergent | ||
918 | +declare spir_func float @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicff12memory_order(float addrspace(1)*, float, i32) local_unnamed_addr #1 | ||
919 | +; CHECK-LLVM-SPV: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float | ||
920 | + | ||
921 | +; Function Attrs: convergent norecurse nounwind | ||
922 | +define dso_local spir_func void @test_double(double addrspace(1)* %a) local_unnamed_addr #0 { | ||
923 | +entry: | ||
924 | + ; CHECK-SPIRV: 7 AtomicFMinEXT [[TYPE_FLOAT_64]] | ||
925 | + ; CHECK-LLVM-CL20: call spir_func double @[[DOUBLE_FUNC_NAME:_Z25atomic_fetch_min_explicit[[:alnum:]]+_Atomicdd[a-zA-Z0-9_]+]]({{.*}}) | ||
926 | + ; CHECK-LLVM-SPV: call spir_func double @[[DOUBLE_FUNC_NAME:_Z21__spirv_AtomicFMinEXT[[:alnum:]]+diid]]({{.*}}) | ||
927 | + %call = tail call spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)* %a, double 0.000000e+00, i32 0) #2 | ||
928 | + ret void | ||
929 | +} | ||
930 | + | ||
931 | +; Function Attrs: convergent | ||
932 | +declare spir_func double @_Z25atomic_fetch_min_explicitPU3AS1VU7_Atomicdd12memory_order(double addrspace(1)*, double, i32) local_unnamed_addr #1 | ||
933 | +; CHECK-LLVM-SPV: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double | ||
934 | + | ||
935 | +; CHECK-LLVM-CL: declare {{.*}}spir_func float @[[FLOAT_FUNC_NAME]](float | ||
936 | +; CHECK-LLVM-CL: declare {{.*}}spir_func double @[[DOUBLE_FUNC_NAME]](double | ||
937 | + | ||
938 | +attributes #0 = { convergent norecurse nounwind "frame-pointer"="none" "min-legal-vector-width"="0" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } | ||
939 | +attributes #1 = { convergent "frame-pointer"="none" "no-trapping-math"="true" "stack-protector-buffer-size"="8" } | ||
940 | +attributes #2 = { convergent nounwind } | ||
941 | + | ||
942 | +!llvm.module.flags = !{!0} | ||
943 | +!opencl.ocl.version = !{!1} | ||
944 | +!opencl.spir.version = !{!1} | ||
945 | +!llvm.ident = !{!2} | ||
946 | + | ||
947 | +!0 = !{i32 1, !"wchar_size", i32 4} | ||
948 | +!1 = !{i32 2, i32 0} | ||
949 | +!2 = !{!"clang version 13.0.0 (https://github.com/llvm/llvm-project.git 94aa388f0ce0723bb15503cf41c2c15b288375b9)"} | ||
950 | diff --git a/test/InvalidAtomicBuiltins.cl b/test/InvalidAtomicBuiltins.cl | ||
951 | index b8ec5b89..2182f070 100644 | ||
952 | --- a/test/InvalidAtomicBuiltins.cl | ||
953 | +++ b/test/InvalidAtomicBuiltins.cl | ||
954 | @@ -41,13 +41,9 @@ float __attribute__((overloadable)) atomic_fetch_xor(volatile generic atomic_flo | ||
955 | double __attribute__((overloadable)) atomic_fetch_and(volatile generic atomic_double *object, double operand, memory_order order); | ||
956 | double __attribute__((overloadable)) atomic_fetch_max(volatile generic atomic_double *object, double operand, memory_order order); | ||
957 | double __attribute__((overloadable)) atomic_fetch_min(volatile generic atomic_double *object, double operand, memory_order order); | ||
958 | -float __attribute__((overloadable)) atomic_fetch_add_explicit(volatile generic atomic_float *object, float operand, memory_order order); | ||
959 | -float __attribute__((overloadable)) atomic_fetch_sub_explicit(volatile generic atomic_float *object, float operand, memory_order order); | ||
960 | float __attribute__((overloadable)) atomic_fetch_or_explicit(volatile generic atomic_float *object, float operand, memory_order order); | ||
961 | float __attribute__((overloadable)) atomic_fetch_xor_explicit(volatile generic atomic_float *object, float operand, memory_order order); | ||
962 | double __attribute__((overloadable)) atomic_fetch_and_explicit(volatile generic atomic_double *object, double operand, memory_order order); | ||
963 | -double __attribute__((overloadable)) atomic_fetch_max_explicit(volatile generic atomic_double *object, double operand, memory_order order); | ||
964 | -double __attribute__((overloadable)) atomic_fetch_min_explicit(volatile generic atomic_double *object, double operand, memory_order order); | ||
965 | |||
966 | __kernel void test_atomic_fn(volatile __global float *p, | ||
967 | volatile __global double *pp, | ||
968 | @@ -86,11 +82,7 @@ __kernel void test_atomic_fn(volatile __global float *p, | ||
969 | d = atomic_fetch_and(pp, val, order); | ||
970 | d = atomic_fetch_min(pp, val, order); | ||
971 | d = atomic_fetch_max(pp, val, order); | ||
972 | - f = atomic_fetch_add_explicit(p, val, order); | ||
973 | - f = atomic_fetch_sub_explicit(p, val, order); | ||
974 | f = atomic_fetch_or_explicit(p, val, order); | ||
975 | f = atomic_fetch_xor_explicit(p, val, order); | ||
976 | d = atomic_fetch_and_explicit(pp, val, order); | ||
977 | - d = atomic_fetch_min_explicit(pp, val, order); | ||
978 | - d = atomic_fetch_max_explicit(pp, val, order); | ||
979 | } | ||
980 | -- | ||
981 | 2.17.1 | ||
982 | |||
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Memory-leak-fix-for-Managed-Static-Mutex.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Memory-leak-fix-for-Managed-Static-Mutex.patch new file mode 100644 index 00000000..3b035f47 --- /dev/null +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0003-Memory-leak-fix-for-Managed-Static-Mutex.patch | |||
@@ -0,0 +1,35 @@ | |||
1 | From cfb18b75e8a353bc7486f337541476a36994b063 Mon Sep 17 00:00:00 2001 | ||
2 | From: juanrod2 <> | ||
3 | Date: Tue, 22 Dec 2020 08:33:08 +0800 | ||
4 | Subject: [PATCH 3/7] Memory leak fix for Managed Static Mutex | ||
5 | |||
6 | Upstream-Status: Backport [Taken from opencl-clang patches; https://github.com/intel/opencl-clang/blob/ocl-open-100/patches/llvm/0001-Memory-leak-fix-for-Managed-Static-Mutex.patch] | ||
7 | |||
8 | Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com> | ||
9 | |||
10 | Cleaning a mutex inside ManagedStatic llvm class. | ||
11 | --- | ||
12 | llvm/lib/Support/ManagedStatic.cpp | 6 +++++- | ||
13 | 1 file changed, 5 insertions(+), 1 deletion(-) | ||
14 | |||
15 | diff --git a/llvm/lib/Support/ManagedStatic.cpp b/llvm/lib/Support/ManagedStatic.cpp | ||
16 | index 053493f72fb5..6571580ccecf 100644 | ||
17 | --- a/llvm/lib/Support/ManagedStatic.cpp | ||
18 | +++ b/llvm/lib/Support/ManagedStatic.cpp | ||
19 | @@ -76,8 +76,12 @@ void ManagedStaticBase::destroy() const { | ||
20 | |||
21 | /// llvm_shutdown - Deallocate and destroy all ManagedStatic variables. | ||
22 | void llvm::llvm_shutdown() { | ||
23 | - std::lock_guard<std::recursive_mutex> Lock(*getManagedStaticMutex()); | ||
24 | + getManagedStaticMutex()->lock(); | ||
25 | |||
26 | while (StaticList) | ||
27 | StaticList->destroy(); | ||
28 | + | ||
29 | + getManagedStaticMutex()->unlock(); | ||
30 | + delete ManagedStaticMutex; | ||
31 | + ManagedStaticMutex = nullptr; | ||
32 | } | ||
33 | -- | ||
34 | 2.17.1 | ||
35 | |||
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-Remove-repo-name-in-LLVM-IR.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0004-Remove-repo-name-in-LLVM-IR.patch index 232ae063..f8dec996 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-Remove-repo-name-in-LLVM-IR.patch +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0004-Remove-repo-name-in-LLVM-IR.patch | |||
@@ -1,18 +1,17 @@ | |||
1 | From b53fd86ffdeacb9b13624bdb110fd25e8c35cb92 Mon Sep 17 00:00:00 2001 | 1 | From b794037bf1f90a93efa4c542855ad569cb13b4c5 Mon Sep 17 00:00:00 2001 |
2 | From: Feng Zou <feng.zou@intel.com> | 2 | From: Feng Zou <feng.zou@intel.com> |
3 | Date: Mon, 19 Oct 2020 14:43:38 +0800 | 3 | Date: Mon, 19 Oct 2020 14:43:38 +0800 |
4 | Subject: [PATCH] Remove repo name in LLVM IR | 4 | Subject: [PATCH 4/7] Remove repo name in LLVM IR |
5 | 5 | ||
6 | Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-100/patches/llvm/0003-Remove-repo-name-in-LLVM-IR.patch] | 6 | Upstream-Status: Backport [Taken from opencl-clang patches, https://github.com/intel/opencl-clang/blob/ocl-open-100/patches/llvm/0003-Remove-repo-name-in-LLVM-IR.patch] |
7 | Signed-off-by: Feng Zou <feng.zou@intel.com> | 7 | Signed-off-by: Feng Zou <feng.zou@intel.com> |
8 | Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com> | 8 | Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com> |
9 | |||
10 | --- | 9 | --- |
11 | llvm/cmake/modules/VersionFromVCS.cmake | 23 ++++++++++++----------- | 10 | llvm/cmake/modules/VersionFromVCS.cmake | 23 ++++++++++++----------- |
12 | 1 file changed, 12 insertions(+), 11 deletions(-) | 11 | 1 file changed, 12 insertions(+), 11 deletions(-) |
13 | 12 | ||
14 | diff --git a/llvm/cmake/modules/VersionFromVCS.cmake b/llvm/cmake/modules/VersionFromVCS.cmake | 13 | diff --git a/llvm/cmake/modules/VersionFromVCS.cmake b/llvm/cmake/modules/VersionFromVCS.cmake |
15 | index 1b6519b4b7c..8fd6b23bb34 100644 | 14 | index 1b6519b4b7c4..8fd6b23bb345 100644 |
16 | --- a/llvm/cmake/modules/VersionFromVCS.cmake | 15 | --- a/llvm/cmake/modules/VersionFromVCS.cmake |
17 | +++ b/llvm/cmake/modules/VersionFromVCS.cmake | 16 | +++ b/llvm/cmake/modules/VersionFromVCS.cmake |
18 | @@ -33,17 +33,18 @@ function(get_source_info path revision repository) | 17 | @@ -33,17 +33,18 @@ function(get_source_info path revision repository) |
@@ -46,5 +45,5 @@ index 1b6519b4b7c..8fd6b23bb34 100644 | |||
46 | endif() | 45 | endif() |
47 | endfunction() | 46 | endfunction() |
48 | -- | 47 | -- |
49 | 2.18.1 | 48 | 2.17.1 |
50 | 49 | ||
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0005-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0005-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch new file mode 100644 index 00000000..f8f177e5 --- /dev/null +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0005-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch | |||
@@ -0,0 +1,47 @@ | |||
1 | From 3dd4766499d25e5978a5d90001f18e657e875da0 Mon Sep 17 00:00:00 2001 | ||
2 | From: haonanya <haonan.yang@intel.com> | ||
3 | Date: Thu, 12 Aug 2021 15:48:34 +0800 | ||
4 | Subject: [PATCH 5/7] Remove __IMAGE_SUPPORT__ macro for SPIR since SPIR | ||
5 | doesn't require image support | ||
6 | |||
7 | Upstream-Status: Backport [Taken from opencl-clang patches; https://github.com/intel/opencl-clang/blob/ocl-open-100/patches/clang/0003-Remove-__IMAGE_SUPPORT__-macro-for-SPIR.patch] | ||
8 | |||
9 | Signed-off-by: haonanya <haonan.yang@intel.com> | ||
10 | Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com> | ||
11 | --- | ||
12 | clang/lib/Frontend/InitPreprocessor.cpp | 3 --- | ||
13 | clang/test/Preprocessor/predefined-macros.c | 4 ---- | ||
14 | 2 files changed, 7 deletions(-) | ||
15 | |||
16 | diff --git a/clang/lib/Frontend/InitPreprocessor.cpp b/clang/lib/Frontend/InitPreprocessor.cpp | ||
17 | index aefd208e6cd3..b4a84636673a 100644 | ||
18 | --- a/clang/lib/Frontend/InitPreprocessor.cpp | ||
19 | +++ b/clang/lib/Frontend/InitPreprocessor.cpp | ||
20 | @@ -1108,9 +1108,6 @@ static void InitializePredefinedMacros(const TargetInfo &TI, | ||
21 | if (TI.getSupportedOpenCLOpts().isSupported(#Ext)) \ | ||
22 | Builder.defineMacro(#Ext); | ||
23 | #include "clang/Basic/OpenCLExtensions.def" | ||
24 | - | ||
25 | - if (TI.getTriple().isSPIR()) | ||
26 | - Builder.defineMacro("__IMAGE_SUPPORT__"); | ||
27 | } | ||
28 | |||
29 | if (TI.hasInt128Type() && LangOpts.CPlusPlus && LangOpts.GNUMode) { | ||
30 | diff --git a/clang/test/Preprocessor/predefined-macros.c b/clang/test/Preprocessor/predefined-macros.c | ||
31 | index b088a37ba665..39a222d02faf 100644 | ||
32 | --- a/clang/test/Preprocessor/predefined-macros.c | ||
33 | +++ b/clang/test/Preprocessor/predefined-macros.c | ||
34 | @@ -184,10 +184,6 @@ | ||
35 | // MSCOPE:#define __OPENCL_MEMORY_SCOPE_WORK_GROUP 1 | ||
36 | // MSCOPE:#define __OPENCL_MEMORY_SCOPE_WORK_ITEM 0 | ||
37 | |||
38 | -// RUN: %clang_cc1 %s -E -dM -o - -x cl -triple spir-unknown-unknown \ | ||
39 | -// RUN: | FileCheck -match-full-lines %s --check-prefix=CHECK-SPIR | ||
40 | -// CHECK-SPIR: #define __IMAGE_SUPPORT__ 1 | ||
41 | - | ||
42 | // RUN: %clang_cc1 %s -E -dM -o - -x hip -triple amdgcn-amd-amdhsa \ | ||
43 | // RUN: | FileCheck -match-full-lines %s --check-prefix=CHECK-HIP | ||
44 | // CHECK-HIP-NOT: #define __CUDA_ARCH__ | ||
45 | -- | ||
46 | 2.17.1 | ||
47 | |||
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0006-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0006-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch new file mode 100644 index 00000000..0b4ee8c7 --- /dev/null +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0006-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch | |||
@@ -0,0 +1,53 @@ | |||
1 | From 2c53abd0008bbecfcfe871c6060f4bbf1c94c74a Mon Sep 17 00:00:00 2001 | ||
2 | From: Raphael Isemann <teemperor@gmail.com> | ||
3 | Date: Thu, 1 Apr 2021 18:41:44 +0200 | ||
4 | Subject: [PATCH 6/7] Avoid calling ParseCommandLineOptions in BackendUtil if | ||
5 | possible | ||
6 | |||
7 | Calling `ParseCommandLineOptions` should only be called from `main` as the | ||
8 | CommandLine setup code isn't thread-safe. As BackendUtil is part of the | ||
9 | generic Clang FrontendAction logic, a process which has several threads executing | ||
10 | Clang FrontendActions will randomly crash in the unsafe setup code. | ||
11 | |||
12 | This patch avoids calling the function unless either the debug-pass option or | ||
13 | limit-float-precision option is set. Without these two options set the | ||
14 | `ParseCommandLineOptions` call doesn't do anything beside parsing | ||
15 | the command line `clang` which doesn't set any options. | ||
16 | |||
17 | See also D99652 where LLDB received a workaround for this crash. | ||
18 | |||
19 | Reviewed By: JDevlieghere | ||
20 | |||
21 | Differential Revision: https://reviews.llvm.org/D99740 | ||
22 | |||
23 | Upstream-Status: Backport [Taken from opencl-clang patches; https://github.com/intel/opencl-clang/blob/ocl-open-100/patches/clang/0004-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch] | ||
24 | |||
25 | Signed-off-by: Raphael Isemann <teemperor@gmail.com> | ||
26 | Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com> | ||
27 | --- | ||
28 | clang/lib/CodeGen/BackendUtil.cpp | 8 ++++++++ | ||
29 | 1 file changed, 8 insertions(+) | ||
30 | |||
31 | diff --git a/clang/lib/CodeGen/BackendUtil.cpp b/clang/lib/CodeGen/BackendUtil.cpp | ||
32 | index 0bfcab88a3a9..db8fd4166d7a 100644 | ||
33 | --- a/clang/lib/CodeGen/BackendUtil.cpp | ||
34 | +++ b/clang/lib/CodeGen/BackendUtil.cpp | ||
35 | @@ -743,7 +743,15 @@ static void setCommandLineOpts(const CodeGenOptions &CodeGenOpts) { | ||
36 | BackendArgs.push_back("-limit-float-precision"); | ||
37 | BackendArgs.push_back(CodeGenOpts.LimitFloatPrecision.c_str()); | ||
38 | } | ||
39 | + // Check for the default "clang" invocation that won't set any cl::opt values. | ||
40 | + // Skip trying to parse the command line invocation to avoid the issues | ||
41 | + // described below. | ||
42 | + if (BackendArgs.size() == 1) | ||
43 | + return; | ||
44 | BackendArgs.push_back(nullptr); | ||
45 | + // FIXME: The command line parser below is not thread-safe and shares a global | ||
46 | + // state, so this call might crash or overwrite the options of another Clang | ||
47 | + // instance in the same process. | ||
48 | llvm::cl::ParseCommandLineOptions(BackendArgs.size() - 1, | ||
49 | BackendArgs.data()); | ||
50 | } | ||
51 | -- | ||
52 | 2.17.1 | ||
53 | |||
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0007-support-cl_ext_float_atomics.patch b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0007-support-cl_ext_float_atomics.patch new file mode 100644 index 00000000..f7d191ff --- /dev/null +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/files/llvm10-0007-support-cl_ext_float_atomics.patch | |||
@@ -0,0 +1,377 @@ | |||
1 | From a685de6fc45afcdbe4a7120e9d5b33e175dd71cd Mon Sep 17 00:00:00 2001 | ||
2 | From: haonanya <haonan.yang@intel.com> | ||
3 | Date: Fri, 13 Aug 2021 10:00:02 +0800 | ||
4 | Subject: [PATCH 7/7] support cl_ext_float_atomics | ||
5 | |||
6 | Upstream-Status: Backport [Taken from opencl-clang patches; https://github.com/intel/opencl-clang/blob/ocl-open-100/patches/clang/0005-OpenCL-support-cl_ext_float_atomics.patch] | ||
7 | |||
8 | Signed-off-by: haonanya <haonan.yang@intel.com> | ||
9 | Signed-off-by: Naveen Saini <naveen.kumar.saini@intel.com> | ||
10 | --- | ||
11 | clang/lib/Headers/opencl-c-base.h | 25 ++++ | ||
12 | clang/lib/Headers/opencl-c.h | 208 ++++++++++++++++++++++++++ | ||
13 | clang/test/Headers/opencl-c-header.cl | 96 ++++++++++++ | ||
14 | 3 files changed, 329 insertions(+) | ||
15 | |||
16 | diff --git a/clang/lib/Headers/opencl-c-base.h b/clang/lib/Headers/opencl-c-base.h | ||
17 | index 2cc688ccc3da..86bbee12fdf8 100644 | ||
18 | --- a/clang/lib/Headers/opencl-c-base.h | ||
19 | +++ b/clang/lib/Headers/opencl-c-base.h | ||
20 | @@ -14,6 +14,31 @@ | ||
21 | #define CL_VERSION_3_0 300 | ||
22 | #endif | ||
23 | |||
24 | +#if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) | ||
25 | +// For SPIR all extensions are supported. | ||
26 | +#if defined(__SPIR__) | ||
27 | +#define cl_ext_float_atomics 1 | ||
28 | +#ifdef cl_khr_fp16 | ||
29 | +#define __opencl_c_ext_fp16_global_atomic_load_store 1 | ||
30 | +#define __opencl_c_ext_fp16_local_atomic_load_store 1 | ||
31 | +#define __opencl_c_ext_fp16_global_atomic_add 1 | ||
32 | +#define __opencl_c_ext_fp16_local_atomic_add 1 | ||
33 | +#define __opencl_c_ext_fp16_global_atomic_min_max 1 | ||
34 | +#define __opencl_c_ext_fp16_local_atomic_min_max 1 | ||
35 | +#endif | ||
36 | +#ifdef __opencl_c_fp64 | ||
37 | +#define __opencl_c_ext_fp64_global_atomic_add 1 | ||
38 | +#define __opencl_c_ext_fp64_local_atomic_add 1 | ||
39 | +#define __opencl_c_ext_fp64_global_atomic_min_max 1 | ||
40 | +#define __opencl_c_ext_fp64_local_atomic_min_max 1 | ||
41 | +#endif | ||
42 | +#define __opencl_c_ext_fp32_global_atomic_add 1 | ||
43 | +#define __opencl_c_ext_fp32_local_atomic_add 1 | ||
44 | +#define __opencl_c_ext_fp32_global_atomic_min_max 1 | ||
45 | +#define __opencl_c_ext_fp32_local_atomic_min_max 1 | ||
46 | +#endif // defined(__SPIR__) | ||
47 | +#endif // (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) | ||
48 | + | ||
49 | // Define features for 2.0 for header backward compatibility | ||
50 | #ifndef __opencl_c_int64 | ||
51 | #define __opencl_c_int64 1 | ||
52 | diff --git a/clang/lib/Headers/opencl-c.h b/clang/lib/Headers/opencl-c.h | ||
53 | index 67d900eb1c3d..b463e702d95e 100644 | ||
54 | --- a/clang/lib/Headers/opencl-c.h | ||
55 | +++ b/clang/lib/Headers/opencl-c.h | ||
56 | @@ -14354,6 +14354,214 @@ intptr_t __ovld atomic_fetch_max_explicit( | ||
57 | // defined(cl_khr_int64_extended_atomics) | ||
58 | #endif // (__OPENCL_C_VERSION__ >= CL_VERSION_3_0) | ||
59 | |||
60 | +#if defined(cl_ext_float_atomics) | ||
61 | + | ||
62 | +#if defined(__opencl_c_ext_fp32_global_atomic_min_max) | ||
63 | +float __ovld atomic_fetch_min(volatile __global atomic_float *object, | ||
64 | + float operand); | ||
65 | +float __ovld atomic_fetch_max(volatile __global atomic_float *object, | ||
66 | + float operand); | ||
67 | +float __ovld atomic_fetch_min_explicit(volatile __global atomic_float *object, | ||
68 | + float operand, memory_order order); | ||
69 | +float __ovld atomic_fetch_max_explicit(volatile __global atomic_float *object, | ||
70 | + float operand, memory_order order); | ||
71 | +float __ovld atomic_fetch_min_explicit(volatile __global atomic_float *object, | ||
72 | + float operand, memory_order order, | ||
73 | + memory_scope scope); | ||
74 | +float __ovld atomic_fetch_max_explicit(volatile __global atomic_float *object, | ||
75 | + float operand, memory_order order, | ||
76 | + memory_scope scope); | ||
77 | +#endif // defined(__opencl_c_ext_fp32_global_atomic_min_max) | ||
78 | + | ||
79 | +#if defined(__opencl_c_ext_fp32_local_atomic_min_max) | ||
80 | +float __ovld atomic_fetch_min(volatile __local atomic_float *object, | ||
81 | + float operand); | ||
82 | +float __ovld atomic_fetch_max(volatile __local atomic_float *object, | ||
83 | + float operand); | ||
84 | +float __ovld atomic_fetch_min_explicit(volatile __local atomic_float *object, | ||
85 | + float operand, memory_order order); | ||
86 | +float __ovld atomic_fetch_max_explicit(volatile __local atomic_float *object, | ||
87 | + float operand, memory_order order); | ||
88 | +float __ovld atomic_fetch_min_explicit(volatile __local atomic_float *object, | ||
89 | + float operand, memory_order order, | ||
90 | + memory_scope scope); | ||
91 | +float __ovld atomic_fetch_max_explicit(volatile __local atomic_float *object, | ||
92 | + float operand, memory_order order, | ||
93 | + memory_scope scope); | ||
94 | +#endif // defined(__opencl_c_ext_fp32_local_atomic_min_max) | ||
95 | + | ||
96 | +#if defined(__opencl_c_ext_fp32_global_atomic_min_max) || \ | ||
97 | + defined(__opencl_c_ext_fp32_local_atomic_min_max) | ||
98 | +float __ovld atomic_fetch_min(volatile atomic_float *object, float operand); | ||
99 | +float __ovld atomic_fetch_max(volatile atomic_float *object, float operand); | ||
100 | +float __ovld atomic_fetch_min_explicit(volatile atomic_float *object, | ||
101 | + float operand, memory_order order); | ||
102 | +float __ovld atomic_fetch_max_explicit(volatile atomic_float *object, | ||
103 | + float operand, memory_order order); | ||
104 | +float __ovld atomic_fetch_min_explicit(volatile atomic_float *object, | ||
105 | + float operand, memory_order order, | ||
106 | + memory_scope scope); | ||
107 | +float __ovld atomic_fetch_max_explicit(volatile atomic_float *object, | ||
108 | + float operand, memory_order order, | ||
109 | + memory_scope scope); | ||
110 | +#endif // defined(__opencl_c_ext_fp32_global_atomic_min_max) || \ | ||
111 | + defined(__opencl_c_ext_fp32_local_atomic_min_max) | ||
112 | + | ||
113 | +#if defined(__opencl_c_ext_fp64_global_atomic_min_max) | ||
114 | +double __ovld atomic_fetch_min(volatile __global atomic_double *object, | ||
115 | + double operand); | ||
116 | +double __ovld atomic_fetch_max(volatile __global atomic_double *object, | ||
117 | + double operand); | ||
118 | +double __ovld atomic_fetch_min_explicit(volatile __global atomic_double *object, | ||
119 | + double operand, memory_order order); | ||
120 | +double __ovld atomic_fetch_max_explicit(volatile __global atomic_double *object, | ||
121 | + double operand, memory_order order); | ||
122 | +double __ovld atomic_fetch_min_explicit(volatile __global atomic_double *object, | ||
123 | + double operand, memory_order order, | ||
124 | + memory_scope scope); | ||
125 | +double __ovld atomic_fetch_max_explicit(volatile __global atomic_double *object, | ||
126 | + double operand, memory_order order, | ||
127 | + memory_scope scope); | ||
128 | +#endif // defined(__opencl_c_ext_fp64_global_atomic_min_max) | ||
129 | + | ||
130 | +#if defined(__opencl_c_ext_fp64_local_atomic_min_max) | ||
131 | +double __ovld atomic_fetch_min(volatile __local atomic_double *object, | ||
132 | + double operand); | ||
133 | +double __ovld atomic_fetch_max(volatile __local atomic_double *object, | ||
134 | + double operand); | ||
135 | +double __ovld atomic_fetch_min_explicit(volatile __local atomic_double *object, | ||
136 | + double operand, memory_order order); | ||
137 | +double __ovld atomic_fetch_max_explicit(volatile __local atomic_double *object, | ||
138 | + double operand, memory_order order); | ||
139 | +double __ovld atomic_fetch_min_explicit(volatile __local atomic_double *object, | ||
140 | + double operand, memory_order order, | ||
141 | + memory_scope scope); | ||
142 | +double __ovld atomic_fetch_max_explicit(volatile __local atomic_double *object, | ||
143 | + double operand, memory_order order, | ||
144 | + memory_scope scope); | ||
145 | +#endif // defined(__opencl_c_ext_fp64_local_atomic_min_max) | ||
146 | + | ||
147 | +#if defined(__opencl_c_ext_fp64_global_atomic_min_max) || \ | ||
148 | + defined(__opencl_c_ext_fp64_local_atomic_min_max) | ||
149 | +double __ovld atomic_fetch_min(volatile atomic_double *object, double operand); | ||
150 | +double __ovld atomic_fetch_max(volatile atomic_double *object, double operand); | ||
151 | +double __ovld atomic_fetch_min_explicit(volatile atomic_double *object, | ||
152 | + double operand, memory_order order); | ||
153 | +double __ovld atomic_fetch_max_explicit(volatile atomic_double *object, | ||
154 | + double operand, memory_order order); | ||
155 | +double __ovld atomic_fetch_min_explicit(volatile atomic_double *object, | ||
156 | + double operand, memory_order order, | ||
157 | + memory_scope scope); | ||
158 | +double __ovld atomic_fetch_max_explicit(volatile atomic_double *object, | ||
159 | + double operand, memory_order order, | ||
160 | + memory_scope scope); | ||
161 | +#endif // defined(__opencl_c_ext_fp64_global_atomic_min_max) || \ | ||
162 | + defined(__opencl_c_ext_fp64_local_atomic_min_max) | ||
163 | + | ||
164 | +#if defined(__opencl_c_ext_fp32_global_atomic_add) | ||
165 | +float __ovld atomic_fetch_add(volatile __global atomic_float *object, | ||
166 | + float operand); | ||
167 | +float __ovld atomic_fetch_sub(volatile __global atomic_float *object, | ||
168 | + float operand); | ||
169 | +float __ovld atomic_fetch_add_explicit(volatile __global atomic_float *object, | ||
170 | + float operand, memory_order order); | ||
171 | +float __ovld atomic_fetch_sub_explicit(volatile __global atomic_float *object, | ||
172 | + float operand, memory_order order); | ||
173 | +float __ovld atomic_fetch_add_explicit(volatile __global atomic_float *object, | ||
174 | + float operand, memory_order order, | ||
175 | + memory_scope scope); | ||
176 | +float __ovld atomic_fetch_sub_explicit(volatile __global atomic_float *object, | ||
177 | + float operand, memory_order order, | ||
178 | + memory_scope scope); | ||
179 | +#endif // defined(__opencl_c_ext_fp32_global_atomic_add) | ||
180 | + | ||
181 | +#if defined(__opencl_c_ext_fp32_local_atomic_add) | ||
182 | +float __ovld atomic_fetch_add(volatile __local atomic_float *object, | ||
183 | + float operand); | ||
184 | +float __ovld atomic_fetch_sub(volatile __local atomic_float *object, | ||
185 | + float operand); | ||
186 | +float __ovld atomic_fetch_add_explicit(volatile __local atomic_float *object, | ||
187 | + float operand, memory_order order); | ||
188 | +float __ovld atomic_fetch_sub_explicit(volatile __local atomic_float *object, | ||
189 | + float operand, memory_order order); | ||
190 | +float __ovld atomic_fetch_add_explicit(volatile __local atomic_float *object, | ||
191 | + float operand, memory_order order, | ||
192 | + memory_scope scope); | ||
193 | +float __ovld atomic_fetch_sub_explicit(volatile __local atomic_float *object, | ||
194 | + float operand, memory_order order, | ||
195 | + memory_scope scope); | ||
196 | +#endif // defined(__opencl_c_ext_fp32_local_atomic_add) | ||
197 | + | ||
198 | +#if defined(__opencl_c_ext_fp32_global_atomic_add) || \ | ||
199 | + defined(__opencl_c_ext_fp32_local_atomic_add) | ||
200 | +float __ovld atomic_fetch_add(volatile atomic_float *object, float operand); | ||
201 | +float __ovld atomic_fetch_sub(volatile atomic_float *object, float operand); | ||
202 | +float __ovld atomic_fetch_add_explicit(volatile atomic_float *object, | ||
203 | + float operand, memory_order order); | ||
204 | +float __ovld atomic_fetch_sub_explicit(volatile atomic_float *object, | ||
205 | + float operand, memory_order order); | ||
206 | +float __ovld atomic_fetch_add_explicit(volatile atomic_float *object, | ||
207 | + float operand, memory_order order, | ||
208 | + memory_scope scope); | ||
209 | +float __ovld atomic_fetch_sub_explicit(volatile atomic_float *object, | ||
210 | + float operand, memory_order order, | ||
211 | + memory_scope scope); | ||
212 | +#endif // defined(__opencl_c_ext_fp32_global_atomic_add) || \ | ||
213 | + defined(__opencl_c_ext_fp32_local_atomic_add) | ||
214 | + | ||
215 | +#if defined(__opencl_c_ext_fp64_global_atomic_add) | ||
216 | +double __ovld atomic_fetch_add(volatile __global atomic_double *object, | ||
217 | + double operand); | ||
218 | +double __ovld atomic_fetch_sub(volatile __global atomic_double *object, | ||
219 | + double operand); | ||
220 | +double __ovld atomic_fetch_add_explicit(volatile __global atomic_double *object, | ||
221 | + double operand, memory_order order); | ||
222 | +double __ovld atomic_fetch_sub_explicit(volatile __global atomic_double *object, | ||
223 | + double operand, memory_order order); | ||
224 | +double __ovld atomic_fetch_add_explicit(volatile __global atomic_double *object, | ||
225 | + double operand, memory_order order, | ||
226 | + memory_scope scope); | ||
227 | +double __ovld atomic_fetch_sub_explicit(volatile __global atomic_double *object, | ||
228 | + double operand, memory_order order, | ||
229 | + memory_scope scope); | ||
230 | +#endif // defined(__opencl_c_ext_fp64_global_atomic_add) | ||
231 | + | ||
232 | +#if defined(__opencl_c_ext_fp64_local_atomic_add) | ||
233 | +double __ovld atomic_fetch_add(volatile __local atomic_double *object, | ||
234 | + double operand); | ||
235 | +double __ovld atomic_fetch_sub(volatile __local atomic_double *object, | ||
236 | + double operand); | ||
237 | +double __ovld atomic_fetch_add_explicit(volatile __local atomic_double *object, | ||
238 | + double operand, memory_order order); | ||
239 | +double __ovld atomic_fetch_sub_explicit(volatile __local atomic_double *object, | ||
240 | + double operand, memory_order order); | ||
241 | +double __ovld atomic_fetch_add_explicit(volatile __local atomic_double *object, | ||
242 | + double operand, memory_order order, | ||
243 | + memory_scope scope); | ||
244 | +double __ovld atomic_fetch_sub_explicit(volatile __local atomic_double *object, | ||
245 | + double operand, memory_order order, | ||
246 | + memory_scope scope); | ||
247 | +#endif // defined(__opencl_c_ext_fp64_local_atomic_add) | ||
248 | + | ||
249 | +#if defined(__opencl_c_ext_fp64_global_atomic_add) || \ | ||
250 | + defined(__opencl_c_ext_fp64_local_atomic_add) | ||
251 | +double __ovld atomic_fetch_add(volatile atomic_double *object, double operand); | ||
252 | +double __ovld atomic_fetch_sub(volatile atomic_double *object, double operand); | ||
253 | +double __ovld atomic_fetch_add_explicit(volatile atomic_double *object, | ||
254 | + double operand, memory_order order); | ||
255 | +double __ovld atomic_fetch_sub_explicit(volatile atomic_double *object, | ||
256 | + double operand, memory_order order); | ||
257 | +double __ovld atomic_fetch_add_explicit(volatile atomic_double *object, | ||
258 | + double operand, memory_order order, | ||
259 | + memory_scope scope); | ||
260 | +double __ovld atomic_fetch_sub_explicit(volatile atomic_double *object, | ||
261 | + double operand, memory_order order, | ||
262 | + memory_scope scope); | ||
263 | +#endif // defined(__opencl_c_ext_fp64_global_atomic_add) || \ | ||
264 | + defined(__opencl_c_ext_fp64_local_atomic_add) | ||
265 | + | ||
266 | +#endif // cl_ext_float_atomics | ||
267 | + | ||
268 | // atomic_store() | ||
269 | |||
270 | #if defined(__opencl_c_atomic_scope_device) && \ | ||
271 | diff --git a/clang/test/Headers/opencl-c-header.cl b/clang/test/Headers/opencl-c-header.cl | ||
272 | index 2716076acdcf..7f720cf28142 100644 | ||
273 | --- a/clang/test/Headers/opencl-c-header.cl | ||
274 | +++ b/clang/test/Headers/opencl-c-header.cl | ||
275 | @@ -98,3 +98,99 @@ global atomic_int z = ATOMIC_VAR_INIT(99); | ||
276 | #pragma OPENCL EXTENSION cl_intel_planar_yuv : enable | ||
277 | |||
278 | // CHECK-MOD: Reading modules | ||
279 | + | ||
280 | +// For SPIR all extensions are supported. | ||
281 | +#if defined(__SPIR__) | ||
282 | + | ||
283 | +#if (defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) | ||
284 | + | ||
285 | +#if __opencl_c_ext_fp16_global_atomic_load_store != 1 | ||
286 | +#error "Incorrectly defined __opencl_c_ext_fp16_global_atomic_load_store" | ||
287 | +#endif | ||
288 | +#if __opencl_c_ext_fp16_local_atomic_load_store != 1 | ||
289 | +#error "Incorrectly defined __opencl_c_ext_fp16_local_atomic_load_store" | ||
290 | +#endif | ||
291 | +#if __opencl_c_ext_fp16_global_atomic_add != 1 | ||
292 | +#error "Incorrectly defined __opencl_c_ext_fp16_global_atomic_add" | ||
293 | +#endif | ||
294 | +#if __opencl_c_ext_fp32_global_atomic_add != 1 | ||
295 | +#error "Incorrectly defined __opencl_c_ext_fp32_global_atomic_add" | ||
296 | +#endif | ||
297 | +#if __opencl_c_ext_fp64_global_atomic_add != 1 | ||
298 | +#error "Incorrectly defined __opencl_c_ext_fp64_global_atomic_add" | ||
299 | +#endif | ||
300 | +#if __opencl_c_ext_fp16_local_atomic_add != 1 | ||
301 | +#error "Incorrectly defined __opencl_c_ext_fp16_local_atomic_add" | ||
302 | +#endif | ||
303 | +#if __opencl_c_ext_fp32_local_atomic_add != 1 | ||
304 | +#error "Incorrectly defined __opencl_c_ext_fp32_local_atomic_add" | ||
305 | +#endif | ||
306 | +#if __opencl_c_ext_fp64_local_atomic_add != 1 | ||
307 | +#error "Incorrectly defined __opencl_c_ext_fp64_local_atomic_add" | ||
308 | +#endif | ||
309 | +#if __opencl_c_ext_fp16_global_atomic_min_max != 1 | ||
310 | +#error "Incorrectly defined __opencl_c_ext_fp16_global_atomic_min_max" | ||
311 | +#endif | ||
312 | +#if __opencl_c_ext_fp32_global_atomic_min_max != 1 | ||
313 | +#error "Incorrectly defined __opencl_c_ext_fp32_global_atomic_min_max" | ||
314 | +#endif | ||
315 | +#if __opencl_c_ext_fp64_global_atomic_min_max != 1 | ||
316 | +#error "Incorrectly defined __opencl_c_ext_fp64_global_atomic_min_max" | ||
317 | +#endif | ||
318 | +#if __opencl_c_ext_fp16_local_atomic_min_max != 1 | ||
319 | +#error "Incorrectly defined __opencl_c_ext_fp16_local_atomic_min_max" | ||
320 | +#endif | ||
321 | +#if __opencl_c_ext_fp32_local_atomic_min_max != 1 | ||
322 | +#error "Incorrectly defined __opencl_c_ext_fp32_local_atomic_min_max" | ||
323 | +#endif | ||
324 | +#if __opencl_c_ext_fp64_local_atomic_min_max != 1 | ||
325 | +#error "Incorrectly defined __opencl_c_ext_fp64_local_atomic_min_max" | ||
326 | +#endif | ||
327 | +#else | ||
328 | + | ||
329 | +#ifdef __opencl_c_ext_fp16_global_atomic_load_store | ||
330 | +#error "Incorrectly __opencl_c_ext_fp16_global_atomic_load_store defined" | ||
331 | +#endif | ||
332 | +#ifdef __opencl_c_ext_fp16_local_atomic_load_store | ||
333 | +#error "Incorrectly __opencl_c_ext_fp16_local_atomic_load_store defined" | ||
334 | +#endif | ||
335 | +#ifdef __opencl_c_ext_fp16_global_atomic_add | ||
336 | +#error "Incorrectly __opencl_c_ext_fp16_global_atomic_add defined" | ||
337 | +#endif | ||
338 | +#ifdef __opencl_c_ext_fp32_global_atomic_add | ||
339 | +#error "Incorrectly __opencl_c_ext_fp32_global_atomic_add defined" | ||
340 | +#endif | ||
341 | +#ifdef __opencl_c_ext_fp64_global_atomic_add | ||
342 | +#error "Incorrectly __opencl_c_ext_fp64_global_atomic_add defined" | ||
343 | +#endif | ||
344 | +#ifdef __opencl_c_ext_fp16_local_atomic_add | ||
345 | +#error "Incorrectly __opencl_c_ext_fp16_local_atomic_add defined" | ||
346 | +#endif | ||
347 | +#ifdef __opencl_c_ext_fp32_local_atomic_add | ||
348 | +#error "Incorrectly __opencl_c_ext_fp32_local_atomic_add defined" | ||
349 | +#endif | ||
350 | +#ifdef __opencl_c_ext_fp64_local_atomic_add | ||
351 | +#error "Incorrectly __opencl_c_ext_fp64_local_atomic_add defined" | ||
352 | +#endif | ||
353 | +#ifdef __opencl_c_ext_fp16_global_atomic_min_max | ||
354 | +#error "Incorrectly __opencl_c_ext_fp16_global_atomic_min_max defined" | ||
355 | +#endif | ||
356 | +#ifdef __opencl_c_ext_fp32_global_atomic_min_max | ||
357 | +#error "Incorrectly __opencl_c_ext_fp32_global_atomic_min_max defined" | ||
358 | +#endif | ||
359 | +#ifdef __opencl_c_ext_fp64_global_atomic_min_max | ||
360 | +#error "Incorrectly __opencl_c_ext_fp64_global_atomic_min_max defined" | ||
361 | +#endif | ||
362 | +#ifdef __opencl_c_ext_fp16_local_atomic_min_max | ||
363 | +#error "Incorrectly __opencl_c_ext_fp16_local_atomic_min_max defined" | ||
364 | +#endif | ||
365 | +#ifdef __opencl_c_ext_fp32_local_atomic_min_max | ||
366 | +#error "Incorrectly __opencl_c_ext_fp32_local_atomic_min_max defined" | ||
367 | +#endif | ||
368 | +#ifdef __opencl_c_ext_fp64_local_atomic_min_max | ||
369 | +#error "Incorrectly __opencl_c_ext_fp64_local_atomic_min_max defined" | ||
370 | +#endif | ||
371 | + | ||
372 | +#endif //(defined(__OPENCL_CPP_VERSION__) || __OPENCL_C_VERSION__ >= 200) | ||
373 | + | ||
374 | +#endif // defined(__SPIR__) | ||
375 | -- | ||
376 | 2.17.1 | ||
377 | |||
diff --git a/dynamic-layers/clang-layer/recipes-devtools/clang/llvm-project-source.bbappend b/dynamic-layers/clang-layer/recipes-devtools/clang/llvm-project-source.bbappend index a09343b3..ac34321c 100644 --- a/dynamic-layers/clang-layer/recipes-devtools/clang/llvm-project-source.bbappend +++ b/dynamic-layers/clang-layer/recipes-devtools/clang/llvm-project-source.bbappend | |||
@@ -1,20 +1,23 @@ | |||
1 | FILESEXTRAPATHS:prepend:intel-x86-common := "${THISDIR}/files:" | 1 | FILESEXTRAPATHS:prepend:intel-x86-common := "${THISDIR}/files:" |
2 | 2 | ||
3 | SPIRV10_SRCREV = "576abae62cecd171992017a4a786e3831221ab8d" | 3 | SPIRV10_SRCREV = "fe4d6b767363a1995ccbfca27f79efb10dcfe110" |
4 | SPIRV11_SRCREV = "2a8c1e6c9778deaa720a23e08c293006dc5d56fd" | 4 | SPIRV11_SRCREV = "2a8c1e6c9778deaa720a23e08c293006dc5d56fd" |
5 | 5 | ||
6 | SPIRV_SRCREV = "${@bb.utils.contains('LLVMVERSION', '10.0.1', '${SPIRV10_SRCREV}', '${SPIRV11_SRCREV}', d)}" | 6 | SPIRV_SRCREV = "${@bb.utils.contains('LLVMVERSION', '10.0.1', '${SPIRV10_SRCREV}', '${SPIRV11_SRCREV}', d)}" |
7 | 7 | ||
8 | SRC_URI_LLVM10_PATCHES = " \ | 8 | SRC_URI_LLVM10_PATCHES = " \ |
9 | file://llvm10-skip-building-tests.patch;patchdir=llvm/projects/llvm-spirv \ | 9 | file://llvm10-0001-llvm-spirv-skip-building-tests.patch;patchdir=llvm/projects/llvm-spirv \ |
10 | file://fix-shared-libs.patch;patchdir=llvm/projects/llvm-spirv \ | 10 | file://llvm10-0002-Fix-building-in-tree-with-cmake-DLLVM_LINK_LLVM_DYLI.patch;patchdir=llvm/projects/llvm-spirv \ |
11 | file://llvm10-0003-Add-support-for-cl_ext_float_atomics-in-SPIRVWriter.patch;patchdir=llvm/projects/llvm-spirv \ | ||
11 | file://BasicBlockUtils-Add-metadata-fixing-in-SplitBlockPre.patch;patchdir=llvm \ | 12 | file://BasicBlockUtils-Add-metadata-fixing-in-SplitBlockPre.patch;patchdir=llvm \ |
12 | file://IndVarSimplify-Do-not-use-SCEV-expander-for-IVCount-.patch;patchdir=llvm \ | 13 | file://IndVarSimplify-Do-not-use-SCEV-expander-for-IVCount-.patch;patchdir=llvm \ |
13 | file://llvm10-OpenCL-3.0-support.patch \ | 14 | file://llvm10-0001-OpenCL-3.0-support.patch \ |
14 | file://0002-Add-cl_khr_extended_subgroup-extensions.patch \ | 15 | file://llvm10-0002-Add-cl_khr_extended_subgroup-extensions.patch \ |
15 | file://0001-Memory-leak-fix-for-Managed-Static-Mutex.patch \ | 16 | file://llvm10-0003-Memory-leak-fix-for-Managed-Static-Mutex.patch \ |
16 | file://llvm10-Remove-repo-name-in-LLVM-IR.patch \ | 17 | file://llvm10-0004-Remove-repo-name-in-LLVM-IR.patch \ |
17 | file://0001-Fix-debug-info-of-work-item-builtin-translation-745.patch;patchdir=llvm/projects/llvm-spirv \ | 18 | file://llvm10-0005-Remove-__IMAGE_SUPPORT__-macro-for-SPIR-since-SPIR-d.patch \ |
19 | file://llvm10-0006-Avoid-calling-ParseCommandLineOptions-in-BackendUtil.patch \ | ||
20 | file://llvm10-0007-support-cl_ext_float_atomics.patch \ | ||
18 | " | 21 | " |
19 | 22 | ||
20 | SRC_URI_LLVM11_PATCHES = " \ | 23 | SRC_URI_LLVM11_PATCHES = " \ |