move nvidia related code from flagcx core

Author: yzhang35

Makefile

@@ -238,6 +238,7 @@ INCLUDEDIR := \
 	$(abspath flagcx/core) \
 	$(abspath flagcx/adaptor) \
 	$(abspath flagcx/adaptor/include) \
+	$(abspath flagcx/adaptor/tuner) \
 	$(abspath flagcx/service)
 
 LIBSRCFILES:= \
@@ -247,6 +248,7 @@ LIBSRCFILES:= \
 	$(wildcard flagcx/adaptor/device/*.cc) \
 	$(wildcard flagcx/adaptor/ccl/*.cc) \
 	$(wildcard flagcx/adaptor/net/*.cc) \
+	$(wildcard flagcx/adaptor/tuner/tuner_util.cc) \
 	$(wildcard flagcx/service/*.cc)
 
 LIBOBJ     := $(LIBSRCFILES:%.cc=$(OBJDIR)/%.o)

flagcx/adaptor/tuner/nccl_param.cc

@@ -28,6 +28,6 @@ static void ncclLoadParam(char const* env, int64_t deftVal, int64_t* value) {
     return value;				 \
   }
 
-// In order to intercept a NCCL environment getter function, add an additional line of NCCL_PRARM here.
+// In order to intercept a NCCL environment getter function, add an additional line of NCCL_PARAM here.
 // Below is an example of intercepting NCCL_P2P_NVL_CHUNKSIZE env.
 NCCL_PARAM(P2pNvlChunkSize, "P2P_NVL_CHUNKSIZE", (1 << 19)); /* 512 kB */

flagcx/adaptor/tuner/tuner_util.cc

@@ -0,0 +1,19 @@
+#include "tuner/tuner_util.h"
+
+#ifdef USE_NVIDIA_ADAPTOR
+static struct flagcxEnvConfig config1 = {
+    "defaultConfig1",
+    1,
+    {FLAGCX_ENV_TYPE_CREATION, "NCCL_P2P_NVL_CHUNKSIZE", "1024", "524288"}};
+static struct flagcxEnvConfig config2 = {
+    "defaultConfig2",
+    1,
+    {FLAGCX_ENV_TYPE_CREATION, "NCCL_P2P_NVL_CHUNKSIZE", "524288", "524288"}};
+
+// demo
+flagcxResult_t loadConfigList(std::vector<struct flagcxEnvConfig> &cfgList) {
+  cfgList.push_back(config1);
+  cfgList.push_back(config2);
+  return flagcxSuccess;
+}
+#endif

flagcx/adaptor/tuner/tuner_util.h

@@ -0,0 +1,10 @@
+#ifndef FLAGCX_TUNER_UTIL_H_
+#define FLAGCX_TUNER_UTIL_H_
+
+#include "tuner.h" // struct flagcxEnvConfig
+#include <vector>
+
+// This is a demonstration function that provide a way to load all config list for a specific GPU.
+flagcxResult_t loadConfigList(std::vector<struct flagcxEnvConfig> &cfgList);
+
+#endif // end include guard

flagcx/core/flagcx_tuner.cc

@@ -2,6 +2,7 @@
 #include "check.h"
 #include "param.h"
 #include "timer.h"
+#include "tuner/tuner_util.h"
 #include <cfloat>
 #include <map>
 #include <sstream>
@@ -113,15 +114,6 @@ struct flagcxTunerContext {
   flagcxTimer<TunerProfileKey> timer;
 };
 
-static struct flagcxEnvConfig config1 = {
-    "defaultConfig1",
-    1,
-    {FLAGCX_ENV_TYPE_CREATION, "NCCL_P2P_NVL_CHUNKSIZE", "1024", "524288"}};
-static struct flagcxEnvConfig config2 = {
-    "defaultConfig2",
-    1,
-    {FLAGCX_ENV_TYPE_CREATION, "NCCL_P2P_NVL_CHUNKSIZE", "524288", "524288"}};
-
 bool operator<(const struct flagcxCommTag &lhs,
                const struct flagcxCommTag &rhs) {
   return strcmp(lhs.tag, rhs.tag) < 0;
@@ -150,8 +142,7 @@ flagcxResult_t flagcxTunerInit(size_t nRanks, size_t nNodes,
                                flagcxDebugLogger_t logFunction,
                                void **context) {
   struct flagcxTunerContext *ctx = new struct flagcxTunerContext;
-  ctx->configList.push_back(config1);
-  ctx->configList.push_back(config2);
+  FLAGCXCHECK(loadConfigList(ctx->configList));
   ctx->logger = logFunction;
   *context = ctx;