Added useAcc as a template parameter to address the performance regression (#1856)

* Added useAcc as a template parameter to address the 2% performance regression in allreduceWithBias
---------

Co-authored-by: Marzieh Berenjkoub <mberenjk@amd.com>
This commit is contained in:
mberenjk
2025-08-14 15:58:54 -05:00
committato da GitHub
parent aaf8613b76
commit c61152baa4
18 ha cambiato i file con 153 aggiunte e 137 eliminazioni
+4 -4
Vedi File
@@ -618,7 +618,7 @@ __device__ __attribute__((noinline)) void reduceCopyPacksWithBias(
thread = warp*WARP_SIZE + lane;
}
template<int Unroll, typename RedFn, typename T,
template<int Unroll, int useAcc, typename RedFn, typename T,
int MultimemSrcs, int MinSrcs, int MaxSrcs,
int MultimemDsts, int MinDsts, int MaxDsts, int PreOpSrcs,
typename IntBytes, typename SrcPtrFn, typename DstPtrFn, typename AccPtrFn>
@@ -641,7 +641,7 @@ __device__ __forceinline__ void reduceCopy(
IntBytes nBytesBehind = 0;
IntBytes nBytesAhead = nElts*sizeof(T);
bool useAcc = accPtrFn() != nullptr;
//bool useAcc = accPtrFn() != nullptr;
#if __cpp_if_constexpr
if constexpr (BigPackSize > sizeof(T)) {
@@ -763,7 +763,7 @@ __device__ __forceinline__ void reduceCopy(
nSrcs, srcPtrFn, nDsts, dstPtrFn, /*&*/nBytesBehind, /*&*/nBytesAhead);
}
template<int Unroll, typename RedFn, typename T,
template<int Unroll, int useAcc, typename RedFn, typename T,
int MultimemSrcs, int MinSrcs, int MaxSrcs,
int MultimemDsts, int MinDsts, int MaxDsts, int PreOpSrcs,
typename IntBytes>
@@ -773,7 +773,7 @@ __device__ __forceinline__ void reduceCopy(
int nSrcs, void** srcPtrs, int nDsts, void** dstPtrs,
IntBytes nElts, void *accPtr = nullptr
) {
reduceCopy<Unroll, RedFn, T,
reduceCopy<Unroll, useAcc, RedFn, T,
MultimemSrcs, MinSrcs, MaxSrcs,
MultimemDsts, MinDsts, MaxDsts, PreOpSrcs, IntBytes>
(thread, nThreads, redArg, preOpArgs, postOp,