1// expected-no-diagnostics 2#ifndef HEADER_INC 3#define HEADER_INC 4 5// This file is regex-heavy and takes a long time to execute the test. To speed 6// testing up, test execution is split over multiple fimes. The RUN commands are 7// in the corresponding .cpp files now. Do not add them here. 8 9// SIMD-ONLY18-NOT: {{__kmpc|__tgt}} 10#ifdef CK19 11 12// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 13// CK19: [[SIZE00:@.+]] = private {{.*}}constant [1 x i64] [i64 4] 14// CK19-USE: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] [i64 32] 15// CK19-NOUSE: [[MTYPE00:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer 16 17// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 18// CK19: [[SIZE00n:@.+]] = private {{.*}}constant [1 x i64] [i64 4] 19// CK19-USE: [[MTYPE00n:@.+]] = private {{.*}}constant [1 x i64] [i64 32] 20// CK19-NOUSE: [[MTYPE00n:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer 21 22// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 23// CK19: [[SIZE01:@.+]] = private {{.*}}constant [1 x i64] [i64 400] 24// CK19-USE: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 33] 25// CK19-NOUSE: [[MTYPE01:@.+]] = private {{.*}}constant [1 x i64] [i64 1] 26 27// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 28// CK19: [[SIZE02:@.+]] = private {{.*}}constant [1 x i64] [i64 240] 29// CK19-USE: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i64] [i64 34] 30// CK19-NOUSE: [[MTYPE02:@.+]] = private {{.*}}constant [1 x i64] [i64 2] 31 32// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 33// CK19: [[SIZE03:@.+]] = private {{.*}}constant [1 x i64] [i64 240] 34// CK19-USE: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 35] 35// CK19-NOUSE: [[MTYPE03:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 36 37// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 38// CK19: [[SIZE04:@.+]] = private {{.*}}constant [1 x i64] [i64 400] 39// CK19-USE: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i64] [i64 32] 40// CK19-NOUSE: [[MTYPE04:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer 41 42// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 43// CK19: [[SIZE05:@.+]] = private {{.*}}constant [1 x i64] [i64 4] 44// CK19-USE: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i64] [i64 33] 45// CK19-NOUSE: [[MTYPE05:@.+]] = private {{.*}}constant [1 x i64] [i64 1] 46 47// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 48// CK19-USE: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i64] [i64 35] 49// CK19-NOUSE: [[MTYPE06:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 50 51// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 52// CK19-USE: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i64] [i64 32] 53// CK19-NOUSE: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer 54 55// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 56// CK19: [[SIZE08:@.+]] = private {{.*}}constant [1 x i64] [i64 4] 57// CK19-USE: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i64] [i64 35] 58// CK19-NOUSE: [[MTYPE08:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 59 60// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 61// CK19: [[SIZE09:@.+]] = private {{.*}}constant [1 x i64] [i64 {{8|4}}] 62// CK19-USE: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 34] 63// CK19-NOUSE: [[MTYPE09:@.+]] = private {{.*}}constant [1 x i64] [i64 2] 64 65// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 66// CK19: [[SIZE10:@.+]] = private {{.*}}constant [1 x i64] [i64 240] 67// CK19-USE: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i64] [i64 35] 68// CK19-NOUSE: [[MTYPE10:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 69 70// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 71// CK19: [[SIZE11:@.+]] = private {{.*}}constant [1 x i64] [i64 240] 72// CK19-USE: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i64] [i64 32] 73// CK19-NOUSE: [[MTYPE11:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer 74 75// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 76// CK19: [[SIZE12:@.+]] = private {{.*}}constant [1 x i64] [i64 4] 77// CK19-USE: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i64] [i64 33] 78// CK19-NOUSE: [[MTYPE12:@.+]] = private {{.*}}constant [1 x i64] [i64 1] 79 80// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 81// CK19-USE: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i64] [i64 32] 82// CK19-NOUSE: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer 83 84// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 85// CK19-USE: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 33] 86// CK19-NOUSE: [[MTYPE14:@.+]] = private {{.*}}constant [1 x i64] [i64 1] 87 88// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 89// CK19: [[SIZE15:@.+]] = private {{.*}}constant [1 x i64] [i64 4] 90// CK19-USE: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i64] [i64 34] 91// CK19-NOUSE: [[MTYPE15:@.+]] = private {{.*}}constant [1 x i64] [i64 2] 92 93// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 94// CK19-USE: [[SIZE16:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 0] 95// CK19-USE: [[MTYPE16:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 33] 96// CK19-NOUSE: [[MTYPE16:@.+]] = private {{.*}}constant [1 x i64] [i64 1] 97 98// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 99// CK19-USE: [[SIZE17:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 240] 100// CK19-USE: [[MTYPE17:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 34] 101// CK19-NOUSE: [[SIZE17:@.+]] = private {{.*}}constant [1 x i64] [i64 240] 102// CK19-NOUSE: [[MTYPE17:@.+]] = private {{.*}}constant [1 x i64] [i64 2] 103 104// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 105// CK19-USE: [[SIZE18:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 240] 106// CK19-USE: [[MTYPE18:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 35] 107// CK19-NOUSE: [[SIZE18:@.+]] = private {{.*}}constant [1 x i64] [i64 240] 108// CK19-NOUSE: [[MTYPE18:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 109 110// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 111// CK19-USE: [[SIZE19:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 0] 112// CK19-USE: [[MTYPE19:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 32] 113// CK19-NOUSE: [[MTYPE19:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer 114 115// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 116// CK19-USE: [[SIZE20:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 4] 117// CK19-USE: [[MTYPE20:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 33] 118// CK19-NOUSE: [[SIZE20:@.+]] = private {{.*}}constant [1 x i64] [i64 4] 119// CK19-NOUSE: [[MTYPE20:@.+]] = private {{.*}}constant [1 x i64] [i64 1] 120 121// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 122// CK19-USE: [[SIZE21:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 0] 123// CK19-USE: [[MTYPE21:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 35] 124// CK19-NOUSE: [[MTYPE21:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 125 126// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 127// CK19-USE: [[SIZE22:@.+]] = private {{.*}}constant [2 x i64] [i64 {{8|4}}, i64 4] 128// CK19-USE: [[MTYPE22:@.+]] = private {{.*}}constant [2 x i64] [i64 800, i64 35] 129// CK19-NOUSE: [[SIZE22:@.+]] = private {{.*}}constant [1 x i64] [i64 4] 130// CK19-NOUSE: [[MTYPE22:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 131 132// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 133// CK19: [[SIZE23:@.+]] = private {{.*}}constant [1 x i64] [i64 4] 134// CK19-USE: [[MTYPE23:@.+]] = private {{.*}}constant [1 x i64] [i64 39] 135// CK19-NOUSE: [[MTYPE23:@.+]] = private {{.*}}constant [1 x i64] [i64 7] 136 137// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 138// CK19: [[SIZE24:@.+]] = private {{.*}}constant [1 x i64] [i64 480] 139// CK19-USE: [[MTYPE24:@.+]] = private {{.*}}constant [1 x i64] [i64 35] 140// CK19-NOUSE: [[MTYPE24:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 141 142// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 143// CK19: [[SIZE25:@.+]] = private {{.*}}constant [1 x i64] [i64 16] 144// CK19-USE: [[MTYPE25:@.+]] = private {{.*}}constant [1 x i64] [i64 35] 145// CK19-NOUSE: [[MTYPE25:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 146 147// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 148// CK19: [[SIZE26:@.+]] = private {{.*}}constant [1 x i64] [i64 24] 149// CK19-USE: [[MTYPE26:@.+]] = private {{.*}}constant [1 x i64] [i64 35] 150// CK19-NOUSE: [[MTYPE26:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 151 152// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 153// CK19: [[SIZE27:@.+]] = private {{.*}}constant [1 x i64] [i64 4] 154// CK19-USE: [[MTYPE27:@.+]] = private {{.*}}constant [1 x i64] [i64 35] 155// CK19-NOUSE: [[MTYPE27:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 156 157// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 158// CK19: [[SIZE28:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 16] 159// CK19-USE: [[MTYPE28:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16, i64 19] 160// CK19-NOUSE: [[MTYPE28:@.+]] = private {{.*}}constant [3 x i64] [i64 3, i64 16, i64 19] 161 162// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 163// CK19: [[SIZE29:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 4] 164// CK19-USE: [[MTYPE29:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16, i64 19] 165// CK19-NOUSE: [[MTYPE29:@.+]] = private {{.*}}constant [3 x i64] [i64 3, i64 16, i64 19] 166 167// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 168// CK19-USE: [[SIZE30:@.+]] = private {{.*}}constant [4 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 {{8|4}}, i64 0] 169// CK19-USE: [[MTYPE30:@.+]] = private {{.*}}constant [4 x i64] [i64 800, i64 800, i64 800, i64 35] 170// CK19-NOUSE: [[MTYPE30:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 171 172// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 173// CK19-USE: [[SIZE31:@.+]] = private {{.*}}constant [4 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 {{8|4}}, i64 40] 174// CK19-USE: [[MTYPE31:@.+]] = private {{.*}}constant [4 x i64] [i64 800, i64 800, i64 800, i64 35] 175// CK19-NOUSE: [[SIZE31:@.+]] = private {{.*}}constant [1 x i64] [i64 40] 176// CK19-NOUSE: [[MTYPE31:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 177 178// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 179// CK19: [[SIZE32:@.+]] = private {{.*}}constant [1 x i64] [i64 13728] 180// CK19-USE: [[MTYPE32:@.+]] = private {{.*}}constant [1 x i64] [i64 35] 181// CK19-NOUSE: [[MTYPE32:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 182 183// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 184// CK19: [[SIZE33:@.+]] = private {{.*}}constant [1 x i64] [i64 13728] 185// CK19-USE: [[MTYPE33:@.+]] = private {{.*}}constant [1 x i64] [i64 35] 186// CK19-NOUSE: [[MTYPE33:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 187 188// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 189// CK19: [[SIZE34:@.+]] = private {{.*}}constant [1 x i64] [i64 13728] 190// CK19-USE: [[MTYPE34:@.+]] = private {{.*}}constant [1 x i64] [i64 35] 191// CK19-NOUSE: [[MTYPE34:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 192 193// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 194// CK19-USE: [[MTYPE35:@.+]] = private {{.*}}constant [1 x i64] [i64 35] 195// CK19-NOUSE: [[MTYPE35:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 196 197// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 198// CK19: [[SIZE36:@.+]] = private {{.*}}constant [1 x i64] [i64 208] 199// CK19-USE: [[MTYPE36:@.+]] = private {{.*}}constant [1 x i64] [i64 35] 200// CK19-NOUSE: [[MTYPE36:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 201 202// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 203// CK19-USE: [[SIZE37:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 0] 204// CK19-USE: [[MTYPE37:@.+]] = private {{.*}}constant [3 x i64] [i64 800, i64 800, i64 35] 205// CK19-NOUSE: [[MTYPE37:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 206 207// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 208// CK19-USE: [[SIZE38:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 0] 209// CK19-USE: [[MTYPE38:@.+]] = private {{.*}}constant [3 x i64] [i64 800, i64 800, i64 35] 210// CK19-NOUSE: [[MTYPE38:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 211 212// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 213// CK19-USE: [[SIZE39:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 0] 214// CK19-USE: [[MTYPE39:@.+]] = private {{.*}}constant [3 x i64] [i64 800, i64 800, i64 35] 215// CK19-NOUSE: [[MTYPE39:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 216 217// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 218// CK19-USE: [[SIZE40:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 0] 219// CK19-USE: [[MTYPE40:@.+]] = private {{.*}}constant [3 x i64] [i64 800, i64 800, i64 35] 220// CK19-NOUSE: [[MTYPE40:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 221 222// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 223// CK19-USE: [[SIZE41:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 208] 224// CK19-USE: [[MTYPE41:@.+]] = private {{.*}}constant [3 x i64] [i64 800, i64 800, i64 35] 225// CK19-NOUSE: [[SIZE41:@.+]] = private {{.*}}constant [1 x i64] [i64 208] 226// CK19-NOUSE: [[MTYPE41:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 227 228// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 229// CK19: [[SIZE42:@.+]] = private {{.*}}constant [3 x i64] [i64 {{8|4}}, i64 {{8|4}}, i64 104] 230// CK19-USE: [[MTYPE42:@.+]] = private {{.*}}constant [3 x i64] [i64 35, i64 16, i64 19] 231// CK19-NOUSE: [[MTYPE42:@.+]] = private {{.*}}constant [3 x i64] [i64 3, i64 16, i64 19] 232 233// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 234// CK19-USE: [[MTYPE43:@.+]] = private {{.*}}constant [1 x i64] [i64 35] 235// CK19-NOUSE: [[MTYPE43:@.+]] = private {{.*}}constant [1 x i64] [i64 3] 236 237// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0 238// CK19: [[SIZE44:@.+]] = private {{.*}}constant [1 x i64] [i64 320] 239// CK19-USE: [[MTYPE44:@.+]] = private {{.*}}constant [1 x i64] [i64 34] 240// CK19-NOUSE: [[MTYPE44:@.+]] = private {{.*}}constant [1 x i64] [i64 2] 241 242// CK19-LABEL: explicit_maps_single{{.*}}( 243void explicit_maps_single (int ii){ 244 // Map of a scalar. 245 int a = ii; 246 247 // Region 00 248 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 249 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 250 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 251 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 252 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 253 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 254 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 255 256 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 257 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 258 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 259 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 260 // CK19-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]] 261 // CK19-DAG: store i32* [[VAR0]], i32** [[CP0]] 262 263 // CK19-USE: call void [[CALL00:@.+]](i32* {{[^,]+}}) 264 // CK19-NOUSE: call void [[CALL00:@.+]]() 265 #pragma omp target map(alloc:a) 266 { 267#ifdef USE 268 ++a; 269#endif 270 } 271 272 // Map of a scalar in nested region. 273 int b = a; 274 275 // Region 00n 276 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 -1, i32 1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 277 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 278 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 279 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 280 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 281 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 282 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 283 284 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 285 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 286 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 287 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 288 // CK19-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]] 289 // CK19-DAG: store i32* [[VAR0]], i32** [[CP0]] 290 291 // CK19-USE: call void [[CALL00n:@.+]](i32* {{[^,]+}}) 292 // CK19-NOUSE: call void [[CALL00n:@.+]]() 293 #pragma omp target map(alloc:b) 294 #pragma omp parallel 295 { 296#ifdef USE 297 ++b; 298#endif 299 } 300 301 // Map of an array. 302 int arra[100]; 303 304 // Region 01 305 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 306 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 307 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 308 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 309 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 310 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 311 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 312 313 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 314 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 315 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [100 x i32]** 316 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [100 x i32]** 317 // CK19-DAG: store [100 x i32]* [[VAR0:%.+]], [100 x i32]** [[CBP0]] 318 // CK19-DAG: store [100 x i32]* [[VAR0]], [100 x i32]** [[CP0]] 319 320 // CK19-USE: call void [[CALL01:@.+]]([100 x i32]* {{[^,]+}}) 321 // CK19-NOUSE: call void [[CALL01:@.+]]() 322 #pragma omp target map(to:arra) 323 { 324#ifdef USE 325 arra[50]++; 326#endif 327 } 328 329 // Region 02 330 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 331 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 332 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 333 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 334 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 335 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 336 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 337 338 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 339 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 340 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [100 x i32]** 341 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 342 // CK19-DAG: store [100 x i32]* [[VAR0:%.+]], [100 x i32]** [[CBP0]] 343 // CK19-DAG: store i32* [[SEC0:%[^,]+]], i32** [[CP0]] 344 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 20 345 346 // CK19-USE: call void [[CALL02:@.+]]([100 x i32]* {{[^,]+}}) 347 // CK19-NOUSE: call void [[CALL02:@.+]]() 348 #pragma omp target map(from:arra[20:60]) 349 { 350#ifdef USE 351 arra[50]++; 352#endif 353 } 354 355 // Region 03 356 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 357 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 358 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 359 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 360 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 361 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 362 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 363 364 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 365 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 366 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [100 x i32]** 367 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 368 // CK19-DAG: store [100 x i32]* [[VAR0:%.+]], [100 x i32]** [[CBP0]] 369 // CK19-DAG: store i32* [[SEC0:%[^,]+]], i32** [[CP0]] 370 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 0 371 372 // CK19-USE: call void [[CALL03:@.+]]([100 x i32]* {{[^,]+}}) 373 // CK19-NOUSE: call void [[CALL03:@.+]]() 374 #pragma omp target map(tofrom:arra[:60]) 375 { 376#ifdef USE 377 arra[50]++; 378#endif 379 } 380 381 // Region 04 382 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 383 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 384 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 385 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 386 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 387 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 388 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 389 390 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 391 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 392 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [100 x i32]** 393 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 394 // CK19-DAG: store [100 x i32]* [[VAR0:%.+]], [100 x i32]** [[CBP0]] 395 // CK19-DAG: store i32* [[SEC0:%[^,]+]], i32** [[CP0]] 396 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 0 397 398 // CK19-USE: call void [[CALL04:@.+]]([100 x i32]* {{[^,]+}}) 399 // CK19-NOUSE: call void [[CALL04:@.+]]() 400 #pragma omp target map(alloc:arra[:]) 401 { 402#ifdef USE 403 arra[50]++; 404#endif 405 } 406 407 // Region 05 408 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 409 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 410 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 411 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 412 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 413 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 414 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 415 416 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 417 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 418 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [100 x i32]** 419 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 420 // CK19-DAG: store [100 x i32]* [[VAR0:%.+]], [100 x i32]** [[CBP0]] 421 // CK19-DAG: store i32* [[SEC0:%[^,]+]], i32** [[CP0]] 422 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 15 423 424 // CK19-USE: call void [[CALL05:@.+]]([100 x i32]* {{[^,]+}}) 425 // CK19-NOUSE: call void [[CALL05:@.+]]() 426 #pragma omp target map(to:arra[15]) 427 { 428#ifdef USE 429 arra[15]++; 430#endif 431 } 432 433 // Region 06 434 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 435 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 436 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 437 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 438 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 439 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 440 // CK19-DAG: store i64* [[SIZES:%.+]], i64** [[SARG]] 441 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 442 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 443 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 444 445 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 446 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 447 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 448 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [100 x i32]** 449 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 450 // CK19-DAG: store [100 x i32]* [[VAR0:%.+]], [100 x i32]** [[CBP0]] 451 // CK19-DAG: store i32* [[SEC0:%[^,]+]], i32** [[CP0]] 452 // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]] 453 // CK19-DAG: [[CSVAL0]] = {{mul nuw i.+ %.*, 4|sext i32 .+ to i64}} 454 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} %{{.*}} 455 456 // CK19-USE: call void [[CALL06:@.+]]([100 x i32]* {{[^,]+}}) 457 // CK19-NOUSE: call void [[CALL06:@.+]]() 458 #pragma omp target map(tofrom:arra[ii:ii+23]) 459 { 460#ifdef USE 461 arra[50]++; 462#endif 463 } 464 465 // Region 07 466 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 467 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 468 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 469 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 470 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 471 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 472 // CK19-DAG: store i64* [[SIZES:%.+]], i64** [[SARG]] 473 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 474 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 475 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 476 477 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 478 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 479 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 480 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [100 x i32]** 481 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 482 // CK19-DAG: store [100 x i32]* [[VAR0:%.+]], [100 x i32]** [[CBP0]] 483 // CK19-DAG: store i32* [[SEC0:%[^,]+]], i32** [[CP0]] 484 // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]] 485 // CK19-DAG: [[CSVAL0]] = {{mul nuw i.+ %.*, 4|sext i32 .+ to i64}} 486 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 0 487 488 // CK19-USE: call void [[CALL07:@.+]]([100 x i32]* {{[^,]+}}) 489 // CK19-NOUSE: call void [[CALL07:@.+]]() 490 #pragma omp target map(alloc:arra[:ii]) 491 { 492#ifdef USE 493 arra[50]++; 494#endif 495 } 496 497 // Region 08 498 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 499 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 500 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 501 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 502 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 503 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 504 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 505 506 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 507 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 508 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [100 x i32]** 509 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 510 // CK19-DAG: store [100 x i32]* [[VAR0:%.+]], [100 x i32]** [[CBP0]] 511 // CK19-DAG: store i32* [[SEC0:%[^,]+]], i32** [[CP0]] 512 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} %{{.*}} 513 514 // CK19-USE: call void [[CALL08:@.+]]([100 x i32]* {{[^,]+}}) 515 // CK19-NOUSE: call void [[CALL08:@.+]]() 516 #pragma omp target map(tofrom:arra[ii]) 517 { 518#ifdef USE 519 arra[15]++; 520#endif 521 } 522 523 // Map of a pointer. 524 int *pa; 525 526 // Region 09 527 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 528 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 529 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 530 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 531 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 532 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 533 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 534 535 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 536 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 537 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32*** 538 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32*** 539 // CK19-DAG: store i32** [[VAR0:%.+]], i32*** [[CBP0]] 540 // CK19-DAG: store i32** [[VAR0]], i32*** [[CP0]] 541 542 // CK19-USE: call void [[CALL09:@.+]](i32** {{[^,]+}}) 543 // CK19-NOUSE: call void [[CALL09:@.+]]() 544 #pragma omp target map(from:pa) 545 { 546#ifdef USE 547 pa[50]++; 548#endif 549 } 550 551 // Region 10 552 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 553 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 554 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 555 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 556 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 557 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 558 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 559 560 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 561 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 562 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 563 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 564 // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]] 565 // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] 566 // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] 567 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 20 568 // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] 569 570 // CK19-USE: call void [[CALL10:@.+]](i32* {{[^,]+}}) 571 // CK19-NOUSE: call void [[CALL10:@.+]]() 572 #pragma omp target map(tofrom:pa[20:60]) 573 { 574#ifdef USE 575 pa[50]++; 576#endif 577 } 578 579 // Region 11 580 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 581 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 582 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 583 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 584 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 585 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 586 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 587 588 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 589 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 590 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 591 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 592 // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]] 593 // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] 594 // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] 595 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0 596 // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] 597 598 // CK19-USE: call void [[CALL11:@.+]](i32* {{[^,]+}}) 599 // CK19-NOUSE: call void [[CALL11:@.+]]() 600 #pragma omp target map(alloc:pa[:60]) 601 { 602#ifdef USE 603 pa[50]++; 604#endif 605 } 606 607 // Region 12 608 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 609 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 610 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 611 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 612 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 613 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 614 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 615 616 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 617 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 618 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 619 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 620 // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]] 621 // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] 622 // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] 623 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 15 624 // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] 625 626 // CK19-USE: call void [[CALL12:@.+]](i32* {{[^,]+}}) 627 // CK19-NOUSE: call void [[CALL12:@.+]]() 628 #pragma omp target map(to:pa[15]) 629 { 630#ifdef USE 631 pa[15]++; 632#endif 633 } 634 635 // Region 13 636 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 637 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 638 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 639 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 640 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 641 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 642 // CK19-DAG: store i64* [[SIZES:%.+]], i64** [[SARG]] 643 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 644 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 645 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 646 647 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 648 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 649 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 650 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 651 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 652 // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]] 653 // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] 654 // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]] 655 // CK19-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}} 656 // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] 657 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} %{{.*}} 658 // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] 659 660 // CK19-USE: call void [[CALL13:@.+]](i32* {{[^,]+}}) 661 // CK19-NOUSE: call void [[CALL13:@.+]]() 662 #pragma omp target map(alloc:pa[ii-23:ii]) 663 { 664#ifdef USE 665 pa[50]++; 666#endif 667 } 668 669 // Region 14 670 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 671 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 672 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 673 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 674 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 675 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 676 // CK19-DAG: store i64* [[SIZES:%.+]], i64** [[SARG]] 677 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 678 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 679 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 680 681 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 682 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 683 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 684 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 685 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 686 // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]] 687 // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] 688 // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]] 689 // CK19-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}} 690 // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] 691 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} 0 692 // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] 693 694 // CK19-USE: call void [[CALL14:@.+]](i32* {{[^,]+}}) 695 // CK19-NOUSE: call void [[CALL14:@.+]]() 696 #pragma omp target map(to:pa[:ii]) 697 { 698#ifdef USE 699 pa[50]++; 700#endif 701 } 702 703 // Region 15 704 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 705 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 706 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 707 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 708 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 709 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 710 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 711 712 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 713 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 714 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 715 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 716 // CK19-DAG: store i32* [[RVAR0:%.+]], i32** [[CBP0]] 717 // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] 718 // CK19-DAG: [[RVAR0]] = load i32*, i32** [[VAR0:%[^,]+]] 719 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[RVAR00:%.+]], i{{.+}} %{{.*}} 720 // CK19-DAG: [[RVAR00]] = load i32*, i32** [[VAR0]] 721 722 // CK19-USE: call void [[CALL15:@.+]](i32* {{[^,]+}}) 723 // CK19-NOUSE: call void [[CALL15:@.+]]() 724 #pragma omp target map(from:pa[ii+12]) 725 { 726#ifdef USE 727 pa[15]++; 728#endif 729 } 730 731 // Map of a variable-size array. 732 int va[ii]; 733 734 // Region 16 735 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 736 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 737 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 738 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 739 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 740 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 741 // CK19-DAG: store i64* [[SIZES:%.+]], i64** [[SARG]] 742 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 743 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 744 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 745 746 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 747 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 748 // CK19-USE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i[[Z:64|32]]* 749 // CK19-USE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i[[Z]]* 750 // CK19-USE-DAG: store i[[Z]] {{%.+}}, i[[Z]]* [[CBP0]] 751 // CK19-USE-DAG: store i[[Z]] {{%.+}}, i[[Z]]* [[CP0]] 752 753 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 754 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 755 // CK19-USE-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 756 // CK19-USE-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32** 757 // CK19-USE-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** 758 // CK19-USE-DAG: store i32* [[VAR1:%.+]], i32** [[CBP1]] 759 // CK19-USE-DAG: store i32* [[VAR1]], i32** [[CP1]] 760 // CK19-USE-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], i{{.+}}* [[S1]] 761 // CK19-USE-DAG: [[CSVAL1]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}} 762 763 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 764 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 765 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 766 // CK19-NOUSE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 767 // CK19-NOUSE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 768 // CK19-NOUSE-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]] 769 // CK19-NOUSE-DAG: store i32* [[VAR0]], i32** [[CP0]] 770 // CK19-NOUSE-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]] 771 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}} 772 773 // CK19-USE: call void [[CALL16:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) 774 // CK19-NOUSE: call void [[CALL16:@.+]]() 775 #pragma omp target map(to:va) 776 { 777#ifdef USE 778 va[50]++; 779#endif 780 } 781 782 // Region 17 783 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 784 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 785 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 786 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 787 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 788 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 789 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 790 791 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 792 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 793 // CK19-USE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i[[Z]]* 794 // CK19-USE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i[[Z]]* 795 // CK19-USE-DAG: store i[[Z]] {{%.+}}, i[[Z]]* [[CBP0]] 796 // CK19-USE-DAG: store i[[Z]] {{%.+}}, i[[Z]]* [[CP0]] 797 798 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 799 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 800 // CK19-USE-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32** 801 // CK19-USE-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** 802 // CK19-USE-DAG: store i32* [[VAR1:%.+]], i32** [[CBP1]] 803 // CK19-USE-DAG: store i32* [[SEC1:%.+]], i32** [[CP1]] 804 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} 20 805 806 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 807 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 808 // CK19-NOUSE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 809 // CK19-NOUSE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 810 // CK19-NOUSE-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]] 811 // CK19-NOUSE-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] 812 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[VAR0]], i{{.+}} 20 813 814 // CK19-USE: call void [[CALL17:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) 815 // CK19-NOUSE: call void [[CALL17:@.+]]() 816 #pragma omp target map(from:va[20:60]) 817 { 818#ifdef USE 819 va[50]++; 820#endif 821 } 822 823 // Region 18 824 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 825 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 826 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 827 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 828 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 829 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 830 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 831 832 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 833 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 834 // CK19-USE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i[[Z]]* 835 // CK19-USE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i[[Z]]* 836 // CK19-USE-DAG: store i[[Z]] {{%.+}}, i[[Z]]* [[CBP0]] 837 // CK19-USE-DAG: store i[[Z]] {{%.+}}, i[[Z]]* [[CP0]] 838 839 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 840 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 841 // CK19-USE-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32** 842 // CK19-USE-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** 843 // CK19-USE-DAG: store i32* [[VAR1:%.+]], i32** [[CBP1]] 844 // CK19-USE-DAG: store i32* [[SEC1:%.+]], i32** [[CP1]] 845 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} 0 846 847 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 848 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 849 // CK19-NOUSE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 850 // CK19-NOUSE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 851 // CK19-NOUSE-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]] 852 // CK19-NOUSE-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] 853 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[VAR0]], i{{.+}} 0 854 855 // CK19-USE: call void [[CALL18:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) 856 // CK19-NOUSE: call void [[CALL18:@.+]]() 857 #pragma omp target map(tofrom:va[:60]) 858 { 859#ifdef USE 860 va[50]++; 861#endif 862 } 863 864 // Region 19 865 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 866 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 867 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 868 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 869 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 870 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 871 // CK19-DAG: store i64* [[SIZES:%.+]], i64** [[SARG]] 872 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 873 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 874 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 875 876 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 877 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 878 // CK19-USE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i[[Z]]* 879 // CK19-USE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i[[Z]]* 880 // CK19-USE-DAG: store i[[Z]] {{%.+}}, i[[Z]]* [[CBP0]] 881 // CK19-USE-DAG: store i[[Z]] {{%.+}}, i[[Z]]* [[CP0]] 882 883 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 884 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 885 // CK19-USE-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 886 // CK19-USE-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32** 887 // CK19-USE-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** 888 // CK19-USE-DAG: store i32* [[VAR1:%.+]], i32** [[CBP1]] 889 // CK19-USE-DAG: store i32* [[SEC1:%.+]], i32** [[CP1]] 890 // CK19-USE-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], i{{.+}}* [[S1]] 891 // CK19-USE-DAG: [[CSVAL1]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}} 892 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} 0 893 894 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 895 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 896 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 897 // CK19-NOUSE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 898 // CK19-NOUSE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 899 // CK19-NOUSE-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]] 900 // CK19-NOUSE-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] 901 // CK19-NOUSE-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]] 902 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}} 903 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[VAR0]], i{{.+}} 0 904 905 // CK19-USE: call void [[CALL19:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) 906 // CK19-NOUSE: call void [[CALL19:@.+]]() 907 #pragma omp target map(alloc:va[:]) 908 { 909#ifdef USE 910 va[50]++; 911#endif 912 } 913 914 // Region 20 915 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 916 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 917 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 918 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 919 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 920 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 921 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 922 923 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 924 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 925 // CK19-USE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i[[Z]]* 926 // CK19-USE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i[[Z]]* 927 // CK19-USE-DAG: store i[[Z]] {{%.+}}, i[[Z]]* [[CBP0]] 928 // CK19-USE-DAG: store i[[Z]] {{%.+}}, i[[Z]]* [[CP0]] 929 930 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 931 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 932 // CK19-USE-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32** 933 // CK19-USE-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** 934 // CK19-USE-DAG: store i32* [[VAR1:%.+]], i32** [[CBP1]] 935 // CK19-USE-DAG: store i32* [[SEC1:%.+]], i32** [[CP1]] 936 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} 15 937 938 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 939 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 940 // CK19-NOUSE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 941 // CK19-NOUSE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 942 // CK19-NOUSE-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]] 943 // CK19-NOUSE-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] 944 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[VAR0]], i{{.+}} 15 945 946 // CK19-USE: call void [[CALL20:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) 947 // CK19-NOUSE: call void [[CALL20:@.+]]() 948 #pragma omp target map(to:va[15]) 949 { 950#ifdef USE 951 va[15]++; 952#endif 953 } 954 955 // Region 21 956 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 957 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 958 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 959 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 960 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 961 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 962 // CK19-DAG: store i64* [[SIZES:%.+]], i64** [[SARG]] 963 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 964 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 965 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 966 967 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 968 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 969 // CK19-USE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i[[Z]]* 970 // CK19-USE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i[[Z]]* 971 // CK19-USE-DAG: store i[[Z]] {{%.+}}, i[[Z]]* [[CBP0]] 972 // CK19-USE-DAG: store i[[Z]] {{%.+}}, i[[Z]]* [[CP0]] 973 974 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 975 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 976 // CK19-USE-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1 977 // CK19-USE-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32** 978 // CK19-USE-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** 979 // CK19-USE-DAG: store i32* [[VAR1:%.+]], i32** [[CBP1]] 980 // CK19-USE-DAG: store i32* [[SEC1:%.+]], i32** [[CP1]] 981 // CK19-USE-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], i{{.+}}* [[S1]] 982 // CK19-USE-DAG: [[CSVAL1]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}} 983 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} %{{.+}} 984 985 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 986 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 987 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 988 // CK19-NOUSE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 989 // CK19-NOUSE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 990 // CK19-NOUSE-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]] 991 // CK19-NOUSE-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] 992 // CK19-NOUSE-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], i{{.+}}* [[S0]] 993 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}} 994 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[VAR0]], i{{.+}} %{{.+}} 995 996 // CK19-USE: call void [[CALL21:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) 997 // CK19-NOUSE: call void [[CALL21:@.+]]() 998 #pragma omp target map(tofrom:va[ii:ii+23]) 999 { 1000#ifdef USE 1001 va[50]++; 1002#endif 1003 } 1004 1005 // Region 22 1006 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1007 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1008 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1009 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1010 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1011 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1012 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1013 1014 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1015 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1016 // CK19-USE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i[[Z]]* 1017 // CK19-USE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i[[Z]]* 1018 // CK19-USE-DAG: store i[[Z]] {{%.+}}, i[[Z]]* [[CBP0]] 1019 // CK19-USE-DAG: store i[[Z]] {{%.+}}, i[[Z]]* [[CP0]] 1020 1021 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 1022 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 1023 // CK19-USE-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32** 1024 // CK19-USE-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32** 1025 // CK19-USE-DAG: store i32* [[VAR1:%.+]], i32** [[CBP1]] 1026 // CK19-USE-DAG: store i32* [[SEC1:%.+]], i32** [[CP1]] 1027 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}i32* [[VAR1]], i{{.+}} %{{.+}} 1028 1029 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1030 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1031 // CK19-NOUSE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 1032 // CK19-NOUSE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 1033 // CK19-NOUSE-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]] 1034 // CK19-NOUSE-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] 1035 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}i32* [[VAR0]], i{{.+}} %{{.+}} 1036 1037 // CK19-USE: call void [[CALL22:@.+]](i{{.+}} {{[^,]+}}, i32* {{[^,]+}}) 1038 // CK19-NOUSE: call void [[CALL22:@.+]]() 1039 #pragma omp target map(tofrom:va[ii]) 1040 { 1041#ifdef USE 1042 va[15]++; 1043#endif 1044 } 1045 1046 // Always. 1047 // Region 23 1048 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1049 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1050 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1051 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1052 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1053 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1054 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1055 1056 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1057 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1058 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** 1059 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 1060 // CK19-DAG: store i32* [[VAR0:%.+]], i32** [[CBP0]] 1061 // CK19-DAG: store i32* [[VAR0]], i32** [[CP0]] 1062 1063 // CK19-USE: call void [[CALL23:@.+]](i32* {{[^,]+}}) 1064 // CK19-NOUSE: call void [[CALL23:@.+]]() 1065 #pragma omp target map(always, tofrom: a) 1066 { 1067#ifdef USE 1068 a++; 1069#endif 1070 } 1071 1072 // Multidimensional arrays. 1073 int marr[4][5][6]; 1074 int ***mptr; 1075 1076 // Region 24 1077 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1078 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1079 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1080 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1081 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1082 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1083 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1084 1085 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1086 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1087 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [4 x [5 x [6 x i32]]]** 1088 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [4 x [5 x [6 x i32]]]** 1089 // CK19-DAG: store [4 x [5 x [6 x i32]]]* [[VAR0:%.+]], [4 x [5 x [6 x i32]]]** [[CBP0]] 1090 // CK19-DAG: store [4 x [5 x [6 x i32]]]* [[VAR0]], [4 x [5 x [6 x i32]]]** [[CP0]] 1091 1092 // CK19-USE: call void [[CALL24:@.+]]([4 x [5 x [6 x i32]]]* {{[^,]+}}) 1093 // CK19-NOUSE: call void [[CALL24:@.+]]() 1094 #pragma omp target map(tofrom: marr) 1095 { 1096#ifdef USE 1097 marr[1][2][3]++; 1098#endif 1099 } 1100 1101 // Region 25 1102 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1103 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1104 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1105 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1106 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1107 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1108 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1109 1110 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1111 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1112 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [4 x [5 x [6 x i32]]]** 1113 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 1114 // CK19-DAG: store [4 x [5 x [6 x i32]]]* [[VAR0:%.+]], [4 x [5 x [6 x i32]]]** [[CBP0]] 1115 // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] 1116 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[6 x i32]* [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 2 1117 // CK19-DAG: [[SEC00]] = getelementptr {{.*}}[5 x [6 x i32]]* [[SEC000:[^,]+]], i{{.+}} 0, i{{.+}} 2 1118 // CK19-DAG: [[SEC000]] = getelementptr {{.*}}[4 x [5 x [6 x i32]]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 1119 1120 // CK19-USE: call void [[CALL25:@.+]]([4 x [5 x [6 x i32]]]* {{[^,]+}}) 1121 // CK19-NOUSE: call void [[CALL25:@.+]]() 1122 #pragma omp target map(tofrom: marr[1][2][2:4]) 1123 { 1124#ifdef USE 1125 marr[1][2][3]++; 1126#endif 1127 } 1128 1129 // Region 26 1130 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1131 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1132 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1133 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1134 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1135 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1136 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1137 1138 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1139 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1140 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [4 x [5 x [6 x i32]]]** 1141 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 1142 // CK19-DAG: store [4 x [5 x [6 x i32]]]* [[VAR0:%.+]], [4 x [5 x [6 x i32]]]** [[CBP0]] 1143 // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] 1144 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[6 x i32]* [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 0 1145 // CK19-DAG: [[SEC00]] = getelementptr {{.*}}[5 x [6 x i32]]* [[SEC000:[^,]+]], i{{.+}} 0, i{{.+}} 2 1146 // CK19-DAG: [[SEC000]] = getelementptr {{.*}}[4 x [5 x [6 x i32]]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 1147 1148 // CK19-USE: call void [[CALL26:@.+]]([4 x [5 x [6 x i32]]]* {{[^,]+}}) 1149 // CK19-NOUSE: call void [[CALL26:@.+]]() 1150 #pragma omp target map(tofrom: marr[1][2][:]) 1151 { 1152#ifdef USE 1153 marr[1][2][3]++; 1154#endif 1155 } 1156 1157 // Region 27 1158 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1159 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1160 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1161 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1162 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1163 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1164 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1165 1166 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1167 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1168 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [4 x [5 x [6 x i32]]]** 1169 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 1170 // CK19-DAG: store [4 x [5 x [6 x i32]]]* [[VAR0:%.+]], [4 x [5 x [6 x i32]]]** [[CBP0]] 1171 // CK19-DAG: store i32* [[SEC0:%.+]], i32** [[CP0]] 1172 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[6 x i32]* [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 3 1173 // CK19-DAG: [[SEC00]] = getelementptr {{.*}}[5 x [6 x i32]]* [[SEC000:[^,]+]], i{{.+}} 0, i{{.+}} 2 1174 // CK19-DAG: [[SEC000]] = getelementptr {{.*}}[4 x [5 x [6 x i32]]]* [[VAR0]], i{{.+}} 0, i{{.+}} 1 1175 1176 // CK19-USE: call void [[CALL27:@.+]]([4 x [5 x [6 x i32]]]* {{[^,]+}}) 1177 // CK19-NOUSE: call void [[CALL27:@.+]]() 1178 #pragma omp target map(tofrom: marr[1][2][3]) 1179 { 1180#ifdef USE 1181 marr[1][2][3]++; 1182#endif 1183 } 1184 1185 // Region 28 1186 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1187 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1188 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1189 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1190 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1191 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1192 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1193 1194 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1195 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1196 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**** 1197 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**** 1198 // CK19-DAG: store i32*** [[VAR0:%.+]], i32**** [[CBP0]] 1199 // CK19-DAG: store i32*** [[SEC0:%.+]], i32**** [[CP0]] 1200 // CK19-DAG: [[VAR0]] = load i32***, i32**** [[PTR:%[^,]+]], 1201 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32*** [[SEC00:[^,]+]], i{{.+}} 1 1202 // CK19-DAG: [[SEC00]] = load i32***, i32**** [[PTR]], 1203 1204 // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 1205 // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 1206 // CK19-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32**** 1207 // CK19-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32*** 1208 // CK19-DAG: store i32*** [[SEC0]], i32**** [[CBP1]] 1209 // CK19-DAG: store i32** [[SEC1:%.+]], i32*** [[CP1]] 1210 // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32** [[SEC11:[^,]+]], i{{.+}} 2 1211 // CK19-DAG: [[SEC11]] = load i32**, i32*** [[SEC111:%[^,]+]], 1212 // CK19-DAG: [[SEC111]] = getelementptr {{.*}}i32*** [[SEC1111:[^,]+]], i{{.+}} 1 1213 // CK19-DAG: [[SEC1111]] = load i32***, i32**** [[PTR]], 1214 1215 // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 1216 // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 1217 // CK19-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to i32*** 1218 // CK19-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32** 1219 // CK19-DAG: store i32** [[SEC1]], i32*** [[CBP2]] 1220 // CK19-DAG: store i32* [[SEC2:%.+]], i32** [[CP2]] 1221 // CK19-DAG: [[SEC2]] = getelementptr {{.*}}i32* [[SEC22:[^,]+]], i{{.+}} 2 1222 // CK19-DAG: [[SEC22]] = load i32*, i32** [[SEC222:%[^,]+]], 1223 // CK19-DAG: [[SEC222]] = getelementptr {{.*}}i32** [[SEC2222:[^,]+]], i{{.+}} 2 1224 // CK19-DAG: [[SEC2222]] = load i32**, i32*** [[SEC22222:%[^,]+]], 1225 // CK19-DAG: [[SEC22222]] = getelementptr {{.*}}i32*** [[SEC222222:[^,]+]], i{{.+}} 1 1226 // CK19-DAG: [[SEC222222]] = load i32***, i32**** [[PTR]], 1227 1228 // CK19-USE: call void [[CALL28:@.+]](i32*** {{[^,]+}}) 1229 // CK19-NOUSE: call void [[CALL28:@.+]]() 1230 #pragma omp target map(tofrom: mptr[1][2][2:4]) 1231 { 1232#ifdef USE 1233 mptr[1][2][3]++; 1234#endif 1235 } 1236 1237 // Region 29 1238 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1239 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1240 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1241 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1242 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1243 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1244 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1245 1246 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1247 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1248 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**** 1249 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**** 1250 // CK19-DAG: store i32*** [[VAR0:%.+]], i32**** [[CBP0]] 1251 // CK19-DAG: store i32*** [[SEC0:%.+]], i32**** [[CP0]] 1252 // CK19-DAG: [[VAR0]] = load i32***, i32**** [[PTR:%[^,]+]], 1253 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}i32*** [[SEC00:[^,]+]], i{{.+}} 1 1254 // CK19-DAG: [[SEC00]] = load i32***, i32**** [[PTR]], 1255 1256 // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 1257 // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 1258 // CK19-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i32**** 1259 // CK19-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i32*** 1260 // CK19-DAG: store i32*** [[SEC0]], i32**** [[CBP1]] 1261 // CK19-DAG: store i32** [[SEC1:%.+]], i32*** [[CP1]] 1262 // CK19-DAG: [[SEC1]] = getelementptr {{.*}}i32** [[SEC11:[^,]+]], i{{.+}} 2 1263 // CK19-DAG: [[SEC11]] = load i32**, i32*** [[SEC111:%[^,]+]], 1264 // CK19-DAG: [[SEC111]] = getelementptr {{.*}}i32*** [[SEC1111:[^,]+]], i{{.+}} 1 1265 // CK19-DAG: [[SEC1111]] = load i32***, i32**** [[PTR]], 1266 1267 // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 1268 // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 1269 // CK19-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to i32*** 1270 // CK19-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i32** 1271 // CK19-DAG: store i32** [[SEC1]], i32*** [[CBP2]] 1272 // CK19-DAG: store i32* [[SEC2:%.+]], i32** [[CP2]] 1273 // CK19-DAG: [[SEC2]] = getelementptr {{.*}}i32* [[SEC22:[^,]+]], i{{.+}} 3 1274 // CK19-DAG: [[SEC22]] = load i32*, i32** [[SEC222:%[^,]+]], 1275 // CK19-DAG: [[SEC222]] = getelementptr {{.*}}i32** [[SEC2222:[^,]+]], i{{.+}} 2 1276 // CK19-DAG: [[SEC2222]] = load i32**, i32*** [[SEC22222:%[^,]+]], 1277 // CK19-DAG: [[SEC22222]] = getelementptr {{.*}}i32*** [[SEC222222:[^,]+]], i{{.+}} 1 1278 // CK19-DAG: [[SEC222222]] = load i32***, i32**** [[PTR]], 1279 1280 // CK19-USE: call void [[CALL29:@.+]](i32*** {{[^,]+}}) 1281 // CK19-NOUSE: call void [[CALL29:@.+]]() 1282 #pragma omp target map(tofrom: mptr[1][2][3]) 1283 { 1284#ifdef USE 1285 mptr[1][2][3]++; 1286#endif 1287 } 1288 1289 // Multidimensional VLA. 1290 double mva[23][ii][ii+5]; 1291 1292 // Region 30 1293 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1294 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1295 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1296 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1297 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1298 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 1299 // CK19-DAG: store i64* [[SIZES:%.+]], i64** [[SARG]] 1300 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1301 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1302 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 1303 // 1304 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1305 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1306 // CK19-USE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i[[Z]]* 1307 // CK19-USE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i[[Z]]* 1308 // CK19-USE-DAG: store i[[Z]] 23, i[[Z]]* [[CBP0]] 1309 // CK19-USE-DAG: store i[[Z]] 23, i[[Z]]* [[CP0]] 1310 // 1311 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 1312 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 1313 // CK19-USE-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[Z]]* 1314 // CK19-USE-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[Z]]* 1315 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], i[[Z]]* [[CBP1]] 1316 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], i[[Z]]* [[CP1]] 1317 // CK19-64-USE-DAG: [[VAR1]] = zext i32 %{{[^,]+}} to i64 1318 // CK19-64-USE-DAG: [[VAR11]] = zext i32 %{{[^,]+}} to i64 1319 // 1320 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 1321 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 1322 // CK19-USE-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to i[[Z]]* 1323 // CK19-USE-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i[[Z]]* 1324 // CK19-USE-DAG: store i[[Z]] [[VAR2:%.+]], i[[Z]]* [[CBP2]] 1325 // CK19-USE-DAG: store i[[Z]] [[VAR22:%.+]], i[[Z]]* [[CP2]] 1326 // CK19-64-USE-DAG: [[VAR2]] = zext i32 %{{[^,]+}} to i64 1327 // CK19-64-USE-DAG: [[VAR22]] = zext i32 %{{[^,]+}} to i64 1328 // 1329 // CK19-USE-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3 1330 // CK19-USE-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3 1331 // CK19-USE-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 3 1332 // CK19-USE-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to double** 1333 // CK19-USE-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to double** 1334 // CK19-USE-DAG: store double* [[VAR3:%.+]], double** [[CBP3]] 1335 // CK19-USE-DAG: store double* [[VAR3]], double** [[CP3]] 1336 // CK19-USE-DAG: store i64 [[CSVAL3:%[^,]+]], i64* [[S3]] 1337 // CK19-USE-DAG: [[CSVAL3]] = {{mul nuw i64 %[^,]+, 8|sext i32 .+ to i64}} 1338 1339 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1340 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1341 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 1342 // CK19-NOUSE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to double** 1343 // CK19-NOUSE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double** 1344 // CK19-NOUSE-DAG: store double* [[VAR0:%.+]], double** [[CBP0]] 1345 // CK19-NOUSE-DAG: store double* [[VAR0]], double** [[CP0]] 1346 // CK19-NOUSE-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]] 1347 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 8|sext i32 .+ to i64}} 1348 1349 // CK19-USE: call void [[CALL30:@.+]](i[[Z]] 23, i[[Z]] %{{[^,]+}}, i[[Z]] %{{[^,]+}}, double* %{{[^,]+}}) 1350 // CK19-NOUSE: call void [[CALL30:@.+]]() 1351 #pragma omp target map(tofrom: mva) 1352 { 1353#ifdef USE 1354 mva[1][2][3]++; 1355#endif 1356 } 1357 1358 // Region 31 1359 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1360 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1361 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1362 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1363 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1364 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1365 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1366 // 1367 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1368 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1369 // CK19-USE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i[[Z]]* 1370 // CK19-USE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i[[Z]]* 1371 // CK19-USE-DAG: store i[[Z]] 23, i[[Z]]* [[CBP0]] 1372 // CK19-USE-DAG: store i[[Z]] 23, i[[Z]]* [[CP0]] 1373 // 1374 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 1375 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 1376 // CK19-USE-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[Z]]* 1377 // CK19-USE-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[Z]]* 1378 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], i[[Z]]* [[CBP1]] 1379 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], i[[Z]]* [[CP1]] 1380 // 1381 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 1382 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 1383 // CK19-USE-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to i[[Z]]* 1384 // CK19-USE-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to i[[Z]]* 1385 // CK19-USE-DAG: store i[[Z]] [[VAR2:%.+]], i[[Z]]* [[CBP2]] 1386 // CK19-USE-DAG: store i[[Z]] [[VAR22:%.+]], i[[Z]]* [[CP2]] 1387 // 1388 // CK19-USE-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 3 1389 // CK19-USE-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 3 1390 // CK19-USE-DAG: [[CBP3:%.+]] = bitcast i8** [[BP3]] to double** 1391 // CK19-USE-DAG: [[CP3:%.+]] = bitcast i8** [[P3]] to double** 1392 // CK19-USE-DAG: store double* [[VAR3:%.+]], double** [[CBP3]] 1393 // CK19-USE-DAG: store double* [[SEC3:%.+]], double** [[CP3]] 1394 // CK19-USE-DAG: [[SEC3]] = getelementptr {{.*}}double* [[SEC33:%.+]], i[[Z]] 0 1395 // CK19-USE-DAG: [[SEC33]] = getelementptr {{.*}}double* [[SEC333:%.+]], i[[Z]] [[IDX3:%.+]] 1396 // CK19-USE-DAG: [[IDX3]] = mul nsw i[[Z]] %{{[^,]+}}, %{{[^,]+}} 1397 // CK19-USE-DAG: [[SEC333]] = getelementptr {{.*}}double* [[VAR3]], i[[Z]] [[IDX33:%.+]] 1398 // CK19-USE-DAG: [[IDX33]] = mul nsw i[[Z]] 1, %{{[^,]+}} 1399 1400 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1401 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1402 // CK19-NOUSE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to double** 1403 // CK19-NOUSE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double** 1404 // CK19-NOUSE-DAG: store double* [[VAR0:%.+]], double** [[CBP0]] 1405 // CK19-NOUSE-DAG: store double* [[SEC0:%.+]], double** [[CP0]] 1406 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}double* [[SEC00:%.+]], i[[Z:64|32]] 0 1407 // CK19-NOUSE-DAG: [[SEC00]] = getelementptr {{.*}}double* [[SEC000:%.+]], i[[Z]] [[IDX0:%.+]] 1408 // CK19-NOUSE-DAG: [[IDX0]] = mul nsw i[[Z]] %{{[^,]+}}, %{{[^,]+}} 1409 // CK19-NOUSE-DAG: [[SEC000]] = getelementptr {{.*}}double* [[VAR0]], i[[Z]] [[IDX00:%.+]] 1410 // CK19-NOUSE-DAG: [[IDX00]] = mul nsw i[[Z]] 1, %{{[^,]+}} 1411 1412 // CK19-USE: call void [[CALL31:@.+]](i[[Z]] 23, i[[Z]] %{{[^,]+}}, i[[Z]] %{{[^,]+}}, double* %{{[^,]+}}) 1413 // CK19-NOUSE: call void [[CALL31:@.+]]() 1414 #pragma omp target map(tofrom: mva[1][ii-2][:5]) 1415 { 1416#ifdef USE 1417 mva[1][2][3]++; 1418#endif 1419 } 1420 1421 // Multidimensional array sections. 1422 double marras[11][12][13]; 1423 double mvlaas[11][ii][13]; 1424 double ***mptras; 1425 1426 // Region 32 1427 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1428 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1429 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1430 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1431 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1432 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1433 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1434 1435 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1436 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1437 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [11 x [12 x [13 x double]]]** 1438 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [11 x [12 x [13 x double]]]** 1439 // CK19-DAG: store [11 x [12 x [13 x double]]]* [[VAR0:%.+]], [11 x [12 x [13 x double]]]** [[CBP0]] 1440 // CK19-DAG: store [11 x [12 x [13 x double]]]* [[VAR0]], [11 x [12 x [13 x double]]]** [[CP0]] 1441 1442 // CK19-USE: call void [[CALL32:@.+]]([11 x [12 x [13 x double]]]* {{[^,]+}}) 1443 // CK19-NOUSE: call void [[CALL32:@.+]]() 1444 #pragma omp target map(marras) 1445 { 1446#ifdef USE 1447 marras[1][2][3]++; 1448#endif 1449 } 1450 1451 // Region 33 1452 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1453 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1454 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1455 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1456 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1457 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1458 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1459 1460 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1461 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1462 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [11 x [12 x [13 x double]]]** 1463 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [12 x [13 x double]]** 1464 // CK19-DAG: store [11 x [12 x [13 x double]]]* [[VAR0:%.+]], [11 x [12 x [13 x double]]]** [[CBP0]] 1465 // CK19-DAG: store [12 x [13 x double]]* [[SEC0:%.+]], [12 x [13 x double]]** [[CP0]] 1466 // CK19-DAG: [[SEC0]] = getelementptr {{.+}}[11 x [12 x [13 x double]]]* [[VAR0]], i[[Z]] 0, i[[Z]] 0 1467 1468 // CK19-USE: call void [[CALL33:@.+]]([11 x [12 x [13 x double]]]* {{[^,]+}}) 1469 // CK19-NOUSE: call void [[CALL33:@.+]]() 1470 #pragma omp target map(marras[:]) 1471 { 1472#ifdef USE 1473 marras[1][2][3]++; 1474#endif 1475 } 1476 1477 // Region 34 1478 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1479 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1480 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1481 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1482 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1483 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1484 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1485 1486 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1487 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1488 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [11 x [12 x [13 x double]]]** 1489 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [12 x [13 x double]]** 1490 // CK19-DAG: store [11 x [12 x [13 x double]]]* [[VAR0:%.+]], [11 x [12 x [13 x double]]]** [[CBP0]] 1491 // CK19-DAG: store [12 x [13 x double]]* [[SEC0:%.+]], [12 x [13 x double]]** [[CP0]] 1492 // CK19-DAG: [[SEC0]] = getelementptr {{.+}}[11 x [12 x [13 x double]]]* [[VAR0]], i[[Z]] 0, i[[Z]] 0 1493 1494 // CK19-USE: call void [[CALL34:@.+]]([11 x [12 x [13 x double]]]* {{[^,]+}}) 1495 // CK19-NOUSE: call void [[CALL34:@.+]]() 1496 #pragma omp target map(marras[:][:][:]) 1497 { 1498#ifdef USE 1499 marras[1][2][3]++; 1500#endif 1501 } 1502 1503 // Region 35 1504 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1505 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1506 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1507 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1508 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1509 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 1510 // CK19-DAG: store i64* [[SIZES:%.+]], i64** [[SARG]] 1511 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1512 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1513 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 1514 // 1515 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1516 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1517 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 1518 1519 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [11 x [12 x [13 x double]]]** 1520 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [13 x double]** 1521 // CK19-DAG: store [11 x [12 x [13 x double]]]* [[VAR0:%.+]], [11 x [12 x [13 x double]]]** [[CBP0]] 1522 // CK19-DAG: store [13 x double]* [[SEC0:%.+]], [13 x double]** [[CP0]] 1523 // CK19-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]] 1524 // CK19-DAG: [[SEC0]] = getelementptr {{.+}}[12 x [13 x double]]* [[SEC00:%[^,]+]], i[[Z]] 0, i[[Z]] 0 1525 // CK19-DAG: [[SEC00]] = getelementptr {{.+}}[11 x [12 x [13 x double]]]* [[VAR0]], i[[Z]] 0, i[[Z]] 1 1526 // CK19-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}} 1527 1528 // CK19-USE: call void [[CALL35:@.+]]([11 x [12 x [13 x double]]]* {{[^,]+}}) 1529 // CK19-NOUSE: call void [[CALL35:@.+]]() 1530 #pragma omp target map(marras[1][:ii][:]) 1531 { 1532#ifdef USE 1533 marras[1][2][3]++; 1534#endif 1535 } 1536 1537 // Region 36 1538 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1539 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1540 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1541 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1542 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1543 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1544 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1545 1546 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1547 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1548 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [11 x [12 x [13 x double]]]** 1549 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [13 x double]** 1550 // CK19-DAG: store [11 x [12 x [13 x double]]]* [[VAR0:%.+]], [11 x [12 x [13 x double]]]** [[CBP0]] 1551 // CK19-DAG: store [13 x double]* [[SEC0:%.+]], [13 x double]** [[CP0]] 1552 // CK19-DAG: [[SEC0]] = getelementptr {{.+}}[13 x double]* [[SEC00:%[^,]+]], i{{.+}} 0 1553 // CK19-DAG: [[SEC00]] = getelementptr {{.+}}[12 x [13 x double]]* [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 0 1554 // CK19-DAG: [[SEC000]] = getelementptr {{.+}}[11 x [12 x [13 x double]]]* [[VAR0]], i{{.+}} 0, i{{.+}} 0 1555 1556 // CK19-USE: call void [[CALL36:@.+]]([11 x [12 x [13 x double]]]* {{[^,]+}}) 1557 // CK19-NOUSE: call void [[CALL36:@.+]]() 1558 #pragma omp target map(marras[:1][:2][:13]) 1559 { 1560#ifdef USE 1561 marras[1][2][3]++; 1562#endif 1563 } 1564 1565 // Region 37 1566 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1567 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1568 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1569 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1570 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1571 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 1572 // CK19-DAG: store i64* [[SIZES:%.+]], i64** [[SARG]] 1573 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1574 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1575 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 1576 // 1577 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1578 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1579 // CK19-USE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i[[Z]]* 1580 // CK19-USE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i[[Z]]* 1581 // CK19-USE-DAG: store i[[Z]] 11, i[[Z]]* [[CBP0]] 1582 // CK19-USE-DAG: store i[[Z]] 11, i[[Z]]* [[CP0]] 1583 // 1584 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 1585 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 1586 // CK19-USE-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[Z]]* 1587 // CK19-USE-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[Z]]* 1588 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], i[[Z]]* [[CBP1]] 1589 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], i[[Z]]* [[CP1]] 1590 // 1591 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 1592 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 1593 // CK19-USE-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2 1594 // CK19-USE-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [13 x double]** 1595 // CK19-USE-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to [13 x double]** 1596 // CK19-USE-DAG: store [13 x double]* [[VAR2:%.+]], [13 x double]** [[CBP2]] 1597 // CK19-USE-DAG: store [13 x double]* [[VAR2]], [13 x double]** [[CP2]] 1598 // CK19-USE-DAG: store i64 [[CSVAL2:%[^,]+]], i64* [[S2]] 1599 // CK19-USE-DAG: [[CSVAL2]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}} 1600 1601 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1602 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1603 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 1604 // CK19-NOUSE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [13 x double]** 1605 // CK19-NOUSE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [13 x double]** 1606 // CK19-NOUSE-DAG: store [13 x double]* [[VAR0:%.+]], [13 x double]** [[CBP0]] 1607 // CK19-NOUSE-DAG: store [13 x double]* [[VAR0]], [13 x double]** [[CP0]] 1608 // CK19-NOUSE-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]] 1609 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}} 1610 1611 // CK19-USE: call void [[CALL37:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, [13 x double]* %{{[^,]+}}) 1612 // CK19-NOUSE: call void [[CALL37:@.+]]() 1613 #pragma omp target map(mvlaas) 1614 { 1615#ifdef USE 1616 mvlaas[1][2][3]++; 1617#endif 1618 } 1619 1620 // Region 38 1621 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1622 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1623 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1624 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1625 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1626 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 1627 // CK19-DAG: store i64* [[SIZES:%.+]], i64** [[SARG]] 1628 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1629 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1630 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 1631 // 1632 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1633 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1634 // CK19-USE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i[[Z]]* 1635 // CK19-USE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i[[Z]]* 1636 // CK19-USE-DAG: store i[[Z]] 11, i[[Z]]* [[CBP0]] 1637 // CK19-USE-DAG: store i[[Z]] 11, i[[Z]]* [[CP0]] 1638 // 1639 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 1640 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 1641 // CK19-USE-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[Z]]* 1642 // CK19-USE-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[Z]]* 1643 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], i[[Z]]* [[CBP1]] 1644 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], i[[Z]]* [[CP1]] 1645 // 1646 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 1647 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 1648 // CK19-USE-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2 1649 // CK19-USE-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [13 x double]** 1650 // CK19-USE-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to [13 x double]** 1651 // CK19-USE-DAG: store [13 x double]* [[VAR2:%.+]], [13 x double]** [[CBP2]] 1652 // CK19-USE-DAG: store [13 x double]* [[SEC2:%.+]], [13 x double]** [[CP2]] 1653 // CK19-USE-DAG: store i64 [[CSVAL2:%[^,]+]], i64* [[S2]] 1654 // CK19-USE-DAG: [[SEC2]] = getelementptr {{.+}}[13 x double]* [[VAR2]], i[[Z]] [[SEC22:%[^,]+]] 1655 // CK19-USE-DAG: [[SEC22]] = mul nsw i[[Z]] 0, %{{[^,]+}} 1656 // CK19-USE-DAG: [[CSVAL2]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}} 1657 1658 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1659 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1660 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 1661 // CK19-NOUSE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [13 x double]** 1662 // CK19-NOUSE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [13 x double]** 1663 // CK19-NOUSE-DAG: store [13 x double]* [[VAR0:%.+]], [13 x double]** [[CBP0]] 1664 // CK19-NOUSE-DAG: store [13 x double]* [[SEC0:%.+]], [13 x double]** [[CP0]] 1665 // CK19-NOUSE-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]] 1666 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.+}}[13 x double]* [[VAR0]], i[[Z]] [[SEC00:%[^,]+]] 1667 // CK19-NOUSE-DAG: [[SEC00]] = mul nsw i[[Z]] 0, %{{[^,]+}} 1668 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}} 1669 1670 // CK19-USE: call void [[CALL38:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, [13 x double]* %{{[^,]+}}) 1671 // CK19-NOUSE: call void [[CALL38:@.+]]() 1672 #pragma omp target map(mvlaas[:]) 1673 { 1674#ifdef USE 1675 mvlaas[1][2][3]++; 1676#endif 1677 } 1678 1679 // Region 39 1680 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1681 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1682 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1683 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1684 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1685 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 1686 // CK19-DAG: store i64* [[SIZES:%.+]], i64** [[SARG]] 1687 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1688 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1689 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 1690 // 1691 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1692 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1693 // CK19-USE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i[[Z]]* 1694 // CK19-USE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i[[Z]]* 1695 // CK19-USE-DAG: store i[[Z]] 11, i[[Z]]* [[CBP0]] 1696 // CK19-USE-DAG: store i[[Z]] 11, i[[Z]]* [[CP0]] 1697 // 1698 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 1699 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 1700 // CK19-USE-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[Z]]* 1701 // CK19-USE-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[Z]]* 1702 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], i[[Z]]* [[CBP1]] 1703 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], i[[Z]]* [[CP1]] 1704 // 1705 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 1706 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 1707 // CK19-USE-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2 1708 // CK19-USE-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [13 x double]** 1709 // CK19-USE-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to [13 x double]** 1710 // CK19-USE-DAG: store [13 x double]* [[VAR2:%.+]], [13 x double]** [[CBP2]] 1711 // CK19-USE-DAG: store [13 x double]* [[SEC2:%.+]], [13 x double]** [[CP2]] 1712 // CK19-USE-DAG: store i64 [[CSVAL2:%[^,]+]], i64* [[S2]] 1713 // CK19-USE-DAG: [[SEC2]] = getelementptr {{.+}}[13 x double]* [[VAR2]], i[[Z]] [[SEC22:%[^,]+]] 1714 // CK19-USE-DAG: [[SEC22]] = mul nsw i[[Z]] 0, %{{[^,]+}} 1715 // CK19-USE-DAG: [[CSVAL2]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}} 1716 1717 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1718 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1719 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 1720 // CK19-NOUSE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [13 x double]** 1721 // CK19-NOUSE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [13 x double]** 1722 // CK19-NOUSE-DAG: store [13 x double]* [[VAR0:%.+]], [13 x double]** [[CBP0]] 1723 // CK19-NOUSE-DAG: store [13 x double]* [[SEC0:%.+]], [13 x double]** [[CP0]] 1724 // CK19-NOUSE-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]] 1725 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.+}}[13 x double]* [[VAR0]], i[[Z]] [[SEC00:%[^,]+]] 1726 // CK19-NOUSE-DAG: [[SEC00]] = mul nsw i[[Z]] 0, %{{[^,]+}} 1727 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}} 1728 1729 // CK19-USE: call void [[CALL39:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, [13 x double]* %{{[^,]+}}) 1730 // CK19-NOUSE: call void [[CALL39:@.+]]() 1731 #pragma omp target map(mvlaas[:][:][:]) 1732 { 1733#ifdef USE 1734 mvlaas[1][2][3]++; 1735#endif 1736 } 1737 1738 // Region 40 1739 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1740 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1741 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1742 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1743 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1744 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 1745 // CK19-DAG: store i64* [[SIZES:%.+]], i64** [[SARG]] 1746 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1747 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1748 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 1749 // 1750 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1751 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1752 // CK19-USE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i[[Z]]* 1753 // CK19-USE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i[[Z]]* 1754 // CK19-USE-DAG: store i[[Z]] 11, i[[Z]]* [[CBP0]] 1755 // CK19-USE-DAG: store i[[Z]] 11, i[[Z]]* [[CP0]] 1756 // 1757 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 1758 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 1759 // CK19-USE-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[Z]]* 1760 // CK19-USE-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[Z]]* 1761 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], i[[Z]]* [[CBP1]] 1762 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], i[[Z]]* [[CP1]] 1763 // 1764 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 1765 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 1766 // CK19-USE-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 2 1767 // CK19-USE-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [13 x double]** 1768 // CK19-USE-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to [13 x double]** 1769 // CK19-USE-DAG: store [13 x double]* [[VAR2:%.+]], [13 x double]** [[CBP2]] 1770 // CK19-USE-DAG: store [13 x double]* [[SEC2:%.+]], [13 x double]** [[CP2]] 1771 // CK19-USE-DAG: store i64 [[CSVAL2:%[^,]+]], i64* [[S2]] 1772 // CK19-USE-DAG: [[SEC2]] = getelementptr {{.+}}[13 x double]* [[SEC22:%[^,]+]], i[[Z]] 0 1773 // CK19-USE-DAG: [[SEC22]] = getelementptr {{.+}}[13 x double]* [[VAR2]], i[[Z]] [[SEC222:%[^,]+]] 1774 // CK19-USE-DAG: [[SEC222]] = mul nsw i[[Z]] 1, %{{[^,]+}} 1775 1776 // CK19-NOUSE-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 1777 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1778 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1779 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 1780 // CK19-NOUSE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [13 x double]** 1781 // CK19-NOUSE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [13 x double]** 1782 // CK19-NOUSE-DAG: store [13 x double]* [[VAR0:%.+]], [13 x double]** [[CBP0]] 1783 // CK19-NOUSE-DAG: store [13 x double]* [[SEC0:%.+]], [13 x double]** [[CP0]] 1784 // CK19-NOUSE-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]] 1785 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.+}}[13 x double]* [[SEC00:%[^,]+]], i[[Z]] 0 1786 // CK19-NOUSE-DAG: [[SEC00]] = getelementptr {{.+}}[13 x double]* [[VAR0]], i[[Z]] [[SEC000:%[^,]+]] 1787 // CK19-NOUSE-DAG: [[SEC000]] = mul nsw i[[Z]] 1, %{{[^,]+}} 1788 1789 // CK19-USE: call void [[CALL40:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, [13 x double]* %{{[^,]+}}) 1790 // CK19-NOUSE: call void [[CALL40:@.+]]() 1791 #pragma omp target map(mvlaas[1][:ii][:]) 1792 { 1793#ifdef USE 1794 mvlaas[1][2][3]++; 1795#endif 1796 } 1797 1798 // Region 41 1799 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1800 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1801 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1802 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1803 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1804 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1805 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1806 // 1807 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1808 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1809 // CK19-USE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i[[Z]]* 1810 // CK19-USE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i[[Z]]* 1811 // CK19-USE-DAG: store i[[Z]] 11, i[[Z]]* [[CBP0]] 1812 // CK19-USE-DAG: store i[[Z]] 11, i[[Z]]* [[CP0]] 1813 // 1814 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 1815 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 1816 // CK19-USE-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to i[[Z]]* 1817 // CK19-USE-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to i[[Z]]* 1818 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], i[[Z]]* [[CBP1]] 1819 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], i[[Z]]* [[CP1]] 1820 // 1821 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 1822 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 1823 // CK19-USE-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to [13 x double]** 1824 // CK19-USE-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to [13 x double]** 1825 // CK19-USE-DAG: store [13 x double]* [[VAR2:%.+]], [13 x double]** [[CBP2]] 1826 // CK19-USE-DAG: store [13 x double]* [[SEC2:%.+]], [13 x double]** [[CP2]] 1827 // CK19-USE-DAG: [[SEC2]] = getelementptr {{.+}}[13 x double]* [[SEC22:%[^,]+]], i[[Z]] 0 1828 // CK19-USE-DAG: [[SEC22]] = getelementptr {{.+}}[13 x double]* [[VAR2]], i[[Z]] [[SEC222:%[^,]+]] 1829 // CK19-USE-DAG: [[SEC222]] = mul nsw i[[Z]] 0, %{{[^,]+}} 1830 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 1831 1832 // CK19-NO-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1833 // CK19-NO-USE-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [13 x double]** 1834 // CK19-NO-USE-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [13 x double]** 1835 // CK19-NO-USE-DAG: store [13 x double]* [[VAR0:%.+]], [13 x double]** [[CBP0]] 1836 // CK19-NO-USE-DAG: store [13 x double]* [[SEC0:%.+]], [13 x double]** [[CP0]] 1837 // CK19-NO-USE-DAG: [[SEC0]] = getelementptr {{.+}}[13 x double]* [[SEC00:%[^,]+]], i[[Z]] 0 1838 // CK19-NO-USE-DAG: [[SEC00]] = getelementptr {{.+}}[13 x double]* [[VAR0]], i[[Z]] [[SEC000:%[^,]+]] 1839 // CK19-NO-USE-DAG: [[SEC000]] = mul nsw i[[Z]] 0, %{{[^,]+}} 1840 1841 // CK19-USE: call void [[CALL41:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, [13 x double]* %{{[^,]+}}) 1842 // CK19-NOUSE: call void [[CALL41:@.+]]() 1843 #pragma omp target map(mvlaas[:1][:2][:13]) 1844 { 1845#ifdef USE 1846 mvlaas[1][2][3]++; 1847#endif 1848 } 1849 1850 // Region 42 1851 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1852 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1853 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1854 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1855 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1856 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1857 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1858 1859 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1860 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1861 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to double**** 1862 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double**** 1863 // CK19-DAG: store double*** [[VAR0:%.+]], double**** [[CBP0]] 1864 // CK19-DAG: store double*** [[SEC0:%.+]], double**** [[CP0]] 1865 // CK19-DAG: [[VAR0]] = load double***, double**** [[PTR:%[^,]+]], 1866 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}double*** [[SEC00:[^,]+]], i{{.+}} 0 1867 // CK19-DAG: [[SEC00]] = load double***, double**** [[PTR]], 1868 1869 // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1 1870 // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1 1871 // CK19-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to double**** 1872 // CK19-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to double*** 1873 // CK19-DAG: store double*** [[SEC0]], double**** [[CBP1]] 1874 // CK19-DAG: store double** [[SEC1:%.+]], double*** [[CP1]] 1875 // CK19-DAG: [[SEC1]] = getelementptr {{.*}}double** [[SEC11:[^,]+]], i{{.+}} 2 1876 // CK19-DAG: [[SEC11]] = load double**, double*** [[SEC111:%[^,]+]], 1877 // CK19-DAG: [[SEC111]] = getelementptr {{.*}}double*** [[SEC1111:[^,]+]], i{{.+}} 0 1878 // CK19-DAG: [[SEC1111]] = load double***, double**** [[PTR]], 1879 1880 // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 2 1881 // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 2 1882 // CK19-DAG: [[CBP2:%.+]] = bitcast i8** [[BP2]] to double*** 1883 // CK19-DAG: [[CP2:%.+]] = bitcast i8** [[P2]] to double** 1884 // CK19-DAG: store double** [[SEC1]], double*** [[CBP2]] 1885 // CK19-DAG: store double* [[SEC2:%.+]], double** [[CP2]] 1886 // CK19-DAG: [[SEC2]] = getelementptr {{.*}}double* [[SEC22:[^,]+]], i{{.+}} 0 1887 // CK19-DAG: [[SEC22]] = load double*, double** [[SEC222:%[^,]+]], 1888 // CK19-DAG: [[SEC222]] = getelementptr {{.*}}double** [[SEC2222:[^,]+]], i{{.+}} 2 1889 // CK19-DAG: [[SEC2222]] = load double**, double*** [[SEC22222:%[^,]+]], 1890 // CK19-DAG: [[SEC22222]] = getelementptr {{.*}}double*** [[SEC222222:[^,]+]], i{{.+}} 0 1891 // CK19-DAG: [[SEC222222]] = load double***, double**** [[PTR]], 1892 1893 // CK19-USE: call void [[CALL42:@.+]](double*** {{[^,]+}}) 1894 // CK19-NOUSE: call void [[CALL42:@.+]]() 1895 #pragma omp target map(mptras[:1][2][:13]) 1896 { 1897#ifdef USE 1898 mptras[1][2][3]++; 1899#endif 1900 } 1901 1902 // Region 43 - the memory is not contiguous for this map - will map the whole last dimension. 1903 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1904 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1905 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1906 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1907 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1908 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4 1909 // CK19-DAG: store i64* [[SIZES:%.+]], i64** [[SARG]] 1910 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1911 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1912 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] 1913 // 1914 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1915 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1916 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 1917 1918 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [11 x [12 x [13 x double]]]** 1919 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [13 x double]** 1920 // CK19-DAG: store [11 x [12 x [13 x double]]]* [[VAR0:%.+]], [11 x [12 x [13 x double]]]** [[CBP0]] 1921 // CK19-DAG: store [13 x double]* [[SEC0:%.+]], [13 x double]** [[CP0]] 1922 // CK19-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]] 1923 // CK19-DAG: [[SEC0]] = getelementptr {{.+}}[12 x [13 x double]]* [[SEC00:%[^,]+]], i[[Z]] 0, i[[Z]] 0 1924 // CK19-DAG: [[SEC00]] = getelementptr {{.+}}[11 x [12 x [13 x double]]]* [[VAR0]], i[[Z]] 0, i[[Z]] 1 1925 // CK19-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}} 1926 1927 // CK19-USE: call void [[CALL43:@.+]]([11 x [12 x [13 x double]]]* {{[^,]+}}) 1928 // CK19-NOUSE: call void [[CALL43:@.+]]() 1929 #pragma omp target map(marras[1][:ii][1:]) 1930 { 1931#ifdef USE 1932 marras[1][2][3]++; 1933#endif 1934 } 1935 1936 // Region 44 1937 // CK19-DAG: call i32 @__tgt_target_kernel(%struct.ident_t* @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, i8* @.{{.+}}.region_id, %struct.__tgt_kernel_arguments* [[ARGS:%.+]]) 1938 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2 1939 // CK19-DAG: store i8** [[BPGEP:%.+]], i8*** [[BPARG]] 1940 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3 1941 // CK19-DAG: store i8** [[PGEP:%.+]], i8*** [[PARG]] 1942 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] 1943 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] 1944 1945 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 1946 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 1947 // CK19-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [100 x i32]** 1948 // CK19-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** 1949 // CK19-DAG: store [100 x i32]* [[VAR0:%.+]], [100 x i32]** [[CBP0]] 1950 // CK19-DAG: store i32* [[SEC0:%[^,]+]], i32** [[CP0]] 1951 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}[100 x i32]* [[VAR0]], i{{.+}} 0, i{{.+}} 20 1952 1953 // CK19-USE: call void [[CALL44:@.+]]([100 x i32]* {{[^,]+}}) 1954 // CK19-NOUSE: call void [[CALL44:@.+]]() 1955 #pragma omp target map(from:arra[20:]) 1956 { 1957#ifdef USE 1958 arra[50]++; 1959#endif 1960 } 1961 1962} 1963 1964// CK19: define {{.+}}[[CALL00]] 1965// CK19: define {{.+}}[[CALL01]] 1966// CK19: define {{.+}}[[CALL02]] 1967// CK19: define {{.+}}[[CALL03]] 1968// CK19: define {{.+}}[[CALL04]] 1969// CK19: define {{.+}}[[CALL05]] 1970// CK19: define {{.+}}[[CALL06]] 1971// CK19: define {{.+}}[[CALL07]] 1972// CK19: define {{.+}}[[CALL08]] 1973// CK19: define {{.+}}[[CALL09]] 1974// CK19: define {{.+}}[[CALL10]] 1975// CK19: define {{.+}}[[CALL11]] 1976// CK19: define {{.+}}[[CALL12]] 1977// CK19: define {{.+}}[[CALL13]] 1978// CK19: define {{.+}}[[CALL14]] 1979// CK19: define {{.+}}[[CALL15]] 1980// CK19: define {{.+}}[[CALL16]] 1981// CK19: define {{.+}}[[CALL17]] 1982// CK19: define {{.+}}[[CALL18]] 1983// CK19: define {{.+}}[[CALL19]] 1984// CK19: define {{.+}}[[CALL20]] 1985// CK19: define {{.+}}[[CALL21]] 1986// CK19: define {{.+}}[[CALL22]] 1987// CK19: define {{.+}}[[CALL23]] 1988// CK19: define {{.+}}[[CALL24]] 1989// CK19: define {{.+}}[[CALL25]] 1990// CK19: define {{.+}}[[CALL26]] 1991// CK19: define {{.+}}[[CALL27]] 1992// CK19: define {{.+}}[[CALL28]] 1993// CK19: define {{.+}}[[CALL29]] 1994// CK19: define {{.+}}[[CALL30]] 1995// CK19: define {{.+}}[[CALL31]] 1996// CK19: define {{.+}}[[CALL32]] 1997// CK19: define {{.+}}[[CALL33]] 1998// CK19: define {{.+}}[[CALL34]] 1999// CK19: define {{.+}}[[CALL35]] 2000// CK19: define {{.+}}[[CALL36]] 2001// CK19: define {{.+}}[[CALL37]] 2002// CK19: define {{.+}}[[CALL38]] 2003// CK19: define {{.+}}[[CALL39]] 2004// CK19: define {{.+}}[[CALL40]] 2005// CK19: define {{.+}}[[CALL41]] 2006// CK19: define {{.+}}[[CALL42]] 2007// CK19: define {{.+}}[[CALL43]] 2008// CK19: define {{.+}}[[CALL44]] 2009 2010#endif // CK19 2011#endif // HEADER_INC 2012