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_coalesced_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_coalesced_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_coalesced_push_stack |
36 | |
37 | // CHECK: define {{.*}}[[BAR]]() |
38 | // CHECK: alloca i32, |
39 | // CHECK: [[A_LOCAL_ADDR:%.+]] = alloca i32, |
40 | // CHECK: [[RES:%.+]] = call i8 @__kmpc_is_spmd_exec_mode() |
41 | // CHECK: [[IS_SPMD:%.+]] = icmp ne i8 [[RES]], 0 |
42 | // CHECK: br i1 [[IS_SPMD]], label |
43 | // CHECK: br label |
44 | // CHECK: [[RES:%.+]] = call i8* @__kmpc_data_sharing_coalesced_push_stack(i64 128, i16 0) |
45 | // CHECK: [[GLOBALS:%.+]] = bitcast i8* [[RES]] to [[GLOBAL_ST:%.+]]* |
46 | // CHECK: br label |
47 | // CHECK: [[ITEMS:%.+]] = phi [[GLOBAL_ST]]* [ null, {{.+}} ], [ [[GLOBALS]], {{.+}} ] |
48 | // CHECK: [[A_ADDR:%.+]] = getelementptr inbounds [[GLOBAL_ST]], [[GLOBAL_ST]]* [[ITEMS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 |
49 | // CHECK: [[TID:%.+]] = call i32 @llvm.nvvm.read.ptx.sreg.tid.x() |
50 | // CHECK: [[LID:%.+]] = and i32 [[TID]], 31 |
51 | // CHECK: [[A_GLOBAL_ADDR:%.+]] = getelementptr inbounds [32 x i32], [32 x i32]* [[A_ADDR]], i32 0, i32 [[LID]] |
52 | // CHECK: [[A_ADDR:%.+]] = select i1 [[IS_SPMD]], i32* [[A_LOCAL_ADDR]], i32* [[A_GLOBAL_ADDR]] |
53 | // CHECK: call {{.*}}[[FOO]](i32* dereferenceable{{.*}} [[A_ADDR]]) |
54 | // CHECK: br i1 [[IS_SPMD]], label |
55 | // CHECK: [[BC:%.+]] = bitcast [[GLOBAL_ST]]* [[ITEMS]] to i8* |
56 | // CHECK: call void @__kmpc_data_sharing_pop_stack(i8* [[BC]]) |
57 | // CHECK: br label |
58 | // CHECK: ret i32 |
59 | |