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

Add 'projects/rccl/' from commit '1f2f9f33bac3e8ecfd84c69af6063d7352c362fc'

Browse Source 
git-subtree-dir: projects/rccl
git-subtree-mainline: 3fd8a0d393
git-subtree-split: 1f2f9f33ba
This commit is contained in:
Ameya Keshava Mallya
2025-12-11 20:46:05 +00:00
parent 3fd8a0d393 1f2f9f33ba
commit 42d84317cf
682 changed files with 293338 additions and 0 deletions
Download Patch File Download Diff File Expand all files Collapse all files
projects/rccl/ext-src/remove-clip.patch
+54
View File
@@ -0,0 +1,54 @@
diff --git a/apps/nccl/src/allreduce.hpp b/apps/nccl/src/allreduce.hpp
index fac105a..9ef93ce 100644
--- a/apps/nccl/src/allreduce.hpp
+++ b/apps/nccl/src/allreduce.hpp
@@ -71,17 +71,29 @@ __forceinline__ __device__ __bfloat162 clip(__bfloat162 val) {
template <typename T>
__forceinline__ __device__ T add_elements(T a, T b) {
- return clip(a + b);
+ #ifdef MSCCLPP_CLIP_ENABLED
+ return clip(a + b);
+ #else
+ return a + b;
+ #endif
}
template <>
__forceinline__ __device__ __half2 add_elements(__half2 a, __half2 b) {
- return clip(__hadd2(a, b));
+ #ifdef MSCCLPP_CLIP_ENABLED
+ return clip(__hadd2(a, b));
+ #else
+ return __hadd2(a, b);
+ #endif
}
template <>
__forceinline__ __device__ __bfloat162 add_elements(__bfloat162 a, __bfloat162 b) {
- return clip(__hadd2(a, b));
+ #ifdef MSCCLPP_CLIP_ENABLED
+ return clip(__hadd2(a, b));
+ #else
+ return __hadd2(a, b);
+ #endif
}
template <typename T>
@@ -558,7 +570,7 @@ __global__ void __launch_bounds__(512, 1)
template <typename T>
-__global__ void __launch_bounds__(512, 1)
+__global__ void __launch_bounds__(1024, 1)
allreduce8(T* buff, T* scratch, T* resultBuff, mscclpp::DeviceHandle<mscclpp::SmChannel>* smChannels,
mscclpp::DeviceHandle<mscclpp::SmChannel>* smOutChannels, size_t channelOutDataOffset,
size_t channelScratchOffset, int rank, int nRanksPerNode, int worldSize, size_t nelems) {
@@ -1045,6 +1057,7 @@ cudaError_t allreduce(T* buff, T* scratch, T* resultBuff, mscclpp::DeviceHandle<
allreduce8Mod<<<nBlocks, nThreadsPerBlock, 0, stream>>>(buff, scratch, resultBuff, smScrChannels,
channelScratchOffset, rank, nRanksPerNode, worldSize, nelems);
} else {
+ nThreadsPerBlock = std::is_same<T,__bfloat16>::value ? 1024 : nThreadsPerBlock;
allreduce8<<<nBlocks, nThreadsPerBlock, 0, stream>>>(buff, scratch, resultBuff, smScrChannels,
smOutChannels, channelOutOffset, channelScratchOffset, rank, nRanksPerNode,
worldSize, nelems);