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