1 // expected-no-diagnostics 2 #ifndef HEADER 3 #define HEADER 4 5 #ifdef CK1 6 7 ///==========================================================================/// 8 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK1 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK1 9 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK1 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 10 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK1 11 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK1 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK1 12 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK1 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 13 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK1 14 15 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK1 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 16 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 17 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 18 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK1 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 19 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK1 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 20 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 21 // SIMD-ONLY10-NOT: {{__kmpc|__tgt}} 22 23 // CK1-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double_complex{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 24 25 // CK1-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] 26 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544 27 // CK1-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544] 28 29 // CK1-LABEL: implicit_maps_double_complex{{.*}}( 30 void implicit_maps_double_complex (int a){ 31 double _Complex dc = (double)a; 32 33 // CK1-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 34 // CK1-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 35 // CK1-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 36 // CK1-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 37 // CK1-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 38 // CK1-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 39 // CK1-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 40 // CK1-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 41 // CK1-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 42 // CK1-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to { double, double }** 43 // CK1-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to { double, double }** 44 // CK1-DAG: store { double, double }* [[PTR:%[^,]+]], { double, double }** [[CBP1]] 45 // CK1-DAG: store { double, double }* [[PTR]], { double, double }** [[CP1]] 46 47 // CK1: call void [[KERNEL:@.+]]({ double, double }* [[PTR]]) 48 #pragma omp target defaultmap(alloc \ 49 : scalar) 50 { 51 dc *= dc; 52 } 53 } 54 55 // CK1: define internal void [[KERNEL]]({ double, double }* {{.*}}[[ARG:%.+]]) 56 // CK1: [[ADDR:%.+]] = alloca { double, double }*, 57 // CK1: store { double, double }* [[ARG]], { double, double }** [[ADDR]], 58 // CK1: [[REF:%.+]] = load { double, double }*, { double, double }** [[ADDR]], 59 // CK1: {{.+}} = getelementptr inbounds { double, double }, { double, double }* [[REF]], i32 0, i32 0 60 #endif 61 ///==========================================================================/// 62 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK2 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK2 63 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK2 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 64 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK2 65 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK2 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK2 66 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK2 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 67 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK2 68 69 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK2 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 70 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK2 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 71 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 72 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK2 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 73 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK2 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 74 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 75 // SIMD-ONLY10-NOT: {{__kmpc|__tgt}} 76 #ifdef CK2 77 78 // CK2-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double_complex{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 79 80 // CK2-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] 81 // Map types: OMP_MAP_TO | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 545 82 // CK2-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 545] 83 84 // CK2-LABEL: implicit_maps_double_complex{{.*}}( 85 void implicit_maps_double_complex (int a){ 86 double _Complex dc = (double)a; 87 88 // CK2-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 89 // CK2-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 90 // CK2-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 91 // CK2-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 92 // CK2-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 93 // CK2-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 94 // CK2-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 95 // CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 96 // CK2-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 97 // CK2-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to { double, double }** 98 // CK2-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to { double, double }** 99 // CK2-DAG: store { double, double }* [[PTR:%[^,]+]], { double, double }** [[CBP1]] 100 // CK2-DAG: store { double, double }* [[PTR]], { double, double }** [[CP1]] 101 102 // CK2: call void [[KERNEL:@.+]]({ double, double }* [[PTR]]) 103 #pragma omp target defaultmap(to \ 104 : scalar) 105 { 106 dc *= dc; 107 } 108 } 109 110 // CK2: define internal void [[KERNEL]]({ double, double }* {{.*}}[[ARG:%.+]]) 111 // CK2: [[ADDR:%.+]] = alloca { double, double }*, 112 // CK2: store { double, double }* [[ARG]], { double, double }** [[ADDR]], 113 // CK2: [[REF:%.+]] = load { double, double }*, { double, double }** [[ADDR]], 114 // CK2: {{.+}} = getelementptr inbounds { double, double }, { double, double }* [[REF]], i32 0, i32 0 115 #endif 116 ///==========================================================================/// 117 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK3 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3 118 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK3 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 119 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3 120 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK3 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3 121 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK3 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 122 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK3 123 124 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK3 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 125 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK3 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 126 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 127 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK3 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 128 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK3 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 129 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY10 %s 130 // SIMD-ONLY10-NOT: {{__kmpc|__tgt}} 131 #ifdef CK3 132 133 // CK3-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double_complex{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 134 135 // CK3-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] 136 // Map types: OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 546 137 // CK3-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 546] 138 139 // CK3-LABEL: implicit_maps_double_complex{{.*}}( 140 void implicit_maps_double_complex (int a){ 141 double _Complex dc = (double)a; 142 143 // CK3-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 144 // CK3-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 145 // CK3-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 146 // CK3-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 147 // CK3-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 148 // CK3-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 149 // CK3-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 150 // CK3-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 151 // CK3-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 152 // CK3-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to { double, double }** 153 // CK3-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to { double, double }** 154 // CK3-DAG: store { double, double }* [[PTR:%[^,]+]], { double, double }** [[CBP1]] 155 // CK3-DAG: store { double, double }* [[PTR]], { double, double }** [[CP1]] 156 157 // CK3: call void [[KERNEL:@.+]]({ double, double }* [[PTR]]) 158 #pragma omp target defaultmap(from \ 159 : scalar) 160 { 161 dc *= dc; 162 } 163 } 164 165 // CK3: define internal void [[KERNEL]]({ double, double }* {{.*}}[[ARG:%.+]]) 166 // CK3: [[ADDR:%.+]] = alloca { double, double }*, 167 // CK3: store { double, double }* [[ARG]], { double, double }** [[ADDR]], 168 // CK3: [[REF:%.+]] = load { double, double }*, { double, double }** [[ADDR]], 169 // CK3: {{.+}} = getelementptr inbounds { double, double }, { double, double }* [[REF]], i32 0, i32 0 170 #endif 171 ///==========================================================================/// 172 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK4 -verify -fopenmp -fopenmp-version=50 -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-64 173 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK4 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 174 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4 --check-prefix CK4-64 175 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK4 -verify -fopenmp -fopenmp-version=50 -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-32 176 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK4 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 177 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK4 --check-prefix CK4-32 178 179 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK4 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 180 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK4 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 181 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s 182 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK4 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 183 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK4 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 184 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s 185 // SIMD-ONLY6-NOT: {{__kmpc|__tgt}} 186 #ifdef CK4 187 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 0 192 193 // CK4-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8] 194 // Map types: OMP_MAP_PRIVATE_VAL | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800 195 // 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 = 673 197 // CK4-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 673] 198 199 // CK4-LABEL: implicit_maps_double{{.*}}( 200 void implicit_maps_double (int a){ 201 double d = (double)a; 202 203 // CK4-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 204 // CK4-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 205 // CK4-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 206 // CK4-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 207 // CK4-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 208 // CK4-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 209 // CK4-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 210 // CK4-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 211 // CK4-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 212 213 // CK4-64-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[sz:64|32]]* 214 // CK4-64-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[sz]]* 215 // CK4-64-DAG: store i[[sz]] [[VAL:%[^,]+]], i[[sz]]* [[CBP1]] 216 // CK4-64-DAG: store i[[sz]] [[VAL]], i[[sz]]* [[CP1]] 217 // CK4-64-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], 218 // CK4-64-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to double* 219 // CK4-64-64-DAG: store double {{.+}}, double* [[CADDR]], 220 221 // CK4-32-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double** 222 // CK4-32-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double** 223 // CK4-32-DAG: store double* [[DECL:%[^,]+]], double** [[CBP1]] 224 // CK4-32-DAG: store double* [[DECL]], double** [[CP1]] 225 226 // CK4-64: call void [[KERNEL:@.+]](i[[sz]] [[VAL]]) 227 // CK4-32: call void [[KERNEL:@.+]](double* [[DECL]]) 228 #pragma omp target defaultmap(firstprivate \ 229 : scalar) 230 { 231 d += 1.0; 232 } 233 } 234 235 // CK4-64: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]]) 236 // CK4-64: [[ADDR:%.+]] = alloca i[[sz]], 237 // CK4-64: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]], 238 // CK4-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to double* 239 // CK4-64: {{.+}} = load double, double* [[CADDR]], 240 241 // CK4-32: define internal void [[KERNEL]](double* {{.+}}[[ARG:%.+]]) 242 // CK4-32: [[ADDR:%.+]] = alloca double*, 243 // CK4-32: store double* [[ARG]], double** [[ADDR]], 244 // CK4-32: [[REF:%.+]] = load double*, double** [[ADDR]], 245 // CK4-32: {{.+}} = load double, double* [[REF]], 246 247 #endif 248 ///==========================================================================/// 249 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK5 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK5 250 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK5 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 251 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK5 252 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK5 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK5 253 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK5 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 254 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK5 255 256 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK5 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 257 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK5 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 258 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 259 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK5 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 260 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK5 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 261 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 262 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 263 #ifdef CK5 264 265 // CK5-LABEL: @.__omp_offloading_{{.*}}implicit_maps_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 266 267 // CK5-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] 268 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544 269 // CK5-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544] 270 271 // CK5-LABEL: implicit_maps_array{{.*}}( 272 void implicit_maps_array (int a){ 273 double darr[2] = {(double)a, (double)a}; 274 275 // CK5-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 276 // CK5-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 277 // CK5-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 278 // CK5-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 279 // CK5-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 280 // CK5-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 281 // CK5-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 282 // CK5-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 283 // CK5-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 284 // CK5-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [2 x double]** 285 // CK5-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [2 x double]** 286 // CK5-DAG: store [2 x double]* [[DECL:%[^,]+]], [2 x double]** [[CBP1]] 287 // CK5-DAG: store [2 x double]* [[DECL]], [2 x double]** [[CP1]] 288 289 // CK5: call void [[KERNEL:@.+]]([2 x double]* [[DECL]]) 290 #pragma omp target defaultmap(alloc \ 291 : aggregate) 292 { 293 darr[0] += 1.0; 294 darr[1] += 1.0; 295 } 296 } 297 298 // CK5: define internal void [[KERNEL]]([2 x double]* {{.+}}[[ARG:%.+]]) 299 // CK5: [[ADDR:%.+]] = alloca [2 x double]*, 300 // CK5: store [2 x double]* [[ARG]], [2 x double]** [[ADDR]], 301 // CK5: [[REF:%.+]] = load [2 x double]*, [2 x double]** [[ADDR]], 302 // CK5: {{.+}} = getelementptr inbounds [2 x double], [2 x double]* [[REF]], i{{64|32}} 0, i{{64|32}} 0 303 #endif 304 ///==========================================================================/// 305 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK6 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK6 306 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK6 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 307 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK6 308 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK6 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK6 309 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK6 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 310 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK6 311 312 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK6 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 313 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK6 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 314 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 315 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK6 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 316 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK6 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 317 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 318 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 319 #ifdef CK6 320 321 // CK6-LABEL: @.__omp_offloading_{{.*}}implicit_maps_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 322 323 // CK6-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] 324 // Map types: OMP_MAP_TO | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 545 325 // CK6-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 545] 326 327 // CK6-LABEL: implicit_maps_array{{.*}}( 328 void implicit_maps_array (int a){ 329 double darr[2] = {(double)a, (double)a}; 330 331 // CK6-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 332 // CK6-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 333 // CK6-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 334 // CK6-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 335 // CK6-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 336 // CK6-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 337 // CK6-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 338 // CK6-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 339 // CK6-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 340 // CK6-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [2 x double]** 341 // CK6-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [2 x double]** 342 // CK6-DAG: store [2 x double]* [[DECL:%[^,]+]], [2 x double]** [[CBP1]] 343 // CK6-DAG: store [2 x double]* [[DECL]], [2 x double]** [[CP1]] 344 345 // CK6: call void [[KERNEL:@.+]]([2 x double]* [[DECL]]) 346 #pragma omp target defaultmap(to) 347 { 348 darr[0] += 1.0; 349 darr[1] += 1.0; 350 } 351 } 352 353 // CK6: define internal void [[KERNEL]]([2 x double]* {{.+}}[[ARG:%.+]]) 354 // CK6: [[ADDR:%.+]] = alloca [2 x double]*, 355 // CK6: store [2 x double]* [[ARG]], [2 x double]** [[ADDR]], 356 // CK6: [[REF:%.+]] = load [2 x double]*, [2 x double]** [[ADDR]], 357 // CK6: {{.+}} = getelementptr inbounds [2 x double], [2 x double]* [[REF]], i{{64|32}} 0, i{{64|32}} 0 358 #endif 359 ///==========================================================================/// 360 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK7 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK7 361 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK7 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 362 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK7 363 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK7 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK7 364 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK7 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 365 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK7 366 367 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK7 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 368 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK7 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 369 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 370 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK7 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 371 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK7 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 372 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 373 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 374 #ifdef CK7 375 376 // CK7-LABEL: @.__omp_offloading_{{.*}}implicit_maps_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 377 378 // CK7-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] 379 // Map types: OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 546 380 // CK7-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 546] 381 382 // CK7-LABEL: implicit_maps_array{{.*}}( 383 void implicit_maps_array (int a){ 384 double darr[2] = {(double)a, (double)a}; 385 386 // CK7-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 387 // CK7-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 388 // CK7-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 389 // CK7-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 390 // CK7-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 391 // CK7-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 392 // CK7-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 393 // CK7-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 394 // CK7-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 395 // CK7-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [2 x double]** 396 // CK7-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [2 x double]** 397 // CK7-DAG: store [2 x double]* [[DECL:%[^,]+]], [2 x double]** [[CBP1]] 398 // CK7-DAG: store [2 x double]* [[DECL]], [2 x double]** [[CP1]] 399 400 // CK7: call void [[KERNEL:@.+]]([2 x double]* [[DECL]]) 401 #pragma omp target defaultmap(from \ 402 : aggregate) 403 { 404 darr[0] += 1.0; 405 darr[1] += 1.0; 406 } 407 } 408 409 // CK7: define internal void [[KERNEL]]([2 x double]* {{.+}}[[ARG:%.+]]) 410 // CK7: [[ADDR:%.+]] = alloca [2 x double]*, 411 // CK7: store [2 x double]* [[ARG]], [2 x double]** [[ADDR]], 412 // CK7: [[REF:%.+]] = load [2 x double]*, [2 x double]** [[ADDR]], 413 // CK7: {{.+}} = getelementptr inbounds [2 x double], [2 x double]* [[REF]], i{{64|32}} 0, i{{64|32}} 0 414 #endif 415 ///==========================================================================/// 416 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK8 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK8 417 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK8 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 418 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK8 419 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK8 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK8 420 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK8 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 421 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK8 422 423 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK8 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 424 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK8 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 425 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 426 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK8 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 427 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK8 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 428 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 429 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 430 #ifdef CK8 431 432 // CK8-LABEL: @.__omp_offloading_{{.*}}implicit_maps_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 433 434 // CK8-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] 435 // Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547 436 // CK8-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 547] 437 438 // CK8-LABEL: implicit_maps_array{{.*}}( 439 void implicit_maps_array (int a){ 440 double darr[2] = {(double)a, (double)a}; 441 442 // CK8-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 443 // CK8-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 444 // CK8-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 445 // CK8-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 446 // CK8-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 447 // CK8-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 448 // CK8-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 449 // CK8-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 450 // CK8-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 451 // CK8-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [2 x double]** 452 // CK8-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [2 x double]** 453 // CK8-DAG: store [2 x double]* [[DECL:%[^,]+]], [2 x double]** [[CBP1]] 454 // CK8-DAG: store [2 x double]* [[DECL]], [2 x double]** [[CP1]] 455 456 // CK8: call void [[KERNEL:@.+]]([2 x double]* [[DECL]]) 457 #pragma omp target defaultmap(tofrom) 458 { 459 darr[0] += 1.0; 460 darr[1] += 1.0; 461 } 462 } 463 464 // CK8: define internal void [[KERNEL]]([2 x double]* {{.+}}[[ARG:%.+]]) 465 // CK8: [[ADDR:%.+]] = alloca [2 x double]*, 466 // CK8: store [2 x double]* [[ARG]], [2 x double]** [[ADDR]], 467 // CK8: [[REF:%.+]] = load [2 x double]*, [2 x double]** [[ADDR]], 468 // CK8: {{.+}} = getelementptr inbounds [2 x double], [2 x double]* [[REF]], i{{64|32}} 0, i{{64|32}} 0 469 #endif 470 471 ///==========================================================================/// 472 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK9 -verify -fopenmp -fopenmp-version=50 -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-64 473 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK9 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 474 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK9 --check-prefix CK9-64 475 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK9 -verify -fopenmp -fopenmp-version=50 -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-32 476 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK9 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 477 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK9 --check-prefix CK9-32 478 479 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK9 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 480 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK9 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 481 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY26 %s 482 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK9 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 483 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK9 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 484 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY26 %s 485 // SIMD-ONLY26-NOT: {{__kmpc|__tgt}} 486 487 #ifdef CK9 488 489 490 // CK9-LABEL: @.__omp_offloading_{{.*}}zero_size_section_and_private_maps{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 491 // CK9: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 40] 492 // CK9: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 673] 493 494 495 // CK9-LABEL: zero_size_section_and_private_maps{{.*}}( 496 void zero_size_section_and_private_maps (int ii){ 497 int pvtArr[10]; 498 499 // Region 09 500 // CK9-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 501 // CK9-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 502 // CK9-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 503 // CK9-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 504 // CK9-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 505 // CK9-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 506 // CK9-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 507 508 // CK9-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 509 // CK9-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 510 // CK9-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [10 x i32]** 511 // CK9-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [10 x i32]** 512 // CK9-DAG: store [10 x i32]* [[VAR0:%.+]], [10 x i32]** [[CBP0]] 513 // CK9-DAG: store [10 x i32]* [[VAR0]], [10 x i32]** [[CP0]] 514 515 // CK9: call void [[CALL09:@.+]]([10 x i32]* {{[^,]+}}) 516 #pragma omp target defaultmap(firstprivate \ 517 : aggregate) 518 { 519 pvtArr[5]++; 520 } 521 522 } 523 524 525 #endif 526 ///==========================================================================/// 527 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK10 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK10 528 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK10 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 529 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK10 530 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK10 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK10 531 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK10 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 532 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK10 533 534 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK10 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 535 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK10 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 536 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 537 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK10 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 538 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK10 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 539 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 540 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 541 #ifdef CK10 542 543 544 // CK10-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 545 546 // CK10: [[SIZE:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}] 547 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544 548 // CK10: [[MTYPE:@.+]] = private {{.*}}constant [1 x i64] [i64 544] 549 550 // CK10-LABEL: explicit_maps_single{{.*}}( 551 void explicit_maps_single (){ 552 int *pa; 553 554 // CK10-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 555 // CK10-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 556 // CK10-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 557 // CK10-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 558 // CK10-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 559 // CK10-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 560 // CK10-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 561 562 // CK10-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 563 // CK10-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 564 // CK10-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** 565 // CK10-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32*** 566 // CK10-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]] 567 // CK10-DAG: store i32** [[VAR0]], i32*** [[CP0]] 568 569 // CK10: call void [[CALL:@.+]](i32** {{[^,]+}}) 570 #pragma omp target defaultmap(alloc \ 571 : pointer) 572 { 573 pa[50]++; 574 } 575 } 576 577 // CK10: define {{.+}}[[CALL]] 578 #endif 579 ///==========================================================================/// 580 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK11 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK11 581 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK11 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 582 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK11 583 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK11 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK11 584 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK11 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 585 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK11 586 587 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK11 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 588 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK11 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 589 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 590 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK11 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 591 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK11 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 592 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 593 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 594 #ifdef CK11 595 596 597 // CK11-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 598 599 // CK11: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}] 600 // Map types: OMP_MAP_TO | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 545 601 // CK11: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 545] 602 603 // CK11-LABEL: explicit_maps_single{{.*}}( 604 void explicit_maps_single (){ 605 int *pa; 606 607 // CK11-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 608 // CK11-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 609 // CK11-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 610 // CK11-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 611 // CK11-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 612 // CK11-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 613 // CK11-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 614 615 // CK11-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 616 // CK11-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 617 // CK11-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** 618 // CK11-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32*** 619 // CK11-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]] 620 // CK11-DAG: store i32** [[VAR0]], i32*** [[CP0]] 621 622 // CK11: call void [[CALL09:@.+]](i32** {{[^,]+}}) 623 #pragma omp target defaultmap(to \ 624 : pointer) 625 { 626 pa[50]++; 627 } 628 } 629 630 // CK11: define {{.+}}[[CALL09]] 631 #endif 632 ///==========================================================================/// 633 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK12 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK12 634 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK12 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 635 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK12 636 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK12 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK12 637 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK12 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 638 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK12 639 640 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK12 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 641 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK12 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 642 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 643 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK12 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 644 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK12 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 645 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 646 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 647 #ifdef CK12 648 649 650 // CK12-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 651 652 // CK12: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}] 653 // Map types: OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 546 654 // CK12: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 546] 655 656 // CK12-LABEL: explicit_maps_single{{.*}}( 657 void explicit_maps_single (){ 658 int *pa; 659 660 // CK12-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 661 // CK12-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 662 // CK12-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 663 // CK12-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 664 // CK12-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 665 // CK12-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 666 // CK12-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 667 668 // CK12-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 669 // CK12-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 670 // CK12-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** 671 // CK12-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32*** 672 // CK12-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]] 673 // CK12-DAG: store i32** [[VAR0]], i32*** [[CP0]] 674 675 // CK12: call void [[CALL09:@.+]](i32** {{[^,]+}}) 676 #pragma omp target defaultmap(from \ 677 : pointer) 678 { 679 pa[50]++; 680 } 681 } 682 683 // CK12: define {{.+}}[[CALL09]] 684 #endif 685 ///==========================================================================/// 686 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK13 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK13 687 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK13 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 688 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK13 689 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK13 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK13 690 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK13 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 691 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK13 692 693 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK13 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 694 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK13 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 695 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 696 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK13 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 697 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK13 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 698 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 699 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 700 #ifdef CK13 701 702 703 // CK13-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 704 705 // CK13: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}] 706 // Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547 707 // CK13: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 547] 708 709 // CK13-LABEL: explicit_maps_single{{.*}}( 710 void explicit_maps_single (){ 711 int *pa; 712 713 // CK13-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 714 // CK13-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 715 // CK13-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 716 // CK13-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 717 // CK13-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 718 // CK13-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 719 // CK13-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 720 721 // CK13-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 722 // CK13-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 723 // CK13-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** 724 // CK13-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32*** 725 // CK13-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]] 726 // CK13-DAG: store i32** [[VAR0]], i32*** [[CP0]] 727 728 // CK13: call void [[CALL09:@.+]](i32** {{[^,]+}}) 729 #pragma omp target defaultmap(tofrom \ 730 : pointer) 731 { 732 pa[50]++; 733 } 734 } 735 736 // CK13: define {{.+}}[[CALL09]] 737 #endif 738 ///==========================================================================/// 739 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK14 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK14 740 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK14 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 741 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK14 742 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK14 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK14 743 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK14 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 744 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK14 745 746 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK14 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 747 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK14 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 748 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 749 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK14 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 750 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK14 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 751 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY8 %s 752 // SIMD-ONLY8-NOT: {{__kmpc|__tgt}} 753 #ifdef CK14 754 755 756 // CK14-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 757 758 // CK14: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer 759 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544 760 // CK14: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 544] 761 762 // CK14-LABEL: explicit_maps_single{{.*}}( 763 void explicit_maps_single (){ 764 int *pa; 765 766 // CK14-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 767 // CK14-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 768 // CK14-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 769 // CK14-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 770 // CK14-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 771 // CK14-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 772 // CK14-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 773 774 // CK14-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 775 // CK14-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 776 // CK14-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 777 // CK14-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 778 // CK14-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]] 779 // CK14-DAG: store i32* [[VAR0]], i32** [[CP0]] 780 781 // CK14: call void [[CALL09:@.+]](i32* {{[^,]+}}) 782 #pragma omp target defaultmap(firstprivate \ 783 : pointer) 784 { 785 pa[50]++; 786 } 787 } 788 789 // CK14: define {{.+}}[[CALL09]] 790 #endif 791 ///==========================================================================/// 792 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK15 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK15 793 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK15 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 794 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK15 795 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK15 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK15 796 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK15 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 797 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK15 798 799 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK15 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 800 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK15 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 801 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY12 %s 802 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK15 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 803 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK15 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 804 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY12 %s 805 // SIMD-ONLY12-NOT: {{__kmpc|__tgt}} 806 #ifdef CK15 807 808 // CK15-LABEL: @.__omp_offloading_{{.*}}implicit_maps_variable_length_array{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 809 810 // CK15-DAG: [[SIZES:@.+]] = {{.+}}constant [3 x i64] [i64 {{4|8}}, i64 {{4|8}}, i64 0] 811 // Map types: 812 // - OMP_MAP_LITERAL + OMP_MAP_TARGET_PARAM + OMP_MAP_IMPLICIT = 800 (vla size) 813 // - OMP_MAP_LITERAL + OMP_MAP_TARGET_PARAM + OMP_MAP_IMPLICIT = 800 (vla size) 814 // - OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544 815 // CK15-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i64] [i64 800, i64 800, i64 544] 816 817 // CK15-LABEL: implicit_maps_variable_length_array{{.*}}( 818 void implicit_maps_variable_length_array (int a){ 819 double vla[2][a]; 820 821 // CK15-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 822 // CK15-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 823 // CK15-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 824 // CK15-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 825 // CK15-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 826 // CK15-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 827 // CK15-DAG: store i64* [[SGEP:%.+]], i64** [[SARG]] 828 // CK15-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 829 // CK15-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 830 // CK15-DAG: [[SGEP]] = getelementptr inbounds {{.+}}[[SS:%[^,]+]], i32 0, i32 0 831 832 // CK15-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 833 // CK15-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 834 // CK15-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i[[sz:64|32]]* 835 // CK15-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i[[sz]]* 836 // CK15-DAG: store i[[sz]] 2, i[[sz]]* [[CBP0]] 837 // CK15-DAG: store i[[sz]] 2, i[[sz]]* [[CP0]] 838 839 // CK15-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1 840 // CK15-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1 841 // CK15-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[sz]]* 842 // CK15-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[sz]]* 843 // CK15-DAG: store i[[sz]] [[VAL:%.+]], i[[sz]]* [[CBP1]] 844 // CK15-DAG: store i[[sz]] [[VAL]], i[[sz]]* [[CP1]] 845 846 // CK15-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2 847 // CK15-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2 848 // CK15-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[SS]], i32 0, i32 2 849 // CK15-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double** 850 // CK15-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to double** 851 // CK15-DAG: store double* [[DECL:%.+]], double** [[CBP2]] 852 // CK15-DAG: store double* [[DECL]], double** [[CP2]] 853 // CK15-DAG: store i64 [[VALS2:%.+]], i64* [[S2]], 854 // CK15-DAG: [[VALS2]] = {{mul nuw i64 %.+, 8|sext i32 %.+ to i64}} 855 856 // CK15: call void [[KERNEL:@.+]](i[[sz]] {{.+}}, i[[sz]] {{.+}}, double* [[DECL]]) 857 #pragma omp target defaultmap(alloc \ 858 : aggregate) 859 { 860 vla[1][3] += 1.0; 861 } 862 } 863 864 // CK15: define internal void [[KERNEL]](i[[sz]] [[VLA0:%.+]], i[[sz]] [[VLA1:%.+]], double* {{.*}}[[ARG:%.+]]) 865 // CK15: [[ADDR0:%.+]] = alloca i[[sz]], 866 // CK15: [[ADDR1:%.+]] = alloca i[[sz]], 867 // CK15: [[ADDR2:%.+]] = alloca double*, 868 // CK15: store i[[sz]] [[VLA0]], i[[sz]]* [[ADDR0]], 869 // CK15: store i[[sz]] [[VLA1]], i[[sz]]* [[ADDR1]], 870 // CK15: store double* [[ARG]], double** [[ADDR2]], 871 // CK15: {{.+}} = load i[[sz]], i[[sz]]* [[ADDR0]], 872 // CK15: {{.+}} = load i[[sz]], i[[sz]]* [[ADDR1]], 873 // CK15: [[REF:%.+]] = load double*, double** [[ADDR2]], 874 // CK15: {{.+}} = getelementptr inbounds double, double* [[REF]], i[[sz]] %{{.+}} 875 #endif 876 ///==========================================================================/// 877 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK16 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK16 878 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK16 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 879 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK16 880 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK16 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK16 881 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK16 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 882 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK16 883 884 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK16 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 885 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK16 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 886 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 887 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK16 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 888 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK16 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 889 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 890 // SIMD-ONLY16-NOT: {{__kmpc|__tgt}} 891 #ifdef CK16 892 893 // CK16-DAG: [[ST:%.+]] = type { i32, double } 894 // CK16-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 895 // CK16-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}] 896 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544 897 // CK16-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544] 898 899 class SSS { 900 public: 901 int a; 902 double b; 903 }; 904 905 // CK16-LABEL: implicit_maps_struct{{.*}}( 906 void implicit_maps_struct (int a){ 907 SSS s = {a, (double)a}; 908 909 // CK16-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 910 // CK16-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 911 // CK16-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 912 // CK16-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 913 // CK16-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 914 // CK16-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 915 // CK16-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 916 // CK16-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 917 // CK16-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 918 // CK16-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]** 919 // CK16-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [[ST]]** 920 // CK16-DAG: store [[ST]]* [[DECL:%.+]], [[ST]]** [[CBP1]] 921 // CK16-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CP1]] 922 923 // CK16: call void [[KERNEL:@.+]]([[ST]]* [[DECL]]) 924 #pragma omp target defaultmap(alloc \ 925 : aggregate) 926 { 927 s.a += 1; 928 s.b += 1.0; 929 } 930 } 931 932 // CK16: define internal void [[KERNEL]]([[ST]]* {{.+}}[[ARG:%.+]]) 933 // CK16: [[ADDR:%.+]] = alloca [[ST]]*, 934 // CK16: store [[ST]]* [[ARG]], [[ST]]** [[ADDR]], 935 // CK16: [[REF:%.+]] = load [[ST]]*, [[ST]]** [[ADDR]], 936 // CK16: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF]], i32 0, i32 0 937 #endif 938 ///==========================================================================/// 939 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK17 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK17 940 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK17 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 941 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK17 942 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK17 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK17 943 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK17 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 944 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK17 945 946 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK17 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 947 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK17 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 948 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 949 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK17 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 950 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK17 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 951 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 952 // SIMD-ONLY16-NOT: {{__kmpc|__tgt}} 953 #ifdef CK17 954 955 // CK17-DAG: [[ST:%.+]] = type { i32, double } 956 // CK17-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 957 // CK17-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}] 958 // Map types: OMP_MAP_TO | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 545 959 // CK17-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 545] 960 961 class SSS { 962 public: 963 int a; 964 double b; 965 }; 966 967 // CK17-LABEL: implicit_maps_struct{{.*}}( 968 void implicit_maps_struct (int a){ 969 SSS s = {a, (double)a}; 970 971 // CK17-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 972 // CK17-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 973 // CK17-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 974 // CK17-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 975 // CK17-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 976 // CK17-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 977 // CK17-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 978 // CK17-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 979 // CK17-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 980 // CK17-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]** 981 // CK17-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [[ST]]** 982 // CK17-DAG: store [[ST]]* [[DECL:%.+]], [[ST]]** [[CBP1]] 983 // CK17-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CP1]] 984 985 // CK17: call void [[KERNEL:@.+]]([[ST]]* [[DECL]]) 986 #pragma omp target defaultmap(to \ 987 : aggregate) 988 { 989 s.a += 1; 990 s.b += 1.0; 991 } 992 } 993 994 // CK17: define internal void [[KERNEL]]([[ST]]* {{.+}}[[ARG:%.+]]) 995 // CK17: [[ADDR:%.+]] = alloca [[ST]]*, 996 // CK17: store [[ST]]* [[ARG]], [[ST]]** [[ADDR]], 997 // CK17: [[REF:%.+]] = load [[ST]]*, [[ST]]** [[ADDR]], 998 // CK17: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF]], i32 0, i32 0 999 #endif 1000 ///==========================================================================/// 1001 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK18 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK18 1002 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK18 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1003 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK18 1004 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK18 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK18 1005 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK18 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1006 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK18 1007 1008 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK18 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 1009 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK18 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1010 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 1011 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK18 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 1012 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK18 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1013 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 1014 // SIMD-ONLY16-NOT: {{__kmpc|__tgt}} 1015 #ifdef CK18 1016 1017 // CK18-DAG: [[ST:%.+]] = type { i32, double } 1018 // CK18-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 1019 // CK18-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}] 1020 // Map types: OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 546 1021 // CK18-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 546] 1022 1023 class SSS { 1024 public: 1025 int a; 1026 double b; 1027 }; 1028 1029 // CK18-LABEL: implicit_maps_struct{{.*}}( 1030 void implicit_maps_struct (int a){ 1031 SSS s = {a, (double)a}; 1032 1033 // CK18-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1034 // CK18-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1035 // CK18-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1036 // CK18-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1037 // CK18-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1038 // CK18-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1039 // CK18-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1040 // CK18-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1041 // CK18-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1042 // CK18-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]** 1043 // CK18-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [[ST]]** 1044 // CK18-DAG: store [[ST]]* [[DECL:%.+]], [[ST]]** [[CBP1]] 1045 // CK18-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CP1]] 1046 1047 // CK18: call void [[KERNEL:@.+]]([[ST]]* [[DECL]]) 1048 #pragma omp target defaultmap(from \ 1049 : aggregate) 1050 { 1051 s.a += 1; 1052 s.b += 1.0; 1053 } 1054 } 1055 1056 // CK18: define internal void [[KERNEL]]([[ST]]* {{.+}}[[ARG:%.+]]) 1057 // CK18: [[ADDR:%.+]] = alloca [[ST]]*, 1058 // CK18: store [[ST]]* [[ARG]], [[ST]]** [[ADDR]], 1059 // CK18: [[REF:%.+]] = load [[ST]]*, [[ST]]** [[ADDR]], 1060 // CK18: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF]], i32 0, i32 0 1061 #endif 1062 ///==========================================================================/// 1063 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK19 -verify -fopenmp -fopenmp-version=50 -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 1064 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK19 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1065 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK19 1066 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK19 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK19 1067 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK19 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1068 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK19 1069 1070 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK19 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 1071 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK19 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1072 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 1073 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK19 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 1074 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK19 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1075 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 1076 // SIMD-ONLY16-NOT: {{__kmpc|__tgt}} 1077 #ifdef CK19 1078 1079 // CK19-DAG: [[ST:%.+]] = type { i32, double } 1080 // CK19-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 1081 // CK19-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}] 1082 // Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547 1083 // CK19-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 547] 1084 1085 class SSS { 1086 public: 1087 int a; 1088 double b; 1089 }; 1090 1091 // CK19-LABEL: implicit_maps_struct{{.*}}( 1092 void implicit_maps_struct (int a){ 1093 SSS s = {a, (double)a}; 1094 1095 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1096 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1097 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1098 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1099 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1100 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1101 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1102 // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1103 // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1104 // CK19-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]** 1105 // CK19-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [[ST]]** 1106 // CK19-DAG: store [[ST]]* [[DECL:%.+]], [[ST]]** [[CBP1]] 1107 // CK19-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CP1]] 1108 1109 // CK19: call void [[KERNEL:@.+]]([[ST]]* [[DECL]]) 1110 #pragma omp target defaultmap(tofrom \ 1111 : aggregate) 1112 { 1113 s.a += 1; 1114 s.b += 1.0; 1115 } 1116 } 1117 1118 // CK19: define internal void [[KERNEL]]([[ST]]* {{.+}}[[ARG:%.+]]) 1119 // CK19: [[ADDR:%.+]] = alloca [[ST]]*, 1120 // CK19: store [[ST]]* [[ARG]], [[ST]]** [[ADDR]], 1121 // CK19: [[REF:%.+]] = load [[ST]]*, [[ST]]** [[ADDR]], 1122 // CK19: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF]], i32 0, i32 0 1123 #endif 1124 ///==========================================================================/// 1125 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK20 -verify -fopenmp -fopenmp-version=50 -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-64 1126 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK20 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1127 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK20 --check-prefix CK20-64 1128 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK20 -verify -fopenmp -fopenmp-version=50 -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-32 1129 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK20 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1130 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK20 --check-prefix CK20-32 1131 1132 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK20 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 1133 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK20 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1134 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s 1135 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK20 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 1136 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK20 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1137 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY6 %s 1138 // SIMD-ONLY6-NOT: {{__kmpc|__tgt}} 1139 #ifdef CK20 1140 1141 // For a 32-bit targets, the value doesn't fit the size of the pointer, 1142 // therefore it is passed by reference with a map 'to' specification. 1143 1144 // CK20-LABEL: @.__omp_offloading_{{.*}}implicit_maps_double{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 1145 1146 // CK20-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8] 1147 // Map types: OMP_MAP_LITERAL | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 800 1148 // CK20-64-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 800] 1149 // Map types: OMP_MAP_TO | OMP_MAP_PRIVATE | OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 673 1150 // CK20-32-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 673] 1151 1152 // CK20-LABEL: implicit_maps_double{{.*}}( 1153 void implicit_maps_double (int a){ 1154 double d = (double)a; 1155 1156 // CK20-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1157 // CK20-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1158 // CK20-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1159 // CK20-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1160 // CK20-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1161 // CK20-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1162 // CK20-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1163 // CK20-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1164 // CK20-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1165 1166 // CK20-64-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[sz:64|32]]* 1167 // CK20-64-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[sz]]* 1168 // CK20-64-DAG: store i[[sz]] [[VAL:%[^,]+]], i[[sz]]* [[CBP1]] 1169 // CK20-64-DAG: store i[[sz]] [[VAL]], i[[sz]]* [[CP1]] 1170 // CK20-64-DAG: [[VAL]] = load i[[sz]], i[[sz]]* [[ADDR:%.+]], 1171 // CK20-64-64-DAG: [[CADDR:%.+]] = bitcast i[[sz]]* [[ADDR]] to double* 1172 // CK20-64-64-DAG: store double {{.+}}, double* [[CADDR]], 1173 1174 // CK20-32-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double** 1175 // CK20-32-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double** 1176 // CK20-32-DAG: store double* [[DECL:%[^,]+]], double** [[CBP1]] 1177 // CK20-32-DAG: store double* [[DECL]], double** [[CP1]] 1178 1179 // CK20-64: call void [[KERNEL:@.+]](i[[sz]] [[VAL]]) 1180 // CK20-32: call void [[KERNEL:@.+]](double* [[DECL]]) 1181 #pragma omp target defaultmap(default \ 1182 : scalar) 1183 { 1184 d += 1.0; 1185 } 1186 } 1187 1188 // CK20-64: define internal void [[KERNEL]](i[[sz]] [[ARG:%.+]]) 1189 // CK20-64: [[ADDR:%.+]] = alloca i[[sz]], 1190 // CK20-64: store i[[sz]] [[ARG]], i[[sz]]* [[ADDR]], 1191 // CK20-64: [[CADDR:%.+]] = bitcast i64* [[ADDR]] to double* 1192 // CK20-64: {{.+}} = load double, double* [[CADDR]], 1193 1194 // CK20-32: define internal void [[KERNEL]](double* {{.+}}[[ARG:%.+]]) 1195 // CK20-32: [[ADDR:%.+]] = alloca double*, 1196 // CK20-32: store double* [[ARG]], double** [[ADDR]], 1197 // CK20-32: [[REF:%.+]] = load double*, double** [[ADDR]], 1198 // CK20-32: {{.+}} = load double, double* [[REF]], 1199 1200 #endif 1201 ///==========================================================================/// 1202 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK21 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK21 1203 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK21 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1204 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK21 1205 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK21 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK21 1206 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK21 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1207 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK21 1208 1209 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK21 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 1210 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK21 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1211 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 1212 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK21 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 1213 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK21 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1214 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY16 %s 1215 // SIMD-ONLY16-NOT: {{__kmpc|__tgt}} 1216 #ifdef CK21 1217 1218 // CK21-DAG: [[ST:%.+]] = type { i32, double } 1219 // CK21-LABEL: @.__omp_offloading_{{.*}}implicit_maps_struct{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 1220 // CK21-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 {{16|12}}] 1221 // Map types: OMP_MAP_TO + OMP_MAP_FROM + OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 547 1222 // CK21-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 547] 1223 1224 class SSS { 1225 public: 1226 int a; 1227 double b; 1228 }; 1229 1230 // CK21-LABEL: implicit_maps_struct{{.*}}( 1231 void implicit_maps_struct (int a){ 1232 SSS s = {a, (double)a}; 1233 1234 // CK21-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1235 // CK21-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1236 // CK21-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1237 // CK21-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1238 // CK21-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1239 // CK21-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1240 // CK21-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1241 // CK21-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1242 // CK21-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1243 // CK21-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to [[ST]]** 1244 // CK21-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to [[ST]]** 1245 // CK21-DAG: store [[ST]]* [[DECL:%.+]], [[ST]]** [[CBP1]] 1246 // CK21-DAG: store [[ST]]* [[DECL]], [[ST]]** [[CP1]] 1247 1248 // CK21: call void [[KERNEL:@.+]]([[ST]]* [[DECL]]) 1249 #pragma omp target defaultmap(default \ 1250 : aggregate) 1251 { 1252 s.a += 1; 1253 s.b += 1.0; 1254 } 1255 } 1256 1257 // CK21: define internal void [[KERNEL]]([[ST]]* {{.+}}[[ARG:%.+]]) 1258 // CK21: [[ADDR:%.+]] = alloca [[ST]]*, 1259 // CK21: store [[ST]]* [[ARG]], [[ST]]** [[ADDR]], 1260 // CK21: [[REF:%.+]] = load [[ST]]*, [[ST]]** [[ADDR]], 1261 // CK21: {{.+}} = getelementptr inbounds [[ST]], [[ST]]* [[REF]], i32 0, i32 0 1262 #endif 1263 ///==========================================================================/// 1264 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK22 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK22 1265 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK22 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1266 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK22 1267 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK22 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK22 1268 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK22 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1269 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK22 1270 1271 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK22 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 1272 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK22 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1273 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY9 %s 1274 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK22 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 1275 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK22 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1276 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY9 %s 1277 // SIMD-ONLY9-NOT: {{__kmpc|__tgt}} 1278 #ifdef CK22 1279 1280 // CK22-LABEL: @.__omp_offloading_{{.*}}implicit_maps_pointer{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 1281 1282 // CK22-DAG: [[SIZES:@.+]] = {{.+}}constant [1 x i64] zeroinitializer 1283 // Map types: OMP_MAP_TARGET_PARAM | OMP_MAP_IMPLICIT = 544 1284 // CK22-DAG: [[TYPES:@.+]] = {{.+}}constant [1 x i64] [i64 544] 1285 1286 // CK22-LABEL: implicit_maps_pointer{{.*}}( 1287 void implicit_maps_pointer (){ 1288 double *ddyn; 1289 1290 // CK22-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1291 // CK22-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1292 // CK22-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1293 // CK22-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1294 // CK22-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1295 // CK22-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1296 // CK22-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1297 // CK22-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1298 // CK22-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1299 // CK22-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double** 1300 // CK22-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double** 1301 // CK22-DAG: store double* [[PTR:%[^,]+]], double** [[CBP1]] 1302 // CK22-DAG: store double* [[PTR]], double** [[CP1]] 1303 1304 // CK22: call void [[KERNEL:@.+]](double* [[PTR]]) 1305 #pragma omp target defaultmap(default \ 1306 : pointer) 1307 { 1308 ddyn[0] += 1.0; 1309 ddyn[1] += 1.0; 1310 } 1311 } 1312 1313 // CK22: define internal void [[KERNEL]](double* {{.*}}[[ARG:%.+]]) 1314 // CK22: [[ADDR:%.+]] = alloca double*, 1315 // CK22: store double* [[ARG]], double** [[ADDR]], 1316 // CK22: [[REF:%.+]] = load double*, double** [[ADDR]], 1317 // CK22: {{.+}} = getelementptr inbounds double, double* [[REF]], i{{64|32}} 0 1318 1319 #endif 1320 ///==========================================================================/// 1321 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK23 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-64 1322 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK23 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1323 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-64 1324 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK23 -verify -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-32 1325 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK23 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1326 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck %s --check-prefix CK23 --check-prefix CK23-32 1327 1328 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK23 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s 1329 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK23 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1330 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s 1331 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK23 -verify -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -emit-llvm %s -o - | FileCheck --check-prefix SIMD-ONLY0 %s 1332 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK23 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1333 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck --check-prefix SIMD-ONLY0 %s 1334 // SIMD-ONLY0-NOT: {{__kmpc|__tgt}} 1335 #ifdef CK23 1336 1337 double *g; 1338 1339 // CK23: @g ={{.*}} global double* 1340 // CK23: [[SIZES00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} {{8|4}}] 1341 // CK23: [[TYPES00:@.+]] = {{.+}}constant [1 x i64] [i64 288] 1342 1343 // CK23: [[SIZES01:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}] 1344 // CK23: [[TYPES01:@.+]] = {{.+}}constant [1 x i64] [i64 288] 1345 1346 // CK23: [[SIZES02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}] 1347 // CK23: [[TYPES02:@.+]] = {{.+}}constant [1 x i64] [i64 288] 1348 1349 // CK23: [[SIZES03:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}] 1350 // CK23: [[TYPES03:@.+]] = {{.+}}constant [1 x i64] [i64 288] 1351 1352 // CK23: [[SIZES04:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}] 1353 // CK23: [[TYPES04:@.+]] = {{.+}}constant [1 x i64] [i64 288] 1354 1355 // CK23: [[SIZES05:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] {{8|4}}] 1356 // CK23: [[TYPES05:@.+]] = {{.+}}constant [1 x i64] [i64 288] 1357 1358 // CK23: [[SIZES06:@.+]] = {{.+}}constant [2 x i[[sz]]] [i[[sz]] {{8|4}}, i[[sz]] {{8|4}}] 1359 // CK23: [[TYPES06:@.+]] = {{.+}}constant [2 x i64] [i64 288, i64 288] 1360 1361 // CK23-LABEL: @_Z3foo{{.*}}( 1362 template<typename T> 1363 void foo(float *&lr, T *&tr) { 1364 float *l; 1365 T *t; 1366 1367 // CK23-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1368 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1369 // CK23-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1370 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1371 // CK23-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1372 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1373 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1374 // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1375 // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1376 // CK23-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double** 1377 // CK23-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double** 1378 // CK23-DAG: store double* [[VAL:%.+]], double** [[CBP1]] 1379 // CK23-DAG: store double* [[VAL]], double** [[CP1]] 1380 // CK23-DAG: [[VAL]] = load double*, double** [[ADDR:@g]], 1381 1382 // CK23: call void [[KERNEL:@.+]](double* [[VAL]]) 1383 #pragma omp target is_device_ptr(g) defaultmap(none \ 1384 : pointer) 1385 { 1386 ++g; 1387 } 1388 1389 // CK23-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1390 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1391 // CK23-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1392 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1393 // CK23-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1394 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1395 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1396 // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1397 // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1398 // CK23-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to float** 1399 // CK23-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to float** 1400 // CK23-DAG: store float* [[VAL:%.+]], float** [[CBP1]] 1401 // CK23-DAG: store float* [[VAL]], float** [[CP1]] 1402 // CK23-DAG: [[VAL]] = load float*, float** [[ADDR:%.+]], 1403 1404 // CK23: call void [[KERNEL:@.+]](float* [[VAL]]) 1405 #pragma omp target is_device_ptr(l) defaultmap(none \ 1406 : pointer) 1407 { 1408 ++l; 1409 } 1410 1411 // CK23-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1412 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1413 // CK23-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1414 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1415 // CK23-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1416 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1417 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1418 // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1419 // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1420 // CK23-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32** 1421 // CK23-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** 1422 // CK23-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]] 1423 // CK23-DAG: store i32* [[VAL]], i32** [[CP1]] 1424 // CK23-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]], 1425 1426 // CK23: call void [[KERNEL:@.+]](i32* [[VAL]]) 1427 #pragma omp target is_device_ptr(t) defaultmap(none \ 1428 : pointer) 1429 { 1430 ++t; 1431 } 1432 1433 // CK23-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1434 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1435 // CK23-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1436 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1437 // CK23-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1438 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1439 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1440 // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1441 // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1442 // CK23-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to float** 1443 // CK23-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to float** 1444 // CK23-DAG: store float* [[VAL:%.+]], float** [[CBP1]] 1445 // CK23-DAG: store float* [[VAL]], float** [[CP1]] 1446 // CK23-DAG: [[VAL]] = load float*, float** [[ADDR:%.+]], 1447 // CK23-DAG: [[ADDR]] = load float**, float*** [[ADDR2:%.+]], 1448 1449 // CK23: call void [[KERNEL:@.+]](float* [[VAL]]) 1450 #pragma omp target is_device_ptr(lr) defaultmap(none \ 1451 : pointer) 1452 { 1453 ++lr; 1454 } 1455 1456 // CK23-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1457 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1458 // CK23-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1459 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1460 // CK23-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1461 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1462 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1463 // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1464 // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1465 // CK23-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32** 1466 // CK23-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** 1467 // CK23-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]] 1468 // CK23-DAG: store i32* [[VAL]], i32** [[CP1]] 1469 // CK23-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]], 1470 // CK23-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]], 1471 1472 // CK23: call void [[KERNEL:@.+]](i32* [[VAL]]) 1473 #pragma omp target is_device_ptr(tr) defaultmap(none \ 1474 : pointer) 1475 { 1476 ++tr; 1477 } 1478 1479 // CK23-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1480 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1481 // CK23-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1482 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1483 // CK23-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1484 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1485 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1486 // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1487 // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1488 // CK23-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32** 1489 // CK23-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** 1490 // CK23-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]] 1491 // CK23-DAG: store i32* [[VAL]], i32** [[CP1]] 1492 // CK23-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]], 1493 // CK23-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]], 1494 1495 // CK23: call void [[KERNEL:@.+]](i32* [[VAL]]) 1496 #pragma omp target is_device_ptr(tr, lr) defaultmap(none \ 1497 : pointer) 1498 { 1499 ++tr; 1500 } 1501 1502 // CK23-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1503 // CK23-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1504 // CK23-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1505 // CK23-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1506 // CK23-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1507 // CK23-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1508 // CK23-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1509 // CK23-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0 1510 // CK23-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0 1511 // CK23-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32** 1512 // CK23-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** 1513 // CK23-DAG: store i32* [[VAL:%.+]], i32** [[CBP1]] 1514 // CK23-DAG: store i32* [[VAL]], i32** [[CP1]] 1515 // CK23-DAG: [[VAL]] = load i32*, i32** [[ADDR:%.+]], 1516 // CK23-DAG: [[ADDR]] = load i32**, i32*** [[ADDR2:%.+]], 1517 1518 // CK23-DAG: [[_BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 1 1519 // CK23-DAG: [[_P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 1 1520 // CK23-DAG: [[_CBP1:%.+]] = bitcast i8** [[_BP1]] to float** 1521 // CK23-DAG: [[_CP1:%.+]] = bitcast i8** [[_P1]] to float** 1522 // CK23-DAG: store float* [[_VAL:%.+]], float** [[_CBP1]] 1523 // CK23-DAG: store float* [[_VAL]], float** [[_CP1]] 1524 // CK23-DAG: [[_VAL]] = load float*, float** [[_ADDR:%.+]], 1525 // CK23-DAG: [[_ADDR]] = load float**, float*** [[_ADDR2:%.+]], 1526 1527 // CK23: call void [[KERNEL:@.+]](i32* [[VAL]], float* [[_VAL]]) 1528 #pragma omp target is_device_ptr(tr, lr) defaultmap(none \ 1529 : pointer) 1530 { 1531 ++tr,++lr; 1532 } 1533 } 1534 1535 void bar(float *&a, int *&b) { 1536 foo<int>(a,b); 1537 } 1538 1539 #endif 1540 ///==========================================================================/// 1541 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK24 -verify -fopenmp -fopenmp-version=50 -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-64 1542 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK24 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1543 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-64 1544 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK24 -verify -fopenmp -fopenmp-version=50 -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-32 1545 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK24 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1546 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK24 --check-prefix CK24-32 1547 1548 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK24 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 1549 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK24 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1550 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1551 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK24 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 1552 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK24 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1553 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1554 // SIMD-ONLY18-NOT: {{__kmpc|__tgt}} 1555 #ifdef CK24 1556 1557 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 1558 // CK24: [[SIZE00:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] 1559 // CK24: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 1059] 1560 1561 // CK24-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 1562 // CK24: [[SIZE01:@.+]] = private {{.*}}constant [1 x i[[Z:64|32]]] [i[[Z:64|32]] 4] 1563 // CK24: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 1063] 1564 1565 // CK24-LABEL: explicit_maps_single{{.*}}( 1566 void explicit_maps_single (int ii){ 1567 // Map of a scalar. 1568 int a = ii; 1569 1570 // Close. 1571 // Region 00 1572 // CK24-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1573 // CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1574 // CK24-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1575 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1576 // CK24-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1577 // CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1578 // CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1579 1580 // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1581 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1582 // CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 1583 // CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 1584 // CK24-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]] 1585 // CK24-DAG: store i32* [[VAR0]], i32** [[CP0]] 1586 1587 // CK24: call void [[CALL00:@.+]](i32* {{[^,]+}}) 1588 #pragma omp target map(close, tofrom \ 1589 : a) defaultmap(none \ 1590 : scalar) 1591 { 1592 a++; 1593 } 1594 1595 // Always Close. 1596 // Region 01 1597 // CK24-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1598 // CK24-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1599 // CK24-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1600 // CK24-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1601 // CK24-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1602 // CK24-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1603 // CK24-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1604 1605 // CK24-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1606 // CK24-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1607 // CK24-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 1608 // CK24-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 1609 // CK24-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]] 1610 // CK24-DAG: store i32* [[VAR0]], i32** [[CP0]] 1611 1612 // CK24: call void [[CALL01:@.+]](i32* {{[^,]+}}) 1613 #pragma omp target map(always close tofrom \ 1614 : a) defaultmap(none \ 1615 : scalar) 1616 { 1617 a++; 1618 } 1619 } 1620 // CK24: define {{.+}}[[CALL00]] 1621 // CK24: define {{.+}}[[CALL01]] 1622 1623 #endif 1624 ///==========================================================================/// 1625 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK25 -verify -fopenmp -fopenmp-version=50 -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-64 1626 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK25 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1627 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK25 --check-prefix CK25-64 1628 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK25 -verify -fopenmp -fopenmp-version=50 -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-32 1629 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK25 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1630 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK25 --check-prefix CK25-32 1631 1632 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK25 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 1633 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK25 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1634 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1635 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK25 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 1636 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK25 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1637 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1638 // SIMD-ONLY18-NOT: {{__kmpc|__tgt}} 1639 #ifdef CK25 1640 1641 extern int x; 1642 #pragma omp declare target to(x) 1643 1644 // CK25-LABEL: @.__omp_offloading_{{.*}}declare_target_to{{.*}}_l{{[0-9]+}}.region_id = weak{{.*}} constant i8 0 1645 1646 void declare_target_to() 1647 { 1648 // CK25: [[C:%.+]] = load i32, i32* @x 1649 // CK25: [[INC:%.+]] = add nsw i32 [[C]], 1 1650 // CK25: store i32 [[INC]], i32* @x 1651 // CK25: ret void 1652 #pragma omp target defaultmap(none : scalar) 1653 { 1654 x++; 1655 } 1656 } 1657 1658 #endif 1659 ///==========================================================================/// 1660 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK26 -verify -fopenmp -fopenmp-version=50 -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-64 1661 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK26 -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1662 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK26 --check-prefix CK26-64 1663 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK26 -verify -fopenmp -fopenmp-version=50 -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-32 1664 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK26 -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1665 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap %s --check-prefix CK26 --check-prefix CK26-32 1666 1667 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK26 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 1668 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK26 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -std=c++11 -triple powerpc64le-unknown-unknown -emit-pch -o %t %s 1669 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1670 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK26 -verify -fopenmp-simd -fopenmp-version=50 -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 %s 1671 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -DCK26 -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -std=c++11 -triple i386-unknown-unknown -emit-pch -o %t %s 1672 // RUN: %clang_cc1 -no-opaque-pointers -no-enable-noundef-analysis -fopenmp-simd -fopenmp-version=50 -fopenmp-targets=i386-pc-linux-gnu -x c++ -triple i386-unknown-unknown -std=c++11 -include-pch %t -verify %s -emit-llvm -o - | FileCheck -allow-deprecated-dag-overlap --check-prefix SIMD-ONLY18 %s 1673 // SIMD-ONLY18-NOT: {{__kmpc|__tgt}} 1674 #ifdef CK26 1675 1676 // CK26-DAG: [[SIZES:@.+]] = {{.+}}constant [3 x i64] [i64 4, i64 4096, i64 {{.+}}] 1677 // Map types: OMP_MAP_TO | OMP_MAP_FROM | OMP_MAP_PTR_AND_OBJ | OMP_MAP_IMPLICIT = 531 1678 // CK26-DAG: [[TYPES:@.+]] = {{.+}}constant [3 x i64] [i64 531, i64 531, i64 531] 1679 1680 float Vector[1024]; 1681 #pragma omp declare target link(Vector) 1682 1683 extern int a; 1684 #pragma omp declare target link(a) 1685 1686 double *ptr; 1687 #pragma omp declare target link(ptr) 1688 1689 void declare_target_link() 1690 { 1691 #pragma omp target defaultmap(none:scalar) defaultmap(none:aggregate) defaultmap(none:pointer) 1692 { 1693 1694 // CK26-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1695 // CK26-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1696 // CK26-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1697 // CK26-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1698 // CK26-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1699 // CK26-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0 1700 // CK26-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0 1701 // CK26-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 2 1702 // CK26-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 2 1703 // CK26-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double**** 1704 // CK26-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double*** 1705 // CK26-DAG: store double*** @ptr_decl_tgt_ref_ptr, double**** [[CBP1]] 1706 // CK26-DAG: store double** @ptr, double*** [[CP1]] 1707 1708 Vector[a]++; 1709 ptr++; 1710 } 1711 } 1712 1713 #endif 1714 #endif 1715