Skip to content

Commit ebe9758

Browse files
hvdijkhvdijk
authored
[SYCL][NativeCPU] Move __spirv_AtomicF{Add,Min,Max}EXT into libclc. (#19365)
1 parent 6ccaafd commit ebe9758

File tree

9 files changed

+132
-123
lines changed

9 files changed

+132
-123
lines changed

‎libclc/libspirv/lib/native_cpu/SOURCES‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -11,6 +11,9 @@ math/native_log2.cl
1111
math/native_sin.cl
1212
math/native_sqrt.cl
1313
math/round.cl
14+
atomic/atomic_add.ll
15+
atomic/atomic_max.ll
16+
atomic/atomic_min.ll
1417
workitem/get_global_id.cl
1518
workitem/get_global_size.cl
1619
workitem/get_num_sub_groups.cl

‎libclc/libspirv/lib/native_cpu/atomic/atomic_add.ll‎

Lines changed: 43 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,43 @@
1+
;;===----------------------------------------------------------------------===;;
2+
;
3+
; Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
; See https://llvm.org/LICENSE.txt for license information.
5+
; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
;
7+
;;===----------------------------------------------------------------------===;;
8+
9+
define float @_Z21__spirv_AtomicFAddEXTPfiif(ptr noundef %p, i32 noundef %scope, i32 noundef %semantics, float noundef %val) nounwind alwaysinline {
10+
entry:
11+
%0 = atomicrmw fadd ptr %p, float %val seq_cst, align 4
12+
ret float %0
13+
}
14+
15+
define float @_Z21__spirv_AtomicFAddEXTPU3AS1fiif(ptr addrspace(1) noundef %p, i32 noundef %scope, i32 noundef %semantics, float noundef %val) nounwind alwaysinline {
16+
entry:
17+
%0 = atomicrmw fadd ptr addrspace(1) %p, float %val seq_cst, align 4
18+
ret float %0
19+
}
20+
21+
define float @_Z21__spirv_AtomicFAddEXTPU3AS3fiif(ptr addrspace(3) noundef %p, i32 noundef %scope, i32 noundef %semantics, float noundef %val) nounwind alwaysinline {
22+
entry:
23+
%0 = atomicrmw fadd ptr addrspace(3) %p, float %val seq_cst, align 4
24+
ret float %0
25+
}
26+
27+
define double @_Z21__spirv_AtomicFAddEXTPdiid(ptr noundef %p, i32 noundef %scope, i32 noundef %semantics, double noundef %val) nounwind alwaysinline {
28+
entry:
29+
%0 = atomicrmw fadd ptr %p, double %val seq_cst, align 8
30+
ret double %0
31+
}
32+
33+
define double @_Z21__spirv_AtomicFAddEXTPU3AS1diid(ptr addrspace(1) noundef %p, i32 noundef %scope, i32 noundef %semantics, double noundef %val) nounwind alwaysinline {
34+
entry:
35+
%0 = atomicrmw fadd ptr addrspace(1) %p, double %val seq_cst, align 8
36+
ret double %0
37+
}
38+
39+
define double @_Z21__spirv_AtomicFAddEXTPU3AS3diid(ptr addrspace(3) noundef %p, i32 noundef %scope, i32 noundef %semantics, double noundef %val) nounwind alwaysinline {
40+
entry:
41+
%0 = atomicrmw fadd ptr addrspace(3) %p, double %val seq_cst, align 8
42+
ret double %0
43+
}

‎libclc/libspirv/lib/native_cpu/atomic/atomic_max.ll‎

Lines changed: 43 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,43 @@
1+
;;===----------------------------------------------------------------------===;;
2+
;
3+
; Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
; See https://llvm.org/LICENSE.txt for license information.
5+
; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
;
7+
;;===----------------------------------------------------------------------===;;
8+
9+
define float @_Z21__spirv_AtomicFMaxEXTPfiif(ptr noundef %p, i32 noundef %scope, i32 noundef %semantics, float noundef %val) nounwind alwaysinline {
10+
entry:
11+
%0 = atomicrmw fmax ptr %p, float %val seq_cst, align 4
12+
ret float %0
13+
}
14+
15+
define float @_Z21__spirv_AtomicFMaxEXTPU3AS1fiif(ptr addrspace(1) noundef %p, i32 noundef %scope, i32 noundef %semantics, float noundef %val) nounwind alwaysinline {
16+
entry:
17+
%0 = atomicrmw fmax ptr addrspace(1) %p, float %val seq_cst, align 4
18+
ret float %0
19+
}
20+
21+
define float @_Z21__spirv_AtomicFMaxEXTPU3AS3fiif(ptr addrspace(3) noundef %p, i32 noundef %scope, i32 noundef %semantics, float noundef %val) nounwind alwaysinline {
22+
entry:
23+
%0 = atomicrmw fmax ptr addrspace(3) %p, float %val seq_cst, align 4
24+
ret float %0
25+
}
26+
27+
define double @_Z21__spirv_AtomicFMaxEXTPdiid(ptr noundef %p, i32 noundef %scope, i32 noundef %semantics, double noundef %val) nounwind alwaysinline {
28+
entry:
29+
%0 = atomicrmw fmax ptr %p, double %val seq_cst, align 8
30+
ret double %0
31+
}
32+
33+
define double @_Z21__spirv_AtomicFMaxEXTPU3AS1diid(ptr addrspace(1) noundef %p, i32 noundef %scope, i32 noundef %semantics, double noundef %val) nounwind alwaysinline {
34+
entry:
35+
%0 = atomicrmw fmax ptr addrspace(1) %p, double %val seq_cst, align 8
36+
ret double %0
37+
}
38+
39+
define double @_Z21__spirv_AtomicFMaxEXTPU3AS3diid(ptr addrspace(3) noundef %p, i32 noundef %scope, i32 noundef %semantics, double noundef %val) nounwind alwaysinline {
40+
entry:
41+
%0 = atomicrmw fmax ptr addrspace(3) %p, double %val seq_cst, align 8
42+
ret double %0
43+
}

‎libclc/libspirv/lib/native_cpu/atomic/atomic_min.ll‎

Lines changed: 43 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,43 @@
1+
;;===----------------------------------------------------------------------===;;
2+
;
3+
; Part of the LLVM Project, under the Apache License v2.0 with LLVM Exceptions.
4+
; See https://llvm.org/LICENSE.txt for license information.
5+
; SPDX-License-Identifier: Apache-2.0 WITH LLVM-exception
6+
;
7+
;;===----------------------------------------------------------------------===;;
8+
9+
define float @_Z21__spirv_AtomicFMinEXTPfiif(ptr noundef %p, i32 noundef %scope, i32 noundef %semantics, float noundef %val) nounwind alwaysinline {
10+
entry:
11+
%0 = atomicrmw fmin ptr %p, float %val seq_cst, align 4
12+
ret float %0
13+
}
14+
15+
define float @_Z21__spirv_AtomicFMinEXTPU3AS1fiif(ptr addrspace(1) noundef %p, i32 noundef %scope, i32 noundef %semantics, float noundef %val) nounwind alwaysinline {
16+
entry:
17+
%0 = atomicrmw fmin ptr addrspace(1) %p, float %val seq_cst, align 4
18+
ret float %0
19+
}
20+
21+
define float @_Z21__spirv_AtomicFMinEXTPU3AS3fiif(ptr addrspace(3) noundef %p, i32 noundef %scope, i32 noundef %semantics, float noundef %val) nounwind alwaysinline {
22+
entry:
23+
%0 = atomicrmw fmin ptr addrspace(3) %p, float %val seq_cst, align 4
24+
ret float %0
25+
}
26+
27+
define double @_Z21__spirv_AtomicFMinEXTPdiid(ptr noundef %p, i32 noundef %scope, i32 noundef %semantics, double noundef %val) nounwind alwaysinline {
28+
entry:
29+
%0 = atomicrmw fmin ptr %p, double %val seq_cst, align 8
30+
ret double %0
31+
}
32+
33+
define double @_Z21__spirv_AtomicFMinEXTPU3AS1diid(ptr addrspace(1) noundef %p, i32 noundef %scope, i32 noundef %semantics, double noundef %val) nounwind alwaysinline {
34+
entry:
35+
%0 = atomicrmw fmin ptr addrspace(1) %p, double %val seq_cst, align 8
36+
ret double %0
37+
}
38+
39+
define double @_Z21__spirv_AtomicFMinEXTPU3AS3diid(ptr addrspace(3) noundef %p, i32 noundef %scope, i32 noundef %semantics, double noundef %val) nounwind alwaysinline {
40+
entry:
41+
%0 = atomicrmw fmin ptr addrspace(3) %p, double %val seq_cst, align 8
42+
ret double %0
43+
}

‎llvm/include/llvm/SYCLLowerIR/FAtomicsNativeCPU.h‎

Lines changed: 0 additions & 28 deletions
This file was deleted.

‎llvm/lib/SYCLNativeCPUUtils/CMakeLists.txt‎

Lines changed: 0 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,7 +3,6 @@ add_llvm_component_library(LLVMSYCLNativeCPUUtils
33
PrepareSYCLNativeCPU.cpp
44
RenameKernelSYCLNativeCPU.cpp
55
ConvertToMuxBuiltinsSYCLNativeCPU.cpp
6-
FAtomicsNativeCPU.cpp
76

87
ADDITIONAL_HEADER_DIRS
98
${LLVM_MAIN_INCLUDE_DIR}/llvm/SYCLLowerIR

‎llvm/lib/SYCLNativeCPUUtils/FAtomicsNativeCPU.cpp‎

Lines changed: 0 additions & 51 deletions
This file was deleted.

‎llvm/lib/SYCLNativeCPUUtils/PipelineSYCLNativeCPU.cpp‎

Lines changed: 0 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -12,7 +12,6 @@
1212
//
1313
//===----------------------------------------------------------------------===//
1414
#include "llvm/SYCLLowerIR/ConvertToMuxBuiltinsSYCLNativeCPU.h"
15-
#include "llvm/SYCLLowerIR/FAtomicsNativeCPU.h"
1615
#include "llvm/SYCLLowerIR/PrepareSYCLNativeCPU.h"
1716
#include "llvm/SYCLLowerIR/RenameKernelSYCLNativeCPU.h"
1817
#include "llvm/SYCLLowerIR/SpecConstants.h"
@@ -70,7 +69,6 @@ void llvm::sycl::utils::addSYCLNativeCPUBackendPasses(
7069
OptimizationLevel OptLevel) {
7170
MPM.addPass(SpecConstantsPass(SpecConstantsPass::HandlingMode::emulation));
7271
MPM.addPass(ConvertToMuxBuiltinsSYCLNativeCPUPass());
73-
MPM.addPass(FAtomicsNativeCPU());
7472
#ifdef NATIVECPU_USE_OCK
7573
MPM.addPass(compiler::utils::PrepareBarriersPass());
7674
MPM.addPass(compiler::utils::TransferKernelMetadataPass());

‎sycl/test/check_device_code/native_cpu/fp_atomic.cpp‎

Lines changed: 0 additions & 41 deletions
This file was deleted.

0 commit comments

Comments
 (0)