[llvm] [offload] Clarify naming for block and thread numbers (PR #197419)
Robert Imschweiler via llvm-commits
llvm-commits at lists.llvm.org
Wed May 13 04:22:35 PDT 2026
https://github.com/ro-i created https://github.com/llvm/llvm-project/pull/197419
As discussed in
https://github.com/llvm/llvm-project/pull/195102#discussion_r3205402105 and on Slack.
>From b65b38a2b9218b21f3badae59e3de64021fdc22b Mon Sep 17 00:00:00 2001
From: Robert Imschweiler <robert.imschweiler at amd.com>
Date: Wed, 13 May 2026 06:20:05 -0500
Subject: [PATCH] [offload] Clarify naming for block and thread numbers
As discussed in
https://github.com/llvm/llvm-project/pull/195102#discussion_r3205402105
and on Slack.
---
offload/include/Shared/APITypes.h | 8 +-
offload/liboffload/src/OffloadImpl.cpp | 12 +--
offload/libomptarget/KernelLanguage/API.cpp | 12 +--
offload/libomptarget/interface.cpp | 18 ++---
offload/libomptarget/omptarget.cpp | 16 ++--
.../common/include/PluginInterface.h | 17 +++--
.../common/src/PluginInterface.cpp | 76 ++++++++++---------
.../common/src/RecordReplay.cpp | 8 +-
8 files changed, 87 insertions(+), 80 deletions(-)
diff --git a/offload/include/Shared/APITypes.h b/offload/include/Shared/APITypes.h
index 5e61bc7c842e7..6c9479af50938 100644
--- a/offload/include/Shared/APITypes.h
+++ b/offload/include/Shared/APITypes.h
@@ -106,10 +106,10 @@ struct KernelArgsTy {
uint64_t DynCGroupMemFallback : 2; // The fallback for dynamic cgroup mem.
uint64_t Unused : 60;
} Flags = {0, 0, 0, 0};
- // The number of teams (for x,y,z dimension).
- uint32_t NumTeams[3] = {0, 0, 0};
- // The number of threads (for x,y,z dimension).
- uint32_t ThreadLimit[3] = {0, 0, 0};
+ // User-requested number of blocks (for x,y,z dimension).
+ uint32_t UserNumBlocks[3] = {0, 0, 0};
+ // User-requested number of threads (for x,y,z dimension).
+ uint32_t UserThreadLimit[3] = {0, 0, 0};
uint32_t DynCGroupMem = 0; // Amount of dynamic cgroup memory requested.
};
static_assert(sizeof(KernelArgsTy().Flags) == sizeof(uint64_t),
diff --git a/offload/liboffload/src/OffloadImpl.cpp b/offload/liboffload/src/OffloadImpl.cpp
index 71712c2bc6b45..d4b9a62bf3bb9 100644
--- a/offload/liboffload/src/OffloadImpl.cpp
+++ b/offload/liboffload/src/OffloadImpl.cpp
@@ -1067,12 +1067,12 @@ Error olLaunchKernel_impl(ol_queue_handle_t Queue, ol_device_handle_t Device,
auto *QueueImpl = Queue ? Queue->AsyncInfo : nullptr;
AsyncInfoWrapperTy AsyncInfoWrapper(*DeviceImpl, QueueImpl);
KernelArgsTy LaunchArgs{};
- LaunchArgs.NumTeams[0] = LaunchSizeArgs->NumGroups.x;
- LaunchArgs.NumTeams[1] = LaunchSizeArgs->NumGroups.y;
- LaunchArgs.NumTeams[2] = LaunchSizeArgs->NumGroups.z;
- LaunchArgs.ThreadLimit[0] = LaunchSizeArgs->GroupSize.x;
- LaunchArgs.ThreadLimit[1] = LaunchSizeArgs->GroupSize.y;
- LaunchArgs.ThreadLimit[2] = LaunchSizeArgs->GroupSize.z;
+ LaunchArgs.UserNumBlocks[0] = LaunchSizeArgs->NumGroups.x;
+ LaunchArgs.UserNumBlocks[1] = LaunchSizeArgs->NumGroups.y;
+ LaunchArgs.UserNumBlocks[2] = LaunchSizeArgs->NumGroups.z;
+ LaunchArgs.UserThreadLimit[0] = LaunchSizeArgs->GroupSize.x;
+ LaunchArgs.UserThreadLimit[1] = LaunchSizeArgs->GroupSize.y;
+ LaunchArgs.UserThreadLimit[2] = LaunchSizeArgs->GroupSize.z;
LaunchArgs.DynCGroupMem = LaunchSizeArgs->DynSharedMemory;
KernelLaunchParamsTy Params;
diff --git a/offload/libomptarget/KernelLanguage/API.cpp b/offload/libomptarget/KernelLanguage/API.cpp
index ef1aad829e7bd..112b27b707e5a 100644
--- a/offload/libomptarget/KernelLanguage/API.cpp
+++ b/offload/libomptarget/KernelLanguage/API.cpp
@@ -60,12 +60,12 @@ unsigned llvmLaunchKernel(const void *func, dim3 gridDim, dim3 blockDim,
void *args, size_t sharedMem, void *stream) {
KernelArgsTy Args = {};
Args.DynCGroupMem = sharedMem;
- Args.NumTeams[0] = gridDim.x;
- Args.NumTeams[1] = gridDim.y;
- Args.NumTeams[2] = gridDim.z;
- Args.ThreadLimit[0] = blockDim.x;
- Args.ThreadLimit[1] = blockDim.y;
- Args.ThreadLimit[2] = blockDim.z;
+ Args.UserNumBlocks[0] = gridDim.x;
+ Args.UserNumBlocks[1] = gridDim.y;
+ Args.UserNumBlocks[2] = gridDim.z;
+ Args.UserThreadLimit[0] = blockDim.x;
+ Args.UserThreadLimit[1] = blockDim.y;
+ Args.UserThreadLimit[2] = blockDim.z;
Args.ArgPtrs = reinterpret_cast<void **>(args);
Args.Flags.IsCUDA = true;
return __tgt_target_kernel(nullptr, 0, gridDim.x, blockDim.x, func, &Args);
diff --git a/offload/libomptarget/interface.cpp b/offload/libomptarget/interface.cpp
index f65ca3cadee8c..64ff078e3ec46 100644
--- a/offload/libomptarget/interface.cpp
+++ b/offload/libomptarget/interface.cpp
@@ -308,12 +308,12 @@ static KernelArgsTy *upgradeKernelArgs(KernelArgsTy *KernelArgs,
LocalKernelArgs.Tripcount = KernelArgs->Tripcount;
LocalKernelArgs.Flags = KernelArgs->Flags;
LocalKernelArgs.DynCGroupMem = 0;
- LocalKernelArgs.NumTeams[0] = NumTeams;
- LocalKernelArgs.NumTeams[1] = 1;
- LocalKernelArgs.NumTeams[2] = 1;
- LocalKernelArgs.ThreadLimit[0] = ThreadLimit;
- LocalKernelArgs.ThreadLimit[1] = 1;
- LocalKernelArgs.ThreadLimit[2] = 1;
+ LocalKernelArgs.UserNumBlocks[0] = NumTeams;
+ LocalKernelArgs.UserNumBlocks[1] = 1;
+ LocalKernelArgs.UserNumBlocks[2] = 1;
+ LocalKernelArgs.UserThreadLimit[0] = ThreadLimit;
+ LocalKernelArgs.UserThreadLimit[1] = 1;
+ LocalKernelArgs.UserThreadLimit[2] = 1;
return &LocalKernelArgs;
}
@@ -325,8 +325,8 @@ static KernelArgsTy *upgradeKernelArgs(KernelArgsTy *KernelArgs,
if (Val[2] == 0)
Val[2] = 1;
};
- CorrectMultiDim(KernelArgs->ThreadLimit);
- CorrectMultiDim(KernelArgs->NumTeams);
+ CorrectMultiDim(KernelArgs->UserThreadLimit);
+ CorrectMultiDim(KernelArgs->UserNumBlocks);
// Version 3 put the implicit argument at the front with no storage.
if (KernelArgs->Version == OMP_KERNEL_ARG_MIN_VERSION_WITH_DYN_PTR) {
@@ -384,7 +384,7 @@ static inline int targetKernel(ident_t *Loc, int64_t DeviceId, int32_t NumTeams,
bool IsTeams = NumTeams != -1;
if (!IsTeams)
- KernelArgs->NumTeams[0] = NumTeams = 1;
+ KernelArgs->UserNumBlocks[0] = NumTeams = 1;
KernelArgsTy LocalKernelArgs;
UpgradedArgBuffersTy UpgradedBufs;
diff --git a/offload/libomptarget/omptarget.cpp b/offload/libomptarget/omptarget.cpp
index 82a0ed73317d5..c2456920ebc1b 100644
--- a/offload/libomptarget/omptarget.cpp
+++ b/offload/libomptarget/omptarget.cpp
@@ -2337,13 +2337,13 @@ int target(ident_t *Loc, DeviceTy &Device, void *HostPtr,
TIMESCOPE_WITH_DETAILS_AND_IDENT(
"Kernel Target",
"NumArguments=" + std::to_string(KernelArgs.NumArgs) +
- ";NumTeams=" + std::to_string(KernelArgs.NumTeams[0]) +
+ ";NumTeams=" + std::to_string(KernelArgs.UserNumBlocks[0]) +
";TripCount=" + std::to_string(KernelArgs.Tripcount),
Loc);
#ifdef OMPT_SUPPORT
/// RAII to establish tool anchors before and after kernel launch
- int32_t NumTeams = KernelArgs.NumTeams[0];
+ int32_t NumTeams = KernelArgs.UserNumBlocks[0];
// No need to guard this with OMPT_IF_BUILT
InterfaceRAII TargetSubmitRAII(
RegionInterface.getCallbacks<ompt_callback_target_submit>(), NumTeams);
@@ -2474,12 +2474,12 @@ int target_replay(ident_t *Loc, DeviceTy &Device, void *HostPtr,
KernelArgs.Version = OMP_KERNEL_ARG_VERSION;
KernelArgs.NumArgs = NumArgs;
KernelArgs.Tripcount = LoopTripCount;
- KernelArgs.NumTeams[0] = NumTeams;
- KernelArgs.NumTeams[1] = 1;
- KernelArgs.NumTeams[2] = 1;
- KernelArgs.ThreadLimit[0] = ThreadLimit;
- KernelArgs.ThreadLimit[1] = 1;
- KernelArgs.ThreadLimit[2] = 1;
+ KernelArgs.UserNumBlocks[0] = NumTeams;
+ KernelArgs.UserNumBlocks[1] = 1;
+ KernelArgs.UserNumBlocks[2] = 1;
+ KernelArgs.UserThreadLimit[0] = ThreadLimit;
+ KernelArgs.UserThreadLimit[1] = 1;
+ KernelArgs.UserThreadLimit[2] = 1;
KernelArgs.DynCGroupMem = SharedMemorySize;
KernelExtraArgsTy KernelExtraArgs{};
diff --git a/offload/plugins-nextgen/common/include/PluginInterface.h b/offload/plugins-nextgen/common/include/PluginInterface.h
index f55fe736cea11..7ba059969bcc7 100644
--- a/offload/plugins-nextgen/common/include/PluginInterface.h
+++ b/offload/plugins-nextgen/common/include/PluginInterface.h
@@ -517,17 +517,20 @@ struct GenericKernelTy {
KernelLaunchEnvironmentTy *KernelLaunchEnvironment,
uint32_t Version) const;
- /// Get the number of threads and blocks for the kernel based on the
- /// user-defined threads and block clauses.
- uint32_t getNumThreads(GenericDeviceTy &GenericDevice,
- uint32_t ThreadLimitClause[3]) const;
+ /// Get the effective number of threads for the kernel based on the
+ /// user-defined number of threads.
+ uint32_t getEffectiveNumThreads(GenericDeviceTy &GenericDevice,
+ uint32_t ThreadLimitClause[3]) const;
+ /// Get the effective number of blocks for the kernel based on the
+ /// user-defined number of blocks and the loop trip count.
/// The number of threads \p NumThreads can be adjusted by this method.
/// \p IsNumThreadsFromUser is true is \p NumThreads is defined by user via
/// thread_limit clause.
- uint32_t getNumBlocks(GenericDeviceTy &GenericDevice,
- uint32_t BlockLimitClause[3], uint64_t LoopTripCount,
- uint32_t &NumThreads, bool IsNumThreadsFromUser) const;
+ uint32_t getEffectiveNumBlocks(GenericDeviceTy &GenericDevice,
+ uint32_t UserNumBlocks[3],
+ uint64_t LoopTripCount, uint32_t &NumThreads,
+ bool IsNumThreadsFromUser) const;
/// Indicate if the kernel works in Generic SPMD, Generic, No-Loop
/// or SPMD mode.
diff --git a/offload/plugins-nextgen/common/src/PluginInterface.cpp b/offload/plugins-nextgen/common/src/PluginInterface.cpp
index 12ccebdfbc57a..77f75285db298 100644
--- a/offload/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/offload/plugins-nextgen/common/src/PluginInterface.cpp
@@ -243,23 +243,27 @@ Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice, void **ArgPtrs,
llvm::SmallVector<void *, 16> Args;
llvm::SmallVector<void *, 16> Ptrs;
- uint32_t NumThreads[3] = {KernelArgs.ThreadLimit[0],
- KernelArgs.ThreadLimit[1],
- KernelArgs.ThreadLimit[2]};
- uint32_t NumBlocks[3] = {KernelArgs.NumTeams[0], KernelArgs.NumTeams[1],
- KernelArgs.NumTeams[2]};
+ uint32_t EffectiveNumThreads[3] = {KernelArgs.UserThreadLimit[0],
+ KernelArgs.UserThreadLimit[1],
+ KernelArgs.UserThreadLimit[2]};
+ uint32_t EffectiveNumBlocks[3] = {KernelArgs.UserNumBlocks[0],
+ KernelArgs.UserNumBlocks[1],
+ KernelArgs.UserNumBlocks[2]};
if (!isBareMode()) {
assert(
- NumThreads[1] == 1 && NumThreads[2] == 1 && NumBlocks[1] == 1 &&
- NumBlocks[2] == 1 &&
+ EffectiveNumThreads[1] == 1 && EffectiveNumThreads[2] == 1 &&
+ EffectiveNumBlocks[1] == 1 && EffectiveNumBlocks[2] == 1 &&
"Non-bare mode should only use the first thread and block dimensions");
- NumThreads[0] = getNumThreads(GenericDevice, NumThreads);
- NumBlocks[0] = getNumBlocks(GenericDevice, NumBlocks, KernelArgs.Tripcount,
- NumThreads[0], KernelArgs.ThreadLimit[0] > 0);
+ EffectiveNumThreads[0] =
+ getEffectiveNumThreads(GenericDevice, EffectiveNumThreads);
+ EffectiveNumBlocks[0] = getEffectiveNumBlocks(
+ GenericDevice, EffectiveNumBlocks, KernelArgs.Tripcount,
+ EffectiveNumThreads[0], KernelArgs.UserThreadLimit[0] > 0);
}
auto DynBlockMemConfOrErr = prepareBlockMemory(
- GenericDevice, KernelArgs, NumBlocks[0] * NumBlocks[1] * NumBlocks[2]);
+ GenericDevice, KernelArgs,
+ EffectiveNumBlocks[0] * EffectiveNumBlocks[1] * EffectiveNumBlocks[2]);
if (!DynBlockMemConfOrErr)
return DynBlockMemConfOrErr.takeError();
@@ -285,8 +289,8 @@ Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice, void **ArgPtrs,
Args, Ptrs, *KernelLaunchEnvOrErr, KernelArgs.Version);
}
- if (auto Err =
- printLaunchInfo(GenericDevice, KernelArgs, NumThreads, NumBlocks))
+ if (auto Err = printLaunchInfo(GenericDevice, KernelArgs, EffectiveNumThreads,
+ EffectiveNumBlocks))
return Err;
RecordReplayTy::HandleTy RRHandle;
@@ -298,16 +302,16 @@ Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice, void **ArgPtrs,
// Record the kernel prologue data before kernel launch.
auto RRHandleOrErr = RecordReplay->recordPrologue(
- *this, KernelArgs, KernelExtraArgs, LaunchParams, NumBlocks, NumThreads,
- DynBlockMemConf.NativeSize);
+ *this, KernelArgs, KernelExtraArgs, LaunchParams, EffectiveNumBlocks,
+ EffectiveNumThreads, DynBlockMemConf.NativeSize);
if (!RRHandleOrErr)
return RRHandleOrErr.takeError();
RRHandle = *RRHandleOrErr;
}
- if (auto Err = launchImpl(GenericDevice, NumThreads, NumBlocks,
- DynBlockMemConf.NativeSize, KernelArgs,
- LaunchParams, AsyncInfoWrapper))
+ if (auto Err = launchImpl(GenericDevice, EffectiveNumThreads,
+ EffectiveNumBlocks, DynBlockMemConf.NativeSize,
+ KernelArgs, LaunchParams, AsyncInfoWrapper))
return Err;
if (RecordReplay) {
@@ -356,7 +360,8 @@ GenericKernelTy::prepareArgs(GenericDeviceTy &GenericDevice, void **ArgPtrs,
return KernelLaunchParamsTy{sizeof(void *) * NumArgs, &Args[0], &Ptrs[0]};
}
-uint32_t GenericKernelTy::getNumThreads(GenericDeviceTy &GenericDevice,
+uint32_t
+GenericKernelTy::getEffectiveNumThreads(GenericDeviceTy &GenericDevice,
uint32_t ThreadLimitClause[3]) const {
assert(!isBareMode() && "bare kernel should not call this function");
@@ -371,24 +376,23 @@ uint32_t GenericKernelTy::getNumThreads(GenericDeviceTy &GenericDevice,
: PreferredNumThreads);
}
-uint32_t GenericKernelTy::getNumBlocks(GenericDeviceTy &GenericDevice,
- uint32_t NumTeamsClause[3],
- uint64_t LoopTripCount,
- uint32_t &NumThreads,
- bool IsNumThreadsFromUser) const {
+uint32_t GenericKernelTy::getEffectiveNumBlocks(
+ GenericDeviceTy &GenericDevice, uint32_t UserNumBlocks[3],
+ uint64_t LoopTripCount, uint32_t &NumThreads,
+ bool IsNumThreadsFromUser) const {
assert(!isBareMode() && "bare kernel should not call this function");
- assert(NumTeamsClause[1] == 1 && NumTeamsClause[2] == 1 &&
+ assert(UserNumBlocks[1] == 1 && UserNumBlocks[2] == 1 &&
"Multi dimensional launch not supported yet.");
- if (NumTeamsClause[0] > 0) {
+ if (UserNumBlocks[0] > 0) {
// TODO: We need to honor any value and consequently allow more than the
// block limit. For this we might need to start multiple kernels or let the
// blocks start again until the requested number has been started.
- return std::min(NumTeamsClause[0], GenericDevice.getBlockLimit());
+ return std::min(UserNumBlocks[0], GenericDevice.getBlockLimit());
}
- // Return the number of teams required to cover the loop iterations.
+ // Return the number of blocks required to cover the loop iterations.
if (isNoLoopMode())
return LoopTripCount > 0 ? (((LoopTripCount - 1) / NumThreads) + 1) : 1;
@@ -397,10 +401,10 @@ uint32_t GenericKernelTy::getNumBlocks(GenericDeviceTy &GenericDevice,
if (LoopTripCount > 0) {
if (isSPMDMode()) {
// We have a combined construct, i.e. `target teams distribute
- // parallel for [simd]`. We launch so many teams so that each thread
+ // parallel for [simd]`. We launch so many blocks so that each thread
// will execute one iteration of the loop; rounded up to the nearest
- // integer. However, if that results in too few teams, we artificially
- // reduce the thread count per team to increase the outer parallelism.
+ // integer. However, if that results in too few blocks, we artificially
+ // reduce the thread count per block to increase the outer parallelism.
auto MinThreads = GenericDevice.getMinThreadsForLowTripCountLoop();
MinThreads = std::min(MinThreads, NumThreads);
@@ -408,13 +412,13 @@ uint32_t GenericKernelTy::getNumBlocks(GenericDeviceTy &GenericDevice,
[[maybe_unused]] auto OldNumThreads = NumThreads;
if (LoopTripCount >= DefaultNumBlocks * NumThreads ||
IsNumThreadsFromUser) {
- // Enough parallelism for teams and threads.
+ // Enough parallelism for blocks and threads.
TripCountNumBlocks = ((LoopTripCount - 1) / NumThreads) + 1;
assert(IsNumThreadsFromUser ||
TripCountNumBlocks >= DefaultNumBlocks &&
"Expected sufficient outer parallelism.");
} else if (LoopTripCount >= DefaultNumBlocks * MinThreads) {
- // Enough parallelism for teams, limit threads.
+ // Enough parallelism for blocks, limit threads.
// This case is hard; for now, we force "full warps":
// First, compute a thread count assuming DefaultNumBlocks.
@@ -429,7 +433,7 @@ uint32_t GenericKernelTy::getNumBlocks(GenericDeviceTy &GenericDevice,
"Expected sufficient inner parallelism.");
TripCountNumBlocks = ((LoopTripCount - 1) / NumThreads) + 1;
} else {
- // Not enough parallelism for teams and threads, limit both.
+ // Not enough parallelism for blocks and threads, limit both.
NumThreads = std::min(NumThreads, MinThreads);
TripCountNumBlocks = ((LoopTripCount - 1) / NumThreads) + 1;
}
@@ -442,7 +446,7 @@ uint32_t GenericKernelTy::getNumBlocks(GenericDeviceTy &GenericDevice,
assert((isGenericMode() || isGenericSPMDMode()) &&
"Unexpected execution mode!");
// If we reach this point, then we have a non-combined construct, i.e.
- // `teams distribute` with a nested `parallel for` and each team is
+ // `teams distribute` with a nested `parallel for` and each block is
// assigned one iteration of the `distribute` loop. E.g.:
//
// #pragma omp target teams distribute
@@ -451,7 +455,7 @@ uint32_t GenericKernelTy::getNumBlocks(GenericDeviceTy &GenericDevice,
// for(...) {}
// }
//
- // Threads within a team will execute the iterations of the `parallel`
+ // Threads within a block will execute the iterations of the `parallel`
// loop.
TripCountNumBlocks = LoopTripCount;
}
diff --git a/offload/plugins-nextgen/common/src/RecordReplay.cpp b/offload/plugins-nextgen/common/src/RecordReplay.cpp
index 8c4b012fdc9e2..1d4acbc5bbd4c 100644
--- a/offload/plugins-nextgen/common/src/RecordReplay.cpp
+++ b/offload/plugins-nextgen/common/src/RecordReplay.cpp
@@ -253,15 +253,15 @@ Error NativeRecordReplayTy::recordDescImpl(
// Add minimum and maximum for allowed number of teams. If zero, it means
// there was no restriction provided by the program.
json::Array JsonTeamsLimits;
- JsonTeamsLimits.push_back(KernelArgs.NumTeams[0]);
- JsonTeamsLimits.push_back(KernelArgs.NumTeams[0]);
+ JsonTeamsLimits.push_back(KernelArgs.UserNumBlocks[0]);
+ JsonTeamsLimits.push_back(KernelArgs.UserNumBlocks[0]);
JsonKernelInfo["TeamsLimits"] = json::Value(std::move(JsonTeamsLimits));
// Add minimum and maximum for allowed number of threads. If zero, it means
// there was no restriction provided by the program.
json::Array JsonThreadsLimits;
- JsonThreadsLimits.push_back(uint32_t(KernelArgs.ThreadLimit[0] > 0));
- JsonThreadsLimits.push_back(KernelArgs.ThreadLimit[0]);
+ JsonThreadsLimits.push_back(uint32_t(KernelArgs.UserThreadLimit[0] > 0));
+ JsonThreadsLimits.push_back(KernelArgs.UserThreadLimit[0]);
JsonKernelInfo["ThreadsLimits"] = json::Value(std::move(JsonThreadsLimits));
json::Array JsonArgPtrs;
More information about the llvm-commits
mailing list