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

removed gfx940 and gfx941 (#1606)

Browse Source 
* removed gfx940 and gfx941

* removed gfx940 and gfx941

* Update "gfx94" to "gfx942" in init.cc

* Updated remaining "gfx94" updates to "gfx942"

* Update filenames and variables from gfx940 to gfx942

---------

Co-authored-by: akolliasAMD <akollias@amd.com>

[ROCm/rccl commit: 6505639cf4]
This commit is contained in:
corey-derochie-amd
2025-03-20 09:34:53 -06:00
committed by GitHub
parent a95b2c9fb7
commit e95578ef4c
30 changed files with 74 additions and 79 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/cmake/scripts/extract_metadata.cmake
+1 -1
View File
@@ -31,7 +31,7 @@ if(list_result EQUAL 0)
## Extract file paths for the selected gfx archs
foreach(line ${cmd_output})
if(line MATCHES "(gfx90a|gfx940|gfx941|gfx942|gfx950)")
if(line MATCHES "(gfx90a|gfx942|gfx950)")
string(REGEX MATCH "\\file://(.*)" file_match ${line})
if(file_match)
list(APPEND file_paths ${file_match})
projects/rccl/src/device/all_gather.h
+1 -1
View File
@@ -11,7 +11,7 @@
namespace {
template<typename T, typename RedOp, typename Proto>
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx940__) && !defined(__gfx941__) && !defined(__gfx942__) && !defined(__gfx950__)
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx942__) && !defined(__gfx950__)
__device__ void runRing(int tid, int nthreads, struct ncclDevWorkColl* work) {
#else
__device__ __attribute__((noinline)) void runRing(int tid, int nthreads, struct ncclDevWorkColl* work) {
projects/rccl/src/device/all_reduce.h
+3 -3
View File
@@ -15,7 +15,7 @@
namespace {
template<typename T, typename RedOp, typename Proto>
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx940__) && !defined(__gfx941__) && !defined(__gfx942__) && !defined(__gfx950__)
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx942__) && !defined(__gfx950__)
__device__ void runRing(int tid, int nthreads, struct ncclDevWorkColl* work) {
#else
__device__ __attribute__((noinline)) void runRing(int tid, int nthreads, struct ncclDevWorkColl* work) {
@@ -209,7 +209,7 @@ namespace {
}
template<typename T, typename RedOp, typename Proto>
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx940__) && !defined(__gfx941__) && !defined(__gfx942__) && !defined(__gfx950__)
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx942__) && !defined(__gfx950__)
__device__ void runTreeUpDown(int tid, int nthreads, struct ncclDevWorkColl* work) {
#else
__device__ __attribute__((noinline)) void runTreeUpDown(int tid, int nthreads, struct ncclDevWorkColl* work) {
@@ -357,7 +357,7 @@ namespace {
}
template<typename T, typename RedOp, typename Proto>
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx940__) && !defined(__gfx941__) && !defined(__gfx942__) && !defined(__gfx950__)
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx942__) && !defined(__gfx950__)
__device__ void runTreeSplit(int tid, int nthreads, struct ncclDevWorkColl* work) {
#else
__device__ __attribute__((noinline)) void runTreeSplit(int tid, int nthreads, struct ncclDevWorkColl* work) {
projects/rccl/src/device/alltoall_pivot.h
+1 -1
View File
@@ -10,7 +10,7 @@
namespace {
template<typename T, typename RedOp, typename Proto>
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx940__) && !defined(__gfx941__) && !defined(__gfx942__) && !defined(__gfx950__)
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx942__) && !defined(__gfx950__)
__device__ void runRing(int tid, int nthreads, struct ncclDevWorkColl* work) {
#else
__device__ __attribute__((noinline)) void runRing(int tid, int nthreads, struct ncclDevWorkColl* work) {
projects/rccl/src/device/broadcast.h
+1 -1
View File
@@ -10,7 +10,7 @@
namespace {
template<typename T, typename RedOp, typename Proto>
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx940__) && !defined(__gfx941__) && !defined(__gfx942__) && !defined(__gfx950__)
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx942__) && !defined(__gfx950__)
__device__ void runRing(int tid, int nthreads, struct ncclDevWorkColl* work) {
#else
__device__ __attribute__((noinline)) void runRing(int tid, int nthreads, struct ncclDevWorkColl* work) {
projects/rccl/src/device/msccl_kernel_impl.h
+1 -1
View File
@@ -152,7 +152,7 @@ __device__ __forceinline__ void mscclRunInterpreter(
int npKitCtxIdx = bid;
int xcc_id = 0;
if (tid == 0) {
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
#if defined(__gfx942__) || defined(__gfx950__)
asm volatile ("s_getreg_b32 %0, hwreg(HW_REG_XCC_ID)" : "=s" (xcc_id));
#endif
}
projects/rccl/src/device/onerank.cu
+1 -1
View File
@@ -11,7 +11,7 @@
#include "common.h"
#include <cuda_runtime.h>
#if defined(__gfx908__) || defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
#if defined(__gfx908__) || defined(__gfx942__) || defined(__gfx950__)
#define COLL_UNROLL 2
#else
#define COLL_UNROLL 4
projects/rccl/src/device/op128.h
+1 -1
View File
@@ -124,7 +124,7 @@ union alignas(16) BytePack<16> {
uint32_t u32[4];
uint64_t u64[2];
ulong2 ul2, native;
#if !defined(USE_INDIRECT_FUNCTION_CALL) || defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
#if !defined(USE_INDIRECT_FUNCTION_CALL) || defined(__gfx942__) || defined(__gfx950__)
inline __device__ BytePack<16>& operator=(BytePack<16> other) {
u64[0] = other.u64[0];
u64[1] = other.u64[1];
projects/rccl/src/device/primitives.h
+1 -1
View File
@@ -15,7 +15,7 @@
#define NCCL_SPINS_BEFORE_CHECK_ABORT 1000000
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
#if defined(__gfx942__) || defined(__gfx950__)
#define __THREAD_FENCE __threadfence_block()
#else
#define __THREAD_FENCE __threadfence()
projects/rccl/src/device/reduce.h
+1 -1
View File
@@ -11,7 +11,7 @@
namespace {
template<typename T, typename RedOp, typename Proto>
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx940__) && !defined(__gfx941__) && !defined(__gfx942__) && !defined(__gfx950__)
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx942__) && !defined(__gfx950__)
__device__ void runRing(int tid, int nthreads, struct ncclDevWorkColl* work) {
#else
__device__ __attribute__((noinline)) void runRing(int tid, int nthreads, struct ncclDevWorkColl* work) {
projects/rccl/src/device/reduce_scatter.h
+2 -2
View File
@@ -11,7 +11,7 @@
namespace {
template<typename T, typename RedOp, typename Proto>
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx940__) && !defined(__gfx941__) && !defined(__gfx942__) && !defined(__gfx950__)
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx942__) && !defined(__gfx950__)
__device__ void runRing(int tid, int nthreads, struct ncclDevWorkColl* work) {
#else
__device__ __attribute__((noinline)) void runRing(int tid, int nthreads, struct ncclDevWorkColl* work) {
@@ -165,7 +165,7 @@ struct RunWorkColl<ncclFuncReduceScatter, T, RedOp, NCCL_ALGO_NVLS, NCCL_PROTO_S
size_t offset;
int nelem;
/* if we are direct NVLS, we only need to allocate 1 warp to scatter for sync;
/* if we are direct NVLS, we only need to allocate 1 warp to scatter for sync;
* if not, based on #ranks, we allocate 7 or 5 warps to reduce to saturate bandwidth
* and the rest are allocated to scatter. */
const int nThreadsReduce = work->regUsed ? (NCCL_MAX_NTHREADS - WARP_SIZE) : (nranks <= 6 ? 7 * WARP_SIZE : 5 * WARP_SIZE);
projects/rccl/src/device/sendrecv.h
+5 -5
View File
@@ -20,7 +20,7 @@ struct RunWorkBatch<ncclFuncSendRecv, T, RedOp, NCCL_ALGO_RING, NCCL_PROTO_SIMPL
__device__ void runSend(int tid, int tn, int group, struct ncclDevWorkP2p* work) {
size_t bytes = work->sendBytes;
int chunkSize = u32fp8Decode(work->sendChunkSize_u32fp8);
#if defined(ENABLE_NPKIT)
bool isNpKitThread = (tid == 0);
int npKitCtxIdx = blockIdx.x + group;
@@ -132,7 +132,7 @@ struct RunWorkBatch<ncclFuncSendRecv, T, RedOp, NCCL_ALGO_RING, NCCL_PROTO_SIMPL
#endif
}
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx940__) && !defined(__gfx941__) && !defined(__gfx942__) && !defined(__gfx950__)
#if defined(USE_INDIRECT_FUNCTION_CALL) && !defined(__gfx942__) && !defined(__gfx950__)
__device__ void run() {
#else
__device__ __attribute__((noinline)) void run() {
@@ -241,7 +241,7 @@ struct RunWorkBatch<ncclFuncSendRecv, T, RedOp, NCCL_ALGO_RING, NCCL_PROTO_SIMPL
}
if (isCopy) {
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
#if defined(__gfx942__) || defined(__gfx950__)
reduceCopy<COLL_UNROLL*2, RedOp, T, 0,1,1, 0,1,1, /*PreOpSrcs=*/0>
(subtid, subtn, 0, nullptr, false, 1, &work->sendAddr, 1, &work->recvAddr, (ssize_t)work->sendBytes);
#else
@@ -254,7 +254,7 @@ struct RunWorkBatch<ncclFuncSendRecv, T, RedOp, NCCL_ALGO_RING, NCCL_PROTO_SIMPL
} else {
#if defined(__gfx90a__)
runSend<ProtoSimple<1,1,8>>(subtid, subtn, group, work);
#elif defined(__gfx908__) || defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
#elif defined(__gfx908__) || defined(__gfx942__) || defined(__gfx950__)
runSend<ProtoSimple<1,1,4>>(subtid, subtn, group, work);
#else
runSend<ProtoSimple<1,1>>(subtid, subtn, group, work);
@@ -266,7 +266,7 @@ struct RunWorkBatch<ncclFuncSendRecv, T, RedOp, NCCL_ALGO_RING, NCCL_PROTO_SIMPL
} else {
#if defined(__gfx90a__)
runRecv<ProtoSimple<1,1,8>>(subtid, subtn, group, work);
#elif defined(__gfx908__) || defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
#elif defined(__gfx908__) || defined(__gfx942__) || defined(__gfx950__)
runRecv<ProtoSimple<1,1,4>>(subtid, subtn, group, work);
#else
runRecv<ProtoSimple<1,1>>(subtid, subtn, group, work);
projects/rccl/src/graph/connect.cc
+2 -2
View File
@@ -286,7 +286,7 @@ static ncclResult_t setTreeDown(struct ncclTree* tree, int* indexes, int d) {
static ncclResult_t connectTrees(struct ncclComm* comm, int* treeToParent, int* treeToChild0, int* treeToChild1, int* treePatterns) {
const int channelLimit = (IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94") || IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) ? 2*CHANNEL_LIMIT : CHANNEL_LIMIT;
const int channelLimit = (IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx942") || IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) ? 2*CHANNEL_LIMIT : CHANNEL_LIMIT;
const int nChannels = (comm->nChannels > channelLimit) ? comm->nChannels / 2 : comm->nChannels;
const int nNodes = comm->nNodes, node = comm->node;
@@ -804,7 +804,7 @@ ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, int* treePa
}
// Only use full MAXCHANNELS for gfx94x and gfx950
int maxChannels = (IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94") || IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) ?
int maxChannels = (IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx942") || IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) ?
((comm->topo->nodes[GPU].nodes[0].gpu.cu == 80 || comm->topo->nodes[GPU].nodes[0].gpu.cu == 20 || comm->topo->nodes[GPU].nodes[0].gpu.cu == 38)
? comm->topo->nodes[GPU].nodes[0].gpu.cu : MAXCHANNELS) : 2*CHANNEL_LIMIT;
projects/rccl/src/graph/paths.cc
+3 -3
View File
@@ -666,7 +666,7 @@ ncclResult_t ncclTopoComputePaths(struct ncclTopoSystem* system, struct ncclComm
int arch, vendor, model;
NCCLCHECK(ncclTopoCpuType(system, &arch, &vendor, &model));
if (arch == NCCL_TOPO_CPU_ARCH_X86 && vendor == NCCL_TOPO_CPU_VENDOR_INTEL &&
(IsArchMatch(system->nodes[GPU].nodes[0].gpu.gcn, "gfx94") || IsArchMatch(system->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) &&
(IsArchMatch(system->nodes[GPU].nodes[0].gpu.gcn, "gfx942") || IsArchMatch(system->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) &&
((system->nodes[GPU].count == 8 && system->nodes[NET].count == 8 && system->nodes[GPU].count == system->nRanks) ||
(system->nodes[GPU].count != system->nRanks))) {
if (!rcclPathOverride(system, 0x100000) && !rcclPathOverride(system, 0x1000))
@@ -843,7 +843,7 @@ static ncclResult_t ncclTopoGetNchannels(struct ncclComm* comm, int g /*local gp
path = system->nodes[GPU].nodes[peer].paths[GPU]+g;
if (path->type == PATH_NVL) {
float nvlBw = ncclTopoXGMISpeed(system->nodes[GPU].nodes[g].gpu.gcn);
*nChannels = ((IsArchMatch(system->nodes[GPU].nodes[0].gpu.gcn, "gfx94") || IsArchMatch(system->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) ? 4 : 2)*std::max(1, (int)(path->bw / nvlBw));
*nChannels = ((IsArchMatch(system->nodes[GPU].nodes[0].gpu.gcn, "gfx942") || IsArchMatch(system->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) ? 4 : 2)*std::max(1, (int)(path->bw / nvlBw));
} else {
*nChannels = 2;
}
@@ -906,7 +906,7 @@ ncclResult_t ncclTopoComputeP2pChannels(struct ncclComm* comm) {
// Round to next pow2 nChannelsPerPeer and nChannels
comm->p2pnChannelsPerPeer = (ncclParamNChannelsPerPeer() == -2 ? pow2Up(minChannels) : ncclParamNChannelsPerPeer());
// Doubling P2P channels per peer on single node
if (comm->topo->nodes[GPU].count == comm->topo->nRanks && (IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94") || IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950"))) comm->p2pnChannelsPerPeer *= 2;
if (comm->topo->nodes[GPU].count == comm->topo->nRanks && (IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx942") || IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950"))) comm->p2pnChannelsPerPeer *= 2;
comm->p2pnChannels = std::min(pow2Up(comm->p2pnChannels), 4*CHANNEL_LIMIT);
// p2pnChannelsPerPeer cannot be greater than MAXCHANNELS
comm->p2pnChannelsPerPeer = std::min(comm->p2pnChannelsPerPeer, MAXCHANNELS);
projects/rccl/src/graph/topo.h
+1 -1
View File
@@ -262,7 +262,7 @@ static ncclResult_t ncclTopoIdToNetDev(struct ncclTopoSystem* system, int64_t id
static float ncclTopoXGMISpeed(const char* gcn) {
if (IsArchMatch(gcn, "gfx90a"))
return MI200_XGMI_WIDTH;
else if (IsArchMatch(gcn, "gfx94"))
else if (IsArchMatch(gcn, "gfx942"))
return GFX94X_XGMI_WIDTH;
else if (IsArchMatch(gcn, "gfx95"))
return GFX95X_XGMI_WIDTH;
projects/rccl/src/graph/tuning.cc
+2 -2
View File
@@ -360,7 +360,7 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom
if (coll == ncclFuncAllGather && a != NCCL_ALGO_RING && a != NCCL_ALGO_NVLS && a != NCCL_ALGO_COLLNET_DIRECT) continue;
for (int p=0; p<NCCL_NUM_PROTOCOLS; p++) {
if (a == NCCL_ALGO_TREE && p == NCCL_PROTO_SIMPLE && (IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94") || IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) && comm->topo->nodes[GPU].count == comm->topo->nRanks) continue;
if (a == NCCL_ALGO_TREE && p == NCCL_PROTO_SIMPLE && (IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx942") || IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) && comm->topo->nodes[GPU].count == comm->topo->nRanks) continue;
if ((a == NCCL_ALGO_NVLS || a == NCCL_ALGO_NVLS_TREE) && p != NCCL_PROTO_SIMPLE) continue;
int collnet = (a == NCCL_ALGO_COLLNET_DIRECT || a == NCCL_ALGO_COLLNET_CHAIN) ? 1 : 0;
float bw = nNodes <= 2 || collnet ? graphs[a]->bwIntra : graphs[a]->bwInter;
@@ -376,7 +376,7 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom
busBw *= rcclTuningModel[comm->topo->tuning].bwRatio[0][a][p];
else
busBw *= rcclTuningModel[comm->topo->tuning].bwRatio[1][a][p];
if (a == NCCL_ALGO_RING && p == NCCL_PROTO_LL && (coll == ncclFuncBroadcast || coll == ncclFuncReduce) && (IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94") || IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) && comm->topo->nodes[GPU].count == comm->topo->nRanks) { busBw = busBw * 1.65; }
if (a == NCCL_ALGO_RING && p == NCCL_PROTO_LL && (coll == ncclFuncBroadcast || coll == ncclFuncReduce) && (IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx942") || IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) && comm->topo->nodes[GPU].count == comm->topo->nRanks) { busBw = busBw * 1.65; }
#else
if (a == NCCL_ALGO_RING && p == NCCL_PROTO_LL) { busBw = std::min(llMaxBw, busBw * .5); }
if (a == NCCL_ALGO_RING && p == NCCL_PROTO_LL128) busBw = std::min(busBw * (ppn < 2 ? 0.7 : 0.92 /*120.0/128.0*/), graphs[a]->nChannels*perChMaxRingLL128Bw);
projects/rccl/src/include/gdrwrap.h
+1 -1
View File
@@ -169,7 +169,7 @@ static gdr_t ncclGdrInit() {
return NULL;
}
GcnArchNameFormat(devProp.gcnArchName, gcnArchNameSubstr);
if (IsArchMatch(gcnArchNameSubstr, "gfx94")) {
if (IsArchMatch(gcnArchNameSubstr, "gfx942")) {
INFO(NCCL_INIT, "Enabled GDRCopy equivalent memory allocation on %s", gcnArchNameSubstr);
return (gdr_t)0x12345678L;
} else {
projects/rccl/src/include/rccl_float8.h
+16 -16
View File
@@ -344,7 +344,7 @@ struct rccl_float8
// default constructor
HIP_HOST_DEVICE rccl_float8() = default;
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
#if defined(__gfx942__) || defined(__gfx950__)
// device specific optimized F8 down-conversion code
template <bool stochastic_rounding = false>
@@ -381,10 +381,10 @@ struct rccl_float8
return i8data;
}
#endif // __gfx940__
#endif // __gfx942__
// constructor from float
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
#if defined(__gfx942__) || defined(__gfx950__)
// NOTE: ON-DEVICE... always optimal bias
explicit HIP_DEVICE rccl_float8(float v,
@@ -402,7 +402,7 @@ struct rccl_float8
// Host only implementation using s/w simulation
explicit HIP_HOST
#else
// both Host and DEVICE for non-gfx940 using s/w simulation
// both Host and DEVICE for non-gfx942 using s/w simulation
explicit HIP_HOST_DEVICE
#endif
rccl_float8(float v,
@@ -446,7 +446,7 @@ struct rccl_float8
}
// convert to float
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
#if defined(__gfx942__) || defined(__gfx950__)
// upcast using device specific intrinsic
explicit inline HIP_DEVICE operator float() const
{
@@ -460,7 +460,7 @@ struct rccl_float8
}
explicit inline HIP_HOST operator float() const
#else // non gfx940
#else // non gfx942
explicit inline HIP_HOST_DEVICE operator float() const
#endif
{
@@ -511,7 +511,7 @@ struct rccl_bfloat8
// default constructor
HIP_HOST_DEVICE rccl_bfloat8() = default;
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
#if defined(__gfx942__) || defined(__gfx950__)
// device specific optimized F8 down-conversion code
template <bool stochastic_rounding = false>
@@ -548,10 +548,10 @@ struct rccl_bfloat8
return i8data;
}
#endif // __gfx940__
#endif // __gfx942__
// constructor from float
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
#if defined(__gfx942__) || defined(__gfx950__)
// NOTE: ON-DEVICE... always optimal bias
explicit HIP_DEVICE rccl_bfloat8(float v,
@@ -569,7 +569,7 @@ struct rccl_bfloat8
// Host only implementation using s/w simulation
explicit HIP_HOST
#else
// both Host and DEVICE for non-gfx940 using s/w simulation
// both Host and DEVICE for non-gfx942 using s/w simulation
explicit HIP_HOST_DEVICE
#endif
rccl_bfloat8(float v,
@@ -613,7 +613,7 @@ struct rccl_bfloat8
}
// convert to float
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
#if defined(__gfx942__) || defined(__gfx950__)
// upcast using device specific intrinsic
explicit inline HIP_DEVICE operator float() const
{
@@ -627,7 +627,7 @@ struct rccl_bfloat8
}
explicit inline HIP_HOST operator float() const
#else // non gfx940
#else // non gfx942
explicit inline HIP_HOST_DEVICE operator float() const
#endif
{
@@ -969,7 +969,7 @@ inline __host__ __device__ T explicit_downcast(Ta a, uint32_t rng = 0)
return a;
}
// Use h/w intrinsic and optimized version when __gfx940__
// Use h/w intrinsic and optimized version when __gfx942__
template <
typename T,
typename Ta,
@@ -980,7 +980,7 @@ template <
= 0>
inline __host__ __device__ T explicit_downcast(Ta a, uint32_t rng)
{
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
#if defined(__gfx942__) || defined(__gfx950__)
// NOTE: we are directly calling cast_to_f8_from_f32 instead of constructor to optimize away one runtime branch
T val;
if(std::is_same<T, rccl_float8>::value)
@@ -988,12 +988,12 @@ inline __host__ __device__ T explicit_downcast(Ta a, uint32_t rng)
else
val.data = rccl_bfloat8::cast_to_bf8_from_f32<stochastic_rounding>(float(a), rng);
return val;
#else // non gfx940
#else // non gfx942
return T(float(a),
stochastic_rounding ? T::rocblas_hip_f8_rounding_mode::stochastic
: T::rocblas_hip_f8_rounding_mode::standard,
rng);
#endif // __gfx940__
#endif // __gfx942__
}
// NOTE NOTE: The above code is good if we don't consider HIP-GEMM code and only consider the quantization
projects/rccl/src/init.cc
+4 -4
View File
@@ -97,7 +97,7 @@ static uint64_t hashUniqueId(ncclUniqueId const &id) {
ncclResult_t commSetUnrollFactor(struct ncclComm* comm) {
hipDeviceProp_t devProp;
CUDACHECK(hipGetDeviceProperties(&devProp, comm->cudaDev));
if(IsArchMatch(devProp.gcnArchName, "gfx908") || ((IsArchMatch(devProp.gcnArchName, "gfx94") || IsArchMatch(devProp.gcnArchName, "gfx950"))
if(IsArchMatch(devProp.gcnArchName, "gfx908") || ((IsArchMatch(devProp.gcnArchName, "gfx942") || IsArchMatch(devProp.gcnArchName, "gfx950"))
&& devProp.multiProcessorCount > 80))
comm->unroll = NCCL_UNROLL_2;
else
@@ -1338,7 +1338,7 @@ static ncclResult_t initTransportsRank(struct ncclComm* comm, struct ncclComm* p
allGather3Data[rank].nc = std::max(allGather3Data[rank].nc, 4/ringGraph->nChannels);
if (ringGraph->nChannels > MAXCHANNELS/2)
allGather3Data[rank].nc = 1;
if (IsArchMatch(comm->topo->nodes[GPU].nodes[idx].gpu.gcn, "gfx94")) {
if (IsArchMatch(comm->topo->nodes[GPU].nodes[idx].gpu.gcn, "gfx942")) {
// Multi-node MI300A
int managed = 0;
CUDACHECK(hipDeviceGetAttribute(&managed, hipDeviceAttributeDirectManagedMemAccessFromHost, 0));
@@ -1873,7 +1873,7 @@ static ncclResult_t ncclCommInitRankFunc(struct ncclAsyncJob* job_) {
// a CUDA memory reconfig on load (c.f. NVSHMEM issue)
#ifdef USE_INDIRECT_FUNCTION_CALL
CUDACHECK(hipGetDeviceProperties(&devProp, 0));
if (ncclParamSetStackSize() == 1 && !IsArchMatch(devProp.gcnArchName,"gfx94") && !IsArchMatch(devProp.gcnArchName,"gfx950")) {
if (ncclParamSetStackSize() == 1 && !IsArchMatch(devProp.gcnArchName,"gfx942") && !IsArchMatch(devProp.gcnArchName,"gfx950")) {
stackSize = rcclParamStackSizeOverride() ? rcclParamStackSizeOverride() : maxLocalSizeBytes;
if (stackSize == 0) {
if (IsArchMatch(devProp.gcnArchName,"gfx906"))
@@ -1939,7 +1939,7 @@ static ncclResult_t ncclCommInitRankFunc(struct ncclAsyncJob* job_) {
if (mscclEnabled() && (comm->topo->mscclEnabled || mscclForceEnabled()) && mscclppCommCompatible(comm)) {
hipDeviceProp_t devProp;
CUDACHECK(hipGetDeviceProperties(&devProp, cudaDev));
comm->mscclppCompatible = IsArchMatch(devProp.gcnArchName, "gfx94") || IsArchMatch(devProp.gcnArchName, "gfx950");
comm->mscclppCompatible = IsArchMatch(devProp.gcnArchName, "gfx942") || IsArchMatch(devProp.gcnArchName, "gfx950");
if (comm->mscclppCompatible) {
bool mapContainsId = (mscclpp_uniqueIdMap.count(job->commId) > 0);
auto& mscclppUniqueId = mscclpp_uniqueIdMap[job->commId];
projects/rccl/src/misc/archinfo.cc
+1 -5
View File
@@ -42,10 +42,6 @@ void convertGcnArchToGcnArchName(const char* gcnArch, const char** gcnArchName)
*gcnArchName = "gfx908";
else if (strcmp(gcnArch, "910") == 0)
*gcnArchName = "gfx90a";
else if (strcmp(gcnArch, "940") == 0)
*gcnArchName = "gfx940";
else if (strcmp(gcnArch, "941") == 0)
*gcnArchName = "gfx941";
else if (strcmp(gcnArch, "942") == 0)
*gcnArchName = "gfx942";
else if (strcmp(gcnArch, "950") == 0)
@@ -65,7 +61,7 @@ int GetGcnArchName(int deviceId, char* out) {
double GetDeviceWallClockRateInKhz(int deviceId) {
char gcn[256];
GetGcnArchName(deviceId, gcn);
if (strncmp("gfx94", gcn, 5) == 0)
if (strncmp("gfx942", gcn, 6) == 0)
return 1.0E5;
else if(strncmp("gfx950", gcn, 6) == 0)
return 1.0E5;
projects/rccl/src/transport/net.cc
+1 -1
View File
@@ -207,7 +207,7 @@ static ncclResult_t sendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph
if (req.netDev < 0) NCCLCHECK(ncclTopoGetNetDev(comm, myInfo->rank, graph, channelId, peerInfo->rank, &netId, &req.netDev, &proxyRank));
NCCLCHECK(ncclTopoCheckGdr(comm->topo, myInfo->busId, netId, 1, &req.useGdr));
send->conn.flags |= req.useGdr ? NCCL_DIRECT_NIC : 0;
if (req.useGdr && !IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx90a") && !IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94") && !IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) {
if (req.useGdr && !IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx90a") && !IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx942") && !IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) {
CUDACHECK(hipDeviceGetAttribute((int*)&req.curr_hdp_reg, hipDeviceAttributeHdpMemFlushCntl, myInfo->cudaDev));
send->conn.curr_hdp_reg = req.curr_hdp_reg;
}
projects/rccl/src/transport/p2p.cc
+1 -1
View File
@@ -368,7 +368,7 @@ ncclResult_t p2pSendSetup(struct ncclComm* comm, struct ncclTopoGraph* graph, st
INFO(NCCL_INIT|NCCL_P2P,"Ring %02d : %d -> %d failed to get link type and hop count", channelId, myInfo->rank, peerInfo->rank);
return ncclInternalError;
}
if (!isXGMI && !IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx90a") && !IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94") && !IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) {
if (!isXGMI && !IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx90a") && !IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx942") && !IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx950")) {
CUDACHECK(hipDeviceGetAttribute((int*)&resources->next_hdp_reg, hipDeviceAttributeHdpMemFlushCntl,peerInfo->cudaDev));
TRACE(NCCL_INIT|NCCL_P2P,"Ring %02d : %d -> %d HDP %p", channelId, myInfo->rank, peerInfo->rank, resources->next_hdp_reg);
}
projects/rccl/tools/JitterBench/Common.hpp
+1 -1
View File
@@ -34,7 +34,7 @@ THE SOFTWARE.
} while (0)
// Macro for collecting HW_REG_XCC_ID
#if defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__)
#if defined(__gfx942__) || defined(__gfx950__)
#define GetXccId(val) \
asm volatile ("s_getreg_b32 %0, hwreg(HW_REG_XCC_ID)" : "=s" (val));
#else
projects/rccl/tools/p2p-latency-test/ll_latency_test.cpp
+4 -4
View File
@@ -147,7 +147,7 @@ int main(int argc, char** argv) {
HIPCHECK(hipStreamCreateWithFlags(&stream[0], hipStreamNonBlocking));
HIPCHECK(hipDeviceEnablePeerAccess(device_id[1], 0));
HIPCHECK(hipGetDeviceProperties(&prop[0], device_id[0]));
HIPCHECK(hipExtMallocWithFlags((void**)&flag[0], HIP_IPC_MEM_MIN_SIZE, strncmp(prop[0].gcnArchName, "gfx94", 5) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipExtMallocWithFlags((void**)&flag[0], HIP_IPC_MEM_MIN_SIZE, strncmp(prop[0].gcnArchName, "gfx942", 6) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipHostMalloc ((void**)&time_delta[0], sizeof(uint64_t), hipHostMallocDefault));
HIPCHECK(hipMalloc((void**)&abortFlag[0], sizeof(uint32_t)));
HIPCHECK(hipMemsetAsync(flag[0], 0, HIP_IPC_MEM_MIN_SIZE, stream[0]));
@@ -158,7 +158,7 @@ int main(int argc, char** argv) {
HIPCHECK(hipStreamCreateWithFlags(&stream[1], hipStreamNonBlocking));
HIPCHECK(hipDeviceEnablePeerAccess(device_id[0], 0));
HIPCHECK(hipGetDeviceProperties(&prop[1], device_id[1]));
HIPCHECK(hipExtMallocWithFlags((void**)&flag[1], HIP_IPC_MEM_MIN_SIZE, strncmp(prop[1].gcnArchName, "gfx94", 5) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipExtMallocWithFlags((void**)&flag[1], HIP_IPC_MEM_MIN_SIZE, strncmp(prop[1].gcnArchName, "gfx942", 6) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipHostMalloc((void**)&time_delta[1], sizeof(uint64_t), hipHostMallocDefault));
HIPCHECK(hipMalloc((void**)&abortFlag[1], sizeof(uint32_t)));
HIPCHECK(hipMemsetAsync(flag[1], 0, HIP_IPC_MEM_MIN_SIZE, stream[1]));
@@ -174,11 +174,11 @@ int main(int argc, char** argv) {
double vega_gpu_rtc_freq;
HIPCHECK(hipStreamSynchronize(stream[0]));
vega_gpu_rtc_freq = strncmp(prop[0].gcnArchName, "gfx94", 5) == 0 ? 1.0E8 : 2.5E7;
vega_gpu_rtc_freq = strncmp(prop[0].gcnArchName, "gfx942", 6) == 0 ? 1.0E8 : 2.5E7;
fprintf(stdout, "One-way latency in us: %g\n", double(*time_delta[0]) * 1e6 / NUM_LOOPS_RUN / vega_gpu_rtc_freq / 2);
HIPCHECK(hipStreamSynchronize(stream[1]));
vega_gpu_rtc_freq = strncmp(prop[1].gcnArchName, "gfx94", 5) == 0 ? 1.0E8 : 2.5E7;
vega_gpu_rtc_freq = strncmp(prop[1].gcnArchName, "gfx942", 6) == 0 ? 1.0E8 : 2.5E7;
fprintf(stdout, "One-way latency in us: %g\n", double(*time_delta[1]) * 1e6 / NUM_LOOPS_RUN / vega_gpu_rtc_freq / 2);
HIPCHECK(hipFree(flag[0]));
projects/rccl/tools/p2p-latency-test/p2p_latency_test.cpp
+4 -4
View File
@@ -86,7 +86,7 @@ int main(int argc, char** argv) {
HIPCHECK(hipStreamCreateWithFlags(&stream[0], hipStreamNonBlocking));
HIPCHECK(hipDeviceEnablePeerAccess(device_id[1], 0));
HIPCHECK(hipGetDeviceProperties(&prop[0], device_id[0]));
HIPCHECK(hipExtMallocWithFlags((void**)&flag[0], HIP_IPC_MEM_MIN_SIZE, strncmp(prop[0].gcnArchName, "gfx94", 5) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipExtMallocWithFlags((void**)&flag[0], HIP_IPC_MEM_MIN_SIZE, strncmp(prop[0].gcnArchName, "gfx942", 6) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipMalloc((void**)&time_delta[0], HIP_IPC_MEM_MIN_SIZE));
HIPCHECK(hipMemsetAsync(flag[0], 0, HIP_IPC_MEM_MIN_SIZE, stream[0]));
HIPCHECK(hipStreamSynchronize(stream[0]));
@@ -95,7 +95,7 @@ int main(int argc, char** argv) {
HIPCHECK(hipStreamCreateWithFlags(&stream[1], hipStreamNonBlocking));
HIPCHECK(hipDeviceEnablePeerAccess(device_id[0], 0));
HIPCHECK(hipGetDeviceProperties(&prop[1], device_id[1]));
HIPCHECK(hipExtMallocWithFlags((void**)&flag[1], HIP_IPC_MEM_MIN_SIZE, strncmp(prop[1].gcnArchName, "gfx94", 5) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipExtMallocWithFlags((void**)&flag[1], HIP_IPC_MEM_MIN_SIZE, strncmp(prop[1].gcnArchName, "gfx942", 6) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipMalloc((void**)&time_delta[1], HIP_IPC_MEM_MIN_SIZE));
HIPCHECK(hipMemsetAsync(flag[1], 0, HIP_IPC_MEM_MIN_SIZE, stream[1]));
HIPCHECK(hipStreamSynchronize(stream[1]));
@@ -109,11 +109,11 @@ int main(int argc, char** argv) {
double vega_gpu_rtc_freq;
HIPCHECK(hipStreamSynchronize(stream[0]));
vega_gpu_rtc_freq = strncmp(prop[0].gcnArchName, "gfx94", 5) == 0 ? 1.0E8 : 2.5E7;
vega_gpu_rtc_freq = strncmp(prop[0].gcnArchName, "gfx942", 6) == 0 ? 1.0E8 : 2.5E7;
fprintf(stdout, "One-way latency in us: %g\n", double(*time_delta[0]) * 1e6 / NUM_LOOPS_RUN / vega_gpu_rtc_freq / 2);
HIPCHECK(hipStreamSynchronize(stream[1]));
vega_gpu_rtc_freq = strncmp(prop[1].gcnArchName, "gfx94", 5) == 0 ? 1.0E8 : 2.5E7;
vega_gpu_rtc_freq = strncmp(prop[1].gcnArchName, "gfx942", 6) == 0 ? 1.0E8 : 2.5E7;
fprintf(stdout, "One-way latency in us: %g\n", double(*time_delta[1]) * 1e6 / NUM_LOOPS_RUN / vega_gpu_rtc_freq / 2);
HIPCHECK(hipFree(flag[0]));
projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp
+7 -7
View File
@@ -426,7 +426,7 @@ int main(int argc,char* argv[])
static const char *ring_4p3l = "0 1 2 3|0 1 3 2|0 2 1 3|0 2 3 1|0 3 1 2|0 3 2 1";
static const char *ring_8p1h = "0 1 3 2 4 5 7 6|6 7 5 4 2 3 1 0|0 1 5 4 6 7 3 2|2 3 7 6 4 5 1 0";
static const char *ring_16p1h = "0 1 3 2 6 7 15 14 10 11 9 8 12 13 5 4|0 1 2 3 7 6 13 12 8 9 10 11 15 14 5 4|0 2 3 7 6 14 15 11 10 8 9 13 12 4 5 1|4 5 13 12 8 9 11 10 14 15 7 6 2 3 1 0|4 5 14 15 11 10 9 8 12 13 6 7 3 2 1 0|1 5 4 12 13 9 8 10 11 15 14 6 7 3 2 0";
static const char *ring_gfx940_8p = "0 1 2 3 4 5 6 7|0 1 2 3 4 5 7 6|0 2 4 1 3 6 5 7|0 2 4 6 1 7 3 5|0 3 1 5 2 7 4 6|0 3 5 1 6 2 7 4|0 4 1 7 3 6 2 5|7 6 5 4 3 2 1 0|6 7 5 4 3 2 1 0|7 5 6 3 1 4 2 0|5 3 7 1 6 4 2 0|6 4 7 2 5 1 3 0|4 7 2 6 1 5 3 0|5 2 6 3 7 1 4 0";
static const char *ring_gfx942_8p = "0 1 2 3 4 5 6 7|0 1 2 3 4 5 7 6|0 2 4 1 3 6 5 7|0 2 4 6 1 7 3 5|0 3 1 5 2 7 4 6|0 3 5 1 6 2 7 4|0 4 1 7 3 6 2 5|7 6 5 4 3 2 1 0|6 7 5 4 3 2 1 0|7 5 6 3 1 4 2 0|5 3 7 1 6 4 2 0|6 4 7 2 5 1 3 0|4 7 2 6 1 5 3 0|5 2 6 3 7 1 4 0";
setupPeers(connection_info, &is_xgmi);
if (!r) {
@@ -436,8 +436,8 @@ int main(int argc,char* argv[])
if (nGpu == 8 && !cr8g) {
hipDeviceProp_t prop;
HIPCHECK(hipGetDeviceProperties(&prop, 0));
if (strncmp(prop.gcnArchName, "gfx94", 5) == 0) {
r = (char *)ring_gfx940_8p;
if (strncmp(prop.gcnArchName, "gfx942", 6) == 0) {
r = (char *)ring_gfx942_8p;
if(!workgroups) workgroups = 28;
} else {
r = (char *)ring_8p1h;
@@ -521,11 +521,11 @@ int main(int argc,char* argv[])
profiling_data[i] = (struct profiling_data_t *)malloc(sizeof(struct profiling_data_t)*iters);
HIPCHECK(hipMalloc((void**) &d_profiling_data[i], sizeof(struct profiling_data_t)*iters));
HIPCHECK(hipExtMallocWithFlags((void**) &transfer_data[i], sizeof(struct transfer_data_t), strncmp(prop.gcnArchName, "gfx94", 5) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipExtMallocWithFlags((void**) &transfer_data[i], sizeof(struct transfer_data_t), strncmp(prop.gcnArchName, "gfx942", 6) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
for (int j = 0; j < workgroups; j++) {
HIPCHECK(hipExtMallocWithFlags((void**) &buff[i*MAX_WORKGROUPS+j], 2*N*sizeof(float), strncmp(prop.gcnArchName, "gfx94", 5) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipExtMallocWithFlags((void**) &buff[i*MAX_WORKGROUPS+j], 2*N*sizeof(float), strncmp(prop.gcnArchName, "gfx942", 6) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
// additional fine grained buffer for local doublecopy, only need 1 buffer (not used by remote)
HIPCHECK(hipExtMallocWithFlags((void**) &buff_fine[i*MAX_WORKGROUPS+j], N*sizeof(float), strncmp(prop.gcnArchName, "gfx94", 5) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipExtMallocWithFlags((void**) &buff_fine[i*MAX_WORKGROUPS+j], N*sizeof(float), strncmp(prop.gcnArchName, "gfx942", 6) == 0 ? hipDeviceMallocUncached : hipDeviceMallocFinegrained));
HIPCHECK(hipMalloc((void**) &buff_coarse[i*MAX_WORKGROUPS+j], 2*N*sizeof(float)));
//randomize test data
hipLaunchKernelGGL(initTestDataKernel,
@@ -670,7 +670,7 @@ int main(int argc,char* argv[])
hipDeviceProp_t prop;
HIPCHECK(hipGetDeviceProperties(&prop, i));
double vega_gpu_rtc_freq, bw_std_dev = 0, mean_write_cycle = 0;
if (strncmp(prop.gcnArchName, "gfx94", 5) == 0)
if (strncmp(prop.gcnArchName, "gfx942", 6) == 0)
vega_gpu_rtc_freq = 1.0E8;
else
vega_gpu_rtc_freq = 2.5E7;
projects/rccl/tools/topo_expl/models/topo_4p_940.xml → projects/rccl/tools/topo_expl/models/topo_4p_942.xml
+4 -4
View File
@@ -1,7 +1,7 @@
<system version="2">
<cpu numaid="2" affinity="000000ff,ffff0000,00000000,000000ff,ffff0000,00000000" arch="x86_64" vendor="AuthenticAMD" familyid="175" modelid="144">
<pci busid="0000:82:00.0" class="0x120000" vendor="0x1002" device="0x74a0" subsystem_vendor="0x1002" subsystem_device="0x74a0" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="0" sm="228" gcn="940" arch="38911" rank="0" gdr="1">
<gpu dev="0" sm="228" gcn="942" arch="38911" rank="0" gdr="1">
<xgmi target="0000:c2:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:02:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:42:00.0" count="1" tclass="0x120000"/>
@@ -15,7 +15,7 @@
</cpu>
<cpu numaid="3" affinity="ffffff00,00000000,00000000,ffffff00,00000000,00000000" arch="x86_64" vendor="AuthenticAMD" familyid="175" modelid="144">
<pci busid="0000:c2:00.0" class="0x120000" vendor="0x1002" device="0x74a0" subsystem_vendor="0x1002" subsystem_device="0x74a0" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="1" sm="228" gcn="940" arch="38911" rank="1" gdr="1">
<gpu dev="1" sm="228" gcn="942" arch="38911" rank="1" gdr="1">
<xgmi target="0000:82:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:02:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:42:00.0" count="1" tclass="0x120000"/>
@@ -29,7 +29,7 @@
</cpu>
<cpu numaid="0" affinity="00000000,00000000,00ffffff,00000000,00000000,00ffffff" arch="x86_64" vendor="AuthenticAMD" familyid="175" modelid="144">
<pci busid="0000:02:00.0" class="0x120000" vendor="0x1002" device="0x74a0" subsystem_vendor="0x1002" subsystem_device="0x74a0" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="2" sm="228" gcn="940" arch="38911" rank="2" gdr="1">
<gpu dev="2" sm="228" gcn="942" arch="38911" rank="2" gdr="1">
<xgmi target="0000:82:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:c2:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:42:00.0" count="1" tclass="0x120000"/>
@@ -43,7 +43,7 @@
</cpu>
<cpu numaid="1" affinity="00000000,0000ffff,ff000000,00000000,0000ffff,ff000000" arch="x86_64" vendor="AuthenticAMD" familyid="175" modelid="144">
<pci busid="0000:42:00.0" class="0x120000" vendor="0x1002" device="0x74a0" subsystem_vendor="0x1002" subsystem_device="0x74a0" link_speed="32.0 GT/s PCIe" link_width="16">
<gpu dev="3" sm="228" gcn="940" arch="38911" rank="3" gdr="1">
<gpu dev="3" sm="228" gcn="942" arch="38911" rank="3" gdr="1">
<xgmi target="0000:82:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:c2:00.0" count="1" tclass="0x120000"/>
<xgmi target="0000:02:00.0" count="1" tclass="0x120000"/>
projects/rccl/tools/topo_expl/models/topo_8p_940.xml → projects/rccl/tools/topo_expl/models/topo_8p_942.xml
View File
projects/rccl/tools/topo_expl/models/topo_8p_940vm.xml → projects/rccl/tools/topo_expl/models/topo_8p_942vm.xml
View File
projects/rccl/tools/topo_expl/topo_expl.cpp
+3 -4
View File
@@ -129,11 +129,10 @@ NodeModelDesc model_descs[] = {
{"topo_8p1h_5.xml", " 8gfx910 2H3XGMI 8NIC 2AMD B"},
{"topo_16p1h.xml", "16gfx910 2H3XGMI 8NIC 4AMD A"},
{"topo_16p1h_vm.xml", "16gfx910 2H3XGMI 8NIC 4AMD B"},
// GFX 940
{"topo_4p_940.xml", " 4gfx940 1H3XGMI 4NIC 4AMD2 A"},
// GFX 942
{"topo_8p_940.xml", " 8gfx942 1H7XGMI 8NIC 2Intel A"},
{"topo_8p_940vm.xml", " 8gfx942 1H7XGMI 8NIC 2Intel B"},
{"topo_4p_942.xml", " 4gfx942 1H3XGMI 4NIC 4AMD2 A"},
{"topo_8p_942.xml", " 8gfx942 1H7XGMI 8NIC 2Intel A"},
{"topo_8p_942vm.xml", " 8gfx942 1H7XGMI 8NIC 2Intel B"},
{"topo_16p_gio-1s-1rp-cascade.xml", "16gfx942 2H7XGMI 1NIC 2AMD A"},
{"topo_16p_gio-3s-1rp-split-flat.xml", "16gfx942 2H7XGMI 1NIC 2AMD B"},
};