Skip to content

[offload][OpenMP] Remove device code for num_threads strict #157893

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 2 commits into from
Sep 11, 2025
Merged

[offload][OpenMP] Remove device code for num_threads strict #157893

Show file tree
Hide file tree
Changes from all commits
Commits
Show all changes
2 commits
Select commit Hold shift + click to select a range
ee83f4f
[offload][OpenMP] Remove device code for num_threads strict
ro-i Sep 10, 2025
113bb4c
remove target from parallel in clang/test/AST/ByteCode/openmp.cpp
ro-i Sep 11, 2025
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
35 changes: 17 additions & 18 deletions clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -1210,9 +1210,9 @@ void CGOpenMPRuntimeGPU::emitParallelCall(
if (!CGF.HaveInsertPoint())
return;

auto &&ParallelGen = [this, Loc, OutlinedFn, CapturedVars, IfCond, NumThreads,
NumThreadsModifier, Severity, Message](
CodeGenFunction &CGF, PrePostActionTy &Action) {
auto &&ParallelGen = [this, Loc, OutlinedFn, CapturedVars, IfCond,
NumThreads](CodeGenFunction &CGF,
PrePostActionTy &Action) {
CGBuilderTy &Bld = CGF.Builder;
llvm::Value *NumThreadsVal = NumThreads;
llvm::Function *WFn = WrapperFunctionsMap[OutlinedFn];
Expand Down Expand Up @@ -1260,22 +1260,21 @@ void CGOpenMPRuntimeGPU::emitParallelCall(
NumThreadsVal = Bld.CreateZExtOrTrunc(NumThreadsVal, CGF.Int32Ty);

assert(IfCondVal && "Expected a value");
RuntimeFunction FnID = OMPRTL___kmpc_parallel_51;
llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);
llvm::SmallVector<llvm::Value *, 10> Args(
{RTLoc, getThreadID(CGF, Loc), IfCondVal, NumThreadsVal,
llvm::ConstantInt::get(CGF.Int32Ty, -1), FnPtr, ID,
Bld.CreateBitOrPointerCast(CapturedVarsAddrs.emitRawPointer(CGF),
CGF.VoidPtrPtrTy),
llvm::ConstantInt::get(CGM.SizeTy, CapturedVars.size())});
if (NumThreadsModifier == OMPC_NUMTHREADS_strict) {
FnID = OMPRTL___kmpc_parallel_60;
Args.append({llvm::ConstantInt::get(CGM.Int32Ty, true),
emitSeverityClause(Severity),
emitMessageClause(CGF, Message)});
}
CGF.EmitRuntimeCall(
OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(), FnID), Args);
llvm::Value *Args[] = {
RTLoc,
getThreadID(CGF, Loc),
IfCondVal,
NumThreadsVal,
llvm::ConstantInt::get(CGF.Int32Ty, -1),
FnPtr,
ID,
Bld.CreateBitOrPointerCast(CapturedVarsAddrs.emitRawPointer(CGF),
CGF.VoidPtrPtrTy),
llvm::ConstantInt::get(CGM.SizeTy, CapturedVars.size())};
CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
CGM.getModule(), OMPRTL___kmpc_parallel_51),
Args);
};

