Page MenuHomePhabricator
Differential D67833 Diff 278291 clang/test/OpenMP/target_exit_data_codegen.cpp

Changeset View

Standalone View

View Options

clang/test/OpenMP/target_exit_data_codegen.cpp

Show First 20 Lines • Show All 45 Lines • ▼ Show 20 Lines
// CK1-LABEL: _Z3fooi // CK1-LABEL: _Z3fooi
void foo(int arg) { void foo(int arg) {
int la; int la;
float lb[arg]; float lb[arg];
// Region 00 // Region 00
// CK1-NOT: __tgt_target_data_begin // CK1-NOT: __tgt_target_data_begin
// CK1-DAG: call void @__tgt_target_data_end_nowait(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}) // CK1-DAG: call void @__tgt_target_data_end_nowait_mapper(i64 [[DEV:%[^,]+]], i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE00]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK1-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64 // CK1-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
// CK1-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}}, // CK1-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK1-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [100 x double]** // CK1-DAG: [[BPC0:%.+]] = bitcast i8** [[BP0]] to [100 x double]**
Show All 9 Linesvoid foo(int arg) {
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
#pragma omp target exit data map(release: la) if(1+3-4) #pragma omp target exit data map(release: la) if(1+3-4)
{++arg;} {++arg;}
// Region 02 // Region 02
// CK1-NOT: __tgt_target_data_begin // CK1-NOT: __tgt_target_data_begin
// CK1: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]] // CK1: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
// CK1: [[IFTHEN]] // CK1: [[IFTHEN]]
// CK1-DAG: call void @__tgt_target_data_end(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}) // CK1-DAG: call void @__tgt_target_data_end_mapper(i64 4, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[SIZE02]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE02]]{{.+}}, i8** null)
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32** // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to i32**
// CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32** // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to i32**
// CK1-DAG: store i32* [[VAL0:%[^,]+]], i32** [[CBP0]] // CK1-DAG: store i32* [[VAL0:%[^,]+]], i32** [[CBP0]]
// CK1-DAG: store i32* [[VAL0]], i32** [[CP0]] // CK1-DAG: store i32* [[VAL0]], i32** [[CP0]]
// CK1: br label %[[IFEND:[^,]+]] // CK1: br label %[[IFEND:[^,]+]]
// CK1: [[IFELSE]] // CK1: [[IFELSE]]
// CK1: br label %[[IFEND]] // CK1: br label %[[IFEND]]
// CK1: [[IFEND]] // CK1: [[IFEND]]
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
#pragma omp target exit data map(release: arg) if(arg) device(4) #pragma omp target exit data map(release: arg) if(arg) device(4)
{++arg;} {++arg;}
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
{++arg;} {++arg;}
// Region 03 // Region 03
// CK1-NOT: __tgt_target_data_begin // CK1-NOT: __tgt_target_data_begin
// CK1-DAG: call void @__tgt_target_data_end(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}) // CK1-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** null)
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
// CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
// CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float** // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
// CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float** // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
// CK1-DAG: store float* [[VAL0:%[^,]+]], float** [[CBP0]] // CK1-DAG: store float* [[VAL0:%[^,]+]], float** [[CBP0]]
// CK1-DAG: store float* [[VAL0]], float** [[CP0]] // CK1-DAG: store float* [[VAL0]], float** [[CP0]]
// CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]] // CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
// CK1-64-DAG: [[CSVAL0]] = mul nuw i[[sz]] %{{[^,]+}}, 4 // CK1-64-DAG: [[CSVAL0]] = mul nuw i[[sz]] %{{[^,]+}}, 4
// CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64 // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
// CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4 // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
#pragma omp target exit data map(always, from: lb) #pragma omp target exit data map(always, from: lb)
{++arg;} {++arg;}
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
{++arg;} {++arg;}
// Region 04 // Region 04
// CK1-NOT: __tgt_target_data_begin // CK1-NOT: __tgt_target_data_begin
// CK1-DAG: call void @__tgt_target_data_end(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}) // CK1-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[SIZE04]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE04]]{{.+}}, i8** null)
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]** // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [[ST]]**
// CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double*** // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double***
// CK1-DAG: store [[ST]]* @gb, [[ST]]** [[CBP0]] // CK1-DAG: store [[ST]]* @gb, [[ST]]** [[CBP0]]
Show All 13 Linesvoid foo(int arg) {
#pragma omp target exit data map(release: gb.b[:3]) #pragma omp target exit data map(release: gb.b[:3])
{++arg;} {++arg;}
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
{++arg;} {++arg;}
// Region 05 // Region 05
// CK1-NOT: __tgt_target_data_begin // CK1-NOT: __tgt_target_data_begin
// CK1-DAG: call void @__tgt_target_data_end(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}) // CK1-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}, i8** null)
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
// CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
// CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float** // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
// CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float** // CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
// CK1-DAG: store float* [[VAL0:%[^,]+]], float** [[CBP0]] // CK1-DAG: store float* [[VAL0:%[^,]+]], float** [[CBP0]]
// CK1-DAG: store float* [[VAL0]], float** [[CP0]] // CK1-DAG: store float* [[VAL0]], float** [[CP0]]
// CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]] // CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
// CK1-64-DAG: [[CSVAL0]] = mul nuw i[[sz]] %{{[^,]+}}, 4 // CK1-64-DAG: [[CSVAL0]] = mul nuw i[[sz]] %{{[^,]+}}, 4
// CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64 // CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
// CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4 // CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
#pragma omp target exit data map(close, from: lb) #pragma omp target exit data map(close, from: lb)
{++arg;} {++arg;}
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1 // CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
{++arg;} {++arg;}
// Region 06 // Region 06
// CK1-NOT: __tgt_target_data_begin // CK1-NOT: __tgt_target_data_begin
// CK1-DAG: call void @__tgt_target_data_end(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}) // CK1-DAG: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}, i8** null)
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] // CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] // CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
// CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
// CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0 // CK1-DAG: [[S0:%.+]] = getelementptr inbounds {{.+}}[[S]], i{{.+}} 0, i{{.+}} 0
// CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float** // CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to float**
▲ Show 20 Lines • Show All 47 Lines • ▼ Show 20 Linesint bar(int arg){
ST<int> A; ST<int> A;
return A.foo(arg); return A.foo(arg);
} }
// Region 00 // Region 00
// CK2-NOT: __tgt_target_data_begin // CK2-NOT: __tgt_target_data_begin
// CK2: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]] // CK2: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
// CK2: [[IFTHEN]] // CK2: [[IFTHEN]]
// CK2-DAG: call void @__tgt_target_data_end(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:.+]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}) // CK2-DAG: call void @__tgt_target_data_end_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:.+]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK2-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64 // CK2-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
// CK2-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}}, // CK2-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
// CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK2-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] // CK2-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK2-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] // CK2-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
// CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK2-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK2-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
▲ Show 20 Lines • Show All 87 Lines • ▼ Show 20 Linesint bar(int arg){
STT<int> A; STT<int> A;
return A.foo(arg); return A.foo(arg);
} }
// Region 00 // Region 00
// CK4-NOT: __tgt_target_data_begin // CK4-NOT: __tgt_target_data_begin
// CK4: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]] // CK4: br i1 %{{[^,]+}}, label %[[IFTHEN:[^,]+]], label %[[IFELSE:[^,]+]]
// CK4: [[IFTHEN]] // CK4: [[IFTHEN]]
// CK4-DAG: call void @__tgt_target_data_end(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:.+]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}) // CK4-DAG: call void @__tgt_target_data_end_mapper(i64 [[DEV:%[^,]+]], i32 2, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz:.+]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[2 x i{{.+}}]* [[MTYPE00]]{{.+}}, i8** null)
// CK4-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64 // CK4-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
// CK4-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}}, // CK4-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
// CK4-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]] // CK4-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
// CK4-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]] // CK4-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
// CK4-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]] // CK4-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
// CK4-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0 // CK4-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
// CK4-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0 // CK4-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
Show All 26 Lines