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
masked()7 void masked() {
8 #pragma target
9 #pragma omp parallel masked
10 {
11 foo();
12 }
13 }
14
maskedFilter()15 void maskedFilter() {
16 const int tid = 1;
17 #pragma target
18 #pragma omp parallel masked filter(tid)
19 {
20 foo();
21 }
22 }
23
master()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 //
113