// 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 _ // RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ \ // RUN: -triple powerpc64le-unknown-unknown \ // RUN: -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o \ // RUN: %t-ppc-host.bc // RUN: %clang_cc1 -verify -fopenmp -fopenmp-cuda-mode -x c++ \ // RUN: -triple nvptx64-unknown-unknown -DCUA \ // RUN: -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s \ // RUN: -fopenmp-is-target-device -fopenmp-host-ir-file-path %t-ppc-host.bc \ // RUN: -o - | FileCheck %s --check-prefix CHECK // expected-no-diagnostics int foo() { int i; int j; _Complex float sum = 0; #pragma omp target teams loop reduction(+:sum) collapse(2) bind(parallel) order(concurrent) lastprivate(j) map(tofrom:sum) for(i=0; i<10; i++) for(j=0; j<10; j++) sum += i; return 0; } // CHECK-LABEL: define {{[^@]+}}@_omp_reduction_shuffle_and_reduce_func // CHECK-SAME: (ptr noundef [[TMP0:%.*]], i16 noundef signext [[TMP1:%.*]], i16 noundef signext [[TMP2:%.*]], i16 noundef signext [[TMP3:%.*]]) #[[ATTR2:[0-9]+]] { // CHECK-NEXT: entry: // CHECK-NEXT: %[[VAL_228:.*]] = alloca ptr, align 8 // CHECK-NEXT: %[[VAL_229:.*]] = alloca i16, align 2 // CHECK-NEXT: %[[VAL_230:.*]] = alloca i16, align 2 // CHECK-NEXT: %[[VAL_231:.*]] = alloca i16, align 2 // CHECK-NEXT: %[[VAL_232:.*]] = alloca [1 x ptr], align 8 // CHECK-NEXT: %[[VAL_233:.*]] = alloca { float, float }, align 8 // CHECK-NEXT: store ptr %[[VAL_234:.*]], ptr %[[VAL_228]], align 8 // CHECK-NEXT: store i16 %[[VAL_235:.*]], ptr %[[VAL_229]], align 2 // CHECK-NEXT: store i16 %[[VAL_236:.*]], ptr %[[VAL_230]], align 2 // CHECK-NEXT: store i16 %[[VAL_237:.*]], ptr %[[VAL_231]], align 2 // CHECK-NEXT: %[[VAL_238:.*]] = load ptr, ptr %[[VAL_228]], align 8 // CHECK-NEXT: %[[VAL_239:.*]] = load i16, ptr %[[VAL_229]], align 2 // CHECK-NEXT: %[[VAL_240:.*]] = load i16, ptr %[[VAL_230]], align 2 // CHECK-NEXT: %[[VAL_241:.*]] = load i16, ptr %[[VAL_231]], align 2 // CHECK-NEXT: %[[VAL_242:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_238]], i64 0, i64 0 // CHECK-NEXT: %[[VAL_243:.*]] = load ptr, ptr %[[VAL_242]], align 8 // CHECK-NEXT: %[[VAL_244:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_232]], i64 0, i64 0 // CHECK-NEXT: %[[VAL_245:.*]] = getelementptr { float, float }, ptr %[[VAL_243]], i64 1 // CHECK-NEXT: %[[VAL_246:.*]] = load i64, ptr %[[VAL_243]], align 8 // CHECK-NEXT: %[[VAL_247:.*]] = call i32 @__kmpc_get_warp_size() // CHECK-NEXT: %[[VAL_248:.*]] = trunc i32 %[[VAL_247]] to i16 // CHECK-NEXT: %[[VAL_249:.*]] = call i64 @__kmpc_shuffle_int64(i64 %[[VAL_246]], i16 %[[VAL_240]], i16 %[[VAL_248]]) // CHECK-NEXT: store i64 %[[VAL_249]], ptr %[[VAL_233]], align 8 // CHECK-NEXT: %[[VAL_250:.*]] = getelementptr i64, ptr %[[VAL_243]], i64 1 // CHECK-NEXT: %[[VAL_251:.*]] = getelementptr i64, ptr %[[VAL_233]], i64 1 // CHECK-NEXT: store ptr %[[VAL_233]], ptr %[[VAL_244]], align 8 // CHECK-NEXT: %[[VAL_252:.*]] = icmp eq i16 %[[VAL_241]], 0 // CHECK-NEXT: %[[VAL_253:.*]] = icmp eq i16 %[[VAL_241]], 1 // CHECK-NEXT: %[[VAL_254:.*]] = icmp ult i16 %[[VAL_239]], %[[VAL_240]] // CHECK-NEXT: %[[VAL_255:.*]] = and i1 %[[VAL_253]], %[[VAL_254]] // CHECK-NEXT: %[[VAL_256:.*]] = icmp eq i16 %[[VAL_241]], 2 // CHECK-NEXT: %[[VAL_257:.*]] = and i16 %[[VAL_239]], 1 // CHECK-NEXT: %[[VAL_258:.*]] = icmp eq i16 %[[VAL_257]], 0 // CHECK-NEXT: %[[VAL_259:.*]] = and i1 %[[VAL_256]], %[[VAL_258]] // CHECK-NEXT: %[[VAL_260:.*]] = icmp sgt i16 %[[VAL_240]], 0 // CHECK-NEXT: %[[VAL_261:.*]] = and i1 %[[VAL_259]], %[[VAL_260]] // CHECK-NEXT: %[[VAL_262:.*]] = or i1 %[[VAL_252]], %[[VAL_255]] // CHECK-NEXT: %[[VAL_263:.*]] = or i1 %[[VAL_262]], %[[VAL_261]] // CHECK-NEXT: br i1 %[[VAL_263]], label %[[VAL_264:.*]], label %[[VAL_265:.*]] // CHECK: then: ; preds = %[[VAL_266:.*]] // 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 // CHECK-NEXT: br label %[[VAL_267:.*]] // CHECK: else: ; preds = %[[VAL_266]] // CHECK-NEXT: br label %[[VAL_267]] // CHECK: ifcont: ; preds = %[[VAL_265]], %[[VAL_264]] // CHECK-NEXT: %[[VAL_268:.*]] = icmp eq i16 %[[VAL_241]], 1 // CHECK-NEXT: %[[VAL_269:.*]] = icmp uge i16 %[[VAL_239]], %[[VAL_240]] // CHECK-NEXT: %[[VAL_270:.*]] = and i1 %[[VAL_268]], %[[VAL_269]] // CHECK-NEXT: br i1 %[[VAL_270]], label %[[VAL_271:.*]], label %[[VAL_272:.*]] // CHECK: then4: ; preds = %[[VAL_267]] // CHECK-NEXT: %[[VAL_273:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_232]], i64 0, i64 0 // CHECK-NEXT: %[[VAL_274:.*]] = load ptr, ptr %[[VAL_273]], align 8 // CHECK-NEXT: %[[VAL_275:.*]] = getelementptr inbounds [1 x ptr], ptr %[[VAL_238]], i64 0, i64 0 // CHECK-NEXT: %[[VAL_276:.*]] = load ptr, ptr %[[VAL_275]], align 8 // CHECK-NEXT: %[[VAL_277:.*]] = getelementptr inbounds { float, float }, ptr %[[VAL_274]], i32 0, i32 0 // CHECK-NEXT: %[[VAL_278:.*]] = load float, ptr %[[VAL_277]], align 4 // CHECK-NEXT: %[[VAL_279:.*]] = getelementptr inbounds { float, float }, ptr %[[VAL_274]], i32 0, i32 1 // CHECK-NEXT: %[[VAL_280:.*]] = load float, ptr %[[VAL_279]], align 4 // CHECK-NEXT: %[[VAL_281:.*]] = getelementptr inbounds { float, float }, ptr %[[VAL_276]], i32 0, i32 0 // CHECK-NEXT: %[[VAL_282:.*]] = getelementptr inbounds { float, float }, ptr %[[VAL_276]], i32 0, i32 1 // CHECK-NEXT: store float %[[VAL_278]], ptr %[[VAL_281]], align 4 // CHECK-NEXT: store float %[[VAL_280]], ptr %[[VAL_282]], align 4 // CHECK-NEXT: br label %[[VAL_283:.*]] // CHECK: else7: ; preds = %[[VAL_267]] // CHECK-NEXT: br label %[[VAL_283]] // CHECK: ifcont8: ; preds = %[[VAL_272]], %[[VAL_271]] // CHECK-NEXT: ret void