| Alexey Bataev | c99042b | 2018-03-15 18:10:54 +0000 | [diff] [blame] | 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 | |
| Alexey Bataev | 9a70017 | 2018-05-08 14:16:57 +0000 | [diff] [blame^] | 12 | // CHECK: define weak void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+5]](i32* dereferenceable{{.*}}) |
| Alexey Bataev | c99042b | 2018-03-15 18:10:54 +0000 | [diff] [blame] | 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: [[RES:%.+]] = call i8* @__kmpc_data_sharing_push_stack(i64 4, i16 0) |
| 28 | // CHECK: [[GLOBALS:%.+]] = bitcast i8* [[RES]] to [[GLOBAL_ST:%struct[.].*]]* |
| 29 | // CHECK: [[B_ADDR:%.+]] = getelementptr inbounds [[GLOBAL_ST]], [[GLOBAL_ST]]* [[GLOBALS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 |
| 30 | // CHECK: call {{.*}}[[FOO:@.*foo.*]](i32* dereferenceable{{.*}} [[B_ADDR]]) |
| 31 | // CHECK: call {{.*}}[[BAR:@.*bar.*]]() |
| 32 | // CHECK: call void @__kmpc_data_sharing_pop_stack(i8* [[RES]]) |
| 33 | // CHECK: ret void |
| 34 | |
| 35 | // CHECK: define {{.*}}[[FOO]](i32* dereferenceable{{.*}}) |
| 36 | // CHECK-NOT: @__kmpc_data_sharing_push_stack |
| 37 | |
| 38 | // CHECK: define {{.*}}[[BAR]]() |
| 39 | // CHECK: [[RES:%.+]] = call i8* @__kmpc_data_sharing_push_stack(i64 4, i16 0) |
| 40 | // CHECK: [[GLOBALS:%.+]] = bitcast i8* [[RES]] to [[GLOBAL_ST:%struct[.].*]]* |
| 41 | // CHECK: [[A_ADDR:%.+]] = getelementptr inbounds [[GLOBAL_ST]], [[GLOBAL_ST]]* [[GLOBALS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 |
| 42 | // CHECK: call {{.*}}[[FOO]](i32* dereferenceable{{.*}} [[A_ADDR]]) |
| 43 | // CHECK: call void @__kmpc_data_sharing_pop_stack(i8* [[RES]]) |
| 44 | // CHECK: ret i32 |