From 06b0e1c4e2538df0ec695420b4a0492e61004f5c Mon Sep 17 00:00:00 2001 From: gilbertlee-amd <44450918+gilbertlee-amd@users.noreply.github.com> Date: Tue, 7 Sep 2021 15:28:16 -0600 Subject: [PATCH 1/6] [TransferBench] ConfigFile parsing fixes, adding additional info (#422) * [TransferBench] Adding GPU to NUMA distance detection, parsing fixes, config file generation fix * [TransferBench] Fixing up NUMA node detection by filtering pools [ROCm/rccl commit: 51d64894ff0733d2e83121bf88839c6079940f70] --- .../TransferBench/GetClosestNumaNode.hpp | 143 ++++++++++++++++++ projects/rccl/tools/TransferBench/Makefile | 2 +- .../tools/TransferBench/TransferBench.cpp | 29 +++- 3 files changed, 165 insertions(+), 9 deletions(-) create mode 100644 projects/rccl/tools/TransferBench/GetClosestNumaNode.hpp diff --git a/projects/rccl/tools/TransferBench/GetClosestNumaNode.hpp b/projects/rccl/tools/TransferBench/GetClosestNumaNode.hpp new file mode 100644 index 0000000000..40bfda07a5 --- /dev/null +++ b/projects/rccl/tools/TransferBench/GetClosestNumaNode.hpp @@ -0,0 +1,143 @@ +/* +Copyright (c) 2021 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. +*/ + +// Helper macro for checking HSA calls +#define HSA_CHECK(cmd) \ + do { \ + hsa_status_t error = (cmd); \ + if (error != HSA_STATUS_SUCCESS) { \ + const char* errString = NULL; \ + hsa_status_string(error, &errString); \ + std::cerr << "Encountered HSA error (" << errString << ") at line " \ + << __LINE__ << " in file " << __FILE__ << "\n"; \ + exit(-1); \ + } \ + } while (0) + +// Structure to hold HSA agent information +struct AgentData +{ + bool isInitialized; + std::vector cpuAgents; + std::vector gpuAgents; + std::vector closestNumaNode; +}; + +// Simple callback function to return any memory pool for an agent +hsa_status_t MemPoolInfoCallback(hsa_amd_memory_pool_t pool, void *data) +{ + hsa_amd_memory_pool_t* poolData = reinterpret_cast(data); + + // Check memory pool flags + uint32_t poolFlags; + HSA_CHECK(hsa_amd_memory_pool_get_info(pool, HSA_AMD_MEMORY_POOL_INFO_GLOBAL_FLAGS, &poolFlags)); + + // Only consider coarse-grained pools + if (!(poolFlags & HSA_AMD_MEMORY_POOL_GLOBAL_FLAG_COARSE_GRAINED)) return HSA_STATUS_SUCCESS; + + *poolData = pool; + return HSA_STATUS_SUCCESS; +} + +// Callback function to gather HSA agent information +hsa_status_t AgentInfoCallback(hsa_agent_t agent, void* data) +{ + AgentData* agentData = reinterpret_cast(data); + + // Get the device type + hsa_device_type_t deviceType; + HSA_CHECK(hsa_agent_get_info(agent, HSA_AGENT_INFO_DEVICE, &deviceType)); + if (deviceType == HSA_DEVICE_TYPE_CPU) + agentData->cpuAgents.push_back(agent); + if (deviceType == HSA_DEVICE_TYPE_GPU) + { + agentData->gpuAgents.push_back(agent); + agentData->closestNumaNode.push_back(0); + } + + return HSA_STATUS_SUCCESS; +} + +AgentData& GetAgentData() +{ + static AgentData agentData = {}; + + if (!agentData.isInitialized) + { + agentData.isInitialized = true; + + // Add all detected agents to the list + HSA_CHECK(hsa_iterate_agents(AgentInfoCallback, &agentData)); + + // Loop over each GPU + for (uint32_t i = 0; i < agentData.gpuAgents.size(); i++) + { + // Collect memory pool + hsa_amd_memory_pool_t pool; + HSA_CHECK(hsa_amd_agent_iterate_memory_pools(agentData.gpuAgents[i], MemPoolInfoCallback, &pool)); + + // Loop over each CPU agent and check distance + int bestDistance = -1; + for (uint32_t j = 0; j < agentData.cpuAgents.size(); j++) + { + // Determine number of hops from GPU memory pool to CPU agent + uint32_t hops = 0; + HSA_CHECK(hsa_amd_agent_memory_pool_get_info(agentData.cpuAgents[j], + pool, + HSA_AMD_AGENT_MEMORY_POOL_INFO_NUM_LINK_HOPS, + &hops)); + // Gather link info + hsa_amd_memory_pool_link_info_t* link_info = + (hsa_amd_memory_pool_link_info_t *)malloc(hops * sizeof(hsa_amd_memory_pool_link_info_t)); + HSA_CHECK(hsa_amd_agent_memory_pool_get_info(agentData.cpuAgents[j], + pool, + HSA_AMD_AGENT_MEMORY_POOL_INFO_LINK_INFO, + link_info)); + int numaDist = 0; + for (int k = 0; k < hops; k++) + { + numaDist += link_info[k].numa_distance; + } + if (bestDistance == -1 || numaDist < bestDistance) + { + agentData.closestNumaNode[i] = j; + bestDistance = numaDist; + } + free(link_info); + } + } + } + return agentData; +} + +// Returns closest CPU NUMA node to provided GPU +// NOTE: This assumes HSA GPU indexing is similar to HIP GPU indexing +int GetClosestNumaNode(int gpuIdx) +{ + AgentData& agentData = GetAgentData(); + if (gpuIdx < 0 || gpuIdx >= agentData.closestNumaNode.size()) + { + printf("[ERROR] GPU index out is out of bounds\n"); + exit(1); + } + return agentData.closestNumaNode[gpuIdx]; +} diff --git a/projects/rccl/tools/TransferBench/Makefile b/projects/rccl/tools/TransferBench/Makefile index c51bd47b4c..62d1046fb7 100644 --- a/projects/rccl/tools/TransferBench/Makefile +++ b/projects/rccl/tools/TransferBench/Makefile @@ -6,7 +6,7 @@ endif HIPCC=$(HIP_PATH)/bin/hipcc EXE=TransferBench -CXXFLAGS = -O3 -I../../src/include -I. -lnuma +CXXFLAGS = -O3 -I../../src/include -I. -lnuma -L$(HIP_PATH)/../hsa/lib -lhsa-runtime64 all: $(EXE) diff --git a/projects/rccl/tools/TransferBench/TransferBench.cpp b/projects/rccl/tools/TransferBench/TransferBench.cpp index 899c71fc55..6fac7df90a 100644 --- a/projects/rccl/tools/TransferBench/TransferBench.cpp +++ b/projects/rccl/tools/TransferBench/TransferBench.cpp @@ -24,6 +24,7 @@ THE SOFTWARE. // on the same node #include "TransferBench.hpp" +#include "GetClosestNumaNode.hpp" #include #include #include @@ -270,10 +271,12 @@ int main(int argc, char **argv) // Report timings totalCpuTime = totalCpuTime / (1.0 * ev.numIterations) * 1000; double totalBandwidthGbs = (numLinks * N * sizeof(float) / 1.0E6) / totalCpuTime; + double maxGpuTime = 0; for (int i = 0; i < numLinks; i++) { double linkDurationMsec = links[i].totalTime / (1.0 * ev.numIterations); double linkBandwidthGbs = (N * sizeof(float) / 1.0E9) / linkDurationMsec * 1000.0f; + maxGpuTime = std::max(maxGpuTime, linkDurationMsec); if (!ev.outputToCsv) { printf(" Link %02d: %c%02d -> [%cPU %02d:%02d] -> %c%02d | %9.3f GB/s | %8.3f ms | %-16s", @@ -310,7 +313,8 @@ int main(int argc, char **argv) // Display aggregate statistics if (!ev.outputToCsv) { - printf(" Aggregate Bandwidth (CPU timed) | %9.3f GB/s | %8.3f ms |\n", totalBandwidthGbs, totalCpuTime); + printf(" Aggregate Bandwidth (CPU timed) | %9.3f GB/s | %8.3f ms | Overhead: %.3f ms\n", totalBandwidthGbs, totalCpuTime, + totalCpuTime - maxGpuTime); } else { @@ -514,7 +518,7 @@ void GenerateConfigFile(char const* cfgFile, int numBlocks) fprintf(fp, "# GPU 0 Gather\n"); fprintf(fp, "%d %d", numGpuDevices-1, numBlocks); for (int i = 1; i < numGpuDevices; i++) - fprintf(fp, " (G%d->G%d->G%d)", 0, i, 0); + fprintf(fp, " (G%d->G%d->G%d)", i, 0, 0); fprintf(fp, "\n\n"); // Full stress test @@ -533,17 +537,16 @@ void GenerateConfigFile(char const* cfgFile, int numBlocks) void DisplayTopology() { - printf("\nDetected topology:\n"); int numGpuDevices; HIP_CALL(hipGetDeviceCount(&numGpuDevices)); - + printf("\nDetected topology: %d CPU NUMA node(s) %d GPU device(s)\n", numa_num_configured_nodes(), numGpuDevices); printf(" |"); for (int j = 0; j < numGpuDevices; j++) printf(" GPU %02d |", j); - printf(" PCIe Bus ID\n"); + printf(" PCIe Bus ID | Closest NUMA\n"); for (int j = 0; j <= numGpuDevices; j++) printf("--------+"); - printf("-------------\n"); + printf("--------------+-------------\n"); char pciBusId[20]; for (int i = 0; i < numGpuDevices; i++) @@ -567,7 +570,7 @@ void DisplayTopology() } } HIP_CALL(hipDeviceGetPCIBusId(pciBusId, 20, i)); - printf(" %s\n", pciBusId); + printf(" %11s | %d \n", pciBusId, GetClosestNumaNode(i)); } } @@ -670,7 +673,7 @@ void ParseLinks(char* line, int numCpus, int numGpus, std::vector& links) // Method 1: Take in triples (srcMem, exeMem, dstMem) int numBlocksToUse; iss >> numBlocksToUse; - if (numBlocksToUse <= 0) + if (numBlocksToUse <= 0 || iss.fail()) { printf("Parsing error: Number of blocks to use (%d) must be greater than 0\n", numBlocksToUse); exit(1); @@ -679,6 +682,11 @@ void ParseLinks(char* line, int numCpus, int numGpus, std::vector& links) for (int i = 0; i < numLinks; i++) { iss >> srcMem >> exeMem >> dstMem; + if (iss.fail()) + { + printf("Parsing error: Unable to read valid Link triplet (possibly missing a SRC or EXE or DST)\n"); + exit(1); + } ParseMemType(srcMem, numCpus, numGpus, &links[i].srcMemType, &links[i].srcIndex); ParseMemType(exeMem, numCpus, numGpus, &links[i].exeMemType, &links[i].exeIndex); ParseMemType(dstMem, numCpus, numGpus, &links[i].dstMemType, &links[i].dstIndex); @@ -699,6 +707,11 @@ void ParseLinks(char* line, int numCpus, int numGpus, std::vector& links) for (int i = 0; i < numLinks; i++) { iss >> srcMem >> exeMem >> dstMem >> links[i].numBlocksToUse; + if (iss.fail()) + { + printf("Parsing error: Unable to read valid Link quadruple (possibly missing a SRC or EXE or DST or #CU)\n"); + exit(1); + } ParseMemType(srcMem, numCpus, numGpus, &links[i].srcMemType, &links[i].srcIndex); ParseMemType(exeMem, numCpus, numGpus, &links[i].exeMemType, &links[i].exeIndex); ParseMemType(dstMem, numCpus, numGpus, &links[i].dstMemType, &links[i].dstIndex); From 4f610a22390f0cc57d9b50380c5c35a61c6d824f Mon Sep 17 00:00:00 2001 From: Wenkai Du Date: Tue, 7 Sep 2021 13:48:59 -0500 Subject: [PATCH 2/6] Revert "rccl-prim-test: add all-to-all benchmark (#185)" This reverts commit e3e1c6b29ca3307e3e285e8d96b1b43e43721019. [ROCm/rccl commit: b22d097524b36eae16df543607ee137fc1fa9642] --- .../rccl/tools/rccl-prim-test/copy_kernel.h | 2 +- .../tools/rccl-prim-test/rccl_prim_test.cpp | 180 ++++++++---------- 2 files changed, 81 insertions(+), 101 deletions(-) diff --git a/projects/rccl/tools/rccl-prim-test/copy_kernel.h b/projects/rccl/tools/rccl-prim-test/copy_kernel.h index 05d96d8406..f14035e8be 100644 --- a/projects/rccl/tools/rccl-prim-test/copy_kernel.h +++ b/projects/rccl/tools/rccl-prim-test/copy_kernel.h @@ -264,4 +264,4 @@ __device__ __forceinline__ void ReduceOrCopyMulti(const int tid, const int nthre ReduceCopyMulti(w, nw, t, nsrcs, srcs, ndsts, dsts, offset, Nrem); } -#endif // COMMON_KERNEL_H_ +#endif \ No newline at end of file diff --git a/projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp b/projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp index 076387f9b9..b050cdf981 100644 --- a/projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp +++ b/projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp @@ -36,13 +36,13 @@ THE SOFTWARE. #define MAX_GPU 8 #define MAX_WORKGROUPS 32 #define THREADS 256 +#define NGPUS 2 #define COPY_UNROLL 4 #define REDUCE_UNROLL 2 #define DOUBLECOPY_UNROLL 2 #define DOUBLECOPYLOCAL_UNROLL 2 #define REDUCECOPY_UNROLL 2 -#define ALL2ALL_UNROLL 2 @@ -59,18 +59,15 @@ THE SOFTWARE. #define RTC_CLOCK_FREQ_DEFAULT 2.7E7 struct transfer_data_t { - // Buffers for all OPs except all to all float *dest0[MAX_WORKGROUPS]; //remote fine grain float *src0[MAX_WORKGROUPS]; //local fine grain float *dest1[MAX_WORKGROUPS]; //local coarse grain float *dest2[MAX_WORKGROUPS]; //local fine grain float *src1[MAX_WORKGROUPS]; //local coarse grain - // Buffers for all to all - const float *srcs[MAX_WORKGROUPS][MAX_GPU]; - float *dsts[MAX_WORKGROUPS][MAX_GPU]; int N; int gpu; int ngpu; + uint64_t *remOpCount; }; struct profiling_data_t { @@ -100,11 +97,10 @@ enum Ops { OP_REDUCE, OP_REDUCECOPY, OP_READ, - OP_ALL2ALL, NUM_OPS, }; -template +template __global__ void flag_sync_kernel(struct transfer_data_t* transfer_data, struct profiling_data_t* profiling_data, uint64_t opCount) { size_t tid = threadIdx.x; uint64_t curr_time; @@ -114,9 +110,19 @@ __global__ void flag_sync_kernel(struct transfer_data_t* transfer_data, struct p const float *srcs[NGPUS]; float *dsts[NGPUS]; + // signal self ready and wait until all GPUs are ready if (tid == 0) { - curr_time = __builtin_amdgcn_s_memrealtime(); + __atomic_fetch_add(&transfer_data->remOpCount[transfer_data->gpu], 1, __ATOMIC_SEQ_CST); + if (sync) { + for (int i = 0; i < transfer_data->ngpu; i++) { + while (LOAD(&transfer_data->remOpCount[i]) < opCount) {}; + } + } } + __syncthreads(); + + if (tid == 0) + curr_time = __builtin_amdgcn_s_memrealtime(); if (op == OP_COPY) { srcs[0] = transfer_data->src0[bid]; @@ -167,36 +173,32 @@ __global__ void flag_sync_kernel(struct transfer_data_t* transfer_data, struct p ReduceOrCopyMulti, float, 1, 1, 1, 1>(threadIdx.x, THREADS, 1, srcs, 1, dsts, n); } - if (op == OP_ALL2ALL) { - for (int i = 0; i < NGPUS; i++) { - srcs[i] = transfer_data->srcs[bid][i]; - dsts[i] = transfer_data->dsts[bid][i]; - } - ReduceOrCopyMulti, float, 1, NGPUS, 1, NGPUS>(tid, THREADS, - NGPUS, srcs, NGPUS, dsts, n); - } __syncthreads(); + if (tid == 0) { __atomic_fetch_add(&(profiling_data->write_cycles[bid]), __builtin_amdgcn_s_memrealtime() - curr_time, __ATOMIC_SEQ_CST); - // for all to all, read and write n itmes to all other GPUs, thus "n * sizeof(float) * (transfer_data->ngpu - 1) * 2" bytes - if (op == OP_ALL2ALL) __atomic_fetch_add(&(profiling_data->bytes_transferred[bid]), n * sizeof(float) * (transfer_data->ngpu - 1) * 2, __ATOMIC_SEQ_CST); - else __atomic_fetch_add(&(profiling_data->bytes_transferred[bid]), n * sizeof(float), __ATOMIC_SEQ_CST); + __atomic_fetch_add(&(profiling_data->bytes_transferred[bid]), n * sizeof(float), __ATOMIC_SEQ_CST); } } typedef void(*flag_sync_kernel_t)(struct transfer_data_t* transfer_data, struct profiling_data_t* profiling_data, uint64_t opCount); -static flag_sync_kernel_t const flagSyncKerns[NUM_OPS+1] = { - flag_sync_kernel, - flag_sync_kernel, - flag_sync_kernel, - flag_sync_kernel, - flag_sync_kernel, - flag_sync_kernel, - flag_sync_kernel, - flag_sync_kernel, - flag_sync_kernel, +static flag_sync_kernel_t const flagSyncKerns[NUM_OPS*2] = { + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, + flag_sync_kernel, }; __global__ void initTestDataKernel(float* data, const size_t N, const int gpu) { @@ -230,11 +232,11 @@ static void setupPeers(uint32_t *info, bool* is_xgmi) { HIPCHECK(hipSetDevice(i)); for (int j = 0; j < deviceCnt; j++) { if (i != j) { - int p2p; + int p2p; HIPCHECK(hipDeviceCanAccessPeer(&p2p, i, j)); if (!p2p) { printf("Cannot enable peer access between device %d and %d. You may use HIP_VISIBLE_DEVICES to limit GPUs.\n", - i, j); + i, j); exit(-1); } HIPCHECK(hipDeviceEnablePeerAccess(j, 0)); @@ -368,7 +370,7 @@ static const char* link_type_name[] = {"HT", "QPI", "PCIE", "IB", "XGMI"}; int main(int argc,char* argv[]) { if (cmdOptionExists(argv, argv + argc, "-h")) { - printf("./rccl_prim_test -w num_workgroups -p copy|localcopy|doublecopy|doublecopylocal|reduce|reducecopy|all2all -i iterations -n bytes -r \"0 1 2 3|3 2 1 0\"\n"); + printf("./rccl_prim_test -w num_workgroups -p copy|localcopy|doublecopy|doublecopylocal|reduce|reducecopy|all -i iterations -n bytes -r \"0 1 2 3|3 2 1 0\"\n"); exit(0); } @@ -391,10 +393,16 @@ int main(int argc,char* argv[]) printf("Benchmarking using %ld bytes\n", nBytes); uint64_t N = nBytes/sizeof(float); + int sync = 1; + char *s = getCmdOption(argv, argv + argc, "-s"); + if (s) + sync = atol(s); + if (sync) printf("Sync all GPUs before operation\n"); + char *r = getCmdOption(argv, argv + argc, "-r"); if (r) printf("User specified ring topology: %s\n", r); - const char *ops[] = {"copy", "localcopy", "doublecopy", "doublecopylocal", "reduce", "reducecopy", "read", "all2all"}; + const char *ops[] = {"copy", "localcopy", "doublecopy", "doublecopylocal", "reduce", "reducecopy", "read", "all"}; char *prim = getCmdOption(argv, argv + argc, "-p"); int op = NUM_OPS, begin_op, end_op; if (prim) { @@ -461,6 +469,10 @@ int main(int argc,char* argv[]) struct profiling_data_t *profiling_data[MAX_GPU], *d_profiling_data[MAX_GPU]; hipStream_t stream[MAX_GPU]; + uint64_t *remOpCount, *d_remOpCount; + HIPCHECK(hipHostMalloc((void**)&remOpCount, sizeof(uint64_t)*MAX_GPU, hipHostMallocMapped)); + HIPCHECK(hipHostGetDevicePointer((void**)&d_remOpCount, (void*)remOpCount, 0)); + // print rings for (int i = 0; i < workgroups; i++) { printRing(i, ring[i], nGpu); @@ -519,16 +531,7 @@ int main(int argc,char* argv[]) h_transfer_data[i].N = N; h_transfer_data[i].gpu = i; h_transfer_data[i].ngpu = nGpu; - } - - for (int i = 0; i < nGpu; i ++) { - for (int j = 0; j < workgroups; j++) { - for (int k = 0; k < nGpu; k++) { - h_transfer_data[i].srcs[j][k] = buff[((i+k)%nGpu)*MAX_WORKGROUPS+j]; - h_transfer_data[i].dsts[j][k] = buff[((i+k)%nGpu)*MAX_WORKGROUPS+j] + N; - //printf("Setup GPU %d bid %d srcs[%d] %p dsts[%d] %p\n", i, j, k, h_transfer_data[i].srcs[j][k], k, h_transfer_data[i].dsts[j][k]); - } - } + h_transfer_data[i].remOpCount = d_remOpCount; } for (int i = 0; i < nGpu; i ++) { @@ -543,13 +546,10 @@ int main(int argc,char* argv[]) hipLaunchParams *launchParamsList= reinterpret_cast( malloc(sizeof(hipLaunchParams)*MAX_GPU)); - uint64_t opCount = 0; + uint64_t opCount = workgroups; for (int op = begin_op; op < end_op; op ++) { - if (op == OP_ALL2ALL && nGpu != 4 && nGpu != 8) { - printf("\n%s only supports 4 or 8 GPUs.\n", ops[op]); - continue; - } - printf("\n[Testing %s]: \n", ops[op]); + const char *OpsName[] = {"Copy", "Local Copy", "Double Copy", "doublecopylocal", "Reduce", "ReduceCopy", "Read"}; + printf("\n[Testing %s]: \n", OpsName[op]); // 4 warm up cycles for (int j = 0; j < 4; j ++) { for (int i = 0; i < nGpu; i ++) { @@ -557,10 +557,8 @@ int main(int argc,char* argv[]) args[i*3] = &transfer_data[i]; args[i*3+1] = &d_profiling_data[i]; args[i*3+2] = &opCount; - if (op == OP_ALL2ALL) - launchParamsList[i].func = reinterpret_cast(flagSyncKerns[op + (nGpu/8)]); - else - launchParamsList[i].func = reinterpret_cast(flagSyncKerns[op]); + launchParamsList[i].func = + reinterpret_cast(flagSyncKerns[op*2 + sync]); launchParamsList[i].gridDim = dim3(workgroups, 1, 1), launchParamsList[i].blockDim = dim3(THREADS, 1, 1), launchParamsList[i].sharedMem = 0; @@ -572,7 +570,7 @@ int main(int argc,char* argv[]) #else HIPCHECK(hipSetDevice(i)); //launch the kernel - hipLaunchKernelGGL(flagSyncKerns[op == OP_ALL2ALL ? op + (nGpu/8) : op], + hipLaunchKernelGGL(flagSyncKerns[op*2 + sync], /*grid dim x,y,z*/ dim3(workgroups, 1, 1), /*block dim x,y,z*/ dim3(THREADS, 1, 1), /*dynamic shared mem*/ 0, @@ -580,7 +578,7 @@ int main(int argc,char* argv[]) /*kernel args*/ transfer_data[i], d_profiling_data[i], opCount); } #endif - opCount++; + opCount+=workgroups; } for (int i = 0; i < nGpu; i ++) { @@ -596,10 +594,8 @@ int main(int argc,char* argv[]) args[i*3] = &transfer_data[i]; args[i*3+1] = &d_profiling_data[i]; args[i*3+2] = &opCount; - if (op == OP_ALL2ALL) - launchParamsList[i].func = reinterpret_cast(flagSyncKerns[op + (nGpu/8)]); - else - launchParamsList[i].func = reinterpret_cast(flagSyncKerns[op]); + launchParamsList[i].func = + reinterpret_cast(flagSyncKerns[op*2 + sync]); launchParamsList[i].gridDim = dim3(workgroups, 1, 1), launchParamsList[i].blockDim = dim3(THREADS, 1, 1), launchParamsList[i].sharedMem = 0; @@ -611,7 +607,7 @@ int main(int argc,char* argv[]) #else HIPCHECK(hipSetDevice(i)); //launch the kernel - hipLaunchKernelGGL(flagSyncKerns[op == OP_ALL2ALL ? op + (nGpu/8) : op], + hipLaunchKernelGGL(flagSyncKerns[op*2 + sync], /*grid dim x,y,z*/ dim3(workgroups, 1, 1), /*block dim x,y,z*/ dim3(THREADS, 1, 1), /*dynamic shared mem*/ 0, @@ -619,7 +615,7 @@ int main(int argc,char* argv[]) /*kernel args*/ transfer_data[i], d_profiling_data[i], opCount); } #endif - opCount++; + opCount+=workgroups; } for (int i = 0; i < nGpu; i ++) { @@ -650,46 +646,24 @@ int main(int argc,char* argv[]) uint32_t hopcount; HIPCHECK(hipExtGetLinkTypeAndHopCount(i, next_gpu , &linktype, &hopcount)); - if (op == OP_ALL2ALL) { - if(prop.gcnArch == 906) { - write_cycle = write_cycle + profiling_data[i]->write_cycles[j]; - bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j]; - double t0 = (double)profiling_data[i]->write_cycles[j]/RTC_CLOCK_FREQ_VEGA20; - fprintf(stderr, "%-20d %-d<->all %-13d %-13s %-13.4f %-20lu %-.2f\n", - i, i, j, link_type_name[linktype], t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9)); - } else if (prop.gcnArch == 908) { - write_cycle = write_cycle + profiling_data[i]->write_cycles[j]; - bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j]; - double t0 = (double)profiling_data[i]->write_cycles[j]/RTC_CLOCK_FREQ_ARCTURUS; - fprintf(stderr, "%-20d %-d<->all %-13d %-13s %-13.4f %-20lu %-.2f\n", - i, i, j, link_type_name[linktype], t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9)); - } else { - write_cycle = write_cycle + profiling_data[i]->write_cycles[j]; - bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j]; - double t0 = (double)profiling_data[i]->write_cycles[j]/RTC_CLOCK_FREQ_DEFAULT; - fprintf(stderr, "%-20d %-d<->all %-13d %-13s %-13.4f %-20lu %-.2f\n", - i, i, j, link_type_name[linktype], t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9)); - } + if(prop.gcnArch == 906) { + write_cycle = write_cycle + profiling_data[i]->write_cycles[j]; + bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j]; + double t0 = (double)profiling_data[i]->write_cycles[j]/((double)RTC_CLOCK_FREQ_VEGA20); + fprintf(stderr, "%-20d %-d->%-10d %-13d %-13s %-13.4f %-20lu %-.2f\n", + i,i, next_gpu,j,link_type_name[linktype],t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9)); + } else if (prop.gcnArch == 908) { + write_cycle = write_cycle + profiling_data[i]->write_cycles[j]; + bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j]; + double t0 = (double)profiling_data[i]->write_cycles[j]/((double)RTC_CLOCK_FREQ_ARCTURUS); + fprintf(stderr, "%-20d %-d->%-10d %-13d %-13s %-13.4f %-20lu %-.2f\n", + i,i, next_gpu,j,link_type_name[linktype],t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9)); } else { - if(prop.gcnArch == 906) { - write_cycle = write_cycle + profiling_data[i]->write_cycles[j]; - bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j]; - double t0 = (double)profiling_data[i]->write_cycles[j]/RTC_CLOCK_FREQ_VEGA20; - fprintf(stderr, "%-20d %-d->%-10d %-13d %-13s %-13.4f %-20lu %-.2f\n", - i, i, next_gpu, j, link_type_name[linktype], t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9)); - } else if (prop.gcnArch == 908) { - write_cycle = write_cycle + profiling_data[i]->write_cycles[j]; - bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j]; - double t0 = (double)profiling_data[i]->write_cycles[j]/RTC_CLOCK_FREQ_ARCTURUS; - fprintf(stderr, "%-20d %-d->%-10d %-13d %-13s %-13.4f %-20lu %-.2f\n", - i, i, next_gpu, j, link_type_name[linktype], t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9)); - } else { - write_cycle = write_cycle + profiling_data[i]->write_cycles[j]; - bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j]; - double t0 = (double)profiling_data[i]->write_cycles[j]/RTC_CLOCK_FREQ_DEFAULT; - fprintf(stderr, "%-20d %-d->%-10d %-13d %-13s %-13.4f %-20lu %-.2f\n", - i, i, next_gpu, j, link_type_name[linktype], t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9)); - } + write_cycle = write_cycle + profiling_data[i]->write_cycles[j]; + bytes_transferred = bytes_transferred + profiling_data[i]->bytes_transferred[j]; + double t0 = (double)profiling_data[i]->write_cycles[j]/((double)RTC_CLOCK_FREQ_DEFAULT); + fprintf(stderr, "%-20d %-d->%-10d %-13d %-13s %-13.4f %-20lu %-.2f\n", + i,i, next_gpu,j,link_type_name[linktype],t0, profiling_data[i]->bytes_transferred[j], (double)profiling_data[i]->bytes_transferred[j]/(t0*1.0E9)); } } print_table_summary_line(); @@ -724,4 +698,10 @@ int main(int argc,char* argv[]) HIPCHECK(hipFree((void*) d_profiling_data[i])); free(profiling_data[i]); } + + printf("opCount: "); + for (int i = 0; i < nGpu; i++) + printf("%ld ", remOpCount[i]); + printf("\n"); + HIPCHECK(hipHostFree((void*)remOpCount)); } From 310d51056ff5f78ccce416d5f571e2d0d7c6acb8 Mon Sep 17 00:00:00 2001 From: Wenkai Du Date: Tue, 7 Sep 2021 11:15:51 -0500 Subject: [PATCH 3/6] rccl-prim-test: enable 8p1h and 16p1h test [ROCm/rccl commit: 7558b5e2bf447da9a9e942dd9d3b7aa2c71ffdd0] --- .../tools/rccl-prim-test/rccl_prim_test.cpp | 56 ++++++++++++++----- 1 file changed, 43 insertions(+), 13 deletions(-) diff --git a/projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp b/projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp index b050cdf981..8c566275cc 100644 --- a/projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp +++ b/projects/rccl/tools/rccl-prim-test/rccl_prim_test.cpp @@ -33,7 +33,7 @@ THE SOFTWARE. #include #include "copy_kernel.h" -#define MAX_GPU 8 +#define MAX_GPU 16 #define MAX_WORKGROUPS 32 #define THREADS 256 #define NGPUS 2 @@ -241,7 +241,9 @@ static void setupPeers(uint32_t *info, bool* is_xgmi) { } HIPCHECK(hipDeviceEnablePeerAccess(j, 0)); uint32_t linktype; - HIPCHECK(hipExtGetLinkTypeAndHopCount(i, j, &linktype, &info[i*deviceCnt+j])); + hipError_t error = hipExtGetLinkTypeAndHopCount(i, j, &linktype, &info[i*deviceCnt+j]); + if (error != hipSuccess) + *is_xgmi = 0; if (linktype != 4 || info[i*deviceCnt+j] != 1) *is_xgmi = 0; } else @@ -268,7 +270,9 @@ static void parseChordalRing(char **str) { int count = 0; for (int n = 0; n= 0 && digit <= 9) { + if (state) + ring[num_rings][j] = ring[num_rings][j]*10 + digit; + else { + ring[num_rings][j] = digit; + state = 1; + } + } + else { + state = 0; + j++; + if (r[n] == ' ') continue; + if (r[n] == '|') { + num_rings ++; + j = 0; + continue; + } } - ring[num_rings][j++] = r[n] - '0'; } while (r[n++] != 0x0); num_rings ++; } else { From d75504e9dc1519e99c712adfb69dfa5092aadccb Mon Sep 17 00:00:00 2001 From: Wenkai Du Date: Wed, 8 Sep 2021 14:20:32 -0500 Subject: [PATCH 4/6] Remove atomic from profiling [ROCm/rccl commit: 31bd4236f1ae59eb1b7b942dc77ed3ef9d54bf1c] --- .../rccl/src/collectives/device/all_reduce.h | 2 +- .../rccl/src/collectives/device/primitives.h | 18 ++++--- projects/rccl/src/include/devcomm.h | 10 ++-- projects/rccl/src/init.cc | 54 ++++++++++++++----- 4 files changed, 60 insertions(+), 24 deletions(-) diff --git a/projects/rccl/src/collectives/device/all_reduce.h b/projects/rccl/src/collectives/device/all_reduce.h index a77b187870..1f54ec2b13 100644 --- a/projects/rccl/src/collectives/device/all_reduce.h +++ b/projects/rccl/src/collectives/device/all_reduce.h @@ -101,7 +101,7 @@ class ncclFunctiontotal_cycle), __builtin_amdgcn_s_memrealtime() - clk, __ATOMIC_SEQ_CST); + if (tid == 0) devProf->elems[blockIdx.x].total_cycle += (__builtin_amdgcn_s_memrealtime() - clk); #endif } }; diff --git a/projects/rccl/src/collectives/device/primitives.h b/projects/rccl/src/collectives/device/primitives.h index 3fc3f45f0f..c64d0bd153 100644 --- a/projects/rccl/src/collectives/device/primitives.h +++ b/projects/rccl/src/collectives/device/primitives.h @@ -146,14 +146,15 @@ class ncclPrimitives { inline __device__ void waitRecv(ssize_t directOffset) { spins = 0; #ifdef ENABLE_PROFILING - uint64_t t0 = __builtin_amdgcn_s_memrealtime(); + uint64_t t0; + if (tid == 0) t0 = __builtin_amdgcn_s_memrealtime(); #endif while (connTailCache < step + SLICESTEPS) { connTailCache = LOAD(connTailPtr); if (checkAbort()) break; } #ifdef ENABLE_PROFILING - if (tid == 0) __atomic_fetch_add(&comm->devProf->wait_recv_cycle[blockIdx.x], __builtin_amdgcn_s_memrealtime() - t0, __ATOMIC_SEQ_CST); + if (tid == 0) comm->devProf->elems[blockIdx.x].wait_recv_cycle += (__builtin_amdgcn_s_memrealtime() - t0); #endif if (connPtrsFifoPtr) srcs[SRC+index] = (const T *)LOAD(connPtrsFifoPtr+step%NCCL_STEPS); else srcs[SRC+index] = directPtr(directOffset); @@ -180,7 +181,8 @@ class ncclPrimitives { for (int slice=0; slice(directOffset+offset, realSize*sizeof(T)); if (realSize > 0) { #ifdef ENABLE_PROFILING - if (tid == 0) __atomic_fetch_add(&comm->devProf->wait_cycle[blockIdx.x], __builtin_amdgcn_s_memrealtime() - t0, __ATOMIC_SEQ_CST); + if (tid == 0) comm->devProf->elems[blockIdx.x].wait_cycle += (__builtin_amdgcn_s_memrealtime() - t0); #endif subBarrier(); ReduceOrCopyMulti(tid, nworkers, RECV*nrecv+SRC, srcs, SEND*nsend+DST, dsts, realSize); @@ -428,12 +430,12 @@ class ncclPrimitives { #ifdef ENABLE_PROFILING #define INIT_COUNTER \ - if (tid == 0) { t0 = __builtin_amdgcn_s_memrealtime(); ws = LOAD(&(devProf->wait_cycle[blockIdx.x])); } + if (tid == 0) { t0 = __builtin_amdgcn_s_memrealtime(); ws = devProf->elems[blockIdx.x].wait_cycle; } #define ACCUMULATE_COUNTER(prim) \ - if (tid == 0) { __atomic_fetch_add(&(devProf->prim##_cycle), __builtin_amdgcn_s_memrealtime() - t0 \ - + ws - LOAD(&(devProf->wait_cycle[blockIdx.x])), __ATOMIC_SEQ_CST); \ - __atomic_fetch_add(&(devProf->prim##_byte), nelem * sizeof(T), __ATOMIC_SEQ_CST); } + if (tid == 0) { devProf->elems[blockIdx.x].prim##_cycle += (__builtin_amdgcn_s_memrealtime() - t0 \ + + ws - devProf->elems[blockIdx.x].wait_cycle); \ + devProf->elems[blockIdx.x].prim##_byte += nelem * sizeof(T); } #else #define INIT_COUNTER #define ACCUMULATE_COUNTER(prim) diff --git a/projects/rccl/src/include/devcomm.h b/projects/rccl/src/include/devcomm.h index 51b57ec790..42976e36c7 100644 --- a/projects/rccl/src/include/devcomm.h +++ b/projects/rccl/src/include/devcomm.h @@ -280,12 +280,12 @@ static_assert(sizeof(struct ncclChannel) == 0x80*sizeof(int), "ncclChannel must #pragma pack(pop) /* restore original alignment from stack */ #ifdef ENABLE_PROFILING -struct ncclProf { +struct ncclProfElem { union { struct { uint64_t total_cycle; - uint64_t wait_cycle[MAXCHANNELS]; // total wait cycle - uint64_t wait_recv_cycle[MAXCHANNELS]; // recv wait cycle + uint64_t wait_cycle; // total wait cycle + uint64_t wait_recv_cycle; // recv wait cycle // primtive cycles uint64_t send_cycle; uint64_t directSend_cycle; @@ -316,6 +316,10 @@ struct ncclProf { int data[0x80]; }; }; + +struct ncclProf { + struct ncclProfElem elems[MAXCHANNELS]; +}; #endif #ifdef ENABLE_COLLTRACE diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc index 73c34219b5..d24b33cb55 100644 --- a/projects/rccl/src/init.cc +++ b/projects/rccl/src/init.cc @@ -298,10 +298,40 @@ static ncclResult_t commFree(ncclComm_t comm) { #ifdef ENABLE_PROFILING struct ncclProf* prof = (struct ncclProf*)malloc(sizeof(struct ncclProf)); CUDACHECK(hipMemcpy(prof, comm->hostDevComm.devProf, sizeof(struct ncclProf), hipMemcpyDeviceToHost)); - uint64_t wait_cycle = 0, wait_recv_cycle = 0; + uint64_t total_cycle = 0, wait_cycle = 0, wait_recv_cycle = 0, send_cycle = 0, directSend_cycle = 0, recv_cycle = 0, \ + directRecv_cycle = 0, copySend_cycle = 0, directCopySend_cycle = 0, recvCopySend_cycle = 0, directRecvCopySend_cycle = 0, \ + recvReduceCopy_cycle = 0, recvReduceSend_cycle = 0, recvReduceCopySend_cycle = 0, directRecvReduceCopySend_cycle = 0, \ + send_byte = 0, directSend_byte = 0, recv_byte = 0, directRecv_byte = 0, copySend_byte = 0, directCopySend_byte = 0, \ + recvCopySend_byte = 0, directRecvCopySend_byte = 0, recvReduceCopy_byte = 0, recvReduceSend_byte = 0, \ + recvReduceCopySend_byte = 0, directRecvReduceCopySend_byte = 0; for (int chan=0; channChannels; chan++) { - wait_cycle += prof->wait_cycle[chan]; - wait_recv_cycle += prof->wait_recv_cycle[chan]; + total_cycle += prof->elems[chan].total_cycle; + wait_cycle += prof->elems[chan].wait_cycle; + wait_recv_cycle += prof->elems[chan].wait_recv_cycle; + send_cycle += prof->elems[chan].send_cycle; + directSend_cycle += prof->elems[chan].directSend_cycle; + recv_cycle += prof->elems[chan].recv_cycle; + directRecv_cycle += prof->elems[chan].directRecv_cycle; + copySend_cycle += prof->elems[chan].copySend_cycle; + directCopySend_cycle += prof->elems[chan].directCopySend_cycle; + recvCopySend_cycle += prof->elems[chan].recvCopySend_cycle; + directRecvCopySend_cycle += prof->elems[chan].directRecvCopySend_cycle; + recvReduceCopy_cycle += prof->elems[chan].recvReduceCopy_cycle; + recvReduceSend_cycle += prof->elems[chan].recvReduceSend_cycle; + recvReduceCopySend_cycle += prof->elems[chan].recvReduceCopySend_cycle; + directRecvReduceCopySend_cycle += prof->elems[chan].directRecvReduceCopySend_cycle; + send_byte += prof->elems[chan].send_byte; + directSend_byte += prof->elems[chan].directSend_byte; + recv_byte += prof->elems[chan].recv_byte; + directRecv_byte += prof->elems[chan].directRecv_byte; + copySend_byte += prof->elems[chan].copySend_byte; + directCopySend_byte += prof->elems[chan].directCopySend_byte; + recvCopySend_byte += prof->elems[chan].recvCopySend_byte; + directRecvCopySend_byte += prof->elems[chan].directRecvCopySend_byte; + recvReduceCopy_byte += prof->elems[chan].recvReduceCopy_byte; + recvReduceSend_byte += prof->elems[chan].recvReduceSend_byte; + recvReduceCopySend_byte += prof->elems[chan].recvReduceCopySend_byte; + directRecvReduceCopySend_byte += prof->elems[chan].directRecvReduceCopySend_byte; } #define VEGA_GPU_RTC_FREQUENCY 2.5E7 if (comm->rank == 0) { @@ -309,17 +339,17 @@ static ncclResult_t commFree(ncclComm_t comm) { INFO(NCCL_INIT, "# %4s %6s %6s %6s %6s %6s %7s %6s %6s %6s %6s %6s", "", "(s)", "(s)", "(s)", "(GB/s)", "(GB/s)", "(GB/s)", "(GB/s)", "(GB/s)", "(GB/s)", "(GB/s)", "(GB/s)"); } INFO(NCCL_INIT, "# %4d %6.4f %6.4f %6.4f %6.2f %6.2f %7.2f %6.2f %6.2f %6.2f %6.2f %6.2f", - comm->rank, (double)prof->total_cycle/VEGA_GPU_RTC_FREQUENCY/comm->nChannels, + comm->rank, (double)total_cycle/VEGA_GPU_RTC_FREQUENCY/comm->nChannels, (double)wait_cycle/VEGA_GPU_RTC_FREQUENCY/comm->nChannels, (double)wait_recv_cycle/VEGA_GPU_RTC_FREQUENCY/comm->nChannels, - (prof->send_cycle) ? (double)prof->send_byte*comm->nChannels/((double)prof->send_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0, - (prof->recvReduceSend_cycle) ? (double)prof->recvReduceSend_byte*comm->nChannels/((double)prof->recvReduceSend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0, - (prof->directRecvReduceCopySend_cycle) ? (double)prof->directRecvReduceCopySend_byte*comm->nChannels/((double)prof->directRecvReduceCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0, - (prof->directRecvCopySend_cycle) ? (double)prof->directRecvCopySend_byte*comm->nChannels/((double)prof->directRecvCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0, - (prof->directRecv_cycle) ? (double)prof->directRecv_byte*comm->nChannels/((double)prof->directRecv_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0, - (prof->copySend_cycle) ? (double)prof->copySend_byte*comm->nChannels/((double)prof->copySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0, - (prof->recv_cycle) ? (double)prof->recv_byte*comm->nChannels/((double)prof->recv_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0, - (prof->recvCopySend_cycle) ? (double)prof->recvCopySend_byte*comm->nChannels/((double)prof->recvCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0); + (send_cycle) ? (double)send_byte*comm->nChannels/((double)send_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0, + (recvReduceSend_cycle) ? (double)recvReduceSend_byte*comm->nChannels/((double)recvReduceSend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0, + (directRecvReduceCopySend_cycle) ? (double)directRecvReduceCopySend_byte*comm->nChannels/((double)directRecvReduceCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0, + (directRecvCopySend_cycle) ? (double)directRecvCopySend_byte*comm->nChannels/((double)directRecvCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0, + (directRecv_cycle) ? (double)directRecv_byte*comm->nChannels/((double)directRecv_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0, + (copySend_cycle) ? (double)copySend_byte*comm->nChannels/((double)copySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0, + (recv_cycle) ? (double)recv_byte*comm->nChannels/((double)recv_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0, + (recvCopySend_cycle) ? (double)recvCopySend_byte*comm->nChannels/((double)recvCopySend_cycle/VEGA_GPU_RTC_FREQUENCY*1.0E9) : 0); free(prof); CUDACHECK(hipFree(comm->hostDevComm.devProf)); From d2580c8cf5d336dedc676ed59a6c5e8687588f4c Mon Sep 17 00:00:00 2001 From: Wenkai Du Date: Wed, 8 Sep 2021 16:14:32 -0500 Subject: [PATCH 5/6] Improve barrier implementation [ROCm/rccl commit: adb8d633528861165252ac53b043d5f96f95dd24] --- .../rccl/src/collectives/device/primitives.h | 16 ++++++++++------ projects/rccl/src/include/devcomm.h | 4 ++-- 2 files changed, 12 insertions(+), 8 deletions(-) diff --git a/projects/rccl/src/collectives/device/primitives.h b/projects/rccl/src/collectives/device/primitives.h index c64d0bd153..c2c2e2cd10 100644 --- a/projects/rccl/src/collectives/device/primitives.h +++ b/projects/rccl/src/collectives/device/primitives.h @@ -33,12 +33,16 @@ } while (0) #define barrier_by_group() do { \ - const int w = threadIdx.x/WARP_SIZE; \ - const int wid = threadIdx.x%WARP_SIZE; \ - if (wid == 0) { \ - barrier_next[w] += nthreads/WARP_SIZE; \ - __atomic_fetch_add(barriers, 1, __ATOMIC_SEQ_CST); \ - while (LOAD(barriers) < barrier_next[w]) /* spin */; \ + if (nthreads == NCCL_MAX_NTHREADS) \ + __syncthreads(); \ + else { \ + const int w = threadIdx.x/WARP_SIZE; \ + const int wid = threadIdx.x%WARP_SIZE; \ + if (wid == 0) { \ + barrier_next[w] += nthreads/WARP_SIZE; \ + __atomic_fetch_add(barriers, 1, __ATOMIC_SEQ_CST); \ + while (LOAD(barriers) < barrier_next[w]) /* spin */; \ + } \ } \ } while (0) diff --git a/projects/rccl/src/include/devcomm.h b/projects/rccl/src/include/devcomm.h index 42976e36c7..b7923f97eb 100644 --- a/projects/rccl/src/include/devcomm.h +++ b/projects/rccl/src/include/devcomm.h @@ -18,8 +18,8 @@ // Convert volatile access to atomic #if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) -#define LOAD(VAR) __atomic_load_n((VAR), __ATOMIC_SEQ_CST) -#define STORE(DST, SRC) __atomic_store_n((DST), (SRC), __ATOMIC_SEQ_CST) +#define LOAD(VAR) __atomic_load_n((VAR), __ATOMIC_ACQUIRE) +#define STORE(DST, SRC) __atomic_store_n((DST), (SRC), __ATOMIC_RELEASE) #else #define LOAD(VAR) *(VAR) #define STORE(DST, SRC) *(DST) = (SRC) From 9ffeb41fe18faefdeb72e75ce701586113269521 Mon Sep 17 00:00:00 2001 From: Wenkai Du <43822138+wenkaidu@users.noreply.github.com> Date: Mon, 13 Sep 2021 08:39:01 -0700 Subject: [PATCH 6/6] Update tuning table (#424) [ROCm/rccl commit: ef432e48e1abf3a9d76cbdc7bc47883d313979fa] --- projects/rccl/src/graph/tuning.cc | 36 ++++++++++++++----------------- projects/rccl/src/include/graph.h | 2 +- projects/rccl/src/init.cc | 2 +- 3 files changed, 18 insertions(+), 22 deletions(-) diff --git a/projects/rccl/src/graph/tuning.cc b/projects/rccl/src/graph/tuning.cc index a63ef585c4..07d6e6df7c 100644 --- a/projects/rccl/src/graph/tuning.cc +++ b/projects/rccl/src/graph/tuning.cc @@ -54,7 +54,7 @@ ncclResult_t parseList(const char* str, const char* elems[], int nelems, int* li // Latencies in us, Bandwidths in GB/s // Tree { LL, LL128, Simple } , Ring { LL, LL128, Simple } -static const float baseLat [NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { { 39.0, 39.0, 49.0 }, { 21.0, 21.0, 30.0 }, { 37.9, 37.9, 40.4 } }; +static const float baseLat [NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { { 12.0, 12.0, 17.0 }, { 12.0, 12.0, 17.0 }, { 12.0, 12.0, 17.0 } }; // NVLink, PCI, Network #define NCCL_HW_NVLINK 0 @@ -63,11 +63,11 @@ static const float baseLat [NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { { 39.0 // Tree/Simple is the latency a 256kB chunk, which is ~ base lat + 256k/12GB/s (+ 256k/12GB/s for the network). static const float hwLat [3][NCCL_NUM_ALGORITHMS][NCCL_NUM_PROTOCOLS] = { /* NVLINK */ - { /* Tree (LL/LL128/Simple)*/ { 2.5, 2.5, 5.5 }, /* Ring (LL/LL128/Simple)*/ { 2.5, 2.5, 5 }, /* CollNet (LL/LL128/Simple)*/ { 1.2, 1.2, 3.8 } }, + { /* Tree (LL/LL128/Simple)*/ { 1.5, 1.5, 4.5 }, /* Ring (LL/LL128/Simple)*/ { 1.5, 1.5, 4.5 }, /* CollNet (LL/LL128/Simple)*/ { 1.5, 1.5, 4.5 } }, /* PCI */ - { /* Tree (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 }, /* Ring (LL/LL128/Simple)*/ { 1.3, 1.3, 1.9 }, /* CollNet (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 } }, + { /* Tree (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 }, /* Ring (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 }, /* CollNet (LL/LL128/Simple)*/ { 2.2, 2.2, 5.7 } }, /* NET */ - { /* Tree (LL/LL128/Simple)*/ { 28.0, 28.0, 66.0 }, /* Ring (LL/LL128/Simple)*/ { 8.5, 8.5, 19.0 }, /* CollNet (LL/LL128/Simple)*/ { 9.8, 9.8, 19.5 } } + { /* Tree (LL/LL128/Simple)*/ { 40.0, 40.0, 50.0 }, /* Ring (LL/LL128/Simple)*/ { 4.0, 4.0, 25.0 }, /* CollNet (LL/LL128/Simple)*/ { 40.0, 40.0, 50.0 } } }; // LL128 max BW (per channel) for the different collectives @@ -76,7 +76,7 @@ static const double ll128MaxBwPerCh[NCCL_NUM_FUNCTIONS] = { 18.8, 12.0, 18.3, 15 static const double llMaxBws[2][3] = { /* Volta-N1/Intel-N2/Intel-N4) */ {39.0, 39.0, 20.4}, /* Ampere-N1/AMD-N2/AMD-N4) */ {87.7, 22.5 /*avg of ring & tree*/, 19.0} }; static const double perChMaxTreeBws[2][3] = { /* Volta (N1/N2/N4) */ {26.5, 18.5, 10.0}, /* Ampere (N1/N2/N4) */ {24.0, 22.5, 16.0} }; -ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCompCap, struct ncclTopoGraph* treeGraph, struct ncclTopoGraph* ringGraph, struct ncclTopoGraph* collNetGraph) { +ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCompCap, struct ncclTopoGraph* treeGraph, struct ncclTopoGraph* ringGraph, struct ncclTopoGraph* collNetGraph, int gcn) { int simpleDefaultThreads = (ringGraph->speedIntra*ringGraph->nChannels <= PCI_WIDTH) ? 256 : NCCL_SIMPLE_MAX_NTHREADS; comm->maxThreads[NCCL_ALGO_RING][NCCL_PROTO_SIMPLE] = #if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) @@ -133,14 +133,10 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom // Various model refinements #if defined(__HIP_PLATFORM_HCC__) || defined(__HCC__) || defined(__HIPCC__) - if (a == NCCL_ALGO_RING && p == NCCL_PROTO_LL) busBw *= 1.0/5.0; - if (a == NCCL_ALGO_RING && p == NCCL_PROTO_LL128) busBw = std::min(busBw * (ppn < 2 ? 0.7 : 0.92 /*120.0/128.0*/), ll128MaxBwPerCh[coll]*graphs[a]->nChannels); - double maxTreeBw = comm->nNodes > 2 ? - compCap80 && p == NCCL_PROTO_LL128 ? 105.0 : 80.0 : - compCap80 && p == NCCL_PROTO_LL128 ? 130.0 : 110.0; - if (a == NCCL_ALGO_TREE) busBw = std::min(busBw*.27, comm->nNodes > 1 ? 70.0 : 90.0); - if (a == NCCL_ALGO_TREE && p == NCCL_PROTO_LL) busBw *= 1.0/2.3; - if (a == NCCL_ALGO_TREE && p == NCCL_PROTO_LL128) busBw = std::min(busBw * (comm->nNodes == 1 ? 7.0/9.0 : 0.915 /*120.0/128.0*/), ll128MaxBwPerCh[coll]*graphs[a]->nChannels*7.0/9.0); + if (a == NCCL_ALGO_RING && (p == NCCL_PROTO_LL || p == NCCL_PROTO_LL128)) busBw *= 0.05; + if (a == NCCL_ALGO_TREE && p == NCCL_PROTO_SIMPLE) (nNodes == 2) ? busBw *= 0.33 : busBw *= 0.11; + if (a == NCCL_ALGO_TREE && (p == NCCL_PROTO_LL || p == NCCL_PROTO_LL128)) busBw *= 0.04; + if (gcn == 910 && p == NCCL_PROTO_LL && nNodes == 1 && nRanks == 16) busBw *= 5.9; #else if (compCap80) busBw = std::min(busBw, 235.0f); if (a == NCCL_ALGO_RING && p == NCCL_PROTO_LL) { busBw = std::min(llMaxBw, busBw * ((nNodes > 1 || coll == ncclFuncAllReduce || coll == ncclFuncReduce) ? 1.0/4.0 : 1.0/3.0)); } @@ -159,7 +155,7 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom comm->latencies[coll][a][p] = baseLat[a][p]; float intraLat = hwLat[intraHw[a]][a][p]; float interLat = hwLat[NCCL_HW_NET][a][p]; - if (nNodes > 1 && p == NCCL_PROTO_LL) intraLat *= 1.8; + //if (nNodes > 1 && p == NCCL_PROTO_LL) intraLat *= 1.8; if (a == NCCL_ALGO_RING) { float lat = hwLat[hw[a]][a][p]; if ((coll == ncclFuncReduce || coll == ncclFuncBroadcast)) { @@ -285,15 +281,15 @@ ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCom // Trees are not perfectly sticking to the model for medium sizes. Applying a static correction // factor is not ideal but works quite well. Powers of two, 64 B to 128MB. static float treeCorrectionFactor[NCCL_NUM_PROTOCOLS][22] = { - { 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, .84, .49, .42, .60, .75, .87, .94, .94, .99, 1.0, 1.0 , 1.0 , 1.0 , 1.0 , 1.0 }, - { 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, .84, .49, .42, .60, .75, .87, .94, .94, .99, 1.0, 1.0 , 1.0 , 1.0 , 1.0 , 1.0 }, - { 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, .41, .27, .25, .39, .46, .72, .76, .87, .92, .97, 1.0, 1.0 , 1.0 , 1.0 , 1.0 , 1.0 } + { 0.7, 0.7, 0.7, 0.6, 0.6, 0.3, 0.9, 0.5, 0.5, 0.6, 0.5, 0.5, 0.8, 0.9, 0.8, 0.8, 0.8, 0.8, 0.8, 0.9, 0.9, 1.0, }, + { 0.7, 0.7, 0.7, 0.6, 0.6, 0.3, 0.9, 0.5, 0.5, 0.6, 0.5, 0.5, 0.8, 0.9, 0.8, 0.8, 0.8, 0.8, 0.8, 0.9, 0.9, 1.0, }, + { 0.4, 0.4, 0.3, 0.3, 0.2, 0.5, 0.5, 0.7, 0.2, 0.2, 0.3, 0.6, 0.7, 1.0, 1.3, 1.0, 1.2, 1.2, 1.1, 1.1, 1.2, 1.2, }, }; static float ringCorrectionFactor[NCCL_NUM_PROTOCOLS][22] = { - { 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, .25, .41, .55, .56, .78, .94, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0 , 1.0 , 1.0 , 1.0 , 1.0 }, - { 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, .25, .41, .55, .56, .78, .94, 1.0, 1.0, 1.0, 1.0, 1.0, 1.0 , 1.0 , 1.0 , 1.0 , 1.0 }, - { 1.0, 1.0, 1.0, 1.0, 1.0, 1.0, .04, .08, .09, .09, .11, .13, .25, .40, .59, .76, .86, 1.0 , 1.0 , 1.0 , 1.0 , 1.0 } + { 0.4, 0.6, 0.6, 0.3, 0.2, 0.2, 0.2, 0.2, 0.4, 0.6, 0.7, 0.9, 1.4, 1.5, 1.0, 0.8, 0.7, 0.8, 0.8, 0.9, 0.9, 0.9, }, + { 0.4, 0.6, 0.6, 0.3, 0.2, 0.2, 0.2, 0.2, 0.4, 0.6, 0.7, 0.9, 1.4, 1.5, 1.0, 0.8, 0.7, 0.8, 0.8, 0.9, 0.9, 0.9, }, + { 0.6, 0.4, 0.4, 0.4, 0.2, 0.3, 0.2, 0.2, 0.2, 0.2, 0.2, 0.2, 0.2, 0.2, 0.2, 0.2, 0.2, 0.3, 0.4, 0.6, 0.8, 0.9, }, }; ncclResult_t ncclTopoGetAlgoTime(struct ncclInfo* info, int algorithm, int protocol, float* time) { diff --git a/projects/rccl/src/include/graph.h b/projects/rccl/src/include/graph.h index 5072ba1249..34b93be4e7 100644 --- a/projects/rccl/src/include/graph.h +++ b/projects/rccl/src/include/graph.h @@ -105,7 +105,7 @@ ncclResult_t ncclTopoPreset(struct ncclComm* comm, ncclResult_t ncclTopoPostset(struct ncclComm* comm, int* firstRanks, int* treePatterns, struct ncclTopoRanks** allTopoRanks, int* rings, struct ncclTopoGraph* collNetGraph, int nc); -ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCompCap, struct ncclTopoGraph* treeGraph, struct ncclTopoGraph* ringGraph, struct ncclTopoGraph* collNetGraph); +ncclResult_t ncclTopoTuneModel(struct ncclComm* comm, int minCompCap, int maxCompCap, struct ncclTopoGraph* treeGraph, struct ncclTopoGraph* ringGraph, struct ncclTopoGraph* collNetGraph, int gcn); #include "info.h" ncclResult_t ncclTopoGetAlgoTime(struct ncclInfo* info, int algorithm, int protocol, float* time); diff --git a/projects/rccl/src/init.cc b/projects/rccl/src/init.cc index d24b33cb55..74ae05e3f4 100644 --- a/projects/rccl/src/init.cc +++ b/projects/rccl/src/init.cc @@ -1114,7 +1114,7 @@ collnet_cleanup: TRACE(NCCL_INIT, "rank %d nranks %d - CONNECTED %d RINGS AND TREES", rank, nranks, comm->nChannels); // Compute time models for algorithm and protocol combinations - NCCLCHECK(ncclTopoTuneModel(comm, minCompCap, maxCompCap, &treeGraph, &ringGraph, &collNetGraph)); + NCCLCHECK(ncclTopoTuneModel(comm, minCompCap, maxCompCap, &treeGraph, &ringGraph, &collNetGraph, comm->topo->nodes[GPU].nodes[0].gpu.gcn)); // Compute nChannels per peer for p2p NCCLCHECK(ncclTopoComputeP2pChannels(comm));