diff --git a/projects/rccl/src/device/msccl_kernel_impl.h b/projects/rccl/src/device/msccl_kernel_impl.h index 98af766a93..03f6f49e3f 100644 --- a/projects/rccl/src/device/msccl_kernel_impl.h +++ b/projects/rccl/src/device/msccl_kernel_impl.h @@ -48,7 +48,7 @@ extern __shared__ struct mscclShmemData mscclShmem; #endif inline __device__ static void barrier(int nthreads) { -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) assert(nthreads == NCCL_MAX_NTHREADS); __asm__ __volatile__("s_waitcnt vmcnt(0) lgkmcnt(0)\ns_barrier"); #else diff --git a/projects/rccl/src/device/prims_ll.h b/projects/rccl/src/device/prims_ll.h index 387b09816b..58005277a3 100644 --- a/projects/rccl/src/device/prims_ll.h +++ b/projects/rccl/src/device/prims_ll.h @@ -81,7 +81,7 @@ private: uint64_t* barrier_next; inline __device__ void barrier() { -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) if (nthreads != WARP_SIZE) barrier_by_group(); #else @@ -159,7 +159,7 @@ private: } #endif -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) union ncclLLFifoLine i4; do { i4.v[0] = LL_LOAD(src->v); @@ -197,7 +197,7 @@ private: for (int i=BeginIx; i < MaxRecv; i++) { if (i < fan.nrecv()) { union ncclLLFifoLine* src = recvPtr(i) + offset; -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) line[i].v[0] = LL_LOAD(src->v); line[i].v[1] = LL_LOAD(src->v+1); #else @@ -219,7 +219,7 @@ private: #endif do { -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) line[i].v[0] = LL_LOAD(src->v); line[i].v[1] = LL_LOAD(src->v+1); #else @@ -243,7 +243,7 @@ private: } __device__ void storeLL(union ncclLLFifoLine* dst, uint64_t val, uint32_t flag) { -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) union ncclLLFifoLine i4; i4.data1 = val & 0xffffffff; i4.flag1 = flag; @@ -267,7 +267,7 @@ private: uint32_t u4; uint64_t u8; }; -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) if(sizeof(U) == 1) u1 = LL_LOAD((uint8_t*)src); else if(sizeof(U) == 2) @@ -299,7 +299,7 @@ private: uint64_t u8; }; elt = val; -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) if(sizeof(U) == 1) LL_STORE(u1, (uint8_t*)dst); else if(sizeof(U) == 2) diff --git a/projects/rccl/src/device/prims_ll128.h b/projects/rccl/src/device/prims_ll128.h index 476fb702a9..6c7674c432 100644 --- a/projects/rccl/src/device/prims_ll128.h +++ b/projects/rccl/src/device/prims_ll128.h @@ -74,7 +74,7 @@ private: #endif inline __device__ void barrier() { -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) if (nthreads != WARP_SIZE) barrier_by_group(); #else diff --git a/projects/rccl/src/enqueue.cc b/projects/rccl/src/enqueue.cc index 0b6dc8578b..98a3d78491 100644 --- a/projects/rccl/src/enqueue.cc +++ b/projects/rccl/src/enqueue.cc @@ -1617,7 +1617,7 @@ static ncclResult_t getChannnelThreadInfo(struct ncclInfo* collInfo) { // Ring/Tree channel tuning while (collInfo->nBytes < nc*nt*threadThreshold) { if (nc >= 2) nc--; -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) // do not reduce threads count on VEGA #else else if ((nt % 128) == 0) nt/=2; @@ -1625,7 +1625,7 @@ static ncclResult_t getChannnelThreadInfo(struct ncclInfo* collInfo) { else break; } } -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) #else if (collInfo->protocol == NCCL_PROTO_SIMPLE) { if (collInfo->algorithm == NCCL_ALGO_RING) nt += WARP_SIZE; // Extra warp for sync diff --git a/projects/rccl/src/graph/paths.cc b/projects/rccl/src/graph/paths.cc index df1eba69e5..6833c4e436 100644 --- a/projects/rccl/src/graph/paths.cc +++ b/projects/rccl/src/graph/paths.cc @@ -304,7 +304,7 @@ compare: // Compute the PCI distance and compare with the p2pLevel. if (path->type <= p2pLevel) *p2p = 1; -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) #else if (*p2p == 1) { // NCCL_IGNORE_DISABLED_P2P=2 is used by unit tests that don't want to diff --git a/projects/rccl/src/graph/search.cc b/projects/rccl/src/graph/search.cc index 8fe666ae48..31d14a74e6 100644 --- a/projects/rccl/src/graph/search.cc +++ b/projects/rccl/src/graph/search.cc @@ -886,7 +886,7 @@ ncclResult_t ncclTopoDupChannels(struct ncclTopoGraph* graph, int ccMin, int ngp return ncclSuccess; } -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) float speedArrayIntra[] = { 48.0, 24.0, 20.0, 18.0, 15.0, 12.0, 10.0, 9.0, 7.0, 6.0, 5.0, 4.0, 3.0, 2.4, 1.2, 0.24, 0.12 }; float speedArrayInter[] = { 48.0, 24.0, 20.0, 18.0, 15.0, 12.0, 10.0, 9.0, 7.0, 6.0, 5.0, 4.0, 3.0, 2.4, 1.2, 0.24, 0.12 }; #define NSPEEDSINTRA (sizeof(speedArrayIntra)/sizeof(float)) diff --git a/projects/rccl/src/graph/topo.cc b/projects/rccl/src/graph/topo.cc index afa179916a..7ad2836cf1 100644 --- a/projects/rccl/src/graph/topo.cc +++ b/projects/rccl/src/graph/topo.cc @@ -21,7 +21,7 @@ #define BUSID_REDUCED_SIZE (sizeof("0000:00")) const char* topoNodeTypeStr[] = { "GPU", "PCI", "NVS", "CPU", "NIC", "NET" }; -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) const char* topoLinkTypeStr[] = { "LOC", "XGMI", "", "PCI", "", "", "", "SYS", "NET" }; const char* topoPathTypeStr[] = { "LOC", "XGMI", "NVB", "PIX", "PXB", "PXN", "PHB", "SYS", "DIS" }; #else @@ -376,7 +376,7 @@ ncclResult_t ncclTopoAddNic(struct ncclXmlNode* xmlNic, struct ncclTopoSystem* s } ncclResult_t ncclTopoAddGpu(struct ncclXmlNode* xmlGpu, struct ncclTopoSystem* system, struct ncclTopoNode* gpu) { -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) // There is no direct mapping between CUDA SM to HIP GFX. Use SM60 as compatibility level. gpu->gpu.cudaCompCap = 60; #else @@ -524,7 +524,7 @@ ncclResult_t ncclTopoAddCpu(struct ncclXmlNode* xmlCpu, struct ncclTopoSystem* s return ncclSuccess; } -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) ncclResult_t ncclTopoAddXGMI(struct ncclXmlNode* node, struct ncclTopoSystem* system, const char* parentBusId) { if (strcmp(node->name, "xgmi") == 0) { struct ncclTopoNode* gpu = NULL; @@ -666,7 +666,7 @@ ncclResult_t ncclTopoGetSystemFromXml(struct ncclXml* xml, struct ncclTopoSystem struct ncclXmlNode* node = topNode->subs[s]; if (strcmp(node->name, "cpu") == 0) NCCLCHECK(ncclTopoAddCpu(node, *topoSystem)); } -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) NCCLCHECK(ncclTopoAddXGMI(topNode, *topoSystem, NULL)); #else NCCLCHECK(ncclTopoAddNvLinks(topNode, *topoSystem, NULL)); diff --git a/projects/rccl/src/graph/tuning.cc b/projects/rccl/src/graph/tuning.cc index ad61051c06..9d3441401e 100644 --- a/projects/rccl/src/graph/tuning.cc +++ b/projects/rccl/src/graph/tuning.cc @@ -303,7 +303,7 @@ static float getNetOverhead(struct ncclComm* comm) { ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCompCap, struct ncclTopoGraph** graphs) { int simpleDefaultThreads = (graphs[NCCL_ALGO_RING]->bwIntra*graphs[NCCL_ALGO_RING]->nChannels <= PCI_BW) ? 256 : NCCL_SIMPLE_MAX_NTHREADS; comm->maxThreads[NCCL_ALGO_RING][NCCL_PROTO_SIMPLE] = -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) getNthreads("NCCL_NTHREADS", ncclParamNthreads(), 4*comm->WarpSize, NCCL_MAX_NTHREADS, simpleDefaultThreads, comm->WarpSize); comm->maxThreads[NCCL_ALGO_TREE][NCCL_PROTO_SIMPLE] = comm->maxThreads[NCCL_ALGO_COLLNET_DIRECT][NCCL_PROTO_SIMPLE] = getNthreads("NCCL_NTHREADS", ncclParamNthreads(), 4*comm->WarpSize, NCCL_MAX_NTHREADS, NCCL_MAX_NTHREADS, comm->WarpSize); @@ -374,7 +374,7 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom if (a == NCCL_ALGO_NVLS_TREE) bw = std::min(graphs[a]->bwIntra, nNodes <= 2 ? graphs[a]->bwInter : graphs[a]->bwInter/2); // Various model refinements -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) if (nNodes <= 2) busBw *= rcclTuningModel[comm->topo->tuning].bwRatio[0][a][p]; else @@ -509,7 +509,7 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom // Disable LL protocol on gfx11xx int pEnable = protoEnable[p]; if (pEnable == 2 && p == NCCL_PROTO_LL128) { -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) #if defined(ENABLE_LL128) // Enable LL128 by default only on gfx90a with available tuning table pEnable = (graphs[a]->typeInter <= PATH_PXB) && graphs[a]->typeIntra <= PATH_NVL && @@ -649,7 +649,7 @@ ncclResult_t ncclTopoGetAlgoTime(struct ncclInfo* info, int algorithm, int proto } int logSize = log2i(info->nBytes>>6); -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) if (algorithm == NCCL_ALGO_TREE) { if (logSize < 27) bw *= rcclTuningModel[info->comm->topo->tuning].treeCorrectionFactor[protocol][logSize]; else bw *= rcclTuningModel[info->comm->topo->tuning].treeCorrectionFactor[protocol][26]; diff --git a/projects/rccl/src/graph/xml.cc b/projects/rccl/src/graph/xml.cc index 6af3d0db27..770438e4fe 100644 --- a/projects/rccl/src/graph/xml.cc +++ b/projects/rccl/src/graph/xml.cc @@ -266,7 +266,7 @@ ncclResult_t ncclTopoXmlLoadC2c(FILE* file, struct ncclXml* xml, struct ncclXmlN return ncclSuccess; } ncclResult_t ncclTopoXmlLoadGpu(FILE* file, struct ncclXml* xml, struct ncclXmlNode* head) { -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) struct xmlHandler handlers[] = { { "xgmi", ncclTopoXmlLoadNvlink } }; #else struct xmlHandler handlers[] = { { "nvlink", ncclTopoXmlLoadNvlink }, { "c2c", ncclTopoXmlLoadC2c } }; @@ -690,7 +690,7 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev struct ncclXmlNode* nvlNode = NULL; NCCLCHECK(xmlGetSub(gpuNode, "nvlink", &nvlNode)); if (nvlNode == NULL) { -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) const char* busId; NCCLCHECK(xmlGetAttr(pciNode, "busid", &busId)); uint32_t deviceCnt; @@ -813,7 +813,7 @@ ncclResult_t ncclTopoGetXmlFromGpu(struct ncclXmlNode* pciNode, uint32_t rocmDev // Fill target classes for (int s=0; snSubs; s++) { struct ncclXmlNode* sub = gpuNode->subs[s]; -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) if (strcmp(sub->name, "xgmi") != 0) continue; #else if (strcmp(sub->name, "nvlink") != 0) continue; @@ -845,7 +845,7 @@ ncclResult_t ncclTopoFillGpu(struct ncclXml* xml, const char* busId, struct nccl NCCLCHECK(ncclTopoGetPciNode(xml, busId, &node)); NCCLCHECK(xmlSetAttrIfUnset(node, "class", "0x03")); NCCLCHECK(ncclTopoGetXmlFromSys(node, xml)); -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) uint32_t devIndex; static int rocmsmiInit = 0; if (rocmsmiInit == 0) { diff --git a/projects/rccl/src/include/comm.h b/projects/rccl/src/include/comm.h index 321f09d52c..dfabb70632 100644 --- a/projects/rccl/src/include/comm.h +++ b/projects/rccl/src/include/comm.h @@ -17,7 +17,7 @@ #include "nccl_net.h" #include "register.h" -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) #define HIPRT_CB #else #if CUDART_VERSION < 9000 diff --git a/projects/rccl/src/include/gdrwrap.h b/projects/rccl/src/include/gdrwrap.h index a791e050f0..669cbace1a 100644 --- a/projects/rccl/src/include/gdrwrap.h +++ b/projects/rccl/src/include/gdrwrap.h @@ -157,7 +157,7 @@ typedef struct gdr_mem_desc { gdr_mh_t gdrMh; } gdr_mem_desc_t; -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) static gdr_t ncclGdrInit() { INFO(NCCL_INIT, "Enabled GDRCopy equivalent memory allocation"); return (gdr_t)0x12345678L; diff --git a/projects/rccl/src/include/nvtx.h b/projects/rccl/src/include/nvtx.h index b09d67c879..ed5741edbc 100644 --- a/projects/rccl/src/include/nvtx.h +++ b/projects/rccl/src/include/nvtx.h @@ -77,7 +77,7 @@ class payload_schema { // @param N schema name // @param S schema (entries) // @param P payload (struct) -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) #define NVTX3_FUNC_WITH_PARAMS(ID, S, P) \ nvtxPayloadData_t nvtx3_bpl__[] = { \ {NVTX_PAYLOAD_ENTRY_TYPE_SCHEMA_ID_STATIC_START + NVTX_SID_##ID, sizeof(P), &(P)}}; \ diff --git a/projects/rccl/src/include/nvtx3/nvtx3.hpp b/projects/rccl/src/include/nvtx3/nvtx3.hpp index a4cef3849f..61bda49ccc 100644 --- a/projects/rccl/src/include/nvtx3/nvtx3.hpp +++ b/projects/rccl/src/include/nvtx3/nvtx3.hpp @@ -2777,7 +2777,7 @@ inline void mark(Args const&... args) noexcept * `domain` to which the `registered_string_in` belongs. Else, * `domain::global` to indicate that the global NVTX domain should be used. */ -#if !defined(__HIP_PLATFORM_HCC__) && !defined(__HCC__) && !defined(__HIPCC__) +#if !defined(__HIP_PLATFORM_AMD__) && !defined(__HCC__) && !defined(__HIPCC__) #define NVTX3_V1_FUNC_RANGE_IN(D) \ static ::nvtx3::v1::registered_string_in const nvtx3_func_name__{__func__}; \ static ::nvtx3::v1::event_attributes const nvtx3_func_attr__{nvtx3_func_name__}; \ diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc index d6deea335a..23abcf8aae 100644 --- a/projects/rccl/src/init.cc +++ b/projects/rccl/src/init.cc @@ -52,7 +52,7 @@ #define STR2(v) #v #define STR(v) STR2(v) -#if CUDART_VERSION >= 9020 || defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if CUDART_VERSION >= 9020 || defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) #define NCCL_GROUP_CUDA_STREAM 0 // CGMD: CUDA 9.2,10.X Don't need to use an internal CUDA stream #else #define NCCL_GROUP_CUDA_STREAM 1 // CGMD: CUDA 9.0,9.1 Need to use an internal CUDA stream @@ -682,7 +682,7 @@ fail: } // Pre-process the string so that running "strings" on the lib can quickly reveal the version. -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) #define VERSION_STRING "RCCL version " STR(NCCL_MAJOR) "." STR(NCCL_MINOR) "." STR(NCCL_PATCH) NCCL_SUFFIX "+hip" STR(HIP_VERSION_MAJOR) "." STR(HIP_VERSION_MINOR) #else #define VERSION_STRING "NCCL version " STR(NCCL_MAJOR) "." STR(NCCL_MINOR) "." STR(NCCL_PATCH) NCCL_SUFFIX "+cuda" STR(CUDA_MAJOR) "." STR(CUDA_MINOR) @@ -733,7 +733,7 @@ static ncclResult_t fillInfo(struct ncclComm* comm, struct ncclPeerInfo* info, u info->comm = comm; info->cudaCompCap = comm->minCompCap = comm->maxCompCap = comm->compCap; -#if !defined(__HIP_PLATFORM_HCC__) && !defined(__HCC__) && !defined(__HIPCC__) +#if !defined(__HIP_PLATFORM_AMD__) && !defined(__HCC__) && !defined(__HIPCC__) // MNNVL support { // MNNVL: Request the fabric UUID and partition info diff --git a/projects/rccl/src/net.cc b/projects/rccl/src/net.cc index 2cb9923266..1326ea50a5 100644 --- a/projects/rccl/src/net.cc +++ b/projects/rccl/src/net.cc @@ -548,7 +548,7 @@ ncclResult_t ncclGpuGdrSupport(struct ncclComm* comm, int* gdrSupport) { ncclNetProperties_t props; NCCLCHECK(comm->ncclNet->getProperties(dev, &props)); if ((props.ptrSupport & NCCL_PTR_CUDA) == 0) continue; -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) gdrSupportMatrix[comm->cudaDev] = 1; break; #endif diff --git a/projects/rccl/src/transport/net.cc b/projects/rccl/src/transport/net.cc index b98ea9567f..2643849522 100644 --- a/projects/rccl/src/transport/net.cc +++ b/projects/rccl/src/transport/net.cc @@ -153,7 +153,7 @@ struct recvNetResources { /* Determine if two peers can communicate with NET */ static ncclResult_t canConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) { *ret = 1; -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) #else if (info1->hostHash == info2->hostHash) { // If on the same host, check intra-node net is not disabled. diff --git a/projects/rccl/src/transport/net_ib.cc b/projects/rccl/src/transport/net_ib.cc index 02fd0cac94..f463e67123 100644 --- a/projects/rccl/src/transport/net_ib.cc +++ b/projects/rccl/src/transport/net_ib.cc @@ -352,7 +352,7 @@ ncclResult_t ncclIbDevices(int* ndev) { ncclResult_t ncclIbGdrSupport() { static int moduleLoaded = -1; if (moduleLoaded == -1) { -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) moduleLoaded = (access("/sys/kernel/mm/memory_peers/amdkfd/version", F_OK) == -1) ? 0 : 1; char strValue[MAX_STR_LEN]; NCCLCHECK(ncclTopoGetStrFromSys("/sys/devices/virtual/dmi/id", "bios_version", strValue)); diff --git a/projects/rccl/src/transport/p2p.cc b/projects/rccl/src/transport/p2p.cc index 0798284195..63a0c99dc0 100644 --- a/projects/rccl/src/transport/p2p.cc +++ b/projects/rccl/src/transport/p2p.cc @@ -106,7 +106,7 @@ static void initCeOperation(); /* Determine if two peers can communicate through p2p */ ncclResult_t p2pCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTopoGraph* graph, struct ncclPeerInfo* info1, struct ncclPeerInfo* info2) { initCeOperation(); -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) if (!info1->hasFineGrain || !info2->hasFineGrain) { *ret = 0; return ncclSuccess; @@ -146,7 +146,7 @@ ncclResult_t p2pCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTop int cudaDev1 = busIdToCudaDev(info1->busId); int cudaDev2 = busIdToCudaDev(info2->busId); if (cudaDev1 == -1 || cudaDev2 == -1) { -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) || CUDART_VERSION >= 10010 +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) || CUDART_VERSION >= 10010 // CUDA 10.1 and later can use P2P with invisible devices. return ncclSuccess; #else @@ -165,7 +165,7 @@ ncclResult_t p2pCanConnect(int* ret, struct ncclTopoSystem* topo, struct ncclTop return ncclSuccess; } -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) #else // This will always fail when using NCCL_CUMEM_ENABLE=1 if (p2p != 0 && !ncclCuMemEnable()) { diff --git a/projects/rccl/tools/rccl-prim-test/copy_kernel.h b/projects/rccl/tools/rccl-prim-test/copy_kernel.h index 62e0bade2a..7fe4ba63cb 100644 --- a/projects/rccl/tools/rccl-prim-test/copy_kernel.h +++ b/projects/rccl/tools/rccl-prim-test/copy_kernel.h @@ -85,7 +85,7 @@ struct MULTI128 { }; inline __device__ void Fetch128(Pack128& v, const Pack128* p) { -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) v.x = __builtin_nontemporal_load(&p->x); v.y = __builtin_nontemporal_load(&p->y); #else @@ -93,7 +93,7 @@ inline __device__ void Fetch128(Pack128& v, const Pack128* p) { #endif } inline __device__ void Store128(Pack128* p, Pack128& v) { -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) __builtin_nontemporal_store(v.x, &p->x); __builtin_nontemporal_store(v.y, &p->y); #else @@ -198,7 +198,7 @@ __device__ int ptrAlign128(T* ptr) { return (uint64_t)ptr % alignof(int32_t); } #define PACKELEMS (sizeof(Pack128) / sizeof(T)) -#if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) +#if defined(__HIP_PLATFORM_AMD__) || defined(__HCC__) || defined(__HIPCC__) // Multiply UNROLL by 2 if single source/single destination #define AUTOUNROLL (UNROLL*((MINSRCS==1 && MINDSTS==1) ? 2 : 1)) #endif