AI/rocm-systems
2
0
フォーク 0
コード イシュー プルリクエスト Actions パッケージ プロジェクト リリース Wiki アクティビティ

Add Compilation Flag for enabling/disabling clipping, and tune number of blocks for mscclpp allreduce8 (#1607)

ソースを参照 
* mscclpp patch apply clip patch and set allreduce8 blocks from 512 to 1024

* add compilation flag for enabling/disabling clipping in mscclpp

* change flag name for consistency, set flag to OFF

* add compilation flag in rccl for enabling clipping in mscclpp

* set 1024 threads for mscclpp allreduce8 only for bfloat16

* fix improper description for ENABLE_MSCCLPP_CLIP flag

* Revert "Merge branch 'clip-patch' of https://github.com/isaki001/rccl into clip-patch"

This reverts commit 6e31857a9db98314b8a748eb024f2c3699ebe2d5, reversing
changes made to 193f4caa8ffa78b4e056893212fd8344aa14e937.

* update clip remove-clip.patch for rebase

[ROCm/rccl commit: 8145c4f3b8]
このコミットが含まれているのは:
isaki001
2025-04-30 16:42:28 -05:00
committed by GitHub
親 33c390b36e
コミット de76d7f649
4個のファイルの変更、78行の追加、5行の削除
Patchファイルをダウンロード Diffファイルをダウンロード すべてのファイルを展開 すべてのファイルを折り畳む
projects/rccl/CMakeLists.txt
+1
ファイルの表示
@@ -27,6 +27,7 @@ option(BUILD_TESTS "Build unit test programs"
option(COLLTRACE "Collective Trace Option" ON)
option(ENABLE_MSCCL_KERNEL "Enable MSCCL while compiling" ON)
option(ENABLE_MSCCLPP "Enable MSCCL++" ON)
option(ENABLE_MSCCLPP_CLIP "Enable MSCCL++" OFF)
option(ENABLE_IFC "Enable indirect function call" OFF)
option(INSTALL_DEPENDENCIES "Force install dependencies" OFF)
option(ROCTX "Enable ROCTX" ON)
projects/rccl/cmake/MSCCLPP.cmake
+15 -4
ファイルの表示
@@ -98,6 +98,11 @@ if(ENABLE_MSCCLPP)
WORKING_DIRECTORY ${MSCCLPP_SOURCE}
)
execute_process(
COMMAND git apply ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/remove-clip.patch
WORKING_DIRECTORY ${MSCCLPP_SOURCE}
)
set(CMAKE_INHERITED_ARGS "")
set(CMAKE_ARGS_LIST "CMAKE_PREFIX_PATH;CMAKE_INSTALL_RPATH_USE_LINK_PATH;HIP_COMPILER")
foreach(arg IN LISTS CMAKE_ARGS_LIST)
@@ -130,8 +135,7 @@ if(ENABLE_MSCCLPP)
#GIT_TAG 4ee15b7ad085daaf74349d4c49c9b8480d28f0dc
INSTALL_DIR ${MSCCLPP_ROOT}
LIST_SEPARATOR %
CMAKE_ARGS "-DGPU_TARGETS=${MSCCLPP_GPU_TARGETS}" -DMSCCLPP_BYPASS_GPU_CHECK=ON -DMSCCLPP_USE_ROCM=ON -DCMAKE_BUILD_TYPE=${CMAKE_BUILD_TYPE} -DMSCCLPP_BUILD_APPS_NCCL=ON -DMSCCLPP_BUILD_PYTHON_BINDINGS=OFF -DMSCCLPP_BUILD_TESTS=OFF -DCMAKE_INSTALL_PREFIX=<INSTALL_DIR> -DCMAKE_VERBOSE_MAKEFILE=1 "${CMAKE_INHERITED_ARGS}" -DFETCHCONTENT_SOURCE_DIR_JSON=${JSON_SOURCE}
LOG_DOWNLOAD FALSE
CMAKE_ARGS "-DGPU_TARGETS=${MSCCLPP_GPU_TARGETS}" -DMSCCLPP_BYPASS_GPU_CHECK=ON -DMSCCLPP_USE_ROCM=ON -DCMAKE_BUILD_TYPE=${CMAKE_BUILD_TYPE} -DMSCCLPP_BUILD_APPS_NCCL=ON -DMSCCLPP_BUILD_PYTHON_BINDINGS=OFF -DMSCCLPP_BUILD_TESTS=OFF -DMSCCLPP_CLIP_ENABLED=${ENABLE_MSCCLPP_CLIP} -DCMAKE_INSTALL_PREFIX=<INSTALL_DIR> -DCMAKE_VERBOSE_MAKEFILE=1 "${CMAKE_INHERITED_ARGS}" -DFETCHCONTENT_SOURCE_DIR_JSON=${JSON_SOURCE} LOG_DOWNLOAD FALSE
LOG_CONFIGURE FALSE
LOG_BUILD FALSE
LOG_INSTALL FALSE
@@ -141,16 +145,23 @@ if(ENABLE_MSCCLPP)
find_package(mscclpp_nccl REQUIRED)
execute_process(
COMMAND git apply --reverse ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/device-flag.patch
execute_process(
COMMAND git apply --reverse ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/remove-clip.patch
WORKING_DIRECTORY ${MSCCLPP_SOURCE}
)
execute_process(
COMMAND git apply --reverse ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/no-cache.patch
WORKING_DIRECTORY ${MSCCLPP_SOURCE}
)
execute_process(
COMMAND git apply --reverse ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/device-flag.patch
WORKING_DIRECTORY ${MSCCLPP_SOURCE}
)
execute_process(
COMMAND git apply --reverse ${CMAKE_CURRENT_SOURCE_DIR}/ext-src/reg-fix.patch
WORKING_DIRECTORY ${MSCCLPP_SOURCE}
projects/rccl/ext-src/remove-clip.patch
+54
ファイルの表示
@@ -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);
projects/rccl/install.sh
+8 -1
ファイルの表示
@@ -26,6 +26,7 @@ install_prefix="${ROCM_PATH}"
log_trace=false
msccl_kernel_enabled=true
mscclpp_enabled=true
enable_mscclpp_clip=false
num_parallel_jobs=$(nproc)
npkit_enabled=false
openmp_test_enabled=false
@@ -48,6 +49,7 @@ function display_help()
echo " --disable-colltrace Build without collective trace"
echo " --disable-msccl-kernel Build without MSCCL kernels"
echo " --disable-mscclpp Build without MSCCL++ support"
echo " --enable-mscclpp-clip Build MSCCL++ with clip wrapper on bfloat16 and half addition routines"
echo " --disable-roctx Build without ROCTX logging"
echo " -f|--fast Quick-build RCCL (local gpu arch only, no backtrace, and collective trace support)"
echo " -h|--help Prints this help message"
@@ -77,7 +79,7 @@ function display_help()
# check if we have a modern version of getopt that can handle whitespace and long parameters
getopt -T
if [[ "$?" -eq 4 ]]; then
GETOPT_PARSE=$(getopt --name "${0}" --options dfhij:lprt --longoptions address-sanitizer,dependencies,debug,enable_backtrace,disable-colltrace,disable-msccl-kernel,disable-mscclpp,fast,help,install,jobs:,local_gpu_only,amdgpu_targets:,no_clean,npkit-enable,log-trace,openmp-test-enable,roctx-enable,package_build,prefix:,rm-legacy-include-dir,run_tests_all,run_tests_quick,static,tests_build,time-trace,verbose -- "$@")
GETOPT_PARSE=$(getopt --name "${0}" --options dfhij:lprt --longoptions address-sanitizer,dependencies,debug,enable_backtrace,disable-colltrace,disable-msccl-kernel,disable-mscclpp,enable-mscclpp-clip,fast,help,install,jobs:,local_gpu_only,amdgpu_targets:,no_clean,npkit-enable,log-trace,openmp-test-enable,roctx-enable,package_build,prefix:,rm-legacy-include-dir,run_tests_all,run_tests_quick,static,tests_build,time-trace,verbose -- "$@")
else
echo "Need a new version of getopt"
exit 1
@@ -99,6 +101,7 @@ while true; do
--disable-colltrace) collective_trace=false; shift ;;
--disable-msccl-kernel) msccl_kernel_enabled=false; shift ;;
--disable-mscclpp) mscclpp_enabled=false; shift ;;
--enable-mscclpp-clip) enable_mscclpp_clip=true; shift ;;
--disable-roctx) roctx_enabled=false; shift ;;
-f | --fast) build_local_gpu_only=true; collective_trace=false; msccl_kernel_enabled=false; shift ;;
-h | --help) display_help; exit 0 ;;
@@ -244,6 +247,10 @@ if [[ "${mscclpp_enabled}" == false ]]; then
cmake_common_options="${cmake_common_options} -DENABLE_MSCCLPP=OFF"
fi
if [[ "${enable_mscclpp_clip}" == true ]]; then
cmake_common_options="${cmake_common_options} -DENABLE_MSCCLPP_CLIP=ON"
fi
# Install dependencies
if [[ "${install_dependencies}" == true ]]; then
cmake_common_options="${cmake_common_options} -DINSTALL_DEPENDENCIES=ON"