diff --git a/clang/lib/CodeGen/CGOpenMPRuntime.cpp b/clang/lib/CodeGen/CGOpenMPRuntime.cpp index 478fda7d913b..3730b9af12fa 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntime.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntime.cpp @@ -2839,12 +2839,12 @@ void CGOpenMPRuntime::emitParallelCall(CodeGenFunction &CGF, SourceLocation Loc, RT.createRuntimeFunction(OMPRTL__kmpc_serialized_parallel), Args); // OutlinedFn(>id, &zero, CapturedStruct); - Address ThreadIDAddr = RT.emitThreadIDAddress(CGF, Loc); Address ZeroAddr = CGF.CreateDefaultAlignTempAlloca(CGF.Int32Ty, /*Name*/ ".zero.addr"); CGF.InitTempAlloca(ZeroAddr, CGF.Builder.getInt32(/*C*/ 0)); llvm::SmallVector OutlinedFnArgs; - OutlinedFnArgs.push_back(ThreadIDAddr.getPointer()); + // ThreadId for serialized parallels is 0. + OutlinedFnArgs.push_back(ZeroAddr.getPointer()); OutlinedFnArgs.push_back(ZeroAddr.getPointer()); OutlinedFnArgs.append(CapturedVars.begin(), CapturedVars.end()); RT.emitOutlinedFunctionCall(CGF, Loc, OutlinedFn, OutlinedFnArgs); diff --git a/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp b/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp index 5aaebff8f76b..036b5371fe0b 100644 --- a/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp +++ b/clang/lib/CodeGen/CGOpenMPRuntimeNVPTX.cpp @@ -1784,8 +1784,9 @@ void CGOpenMPRuntimeNVPTX::emitNonSPMDParallelCall( /*DestWidth=*/32, /*Signed=*/1), ".zero.addr"); CGF.InitTempAlloca(ZeroAddr, CGF.Builder.getInt32(/*C*/ 0)); - Address ThreadIDAddr = emitThreadIDAddress(CGF, Loc); - auto &&CodeGen = [this, Fn, CapturedVars, Loc, ZeroAddr, ThreadIDAddr]( + // ThreadId for serialized parallels is 0. + Address ThreadIDAddr = ZeroAddr; + auto &&CodeGen = [this, Fn, CapturedVars, Loc, ZeroAddr, &ThreadIDAddr]( CodeGenFunction &CGF, PrePostActionTy &Action) { Action.Enter(CGF); @@ -1883,8 +1884,9 @@ void CGOpenMPRuntimeNVPTX::emitNonSPMDParallelCall( Work.emplace_back(WFn); }; - auto &&LNParallelGen = [this, Loc, &SeqGen, &L0ParallelGen, &CodeGen]( - CodeGenFunction &CGF, PrePostActionTy &Action) { + auto &&LNParallelGen = [this, Loc, &SeqGen, &L0ParallelGen, &CodeGen, + &ThreadIDAddr](CodeGenFunction &CGF, + PrePostActionTy &Action) { RegionCodeGenTy RCG(CodeGen); if (IsInParallelRegion) { SeqGen(CGF, Action); @@ -1936,6 +1938,8 @@ void CGOpenMPRuntimeNVPTX::emitNonSPMDParallelCall( // There is no need to emit line number for unconditional branch. (void)ApplyDebugLocation::CreateEmpty(CGF); CGF.EmitBlock(ElseBlock); + // In the worker need to use the real thread id. + ThreadIDAddr = emitThreadIDAddress(CGF, Loc); RCG(CGF); // There is no need to emit line number for unconditional branch. (void)ApplyDebugLocation::CreateEmpty(CGF); @@ -1965,10 +1969,11 @@ void CGOpenMPRuntimeNVPTX::emitSPMDParallelCall( /*DestWidth=*/32, /*Signed=*/1), ".zero.addr"); CGF.InitTempAlloca(ZeroAddr, CGF.Builder.getInt32(/*C*/ 0)); - Address ThreadIDAddr = emitThreadIDAddress(CGF, Loc); + // ThreadId for serialized parallels is 0. + Address ThreadIDAddr = ZeroAddr; auto &&CodeGen = [this, OutlinedFn, CapturedVars, Loc, ZeroAddr, - ThreadIDAddr](CodeGenFunction &CGF, - PrePostActionTy &Action) { + &ThreadIDAddr](CodeGenFunction &CGF, + PrePostActionTy &Action) { Action.Enter(CGF); llvm::SmallVector OutlinedFnArgs; @@ -1995,6 +2000,8 @@ void CGOpenMPRuntimeNVPTX::emitSPMDParallelCall( }; if (IsInTargetMasterThreadRegion) { + // In the worker need to use the real thread id. + ThreadIDAddr = emitThreadIDAddress(CGF, Loc); RegionCodeGenTy RCG(CodeGen); RCG(CGF); } else { diff --git a/clang/test/OpenMP/nvptx_target_codegen.cpp b/clang/test/OpenMP/nvptx_target_codegen.cpp index d4d9941bc431..5f9b3bd32808 100644 --- a/clang/test/OpenMP/nvptx_target_codegen.cpp +++ b/clang/test/OpenMP/nvptx_target_codegen.cpp @@ -562,7 +562,6 @@ int baz(int f, double &a) { // CHECK: [[REC_ADDR:%.+]] = bitcast i8* [[PTR]] to %struct._globalized_locals_ty* // CHECK: [[F_PTR:%.+]] = getelementptr inbounds %struct._globalized_locals_ty, %struct._globalized_locals_ty* [[REC_ADDR]], i32 0, i32 0 // CHECK: store i32 %{{.+}}, i32* [[F_PTR]], - // CHECK: store i32 [[GTID]], i32* [[GTID_ADDR]], // CHECK: [[RES:%.+]] = call i8 @__kmpc_is_spmd_exec_mode() // CHECK: icmp ne i8 [[RES]], 0 @@ -573,7 +572,7 @@ int baz(int f, double &a) { // CHECK: br i1 // CHECK: call void @__kmpc_serialized_parallel(%struct.ident_t* @{{.+}}, i32 [[GTID]]) - // CHECK: call void [[OUTLINED:@.+]](i32* [[GTID_ADDR]], i32* [[ZERO_ADDR]], i32* [[F_PTR]], double* %{{.+}}) + // CHECK: call void [[OUTLINED:@.+]](i32* [[ZERO_ADDR]], i32* [[ZERO_ADDR]], i32* [[F_PTR]], double* %{{.+}}) // CHECK: call void @__kmpc_end_serialized_parallel(%struct.ident_t* @{{.+}}, i32 [[GTID]]) // CHECK: br label @@ -591,6 +590,7 @@ int baz(int f, double &a) { // CHECK: call void @__kmpc_end_sharing_variables() // CHECK: br label + // CHECK: store i32 [[GTID]], i32* [[GTID_ADDR]], // CHECK: call void [[OUTLINED]](i32* [[GTID_ADDR]], i32* [[ZERO_ADDR]], i32* [[F_PTR]], double* %{{.+}}) // CHECK: br label diff --git a/clang/test/OpenMP/parallel_if_codegen.cpp b/clang/test/OpenMP/parallel_if_codegen.cpp index e59415a9d295..49e4bd402ce8 100644 --- a/clang/test/OpenMP/parallel_if_codegen.cpp +++ b/clang/test/OpenMP/parallel_if_codegen.cpp @@ -29,12 +29,12 @@ void gtid_test() { } // CHECK: define internal {{.*}}void [[GTID_TEST_REGION1]](i{{.+}}* noalias [[GTID_PARAM:%.+]], i32* noalias +// CHECK: store i32 0, i32* [[ZERO_ADDR:%.+]], // CHECK: store i{{[0-9]+}}* [[GTID_PARAM]], i{{[0-9]+}}** [[GTID_ADDR_REF:%.+]], // CHECK: [[GTID_ADDR:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[GTID_ADDR_REF]] // CHECK: [[GTID:%.+]] = load i{{[0-9]+}}, i{{[0-9]+}}* [[GTID_ADDR]] // CHECK: call {{.*}}void @__kmpc_serialized_parallel(%{{.+}}* @{{.+}}, i{{.+}} [[GTID]]) -// CHECK: [[GTID_ADDR:%.+]] = load i{{[0-9]+}}*, i{{[0-9]+}}** [[GTID_ADDR_REF]] -// CHECK: call void [[GTID_TEST_REGION2:@.+]](i{{[0-9]+}}* [[GTID_ADDR]] +// CHECK: call void [[GTID_TEST_REGION2:@.+]](i{{[0-9]+}}* [[ZERO_ADDR]] // CHECK: call {{.*}}void @__kmpc_end_serialized_parallel(%{{.+}}* @{{.+}}, i{{.+}} [[GTID]]) // CHECK: ret void @@ -56,12 +56,13 @@ int tmain(T Arg) { // CHECK-LABEL: define {{.*}}i{{[0-9]+}} @main() int main() { // CHECK: [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num( +// CHECK: store i32 0, i32* [[ZERO_ADDR2:%.+]], +// CHECK: store i32 0, i32* [[ZERO_ADDR1:%.+]], // CHECK: call {{.*}}void {{.+}} @__kmpc_fork_call(%{{.+}}* @{{.+}}, i{{.+}} 0, void {{.+}}* [[CAP_FN4:@.+]] to void #pragma omp parallel if (true) fn4(); // CHECK: call {{.*}}void @__kmpc_serialized_parallel(%{{.+}}* @{{.+}}, i32 [[GTID]]) -// CHECK: store i32 [[GTID]], i32* [[GTID_ADDR:%.+]], -// CHECK: call void [[CAP_FN5:@.+]](i32* [[GTID_ADDR]], +// CHECK: call void [[CAP_FN5:@.+]](i32* [[ZERO_ADDR1]], i32* [[ZERO_ADDR1]]) // CHECK: call {{.*}}void @__kmpc_end_serialized_parallel(%{{.+}}* @{{.+}}, i32 [[GTID]]) #pragma omp parallel if (false) fn5(); @@ -72,8 +73,7 @@ int main() { // CHECK: br label %[[OMP_END:.+]] // CHECK: [[OMP_ELSE]] // CHECK: call {{.*}}void @__kmpc_serialized_parallel(%{{.+}}* @{{.+}}, i32 [[GTID]]) -// CHECK: store i32 [[GTID]], i32* [[GTID_ADDR:%.+]], -// CHECK: call void [[CAP_FN6]](i32* [[GTID_ADDR]], +// CHECK: call void [[CAP_FN6]](i32* [[ZERO_ADDR2]], i32* [[ZERO_ADDR2]]) // CHECK: call {{.*}}void @__kmpc_end_serialized_parallel(%{{.+}}* @{{.+}}, i32 [[GTID]]) // CHECK: br label %[[OMP_END]] // CHECK: [[OMP_END]] @@ -97,10 +97,11 @@ int main() { // CHECK-LABEL: define {{.+}} @{{.+}}tmain // CHECK: [[GTID:%.+]] = call {{.*}}i32 @__kmpc_global_thread_num( +// CHECK: store i32 0, i32* [[ZERO_ADDR2:%.+]], +// CHECK: store i32 0, i32* [[ZERO_ADDR1:%.+]], // CHECK: call {{.*}}void {{.+}} @__kmpc_fork_call(%{{.+}}* @{{.+}}, i{{.+}} 0, void {{.+}}* [[CAP_FN1:@.+]] to void // CHECK: call {{.*}}void @__kmpc_serialized_parallel(%{{.+}}* @{{.+}}, i32 [[GTID]]) -// CHECK: store i32 [[GTID]], i32* [[GTID_ADDR:%.+]], -// CHECK: call void [[CAP_FN2:@.+]](i32* [[GTID_ADDR]], +// CHECK: call void [[CAP_FN2:@.+]](i32* [[ZERO_ADDR1]], i32* [[ZERO_ADDR1]]) // CHECK: call {{.*}}void @__kmpc_end_serialized_parallel(%{{.+}}* @{{.+}}, i32 [[GTID]]) // CHECK: br i1 %{{.+}}, label %[[OMP_THEN:.+]], label %[[OMP_ELSE:.+]] // CHECK: [[OMP_THEN]] @@ -108,8 +109,7 @@ int main() { // CHECK: br label %[[OMP_END:.+]] // CHECK: [[OMP_ELSE]] // CHECK: call {{.*}}void @__kmpc_serialized_parallel(%{{.+}}* @{{.+}}, i32 [[GTID]]) -// CHECK: store i32 [[GTID]], i32* [[GTID_ADDR:%.+]], -// CHECK: call void [[CAP_FN3]](i32* [[GTID_ADDR]], +// CHECK: call void [[CAP_FN3]](i32* [[ZERO_ADDR2]], i32* [[ZERO_ADDR2]]) // CHECK: call {{.*}}void @__kmpc_end_serialized_parallel(%{{.+}}* @{{.+}}, i32 [[GTID]]) // CHECK: br label %[[OMP_END]] // CHECK: [[OMP_END]]