[llvm-branch-commits] [llvm] [offload][omp] Move strict threads & groups computation to libomptarget (PR #222607)
Alex Duran via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Sep 10 04:53:22 PDT 2026
https://github.com/adurang updated https://github.com/llvm/llvm-project/pull/222607
>From 05d870b4cac9009825b61a18694d5f555aeaf60d Mon Sep 17 00:00:00 2001
From: "Duran, Alex" <alejandro.duran at intel.com>
Date: Thu, 10 Sep 2026 03:40:36 -0700
Subject: [PATCH] [offload][omp] Move strict threads & groups computation to
libomptarget
---
offload/include/device.h | 41 +++-
offload/liboffload/src/OffloadImpl.cpp | 2 -
offload/libompaccsupport/device.cpp | 180 +++++++++++++++++-
.../common/include/PluginInterface.h | 70 +------
.../common/src/PluginInterface.cpp | 147 --------------
offload/test/offloading/ompx_bare.c | 2 +-
offload/test/offloading/ompx_bare_gridsize.c | 18 +-
.../test/offloading/ompx_bare_multi_dim.cpp | 2 +-
8 files changed, 235 insertions(+), 227 deletions(-)
diff --git a/offload/include/device.h b/offload/include/device.h
index 3d76b742a59b9..9c5fa061ec608 100644
--- a/offload/include/device.h
+++ b/offload/include/device.h
@@ -39,12 +39,51 @@
using GenericPluginTy = llvm::omp::target::plugin::GenericPluginTy;
using DeviceInfo = llvm::omp::target::plugin::DeviceInfo;
using InfoTreeNode = llvm::omp::target::plugin::InfoTreeNode;
-using KernelLaunchInfoTy = llvm::omp::target::plugin::KernelLaunchInfoTy;
// Forward declarations.
struct __tgt_bin_desc;
struct __tgt_target_table;
+/// Kernel launch-geometry properties.
+struct KernelLaunchInfoTy {
+ uint32_t MaxNumThreads = 0;
+ uint32_t PreferredNumThreads = 0;
+ uint32_t ReductionDataSize = 0;
+ llvm::omp::OMPTgtExecModeFlags Mode = llvm::omp::OMP_TGT_EXEC_MODE_BARE;
+
+ bool isBareMode() const { return Mode == llvm::omp::OMP_TGT_EXEC_MODE_BARE; }
+ bool isGenericMode() const {
+ return Mode == llvm::omp::OMP_TGT_EXEC_MODE_GENERIC;
+ }
+ bool isGenericSPMDMode() const {
+ return Mode == llvm::omp::OMP_TGT_EXEC_MODE_GENERIC_SPMD;
+ }
+ bool isSPMDMode() const { return Mode == llvm::omp::OMP_TGT_EXEC_MODE_SPMD; }
+ bool isNoLoopMode() const {
+ return Mode == llvm::omp::OMP_TGT_EXEC_MODE_SPMD_NO_LOOP;
+ }
+
+ static const char *getExecutionModeName(llvm::omp::OMPTgtExecModeFlags Mode) {
+ switch (Mode) {
+ case llvm::omp::OMP_TGT_EXEC_MODE_BARE:
+ return "BARE";
+ case llvm::omp::OMP_TGT_EXEC_MODE_SPMD:
+ return "SPMD";
+ case llvm::omp::OMP_TGT_EXEC_MODE_GENERIC:
+ return "Generic";
+ case llvm::omp::OMP_TGT_EXEC_MODE_GENERIC_SPMD:
+ return "Generic-SPMD";
+ case llvm::omp::OMP_TGT_EXEC_MODE_SPMD_NO_LOOP:
+ return "SPMD-No-Loop";
+ }
+ return "Unknown";
+ }
+
+ const char *getExecutionModeName() const {
+ return getExecutionModeName(Mode);
+ }
+};
+
struct DeviceTy {
int32_t DeviceID;
GenericPluginTy *RTL;
diff --git a/offload/liboffload/src/OffloadImpl.cpp b/offload/liboffload/src/OffloadImpl.cpp
index 11c07483b38aa..11f1c11dc606d 100644
--- a/offload/liboffload/src/OffloadImpl.cpp
+++ b/offload/liboffload/src/OffloadImpl.cpp
@@ -1279,8 +1279,6 @@ Error olLaunchKernel_impl(ol_queue_handle_t Queue, ol_device_handle_t Device,
LaunchArgs.UserThreadLimit[1] = LaunchSizeArgs->GroupSize.y;
LaunchArgs.UserThreadLimit[2] = LaunchSizeArgs->GroupSize.z;
LaunchArgs.DynCGroupMem = LaunchSizeArgs->DynSharedMemory;
- LaunchArgs.Flags.StrictBlocks = true;
- LaunchArgs.Flags.StrictThreads = true;
while (Properties && Properties->type != OL_KERNEL_LAUNCH_PROP_TYPE_NONE) {
switch (Properties->type) {
diff --git a/offload/libompaccsupport/device.cpp b/offload/libompaccsupport/device.cpp
index ff51d346436a9..bea49e6e4466d 100644
--- a/offload/libompaccsupport/device.cpp
+++ b/offload/libompaccsupport/device.cpp
@@ -24,12 +24,14 @@
#include "Shared/EnvironmentVar.h"
#include "llvm/Frontend/OpenMP/OMPConstants.h"
#include "llvm/Support/Error.h"
+#include "llvm/Support/MathExtras.h"
#include <algorithm>
#include <cassert>
#include <climits>
#include <cstdint>
#include <cstdio>
+#include <limits>
#include <mutex>
#include <string>
#include <thread>
@@ -387,6 +389,122 @@ static void resolveKernelLaunchParams(void **const TgtArgs,
LaunchArgs.Args = &Ptrs[0];
}
+/// Get the effective number of threads for the kernel based on the
+/// user-defined number of threads.
+static uint32_t getEffectiveNumThreads(GenericDeviceTy &GenericDevice,
+ uint32_t UserThreadLimit,
+ const KernelLaunchInfoTy &KernelEnv) {
+ assert(!KernelEnv.isBareMode() &&
+ "bare kernel should not call this function");
+
+ if (UserThreadLimit > 0 && KernelEnv.isGenericMode())
+ UserThreadLimit += GenericDevice.getWarpSize();
+
+ return std::min(KernelEnv.MaxNumThreads, (UserThreadLimit > 0)
+ ? UserThreadLimit
+ : KernelEnv.PreferredNumThreads);
+}
+
+/// 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 EffectiveNumThreads can be adjusted by this
+/// method. \p IsNumThreadsFromUser is true if \p EffectiveNumThreads is
+/// defined by the user via the thread_limit clause.
+static uint32_t
+getEffectiveNumBlocks(GenericDeviceTy &GenericDevice, uint32_t UserNumBlocks,
+ uint64_t LoopTripCount, uint32_t &EffectiveNumThreads,
+ bool IsNumThreadsStrict, bool IsNumThreadsFromUser,
+ const KernelLaunchInfoTy &KernelEnv) {
+ assert(!KernelEnv.isBareMode() &&
+ "bare kernel should not call this function");
+
+ // NOTE: This clamps the user-requested number of blocks to the device limit
+ // rather than honoring it exactly, which is non-standard behavior. Truly
+ // honoring an arbitrary value would require launching multiple kernels or
+ // reusing blocks until the requested count has been served.
+ if (UserNumBlocks > 0)
+ return std::min(UserNumBlocks,
+ GenericDevice.getBlockLimit(EffectiveNumThreads));
+
+ // Return the number of blocks required to cover the loop iterations.
+ if (KernelEnv.isNoLoopMode())
+ return LoopTripCount > 0 ? (((LoopTripCount - 1) / EffectiveNumThreads) + 1)
+ : 1;
+
+ uint64_t DefaultNumBlocks = GenericDevice.getDefaultNumBlocks();
+ uint64_t TripCountNumBlocks = std::numeric_limits<uint64_t>::max();
+ if (LoopTripCount > 0) {
+ if (KernelEnv.isSPMDMode()) {
+ // We have a combined construct, i.e. `target teams distribute
+ // 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 blocks, we artificially
+ // reduce the thread count per block to increase the outer parallelism.
+ auto MinThreads = GenericDevice.getMinThreadsForLowTripCountLoop();
+ MinThreads = std::min(MinThreads, EffectiveNumThreads);
+
+ // Honor the thread_limit clause; only lower the number of threads.
+ [[maybe_unused]] auto OldNumThreads = EffectiveNumThreads;
+ if (LoopTripCount >= DefaultNumBlocks * EffectiveNumThreads ||
+ IsNumThreadsFromUser || IsNumThreadsStrict) {
+ // Enough parallelism for blocks and threads.
+ TripCountNumBlocks = ((LoopTripCount - 1) / EffectiveNumThreads) + 1;
+ assert(IsNumThreadsFromUser ||
+ TripCountNumBlocks >= DefaultNumBlocks &&
+ "Expected sufficient outer parallelism.");
+ } else if (LoopTripCount >= DefaultNumBlocks * MinThreads) {
+ // Enough parallelism for blocks, limit threads.
+
+ // This case is hard; for now, we force "full warps":
+ // First, compute a thread count assuming DefaultNumBlocks.
+ auto NumThreadsDefaultBlocks =
+ (LoopTripCount + DefaultNumBlocks - 1) / DefaultNumBlocks;
+ // Now get a power of two that is larger or equal.
+ auto NumThreadsDefaultBlocksP2 =
+ llvm::PowerOf2Ceil(NumThreadsDefaultBlocks);
+ // Do not increase a thread limit given be the user.
+ EffectiveNumThreads =
+ std::min(EffectiveNumThreads, uint32_t(NumThreadsDefaultBlocksP2));
+ assert(EffectiveNumThreads >= MinThreads &&
+ "Expected sufficient inner parallelism.");
+ TripCountNumBlocks = ((LoopTripCount - 1) / EffectiveNumThreads) + 1;
+ } else {
+ // Not enough parallelism for blocks and threads, limit both.
+ EffectiveNumThreads = std::min(EffectiveNumThreads, MinThreads);
+ TripCountNumBlocks = ((LoopTripCount - 1) / EffectiveNumThreads) + 1;
+ }
+
+ assert(EffectiveNumThreads * TripCountNumBlocks >= LoopTripCount &&
+ "Expected sufficient parallelism");
+ assert(OldNumThreads >= EffectiveNumThreads &&
+ "Number of threads cannot be increased!");
+ } else {
+ assert((KernelEnv.isGenericMode() || KernelEnv.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 block is
+ // assigned one iteration of the `distribute` loop. E.g.:
+ //
+ // #pragma omp target teams distribute
+ // for(...loop_tripcount...) {
+ // #pragma omp parallel for
+ // for(...) {}
+ // }
+ //
+ // Threads within a block will execute the iterations of the `parallel`
+ // loop.
+ TripCountNumBlocks = LoopTripCount;
+ }
+ }
+
+ uint32_t PreferredNumBlocks = TripCountNumBlocks;
+ // If the loops are long running we rather reuse blocks than spawn too many.
+ if (GenericDevice.getReuseBlocksForHighTripCount())
+ PreferredNumBlocks = std::min(TripCountNumBlocks, DefaultNumBlocks);
+ return std::min(PreferredNumBlocks,
+ GenericDevice.getBlockLimit(EffectiveNumThreads));
+}
+
// Run region on device
int32_t DeviceTy::launchKernel(void *TgtEntryPtr, void **TgtVarsPtr,
ptrdiff_t *TgtOffsets, KernelArgsTy &KernelArgs,
@@ -404,10 +522,57 @@ int32_t DeviceTy::launchKernel(void *TgtEntryPtr, void **TgtVarsPtr,
llvm::copy(KernelArgs.UserNumBlocks, LaunchArgs.UserNumBlocks);
llvm::copy(KernelArgs.UserThreadLimit, LaunchArgs.UserThreadLimit);
LaunchArgs.Flags.Cooperative = KernelArgs.Flags.Cooperative;
- LaunchArgs.Flags.StrictBlocks = KernelArgs.Flags.StrictBlocks;
- LaunchArgs.Flags.StrictThreads = KernelArgs.Flags.StrictThreads;
LaunchArgs.Flags.DynCGroupMemFallback = KernelArgs.Flags.DynCGroupMemFallback;
- LaunchArgs.KernelEnvironment = getKernelLaunchInfo(TgtEntryPtr);
+
+ KernelLaunchInfoTy KernelEnv = getKernelLaunchInfo(TgtEntryPtr);
+ LaunchArgs.KernelEnvironment.ReductionDataSize = KernelEnv.ReductionDataSize;
+
+ const bool StrictBlocks = KernelArgs.Flags.StrictBlocks;
+ const bool StrictThreads = KernelArgs.Flags.StrictThreads;
+
+ // Multidimensional is only supported with bare mode for now.
+ assert(KernelEnv.isBareMode() ||
+ LaunchArgs.UserThreadLimit[1] == 1 &&
+ LaunchArgs.UserThreadLimit[2] == 1 &&
+ LaunchArgs.UserNumBlocks[1] == 1 &&
+ LaunchArgs.UserNumBlocks[2] == 1 &&
+ "Non-bare mode should only use the first thread and block "
+ "dimensions");
+
+ assert(!StrictBlocks ||
+ LaunchArgs.UserNumBlocks[0] > 0 && LaunchArgs.UserNumBlocks[1] > 0 &&
+ LaunchArgs.UserNumBlocks[2] > 0 &&
+ "Strict requires number of blocks greater than zero");
+ assert(!StrictThreads ||
+ LaunchArgs.UserThreadLimit[0] > 0 &&
+ LaunchArgs.UserThreadLimit[1] > 0 &&
+ LaunchArgs.UserThreadLimit[2] > 0 &&
+ "Strict requires number of threads greater than zero");
+
+ // Record whether the user actually requested a thread limit (thread_limit
+ // clause) before possibly overwriting UserThreadLimit[0] below with the
+ // computed effective value.
+ const bool ThreadLimitFromUser = LaunchArgs.UserThreadLimit[0] > 0;
+
+ // Calculate or adjust the effective number of threads and blocks for the
+ // first dimension, if the caller didn't request strict counts.
+ if (!StrictThreads || !StrictBlocks) {
+ assert(!KernelEnv.isBareMode() &&
+ "bare kernel launches must request strict thread/block counts");
+
+ GenericDeviceTy &GenericDevice = RTL->getDevice(RTLDeviceID);
+ uint32_t EffectiveNumThreads = LaunchArgs.UserThreadLimit[0];
+ if (!StrictThreads)
+ EffectiveNumThreads =
+ getEffectiveNumThreads(GenericDevice, EffectiveNumThreads, KernelEnv);
+
+ if (!StrictBlocks)
+ LaunchArgs.UserNumBlocks[0] = getEffectiveNumBlocks(
+ GenericDevice, LaunchArgs.UserNumBlocks[0], LaunchArgs.Tripcount,
+ EffectiveNumThreads, StrictThreads, ThreadLimitFromUser, KernelEnv);
+
+ LaunchArgs.UserThreadLimit[0] = EffectiveNumThreads;
+ }
if (KernelArgs.Flags.IsCUDA) {
// Kernel languages (CUDA/HIP) pass an already-flattened argument-pointer
@@ -443,6 +608,15 @@ int32_t DeviceTy::launchKernel(void *TgtEntryPtr, void **TgtVarsPtr,
}
}
+ auto *Kernel = reinterpret_cast<GenericKernelTy *>(TgtEntryPtr);
+ INFO(OMP_INFOTYPE_PLUGIN_KERNEL, RTL->getDevice(RTLDeviceID).getDeviceId(),
+ "Launching kernel %s with [%u,%u,%u] blocks and [%u,%u,%u] threads in "
+ "%s mode\n",
+ Kernel->getName(), LaunchArgs.UserNumBlocks[0],
+ LaunchArgs.UserNumBlocks[1], LaunchArgs.UserNumBlocks[2],
+ LaunchArgs.UserThreadLimit[0], LaunchArgs.UserThreadLimit[1],
+ LaunchArgs.UserThreadLimit[2], KernelEnv.getExecutionModeName());
+
return RTL->launch_kernel(RTLDeviceID, TgtEntryPtr, LaunchArgs, AsyncInfo);
}
diff --git a/offload/plugins-nextgen/common/include/PluginInterface.h b/offload/plugins-nextgen/common/include/PluginInterface.h
index 2247a47d3ee96..a350f23905cac 100644
--- a/offload/plugins-nextgen/common/include/PluginInterface.h
+++ b/offload/plugins-nextgen/common/include/PluginInterface.h
@@ -425,46 +425,6 @@ class DeviceImageTy {
}
};
-struct KernelLaunchInfoTy {
- uint32_t MaxNumThreads = 0;
- uint32_t PreferredNumThreads = 0;
- uint32_t ReductionDataSize = 0;
- /// Defaults to OMP_TGT_EXEC_MODE_BARE.
- OMPTgtExecModeFlags Mode = OMP_TGT_EXEC_MODE_BARE;
-
- /// Indicate if the kernel works in Bare, Generic SPMD, Generic, No-Loop
- /// or SPMD mode.
- bool isBareMode() const { return Mode == OMP_TGT_EXEC_MODE_BARE; }
- bool isGenericMode() const { return Mode == OMP_TGT_EXEC_MODE_GENERIC; }
- bool isGenericSPMDMode() const {
- return Mode == OMP_TGT_EXEC_MODE_GENERIC_SPMD;
- }
- bool isSPMDMode() const { return Mode == OMP_TGT_EXEC_MODE_SPMD; }
- bool isNoLoopMode() const { return Mode == OMP_TGT_EXEC_MODE_SPMD_NO_LOOP; }
-
- static const char *getExecutionModeName(OMPTgtExecModeFlags Mode) {
- switch (Mode) {
- case OMP_TGT_EXEC_MODE_BARE:
- return "BARE";
- case OMP_TGT_EXEC_MODE_SPMD:
- return "SPMD";
- case OMP_TGT_EXEC_MODE_GENERIC:
- return "Generic";
- case OMP_TGT_EXEC_MODE_GENERIC_SPMD:
- return "Generic-SPMD";
- case OMP_TGT_EXEC_MODE_SPMD_NO_LOOP:
- return "SPMD-No-Loop";
- }
- return "Unknown";
- }
-
- /// Return the display name of this kernel's execution mode, for
- /// debug/info logging only.
- const char *getExecutionModeName() const {
- return getExecutionModeName(Mode);
- }
-};
-
/// The subset of KernelArgsTy fields the plugin interface needs to launch a
/// kernel, plus the resolved argument-pointer array. Unlike KernelArgsTy,
/// this struct is populated by libomptarget on the stack for every launch,
@@ -495,15 +455,16 @@ struct KernelLaunchArgsTy {
uint32_t UserNumBlocks[3] = {0, 0, 0};
/// User-requested number of threads (for x,y,z dimension).
uint32_t UserThreadLimit[3] = {0, 0, 0};
- KernelLaunchInfoTy KernelEnvironment;
+ struct {
+ /// Size in bytes of a single cross-team reduction buffer element for
+ /// this kernel, or 0 if the kernel does not need a reduction buffer.
+ uint32_t ReductionDataSize = 0;
+ } KernelEnvironment;
struct {
uint64_t Cooperative : 1; // Was this kernel spawned as cooperative.
- uint64_t StrictBlocks : 1; // The user-requested number of blocks is strict.
- uint64_t StrictThreads
- : 1; // The user-requested number of threads is strict.
uint64_t DynCGroupMemFallback : 2; // The fallback for dynamic cgroup mem.
- uint64_t Unused : 60;
- } Flags = {0, 0, 0, 0, 0};
+ uint64_t Unused : 61;
+ } Flags = {0, 0, 0};
/// Set by the caller when replaying a previously recorded kernel launch, so
/// the plugin can report the outcome back; null for a normal launch.
KernelReplayOutcomeTy *ReplayOutcome = nullptr;
@@ -610,23 +571,6 @@ struct GenericKernelTy {
const KernelLaunchArgsTy &LaunchArgs,
uint32_t NumBlocks) const;
- /// Get the effective number of threads for the kernel based on the
- /// user-defined number of threads.
- static uint32_t getEffectiveNumThreads(GenericDeviceTy &GenericDevice,
- uint32_t UserThreadLimit,
- const KernelLaunchArgsTy &LaunchArgs);
-
- /// 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.
- static uint32_t
- getEffectiveNumBlocks(GenericDeviceTy &GenericDevice, uint32_t UserNumBlocks,
- uint64_t LoopTripCount, uint32_t &EffectiveNumThreads,
- bool IsNumThreadsStrict, bool IsNumThreadsFromUser,
- const KernelLaunchArgsTy &LaunchArgs);
-
/// The kernel name.
std::string Name;
diff --git a/offload/plugins-nextgen/common/src/PluginInterface.cpp b/offload/plugins-nextgen/common/src/PluginInterface.cpp
index c7103a701c531..cb8dcdc857cdc 100644
--- a/offload/plugins-nextgen/common/src/PluginInterface.cpp
+++ b/offload/plugins-nextgen/common/src/PluginInterface.cpp
@@ -29,13 +29,11 @@
#include "llvm/Bitcode/BitcodeReader.h"
#include "llvm/Frontend/OpenMP/OMPConstants.h"
#include "llvm/Support/Error.h"
-#include "llvm/Support/MathExtras.h"
#include "llvm/Support/MemoryBuffer.h"
#include "llvm/Support/Signals.h"
#include "llvm/Support/raw_ostream.h"
#include <cstdint>
-#include <limits>
using namespace llvm;
using namespace omp;
@@ -153,12 +151,6 @@ Error GenericKernelTy::printLaunchInfo(GenericDeviceTy &GenericDevice,
const KernelLaunchArgsTy &LaunchArgs,
uint32_t NumThreads[3],
uint32_t NumBlocks[3]) const {
- INFO(OMP_INFOTYPE_PLUGIN_KERNEL, GenericDevice.getDeviceId(),
- "Launching kernel %s with [%u,%u,%u] blocks and [%u,%u,%u] threads in "
- "%s mode\n",
- getName(), NumBlocks[0], NumBlocks[1], NumBlocks[2], NumThreads[0],
- NumThreads[1], NumThreads[2],
- LaunchArgs.KernelEnvironment.getExecutionModeName());
return printLaunchInfoDetails(GenericDevice, LaunchArgs, NumThreads,
NumBlocks);
}
@@ -226,33 +218,6 @@ Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice,
LaunchArgs.UserNumBlocks[1],
LaunchArgs.UserNumBlocks[2]};
- // Multidimensional is only supported with bare mode for now.
- assert(LaunchArgs.KernelEnvironment.isBareMode() ||
- EffectiveNumThreads[1] == 1 && EffectiveNumThreads[2] == 1 &&
- EffectiveNumBlocks[1] == 1 && EffectiveNumBlocks[2] == 1 &&
- "Non-bare mode should only use the first thread and block "
- "dimensions");
-
- assert(!LaunchArgs.Flags.StrictBlocks ||
- EffectiveNumBlocks[0] > 0 && EffectiveNumBlocks[1] > 0 &&
- EffectiveNumBlocks[2] > 0 &&
- "Strict requires number of blocks greater than zero");
- assert(!LaunchArgs.Flags.StrictThreads ||
- EffectiveNumThreads[0] > 0 && EffectiveNumThreads[1] > 0 &&
- EffectiveNumThreads[2] > 0 &&
- "Strict requires number of threads greater than zero");
-
- // Calculate or adjust the effective number of threads and blocks if needed.
- if (!LaunchArgs.Flags.StrictThreads)
- EffectiveNumThreads[0] = getEffectiveNumThreads(
- GenericDevice, EffectiveNumThreads[0], LaunchArgs);
-
- if (!LaunchArgs.Flags.StrictBlocks)
- EffectiveNumBlocks[0] = getEffectiveNumBlocks(
- GenericDevice, EffectiveNumBlocks[0], LaunchArgs.Tripcount,
- EffectiveNumThreads[0], LaunchArgs.Flags.StrictThreads,
- LaunchArgs.UserThreadLimit[0] > 0, LaunchArgs);
-
auto DynBlockMemConfOrErr = prepareBlockMemory(
GenericDevice, LaunchArgs,
EffectiveNumBlocks[0] * EffectiveNumBlocks[1] * EffectiveNumBlocks[2]);
@@ -312,118 +277,6 @@ Error GenericKernelTy::launch(GenericDeviceTy &GenericDevice,
return Plugin::success();
}
-uint32_t
-GenericKernelTy::getEffectiveNumThreads(GenericDeviceTy &GenericDevice,
- uint32_t UserThreadLimit,
- const KernelLaunchArgsTy &LaunchArgs) {
- assert(!LaunchArgs.KernelEnvironment.isBareMode() &&
- "bare kernel should not call this function");
-
- if (UserThreadLimit > 0 && LaunchArgs.KernelEnvironment.isGenericMode())
- UserThreadLimit += GenericDevice.getWarpSize();
-
- return std::min(LaunchArgs.KernelEnvironment.MaxNumThreads,
- (UserThreadLimit > 0)
- ? UserThreadLimit
- : LaunchArgs.KernelEnvironment.PreferredNumThreads);
-}
-
-uint32_t GenericKernelTy::getEffectiveNumBlocks(
- GenericDeviceTy &GenericDevice, uint32_t UserNumBlocks,
- uint64_t LoopTripCount, uint32_t &EffectiveNumThreads,
- bool IsNumThreadsStrict, bool IsNumThreadsFromUser,
- const KernelLaunchArgsTy &LaunchArgs) {
- assert(!LaunchArgs.KernelEnvironment.isBareMode() &&
- "bare kernel should not call this function");
-
- // NOTE: This clamps the user-requested number of blocks to the device limit
- // rather than honoring it exactly, which is non-standard behavior. Truly
- // honoring an arbitrary value would require launching multiple kernels or
- // reusing blocks until the requested count has been served.
- if (UserNumBlocks > 0)
- return std::min(UserNumBlocks,
- GenericDevice.getBlockLimit(EffectiveNumThreads));
-
- // Return the number of blocks required to cover the loop iterations.
- if (LaunchArgs.KernelEnvironment.isNoLoopMode())
- return LoopTripCount > 0 ? (((LoopTripCount - 1) / EffectiveNumThreads) + 1)
- : 1;
-
- uint64_t DefaultNumBlocks = GenericDevice.getDefaultNumBlocks();
- uint64_t TripCountNumBlocks = std::numeric_limits<uint64_t>::max();
- if (LoopTripCount > 0) {
- if (LaunchArgs.KernelEnvironment.isSPMDMode()) {
- // We have a combined construct, i.e. `target teams distribute
- // 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 blocks, we artificially
- // reduce the thread count per block to increase the outer parallelism.
- auto MinThreads = GenericDevice.getMinThreadsForLowTripCountLoop();
- MinThreads = std::min(MinThreads, EffectiveNumThreads);
-
- // Honor the thread_limit clause; only lower the number of threads.
- [[maybe_unused]] auto OldNumThreads = EffectiveNumThreads;
- if (LoopTripCount >= DefaultNumBlocks * EffectiveNumThreads ||
- IsNumThreadsFromUser || IsNumThreadsStrict) {
- // Enough parallelism for blocks and threads.
- TripCountNumBlocks = ((LoopTripCount - 1) / EffectiveNumThreads) + 1;
- assert(IsNumThreadsFromUser ||
- TripCountNumBlocks >= DefaultNumBlocks &&
- "Expected sufficient outer parallelism.");
- } else if (LoopTripCount >= DefaultNumBlocks * MinThreads) {
- // Enough parallelism for blocks, limit threads.
-
- // This case is hard; for now, we force "full warps":
- // First, compute a thread count assuming DefaultNumBlocks.
- auto NumThreadsDefaultBlocks =
- (LoopTripCount + DefaultNumBlocks - 1) / DefaultNumBlocks;
- // Now get a power of two that is larger or equal.
- auto NumThreadsDefaultBlocksP2 =
- llvm::PowerOf2Ceil(NumThreadsDefaultBlocks);
- // Do not increase a thread limit given be the user.
- EffectiveNumThreads =
- std::min(EffectiveNumThreads, uint32_t(NumThreadsDefaultBlocksP2));
- assert(EffectiveNumThreads >= MinThreads &&
- "Expected sufficient inner parallelism.");
- TripCountNumBlocks = ((LoopTripCount - 1) / EffectiveNumThreads) + 1;
- } else {
- // Not enough parallelism for blocks and threads, limit both.
- EffectiveNumThreads = std::min(EffectiveNumThreads, MinThreads);
- TripCountNumBlocks = ((LoopTripCount - 1) / EffectiveNumThreads) + 1;
- }
-
- assert(EffectiveNumThreads * TripCountNumBlocks >= LoopTripCount &&
- "Expected sufficient parallelism");
- assert(OldNumThreads >= EffectiveNumThreads &&
- "Number of threads cannot be increased!");
- } else {
- assert((LaunchArgs.KernelEnvironment.isGenericMode() ||
- LaunchArgs.KernelEnvironment.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 block is
- // assigned one iteration of the `distribute` loop. E.g.:
- //
- // #pragma omp target teams distribute
- // for(...loop_tripcount...) {
- // #pragma omp parallel for
- // for(...) {}
- // }
- //
- // Threads within a block will execute the iterations of the `parallel`
- // loop.
- TripCountNumBlocks = LoopTripCount;
- }
- }
-
- uint32_t PreferredNumBlocks = TripCountNumBlocks;
- // If the loops are long running we rather reuse blocks than spawn too many.
- if (GenericDevice.getReuseBlocksForHighTripCount())
- PreferredNumBlocks = std::min(TripCountNumBlocks, DefaultNumBlocks);
- return std::min(PreferredNumBlocks,
- GenericDevice.getBlockLimit(EffectiveNumThreads));
-}
-
GenericDeviceTy::GenericDeviceTy(GenericPluginTy &Plugin, int32_t DeviceId,
int32_t NumDevices,
const llvm::omp::GV &OMPGridValues)
diff --git a/offload/test/offloading/ompx_bare.c b/offload/test/offloading/ompx_bare.c
index 35e335e062a2a..6c690ae086a13 100644
--- a/offload/test/offloading/ompx_bare.c
+++ b/offload/test/offloading/ompx_bare.c
@@ -15,7 +15,7 @@ int main(int argc, char *argv[]) {
const int N = num_blocks * block_size;
int *data = (int *)malloc(N * sizeof(int));
- // CHECK: PluginInterface device 0 info: Launching kernel
+ // CHECK: omptarget device 0 info: Launching kernel
// CHECK-SAME: __omp_offloading_{{.*}} with [64,1,1] blocks and [64,1,1]
// CHECK-SAME: threads in BARE mode
diff --git a/offload/test/offloading/ompx_bare_gridsize.c b/offload/test/offloading/ompx_bare_gridsize.c
index 5f295d785bf52..9081102da1e56 100644
--- a/offload/test/offloading/ompx_bare_gridsize.c
+++ b/offload/test/offloading/ompx_bare_gridsize.c
@@ -23,7 +23,7 @@ void get_gridsizes(int *nblocks, int *nthreads) {
int main(int argc, char *argv[]) {
int nblocks[3], nthreads[3];
- // CHECK: PluginInterface device 0 info: Launching kernel
+ // CHECK: omptarget device 0 info: Launching kernel
// CHECK-SAME: __omp_offloading_{{.*}} with [64,1,1] blocks and [32,1,1]
// CHECK-SAME: threads in BARE mode
nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
@@ -37,7 +37,7 @@ int main(int argc, char *argv[]) {
fprintf(stderr, "nblocks: %d %d %d, nthreads: %d %d %d\n", nblocks[0],
nblocks[1], nblocks[2], nthreads[0], nthreads[1], nthreads[2]);
- // CHECK: PluginInterface device 0 info: Launching kernel
+ // CHECK: omptarget device 0 info: Launching kernel
// CHECK-SAME: __omp_offloading_{{.*}} with [64,1,1] blocks and [32,4,1]
// CHECK-SAME: threads in BARE mode
nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
@@ -51,7 +51,7 @@ int main(int argc, char *argv[]) {
fprintf(stderr, "nblocks: %d %d %d, nthreads: %d %d %d\n", nblocks[0],
nblocks[1], nblocks[2], nthreads[0], nthreads[1], nthreads[2]);
- // CHECK: PluginInterface device 0 info: Launching kernel
+ // CHECK: omptarget device 0 info: Launching kernel
// CHECK-SAME: __omp_offloading_{{.*}} with [64,1,1] blocks and [32,4,2]
// CHECK-SAME: threads in BARE mode
nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
@@ -65,7 +65,7 @@ int main(int argc, char *argv[]) {
fprintf(stderr, "nblocks: %d %d %d, nthreads: %d %d %d\n", nblocks[0],
nblocks[1], nblocks[2], nthreads[0], nthreads[1], nthreads[2]);
- // CHECK: PluginInterface device 0 info: Launching kernel
+ // CHECK: omptarget device 0 info: Launching kernel
// CHECK-SAME: __omp_offloading_{{.*}} with [64,16,1] blocks and [32,1,1]
// CHECK-SAME: threads in BARE mode
nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
@@ -79,7 +79,7 @@ int main(int argc, char *argv[]) {
fprintf(stderr, "nblocks: %d %d %d, nthreads: %d %d %d\n", nblocks[0],
nblocks[1], nblocks[2], nthreads[0], nthreads[1], nthreads[2]);
- // CHECK: PluginInterface device 0 info: Launching kernel
+ // CHECK: omptarget device 0 info: Launching kernel
// CHECK-SAME: __omp_offloading_{{.*}} with [64,16,1] blocks and [32,4,1]
// CHECK-SAME: threads in BARE mode
nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
@@ -93,7 +93,7 @@ int main(int argc, char *argv[]) {
fprintf(stderr, "nblocks: %d %d %d, nthreads: %d %d %d\n", nblocks[0],
nblocks[1], nblocks[2], nthreads[0], nthreads[1], nthreads[2]);
- // CHECK: PluginInterface device 0 info: Launching kernel
+ // CHECK: omptarget device 0 info: Launching kernel
// CHECK-SAME: __omp_offloading_{{.*}} with [64,16,1] blocks and [32,4,2]
// CHECK-SAME: threads in BARE mode
nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
@@ -107,7 +107,7 @@ int main(int argc, char *argv[]) {
fprintf(stderr, "nblocks: %d %d %d, nthreads: %d %d %d\n", nblocks[0],
nblocks[1], nblocks[2], nthreads[0], nthreads[1], nthreads[2]);
- // CHECK: PluginInterface device 0 info: Launching kernel
+ // CHECK: omptarget device 0 info: Launching kernel
// CHECK-SAME: __omp_offloading_{{.*}} with [64,16,8] blocks and [32,1,1]
// CHECK-SAME: threads in BARE mode
nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
@@ -121,7 +121,7 @@ int main(int argc, char *argv[]) {
fprintf(stderr, "nblocks: %d %d %d, nthreads: %d %d %d\n", nblocks[0],
nblocks[1], nblocks[2], nthreads[0], nthreads[1], nthreads[2]);
- // CHECK: PluginInterface device 0 info: Launching kernel
+ // CHECK: omptarget device 0 info: Launching kernel
// CHECK-SAME: __omp_offloading_{{.*}} with [64,16,8] blocks and [32,4,1]
// CHECK-SAME: threads in BARE mode
nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
@@ -135,7 +135,7 @@ int main(int argc, char *argv[]) {
fprintf(stderr, "nblocks: %d %d %d, nthreads: %d %d %d\n", nblocks[0],
nblocks[1], nblocks[2], nthreads[0], nthreads[1], nthreads[2]);
- // CHECK: PluginInterface device 0 info: Launching kernel
+ // CHECK: omptarget device 0 info: Launching kernel
// CHECK-SAME: __omp_offloading_{{.*}} with [64,16,8] blocks and [32,4,2]
// CHECK-SAME: threads in BARE mode
nblocks[0] = nblocks[1] = nblocks[2] = nthreads[0] = nthreads[1] =
diff --git a/offload/test/offloading/ompx_bare_multi_dim.cpp b/offload/test/offloading/ompx_bare_multi_dim.cpp
index 93a7b3604e411..6d89f9c27faaa 100644
--- a/offload/test/offloading/ompx_bare_multi_dim.cpp
+++ b/offload/test/offloading/ompx_bare_multi_dim.cpp
@@ -10,7 +10,7 @@
#include <cassert>
#include <vector>
-// CHECK: PluginInterface device 0 info: Launching kernel
+// CHECK: omptarget device 0 info: Launching kernel
// CHECK-SAME: __omp_offloading_{{.*}} with [2,4,6] blocks and [32,4,2] threads
// CHECK-SAME: in BARE mode
More information about the llvm-branch-commits
mailing list