113 lines · 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 %s3// expected-no-diagnostics4 5void foo();6 7void masked() {8 #pragma target9 #pragma omp parallel masked10 {11 foo();12 }13}14 15void maskedFilter() {16 const int tid = 1;17 #pragma target18 #pragma omp parallel masked filter(tid)19 {20 foo();21 }22}23 24void master() {25 #pragma target26 #pragma omp parallel master27 {28 foo();29 }30}31// CHECK-LABEL: define {{[^@]+}}@masked32// 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 void36//37//38// CHECK-LABEL: define {{[^@]+}}@masked.omp_outlined39// 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 842// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 843// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 844// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 845// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 846// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 447// CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_masked(ptr @[[GLOB1]], i32 [[TMP1]], i32 0)48// CHECK-NEXT: [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 049// 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 void56//57//58// CHECK-LABEL: define {{[^@]+}}@maskedFilter59// CHECK-SAME: () #[[ATTR0]] {60// CHECK-NEXT: entry:61// CHECK-NEXT: [[TID:%.*]] = alloca i32, align 462// CHECK-NEXT: store i32 1, ptr [[TID]], align 463// CHECK-NEXT: call void (ptr, i32, ptr, ...) @__kmpc_fork_call(ptr @[[GLOB1]], i32 0, ptr @maskedFilter.omp_outlined)64// CHECK-NEXT: ret void65//66//67// CHECK-LABEL: define {{[^@]+}}@maskedFilter.omp_outlined68// 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 871// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 872// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 873// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 874// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 875// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 476// CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_masked(ptr @[[GLOB1]], i32 [[TMP1]], i32 1)77// CHECK-NEXT: [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 078// 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 void85//86//87// CHECK-LABEL: define {{[^@]+}}@master88// 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 void92//93//94// CHECK-LABEL: define {{[^@]+}}@master.omp_outlined95// 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 898// CHECK-NEXT: [[DOTBOUND_TID__ADDR:%.*]] = alloca ptr, align 899// CHECK-NEXT: store ptr [[DOTGLOBAL_TID_]], ptr [[DOTGLOBAL_TID__ADDR]], align 8100// CHECK-NEXT: store ptr [[DOTBOUND_TID_]], ptr [[DOTBOUND_TID__ADDR]], align 8101// CHECK-NEXT: [[TMP0:%.*]] = load ptr, ptr [[DOTGLOBAL_TID__ADDR]], align 8102// CHECK-NEXT: [[TMP1:%.*]] = load i32, ptr [[TMP0]], align 4103// CHECK-NEXT: [[TMP2:%.*]] = call i32 @__kmpc_master(ptr @[[GLOB1]], i32 [[TMP1]])104// CHECK-NEXT: [[TMP3:%.*]] = icmp ne i32 [[TMP2]], 0105// 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 void112//113