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" <[email protected]> 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" <[email protected]> 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. _______________________________________________ llvm-branch-commits mailing list [email protected] https://lists.llvm.org/cgi-bin/mailman/listinfo/llvm-branch-commits
