From e95578ef4c40492dcb66c4325cc97f5b0448950f Mon Sep 17 00:00:00 2001 From: corey-derochie-amd <161367113+corey-derochie-amd@users.noreply.github.com> Date: Thu, 20 Mar 2025 09:34:53 -0600 Subject: [PATCH] removed gfx940 and gfx941 (#1606) * 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 [ROCm/rccl commit: 6505639cf488ddf65d3b086c34ae98b776d34c3e] --- .../rccl/cmake/scripts/extract_metadata.cmake | 2 +- projects/rccl/src/device/all_gather.h | 2 +- projects/rccl/src/device/all_reduce.h | 6 ++-- projects/rccl/src/device/alltoall_pivot.h | 2 +- projects/rccl/src/device/broadcast.h | 2 +- projects/rccl/src/device/msccl_kernel_impl.h | 2 +- projects/rccl/src/device/onerank.cu | 2 +- projects/rccl/src/device/op128.h | 2 +- projects/rccl/src/device/primitives.h | 2 +- projects/rccl/src/device/reduce.h | 2 +- projects/rccl/src/device/reduce_scatter.h | 4 +-- projects/rccl/src/device/sendrecv.h | 10 +++--- projects/rccl/src/graph/connect.cc | 4 +-- projects/rccl/src/graph/paths.cc | 6 ++-- projects/rccl/src/graph/topo.h | 2 +- projects/rccl/src/graph/tuning.cc | 4 +-- projects/rccl/src/include/gdrwrap.h | 2 +- projects/rccl/src/include/rccl_float8.h | 32 +++++++++---------- projects/rccl/src/init.cc | 8 ++--- projects/rccl/src/misc/archinfo.cc | 6 +--- projects/rccl/src/transport/net.cc | 2 +- projects/rccl/src/transport/p2p.cc | 2 +- projects/rccl/tools/JitterBench/Common.hpp | 2 +- .../p2p-latency-test/ll_latency_test.cpp | 8 ++--- .../p2p-latency-test/p2p_latency_test.cpp | 8 ++--- .../tools/rccl-prim-test/rccl_prim_test.cpp | 14 ++++---- .../{topo_4p_940.xml => topo_4p_942.xml} | 8 ++--- .../{topo_8p_940.xml => topo_8p_942.xml} | 0 .../{topo_8p_940vm.xml => topo_8p_942vm.xml} | 0 projects/rccl/tools/topo_expl/topo_expl.cpp | 7 ++-- 30 files changed, 74 insertions(+), 79 deletions(-) rename projects/rccl/tools/topo_expl/models/{topo_4p_940.xml => topo_4p_942.xml} (93%) rename projects/rccl/tools/topo_expl/models/{topo_8p_940.xml => topo_8p_942.xml} (100%) rename projects/rccl/tools/topo_expl/models/{topo_8p_940vm.xml => topo_8p_942vm.xml} (100%) diff --git a/projects/rccl/cmake/scripts/extract_metadata.cmake b/projects/rccl/cmake/scripts/extract_metadata.cmake index 46dba319e1..9e0e7d10d1 100644 --- a/projects/rccl/cmake/scripts/extract_metadata.cmake +++ b/projects/rccl/cmake/scripts/extract_metadata.cmake @@ -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}) diff --git a/projects/rccl/src/device/all_gather.h b/projects/rccl/src/device/all_gather.h index 88700ada79..8468948610 100644 --- a/projects/rccl/src/device/all_gather.h +++ b/projects/rccl/src/device/all_gather.h @@ -11,7 +11,7 @@ namespace { template -#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) { diff --git a/projects/rccl/src/device/all_reduce.h b/projects/rccl/src/device/all_reduce.h index cabf12ea74..6a394fb2bd 100644 --- a/projects/rccl/src/device/all_reduce.h +++ b/projects/rccl/src/device/all_reduce.h @@ -15,7 +15,7 @@ namespace { template -#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 -#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 -#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) { diff --git a/projects/rccl/src/device/alltoall_pivot.h b/projects/rccl/src/device/alltoall_pivot.h index 6b5a6392ff..d0b12cf50d 100644 --- a/projects/rccl/src/device/alltoall_pivot.h +++ b/projects/rccl/src/device/alltoall_pivot.h @@ -10,7 +10,7 @@ namespace { template -#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) { diff --git a/projects/rccl/src/device/broadcast.h b/projects/rccl/src/device/broadcast.h index c53f1553c8..50150e5b79 100644 --- a/projects/rccl/src/device/broadcast.h +++ b/projects/rccl/src/device/broadcast.h @@ -10,7 +10,7 @@ namespace { template -#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) { diff --git a/projects/rccl/src/device/msccl_kernel_impl.h b/projects/rccl/src/device/msccl_kernel_impl.h index 75ae68890f..1daf10e37d 100644 --- a/projects/rccl/src/device/msccl_kernel_impl.h +++ b/projects/rccl/src/device/msccl_kernel_impl.h @@ -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 } diff --git a/projects/rccl/src/device/onerank.cu b/projects/rccl/src/device/onerank.cu index 809e6af9ca..e767f0a69a 100644 --- a/projects/rccl/src/device/onerank.cu +++ b/projects/rccl/src/device/onerank.cu @@ -11,7 +11,7 @@ #include "common.h" #include -#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 diff --git a/projects/rccl/src/device/op128.h b/projects/rccl/src/device/op128.h index 99e4fa268e..9311edb33a 100644 --- a/projects/rccl/src/device/op128.h +++ b/projects/rccl/src/device/op128.h @@ -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]; diff --git a/projects/rccl/src/device/primitives.h b/projects/rccl/src/device/primitives.h index c49c28d582..f57d95244a 100644 --- a/projects/rccl/src/device/primitives.h +++ b/projects/rccl/src/device/primitives.h @@ -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() diff --git a/projects/rccl/src/device/reduce.h b/projects/rccl/src/device/reduce.h index 689e4261f7..0fa6395983 100644 --- a/projects/rccl/src/device/reduce.h +++ b/projects/rccl/src/device/reduce.h @@ -11,7 +11,7 @@ namespace { template -#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) { diff --git a/projects/rccl/src/device/reduce_scatter.h b/projects/rccl/src/device/reduce_scatter.h index e618b8c3b2..b23975ef57 100644 --- a/projects/rccl/src/device/reduce_scatter.h +++ b/projects/rccl/src/device/reduce_scatter.h @@ -11,7 +11,7 @@ namespace { template -#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 RunWorkCollregUsed ? (NCCL_MAX_NTHREADS - WARP_SIZE) : (nranks <= 6 ? 7 * WARP_SIZE : 5 * WARP_SIZE); diff --git a/projects/rccl/src/device/sendrecv.h b/projects/rccl/src/device/sendrecv.h index 91ed574e02..e45e959812 100644 --- a/projects/rccl/src/device/sendrecv.h +++ b/projects/rccl/src/device/sendrecv.h @@ -20,7 +20,7 @@ struct RunWorkBatchsendBytes; 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 (subtid, subtn, 0, nullptr, false, 1, &work->sendAddr, 1, &work->recvAddr, (ssize_t)work->sendBytes); #else @@ -254,7 +254,7 @@ struct RunWorkBatch>(subtid, subtn, group, work); -#elif defined(__gfx908__) || defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__) +#elif defined(__gfx908__) || defined(__gfx942__) || defined(__gfx950__) runSend>(subtid, subtn, group, work); #else runSend>(subtid, subtn, group, work); @@ -266,7 +266,7 @@ struct RunWorkBatch>(subtid, subtn, group, work); -#elif defined(__gfx908__) || defined(__gfx940__) || defined(__gfx941__) || defined(__gfx942__) || defined(__gfx950__) +#elif defined(__gfx908__) || defined(__gfx942__) || defined(__gfx950__) runRecv>(subtid, subtn, group, work); #else runRecv>(subtid, subtn, group, work); diff --git a/projects/rccl/src/graph/connect.cc b/projects/rccl/src/graph/connect.cc index 6c6a743542..13caa378bd 100644 --- a/projects/rccl/src/graph/connect.cc +++ b/projects/rccl/src/graph/connect.cc @@ -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; diff --git a/projects/rccl/src/graph/paths.cc b/projects/rccl/src/graph/paths.cc index a09df907dd..dca35eba4a 100644 --- a/projects/rccl/src/graph/paths.cc +++ b/projects/rccl/src/graph/paths.cc @@ -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); diff --git a/projects/rccl/src/graph/topo.h b/projects/rccl/src/graph/topo.h index 47713db7ce..17d1906472 100644 --- a/projects/rccl/src/graph/topo.h +++ b/projects/rccl/src/graph/topo.h @@ -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; diff --git a/projects/rccl/src/graph/tuning.cc b/projects/rccl/src/graph/tuning.cc index 3c90705fa5..494e683f6c 100644 --- a/projects/rccl/src/graph/tuning.cc +++ b/projects/rccl/src/graph/tuning.cc @@ -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; ptopo->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); diff --git a/projects/rccl/src/include/gdrwrap.h b/projects/rccl/src/include/gdrwrap.h index d305690856..281d30f9f9 100644 --- a/projects/rccl/src/include/gdrwrap.h +++ b/projects/rccl/src/include/gdrwrap.h @@ -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 { diff --git a/projects/rccl/src/include/rccl_float8.h b/projects/rccl/src/include/rccl_float8.h index 944d781a11..19304c74b0 100644 --- a/projects/rccl/src/include/rccl_float8.h +++ b/projects/rccl/src/include/rccl_float8.h @@ -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 @@ -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 @@ -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::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(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 diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc index 3db87c9296..d7facdc439 100644 --- a/projects/rccl/src/init.cc +++ b/projects/rccl/src/init.cc @@ -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]; diff --git a/projects/rccl/src/misc/archinfo.cc b/projects/rccl/src/misc/archinfo.cc index 547bcccd92..e9ee4dc4b1 100644 --- a/projects/rccl/src/misc/archinfo.cc +++ b/projects/rccl/src/misc/archinfo.cc @@ -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; diff --git a/projects/rccl/src/transport/net.cc b/projects/rccl/src/transport/net.cc index 5fd36f6f71..8018c55652 100644 --- a/projects/rccl/src/transport/net.cc +++ b/projects/rccl/src/transport/net.cc @@ -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; } diff --git a/projects/rccl/src/transport/p2p.cc b/projects/rccl/src/transport/p2p.cc index 5c3871163a..240ac59f0e 100644 --- a/projects/rccl/src/transport/p2p.cc +++ b/projects/rccl/src/transport/p2p.cc @@ -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); } diff --git a/projects/rccl/tools/JitterBench/Common.hpp b/projects/rccl/tools/JitterBench/Common.hpp index 16aa82c4ce..bad12a1b6f 100644 --- a/projects/rccl/tools/JitterBench/Common.hpp +++ b/projects/rccl/tools/JitterBench/Common.hpp @@ -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 diff --git a/projects/rccl/tools/p2p-latency-test/ll_latency_test.cpp b/projects/rccl/tools/p2p-latency-test/ll_latency_test.cpp index 8a20e55615..f049c9db73 100644 --- a/projects/rccl/tools/p2p-latency-test/ll_latency_test.cpp +++ b/projects/rccl/tools/p2p-latency-test/ll_latency_test.cpp @@ -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])); diff --git a/projects/rccl/tools/p2p-latency-test/p2p_latency_test.cpp b/projects/rccl/tools/p2p-latency-test/p2p_latency_test.cpp index 0ed3b5e982..da98c8b2bc 100644 --- a/projects/rccl/tools/p2p-latency-test/p2p_latency_test.cpp +++ b/projects/rccl/tools/p2p-latency-test/p2p_latency_test.cpp @@ -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])); diff --git a/projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp b/projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp index 7f3826199c..633eef2dc9 100644 --- a/projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp +++ b/projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp @@ -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; diff --git a/projects/rccl/tools/topo_expl/models/topo_4p_940.xml b/projects/rccl/tools/topo_expl/models/topo_4p_942.xml similarity index 93% rename from projects/rccl/tools/topo_expl/models/topo_4p_940.xml rename to projects/rccl/tools/topo_expl/models/topo_4p_942.xml index 03665513c5..b0eefe4747 100644 --- a/projects/rccl/tools/topo_expl/models/topo_4p_940.xml +++ b/projects/rccl/tools/topo_expl/models/topo_4p_942.xml @@ -1,7 +1,7 @@ - + @@ -15,7 +15,7 @@ - + @@ -29,7 +29,7 @@ - + @@ -43,7 +43,7 @@ - + diff --git a/projects/rccl/tools/topo_expl/models/topo_8p_940.xml b/projects/rccl/tools/topo_expl/models/topo_8p_942.xml similarity index 100% rename from projects/rccl/tools/topo_expl/models/topo_8p_940.xml rename to projects/rccl/tools/topo_expl/models/topo_8p_942.xml diff --git a/projects/rccl/tools/topo_expl/models/topo_8p_940vm.xml b/projects/rccl/tools/topo_expl/models/topo_8p_942vm.xml similarity index 100% rename from projects/rccl/tools/topo_expl/models/topo_8p_940vm.xml rename to projects/rccl/tools/topo_expl/models/topo_8p_942vm.xml diff --git a/projects/rccl/tools/topo_expl/topo_expl.cpp b/projects/rccl/tools/topo_expl/topo_expl.cpp index dde0de644a..d57388136c 100644 --- a/projects/rccl/tools/topo_expl/topo_expl.cpp +++ b/projects/rccl/tools/topo_expl/topo_expl.cpp @@ -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"}, };