AI/rocm-systems
2
0
Garpu 0
Kode Masalah Tarik Permintaan Actions Packages Projects Rilis Wiki Kegiatan

Fix gfx940 CPX mode (#1290)

Telusuri Sumber
This commit is contained in:
Wenkai Du
2024-08-16 08:46:06 +08:00
committed by GitHub
orang tua eff56735b0
melakukan d3171b51b7
4 mengubah file dengan 10 tambahan dan 5 penghapusan
Download Patch File Download Diff File Expand all files Collapse all files
src/graph/connect.cc
+2 -1
Melihat File
@@ -626,7 +626,8 @@ ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, int* treePa
// Only use full MAXCHANNELS for gfx94x
int maxChannels = IsArchMatch(comm->topo->nodes[GPU].nodes[0].gpu.gcn, "gfx94") ?
(comm->topo->nodes[GPU].nodes[0].gpu.cu == 80 ? 80 : MAXCHANNELS) : 2*CHANNEL_LIMIT;
((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 : MAXCHANNELS) : 2*CHANNEL_LIMIT;
if (graphs[NCCL_ALGO_RING]->nIntraChannels > 0 || comm->nNodes > 1) {
maxChannels = std::min(64, maxChannels);
src/graph/xml.cc
+3 -1
Melihat File
@@ -347,7 +347,9 @@ static void memcpylower(char* dst, const char* src, const size_t size) {
static ncclResult_t getPciPath(const char* busId, char** path) {
char busPath[] = "/sys/class/pci_bus/0000:00/../../0000:00:00.0";
memcpylower(busPath+sizeof("/sys/class/pci_bus/")-1, busId, BUSID_REDUCED_SIZE-1);
memcpylower(busPath+sizeof("/sys/class/pci_bus/0000:00/../../")-1, busId, BUSID_SIZE-1);
memcpylower(busPath+sizeof("/sys/class/pci_bus/0000:00/../../")-1, busId, BUSID_SIZE-2);
// override PCIe device function ID in CPX mode
busPath[sizeof("/sys/class/pci_bus/0000:00/../../")+BUSID_SIZE-3] = '0';
*path = realpath(busPath, NULL);
if (*path == NULL) {
WARN("Could not find real path of %s", busPath);
src/graph/xml.h
+1 -1
Melihat File
@@ -18,7 +18,7 @@
#define MAX_STR_LEN 255
#define MAX_ATTR_COUNT 16
#define MAX_SUBS 64 //Changed the value from 32 to 64 for CPX mode
#define MAX_NODES 1024
#define MAX_NODES 4096 //Changed the value from 1024 to 4096 for CPX mode
#define NODE_TYPE_NONE 0
#define NODE_TYPE_OPEN 1
src/misc/rocm_smi_wrap.cc
+4 -2
Melihat File
@@ -172,7 +172,8 @@ ncclResult_t rocm_smi_getLinkInfo(int srcIndex, int dstIndex, RSMI_IO_LINK_TYPE*
if (rcclParamUseRocmSmiLib()) {
ROCMSMICHECK(rsmi_topo_get_link_type(srcIndex, dstIndex, &rsmi_hops, rsmi_type));
ROCMSMICHECK(rsmi_topo_get_link_weight(srcIndex, dstIndex, &rsmi_weight));
if (*rsmi_type == RSMI_IOLINK_TYPE_XGMI && rsmi_weight == 15) {
if (*rsmi_type == RSMI_IOLINK_TYPE_XGMI && (rsmi_weight == 15 ||
rsmi_weight == 41 || rsmi_weight == 13)) {
uint64_t min_bw = 0, max_bw = 0;
*hops = 1;
#if defined USE_ROCM_SMI64CONFIG && rocm_smi_VERSION_MAJOR >= 5
@@ -189,7 +190,8 @@ ncclResult_t rocm_smi_getLinkInfo(int srcIndex, int dstIndex, RSMI_IO_LINK_TYPE*
ARSMICHECK(ARSMI_topo_get_link_info(srcIndex, dstIndex, &tinfo));
*rsmi_type = (RSMI_IO_LINK_TYPE) tinfo.type;
if (*rsmi_type == RSMI_IOLINK_TYPE_XGMI && tinfo.weight == 15) {
if (*rsmi_type == RSMI_IOLINK_TYPE_XGMI && (tinfo.weight == 15 ||
tinfo.weight == 41 || tinfo.weight == 13)) {
*hops = 1;
if (tinfo.max_bandwidth && tinfo.min_bandwidth)
*count = tinfo.max_bandwidth/tinfo.min_bandwidth;