Skip to content

Commit 6b1c51b

Browse files
TIFitisjsjodin
TIFitis
and
jsjodin
authored
[OpenMP] Migrate GPU Reductions CodeGen from Clang to OMPIRBuilder (llvm#80343)
This patch migrates the CGOpenMPRuntimeGPU::emitReduction and related functions to the OpenMPIRBUilder. In future patches MLIR OpenMP translation would be making use of these functions. Co-authored-by: Jan Leyonberg <[email protected]>
1 parent 22b36bf commit 6b1c51b

File tree

12 files changed

+2716
-1834
lines changed

12 files changed

+2716
-1834
lines changed

‎clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp

Lines changed: 74 additions & 1296 deletions
Large diffs are not rendered by default.

‎clang/lib/CodeGen/CGOpenMPRuntimeGPU.h

Lines changed: 0 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -150,9 +150,6 @@ class CGOpenMPRuntimeGPU : public CGOpenMPRuntime {
150150
CodeGenFunction &CGF,
151151
const std::pair<llvm::Value *, llvm::Value *> &AddrSizePair) override;
152152

153-
/// Get the GPU warp size.
154-
llvm::Value *getGPUWarpSize(CodeGenFunction &CGF);
155-
156153
/// Get the id of the current thread on the GPU.
157154
llvm::Value *getGPUThreadID(CodeGenFunction &CGF);
158155

‎clang/test/OpenMP/nvptx_target_parallel_reduction_codegen.cpp

Lines changed: 93 additions & 72 deletions
Large diffs are not rendered by default.

‎clang/test/OpenMP/nvptx_target_parallel_reduction_codegen_tbaa_PR46146.cpp

Lines changed: 283 additions & 281 deletions
Large diffs are not rendered by default.

‎clang/test/OpenMP/nvptx_teams_reduction_codegen.cpp

Lines changed: 144 additions & 114 deletions
Large diffs are not rendered by default.

‎clang/test/OpenMP/reduction_complex.c

Lines changed: 96 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,96 @@
1+
// NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --replace-value-regex "__omp_offloading_[0-9a-z]+_[0-9a-z]+" "reduction_size[.].+[.]" "pl_cond[.].+[.|,]" --prefix-filecheck-ir-name _
2+
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ \
3+
// RUN: -triple powerpc64le-unknown-unknown \
4+
// RUN: -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o \
5+
// RUN: %t-ppc-host.bc
6+
7+
// RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ \
8+
// RUN: -triple nvptx64-unknown-unknown -DCUA \
9+
// RUN: -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s \
10+
// RUN: -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc \
11+
// RUN: -o - | FileCheck %s --check-prefix CHECK
12+
13+
// expected-no-diagnostics
14+
int foo() {
15+
int i;
16+
int j;
17+
_Complex float sum = 0;
18+
19+
#pragma omp target teams loop reduction(+:sum) collapse(2) bind(parallel) order(concurrent) lastprivate(j) map(tofrom:sum)
20+
21+
for(i=0; i<10; i++)
22+
for(j=0; j<10; j++)
23+
sum += i;
24+
25+
return 0;
26+
}
27+
// CHECK-LABEL: define {{[^@]+}}@_omp_reduction_shuffle_and_reduce_func
28+
// CHECK-SAME: (ptr noundef [[TMP0:%.*]], i16 noundef signext [[TMP1:%.*]], i16 noundef signext [[TMP2:%.*]], i16 noundef signext [[TMP3:%.*]]) #[[ATTR2:[0-9]+]] {
29+
// CHECK-NEXT: entry:
30+
// CHECK-NEXT: %[[VAL_228:.*]] = alloca ptr, align 8
31+
// CHECK-NEXT: %[[VAL_229:.*]] = alloca i16, align 2
32+
// CHECK-NEXT: %[[VAL_230:.*]] = alloca i16, align 2
33+
// CHECK-NEXT: %[[VAL_231:.*]] = alloca i16, align 2
34+
// CHECK-NEXT: %[[VAL_232:.*]] = alloca [1 x ptr], align 8
35+
// CHECK-NEXT: %[[VAL_233:.*]] = alloca { float, float }, align 8
36+
// CHECK-NEXT: store ptr %[[VAL_234:.*]], ptr %[[VAL_228]], align 8
37+
// CHECK-NEXT: store i16 %[[VAL_235:.*]], ptr %[[VAL_229]], align 2
38+
// CHECK-NEXT: store i16 %[[VAL_236:.*]], ptr %[[VAL_230]], align 2
39+
// CHECK-NEXT: store i16 %[[VAL_237:.*]], ptr %[[VAL_231]], align 2
40+
// CHECK-NEXT: %[[VAL_238:.*]] = load ptr, ptr %[[VAL_228]], align 8
41+
// CHECK-NEXT: %[[VAL_239:.*]] = load i16, ptr %[[VAL_229]], align 2
42+
// CHECK-NEXT: %[[VAL_240:.*]] = load i16, ptr %[[VAL_230]], align 2
43+
// CHECK-NEXT: %[[VAL_241:.*]] = load i16, ptr %[[VAL_231]], align 2
44+
// CHECK-NEXT: %[[VAL_242:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_238]], i64 0, i64 0
45+
// CHECK-NEXT: %[[VAL_243:.*]] = load ptr, ptr %[[VAL_242]], align 8
46+
// CHECK-NEXT: %[[VAL_244:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_232]], i64 0, i64 0
47+
// CHECK-NEXT: %[[VAL_245:.*]] = getelementptr { float, float }, ptr %[[VAL_243]], i64 1
48+
// CHECK-NEXT: %[[VAL_246:.*]] = load i64, ptr %[[VAL_243]], align 8
49+
// CHECK-NEXT: %[[VAL_247:.*]] = call i32 @__kmpc_get_warp_size()
50+
// CHECK-NEXT: %[[VAL_248:.*]] = trunc i32 %[[VAL_247]] to i16
51+
// CHECK-NEXT: %[[VAL_249:.*]] = call i64 @__kmpc_shuffle_int64(i64 %[[VAL_246]], i16 %[[VAL_240]], i16 %[[VAL_248]])
52+
// CHECK-NEXT: store i64 %[[VAL_249]], ptr %[[VAL_233]], align 8
53+
// CHECK-NEXT: %[[VAL_250:.*]] = getelementptr i64, ptr %[[VAL_243]], i64 1
54+
// CHECK-NEXT: %[[VAL_251:.*]] = getelementptr i64, ptr %[[VAL_233]], i64 1
55+
// CHECK-NEXT: store ptr %[[VAL_233]], ptr %[[VAL_244]], align 8
56+
// CHECK-NEXT: %[[VAL_252:.*]] = icmp eq i16 %[[VAL_241]], 0
57+
// CHECK-NEXT: %[[VAL_253:.*]] = icmp eq i16 %[[VAL_241]], 1
58+
// CHECK-NEXT: %[[VAL_254:.*]] = icmp ult i16 %[[VAL_239]], %[[VAL_240]]
59+
// CHECK-NEXT: %[[VAL_255:.*]] = and i1 %[[VAL_253]], %[[VAL_254]]
60+
// CHECK-NEXT: %[[VAL_256:.*]] = icmp eq i16 %[[VAL_241]], 2
61+
// CHECK-NEXT: %[[VAL_257:.*]] = and i16 %[[VAL_239]], 1
62+
// CHECK-NEXT: %[[VAL_258:.*]] = icmp eq i16 %[[VAL_257]], 0
63+
// CHECK-NEXT: %[[VAL_259:.*]] = and i1 %[[VAL_256]], %[[VAL_258]]
64+
// CHECK-NEXT: %[[VAL_260:.*]] = icmp sgt i16 %[[VAL_240]], 0
65+
// CHECK-NEXT: %[[VAL_261:.*]] = and i1 %[[VAL_259]], %[[VAL_260]]
66+
// CHECK-NEXT: %[[VAL_262:.*]] = or i1 %[[VAL_252]], %[[VAL_255]]
67+
// CHECK-NEXT: %[[VAL_263:.*]] = or i1 %[[VAL_262]], %[[VAL_261]]
68+
// CHECK-NEXT: br i1 %[[VAL_263]], label %[[VAL_264:.*]], label %[[VAL_265:.*]]
69+
// CHECK: then: ; preds = %[[VAL_266:.*]]
70+
// CHECK-NEXT: call void @"{{__omp_offloading_[0-9a-z]+_[0-9a-z]+}}__Z3foov_l{{[0-9]+}}_omp_outlined_omp_outlined_omp$reduction$reduction_func"(ptr %[[VAL_238]], ptr %[[VAL_232]]) #2
71+
// CHECK-NEXT: br label %[[VAL_267:.*]]
72+
// CHECK: else: ; preds = %[[VAL_266]]
73+
// CHECK-NEXT: br label %[[VAL_267]]
74+
// CHECK: ifcont: ; preds = %[[VAL_265]], %[[VAL_264]]
75+
// CHECK-NEXT: %[[VAL_268:.*]] = icmp eq i16 %[[VAL_241]], 1
76+
// CHECK-NEXT: %[[VAL_269:.*]] = icmp uge i16 %[[VAL_239]], %[[VAL_240]]
77+
// CHECK-NEXT: %[[VAL_270:.*]] = and i1 %[[VAL_268]], %[[VAL_269]]
78+
// CHECK-NEXT: br i1 %[[VAL_270]], label %[[VAL_271:.*]], label %[[VAL_272:.*]]
79+
// CHECK: then4: ; preds = %[[VAL_267]]
80+
// CHECK-NEXT: %[[VAL_273:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_232]], i64 0, i64 0
81+
// CHECK-NEXT: %[[VAL_274:.*]] = load ptr, ptr %[[VAL_273]], align 8
82+
// CHECK-NEXT: %[[VAL_275:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_238]], i64 0, i64 0
83+
// CHECK-NEXT: %[[VAL_276:.*]] = load ptr, ptr %[[VAL_275]], align 8
84+
// CHECK-NEXT: %[[VAL_277:.*]] = getelementptr inbounds { float, float }, ptr %[[VAL_274]], i32 0, i32 0
85+
// CHECK-NEXT: %[[VAL_278:.*]] = load float, ptr %[[VAL_277]], align 4
86+
// CHECK-NEXT: %[[VAL_279:.*]] = getelementptr inbounds { float, float }, ptr %[[VAL_274]], i32 0, i32 1
87+
// CHECK-NEXT: %[[VAL_280:.*]] = load float, ptr %[[VAL_279]], align 4
88+
// CHECK-NEXT: %[[VAL_281:.*]] = getelementptr inbounds { float, float }, ptr %[[VAL_276]], i32 0, i32 0
89+
// CHECK-NEXT: %[[VAL_282:.*]] = getelementptr inbounds { float, float }, ptr %[[VAL_276]], i32 0, i32 1
90+
// CHECK-NEXT: store float %[[VAL_278]], ptr %[[VAL_281]], align 4
91+
// CHECK-NEXT: store float %[[VAL_280]], ptr %[[VAL_282]], align 4
92+
// CHECK-NEXT: br label %[[VAL_283:.*]]
93+
// CHECK: else7: ; preds = %[[VAL_267]]
94+
// CHECK-NEXT: br label %[[VAL_283]]
95+
// CHECK: ifcont8: ; preds = %[[VAL_272]], %[[VAL_271]]
96+
// CHECK-NEXT: ret void

‎clang/test/OpenMP/reduction_implicit_map.cpp

Lines changed: 5 additions & 4 deletions
Original file line numberDiff line numberDiff line change
@@ -233,7 +233,6 @@ int main()
233233
// CHECK-NEXT: [[DOTADDR:%.*]] = alloca ptr, align 8
234234
// CHECK-NEXT: [[DOTADDR1:%.*]] = alloca i32, align 4
235235
// CHECK-NEXT: [[DOTCNT_ADDR:%.*]] = alloca i32, align 4
236-
// CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
237236
// CHECK-NEXT: store ptr [[TMP0]], ptr [[DOTADDR]], align 8
238237
// CHECK-NEXT: store i32 [[TMP1]], ptr [[DOTADDR1]], align 4
239238
// CHECK-NEXT: [[TMP3:%.*]] = call i32 @__kmpc_get_hardware_thread_id_in_block()
@@ -249,6 +248,7 @@ int main()
249248
// CHECK-NEXT: [[TMP8:%.*]] = icmp ult i32 [[TMP7]], 2
250249
// CHECK-NEXT: br i1 [[TMP8]], label [[BODY:%.*]], label [[EXIT:%.*]]
251250
// CHECK: body:
251+
// CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
252252
// CHECK-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2:[0-9]+]], i32 [[TMP2]])
253253
// CHECK-NEXT: [[WARP_MASTER:%.*]] = icmp eq i32 [[NVPTX_LANE_ID]], 0
254254
// CHECK-NEXT: br i1 [[WARP_MASTER]], label [[THEN:%.*]], label [[ELSE:%.*]]
@@ -263,21 +263,22 @@ int main()
263263
// CHECK: else:
264264
// CHECK-NEXT: br label [[IFCONT]]
265265
// CHECK: ifcont:
266+
// CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_global_thread_num(ptr @[[GLOB1]])
266267
// CHECK-NEXT: call void @__kmpc_barrier(ptr @[[GLOB2]], i32 [[TMP2]])
267268
// CHECK-NEXT: [[TMP14:%.*]] = load i32, ptr [[DOTADDR1]], align 4
268269
// CHECK-NEXT: [[IS_ACTIVE_THREAD:%.*]] = icmp ult i32 [[TMP3]], [[TMP14]]
269270
// CHECK-NEXT: br i1 [[IS_ACTIVE_THREAD]], label [[THEN2:%.*]], label [[ELSE3:%.*]]
270-
// CHECK: then2:
271+
// CHECK: then3:
271272
// CHECK-NEXT: [[TMP15:%.*]] = getelementptr inbounds [32 x i32], ptr addrspace(3) @__openmp_nvptx_data_transfer_temporary_storage, i64 0, i32 [[TMP3]]
272273
// CHECK-NEXT: [[TMP16:%.*]] = getelementptr inbounds [1 x ptr], ptr [[TMP6]], i64 0, i64 0
273274
// CHECK-NEXT: [[TMP17:%.*]] = load ptr, ptr [[TMP16]], align 8
274275
// CHECK-NEXT: [[TMP18:%.*]] = getelementptr i32, ptr [[TMP17]], i32 [[TMP7]]
275276
// CHECK-NEXT: [[TMP19:%.*]] = load volatile i32, ptr addrspace(3) [[TMP15]], align 4
276277
// CHECK-NEXT: store i32 [[TMP19]], ptr [[TMP18]], align 4
277278
// CHECK-NEXT: br label [[IFCONT4:%.*]]
278-
// CHECK: else3:
279+
// CHECK: else4:
279280
// CHECK-NEXT: br label [[IFCONT4]]
280-
// CHECK: ifcont4:
281+
// CHECK: ifcont5:
281282
// CHECK-NEXT: [[TMP20:%.*]] = add nsw i32 [[TMP7]], 1
282283
// CHECK-NEXT: store i32 [[TMP20]], ptr [[DOTCNT_ADDR]], align 4
283284
// CHECK-NEXT: br label [[PRECOND]]

0 commit comments

Comments
 (0)