RegionCodeGenTy RCG(ParallelGen);
Expand Down
11 changes: 3 additions & 8 deletions clang/lib/CodeGen/CGOpenMPRuntimeGPU.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -165,11 +165,6 @@ class CGOpenMPRuntimeGPU : public CGOpenMPRuntime {
/// Emits call to void __kmpc_push_num_threads(ident_t *loc, kmp_int32
/// global_tid, kmp_int32 num_threads) to generate code for 'num_threads'
/// clause.
/// If the modifier 'strict' is given:
/// Emits call to void __kmpc_push_num_threads_strict(ident_t *loc, kmp_int32
/// global_tid, kmp_int32 num_threads, int severity, const char *message) to
/// generate code for 'num_threads' clause with 'strict' modifier.
/// \param NumThreads An integer value of threads.
void emitNumThreadsClause(
CodeGenFunction &CGF, llvm::Value *NumThreads, SourceLocation Loc,
OpenMPNumThreadsClauseModifier Modifier = OMPC_NUMTHREADS_unknown,
Expand Down Expand Up @@ -238,11 +233,11 @@ class CGOpenMPRuntimeGPU : public CGOpenMPRuntime {
/// \param NumThreads The value corresponding to the num_threads clause, if
/// any, or nullptr.
/// \param NumThreadsModifier The modifier of the num_threads clause, if
/// any, ignored otherwise.
/// any, ignored otherwise. Currently unused on the device.
/// \param Severity The severity corresponding to the num_threads clause, if
/// any, ignored otherwise.
/// any, ignored otherwise. Currently unused on the device.
/// \param Message The message string corresponding to the num_threads clause,
/// if any, or nullptr.
/// if any, or nullptr. Currently unused on the device.
void emitParallelCall(CodeGenFunction &CGF, SourceLocation Loc,
llvm::Function *OutlinedFn,
ArrayRef<llvm::Value *> CapturedVars,
Expand Down
6 changes: 3 additions & 3 deletions clang/test/AST/ByteCode/openmp.cpp
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -17,12 +17,12 @@ extern int omp_get_thread_num(void);

int test2() {
int x = 0;
int device_result[N] = {0};
int result[N] = {0};

#pragma omp target parallel loop num_threads(strict: N) severity(warning) message("msg")
#pragma omp parallel loop num_threads(strict: N) severity(warning) message("msg")
for (int i = 0; i < N; i++) {
x = omp_get_thread_num();
device_result[i] = i + x;
result[i] = i + x;
}
}

Expand Down
1,095 changes: 0 additions & 1,095 deletions clang/test/OpenMP/amdgcn_target_parallel_num_threads_codegen.cpp
View file
Open in desktop

This file was deleted.

913 changes: 66 additions & 847 deletions clang/test/OpenMP/nvptx_target_codegen.cpp
View file
Open in desktop

Large diffs are not rendered by default.

760 changes: 35 additions & 725 deletions clang/test/OpenMP/nvptx_target_parallel_num_threads_codegen.cpp
View file
Open in desktop

Large diffs are not rendered by default.

436 changes: 18 additions & 418 deletions clang/test/OpenMP/target_parallel_generic_loop_codegen.cpp
View file
Open in desktop

Large diffs are not rendered by default.

2,956 changes: 0 additions & 2,956 deletions clang/test/OpenMP/target_parallel_num_threads_strict_codegen.cpp
View file
Open in desktop

This file was deleted.

22 changes: 0 additions & 22 deletions llvm/include/llvm/Frontend/OpenMP/OMP.td
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2061,11 +2061,9 @@ def OMP_TargetParallel : Directive<[Spelling<"target parallel">]> {
let allowedOnceClauses = [
VersionedClause<OMPC_DefaultMap>,
VersionedClause<OMPC_Device>,
VersionedClause<OMPC_Message, 60>,
VersionedClause<OMPC_NumThreads>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
VersionedClause<OMPC_ProcBind>,
VersionedClause<OMPC_Severity, 60>,
VersionedClause<OMPC_ThreadLimit, 51>,
];
let leafConstructs = [OMP_Target, OMP_Parallel];
Expand Down Expand Up @@ -2093,14 +2091,12 @@ def OMP_TargetParallelDo : Directive<[Spelling<"target parallel do">]> {
VersionedClause<OMPC_Collapse>,
VersionedClause<OMPC_DefaultMap>,
VersionedClause<OMPC_Device>,
VersionedClause<OMPC_Message, 60>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_NumThreads>,
VersionedClause<OMPC_Order, 50>,
VersionedClause<OMPC_Ordered>,
VersionedClause<OMPC_ProcBind>,
VersionedClause<OMPC_Schedule>,
VersionedClause<OMPC_Severity, 60>,
];
let leafConstructs = [OMP_Target, OMP_Parallel, OMP_Do];
let category = CA_Executable;
Expand All @@ -2124,7 +2120,6 @@ def OMP_TargetParallelDoSimd
VersionedClause<OMPC_LastPrivate>,
VersionedClause<OMPC_Linear>,
VersionedClause<OMPC_Map>,
VersionedClause<OMPC_Message, 60>,
VersionedClause<OMPC_NonTemporal>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_NumThreads>,
Expand All @@ -2135,7 +2130,6 @@ def OMP_TargetParallelDoSimd
VersionedClause<OMPC_Reduction>,
VersionedClause<OMPC_SafeLen>,
VersionedClause<OMPC_Schedule>,
VersionedClause<OMPC_Severity, 60>,
VersionedClause<OMPC_Shared>,
VersionedClause<OMPC_SimdLen>,
VersionedClause<OMPC_UsesAllocators>,
Expand All @@ -2160,7 +2154,6 @@ def OMP_TargetParallelFor : Directive<[Spelling<"target parallel for">]> {
VersionedClause<OMPC_LastPrivate>,
VersionedClause<OMPC_Linear>,
VersionedClause<OMPC_Map>,
VersionedClause<OMPC_Message, 60>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_NumThreads>,
VersionedClause<OMPC_OMPX_Attribute>,
Expand All @@ -2170,7 +2163,6 @@ def OMP_TargetParallelFor : Directive<[Spelling<"target parallel for">]> {
VersionedClause<OMPC_ProcBind>,
VersionedClause<OMPC_Reduction>,
VersionedClause<OMPC_Schedule>,
VersionedClause<OMPC_Severity, 60>,
VersionedClause<OMPC_Shared>,
VersionedClause<OMPC_UsesAllocators, 50>,
];
Expand Down Expand Up @@ -2200,7 +2192,6 @@ def OMP_TargetParallelForSimd
VersionedClause<OMPC_LastPrivate>,
VersionedClause<OMPC_Linear>,
VersionedClause<OMPC_Map>,
VersionedClause<OMPC_Message, 60>,
VersionedClause<OMPC_NonTemporal, 50>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_NumThreads>,
Expand All @@ -2212,7 +2203,6 @@ def OMP_TargetParallelForSimd
VersionedClause<OMPC_Reduction>,
VersionedClause<OMPC_SafeLen>,
VersionedClause<OMPC_Schedule>,
VersionedClause<OMPC_Severity, 60>,
VersionedClause<OMPC_Shared>,
VersionedClause<OMPC_SimdLen>,
VersionedClause<OMPC_UsesAllocators, 50>,
Expand Down Expand Up @@ -2248,13 +2238,11 @@ def OMP_target_parallel_loop : Directive<[Spelling<"target parallel loop">]> {
VersionedClause<OMPC_Collapse>,
VersionedClause<OMPC_Default>,
VersionedClause<OMPC_DefaultMap>,
VersionedClause<OMPC_Message, 60>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_NumThreads>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
VersionedClause<OMPC_Order>,
VersionedClause<OMPC_ProcBind>,
VersionedClause<OMPC_Severity, 60>,
VersionedClause<OMPC_ThreadLimit, 51>,
];
let leafConstructs = [OMP_Target, OMP_Parallel, OMP_loop];
Expand Down Expand Up @@ -2285,14 +2273,12 @@ def OMP_TargetSimd : Directive<[Spelling<"target simd">]> {
VersionedClause<OMPC_Collapse>,
VersionedClause<OMPC_DefaultMap>,
VersionedClause<OMPC_Device>,
VersionedClause<OMPC_Message, 60>,
VersionedClause<OMPC_NumThreads>,
VersionedClause<OMPC_OMPX_DynCGroupMem>,
VersionedClause<OMPC_Order, 50>,
VersionedClause<OMPC_ProcBind>,
VersionedClause<OMPC_SafeLen>,
VersionedClause<OMPC_Schedule>,
VersionedClause<OMPC_Severity, 60>,
VersionedClause<OMPC_SimdLen>,
VersionedClause<OMPC_ThreadLimit, 51>,
];
Expand Down Expand Up @@ -2385,14 +2371,12 @@ def OMP_TargetTeamsDistributeParallelDo
VersionedClause<OMPC_DefaultMap>,
VersionedClause<OMPC_Device>,
VersionedClause<OMPC_DistSchedule>,
VersionedClause<OMPC_Message, 60>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_NumTeams>,
VersionedClause<OMPC_NumThreads>,
VersionedClause<OMPC_Order, 50>,
VersionedClause<OMPC_ProcBind>,
VersionedClause<OMPC_Schedule>,
VersionedClause<OMPC_Severity, 60>,
VersionedClause<OMPC_ThreadLimit>,
];
let leafConstructs =
Expand Down Expand Up @@ -2426,15 +2410,13 @@ def OMP_TargetTeamsDistributeParallelDoSimd
VersionedClause<OMPC_DefaultMap>,
VersionedClause<OMPC_Device>,
VersionedClause<OMPC_DistSchedule>,
VersionedClause<OMPC_Message, 60>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_NumTeams>,
VersionedClause<OMPC_NumThreads>,
VersionedClause<OMPC_Order, 50>,
VersionedClause<OMPC_ProcBind>,
VersionedClause<OMPC_SafeLen>,
VersionedClause<OMPC_Schedule>,
VersionedClause<OMPC_Severity, 60>,
VersionedClause<OMPC_SimdLen>,
VersionedClause<OMPC_ThreadLimit>,
];
Expand All @@ -2460,7 +2442,6 @@ def OMP_TargetTeamsDistributeParallelFor
VersionedClause<OMPC_IsDevicePtr>,
VersionedClause<OMPC_LastPrivate>,
VersionedClause<OMPC_Map>,
VersionedClause<OMPC_Message, 60>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_NumTeams>,
VersionedClause<OMPC_NumThreads>,
Expand All @@ -2470,7 +2451,6 @@ def OMP_TargetTeamsDistributeParallelFor
VersionedClause<OMPC_ProcBind>,
VersionedClause<OMPC_Reduction>,
VersionedClause<OMPC_Schedule>,
VersionedClause<OMPC_Severity, 60>,
VersionedClause<OMPC_Shared>,
VersionedClause<OMPC_ThreadLimit>,
VersionedClause<OMPC_UsesAllocators, 50>,
Expand Down Expand Up @@ -2502,7 +2482,6 @@ def OMP_TargetTeamsDistributeParallelForSimd
VersionedClause<OMPC_LastPrivate>,
VersionedClause<OMPC_Linear>,
VersionedClause<OMPC_Map>,
VersionedClause<OMPC_Message, 60>,
VersionedClause<OMPC_NonTemporal, 50>,
VersionedClause<OMPC_NoWait>,
VersionedClause<OMPC_NumTeams>,
Expand All @@ -2514,7 +2493,6 @@ def OMP_TargetTeamsDistributeParallelForSimd
VersionedClause<OMPC_Reduction>,
VersionedClause<OMPC_SafeLen>,
VersionedClause<OMPC_Schedule>,
VersionedClause<OMPC_Severity, 60>,
VersionedClause<OMPC_Shared>,
VersionedClause<OMPC_SimdLen>,
VersionedClause<OMPC_ThreadLimit>,
Expand Down
6 changes: 0 additions & 6 deletions llvm/include/llvm/Frontend/OpenMP/OMPKinds.def
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -472,8 +472,6 @@ __OMP_RTL(__kmpc_target_deinit, false, Void,)
__OMP_RTL(__kmpc_kernel_prepare_parallel, false, Void, VoidPtr)
__OMP_RTL(__kmpc_parallel_51, false, Void, IdentPtr, Int32, Int32, Int32, Int32,
VoidPtr, VoidPtr, VoidPtrPtr, SizeTy)
__OMP_RTL(__kmpc_parallel_60, false, Void, IdentPtr, Int32, Int32, Int32, Int32,
VoidPtr, VoidPtr, VoidPtrPtr, SizeTy, Int32, Int32, Int8Ptr)
__OMP_RTL(__kmpc_for_static_loop_4, false, Void, IdentPtr, VoidPtr, VoidPtr, Int32, Int32, Int32, Int8)
__OMP_RTL(__kmpc_for_static_loop_4u, false, Void, IdentPtr, VoidPtr, VoidPtr, Int32, Int32, Int32, Int8)
__OMP_RTL(__kmpc_for_static_loop_8, false, Void, IdentPtr, VoidPtr, VoidPtr, Int64, Int64, Int64, Int8)
Expand Down Expand Up @@ -1087,10 +1085,6 @@ __OMP_RTL_ATTRS(__kmpc_parallel_51, AlwaysInlineAttrs, AttributeSet(),
ParamAttrs(AttributeSet(), SExt, SExt, SExt, SExt,
AttributeSet(), AttributeSet(), AttributeSet(),
SizeTyExt))
__OMP_RTL_ATTRS(__kmpc_parallel_60, AlwaysInlineAttrs, AttributeSet(),
ParamAttrs(AttributeSet(), SExt, SExt, SExt, SExt,
AttributeSet(), AttributeSet(), AttributeSet(),
SizeTyExt, SExt, SExt, AttributeSet()))
__OMP_RTL_ATTRS(__kmpc_serialized_parallel, InaccessibleArgOnlyAttrs,
AttributeSet(), ParamAttrs(ReadOnlyPtrAttrs, SExt))
__OMP_RTL_ATTRS(__kmpc_end_serialized_parallel, InaccessibleArgOnlyAttrs,
Expand Down
6 changes: 0 additions & 6 deletions openmp/device/include/DeviceTypes.h
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -136,12 +136,6 @@ struct omp_lock_t {
void *Lock;
};

// see definition in openmp/runtime kmp.h
typedef enum omp_severity_t {
severity_warning = 1,
severity_fatal = 2
} omp_severity_t;

using InterWarpCopyFnTy = void (*)(void *src, int32_t warp_num);
using ShuffleReductFnTy = void (*)(void *rhsData, int16_t lane_id,
int16_t lane_offset, int16_t shortCircuit);
Expand Down
Loading