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 - -disable-llvm-optzns | 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* nonnull align {{[0-9]+}} dereferenceable({{[0-9]+}}) %{{.*}}) 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* nonnull align {{[0-9]+}} dereferenceable({{[0-9]+}}) %{{.*}}) 27 // CHECK-NOT: call i8* @__kmpc_data_sharing_coalesced_push_stack( 28 // CHECK: [[B_ADDR:%.+]] = alloca i32, 29 // CHECK: call {{.*}}[[FOO:@.*foo.*]](i32* nonnull align {{[0-9]+}} dereferenceable({{[0-9]+}}) [[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* nonnull align {{[0-9]+}} 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* nonnull align {{[0-9]+}} 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