1; NOTE: Assertions have been autogenerated by utils/update_test_checks.py UTC_ARGS: -p --function-signature --scrub-attributes 2; RUN: opt -S -passes=openmp-opt-cgscc -aa-pipeline=basic-aa -openmp-hide-memory-transfer-latency < %s | FileCheck %s 3target datalayout = "e-m:e-p270:32:32-p271:32:32-p272:64:64-i64:64-f80:128-n8:16:32:64-S128" 4 5; CHECK: %struct.__tgt_async_info = type { i8* } 6 7%struct.ident_t = type { i32, i32, i32, i32, i8* } 8%struct.__tgt_offload_entry = type { i8*, i8*, i64, i32, i32 } 9 10@.offload_maptypes = private unnamed_addr constant [1 x i64] [i64 35] 11@.__omp_offloading_heavyComputation1.region_id = weak constant i8 0 12@.offload_sizes.1 = private unnamed_addr constant [1 x i64] [i64 8] 13@.offload_maptypes.2 = private unnamed_addr constant [1 x i64] [i64 800] 14 15@.__omp_offloading_heavyComputation2.region_id = weak constant i8 0 16@.offload_maptypes.3 = private unnamed_addr constant [2 x i64] [i64 35, i64 35] 17 18@.__omp_offloading_heavyComputation3.region_id = weak constant i8 0 19@.offload_sizes.2 = private unnamed_addr constant [2 x i64] [i64 4, i64 0] 20@.offload_maptypes.4 = private unnamed_addr constant [2 x i64] [i64 800, i64 544] 21 22@.offload_maptypes.5 = private unnamed_addr constant [1 x i64] [i64 33] 23 24@0 = private unnamed_addr global %struct.ident_t { i32 0, i32 34, i32 0, i32 0, i8* getelementptr inbounds ([23 x i8], [23 x i8]* @.str0, i32 0, i32 0) }, align 8 25@.str0 = private unnamed_addr constant [23 x i8] c";unknown;unknown;0;0;;\00", align 1 26 27;double heavyComputation1() { 28; double a = rand() % 777; 29; double random = rand(); 30; 31; //#pragma omp target data map(a) 32; void* args[1]; 33; args[0] = &a; 34; __tgt_target_data_begin(..., args, ...) 35; 36; #pragma omp target teams 37; for (int i = 0; i < 1000; ++i) { 38; a *= i*i / 2; 39; } 40; 41; return random + a; 42;} 43define dso_local double @heavyComputation1() { 44; CHECK-LABEL: define {{[^@]+}}@heavyComputation1() { 45; CHECK-NEXT: entry: 46 47; CHECK-NEXT: %handle = alloca %struct.__tgt_async_info, align 8 48 49; CHECK-NEXT: %a = alloca double, align 8 50; CHECK-NEXT: %.offload_baseptrs = alloca [1 x i8*], align 8 51; CHECK-NEXT: %.offload_ptrs = alloca [1 x i8*], align 8 52; CHECK-NEXT: %.offload_baseptrs4 = alloca [1 x i8*], align 8 53; CHECK-NEXT: %.offload_ptrs5 = alloca [1 x i8*], align 8 54; CHECK-NEXT: %0 = bitcast double* %a to i8* 55; CHECK-NEXT: %call = tail call i32 (...) @rand() 56; CHECK-NEXT: %rem = srem i32 %call, 777 57; CHECK-NEXT: %conv = sitofp i32 %rem to double 58; CHECK-NEXT: store double %conv, double* %a, align 8 59; CHECK-NEXT: %call1 = tail call i32 (...) @rand() 60; CHECK-NEXT: %1 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_baseptrs, i64 0, i64 0 61; CHECK-NEXT: %2 = bitcast [1 x i8*]* %.offload_baseptrs to double** 62; CHECK-NEXT: store double* %a, double** %2, align 8 63; CHECK-NEXT: %3 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i64 0, i64 0 64; CHECK-NEXT: %4 = bitcast [1 x i8*]* %.offload_ptrs to double** 65; CHECK-NEXT: store double* %a, double** %4, align 8 66 67; CHECK-NEXT: call void @__tgt_target_data_begin_mapper_issue(%struct.ident_t* @0, i64 -1, i32 1, i8** %1, i8** %3, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i64 0, i64 0), i8** null, i8** null, %struct.__tgt_async_info* %handle) 68 69; CHECK-NEXT: %5 = bitcast double* %a to i64* 70 71; CHECK-NEXT: call void @__tgt_target_data_begin_mapper_wait(i64 -1, %struct.__tgt_async_info* %handle) 72 73; CHECK-NEXT: %6 = load i64, i64* %5, align 8 74; CHECK-NEXT: %7 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_baseptrs4, i64 0, i64 0 75; CHECK-NEXT: %8 = bitcast [1 x i8*]* %.offload_baseptrs4 to i64* 76; CHECK-NEXT: store i64 %6, i64* %8, align 8 77; CHECK-NEXT: %9 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs5, i64 0, i64 0 78; CHECK-NEXT: %10 = bitcast [1 x i8*]* %.offload_ptrs5 to i64* 79; CHECK-NEXT: store i64 %6, i64* %10, align 8 80; CHECK-NEXT: %11 = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @0, i64 -1, i8* nonnull @.__omp_offloading_heavyComputation1.region_id, i32 1, i8** nonnull %7, i8** nonnull %9, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.2, i64 0, i64 0), i8** null, i8** null, i32 0, i32 0) 81; CHECK-NEXT: %.not = icmp eq i32 %11, 0 82; CHECK-NEXT: br i1 %.not, label %omp_offload.cont, label %omp_offload.failed 83; CHECK: omp_offload.failed: 84; CHECK-NEXT: call void @heavyComputation1FallBack(i64 %6) 85; CHECK-NEXT: br label %omp_offload.cont 86; CHECK: omp_offload.cont: 87; CHECK-NEXT: %conv2 = sitofp i32 %call1 to double 88; CHECK-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @0, i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i64 0, i64 0), i8** null, i8** null) 89; CHECK-NEXT: %12 = load double, double* %a, align 8 90; CHECK-NEXT: %add = fadd double %12, %conv2 91; CHECK-NEXT: ret double %add 92; 93entry: 94 %a = alloca double, align 8 95 %.offload_baseptrs = alloca [1 x i8*], align 8 96 %.offload_ptrs = alloca [1 x i8*], align 8 97 %.offload_baseptrs4 = alloca [1 x i8*], align 8 98 %.offload_ptrs5 = alloca [1 x i8*], align 8 99 100 %0 = bitcast double* %a to i8* 101 %call = tail call i32 (...) @rand() 102 %rem = srem i32 %call, 777 103 %conv = sitofp i32 %rem to double 104 store double %conv, double* %a, align 8 105 106 ; FIXME: call to @__tgt_target_data_begin_mapper_issue(%struct.ident_t* @0, ...) should be moved here. 107 %call1 = tail call i32 (...) @rand() 108 109 %1 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_baseptrs, i64 0, i64 0 110 %2 = bitcast [1 x i8*]* %.offload_baseptrs to double** 111 store double* %a, double** %2, align 8 112 %3 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i64 0, i64 0 113 %4 = bitcast [1 x i8*]* %.offload_ptrs to double** 114 store double* %a, double** %4, align 8 115 call void @__tgt_target_data_begin_mapper(%struct.ident_t* @0, i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i64 0, i64 0), i8** null, i8** null) 116 117 %5 = bitcast double* %a to i64* 118 %6 = load i64, i64* %5, align 8 119 %7 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_baseptrs4, i64 0, i64 0 120 %8 = bitcast [1 x i8*]* %.offload_baseptrs4 to i64* 121 store i64 %6, i64* %8, align 8 122 %9 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs5, i64 0, i64 0 123 %10 = bitcast [1 x i8*]* %.offload_ptrs5 to i64* 124 store i64 %6, i64* %10, align 8 125 126 ; FIXME: call to @__tgt_target_data_begin_mapper_wait(...) should be moved here. 127 %11 = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @0, i64 -1, i8* nonnull @.__omp_offloading_heavyComputation1.region_id, i32 1, i8** nonnull %7, i8** nonnull %9, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.2, i64 0, i64 0), i8** null, i8** null, i32 0, i32 0) 128 %.not = icmp eq i32 %11, 0 129 br i1 %.not, label %omp_offload.cont, label %omp_offload.failed 130 131omp_offload.failed: ; preds = %entry 132 call void @heavyComputation1FallBack(i64 %6) 133 br label %omp_offload.cont 134 135omp_offload.cont: ; preds = %omp_offload.failed, %entry 136 %conv2 = sitofp i32 %call1 to double 137 call void @__tgt_target_data_end_mapper(%struct.ident_t* @0, i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_sizes.1, i64 0, i64 0), i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes, i64 0, i64 0), i8** null, i8** null) 138 %12 = load double, double* %a, align 8 139 %add = fadd double %12, %conv2 140 ret double %add 141} 142 143define internal void @heavyComputation1FallBack(i64 %a) { 144entry: 145 ; Fallback for offloading function heavyComputation1. 146 ret void 147} 148 149;int heavyComputation2(double* a, unsigned size) { 150; int random = rand() % 7; 151; 152; //#pragma omp target data map(a[0:size], size) 153; void* args[2]; 154; args[0] = &a; 155; args[1] = &size; 156; __tgt_target_data_begin(..., args, ...) 157; 158; #pragma omp target teams 159; for (int i = 0; i < size; ++i) { 160; a[i] = ++a[i] * 3.141624; 161; } 162; 163; return random; 164;} 165define dso_local i32 @heavyComputation2(double* %a, i32 %size) { 166; CHECK-LABEL: define {{[^@]+}}@heavyComputation2(double* %a, i32 %size) { 167; CHECK-NEXT: entry: 168; CHECK-NEXT: %size.addr = alloca i32, align 4 169; CHECK-NEXT: %.offload_baseptrs = alloca [2 x i8*], align 8 170; CHECK-NEXT: %.offload_ptrs = alloca [2 x i8*], align 8 171; CHECK-NEXT: %.offload_sizes = alloca [2 x i64], align 8 172; CHECK-NEXT: %.offload_baseptrs2 = alloca [2 x i8*], align 8 173; CHECK-NEXT: %.offload_ptrs3 = alloca [2 x i8*], align 8 174; CHECK-NEXT: store i32 %size, i32* %size.addr, align 4 175; CHECK-NEXT: %call = tail call i32 (...) @rand() 176; CHECK-NEXT: %conv = zext i32 %size to i64 177; CHECK-NEXT: %0 = shl nuw nsw i64 %conv, 3 178; CHECK-NEXT: %1 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs, i64 0, i64 0 179; CHECK-NEXT: %2 = bitcast [2 x i8*]* %.offload_baseptrs to double** 180; CHECK-NEXT: store double* %a, double** %2, align 8 181; CHECK-NEXT: %3 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i64 0, i64 0 182; CHECK-NEXT: %4 = bitcast [2 x i8*]* %.offload_ptrs to double** 183; CHECK-NEXT: store double* %a, double** %4, align 8 184; CHECK-NEXT: %5 = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i64 0, i64 0 185; CHECK-NEXT: store i64 %0, i64* %5, align 8 186; CHECK-NEXT: %6 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs, i64 0, i64 1 187; CHECK-NEXT: %7 = bitcast i8** %6 to i32** 188; CHECK-NEXT: store i32* %size.addr, i32** %7, align 8 189; CHECK-NEXT: %8 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i64 0, i64 1 190; CHECK-NEXT: %9 = bitcast i8** %8 to i32** 191; CHECK-NEXT: store i32* %size.addr, i32** %9, align 8 192; CHECK-NEXT: %10 = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i64 0, i64 1 193; CHECK-NEXT: store i64 4, i64* %10, align 8 194 195; CHECK-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @0, i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0), i8** null, i8** null) 196 197; CHECK-NEXT: %11 = load i32, i32* %size.addr, align 4 198; CHECK-NEXT: %size.casted = zext i32 %11 to i64 199; CHECK-NEXT: %12 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs2, i64 0, i64 0 200; CHECK-NEXT: %13 = bitcast [2 x i8*]* %.offload_baseptrs2 to i64* 201; CHECK-NEXT: store i64 %size.casted, i64* %13, align 8 202; CHECK-NEXT: %14 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs3, i64 0, i64 0 203; CHECK-NEXT: %15 = bitcast [2 x i8*]* %.offload_ptrs3 to i64* 204; CHECK-NEXT: store i64 %size.casted, i64* %15, align 8 205; CHECK-NEXT: %16 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs2, i64 0, i64 1 206; CHECK-NEXT: %17 = bitcast i8** %16 to double** 207; CHECK-NEXT: store double* %a, double** %17, align 8 208; CHECK-NEXT: %18 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs3, i64 0, i64 1 209; CHECK-NEXT: %19 = bitcast i8** %18 to double** 210; CHECK-NEXT: store double* %a, double** %19, align 8 211; CHECK-NEXT: %20 = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @0, i64 -1, i8* nonnull @.__omp_offloading_heavyComputation2.region_id, i32 2, i8** nonnull %12, i8** nonnull %14, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.2, i64 0, i64 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.4, i64 0, i64 0), i8** null, i8** null, i32 0, i32 0) 212; CHECK-NEXT: %.not = icmp eq i32 %20, 0 213; CHECK-NEXT: br i1 %.not, label %omp_offload.cont, label %omp_offload.failed 214; CHECK: omp_offload.failed: 215; CHECK-NEXT: call void @heavyComputation2FallBack(i64 %size.casted, double* %a) 216; CHECK-NEXT: br label %omp_offload.cont 217; CHECK: omp_offload.cont: 218; CHECK-NEXT: %rem = srem i32 %call, 7 219; CHECK-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @0, i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0), i8** null, i8** null) 220; CHECK-NEXT: ret i32 %rem 221; 222entry: 223 %size.addr = alloca i32, align 4 224 %.offload_baseptrs = alloca [2 x i8*], align 8 225 %.offload_ptrs = alloca [2 x i8*], align 8 226 %.offload_sizes = alloca [2 x i64], align 8 227 %.offload_baseptrs2 = alloca [2 x i8*], align 8 228 %.offload_ptrs3 = alloca [2 x i8*], align 8 229 230 store i32 %size, i32* %size.addr, align 4 231 %call = tail call i32 (...) @rand() 232 233 %conv = zext i32 %size to i64 234 %0 = shl nuw nsw i64 %conv, 3 235 %1 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs, i64 0, i64 0 236 %2 = bitcast [2 x i8*]* %.offload_baseptrs to double** 237 store double* %a, double** %2, align 8 238 %3 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i64 0, i64 0 239 %4 = bitcast [2 x i8*]* %.offload_ptrs to double** 240 store double* %a, double** %4, align 8 241 %5 = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i64 0, i64 0 242 store i64 %0, i64* %5, align 8 243 %6 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs, i64 0, i64 1 244 %7 = bitcast i8** %6 to i32** 245 store i32* %size.addr, i32** %7, align 8 246 %8 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i64 0, i64 1 247 %9 = bitcast i8** %8 to i32** 248 store i32* %size.addr, i32** %9, align 8 249 %10 = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i64 0, i64 1 250 store i64 4, i64* %10, align 8 251 call void @__tgt_target_data_begin_mapper(%struct.ident_t* @0, i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0), i8** null, i8** null) 252 253 %11 = load i32, i32* %size.addr, align 4 254 %size.casted = zext i32 %11 to i64 255 %12 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs2, i64 0, i64 0 256 %13 = bitcast [2 x i8*]* %.offload_baseptrs2 to i64* 257 store i64 %size.casted, i64* %13, align 8 258 %14 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs3, i64 0, i64 0 259 %15 = bitcast [2 x i8*]* %.offload_ptrs3 to i64* 260 store i64 %size.casted, i64* %15, align 8 261 %16 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs2, i64 0, i64 1 262 %17 = bitcast i8** %16 to double** 263 store double* %a, double** %17, align 8 264 %18 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs3, i64 0, i64 1 265 %19 = bitcast i8** %18 to double** 266 store double* %a, double** %19, align 8 267 268 ; FIXME: call to @__tgt_target_data_begin_mapper_wait(...) should be moved here. 269 %20 = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @0, i64 -1, i8* nonnull @.__omp_offloading_heavyComputation2.region_id, i32 2, i8** nonnull %12, i8** nonnull %14, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.2, i64 0, i64 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.4, i64 0, i64 0), i8** null, i8** null, i32 0, i32 0) 270 %.not = icmp eq i32 %20, 0 271 br i1 %.not, label %omp_offload.cont, label %omp_offload.failed 272 273omp_offload.failed: ; preds = %entry 274 call void @heavyComputation2FallBack(i64 %size.casted, double* %a) 275 br label %omp_offload.cont 276 277omp_offload.cont: ; preds = %omp_offload.failed, %entry 278 %rem = srem i32 %call, 7 279 call void @__tgt_target_data_end_mapper(%struct.ident_t* @0, i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0), i8** null, i8** null) 280 ret i32 %rem 281} 282 283define internal void @heavyComputation2FallBack(i64 %size, double* %a) { 284entry: 285 ; Fallback for offloading function heavyComputation2. 286 ret void 287} 288 289;int heavyComputation3(double* restrict a, unsigned size) { 290; int random = rand() % 7; 291; 292; //#pragma omp target data map(a[0:size], size) 293; void* args[2]; 294; args[0] = &a; 295; args[1] = &size; 296; __tgt_target_data_begin(..., args, ...) 297; 298; #pragma omp target teams 299; for (int i = 0; i < size; ++i) { 300; a[i] = ++a[i] * 3.141624; 301; } 302; 303; return random; 304;} 305define dso_local i32 @heavyComputation3(double* noalias %a, i32 %size) { 306; CHECK-LABEL: define {{[^@]+}}@heavyComputation3(double* noalias %a, i32 %size) { 307; CHECK-NEXT: entry: 308; CHECK-NEXT: %size.addr = alloca i32, align 4 309; CHECK-NEXT: %.offload_baseptrs = alloca [2 x i8*], align 8 310; CHECK-NEXT: %.offload_ptrs = alloca [2 x i8*], align 8 311; CHECK-NEXT: %.offload_sizes = alloca [2 x i64], align 8 312; CHECK-NEXT: %.offload_baseptrs2 = alloca [2 x i8*], align 8 313; CHECK-NEXT: %.offload_ptrs3 = alloca [2 x i8*], align 8 314; CHECK-NEXT: store i32 %size, i32* %size.addr, align 4 315; CHECK-NEXT: %call = tail call i32 (...) @rand() 316; CHECK-NEXT: %conv = zext i32 %size to i64 317; CHECK-NEXT: %0 = shl nuw nsw i64 %conv, 3 318; CHECK-NEXT: %1 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs, i64 0, i64 0 319; CHECK-NEXT: %2 = bitcast [2 x i8*]* %.offload_baseptrs to double** 320; CHECK-NEXT: store double* %a, double** %2, align 8 321; CHECK-NEXT: %3 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i64 0, i64 0 322; CHECK-NEXT: %4 = bitcast [2 x i8*]* %.offload_ptrs to double** 323; CHECK-NEXT: store double* %a, double** %4, align 8 324; CHECK-NEXT: %5 = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i64 0, i64 0 325; CHECK-NEXT: store i64 %0, i64* %5, align 8 326; CHECK-NEXT: %6 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs, i64 0, i64 1 327; CHECK-NEXT: %7 = bitcast i8** %6 to i32** 328; CHECK-NEXT: store i32* %size.addr, i32** %7, align 8 329; CHECK-NEXT: %8 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i64 0, i64 1 330; CHECK-NEXT: %9 = bitcast i8** %8 to i32** 331; CHECK-NEXT: store i32* %size.addr, i32** %9, align 8 332; CHECK-NEXT: %10 = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i64 0, i64 1 333; CHECK-NEXT: store i64 4, i64* %10, align 8 334 335; CHECK-NEXT: call void @__tgt_target_data_begin_mapper(%struct.ident_t* @0, i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0), i8** null, i8** null) 336 337; CHECK-NEXT: %11 = load i32, i32* %size.addr, align 4 338; CHECK-NEXT: %size.casted = zext i32 %11 to i64 339; CHECK-NEXT: %12 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs2, i64 0, i64 0 340; CHECK-NEXT: %13 = bitcast [2 x i8*]* %.offload_baseptrs2 to i64* 341; CHECK-NEXT: store i64 %size.casted, i64* %13, align 8 342; CHECK-NEXT: %14 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs3, i64 0, i64 0 343; CHECK-NEXT: %15 = bitcast [2 x i8*]* %.offload_ptrs3 to i64* 344; CHECK-NEXT: store i64 %size.casted, i64* %15, align 8 345; CHECK-NEXT: %16 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs2, i64 0, i64 1 346; CHECK-NEXT: %17 = bitcast i8** %16 to double** 347; CHECK-NEXT: store double* %a, double** %17, align 8 348; CHECK-NEXT: %18 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs3, i64 0, i64 1 349; CHECK-NEXT: %19 = bitcast i8** %18 to double** 350; CHECK-NEXT: store double* %a, double** %19, align 8 351; CHECK-NEXT: %20 = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @0, i64 -1, i8* nonnull @.__omp_offloading_heavyComputation3.region_id, i32 2, i8** nonnull %12, i8** nonnull %14, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.2, i64 0, i64 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.4, i64 0, i64 0), i8** null, i8** null, i32 0, i32 0) 352; CHECK-NEXT: %.not = icmp eq i32 %20, 0 353; CHECK-NEXT: br i1 %.not, label %omp_offload.cont, label %omp_offload.failed 354; CHECK: omp_offload.failed: 355; CHECK-NEXT: call void @heavyComputation3FallBack(i64 %size.casted, double* %a) 356; CHECK-NEXT: br label %omp_offload.cont 357; CHECK: omp_offload.cont: 358; CHECK-NEXT: %rem = srem i32 %call, 7 359; CHECK-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @0, i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0), i8** null, i8** null) 360; CHECK-NEXT: ret i32 %rem 361; 362entry: 363 %size.addr = alloca i32, align 4 364 %.offload_baseptrs = alloca [2 x i8*], align 8 365 %.offload_ptrs = alloca [2 x i8*], align 8 366 %.offload_sizes = alloca [2 x i64], align 8 367 %.offload_baseptrs2 = alloca [2 x i8*], align 8 368 %.offload_ptrs3 = alloca [2 x i8*], align 8 369 store i32 %size, i32* %size.addr, align 4 370 371 ; FIXME: call to @__tgt_target_data_begin_mapper_issue(%struct.ident_t* @0, ...) should be moved here. 372 %call = tail call i32 (...) @rand() 373 374 %conv = zext i32 %size to i64 375 %0 = shl nuw nsw i64 %conv, 3 376 %1 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs, i64 0, i64 0 377 %2 = bitcast [2 x i8*]* %.offload_baseptrs to double** 378 store double* %a, double** %2, align 8 379 %3 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i64 0, i64 0 380 %4 = bitcast [2 x i8*]* %.offload_ptrs to double** 381 store double* %a, double** %4, align 8 382 %5 = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i64 0, i64 0 383 store i64 %0, i64* %5, align 8 384 %6 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs, i64 0, i64 1 385 %7 = bitcast i8** %6 to i32** 386 store i32* %size.addr, i32** %7, align 8 387 %8 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs, i64 0, i64 1 388 %9 = bitcast i8** %8 to i32** 389 store i32* %size.addr, i32** %9, align 8 390 %10 = getelementptr inbounds [2 x i64], [2 x i64]* %.offload_sizes, i64 0, i64 1 391 store i64 4, i64* %10, align 8 392 call void @__tgt_target_data_begin_mapper(%struct.ident_t* @0, i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0), i8** null, i8** null) 393 394 %11 = load i32, i32* %size.addr, align 4 395 %size.casted = zext i32 %11 to i64 396 %12 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs2, i64 0, i64 0 397 %13 = bitcast [2 x i8*]* %.offload_baseptrs2 to i64* 398 store i64 %size.casted, i64* %13, align 8 399 %14 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs3, i64 0, i64 0 400 %15 = bitcast [2 x i8*]* %.offload_ptrs3 to i64* 401 store i64 %size.casted, i64* %15, align 8 402 %16 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_baseptrs2, i64 0, i64 1 403 %17 = bitcast i8** %16 to double** 404 store double* %a, double** %17, align 8 405 %18 = getelementptr inbounds [2 x i8*], [2 x i8*]* %.offload_ptrs3, i64 0, i64 1 406 %19 = bitcast i8** %18 to double** 407 store double* %a, double** %19, align 8 408 409 ; FIXME: call to @__tgt_target_data_begin_mapper_wait(...) should be moved here. 410 %20 = call i32 @__tgt_target_teams_mapper(%struct.ident_t* @0, i64 -1, i8* nonnull @.__omp_offloading_heavyComputation3.region_id, i32 2, i8** nonnull %12, i8** nonnull %14, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_sizes.2, i64 0, i64 0), i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.4, i64 0, i64 0), i8** null, i8** null, i32 0, i32 0) 411 %.not = icmp eq i32 %20, 0 412 br i1 %.not, label %omp_offload.cont, label %omp_offload.failed 413 414omp_offload.failed: ; preds = %entry 415 call void @heavyComputation3FallBack(i64 %size.casted, double* %a) 416 br label %omp_offload.cont 417 418omp_offload.cont: ; preds = %omp_offload.failed, %entry 419 %rem = srem i32 %call, 7 420 call void @__tgt_target_data_end_mapper(%struct.ident_t* @0, i64 -1, i32 2, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([2 x i64], [2 x i64]* @.offload_maptypes.3, i64 0, i64 0), i8** null, i8** null) 421 ret i32 %rem 422} 423 424define internal void @heavyComputation3FallBack(i64 %size, double* %a) { 425entry: 426 ; Fallback for offloading function heavyComputation3. 427 ret void 428} 429 430;int dataTransferOnly1(double* restrict a, unsigned size) { 431; // Random computation. 432; int random = rand(); 433; 434; //#pragma omp target data map(to:a[0:size]) 435; void* args[1]; 436; args[0] = &a; 437; __tgt_target_data_begin(..., args, ...) 438; 439; // Random computation. 440; random %= size; 441; return random; 442;} 443define dso_local i32 @dataTransferOnly1(double* noalias %a, i32 %size) { 444; CHECK-LABEL: define {{[^@]+}}@dataTransferOnly1(double* noalias %a, i32 %size) { 445; CHECK-NEXT: entry: 446 447; CHECK-NEXT: %handle = alloca %struct.__tgt_async_info, align 8 448 449; CHECK-NEXT: %.offload_baseptrs = alloca [1 x i8*], align 8 450; CHECK-NEXT: %.offload_ptrs = alloca [1 x i8*], align 8 451; CHECK-NEXT: %.offload_sizes = alloca [1 x i64], align 8 452; CHECK-NEXT: %call = tail call i32 (...) @rand() 453; CHECK-NEXT: %conv = zext i32 %size to i64 454; CHECK-NEXT: %0 = shl nuw nsw i64 %conv, 3 455; CHECK-NEXT: %1 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_baseptrs, i64 0, i64 0 456; CHECK-NEXT: %2 = bitcast [1 x i8*]* %.offload_baseptrs to double** 457; CHECK-NEXT: store double* %a, double** %2, align 8 458; CHECK-NEXT: %3 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i64 0, i64 0 459; CHECK-NEXT: %4 = bitcast [1 x i8*]* %.offload_ptrs to double** 460; CHECK-NEXT: store double* %a, double** %4, align 8 461; CHECK-NEXT: %5 = getelementptr inbounds [1 x i64], [1 x i64]* %.offload_sizes, i64 0, i64 0 462; CHECK-NEXT: store i64 %0, i64* %5, align 8 463 464; CHECK-NEXT: call void @__tgt_target_data_begin_mapper_issue(%struct.ident_t* @0, i64 -1, i32 1, i8** %1, i8** %3, i64* %5, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.5, i64 0, i64 0), i8** null, i8** null, %struct.__tgt_async_info* %handle) 465 466; CHECK-NEXT: %rem = urem i32 %call, %size 467 468; CHECK-NEXT: call void @__tgt_target_data_begin_mapper_wait(i64 -1, %struct.__tgt_async_info* %handle) 469 470; CHECK-NEXT: call void @__tgt_target_data_end_mapper(%struct.ident_t* @0, i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.5, i64 0, i64 0), i8** null, i8** null) 471; CHECK-NEXT: ret i32 %rem 472; 473entry: 474 %.offload_baseptrs = alloca [1 x i8*], align 8 475 %.offload_ptrs = alloca [1 x i8*], align 8 476 %.offload_sizes = alloca [1 x i64], align 8 477 478 ; FIXME: call to @__tgt_target_data_begin_issue_mapper(...) should be moved here. 479 %call = tail call i32 (...) @rand() 480 481 %conv = zext i32 %size to i64 482 %0 = shl nuw nsw i64 %conv, 3 483 %1 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_baseptrs, i64 0, i64 0 484 %2 = bitcast [1 x i8*]* %.offload_baseptrs to double** 485 store double* %a, double** %2, align 8 486 %3 = getelementptr inbounds [1 x i8*], [1 x i8*]* %.offload_ptrs, i64 0, i64 0 487 %4 = bitcast [1 x i8*]* %.offload_ptrs to double** 488 store double* %a, double** %4, align 8 489 %5 = getelementptr inbounds [1 x i64], [1 x i64]* %.offload_sizes, i64 0, i64 0 490 store i64 %0, i64* %5, align 8 491 call void @__tgt_target_data_begin_mapper(%struct.ident_t* @0, i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.5, i64 0, i64 0), i8** null, i8** null) 492 493 %rem = urem i32 %call, %size 494 495 call void @__tgt_target_data_end_mapper(%struct.ident_t* @0, i64 -1, i32 1, i8** nonnull %1, i8** nonnull %3, i64* nonnull %5, i64* getelementptr inbounds ([1 x i64], [1 x i64]* @.offload_maptypes.5, i64 0, i64 0), i8** null, i8** null) 496 ret i32 %rem 497} 498 499declare void @__tgt_target_data_begin_mapper(%struct.ident_t*, i64, i32, i8**, i8**, i64*, i64*, i8**, i8**) 500declare i32 @__tgt_target_teams_mapper(%struct.ident_t*, i64, i8*, i32, i8**, i8**, i64*, i64*, i8**, i8**, i32, i32) 501declare void @__tgt_target_data_end_mapper(%struct.ident_t*, i64, i32, i8**, i8**, i64*, i64*, i8**, i8**) 502 503declare dso_local i32 @rand(...) 504 505; CHECK: declare void @__tgt_target_data_begin_mapper_issue(%struct.ident_t*, i64, i32, i8**, i8**, i64*, i64*, i8**, i8**, %struct.__tgt_async_info*) 506; CHECK: declare void @__tgt_target_data_begin_mapper_wait(i64, %struct.__tgt_async_info*) 507