[llvm-branch-commits] [llvm] [offload][omp] Move strict threads & groups computation to libomptarget (PR #222607)
via llvm-branch-commits
llvm-branch-commits at lists.llvm.org
Thu Sep 10 05:34:47 PDT 2026
llvmorg-github-actions[bot] wrote:
<!--LLVM PR SUMMARY COMMENT-->
@llvm/pr-subscribers-offload
Author: Alex Duran (adurang)
<details>
<summary>Changes</summary>
Move all computation related to strict threads and strict groups to libomptarget as it's specific to OpenMP. This also removes all refrences to ExecutionModes from the Plugin Interface.
With this, the data from the OpenMP kernel environment being used is ReductionDataSize used to create the KLE, and maxNumThreads used by RecordReplay. They're are purposely left for a future PR.
Assisted by Claude.
---
Patch is 31.65 KiB, truncated to 20.00 KiB below, full version: https://github.com/llvm/llvm-project/pull/222607.diff
8 Files Affected:
- (modified) offload/include/device.h (+40-1)
- (modified) offload/liboffload/src/OffloadImpl.cpp (-2)
- (modified) offload/libompaccsupport/device.cpp (+178-3)
- (modified) offload/plugins-nextgen/common/include/PluginInterface.h (+9-63)
- (modified) offload/plugins-nextgen/common/src/PluginInterface.cpp (-147)
- (modified) offload/test/offloading/ompx_bare.c (+1-1)
- (modified) offload/test/offloading/ompx_bare_gridsize.c (+9-9)
- (modified) offload/test/offloading/ompx_bare_multi_dim.cpp (+1-1)
``````````diff
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);...
[truncated]
``````````
</details>
https://github.com/llvm/llvm-project/pull/222607
More information about the llvm-branch-commits
mailing list