1*c99042baSAlexey Bataev // 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*c99042baSAlexey Bataev // 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*c99042baSAlexey Bataev // expected-no-diagnostics 4*c99042baSAlexey Bataev 5*c99042baSAlexey Bataev int foo(int &a) { return a; } 6*c99042baSAlexey Bataev 7*c99042baSAlexey Bataev int bar() { 8*c99042baSAlexey Bataev int a; 9*c99042baSAlexey Bataev return foo(a); 10*c99042baSAlexey Bataev } 11*c99042baSAlexey Bataev 12*c99042baSAlexey Bataev // CHECK: define void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+5]](i32* dereferenceable{{.*}}) 13*c99042baSAlexey Bataev // CHECK-NOT: @__kmpc_data_sharing_push_stack 14*c99042baSAlexey Bataev 15*c99042baSAlexey Bataev int maini1() { 16*c99042baSAlexey Bataev int a; 17*c99042baSAlexey Bataev #pragma omp target parallel map(from:a) 18*c99042baSAlexey Bataev { 19*c99042baSAlexey Bataev int b; 20*c99042baSAlexey Bataev a = foo(b) + bar(); 21*c99042baSAlexey Bataev } 22*c99042baSAlexey Bataev return a; 23*c99042baSAlexey Bataev } 24*c99042baSAlexey Bataev 25*c99042baSAlexey Bataev // parallel region 26*c99042baSAlexey Bataev // CHECK: define {{.*}}void @{{.*}}(i32* noalias {{.*}}, i32* noalias {{.*}}, i32* dereferenceable{{.*}}) 27*c99042baSAlexey Bataev // CHECK: [[RES:%.+]] = call i8* @__kmpc_data_sharing_push_stack(i64 4, i16 0) 28*c99042baSAlexey Bataev // CHECK: [[GLOBALS:%.+]] = bitcast i8* [[RES]] to [[GLOBAL_ST:%struct[.].*]]* 29*c99042baSAlexey Bataev // CHECK: [[B_ADDR:%.+]] = getelementptr inbounds [[GLOBAL_ST]], [[GLOBAL_ST]]* [[GLOBALS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 30*c99042baSAlexey Bataev // CHECK: call {{.*}}[[FOO:@.*foo.*]](i32* dereferenceable{{.*}} [[B_ADDR]]) 31*c99042baSAlexey Bataev // CHECK: call {{.*}}[[BAR:@.*bar.*]]() 32*c99042baSAlexey Bataev // CHECK: call void @__kmpc_data_sharing_pop_stack(i8* [[RES]]) 33*c99042baSAlexey Bataev // CHECK: ret void 34*c99042baSAlexey Bataev 35*c99042baSAlexey Bataev // CHECK: define {{.*}}[[FOO]](i32* dereferenceable{{.*}}) 36*c99042baSAlexey Bataev // CHECK-NOT: @__kmpc_data_sharing_push_stack 37*c99042baSAlexey Bataev 38*c99042baSAlexey Bataev // CHECK: define {{.*}}[[BAR]]() 39*c99042baSAlexey Bataev // CHECK: [[RES:%.+]] = call i8* @__kmpc_data_sharing_push_stack(i64 4, i16 0) 40*c99042baSAlexey Bataev // CHECK: [[GLOBALS:%.+]] = bitcast i8* [[RES]] to [[GLOBAL_ST:%struct[.].*]]* 41*c99042baSAlexey Bataev // CHECK: [[A_ADDR:%.+]] = getelementptr inbounds [[GLOBAL_ST]], [[GLOBAL_ST]]* [[GLOBALS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0 42*c99042baSAlexey Bataev // CHECK: call {{.*}}[[FOO]](i32* dereferenceable{{.*}} [[A_ADDR]]) 43*c99042baSAlexey Bataev // CHECK: call void @__kmpc_data_sharing_pop_stack(i8* [[RES]]) 44*c99042baSAlexey Bataev // CHECK: ret i32 45