forked from OSchip/llvm-project
[Clang][OpenMP] Added the support for target data nowait
Previously we added support for target nowait, but target data nowait has not been supported yet. In this patch, target data nowait will also be wrapped into a task. Reviewed By: jdoerfert Differential Revision: https://reviews.llvm.org/D90099
This commit is contained in:
Shilei Tian
parent
ba78cae20f
commit
0661328d7e
|
|
@ -10513,7 +10513,8 @@ void CGOpenMPRuntime::emitTargetDataStandAloneCall(
|
|||
TargetDataInfo Info;
|
||||
// Fill up the arrays and create the arguments.
|
||||
emitOffloadingArrays(CGF, CombinedInfo, Info);
|
||||
bool HasDependClauses = D.hasClausesOfKind<OMPDependClause>();
|
||||
bool RequiresOuterTask = D.hasClausesOfKind<OMPDependClause>() ||
|
||||
D.hasClausesOfKind<OMPNowaitClause>();
|
||||
emitOffloadingArraysArgument(
|
||||
CGF, Info.BasePointersArray, Info.PointersArray, Info.SizesArray,
|
||||
Info.MapTypesArray, Info.MappersArray, Info, {/*ForEndTask=*/false});
|
||||
|
|
@ -10526,7 +10527,7 @@ void CGOpenMPRuntime::emitTargetDataStandAloneCall(
|
|||
Address(Info.SizesArray, CGM.getPointerAlign());
|
||||
InputInfo.MappersArray = Address(Info.MappersArray, CGM.getPointerAlign());
|
||||
MapTypesArray = Info.MapTypesArray;
|
||||
if (HasDependClauses)
|
||||
if (RequiresOuterTask)
|
||||
CGF.EmitOMPTargetTaskBasedDirective(D, ThenGen, InputInfo);
|
||||
else
|
||||
emitInlinedDirective(CGF, D.getDirectiveKind(), ThenGen);
|
||||
|
|
|
|||
|
|
@ -33,6 +33,15 @@
|
|||
// CK0: [[KMP_TASK_T_WITH_PRIVATES_1:%.+]] = type { [[KMP_TASK_T]], [[KMP_PRIVATES_T_2:%.+]] }
|
||||
// CK0-32: [[KMP_PRIVATES_T_2]] = type { [1 x i64], [1 x i8*], [1 x i8*], [1 x i8*] }
|
||||
// CK0-64: [[KMP_PRIVATES_T_2]] = type { [1 x i8*], [1 x i8*], [1 x i64], [1 x i8*] }
|
||||
// CK0: [[KMP_TASK_T_WITH_PRIVATES_4:%.+]] = type { [[KMP_TASK_T]], [[KMP_PRIVATES_T_5:%.+]] }
|
||||
// CK0-32: [[KMP_PRIVATES_T_5]] = type { [1 x i64], [1 x i8*], [1 x i8*], [1 x i8*] }
|
||||
// CK0-64: [[KMP_PRIVATES_T_5]] = type { [1 x i8*], [1 x i8*], [1 x i64], [1 x i8*] }
|
||||
// CK0: [[KMP_TASK_T_WITH_PRIVATES_7:%.+]] = type { [[KMP_TASK_T]], [[KMP_PRIVATES_T_8:%.+]] }
|
||||
// CK0-32: [[KMP_PRIVATES_T_8]] = type { [1 x i64], [1 x i8*], [1 x i8*], [1 x i8*] }
|
||||
// CK0-64: [[KMP_PRIVATES_T_8]] = type { [1 x i8*], [1 x i8*], [1 x i64], [1 x i8*] }
|
||||
// CK0: [[KMP_TASK_T_WITH_PRIVATES_10:%.+]] = type { [[KMP_TASK_T]], [[KMP_PRIVATES_T_11:%.+]] }
|
||||
// CK0-32: [[KMP_PRIVATES_T_11]] = type { [1 x i64], [1 x i8*], [1 x i8*], [1 x i8*] }
|
||||
// CK0-64: [[KMP_PRIVATES_T_11]] = type { [1 x i8*], [1 x i8*], [1 x i64], [1 x i8*] }
|
||||
|
||||
// CK0-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}.region_id = weak constant i8 0
|
||||
// CK0-64: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16]
|
||||
|
|
@ -330,18 +339,39 @@ void foo(int a){
|
|||
// CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
|
||||
#pragma omp target enter data map(mapper(id),to: c)
|
||||
|
||||
// CK0-DAG: call void @__tgt_target_data_begin_nowait_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EDNWSIZES]]{{.+}}, {{.+}}[[EDNWTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
|
||||
// CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
|
||||
// CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
|
||||
// CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%[^,]+]] to i8**
|
||||
// CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
|
||||
// CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
|
||||
// CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
|
||||
// CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
|
||||
// CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
|
||||
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
|
||||
// CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
|
||||
// CK0-DAG: call i32 @__kmpc_omp_task([[IDENT_T]]* @{{[^,]+}}, i32 %{{[^,]+}}, i8* [[TASK_2:%.+]])
|
||||
// CK0-DAG: [[TASK_2]] = call i8* @__kmpc_omp_target_task_alloc([[IDENT_T]]* @{{[^,]+}}, i32 %{{[^,]+}}, i32 1, i[[sz]] {{40|72}}, i[[sz]] 1, i32 (i32, i8*)* bitcast (i32 (i32, [[KMP_TASK_T_WITH_PRIVATES_4]]*)* [[OMP_TASK_ENTRY_18:@[^,]+]] to i32 (i32, i8*)*), i64 -1)
|
||||
// CK0-DAG: [[TASK_WITH_PRIVATES:%.+]] = bitcast i8* [[TASK_2]] to [[KMP_TASK_T_WITH_PRIVATES_4]]*
|
||||
// CK0-DAG: [[PRIVATES:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES_4]], [[KMP_TASK_T_WITH_PRIVATES_4]]* [[TASK_WITH_PRIVATES]], i32 0, i32 1
|
||||
// CK0-32-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_5]], [[KMP_PRIVATES_T_5]]* [[PRIVATES]], i32 0, i32 1
|
||||
// CK0-64-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_5]], [[KMP_PRIVATES_T_5]]* [[PRIVATES]], i32 0, i32 0
|
||||
// CK0-DAG: [[FPBPADDR:%.+]] = bitcast [1 x i8*]* [[FPBPGEP]] to i8*
|
||||
// CK0-DAG: [[BPADDR:%.+]] = bitcast i8** [[BPGEP:%.+]] to i8*
|
||||
// CK0-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPBPADDR]], i8* align {{4|8}} [[BPADDR]], i[[sz]] {{4|8}}, i1 false)
|
||||
// CK0-DAG: [[BPGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], i32 0, i32 0
|
||||
// CK0-DAG: [[BPGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP]], i32 0, i32 0
|
||||
// CK0-DAG: [[BPADDR:%.+]] = bitcast i8** [[BPGEP]] to %class.C**
|
||||
// CK0-DAG: store %class.C* [[C:%[^,]+]], %class.C** [[BPADDR]], align
|
||||
// CK0-32-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_5]], [[KMP_PRIVATES_T_5]]* [[PRIVATES]], i32 0, i32 2
|
||||
// CK0-64-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_5]], [[KMP_PRIVATES_T_5]]* [[PRIVATES]], i32 0, i32 1
|
||||
// CK0-DAG: [[FPPADDR:%.+]] = bitcast [1 x i8*]* [[FPPGEP]] to i8*
|
||||
// CK0-DAG: [[PADDR:%.+]] = bitcast i8** [[PGEP:%.+]] to i8*
|
||||
// CK0-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPPADDR]], i8* align {{4|8}} [[PADDR]], i[[sz]] {{4|8}}, i1 false)
|
||||
// CK0-DAG: [[PGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], i32 0, i32 0
|
||||
// CK0-DAG: [[PGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P]], i32 0, i32 0
|
||||
// CK0-DAG: [[PADDR:%.+]] = bitcast i8** [[PGEP]] to %class.C**
|
||||
// CK0-DAG: store %class.C* [[C]], %class.C** [[PADDR]], align
|
||||
// CK0-32-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_5]], [[KMP_PRIVATES_T_5]]* [[PRIVATES]], i32 0, i32 0
|
||||
// CK0-64-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_5]], [[KMP_PRIVATES_T_5]]* [[PRIVATES]], i32 0, i32 2
|
||||
// CK0-DAG: [[FPSZADDR:%.+]] = bitcast [1 x i64]* [[FPSZGEP]] to i8*
|
||||
// CK0-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPSZADDR]], i8* align {{4|8}} bitcast ([1 x i64]* [[EDNWSIZES]] to i8*), i[[sz]] {{4|8}}, i1 false)
|
||||
// CK0-DAG: [[FPMPRGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_5]], [[KMP_PRIVATES_T_5]]* [[PRIVATES]], i32 0, i32 3
|
||||
// CK0-DAG: [[FPMPRADDR:%.+]] = bitcast [1 x i8*]* [[FPMPRGEP]] to i8*
|
||||
// CK0-DAG: [[MPRADDR:%.+]] = bitcast i8** [[MPRGEP:%.+]] to i8*
|
||||
// CK0-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPMPRADDR]], i8* align {{4|8}} [[MPRADDR]], i[[sz]] {{4|8}}, i1 false)
|
||||
// CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%.+]] to i8**
|
||||
// CK0-DAG: [[MPRGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[MPR]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPRGEP]], align
|
||||
#pragma omp target enter data map(mapper(id),to: c) nowait
|
||||
|
||||
// CK0-DAG: call void @__tgt_target_data_end_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EXDSIZES]]{{.+}}, {{.+}}[[EXDTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
|
||||
|
|
@ -358,18 +388,39 @@ void foo(int a){
|
|||
// CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
|
||||
#pragma omp target exit data map(mapper(id),from: c)
|
||||
|
||||
// CK0-DAG: call void @__tgt_target_data_end_nowait_mapper(i64 {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[EXDNWSIZES]]{{.+}}, {{.+}}[[EXDNWTYPES]]{{.+}}, i8** [[MPRGEP:%.+]])
|
||||
// CK0-DAG: [[BPGEP]] = getelementptr inbounds {{.+}}[[BPS:%[^,]+]], i32 0, i32 0
|
||||
// CK0-DAG: [[PGEP]] = getelementptr inbounds {{.+}}[[PS:%[^,]+]], i32 0, i32 0
|
||||
// CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%[^,]+]] to i8**
|
||||
// CK0-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BPS]], i32 0, i32 0
|
||||
// CK0-DAG: [[P1:%.+]] = getelementptr inbounds {{.+}}[[PS]], i32 0, i32 0
|
||||
// CK0-DAG: [[MPR1:%.+]] = getelementptr inbounds {{.+}}[[MPR]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: [[CBP1:%.+]] = bitcast i8** [[BP1]] to %class.C**
|
||||
// CK0-DAG: [[CP1:%.+]] = bitcast i8** [[P1]] to %class.C**
|
||||
// CK0-DAG: store %class.C* [[VAL:%[^,]+]], %class.C** [[CBP1]]
|
||||
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[CP1]]
|
||||
// CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPR1]]
|
||||
// CK0-DAG: call i32 @__kmpc_omp_task([[IDENT_T]]* @{{[^,]+}}, i32 %{{[^,]+}}, i8* [[TASK_3:%.+]])
|
||||
// CK0-DAG: [[TASK_3]] = call i8* @__kmpc_omp_target_task_alloc([[IDENT_T]]* @{{[^,]+}}, i32 %{{[^,]+}}, i32 1, i[[sz]] {{40|72}}, i[[sz]] 1, i32 (i32, i8*)* bitcast (i32 (i32, [[KMP_TASK_T_WITH_PRIVATES_7]]*)* [[OMP_TASK_ENTRY_25:@[^,]+]] to i32 (i32, i8*)*), i64 -1)
|
||||
// CK0-DAG: [[TASK_WITH_PRIVATES:%.+]] = bitcast i8* [[TASK_3]] to [[KMP_TASK_T_WITH_PRIVATES_7]]*
|
||||
// CK0-DAG: [[PRIVATES:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES_7]], [[KMP_TASK_T_WITH_PRIVATES_7]]* [[TASK_WITH_PRIVATES]], i32 0, i32 1
|
||||
// CK0-32-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_8]], [[KMP_PRIVATES_T_8]]* [[PRIVATES]], i32 0, i32 1
|
||||
// CK0-64-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_8]], [[KMP_PRIVATES_T_8]]* [[PRIVATES]], i32 0, i32 0
|
||||
// CK0-DAG: [[FPBPADDR:%.+]] = bitcast [1 x i8*]* [[FPBPGEP]] to i8*
|
||||
// CK0-DAG: [[BPADDR:%.+]] = bitcast i8** [[BPGEP:%.+]] to i8*
|
||||
// CK0-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPBPADDR]], i8* align {{4|8}} [[BPADDR]], i[[sz]] {{4|8}}, i1 false)
|
||||
// CK0-DAG: [[BPGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], i32 0, i32 0
|
||||
// CK0-DAG: [[BPGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP]], i32 0, i32 0
|
||||
// CK0-DAG: [[BPADDR:%.+]] = bitcast i8** [[BPGEP]] to %class.C**
|
||||
// CK0-DAG: store %class.C* [[C:%[^,]+]], %class.C** [[BPADDR]], align
|
||||
// CK0-32-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_8]], [[KMP_PRIVATES_T_8]]* [[PRIVATES]], i32 0, i32 2
|
||||
// CK0-64-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_8]], [[KMP_PRIVATES_T_8]]* [[PRIVATES]], i32 0, i32 1
|
||||
// CK0-DAG: [[FPPADDR:%.+]] = bitcast [1 x i8*]* [[FPPGEP]] to i8*
|
||||
// CK0-DAG: [[PADDR:%.+]] = bitcast i8** [[PGEP:%.+]] to i8*
|
||||
// CK0-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPPADDR]], i8* align {{4|8}} [[PADDR]], i[[sz]] {{4|8}}, i1 false)
|
||||
// CK0-DAG: [[PGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], i32 0, i32 0
|
||||
// CK0-DAG: [[PGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P]], i32 0, i32 0
|
||||
// CK0-DAG: [[PADDR:%.+]] = bitcast i8** [[PGEP]] to %class.C**
|
||||
// CK0-DAG: store %class.C* [[C]], %class.C** [[PADDR]], align
|
||||
// CK0-32-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_8]], [[KMP_PRIVATES_T_8]]* [[PRIVATES]], i32 0, i32 0
|
||||
// CK0-64-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_8]], [[KMP_PRIVATES_T_8]]* [[PRIVATES]], i32 0, i32 2
|
||||
// CK0-DAG: [[FPSZADDR:%.+]] = bitcast [1 x i64]* [[FPSZGEP]] to i8*
|
||||
// CK0-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPSZADDR]], i8* align {{4|8}} bitcast ([1 x i64]* [[EXDNWSIZES]] to i8*), i[[sz]] {{4|8}}, i1 false)
|
||||
// CK0-DAG: [[FPMPRGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_8]], [[KMP_PRIVATES_T_8]]* [[PRIVATES]], i32 0, i32 3
|
||||
// CK0-DAG: [[FPMPRADDR:%.+]] = bitcast [1 x i8*]* [[FPMPRGEP]] to i8*
|
||||
// CK0-DAG: [[MPRADDR:%.+]] = bitcast i8** [[MPRGEP:%.+]] to i8*
|
||||
// CK0-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPMPRADDR]], i8* align {{4|8}} [[MPRADDR]], i[[sz]] {{4|8}}, i1 false)
|
||||
// CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%.+]] to i8**
|
||||
// CK0-DAG: [[MPRGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[MPR]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPRGEP]], align
|
||||
#pragma omp target exit data map(mapper(id),from: c) nowait
|
||||
|
||||
// CK0-DAG: call void @__tgt_target_data_update_mapper(i64 -1, i32 1, i8** [[TGEPBP:%.+]], i8** [[TGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[TSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[TTYPES]]{{.+}}, i8** [[TMPRGEP:%.+]])
|
||||
|
|
@ -400,18 +451,39 @@ void foo(int a){
|
|||
// CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[FMPR1]]
|
||||
#pragma omp target update from(mapper(id): c)
|
||||
|
||||
// CK0-DAG: call void @__tgt_target_data_update_nowait_mapper(i64 -1, i32 1, i8** [[FGEPBP:%.+]], i8** [[FGEPP:%.+]], i64* getelementptr {{.+}}[1 x i64]* [[FNWSIZES]], i32 0, i32 0), {{.+}}getelementptr {{.+}}[1 x i64]* [[FNWTYPES]]{{.+}}, i8** [[FMPRGEP:%.+]])
|
||||
// CK0-DAG: [[FGEPBP]] = getelementptr inbounds {{.+}}[[FBP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK0-DAG: [[FGEPP]] = getelementptr inbounds {{.+}}[[FP:%[^,]+]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK0-DAG: [[FMPRGEP]] = bitcast [1 x i8*]* [[FMPR:%[^,]+]] to i8**
|
||||
// CK0-DAG: [[FBP0:%.+]] = getelementptr inbounds {{.+}}[[FBP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK0-DAG: [[FP0:%.+]] = getelementptr inbounds {{.+}}[[FP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK0-DAG: [[FMPR1:%.+]] = getelementptr inbounds {{.+}}[[FMPR]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: [[FCBP0:%.+]] = bitcast i8** [[FBP0]] to %class.C**
|
||||
// CK0-DAG: [[FCP0:%.+]] = bitcast i8** [[FP0]] to %class.C**
|
||||
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCBP0]]
|
||||
// CK0-DAG: store %class.C* [[VAL]], %class.C** [[FCP0]]
|
||||
// CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[FMPR1]]
|
||||
// CK0-DAG: call i32 @__kmpc_omp_task([[IDENT_T]]* @{{[^,]+}}, i32 %{{[^,]+}}, i8* [[TASK_4:%.+]])
|
||||
// CK0-DAG: [[TASK_4]] = call i8* @__kmpc_omp_target_task_alloc([[IDENT_T]]* @{{[^,]+}}, i32 %{{[^,]+}}, i32 1, i[[sz]] {{40|72}}, i[[sz]] 1, i32 (i32, i8*)* bitcast (i32 (i32, [[KMP_TASK_T_WITH_PRIVATES_10]]*)* [[OMP_TASK_ENTRY_34:@[^,]+]] to i32 (i32, i8*)*), i64 -1)
|
||||
// CK0-DAG: [[TASK_WITH_PRIVATES:%.+]] = bitcast i8* [[TASK_4]] to [[KMP_TASK_T_WITH_PRIVATES_10]]*
|
||||
// CK0-DAG: [[PRIVATES:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES_10]], [[KMP_TASK_T_WITH_PRIVATES_10]]* [[TASK_WITH_PRIVATES]], i32 0, i32 1
|
||||
// CK0-32-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_11]], [[KMP_PRIVATES_T_11]]* [[PRIVATES]], i32 0, i32 1
|
||||
// CK0-64-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_11]], [[KMP_PRIVATES_T_11]]* [[PRIVATES]], i32 0, i32 0
|
||||
// CK0-DAG: [[FPBPADDR:%.+]] = bitcast [1 x i8*]* [[FPBPGEP]] to i8*
|
||||
// CK0-DAG: [[BPADDR:%.+]] = bitcast i8** [[BPGEP:%.+]] to i8*
|
||||
// CK0-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPBPADDR]], i8* align {{4|8}} [[BPADDR]], i[[sz]] {{4|8}}, i1 false)
|
||||
// CK0-DAG: [[BPGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], i32 0, i32 0
|
||||
// CK0-DAG: [[BPGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP]], i32 0, i32 0
|
||||
// CK0-DAG: [[BPADDR:%.+]] = bitcast i8** [[BPGEP]] to %class.C**
|
||||
// CK0-DAG: store %class.C* [[C:%[^,]+]], %class.C** [[BPADDR]], align
|
||||
// CK0-32-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_11]], [[KMP_PRIVATES_T_11]]* [[PRIVATES]], i32 0, i32 2
|
||||
// CK0-64-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_11]], [[KMP_PRIVATES_T_11]]* [[PRIVATES]], i32 0, i32 1
|
||||
// CK0-DAG: [[FPPADDR:%.+]] = bitcast [1 x i8*]* [[FPPGEP]] to i8*
|
||||
// CK0-DAG: [[PADDR:%.+]] = bitcast i8** [[PGEP:%.+]] to i8*
|
||||
// CK0-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPPADDR]], i8* align {{4|8}} [[PADDR]], i[[sz]] {{4|8}}, i1 false)
|
||||
// CK0-DAG: [[PGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], i32 0, i32 0
|
||||
// CK0-DAG: [[PGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P]], i32 0, i32 0
|
||||
// CK0-DAG: [[PADDR:%.+]] = bitcast i8** [[PGEP]] to %class.C**
|
||||
// CK0-DAG: store %class.C* [[C]], %class.C** [[PADDR]], align
|
||||
// CK0-32-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_11]], [[KMP_PRIVATES_T_11]]* [[PRIVATES]], i32 0, i32 0
|
||||
// CK0-64-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_11]], [[KMP_PRIVATES_T_11]]* [[PRIVATES]], i32 0, i32 2
|
||||
// CK0-DAG: [[FPSZADDR:%.+]] = bitcast [1 x i64]* [[FPSZGEP]] to i8*
|
||||
// CK0-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPSZADDR]], i8* align {{4|8}} bitcast ([1 x i64]* [[FNWSIZES]] to i8*), i[[sz]] {{4|8}}, i1 false)
|
||||
// CK0-DAG: [[FPMPRGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T_11]], [[KMP_PRIVATES_T_11]]* [[PRIVATES]], i32 0, i32 3
|
||||
// CK0-DAG: [[FPMPRADDR:%.+]] = bitcast [1 x i8*]* [[FPMPRGEP]] to i8*
|
||||
// CK0-DAG: [[MPRADDR:%.+]] = bitcast i8** [[MPRGEP:%.+]] to i8*
|
||||
// CK0-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPMPRADDR]], i8* align {{4|8}} [[MPRADDR]], i[[sz]] {{4|8}}, i1 false)
|
||||
// CK0-DAG: [[MPRGEP]] = bitcast [1 x i8*]* [[MPR:%.+]] to i8**
|
||||
// CK0-DAG: [[MPRGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[MPR]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: store i8* bitcast (void (i8*, i8*, i8*, i64, i64)* [[MPRFUNC]] to i8*), i8** [[MPRGEP]], align
|
||||
#pragma omp target update from(mapper(id): c) nowait
|
||||
}
|
||||
|
||||
|
|
@ -496,6 +568,61 @@ void foo(int a){
|
|||
// CK0: call void [[OUTLINE_1]](i32 {{%.+}}, i32* {{%.+}}, i8* [[PRIVATES]], {{.+}}, i8* [[TASK_WITH_PRIVATES]], [[ANON_T_0]]* [[ANON]])
|
||||
// CK0: }
|
||||
|
||||
// CK0: define internal void [[OMP_OUTLINED_16:@.+]](i32{{.*}} %{{[^,]+}}, i32* noalias %{{[^,]+}}, i8* noalias %{{[^,]+}}
|
||||
// CK0-DAG: call void @__tgt_target_data_begin_nowait_mapper(i64 -1, i32 1, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i64* [[SZ:%[^,]+]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[EDNWTYPES]], i32 0, i32 0), i8** [[MPR:%.+]])
|
||||
// CK0-DAG: [[BP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BPADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: [[P]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[PADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: [[SZ]] = getelementptr inbounds [1 x i64], [1 x i64]* [[SZADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: [[MPR]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[MPRADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: [[BPADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPBPADDR:%[^,]+]], align
|
||||
// CK0-DAG: [[PADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPPADDR:%[^,]+]], align
|
||||
// CK0-DAG: [[SZADDR]] = load [1 x i64]*, [1 x i64]** [[FPSZADDR:%[^,]+]], align
|
||||
// CK0-DAG: [[MPRADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPMPRADDR:%[^,]+]], align
|
||||
// CK0-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [1 x i8*]** [[FPBPADDR]], [1 x i8*]** [[FPPADDR]], [1 x i64]** [[FPSZADDR]], [1 x i8*]** [[FPMPRADDR]])
|
||||
// CK0: ret void
|
||||
// CK0: }
|
||||
|
||||
// CK0: define internal {{.*}}i32 [[OMP_TASK_ENTRY_18]](i32 {{.*}}%{{[^,]+}}, [[KMP_TASK_T_WITH_PRIVATES_4]]* noalias %{{[^,]+}})
|
||||
// CK0: call void [[OMP_OUTLINED_16]]
|
||||
// CK0: ret i32 0
|
||||
// CK0: }
|
||||
|
||||
// CK0: define internal void [[OMP_OUTLINED_23:@.+]](i32{{.*}} %{{[^,]+}}, i32* noalias %{{[^,]+}}, i8* noalias %{{[^,]+}}
|
||||
// CK0-DAG: call void @__tgt_target_data_end_nowait_mapper(i64 -1, i32 1, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i64* [[SZ:%[^,]+]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[EXDNWTYPES]], i32 0, i32 0), i8** [[MPR:%.+]])
|
||||
// CK0-DAG: [[BP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BPADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: [[P]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[PADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: [[SZ]] = getelementptr inbounds [1 x i64], [1 x i64]* [[SZADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: [[MPR]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[MPRADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: [[BPADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPBPADDR:%[^,]+]], align
|
||||
// CK0-DAG: [[PADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPPADDR:%[^,]+]], align
|
||||
// CK0-DAG: [[SZADDR]] = load [1 x i64]*, [1 x i64]** [[FPSZADDR:%[^,]+]], align
|
||||
// CK0-DAG: [[MPRADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPMPRADDR:%[^,]+]], align
|
||||
// CK0-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [1 x i8*]** [[FPBPADDR]], [1 x i8*]** [[FPPADDR]], [1 x i64]** [[FPSZADDR]], [1 x i8*]** [[FPMPRADDR]])
|
||||
// CK0: }
|
||||
|
||||
// CK0: define internal {{.*}}i32 [[OMP_TASK_ENTRY_25]](i32 {{.*}}%{{[^,]+}}, [[KMP_TASK_T_WITH_PRIVATES_7]]* noalias %{{[^,]+}})
|
||||
// CK0: call void [[OMP_OUTLINED_23]]
|
||||
// CK0: ret i32 0
|
||||
// CK0: }
|
||||
|
||||
// CK0: define internal void [[OMP_OUTLINED_32:@.+]](i32{{.*}} %{{[^,]+}}, i32* noalias %{{[^,]+}}, i8* noalias %{{[^,]+}}
|
||||
// CK0-DAG: call void @__tgt_target_data_update_nowait_mapper(i64 -1, i32 1, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i64* [[SZ:%[^,]+]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[FNWTYPES]], i32 0, i32 0), i8** [[MPR:%.+]])
|
||||
// CK0-DAG: [[BP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BPADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: [[P]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[PADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: [[SZ]] = getelementptr inbounds [1 x i64], [1 x i64]* [[SZADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: [[MPR]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[MPRADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK0-DAG: [[BPADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPBPADDR:%[^,]+]], align
|
||||
// CK0-DAG: [[PADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPPADDR:%[^,]+]], align
|
||||
// CK0-DAG: [[SZADDR]] = load [1 x i64]*, [1 x i64]** [[FPSZADDR:%[^,]+]], align
|
||||
// CK0-DAG: [[MPRADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPMPRADDR:%[^,]+]], align
|
||||
// CK0-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [1 x i8*]** [[FPBPADDR]], [1 x i8*]** [[FPPADDR]], [1 x i64]** [[FPSZADDR]], [1 x i8*]** [[FPMPRADDR]])
|
||||
// CK0: }
|
||||
|
||||
// CK0: define internal {{.*}}i32 [[OMP_TASK_ENTRY_34]](i32 {{.*}}%{{[^,]+}}, [[KMP_TASK_T_WITH_PRIVATES_10]]* noalias %{{[^,]+}})
|
||||
// CK0: call void [[OMP_OUTLINED_32]]
|
||||
// CK0: ret i32 0
|
||||
// CK0: }
|
||||
|
||||
#endif // CK0
|
||||
|
||||
|
||||
|
|
|
|||
|
|
@ -29,10 +29,17 @@ struct ST {
|
|||
ST<int> gb;
|
||||
double gc[100];
|
||||
|
||||
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800]
|
||||
// CK1: [[IDENT_T:%.+]] = type { i32, i32, i32, i32, i8* }
|
||||
|
||||
// CK1: [[KMP_TASK_T_WITH_PRIVATES:%.+]] = type { [[TASK_T:%[^,]+]], [[KMP_PRIVATES_T:%.+]] }
|
||||
// CK1: [[TASK_T]] = type { i8*, i32 (i32, i8*)*, i32, %{{[^,]+}}, %{{[^,]+}} }
|
||||
// CK1-32: [[KMP_PRIVATES_T]] = type { [1 x i64], [1 x i8*], [1 x i8*] }
|
||||
// CK1-64: [[KMP_PRIVATES_T]] = type { [1 x i8*], [1 x i8*], [1 x i64] }
|
||||
|
||||
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 800]
|
||||
// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 32]
|
||||
|
||||
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
|
||||
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i64] [i64 4]
|
||||
// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 33]
|
||||
|
||||
// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 37]
|
||||
|
|
@ -50,21 +57,35 @@ void foo(int arg) {
|
|||
float lb[arg];
|
||||
|
||||
// Region 00
|
||||
// CK1-DAG: call void @__tgt_target_data_begin_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: [[DEVi32]] = load i32, i32* %{{[^,]+}},
|
||||
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
|
||||
// CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[P0:%.+]] = getelementptr inbounds {{.+}}[[P]], i{{.+}} 0, i{{.+}} 0
|
||||
// CK1-DAG: [[CBP0:%.+]] = bitcast i8** [[BP0]] to [100 x double]**
|
||||
// CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to [100 x double]**
|
||||
// CK1-DAG: store [100 x double]* @gc, [100 x double]** [[CBP0]]
|
||||
// CK1-DAG: store [100 x double]* @gc, [100 x double]** [[CP0]]
|
||||
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
// CK1-NOT: __tgt_target_data_end
|
||||
// CK1-DAG: call i32 @__kmpc_omp_task([[IDENT_T]]* @{{[^,]+}}, i32 %{{[^,]+}}, i8* [[TASK:%.+]])
|
||||
// CK1-DAG: [[TASK]] = call i8* @__kmpc_omp_target_task_alloc([[IDENT_T]]* @{{[^,]+}}, i32 %{{[^,]+}}, i32 1, i{{32|64}} {{36|64}}, i{{32|64}} 4, i32 (i32, i8*)* bitcast (i32 (i32, [[KMP_TASK_T_WITH_PRIVATES]]*)* [[OMP_TASK_ENTRY:@[^,]+]] to i32 (i32, i8*)*), i64 [[DEV:%.+]])
|
||||
// CK1-DAG: [[DEV]] = sext i32 [[DEVi32:%[^,]+]] to i64
|
||||
// CK1-DAG: [[DEVi32]] = load i32, i32* %{{[^,]+}},
|
||||
// CK1-DAG: [[TASK_CAST:%.+]] = bitcast i8* [[TASK]] to [[KMP_TASK_T_WITH_PRIVATES]]*
|
||||
// CK1-DAG: [[TASK_WITH_PRIVATES:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES]], [[KMP_TASK_T_WITH_PRIVATES]]* [[TASK_CAST]], i32 0, i32 1
|
||||
// CK1-32-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[TASK_WITH_PRIVATES]], i32 0, i32 1
|
||||
// CK1-64-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[TASK_WITH_PRIVATES]], i32 0, i32 0
|
||||
// CK1-DAG: [[FPBPADDR:%.+]] = bitcast [1 x i8*]* [[FPBPGEP]] to i8*
|
||||
// CK1-DAG: [[BPADDR:%.+]] = bitcast i8** [[BPGEP:%.+]] to i8*
|
||||
// CK1-DAG: call void @llvm.memcpy.p0i8.p0i8.i{{32|64}}(i8* align {{4|8}} [[FPBPADDR]], i8* align {{4|8}} [[BPADDR]], i{{32|64}} {{4|8}}, i1 false)
|
||||
// CK1-DAG: [[BPGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], i32 0, i32 0
|
||||
// CK1-DAG: [[BPGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP]], i32 0, i32 0
|
||||
// CK1-DAG: [[BPCAST:%.+]] = bitcast i8** [[BPGEP]] to [100 x double]**
|
||||
// CK1-DAG: store [100 x double]* @gc, [100 x double]** [[BPCAST]], align
|
||||
// CK1-32-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[TASK_WITH_PRIVATES]], i32 0, i32 2
|
||||
// CK1-64-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[TASK_WITH_PRIVATES]], i32 0, i32 1
|
||||
// CK1-DAG: [[FPPADDR:%.+]] = bitcast [1 x i8*]* [[FPPGEP]] to i8*
|
||||
// CK1-DAG: [[PADDR:%.+]] = bitcast i8** [[PGEP:%.+]] to i8*
|
||||
// CK1-DAG: call void @llvm.memcpy.p0i8.p0i8.i{{32|64}}(i8* align {{4|8}} [[FPPADDR]], i8* align {{4|8}} [[PADDR]], i{{32|64}} {{4|8}}, i1 false)
|
||||
// CK1-DAG: [[PGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], i32 0, i32 0
|
||||
// CK1-DAG: [[PGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P]], i32 0, i32 0
|
||||
// CK1-DAG: [[PCAST:%.+]] = bitcast i8** [[PGEP]] to [100 x double]**
|
||||
// CK1-DAG: store [100 x double]* @gc, [100 x double]** [[PCAST]], align
|
||||
// CK1-32-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[TASK_WITH_PRIVATES]], i32 0, i32 0
|
||||
// CK1-64-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[TASK_WITH_PRIVATES]], i32 0, i32 2
|
||||
// CK1-DAG: [[FPSZADDR:%.+]] = bitcast [1 x i64]* [[FPSZGEP]] to i8*
|
||||
// CK1-DAG: call void @llvm.memcpy.p0i8.p0i8.i{{32|64}}(i8* align {{4|8}} [[FPSZADDR]], i8* align {{4|8}} bitcast ([1 x i64]* [[SIZE00]] to i8*), i{{32|64}} 8, i1 false)
|
||||
// CK1-NOT: __tgt_target_data_end
|
||||
#pragma omp target enter data if(1+3-5) device(arg) map(alloc: gc) nowait
|
||||
{++arg;}
|
||||
|
||||
|
|
@ -155,7 +176,7 @@ void foo(int arg) {
|
|||
{++arg;}
|
||||
|
||||
// Region 05
|
||||
// CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}, i8** null)
|
||||
// CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}, i8** null)
|
||||
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
|
||||
|
|
@ -167,7 +188,7 @@ void foo(int arg) {
|
|||
// CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
|
||||
// CK1-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
|
||||
// CK1-DAG: store float* [[VAR0]], float** [[CP0]]
|
||||
// CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
|
||||
// CK1-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]]
|
||||
// CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
|
||||
// CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
|
||||
// CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
|
||||
|
|
@ -180,7 +201,7 @@ void foo(int arg) {
|
|||
{++arg;}
|
||||
|
||||
// Region 06
|
||||
// CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}, i8** null)
|
||||
// CK1-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}, i8** null)
|
||||
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
|
||||
|
|
@ -192,7 +213,7 @@ void foo(int arg) {
|
|||
// CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
|
||||
// CK1-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
|
||||
// CK1-DAG: store float* [[VAR0]], float** [[CP0]]
|
||||
// CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
|
||||
// CK1-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]]
|
||||
// CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
|
||||
// CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
|
||||
// CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
|
||||
|
|
@ -201,6 +222,20 @@ void foo(int arg) {
|
|||
#pragma omp target enter data map(always close, to: lb)
|
||||
{++arg;}
|
||||
}
|
||||
|
||||
|
||||
// CK1: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%0, [[KMP_TASK_T_WITH_PRIVATES]]* noalias %1)
|
||||
// CK1-DAG: call void @__tgt_target_data_begin_nowait_mapper(i64 %{{[^,]+}}, i32 1, i8** [[BPADDR:%[^,]+]], i8** [[PADDR:%[^,]+]], i64* [[SZADDR:%[^,]+]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[MTYPE00]], i32 0, i32 0), i8** null)
|
||||
// CK1-DAG: [[BPADDR]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[FPBPADDR:%[^,]+]], i{{32|64}} 0, i{{32|64}} 0
|
||||
// CK1-DAG: [[PADDR]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[FPPADDR:%[^,]+]], i{{32|64}} 0, i{{32|64}} 0
|
||||
// CK1-DAG: [[SZADDR]] = getelementptr inbounds [1 x i64], [1 x i64]* [[FPSZADDR:%[^,]+]], i{{32|64}} 0, i{{32|64}} 0
|
||||
// CK1-DAG: [[FPBPADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPBP:%[^,]+]], align
|
||||
// CK1-DAG: [[FPPADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPP:%[^,]+]], align
|
||||
// CK1-DAG: [[FPSZADDR]] = load [1 x i64]*, [1 x i64]** [[FPSZ:%[^,]+]], align
|
||||
// CK1-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [1 x i8*]** [[FPBP]], [1 x i8*]** [[FPP]], [1 x i64]** [[FPSZ]])
|
||||
// CK1: ret i32 0
|
||||
// CK1: }
|
||||
|
||||
#endif
|
||||
///==========================================================================///
|
||||
// RUN: %clang_cc1 -DCK1A -verify -fopenmp -fopenmp-version=51 -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK1A --check-prefix CK1A-64
|
||||
|
|
@ -253,7 +288,7 @@ void foo(int arg) {
|
|||
// CK1A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
|
||||
// CK1A-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
|
||||
// CK1A-DAG: store float* [[VAR0]], float** [[CP0]]
|
||||
// CK1A-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
|
||||
// CK1A-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]]
|
||||
// CK1A-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
|
||||
// CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
|
||||
// CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
|
||||
|
|
@ -266,7 +301,7 @@ void foo(int arg) {
|
|||
{++arg;}
|
||||
|
||||
// Region 01
|
||||
// CK1A-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i[[sz]]* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
|
||||
// CK1A-DAG: call void @__tgt_target_data_begin_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE01]]{{.+}})
|
||||
// CK1A-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK1A-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
// CK1A-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
|
||||
|
|
@ -278,7 +313,7 @@ void foo(int arg) {
|
|||
// CK1A-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
|
||||
// CK1A-DAG: store float* [[VAR0:%.+]], float** [[CBP0]]
|
||||
// CK1A-DAG: store float* [[VAR0]], float** [[CP0]]
|
||||
// CK1A-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
|
||||
// CK1A-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]]
|
||||
// CK1A-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
|
||||
// CK1A-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
|
||||
// CK1A-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
|
||||
|
|
@ -344,7 +379,7 @@ int bar(int arg){
|
|||
// CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double***
|
||||
// CK2-DAG: store [[ST]]* [[VAR0:%.+]], [[ST]]** [[CBP0]]
|
||||
// CK2-DAG: store double** [[SEC0:%.+]], double*** [[CP0]]
|
||||
// CK2-DAG: store i[[sz]] {{%.+}}, i[[sz]]* [[S0]]
|
||||
// CK2-DAG: store i64 {{%.+}}, i64* [[S0]]
|
||||
// CK2-DAG: [[SEC0]] = getelementptr inbounds {{.*}}[[ST]]* [[VAR0]], i32 0, i32 1
|
||||
|
||||
// CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
|
||||
|
|
@ -494,7 +529,7 @@ int bar(int arg){
|
|||
// CK5-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double***
|
||||
// CK5-DAG: store [[STT]]* [[VAR0:%.+]], [[STT]]** [[CBP0]]
|
||||
// CK5-DAG: store double** [[SEC0:%.+]], double*** [[CP0]]
|
||||
// CK5-DAG: store i[[sz]] {{%.+}}, i[[sz]]* [[S0]]
|
||||
// CK5-DAG: store i64 {{%.+}}, i64* [[S0]]
|
||||
// CK5-DAG: [[SEC0]] = getelementptr inbounds {{.*}}[[STT]]* [[VAR0]], i32 0, i32 1
|
||||
|
||||
// CK5-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
|
||||
|
|
|
|||
|
|
@ -29,10 +29,16 @@ struct ST {
|
|||
ST<int> gb;
|
||||
double gc[100];
|
||||
|
||||
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800]
|
||||
// CK1: [[IDENT_T:%.+]] = type { i32, i32, i32, i32, i8* }
|
||||
// CK1: [[KMP_TASK_T_WITH_PRIVATES:%.+]] = type { [[KMP_TASK_T:%[^,]+]], [[KMP_PRIVATES_T:%.+]] }
|
||||
// CK1: [[KMP_TASK_T]] = type { i8*, i32 (i32, i8*)*, i32, %{{[^,]+}}, %{{[^,]+}} }
|
||||
// CK1-32: [[KMP_PRIVATES_T]] = type { [1 x i64], [1 x i8*], [1 x i8*] }
|
||||
// CK1-64: [[KMP_PRIVATES_T]] = type { [1 x i8*], [1 x i8*], [1 x i64] }
|
||||
|
||||
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i64] [i64 800]
|
||||
// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 34]
|
||||
|
||||
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i[[sz]]] [i[[sz]] 4]
|
||||
// CK1: [[SIZE02:@.+]] = {{.+}}constant [1 x i64] [i64 4]
|
||||
// CK1: [[MTYPE02:@.+]] = {{.+}}constant [1 x i64] [i64 32]
|
||||
|
||||
// CK1: [[MTYPE03:@.+]] = {{.+}}constant [1 x i64] [i64 38]
|
||||
|
|
@ -51,20 +57,33 @@ void foo(int arg) {
|
|||
|
||||
// Region 00
|
||||
// CK1-NOT: __tgt_target_data_begin
|
||||
// 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: [[DEVi32]] = load i32, i32* %{{[^,]+}},
|
||||
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
|
||||
// CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], 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: [[PC0:%.+]] = bitcast i8** [[P0]] to [100 x double]**
|
||||
// CK1-DAG: store [100 x double]* @gc, [100 x double]** [[BPC0]]
|
||||
// CK1-DAG: store [100 x double]* @gc, [100 x double]** [[PC0]]
|
||||
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
// CK1-DAG: call i32 @__kmpc_omp_task([[IDENT_T]]* @{{[^,]+}}, i32 %{{[^,]+}}, i8* [[TASK:%.+]])
|
||||
// CK1-DAG: [[TASK]] = call i8* @__kmpc_omp_target_task_alloc([[IDENT_T]]* @{{[^,]+}}, i32 %{{[^,]+}}, i32 1, i[[sz:32|64]] {{36|64}}, i{{32|64}} 4, i32 (i32, i8*)* bitcast (i32 (i32, [[KMP_TASK_T_WITH_PRIVATES]]*)* [[OMP_TASK_ENTRY:@[^,]+]] to i32 (i32, i8*)*), i64 [[DEV:%.+]])
|
||||
// CK1-DAG: [[DEV]] = sext i32 [[DEV32:%.+]] to i64
|
||||
// CK1-DAG: [[TASK_WITH_PRIVATES:%.+]] = bitcast i8* [[TASK]] to [[KMP_TASK_T_WITH_PRIVATES]]*
|
||||
// CK1-DAG: [[PRIVATES:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES]], [[KMP_TASK_T_WITH_PRIVATES]]* [[TASK_WITH_PRIVATES]], i32 0, i32 1
|
||||
// CK1-32-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[PRIVATES]], i32 0, i32 1
|
||||
// CK1-64-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[PRIVATES]], i32 0, i32 0
|
||||
// CK1-DAG: [[FPBPADDR:%.+]] = bitcast [1 x i8*]* [[FPBPGEP]] to i8*
|
||||
// CK1-DAG: [[BPADDR:%.+]] = bitcast i8** [[BPGEP:%.+]] to i8*
|
||||
// CK1-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPBPADDR]], i8* align {{4|8}} [[BPADDR]], i[[sz]] {{4|8}}, i1 false)
|
||||
// CK1-DAG: [[BPGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], i32 0, i32 0
|
||||
// CK1-DAG: [[BPGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP]], i32 0, i32 0
|
||||
// CK1-DAG: [[BPADDR:%.+]] = bitcast i8** [[BPGEP]] to [100 x double]**
|
||||
// CK1-DAG: store [100 x double]* [[GC:@[^,]+]], [100 x double]** [[BPADDR]], align
|
||||
// CK1-32-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[PRIVATES]], i32 0, i32 2
|
||||
// CK1-64-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[PRIVATES]], i32 0, i32 1
|
||||
// CK1-DAG: [[FPPADDR:%.+]] = bitcast [1 x i8*]* [[FPPGEP]] to i8*
|
||||
// CK1-DAG: [[PADDR:%.+]] = bitcast i8** [[PGEP:%.+]] to i8*
|
||||
// CK1-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPPADDR]], i8* align {{4|8}} [[PADDR]], i[[sz]] {{4|8}}, i1 false)
|
||||
// CK1-DAG: [[PGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], i32 0, i32 0
|
||||
// CK1-DAG: [[PGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P]], i32 0, i32 0
|
||||
// CK1-DAG: [[PADDR:%.+]] = bitcast i8** [[PGEP]] to [100 x double]**
|
||||
// CK1-DAG: store [100 x double]* [[GC]], [100 x double]** [[PADDR]], align
|
||||
// CK1-32-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[PRIVATES]], i32 0, i32 0
|
||||
// CK1-64-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[PRIVATES]], i32 0, i32 2
|
||||
// CK1-DAG: [[FPSZADDR:%.+]] = bitcast [1 x i64]* [[FPSZGEP]] to i8*
|
||||
// CK1-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPSZADDR]], i8* align {{4|8}} bitcast ([1 x i64]* [[SIZE00]] to i8*), i[[sz]] {{4|8}}, i1 false)
|
||||
#pragma omp target exit data if(1+3-5) device(arg) map(from: gc) nowait
|
||||
{++arg;}
|
||||
|
||||
|
|
@ -101,7 +120,7 @@ void foo(int arg) {
|
|||
|
||||
// Region 03
|
||||
// CK1-NOT: __tgt_target_data_begin
|
||||
// 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: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE03]]{{.+}}, i8** null)
|
||||
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
|
||||
|
|
@ -113,8 +132,8 @@ void foo(int arg) {
|
|||
// CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
|
||||
// CK1-DAG: store float* [[VAL0:%[^,]+]], float** [[CBP0]]
|
||||
// CK1-DAG: store float* [[VAL0]], float** [[CP0]]
|
||||
// CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
|
||||
// CK1-64-DAG: [[CSVAL0]] = mul nuw i[[sz]] %{{[^,]+}}, 4
|
||||
// CK1-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]]
|
||||
// CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
|
||||
// CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
|
||||
// CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
|
|
@ -156,7 +175,7 @@ void foo(int arg) {
|
|||
|
||||
// Region 05
|
||||
// CK1-NOT: __tgt_target_data_begin
|
||||
// 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: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE05]]{{.+}}, i8** null)
|
||||
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
|
||||
|
|
@ -168,8 +187,8 @@ void foo(int arg) {
|
|||
// CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
|
||||
// CK1-DAG: store float* [[VAL0:%[^,]+]], float** [[CBP0]]
|
||||
// CK1-DAG: store float* [[VAL0]], float** [[CP0]]
|
||||
// CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
|
||||
// CK1-64-DAG: [[CSVAL0]] = mul nuw i[[sz]] %{{[^,]+}}, 4
|
||||
// CK1-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]]
|
||||
// CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
|
||||
// CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
|
||||
// CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
|
|
@ -181,7 +200,7 @@ void foo(int arg) {
|
|||
|
||||
// Region 06
|
||||
// CK1-NOT: __tgt_target_data_begin
|
||||
// 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: call void @__tgt_target_data_end_mapper(i64 -1, i32 1, i8** [[GEPBP:%.+]], i8** [[GEPP:%.+]], i64* [[GEPS:%.+]], {{.+}}getelementptr {{.+}}[1 x i{{.+}}]* [[MTYPE06]]{{.+}}, i8** null)
|
||||
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
// CK1-DAG: [[GEPS]] = getelementptr inbounds {{.+}}[[S:%[^,]+]]
|
||||
|
|
@ -193,14 +212,27 @@ void foo(int arg) {
|
|||
// CK1-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to float**
|
||||
// CK1-DAG: store float* [[VAL0:%[^,]+]], float** [[CBP0]]
|
||||
// CK1-DAG: store float* [[VAL0]], float** [[CP0]]
|
||||
// CK1-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
|
||||
// CK1-64-DAG: [[CSVAL0]] = mul nuw i[[sz]] %{{[^,]+}}, 4
|
||||
// CK1-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]]
|
||||
// CK1-64-DAG: [[CSVAL0]] = mul nuw i64 %{{[^,]+}}, 4
|
||||
// CK1-32-DAG: [[CSVAL0]] = sext i32 [[CSVAL032:%.+]] to i64
|
||||
// CK1-32-DAG: [[CSVAL032]] = mul nuw i32 %{{[^,]+}}, 4
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
#pragma omp target exit data map(always close, from: lb)
|
||||
{++arg;}
|
||||
}
|
||||
|
||||
// CK1: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%{{[^,]+}}, [[KMP_TASK_T_WITH_PRIVATES]]* noalias %{{[^,]+}})
|
||||
// CK1-DAG: call void @__tgt_target_data_end_nowait_mapper(i64 %{{[^,]+}}, i32 1, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i64* [[SZ:%[^,]+]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[MTYPE00]], i32 0, i32 0), i8** null)
|
||||
// CK1-DAG: [[BP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BPADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK1-DAG: [[P]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[PADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK1-DAG: [[SZ]] = getelementptr inbounds [1 x i64], [1 x i64]* [[SZADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK1-DAG: [[BPADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPBPADDR:%[^,]+]], align
|
||||
// CK1-DAG: [[PADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPPADDR:%[^,]+]], align
|
||||
// CK1-DAG: [[SZADDR]] = load [1 x i64]*, [1 x i64]** [[FPSZADDR:%[^,]+]], align
|
||||
// CK1-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [1 x i8*]** [[FPBPADDR]], [1 x i8*]** [[FPPADDR]], [1 x i64]** [[FPSZADDR]])
|
||||
// CK1: ret i32 0
|
||||
// CK1: }
|
||||
|
||||
#endif
|
||||
///==========================================================================///
|
||||
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
|
||||
|
|
@ -259,7 +291,7 @@ int bar(int arg){
|
|||
// CK2-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double***
|
||||
// CK2-DAG: store [[ST]]* [[VAR0:%[^,]+]], [[ST]]** [[CBP0]]
|
||||
// CK2-DAG: store double** [[SEC0:%[^,]+]], double*** [[CP0]]
|
||||
// CK2-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
|
||||
// CK2-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]]
|
||||
// CK2-DAG: [[SEC0]] = getelementptr inbounds {{.*}}[[ST]]* [[VAR0]], i32 0, i32 1
|
||||
|
||||
// CK2-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
|
||||
|
|
@ -363,7 +395,7 @@ int bar(int arg){
|
|||
// CK4-DAG: [[CP0:%.+]] = bitcast i8** [[P0]] to double***
|
||||
// CK4-DAG: store [[STT]]* [[VAR0:%[^,]+]], [[STT]]** [[CBP0]]
|
||||
// CK4-DAG: store double** [[SEC0:%[^,]+]], double*** [[CP0]]
|
||||
// CK4-DAG: store i[[sz]] [[CSVAL0:%[^,]+]], i[[sz]]* [[S0]]
|
||||
// CK4-DAG: store i64 [[CSVAL0:%[^,]+]], i64* [[S0]]
|
||||
// CK4-DAG: [[SEC0]] = getelementptr inbounds {{.*}}[[STT]]* [[VAR0]], i32 0, i32 1
|
||||
|
||||
// CK4-DAG: [[BP1:%.+]] = getelementptr inbounds {{.+}}[[BP]], i{{.+}} 0, i{{.+}} 1
|
||||
|
|
|
|||
|
|
@ -29,6 +29,12 @@ struct ST {
|
|||
ST<int> gb;
|
||||
double gc[100];
|
||||
|
||||
// CK1: [[IDENT_T:%.+]] = type { i32, i32, i32, i32, i8* }
|
||||
// CK1: [[KMP_TASK_T_WITH_PRIVATES:%.+]] = type { [[KMP_TASK_T:%[^,]+]], [[KMP_PRIVATES_T:%.+]] }
|
||||
// CK1: [[KMP_TASK_T]] = type { i8*, i32 (i32, i8*)*, i32, %{{[^,]+}}, %{{[^,]+}} }
|
||||
// CK1-32: [[KMP_PRIVATES_T]] = type { [1 x i64], [1 x i8*], [1 x i8*] }
|
||||
// CK1-64: [[KMP_PRIVATES_T]] = type { [1 x i8*], [1 x i8*], [1 x i64] }
|
||||
|
||||
// CK1: [[SIZE00:@.+]] = {{.+}}constant [1 x i[[sz:64|32]]] [i{{64|32}} 800]
|
||||
// CK1: [[MTYPE00:@.+]] = {{.+}}constant [1 x i64] [i64 34]
|
||||
|
||||
|
|
@ -46,20 +52,33 @@ void foo(int arg) {
|
|||
float lb[arg];
|
||||
|
||||
// Region 00
|
||||
// CK1-DAG: call void @__tgt_target_data_update_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: [[DEVi32]] = load i32, i32* %{{[^,]+}},
|
||||
// CK1-DAG: [[GEPBP]] = getelementptr inbounds {{.+}}[[BP:%[^,]+]]
|
||||
// CK1-DAG: [[GEPP]] = getelementptr inbounds {{.+}}[[P:%[^,]+]]
|
||||
|
||||
// CK1-DAG: [[BP0:%.+]] = getelementptr inbounds {{.+}}[[BP]], 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: [[PC0:%.+]] = bitcast i8** [[P0]] to [100 x double]**
|
||||
// CK1-DAG: store [100 x double]* @gc, [100 x double]** [[BPC0]]
|
||||
// CK1-DAG: store [100 x double]* @gc, [100 x double]** [[PC0]]
|
||||
|
||||
// CK1: %{{.+}} = add nsw i32 %{{[^,]+}}, 1
|
||||
// CK1-DAG: call i32 @__kmpc_omp_task([[IDENT_T]]* @{{[^,]+}}, i32 %{{[^,]+}}, i8* [[TASK:%.+]])
|
||||
// CK1-DAG: [[TASK]] = call i8* @__kmpc_omp_target_task_alloc([[IDENT_T]]* @{{[^,]+}}, i32 %{{[^,]+}}, i32 1, i[[sz:32|64]] {{36|64}}, i{{32|64}} 4, i32 (i32, i8*)* bitcast (i32 (i32, [[KMP_TASK_T_WITH_PRIVATES]]*)* [[OMP_TASK_ENTRY:@[^,]+]] to i32 (i32, i8*)*), i64 [[DEV:%.+]])
|
||||
// CK1-DAG: [[DEV]] = sext i32 [[DEV32:%.+]] to i64
|
||||
// CK1-DAG: [[TASK_WITH_PRIVATES:%.+]] = bitcast i8* [[TASK]] to [[KMP_TASK_T_WITH_PRIVATES]]*
|
||||
// CK1-DAG: [[PRIVATES:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES]], [[KMP_TASK_T_WITH_PRIVATES]]* [[TASK_WITH_PRIVATES]], i32 0, i32 1
|
||||
// CK1-32-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[PRIVATES]], i32 0, i32 1
|
||||
// CK1-64-DAG: [[FPBPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[PRIVATES]], i32 0, i32 0
|
||||
// CK1-DAG: [[FPBPADDR:%.+]] = bitcast [1 x i8*]* [[FPBPGEP]] to i8*
|
||||
// CK1-DAG: [[BPADDR:%.+]] = bitcast i8** [[BPGEP:%.+]] to i8*
|
||||
// CK1-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPBPADDR]], i8* align {{4|8}} [[BPADDR]], i[[sz]] {{4|8}}, i1 false)
|
||||
// CK1-DAG: [[BPGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP:%.+]], i32 0, i32 0
|
||||
// CK1-DAG: [[BPGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BP]], i32 0, i32 0
|
||||
// CK1-DAG: [[BPADDR:%.+]] = bitcast i8** [[BPGEP]] to [100 x double]**
|
||||
// CK1-DAG: store [100 x double]* [[GC:@[^,]+]], [100 x double]** [[BPADDR]], align
|
||||
// CK1-32-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[PRIVATES]], i32 0, i32 2
|
||||
// CK1-64-DAG: [[FPPGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[PRIVATES]], i32 0, i32 1
|
||||
// CK1-DAG: [[FPPADDR:%.+]] = bitcast [1 x i8*]* [[FPPGEP]] to i8*
|
||||
// CK1-DAG: [[PADDR:%.+]] = bitcast i8** [[PGEP:%.+]] to i8*
|
||||
// CK1-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPPADDR]], i8* align {{4|8}} [[PADDR]], i[[sz]] {{4|8}}, i1 false)
|
||||
// CK1-DAG: [[PGEP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P:%.+]], i32 0, i32 0
|
||||
// CK1-DAG: [[PGEP:%.+]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[P]], i32 0, i32 0
|
||||
// CK1-DAG: [[PADDR:%.+]] = bitcast i8** [[PGEP]] to [100 x double]**
|
||||
// CK1-DAG: store [100 x double]* [[GC]], [100 x double]** [[PADDR]], align
|
||||
// CK1-32-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[PRIVATES]], i32 0, i32 0
|
||||
// CK1-64-DAG: [[FPSZGEP:%.+]] = getelementptr inbounds [[KMP_PRIVATES_T]], [[KMP_PRIVATES_T]]* [[PRIVATES]], i32 0, i32 2
|
||||
// CK1-DAG: [[FPSZADDR:%.+]] = bitcast [1 x i64]* [[FPSZGEP]] to i8*
|
||||
// CK1-DAG: call void @llvm.memcpy.p0i8.p0i8.i[[sz]](i8* align {{4|8}} [[FPSZADDR]], i8* align {{4|8}} bitcast ([1 x i64]* [[SIZE00]] to i8*), i[[sz]] {{4|8}}, i1 false)
|
||||
#pragma omp target update if(1+3-5) device(arg) from(gc) nowait
|
||||
{++arg;}
|
||||
|
||||
|
|
@ -142,6 +161,18 @@ void foo(int arg) {
|
|||
#pragma omp target update to(gb.b[:3])
|
||||
{++arg;}
|
||||
}
|
||||
|
||||
// CK1: define internal {{.*}}i32 [[OMP_TASK_ENTRY]](i32 {{.*}}%{{[^,]+}}, [[KMP_TASK_T_WITH_PRIVATES]]* noalias %{{[^,]+}})
|
||||
// CK1-DAG: call void @__tgt_target_data_update_nowait_mapper(i64 %{{[^,]+}}, i32 1, i8** [[BP:%[^,]+]], i8** [[P:%[^,]+]], i64* [[SZ:%[^,]+]], i64* getelementptr inbounds ([1 x i64], [1 x i64]* [[MTYPE00]], i32 0, i32 0), i8** null)
|
||||
// CK1-DAG: [[BP]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[BPADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK1-DAG: [[P]] = getelementptr inbounds [1 x i8*], [1 x i8*]* [[PADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK1-DAG: [[SZ]] = getelementptr inbounds [1 x i64], [1 x i64]* [[SZADDR:%[^,]+]], i[[sz]] 0, i[[sz]] 0
|
||||
// CK1-DAG: [[BPADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPBPADDR:%[^,]+]], align
|
||||
// CK1-DAG: [[PADDR]] = load [1 x i8*]*, [1 x i8*]** [[FPPADDR:%[^,]+]], align
|
||||
// CK1-DAG: [[SZADDR]] = load [1 x i64]*, [1 x i64]** [[FPSZADDR:%[^,]+]], align
|
||||
// CK1-DAG: call void (i8*, ...) %{{.+}}(i8* %{{[^,]+}}, [1 x i8*]** [[FPBPADDR]], [1 x i8*]** [[FPPADDR]], [1 x i64]** [[FPSZADDR]])
|
||||
// CK1: ret i32 0
|
||||
// CK1: }
|
||||
#endif
|
||||
///==========================================================================///
|
||||
// RUN: %clang_cc1 -DCK2 -verify -fopenmp -fopenmp-targets=powerpc64le-ibm-linux-gnu -x c++ -triple powerpc64le-unknown-unknown -emit-llvm %s -o - | FileCheck %s --check-prefix CK2 --check-prefix CK2-64
|
||||
|
|
|
|||
Loading…
Reference in New Issue