Skip to content

Commit 77da163

Browse files
ronliebronlieb
authored
[offload][OpenMP] Remove device code for num_threads strict (llvm#157893) (llvm#4022)
2 parents bc34693 + f700833 commit 77da163

File tree

12 files changed

+160
-6188
lines changed

12 files changed

+160
-6188
lines changed

‎clang/lib/CodeGen/CGOpenMPRuntimeGPU.cpp‎

Lines changed: 17 additions & 18 deletions
Original file line numberDiff line numberDiff line change
@@ -1429,9 +1429,9 @@ void CGOpenMPRuntimeGPU::emitParallelCall(
14291429
if (!CGF.HaveInsertPoint())
14301430
return;
14311431

1432-
auto &&ParallelGen = [this, Loc, OutlinedFn, CapturedVars, IfCond, NumThreads,
1433-
NumThreadsModifier, Severity, Message](
1434-
CodeGenFunction &CGF, PrePostActionTy &Action) {
1432+
auto &&ParallelGen = [this, Loc, OutlinedFn, CapturedVars, IfCond,
1433+
NumThreads](CodeGenFunction &CGF,
1434+
PrePostActionTy &Action) {
14351435
CGBuilderTy &Bld = CGF.Builder;
14361436
llvm::Value *NumThreadsVal = NumThreads;
14371437
llvm::Function *WFn = WrapperFunctionsMap[OutlinedFn];
@@ -1479,22 +1479,21 @@ void CGOpenMPRuntimeGPU::emitParallelCall(
14791479
NumThreadsVal = Bld.CreateZExtOrTrunc(NumThreadsVal, CGF.Int32Ty);
14801480

14811481
assert(IfCondVal && "Expected a value");
1482-
RuntimeFunction FnID = OMPRTL___kmpc_parallel_51;
14831482
llvm::Value *RTLoc = emitUpdateLocation(CGF, Loc);
1484-
llvm::SmallVector<llvm::Value *, 10> Args(
1485-
{RTLoc, getThreadID(CGF, Loc), IfCondVal, NumThreadsVal,
1486-
llvm::ConstantInt::get(CGF.Int32Ty, -1), FnPtr, ID,
1487-
Bld.CreateBitOrPointerCast(CapturedVarsAddrs.emitRawPointer(CGF),
1488-
CGF.VoidPtrPtrTy),
1489-
llvm::ConstantInt::get(CGM.SizeTy, CapturedVars.size())});
1490-
if (NumThreadsModifier == OMPC_NUMTHREADS_strict) {
1491-
FnID = OMPRTL___kmpc_parallel_60;
1492-
Args.append({llvm::ConstantInt::get(CGM.Int32Ty, true),
1493-
emitSeverityClause(Severity),
1494-
emitMessageClause(CGF, Message)});
1495-
}
1496-
CGF.EmitRuntimeCall(
1497-
OMPBuilder.getOrCreateRuntimeFunction(CGM.getModule(), FnID), Args);
1483+
llvm::Value *Args[] = {
1484+
RTLoc,
1485+
getThreadID(CGF, Loc),
1486+
IfCondVal,
1487+
NumThreadsVal,
1488+
llvm::ConstantInt::get(CGF.Int32Ty, -1),
1489+
FnPtr,
1490+
ID,
1491+
Bld.CreateBitOrPointerCast(CapturedVarsAddrs.emitRawPointer(CGF),
1492+
CGF.VoidPtrPtrTy),
1493+
llvm::ConstantInt::get(CGM.SizeTy, CapturedVars.size())};
1494+
CGF.EmitRuntimeCall(OMPBuilder.getOrCreateRuntimeFunction(
1495+
CGM.getModule(), OMPRTL___kmpc_parallel_51),
1496+
Args);
14981497
};
14991498

15001499
RegionCodeGenTy RCG(ParallelGen);

‎clang/lib/CodeGen/CGOpenMPRuntimeGPU.h‎

Lines changed: 3 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -224,11 +224,6 @@ class CGOpenMPRuntimeGPU : public CGOpenMPRuntime {
224224
/// Emits call to void __kmpc_push_num_threads(ident_t *loc, kmp_int32
225225
/// global_tid, kmp_int32 num_threads) to generate code for 'num_threads'
226226
/// clause.
227-
/// If the modifier 'strict' is given:
228-
/// Emits call to void __kmpc_push_num_threads_strict(ident_t *loc, kmp_int32
229-
/// global_tid, kmp_int32 num_threads, int severity, const char *message) to
230-
/// generate code for 'num_threads' clause with 'strict' modifier.
231-
/// \param NumThreads An integer value of threads.
232227
void emitNumThreadsClause(
233228
CodeGenFunction &CGF, llvm::Value *NumThreads, SourceLocation Loc,
234229
OpenMPNumThreadsClauseModifier Modifier = OMPC_NUMTHREADS_unknown,
@@ -297,11 +292,11 @@ class CGOpenMPRuntimeGPU : public CGOpenMPRuntime {
297292
/// \param NumThreads The value corresponding to the num_threads clause, if
298293
/// any, or nullptr.
299294
/// \param NumThreadsModifier The modifier of the num_threads clause, if
300-
/// any, ignored otherwise.
295+
/// any, ignored otherwise. Currently unused on the device.
301296
/// \param Severity The severity corresponding to the num_threads clause, if
302-
/// any, ignored otherwise.
297+
/// any, ignored otherwise. Currently unused on the device.
303298
/// \param Message The message string corresponding to the num_threads clause,
304-
/// if any, or nullptr.
299+
/// if any, or nullptr. Currently unused on the device.
305300
void emitParallelCall(CodeGenFunction &CGF, SourceLocation Loc,
306301
llvm::Function *OutlinedFn,
307302
ArrayRef<llvm::Value *> CapturedVars,

‎clang/test/AST/ByteCode/openmp.cpp‎

Lines changed: 3 additions & 3 deletions
Original file line numberDiff line numberDiff line change
@@ -17,12 +17,12 @@ extern int omp_get_thread_num(void);
1717

1818
int test2() {
1919
int x = 0;
20-
int device_result[N] = {0};
20+
int result[N] = {0};
2121

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

‎clang/test/OpenMP/amdgcn_target_parallel_num_threads_codegen.cpp‎

Lines changed: 0 additions & 1095 deletions
This file was deleted.

‎clang/test/OpenMP/nvptx_target_codegen.cpp‎

Lines changed: 66 additions & 847 deletions
Large diffs are not rendered by default.

‎clang/test/OpenMP/nvptx_target_parallel_num_threads_codegen.cpp‎

Lines changed: 35 additions & 725 deletions
Large diffs are not rendered by default.

‎clang/test/OpenMP/target_parallel_generic_loop_codegen.cpp‎

Lines changed: 18 additions & 418 deletions
Large diffs are not rendered by default.

‎clang/test/OpenMP/target_parallel_num_threads_strict_codegen.cpp‎

Lines changed: 0 additions & 2956 deletions
This file was deleted.

‎llvm/include/llvm/Frontend/OpenMP/OMP.td‎

Lines changed: 0 additions & 22 deletions
Original file line numberDiff line numberDiff line change
@@ -2064,11 +2064,9 @@ def OMP_TargetParallel : Directive<[Spelling<"target parallel">]> {
20642064
let allowedOnceClauses = [
20652065
VersionedClause<OMPC_DefaultMap>,
20662066
VersionedClause<OMPC_Device>,
2067-
VersionedClause<OMPC_Message, 60>,
20682067
VersionedClause<OMPC_NumThreads>,
20692068
VersionedClause<OMPC_OMPX_DynCGroupMem>,
20702069
VersionedClause<OMPC_ProcBind>,
2071-
VersionedClause<OMPC_Severity, 60>,
20722070
VersionedClause<OMPC_ThreadLimit, 51>,
20732071
];
20742072
let leafConstructs = [OMP_Target, OMP_Parallel];
@@ -2096,14 +2094,12 @@ def OMP_TargetParallelDo : Directive<[Spelling<"target parallel do">]> {
20962094
VersionedClause<OMPC_Collapse>,
20972095
VersionedClause<OMPC_DefaultMap>,
20982096
VersionedClause<OMPC_Device>,
2099-
VersionedClause<OMPC_Message, 60>,
21002097
VersionedClause<OMPC_NoWait>,
21012098
VersionedClause<OMPC_NumThreads>,
21022099
VersionedClause<OMPC_Order, 50>,
21032100
VersionedClause<OMPC_Ordered>,
21042101
VersionedClause<OMPC_ProcBind>,
21052102
VersionedClause<OMPC_Schedule>,
2106-
VersionedClause<OMPC_Severity, 60>,
21072103
];
21082104
let leafConstructs = [OMP_Target, OMP_Parallel, OMP_Do];
21092105
let category = CA_Executable;
@@ -2127,7 +2123,6 @@ def OMP_TargetParallelDoSimd
21272123
VersionedClause<OMPC_LastPrivate>,
21282124
VersionedClause<OMPC_Linear>,
21292125
VersionedClause<OMPC_Map>,
2130-
VersionedClause<OMPC_Message, 60>,
21312126
VersionedClause<OMPC_NonTemporal>,
21322127
VersionedClause<OMPC_NoWait>,
21332128
VersionedClause<OMPC_NumThreads>,
@@ -2138,7 +2133,6 @@ def OMP_TargetParallelDoSimd
21382133
VersionedClause<OMPC_Reduction>,
21392134
VersionedClause<OMPC_SafeLen>,
21402135
VersionedClause<OMPC_Schedule>,
2141-
VersionedClause<OMPC_Severity, 60>,
21422136
VersionedClause<OMPC_Shared>,
21432137
VersionedClause<OMPC_SimdLen>,
21442138
VersionedClause<OMPC_UsesAllocators>,
@@ -2163,7 +2157,6 @@ def OMP_TargetParallelFor : Directive<[Spelling<"target parallel for">]> {
21632157
VersionedClause<OMPC_LastPrivate>,
21642158
VersionedClause<OMPC_Linear>,
21652159
VersionedClause<OMPC_Map>,
2166-
VersionedClause<OMPC_Message, 60>,
21672160
VersionedClause<OMPC_NoWait>,
21682161
VersionedClause<OMPC_NumThreads>,
21692162
VersionedClause<OMPC_OMPX_Attribute>,
@@ -2173,7 +2166,6 @@ def OMP_TargetParallelFor : Directive<[Spelling<"target parallel for">]> {
21732166
VersionedClause<OMPC_ProcBind>,
21742167
VersionedClause<OMPC_Reduction>,
21752168
VersionedClause<OMPC_Schedule>,
2176-
VersionedClause<OMPC_Severity, 60>,
21772169
VersionedClause<OMPC_Shared>,
21782170
VersionedClause<OMPC_UsesAllocators, 50>,
21792171
];
@@ -2203,7 +2195,6 @@ def OMP_TargetParallelForSimd
22032195
VersionedClause<OMPC_LastPrivate>,
22042196
VersionedClause<OMPC_Linear>,
22052197
VersionedClause<OMPC_Map>,
2206-
VersionedClause<OMPC_Message, 60>,
22072198
VersionedClause<OMPC_NonTemporal, 50>,
22082199
VersionedClause<OMPC_NoWait>,
22092200
VersionedClause<OMPC_NumThreads>,
@@ -2215,7 +2206,6 @@ def OMP_TargetParallelForSimd
22152206
VersionedClause<OMPC_Reduction>,
22162207
VersionedClause<OMPC_SafeLen>,
22172208
VersionedClause<OMPC_Schedule>,
2218-
VersionedClause<OMPC_Severity, 60>,
22192209
VersionedClause<OMPC_Shared>,
22202210
VersionedClause<OMPC_SimdLen>,
22212211
VersionedClause<OMPC_UsesAllocators, 50>,
@@ -2251,13 +2241,11 @@ def OMP_target_parallel_loop : Directive<[Spelling<"target parallel loop">]> {
22512241
VersionedClause<OMPC_Collapse>,
22522242
VersionedClause<OMPC_Default>,
22532243
VersionedClause<OMPC_DefaultMap>,
2254-
VersionedClause<OMPC_Message, 60>,
22552244
VersionedClause<OMPC_NoWait>,
22562245
VersionedClause<OMPC_NumThreads>,
22572246
VersionedClause<OMPC_OMPX_DynCGroupMem>,
22582247
VersionedClause<OMPC_Order>,
22592248
VersionedClause<OMPC_ProcBind>,
2260-
VersionedClause<OMPC_Severity, 60>,
22612249
VersionedClause<OMPC_ThreadLimit, 51>,
22622250
];
22632251
let leafConstructs = [OMP_Target, OMP_Parallel, OMP_loop];
@@ -2288,14 +2276,12 @@ def OMP_TargetSimd : Directive<[Spelling<"target simd">]> {
22882276
VersionedClause<OMPC_Collapse>,
22892277
VersionedClause<OMPC_DefaultMap>,
22902278
VersionedClause<OMPC_Device>,
2291-
VersionedClause<OMPC_Message, 60>,
22922279
VersionedClause<OMPC_NumThreads>,
22932280
VersionedClause<OMPC_OMPX_DynCGroupMem>,
22942281
VersionedClause<OMPC_Order, 50>,
22952282
VersionedClause<OMPC_ProcBind>,
22962283
VersionedClause<OMPC_SafeLen>,
22972284
VersionedClause<OMPC_Schedule>,
2298-
VersionedClause<OMPC_Severity, 60>,
22992285
VersionedClause<OMPC_SimdLen>,
23002286
VersionedClause<OMPC_ThreadLimit, 51>,
23012287
];
@@ -2388,14 +2374,12 @@ def OMP_TargetTeamsDistributeParallelDo
23882374
VersionedClause<OMPC_DefaultMap>,
23892375
VersionedClause<OMPC_Device>,
23902376
VersionedClause<OMPC_DistSchedule>,
2391-
VersionedClause<OMPC_Message, 60>,
23922377
VersionedClause<OMPC_NoWait>,
23932378
VersionedClause<OMPC_NumTeams>,
23942379
VersionedClause<OMPC_NumThreads>,
23952380
VersionedClause<OMPC_Order, 50>,
23962381
VersionedClause<OMPC_ProcBind>,
23972382
VersionedClause<OMPC_Schedule>,
2398-
VersionedClause<OMPC_Severity, 60>,
23992383
VersionedClause<OMPC_ThreadLimit>,
24002384
];
24012385
let leafConstructs =
@@ -2429,15 +2413,13 @@ def OMP_TargetTeamsDistributeParallelDoSimd
24292413
VersionedClause<OMPC_DefaultMap>,
24302414
VersionedClause<OMPC_Device>,
24312415
VersionedClause<OMPC_DistSchedule>,
2432-
VersionedClause<OMPC_Message, 60>,
24332416
VersionedClause<OMPC_NoWait>,
24342417
VersionedClause<OMPC_NumTeams>,
24352418
VersionedClause<OMPC_NumThreads>,
24362419
VersionedClause<OMPC_Order, 50>,
24372420
VersionedClause<OMPC_ProcBind>,
24382421
VersionedClause<OMPC_SafeLen>,
24392422
VersionedClause<OMPC_Schedule>,
2440-
VersionedClause<OMPC_Severity, 60>,
24412423
VersionedClause<OMPC_SimdLen>,
24422424
VersionedClause<OMPC_ThreadLimit>,
24432425
];
@@ -2463,7 +2445,6 @@ def OMP_TargetTeamsDistributeParallelFor
24632445
VersionedClause<OMPC_IsDevicePtr>,
24642446
VersionedClause<OMPC_LastPrivate>,
24652447
VersionedClause<OMPC_Map>,
2466-
VersionedClause<OMPC_Message, 60>,
24672448
VersionedClause<OMPC_NoWait>,
24682449
VersionedClause<OMPC_NumTeams>,
24692450
VersionedClause<OMPC_NumThreads>,
@@ -2473,7 +2454,6 @@ def OMP_TargetTeamsDistributeParallelFor
24732454
VersionedClause<OMPC_ProcBind>,
24742455
VersionedClause<OMPC_Reduction>,
24752456
VersionedClause<OMPC_Schedule>,
2476-
VersionedClause<OMPC_Severity, 60>,
24772457
VersionedClause<OMPC_Shared>,
24782458
VersionedClause<OMPC_ThreadLimit>,
24792459
VersionedClause<OMPC_UsesAllocators, 50>,
@@ -2505,7 +2485,6 @@ def OMP_TargetTeamsDistributeParallelForSimd
25052485
VersionedClause<OMPC_LastPrivate>,
25062486
VersionedClause<OMPC_Linear>,
25072487
VersionedClause<OMPC_Map>,
2508-
VersionedClause<OMPC_Message, 60>,
25092488
VersionedClause<OMPC_NonTemporal, 50>,
25102489
VersionedClause<OMPC_NoWait>,
25112490
VersionedClause<OMPC_NumTeams>,
@@ -2517,7 +2496,6 @@ def OMP_TargetTeamsDistributeParallelForSimd
25172496
VersionedClause<OMPC_Reduction>,
25182497
VersionedClause<OMPC_SafeLen>,
25192498
VersionedClause<OMPC_Schedule>,
2520-
VersionedClause<OMPC_Severity, 60>,
25212499
VersionedClause<OMPC_Shared>,
25222500
VersionedClause<OMPC_SimdLen>,
25232501
VersionedClause<OMPC_ThreadLimit>,

‎llvm/include/llvm/Frontend/OpenMP/OMPKinds.def‎

Lines changed: 0 additions & 6 deletions
Original file line numberDiff line numberDiff line change
@@ -499,8 +499,6 @@ __OMP_RTL(__kmpc_kernel_prepare_parallel, false, Void, VoidPtr)
499499
__OMP_RTL(__kmpc_parallel_51, false, Void, IdentPtr, Int32, Int32, Int32, Int32,
500500
VoidPtr, VoidPtr, VoidPtrPtr, SizeTy)
501501
__OMP_RTL(__kmpc_parallel_spmd, false, Void, IdentPtr, Int32, VoidPtr, VoidPtrPtr, SizeTy)
502-
__OMP_RTL(__kmpc_parallel_60, false, Void, IdentPtr, Int32, Int32, Int32, Int32,
503-
VoidPtr, VoidPtr, VoidPtrPtr, SizeTy, Int32, Int32, Int8Ptr)
504502
__OMP_RTL(__kmpc_for_static_loop_4, false, Void, IdentPtr, VoidPtr, VoidPtr, Int32, Int32, Int32, Int8)
505503
__OMP_RTL(__kmpc_for_static_loop_4u, false, Void, IdentPtr, VoidPtr, VoidPtr, Int32, Int32, Int32, Int8)
506504
__OMP_RTL(__kmpc_for_static_loop_8, false, Void, IdentPtr, VoidPtr, VoidPtr, Int64, Int64, Int64, Int8)
@@ -1365,10 +1363,6 @@ __OMP_RTL_ATTRS(__kmpc_parallel_51, AlwaysInlineAttrs, AttributeSet(),
13651363
ParamAttrs(AttributeSet(), SExt, SExt, SExt, SExt,
13661364
AttributeSet(), AttributeSet(), AttributeSet(),
13671365
SizeTyExt))
1368-
__OMP_RTL_ATTRS(__kmpc_parallel_60, AlwaysInlineAttrs, AttributeSet(),
1369-
ParamAttrs(AttributeSet(), SExt, SExt, SExt, SExt,
1370-
AttributeSet(), AttributeSet(), AttributeSet(),
1371-
SizeTyExt, SExt, SExt, AttributeSet()))
13721366
__OMP_RTL_ATTRS(__kmpc_serialized_parallel, InaccessibleArgOnlyAttrs,
13731367
AttributeSet(), ParamAttrs(ReadOnlyPtrAttrs, SExt))
13741368
__OMP_RTL_ATTRS(__kmpc_end_serialized_parallel, InaccessibleArgOnlyAttrs,

0 commit comments

Comments
 (0)