191 lines · cpp
1// expected-no-diagnostics2#ifndef HEADER3#define HEADER4 5///==========================================================================///6// RUN: %clang_cc1 -DUSE -DCK31A -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31A,CK31A-64,CK31A-USE7// RUN: %clang_cc1 -DUSE -DCK31A -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s8// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31A,CK31A-64,CK31A-USE9// RUN: %clang_cc1 -DUSE -DCK31A -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31A,CK31A-32,CK31A-USE10// RUN: %clang_cc1 -DUSE -DCK31A -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s11// RUN: %clang_cc1 -DUSE -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31A,CK31A-32,CK31A-USE12 13// RUN: %clang_cc1 -DUSE -DCK31A -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s14// RUN: %clang_cc1 -DUSE -DCK31A -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s15// RUN: %clang_cc1 -DUSE -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s16// RUN: %clang_cc1 -DUSE -DCK31A -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s17// RUN: %clang_cc1 -DUSE -DCK31A -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s18// RUN: %clang_cc1 -DUSE -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s19 20// RUN: %clang_cc1 -DCK31A -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31A,CK31A-64,CK31A-NOUSE21// RUN: %clang_cc1 -DCK31A -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s22// RUN: %clang_cc1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31A,CK31A-64,CK31A-NOUSE23// RUN: %clang_cc1 -DCK31A -verify -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31A,CK31A-32,CK31A-NOUSE24// RUN: %clang_cc1 -DCK31A -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s25// RUN: %clang_cc1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefixes=CK31A,CK31A-32,CK31A-NOUSE26 27// RUN: %clang_cc1 -DCK31A -verify -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s28// RUN: %clang_cc1 -DCK31A -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s29// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s30// RUN: %clang_cc1 -DCK31A -verify -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s31// RUN: %clang_cc1 -DCK31A -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s32// RUN: %clang_cc1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s33// SIMD-ONLY18-NOT: {{__kmpc|__tgt}}34#ifdef CK31A35 36// CK31A: [[ST:%.+]] = type { i32, i32 }37struct ST {38 int i;39 int j;40};41 42// CK31A-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 043//44// CK31A: [[SIZE00:@.+]] = private {{.*}}constant [7 x i64] [i64 0, i64 4, i64 4, i64 4, i64 0, i64 4, i64 4]45// PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x102046// CK31A-USE: [[MTYPE00:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x1020]],47// CK31A-NOUSE: [[MTYPE00:@.+]] = private {{.*}}constant [7 x i64] [i64 [[#0x1000]],48//49// MEMBER_OF_1=0x1000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x100000000100350// MEMBER_OF_1=0x1000000000000 | FROM=0x2 | TO=0x1 = 0x100000000000351// PRESENT=0x1000 | TARGET_PARAM=0x20 | FROM=0x2 | TO=0x1 = 0x102352// CK31A-USE-SAME: {{^}} i64 [[#0x1000000001003]], i64 [[#0x1000000000003]], i64 [[#0x1023]],53// CK31A-NOUSE-SAME: {{^}} i64 [[#0x1000000001003]], i64 [[#0x1000000000003]], i64 [[#0x1003]],54//55// PRESENT=0x1000 | TARGET_PARAM=0x20 = 0x102056// MEMBER_OF_5=0x5000000000000 | PRESENT=0x1000 | FROM=0x2 | TO=0x1 = 0x500000000100357// MEMBER_OF_5=0x5000000000000 | FROM=0x2 | TO=0x1 = 0x500000000000358// CK31A-USE-SAME: {{^}} i64 [[#0x1020]], i64 [[#0x5000000001003]], i64 [[#0x5000000000003]]]59// CK31A-NOUSE-SAME: {{^}} i64 [[#0x1000]], i64 [[#0x5000000001003]], i64 [[#0x5000000000003]]]60 61// CK31A-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 062// CK31A: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]63//64// PRESENT=0x1000 | CLOSE=0x400 | TARGET_PARAM=0x20 | ALWAYS=0x4 | FROM=0x2 | TO=0x1 = 0x142765// CK31A-USE: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 [[#0x1427]]]66// CK31A-NOUSE: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 [[#0x1407]]]67 68// CK31A-LABEL: explicit_maps_single{{.*}}(69void explicit_maps_single (int ii){70 // CK31A: alloca i3271 72 // Map of a scalar.73 // CK31A: [[A:%.+]] = alloca i3274 int a = ii;75 76 // CK31A: [[ST1:%.+]] = alloca [[ST]]77 // CK31A: [[ST2:%.+]] = alloca [[ST]]78 struct ST st1;79 struct ST st2;80 81// Make sure the struct picks up present even if another element of the struct82// doesn't have present.83// Region 0084// CK31A: [[ST1_J:%.+]] = getelementptr inbounds nuw [[ST]], ptr [[ST1]], i{{.+}} 0, i{{.+}} 185// CK31A: [[ST1_I:%.+]] = getelementptr inbounds nuw [[ST]], ptr [[ST1]], i{{.+}} 0, i{{.+}} 086// CK31A: [[ST2_I:%.+]] = getelementptr inbounds nuw [[ST]], ptr [[ST2]], i{{.+}} 0, i{{.+}} 087// CK31A: [[ST2_J:%.+]] = getelementptr inbounds nuw [[ST]], ptr [[ST2]], i{{.+}} 0, i{{.+}} 188// CK31A-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])89// CK31A-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 290// CK31A-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]91// CK31A-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 392// CK31A-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]93// CK31A-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 494// CK31A-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]95// CK31A-DAG: [[SIZES]] = getelementptr inbounds [7 x i64], ptr [[S:%.+]], i{{.+}} 0, i{{.+}} 096// CK31A-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]97// CK31A-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]98 99// st1100// CK31A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0101// CK31A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0102// CK31A-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0103// CK31A-DAG: store ptr [[ST1]], ptr [[BP0]]104// CK31A-DAG: store ptr [[ST1_I]], ptr [[P0]]105// CK31A-DAG: store i64 %{{.+}}, ptr [[S0]]106 107// st1.j108// CK31A-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1109// CK31A-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1110// CK31A-DAG: store ptr [[ST1]], ptr [[BP1]]111// CK31A-DAG: store ptr [[ST1_J]], ptr [[P1]]112 113// st1.i114// CK31A-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2115// CK31A-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2116// CK31A-DAG: store ptr [[ST1]], ptr [[BP2]]117// CK31A-DAG: store ptr [[ST1_I]], ptr [[P2]]118 119// a120// CK31A-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3121// CK31A-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3122// CK31A-DAG: store ptr [[A]], ptr [[BP3]]123// CK31A-DAG: store ptr [[A]], ptr [[P3]]124 125// st2126// CK31A-DAG: [[BP4:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 4127// CK31A-DAG: [[P4:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 4128// CK31A-DAG: [[S4:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 4129// CK31A-DAG: store ptr [[ST2]], ptr [[BP4]]130// CK31A-DAG: store ptr [[ST2_I]], ptr [[P4]]131// CK31A-DAG: store i64 %{{.+}}, ptr [[S4]]132 133// st2.i134// CK31A-DAG: [[BP5:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 5135// CK31A-DAG: [[P5:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 5136// CK31A-DAG: store ptr [[ST2]], ptr [[BP5]]137// CK31A-DAG: store ptr [[ST2_I]], ptr [[P5]]138 139// st2.j140// CK31A-DAG: [[BP6:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 6141// CK31A-DAG: [[P6:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 6142// CK31A-DAG: store ptr [[ST2]], ptr [[BP6]]143// CK31A-DAG: store ptr [[ST2_J]], ptr [[P6]]144 145// CK31A-USE: call void [[CALL00:@.+]](ptr [[ST1]], ptr [[A]], ptr [[ST2]])146// CK31A-NOUSE: call void [[CALL00:@.+]]()147#pragma omp target map(tofrom \148 : st1.i) map(present, tofrom \149 : a, st1.j, st2.i) map(tofrom \150 : st2.j)151 {152#ifdef USE153 st1.i++;154 a++;155 st1.j++;156 st2.i++;157 st2.j++;158#endif159 }160 161// Always Close Present.162// Region 01163// CK31A-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])164// CK31A-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2165// CK31A-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]166// CK31A-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3167// CK31A-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]168// CK31A-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]169// CK31A-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]170 171// CK31A-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0172// CK31A-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0173// CK31A-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]174// CK31A-DAG: store ptr [[VAR0]], ptr [[P0]]175 176// CK31A-USE: call void [[CALL01:@.+]](ptr {{[^,]+}})177// CK31A-NOUSE: call void [[CALL01:@.+]]()178#pragma omp target map(always close present tofrom \179 : a)180 {181#ifdef USE182 a++;183#endif184 }185}186// CK31A: define {{.+}}[[CALL00]]187// CK31A: define {{.+}}[[CALL01]]188 189#endif // CK31A190#endif191