diff --git a/projects/rccl/CMakeLists.txt b/projects/rccl/CMakeLists.txt index 1fa904cbbf..1c8ee6e26a 100644 --- a/projects/rccl/CMakeLists.txt +++ b/projects/rccl/CMakeLists.txt @@ -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) diff --git a/projects/rccl/cmake/MSCCLPP.cmake b/projects/rccl/cmake/MSCCLPP.cmake index 4ddd62422c..ef9beb2d43 100644 --- a/projects/rccl/cmake/MSCCLPP.cmake +++ b/projects/rccl/cmake/MSCCLPP.cmake @@ -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= -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= -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} diff --git a/projects/rccl/ext-src/remove-clip.patch b/projects/rccl/ext-src/remove-clip.patch new file mode 100644 index 0000000000..902881ffdf --- /dev/null +++ b/projects/rccl/ext-src/remove-clip.patch @@ -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 + __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 +@@ -558,7 +570,7 @@ __global__ void __launch_bounds__(512, 1) + + + template +-__global__ void __launch_bounds__(512, 1) ++__global__ void __launch_bounds__(1024, 1) + allreduce8(T* buff, T* scratch, T* resultBuff, mscclpp::DeviceHandle* smChannels, + mscclpp::DeviceHandle* 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<<>>(buff, scratch, resultBuff, smScrChannels, + channelScratchOffset, rank, nRanksPerNode, worldSize, nelems); + } else { ++ nThreadsPerBlock = std::is_same::value ? 1024 : nThreadsPerBlock; + allreduce8<<>>(buff, scratch, resultBuff, smScrChannels, + smOutChannels, channelOutOffset, channelScratchOffset, rank, nRanksPerNode, + worldSize, nelems); diff --git a/projects/rccl/install.sh b/projects/rccl/install.sh index fd1a780437..b72c2eaf97 100755 --- a/projects/rccl/install.sh +++ b/projects/rccl/install.sh @@ -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"