From f7fb83c79b9e24e2ca0836030e148f88c8a44829 Mon Sep 17 00:00:00 2001 From: RoastSea8 Date: Mon, 25 Mar 2024 20:33:13 -0700 Subject: [PATCH 01/16] add cudaDeviceSynchronize for NCCL --- allgather.cu | 2 ++ 1 file changed, 2 insertions(+) diff --git a/allgather.cu b/allgather.cu index cf1a882..5953041 100644 --- a/allgather.cu +++ b/allgather.cu @@ -158,6 +158,7 @@ int main(int argc, char *argv[]) { MPI_CHECK(MPI_Wait(&request, &status)); #elif defined(USE_NCCL) NCCL_CHECK(ncclAllGather((const void*)d_local_data, (void*)d_global_data, msg_count, ncclBfloat16, nccl_comm, NULL)); + cudaDeviceSynchronize(); #elif defined(USE_RCCL) // TODO: fix later rcclAllReduce((const void*)d_local_data, (void*)d_global_data, global_data_size, rcclInt, rcclSum, comm, NULL); @@ -177,6 +178,7 @@ int main(int argc, char *argv[]) { MPI_CHECK(MPI_Wait(&request, &status)); #elif defined(USE_NCCL) NCCL_CHECK(ncclAllGather((const void*)d_local_data, (void*)d_global_data, msg_count, ncclBfloat16, nccl_comm, NULL)); + cudaDeviceSynchronize(); #elif defined(USE_RCCL) // TODO: fix later rcclAllReduce((const void*)d_local_data, (void*)d_global_data, global_data_size, rcclInt, rcclSum, comm, NULL); From 453a397d2c609f41eac92012e9564ef6e1fa8ed6 Mon Sep 17 00:00:00 2001 From: RoastSea8 Date: Sat, 30 Mar 2024 14:15:33 -0700 Subject: [PATCH 02/16] add allreduce code --- Makefile | 9 ++- allreduce.cu | 219 +++++++++++++++++++++++++++++++++++++++++++++++++++ 2 files changed, 225 insertions(+), 3 deletions(-) create mode 100644 allreduce.cu diff --git a/Makefile b/Makefile index df453b4..231e499 100644 --- a/Makefile +++ b/Makefile @@ -5,14 +5,17 @@ CC = cc INC = -I/global/common/software/nersc9/nccl/2.19.4/include -CFLAGS = -std=c++11 -O2 -target-accel=nvidia80 --cuda-gpu-arch=sm_80 -DUSE_CUDA -DUSE_NCCL +CFLAGS = -std=c++11 -O2 -target-accel=nvidia80 --cuda-gpu-arch=sm_80 -DUSE_CUDA -DUSE_MPI LDFLAGS = -L/global/common/software/nersc9/nccl/2.19.4/lib -lnccl -all: allgather.x +all: allgather.x allreduce.x allgather.x: allgather.cu ${CC} ${CFLAGS} ${INC} ${LDFLAGS} -o allgather.x allgather.cu +allreduce.x: allreduce.cu + ${CC} ${CFLAGS} ${INC} ${LDFLAGS} -o allreduce.x allreduce.cu + clean: - rm -f allgather.x + rm -f allgather.x allreduce.x diff --git a/allreduce.cu b/allreduce.cu new file mode 100644 index 0000000..062b120 --- /dev/null +++ b/allreduce.cu @@ -0,0 +1,219 @@ +/* \file allreduce.cu + * Copyright 2024 Parallel Software and Systems Group, University of Maryland. + * See the top-level LICENSE file for details. + * + * SPDX-License-Identifier: MIT + */ + +#include +#include +#include + +#ifdef USE_CUDA + #include + #include +#endif + +#ifdef USE_NCCL + #include "nccl.h" +#elif defined(USE_RCCL) + #include "rccl.h" +#endif + +#define NUM_WARMUP_ITERATIONS 5 + +#define MPI_CHECK(cmd) do { \ + int e = cmd; \ + if( e != MPI_SUCCESS ) { \ + printf("Failed: MPI error %s:%d '%d'\n", \ + __FILE__,__LINE__, e); \ + exit(EXIT_FAILURE); \ + } \ +} while(0) + +#define CUDA_CHECK(cmd) do { \ + cudaError_t e = cmd; \ + if(e != cudaSuccess) { \ + printf("CUDA error %s:%d: %s\n", \ + __FILE__, __LINE__, cudaGetErrorString(e)); \ + exit(EXIT_FAILURE); \ + } \ +} while(0) + +#define NCCL_CHECK(cmd) do { \ + ncclResult_t e = cmd; \ + if (e != ncclSuccess) { \ + printf("NCCL error %s:%d %s\n", \ + __FILE__, __LINE__, ncclGetErrorString(e)); \ + exit(EXIT_FAILURE); \ + } \ +} while(0) + +void initializeData(nv_bfloat16 *data, int size) { + for (int i = 0; i < (size / sizeof(nv_bfloat16)); ++i) { + data[i] = __float2bfloat16((float)i); + } +} + +void custom_bf16_sum(void *invec, void *inoutvec, int *len, MPI_Datatype *datatype) { + nv_bfloat16* in = (nv_bfloat16*) invec; + nv_bfloat16* inout = (nv_bfloat16*) inoutvec; + for (int i = 0; i < *len; i++) + inout[i] = __hadd(in[i], inout[i]); +} + +int main(int argc, char *argv[]) { + if (argc != 5) { + fprintf(stderr, "Usage: %s \n", argv[0]); + return EXIT_FAILURE; + } + + int num_gpus = atoi(argv[1]); + int min_msg_size = atoi(argv[2]); + int max_msg_size = atoi(argv[3]); + int iterations = atoi(argv[4]); + + if (num_gpus < 2 || min_msg_size <= 0 || max_msg_size <= 0 || min_msg_size > max_msg_size || iterations <= 0) { + fprintf(stderr, "Invalid input parameters.\n"); + return EXIT_FAILURE; + } + + int my_rank, num_pes; + int num_gpus_per_node; + int msg_count; + + MPI_Init(&argc, &argv); + MPI_Comm_rank(MPI_COMM_WORLD, &my_rank); + MPI_Comm_size(MPI_COMM_WORLD, &num_pes); + + if (num_pes != num_gpus) { + fprintf(stderr, "Number of processes must match number of GPUs.\n"); + MPI_Finalize(); + return EXIT_FAILURE; + } + + // Initialize GPU context + cudaGetDeviceCount(&num_gpus_per_node); + cudaSetDevice((my_rank % num_gpus_per_node)); + + int local_data_size = max_msg_size; // Size of local data + int global_data_size = local_data_size; // Size of global data + + nv_bfloat16 *local_data = (nv_bfloat16*)malloc(local_data_size); + nv_bfloat16 *global_data = (nv_bfloat16*)malloc(global_data_size); + + // Initialize local data + initializeData(local_data, local_data_size); + + // Allocate memory on GPU + nv_bfloat16 *d_local_data, *d_global_data; + CUDA_CHECK(cudaMalloc(&d_local_data, local_data_size)); + CUDA_CHECK(cudaMalloc(&d_global_data, global_data_size)); + + // Copy local data to GPU + CUDA_CHECK(cudaMemcpy(d_local_data, local_data, local_data_size, cudaMemcpyHostToDevice)); + + #ifdef USE_MPI + // create 2-byte datatype (send raw, un-interpreted bytes) + MPI_Datatype mpi_type_bfloat16; + MPI_Type_contiguous(2, MPI_BYTE, &mpi_type_bfloat16); + MPI_Type_commit(&mpi_type_bfloat16); + + // define custom reduce operation for nv_bfloat16 types + MPI_Op CUSTOM_SUM; + MPI_Op_create(&custom_bf16_sum, 1, &CUSTOM_SUM); + + #elif USE_NCCL + ncclUniqueId nccl_comm_id; + ncclComm_t nccl_comm; + + if (my_rank == 0) { + /* Generates an Id to be used in ncclCommInitRank. */ + ncclGetUniqueId(&nccl_comm_id); + } + + /* distribute nccl_comm_id to all ranks */ + MPI_CHECK(MPI_Bcast((void *)&nccl_comm_id, sizeof(nccl_comm_id), MPI_BYTE, + 0, MPI_COMM_WORLD)); + + /* Create a new NCCL communicator */ + NCCL_CHECK(ncclCommInitRank(&nccl_comm, num_pes, nccl_comm_id, my_rank)); + + #elif defined(USE_RCCL) + // TODO: fix later + rcclComm_t rccl_comm; + rcclCommInitRank(&comm, num_gpus, 0, rccl_root); + #endif + + // Perform MPI_Iallgather, NCCL allgather, or RCCL allgather + double total_time, start_time; + MPI_Request request; + MPI_Status status; + + // Print benchmark results + if (my_rank == 0) { + printf("Number of GPUs: %d\n", num_gpus); + printf("Message size range: %d - %d\n", min_msg_size, max_msg_size); + printf("Number of iterations: %d\n", iterations); + } + fflush(NULL); + + for (int msg_size = min_msg_size; msg_size <= max_msg_size; msg_size *= 2) { + msg_count = msg_size / sizeof(nv_bfloat16); + // warmup iterations + for (int i = 0; i < NUM_WARMUP_ITERATIONS; ++i) { + #ifdef USE_MPI + MPI_CHECK(MPI_Iallreduce(d_local_data, d_global_data, msg_count, mpi_type_bfloat16, + CUSTOM_SUM, MPI_COMM_WORLD, &request)); + + MPI_CHECK(MPI_Wait(&request, &status)); + #elif defined(USE_NCCL) + NCCL_CHECK(ncclAllReduce((const void*)d_local_data, (void*)d_global_data, msg_count, ncclBfloat16, ncclSum, nccl_comm, NULL)); + cudaDeviceSynchronize(); + #elif defined(USE_RCCL) + // TODO: fix later + rcclAllReduce((const void*)d_local_data, (void*)d_global_data, global_data_size, rcclInt, rcclSum, comm, NULL); + #endif + } + + if(msg_size >= 8388608) + iterations = 20; + + MPI_Barrier(MPI_COMM_WORLD); + start_time = MPI_Wtime(); + for (int i = 0; i < iterations; ++i) { + #ifdef USE_MPI + MPI_CHECK(MPI_Iallreduce(d_local_data, d_global_data, msg_count, mpi_type_bfloat16, + CUSTOM_SUM, MPI_COMM_WORLD, &request)); + + MPI_CHECK(MPI_Wait(&request, &status)); + #elif defined(USE_NCCL) + NCCL_CHECK(ncclAllReduce((const void*)d_local_data, (void*)d_global_data, msg_count, ncclBfloat16, ncclSum, nccl_comm, NULL)); + cudaDeviceSynchronize(); + #elif defined(USE_RCCL) + // TODO: fix later + rcclAllReduce((const void*)d_local_data, (void*)d_global_data, global_data_size, rcclInt, rcclSum, comm, NULL); + #endif + } + MPI_Barrier(MPI_COMM_WORLD); + total_time = MPI_Wtime() - start_time; + if (my_rank == 0) + printf("%d %.6f seconds\n", msg_size, (total_time / iterations)); + } + + // Cleanup + free(local_data); + free(global_data); + CUDA_CHECK(cudaFree(d_local_data)); + CUDA_CHECK(cudaFree(d_global_data)); + + #ifdef USE_NCCL + ncclCommDestroy(nccl_comm); + #elif defined(USE_RCCL) + rcclCommDestroy(rccl_comm); + #endif + + MPI_Finalize(); + return EXIT_SUCCESS; +} + From 982ccaf7c6d265bf9ed21962432f8d462621dfbe Mon Sep 17 00:00:00 2001 From: Aditya Tomar Date: Sun, 31 Mar 2024 23:43:59 -0400 Subject: [PATCH 03/16] add code for ROCm and RCCL --- Makefile | 16 +++++--- allgather.cu | 101 ++++++++++++++++++++++++++++++++++----------------- 2 files changed, 78 insertions(+), 39 deletions(-) diff --git a/Makefile b/Makefile index df453b4..9f6d40a 100644 --- a/Makefile +++ b/Makefile @@ -4,15 +4,21 @@ # SPDX-License-Identifier: MIT CC = cc -INC = -I/global/common/software/nersc9/nccl/2.19.4/include -CFLAGS = -std=c++11 -O2 -target-accel=nvidia80 --cuda-gpu-arch=sm_80 -DUSE_CUDA -DUSE_NCCL -LDFLAGS = -L/global/common/software/nersc9/nccl/2.19.4/lib -lnccl +# perlmutter flags +# INC = -I/global/common/software/nersc9/nccl/2.19.4/include +# CFLAGS = -std=c++11 -O2 -target-accel=nvidia80 --cuda-gpu-arch=sm_80 -DUSE_CUDA -DUSE_NCCL +# LDFLAGS = -L/global/common/software/nersc9/nccl/2.19.4/lib -lnccl + +# frontier flags +INC = -L${ROCM_PATH}/lib -lamdhip64 +CFLAGS = -std=c++11 -D__HIP_ROCclr__ -D__HIP_ARCH_GFX90A__=1 --rocm-path=${ROCM_PATH} --offload-arch=gfx90a -x hip -DUSE_ROCM -DUSE_RCCL +LDFLAGS = --rocm-path=${ROCM_PATH} -lrccl all: allgather.x -allgather.x: allgather.cu +allgather.x: allgather.cu ${CC} ${CFLAGS} ${INC} ${LDFLAGS} -o allgather.x allgather.cu clean: - rm -f allgather.x + rm -f allgather.x diff --git a/allgather.cu b/allgather.cu index 5953041..b6fd992 100644 --- a/allgather.cu +++ b/allgather.cu @@ -8,16 +8,21 @@ #include #include #include - +#include #ifdef USE_CUDA - #include #include + #define bfloat16 nv_bfloat16 +#elif USE_ROCM + #include + #include + #include + #define bfloat16 hip_bfloat16 #endif #ifdef USE_NCCL #include "nccl.h" -#elif defined(USE_RCCL) - #include "rccl.h" +#elif USE_RCCL + #include #endif #define NUM_WARMUP_ITERATIONS 5 @@ -40,6 +45,16 @@ } \ } while(0) +#define HIP_CHECK(cmd) do { \ + hipError_t e = cmd; \ + if(e != hipSuccess) { \ + printf("HIP error %s:%d: %s\n", \ + __FILE__, __LINE__, hipGetErrorString(e)); \ + exit(EXIT_FAILURE); \ + } \ +} while(0) + +// NCCL_CHECK is used to validate RCCL functions as well #define NCCL_CHECK(cmd) do { \ ncclResult_t e = cmd; \ if (e != ncclSuccess) { \ @@ -49,9 +64,14 @@ } \ } while(0) -void initializeData(nv_bfloat16 *data, int size) { - for (int i = 0; i < (size / sizeof(nv_bfloat16)); ++i) { +void initializeData(bfloat16 *data, int size) { + for (int i = 0; i < (size / sizeof(bfloat16)); ++i) { + #ifdef USE_CUDA data[i] = __float2bfloat16((float)i); + #elif USE_ROCM + // ROCm doesn't have a float2bfloat16 method + data[i] = (bfloat16) ((float) i); + #endif } } @@ -86,33 +106,44 @@ int main(int argc, char *argv[]) { } // Initialize GPU context + #if USE_CUDA cudaGetDeviceCount(&num_gpus_per_node); cudaSetDevice((my_rank % num_gpus_per_node)); + #elif USE_ROCM + hipGetDeviceCount(&num_gpus_per_node); + hipSetDevice((my_rank % num_gpus_per_node)); + #endif int local_data_size = max_msg_size; // Size of local data int global_data_size = local_data_size * num_gpus; // Size of global data - nv_bfloat16 *local_data = (nv_bfloat16*)malloc(local_data_size); - nv_bfloat16 *global_data = (nv_bfloat16*)malloc(global_data_size); + bfloat16 *local_data = (bfloat16*)malloc(local_data_size); + bfloat16 *global_data = (bfloat16*)malloc(global_data_size); // Initialize local data initializeData(local_data, local_data_size); // Allocate memory on GPU - nv_bfloat16 *d_local_data, *d_global_data; + bfloat16 *d_local_data, *d_global_data; + #ifdef USE_CUDA CUDA_CHECK(cudaMalloc(&d_local_data, local_data_size)); CUDA_CHECK(cudaMalloc(&d_global_data, global_data_size)); - // Copy local data to GPU CUDA_CHECK(cudaMemcpy(d_local_data, local_data, local_data_size, cudaMemcpyHostToDevice)); + #elif USE_ROCM + HIP_CHECK(hipMalloc(&d_local_data, local_data_size)); + HIP_CHECK(hipMalloc(&d_global_data, global_data_size)); + HIP_CHECK(hipMemcpy(d_local_data, local_data, local_data_size, hipMemcpyHostToDevice)); + #endif + #ifdef USE_MPI // create 2-byte datatype (send raw, un-interpreted bytes) MPI_Datatype mpi_type_bfloat16; MPI_Type_contiguous(2, MPI_BYTE, &mpi_type_bfloat16); MPI_Type_commit(&mpi_type_bfloat16); - #elif USE_NCCL + #elif defined(USE_NCCL) || defined(USE_RCCL) ncclUniqueId nccl_comm_id; ncclComm_t nccl_comm; @@ -125,13 +156,8 @@ int main(int argc, char *argv[]) { MPI_CHECK(MPI_Bcast((void *)&nccl_comm_id, sizeof(nccl_comm_id), MPI_BYTE, 0, MPI_COMM_WORLD)); - /* Create a new NCCL communicator */ + /* Create a new NCCL/RCCL communicator */ NCCL_CHECK(ncclCommInitRank(&nccl_comm, num_pes, nccl_comm_id, my_rank)); - - #elif defined(USE_RCCL) - // TODO: fix later - rcclComm_t rccl_comm; - rcclCommInitRank(&comm, num_gpus, 0, rccl_root); #endif // Perform MPI_Iallgather, NCCL allgather, or RCCL allgather @@ -148,7 +174,7 @@ int main(int argc, char *argv[]) { fflush(NULL); for (int msg_size = min_msg_size; msg_size <= max_msg_size; msg_size *= 2) { - msg_count = msg_size / sizeof(nv_bfloat16); + msg_count = msg_size / sizeof(bfloat16); // warmup iterations for (int i = 0; i < NUM_WARMUP_ITERATIONS; ++i) { #ifdef USE_MPI @@ -156,12 +182,14 @@ int main(int argc, char *argv[]) { d_global_data, msg_count, mpi_type_bfloat16, MPI_COMM_WORLD, &request)); MPI_CHECK(MPI_Wait(&request, &status)); - #elif defined(USE_NCCL) + #elif defined(USE_NCCL) || defined(USE_RCCL) NCCL_CHECK(ncclAllGather((const void*)d_local_data, (void*)d_global_data, msg_count, ncclBfloat16, nccl_comm, NULL)); - cudaDeviceSynchronize(); - #elif defined(USE_RCCL) - // TODO: fix later - rcclAllReduce((const void*)d_local_data, (void*)d_global_data, global_data_size, rcclInt, rcclSum, comm, NULL); + #endif + + #ifdef USE_CUDA + cudaDeviceSynchronize(); + #elif USE_ROCM + hipDeviceSynchronize(); #endif } @@ -172,16 +200,18 @@ int main(int argc, char *argv[]) { start_time = MPI_Wtime(); for (int i = 0; i < iterations; ++i) { #ifdef USE_MPI - MPI_CHECK(MPI_Iallgather(d_local_data, msg_count, mpi_type_bfloat16, - d_global_data, msg_count, mpi_type_bfloat16, MPI_COMM_WORLD, &request)); - + MPI_CHECK(MPI_Iallgather(d_local_data, msg_count, mpi_type_bfloat16, + d_global_data, msg_count, mpi_type_bfloat16, MPI_COMM_WORLD, &request)); + MPI_CHECK(MPI_Wait(&request, &status)); - #elif defined(USE_NCCL) + #elif defined(USE_NCCL) || defined(USE_RCCL) NCCL_CHECK(ncclAllGather((const void*)d_local_data, (void*)d_global_data, msg_count, ncclBfloat16, nccl_comm, NULL)); - cudaDeviceSynchronize(); - #elif defined(USE_RCCL) - // TODO: fix later - rcclAllReduce((const void*)d_local_data, (void*)d_global_data, global_data_size, rcclInt, rcclSum, comm, NULL); + #endif + + #ifdef USE_CUDA + cudaDeviceSynchronize(); + #elif USE_ROCM + hipDeviceSynchronize(); #endif } MPI_Barrier(MPI_COMM_WORLD); @@ -193,13 +223,16 @@ int main(int argc, char *argv[]) { // Cleanup free(local_data); free(global_data); + #ifdef USE_CUDA CUDA_CHECK(cudaFree(d_local_data)); CUDA_CHECK(cudaFree(d_global_data)); + #elif USE_ROCM + HIP_CHECK(hipFree(d_local_data)); + HIP_CHECK(hipFree(d_global_data)); + #endif - #ifdef USE_NCCL + #ifdef defined(USE_NCCL) || defined(USE_RCCL) ncclCommDestroy(nccl_comm); - #elif defined(USE_RCCL) - rcclCommDestroy(rccl_comm); #endif MPI_Finalize(); From 8ab25d10e8dce0cce796837d5813b47a8425ebd7 Mon Sep 17 00:00:00 2001 From: Aditya Tomar Date: Sun, 31 Mar 2024 23:59:33 -0400 Subject: [PATCH 04/16] revert Makefile to original --- Makefile | 14 ++++---------- 1 file changed, 4 insertions(+), 10 deletions(-) diff --git a/Makefile b/Makefile index 9f6d40a..11f5145 100644 --- a/Makefile +++ b/Makefile @@ -4,20 +4,14 @@ # SPDX-License-Identifier: MIT CC = cc +INC = -I/global/common/software/nersc9/nccl/2.19.4/include +CFLAGS = -std=c++11 -O2 -target-accel=nvidia80 --cuda-gpu-arch=sm_80 -DUSE_CUDA -DUSE_NCCL +LDFLAGS = -L/global/common/software/nersc9/nccl/2.19.4/lib -lnccl -# perlmutter flags -# INC = -I/global/common/software/nersc9/nccl/2.19.4/include -# CFLAGS = -std=c++11 -O2 -target-accel=nvidia80 --cuda-gpu-arch=sm_80 -DUSE_CUDA -DUSE_NCCL -# LDFLAGS = -L/global/common/software/nersc9/nccl/2.19.4/lib -lnccl - -# frontier flags -INC = -L${ROCM_PATH}/lib -lamdhip64 -CFLAGS = -std=c++11 -D__HIP_ROCclr__ -D__HIP_ARCH_GFX90A__=1 --rocm-path=${ROCM_PATH} --offload-arch=gfx90a -x hip -DUSE_ROCM -DUSE_RCCL -LDFLAGS = --rocm-path=${ROCM_PATH} -lrccl all: allgather.x -allgather.x: allgather.cu +allgather.x: allgather.cu ${CC} ${CFLAGS} ${INC} ${LDFLAGS} -o allgather.x allgather.cu clean: From ef65ccde73efcf113fa3799462638dc3c8fcbf5e Mon Sep 17 00:00:00 2001 From: Aditya Tomar Date: Mon, 1 Apr 2024 00:12:38 -0400 Subject: [PATCH 05/16] remove unneeded import --- allgather.cu | 1 - 1 file changed, 1 deletion(-) diff --git a/allgather.cu b/allgather.cu index b6fd992..698e425 100644 --- a/allgather.cu +++ b/allgather.cu @@ -8,7 +8,6 @@ #include #include #include -#include #ifdef USE_CUDA #include #define bfloat16 nv_bfloat16 From 795a6d3e4f5323b3fa8e706ae4f159b58c0dbfca Mon Sep 17 00:00:00 2001 From: Aditya Tomar Date: Mon, 1 Apr 2024 00:28:31 -0400 Subject: [PATCH 06/16] add ROCm and RCCL code for all-reduce --- allreduce.cu | 95 ++++++++++++++++++++++++++++++++++------------------ 1 file changed, 63 insertions(+), 32 deletions(-) diff --git a/allreduce.cu b/allreduce.cu index 062b120..4394249 100644 --- a/allreduce.cu +++ b/allreduce.cu @@ -8,16 +8,20 @@ #include #include #include - #ifdef USE_CUDA - #include #include + #define bfloat16 nv_bfloat16 +#elif USE_ROCM + #include + #include + #include + #define bfloat16 hip_bfloat16 #endif #ifdef USE_NCCL #include "nccl.h" -#elif defined(USE_RCCL) - #include "rccl.h" +#elif USE_RCCL + #include #endif #define NUM_WARMUP_ITERATIONS 5 @@ -40,6 +44,16 @@ } \ } while(0) +#define HIP_CHECK(cmd) do { \ + hipError_t e = cmd; \ + if(e != hipSuccess) { \ + printf("HIP error %s:%d: %s\n", \ + __FILE__, __LINE__, hipGetErrorString(e)); \ + exit(EXIT_FAILURE); \ + } \ +} while(0) + +// NCCL_CHECK is used to validate RCCL functions as well #define NCCL_CHECK(cmd) do { \ ncclResult_t e = cmd; \ if (e != ncclSuccess) { \ @@ -49,9 +63,14 @@ } \ } while(0) -void initializeData(nv_bfloat16 *data, int size) { - for (int i = 0; i < (size / sizeof(nv_bfloat16)); ++i) { +void initializeData(bfloat16 *data, int size) { + for (int i = 0; i < (size / sizeof(bfloat16)); ++i) { + #ifdef USE_CUDA data[i] = __float2bfloat16((float)i); + #elif USE_ROCM + // ROCm doesn't have a float2bfloat16 method + data[i] = (bfloat16) ((float) i); + #endif } } @@ -93,26 +112,36 @@ int main(int argc, char *argv[]) { } // Initialize GPU context + #if USE_CUDA cudaGetDeviceCount(&num_gpus_per_node); cudaSetDevice((my_rank % num_gpus_per_node)); + #elif USE_ROCM + hipGetDeviceCount(&num_gpus_per_node); + hipSetDevice((my_rank % num_gpus_per_node)); + #endif int local_data_size = max_msg_size; // Size of local data int global_data_size = local_data_size; // Size of global data - nv_bfloat16 *local_data = (nv_bfloat16*)malloc(local_data_size); - nv_bfloat16 *global_data = (nv_bfloat16*)malloc(global_data_size); + bfloat16 *local_data = (bfloat16*)malloc(local_data_size); + bfloat16 *global_data = (bfloat16*)malloc(global_data_size); // Initialize local data initializeData(local_data, local_data_size); - // Allocate memory on GPU - nv_bfloat16 *d_local_data, *d_global_data; + bfloat16 *d_local_data, *d_global_data; + #ifdef USE_CUDA CUDA_CHECK(cudaMalloc(&d_local_data, local_data_size)); CUDA_CHECK(cudaMalloc(&d_global_data, global_data_size)); - // Copy local data to GPU CUDA_CHECK(cudaMemcpy(d_local_data, local_data, local_data_size, cudaMemcpyHostToDevice)); + #elif USE_ROCM + HIP_CHECK(hipMalloc(&d_local_data, local_data_size)); + HIP_CHECK(hipMalloc(&d_global_data, global_data_size)); + HIP_CHECK(hipMemcpy(d_local_data, local_data, local_data_size, hipMemcpyHostToDevice)); + #endif + #ifdef USE_MPI // create 2-byte datatype (send raw, un-interpreted bytes) MPI_Datatype mpi_type_bfloat16; @@ -123,7 +152,7 @@ int main(int argc, char *argv[]) { MPI_Op CUSTOM_SUM; MPI_Op_create(&custom_bf16_sum, 1, &CUSTOM_SUM); - #elif USE_NCCL + #elif defined(USE_NCCL) || defined(USE_RCCL) ncclUniqueId nccl_comm_id; ncclComm_t nccl_comm; @@ -136,13 +165,8 @@ int main(int argc, char *argv[]) { MPI_CHECK(MPI_Bcast((void *)&nccl_comm_id, sizeof(nccl_comm_id), MPI_BYTE, 0, MPI_COMM_WORLD)); - /* Create a new NCCL communicator */ + /* Create a new NCCL/RCCL communicator */ NCCL_CHECK(ncclCommInitRank(&nccl_comm, num_pes, nccl_comm_id, my_rank)); - - #elif defined(USE_RCCL) - // TODO: fix later - rcclComm_t rccl_comm; - rcclCommInitRank(&comm, num_gpus, 0, rccl_root); #endif // Perform MPI_Iallgather, NCCL allgather, or RCCL allgather @@ -159,7 +183,7 @@ int main(int argc, char *argv[]) { fflush(NULL); for (int msg_size = min_msg_size; msg_size <= max_msg_size; msg_size *= 2) { - msg_count = msg_size / sizeof(nv_bfloat16); + msg_count = msg_size / sizeof(bfloat16); // warmup iterations for (int i = 0; i < NUM_WARMUP_ITERATIONS; ++i) { #ifdef USE_MPI @@ -167,12 +191,14 @@ int main(int argc, char *argv[]) { CUSTOM_SUM, MPI_COMM_WORLD, &request)); MPI_CHECK(MPI_Wait(&request, &status)); - #elif defined(USE_NCCL) + #elif defined(USE_NCCL) || defined(USE_RCCL) NCCL_CHECK(ncclAllReduce((const void*)d_local_data, (void*)d_global_data, msg_count, ncclBfloat16, ncclSum, nccl_comm, NULL)); - cudaDeviceSynchronize(); - #elif defined(USE_RCCL) - // TODO: fix later - rcclAllReduce((const void*)d_local_data, (void*)d_global_data, global_data_size, rcclInt, rcclSum, comm, NULL); + #endif + + #ifdef USE_CUDA + cudaDeviceSynchronize(); + #elif USE_ROCM + hipDeviceSynchronize(); #endif } @@ -187,12 +213,14 @@ int main(int argc, char *argv[]) { CUSTOM_SUM, MPI_COMM_WORLD, &request)); MPI_CHECK(MPI_Wait(&request, &status)); - #elif defined(USE_NCCL) + #elif defined(USE_NCCL) || defined(USE_RCCL) NCCL_CHECK(ncclAllReduce((const void*)d_local_data, (void*)d_global_data, msg_count, ncclBfloat16, ncclSum, nccl_comm, NULL)); - cudaDeviceSynchronize(); - #elif defined(USE_RCCL) - // TODO: fix later - rcclAllReduce((const void*)d_local_data, (void*)d_global_data, global_data_size, rcclInt, rcclSum, comm, NULL); + #endif + + #ifdef USE_CUDA + cudaDeviceSynchronize(); + #elif USE_ROCM + hipDeviceSynchronize(); #endif } MPI_Barrier(MPI_COMM_WORLD); @@ -204,13 +232,16 @@ int main(int argc, char *argv[]) { // Cleanup free(local_data); free(global_data); + #ifdef USE_CUDA CUDA_CHECK(cudaFree(d_local_data)); CUDA_CHECK(cudaFree(d_global_data)); + #elif USE_ROCM + HIP_CHECK(hipFree(d_local_data)); + HIP_CHECK(hipFree(d_global_data)); + #endif - #ifdef USE_NCCL + #ifdef defined(USE_NCCL) || defined(USE_RCCL) ncclCommDestroy(nccl_comm); - #elif defined(USE_RCCL) - rcclCommDestroy(rccl_comm); #endif MPI_Finalize(); From 4a87dfc7a3084f57f1727865ab4889b3441fcdfd Mon Sep 17 00:00:00 2001 From: Aditya Tomar Date: Mon, 1 Apr 2024 00:59:47 -0400 Subject: [PATCH 07/16] revert Makefile to original --- Makefile | 9 +++------ 1 file changed, 3 insertions(+), 6 deletions(-) diff --git a/Makefile b/Makefile index 231e499..df453b4 100644 --- a/Makefile +++ b/Makefile @@ -5,17 +5,14 @@ CC = cc INC = -I/global/common/software/nersc9/nccl/2.19.4/include -CFLAGS = -std=c++11 -O2 -target-accel=nvidia80 --cuda-gpu-arch=sm_80 -DUSE_CUDA -DUSE_MPI +CFLAGS = -std=c++11 -O2 -target-accel=nvidia80 --cuda-gpu-arch=sm_80 -DUSE_CUDA -DUSE_NCCL LDFLAGS = -L/global/common/software/nersc9/nccl/2.19.4/lib -lnccl -all: allgather.x allreduce.x +all: allgather.x allgather.x: allgather.cu ${CC} ${CFLAGS} ${INC} ${LDFLAGS} -o allgather.x allgather.cu -allreduce.x: allreduce.cu - ${CC} ${CFLAGS} ${INC} ${LDFLAGS} -o allreduce.x allreduce.cu - clean: - rm -f allgather.x allreduce.x + rm -f allgather.x From b6083d1cea284f619120d27e1a9c62018cf1464a Mon Sep 17 00:00:00 2001 From: Aditya Tomar Date: Mon, 1 Apr 2024 01:00:41 -0400 Subject: [PATCH 08/16] revert Makefile to original --- Makefile | 2 +- 1 file changed, 1 insertion(+), 1 deletion(-) diff --git a/Makefile b/Makefile index 11f5145..df453b4 100644 --- a/Makefile +++ b/Makefile @@ -15,4 +15,4 @@ allgather.x: allgather.cu ${CC} ${CFLAGS} ${INC} ${LDFLAGS} -o allgather.x allgather.cu clean: - rm -f allgather.x + rm -f allgather.x From 79b2fb96a335ff798f63160ab36d530753665114 Mon Sep 17 00:00:00 2001 From: Aditya Tomar Date: Mon, 1 Apr 2024 01:34:35 -0400 Subject: [PATCH 09/16] update custom bf16 sum function --- allreduce.cu | 8 ++++++-- 1 file changed, 6 insertions(+), 2 deletions(-) diff --git a/allreduce.cu b/allreduce.cu index 4394249..7bf1031 100644 --- a/allreduce.cu +++ b/allreduce.cu @@ -75,10 +75,14 @@ void initializeData(bfloat16 *data, int size) { } void custom_bf16_sum(void *invec, void *inoutvec, int *len, MPI_Datatype *datatype) { - nv_bfloat16* in = (nv_bfloat16*) invec; - nv_bfloat16* inout = (nv_bfloat16*) inoutvec; + bfloat16* in = (bfloat16*) invec; + bfloat16* inout = (bfloat16*) inoutvec; for (int i = 0; i < *len; i++) + #ifdef USE_CUDA inout[i] = __hadd(in[i], inout[i]); + #elif USE_ROCM + inout[i] = in[i] + inout[i]; + #endif } int main(int argc, char *argv[]) { From ef6fb0d6d22e9e0dac426da032cd678df12875e3 Mon Sep 17 00:00:00 2001 From: Aditya Tomar Date: Mon, 1 Apr 2024 01:41:28 -0400 Subject: [PATCH 10/16] fix custom bf16 sum function --- allreduce.cu | 3 ++- 1 file changed, 2 insertions(+), 1 deletion(-) diff --git a/allreduce.cu b/allreduce.cu index 7bf1031..9f017db 100644 --- a/allreduce.cu +++ b/allreduce.cu @@ -77,12 +77,13 @@ void initializeData(bfloat16 *data, int size) { void custom_bf16_sum(void *invec, void *inoutvec, int *len, MPI_Datatype *datatype) { bfloat16* in = (bfloat16*) invec; bfloat16* inout = (bfloat16*) inoutvec; - for (int i = 0; i < *len; i++) + for (int i = 0; i < *len; i++) { #ifdef USE_CUDA inout[i] = __hadd(in[i], inout[i]); #elif USE_ROCM inout[i] = in[i] + inout[i]; #endif + } } int main(int argc, char *argv[]) { From 8052ca765fc87da0cb8f9de17bcc8d252b4dad56 Mon Sep 17 00:00:00 2001 From: Aditya Tomar Date: Mon, 1 Apr 2024 01:45:14 -0400 Subject: [PATCH 11/16] fix indents --- allreduce.cu | 8 ++++---- 1 file changed, 4 insertions(+), 4 deletions(-) diff --git a/allreduce.cu b/allreduce.cu index 9f017db..2ffac86 100644 --- a/allreduce.cu +++ b/allreduce.cu @@ -75,11 +75,11 @@ void initializeData(bfloat16 *data, int size) { } void custom_bf16_sum(void *invec, void *inoutvec, int *len, MPI_Datatype *datatype) { - bfloat16* in = (bfloat16*) invec; - bfloat16* inout = (bfloat16*) inoutvec; - for (int i = 0; i < *len; i++) { + bfloat16* in = (bfloat16*) invec; + bfloat16* inout = (bfloat16*) inoutvec; + for (int i = 0; i < *len; i++) { #ifdef USE_CUDA - inout[i] = __hadd(in[i], inout[i]); + inout[i] = __hadd(in[i], inout[i]); #elif USE_ROCM inout[i] = in[i] + inout[i]; #endif From abc811fa0c1a5db94b245e7c78eaa2405b9d30bb Mon Sep 17 00:00:00 2001 From: RoastSea8 Date: Sat, 13 Apr 2024 13:24:59 -0700 Subject: [PATCH 12/16] change to int64_t for global/local data size --- allreduce.cu | 33 ++++++++++++++++++++------------- 1 file changed, 20 insertions(+), 13 deletions(-) diff --git a/allreduce.cu b/allreduce.cu index 2ffac86..6e1cbef 100644 --- a/allreduce.cu +++ b/allreduce.cu @@ -8,6 +8,8 @@ #include #include #include +#include + #ifdef USE_CUDA #include #define bfloat16 nv_bfloat16 @@ -27,9 +29,9 @@ #define NUM_WARMUP_ITERATIONS 5 #define MPI_CHECK(cmd) do { \ - int e = cmd; \ + int64_t e = cmd; \ if( e != MPI_SUCCESS ) { \ - printf("Failed: MPI error %s:%d '%d'\n", \ + printf("Failed: MPI error %s:%d '%ld'\n", \ __FILE__,__LINE__, e); \ exit(EXIT_FAILURE); \ } \ @@ -63,8 +65,8 @@ } \ } while(0) -void initializeData(bfloat16 *data, int size) { - for (int i = 0; i < (size / sizeof(bfloat16)); ++i) { +void initializeData(bfloat16 *data, int64_t size) { + for (int64_t i = 0; i < (size / sizeof(bfloat16)); ++i) { #ifdef USE_CUDA data[i] = __float2bfloat16((float)i); #elif USE_ROCM @@ -74,10 +76,10 @@ void initializeData(bfloat16 *data, int size) { } } -void custom_bf16_sum(void *invec, void *inoutvec, int *len, MPI_Datatype *datatype) { +void custom_bf16_sum(void *invec, void *inoutvec, int64_t *len, MPI_Datatype *datatype) { bfloat16* in = (bfloat16*) invec; bfloat16* inout = (bfloat16*) inoutvec; - for (int i = 0; i < *len; i++) { + for (int64_t i = 0; i < *len; i++) { #ifdef USE_CUDA inout[i] = __hadd(in[i], inout[i]); #elif USE_ROCM @@ -93,8 +95,8 @@ int main(int argc, char *argv[]) { } int num_gpus = atoi(argv[1]); - int min_msg_size = atoi(argv[2]); - int max_msg_size = atoi(argv[3]); + int64_t min_msg_size = atoi(argv[2]); + int64_t max_msg_size = atoi(argv[3]); int iterations = atoi(argv[4]); if (num_gpus < 2 || min_msg_size <= 0 || max_msg_size <= 0 || min_msg_size > max_msg_size || iterations <= 0) { @@ -125,8 +127,13 @@ int main(int argc, char *argv[]) { hipSetDevice((my_rank % num_gpus_per_node)); #endif - int local_data_size = max_msg_size; // Size of local data - int global_data_size = local_data_size; // Size of global data + int64_t local_data_size = max_msg_size; // Size of local data + int64_t global_data_size = local_data_size; // Size of global data + + if (my_rank == 0) { + fprintf(stdout, "Local data size: %ld\n", (local_data_size / 1024) / 1024); + fprintf(stdout, "Global data size: %ld\n", (global_data_size / 1024) / 1024); + } bfloat16 *local_data = (bfloat16*)malloc(local_data_size); bfloat16 *global_data = (bfloat16*)malloc(global_data_size); @@ -182,12 +189,12 @@ int main(int argc, char *argv[]) { // Print benchmark results if (my_rank == 0) { printf("Number of GPUs: %d\n", num_gpus); - printf("Message size range: %d - %d\n", min_msg_size, max_msg_size); + printf("Message size range: %ld - %ld\n", min_msg_size, max_msg_size); printf("Number of iterations: %d\n", iterations); } fflush(NULL); - for (int msg_size = min_msg_size; msg_size <= max_msg_size; msg_size *= 2) { + for (int64_t msg_size = min_msg_size; msg_size <= max_msg_size; msg_size *= 2) { msg_count = msg_size / sizeof(bfloat16); // warmup iterations for (int i = 0; i < NUM_WARMUP_ITERATIONS; ++i) { @@ -231,7 +238,7 @@ int main(int argc, char *argv[]) { MPI_Barrier(MPI_COMM_WORLD); total_time = MPI_Wtime() - start_time; if (my_rank == 0) - printf("%d %.6f seconds\n", msg_size, (total_time / iterations)); + printf("%ld %.6f seconds\n", msg_size, (total_time / iterations)); } // Cleanup From b25520e72aa39a546c016e6a3cb89c0e03ae04c5 Mon Sep 17 00:00:00 2001 From: RoastSea8 Date: Sat, 13 Apr 2024 14:30:39 -0700 Subject: [PATCH 13/16] revert type change for custom sum --- allreduce.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/allreduce.cu b/allreduce.cu index 6e1cbef..1f9a9b9 100644 --- a/allreduce.cu +++ b/allreduce.cu @@ -76,10 +76,10 @@ void initializeData(bfloat16 *data, int64_t size) { } } -void custom_bf16_sum(void *invec, void *inoutvec, int64_t *len, MPI_Datatype *datatype) { +void custom_bf16_sum(void *invec, void *inoutvec, int *len, MPI_Datatype *datatype) { bfloat16* in = (bfloat16*) invec; bfloat16* inout = (bfloat16*) inoutvec; - for (int64_t i = 0; i < *len; i++) { + for (int i = 0; i < *len; i++) { #ifdef USE_CUDA inout[i] = __hadd(in[i], inout[i]); #elif USE_ROCM From d2954f98a95493bd35e3becad8563f2c2589e315 Mon Sep 17 00:00:00 2001 From: RoastSea8 Date: Sat, 13 Apr 2024 17:58:11 -0700 Subject: [PATCH 14/16] change atoi to strtoll --- allreduce.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/allreduce.cu b/allreduce.cu index 1f9a9b9..d8a183a 100644 --- a/allreduce.cu +++ b/allreduce.cu @@ -95,8 +95,8 @@ int main(int argc, char *argv[]) { } int num_gpus = atoi(argv[1]); - int64_t min_msg_size = atoi(argv[2]); - int64_t max_msg_size = atoi(argv[3]); + int64_t min_msg_size = strtoll(argv[2], NULL, 10); + int64_t max_msg_size = strtoll(argv[3], NULL, 10); int iterations = atoi(argv[4]); if (num_gpus < 2 || min_msg_size <= 0 || max_msg_size <= 0 || min_msg_size > max_msg_size || iterations <= 0) { From aa8e516572abc35c191df3f256982f068d6f24fc Mon Sep 17 00:00:00 2001 From: Aditya Tomar Date: Sun, 14 Apr 2024 04:44:16 -0700 Subject: [PATCH 15/16] add hip directive --- allreduce.cu | 1 + 1 file changed, 1 insertion(+) diff --git a/allreduce.cu b/allreduce.cu index d8a183a..3b31238 100644 --- a/allreduce.cu +++ b/allreduce.cu @@ -14,6 +14,7 @@ #include #define bfloat16 nv_bfloat16 #elif USE_ROCM + #define __HIP_PLATFORM_AMD__ #include #include #include From e841fe4b37eff6ab91d61864ea8fece4fd9c4952 Mon Sep 17 00:00:00 2001 From: Aditya Tomar Date: Fri, 26 Apr 2024 10:53:36 -0700 Subject: [PATCH 16/16] update --- allreduce.cu | 22 +++++++++------------- 1 file changed, 9 insertions(+), 13 deletions(-) diff --git a/allreduce.cu b/allreduce.cu index 3b31238..da62eb1 100644 --- a/allreduce.cu +++ b/allreduce.cu @@ -30,9 +30,9 @@ #define NUM_WARMUP_ITERATIONS 5 #define MPI_CHECK(cmd) do { \ - int64_t e = cmd; \ + int64_t e = cmd; \ if( e != MPI_SUCCESS ) { \ - printf("Failed: MPI error %s:%d '%ld'\n", \ + printf("Failed: MPI error %s:%d '%ld'\n", \ __FILE__,__LINE__, e); \ exit(EXIT_FAILURE); \ } \ @@ -47,11 +47,11 @@ } \ } while(0) -#define HIP_CHECK(cmd) do { \ - hipError_t e = cmd; \ - if(e != hipSuccess) { \ - printf("HIP error %s:%d: %s\n", \ - __FILE__, __LINE__, hipGetErrorString(e)); \ +#define HIP_CHECK(cmd) do { \ + hipError_t e = cmd; \ + if(e != hipSuccess) { \ + printf("HIP error %s:%d: %s\n", \ + __FILE__, __LINE__, hipGetErrorString(e)); \ exit(EXIT_FAILURE); \ } \ } while(0) @@ -72,8 +72,7 @@ void initializeData(bfloat16 *data, int64_t size) { data[i] = __float2bfloat16((float)i); #elif USE_ROCM // ROCm doesn't have a float2bfloat16 method - data[i] = (bfloat16) ((float) i); - #endif + data[i] = (bfloat16) ((float) i); #endif } } @@ -182,7 +181,7 @@ int main(int argc, char *argv[]) { NCCL_CHECK(ncclCommInitRank(&nccl_comm, num_pes, nccl_comm_id, my_rank)); #endif - // Perform MPI_Iallgather, NCCL allgather, or RCCL allgather + // Perform MPI_Iallreduce, NCCL allreduce, or RCCL allreduce double total_time, start_time; MPI_Request request; MPI_Status status; @@ -215,9 +214,6 @@ int main(int argc, char *argv[]) { #endif } - if(msg_size >= 8388608) - iterations = 20; - MPI_Barrier(MPI_COMM_WORLD); start_time = MPI_Wtime(); for (int i = 0; i < iterations; ++i) {