clang/test/OpenMP/nvptx_lambda_pointer_capturing.cpp

   1 // Test host codegen only.
   2 // RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
   3 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -emit-pch -o %t %s
   4 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -fopenmp-targets=powerpc64le-ibm-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-64
   5 // RUN: %clang_cc1 -verify -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-llvm %s -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
   6 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -emit-pch -o %t %s
   7 // RUN: %clang_cc1 -fopenmp -x c++ -std=c++11 -triple i386-unknown-unknown -fopenmp-targets=i386-pc-linux-gnu -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CHECK --check-prefix CHECK-32
   8
   9 // expected-no-diagnostics
  10 #ifndef HEADER
  11 #define HEADER
  12
  13 // CHECK: [[ANON_T:%.+]] = type { ptr, ptr }
  14 // CHECK-DAG: [[SIZES_TEMPLATE:@.+]] = private {{.+}} constant [5 x i[[PTRSZ:32|64]]] [i{{32|64}} 4, i{{32|64}} 4, i{{32|64}} {{8|16}}, i{{32|64}} 0, i{{32|64}} 0]
  15 // CHECK-DAG: [[TYPES_TEMPLATE:@.+]] = private {{.+}} constant [5 x i64] [i64 800, i64 800, i64 673, i64 844424930132752, i64 844424930132752]
  16 // CHECK-DAG: [[SIZES:@.+]] = private {{.+}} constant [3 x i[[PTRSZ:32|64]]] [i{{32|64}} {{8|16}}, i{{32|64}} 0, i{{32|64}} 0]
  17 // CHECK-DAG: [[TYPES:@.+]] = private {{.+}} constant [3 x i64] [i64 673, i64 281474976711440, i64 281474976711440]
  18 // CHECK-DAG: [[TYPES3:@.+]] = private {{.+}} constant [3 x i64] [i64 545, i64 281474976711440, i64 800]
  19 // CHECK-DAG: [[TYPES11:@.+]] = private {{.+}} constant [5 x i64] [i64 800, i64 800, i64 549, i64 844424930132752, i64 844424930132752]
  20 // CHECK-DAG: [[TYPES13:@.+]] = private {{.+}} constant [2 x i64] [i64 545, i64 281474976711440]
  21 // CHECK-DAG: [[TYPES15:@.+]] = private {{.+}} constant [2 x i64] [i64 673, i64 281474976711440]
  22
  23 template <typename F>
  24 void omp_loop(int start, int end, F body) {
  25 #pragma omp target teams distribute parallel for
  26   for (int i = start; i < end; ++i) {
  27     body(i);
  28   }
  29 }
  30
  31 template <typename F>
  32 void omp_loop_ref(int start, int end, F& body) {
  33 #pragma omp target teams distribute parallel for map(always, to: body)
  34   for (int i = start; i < end; ++i) {
  35     body(i);
  36   }
  37   int *p;
  38   const auto &body_ref = [=](int i) {p[i]=0;};
  39   #pragma omp target map(to: body_ref)
  40   body_ref(10);
  41   #pragma omp target
  42   body_ref(10);
  43 }
  44
  45 template <class FTy>
  46 struct C {
  47   static void xoo(const FTy& f) {
  48     int x = 10;
  49     #pragma omp target map(to:f)
  50       f(x);
  51   }
  52 };
  53
  54 template <class FTy>
  55 void zoo(const FTy &functor) {
  56   C<FTy>::xoo(functor);
  57 }
  58
  59 // CHECK: define {{.*}}[[MAIN:@.+]](
  60 int main()
  61 {
  62   int* p = new int[100];
  63   int* q = new int[100];
  64   auto body = [=](int i){
  65     p[i] = q[i];
  66   };
  67   zoo([=](int i){p[i] = 0;});
  68
  69 #pragma omp target teams distribute parallel for
  70   for (int i = 0; i < 100; ++i) {
  71     body(i);
  72   }
  73
  74 // CHECK: [[BASE_PTRS:%.+]] = alloca [3 x ptr]{{.+}}
  75 // CHECK: [[PTRS:%.+]] = alloca [3 x ptr]{{.+}}
  76
  77 // First gep of pointers inside lambdas to store the values across function call need to be ignored
  78 // CHECK: {{%.+}} = getelementptr inbounds nuw [[ANON_T]], ptr %{{.+}}, i{{.+}} 0, i{{.+}} 0
  79 // CHECK: {{%.+}} = getelementptr inbounds nuw [[ANON_T]], ptr %{{.+}}, i{{.+}} 0, i{{.+}} 1
  80
  81 // access of pointers inside lambdas
  82 // CHECK: [[BASE_PTR1:%.+]] = getelementptr inbounds nuw [[ANON_T]], ptr %{{.+}}, i{{.+}} 0, i{{.+}} 0
  83 // CHECK: [[PTR1:%.+]] = load ptr, ptr [[BASE_PTR1]]
  84 // CHECK: [[BASE_PTR2:%.+]] = getelementptr inbounds nuw [[ANON_T]], ptr %{{.+}}, i{{.+}} 0, i{{.+}} 1
  85 // CHECK: [[PTR2:%.+]] = load ptr, ptr [[BASE_PTR2]]
  86
  87 // storage of pointers in baseptrs and ptrs arrays
  88 // CHECK: [[LOC_LAMBDA:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BASE_PTRS]], i{{.+}} 0, i{{.+}} 0
  89 // CHECK: store ptr %{{.+}}, ptr [[LOC_LAMBDA]]{{.+}}
  90 // CHECK: [[LOC_LAMBDA:%.+]] = getelementptr inbounds [3 x ptr], ptr [[PTRS]], i{{.+}} 0, i{{.+}} 0
  91 // CHECK: store ptr %{{.+}}, ptr [[LOC_LAMBDA]]{{.+}}
  92
  93 // CHECK: [[LOC_PTR1:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BASE_PTRS]], i{{.+}} 0, i{{.+}} 1
  94 // CHECK: store ptr [[BASE_PTR1]], ptr [[LOC_PTR1]]{{.+}}
  95 // CHECK: [[LOC_PTR1:%.+]] = getelementptr inbounds [3 x ptr], ptr [[PTRS]], i{{.+}} 0, i{{.+}} 1
  96 // CHECK: store ptr [[PTR1]], ptr [[LOC_PTR1]]{{.+}}
  97
  98
  99 // CHECK: [[LOC_PTR2:%.+]] = getelementptr inbounds [3 x ptr], ptr [[BASE_PTRS]], i{{.+}} 0, i{{.+}} 2
 100 // CHECK: store ptr [[BASE_PTR2]], ptr [[LOC_PTR2]]{{.+}}
 101 // CHECK: [[LOC_PTR2:%.+]] = getelementptr inbounds [3 x ptr], ptr [[PTRS]], i{{.+}} 0, i{{.+}} 2
 102 // CHECK: store ptr [[PTR2]], ptr [[LOC_PTR2]]{{.+}}
 103
 104   // actual target invocation
 105   // CHECK: [[BASES_GEP:%.+]] = getelementptr {{.+}} [3 x ptr], ptr [[BASE_PTRS]], {{.+}} 0, {{.+}} 0
 106   // CHECK: [[PTRS_GEP:%.+]] = getelementptr {{.+}} [3 x ptr], ptr [[PTRS]], {{.+}} 0, {{.+}} 0
 107   // CHECK: {{%.+}} = call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 0, i32 0, ptr @.{{.+}}.region_id, ptr %{{.+}})
 108
 109   omp_loop(0,100,body);
 110   omp_loop_ref(0,100,body);
 111 }
 112
 113 // CHECK: [[BASE_PTRS:%.+]] = alloca [5 x ptr]{{.+}}
 114 // CHECK: [[PTRS:%.+]] = alloca [5 x ptr]{{.+}}
 115
 116 // access of pointers inside lambdas
 117 // CHECK: [[BASE_PTR1:%.+]] = getelementptr inbounds nuw [[ANON_T]], ptr %{{.+}}, i{{.+}} 0, i{{.+}} 0
 118 // CHECK: [[PTR1:%.+]] = load ptr, ptr [[BASE_PTR1]]
 119 // CHECK: [[BASE_PTR2:%.+]] = getelementptr inbounds nuw [[ANON_T]], ptr %{{.+}}, i{{.+}} 0, i{{.+}} 1
 120 // CHECK: [[PTR2:%.+]] = load ptr, ptr [[BASE_PTR2]]
 121
 122 // storage of pointers in baseptrs and ptrs arrays
 123 // CHECK: [[LOC_LAMBDA:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BASE_PTRS]], i{{.+}} 0, i{{.+}} 2
 124 // CHECK: store ptr %{{.+}}, ptr [[LOC_LAMBDA]]{{.+}}
 125 // CHECK: [[LOC_LAMBDA:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i{{.+}} 0, i{{.+}} 2
 126 // CHECK: store ptr %{{.+}}, ptr [[LOC_LAMBDA]]{{.+}}
 127
 128 // CHECK: [[LOC_PTR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BASE_PTRS]], i{{.+}} 0, i{{.+}} 3
 129 // CHECK: store ptr [[BASE_PTR1]], ptr [[LOC_PTR1]]{{.+}}
 130 // CHECK: [[LOC_PTR1:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i{{.+}} 0, i{{.+}} 3
 131 // CHECK: store ptr [[PTR1]], ptr [[LOC_PTR1]]{{.+}}
 132
 133
 134 // CHECK: [[LOC_PTR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[BASE_PTRS]], i{{.+}} 0, i{{.+}} 4
 135 // CHECK: store ptr [[BASE_PTR2]], ptr [[LOC_PTR2]]{{.+}}
 136 // CHECK: [[LOC_PTR2:%.+]] = getelementptr inbounds [5 x ptr], ptr [[PTRS]], i{{.+}} 0, i{{.+}} 4
 137 // CHECK: store ptr [[PTR2]], ptr [[LOC_PTR2]]{{.+}}
 138
 139 // actual target invocation
 140 // CHECK: [[BASES_GEP:%.+]] = getelementptr {{.+}} [5 x ptr], ptr [[BASE_PTRS]], {{.+}} 0, {{.+}} 0
 141 // CHECK: [[PTRS_GEP:%.+]] = getelementptr {{.+}} [5 x ptr], ptr [[PTRS]], {{.+}} 0, {{.+}} 0
 142
 143 // CHECK: define internal void @{{.+}}omp_loop_ref{{.+}}(
 144 // CHECK: [[BODY:%body.addr]] = alloca ptr
 145 // CHECK: [[TMP:%tmp]] = alloca ptr
 146 // CHECK: [[BODY_REF:%body_ref]] = alloca ptr
 147 // CHECK: [[REF_TMP:%ref.tmp]] = alloca %class.anon.1
 148 // CHECK: [[TMP8:%tmp.+]] = alloca ptr
 149 // CHECK: [[L0:%.+]] = load ptr, ptr [[BODY]]
 150 // CHECK: store ptr [[L0]], ptr [[TMP]]
 151 // CHECK: [[L5:%.+]] = load ptr, ptr [[TMP]]
 152 // CHECK-NOT [[L6:%.+]] = load ptr, ptr [[TMP]]
 153 // CHECK-NOT [[L7:%.+]] = load ptr, ptr [[TMP]]
 154 // CHECK: store ptr [[REF_TMP]], ptr [[BODY_REF]]
 155 // CHECK:[[L47:%.+]] =  load ptr, ptr [[BODY_REF]]
 156 // CHECK: store ptr [[L47]], ptr [[TMP8]]
 157 // CHECK: [[L48:%.+]] = load ptr, ptr [[TMP8]]
 158 // CHECK-NOT: [[L49:%.+]] = load ptr, ptr [[TMP8]]
 159 // CHECK-NOT: [[L50:%.+]] = load ptr, ptr [[TMP8]]
 160 // CHECK: ret void
 161
 162 // CHECK: define internal void @{{.+}}xoo{{.+}}(
 163 // CHECK: [[FADDR:%f.addr]] = alloca ptr
 164 // CHECK: [[L0:%.+]] = load ptr, ptr [[FADDR]]
 165 // CHECK: store ptr [[L0]], ptr [[TMP:%tmp]]
 166 // CHECK: [[L1:%.+]] = load ptr, ptr [[TMP]]
 167 // CHECK-NOT: %4 = load ptr, ptr [[TMP]]
 168 // CHECK-NOT: %5 = load ptr, ptr [[TMP]]
 169 // CHECK: [[L4:%.+]] = getelementptr inbounds nuw %class.anon.0, ptr [[L1]], i32 0, i32 0
 170 // CHECK: [[L5:%.+]] = load ptr, ptr [[L4]]
 171 // CHECK: ret void
 172
 173 #endif