Skip to content

Commit ee8b2b7

Browse files
cq-engchenqingyzhang35
authored
Add the feature of automatically generating candidates (#274)
--------- Co-authored-by: chenqing <chenqing@infrawaves.com> Co-authored-by: Yan Zhang <yzhang35@gmail.com>
1 parent 00873c2 commit ee8b2b7

9 files changed

Lines changed: 153 additions & 23 deletions

File tree

Makefile

100644100755
Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -248,7 +248,7 @@ LIBSRCFILES:= \
248248
$(wildcard flagcx/adaptor/device/*.cc) \
249249
$(wildcard flagcx/adaptor/ccl/*.cc) \
250250
$(wildcard flagcx/adaptor/net/*.cc) \
251-
$(wildcard flagcx/adaptor/tuner/tuner_util.cc) \
251+
$(wildcard flagcx/adaptor/tuner/*.cc) \
252252
$(wildcard flagcx/service/*.cc)
253253

254254
LIBOBJ := $(LIBSRCFILES:%.cc=$(OBJDIR)/%.o)

flagcx/adaptor/device/cuda_adaptor.cc

100644100755
Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -238,7 +238,7 @@ flagcxResult_t cudaAdaptorEventCreate(flagcxEvent_t *event) {
238238
(*event) = NULL;
239239
flagcxCalloc(event, 1);
240240
DEVCHECK(cudaEventCreateWithFlags((cudaEvent_t *)(*event),
241-
cudaEventDisableTiming));
241+
cudaEventDefault));
242242
return flagcxSuccess;
243243
}
244244

flagcx/adaptor/tuner/nccl_tuner.cc

Lines changed: 38 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -0,0 +1,38 @@
1+
#include "tuner/tuner_util.h"
2+
3+
4+
#ifdef USE_NVIDIA_ADAPTOR
5+
6+
static EnvVar algo(
7+
"NCCL_ALGO",
8+
{"ring", "tree"},
9+
"ring"
10+
);
11+
12+
static EnvVar proto(
13+
"NCCL_PROTO",
14+
{"LL", "LL128", "Simple"},
15+
"Simple"
16+
);
17+
18+
static EnvVar thread(
19+
"NCCL_NTHREADS",
20+
{"128", "256"},
21+
"256"
22+
);
23+
24+
static EnvVar minChannel(
25+
"NCCL_MIN_NCHANNELS",
26+
{"16", "32"},
27+
"16"
28+
);
29+
30+
static EnvVar chunkSize(
31+
"NCCL_P2P_NVL_CHUNKSIZE",
32+
{"1024", "2048"},
33+
"1024"
34+
);
35+
36+
std::vector<EnvVar> vars = {algo, proto, thread, minChannel, chunkSize};
37+
38+
#endif
Lines changed: 2 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,8 +3,9 @@ CXX = g++
33
CXXFLAGS = -fPIC -Wall -Wextra -O2 -std=c++11
44
LDFLAGS = -shared -pthread
55

6+
67
# Target shared library name
7-
TARGET = libnccl_param.so
8+
TARGET = libflagcx_param.so
89

910
# Source files
1011
SOURCES = nccl_param.cc
Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -29,5 +29,8 @@ static void ncclLoadParam(char const* env, int64_t deftVal, int64_t* value) {
2929
}
3030

3131
// In order to intercept a NCCL environment getter function, add an additional line of NCCL_PARAM here.
32-
// Below is an example of intercepting NCCL_P2P_NVL_CHUNKSIZE env.
32+
// Below is an example of intercepting NCCL_P2P_NVL_CHUNKSIZE / NTHREADS / MIN_NCHANNELS env.
33+
3334
NCCL_PARAM(P2pNvlChunkSize, "P2P_NVL_CHUNKSIZE", (1 << 19)); /* 512 kB */
35+
NCCL_PARAM(Nthreads, "NTHREADS", (1 << 8)); /*256*/
36+
NCCL_PARAM(MinNchannels, "MIN_NCHANNELS", (1 << 10)); /*1024*/

flagcx/adaptor/tuner/tuner_util.cc

100644100755
Lines changed: 87 additions & 16 deletions
Original file line numberDiff line numberDiff line change
@@ -1,19 +1,90 @@
11
#include "tuner/tuner_util.h"
22

3-
#ifdef USE_NVIDIA_ADAPTOR
4-
static struct flagcxEnvConfig config1 = {
5-
"defaultConfig1",
6-
1,
7-
{FLAGCX_ENV_TYPE_CREATION, "NCCL_P2P_NVL_CHUNKSIZE", "1024", "524288"}};
8-
static struct flagcxEnvConfig config2 = {
9-
"defaultConfig2",
10-
1,
11-
{FLAGCX_ENV_TYPE_CREATION, "NCCL_P2P_NVL_CHUNKSIZE", "524288", "524288"}};
12-
13-
// demo
14-
flagcxResult_t loadConfigList(std::vector<struct flagcxEnvConfig> &cfgList) {
15-
cfgList.push_back(config1);
16-
cfgList.push_back(config2);
17-
return flagcxSuccess;
3+
// Safely copy std::string to char buffer, ensuring NUL termination and truncation
4+
static void safeStrCopy(char *dst, size_t dstSize, const std::string &src) {
5+
if (dstSize == 0) return;
6+
size_t copyLen = std::min(dstSize - 1, src.size());
7+
if (copyLen > 0) memcpy(dst, src.data(), copyLen);
8+
dst[copyLen] = '\0';
9+
}
10+
11+
// Generate all combinations and return a vector of flagcxEnvConfig
12+
flagcxResult_t generateCandidate(std::vector<struct flagcxEnvConfig> &cfgList) {
13+
14+
// Return empty if there are no environment variables
15+
if (vars.empty()){
16+
INFO(FLAGCX_INIT, "Invalid number of environment variables: 0");
17+
return flagcxInvalidArgument;
18+
}
19+
20+
// If the number of variables exceeds the structure capacity, truncate
21+
if (vars.size() > (size_t)FLAGCX_ENV_LIST_MAX_LENGTH) {
22+
INFO(FLAGCX_INIT, "The number of environment variables exceeds the maximum length defined by FLAGCX_ENV_LIST_MAX_LENGTH");
23+
vars.resize(FLAGCX_ENV_LIST_MAX_LENGTH); // Truncate the vars vector
24+
INFO(FLAGCX_INIT, "The number of environment variables has been truncated to FLAGCX_ENV_LIST_MAX_LENGTH (%d)",FLAGCX_ENV_LIST_MAX_LENGTH);
25+
return flagcxSuccess;
26+
}
27+
28+
// Prepare candidate value lists for each variable (at least one empty string to ensure uniform combination logic)
29+
std::vector<std::vector<std::string>> lists;
30+
lists.reserve(vars.size());
31+
for (const auto &v : vars) {
32+
if (v.choices.empty()) {
33+
lists.emplace_back(std::vector<std::string>{""});
34+
} else {
35+
lists.emplace_back(v.choices);
36+
}
37+
}
38+
39+
// Use an index vector to iterate through the Cartesian product (multi-dimensional counter)
40+
size_t nvars = lists.size();
41+
std::vector<size_t> idx(nvars, 0);
42+
bool done = (nvars == 0);
43+
unsigned long numCandidate = 0;
44+
45+
while (!done) {
46+
// Construct a flagcxEnvConfig and zero-initialize
47+
flagcxEnvConfig cfg;
48+
memset(&cfg, 0, sizeof(cfg)); // this zeroes commTag and all fields; adjust if you want non-zero defaults
49+
50+
std::string tagStr = "Config " + std::to_string(numCandidate);
51+
if (tagStr.size() < sizeof(cfg.commTag.tag)) {
52+
safeStrCopy(cfg.commTag.tag, sizeof(cfg.commTag.tag), tagStr);
53+
} else {
54+
INFO(FLAGCX_INIT, "Tag string too long, potential buffer overflow");
55+
return flagcxInvalidArgument;
56+
}
57+
cfg.envCount = 0;
58+
59+
// Fill envs
60+
for (size_t i = 0; i < nvars; ++i) {
61+
flagcxEnvEntity &ent = cfg.envs[i];
62+
// type
63+
ent.type = FLAGCX_ENV_TYPE_CREATION;
64+
// name
65+
safeStrCopy(ent.name, sizeof(ent.name), vars[i].name);
66+
// value
67+
const std::string &val = lists[i][idx[i]];
68+
safeStrCopy(ent.value, sizeof(ent.value), val);
69+
// defaultValue
70+
safeStrCopy(ent.defaultValue, sizeof(ent.defaultValue), vars[i].defaultValue);
71+
72+
cfg.envCount++;
73+
// Stop if exceeding the maximum allowed envs (should not happen since we truncated vars earlier)
74+
if (cfg.envCount >= FLAGCX_ENV_LIST_MAX_LENGTH) break;
75+
}
76+
77+
cfgList.push_back(cfg);
78+
79+
// Increment counter (from least significant to most significant)
80+
for (int i = (int)nvars - 1; i >= 0; --i) {
81+
idx[i]++;
82+
if (idx[i] < lists[i].size()) break;
83+
idx[i] = 0;
84+
if (i == 0) done = true;
85+
}
86+
numCandidate += 1;
87+
}
88+
89+
return flagcxSuccess;
1890
}
19-
#endif

flagcx/adaptor/tuner/tuner_util.h

100644100755
Lines changed: 15 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -3,8 +3,22 @@
33

44
#include "tuner.h" // struct flagcxEnvConfig
55
#include <vector>
6+
#include <string>
67

78
// This is a demonstration function that provide a way to load all config list for a specific GPU.
8-
flagcxResult_t loadConfigList(std::vector<struct flagcxEnvConfig> &cfgList);
9+
10+
struct EnvVar {
11+
std::string name;
12+
std::vector<std::string> choices;
13+
std::string defaultValue;
14+
EnvVar(std::string n="") : name(std::move(n)) {}
15+
EnvVar(std::string n, std::vector<std::string> c, std::string d = "")
16+
: name(std::move(n)), choices(std::move(c)), defaultValue(std::move(d)) {}
17+
};
18+
19+
flagcxResult_t generateCandidate(std::vector<struct flagcxEnvConfig> &cfgList);
20+
static void safeStrCopy(char *dst, size_t dstSize, const std::string &src);
21+
22+
extern std::vector<EnvVar> vars;
923

1024
#endif // end include guard

flagcx/core/flagcx_tuner.cc

100644100755
Lines changed: 4 additions & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -8,6 +8,7 @@
88
#include <sstream>
99
#include <string>
1010
#include <vector>
11+
#include <iostream>
1112

1213
// A category of collective operation. the minimal unit for tuning.
1314
struct TunerCollCategory {
@@ -142,7 +143,9 @@ flagcxResult_t flagcxTunerInit(size_t nRanks, size_t nNodes,
142143
flagcxDebugLogger_t logFunction,
143144
void **context) {
144145
struct flagcxTunerContext *ctx = new struct flagcxTunerContext;
145-
FLAGCXCHECK(loadConfigList(ctx->configList));
146+
FLAGCXCHECK(generateCandidate(ctx->configList));
147+
INFO(FLAGCX_TUNING,
148+
"Candidate number: %ld.", ctx->configList.size());
146149
ctx->logger = logFunction;
147150
*context = ctx;
148151

flagcx/flagcx.cc

100644100755
Lines changed: 1 addition & 1 deletion
Original file line numberDiff line numberDiff line change
@@ -440,7 +440,7 @@ flagcxResult_t flagcxCommInitRank(flagcxComm_t *comm, int nranks,
440440
for (uint32_t i = 0; i < nConfigs; ++i) {
441441
struct flagcxCommTag tag = {""};
442442
FLAGCXCHECK((*comm)->tuner->setCandidate((*comm)->tunerContext, i, &tag));
443-
INFO(FLAGCX_INIT, "start to prepare communicator tag=%s(%u/%u)", tag.tag,
443+
INFO(FLAGCX_INIT | FLAGCX_TUNING, "start to prepare communicator tag=%s(%u/%u)", tag.tag,
444444
i, nConfigs);
445445

446446
flagcxInnerComm_t innerComm = NULL;

0 commit comments

Comments
 (0)