diff --git a/offload/include/device.h b/offload/include/device.h index 96679b1656e778..1d0d35202df0f0 100644 --- a/offload/include/device.h +++ b/offload/include/device.h @@ -42,8 +42,8 @@ using GenericPluginTy = llvm::omp::target::plugin::GenericPluginTy; using DeviceInfo = llvm::omp::target::plugin::DeviceInfo; using InfoTreeNode = llvm::omp::target::plugin::InfoTreeNode; // Downstream, the kernel launch-geometry properties live in the plugin -// interface: the AMDGPU plugin computes the launch geometry itself and needs -// them, including the AMD-only execution modes. +// interface: the AMDGPU plugin reads them, including the AMD-only execution +// modes, e.g., for its kernel trace. using KernelLaunchInfoTy = llvm::omp::target::plugin::KernelLaunchInfoTy; // Forward declarations. @@ -63,6 +63,10 @@ struct DeviceGeometryPropsTy { uint32_t DefaultNumBlocks = 0; /// Maximum amount of shared memory, in bytes, a block can use in a kernel. uint64_t MaxBlockSharedMemSize = 0; + /// Number of compute units, or zero if the device does not report it. + uint32_t NumComputeUnits = 0; + /// Backend of the platform of the device. + ol_platform_backend_t Backend = OL_PLATFORM_BACKEND_UNKNOWN; }; struct PluginManager; diff --git a/offload/libompaccsupport/device.cpp b/offload/libompaccsupport/device.cpp index 1a1ef26f64ba0a..28bc9e23889c57 100644 --- a/offload/libompaccsupport/device.cpp +++ b/offload/libompaccsupport/device.cpp @@ -127,6 +127,18 @@ llvm::Error DeviceTy::initGeometryProps() { if (auto Err = Query(DeviceInfo::WORK_GROUP_LOCAL_MEM_SIZE, GeometryProps.MaxBlockSharedMemSize)) return Err; + // Not every device reports its number of compute units; zero turns off the + // heuristics that need it. + GeometryProps.NumComputeUnits = + getInfo(DeviceInfo::NUM_COMPUTE_UNITS); + // Some defaults depend on the backend. If it is unknown, the generic ones + // apply. + if (auto Platform = getInfo(DeviceInfo::PLATFORM)) { + if (olGetPlatformInfo(Platform, OL_PLATFORM_INFO_BACKEND, + sizeof(GeometryProps.Backend), + &GeometryProps.Backend)) + GeometryProps.Backend = OL_PLATFORM_BACKEND_UNKNOWN; + } return llvm::Error::success(); } @@ -846,19 +858,96 @@ static uint32_t getEffectiveNumThreads(const DeviceGeometryPropsTy &Props, : KernelEnv.PreferredNumThreads); } +/// Get the number of teams the user set with OMP_NUM_TEAMS, or zero. +static uint32_t getOMPNumTeams() { + static Int32Envar OMP_NumTeams("OMP_NUM_TEAMS"); + return OMP_NumTeams > 0 ? uint32_t(OMP_NumTeams) : 0; +} + /// Get the maximum number of blocks the device can launch for a kernel using /// \p NumThreads threads per block, further limited by OMP_NUM_TEAMS if the /// user set it. static uint32_t getDeviceBlockLimit(GenericDeviceTy &GenericDevice, uint32_t NumThreads) { - static Int32Envar OMP_NumTeams("OMP_NUM_TEAMS"); - uint32_t BlockLimit = GenericDevice.getBlockLimit(NumThreads); - if (OMP_NumTeams > 0) - return std::min(BlockLimit, uint32_t(OMP_NumTeams)); + if (uint32_t NumTeams = getOMPNumTeams()) + return std::min(BlockLimit, NumTeams); return BlockLimit; } +/// Threshold for the reduction data size (in bytes) above which a cross-team +/// reduction classifies as "large". +static constexpr uint32_t LargeTeamsReductionBytes = 32; + +/// Downstream: maximum number of blocks of a big-jump-loop kernel. +static constexpr uint64_t BigJumpLoopMaxNumBlocks = 1024 * 1024; + +/// Get the number of blocks (each having \p NumThreads threads) to launch a +/// kernel with a cross-team reduction, or zero to handle it like a +/// non-reduction kernel. +/// +/// Every block adds to the cost of combining the partial results of the blocks. +/// For a "large" reduction (see \p LargeTeamsReductionBytes), this cost +/// dominates -> one block / compute unit. A small/scalar reduction prefers more +/// blocks -> as many as needed given the device's number of waves per compute +/// unit, but at least one / compute unit. +static uint64_t +getXTeamReductionNumBlocks(const DeviceGeometryPropsTy &Props, + uint32_t NumThreads, uint64_t LoopTripCount, + const KernelLaunchInfoTy &KernelEnv) { + if (!KernelEnv.ReductionDataSize || !NumThreads) + return 0; + + // In Generic and Generic-SPMD mode, each block runs one iteration of the + // distribute loop. For these modes, there currently is only performance data + // for large reductions. + bool IsLarge = KernelEnv.ReductionDataSize > LargeTeamsReductionBytes; + bool IsSPMD = KernelEnv.isSPMDMode(); + if (!IsSPMD && !(IsLarge && (KernelEnv.isGenericMode() || + KernelEnv.isGenericSPMDMode()))) + return 0; + + // Envar to override the device's number of waves per compute unit for a + // kernel with a scalar cross-team reduction. Zero keeps the device's value. + static UInt32Envar OMPX_TeamsReductionWavesPerCU( + "LIBOMPTARGET_TEAMS_REDUCTION_WAVES_PER_CU", 0); + uint32_t ScalarWavesPerCU = OMPX_TeamsReductionWavesPerCU.get(); + // Measured on MI210, MI300A, MI300X, MI350X, and MI450: scalar reductions + // with many iterations prefer more waves, ones with fewer iterations or more + // memory traffic prefer less waves. 16 waves were never far from the best. + if (!ScalarWavesPerCU && Props.Backend == OL_PLATFORM_BACKEND_AMDGPU) + ScalarWavesPerCU = 16; + uint64_t NumCUs = Props.NumComputeUnits; + if (!ScalarWavesPerCU || !NumCUs) + return 0; + + uint32_t WavesPerBlock = llvm::divideCeil(NumThreads, Props.WarpSize); + uint64_t BlocksPerCU = + IsLarge ? 1 : std::max(1u, ScalarWavesPerCU / WavesPerBlock); + uint64_t NumBlocks = NumCUs * BlocksPerCU; + // No more blocks than there is work for. + if (LoopTripCount > 0) + NumBlocks = + std::min(NumBlocks, IsSPMD ? llvm::divideCeil(LoopTripCount, NumThreads) + : LoopTripCount); + return NumBlocks; +} + +/// Get the default number of blocks for a kernel in Generic-SPMD mode. +static uint32_t +getDefaultNumBlocksGenericSPMD(const DeviceGeometryPropsTy &Props) { + // Envar to set the default number of blocks per compute unit for a kernel in + // Generic-SPMD mode on AMDGPU devices. Zero uses the default number of blocks + // of other kernels. + static UInt32Envar OMPX_GenericSpmdTeamsPerCU( + "LIBOMPTARGET_AMDGPU_GENERIC_SPMD_TEAMS_PER_CU", 6); + + if (Props.Backend != OL_PLATFORM_BACKEND_AMDGPU || + !OMPX_GenericSpmdTeamsPerCU.get() || !Props.NumComputeUnits) + return Props.DefaultNumBlocks; + return Props.NumComputeUnits * OMPX_GenericSpmdTeamsPerCU.get(); +} + /// 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 @@ -881,11 +970,28 @@ static uint32_t getEffectiveNumBlocks( getDeviceBlockLimit(GenericDevice, EffectiveNumThreads)); // Return the number of blocks required to cover the loop iterations. - if (KernelEnv.isNoLoopMode()) - return LoopTripCount > 0 ? (((LoopTripCount - 1) / EffectiveNumThreads) + 1) - : 1; + // Downstream: a big-jump-loop kernel has no per-block overhead, so it also + // gets one thread per iteration. It loops over the grid, so the number of + // blocks may be capped. Only AMDGPU devices run big-jump-loop kernels. + if (KernelEnv.isNoLoopMode() || KernelEnv.isBigJumpLoopMode()) { + uint64_t NumBlocks = + LoopTripCount > 0 ? llvm::divideCeil(LoopTripCount, EffectiveNumThreads) + : 1; + if (KernelEnv.isBigJumpLoopMode()) + NumBlocks = std::min( + {NumBlocks, BigJumpLoopMaxNumBlocks, + uint64_t(getDeviceBlockLimit(GenericDevice, EffectiveNumThreads))}); + return NumBlocks; + } + + if (uint64_t NumBlocks = getXTeamReductionNumBlocks( + Props, EffectiveNumThreads, LoopTripCount, KernelEnv)) + return std::min(NumBlocks, uint64_t(getDeviceBlockLimit( + GenericDevice, EffectiveNumThreads))); - uint64_t DefaultNumBlocks = Props.DefaultNumBlocks; + uint64_t DefaultNumBlocks = KernelEnv.isGenericSPMDMode() + ? getDefaultNumBlocksGenericSPMD(Props) + : Props.DefaultNumBlocks; uint64_t TripCountNumBlocks = std::numeric_limits::max(); if (LoopTripCount > 0) { if (KernelEnv.isSPMDMode()) { @@ -961,6 +1067,22 @@ static uint32_t getEffectiveNumBlocks( "LIBOMPTARGET_REUSE_BLOCKS_FOR_HIGH_TRIP_COUNT", true); if (OMPX_ReuseBlocksForHighTripCount.get()) PreferredNumBlocks = std::min(TripCountNumBlocks, DefaultNumBlocks); + + // In Generic-SPMD mode, the threads of a block run the iterations of the + // inner parallel loops, whose trip counts are unknown. If they are lower than + // the number of threads, most threads idle. So trade threads for blocks, down + // to one warp per block, as long as the distribute loop has an iteration for + // every block. This keeps the device as busy if the inner trip counts are + // high. Neither is done if the user set the number of threads or teams. + if (KernelEnv.isGenericSPMDMode() && !IsNumThreadsFromUser && + !IsNumThreadsStrict && !getOMPNumTeams()) { + while (uint64_t(PreferredNumBlocks) * 2 <= LoopTripCount && + EffectiveNumThreads > Props.WarpSize) { + EffectiveNumThreads /= 2; + PreferredNumBlocks *= 2; + } + } + return std::min(PreferredNumBlocks, getDeviceBlockLimit(GenericDevice, EffectiveNumThreads)); } @@ -1027,7 +1149,6 @@ static void checkLaunchInvariants(const KernelLaunchArgsTy &LaunchArgs, /// Calculate or adjust, in place, the effective number of threads and blocks /// for the first dimension, unless the caller requested strict counts. static void adjustEffectiveGeometry(GenericDeviceTy &GenericDevice, - const GenericKernelTy &Kernel, const DeviceGeometryPropsTy &Props, KernelLaunchArgsTy &LaunchArgs, const KernelArgsTy &KernelArgs, @@ -1046,26 +1167,16 @@ static void adjustEffectiveGeometry(GenericDeviceTy &GenericDevice, const bool ThreadLimitFromUser = LaunchArgs.UserThreadLimit[0] > 0; uint32_t EffectiveNumThreads = LaunchArgs.UserThreadLimit[0]; - uint32_t EffectiveNumBlocks = LaunchArgs.UserNumBlocks[0]; - - // Downstream: the plugin may compute the launch geometry itself (the AMDGPU - // plugin does, using occupancy and other device-specific heuristics). - // Otherwise, fall back to the generic computation. - if (!Kernel.computeLaunchGeometry(GenericDevice, LaunchArgs, StrictThreads, - StrictBlocks, EffectiveNumThreads, - EffectiveNumBlocks)) { - if (!StrictThreads) - EffectiveNumThreads = - getEffectiveNumThreads(Props, EffectiveNumThreads, KernelEnv); - - if (!StrictBlocks) - EffectiveNumBlocks = getEffectiveNumBlocks( - GenericDevice, Props, EffectiveNumBlocks, LaunchArgs.Tripcount, - EffectiveNumThreads, StrictThreads, ThreadLimitFromUser, KernelEnv); - } + if (!StrictThreads) + EffectiveNumThreads = + getEffectiveNumThreads(Props, EffectiveNumThreads, KernelEnv); + + if (!StrictBlocks) + LaunchArgs.UserNumBlocks[0] = getEffectiveNumBlocks( + GenericDevice, Props, LaunchArgs.UserNumBlocks[0], LaunchArgs.Tripcount, + EffectiveNumThreads, StrictThreads, ThreadLimitFromUser, KernelEnv); LaunchArgs.UserThreadLimit[0] = EffectiveNumThreads; - LaunchArgs.UserNumBlocks[0] = EffectiveNumBlocks; } /// Flatten the kernel arguments into \p LaunchArgs.Args. Returns the address @@ -1176,8 +1287,8 @@ int32_t DeviceTy::launchKernel(void *TgtEntryPtr, void **TgtVarsPtr, } checkLaunchInvariants(LaunchArgs, KernelArgs, KernelLaunchInfo); - adjustEffectiveGeometry(GenericDevice, *Kernel, GeometryProps, LaunchArgs, - KernelArgs, KernelLaunchInfo); + adjustEffectiveGeometry(GenericDevice, GeometryProps, LaunchArgs, KernelArgs, + KernelLaunchInfo); void **DynPtrSlot = resolveArgsAndDynPtrSlot( KernelArgs, TgtVarsPtr, TgtOffsets, Args, Ptrs, ArgSizes, LaunchArgs); diff --git a/offload/plugins-nextgen/amdgpu/src/rtl.cpp b/offload/plugins-nextgen/amdgpu/src/rtl.cpp index 2192dccd4147fd..6b13c7f8cdd877 100644 --- a/offload/plugins-nextgen/amdgpu/src/rtl.cpp +++ b/offload/plugins-nextgen/amdgpu/src/rtl.cpp @@ -767,38 +767,12 @@ struct AMDGPUDeviceImageTy : public DeviceImageTy { uint16_t ELFABIVersion; }; -/// OpenMP environment variables limiting the launch geometry. Upstream, only -/// libomptarget's generic launch-geometry computation applies them. -/// Downstream, the AMDGPU launch-geometry heuristics below replace that -/// computation and apply them themselves. -static int32_t getOMPNumTeams() { - static Int32Envar OMP_NumTeams("OMP_NUM_TEAMS"); - return OMP_NumTeams; -} -static int32_t getOMPTeamsThreadLimit() { - static Int32Envar OMP_TeamsThreadLimit("OMP_TEAMS_THREAD_LIMIT"); - return OMP_TeamsThreadLimit; -} - -/// Whether to reuse blocks rather than spawn many for long-running loops. -/// Upstream, only libomptarget's generic launch-geometry computation reads it. -static bool getReuseBlocksForHighTripCount() { - static BoolEnvar OMPX_ReuseBlocksForHighTripCount( - "LIBOMPTARGET_REUSE_BLOCKS_FOR_HIGH_TRIP_COUNT", true); - return OMPX_ReuseBlocksForHighTripCount; -} - /// Class implementing the AMDGPU kernel functionalities which derives from the /// generic kernel class. struct AMDGPUKernelTy : public GenericKernelTy { /// Create an AMDGPU kernel with a name and an execution mode. AMDGPUKernelTy(StringRef Name, GenericGlobalHandlerTy &Handler) - : GenericKernelTy(Name), - OMPX_SPMDOccupancyBasedOpt("OMPX_SPMD_OCCUPANCY_BASED_OPT", false), - OMPX_GenericSPMDOccupancyBasedOpt( - "OMPX_GENERIC_SPMD_OCCUPANCY_BASED_OPT", false), - OMPX_BigJumpLoopOccupancyBasedOpt( - "OMPX_BIGJUMPLOOP_OCCUPANCY_BASED_OPT", false) {} + : GenericKernelTy(Name) {} /// Initialize the AMDGPU kernel. Error initImpl(GenericDeviceTy &Device, DeviceImageTy &Image) override { @@ -909,34 +883,6 @@ struct AMDGPUKernelTy : public GenericKernelTy { uint32_t getPrivateSize() const { return PrivateSize; } uint16_t getConstWGSize() const { return ConstWGSize; } - /// Compute the launch geometry with the AMDGPU-specific heuristics. - bool computeLaunchGeometry(GenericDeviceTy &GenericDevice, - const KernelLaunchArgsTy &LaunchArgs, - bool StrictThreads, bool StrictBlocks, - uint32_t &NumThreads, - uint32_t &NumBlocks) const override { - // The occupancy-based heuristics need the kernel's max occupancy. - computeMaxOccupancy(GenericDevice); - - if (!StrictThreads) { - NumThreads = - getEffectiveNumThreads(GenericDevice, NumThreads, LaunchArgs); - - std::pair AdjustInfo = adjustNumThreadsForLowTripCount( - GenericDevice, NumThreads, LaunchArgs.Tripcount, - LaunchArgs.UserThreadLimit, LaunchArgs); - if (AdjustInfo.first) - NumThreads = AdjustInfo.second; - } - - if (!StrictBlocks) - NumBlocks = getEffectiveNumBlocks( - GenericDevice, NumBlocks, LaunchArgs.Tripcount, NumThreads, - StrictThreads, LaunchArgs.UserThreadLimit[0] > 0, LaunchArgs); - - return true; - } - /// Get the HSA kernel object representing the kernel function. uint64_t getKernelObject() const { return KernelObject; } @@ -950,15 +896,6 @@ struct AMDGPUKernelTy : public GenericKernelTy { void setKernelLaunchId(uint32_t Id) const { KernelLaunchId = Id; } - /// Envar to enable occupancy-based optimization for SPMD kernel. - BoolEnvar OMPX_SPMDOccupancyBasedOpt; - - /// Envar to enable occupancy-based optimization for generic SPMD kernel. - BoolEnvar OMPX_GenericSPMDOccupancyBasedOpt; - - /// Envar to enable occupancy-based optimization for big jump loop. - BoolEnvar OMPX_BigJumpLoopOccupancyBasedOpt; - private: /// The kernel object to execute. uint64_t KernelObject; @@ -983,461 +920,6 @@ struct AMDGPUKernelTy : public GenericKernelTy { static thread_local uint32_t KernelLaunchId; - /// Lower number of threads if tripcount is low. This should produce - /// a larger number of teams if allowed by other constraints. - std::pair - adjustNumThreadsForLowTripCount(GenericDeviceTy &GenericDevice, - uint32_t BlockSize, uint64_t LoopTripCount, - const uint32_t ThreadLimitClause[3], - const KernelLaunchArgsTy &LaunchArgs) const { - const KernelLaunchInfoTy &Info = LaunchArgs.KernelLaunchInfo; - uint32_t NumThreads = BlockSize; - - // If there is an override already, do nothing. Note the different - // default for cross-team reductions. - const bool IsTeamsReduction = - Info.isSPMDMode() && Info.doesTeamsReduction(); - - if (!IsTeamsReduction && - NumThreads != GenericDevice.getDefaultNumThreads() && - NumThreads != ConstWGSize) - return std::make_pair(false, NumThreads); - - if (IsTeamsReduction && - NumThreads != llvm::omp::xteam_red::DefaultBlockSize && - NumThreads != ConstWGSize) - return std::make_pair(false, NumThreads); - - // If tripcount not set or not low, do nothing. - if ((LoopTripCount == 0) || - (LoopTripCount > GenericDevice.getOMPXLowTripCount())) - return std::make_pair(false, NumThreads); - - // Environment variable present, do nothing. - if (getOMPTeamsThreadLimit() > 0) - return std::make_pair(false, NumThreads); - - // num_threads clause present, do nothing. - if ((ThreadLimitClause[0] > 0) && (ThreadLimitClause[0] != (uint32_t)-1)) - return std::make_pair(false, NumThreads); - - // If generic or generic-SPMD kernel, do nothing. - if (Info.isGenericMode() || Info.isGenericSPMDMode()) - return std::make_pair(false, NumThreads); - - // Reduce the blocksize as long as it is above the tunable limit. - while (NumThreads > GenericDevice.getOMPXSmallBlockSize()) - NumThreads >>= 1; - - if (NumThreads == 0) - return std::make_pair(false, BlockSize); - - // The repeated halving above can leave an awkward block size when the - // CodeGen-provided one is not a power of two. Round the value the plugin - // picked itself; a block size the user asked for is honored verbatim, see - // getEffectiveNumThreads(). - if (IsTeamsReduction) - return std::make_pair(true, - llvm::omp::getBlockSizeAsPowerOfTwo(NumThreads)); - - return std::make_pair(true, NumThreads); - } - - /// Optimize the number of teams based on the max occupancy value. - uint64_t OptimizeNumTeamsBaseOccupancy(GenericDeviceTy &GenericDevice, - uint32_t NumThreads) const { - unsigned NumWavesPerTeam = - divideCeil(NumThreads, GenericDevice.getWarpSize()); - unsigned TotalWavesPerCU = MaxOccupancy * amdgpu_arch::SIMDPerCU; - // Per device - unsigned TotalWavesPerDevice = - TotalWavesPerCU * GenericDevice.getNumComputeUnits(); - unsigned NumTeams = divideCeil(TotalWavesPerDevice, NumWavesPerTeam); - - return static_cast(NumTeams); - } - - /// Get the number of threads and blocks for the kernel based on the - /// user-defined threads and block clauses. - uint32_t getEffectiveNumThreads(GenericDeviceTy &GenericDevice, - uint32_t UserThreadLimit, - const KernelLaunchArgsTy &LaunchArgs) const { - const KernelLaunchInfoTy &Info = LaunchArgs.KernelLaunchInfo; - assert(!Info.isBareMode() && "bare kernel should not call this function"); - - // Honor OMP_TEAMS_THREAD_LIMIT environment variable and - // num_threads/thread_limit clause for BigJumpLoop and NoLoop kernel types. - int32_t TeamsThreadLimitEnvVar = getOMPTeamsThreadLimit(); - if (Info.isBigJumpLoopMode() || Info.isNoLoopMode()) { - if (TeamsThreadLimitEnvVar > 0) - return std::min(static_cast(ConstWGSize), - TeamsThreadLimitEnvVar); - if ((UserThreadLimit > 0) && (UserThreadLimit != (uint32_t)-1)) - return std::min(static_cast(ConstWGSize), UserThreadLimit); - return ConstWGSize; - } - - // Cross-team reduction kernels default to the block size CodeGen picked - // for them, which is larger than the generic SPMD default. A user-provided - // value is honored as given, as long as it fits: the upstream cross-team - // reduction works with any block size, unlike the removed Xteamr helpers - // that required a power of two. - if (Info.isSPMDMode() && Info.doesTeamsReduction()) { - if (TeamsThreadLimitEnvVar > 0 && - TeamsThreadLimitEnvVar <= static_cast(ConstWGSize)) - return TeamsThreadLimitEnvVar; - if (UserThreadLimit > 0 && UserThreadLimit != (uint32_t)-1 && - UserThreadLimit <= static_cast(ConstWGSize)) - return UserThreadLimit; - uint32_t BlockSizeOverride = GenericDevice.getOMPXXteamBlockSize(); - if (BlockSizeOverride > 0 && - BlockSizeOverride <= static_cast(ConstWGSize)) - return BlockSizeOverride; - return ConstWGSize; - } - - if (UserThreadLimit > 0 && Info.isGenericMode()) { - if (UserThreadLimit == (uint32_t)-1) - UserThreadLimit = Info.PreferredNumThreads; - else - UserThreadLimit += GenericDevice.getWarpSize(); - } - - // Limit number of threads taking into consideration the user - // environment variable OMP_TEAMS_THREAD_LIMIT if provided. - uint32_t CurrentMaxNumThreads = Info.MaxNumThreads; - if (TeamsThreadLimitEnvVar > 0) - CurrentMaxNumThreads = std::min( - static_cast(TeamsThreadLimitEnvVar), CurrentMaxNumThreads); - - return std::min(CurrentMaxNumThreads, (UserThreadLimit > 0) - ? UserThreadLimit - : Info.PreferredNumThreads); - } - uint32_t getEffectiveNumBlocks(GenericDeviceTy &GenericDevice, - uint32_t UserNumBlocks, uint64_t LoopTripCount, - uint32_t &EffectiveNumThreads, - bool IsNumThreadsStrict, - bool IsNumThreadsFromUser, - const KernelLaunchArgsTy &LaunchArgs) const { - const KernelLaunchInfoTy &Info = LaunchArgs.KernelLaunchInfo; - assert(!Info.isBareMode() && "bare kernel should not call this function"); - - const auto getNumGroupsFromThreadsAndTripCount = - [](const uint64_t TripCount, const uint32_t NumThreads) { - return ((TripCount - 1) / NumThreads) + 1; - }; - uint64_t DeviceNumCUs = GenericDevice.getNumComputeUnits(); // FIXME - - if (Info.isNoLoopMode()) { - return LoopTripCount > 0 ? getNumGroupsFromThreadsAndTripCount( - LoopTripCount, EffectiveNumThreads) - : 1; - } - - uint64_t NumWavesInGroup = - (EffectiveNumThreads - 1) / GenericDevice.getWarpSize() + 1; - - if (Info.isBigJumpLoopMode()) { - int32_t NumTeamsEnvVar = getOMPNumTeams(); - uint64_t NumGroups = 1; - // Cannot assert a non-zero tripcount. Instead, launch with 1 team if the - // tripcount is indeed zero. - if (LoopTripCount > 0) - NumGroups = getNumGroupsFromThreadsAndTripCount(LoopTripCount, - EffectiveNumThreads); - - // Honor num_teams clause but lower it if tripcount dictates. - if (UserNumBlocks > 0 && - UserNumBlocks <= GenericDevice.getBlockLimit(EffectiveNumThreads)) { - NumGroups = std::min(static_cast(UserNumBlocks), NumGroups); - } - // Without a num_teams clause, honor OMP_NUM_TEAMS environment variable - // for BigJumpLoop kernel type. - else if (UserNumBlocks == 0 && NumTeamsEnvVar > 0 && - static_cast(NumTeamsEnvVar) <= - GenericDevice.getBlockLimit(EffectiveNumThreads)) { - NumGroups = std::min(static_cast(NumTeamsEnvVar), NumGroups); - } else { - // num_teams clause is not specified. Choose lower of tripcount-based - // NumGroups and a value determined as follows: - // - If the number of teams per CU is specified by the user with the - // envar LIBOMPTARGET_AMDGPU_BIG_JUMP_LOOP_TEAMS_PER_CU, compute - // NumGroups from that specified value. This envar is OFF by default. - // - Otherwise, use the max total teams specified with the envar - /// LIBOMPTARGET_AMDGPU_BIG_JUMP_LOOP_MAX_TOTAL_TEAMS. - // This envar is used by default with 1M as the default value. - if (GenericDevice.getOMPXBigJumpLoopTeamsPerCU() > 0) { - NumGroups = - std::min(NumGroups, GenericDevice.getOMPXBigJumpLoopTeamsPerCU() * - DeviceNumCUs); - } else { - NumGroups = std::min( - NumGroups, static_cast( - GenericDevice.getOMPXBigJumpLoopMaxTotalTeams())); - } - - // If the user specifies a number of teams for low trip count loops, - // honor it. - uint64_t LowTripCountBlocks = - GenericDevice.getOMPXNumBlocksForLowTripcount(LoopTripCount); - if (LowTripCountBlocks) { - NumGroups = LowTripCountBlocks; - } - } - // If envar OMPX_BIGJUMPLOOP_OCCUPANCY_BASED_OPT is set and no num_teams - // clause or OMP_NUM_TEAMS is specified, optimize the number of teams - // based on occupancy value. - if (OMPX_BigJumpLoopOccupancyBasedOpt && NumTeamsEnvVar == 0 && - UserNumBlocks == 0) { - return std::min(NumGroups, OptimizeNumTeamsBaseOccupancy( - GenericDevice, EffectiveNumThreads)); - } - return std::min(NumGroups, - static_cast( - GenericDevice.getBlockLimit(EffectiveNumThreads))); - } - - // A cross-team reduction kernel is now emitted by CodeGen as a plain SPMD - // kernel using the upstream reduction path (the downstream Xteam reduction - // execution mode has been removed). Recognize it via the reduction data - // size so the AMDGPU reduction grid-size heuristic (which upstream has no - // equivalent for) still applies. - if (Info.isSPMDMode() && Info.doesTeamsReduction()) { - // Here's the default number of teams. - uint64_t NumGroups = DeviceNumCUs; - // The number of teams must not exceed this upper limit. - uint64_t MaxNumGroups = NumGroups; - // Honor OMP_NUM_TEAMS environment variable for XteamReduction kernel - // type, if possible. - int32_t NumTeamsEnvVar = getOMPNumTeams(); - // CU mulitiplier from envar. - uint32_t EnvarCUMultiplier = GenericDevice.getXTeamRedTeamsPerCU(); - - if (GenericDevice.isFastReductionEnabled()) { - // When fast reduction is enabled, the number of teams is capped by - // the MaxCUMultiplier constant. - MaxNumGroups = DeviceNumCUs * llvm::omp::xteam_red::MaxCUMultiplier; - } else { - // When fast reduction is not enabled, the number of teams is capped - // by the metadata that clang CodeGen created. The number of teams - // used here must not exceed the upper limit determined during - // CodeGen. This upper limit is not currently communicated from - // CodeGen to the plugin. So it is re-computed here. - - // ConstWGSize is the block size that CodeGen used. - uint32_t CUMultiplier = - llvm::omp::xteam_red::getXteamRedCUMultiplier(ConstWGSize); - MaxNumGroups = DeviceNumCUs * CUMultiplier; - } - - // If envar OMPX_XTEAMREDUCTION_OCCUPANCY_BASED_OPT is set and no - // OMP_NUM_TEAMS or num_teams clause is specified, optimize the num of - // teams based on occupancy value. We apply this optimization only when - // the MaxOccupancy equals or exceeds the desirable waves per CU. The - // assumption is that anything lower is probably resource constrained - // already and this optimization may not be beneficial. - if (GenericDevice.getOMPXXTeamReductionOccupancyBasedOpt() && - NumTeamsEnvVar == 0 && UserNumBlocks == 0 && - (MaxOccupancy * amdgpu_arch::SIMDPerCU >= - llvm::omp::xteam_red::DesiredWavesPerCU)) { - NumGroups = std::min( - OptimizeNumTeamsBaseOccupancy(GenericDevice, EffectiveNumThreads), - MaxNumGroups); - // Saturating the device is only worth it if there are enough - // iterations to go around: a low trip count would otherwise launch - // teams that have nothing to execute and only add to the reduction. - // The non-occupancy path below applies the same bound. - if (LoopTripCount > 0) - NumGroups = - std::min(NumGroups, getNumGroupsFromThreadsAndTripCount( - LoopTripCount, EffectiveNumThreads)); - ODBG(ODT_Tool) << "xteam-red:NumCUs=" << DeviceNumCUs - << " xteam-red:NumGroups=" << NumGroups; - return NumGroups; - } - - // Prefer num_teams clause over environment variable. There is a corner - // case where inspite of the presence of a num_teams clause, CodeGen - // may fail to extract it, instead using the alternative computation of - // the number of teams. But the runtime here will still see the value - // of the clause, so we need to check against the upper limit. - if (UserNumBlocks > 0 && - UserNumBlocks <= GenericDevice.getBlockLimit(EffectiveNumThreads)) { - NumGroups = - std::min(static_cast(UserNumBlocks), MaxNumGroups); - } else if (UserNumBlocks == 0 && NumTeamsEnvVar > 0 && - static_cast(NumTeamsEnvVar) <= - GenericDevice.getBlockLimit(EffectiveNumThreads)) { - NumGroups = - std::min(static_cast(NumTeamsEnvVar), MaxNumGroups); - } else { - // Ensure we don't have a large number of teams running if the tripcount - // is low - uint64_t NumGroupsFromTripCount = 1; - if (LoopTripCount > 0) - NumGroupsFromTripCount = getNumGroupsFromThreadsAndTripCount( - LoopTripCount, EffectiveNumThreads); - - // Compute desired number of groups in the absence of user input - // based on a factor controlled by an integer env-var. - // Note that the upper bound is MaxNumGroups. - uint32_t AdjustFactor = - GenericDevice.getOMPXAdjustNumTeamsForXteamRedSmallBlockSize(); - if (EffectiveNumThreads > 0 && AdjustFactor > 0) { - uint64_t DesiredNumGroups = NumGroups; - if (AdjustFactor == 1) { - DesiredNumGroups = - DeviceNumCUs * - (llvm::omp::xteam_red::DesiredWavesPerCU / NumWavesInGroup); - } else { - DesiredNumGroups = DeviceNumCUs * AdjustFactor; - } - NumGroups = DesiredNumGroups; - } - - // Prefer OMPX_AdjustNumTeamsForXteamRedSmallBlockSize over - // OMPX_XTeamRedTeamsPerCU. - if (AdjustFactor == 0 && EnvarCUMultiplier > 0) - NumGroups = DeviceNumCUs * EnvarCUMultiplier; - - NumGroups = std::min(NumGroups, MaxNumGroups); - NumGroups = std::min(NumGroups, NumGroupsFromTripCount); - - // If the user specifies a number of teams for low trip count loops, - // and no num_teams clause was used, honor it. - uint64_t LowTripCountBlocks = - GenericDevice.getOMPXNumBlocksForLowTripcount(LoopTripCount); - if (LowTripCountBlocks) { - NumGroups = std::min(MaxNumGroups, LowTripCountBlocks); - } - } - ODBG(ODT_Tool) << "xteam-red:NumCUs=" << DeviceNumCUs - << " xteam-red:NumGroups=" << NumGroups; - return NumGroups; - } - - if (UserNumBlocks > 0) { - // TODO: We need to honor any value and consequently allow more than the - // block limit. For this we might need to start multiple kernels or let - // the blocks start again until the requested number has been started. - return std::min(UserNumBlocks, - GenericDevice.getBlockLimit(EffectiveNumThreads)); - } - - // If envar OMPX_SPMD_OCCUPANCY_BASED_OPT is set and no OMP_NUM_TEAMS is - // specified, optimize the num of teams based on occupancy value. - int32_t NumTeamsEnvVar = getOMPNumTeams(); - uint64_t TripCountNumBlocks = std::numeric_limits::max(); - if (LoopTripCount > 0) { - if (Info.isSPMDMode()) { - // We have a combined construct, i.e. `target teams distribute - // parallel for [simd]`. We launch so many teams so that each thread - // will execute one iteration of the loop. round up to the nearest - // integer - TripCountNumBlocks = ((LoopTripCount - 1) / EffectiveNumThreads) + 1; - } else { - assert((Info.isGenericMode() || Info.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 team 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 team will execute the iterations of the `parallel` - // loop. - TripCountNumBlocks = LoopTripCount; - } - } - - if (Info.isSPMDMode() && OMPX_SPMDOccupancyBasedOpt && - NumTeamsEnvVar == 0 && UserNumBlocks == 0) { - return std::min( - TripCountNumBlocks, - OptimizeNumTeamsBaseOccupancy(GenericDevice, EffectiveNumThreads)); - } - - auto getAdjustedDefaultNumBlocks = - [&Info](GenericDeviceTy &GenericDevice, - uint64_t DeviceNumCUs) -> uint64_t { - if (!Info.isGenericSPMDMode() || - GenericDevice.getOMPXGenericSpmdTeamsPerCU() == 0) - return static_cast(GenericDevice.getDefaultNumBlocks()); - return DeviceNumCUs * static_cast( - GenericDevice.getOMPXGenericSpmdTeamsPerCU()); - }; - - // If the loops are long running we rather reuse blocks than spawn too many. - // Additionally, under an env-var, adjust the number of teams based on the - // number of wave-slots in a CU that we aim to occupy. - uint64_t AdjustedNumBlocks = - getAdjustedDefaultNumBlocks(GenericDevice, DeviceNumCUs); - if (GenericDevice.getOMPXAdjustNumTeamsForSmallBlockSize()) { - uint64_t DefaultNumWavesInGroup = - (GenericDevice.getDefaultNumThreads() - 1) / - GenericDevice.getWarpSize() + - 1; - AdjustedNumBlocks = - (AdjustedNumBlocks * DefaultNumWavesInGroup) / NumWavesInGroup; - } - - // If the user specifies a number of teams for low trip count loops, honor - // it. - uint64_t LowTripCountBlocks = - GenericDevice.getOMPXNumBlocksForLowTripcount(LoopTripCount); - if (LowTripCountBlocks) { - return LowTripCountBlocks; - } - - uint64_t PreferredNumBlocks = TripCountNumBlocks; - // Occupancy-based setting overrides block reuse. - if (OMPX_GenericSPMDOccupancyBasedOpt && NumTeamsEnvVar == 0 && - UserNumBlocks == 0) { - PreferredNumBlocks = std::min( - PreferredNumBlocks, - OptimizeNumTeamsBaseOccupancy(GenericDevice, EffectiveNumThreads)); - } else if (getReuseBlocksForHighTripCount()) { - // If the loops are long running we rather reuse blocks than spawn too - // many. - PreferredNumBlocks = std::min(TripCountNumBlocks, AdjustedNumBlocks); - } - - // For most generic-SPMD kernels, the tripcount of the outer distribute-loop - // determines the number of teams launched. The tripcounts of the inner - // parallel loops should determine the number of threads launched. However, - // the inner loop tripcounts are unknown, so the runtime just launches 256 - // threads by default. But if the inner loop tripcount is lower than 256, - // many of the threads in every workgroup are idle and just waste resources. - // In order to reduce this wastage, we reduce the blocksize upto the - // wavefront size if the tripcount is large enough to proportionally - // increase the number of teams. The increase in the number of teams is - // required to preserve the occupancy in case the inner loop tripcounts are - // larger than the blocksize. This change is done only when the user has not - // specified the number of teams or threads. - if (Info.isGenericSPMDMode() && !IsNumThreadsFromUser && - !IsNumThreadsStrict && UserNumBlocks == 0 && NumTeamsEnvVar == 0 && - GenericDevice.getOMPXGenericSpmdUseSmallBlockSize()) { - uint64_t TmpPreferredNumBlocks = PreferredNumBlocks << 1; - while (TmpPreferredNumBlocks <= LoopTripCount && - EffectiveNumThreads > GenericDevice.getWarpSize()) { - EffectiveNumThreads >>= 1; - PreferredNumBlocks = TmpPreferredNumBlocks; - TmpPreferredNumBlocks <<= 1; - } - } - // Without a num_teams clause, OMP_NUM_TEAMS is an upper bound on the - // number of teams. - uint64_t BlockLimit = GenericDevice.getBlockLimit(EffectiveNumThreads); - if (NumTeamsEnvVar > 0) - BlockLimit = std::min(BlockLimit, static_cast(NumTeamsEnvVar)); - return std::min(PreferredNumBlocks, BlockLimit); - } - /// Compute the occupancy with the constraint on the number of SGPRs /// Follow the logic on the backend /// Ref: @@ -3318,29 +2800,6 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { OMPX_NumQueues("LIBOMPTARGET_AMDGPU_NUM_HSA_QUEUES", 4), OMPX_QueueSize("LIBOMPTARGET_AMDGPU_HSA_QUEUE_SIZE", 512), OMPX_DefaultTeamsPerCU("LIBOMPTARGET_AMDGPU_TEAMS_PER_CU", 4), - OMPX_GenericSpmdTeamsPerCU( - "LIBOMPTARGET_AMDGPU_GENERIC_SPMD_TEAMS_PER_CU", 6), - OMPX_BigJumpLoopTeamsPerCU( - "LIBOMPTARGET_AMDGPU_BIG_JUMP_LOOP_TEAMS_PER_CU", 0), - OMPX_XTeamRedTeamsPerCU("LIBOMPTARGET_AMDGPU_XTEAM_RED_TEAMS_PER_CU", - 0), - OMPX_BigJumpLoopMaxTotalTeams( - "LIBOMPTARGET_AMDGPU_BIG_JUMP_LOOP_MAX_TOTAL_TEAMS", 1024 * 1024), - OMPX_LowTripCount("LIBOMPTARGET_AMDGPU_LOW_TRIPCOUNT", 9000), - OMPX_SmallBlockSize("LIBOMPTARGET_MIN_THREADS_FOR_LOW_TRIP_COUNT", 32), - OMPX_NumBlocksForLowTripcount("LIBOMPTARGET_BLOCKS_FOR_LOW_TRIP_COUNT", - 0), - OMPX_WavesPerCUForLowTripcount( - "LIBOMPTARGET_WAVES_PER_CU_FOR_LOW_TRIP_COUNT", 0), - OMPX_AdjustNumTeamsForSmallBlockSize("LIBOMPTARGET_AMDGPU_ADJUST_TEAMS", - 0), - OMPX_XTeamReductionOccupancyBasedOpt( - "OMPX_XTEAMREDUCTION_OCCUPANCY_BASED_OPT", false), - OMPX_AdjustNumTeamsForXteamRedSmallBlockSize( - "LIBOMPTARGET_AMDGPU_ADJUST_XTEAM_RED_TEAMS", 1), - OMPX_GenericSpmdUseSmallBlockSize( - "LIBOMPTARGET_AMDGPU_GENERIC_SPMD_USE_SMALL_BLOCKSIZE", 1), - OMPX_XteamBlockSize("LIBOMPTARGET_AMDGPU_XTEAM_BLOCKSIZE", 0), OMPX_MaxAsyncCopyBytes("LIBOMPTARGET_AMDGPU_MAX_ASYNC_COPY_BYTES", 64 * 1024), OMPX_InitialNumSignals("LIBOMPTARGET_AMDGPU_NUM_INITIAL_HSA_SIGNALS", @@ -3371,29 +2830,10 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { if (!OMPX_UseMultipleSdmaEngines.isPresent()) { OMPX_UseMultipleSdmaEngines = EnvarConfig.OMPX_UseMultipleSdmaEngines; } - if (!OMPX_AdjustNumTeamsForXteamRedSmallBlockSize.isPresent()) { - OMPX_AdjustNumTeamsForXteamRedSmallBlockSize = - EnvarConfig.OMPX_AdjustNumTeamsForXteamRedSmallBlockSize; - } - if (!OMPX_XteamBlockSize.isPresent()) { - OMPX_XteamBlockSize = - EnvarConfig.OMPX_XteamBlockSize; - } - if (!OMPX_XTeamReductionOccupancyBasedOpt.isPresent()) { - OMPX_XTeamReductionOccupancyBasedOpt = - EnvarConfig.OMPX_XTeamReductionOccupancyBasedOpt; - } // Print potential GPU envars. - ODBG(ODT_Tool) - << "Loaded per GPU envars:" - << " OMPX_UseMultipleSdmaEngines=%d\n" - << EnvarConfig.OMPX_UseMultipleSdmaEngines - << " OMPX_AdjustNumTeamsForXteamRedSmallBlockSize=" - << EnvarConfig.OMPX_AdjustNumTeamsForXteamRedSmallBlockSize - << " OMPX_XteamBlockSize=" - << EnvarConfig.OMPX_XteamBlockSize - << " OMPX_XTeamReductionOccupancyBasedOpt=" - << EnvarConfig.OMPX_XTeamReductionOccupancyBasedOpt; + ODBG(ODT_Tool) << "Loaded per GPU envars:" + << " OMPX_UseMultipleSdmaEngines=" + << EnvarConfig.OMPX_UseMultipleSdmaEngines; } ~AMDGPUDeviceTy() {} @@ -3418,98 +2858,6 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { return OMPX_NumQueues; } - virtual uint32_t getOMPXGenericSpmdTeamsPerCU() const override { - return OMPX_GenericSpmdTeamsPerCU; - } - virtual uint32_t getOMPXBigJumpLoopTeamsPerCU() const override { - return OMPX_BigJumpLoopTeamsPerCU; - } - virtual uint32_t getXTeamRedTeamsPerCU() const override { - return OMPX_XTeamRedTeamsPerCU; - } - virtual uint32_t getOMPXBigJumpLoopMaxTotalTeams() const override { - return OMPX_BigJumpLoopMaxTotalTeams; - } - virtual uint32_t getOMPXLowTripCount() const override { - return OMPX_LowTripCount; - } - virtual uint32_t getOMPXSmallBlockSize() const override { - return OMPX_SmallBlockSize; - } - virtual uint32_t - getOMPXNumBlocksForLowTripcount(uint64_t LoopTripCount) const override { - uint32_t NumBlocks = 0; - - if (LoopTripCount > OMPX_LowTripCount) - return NumBlocks; - - // if NumBlocksForLowTripcount is set, it has the highest priority. - if (OMPX_NumBlocksForLowTripcount > 0) { - NumBlocks = OMPX_NumBlocksForLowTripcount; - ODBG(ODT_Tool) << "Small trip count loop: Using " - << NumBlocks - << " blocks"; - } - - // Next, check if the waves per CU is set. This will launch a number of - // blocks such that we only have at most OMPX_WavesPerCUForLowTripcount - // waves per CU. - if (OMPX_WavesPerCUForLowTripcount > 0) { - // Compute the number of waves per block. For sizes smaller than a full - // wave the size is 1. - uint32_t WavesPerBlock = (uint32_t)((OMPX_SmallBlockSize - 1) / 64) + 1; - ODBG(ODT_Tool) << "Small trip count loop: Using " - << WavesPerBlock - << " waves per block"; - - // We cannot return less than the number of CUs: - if (WavesPerBlock >= OMPX_WavesPerCUForLowTripcount) { - NumBlocks = NumComputeUnits; - ODBG(ODT_Tool) << "Small trip count loop: Using 1 block per CU"; - } else { - uint32_t BlocksPerCU = - (uint32_t)(OMPX_WavesPerCUForLowTripcount / WavesPerBlock); - ODBG(ODT_Tool) << "Small trip count loop: Using " - << BlocksPerCU - << " blocks per CU"; - NumBlocks = (uint32_t)(BlocksPerCU * NumComputeUnits); - } - } - - // Adjust the number of blocks to the trip count if number of blocks x - // threads is much larger than the loop trip count. - if (NumBlocks) { - if (LoopTripCount <= OMPX_SmallBlockSize) - NumBlocks = 1; - - uint32_t MaxBlocks = - (uint32_t)((LoopTripCount - 1) / OMPX_SmallBlockSize) + 1; - if (NumBlocks > MaxBlocks) { - NumBlocks = MaxBlocks; - ODBG(ODT_Tool) << "Small trip count loop: number of blocks capped to " - << NumBlocks - << "trip count"; - } - } - return NumBlocks; - } - virtual uint32_t getOMPXAdjustNumTeamsForSmallBlockSize() const override { - return OMPX_AdjustNumTeamsForSmallBlockSize; - } - virtual uint32_t - getOMPXAdjustNumTeamsForXteamRedSmallBlockSize() const override { - return OMPX_AdjustNumTeamsForXteamRedSmallBlockSize; - } - virtual bool getOMPXXTeamReductionOccupancyBasedOpt() const override { - return OMPX_XTeamReductionOccupancyBasedOpt; - } - virtual bool getOMPXGenericSpmdUseSmallBlockSize() const override { - return OMPX_GenericSpmdUseSmallBlockSize; - } - virtual uint32_t getOMPXXteamBlockSize() const override { - return OMPX_XteamBlockSize; - } - uint64_t getDeviceTimeStamp() override { return getSystemTimestampInNs(); } /// Initialize the device, its resources and get its properties. @@ -5317,80 +4665,6 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { /// #default_teams = OMPX_DefaultTeamsPerCU * #CUs. UInt32Envar OMPX_DefaultTeamsPerCU; - /// Envar for controlling the number of teams relative to the number of - /// compute units (CUs) for generic-SPMD kernels. 0 indicates that this value - /// is not specified, so instead OMPX_DefaultTeamsPerCU should be used. If - /// non-zero, the number of teams = OMPX_GenericSpmdTeamsPerCU * #CUs. - UInt32Envar OMPX_GenericSpmdTeamsPerCU; - - /// Envar for controlling the number of teams relative to the number of - /// compute units (CUs) for Big-Jump-Loop kernels. 0 indicates that this value - /// is not specified. If non-zero, the number of teams = - /// OMPX_BigJumpLoopTeamsPerCU * #CUs. - UInt32Envar OMPX_BigJumpLoopTeamsPerCU; - - /// Envar for controlling the number of teams relative to the number of - /// compute units (CUs) for cross-team-reduction kernels. 0 indicates that - /// this value is not specified. If non-zero, the number of teams = - /// OMPX_XTeamRedTeamsPerCU * #CUs. - UInt32Envar OMPX_XTeamRedTeamsPerCU; - - /// Envar controlling the maximum number of teams per device for - /// Big-Jump-Loop kernels. - UInt32Envar OMPX_BigJumpLoopMaxTotalTeams; - - /// Envar specifying tripcount below which the blocksize should be adjusted. - UInt32Envar OMPX_LowTripCount; - - /// Envar specifying a value till which the blocksize can be adjusted if the - /// tripcount is low. - UInt32Envar OMPX_SmallBlockSize; - - /// Envar for the number of blocks when the loop trip count is under the small - /// trip count limit. - /// The default value of 0 means that the number of blocks will be inferred by - /// the existing getEffectiveNumBlocks logic. - UInt32Envar OMPX_NumBlocksForLowTripcount; - - /// Envar to set the number of waves per CU for small trip count loops. The - /// number of blocks will be adjusted such that there are no more than the - /// specified number of blocks per CU than this variable specifies. For - /// example: - /// Given: - // a GPU with CUs = 100 - /// and OMPX_WavesPerCUForLowTripcount = 8 - /// and a waves per block number of 4 (256 threads) - /// The total number of blocks will be: 200 - UInt32Envar OMPX_WavesPerCUForLowTripcount; - - /// Envar to allow adjusting number of teams after small tripcount - /// optimization. The default 0 means no adjustment of number of teams is - /// done. - UInt32Envar OMPX_AdjustNumTeamsForSmallBlockSize; - - /// Envar to enable occupancy-based optimization for cross team reduction. - /// Defaults come from the per-device envar config table. - BoolEnvar OMPX_XTeamReductionOccupancyBasedOpt; - - /// Envar to allow scaling up the number of teams for Xteam-Reduction, - /// whenever the blocksize has been reduced from the max. The value 0 - /// indicates that this functionality is disabled. The default value is 1, - /// indicating that if the number of waves is lower than the max, increase the - /// number of teams proportionally. A value greater than 1 indicates that the - /// value should be used as the scaling factor for the number of teams. - UInt32Envar OMPX_AdjustNumTeamsForXteamRedSmallBlockSize; - - /// Envar indicating whether, for generic-SPMD kernels, the blocksize should - /// be reduced and the corresponding number of teams adjusted. - BoolEnvar OMPX_GenericSpmdUseSmallBlockSize; - - /// Envar indicating the blocksize to be used for Xteam reduction kernels. The - /// default of 0 indicates that there is no runtime override and the value - /// indicated by CodeGen will be used. If a non-zero value is specified, the - /// runtime will attempt to use it as an override if other constraints are - /// satisfied. - UInt32Envar OMPX_XteamBlockSize; - /// Envar specifying the maximum size in bytes where the memory copies are /// asynchronous operations. Up to this transfer size, the memory copies are /// asynchronous operations pushed to the corresponding stream. For larger @@ -5570,65 +4844,19 @@ struct AMDGPUDeviceTy : public GenericDeviceTy, AMDGenericDeviceTy { struct DeviceEnvarConfigTy { bool OMPX_UseMultipleSdmaEngines; // LIBOMPTARGET_AMDGPU_USE_MULTIPLE_SDMA_ENGINES - bool - OMPX_AdjustNumTeamsForXteamRedSmallBlockSize; - int - OMPX_XteamBlockSize; - bool - OMPX_XTeamReductionOccupancyBasedOpt; }; - // OMPX_XTeamReductionOccupancyBasedOpt is disabled everywhere: the member - // shadowing that kept it from taking effect was only fixed recently, so the - // entries that requested it have never actually run. Enable it per device - // once the occupancy-based team count has been evaluated there. static inline const std::unordered_map - EnvarConfigs = {{"MI210", - {.OMPX_UseMultipleSdmaEngines = true, - .OMPX_XteamBlockSize = 512, - .OMPX_XTeamReductionOccupancyBasedOpt = false, - .OMPX_AdjustNumTeamsForXteamRedSmallBlockSize = 0}}, - {"MI250X", - {.OMPX_UseMultipleSdmaEngines = true, - .OMPX_XteamBlockSize = 512, - .OMPX_XTeamReductionOccupancyBasedOpt = false, - .OMPX_AdjustNumTeamsForXteamRedSmallBlockSize = 0}}, - {"MI250X/MI250", - {.OMPX_UseMultipleSdmaEngines = true, - .OMPX_XteamBlockSize = 512, - .OMPX_XTeamReductionOccupancyBasedOpt = false, - .OMPX_AdjustNumTeamsForXteamRedSmallBlockSize = 0}}, - {"MI300A", - {.OMPX_UseMultipleSdmaEngines = false, - .OMPX_XteamBlockSize = 512, - .OMPX_XTeamReductionOccupancyBasedOpt = false, - .OMPX_AdjustNumTeamsForXteamRedSmallBlockSize = 1}}, - {"MI300X", - {.OMPX_UseMultipleSdmaEngines = true, - .OMPX_XteamBlockSize = 512, - .OMPX_XTeamReductionOccupancyBasedOpt = false, - .OMPX_AdjustNumTeamsForXteamRedSmallBlockSize = 1}}, - {"MI308X", - {.OMPX_UseMultipleSdmaEngines = true, - .OMPX_XteamBlockSize = 256, - .OMPX_XTeamReductionOccupancyBasedOpt = false, - .OMPX_AdjustNumTeamsForXteamRedSmallBlockSize = 0}}, - {"MI350X", - {.OMPX_UseMultipleSdmaEngines = true, - .OMPX_XteamBlockSize = 512, - .OMPX_XTeamReductionOccupancyBasedOpt = false, - .OMPX_AdjustNumTeamsForXteamRedSmallBlockSize = 1}}, - {"MI355X", - {.OMPX_UseMultipleSdmaEngines = true, - .OMPX_XteamBlockSize = 512, - .OMPX_XTeamReductionOccupancyBasedOpt = false, - .OMPX_AdjustNumTeamsForXteamRedSmallBlockSize = 1}}, + EnvarConfigs = {{"MI210", {.OMPX_UseMultipleSdmaEngines = true}}, + {"MI250X", {.OMPX_UseMultipleSdmaEngines = true}}, + {"MI250X/MI250", {.OMPX_UseMultipleSdmaEngines = true}}, + {"MI300A", {.OMPX_UseMultipleSdmaEngines = false}}, + {"MI300X", {.OMPX_UseMultipleSdmaEngines = true}}, + {"MI308X", {.OMPX_UseMultipleSdmaEngines = true}}, + {"MI350X", {.OMPX_UseMultipleSdmaEngines = true}}, + {"MI355X", {.OMPX_UseMultipleSdmaEngines = true}}, // Default config for unknown devices. - {"DEFAULT", - {.OMPX_UseMultipleSdmaEngines = false, - .OMPX_XteamBlockSize = 512, - .OMPX_XTeamReductionOccupancyBasedOpt = false, - .OMPX_AdjustNumTeamsForXteamRedSmallBlockSize = 1}}}; + {"DEFAULT", {.OMPX_UseMultipleSdmaEngines = false}}}; const DeviceEnvarConfigTy &getEnvarConfig() const { std::string DeviceMarketingName = getNormMarketingName(); diff --git a/offload/plugins-nextgen/common/include/PluginInterface.h b/offload/plugins-nextgen/common/include/PluginInterface.h index c6ac570c9f153d..039b9eaff121f1 100644 --- a/offload/plugins-nextgen/common/include/PluginInterface.h +++ b/offload/plugins-nextgen/common/include/PluginInterface.h @@ -545,21 +545,6 @@ struct GenericKernelTy { return AchievedOccupancy; } - /// Downstream: let the plugin compute the effective launch geometry of a - /// non-bare OpenMP kernel launch, i.e., the number of threads and blocks of - /// the first dimension. On entry, \p NumThreads and \p NumBlocks hold the - /// user-requested values; the ones flagged by \p StrictThreads and - /// \p StrictBlocks must be kept. Returns false if the plugin does not - /// provide its own computation, in which case libomptarget uses the - /// generic one. - virtual bool computeLaunchGeometry(GenericDeviceTy &GenericDevice, - const KernelLaunchArgsTy &LaunchArgs, - bool StrictThreads, bool StrictBlocks, - uint32_t &NumThreads, - uint32_t &NumBlocks) const { - return false; - } - protected: /// Prints generic kernel launch information. Error printLaunchInfo(GenericDeviceTy &GenericDevice, @@ -1276,44 +1261,6 @@ struct GenericDeviceTy : public DeviceAllocatorTy { } virtual uint64_t getClockFrequency() const { return CLOCKS_PER_SEC; } - virtual uint32_t getOMPXGenericSpmdTeamsPerCU() const { - llvm_unreachable("Unimplemented"); - } - virtual uint32_t getOMPXBigJumpLoopTeamsPerCU() const { - llvm_unreachable("Unimplemented"); - } - virtual uint32_t getXTeamRedTeamsPerCU() const { - llvm_unreachable("Unimplemented"); - } - virtual uint32_t getOMPXBigJumpLoopMaxTotalTeams() const { - llvm_unreachable("Unimplemented"); - } - virtual uint32_t getOMPXLowTripCount() const { - llvm_unreachable("Unimplemented"); - } - virtual uint32_t getOMPXSmallBlockSize() const { - llvm_unreachable("Unimplemented"); - } - virtual uint32_t - getOMPXNumBlocksForLowTripcount(uint64_t LoopTripCount) const { - llvm_unreachable("Unimplemented"); - } - virtual uint32_t getOMPXAdjustNumTeamsForSmallBlockSize() const { - llvm_unreachable("Unimplemented"); - } - virtual uint32_t getOMPXAdjustNumTeamsForXteamRedSmallBlockSize() const { - llvm_unreachable("Unimplemented"); - } - virtual bool getOMPXXTeamReductionOccupancyBasedOpt() const { - llvm_unreachable("Unimplemented"); - } - virtual bool getOMPXGenericSpmdUseSmallBlockSize() const { - llvm_unreachable("Unimplemented"); - } - virtual uint32_t getOMPXXteamBlockSize() const { - llvm_unreachable("Unimplemented"); - } - /// Get a device-specific timestamp in nanoseconds, used by the profiler /// for timing device operations. Subclasses should override this to provide /// hardware-accurate timestamps (e.g., via HSA system info). diff --git a/offload/test/offloading/teams_reduction_grid.c b/offload/test/offloading/teams_reduction_grid.c new file mode 100644 index 00000000000000..ff322fc59433c1 --- /dev/null +++ b/offload/test/offloading/teams_reduction_grid.c @@ -0,0 +1,70 @@ +// clang-format off +// This test verifies the grid size of kernels with a cross-team reduction: +// one block per compute unit for a large reduction, the given number of waves +// per compute unit for a scalar one, no more blocks than there are iterations, +// and the num_teams clause honored. The number of compute units is taken from +// the first kernel. +// +// RUN: %libomptarget-compile-generic +// RUN: env LIBOMPTARGET_KERNEL_TRACE=1 LIBOMPTARGET_TEAMS_REDUCTION_WAVES_PER_CU=16 \ +// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic + +// UNSUPPORTED: nvptx64-nvidia-cuda +// UNSUPPORTED: nvptx64-nvidia-cuda-LTO +// UNSUPPORTED: aarch64-unknown-linux-gnu +// UNSUPPORTED: aarch64-unknown-linux-gnu-LTO +// UNSUPPORTED: x86_64-unknown-linux-gnu +// UNSUPPORTED: x86_64-unknown-linux-gnu-LTO + +// clang-format on +#include + +#define N (1 << 20) +#define M 8 + +int main() { + double Sum = 0; + double Sums[M] = {0}; + + // Large reduction (64 bytes per team): one block per compute unit. +#pragma omp target teams distribute parallel for thread_limit(256) \ + reduction(+ : Sums[0 : M]) + for (int I = 0; I < N; ++I) + Sums[I % M] += 1; + // CHECK: SGN:2 + // CHECK-SAME: teamsXthrds:({{ *}}[[#CUS:]]X 256) + + // Scalar reduction: 16 waves, that is four blocks of 256 threads, per + // compute unit. +#pragma omp target teams distribute parallel for thread_limit(256) \ + reduction(+ : Sum) + for (int I = 0; I < N; ++I) + Sum += 1; + // CHECK: SGN:2 + // CHECK-SAME: teamsXthrds:({{ *}}[[#mul(CUS,4)]]X 256) + double Sum1 = Sum; + + // No more blocks than there are iterations. + Sum = 0; +#pragma omp target teams distribute parallel for reduction(+ : Sum) + for (int I = 0; I < 16; ++I) + Sum += 1; + // CHECK: SGN:2 + // CHECK-SAME: teamsXthrds:({{ *}}1X + double Sum2 = Sum; + + // The num_teams clause is honored. + Sum = 0; +#pragma omp target teams distribute parallel for num_teams(3) \ + thread_limit(256) reduction(+ : Sum) + for (int I = 0; I < N; ++I) + Sum += 1; + // CHECK: SGN:2 + // CHECK-SAME: teamsXthrds:({{ *}}3X 256) + + // The kernel trace goes to stderr, so check the results after all of it. + // CHECK: Sums[0]=131072 Sum1=1048576 Sum2=16 Sum3=1048576 + printf("Sums[0]=%.0f Sum1=%.0f Sum2=%.0f Sum3=%.0f\n", Sums[0], Sum1, Sum2, + Sum); + return 0; +} diff --git a/offload/test/offloading/xteam_red_1.c b/offload/test/offloading/xteam_red_1.c index c201a6490be4dc..853bfb8a597452 100644 --- a/offload/test/offloading/xteam_red_1.c +++ b/offload/test/offloading/xteam_red_1.c @@ -1,15 +1,13 @@ // clang-format off -// This test verifies the AMDGPU grid heuristic for cross-team reduction -// kernels under a low trip count: the kernel is launched with 32 threads per -// team and the team count is capped at 4 x #CUs (2048 max threads per CU -// divided by the 512 block size). The expected team count below therefore -// depends on the device: 416 == 4 x 104 CUs on gfx90a. +// This test verifies the result of a fast cross-team reduction kernel under a +// low trip count. A fast reduction has no reduction data, so its grid is that +// of a plain SPMD kernel, which depends on the device. // // Cross-team reductions are emitted as plain SPMD kernels (SGN:2) since the // downstream Xteam reduction execution mode (SGN:8) was removed. // // RUN: %libomptarget-compile-generic -fopenmp-target-fast -fopenmp-target-fast-reduction -// RUN: env LIBOMPTARGET_KERNEL_TRACE=1 LIBOMPTARGET_AMDGPU_LOW_TRIPCOUNT=15360 LIBOMPTARGET_AMDGPU_ADJUST_XTEAM_RED_TEAMS=32 \ +// RUN: env LIBOMPTARGET_KERNEL_TRACE=1 \ // RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic // UNSUPPORTED: nvptx64-nvidia-cuda @@ -43,7 +41,6 @@ int main() { } // clang-format off /// CHECK: DEVID:[[S:[ ]*]][[DEVID:[0-9]+]] SGN:2 -/// CHECK-SAME: teamsXthrds:( 416X 32) /// CHECK-SAME: tripcount:15360 /// CHECK: sum1=117957120.000000 diff --git a/offload/test/offloading/xteam_red_2.c b/offload/test/offloading/xteam_red_2.c index 9427c73b991665..2f516defaf562b 100644 --- a/offload/test/offloading/xteam_red_2.c +++ b/offload/test/offloading/xteam_red_2.c @@ -1,17 +1,8 @@ // clang-format off -// This test verifies the AMDGPU grid heuristic for cross-team reduction -// kernels: the team count saturates the device with the desired number of -// waves per CU. A 512-thread team is 8 waves, so with the desired 16 waves -// per CU the heuristic picks two teams per CU. -// Which heuristic computes the team count is selected per device by the envar -// config table, so pin both knobs to keep this test on the path it is written -// for, independently of the device it runs on: the occupancy-based team count -// is a different heuristic that replaces this one, and without the adjust factor -// the team count stays at the one-team-per-CU default. +// This test verifies the result of a cross-team reduction kernel with the +// default grid size. teams_reduction_grid.c verifies the grid size. // RUN: %libomptarget-compile-generic -fopenmp-target-fast -// RUN: env LIBOMPTARGET_DEBUG=1 OMPX_XTEAMREDUCTION_OCCUPANCY_BASED_OPT=0 \ -// RUN: LIBOMPTARGET_AMDGPU_ADJUST_XTEAM_RED_TEAMS=1 \ -// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic +// RUN: %libomptarget-run-generic 2>&1 | %fcheck-generic // UNSUPPORTED: nvptx64-nvidia-cuda // UNSUPPORTED: nvptx64-nvidia-cuda-LTO @@ -43,8 +34,5 @@ int main() { return 0; } // clang-format off -/// CHECK: xteam-red:NumCUs=[[#CU_COUNT:]] -/// CHECK-SAME: xteam-red:NumGroups=[[#CU_COUNT+CU_COUNT]] -/// CHECK: Launching kernel {{.*}} with {{\[}}[[#CU_COUNT+CU_COUNT]],1,1] blocks and [512,1,1] threads in SPMD mode /// CHECK: sum1=499999500000.000000 diff --git a/openmp/docs/design/Runtimes.md b/openmp/docs/design/Runtimes.md index 6fae6a84a549a6..37136938cf6323 100644 --- a/openmp/docs/design/Runtimes.md +++ b/openmp/docs/design/Runtimes.md @@ -717,6 +717,7 @@ variables is defined below. > - `LIBOMPTARGET_JIT_SAVE_IMAGE_FILENAME= (device image file)` > - `LIBOMPTARGET_MIN_THREADS_FOR_LOW_TRIP_COUNT= (default: 32)` > - `LIBOMPTARGET_REUSE_BLOCKS_FOR_HIGH_TRIP_COUNT=[TRUE/FALSE] (default TRUE)` +> - `LIBOMPTARGET_TEAMS_REDUCTION_WAVES_PER_CU= (default: per device)` > - `OFFLOAD_TRACK_ALLOCATION_TRACES=[TRUE/FALSE] (default FALSE)` > - `OFFLOAD_TRACK_NUM_KERNEL_LAUNCH_TRACES= (default 0)` @@ -1139,6 +1140,20 @@ This environment variable can be used to control how the OpenMP runtime assigns blocks to loops with high trip counts. By default we reuse existing blocks rather than spawning new blocks. +#### LIBOMPTARGET_TEAMS_REDUCTION_WAVES_PER_CU + +This environment variable sets the number of waves per compute unit for a +kernel with a scalar cross-team reduction, e.g., `target teams distribute +parallel for reduction(...)`. The number of blocks of a kernel with a +cross-team reduction follows from the size of the reduction: if the data +reduced per team is larger than 32 bytes, one block per compute unit; +otherwise, as many blocks as fill this number of waves per compute unit, but +at least one per compute unit. No more blocks are launched than the loop has +iterations for. In Generic mode, only large reductions are sized this way. The +`num_teams` clause sets the number of blocks, and `OMP_NUM_TEAMS` caps it. The +default on AMDGPU devices is `16`. On other devices, the default is `0`: kernels +with a cross-team reduction get the same number of blocks as other kernels. + #### OFFLOAD_TRACK_ALLOCATION_TRACES This environment variable determines if the stack traces of allocations and