[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 Oct 1 06:39:10 PDT 2026
https://github.com/adurang updated https://github.com/llvm/llvm-project/pull/222607
>From 0c7231ebe7c87c70a8a8701d12e30b9f4eb9c157 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 1/2] [offload][omp] Move strict threads & groups computation
to libomptarget
---
offload/include/device.h | 41 +++-
offload/liboffload/src/OffloadImpl.cpp | 2 -
offload/libompaccsupport/device.cpp | 183 +++++++++++++++++-
.../common/include/PluginInterface.h | 72 +------
.../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, 240 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 7f804c53142a4..a0a6cbdd03378 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 b892aaf4a93ce..c3ac08dcd58f4 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,60 @@ 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.KernelLaunchInfo = getKernelLaunchInfo(TgtEntryPtr);
+
+ KernelLaunchInfoTy KernelLaunchInfo = getKernelLaunchInfo(TgtEntryPtr);
+ LaunchArgs.KernelLaunchInfo.ReductionDataSize =
+ KernelLaunchInfo.ReductionDataSize;
+ LaunchArgs.KernelLaunchInfo.MaxNumThreads = KernelLaunchInfo.MaxNumThreads;
+
+ const bool StrictBlocks = KernelArgs.Flags.StrictBlocks;
+ const bool StrictThreads = KernelArgs.Flags.StrictThreads;
+
+ // Multidimensional is only supported with bare mode for now.
+ assert(KernelLaunchInfo.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(!KernelLaunchInfo.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, KernelLaunchInfo);
+
+ if (!StrictBlocks)
+ LaunchArgs.UserNumBlocks[0] = getEffectiveNumBlocks(
+ GenericDevice, LaunchArgs.UserNumBlocks[0], LaunchArgs.Tripcount,
+ EffectiveNumThreads, StrictThreads, ThreadLimitFromUser,
+ KernelLaunchInfo);
+
+ LaunchArgs.UserThreadLimit[0] = EffectiveNumThreads;
+ }
if (KernelArgs.Flags.IsCUDA) {
// Kernel languages (CUDA/HIP) pass an already-flattened argument-pointer
@@ -443,6 +611,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], KernelLaunchInfo.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 86efc085efa03..2909336870425 100644
--- a/offload/plugins-nextgen/common/include/PluginInterface.h
+++ b/offload/plugins-nextgen/common/include/PluginInterface.h
@@ -429,46 +429,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,
@@ -499,15 +459,18 @@ 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 KernelLaunchInfo;
+ 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;
+ /// Maximum number of threads per block that this kernel may use.
+ uint32_t MaxNumThreads = 0;
+ } KernelLaunchInfo;
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;
@@ -614,23 +577,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 0153a9daeb363..48f238541b47c 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.KernelLaunchInfo.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.KernelLaunchInfo.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.KernelLaunchInfo.isBareMode() &&
- "bare kernel should not call this function");
-
- if (UserThreadLimit > 0 && LaunchArgs.KernelLaunchInfo.isGenericMode())
- UserThreadLimit += GenericDevice.getWarpSize();
-
- return std::min(LaunchArgs.KernelLaunchInfo.MaxNumThreads,
- (UserThreadLimit > 0)
- ? UserThreadLimit
- : LaunchArgs.KernelLaunchInfo.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.KernelLaunchInfo.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.KernelLaunchInfo.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.KernelLaunchInfo.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.KernelLaunchInfo.isGenericMode() ||
- LaunchArgs.KernelLaunchInfo.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
>From 84bd6a86a9f90835f83dad98a459806241dc2aab Mon Sep 17 00:00:00 2001
From: "Duran, Alex" <alejandro.duran at intel.com>
Date: Thu, 1 Oct 2026 06:38:47 -0700
Subject: [PATCH 2/2] update fortran tests
---
.../offloading/fortran/basic-target-parallel-do.f90 | 2 +-
.../fortran/basic-target-parallel-reduction.f90 | 2 +-
.../basic-target-teams-parallel-reduction.f90 | 2 +-
.../fortran/default-mapper-nested-derived-type.f90 | 2 +-
.../fortran/do-concurrent-to-omp-min-reduce.f90 | 2 +-
.../do-concurrent-to-omp-nested-derived-type.f90 | 4 ++--
.../fortran/do-concurrent-to-omp-saxpy-2d.f90 | 2 +-
.../fortran/do-concurrent-to-omp-saxpy.f90 | 2 +-
.../explicit-and-implicit-record-field-mapping.f90 | 2 +-
.../fortran/implicit-record-field-mapping.f90 | 2 +-
.../fortran/target-custom-reduction-derivedtype.f90 | 2 +-
offload/test/offloading/fortran/target-no-loop.f90 | 12 ++++++------
.../fortran/target-parallel-do-collapse.f90 | 2 +-
.../fortran/usm_derived_type_allocatable_member.f90 | 2 +-
14 files changed, 20 insertions(+), 20 deletions(-)
diff --git a/offload/test/offloading/fortran/basic-target-parallel-do.f90 b/offload/test/offloading/fortran/basic-target-parallel-do.f90
index ee86b69a279f3..425635ea5d8b1 100644
--- a/offload/test/offloading/fortran/basic-target-parallel-do.f90
+++ b/offload/test/offloading/fortran/basic-target-parallel-do.f90
@@ -25,5 +25,5 @@ program main
end program main
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}}
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}}
! CHECK: number of errors: 0
diff --git a/offload/test/offloading/fortran/basic-target-parallel-reduction.f90 b/offload/test/offloading/fortran/basic-target-parallel-reduction.f90
index 5300eb0df49c6..ec84af3bcd32c 100644
--- a/offload/test/offloading/fortran/basic-target-parallel-reduction.f90
+++ b/offload/test/offloading/fortran/basic-target-parallel-reduction.f90
@@ -23,5 +23,5 @@ program main
end program main
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}}
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}}
! CHECK: number of errors: 0
diff --git a/offload/test/offloading/fortran/basic-target-teams-parallel-reduction.f90 b/offload/test/offloading/fortran/basic-target-teams-parallel-reduction.f90
index 496449dc714ae..37f7888ffe8c1 100644
--- a/offload/test/offloading/fortran/basic-target-teams-parallel-reduction.f90
+++ b/offload/test/offloading/fortran/basic-target-teams-parallel-reduction.f90
@@ -23,5 +23,5 @@ program main
end program main
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}}
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}}
! CHECK: number of errors: 0
diff --git a/offload/test/offloading/fortran/default-mapper-nested-derived-type.f90 b/offload/test/offloading/fortran/default-mapper-nested-derived-type.f90
index e54110f7e5d60..9ea1579fe765a 100644
--- a/offload/test/offloading/fortran/default-mapper-nested-derived-type.f90
+++ b/offload/test/offloading/fortran/default-mapper-nested-derived-type.f90
@@ -31,4 +31,4 @@ program test_omp_target_map_bug_v5
deallocate(config%derived_field%alloc_field)
end program test_omp_target_map_bug_v5
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}}
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}}
diff --git a/offload/test/offloading/fortran/do-concurrent-to-omp-min-reduce.f90 b/offload/test/offloading/fortran/do-concurrent-to-omp-min-reduce.f90
index d54f75e4714cb..b203697dac097 100644
--- a/offload/test/offloading/fortran/do-concurrent-to-omp-min-reduce.f90
+++ b/offload/test/offloading/fortran/do-concurrent-to-omp-min-reduce.f90
@@ -44,5 +44,5 @@ program main
end if
end program main
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}}
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}}
! CHECK: PASS
diff --git a/offload/test/offloading/fortran/do-concurrent-to-omp-nested-derived-type.f90 b/offload/test/offloading/fortran/do-concurrent-to-omp-nested-derived-type.f90
index e8e91fd7ef07e..b0ecfea0b8a13 100644
--- a/offload/test/offloading/fortran/do-concurrent-to-omp-nested-derived-type.f90
+++ b/offload/test/offloading/fortran/do-concurrent-to-omp-nested-derived-type.f90
@@ -54,6 +54,6 @@ program main
deallocate(aad%buf(2)%data)
end program main
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}}
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}}
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}}
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}}
! CHECK: PASS
diff --git a/offload/test/offloading/fortran/do-concurrent-to-omp-saxpy-2d.f90 b/offload/test/offloading/fortran/do-concurrent-to-omp-saxpy-2d.f90
index 2f1526c126e97..d5addcd90ccb5 100644
--- a/offload/test/offloading/fortran/do-concurrent-to-omp-saxpy-2d.f90
+++ b/offload/test/offloading/fortran/do-concurrent-to-omp-saxpy-2d.f90
@@ -47,7 +47,7 @@ program main
deallocate(x,y)
end program main
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}}
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}}
! CHECK: plausibility check:
! CHECK: y(1,1) 4.0
! CHECK: y(n,m) 4.0
diff --git a/offload/test/offloading/fortran/do-concurrent-to-omp-saxpy.f90 b/offload/test/offloading/fortran/do-concurrent-to-omp-saxpy.f90
index 91eedd3afae3c..e7d95711fee39 100644
--- a/offload/test/offloading/fortran/do-concurrent-to-omp-saxpy.f90
+++ b/offload/test/offloading/fortran/do-concurrent-to-omp-saxpy.f90
@@ -47,7 +47,7 @@ program main
deallocate(x,y)
end program main
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}}
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}}
! CHECK: plausibility check:
! CHECK: y(1) 4.0
! CHECK: y(n) 4.0
diff --git a/offload/test/offloading/fortran/explicit-and-implicit-record-field-mapping.f90 b/offload/test/offloading/fortran/explicit-and-implicit-record-field-mapping.f90
index 6ec9ffd0264bf..2f71e007a5c3d 100644
--- a/offload/test/offloading/fortran/explicit-and-implicit-record-field-mapping.f90
+++ b/offload/test/offloading/fortran/explicit-and-implicit-record-field-mapping.f90
@@ -79,5 +79,5 @@ program reproducer
print *, "======= Test Passed! ======="
end program reproducer
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}}
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}}
! CHECK: ======= Test Passed! =======
diff --git a/offload/test/offloading/fortran/implicit-record-field-mapping.f90 b/offload/test/offloading/fortran/implicit-record-field-mapping.f90
index a49c02a0da4e2..60e8f0326af73 100644
--- a/offload/test/offloading/fortran/implicit-record-field-mapping.f90
+++ b/offload/test/offloading/fortran/implicit-record-field-mapping.f90
@@ -44,5 +44,5 @@ program test_implicit_field_mapping
endif
end program
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}}
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}}
! CHECK: Test succeeded!
diff --git a/offload/test/offloading/fortran/target-custom-reduction-derivedtype.f90 b/offload/test/offloading/fortran/target-custom-reduction-derivedtype.f90
index c0aaa357eb746..3979baa28876f 100644
--- a/offload/test/offloading/fortran/target-custom-reduction-derivedtype.f90
+++ b/offload/test/offloading/fortran/target-custom-reduction-derivedtype.f90
@@ -83,6 +83,6 @@ program main
end program main
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}}
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}}
! CHECK: PASSED
diff --git a/offload/test/offloading/fortran/target-no-loop.f90 b/offload/test/offloading/fortran/target-no-loop.f90
index e3e1be5e4984a..99c02acdc2567 100644
--- a/offload/test/offloading/fortran/target-no-loop.f90
+++ b/offload/test/offloading/fortran/target-no-loop.f90
@@ -81,17 +81,17 @@ program main
end program main
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode
! CHECK: info: #Args: 3 Teams x Thrds: 64x 16
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode
! CHECK: info: #Args: 3 Teams x Thrds: 3x 16 {{.*}}
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD-No-Loop mode
! CHECK: info: #Args: 3 Teams x Thrds: 64x 16 {{.*}}
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode
! CHECK: info: #Args: 3 Teams x Thrds: 1x 16
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} Generic mode
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} Generic mode
! CHECK: info: #Args: 3 Teams x Thrds: 16x 16 {{.*}}
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}} SPMD mode
! CHECK: info: #Args: 4 Teams x Thrds: 16x 16 {{.*}}
! CHECK: number of errors: 0
diff --git a/offload/test/offloading/fortran/target-parallel-do-collapse.f90 b/offload/test/offloading/fortran/target-parallel-do-collapse.f90
index 619ba340d5f27..ce9229c40a9b4 100644
--- a/offload/test/offloading/fortran/target-parallel-do-collapse.f90
+++ b/offload/test/offloading/fortran/target-parallel-do-collapse.f90
@@ -34,6 +34,6 @@ program main
end program main
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel {{.*}}
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel {{.*}}
! CHECK: number of errors: 0
diff --git a/offload/test/offloading/fortran/usm_derived_type_allocatable_member.f90 b/offload/test/offloading/fortran/usm_derived_type_allocatable_member.f90
index 1653fc38f3a41..b3084c56bae2d 100644
--- a/offload/test/offloading/fortran/usm_derived_type_allocatable_member.f90
+++ b/offload/test/offloading/fortran/usm_derived_type_allocatable_member.f90
@@ -31,5 +31,5 @@ PROGRAM reproducer
PRINT *, GRID%DZ(1)
END PROGRAM reproducer
-! CHECK: PluginInterface device {{[0-9]+}} info: Launching kernel
+! CHECK: omptarget device {{[0-9]+}} info: Launching kernel
! CHECK: 3.
More information about the llvm-branch-commits
mailing list