blob: 04882d346cc46247bd3e014eaf9141616c2313b9 [file] [log] [blame]
Alexey Bataevc99042b2018-03-15 18:10:54 +00001// 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
Alexey Bataev8c5555c2019-05-21 15:11:58 +00002// 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
Alexey Bataevc99042b2018-03-15 18:10:54 +00003// expected-no-diagnostics
4
5int foo(int &a) { return a; }
6
7int bar() {
8 int a;
9 return foo(a);
10}
11
Alexey Bataev9a700172018-05-08 14:16:57 +000012// CHECK: define weak void @__omp_offloading_{{.*}}maini1{{.*}}_l[[@LINE+5]](i32* dereferenceable{{.*}})
Alexey Bataev1fc1f8e2018-11-02 16:08:31 +000013// CHECK-NOT: @__kmpc_data_sharing_coalesced_push_stack
Alexey Bataevc99042b2018-03-15 18:10:54 +000014
15int 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{{.*}})
Alexey Bataev1fc1f8e2018-11-02 16:08:31 +000027// CHECK-NOT: call i8* @__kmpc_data_sharing_coalesced_push_stack(
Alexey Bataevb99dcb52018-07-09 17:43:58 +000028// CHECK: [[B_ADDR:%.+]] = alloca i32,
Alexey Bataevc99042b2018-03-15 18:10:54 +000029// CHECK: call {{.*}}[[FOO:@.*foo.*]](i32* dereferenceable{{.*}} [[B_ADDR]])
30// CHECK: call {{.*}}[[BAR:@.*bar.*]]()
Alexey Bataevb99dcb52018-07-09 17:43:58 +000031// CHECK-NOT: call void @__kmpc_data_sharing_pop_stack(
Alexey Bataevc99042b2018-03-15 18:10:54 +000032// CHECK: ret void
33
34// CHECK: define {{.*}}[[FOO]](i32* dereferenceable{{.*}})
Alexey Bataev1fc1f8e2018-11-02 16:08:31 +000035// CHECK-NOT: @__kmpc_data_sharing_coalesced_push_stack
Alexey Bataevc99042b2018-03-15 18:10:54 +000036
37// CHECK: define {{.*}}[[BAR]]()
Alexey Bataev9ea3c382018-10-09 14:49:00 +000038// CHECK: alloca i32,
39// CHECK: [[A_LOCAL_ADDR:%.+]] = alloca i32,
Alexey Bataev8d8e1232018-08-29 18:32:21 +000040// 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
Alexey Bataev1fc1f8e2018-11-02 16:08:31 +000044// CHECK: [[RES:%.+]] = call i8* @__kmpc_data_sharing_coalesced_push_stack(i64 128, i16 0)
Alexey Bataev9ea3c382018-10-09 14:49:00 +000045// CHECK: [[GLOBALS:%.+]] = bitcast i8* [[RES]] to [[GLOBAL_ST:%.+]]*
Alexey Bataev8d8e1232018-08-29 18:32:21 +000046// CHECK: br label
Alexey Bataev9ea3c382018-10-09 14:49:00 +000047// CHECK: [[ITEMS:%.+]] = phi [[GLOBAL_ST]]* [ null, {{.+}} ], [ [[GLOBALS]], {{.+}} ]
Alexey Bataev8d8e1232018-08-29 18:32:21 +000048// CHECK: [[A_ADDR:%.+]] = getelementptr inbounds [[GLOBAL_ST]], [[GLOBAL_ST]]* [[ITEMS]], i{{[0-9]+}} 0, i{{[0-9]+}} 0
Alexey Bataev9ea3c382018-10-09 14:49:00 +000049// 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]]
Alexey Bataevc99042b2018-03-15 18:10:54 +000053// CHECK: call {{.*}}[[FOO]](i32* dereferenceable{{.*}} [[A_ADDR]])
Alexey Bataev8d8e1232018-08-29 18:32:21 +000054// 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
Alexey Bataevc99042b2018-03-15 18:10:54 +000058// CHECK: ret i32