From 85bdda38120875455475b0ec5269b395d38fb45f Mon Sep 17 00:00:00 2001 From: Corey Derochie Date: Wed, 1 May 2024 20:41:12 -0500 Subject: [PATCH 1/2] Wrapped the warmup iters in captures when doing graph mode to do a proper warmup. [ROCm/rccl-tests commit: 0c762d210c29988c4572dc7ffb738103bec5c759] --- projects/rccl-tests/src/common.cu | 88 +++++++++++++++++++++++++++++++ 1 file changed, 88 insertions(+) diff --git a/projects/rccl-tests/src/common.cu b/projects/rccl-tests/src/common.cu index 6e1306830f..bae774abc9 100644 --- a/projects/rccl-tests/src/common.cu +++ b/projects/rccl-tests/src/common.cu @@ -647,18 +647,106 @@ testResult_t TimeTest(struct threadArgs* args, ncclDataType_t type, const char* // Warm-up for large size setupArgs(args->maxbytes, type, args); +#if HIP_VERSION >= 50221310 + cudaGraph_t graphs[args->nGpus]; + cudaGraphExec_t graphExec[args->nGpus]; + if (cudaGraphLaunches >= 1) { + // Begin cuda graph capture + for (int i=0; inGpus; i++) { + // Thread local mdoe is needed for: + // - Multi-thread mode: where graph capture and instantiation can happen concurrently across threads + // - P2P pre-connect: when there is no warm-up, P2P pre-connect is done during graph capture. + // Since pre-connect calls cudaMalloc, we cannot use global capture mode + CUDACHECK(cudaStreamBeginCapture(args->streams[i], cudaStreamCaptureModeThreadLocal)); + } + } +#endif for (int iter = 0; iter < warmup_iters; iter++) { TESTCHECK(startColl(args, type, op, root, 0, iter)); } + +#if HIP_VERSION >= 50221310 + if (cudaGraphLaunches >= 1) { + // End cuda graph capture + for (int i=0; inGpus; i++) { + CUDACHECK(cudaStreamEndCapture(args->streams[i], graphs+i)); + } + // Instantiate cuda graph + for (int i=0; inGpus; i++) { + CUDACHECK(cudaGraphInstantiate(graphExec+i, graphs[i], NULL, NULL, 0)); + } + // Resync CPU, restart timing, launch cuda graph + Barrier(args); + for (int l=0; lnGpus; i++) { + CUDACHECK(cudaGraphLaunch(graphExec[i], args->streams[i])); + } + } + } +#endif + TESTCHECK(completeColl(args)); +#if HIP_VERSION >= 50221310 + if (cudaGraphLaunches >= 1) { + //destroy cuda graph + for (int i=0; inGpus; i++) { + CUDACHECK(cudaGraphExecDestroy(graphExec[i])); + CUDACHECK(cudaGraphDestroy(graphs[i])); + } + } +#endif + // Warm-up for small size setupArgs(args->minbytes, type, args); +#if HIP_VERSION >= 50221310 + if (cudaGraphLaunches >= 1) { + // Begin cuda graph capture + for (int i=0; inGpus; i++) { + // Thread local mdoe is needed for: + // - Multi-thread mode: where graph capture and instantiation can happen concurrently across threads + // - P2P pre-connect: when there is no warm-up, P2P pre-connect is done during graph capture. + // Since pre-connect calls cudaMalloc, we cannot use global capture mode + CUDACHECK(cudaStreamBeginCapture(args->streams[i], cudaStreamCaptureModeThreadLocal)); + } + } +#endif for (int iter = 0; iter < warmup_iters; iter++) { TESTCHECK(startColl(args, type, op, root, iter < warmup_iters/2 ? 0 : 1, iter)); } + +#if HIP_VERSION >= 50221310 + if (cudaGraphLaunches >= 1) { + // End cuda graph capture + for (int i=0; inGpus; i++) { + CUDACHECK(cudaStreamEndCapture(args->streams[i], graphs+i)); + } + // Instantiate cuda graph + for (int i=0; inGpus; i++) { + CUDACHECK(cudaGraphInstantiate(graphExec+i, graphs[i], NULL, NULL, 0)); + } + // Resync CPU, restart timing, launch cuda graph + Barrier(args); + for (int l=0; lnGpus; i++) { + CUDACHECK(cudaGraphLaunch(graphExec[i], args->streams[i])); + } + } + } +#endif + TESTCHECK(completeColl(args)); +#if HIP_VERSION >= 50221310 + if (cudaGraphLaunches >= 1) { + //destroy cuda graph + for (int i=0; inGpus; i++) { + CUDACHECK(cudaGraphExecDestroy(graphExec[i])); + CUDACHECK(cudaGraphDestroy(graphs[i])); + } + } +#endif + for (size_t iter = 0; iter < stress_cycles; iter++) { if (iter > 0) PRINT("# Testing %lu cycle.\n", iter+1); // Benchmark From fe151f517b657ff6347e98f2da04379d13fb6e60 Mon Sep 17 00:00:00 2001 From: corey-derochie-amd <161367113+corey-derochie-amd@users.noreply.github.com> Date: Thu, 2 May 2024 09:18:25 -0600 Subject: [PATCH 2/2] Fixed spelling [ROCm/rccl-tests commit: f74c04b6863ff727ac8b000ed68bbb0df0c79f0c] --- projects/rccl-tests/src/common.cu | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/projects/rccl-tests/src/common.cu b/projects/rccl-tests/src/common.cu index bae774abc9..def6ca029e 100644 --- a/projects/rccl-tests/src/common.cu +++ b/projects/rccl-tests/src/common.cu @@ -653,7 +653,7 @@ testResult_t TimeTest(struct threadArgs* args, ncclDataType_t type, const char* if (cudaGraphLaunches >= 1) { // Begin cuda graph capture for (int i=0; inGpus; i++) { - // Thread local mdoe is needed for: + // Thread local mode is needed for: // - Multi-thread mode: where graph capture and instantiation can happen concurrently across threads // - P2P pre-connect: when there is no warm-up, P2P pre-connect is done during graph capture. // Since pre-connect calls cudaMalloc, we cannot use global capture mode @@ -703,7 +703,7 @@ testResult_t TimeTest(struct threadArgs* args, ncclDataType_t type, const char* if (cudaGraphLaunches >= 1) { // Begin cuda graph capture for (int i=0; inGpus; i++) { - // Thread local mdoe is needed for: + // Thread local mode is needed for: // - Multi-thread mode: where graph capture and instantiation can happen concurrently across threads // - P2P pre-connect: when there is no warm-up, P2P pre-connect is done during graph capture. // Since pre-connect calls cudaMalloc, we cannot use global capture mode