1645 lines · cpp
1// expected-no-diagnostics2#ifndef HEADER3#define HEADER4 5#ifdef CK16 7///==========================================================================///8// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK19// Verify implicit-behavior 'storage' as alias for 'alloc' in OpenMP 6.010// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -DOMP6 -verify -Wno-vla -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK111// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s12// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK113// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK114// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s15// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK116 17// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -verify -Wno-vla -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-ONLY10 %s18// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s19// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s20// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -verify -Wno-vla -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-ONLY10 %s21// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK1 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s22// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s23// SIMD-ONLY10-NOT: {{__kmpc|__tgt}}24 25// CK1-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double_complex{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 026 27// CK1-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]28// Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 54429// CK1-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544]30 31// CK1-LABEL: implicit_maps_double_complex{{.*}}(32void implicit_maps_double_complex (int a){33 double _Complex dc = (double)a;34 35// CK1-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])36// CK1-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 237// CK1-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]38// CK1-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 339// CK1-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]40// CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 041// CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 042// CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 043// CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 044// CK1-DAG: store ptr [[PTR:%[^,]+]], ptr [[BP1]]45// CK1-DAG: store ptr [[PTR]], ptr [[P1]]46 47// CK1: call void [[KERNEL:@.+]](ptr [[PTR]])48#ifdef OMP649// 'storage' is an alias for 'alloc' in OpenMP 6.050#pragma omp target defaultmap(storage : scalar)51#else52#pragma omp target defaultmap(alloc : scalar)53#endif // OMP654 {55 dc *= dc;56 }57}58 59// CK1: define internal void [[KERNEL]](ptr {{.*}}[[ARG:%.+]])60// CK1: [[ADDR:%.+]] = alloca ptr,61// CK1: store ptr [[ARG]], ptr [[ADDR]],62// CK1: [[REF:%.+]] = load ptr, ptr [[ADDR]],63// CK1: {{.+}} = getelementptr inbounds nuw { double, double }, ptr [[REF]], i32 0, i32 064#endif65///==========================================================================///66// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK267// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s68// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK269// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK270// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s71// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK272 73// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -verify -Wno-vla -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-ONLY10 %s74// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s75// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s76// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -verify -Wno-vla -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-ONLY10 %s77// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK2 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s78// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s79// SIMD-ONLY10-NOT: {{__kmpc|__tgt}}80#ifdef CK281 82// CK2-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double_complex{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 083 84// CK2-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]85// Map types: OMP_MAP_TO | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 54586// CK2-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 545]87 88// CK2-LABEL: implicit_maps_double_complex{{.*}}(89void implicit_maps_double_complex (int a){90 double _Complex dc = (double)a;91 92// CK2-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])93// CK2-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 294// CK2-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]95// CK2-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 396// CK2-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]97// CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 098// CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 099// CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0100// CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0101// CK2-DAG: store ptr [[PTR:%[^,]+]], ptr [[BP1]]102// CK2-DAG: store ptr [[PTR]], ptr [[P1]]103 104// CK2: call void [[KERNEL:@.+]](ptr [[PTR]])105#pragma omp target defaultmap(to \106 : scalar)107 {108 dc *= dc;109 }110}111 112// CK2: define internal void [[KERNEL]](ptr {{.*}}[[ARG:%.+]])113// CK2: [[ADDR:%.+]] = alloca ptr,114// CK2: store ptr [[ARG]], ptr [[ADDR]],115// CK2: [[REF:%.+]] = load ptr, ptr [[ADDR]],116// CK2: {{.+}} = getelementptr inbounds nuw { double, double }, ptr [[REF]], i32 0, i32 0117#endif118///==========================================================================///119// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3120// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s121// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3122// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3123// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s124// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3125 126// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -verify -Wno-vla -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-ONLY10 %s127// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s128// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s129// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -verify -Wno-vla -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-ONLY10 %s130// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK3 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s131// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s132// SIMD-ONLY10-NOT: {{__kmpc|__tgt}}133#ifdef CK3134 135// CK3-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double_complex{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0136 137// CK3-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]138// Map types: OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 546139// CK3-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 546]140 141// CK3-LABEL: implicit_maps_double_complex{{.*}}(142void implicit_maps_double_complex (int a){143 double _Complex dc = (double)a;144 145// CK3-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])146// CK3-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2147// CK3-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]148// CK3-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3149// CK3-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]150// CK3-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0151// CK3-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0152// CK3-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0153// CK3-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0154// CK3-DAG: store ptr [[PTR:%[^,]+]], ptr [[BP1]]155// CK3-DAG: store ptr [[PTR]], ptr [[P1]]156 157// CK3: call void [[KERNEL:@.+]](ptr [[PTR]])158#pragma omp target defaultmap(from \159 : scalar)160 {161 dc *= dc;162 }163}164 165// CK3: define internal void [[KERNEL]](ptr {{.*}}[[ARG:%.+]])166// CK3: [[ADDR:%.+]] = alloca ptr,167// CK3: store ptr [[ARG]], ptr [[ADDR]],168// CK3: [[REF:%.+]] = load ptr, ptr [[ADDR]],169// CK3: {{.+}} = getelementptr inbounds nuw { double, double }, ptr [[REF]], i32 0, i32 0170#endif171///==========================================================================///172// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4 --check-prefix CK4-64173// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s174// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4 --check-prefix CK4-64175// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4 --check-prefix CK4-32176// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s177// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4 --check-prefix CK4-32178 179// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -verify -Wno-vla -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-ONLY6 %s180// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s181// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s182// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -verify -Wno-vla -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-ONLY6 %s183// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK4 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s184// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s185// SIMD-ONLY6-NOT: {{__kmpc|__tgt}}186#ifdef CK4187 188// For a 32-bit targets, the value doesn't fit the size of the pointer,189// therefore it is passed by reference with a map 'to' specification.190 191// CK4-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0192 193// CK4-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]194// Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800195// CK4-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 800]196// Map types: OMP_MAP_TO | OMP_MAP_PRIVATE | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 673197// CK4-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 673]198 199// CK4-LABEL: implicit_maps_double{{.*}}(200void implicit_maps_double (int a){201 double d = (double)a;202 203// CK4-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])204// CK4-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2205// CK4-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]206// CK4-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3207// CK4-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]208// CK4-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0209// CK4-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0210// CK4-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0211// CK4-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0212 213// CK4-64-DAG: store i[[sz:64|32]] [[VAL:%[^,]+]], ptr [[BP1]]214// CK4-64-DAG: store i[[sz]] [[VAL]], ptr [[P1]]215// CK4-64-DAG: [[VAL]] = load i[[sz]], ptr [[ADDR:%.+]],216// CK4-64-64-DAG: store double {{.+}}, ptr [[ADDR]],217 218// CK4-32-DAG: store ptr [[DECL:%[^,]+]], ptr [[BP1]]219// CK4-32-DAG: store ptr [[DECL]], ptr [[P1]]220 221// CK4-64: call void [[KERNEL:@.+]](i[[sz]] [[VAL]])222// CK4-32: call void [[KERNEL:@.+]](ptr [[DECL]])223#pragma omp target defaultmap(firstprivate \224 : scalar)225 {226 d += 1.0;227 }228}229 230// CK4-64: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]])231// CK4-64: [[ADDR:%.+]] = alloca i[[sz]],232// CK4-64: store i[[sz]] [[ARG]], ptr [[ADDR]],233// CK4-64: {{.+}} = load double, ptr [[ADDR]],234 235// CK4-32: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]])236// CK4-32: [[ADDR:%.+]] = alloca ptr,237// CK4-32: store ptr [[ARG]], ptr [[ADDR]],238// CK4-32: [[REF:%.+]] = load ptr, ptr [[ADDR]],239// CK4-32: {{.+}} = load double, ptr [[REF]],240 241#endif242///==========================================================================///243// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK5 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK5244// Verify implicit-behavior 'alloc' still accepted in OpenMP 6.0245// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK5 -verify -Wno-vla -fopenmp -fopenmp-version=60 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK5246// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK5 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s247// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK5248// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK5 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK5249// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK5 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s250// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK5251 252// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK5 -verify -Wno-vla -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-ONLY8 %s253// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK5 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s254// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s255// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK5 -verify -Wno-vla -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-ONLY8 %s256// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK5 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s257// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s258// SIMD-ONLY8-NOT: {{__kmpc|__tgt}}259#ifdef CK5260 261// CK5-LABEL: @.__omp_offloading_{{.*}}implicit_maps_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0262 263// CK5-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]264// Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544265// CK5-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544]266 267// CK5-LABEL: implicit_maps_array{{.*}}(268void implicit_maps_array (int a){269 double darr[2] = {(double)a, (double)a};270 271// CK5-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])272// CK5-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2273// CK5-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]274// CK5-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3275// CK5-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]276// CK5-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0277// CK5-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0278// CK5-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0279// CK5-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0280// CK5-DAG: store ptr [[DECL:%[^,]+]], ptr [[BP1]]281// CK5-DAG: store ptr [[DECL]], ptr [[P1]]282 283// CK5: call void [[KERNEL:@.+]](ptr [[DECL]])284#pragma omp target defaultmap(alloc \285 : aggregate)286 {287 darr[0] += 1.0;288 darr[1] += 1.0;289 }290}291 292// CK5: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]])293// CK5: [[ADDR:%.+]] = alloca ptr,294// CK5: store ptr [[ARG]], ptr [[ADDR]],295// CK5: [[REF:%.+]] = load ptr, ptr [[ADDR]],296// CK5: {{.+}} = getelementptr inbounds [2 x double], ptr [[REF]], i{{64|32}} 0, i{{64|32}} 0297#endif298///==========================================================================///299// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK6 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK6300// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK6 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s301// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK6302// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK6 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK6303// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK6 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s304// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK6305 306// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK6 -verify -Wno-vla -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-ONLY8 %s307// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK6 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s308// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s309// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK6 -verify -Wno-vla -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-ONLY8 %s310// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK6 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s311// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s312// SIMD-ONLY8-NOT: {{__kmpc|__tgt}}313#ifdef CK6314 315// CK6-LABEL: @.__omp_offloading_{{.*}}implicit_maps_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0316 317// CK6-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]318// Map types: OMP_MAP_TO | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 545319// CK6-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 545]320 321// CK6-LABEL: implicit_maps_array{{.*}}(322void implicit_maps_array (int a){323 double darr[2] = {(double)a, (double)a};324 325 // CK6-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])326 // CK6-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2327 // CK6-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]328 // CK6-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3329 // CK6-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]330 // CK6-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0331 // CK6-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0332 // CK6-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0333 // CK6-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0334 // CK6-DAG: store ptr [[DECL:%[^,]+]], ptr [[BP1]]335 // CK6-DAG: store ptr [[DECL]], ptr [[P1]]336 337 // CK6: call void [[KERNEL:@.+]](ptr [[DECL]])338#pragma omp target defaultmap(to)339 {340 darr[0] += 1.0;341 darr[1] += 1.0;342 }343}344 345// CK6: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]])346// CK6: [[ADDR:%.+]] = alloca ptr,347// CK6: store ptr [[ARG]], ptr [[ADDR]],348// CK6: [[REF:%.+]] = load ptr, ptr [[ADDR]],349// CK6: {{.+}} = getelementptr inbounds [2 x double], ptr [[REF]], i{{64|32}} 0, i{{64|32}} 0350#endif351///==========================================================================///352// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK7 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK7353// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK7 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s354// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK7355// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK7 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK7356// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK7 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s357// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK7358 359// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK7 -verify -Wno-vla -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-ONLY8 %s360// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK7 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s361// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s362// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK7 -verify -Wno-vla -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-ONLY8 %s363// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK7 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s364// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s365// SIMD-ONLY8-NOT: {{__kmpc|__tgt}}366#ifdef CK7367 368// CK7-LABEL: @.__omp_offloading_{{.*}}implicit_maps_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0369 370// CK7-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]371// Map types: OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 546372// CK7-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 546]373 374// CK7-LABEL: implicit_maps_array{{.*}}(375void implicit_maps_array (int a){376 double darr[2] = {(double)a, (double)a};377 378// CK7-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])379// CK7-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2380// CK7-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]381// CK7-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3382// CK7-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]383// CK7-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0384// CK7-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0385// CK7-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0386// CK7-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0387// CK7-DAG: store ptr [[DECL:%[^,]+]], ptr [[BP1]]388// CK7-DAG: store ptr [[DECL]], ptr [[P1]]389 390// CK7: call void [[KERNEL:@.+]](ptr [[DECL]])391#pragma omp target defaultmap(from \392 : aggregate)393 {394 darr[0] += 1.0;395 darr[1] += 1.0;396 }397}398 399// CK7: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]])400// CK7: [[ADDR:%.+]] = alloca ptr,401// CK7: store ptr [[ARG]], ptr [[ADDR]],402// CK7: [[REF:%.+]] = load ptr, ptr [[ADDR]],403// CK7: {{.+}} = getelementptr inbounds [2 x double], ptr [[REF]], i{{64|32}} 0, i{{64|32}} 0404#endif405///==========================================================================///406// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK8 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK8407// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK8 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s408// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK8409// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK8 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK8410// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK8 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s411// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK8412 413// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK8 -verify -Wno-vla -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-ONLY8 %s414// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK8 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s415// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s416// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK8 -verify -Wno-vla -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-ONLY8 %s417// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK8 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s418// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s419// SIMD-ONLY8-NOT: {{__kmpc|__tgt}}420#ifdef CK8421 422// CK8-LABEL: @.__omp_offloading_{{.*}}implicit_maps_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0423 424// CK8-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]425// Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547426// CK8-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 547]427 428// CK8-LABEL: implicit_maps_array{{.*}}(429void implicit_maps_array (int a){430 double darr[2] = {(double)a, (double)a};431 432 // CK8-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])433 // CK8-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2434 // CK8-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]435 // CK8-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3436 // CK8-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]437 // CK8-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0438 // CK8-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0439 // CK8-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0440 // CK8-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0441 // CK8-DAG: store ptr [[DECL:%[^,]+]], ptr [[BP1]]442 // CK8-DAG: store ptr [[DECL]], ptr [[P1]]443 444 // CK8: call void [[KERNEL:@.+]](ptr [[DECL]])445#pragma omp target defaultmap(tofrom)446 {447 darr[0] += 1.0;448 darr[1] += 1.0;449 }450}451 452// CK8: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]])453// CK8: [[ADDR:%.+]] = alloca ptr,454// CK8: store ptr [[ARG]], ptr [[ADDR]],455// CK8: [[REF:%.+]] = load ptr, ptr [[ADDR]],456// CK8: {{.+}} = getelementptr inbounds [2 x double], ptr [[REF]], i{{64|32}} 0, i{{64|32}} 0457#endif458 459///==========================================================================///460// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK9 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK9 --check-prefix CK9-64461// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK9 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s462// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK9 --check-prefix CK9-64463// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK9 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK9 --check-prefix CK9-32464// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK9 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s465// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK9 --check-prefix CK9-32466 467// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK9 -verify -Wno-vla -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-ONLY26 %s468// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK9 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s469// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY26 %s470// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK9 -verify -Wno-vla -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-ONLY26 %s471// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK9 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s472// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY26 %s473// SIMD-ONLY26-NOT: {{__kmpc|__tgt}}474 475#ifdef CK9476 477 478// CK9-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0479// CK9: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 40]480// CK9: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 673]481 482 483// CK9-LABEL: zero_size_section_and_private_maps{{.*}}(484void zero_size_section_and_private_maps (int ii){485 int pvtArr[10];486 487// Region 09488// CK9-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])489// CK9-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2490// CK9-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]491// CK9-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3492// CK9-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]493// CK9-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]494// CK9-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]495 496// CK9-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0497// CK9-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0498// CK9-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]499// CK9-DAG: store ptr [[VAR0]], ptr [[P0]]500 501// CK9: call void [[CALL09:@.+]](ptr {{[^,]+}})502#pragma omp target defaultmap(firstprivate \503 : aggregate)504 {505 pvtArr[5]++;506 }507 508}509 510 511#endif512///==========================================================================///513// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK10 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK10514// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK10 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s515// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK10516// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK10 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK10517// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK10 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s518// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK10519 520// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK10 -verify -Wno-vla -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-ONLY8 %s521// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK10 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s522// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s523// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK10 -verify -Wno-vla -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-ONLY8 %s524// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK10 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s525// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s526// SIMD-ONLY8-NOT: {{__kmpc|__tgt}}527#ifdef CK10528 529 530// CK10-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0531 532// CK10: [[SIZE:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}]533// Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544534// CK10: [[MTYPE:@.+]] = private {{.*}}constant [1 x i64] [i64 544]535 536// CK10-LABEL: explicit_maps_single{{.*}}(537void explicit_maps_single (){538 int *pa;539 540// CK10-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])541// CK10-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2542// CK10-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]543// CK10-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3544// CK10-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]545// CK10-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]546// CK10-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]547 548// CK10-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0549// CK10-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0550// CK10-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]551// CK10-DAG: store ptr [[VAR0]], ptr [[P0]]552 553// CK10: call void [[CALL:@.+]](ptr {{[^,]+}})554#pragma omp target defaultmap(alloc \555 : pointer)556 {557 pa[50]++;558 }559}560 561 // CK10: define {{.+}}[[CALL]]562#endif563///==========================================================================///564// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK11 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK11565// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK11 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s566// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK11567// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK11 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK11568// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK11 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s569// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK11570 571// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK11 -verify -Wno-vla -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-ONLY8 %s572// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK11 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s573// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s574// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK11 -verify -Wno-vla -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-ONLY8 %s575// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK11 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s576// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s577// SIMD-ONLY8-NOT: {{__kmpc|__tgt}}578#ifdef CK11579 580 581// CK11-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0582 583// CK11: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}]584// Map types: OMP_MAP_TO | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 545585// CK11: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 545]586 587// CK11-LABEL: explicit_maps_single{{.*}}(588void explicit_maps_single (){589 int *pa;590 591// CK11-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])592// CK11-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2593// CK11-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]594// CK11-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3595// CK11-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]596// CK11-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]597// CK11-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]598 599// CK11-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0600// CK11-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0601// CK11-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]602// CK11-DAG: store ptr [[VAR0]], ptr [[P0]]603 604// CK11: call void [[CALL09:@.+]](ptr {{[^,]+}})605#pragma omp target defaultmap(to \606 : pointer)607 {608 pa[50]++;609 }610}611 612 // CK11: define {{.+}}[[CALL09]]613#endif614///==========================================================================///615// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK12 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK12616// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK12 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s617// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK12618// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK12 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK12619// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK12 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s620// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK12621 622// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK12 -verify -Wno-vla -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-ONLY8 %s623// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK12 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s624// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s625// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK12 -verify -Wno-vla -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-ONLY8 %s626// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK12 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s627// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s628// SIMD-ONLY8-NOT: {{__kmpc|__tgt}}629#ifdef CK12630 631 632// CK12-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0633 634// CK12: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}]635// Map types: OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 546636// CK12: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 546]637 638// CK12-LABEL: explicit_maps_single{{.*}}(639void explicit_maps_single (){640 int *pa;641 642// CK12-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])643// CK12-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2644// CK12-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]645// CK12-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3646// CK12-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]647// CK12-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]648// CK12-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]649 650// CK12-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0651// CK12-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0652// CK12-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]653// CK12-DAG: store ptr [[VAR0]], ptr [[P0]]654 655// CK12: call void [[CALL09:@.+]](ptr {{[^,]+}})656#pragma omp target defaultmap(from \657 : pointer)658 {659 pa[50]++;660 }661}662 663 // CK12: define {{.+}}[[CALL09]]664#endif665///==========================================================================///666// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK13 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK13667// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK13 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s668// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK13669// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK13 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK13670// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK13 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s671// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK13672 673// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK13 -verify -Wno-vla -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-ONLY8 %s674// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK13 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s675// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s676// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK13 -verify -Wno-vla -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-ONLY8 %s677// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK13 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s678// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s679// SIMD-ONLY8-NOT: {{__kmpc|__tgt}}680#ifdef CK13681 682 683// CK13-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0684 685// CK13: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}]686// Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547687// CK13: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 547]688 689// CK13-LABEL: explicit_maps_single{{.*}}(690void explicit_maps_single (){691 int *pa;692 693// CK13-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])694// CK13-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2695// CK13-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]696// CK13-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3697// CK13-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]698// CK13-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]699// CK13-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]700 701// CK13-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0702// CK13-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0703// CK13-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]704// CK13-DAG: store ptr [[VAR0]], ptr [[P0]]705 706// CK13: call void [[CALL09:@.+]](ptr {{[^,]+}})707#pragma omp target defaultmap(tofrom \708 : pointer)709 {710 pa[50]++;711 }712}713 714 // CK13: define {{.+}}[[CALL09]]715#endif716///==========================================================================///717// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK14 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK14718// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK14 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s719// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK14720// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK14 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK14721// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK14 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s722// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK14723 724// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK14 -verify -Wno-vla -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-ONLY8 %s725// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK14 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s726// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s727// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK14 -verify -Wno-vla -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-ONLY8 %s728// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK14 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s729// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s730// SIMD-ONLY8-NOT: {{__kmpc|__tgt}}731#ifdef CK14732 733 734// CK14-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0735 736// CK14: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer737// Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544738// CK14: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 544]739 740// CK14-LABEL: explicit_maps_single{{.*}}(741void explicit_maps_single (){742 int *pa;743 744// CK14-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])745// CK14-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2746// CK14-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]747// CK14-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3748// CK14-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]749// CK14-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]750// CK14-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]751 752// CK14-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0753// CK14-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0754// CK14-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]755// CK14-DAG: store ptr [[VAR0]], ptr [[P0]]756 757// CK14: call void [[CALL09:@.+]](ptr {{[^,]+}})758#pragma omp target defaultmap(firstprivate \759 : pointer)760 {761 pa[50]++;762 }763}764 765 // CK14: define {{.+}}[[CALL09]]766#endif767///==========================================================================///768// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK15 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK15769// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK15 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s770// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK15771// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK15 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK15772// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK15 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s773// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK15774 775// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK15 -verify -Wno-vla -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-ONLY12 %s776// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK15 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s777// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY12 %s778// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK15 -verify -Wno-vla -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-ONLY12 %s779// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK15 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s780// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY12 %s781// SIMD-ONLY12-NOT: {{__kmpc|__tgt}}782#ifdef CK15783 784// CK15-LABEL: @.__omp_offloading_{{.*}}implicit_maps_variable_length_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0785 786// CK15-DAG: [[SIZES:@.+]] = {{.+}}constant [3 x i64] [i64 {{4|8}}, i64 {{4|8}}, i64 0]787// Map types:788// - OMP_MAP_LITERAL + OMP_MAP_TARGET_PARAM + OMP_MAP_IMPLICIT = 800 (vla size)789// - OMP_MAP_LITERAL + OMP_MAP_TARGET_PARAM + OMP_MAP_IMPLICIT = 800 (vla size)790// - OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544791// CK15-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i64] [i64 800, i64 800, i64 544]792 793// CK15-LABEL: implicit_maps_variable_length_array{{.*}}(794void implicit_maps_variable_length_array (int a){795 double vla[2][a];796 797// CK15-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])798// CK15-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2799// CK15-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]800// CK15-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3801// CK15-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]802// CK15-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4803// CK15-DAG: store ptr [[SGEP:%.+]], ptr [[SARG]]804// CK15-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0805// CK15-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0806// CK15-DAG: [[SGEP]] = getelementptr inbounds {{.+}}[[SS:%[^,]+]], i32 0, i32 0807 808// CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0809// CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0810// CK15-DAG: store i[[sz:64|32]] 2, ptr [[BP0]]811// CK15-DAG: store i[[sz]] 2, ptr [[P0]]812 813// CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1814// CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1815// CK15-DAG: store i[[sz]] [[VAL:%.+]], ptr [[BP1]]816// CK15-DAG: store i[[sz]] [[VAL]], ptr [[P1]]817 818// CK15-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2819// CK15-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2820// CK15-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[SS]], i32 0, i32 2821// CK15-DAG: store ptr [[DECL:%.+]], ptr [[BP2]]822// CK15-DAG: store ptr [[DECL]], ptr [[P2]]823// CK15-DAG: store i64 [[VALS2:%.+]], ptr [[S2]],824// CK15-DAG: [[VALS2]] = {{mul nuw i64 %.+, 8|sext i32 %.+ to i64}}825 826// CK15: call void [[KERNEL:@.+]](i[[sz]] {{.+}}, i[[sz]] {{.+}}, ptr [[DECL]])827#pragma omp target defaultmap(alloc \828 : aggregate)829 {830 vla[1][3] += 1.0;831 }832}833 834// CK15: define internal void [[KERNEL]](i[[sz]] [[VLA0:%.+]], i[[sz]] [[VLA1:%.+]], ptr {{.*}}[[ARG:%.+]])835// CK15: [[ADDR0:%.+]] = alloca i[[sz]],836// CK15: [[ADDR1:%.+]] = alloca i[[sz]],837// CK15: [[ADDR2:%.+]] = alloca ptr,838// CK15: store i[[sz]] [[VLA0]], ptr [[ADDR0]],839// CK15: store i[[sz]] [[VLA1]], ptr [[ADDR1]],840// CK15: store ptr [[ARG]], ptr [[ADDR2]],841// CK15: {{.+}} = load i[[sz]], ptr [[ADDR0]],842// CK15: {{.+}} = load i[[sz]], ptr [[ADDR1]],843// CK15: [[REF:%.+]] = load ptr, ptr [[ADDR2]],844// CK15: {{.+}} = getelementptr inbounds double, ptr [[REF]], i[[sz]] %{{.+}}845#endif846///==========================================================================///847// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK16 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK16848// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK16 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s849// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK16850// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK16 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK16851// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK16 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s852// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK16853 854// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK16 -verify -Wno-vla -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-ONLY16 %s855// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK16 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s856// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s857// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK16 -verify -Wno-vla -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-ONLY16 %s858// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK16 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s859// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s860// SIMD-ONLY16-NOT: {{__kmpc|__tgt}}861#ifdef CK16862 863// CK16-DAG: [[ST:%.+]] = type { i32, double }864// CK16-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0865// CK16-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}]866// Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544867// CK16-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544]868 869class SSS {870public:871 int a;872 double b;873};874 875// CK16-LABEL: implicit_maps_struct{{.*}}(876void implicit_maps_struct (int a){877 SSS s = {a, (double)a};878 879// CK16-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])880// CK16-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2881// CK16-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]882// CK16-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3883// CK16-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]884// CK16-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0885// CK16-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0886// CK16-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0887// CK16-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0888// CK16-DAG: store ptr [[DECL:%.+]], ptr [[BP1]]889// CK16-DAG: store ptr [[DECL]], ptr [[P1]]890 891// CK16: call void [[KERNEL:@.+]](ptr [[DECL]])892#pragma omp target defaultmap(alloc \893 : aggregate)894 {895 s.a += 1;896 s.b += 1.0;897 }898}899 900// CK16: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]])901// CK16: [[ADDR:%.+]] = alloca ptr,902// CK16: store ptr [[ARG]], ptr [[ADDR]],903// CK16: [[REF:%.+]] = load ptr, ptr [[ADDR]],904// CK16: {{.+}} = getelementptr inbounds nuw [[ST]], ptr [[REF]], i32 0, i32 0905#endif906///==========================================================================///907// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK17 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK17908// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK17 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s909// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK17910// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK17 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK17911// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK17 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s912// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK17913 914// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK17 -verify -Wno-vla -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-ONLY16 %s915// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK17 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s916// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s917// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK17 -verify -Wno-vla -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-ONLY16 %s918// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK17 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s919// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s920// SIMD-ONLY16-NOT: {{__kmpc|__tgt}}921#ifdef CK17922 923// CK17-DAG: [[ST:%.+]] = type { i32, double }924// CK17-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0925// CK17-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}]926// Map types: OMP_MAP_TO | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 545927// CK17-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 545]928 929class SSS {930public:931 int a;932 double b;933};934 935// CK17-LABEL: implicit_maps_struct{{.*}}(936void implicit_maps_struct (int a){937 SSS s = {a, (double)a};938 939// CK17-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])940// CK17-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2941// CK17-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]942// CK17-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3943// CK17-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]944// CK17-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0945// CK17-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0946// CK17-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0947// CK17-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0948// CK17-DAG: store ptr [[DECL:%.+]], ptr [[BP1]]949// CK17-DAG: store ptr [[DECL]], ptr [[P1]]950 951// CK17: call void [[KERNEL:@.+]](ptr [[DECL]])952#pragma omp target defaultmap(to \953 : aggregate)954 {955 s.a += 1;956 s.b += 1.0;957 }958}959 960// CK17: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]])961// CK17: [[ADDR:%.+]] = alloca ptr,962// CK17: store ptr [[ARG]], ptr [[ADDR]],963// CK17: [[REF:%.+]] = load ptr, ptr [[ADDR]],964// CK17: {{.+}} = getelementptr inbounds nuw [[ST]], ptr [[REF]], i32 0, i32 0965#endif966///==========================================================================///967// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK18 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK18968// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK18 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s969// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK18970// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK18 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK18971// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK18 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s972// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK18973 974// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK18 -verify -Wno-vla -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-ONLY16 %s975// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK18 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s976// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s977// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK18 -verify -Wno-vla -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-ONLY16 %s978// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK18 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s979// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s980// SIMD-ONLY16-NOT: {{__kmpc|__tgt}}981#ifdef CK18982 983// CK18-DAG: [[ST:%.+]] = type { i32, double }984// CK18-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0985// CK18-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}]986// Map types: OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 546987// CK18-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 546]988 989class SSS {990public:991 int a;992 double b;993};994 995// CK18-LABEL: implicit_maps_struct{{.*}}(996void implicit_maps_struct (int a){997 SSS s = {a, (double)a};998 999// CK18-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1000// CK18-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21001// CK18-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1002// CK18-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31003// CK18-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1004// CK18-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 01005// CK18-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 01006// CK18-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 01007// CK18-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 01008// CK18-DAG: store ptr [[DECL:%.+]], ptr [[BP1]]1009// CK18-DAG: store ptr [[DECL]], ptr [[P1]]1010 1011// CK18: call void [[KERNEL:@.+]](ptr [[DECL]])1012#pragma omp target defaultmap(from \1013 : aggregate)1014 {1015 s.a += 1;1016 s.b += 1.0;1017 }1018}1019 1020// CK18: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]])1021// CK18: [[ADDR:%.+]] = alloca ptr,1022// CK18: store ptr [[ARG]], ptr [[ADDR]],1023// CK18: [[REF:%.+]] = load ptr, ptr [[ADDR]],1024// CK18: {{.+}} = getelementptr inbounds nuw [[ST]], ptr [[REF]], i32 0, i32 01025#endif1026///==========================================================================///1027// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK19 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK191028// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK19 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s1029// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK191030// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK19 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK191031// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK19 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s1032// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK191033 1034// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK19 -verify -Wno-vla -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-ONLY16 %s1035// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK19 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s1036// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s1037// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK19 -verify -Wno-vla -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-ONLY16 %s1038// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK19 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s1039// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s1040// SIMD-ONLY16-NOT: {{__kmpc|__tgt}}1041#ifdef CK191042 1043// CK19-DAG: [[ST:%.+]] = type { i32, double }1044// CK19-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 01045// CK19-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}]1046// Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 5471047// CK19-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 547]1048 1049class SSS {1050public:1051 int a;1052 double b;1053};1054 1055// CK19-LABEL: implicit_maps_struct{{.*}}(1056void implicit_maps_struct (int a){1057 SSS s = {a, (double)a};1058 1059// CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1060// CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21061// CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1062// CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31063// CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1064// CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 01065// CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 01066// CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 01067// CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 01068// CK19-DAG: store ptr [[DECL:%.+]], ptr [[BP1]]1069// CK19-DAG: store ptr [[DECL]], ptr [[P1]]1070 1071// CK19: call void [[KERNEL:@.+]](ptr [[DECL]])1072#pragma omp target defaultmap(tofrom \1073 : aggregate)1074 {1075 s.a += 1;1076 s.b += 1.0;1077 }1078}1079 1080// CK19: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]])1081// CK19: [[ADDR:%.+]] = alloca ptr,1082// CK19: store ptr [[ARG]], ptr [[ADDR]],1083// CK19: [[REF:%.+]] = load ptr, ptr [[ADDR]],1084// CK19: {{.+}} = getelementptr inbounds nuw [[ST]], ptr [[REF]], i32 0, i32 01085#endif1086///==========================================================================///1087// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK20 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK20 --check-prefix CK20-641088// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK20 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s1089// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK20 --check-prefix CK20-641090// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK20 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK20 --check-prefix CK20-321091// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK20 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s1092// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK20 --check-prefix CK20-321093 1094// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK20 -verify -Wno-vla -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-ONLY6 %s1095// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK20 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s1096// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s1097// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK20 -verify -Wno-vla -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-ONLY6 %s1098// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK20 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s1099// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s1100// SIMD-ONLY6-NOT: {{__kmpc|__tgt}}1101#ifdef CK201102 1103// For a 32-bit targets, the value doesn't fit the size of the pointer,1104// therefore it is passed by reference with a map 'to' specification.1105 1106// CK20-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 01107 1108// CK20-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8]1109// Map types: OMP_MAP_LITERAL | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 8001110// CK20-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 800]1111// Map types: OMP_MAP_TO | OMP_MAP_PRIVATE | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 6731112// CK20-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 673]1113 1114// CK20-LABEL: implicit_maps_double{{.*}}(1115void implicit_maps_double (int a){1116 double d = (double)a;1117 1118// CK20-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1119// CK20-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21120// CK20-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1121// CK20-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31122// CK20-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1123// CK20-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 01124// CK20-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 01125// CK20-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 01126// CK20-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 01127 1128// CK20-64-DAG: store i[[sz:64|32]] [[VAL:%[^,]+]], ptr [[BP1]]1129// CK20-64-DAG: store i[[sz]] [[VAL]], ptr [[P1]]1130// CK20-64-DAG: [[VAL]] = load i[[sz]], ptr [[ADDR:%.+]],1131// CK20-64-64-DAG: store double {{.+}}, ptr [[ADDR]],1132 1133// CK20-32-DAG: store ptr [[DECL:%[^,]+]], ptr [[BP1]]1134// CK20-32-DAG: store ptr [[DECL]], ptr [[P1]]1135 1136// CK20-64: call void [[KERNEL:@.+]](i[[sz]] [[VAL]])1137// CK20-32: call void [[KERNEL:@.+]](ptr [[DECL]])1138#pragma omp target defaultmap(default \1139 : scalar)1140 {1141 d += 1.0;1142 }1143}1144 1145// CK20-64: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]])1146// CK20-64: [[ADDR:%.+]] = alloca i[[sz]],1147// CK20-64: store i[[sz]] [[ARG]], ptr [[ADDR]],1148// CK20-64: {{.+}} = load double, ptr [[ADDR]],1149 1150// CK20-32: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]])1151// CK20-32: [[ADDR:%.+]] = alloca ptr,1152// CK20-32: store ptr [[ARG]], ptr [[ADDR]],1153// CK20-32: [[REF:%.+]] = load ptr, ptr [[ADDR]],1154// CK20-32: {{.+}} = load double, ptr [[REF]],1155 1156#endif1157///==========================================================================///1158// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK21 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK211159// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK21 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s1160// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK211161// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK21 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK211162// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK21 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s1163// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK211164 1165// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK21 -verify -Wno-vla -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-ONLY16 %s1166// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK21 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s1167// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s1168// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK21 -verify -Wno-vla -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-ONLY16 %s1169// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK21 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s1170// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s1171// SIMD-ONLY16-NOT: {{__kmpc|__tgt}}1172#ifdef CK211173 1174// CK21-DAG: [[ST:%.+]] = type { i32, double }1175// CK21-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 01176// CK21-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}]1177// Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 5471178// CK21-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 547]1179 1180class SSS {1181public:1182 int a;1183 double b;1184};1185 1186// CK21-LABEL: implicit_maps_struct{{.*}}(1187void implicit_maps_struct (int a){1188 SSS s = {a, (double)a};1189 1190// CK21-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1191// CK21-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21192// CK21-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1193// CK21-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31194// CK21-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1195// CK21-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 01196// CK21-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 01197// CK21-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 01198// CK21-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 01199// CK21-DAG: store ptr [[DECL:%.+]], ptr [[BP1]]1200// CK21-DAG: store ptr [[DECL]], ptr [[P1]]1201 1202// CK21: call void [[KERNEL:@.+]](ptr [[DECL]])1203#pragma omp target defaultmap(default \1204 : aggregate)1205 {1206 s.a += 1;1207 s.b += 1.0;1208 }1209}1210 1211// CK21: define internal void [[KERNEL]](ptr {{.+}}[[ARG:%.+]])1212// CK21: [[ADDR:%.+]] = alloca ptr,1213// CK21: store ptr [[ARG]], ptr [[ADDR]],1214// CK21: [[REF:%.+]] = load ptr, ptr [[ADDR]],1215// CK21: {{.+}} = getelementptr inbounds nuw [[ST]], ptr [[REF]], i32 0, i32 01216#endif1217///==========================================================================///1218// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK22 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK221219// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK22 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s1220// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK221221// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK22 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK221222// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK22 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s1223// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK221224 1225// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK22 -verify -Wno-vla -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-ONLY9 %s1226// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK22 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s1227// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY9 %s1228// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK22 -verify -Wno-vla -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-ONLY9 %s1229// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK22 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s1230// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY9 %s1231// SIMD-ONLY9-NOT: {{__kmpc|__tgt}}1232#ifdef CK221233 1234// CK22-LABEL: @.__omp_offloading_{{.*}}implicit_maps_pointer{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 01235 1236// CK22-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] zeroinitializer1237// Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 5441238// CK22-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544]1239 1240// CK22-LABEL: implicit_maps_pointer{{.*}}(1241void implicit_maps_pointer (){1242 double *ddyn;1243 1244// CK22-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1245// CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21246// CK22-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1247// CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31248// CK22-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1249// CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 01250// CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 01251// CK22-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 01252// CK22-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 01253// CK22-DAG: store ptr [[PTR:%[^,]+]], ptr [[BP1]]1254// CK22-DAG: store ptr [[PTR]], ptr [[P1]]1255 1256// CK22: call void [[KERNEL:@.+]](ptr [[PTR]])1257#pragma omp target defaultmap(default \1258 : pointer)1259 {1260 ddyn[0] += 1.0;1261 ddyn[1] += 1.0;1262 }1263}1264 1265// CK22: define internal void [[KERNEL]](ptr {{.*}}[[ARG:%.+]])1266// CK22: [[ADDR:%.+]] = alloca ptr,1267// CK22: store ptr [[ARG]], ptr [[ADDR]],1268// CK22: [[REF:%.+]] = load ptr, ptr [[ADDR]],1269// CK22: {{.+}} = getelementptr inbounds double, ptr [[REF]], i{{64|32}} 01270 1271#endif1272///==========================================================================///1273// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK23 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-641274// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK23 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s1275// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-641276// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK23 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-321277// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK23 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s1278// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-321279 1280// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK23 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s1281// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK23 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s1282// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s1283// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK23 -verify -Wno-vla -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s1284// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK23 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s1285// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s1286// SIMD-ONLY0-NOT: {{__kmpc|__tgt}}1287#ifdef CK231288 1289double *g;1290 1291// CK23: @g ={{.*}} global ptr1292// CK23: [[SIZES00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}]1293// CK23: [[TYPES00:@.+]] = {{.+}}constant [1 x i64] [i64 288]1294 1295// CK23: [[SIZES01:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]1296// CK23: [[TYPES01:@.+]] = {{.+}}constant [1 x i64] [i64 288]1297 1298// CK23: [[SIZES02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]1299// CK23: [[TYPES02:@.+]] = {{.+}}constant [1 x i64] [i64 288]1300 1301// CK23: [[SIZES03:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]1302// CK23: [[TYPES03:@.+]] = {{.+}}constant [1 x i64] [i64 288]1303 1304// CK23: [[SIZES04:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]1305// CK23: [[TYPES04:@.+]] = {{.+}}constant [1 x i64] [i64 288]1306 1307// CK23: [[SIZES05:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}]1308// CK23: [[TYPES05:@.+]] = {{.+}}constant [1 x i64] [i64 288]1309 1310// CK23: [[SIZES06:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] {{8|4}}]1311// CK23: [[TYPES06:@.+]] = {{.+}}constant [2 x i64] [i64 288, i64 288]1312 1313// CK23-LABEL: @_Z3foo{{.*}}(1314template<typename T>1315void foo(float *&lr, T *&tr) {1316 float *l;1317 T *t;1318 1319// CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1320// CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21321// CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1322// CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31323// CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1324// CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 01325// CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 01326// CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 01327// CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 01328// CK23-DAG: store ptr [[VAL:%.+]], ptr [[BP1]]1329// CK23-DAG: store ptr [[VAL]], ptr [[P1]]1330// CK23-DAG: [[VAL]] = load ptr, ptr [[ADDR:@g]],1331 1332// CK23: call void [[KERNEL:@.+]](ptr [[VAL]])1333#pragma omp target is_device_ptr(g) defaultmap(none \1334 : pointer)1335 {1336 ++g;1337 }1338 1339// CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1340// CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21341// CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1342// CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31343// CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1344// CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 01345// CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 01346// CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 01347// CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 01348// CK23-DAG: store ptr [[VAL:%.+]], ptr [[BP1]]1349// CK23-DAG: store ptr [[VAL]], ptr [[P1]]1350// CK23-DAG: [[VAL]] = load ptr, ptr [[ADDR:%.+]],1351 1352// CK23: call void [[KERNEL:@.+]](ptr [[VAL]])1353#pragma omp target is_device_ptr(l) defaultmap(none \1354 : pointer)1355 {1356 ++l;1357 }1358 1359// CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1360// CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21361// CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1362// CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31363// CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1364// CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 01365// CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 01366// CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 01367// CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 01368// CK23-DAG: store ptr [[VAL:%.+]], ptr [[BP1]]1369// CK23-DAG: store ptr [[VAL]], ptr [[P1]]1370// CK23-DAG: [[VAL]] = load ptr, ptr [[ADDR:%.+]],1371 1372// CK23: call void [[KERNEL:@.+]](ptr [[VAL]])1373#pragma omp target is_device_ptr(t) defaultmap(none \1374 : pointer)1375 {1376 ++t;1377 }1378 1379// CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1380// CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21381// CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1382// CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31383// CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1384// CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 01385// CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 01386// CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 01387// CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 01388// CK23-DAG: store ptr [[VAL:%.+]], ptr [[BP1]]1389// CK23-DAG: store ptr [[VAL]], ptr [[P1]]1390// CK23-DAG: [[VAL]] = load ptr, ptr [[ADDR:%.+]],1391// CK23-DAG: [[ADDR]] = load ptr, ptr [[ADDR2:%.+]],1392 1393// CK23: call void [[KERNEL:@.+]](ptr [[VAL]])1394#pragma omp target is_device_ptr(lr) defaultmap(none \1395 : pointer)1396 {1397 ++lr;1398 }1399 1400// CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1401// CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21402// CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1403// CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31404// CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1405// CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 01406// CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 01407// CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 01408// CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 01409// CK23-DAG: store ptr [[VAL:%.+]], ptr [[BP1]]1410// CK23-DAG: store ptr [[VAL]], ptr [[P1]]1411// CK23-DAG: [[VAL]] = load ptr, ptr [[ADDR:%.+]],1412// CK23-DAG: [[ADDR]] = load ptr, ptr [[ADDR2:%.+]],1413 1414// CK23: call void [[KERNEL:@.+]](ptr [[VAL]])1415#pragma omp target is_device_ptr(tr) defaultmap(none \1416 : pointer)1417 {1418 ++tr;1419 }1420 1421// CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1422// CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21423// CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1424// CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31425// CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1426// CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 01427// CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 01428// CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 01429// CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 01430// CK23-DAG: store ptr [[VAL:%.+]], ptr [[BP1]]1431// CK23-DAG: store ptr [[VAL]], ptr [[P1]]1432// CK23-DAG: [[VAL]] = load ptr, ptr [[ADDR:%.+]],1433// CK23-DAG: [[ADDR]] = load ptr, ptr [[ADDR2:%.+]],1434 1435// CK23: call void [[KERNEL:@.+]](ptr [[VAL]])1436#pragma omp target is_device_ptr(tr, lr) defaultmap(none \1437 : pointer)1438 {1439 ++tr;1440 }1441 1442// CK23-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1443// CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21444// CK23-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1445// CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31446// CK23-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1447// CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 01448// CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 01449// CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 01450// CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 01451// CK23-DAG: store ptr [[VAL:%.+]], ptr [[BP1]]1452// CK23-DAG: store ptr [[VAL]], ptr [[P1]]1453// CK23-DAG: [[VAL]] = load ptr, ptr [[ADDR:%.+]],1454// CK23-DAG: [[ADDR]] = load ptr, ptr [[ADDR2:%.+]],1455 1456// CK23-DAG: [[_BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 11457// CK23-DAG: [[_P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 11458// CK23-DAG: store ptr [[_VAL:%.+]], ptr [[_BP1]]1459// CK23-DAG: store ptr [[_VAL]], ptr [[_P1]]1460// CK23-DAG: [[_VAL]] = load ptr, ptr [[_ADDR:%.+]],1461// CK23-DAG: [[_ADDR]] = load ptr, ptr [[_ADDR2:%.+]],1462 1463// CK23: call void [[KERNEL:@.+]](ptr [[VAL]], ptr [[_VAL]])1464#pragma omp target is_device_ptr(tr, lr) defaultmap(none \1465 : pointer)1466 {1467 ++tr,++lr;1468 }1469}1470 1471void bar(float *&a, int *&b) {1472 foo<int>(a,b);1473}1474 1475#endif1476///==========================================================================///1477// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK24 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-641478// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK24 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s1479// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-641480// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK24 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-321481// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK24 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s1482// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-321483 1484// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK24 -verify -Wno-vla -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 %s1485// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK24 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s1486// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s1487// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK24 -verify -Wno-vla -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 %s1488// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK24 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s1489// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s1490// SIMD-ONLY18-NOT: {{__kmpc|__tgt}}1491#ifdef CK241492 1493// CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 01494// CK24: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]1495// CK24: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 1059]1496 1497// CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 01498// CK24: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4]1499// CK24: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 1063]1500 1501// CK24-LABEL: explicit_maps_single{{.*}}(1502void explicit_maps_single (int ii){1503 // Map of a scalar.1504 int a = ii;1505 1506// Close.1507// Region 001508// CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1509// CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21510// CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1511// CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31512// CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1513// CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1514// CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1515 1516// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01517// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01518// CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1519// CK24-DAG: store ptr [[VAR0]], ptr [[P0]]1520 1521// CK24: call void [[CALL00:@.+]](ptr {{[^,]+}})1522#pragma omp target map(close, tofrom \1523 : a) defaultmap(none \1524 : scalar)1525 {1526 a++;1527 }1528 1529// Always Close.1530// Region 011531// CK24-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1532// CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21533// CK24-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1534// CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31535// CK24-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1536// CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1537// CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1538 1539// CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01540// CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01541// CK24-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1542// CK24-DAG: store ptr [[VAR0]], ptr [[P0]]1543 1544// CK24: call void [[CALL01:@.+]](ptr {{[^,]+}})1545#pragma omp target map(always close tofrom \1546 : a) defaultmap(none \1547 : scalar)1548 {1549 a++;1550 }1551}1552// CK24: define {{.+}}[[CALL00]]1553// CK24: define {{.+}}[[CALL01]]1554 1555#endif1556///==========================================================================///1557// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK25 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK25 --check-prefix CK25-641558// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK25 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s1559// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK25 --check-prefix CK25-641560// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK25 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK25 --check-prefix CK25-321561// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK25 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s1562// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK25 --check-prefix CK25-321563 1564// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK25 -verify -Wno-vla -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 %s1565// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK25 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s1566// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s1567// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK25 -verify -Wno-vla -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 %s1568// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK25 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s1569// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s1570// SIMD-ONLY18-NOT: {{__kmpc|__tgt}}1571#ifdef CK251572 1573extern int x;1574#pragma omp declare target to(x)1575 1576// CK25-LABEL: @.__omp_offloading_{{.*}}declare_target_to{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 01577 1578void declare_target_to()1579{1580 // CK25: [[C:%.+]] = load i32, ptr @x1581 // CK25: [[INC:%.+]] = add nsw i32 [[C]], 11582 // CK25: store i32 [[INC]], ptr @x1583 // CK25: ret void1584 #pragma omp target defaultmap(none : scalar)1585 {1586 x++;1587 }1588}1589 1590#endif1591///==========================================================================///1592// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK26 -verify -Wno-vla -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK26 --check-prefix CK26-641593// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK26 -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s1594// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK26 --check-prefix CK26-641595// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK26 -verify -Wno-vla -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK26 --check-prefix CK26-321596// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK26 -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s1597// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK26 --check-prefix CK26-321598 1599// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK26 -verify -Wno-vla -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 %s1600// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK26 -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s1601// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s1602// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK26 -verify -Wno-vla -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 %s1603// RUN: %clang_cc1 -no-enable-noundef-analysis -DCK26 -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s1604// RUN: %clang_cc1 -no-enable-noundef-analysis -fopenmp-simd -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify -Wno-vla %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s1605// SIMD-ONLY18-NOT: {{__kmpc|__tgt}}1606#ifdef CK261607 1608// CK26-DAG: [[SIZES:@.+]] = {{.+}}constant [3 x i64] [i64 4, i64 4096, i64 {{.+}}]1609// Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_PTR_AND_OBJ | OMP_MAP_IMPLICIT = 5311610// CK26-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i64] [i64 531, i64 531, i64 531]1611 1612float Vector[1024];1613#pragma omp declare target link(Vector)1614 1615extern int a;1616#pragma omp declare target link(a)1617 1618double *ptr;1619#pragma omp declare target link(ptr)1620 1621void declare_target_link()1622{1623#pragma omp target defaultmap(none:scalar) defaultmap(none:aggregate) defaultmap(none:pointer)1624 {1625 1626 // CK26-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1627 // CK26-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21628 // CK26-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1629 // CK26-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31630 // CK26-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1631 // CK26-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 01632 // CK26-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 01633 // CK26-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 21634 // CK26-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 21635 // CK26-DAG: store ptr @ptr_decl_tgt_ref_ptr, ptr [[BP1]]1636 // CK26-DAG: store ptr @ptr, ptr [[P1]]1637 1638 Vector[a]++;1639 ptr++;1640 }1641}1642 1643#endif1644#endif1645