diff --git a/CMakeLists.txt b/CMakeLists.txt index d0a316ca69..90786f8600 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -522,7 +522,7 @@ list(APPEND HIP_SOURCES ${CMAKE_CURRENT_BINARY_DIR}/git_version.cpp) # Create a custom target that updates git_version.cpp and executes whenever rccl is built add_custom_target(git_version_check COMMENT "Updating git_version.cpp if necessary" - COMMAND ${CMAKE_COMMAND} -P ${CMAKE_CURRENT_SOURCE_DIR}/cmake/git_version.cmake + COMMAND ${CMAKE_COMMAND} -P ${CMAKE_CURRENT_SOURCE_DIR}/cmake/scripts/git_version.cmake VERBATIM ) @@ -602,6 +602,7 @@ if(DEMANGLE_DIR) target_compile_definitions(rccl PRIVATE "HAVE_DECL_BASENAME=1") endif() if(${hipcc_version_string} VERSION_GREATER_EQUAL "6.1.33591") + set(LL128_ENABLED ON) target_compile_definitions(rccl PRIVATE ENABLE_LL128) message(STATUS "RCCL LL128 protocol enabled") endif() @@ -749,6 +750,12 @@ if(BUILD_TESTS) rocm_package_setup_component(clients) rocm_package_setup_client_component(tests PACKAGE_NAME unittests) add_subdirectory(test) + + add_custom_command(TARGET rccl POST_BUILD + COMMENT "Extracting metadata from librccl.so" + COMMAND COMMAND ${CMAKE_COMMAND} -P ${CMAKE_CURRENT_SOURCE_DIR}/cmake/scripts/extract_metadata.cmake + VERBATIM + ) endif() rocm_create_package( diff --git a/cmake/scripts/extract_metadata.cmake b/cmake/scripts/extract_metadata.cmake new file mode 100644 index 0000000000..ca025ab033 --- /dev/null +++ b/cmake/scripts/extract_metadata.cmake @@ -0,0 +1,55 @@ +# Copyright (c) 2024 Advanced Micro Devices, Inc. All rights reserved. +# +# Permission is hereby granted, free of charge, to any person obtaining a copy +# of this software and associated documentation files (the "Software"), to deal +# in the Software without restriction, including without limitation the rights +# to use, copy, modify, merge, publish, distribute, sublicense, and/or sell +# copies of the Software, and to permit persons to whom the Software is +# furnished to do so, subject to the following conditions: +# +# The above copyright notice and this permission notice shall be included in all +# copies or substantial portions of the Software. +# +# THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR +# IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY, +# FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE +# AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER +# LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM, +# OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE +# SOFTWARE. + +## List the objects for each gfx architecture +execute_process( COMMAND roc-obj-ls librccl.so + RESULT_VARIABLE list_result + OUTPUT_VARIABLE cmd_output +) + +if(list_result EQUAL 0) + ## Convert cmd output to list of lines + string(REGEX REPLACE "\n$" "" cmd_output "${cmd_output}") + string(REPLACE "\n" ";" cmd_output "${cmd_output}") + + ## Extract file paths for the selected gfx archs + foreach(line ${cmd_output}) + if(line MATCHES "(gfx90a|gfx940|gfx941|gfx942)") + string(REGEX MATCH "\\file://(.*)" file_match ${line}) + if(file_match) + list(APPEND file_paths ${file_match}) + endif() + endif() + endforeach() + + ## Extract objects from files + foreach(file ${file_paths}) + execute_process( + COMMAND roc-obj-extract ${file} + RESULT_VARIABLE extraction_result + ) + if(NOT extraction_result EQUAL 0) + message(WARNING "Could not extract objects from ${file}") + endif() + endforeach() +else() + ## We don't want to stop building unit-tests if this command fails. + message(WARNING "Command failed with error code ${result}") +endif() \ No newline at end of file diff --git a/cmake/git_version.cmake b/cmake/scripts/git_version.cmake similarity index 100% rename from cmake/git_version.cmake rename to cmake/scripts/git_version.cmake diff --git a/src/device/reduce_kernel.h b/src/device/reduce_kernel.h index c04b7683af..ca372a26ba 100644 --- a/src/device/reduce_kernel.h +++ b/src/device/reduce_kernel.h @@ -239,22 +239,7 @@ struct Apply_Reduce, /*EltPerPack=*/4> { SPECIALIZE_REDUCE(FuncMinMax, float, 1, float, fn.isMinNotMax ? fminf(x, y) : fmaxf(x, y)) SPECIALIZE_REDUCE(FuncMinMax, double, 1, double, fn.isMinNotMax ? fmin(x, y) : fmax(x, y)) -#if __CUDA_ARCH__ >= 530 && __CUDA_ARCH__ != 610 - SPECIALIZE_REDUCE(FuncSum, half, 1, half, __hadd(x, y)) - SPECIALIZE_REDUCE(FuncSum, half, 2, half2, __hadd2(x, y)) - SPECIALIZE_REDUCE(FuncProd, half, 1, half, __hmul(x, y)) - SPECIALIZE_REDUCE(FuncProd, half, 2, half2, __hmul2(x, y)) -#else - SPECIALIZE_REDUCE(FuncSum, half, 1, half, __float2half(__half2float(x) + __half2float(y))) - SPECIALIZE_REDUCE(FuncProd, half, 1, half, __float2half(__half2float(x) * __half2float(y))) -#endif - -#if __CUDA_ARCH__ >= 800 - SPECIALIZE_REDUCE(FuncMinMax, half, 1, half, fn.isMinNotMax ? __hmin(x, y) : __hmax(x, y)) - SPECIALIZE_REDUCE(FuncMinMax, half, 2, half2, fn.isMinNotMax ? __hmin2(x, y) : __hmax2(x, y)) -#else - SPECIALIZE_REDUCE(FuncMinMax, half, 1, half, __float2half(fn.isMinNotMax ? fminf(__half2float(x), __half2float(y)) : fmaxf(__half2float(x), __half2float(y)))) -#endif +SPECIALIZE_REDUCE(FuncMinMax, half, 1, half, fn.isMinNotMax ? __hmin(x, y) : __hmax(x, y)) #if defined(RCCL_BFLOAT16) #if __CUDA_ARCH__ >= 800 @@ -374,7 +359,6 @@ struct FuncPreMulSum { template<> struct FuncPreMulSum { using EltType = half; -#if __CUDA_ARCH__ >= 530 && __CUDA_ARCH__ != 610 half2 scalar; __device__ FuncPreMulSum(uint64_t opArg=0) { union { uint64_t u64; half val; }; @@ -382,14 +366,6 @@ struct FuncPreMulSum { scalar.x = val; scalar.y = val; } -#else - float scalar; - __device__ FuncPreMulSum(uint64_t opArg=0) { - union { uint64_t u64; half val; }; - u64 = opArg; - scalar = __half2float(val); - } -#endif }; #if defined(RCCL_BFLOAT16) @@ -439,11 +415,7 @@ template<> struct Apply_PreOp, /*EltPerPack=*/1> { static constexpr bool IsIdentity = false; __device__ static BytePack preOp(FuncPreMulSum fn, BytePack a) { - #if __CUDA_ARCH__ >= 530 && __CUDA_ARCH__ != 610 return toPack(__hmul(fromPack(a), fn.scalar.x)); - #else - return toPack(__float2half(__half2float(fromPack(a)) * fn.scalar)); - #endif } }; #if __CUDA_ARCH__ >= 530 && __CUDA_ARCH__ != 610 diff --git a/src/init_nvtx.cc b/src/init_nvtx.cc index 44face6814..dbd54a51a3 100644 --- a/src/init_nvtx.cc +++ b/src/init_nvtx.cc @@ -11,6 +11,7 @@ static constexpr const nvtxPayloadEnum_t NvtxEnumRedSchema[] = { // Must be called before the first call to any reduction operation. void initNvtxRegisteredEnums() { +#ifndef NVTX_NO_IMPL // Register schemas and strings constexpr const nvtxPayloadEnumAttr_t eAttr { .fieldMask = NVTX_PAYLOAD_ENUM_ATTR_ENTRIES | NVTX_PAYLOAD_ENUM_ATTR_NUM_ENTRIES | @@ -23,4 +24,5 @@ void initNvtxRegisteredEnums() { }; nvtxPayloadEnumRegister(nvtx3::domain::get(), &eAttr); +#endif } diff --git a/test/CMakeLists.txt b/test/CMakeLists.txt index 5108e4dce4..48e57ac0ee 100644 --- a/test/CMakeLists.txt +++ b/test/CMakeLists.txt @@ -68,9 +68,19 @@ if(BUILD_TESTS) endif() add_executable(rccl-UnitTests ${COMMON_SOURCE_FILES} ${TEST_SOURCE_FILES}) + + ## Set rccl-UnitTests include directories target_include_directories(rccl-UnitTests PRIVATE ${ROCM_PATH} ${GTEST_INCLUDE_DIRS}) - target_include_directories(rccl-UnitTests PRIVATE ${PROJECT_BINARY_DIR}/include) # for generated rccl.h header + target_include_directories(rccl-UnitTests PRIVATE ${PROJECT_BINARY_DIR}/include) # for generated rccl.h header target_include_directories(rccl-UnitTests PRIVATE ${PROJECT_BINARY_DIR}/hipify/src/include) # for rccl_bfloat16.h + + ## Set rccl-UnitTests compile definitions + if(LL128_ENABLED) + target_compile_definitions(rccl-UnitTests PRIVATE ENABLE_LL128) + endif() + target_compile_definitions(rccl-UnitTests PRIVATE ROCM_PATH="${ROCM_PATH}") + + ## Set rccl-UnitTests linked libraries target_link_libraries(rccl-UnitTests PRIVATE ${GTEST_BOTH_LIBRARIES}) target_link_libraries(rccl-UnitTests PRIVATE hip::host hip::device hsa-runtime64::hsa-runtime64) target_link_libraries(rccl-UnitTests PRIVATE Threads::Threads) diff --git a/test/StandaloneTests.cpp b/test/StandaloneTests.cpp index 4db1dcacf7..887c15204d 100644 --- a/test/StandaloneTests.cpp +++ b/test/StandaloneTests.cpp @@ -9,7 +9,11 @@ #include "StandaloneUtils.hpp" -namespace RcclUnitTesting { +namespace RcclUnitTesting +{ + /** + * \brief Verify that each device is assigned to the right rank using ncclCommSplit API. + * ******************************************************************************************/ TEST(Standalone, SplitComms_RankCheck) { // Check for multi-gpu @@ -52,6 +56,9 @@ namespace RcclUnitTesting { NCCLCHECK(ncclCommDestroy(comm)); } + /** + * \brief Creates a communicator for each device and gathers them all in one rank. + * ******************************************************************************************/ TEST(Standalone, SplitComms_OneColor) { // Check for multi-gpu @@ -93,6 +100,9 @@ namespace RcclUnitTesting { NCCLCHECK(ncclCommDestroy(comm)); } + /** + * \brief Creates a communicator for each device and reduces them into (numDevices / 2) ranks. + * ******************************************************************************************/ TEST(Standalone, SplitComms_Reduce) { // Check for multi-gpu @@ -140,7 +150,10 @@ namespace RcclUnitTesting { for (auto& comm : comms) NCCLCHECK(ncclCommDestroy(comm)); } - + + /** + * \brief Verify there is no regression in timing for each protocol [LL, LL128, Simple] + * ******************************************************************************************/ TEST(Standalone, RegressionTiming) { // timing @@ -241,4 +254,41 @@ namespace RcclUnitTesting { else unsetenv("NCCL_PROTO"); } + + /** + * \brief Verify rccl generic kernel stack size for each gfx architecture is less than the + * expected MAX_STACK_SIZE. + * ******************************************************************************************/ + TEST(Standalone, StackSize) { + const char* mainKernel = "ncclDevKernel"; + + // Look for the .co files + std::vector coFileList = splitString(executeCommand("find ../ -type f -name \"*.co\""), '\n'); + + // Check if the .co files exist in the build directory + if (coFileList.empty()) + GTEST_SKIP() << "Skipping... Could not found required files in the build directory."; + + for (const auto& file : coFileList) { + // Store the output in a list + std::string cmd = std::string(ROCM_PATH) + "/llvm/bin/llvm-readelf --notes " + file; + std::vector metadata = splitString(executeCommand(cmd.c_str()), '\n'); + + // Skip if llvm is not installed + if (metadata.empty()) + GTEST_SKIP() << "Skipping... llvm is not found."; + + // Parse metadata from file and store it for each arch + ArchInfo archInfo = parseMetadata(metadata); + + // iterate over each archs kernels + for (const auto& kernel : archInfo.kernels) { + if (kernel.name.find(mainKernel) != std::string::npos) { + // Kernel stack size should be less than or equal to the maxStackSize value + printf("[ INFO ] Arch: %s Kernel: %s Size: %d\n", archInfo.archName.c_str(), kernel.name.c_str(), kernel.privateSegmentFixedSize); + EXPECT_LE(kernel.privateSegmentFixedSize, archInfo.archName == "gfx90a" ? MAX_STACK_SIZE_gfx90a : MAX_STACK_SIZE); + } + } + } + } } diff --git a/test/common/StandaloneUtils.hpp b/test/common/StandaloneUtils.hpp index 5be85c0c33..7fc63e70df 100644 --- a/test/common/StandaloneUtils.hpp +++ b/test/common/StandaloneUtils.hpp @@ -1,6 +1,10 @@ #ifndef STANDALONE_UTILS_H #define STANDALONE_UTILS_H +#include +#include +#include + #define HIPCALL(cmd) \ do { \ hipError_t error = (cmd); \ @@ -20,4 +24,83 @@ } \ } while(0) +#define MAX_STACK_SIZE 112 + +#ifdef ENABLE_LL128 +#define MAX_STACK_SIZE_gfx90a 288 +#else +#define MAX_STACK_SIZE_gfx90a MAX_STACK_SIZE +#endif + +struct KernelInfo { + std::string name; + int privateSegmentFixedSize = 0; +}; + +struct ArchInfo { + std::string archName; + std::vector kernels; +}; + +std::string executeCommand(const char* cmd) { + std::string result; + FILE* pipe = popen(cmd, "r"); + + if (!pipe) { + std::cerr << "Error executing command: " << cmd << std::endl; + return result; + } + + char buffer[128]; + while (!feof(pipe)) { + if (fgets(buffer, 128, pipe) != NULL) { + result += buffer; + } + } + + pclose(pipe); + return result; +} + +std::vector splitString(const std::string& str, char delimiter) { + std::vector result; + std::istringstream iss(str); + + std::string line; + while(std::getline(iss, line, delimiter)) { + result.push_back(line); + } + + return result; +} + + +ArchInfo parseMetadata(const std::vector& list) { + ArchInfo archInfo; + KernelInfo currKernelInfo; + + std::regex amdhsaTargetRegex("amdhsa.target:\\s+(?:'?)amdgcn-amd-amdhsa--(\\w+)(?:'?)"); + std::regex kernelNameRegex("\\.name:\\s+(\\w+)"); + std::regex privateSegmentSizeRegex("\\.private_segment_fixed_size:\\s+(\\d+)"); + + for (const auto& line : list) { + std::smatch match; + + if (std::regex_search(line, match, amdhsaTargetRegex)) { + archInfo.archName = match[1]; + } else if (std::regex_search(line, match, kernelNameRegex)) { + currKernelInfo.name = match[1]; + } else if (std::regex_search(line, match, privateSegmentSizeRegex)) { + currKernelInfo.privateSegmentFixedSize = std::stoi(match[1]); + } + + if (!currKernelInfo.name.empty() && currKernelInfo.privateSegmentFixedSize != 0) { + archInfo.kernels.push_back(currKernelInfo); + currKernelInfo = {}; // Empty kernelInfo + } + } + + return archInfo; +} + #endif \ No newline at end of file diff --git a/tools/scripts/npkit_trace_generator.py b/tools/scripts/npkit_trace_generator.py index 9ac472ad3c..d09b5b9ac0 100644 --- a/tools/scripts/npkit_trace_generator.py +++ b/tools/scripts/npkit_trace_generator.py @@ -55,7 +55,7 @@ def parse_cpu_event(event_bytes): 'timestamp': int.from_bytes(event_bytes[8:16], byteorder='little', signed=False) } -def parse_gpu_event_file_time(sync_dictionary, npkit_dump_dir, npkit_event_def, rank, buf_idx, gpu_clock_scale, cpu_clock_scale, dictionary_of_stats, warmup_runs=5): +def parse_gpu_event_file_time(sync_dictionary, npkit_dump_dir, npkit_event_def, rank, buf_idx, gpu_clock_scale, cpu_clock_scale, dictionary_of_stats, warmup_runs): gpu_event_file_path = os.path.join(npkit_dump_dir, 'gpu_events_rank_%d_buf_%d' % (rank, buf_idx)) stats_key = 'gpu_rank_%d' % (rank) channel_stats = {} @@ -286,7 +286,7 @@ def parse_cpu_event_file(npkit_dump_dir, npkit_event_def, rank, channel, cpu_clo -def convert_npkit_dump_to_trace(npkit_dump_dir, output_dir, npkit_event_def, gpu_statistics): +def convert_npkit_dump_to_trace(npkit_dump_dir, output_dir, npkit_event_def, gpu_statistics, warmup_runs=0): files_in_dump_dir = next(os.walk(npkit_dump_dir))[2] gpu_event_files = [x for x in files_in_dump_dir if x.startswith('gpu_events_rank_')] cpu_event_files = [x for x in files_in_dump_dir if x.startswith('cpu_events_rank_')] @@ -308,7 +308,7 @@ def convert_npkit_dump_to_trace(npkit_dump_dir, output_dir, npkit_event_def, gpu avg_time = {} number_events=0 for buf_idx in buf_indices: # get the avg time - parse_gpu_event_file_time(sync_dictionary, npkit_dump_dir, npkit_event_def, rank, buf_idx, gpu_clock_scale, cpu_clock_scale, dictionary_of_stats) + parse_gpu_event_file_time(sync_dictionary, npkit_dump_dir, npkit_event_def, rank, buf_idx, gpu_clock_scale, cpu_clock_scale, dictionary_of_stats, warmup_runs) for key in sync_dictionary: avg_time[key] = 0 @@ -317,7 +317,7 @@ def convert_npkit_dump_to_trace(npkit_dump_dir, output_dir, npkit_event_def, gpu avg_time[key] = avg_time[key] + (event['timestamp']/number_events) for buf_idx in buf_indices: - gpu_events = parse_gpu_event_file(avg_time, npkit_dump_dir, npkit_event_def, rank, buf_idx, gpu_clock_scale, cpu_clock_scale, dictionary_of_stats) + gpu_events = parse_gpu_event_file(avg_time, npkit_dump_dir, npkit_event_def, rank, buf_idx, gpu_clock_scale, cpu_clock_scale, dictionary_of_stats, warmup_runs) trace['traceEvents'].extend(gpu_events) @@ -345,9 +345,10 @@ if __name__ == '__main__': parser.add_argument('--npkit_event_header_path', type=str, required=True, help='Path to npkit_event.h.') parser.add_argument('--output_dir', type=str, required=True, help='Path to output directory.') parser.add_argument('--gpu_run_stats', type=bool, nargs='?', const=True, default=False, help="print stats instead.") + parser.add_argument('--warmup_runs', type=int, required=False, default=0, help="amount of warmup_runs on rccl.") args = parser.parse_args() gpu_statistics = False if args.gpu_run_stats is not None: gpu_statistics = args.gpu_run_stats npkit_event_def = parse_npkit_event_header(args.npkit_event_header_path) - convert_npkit_dump_to_trace(args.npkit_dump_dir, args.output_dir, npkit_event_def, gpu_statistics) \ No newline at end of file + convert_npkit_dump_to_trace(args.npkit_dump_dir, args.output_dir, npkit_event_def, gpu_statistics, args.warmup_runs) \ No newline at end of file