https://github.com/adurang updated https://github.com/llvm/llvm-project/pull/222607
>From 821de943eeec77576d9fd5a17f40bb532e6b7e3b Mon Sep 17 00:00:00 2001 From: "Duran, Alex" <[email protected]> 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 | 181 +++++++++++++++++- .../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, 238 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..14ab45709d0d5 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,58 @@ 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; + LaunchArgs.KernelEnvironment.MaxNumThreads = KernelEnv.MaxNumThreads; + + 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 +609,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..da2b09cfbc7ea 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,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 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; + /// Maximum number of threads per block that this kernel may use. + uint32_t MaxNumThreads = 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 +573,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 _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
