clang/test/OpenMP/reduction_complex.c

   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