diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -9385,7 +9385,8 @@ assert(OutlinedFn && "Invalid outlined function!"); - const bool RequiresOuterTask = D.hasClausesOfKind(); + const bool RequiresOuterTask = D.hasClausesOfKind() || + D.hasClausesOfKind(); llvm::SmallVector CapturedVars; const CapturedStmt &CS = *D.getCapturedStmt(OMPD_target); auto &&ArgsCodegen = [&CS, &CapturedVars](CodeGenFunction &CGF, diff --git a/clang/test/OpenMP/declare_mapper_codegen.cpp b/clang/test/OpenMP/declare_mapper_codegen.cpp --- a/clang/test/OpenMP/declare_mapper_codegen.cpp +++ b/clang/test/OpenMP/declare_mapper_codegen.cpp @@ -22,6 +22,13 @@ #ifdef CK0 // Mapper function code generation and runtime interface. +// CK0: [[IDENT_T:%.+]] = type { i32, i32, i32, i32, i8* } +// CK0: [[KMP_TASK_T_WITH_PRIVATES:%.+]] = type { [[KMP_TASK_T:%.+]], [[KMP_PRIVATES:%.+]] } +// CK0: [[KMP_TASK_T]] = type { i8*, i32 (i32, i8*)*, i32, %union{{.+}}, %union{{.+}} } +// CK0: [[KMP_PRIVATES]] = type { [1 x i64], [1 x i8*], [1 x i8*], [1 x i8*] } +// CK0: [[KMP_TASK_T_WITH_PRIVATES_1:%.+]] = type { [[KMP_TASK_T]], [[KMP_PRIVATES_1:%.+]] } +// CK0: [[KMP_PRIVATES_1]] = type { [1 x i64], [1 x i8*], [1 x i8*], [1 x i8*] } + // CK0-LABEL: @.__omp_offloading_{{.*}}foo{{.*}}.region_id = weak constant i8 0 // CK0-64: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 16] // CK0-32: [[SIZES:@.+]] = {{.+}}constant [1 x i64] [i64 8] @@ -254,19 +261,11 @@ ++c.a; } - // CK0-DAG: call i32 @__tgt_target_nowait_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[NWSIZES]]{{.+}}, {{.+}}[[NWTYPES]]{{.+}}, 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: call void [[KERNEL:@.+]](%class.C* [[VAL]]) + // CK0: [[TASK:%.+]] = call i8* @__kmpc_omp_target_task_alloc([[IDENT_T]]* {{.+}}, i32 {{.+}}, i32 1, i32 40, i32 4, i32 (i32, i8*)* bitcast (i32 (i32, [[KMP_TASK_T_WITH_PRIVATES]]*)* [[TASK_ENTRY:@.+]] to i32 (i32, i8*)*), i64 -1) + // CK0: [[TASK_CAST:%.+]] = bitcast i8* [[TASK]] to [[KMP_TASK_T_WITH_PRIVATES]]* + // CK0: [[TASK_CAST_GET:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES]], [[KMP_TASK_T_WITH_PRIVATES]]* [[TASK_CAST]], i32 0, i32 0 + // CK0: {{.+}} = getelementptr inbounds [[KMP_TASK_T]], [[KMP_TASK_T]]* [[TASK_CAST_GET]], i32 0, i32 0 + // CK0: {{.+}} = call i32 @__kmpc_omp_task([[IDENT_T]]* @1, i32 {{.+}}, i8* [[TASK]]) #pragma omp target map(mapper(id),tofrom: c) nowait { ++c.a; @@ -290,19 +289,11 @@ ++c.a; } - // CK0-DAG: call i32 @__tgt_target_teams_nowait_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[TEAMNWSIZES]]{{.+}}, {{.+}}[[TEAMNWTYPES]]{{.+}}, i8** [[MPRGEP:%.+]], i32 0, i32 0) - // 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: call void [[KERNEL:@.+]](%class.C* [[VAL]]) + // CK0: [[TASK_1:%.+]] = call i8* @__kmpc_omp_target_task_alloc([[IDENT_T]]* {{.+}}, i32 {{.+}}, i32 1, i32 40, i32 4, i32 (i32, i8*)* bitcast (i32 (i32, [[KMP_TASK_T_WITH_PRIVATES_1]]*)* [[TASK_ENTRY_1:@.+]] to i32 (i32, i8*)*), i64 -1) + // CK0: [[TASK_CAST_1:%.+]] = bitcast i8* [[TASK_1]] to [[KMP_TASK_T_WITH_PRIVATES_1]]* + // CK0: [[TASK_CAST_GET_1:%.+]] = getelementptr inbounds [[KMP_TASK_T_WITH_PRIVATES_1]], [[KMP_TASK_T_WITH_PRIVATES_1]]* [[TASK_CAST_1]], i32 0, i32 0 + // CK0: {{.+}} = getelementptr inbounds [[KMP_TASK_T]], [[KMP_TASK_T]]* [[TASK_CAST_GET_1]], i32 0, i32 0 + // CK0: {{.+}} = call i32 @__kmpc_omp_task([[IDENT_T]]* @1, i32 {{.+}}, i8* [[TASK_1]]) #pragma omp target teams map(mapper(id),to: c) nowait { ++c.a; @@ -417,6 +408,46 @@ // CK0: {{.+}} = add nsw i32 [[VAL]], 1 // CK0: } +// CK0: define internal void [[OUTLINED:@.+]](i32 {{%.+}}, {{.+}}) +// CK0-DAG: call i32 @__tgt_target_nowait_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[NWSIZES]]{{.+}}, {{.+}}[[NWTYPES]]{{.+}}, 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: call void [[KERNEL:@.+]](%class.C* [[VAL]]) +// CK0: } + +// CK0: define internal i32 [[TASK_ENTRY]](i32 %0, [[TASK_T_WITH_PRIVATES]]* noalias %1) +// CK0: call void [[OUTLINED]]({{.+}}) +// CK0: } + +// CK0: define internal void [[OUTLINE_1:@.+]]({{.+}}) +// CK0-DAG: call i32 @__tgt_target_teams_nowait_mapper(i64 {{.+}}, i8* {{.+}}, i32 1, i8** [[BPGEP:%[0-9]+]], i8** [[PGEP:%[0-9]+]], {{.+}}[[TEAMNWSIZES]]{{.+}}, {{.+}}[[TEAMNWTYPES]]{{.+}}, i8** [[MPRGEP:%.+]], i32 0, i32 0) +// 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: call void [[KERNEL:@.+]](%class.C* [[VAL]]) +// CK0: } + +// CK0: define internal i32 [[TASK_ENTRY_1]](i32 %0, [[TASK_T_WITH_PRIVATES_1]]* noalias %1) +// CK0: call void [[OUTLINED_1]]({{.+}}) +// CK0: } + #endif // CK0