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

update mscclpp (#1488)

Browse Source 
* update commit hash for mscclpp submodule

* update mscclpp submodule

* remove print messages in cmake

* add back some print messages, update MSCLPP CMAKE_ARGS

* enable MSCCL++ patches regardless of finding mscclpp_nccl package
This commit is contained in:
isaki001
2025-01-20 08:06:43 -06:00
committed by GitHub
parent c68b558ed5
commit d89432e8c8
6 changed files with 81 additions and 51 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ext-src/read-allred.patch
+45 -26
View File
@@ -1,8 +1,8 @@
diff --git a/apps/nccl/src/allreduce.hpp b/apps/nccl/src/allreduce.hpp
index 1b85136..ee90c2f 100644
index 4134241..d65be4b 100644
--- a/apps/nccl/src/allreduce.hpp
+++ b/apps/nccl/src/allreduce.hpp
@@ -386,24 +386,353 @@ __global__ void __launch_bounds__(512, 1)
@@ -495,24 +495,348 @@ __global__ void __launch_bounds__(512, 1)
}
}
@@ -311,9 +311,6 @@ index 1b85136..ee90c2f 100644
+ }
+
+}
+
+
+
+
template <typename T>
cudaError_t allreduce(T* buff, T* scratch, T* resultBuff, mscclpp::DeviceHandle<mscclpp::SmChannel>* smChannels,
@@ -326,7 +323,7 @@ index 1b85136..ee90c2f 100644
+ int readAllred = 0, hieAllred = 0;
+ char* envValue = nullptr;
+ char* envValue1 = nullptr;
+
+ nRanksPerNode = (worldSize < nRanksPerNode) ? worldSize : nRanksPerNode;
+
+ envValue = std::getenv("MSCCLPP_READ_ALLRED");
@@ -342,8 +339,6 @@ index 1b85136..ee90c2f 100644
+ hieAllred = 1;
+ }
+ }
+
if (sizeof(T) * nelems < worldSize * sizeof(int)) {
int nBlocks = 7;
int nThreadsPerBlock = 32;
@@ -361,8 +356,8 @@ index 1b85136..ee90c2f 100644
+ nBlocks = 8*(nRanksPerNode - 1);
nThreadsPerBlock = (nelems <= 76800) ? 512 : 1024;
}
allreduce7<<<nBlocks, nThreadsPerBlock, 0, stream>>>(buff, scratch, resultBuff, smChannels, channelInOffset,
@@ -412,9 +741,21 @@ cudaError_t allreduce(T* buff, T* scratch, T* resultBuff, mscclpp::DeviceHandle<
#if defined(ENABLE_NPKIT)
@@ -528,9 +852,21 @@ cudaError_t allreduce(T* buff, T* scratch, T* resultBuff, mscclpp::DeviceHandle<
} else {
int nBlocks = 35;
int nThreadsPerBlock = 512;
@@ -388,24 +383,25 @@ index 1b85136..ee90c2f 100644
return cudaGetLastError();
diff --git a/apps/nccl/src/common.hpp b/apps/nccl/src/common.hpp
index 25c74e7..5e85468 100644
index 015e0a2..f8ba6d6 100644
--- a/apps/nccl/src/common.hpp
+++ b/apps/nccl/src/common.hpp
@@ -11,7 +11,9 @@
@@ -13,8 +13,10 @@
#define WARP_SIZE 32
#endif
+constexpr int NRANKS1_PER_NODE = 4;
constexpr int NRANKS_PER_NODE = 8;
constexpr int SCRATCH_SIZE = 2 * 1024 * 1024 * 70; // double buffer * 35 thread-blocks * 8 ranks * 256KB = 70MB
constexpr int NPEERS = 7;
+constexpr int NPEER = 7;
#endif // NCCL_COMMON_HPP_
constexpr int SCRATCH_SIZE = 2 * 1024 * 1024 * 70; // double buffer * 35 thread-blocks * 8 ranks * 256KB = 70MB
diff --git a/apps/nccl/src/nccl.cu b/apps/nccl/src/nccl.cu
index cb0e7d5..a697be2 100644
index f91d15e..022d398 100644
--- a/apps/nccl/src/nccl.cu
+++ b/apps/nccl/src/nccl.cu
@@ -49,7 +49,9 @@ struct hash<channelKey> {
@@ -70,7 +70,9 @@ struct hash<channelKey> {
struct ChannelInfo {
std::vector<mscclpp::SmChannel> smChannels;
@@ -415,18 +411,15 @@ index cb0e7d5..a697be2 100644
};
struct ncclComm {
@@ -212,8 +214,10 @@ static ncclResult_t ncclAllReduceFallback(const void* sendbuff, void* recvbuff,
int rank = comm->comm->bootstrap()->getRank();
channelKey sendKey{(void*)sendBasePtr, sendBytes};
@@ -213,6 +215,7 @@ static ncclResult_t ncclAllReduceFallback(const void* sendbuff, void* recvbuff,
channelKey recvKey{(void*)recvBasePtr, recvBytes};
+
mscclpp::DeviceHandle<mscclpp::SmChannel>* smChannels = nullptr;
mscclpp::DeviceHandle<mscclpp::SmChannel>* smOutChannels = nullptr;
+ mscclpp::DeviceHandle<mscclpp::SmChannel>* smScrChannels = nullptr;
// Creating the channels
if (count * ncclTypeSize(datatype) <= comm->largeMessageSizeBoundary) {
@@ -221,19 +225,25 @@ static ncclResult_t ncclAllReduceFallback(const void* sendbuff, void* recvbuff,
if (count * ncclTypeSize(datatype) <= (1 << 20)) {
@@ -220,19 +223,24 @@ static ncclResult_t ncclAllReduceFallback(const void* sendbuff, void* recvbuff,
if (sendIt == comm->channelScratchInfos.end()) {
std::vector<mscclpp::SmChannel> channels =
setupSmChannels(comm, comm->remoteScratchRegMemories, const_cast<void*>((void*)sendBasePtr));
@@ -438,8 +431,8 @@ index cb0e7d5..a697be2 100644
smChannels = sendIt->second.smChannelDeviceHandles.get();
} else {
std::vector<mscclpp::RegisteredMemory> remoteMemories;
-
+ std::vector<mscclpp::RegisteredMemory> remoteMemories1;
auto sendIt = comm->channelInInfos.find(sendKey);
if (sendIt == comm->channelInInfos.end()) {
std::vector<mscclpp::SmChannel> channels =
@@ -454,7 +447,7 @@ index cb0e7d5..a697be2 100644
sendIt = comm->channelInInfos.emplace(sendKey, channelInfo).first;
}
@@ -243,35 +253,36 @@ static ncclResult_t ncclAllReduceFallback(const void* sendbuff, void* recvbuff,
@@ -242,35 +250,36 @@ static ncclResult_t ncclAllReduceFallback(const void* sendbuff, void* recvbuff,
setupRemoteMemories(comm->comm, rank, (void*)recvBasePtr, recvBytes, mscclpp::Transport::CudaIpc);
std::vector<mscclpp::SmChannel> outChannels =
setupSmChannels(comm, remoteMemories, const_cast<void*>((void*)recvBasePtr));
@@ -501,8 +494,8 @@ index cb0e7d5..a697be2 100644
+ NRANKS_PER_NODE, comm->comm->bootstrap()->getNranks(), count, stream));
break;
default:
return ncclInvalidArgument;
@@ -550,7 +561,7 @@ NCCL_API ncclResult_t ncclAllGather(const void* sendbuff, void* recvbuff, size_t
WARN("datatype is invalid");
@@ -315,7 +324,7 @@ static ncclResult_t ncclAllGatherFallback(const void* sendbuff, void* recvbuff,
std::vector<mscclpp::DeviceHandle<mscclpp::SmChannel>> smChannelDeviceHandles;
std::transform(channels.begin(), channels.end(), std::back_inserter(smChannelDeviceHandles),
[](const mscclpp::SmChannel& smChannel) { return mscclpp::deviceHandle(smChannel); });
@@ -511,3 +504,29 @@ index cb0e7d5..a697be2 100644
it = comm->channelOutInfos.emplace(recvKey, channelInfo).first;
}
@@ -597,7 +606,7 @@ NCCL_API ncclResult_t ncclBroadcastFallback(const void* sendbuff, void* recvbuff
std::vector<mscclpp::DeviceHandle<mscclpp::SmChannel>> smChannelDeviceHandles;
std::transform(channels.begin(), channels.end(), std::back_inserter(smChannelDeviceHandles),
[](const mscclpp::SmChannel& smChannel) { return mscclpp::deviceHandle(smChannel); });
- ChannelInfo channelInfo{channels, setupSmChannelDeviceHandles(channels)};
+ ChannelInfo channelInfo{channels, channels, setupSmChannelDeviceHandles(channels), setupSmChannelDeviceHandles(channels)};
it = comm->channelOutInfos.emplace(recvKey, channelInfo).first;
}
@@ -805,16 +814,6 @@ NCCL_API ncclResult_t ncclGroupEnd() {
return ncclSuccess;
}
-NCCL_API ncclResult_t ncclCommRegister(const ncclComm_t, void*, size_t, void**) {
- // TODO: Implementation
- return ncclSuccess;
-}
-
-NCCL_API ncclResult_t ncclCommDeregister(const ncclComm_t, void*) {
- // TODO: Implementation
- return ncclSuccess;
-}
-
ncclResult_t ncclMemAlloc(void** ptr, size_t size) {
if (ptr == nullptr || size == 0) {
WARN("ptr is nullptr or size is 0");