clang/test/OpenMP/parallel_masked_target.cpp

   1 // NOTE: Assertions have been autogenerated by utils/update_cc_test_checks.py UTC_ARGS: --function-signature --include-generated-funcs --prefix-filecheck-ir-name _
   2 // RUN: %clang_cc1 -triple x86_64-unknown-unknown -fopenmp -fopenmp-version=52 -fopenmp-targets=nvptx64 -offload-device-only -x c -emit-llvm %s -o - | FileCheck %s
   3 // expected-no-diagnostics
   4
   5 void foo();
   6
   7 void masked() {
   8     #pragma target
   9     #pragma omp parallel masked
  10     {
  11         foo();
  12     }
  13 }
  14
  15 void maskedFilter() {
  16     const int tid = 1;
  17     #pragma target
  18     #pragma omp parallel masked filter(tid)
  19     {
  20         foo();
  21     }
  22 }
  23
  24 void master() {
  25     #pragma target
  26     #pragma omp parallel master
  27     {
  28         foo();
  29     }
  30 }
  31 // CHECK-LABEL: define {{[^@]+}}@masked
  32 // CHECK-SAME: () #[[ATTR0:[0-9]+]] {
  33 // CHECK-NEXT:  entry:
  34 // CHECK-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1:[0-9]+]], i32 0, ptr @masked.omp_outlined)
  35 // CHECK-NEXT:    ret void
  36 //
  37 //
  38 // CHECK-LABEL: define {{[^@]+}}@masked.omp_outlined
  39 // CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1:[0-9]+]] {
  40 // CHECK-NEXT:  entry:
  41 // CHECK-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
  42 // CHECK-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
  43 // CHECK-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
  44 // CHECK-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
  45 // CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
  46 // CHECK-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
  47 // CHECK-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_masked(ptr @[[GLOB1]], i32 [[TMP1]], i32 0)
  48 // CHECK-NEXT:    [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0
  49 // CHECK-NEXT:    br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]]
  50 // CHECK:       omp_if.then:
  51 // CHECK-NEXT:    call void (...) @foo()
  52 // CHECK-NEXT:    call void @__kmpc_end_masked(ptr @[[GLOB1]], i32 [[TMP1]])
  53 // CHECK-NEXT:    br label [[OMP_IF_END]]
  54 // CHECK:       omp_if.end:
  55 // CHECK-NEXT:    ret void
  56 //
  57 //
  58 // CHECK-LABEL: define {{[^@]+}}@maskedFilter
  59 // CHECK-SAME: () #[[ATTR0]] {
  60 // CHECK-NEXT:  entry:
  61 // CHECK-NEXT:    [[TID:%.*]] = alloca i32, align 4
  62 // CHECK-NEXT:    store i32 1, ptr [[TID]], align 4
  63 // CHECK-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @maskedFilter.omp_outlined)
  64 // CHECK-NEXT:    ret void
  65 //
  66 //
  67 // CHECK-LABEL: define {{[^@]+}}@maskedFilter.omp_outlined
  68 // CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {
  69 // CHECK-NEXT:  entry:
  70 // CHECK-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
  71 // CHECK-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
  72 // CHECK-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
  73 // CHECK-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
  74 // CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
  75 // CHECK-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
  76 // CHECK-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_masked(ptr @[[GLOB1]], i32 [[TMP1]], i32 1)
  77 // CHECK-NEXT:    [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0
  78 // CHECK-NEXT:    br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]]
  79 // CHECK:       omp_if.then:
  80 // CHECK-NEXT:    call void (...) @foo()
  81 // CHECK-NEXT:    call void @__kmpc_end_masked(ptr @[[GLOB1]], i32 [[TMP1]])
  82 // CHECK-NEXT:    br label [[OMP_IF_END]]
  83 // CHECK:       omp_if.end:
  84 // CHECK-NEXT:    ret void
  85 //
  86 //
  87 // CHECK-LABEL: define {{[^@]+}}@master
  88 // CHECK-SAME: () #[[ATTR0]] {
  89 // CHECK-NEXT:  entry:
  90 // CHECK-NEXT:    call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @master.omp_outlined)
  91 // CHECK-NEXT:    ret void
  92 //
  93 //
  94 // CHECK-LABEL: define {{[^@]+}}@master.omp_outlined
  95 // CHECK-SAME: (ptr noalias noundef [[DOTGLOBAL_TID_:%.*]], ptr noalias noundef [[DOTBOUND_TID_:%.*]]) #[[ATTR1]] {
  96 // CHECK-NEXT:  entry:
  97 // CHECK-NEXT:    [[DOTGLOBAL_TID__ADDR:%.*]] = alloca ptr, align 8
  98 // CHECK-NEXT:    [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 8
  99 // CHECK-NEXT:    store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8
 100 // CHECK-NEXT:    store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8
 101 // CHECK-NEXT:    [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8
 102 // CHECK-NEXT:    [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4
 103 // CHECK-NEXT:    [[TMP2:%.*]] = call i32 @__kmpc_master(ptr @[[GLOB1]], i32 [[TMP1]])
 104 // CHECK-NEXT:    [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0
 105 // CHECK-NEXT:    br i1 [[TMP3]], label [[OMP_IF_THEN:%.*]], label [[OMP_IF_END:%.*]]
 106 // CHECK:       omp_if.then:
 107 // CHECK-NEXT:    call void (...) @foo()
 108 // CHECK-NEXT:    call void @__kmpc_end_master(ptr @[[GLOB1]], i32 [[TMP1]])
 109 // CHECK-NEXT:    br label [[OMP_IF_END]]
 110 // CHECK:       omp_if.end:
 111 // CHECK-NEXT:    ret void
 112 //