1834 lines · plain
1// expected-no-diagnostics2#ifndef HEADER_INC3#define HEADER_INC4 5// This file is regex-heavy and takes a long time to execute the test. To speed6// testing up, test execution is split over multiple fimes. The RUN commands are7// in the corresponding .cpp files now. Do not add them here.8 9// SIMD-ONLY18-NOT: {{__kmpc|__tgt}}10#ifdef CK1911 12// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 013// 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] zeroinitializer16 17// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 018// 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] zeroinitializer21 22// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 023// 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 028// 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 033// 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 038// 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] zeroinitializer41 42// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 043// 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 048// 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 052// CK19-USE: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i64] [i64 32]53// CK19-NOUSE: [[MTYPE07:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer54 55// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 056// 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 061// 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 066// 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 071// 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] zeroinitializer74 75// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 076// 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 081// CK19-USE: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i64] [i64 32]82// CK19-NOUSE: [[MTYPE13:@.+]] = private {{.*}}constant [1 x i64] zeroinitializer83 84// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 085// 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 089// 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 094// 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 099// 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 0105// 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 0111// 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] zeroinitializer114 115// CK19-LABEL: @.__omp_offloading_{{.*}}explicit_maps_single{{.*}}_l{{[0-9]+}}.region_id = weak constant i8 0116// 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 0122// 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 0127// 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 0133// 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 0138// 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 0143// 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 0148// 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 0153// 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 0158// 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 0163// 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 0168// 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 0173// 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 0179// 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 0184// 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 0189// 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 0194// 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 0198// 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 0203// 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 0208// 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 0213// 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 0218// 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 0223// 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 0229// 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 0234// 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 0238// 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 00248 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])249 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2250 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]251 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3252 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[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{{.+}} 0257 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0258 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]259 // CK19-DAG: store ptr [[VAR0]], ptr [[P0]]260 261 // CK19-USE: call void [[CALL00:@.+]](ptr {{[^,]+}})262 // CK19-NOUSE: call void [[CALL00:@.+]]()263 #pragma omp target map(alloc:a)264 {265#ifdef USE266 ++a;267#endif268 }269 270 // Map of a scalar in nested region.271 int b = a;272 273 // Region 00n274 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 -1, i32 1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])275 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2276 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]277 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3278 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]279 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]280 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]281 282 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0283 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0284 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]285 // CK19-DAG: store ptr [[VAR0]], ptr [[P0]]286 287 // CK19-USE: call void [[CALL00n:@.+]](ptr {{[^,]+}})288 // CK19-NOUSE: call void [[CALL00n:@.+]]()289 #pragma omp target map(alloc:b)290 #pragma omp parallel291 {292#ifdef USE293 ++b;294#endif295 }296 297 // Map of an array.298 int arra[100];299 300 // Region 01301 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])302 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2303 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]304 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3305 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]306 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]307 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]308 309 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0310 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0311 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]312 // CK19-DAG: store ptr [[VAR0]], ptr [[P0]]313 314 // CK19-USE: call void [[CALL01:@.+]](ptr {{[^,]+}})315 // CK19-NOUSE: call void [[CALL01:@.+]]()316 #pragma omp target map(to:arra)317 {318#ifdef USE319 arra[50]++;320#endif321 }322 323 // Region 02324 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])325 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2326 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]327 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3328 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]329 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]330 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]331 332 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0333 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0334 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]335 // CK19-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]]336 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 20337 338 // CK19-USE: call void [[CALL02:@.+]](ptr {{[^,]+}})339 // CK19-NOUSE: call void [[CALL02:@.+]]()340 #pragma omp target map(from:arra[20:60])341 {342#ifdef USE343 arra[50]++;344#endif345 }346 347 // Region 03348 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])349 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2350 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]351 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3352 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]353 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]354 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]355 356 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0357 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0358 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]359 // CK19-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]]360 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 0361 362 // CK19-USE: call void [[CALL03:@.+]](ptr {{[^,]+}})363 // CK19-NOUSE: call void [[CALL03:@.+]]()364 #pragma omp target map(tofrom:arra[:60])365 {366#ifdef USE367 arra[50]++;368#endif369 }370 371 // Region 04372 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])373 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2374 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]375 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3376 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]377 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]378 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]379 380 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0381 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0382 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]383 // CK19-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]]384 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 0385 386 // CK19-USE: call void [[CALL04:@.+]](ptr {{[^,]+}})387 // CK19-NOUSE: call void [[CALL04:@.+]]()388 #pragma omp target map(alloc:arra[:])389 {390#ifdef USE391 arra[50]++;392#endif393 }394 395 // Region 05396 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])397 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2398 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]399 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3400 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]401 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]402 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]403 404 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0405 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0406 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]407 // CK19-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]]408 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 15409 410 // CK19-USE: call void [[CALL05:@.+]](ptr {{[^,]+}})411 // CK19-NOUSE: call void [[CALL05:@.+]]()412 #pragma omp target map(to:arra[15])413 {414#ifdef USE415 arra[15]++;416#endif417 }418 419 // Region 06420 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])421 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2422 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]423 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3424 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]425 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4426 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]427 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]428 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]429 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]430 431 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0432 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0433 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0434 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]435 // CK19-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]]436 // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], ptr [[S0]]437 // CK19-DAG: [[CSVAL0]] = {{mul nuw i.+ %.*, 4|sext i32 .+ to i64}}438 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} %{{.*}}439 440 // CK19-USE: call void [[CALL06:@.+]](ptr {{[^,]+}})441 // CK19-NOUSE: call void [[CALL06:@.+]]()442 #pragma omp target map(tofrom:arra[ii:ii+23])443 {444#ifdef USE445 arra[50]++;446#endif447 }448 449 // Region 07450 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])451 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2452 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]453 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3454 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]455 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4456 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]457 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]458 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]459 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]460 461 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0462 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0463 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0464 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]465 // CK19-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]]466 // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], ptr [[S0]]467 // CK19-DAG: [[CSVAL0]] = {{mul nuw i.+ %.*, 4|sext i32 .+ to i64}}468 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 0469 470 // CK19-USE: call void [[CALL07:@.+]](ptr {{[^,]+}})471 // CK19-NOUSE: call void [[CALL07:@.+]]()472 #pragma omp target map(alloc:arra[:ii])473 {474#ifdef USE475 arra[50]++;476#endif477 }478 479 // Region 08480 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])481 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2482 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]483 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3484 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]485 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]486 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]487 488 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0489 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0490 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]491 // CK19-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]]492 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} %{{.*}}493 494 // CK19-USE: call void [[CALL08:@.+]](ptr {{[^,]+}})495 // CK19-NOUSE: call void [[CALL08:@.+]]()496 #pragma omp target map(tofrom:arra[ii])497 {498#ifdef USE499 arra[15]++;500#endif501 }502 503 // Map of a pointer.504 int *pa;505 506 // Region 09507 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])508 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2509 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]510 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3511 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]512 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]513 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]514 515 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0516 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0517 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]518 // CK19-DAG: store ptr [[VAR0]], ptr [[P0]]519 520 // CK19-USE: call void [[CALL09:@.+]](ptr {{[^,]+}})521 // CK19-NOUSE: call void [[CALL09:@.+]]()522 #pragma omp target map(from:pa)523 {524#ifdef USE525 pa[50]++;526#endif527 }528 529 // Region 10530 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])531 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2532 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]533 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3534 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]535 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]536 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]537 538 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0539 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0540 // CK19-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]541 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]542 // CK19-DAG: [[RVAR0]] = load ptr, ptr [[VAR0:%[^,]+]]543 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 20544 // CK19-DAG: [[RVAR00]] = load ptr, ptr [[VAR0]]545 546 // CK19-USE: call void [[CALL10:@.+]](ptr {{[^,]+}})547 // CK19-NOUSE: call void [[CALL10:@.+]]()548 #pragma omp target map(tofrom:pa[20:60])549 {550#ifdef USE551 pa[50]++;552#endif553 }554 555 // Region 11556 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])557 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2558 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]559 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3560 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]561 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]562 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]563 564 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0565 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0566 // CK19-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]567 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]568 // CK19-DAG: [[RVAR0]] = load ptr, ptr [[VAR0:%[^,]+]]569 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 0570 // CK19-DAG: [[RVAR00]] = load ptr, ptr [[VAR0]]571 572 // CK19-USE: call void [[CALL11:@.+]](ptr {{[^,]+}})573 // CK19-NOUSE: call void [[CALL11:@.+]]()574 #pragma omp target map(alloc:pa[:60])575 {576#ifdef USE577 pa[50]++;578#endif579 }580 581 // Region 12582 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])583 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2584 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]585 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3586 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]587 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]588 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]589 590 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0591 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0592 // CK19-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]593 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]594 // CK19-DAG: [[RVAR0]] = load ptr, ptr [[VAR0:%[^,]+]]595 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 15596 // CK19-DAG: [[RVAR00]] = load ptr, ptr [[VAR0]]597 598 // CK19-USE: call void [[CALL12:@.+]](ptr {{[^,]+}})599 // CK19-NOUSE: call void [[CALL12:@.+]]()600 #pragma omp target map(to:pa[15])601 {602#ifdef USE603 pa[15]++;604#endif605 }606 607 // Region 13608 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])609 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2610 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]611 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3612 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]613 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4614 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]615 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]616 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]617 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]618 619 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0620 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0621 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0622 // CK19-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]623 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]624 // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], ptr [[S0]]625 // CK19-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}626 // CK19-DAG: [[RVAR0]] = load ptr, ptr [[VAR0:%[^,]+]]627 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} %{{.*}}628 // CK19-DAG: [[RVAR00]] = load ptr, ptr [[VAR0]]629 630 // CK19-USE: call void [[CALL13:@.+]](ptr {{[^,]+}})631 // CK19-NOUSE: call void [[CALL13:@.+]]()632 #pragma omp target map(alloc:pa[ii-23:ii])633 {634#ifdef USE635 pa[50]++;636#endif637 }638 639 // Region 14640 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])641 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2642 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]643 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3644 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]645 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4646 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]647 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]648 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]649 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]650 651 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0652 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0653 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0654 // CK19-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]655 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]656 // CK19-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], ptr [[S0]]657 // CK19-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}658 // CK19-DAG: [[RVAR0]] = load ptr, ptr [[VAR0:%[^,]+]]659 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} 0660 // CK19-DAG: [[RVAR00]] = load ptr, ptr [[VAR0]]661 662 // CK19-USE: call void [[CALL14:@.+]](ptr {{[^,]+}})663 // CK19-NOUSE: call void [[CALL14:@.+]]()664 #pragma omp target map(to:pa[:ii])665 {666#ifdef USE667 pa[50]++;668#endif669 }670 671 // Region 15672 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])673 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2674 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]675 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3676 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]677 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]678 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]679 680 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0681 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0682 // CK19-DAG: store ptr [[RVAR0:%.+]], ptr [[BP0]]683 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]684 // CK19-DAG: [[RVAR0]] = load ptr, ptr [[VAR0:%[^,]+]]685 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[RVAR00:%.+]], i{{.+}} %{{.*}}686 // CK19-DAG: [[RVAR00]] = load ptr, ptr [[VAR0]]687 688 // CK19-USE: call void [[CALL15:@.+]](ptr {{[^,]+}})689 // CK19-NOUSE: call void [[CALL15:@.+]]()690 #pragma omp target map(from:pa[ii+12])691 {692#ifdef USE693 pa[15]++;694#endif695 }696 697 // Map of a variable-size array.698 int va[ii];699 700 // Region 16701 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])702 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2703 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]704 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3705 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]706 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4707 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]708 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]709 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]710 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]711 712 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0713 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0714 // CK19-USE-DAG: store i[[Z:64|32]] {{%.+}}, ptr [[BP0]]715 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[P0]]716 717 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1718 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1719 // CK19-USE-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1720 // CK19-USE-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]721 // CK19-USE-DAG: store ptr [[VAR1]], ptr [[P1]]722 // CK19-USE-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], ptr [[S1]]723 // CK19-USE-DAG: [[CSVAL1]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}724 725 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0726 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0727 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0728 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]729 // CK19-NOUSE-DAG: store ptr [[VAR0]], ptr [[P0]]730 // CK19-NOUSE-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], ptr [[S0]]731 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}732 733 // CK19-USE: call void [[CALL16:@.+]](i{{.+}} {{[^,]+}}, ptr {{[^,]+}})734 // CK19-NOUSE: call void [[CALL16:@.+]]()735 #pragma omp target map(to:va)736 {737#ifdef USE738 va[50]++;739#endif740 }741 742 // Region 17743 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])744 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2745 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]746 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3747 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]748 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]749 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]750 751 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0752 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0753 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[BP0]]754 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[P0]]755 756 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1757 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1758 // CK19-USE-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]759 // CK19-USE-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]760 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[VAR1]], i{{.+}} 20761 762 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0763 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0764 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]765 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]766 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 20767 768 // CK19-USE: call void [[CALL17:@.+]](i{{.+}} {{[^,]+}}, ptr {{[^,]+}})769 // CK19-NOUSE: call void [[CALL17:@.+]]()770 #pragma omp target map(from:va[20:60])771 {772#ifdef USE773 va[50]++;774#endif775 }776 777 // Region 18778 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])779 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2780 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]781 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3782 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]783 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]784 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]785 786 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0787 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0788 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[BP0]]789 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[P0]]790 791 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1792 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1793 // CK19-USE-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]794 // CK19-USE-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]795 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[VAR1]], i{{.+}} 0796 797 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0798 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0799 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]800 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]801 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0802 803 // CK19-USE: call void [[CALL18:@.+]](i{{.+}} {{[^,]+}}, ptr {{[^,]+}})804 // CK19-NOUSE: call void [[CALL18:@.+]]()805 #pragma omp target map(tofrom:va[:60])806 {807#ifdef USE808 va[50]++;809#endif810 }811 812 // Region 19813 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])814 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2815 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]816 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3817 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]818 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4819 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]820 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]821 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]822 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]823 824 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0825 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0826 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[BP0]]827 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[P0]]828 829 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1830 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1831 // CK19-USE-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1832 // CK19-USE-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]833 // CK19-USE-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]834 // CK19-USE-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], ptr [[S1]]835 // CK19-USE-DAG: [[CSVAL1]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}836 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[VAR1]], i{{.+}} 0837 838 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0839 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0840 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0841 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]842 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]843 // CK19-NOUSE-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], ptr [[S0]]844 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}845 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0846 847 // CK19-USE: call void [[CALL19:@.+]](i{{.+}} {{[^,]+}}, ptr {{[^,]+}})848 // CK19-NOUSE: call void [[CALL19:@.+]]()849 #pragma omp target map(alloc:va[:])850 {851#ifdef USE852 va[50]++;853#endif854 }855 856 // Region 20857 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])858 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2859 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]860 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3861 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]862 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]863 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]864 865 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0866 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0867 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[BP0]]868 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[P0]]869 870 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1871 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1872 // CK19-USE-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]873 // CK19-USE-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]874 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[VAR1]], i{{.+}} 15875 876 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0877 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0878 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]879 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]880 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 15881 882 // CK19-USE: call void [[CALL20:@.+]](i{{.+}} {{[^,]+}}, ptr {{[^,]+}})883 // CK19-NOUSE: call void [[CALL20:@.+]]()884 #pragma omp target map(to:va[15])885 {886#ifdef USE887 va[15]++;888#endif889 }890 891 // Region 21892 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])893 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2894 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]895 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3896 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]897 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 4898 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]899 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]900 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]901 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]902 903 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0904 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0905 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[BP0]]906 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[P0]]907 908 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1909 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1910 // CK19-USE-DAG: [[S1:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 1911 // CK19-USE-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]912 // CK19-USE-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]913 // CK19-USE-DAG: store i{{.+}} [[CSVAL1:%[^,]+]], ptr [[S1]]914 // CK19-USE-DAG: [[CSVAL1]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}915 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[VAR1]], i{{.+}} %{{.+}}916 917 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0918 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0919 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0920 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]921 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]922 // CK19-NOUSE-DAG: store i{{.+}} [[CSVAL0:%[^,]+]], ptr [[S0]]923 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %.*, 4|sext i32 .+ to i64}}924 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} %{{.+}}925 926 // CK19-USE: call void [[CALL21:@.+]](i{{.+}} {{[^,]+}}, ptr {{[^,]+}})927 // CK19-NOUSE: call void [[CALL21:@.+]]()928 #pragma omp target map(tofrom:va[ii:ii+23])929 {930#ifdef USE931 va[50]++;932#endif933 }934 935 // Region 22936 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])937 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2938 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]939 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3940 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]941 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]942 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]943 944 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0945 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0946 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[BP0]]947 // CK19-USE-DAG: store i[[Z]] {{%.+}}, ptr [[P0]]948 949 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1950 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 1951 // CK19-USE-DAG: store ptr [[VAR1:%.+]], ptr [[BP1]]952 // CK19-USE-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]953 // CK19-USE-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[VAR1]], i{{.+}} %{{.+}}954 955 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0956 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0957 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]958 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]959 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} %{{.+}}960 961 // CK19-USE: call void [[CALL22:@.+]](i{{.+}} {{[^,]+}}, ptr {{[^,]+}})962 // CK19-NOUSE: call void [[CALL22:@.+]]()963 #pragma omp target map(tofrom:va[ii])964 {965#ifdef USE966 va[15]++;967#endif968 }969 970 // Always.971 // Region 23972 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])973 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 2974 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]975 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 3976 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]977 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]978 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]979 980 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0981 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0982 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]983 // CK19-DAG: store ptr [[VAR0]], ptr [[P0]]984 985 // CK19-USE: call void [[CALL23:@.+]](ptr {{[^,]+}})986 // CK19-NOUSE: call void [[CALL23:@.+]]()987 #pragma omp target map(always, tofrom: a)988 {989#ifdef USE990 a++;991#endif992 }993 994 // Multidimensional arrays.995 int marr[4][5][6];996 int ***mptr;997 998 // Region 24999 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1000 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21001 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1002 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31003 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1004 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1005 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1006 1007 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01008 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01009 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1010 // CK19-DAG: store ptr [[VAR0]], ptr [[P0]]1011 1012 // CK19-USE: call void [[CALL24:@.+]](ptr {{[^,]+}})1013 // CK19-NOUSE: call void [[CALL24:@.+]]()1014 #pragma omp target map(tofrom: marr)1015 {1016#ifdef USE1017 marr[1][2][3]++;1018#endif1019 }1020 1021 // Region 251022 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1023 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21024 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1025 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31026 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1027 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1028 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1029 1030 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01031 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01032 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1033 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]1034 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 21035 // CK19-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:[^,]+]], i{{.+}} 0, i{{.+}} 21036 // CK19-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 11037 1038 // CK19-USE: call void [[CALL25:@.+]](ptr {{[^,]+}})1039 // CK19-NOUSE: call void [[CALL25:@.+]]()1040 #pragma omp target map(tofrom: marr[1][2][2:4])1041 {1042#ifdef USE1043 marr[1][2][3]++;1044#endif1045 }1046 1047 // Region 261048 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1049 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21050 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1051 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31052 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[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{{.+}} 01057 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01058 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1059 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]1060 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 01061 // CK19-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:[^,]+]], i{{.+}} 0, i{{.+}} 21062 // CK19-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 11063 1064 // CK19-USE: call void [[CALL26:@.+]](ptr {{[^,]+}})1065 // CK19-NOUSE: call void [[CALL26:@.+]]()1066 #pragma omp target map(tofrom: marr[1][2][:])1067 {1068#ifdef USE1069 marr[1][2][3]++;1070#endif1071 }1072 1073 // Region 271074 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1075 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21076 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1077 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31078 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1079 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1080 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1081 1082 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01083 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01084 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1085 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]1086 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:[^,]+]], i{{.+}} 0, i{{.+}} 31087 // CK19-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:[^,]+]], i{{.+}} 0, i{{.+}} 21088 // CK19-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 11089 1090 // CK19-USE: call void [[CALL27:@.+]](ptr {{[^,]+}})1091 // CK19-NOUSE: call void [[CALL27:@.+]]()1092 #pragma omp target map(tofrom: marr[1][2][3])1093 {1094#ifdef USE1095 marr[1][2][3]++;1096#endif1097 }1098 1099 // Region 281100 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1101 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21102 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1103 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31104 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1105 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1106 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1107 1108 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01109 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01110 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1111 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]1112 // CK19-DAG: [[VAR0]] = load ptr, ptr [[PTR:%[^,]+]],1113 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:[^,]+]], i{{.+}} 11114 // CK19-DAG: [[SEC00]] = load ptr, ptr [[PTR]],1115 1116 // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 11117 // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 11118 // CK19-DAG: store ptr [[SEC0]], ptr [[BP1]]1119 // CK19-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]1120 // CK19-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:[^,]+]], i{{.+}} 21121 // CK19-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],1122 // CK19-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[SEC1111:[^,]+]], i{{.+}} 11123 // CK19-DAG: [[SEC1111]] = load ptr, ptr [[PTR]],1124 1125 // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 21126 // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 21127 // CK19-DAG: store ptr [[SEC1]], ptr [[BP2]]1128 // CK19-DAG: store ptr [[SEC2:%.+]], ptr [[P2]]1129 // CK19-DAG: [[SEC2]] = getelementptr {{.*}}ptr [[SEC22:[^,]+]], i{{.+}} 21130 // CK19-DAG: [[SEC22]] = load ptr, ptr [[SEC222:%[^,]+]],1131 // CK19-DAG: [[SEC222]] = getelementptr {{.*}}ptr [[SEC2222:[^,]+]], i{{.+}} 21132 // CK19-DAG: [[SEC2222]] = load ptr, ptr [[SEC22222:%[^,]+]],1133 // CK19-DAG: [[SEC22222]] = getelementptr {{.*}}ptr [[SEC222222:[^,]+]], i{{.+}} 11134 // CK19-DAG: [[SEC222222]] = load ptr, ptr [[PTR]],1135 1136 // CK19-USE: call void [[CALL28:@.+]](ptr {{[^,]+}})1137 // CK19-NOUSE: call void [[CALL28:@.+]]()1138 #pragma omp target map(tofrom: mptr[1][2][2:4])1139 {1140#ifdef USE1141 mptr[1][2][3]++;1142#endif1143 }1144 1145 // Region 291146 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1147 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21148 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1149 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31150 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1151 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1152 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1153 1154 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01155 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01156 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1157 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]1158 // CK19-DAG: [[VAR0]] = load ptr, ptr [[PTR:%[^,]+]],1159 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:[^,]+]], i{{.+}} 11160 // CK19-DAG: [[SEC00]] = load ptr, ptr [[PTR]],1161 1162 // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 11163 // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 11164 // CK19-DAG: store ptr [[SEC0]], ptr [[BP1]]1165 // CK19-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]1166 // CK19-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:[^,]+]], i{{.+}} 21167 // CK19-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],1168 // CK19-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[SEC1111:[^,]+]], i{{.+}} 11169 // CK19-DAG: [[SEC1111]] = load ptr, ptr [[PTR]],1170 1171 // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 21172 // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 21173 // CK19-DAG: store ptr [[SEC1]], ptr [[BP2]]1174 // CK19-DAG: store ptr [[SEC2:%.+]], ptr [[P2]]1175 // CK19-DAG: [[SEC2]] = getelementptr {{.*}}ptr [[SEC22:[^,]+]], i{{.+}} 31176 // CK19-DAG: [[SEC22]] = load ptr, ptr [[SEC222:%[^,]+]],1177 // CK19-DAG: [[SEC222]] = getelementptr {{.*}}ptr [[SEC2222:[^,]+]], i{{.+}} 21178 // CK19-DAG: [[SEC2222]] = load ptr, ptr [[SEC22222:%[^,]+]],1179 // CK19-DAG: [[SEC22222]] = getelementptr {{.*}}ptr [[SEC222222:[^,]+]], i{{.+}} 11180 // CK19-DAG: [[SEC222222]] = load ptr, ptr [[PTR]],1181 1182 // CK19-USE: call void [[CALL29:@.+]](ptr {{[^,]+}})1183 // CK19-NOUSE: call void [[CALL29:@.+]]()1184 #pragma omp target map(tofrom: mptr[1][2][3])1185 {1186#ifdef USE1187 mptr[1][2][3]++;1188#endif1189 }1190 1191 // Multidimensional VLA.1192 double mva[23][ii][ii+5];1193 1194 // Region 301195 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1196 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21197 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1198 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31199 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1200 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 41201 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]1202 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1203 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1204 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]1205 //1206 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01207 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01208 // CK19-USE-DAG: store i[[Z]] 23, ptr [[BP0]]1209 // CK19-USE-DAG: store i[[Z]] 23, ptr [[P0]]1210 //1211 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 11212 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 11213 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], ptr [[BP1]]1214 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], ptr [[P1]]1215 // CK19-64-USE-DAG: [[VAR1]] = zext i32 %{{[^,]+}} to i641216 // CK19-64-USE-DAG: [[VAR11]] = zext i32 %{{[^,]+}} to i641217 //1218 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 21219 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 21220 // CK19-USE-DAG: store i[[Z]] [[VAR2:%.+]], ptr [[BP2]]1221 // CK19-USE-DAG: store i[[Z]] [[VAR22:%.+]], ptr [[P2]]1222 // CK19-64-USE-DAG: [[VAR2]] = zext i32 %{{[^,]+}} to i641223 // CK19-64-USE-DAG: [[VAR22]] = zext i32 %{{[^,]+}} to i641224 //1225 // CK19-USE-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 31226 // CK19-USE-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 31227 // CK19-USE-DAG: [[S3:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 31228 // CK19-USE-DAG: store ptr [[VAR3:%.+]], ptr [[BP3]]1229 // CK19-USE-DAG: store ptr [[VAR3]], ptr [[P3]]1230 // CK19-USE-DAG: store i64 [[CSVAL3:%[^,]+]], ptr [[S3]]1231 // CK19-USE-DAG: [[CSVAL3]] = {{mul nuw i64 %[^,]+, 8|sext i32 .+ to i64}}1232 1233 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01234 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01235 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 01236 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1237 // CK19-NOUSE-DAG: store ptr [[VAR0]], ptr [[P0]]1238 // CK19-NOUSE-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[S0]]1239 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 8|sext i32 .+ to i64}}1240 1241 // CK19-USE: call void [[CALL30:@.+]](i[[Z]] 23, i[[Z]] %{{[^,]+}}, i[[Z]] %{{[^,]+}}, ptr %{{[^,]+}})1242 // CK19-NOUSE: call void [[CALL30:@.+]]()1243 #pragma omp target map(tofrom: mva)1244 {1245#ifdef USE1246 mva[1][2][3]++;1247#endif1248 }1249 1250 // Region 311251 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1252 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21253 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1254 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31255 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1256 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1257 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1258 //1259 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01260 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01261 // CK19-USE-DAG: store i[[Z]] 23, ptr [[BP0]]1262 // CK19-USE-DAG: store i[[Z]] 23, ptr [[P0]]1263 //1264 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 11265 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 11266 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], ptr [[BP1]]1267 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], ptr [[P1]]1268 //1269 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 21270 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 21271 // CK19-USE-DAG: store i[[Z]] [[VAR2:%.+]], ptr [[BP2]]1272 // CK19-USE-DAG: store i[[Z]] [[VAR22:%.+]], ptr [[P2]]1273 //1274 // CK19-USE-DAG: [[BP3:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 31275 // CK19-USE-DAG: [[P3:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 31276 // CK19-USE-DAG: store ptr [[VAR3:%.+]], ptr [[BP3]]1277 // CK19-USE-DAG: store ptr [[SEC3:%.+]], ptr [[P3]]1278 // CK19-USE-DAG: [[SEC3]] = getelementptr {{.*}}ptr [[SEC33:%.+]], i[[Z]] 01279 // CK19-USE-DAG: [[SEC33]] = getelementptr {{.*}}ptr [[SEC333:%.+]], i[[Z]] [[IDX3:%.+]]1280 // CK19-USE-DAG: [[IDX3]] = mul nsw i[[Z]] %{{[^,]+}}, %{{[^,]+}}1281 // CK19-USE-DAG: [[SEC333]] = getelementptr {{.*}}ptr [[VAR3]], i[[Z]] [[IDX33:%.+]]1282 // CK19-USE-DAG: [[IDX33]] = mul nsw i[[Z]] 1, %{{[^,]+}}1283 1284 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01285 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01286 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1287 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]1288 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:%.+]], i[[Z:64|32]] 01289 // CK19-NOUSE-DAG: [[SEC00]] = getelementptr {{.*}}ptr [[SEC000:%.+]], i[[Z]] [[IDX0:%.+]]1290 // CK19-NOUSE-DAG: [[IDX0]] = mul nsw i[[Z]] %{{[^,]+}}, %{{[^,]+}}1291 // CK19-NOUSE-DAG: [[SEC000]] = getelementptr {{.*}}ptr [[VAR0]], i[[Z]] [[IDX00:%.+]]1292 // CK19-NOUSE-DAG: [[IDX00]] = mul nsw i[[Z]] 1, %{{[^,]+}}1293 1294 // CK19-USE: call void [[CALL31:@.+]](i[[Z]] 23, i[[Z]] %{{[^,]+}}, i[[Z]] %{{[^,]+}}, ptr %{{[^,]+}})1295 // CK19-NOUSE: call void [[CALL31:@.+]]()1296 #pragma omp target map(tofrom: mva[1][ii-2][:5])1297 {1298#ifdef USE1299 mva[1][2][3]++;1300#endif1301 }1302 1303 // Multidimensional array sections.1304 double marras[11][12][13];1305 double mvlaas[11][ii][13];1306 double ***mptras;1307 1308 // Region 321309 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1310 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21311 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1312 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31313 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1314 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1315 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1316 1317 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01318 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01319 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1320 // CK19-DAG: store ptr [[VAR0]], ptr [[P0]]1321 1322 // CK19-USE: call void [[CALL32:@.+]](ptr {{[^,]+}})1323 // CK19-NOUSE: call void [[CALL32:@.+]]()1324 #pragma omp target map(marras)1325 {1326#ifdef USE1327 marras[1][2][3]++;1328#endif1329 }1330 1331 // Region 331332 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1333 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21334 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1335 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31336 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1337 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1338 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1339 1340 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01341 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01342 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1343 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]1344 // CK19-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[VAR0]], i[[Z]] 0, i[[Z]] 01345 1346 // CK19-USE: call void [[CALL33:@.+]](ptr {{[^,]+}})1347 // CK19-NOUSE: call void [[CALL33:@.+]]()1348 #pragma omp target map(marras[:])1349 {1350#ifdef USE1351 marras[1][2][3]++;1352#endif1353 }1354 1355 // Region 341356 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1357 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21358 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1359 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31360 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1361 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1362 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1363 1364 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01365 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01366 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1367 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]1368 // CK19-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[VAR0]], i[[Z]] 0, i[[Z]] 01369 1370 // CK19-USE: call void [[CALL34:@.+]](ptr {{[^,]+}})1371 // CK19-NOUSE: call void [[CALL34:@.+]]()1372 #pragma omp target map(marras[:][:][:])1373 {1374#ifdef USE1375 marras[1][2][3]++;1376#endif1377 }1378 1379 // Region 351380 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1381 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21382 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1383 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31384 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1385 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 41386 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]1387 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1388 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1389 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]1390 //1391 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01392 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01393 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 01394 1395 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1396 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]1397 // CK19-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[S0]]1398 // CK19-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[SEC00:%[^,]+]], i[[Z]] 0, i[[Z]] 01399 // CK19-DAG: [[SEC00]] = getelementptr {{.+}}ptr [[VAR0]], i[[Z]] 0, i[[Z]] 11400 // CK19-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}}1401 1402 // CK19-USE: call void [[CALL35:@.+]](ptr {{[^,]+}})1403 // CK19-NOUSE: call void [[CALL35:@.+]]()1404 #pragma omp target map(marras[1][:ii][:])1405 {1406#ifdef USE1407 marras[1][2][3]++;1408#endif1409 }1410 1411 // Region 361412 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1413 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21414 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1415 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31416 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1417 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1418 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1419 1420 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01421 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01422 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1423 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]1424 // CK19-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[SEC00:%[^,]+]], i{{.+}} 01425 // CK19-DAG: [[SEC00]] = getelementptr {{.+}}ptr [[SEC000:%[^,]+]], i{{.+}} 0, i{{.+}} 01426 // CK19-DAG: [[SEC000]] = getelementptr {{.+}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 01427 1428 // CK19-USE: call void [[CALL36:@.+]](ptr {{[^,]+}})1429 // CK19-NOUSE: call void [[CALL36:@.+]]()1430 #pragma omp target map(marras[:1][:2][:13])1431 {1432#ifdef USE1433 marras[1][2][3]++;1434#endif1435 }1436 1437 // Region 371438 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1439 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21440 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1441 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31442 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1443 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 41444 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]1445 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1446 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1447 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]1448 //1449 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01450 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01451 // CK19-USE-DAG: store i[[Z]] 11, ptr [[BP0]]1452 // CK19-USE-DAG: store i[[Z]] 11, ptr [[P0]]1453 //1454 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 11455 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 11456 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], ptr [[BP1]]1457 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], ptr [[P1]]1458 //1459 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 21460 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 21461 // CK19-USE-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 21462 // CK19-USE-DAG: store ptr [[VAR2:%.+]], ptr [[BP2]]1463 // CK19-USE-DAG: store ptr [[VAR2]], ptr [[P2]]1464 // CK19-USE-DAG: store i64 [[CSVAL2:%[^,]+]], ptr [[S2]]1465 // CK19-USE-DAG: [[CSVAL2]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}}1466 1467 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01468 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01469 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 01470 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1471 // CK19-NOUSE-DAG: store ptr [[VAR0]], ptr [[P0]]1472 // CK19-NOUSE-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[S0]]1473 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}}1474 1475 // CK19-USE: call void [[CALL37:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, ptr %{{[^,]+}})1476 // CK19-NOUSE: call void [[CALL37:@.+]]()1477 #pragma omp target map(mvlaas)1478 {1479#ifdef USE1480 mvlaas[1][2][3]++;1481#endif1482 }1483 1484 // Region 381485 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1486 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21487 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1488 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31489 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1490 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 41491 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]1492 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1493 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1494 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]1495 //1496 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01497 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01498 // CK19-USE-DAG: store i[[Z]] 11, ptr [[BP0]]1499 // CK19-USE-DAG: store i[[Z]] 11, ptr [[P0]]1500 //1501 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 11502 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 11503 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], ptr [[BP1]]1504 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], ptr [[P1]]1505 //1506 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 21507 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 21508 // CK19-USE-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 21509 // CK19-USE-DAG: store ptr [[VAR2:%.+]], ptr [[BP2]]1510 // CK19-USE-DAG: store ptr [[SEC2:%.+]], ptr [[P2]]1511 // CK19-USE-DAG: store i64 [[CSVAL2:%[^,]+]], ptr [[S2]]1512 // CK19-USE-DAG: [[SEC2]] = getelementptr {{.+}}ptr [[VAR2]], i[[Z]] [[SEC22:%[^,]+]]1513 // CK19-USE-DAG: [[SEC22]] = mul nsw i[[Z]] 0, %{{[^,]+}}1514 // CK19-USE-DAG: [[CSVAL2]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}}1515 1516 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01517 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01518 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 01519 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1520 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]1521 // CK19-NOUSE-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[S0]]1522 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[VAR0]], i[[Z]] [[SEC00:%[^,]+]]1523 // CK19-NOUSE-DAG: [[SEC00]] = mul nsw i[[Z]] 0, %{{[^,]+}}1524 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}}1525 1526 // CK19-USE: call void [[CALL38:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, ptr %{{[^,]+}})1527 // CK19-NOUSE: call void [[CALL38:@.+]]()1528 #pragma omp target map(mvlaas[:])1529 {1530#ifdef USE1531 mvlaas[1][2][3]++;1532#endif1533 }1534 1535 // Region 391536 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1537 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21538 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1539 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31540 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1541 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 41542 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]1543 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1544 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1545 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]1546 //1547 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01548 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01549 // CK19-USE-DAG: store i[[Z]] 11, ptr [[BP0]]1550 // CK19-USE-DAG: store i[[Z]] 11, ptr [[P0]]1551 //1552 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 11553 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 11554 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], ptr [[BP1]]1555 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], ptr [[P1]]1556 //1557 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 21558 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 21559 // CK19-USE-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 21560 // CK19-USE-DAG: store ptr [[VAR2:%.+]], ptr [[BP2]]1561 // CK19-USE-DAG: store ptr [[SEC2:%.+]], ptr [[P2]]1562 // CK19-USE-DAG: store i64 [[CSVAL2:%[^,]+]], ptr [[S2]]1563 // CK19-USE-DAG: [[SEC2]] = getelementptr {{.+}}ptr [[VAR2]], i[[Z]] [[SEC22:%[^,]+]]1564 // CK19-USE-DAG: [[SEC22]] = mul nsw i[[Z]] 0, %{{[^,]+}}1565 // CK19-USE-DAG: [[CSVAL2]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}}1566 1567 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01568 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01569 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 01570 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1571 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]1572 // CK19-NOUSE-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[S0]]1573 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[VAR0]], i[[Z]] [[SEC00:%[^,]+]]1574 // CK19-NOUSE-DAG: [[SEC00]] = mul nsw i[[Z]] 0, %{{[^,]+}}1575 // CK19-NOUSE-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}}1576 1577 // CK19-USE: call void [[CALL39:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, ptr %{{[^,]+}})1578 // CK19-NOUSE: call void [[CALL39:@.+]]()1579 #pragma omp target map(mvlaas[:][:][:])1580 {1581#ifdef USE1582 mvlaas[1][2][3]++;1583#endif1584 }1585 1586 // Region 401587 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1588 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21589 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1590 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31591 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1592 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 41593 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]1594 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1595 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1596 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]1597 //1598 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01599 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01600 // CK19-USE-DAG: store i[[Z]] 11, ptr [[BP0]]1601 // CK19-USE-DAG: store i[[Z]] 11, ptr [[P0]]1602 //1603 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 11604 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 11605 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], ptr [[BP1]]1606 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], ptr [[P1]]1607 //1608 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 21609 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 21610 // CK19-USE-DAG: [[S2:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 21611 // CK19-USE-DAG: store ptr [[VAR2:%.+]], ptr [[BP2]]1612 // CK19-USE-DAG: store ptr [[SEC2:%.+]], ptr [[P2]]1613 // CK19-USE-DAG: store i64 [[CSVAL2:%[^,]+]], ptr [[S2]]1614 // CK19-USE-DAG: [[SEC2]] = getelementptr {{.+}}ptr [[SEC22:%[^,]+]], i[[Z]] 01615 // CK19-USE-DAG: [[SEC22]] = getelementptr {{.+}}ptr [[VAR2]], i[[Z]] [[SEC222:%[^,]+]]1616 // CK19-USE-DAG: [[SEC222]] = mul nsw i[[Z]] 1, %{{[^,]+}}1617 1618 // CK19-NOUSE-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]1619 // CK19-NOUSE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01620 // CK19-NOUSE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01621 // CK19-NOUSE-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 01622 // CK19-NOUSE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1623 // CK19-NOUSE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]1624 // CK19-NOUSE-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[S0]]1625 // CK19-NOUSE-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[SEC00:%[^,]+]], i[[Z]] 01626 // CK19-NOUSE-DAG: [[SEC00]] = getelementptr {{.+}}ptr [[VAR0]], i[[Z]] [[SEC000:%[^,]+]]1627 // CK19-NOUSE-DAG: [[SEC000]] = mul nsw i[[Z]] 1, %{{[^,]+}}1628 1629 // CK19-USE: call void [[CALL40:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, ptr %{{[^,]+}})1630 // CK19-NOUSE: call void [[CALL40:@.+]]()1631 #pragma omp target map(mvlaas[1][:ii][:])1632 {1633#ifdef USE1634 mvlaas[1][2][3]++;1635#endif1636 }1637 1638 // Region 411639 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1640 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21641 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1642 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31643 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1644 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1645 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1646 //1647 // CK19-USE-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01648 // CK19-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01649 // CK19-USE-DAG: store i[[Z]] 11, ptr [[BP0]]1650 // CK19-USE-DAG: store i[[Z]] 11, ptr [[P0]]1651 //1652 // CK19-USE-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 11653 // CK19-USE-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 11654 // CK19-USE-DAG: store i[[Z]] [[VAR1:%.+]], ptr [[BP1]]1655 // CK19-USE-DAG: store i[[Z]] [[VAR11:%.+]], ptr [[P1]]1656 //1657 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 21658 // CK19-USE-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 21659 // CK19-USE-DAG: store ptr [[VAR2:%.+]], ptr [[BP2]]1660 // CK19-USE-DAG: store ptr [[SEC2:%.+]], ptr [[P2]]1661 // CK19-USE-DAG: [[SEC2]] = getelementptr {{.+}}ptr [[SEC22:%[^,]+]], i[[Z]] 01662 // CK19-USE-DAG: [[SEC22]] = getelementptr {{.+}}ptr [[VAR2]], i[[Z]] [[SEC222:%[^,]+]]1663 // CK19-USE-DAG: [[SEC222]] = mul nsw i[[Z]] 0, %{{[^,]+}}1664 // CK19-USE-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 21665 1666 // CK19-NO-USE-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01667 // CK19-NO-USE-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1668 // CK19-NO-USE-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]1669 // CK19-NO-USE-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[SEC00:%[^,]+]], i[[Z]] 01670 // CK19-NO-USE-DAG: [[SEC00]] = getelementptr {{.+}}ptr [[VAR0]], i[[Z]] [[SEC000:%[^,]+]]1671 // CK19-NO-USE-DAG: [[SEC000]] = mul nsw i[[Z]] 0, %{{[^,]+}}1672 1673 // CK19-USE: call void [[CALL41:@.+]](i[[Z]] 11, i[[Z]] %{{[^,]+}}, ptr %{{[^,]+}})1674 // CK19-NOUSE: call void [[CALL41:@.+]]()1675 #pragma omp target map(mvlaas[:1][:2][:13])1676 {1677#ifdef USE1678 mvlaas[1][2][3]++;1679#endif1680 }1681 1682 // Region 421683 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1684 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21685 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1686 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31687 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1688 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1689 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1690 1691 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01692 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01693 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1694 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]1695 // CK19-DAG: [[VAR0]] = load ptr, ptr [[PTR:%[^,]+]],1696 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[SEC00:[^,]+]], i{{.+}} 01697 // CK19-DAG: [[SEC00]] = load ptr, ptr [[PTR]],1698 1699 // CK19-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 11700 // CK19-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 11701 // CK19-DAG: store ptr [[SEC0]], ptr [[BP1]]1702 // CK19-DAG: store ptr [[SEC1:%.+]], ptr [[P1]]1703 // CK19-DAG: [[SEC1]] = getelementptr {{.*}}ptr [[SEC11:[^,]+]], i{{.+}} 21704 // CK19-DAG: [[SEC11]] = load ptr, ptr [[SEC111:%[^,]+]],1705 // CK19-DAG: [[SEC111]] = getelementptr {{.*}}ptr [[SEC1111:[^,]+]], i{{.+}} 01706 // CK19-DAG: [[SEC1111]] = load ptr, ptr [[PTR]],1707 1708 // CK19-DAG: [[BP2:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 21709 // CK19-DAG: [[P2:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 21710 // CK19-DAG: store ptr [[SEC1]], ptr [[BP2]]1711 // CK19-DAG: store ptr [[SEC2:%.+]], ptr [[P2]]1712 // CK19-DAG: [[SEC2]] = getelementptr {{.*}}ptr [[SEC22:[^,]+]], i{{.+}} 01713 // CK19-DAG: [[SEC22]] = load ptr, ptr [[SEC222:%[^,]+]],1714 // CK19-DAG: [[SEC222]] = getelementptr {{.*}}ptr [[SEC2222:[^,]+]], i{{.+}} 21715 // CK19-DAG: [[SEC2222]] = load ptr, ptr [[SEC22222:%[^,]+]],1716 // CK19-DAG: [[SEC22222]] = getelementptr {{.*}}ptr [[SEC222222:[^,]+]], i{{.+}} 01717 // CK19-DAG: [[SEC222222]] = load ptr, ptr [[PTR]],1718 1719 // CK19-USE: call void [[CALL42:@.+]](ptr {{[^,]+}})1720 // CK19-NOUSE: call void [[CALL42:@.+]]()1721 #pragma omp target map(mptras[:1][2][:13])1722 {1723#ifdef USE1724 mptras[1][2][3]++;1725#endif1726 }1727 1728 // Region 43 - the memory is not contiguous for this map - will map the whole last dimension.1729 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1730 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21731 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1732 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31733 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1734 // CK19-DAG: [[SARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 41735 // CK19-DAG: store ptr [[SIZES:%.+]], ptr [[SARG]]1736 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1737 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1738 // CK19-DAG: [[SIZES]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]1739 //1740 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01741 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01742 // CK19-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 01743 1744 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1745 // CK19-DAG: store ptr [[SEC0:%.+]], ptr [[P0]]1746 // CK19-DAG: store i64 [[CSVAL0:%[^,]+]], ptr [[S0]]1747 // CK19-DAG: [[SEC0]] = getelementptr {{.+}}ptr [[SEC00:%[^,]+]], i[[Z]] 0, i[[Z]] 01748 // CK19-DAG: [[SEC00]] = getelementptr {{.+}}ptr [[VAR0]], i[[Z]] 0, i[[Z]] 11749 // CK19-DAG: [[CSVAL0]] = {{mul nuw i64 %[^,]+, 104|sext i32 .+ to i64}}1750 1751 // CK19-USE: call void [[CALL43:@.+]](ptr {{[^,]+}})1752 // CK19-NOUSE: call void [[CALL43:@.+]]()1753 #pragma omp target map(marras[1][:ii][1:])1754 {1755#ifdef USE1756 marras[1][2][3]++;1757#endif1758 }1759 1760 // Region 441761 // CK19-DAG: call i32 @__tgt_target_kernel(ptr @{{.+}}, i64 [[DEVICE:.+]], i32 -1, i32 0, ptr @.{{.+}}.region_id, ptr [[ARGS:%.+]])1762 // CK19-DAG: [[BPARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 21763 // CK19-DAG: store ptr [[BPGEP:%.+]], ptr [[BPARG]]1764 // CK19-DAG: [[PARG:%.+]] = getelementptr inbounds {{.+}}[[ARGS]], i32 0, i32 31765 // CK19-DAG: store ptr [[PGEP:%.+]], ptr [[PARG]]1766 // CK19-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]1767 // CK19-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]1768 1769 // CK19-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 01770 // CK19-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 01771 // CK19-DAG: store ptr [[VAR0:%.+]], ptr [[BP0]]1772 // CK19-DAG: store ptr [[SEC0:%[^,]+]], ptr [[P0]]1773 // CK19-DAG: [[SEC0]] = getelementptr {{.*}}ptr [[VAR0]], i{{.+}} 0, i{{.+}} 201774 1775 // CK19-USE: call void [[CALL44:@.+]](ptr {{[^,]+}})1776 // CK19-NOUSE: call void [[CALL44:@.+]]()1777 #pragma omp target map(from:arra[20:])1778 {1779#ifdef USE1780 arra[50]++;1781#endif1782 }1783 1784}1785 1786// CK19: define {{.+}}[[CALL00]]1787// CK19: define {{.+}}[[CALL01]]1788// CK19: define {{.+}}[[CALL02]]1789// CK19: define {{.+}}[[CALL03]]1790// CK19: define {{.+}}[[CALL04]]1791// CK19: define {{.+}}[[CALL05]]1792// CK19: define {{.+}}[[CALL06]]1793// CK19: define {{.+}}[[CALL07]]1794// CK19: define {{.+}}[[CALL08]]1795// CK19: define {{.+}}[[CALL09]]1796// CK19: define {{.+}}[[CALL10]]1797// CK19: define {{.+}}[[CALL11]]1798// CK19: define {{.+}}[[CALL12]]1799// CK19: define {{.+}}[[CALL13]]1800// CK19: define {{.+}}[[CALL14]]1801// CK19: define {{.+}}[[CALL15]]1802// CK19: define {{.+}}[[CALL16]]1803// CK19: define {{.+}}[[CALL17]]1804// CK19: define {{.+}}[[CALL18]]1805// CK19: define {{.+}}[[CALL19]]1806// CK19: define {{.+}}[[CALL20]]1807// CK19: define {{.+}}[[CALL21]]1808// CK19: define {{.+}}[[CALL22]]1809// CK19: define {{.+}}[[CALL23]]1810// CK19: define {{.+}}[[CALL24]]1811// CK19: define {{.+}}[[CALL25]]1812// CK19: define {{.+}}[[CALL26]]1813// CK19: define {{.+}}[[CALL27]]1814// CK19: define {{.+}}[[CALL28]]1815// CK19: define {{.+}}[[CALL29]]1816// CK19: define {{.+}}[[CALL30]]1817// CK19: define {{.+}}[[CALL31]]1818// CK19: define {{.+}}[[CALL32]]1819// CK19: define {{.+}}[[CALL33]]1820// CK19: define {{.+}}[[CALL34]]1821// CK19: define {{.+}}[[CALL35]]1822// CK19: define {{.+}}[[CALL36]]1823// CK19: define {{.+}}[[CALL37]]1824// CK19: define {{.+}}[[CALL38]]1825// CK19: define {{.+}}[[CALL39]]1826// CK19: define {{.+}}[[CALL40]]1827// CK19: define {{.+}}[[CALL41]]1828// CK19: define {{.+}}[[CALL42]]1829// CK19: define {{.+}}[[CALL43]]1830// CK19: define {{.+}}[[CALL44]]1831 1832#endif // CK191833#endif // HEADER_INC1834