AI/rocm-systems
2
0
Fork 0
Code Issues Pull Requests Actions Packages Projects Releases Wiki Activity

2.22.3-1

Browse Source 
Rework core for NVIDIA Trusted Computing
 * Compress work structs so that they are shared between channels
 * Utilize the full amount of kernel argument space permitted (4k)
   before resorting to work fifo.
 * Rework the task preprocessing phase.
 * Use a separate abortDevFlag which is kept in sync with abortFlag
   using cudaMemcpy operations.
 * Rename src/include/align.h to src/include/bitops.h

Add lazy connection establishment for collective operations
 * Move buffer allocation and connection establishment to the first
   collective operation using that algorithm.
 * Accelerate init time and reduce memory usage.
 * Avoid allocating NVLS buffers if all calls are registered.
 * Compute algo/proto in ncclLaunchCollTasksInfo early on.
 * Connect peers in ncclCollPreconnectFunc if not connected already.
 * Also move shared buffer creation to the first send/recv call.

Accelerate intra-node NVLink detection
 * Make each rank only detect NVLinks attached to its GPU.
 * Fuse XMLs to reconstruct the full NVLink topology

Add init profiling to report time spend in different init phases.
 * Report timings of bootstrap, allgather, search, connect, etc.
 * Add new "PROFILE" category for NCCL_DEBUG_SUBSYS.

Add support for PCI p2p on split PCI switches
 * Detect split PCI switches through a kernel module exposing
   switch information.
 * Update the topology XML and graph to add those inter-switch
   connections.

Add cost estimation API
 * Add a new ncclGroupEndSimulate primitive to return the estimated
   time a group would take.

Net/IB: Add separate traffic class for fifo messages
 * Add NCCL_IB_FIFO_TC to control the traffic class of fifo messages
   independently from NCCL_IB_TC.
   Merges PR #1194

Net/IB: Add support for IB router
 * Use flid instead of lid if subnets do not match
 * Warn if flid is 0

Optimizations and fixes for device network offload (unpack)
 * Double the default number of channels
 * Cache netDeviceType
 * Fix save/increment head logic to enable Tree support.

Support ncclGroupStart/End for ncclCommAbort/Destroy
 * Allow Abort/Destroy to be called within a group when managing
   multiple GPUs with a single process.

Improve Tuner API
 * Provide to the plugin the original cost table so that the plugin
   can leave unknown or disabled algo/proto combinations untouched.
 * Remove nvlsSupport and collnetSupport.

Do not print version to stdout when using a debug file
 * Also print version from all processes with INFO debug level.
   Fixes issue #1271

Fix clang warnings in NVTX headers
 * Update NVTX headers to the latest version
   Fixes issue #1270

Disable port fusion in heterogeneous systems
 * Do not fuse ports if a mix of multi-port and single port are detected.

Fix NVLS graphs search for dual NICs.
 * Fix NVLS graph search when we have more than one NIC per GPU.

Fix crash with collnetDirect
 * Add separate graph search for collnetDirect, testing alltoall paths
   and working similarly to the NVLS search.

Fix hang when nodes have different CPU types
 * Add the CPU type to the rank peer info.
 * Align all ranks on the CPU type after the first allgather.
 * Only use the aligned CPU type for all tuning operations.
   Fixes issue #1136
   Fixes issue #1184

Fix performance of registered send/recv operations
 * Allow for single full size operations
 * Add INFO to confirm the registration of send/recv buffers.

Move all sync ops to finalize stage
 * Ensure ncclCommDestroy is non-blocking if ncclCommFinalize has
   been called.

Improve error reporting during SHM segment creation

Improve support of various compilers
   Merges PR #1177
   Merges PR #1228

Allow net and tuner plugins to be statically linked
 * Search for ncclNet or ncclTuner symbols in the main binary.
   Merges PR #979

Plugin examples includes cleanup
 * Harmonize err.h and common.h usage.
 * Add mixed plugin with both net and tuner.


