1 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple powerpc64le-unknown-unknown -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm-bc %s -o %t-ppc-host.bc
2 // RUN: %clang_cc1 -verify -fopenmp -x c++ -triple nvptx64-nvidia-cuda -fopenmp-targets=nvptx64-nvidia-cuda -emit-llvm %s -fopenmp-is-device -fopenmp-host-ir-file-path %t-ppc-host.bc -o - | FileCheck %s
3 // expected-no-diagnostics
4 
5 int foo(int &a) { return a; }
6 
7 int bar() {
8   int a;
9   return foo(a);
10 }
11 
12 // CHECK: define weak void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+5]](i32* dereferenceable{{.*}})
13 // CHECK-NOT: @__kmpc_data_sharing_push_stack
14 
15 int maini1() {
16   int a;
17 #pragma omp target parallel map(from:a)
18   {
19     int b;
20     a = foo(b) + bar();
21   }
22   return a;
23 }
24 
25 // parallel region
26 // CHECK: define {{.*}}void @{{.*}}(i32* noalias {{.*}}, i32* noalias {{.*}}, i32* dereferenceable{{.*}})
27 // CHECK-NOT: call i8* @__kmpc_data_sharing_push_stack(
28 // CHECK: [[B_ADDR:%.+]] = alloca i32,
29 // CHECK: call {{.*}}[[FOO:@.*foo.*]](i32* dereferenceable{{.*}} [[B_ADDR]])
30 // CHECK: call {{.*}}[[BAR:@.*bar.*]]()
31 // CHECK-NOT: call void @__kmpc_data_sharing_pop_stack(
32 // CHECK: ret void
33 
34 // CHECK: define {{.*}}[[FOO]](i32* dereferenceable{{.*}})
35 // CHECK-NOT: @__kmpc_data_sharing_push_stack
36 
37 // CHECK: define {{.*}}[[BAR]]()
38 // CHECK: [[STACK:%.+]] = alloca [[GLOBAL_ST:%.+]],
39 // CHECK: [[RES:%.+]] = call i8 @__kmpc_is_spmd_exec_mode()
40 // CHECK: [[IS_SPMD:%.+]] = icmp ne i8 [[RES]], 0
41 // CHECK: br i1 [[IS_SPMD]], label
42 // CHECK: br label
43 // CHECK: [[RES:%.+]] = call i8* @__kmpc_data_sharing_push_stack(i64 4, i16 0)
44 // CHECK: [[GLOBALS:%.+]] = bitcast i8* [[RES]] to [[GLOBAL_ST]]*
45 // CHECK: br label
46 // CHECK: [[ITEMS:%.+]] = phi [[GLOBAL_ST]]* [ [[STACK]], {{.+}} ], [ [[GLOBALS]], {{.+}} ]
47 // CHECK: [[A_ADDR:%.+]] = getelementptr inbounds [[GLOBAL_ST]], [[GLOBAL_ST]]* [[ITEMS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
48 // CHECK: call {{.*}}[[FOO]](i32* dereferenceable{{.*}} [[A_ADDR]])
49 // CHECK: br i1 [[IS_SPMD]], label
50 // CHECK: [[BC:%.+]] = bitcast [[GLOBAL_ST]]* [[ITEMS]] to i8*
51 // CHECK: call void @__kmpc_data_sharing_pop_stack(i8* [[BC]])
52 // CHECK: br label
53 // CHECK: ret i32
54