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

Compiler warnings fix 2 (#1801)

Browse Source 
* Changes to device code

* Changes to src/misc

* Changes to graph

* src/include changes

* src/transport changes

* changes in init, enqueue, proxy

* Changes to CMakeLists.txt

* Additional changes to device code

* Additional changes to net.cc

* adding 'compiler warning' tag to ease upstream merge'

* typo correction

* Addessing comments

* Additional changes for new commits
This commit is contained in:
Avinash
2025-08-05 17:36:23 -05:00
committed by GitHub
parent 6fc228e247
commit 3f8cac388e
28 changed files with 104 additions and 106 deletions
Download Patch File Download Diff File Expand all files Collapse all files
src/graph/connect.cc
+2 -2
View File
@@ -119,14 +119,14 @@ bool isRankHere(const char* s, int start, int end, int rank) {
ncclResult_t ncclTreeBasePostset(struct ncclComm* comm,
struct ncclTopoGraph* treeGraph) {
int x=0, y=0;
int x=0;
for (int i=0; treeGraph->treeBase[i][0]!=0; i++)
{
x=i+1;
}
if( treeGraph->treeBase[0][0] == 0) return ncclSuccess;
int nChannels = comm->nChannels;
int localRanks = comm->topo->nodes[GPU].count;
//int localRanks = comm->topo->nodes[GPU].count; // unused variable - compiler warning
//new tree
for (int c=0; c<nChannels; c++) { // in here
int buff = c%x;
src/graph/paths.cc
+6 -4
View File
@@ -286,14 +286,17 @@ ncclResult_t ncclTopoCheckP2p(struct ncclComm* comm, struct ncclTopoSystem* syst
// GPU not found, we can't use p2p.
return ncclSuccess;
}
#if !defined(__HIP_PLATFORM_AMD__) && !defined(__HIPCC__)
int intermediateIndex = -1;
#endif
// Set intermediate GPU rank, if routing through an intermediate GPU.
struct ncclTopoLinkList* path = gpu1->paths[GPU]+g2;
if (path->count == 2) {
struct ncclTopoNode* intermediateNode = path->list[0]->remNode;
if (intermediateNode->type == GPU) {
#if !defined(__HIP_PLATFORM_AMD__) && !defined(__HIPCC__)
intermediateIndex = intermediateNode - system->nodes[GPU].nodes;
#endif
if (intermediateRank) *intermediateRank = intermediateNode->gpu.rank;
}
}
@@ -324,8 +327,7 @@ compare:
// Compute the PCI distance and compare with the p2pLevel.
if (path->type <= p2pLevel) *p2p = 1;
#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
#else
#if !defined(__HIP_PLATFORM_AMD__) && !defined(__HIPCC__)
if (*p2p == 1) {
// NCCL_IGNORE_DISABLED_P2P=2 is used by unit tests that don't want to
// validate against NVML at all since they are pretending to be on other hw.
@@ -515,10 +517,10 @@ ncclResult_t ncclTopoNeedFlush(struct ncclComm* comm, int64_t netId, int netDev,
int g;
struct ncclTopoSystem* system = comm->topo;
NCCLCHECK(ncclTopoRankToIndex(system, rank, &g));
struct ncclTopoNode* gpu = system->nodes[GPU].nodes+g;
#if defined(__HIP_PLATFORM_AMD__) || defined(__HIPCC__)
*flush = 1;
#else
struct ncclTopoNode* gpu = system->nodes[GPU].nodes+g; // unused variable - compiler warning
// Flush is required on Ampere and earlier
if (gpu->gpu.cudaCompCap >= 90) *flush = 0;
// On C2C platforms, data could go through a PCI switch while completions and
src/graph/rome_models.cc
+1 -10
View File
@@ -1338,7 +1338,6 @@ end:
*/
ncclResult_t parseGraphLight(const char* str, struct ncclTopoSystem* system, struct ncclTopoGraph* graph, int* gpu_map) {
int gpus[NCCL_TOPO_MAX_NODES]; //transcribe/change according to gpu_map
int nChannels = 0;
int gpu = 0;
int offset = 0;
int start_offset = offset;
@@ -1348,7 +1347,7 @@ ncclResult_t parseGraphLight(const char* str, struct ncclTopoSystem* system, str
}
int status = 0; // 0 : between numbers, 1 : inside number
int ngpus = system->nodes[GPU].count;
int x=0, y=0;
int x=0;
do {
int digit = str[offset] - '0';
if (digit >= 0 && digit <= 9) {
@@ -1855,7 +1854,6 @@ ncclResult_t parseA2a8P(struct ncclTopoSystem* system, struct ncclTopoGraph* gra
bool isAlltoall = checkAlltoallWidth(&romeTopo);
if (!isAlltoall) return ncclSuccess;
int gcnt = 0;
int *g8, n[NCCL_TOPO_MAX_NODES];
int *all_gpu_permutations = (int *)malloc(TOTAL_PERMUTE_COUNT*NUMA_CPUS*NUMA_GPUS*sizeof(int));
struct timeval tvs, tve;
@@ -1878,7 +1876,6 @@ ncclResult_t parseA2a8P(struct ncclTopoSystem* system, struct ncclTopoGraph* gra
}
if (ngpusPerNuma == 0) continue;
if (ngpusPerNuma != NUMA_GPUS) break;
gcnt++;
// init GPU mapping
for (int k = 0; k < ngpus; k++) {
if (romeTopo.gpuNuma[k] != j) continue;
@@ -1927,7 +1924,6 @@ ncclResult_t parseA2a8P(struct ncclTopoSystem* system, struct ncclTopoGraph* gra
if (p < TOTAL_PERMUTE_COUNT) break;
}
gettimeofday(&tve, NULL);
float t = (tve.tv_sec - tvs.tv_sec)*1E3 + (tve.tv_usec - tvs.tv_usec)/1E3;
if (i >= sizeof(romeTopoModels)/sizeof(romeTopoModels[0])) {
//printf("No solution in %.2fms\n", t);
return ncclSuccess;
@@ -2035,7 +2031,6 @@ ncclResult_t parseRome4P2H(struct ncclTopoSystem* system, struct ncclTopoGraph*
int i;
int ngpus = system->nodes[GPU].count;
int ncpus = system->nodes[CPU].count;
int nnets = system->nodes[NET].count;
// Only support ring and tree graphs
@@ -2129,7 +2124,6 @@ ncclResult_t parseRome4P2H(struct ncclTopoSystem* system, struct ncclTopoGraph*
}
}
gettimeofday(&tve, NULL);
float t = (tve.tv_sec - tvs.tv_sec)*1E3 + (tve.tv_usec - tvs.tv_usec)/1E3;
if (i >= sizeof(romeTopoModels)/sizeof(romeTopoModels[0])) {
//printf("No solution in %.2fms (%d iter)\n", t, time);
return ncclSuccess;
@@ -2241,7 +2235,6 @@ ncclResult_t parse1H16P(struct ncclTopoSystem* system, struct ncclTopoGraph* gra
// only match for system with 16 GPUs
if (ngpus != 16 || ncpus != NUMA_CPUS) return ncclSuccess;
int gcnt = 0;
int *g16, n[NCCL_TOPO_MAX_NODES], rdm[NUMA_GPUS*NUMA_CPUS];
int *all_gpu_permutations = (int *)malloc(TOTAL_PERMUTE_COUNT*NUMA_CPUS*NUMA_GPUS*sizeof(int));
struct timeval tvs, tve;
@@ -2262,7 +2255,6 @@ ncclResult_t parse1H16P(struct ncclTopoSystem* system, struct ncclTopoGraph* gra
}
if (ngpusPerNuma == 0) continue;
if (ngpusPerNuma != NUMA_GPUS) break;
gcnt++;
// init GPU mapping
for (int k = 0; k < ngpus; k++) {
if (romeTopo.gpuNuma[k] != j) continue;
@@ -2317,7 +2309,6 @@ ncclResult_t parse1H16P(struct ncclTopoSystem* system, struct ncclTopoGraph* gra
if (p < TOTAL_PERMUTE_COUNT) break;
}
gettimeofday(&tve, NULL);
float t = (tve.tv_sec - tvs.tv_sec)*1E3 + (tve.tv_usec - tvs.tv_usec)/1E3;
if (i >= sizeof(romeTopoModels)/sizeof(romeTopoModels[0])) {
//printf("No solution in %.2fms\n", t);
return ncclSuccess;
src/graph/tuning.cc
+11 -6
View File
@@ -168,7 +168,7 @@ static struct tuningModel tuning_model_0 {
{ 1.0, 0.8, 0.2, 1.0, 1.0, 0.3, 1.0, 0.1, 0.1, 0.2, 0.2, 0.1, 0.5, 1.0, 0.8, 0.8, 1.0, 0.9, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, },
},
.llProtoRanges = {RCCL_LL_LIMITS_UNDEFINED},
.llProtoRanges = {{{RCCL_LL_LIMITS_UNDEFINED}}},
};
static struct tuningModel tuning_model_1 {
@@ -200,7 +200,7 @@ static struct tuningModel tuning_model_1 {
{ 0.3, 1.0, 0.3, 0.1, 0.1, 0.1, 0.3, 0.7, 1.0, 0.2, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.2, 0.3, 0.5, 0.9, 1.0, 1.0, 1.0, 1.0, },
},
.llProtoRanges = {RCCL_LL_LIMITS_UNDEFINED},
.llProtoRanges = {{{RCCL_LL_LIMITS_UNDEFINED}}},
};
static struct tuningModel tuning_model_2 {
@@ -232,7 +232,7 @@ static struct tuningModel tuning_model_2 {
{ 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 1.0, 0.2, 0.2, 0.1, 0.1, 0.1, 0.1, 0.1, 0.2, 0.4, 0.5, 0.6, 0.9, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, },
},
.llProtoRanges = {RCCL_LL_LIMITS_UNDEFINED},
.llProtoRanges = {{{RCCL_LL_LIMITS_UNDEFINED}}},
};
static struct tuningModel tuning_model_3 {
@@ -264,7 +264,7 @@ static struct tuningModel tuning_model_3 {
{ 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.5, 1.0, 0.1, 0.3, 0.1, 0.1, 0.1, 0.2, 0.2, 0.2, 0.3, 0.4, 0.7, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, },
},
.llProtoRanges = {RCCL_LL_LIMITS_UNDEFINED},
.llProtoRanges = {{{RCCL_LL_LIMITS_UNDEFINED}}},
};
static struct tuningModel tuning_model_4 {
@@ -296,7 +296,7 @@ static struct tuningModel tuning_model_4 {
{ 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 0.1, 1.0, 1.0, 0.8, 0.5, 0.1, 0.7, 0.2, 0.4, 0.4, 0.6, 0.7, 0.9, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, },
},
.llProtoRanges = {RCCL_LL_LIMITS_UNDEFINED},
.llProtoRanges = {{{RCCL_LL_LIMITS_UNDEFINED}}},
};
static struct tuningModel tuning_model_5 {
@@ -394,6 +394,7 @@ static struct tuningModel rcclTuningModel[] = {
#define HOPPER_COMPCAP_IDX 2
#define BLACKWELL_COMPCAP_IDX 3
#if !defined(__HIP_PLATFORM_AMD__) && !defined(__HIPCC__)
// LL128 max BW per channel
static const double llMaxBws[][3] = {
/* Volta-N1/Intel-N2/Intel-N4) */ {39.0, 39.0, 20.4},
@@ -420,6 +421,7 @@ static const double perChMaxTreeBws[][3] = {
/* Hopper (N1/N2/N4) */ {38.7, 41.4, 36.0},
/* Blackwell (N1/N2/N4) */ {2*38.7, 2*41.4, 2*36.0},
};
#endif
NCCL_PARAM(PatEnable, "PAT_ENABLE", 0);
static int ncclPatEnable(struct ncclComm* comm) {
@@ -470,7 +472,7 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom
int nNodes = comm->nNodes;
int nRanks = comm->nRanks;
if (nRanks <= 1) return ncclSuccess;
#if !defined(__HIP_PLATFORM_AMD__) && !defined(__HIPCC__)
int compCapIndex = minCompCap >= 100 ? BLACKWELL_COMPCAP_IDX : (minCompCap >= 90 ? HOPPER_COMPCAP_IDX : minCompCap >= 80 ? AMPERE_COMPCAP_IDX : VOLTA_COMPCAP_IDX);
int index2 = nNodes <= 2 ? nNodes-1 : 2;
// LL: for single node, we look at GPU type; for multi-node, we look at CPU type
@@ -480,6 +482,7 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom
double perChMaxTreeBw = perChMaxTreeBws[compCapIndex][index2];
double perChMaxRingLL128Bw = perChMaxRingLL128Bws[compCapIndex][index2];
double perChMaxTreeLL128Bw = perChMaxTreeLL128Bws[compCapIndex][index2];
#endif
// De-penalize Tree/Simple latency on Power systems to favor Tree than Ring
//if (comm->cpuArch == NCCL_TOPO_CPU_ARCH_POWER) hwLat[NCCL_HW_PCI][NCCL_ALGO_TREE][NCCL_PROTO_SIMPLE] = hwLat[NCCL_HW_PCI][NCCL_ALGO_RING][NCCL_PROTO_SIMPLE];
float ppn = (float)nRanks / nNodes;
@@ -805,11 +808,13 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom
// Trees are not perfectly sticking to the model for medium sizes. Applying a static correction
// factor is not ideal but works quite well. Powers of two, 64 B to 256MB.
#if !defined(__HIP_PLATFORM_AMD__) && !defined(__HIPCC__)
static float treeCorrectionFactor[NCCL_NUM_PROTOCOLS][23] = {
{ 1.0, 1.0, 1.0, 1.0, .9, .8, .7, .7, .7, .7, .6, .5, .4, .4, .5, .6, .7, .8, .9, 1.0, 1.0, 1.0, 1.0 },
{ 1.0, 1.0, 1.0, 1.0, 1.0, .9, .8, .8, .8, .7, .6, .6, .6, .6, .6, .6, .8, .9, .9, .9, .9, 1.0, 1.0 },
{ .9, .9, .9, .9, .9, .9, .9, .8, .7, .6, .6, .5, .5, .5, .5, .6, .7, .8, .7, .7, .8, .9, .9 }
};
#endif
ncclResult_t ncclTopoGetAlgoTime(struct ncclComm* comm, int coll, int algorithm, int protocol, size_t nBytes, int numPipeOps, float* time) {
float bw = comm->bandwidths[coll][algorithm][protocol];