[ROCm/rccl commit: 178b6b7590]
This commit is contained in:
Sylvain Jeaugey
2024-06-11 01:28:01 -07:00
parent 1cabd1d638
commit 5ca1b6c160
115 changed files with 8672 additions and 4403 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/src/include/device.h
+182 -95
View File
@@ -9,8 +9,10 @@
#include "nccl.h"
#include "nccl_common.h"
#include "align.h"
#include "bitops.h"
#include <algorithm>
#include <stdint.h>
#include <sys/types.h>
extern const char* ncclFuncStr[NCCL_NUM_FUNCTIONS];
@@ -21,6 +23,12 @@ extern const char* ncclProtoStr[NCCL_NUM_PROTOCOLS];
#define NCCL_MAX_OPS 2048
#define NCCL_STEPS 8
#ifdef __CUDA_ARCH__
#define NCCL_CUDA_ARCH __CUDA_ARCH__
#else
#define NCCL_CUDA_ARCH 0
#endif
#include "net_device.h"
enum ncclDevRedOp_t {
@@ -52,8 +60,11 @@ union ncclLLFifoLine {
#define WARP_SIZE 32
#define MAXCHANNELS 32
#define NCCL_MAX_LOCAL_RANKS 64
#define NCCL_MAX_NTHREADS 640
#define NCCL_MIN_NTHREADS (4*WARP_SIZE)
#define NCCL_SIMPLE_MAX_NTHREADS 512
#define NCCL_SIMPLE_EXTRA_GROUP_IF_NTHREADS_GE (3*WARP_SIZE)
#define NCCL_LL_MAX_NTHREADS 512
#define NCCL_LL_LINES_PER_THREAD 8
#ifdef TEST_LL_CLEANUP
@@ -84,6 +95,9 @@ static_assert(NCCL_LL_CLEAN_MASK % NCCL_STEPS == 0, "Invalid NCCL_LL_CLEAN_MASK
#define NCCL_IPC_READ 0x10
#define NCCL_NVLS_MIN_POLL 0x20
// Number of named barriers supported by CUDA
#define NCCL_MAX_GROUPS 16
#define NCCL_MAX_COLLNET_SIZE (1L << 29)
enum ncclRegBufferType {
@@ -196,112 +210,155 @@ struct ncclChannelPeer {
struct ncclDevComm;
/* ncclWork is to be a power of two, currently 8x64 bytes, */
/* to make sure reads to host from the CUDA kernel are aligned. */
/* Make sure to adjust padding at the end of ncclWorkElem. */
#define NCCL_WORK_SIZE 512
struct alignas(16) ncclDevWorkP2p {
void *sendAddr, *recvAddr;
size_t sendBytes, recvBytes;
int sendRank, recvRank;
// From the part index, nP2pChannels, and channelBase the device code can
// calculate which part of the transfer a channel is responsible for.
uint8_t nP2pChannels; // Always equal to comm->p2pnChannels
uint8_t channelBase; // Channel owning first part.
// Zero channels indicates no work in that direction.
uint8_t nSendChannels, nRecvChannels;
// Chunk size stored in 8 bits via u32fp8Encode/Decode.
uint8_t sendChunkSize_u32fp8, recvChunkSize_u32fp8;
enum ncclWorkType : uint8_t {
ncclWorkTypeUnused=0,
ncclWorkTypeColl=1,
ncclWorkTypeP2p=2,
ncclWorkTypeRegColl=3
};
enum ncclWorkP2PType : uint8_t {
ncclWorkP2pTypeUnused=0,
ncclWorkP2pTypeSend,
ncclWorkP2pTypeRecv
uint8_t sendProtoLL:1, recvProtoLL:1;
uint8_t sendRegistered:1, recvRegistered:1;
};
struct ncclWorkHeader {
union {
int32_t workNext; // when isLast=0: Offset from kernel argument workHead
uint32_t doneAcks; // when isLast=1: Monotonic (mod 1<<32) ack value to send back.
};
uint16_t funcIndex;
uint8_t isLast:1; // last work for this kernel
uint8_t inFifo:1; // is this work in the fifo
enum ncclWorkType type;
};
// Compute the subset of the data transfer corresponding to the given part index.
inline __host__ __device__ void ncclP2pPartBounds(int nParts, int part, size_t bytes, size_t* partBeg, size_t* partEnd) {
size_t partBytes = alignUp(divUp(bytes, nParts), 4<<10);
#if __CUDA_ARCH__
*partBeg = min((part+0)*partBytes, bytes);
*partEnd = min((part+1)*partBytes, bytes);
#else
*partBeg = std::min<size_t>((part+0)*partBytes, bytes);
*partEnd = std::min<size_t>((part+1)*partBytes, bytes);
#endif
}
struct ncclWorkElem {
union {
uint8_t flagBits;
struct {
uint8_t isUsed:1, redOpArgIsPtr:1, oneNode:1;
};
};
uint8_t regUsed;
uint8_t nWarps;
uint8_t direct;
// implemented in channel.h
inline __host__ uint8_t ncclP2pChannelBaseForRound(struct ncclComm* comm, int p2pRound);
// ncclP2pChannelToPart and ncclP2pChannelForPart are inverses. The device code
// uses ncclP2pChannelToPart to determine which part "this" channel is responsible for.
inline __host__ int ncclP2pChannelForPart(int nP2pChannels, int base, int part) {
// Only works because nP2pChannels is pow2
int nChannelsLog2 = countOneBits(nP2pChannels-1);
int delta = reverseBits(part, nChannelsLog2);
return (base + delta) & (nP2pChannels-1);
}
inline __device__ int ncclP2pChannelToPart(int nP2pChannels, int base, int channel) {
// Only works because nP2pChannels is pow2
int nChannelsLog2 = countOneBits(nP2pChannels-1);
int delta = (channel-base) & (nP2pChannels-1);
return reverseBits(delta, nChannelsLog2);
}
struct alignas(16) ncclDevWorkColl {
// Running on channels [channelLo..channelHi], hi is inclusive.
// nChannels == (channelHi - channelLo) + 1
uint32_t channelLo:8, channelHi:8;
uint32_t nWarps:8;
uint32_t redOpArgIsPtr:1, regUsed:2, oneNode:1, direct:4;
uint32_t root;
const void *sendbuff;
void *recvbuff;
size_t count;
uint64_t redOpArg;
uint64_t chunkCount:25, workCount:39;
void* recvbuff;
void* sendbuff;
union {
// Continuous-byte-distribution scheduling. The lo and hi channels are of
// different size than the channels in the middle.
struct {
uint64_t lastChunkCount:25;
uint64_t workOffset:39;
};
size_t countLo, countMid, countHi;
// Chunk counts where units are ncclProtoGrainSize(protocol) bytes
uint64_t chunkGrainsLo:21, chunkGrainsMid:21, chunkGrainsHi:21;
} cbd;
// Collnet scheduling. All channels divide work evenly.
struct {
uint64_t bid:32;
uint64_t nChannels:32;
};
size_t count; // Total size, not divided per channel.
uint32_t chunkCount;
} collnet;
};
uint64_t redOpArg;
};
#define NCCL_MAX_WORK_ELEMENTS ((NCCL_WORK_SIZE - alignUp(sizeof(ncclWorkHeader), alignof(ncclWorkElem)))/sizeof(ncclWorkElem))
static_assert(NCCL_MAX_WORK_ELEMENTS == 9, "Sanity check: NCCL_MAX_WORK_ELEMENTS == 9");
struct ncclWorkElemP2p {
int peer : 30;
int proto : 2;
__host__ __device__ constexpr int ncclProtoGrainSize(int proto) {
return proto == NCCL_PROTO_LL ? 16 :
proto == NCCL_PROTO_LL128 ? WARP_SIZE*NCCL_LL128_SHMEM_ELEMS_PER_THREAD/NCCL_LL128_LINEELEMS*NCCL_LL128_DATAELEMS*sizeof(uint64_t) :
proto == NCCL_PROTO_SIMPLE ? 512 :
-1;
}
enum ncclWorkP2PType p2pType;
uint8_t reg:1;
uint8_t nWarps:5;
uint8_t warpStart;
uint8_t ngroups;
// Important not to use any fields with greater than 4-byte alignment since
// we need sizeof(ncclWorkElemP2p)==28, but that would be padded up to 32 if
// there were 8-byte fields.
//void* buff;
uint32_t buffHi32, buffLo32; // buff = buffHi32<<32 | buffLo32;
//size_t count;
uint32_t countHi32, countLo32; // count = countHi32<<32 | countLo32;
int chunkSize;
};
template<typename Int>
__host__ __device__ inline void ncclCollCbdPart(
struct ncclDevWorkColl* work, uint32_t channelId, int proto, int eltSize,
Int* count, Int* partOffset, Int* partCount, Int* chunkCount
) {
int eltPerGrain = ncclProtoGrainSize(proto)/eltSize;
int nMidChannels = work->channelHi - work->channelLo - 1;
// We can assum that nMidChannels<0 implies countMid==0, which let's us assume
// that countMid*nMidChannels == 0.
if (count != nullptr) {
*count = work->cbd.countLo + work->cbd.countMid*nMidChannels + work->cbd.countHi;
}
if (channelId == work->channelLo) {
*partOffset = 0;
*partCount = work->cbd.countLo;
*chunkCount = work->cbd.chunkGrainsLo*eltPerGrain;
} else if (channelId == work->channelHi) {
*partOffset = work->cbd.countLo + nMidChannels*work->cbd.countMid;
*partCount = work->cbd.countHi;
*chunkCount = work->cbd.chunkGrainsHi*eltPerGrain;
} else {
int mid = channelId - work->channelLo - 1;
*partOffset = work->cbd.countLo + mid*work->cbd.countMid;
*partCount = work->cbd.countMid;
*chunkCount = work->cbd.chunkGrainsMid*eltPerGrain;
}
}
static_assert(((NCCL_WORK_SIZE - alignUp(sizeof(ncclWorkHeader), alignof(ncclWorkElemP2p)))/sizeof(ncclWorkElemP2p)) >= 16, "Sanity check: NCCL_MAX_WORK_ELEMENTS_P2P == 16");
#define NCCL_MAX_WORK_ELEMENTS_P2P 16
struct ncclWorkElemReg {
struct ncclWorkElem elem;
struct alignas(16) ncclDevWorkCollReg {
struct ncclDevWorkColl coll;
void* dnInputs[NCCL_MAX_DIRECT_ARITY+1];
void* dnOutputs[NCCL_MAX_DIRECT_ARITY+1];
void* upOutputs[NCCL_MAX_DIRECT_ARITY+1];
};
#define NCCL_MAX_WORK_ELEMENTS_REG ((NCCL_WORK_SIZE - alignUp(sizeof(ncclWorkHeader), alignof(ncclWorkElemReg)))/sizeof(ncclWorkElemReg))
static_assert(NCCL_MAX_WORK_ELEMENTS_REG == 2, "Sanity check: NCCL_MAX_WORK_ELEMENTS_REG == 2");
// Number of named barriers supported by CUDA
#define NCCL_MAX_GROUPS 16
struct ncclWork {
struct ncclWorkHeader header;
union {
char pad[NCCL_WORK_SIZE - sizeof(struct ncclWorkHeader)];
struct ncclWorkElem elems[NCCL_MAX_WORK_ELEMENTS];
struct ncclWorkElemP2p p2pElems[NCCL_MAX_WORK_ELEMENTS_P2P];
struct ncclWorkElemReg regElems[NCCL_MAX_WORK_ELEMENTS_REG];
};
enum ncclDevWorkType: uint8_t {
ncclDevWorkTypeP2p,
ncclDevWorkTypeColl,
ncclDevWorkTypeCollReg
};
constexpr size_t ncclDevWorkSize(enum ncclDevWorkType type) {
return type == ncclDevWorkTypeP2p ? sizeof(ncclDevWorkP2p) :
type == ncclDevWorkTypeColl ? sizeof(ncclDevWorkColl) : sizeof(ncclDevWorkCollReg);
}
#define NCCL_MAX_DEV_WORK_BATCH_BYTES 1024
#define NCCL_MAX_DEV_WORK_BATCH_COLLS (NCCL_MAX_DEV_WORK_BATCH_BYTES/sizeof(ncclDevWorkColl))
#define NCCL_MAX_DEV_WORK_P2P_PER_BATCH 8
struct alignas(16) ncclDevWorkBatch {
union {
struct {
// nextExtends: should next one be merged into this one.
// nextJump=0: end of this channel's batch list
// nextJump>0: batches[thisIndex+nextJump] is next batch in this list
uint32_t nextJump:14, nextExtends:1;
uint32_t workType:2, funcId:15;
};
// Unioning bitfields with underlying type hints compiler to emit the best
// SASS LD/ST accesses.
uint32_t flags;
};
// Rolling offset in fifo where this batch's work structs begin
uint32_t offsetBase;
// Set of relative offsets from offsetBase for this channel's subset of the batch:
// For each bit index i in offsetMask, find work at fifo offset: offsetBase + i*sizeof(WorkStructType)
uint64_t offsetBitset;
};
static_assert(sizeof(struct ncclWork) == NCCL_WORK_SIZE, "Sanity check: sizeof(struct ncclWork) == NCCL_WORK_SIZE");
static_assert(sizeof(struct ncclWork)%16 == 0, "Sanity check: sizeof(struct ncclWork)%16 == 0");
struct ncclDevChannelPeer {
// Stripped version of ncclChannelPeer where we only keep the ncclConnInfo
@@ -328,9 +385,8 @@ struct ncclDevComm {
int buffSizes[NCCL_NUM_PROTOCOLS];
int p2pChunkSize;
// Operation list for aggregation
int workFifoDepth;
struct ncclWork* workFifoHeap; // may be cudaHost or GDR memory
// Work fifo return credits
uint32_t* workConsumed/*[MAXCHANNELS]*/;
int* collNetDenseToUserRank;
@@ -346,11 +402,37 @@ struct alignas(16) ncclDevCommAndChannels {
struct ncclDevChannel channels[MAXCHANNELS];
};
#ifdef __CUDA_ARCH__
#define NCCL_CUDA_ARCH __CUDA_ARCH__
#else
#define NCCL_CUDA_ARCH 0
#endif
enum ncclDevWorkStorageType: uint8_t {
ncclDevWorkStorageTypeArgs=0,
ncclDevWorkStorageTypeFifo=1,
ncclDevWorkStorageTypePersistent=2
};
struct alignas(16) ncclDevKernelArgs {
struct ncclDevComm* comm;
uint64_t channelMask;
enum ncclDevWorkStorageType workStorageType;
uint32_t workMask;
void* workBuf;
// A channel's first batch is at `blockIdx.x`. Use `nextJump` to follow rest of list.
// struct ncclDevWorkBatch batches[];
};
__host__ __device__ constexpr int ncclMaxKernelArgsSize(/*int cudaDriver, */int cudaArch=NCCL_CUDA_ARCH) {
//return (cudaArch < 700 || cudaDriver < 12010) ? 4<<10 : (32<<10)-4;
return 4<<10;
}
template<size_t capacity>
struct alignas(16) ncclDevKernelArgsStorage {
union {
struct ncclDevKernelArgs args;
ulong2 storage[capacity/sizeof(ulong2)];
};
};
typedef ncclDevKernelArgsStorage<(4<<10)> ncclDevKernelArgs4K;
//typedef ncclDevKernelArgsStorage<(32<<10)-4> ncclDevKernelArgs31K;
template<typename T>
__host__ __device__ constexpr T min_constexpr(T a) { return a; }
@@ -366,6 +448,10 @@ __host__ __device__ constexpr T max_constexpr(T a, T b, Ts ...c) {
return max_constexpr<T>((a > b ? a : b), c...);
}
constexpr int ncclDevMaxChannelsForArgsBytes(size_t argsBytes) {
return min_constexpr<size_t>(MAXCHANNELS, (argsBytes - sizeof(struct ncclDevKernelArgs))/sizeof(struct ncclDevWorkBatch));
}
// Calculate the unroll factor given:
// * bytePerPack: number of bytes accessed per instruction
// * insns: max permissible unroll value
@@ -412,6 +498,7 @@ extern int const ncclDevKernelCount;
extern void* const ncclDevKernelList[/*ncclDevKernelCount*/];
// Table of most specialized kernel function to run given func index.
extern int const ncclDevFuncIdCount;
extern int const ncclDevFuncRowToId[];
extern void* const ncclDevKernelForFunc[/*funcIndex*/];
extern bool const ncclDevKernelForFuncIsSpecialized[/*funcIndex*/];