D2H background traffic interferes with 4-GPU allpairs cudaMemcpyPeerAsync bandwidth

Accelerated Computing CUDA CUDA Programming and Performance
1

CUDA-only D2D Peer Copy Interference from D2H Traffic

I am seeing a large D2D peer-copy bandwidth drop only in one CUDA-only case. The tests use cudaMemcpyPeerAsync for D2D, cudaMemcpyAsync with pinned host memory for D2H/H2D, and CUDA events for D2D timing.

Test machine: 4x Tesla V100-SXM2-32GB, all tested GPU pairs are NV2, driver 580.159.03, CUDA 12.6.

The system has 4 Tesla V100-SXM2-32GB GPUs, and every GPU pair is connected by NV2 links.

Cases tested

  1. 4-GPU allpairs D2D

    Every ordered GPU pair copies to every other GPU (GPU i -> GPU j, i != j).

    Initial result: baseline is about 180.97 GB/s. With D2H background traffic, D2D drops to about 127.97 GB/s; D2H background is about 10.6-13.4 GB/s aggregate. With H2D background traffic using the same 255 MiB copy size, there is no noticeable D2D bandwidth drop.

  2. 4-GPU ring D2D

    Ring pattern: GPU i -> GPU (i + 1) % 4.

    Initial result: no noticeable D2D bandwidth drop with either D2H or H2D background traffic.

  3. 2-GPU D2D

    Two directed peer copies are used: GPU0 -> GPU1 and GPU1 -> GPU0.

    Initial result: no noticeable D2D bandwidth drop with either D2H or H2D background traffic.

Minimal commands

 # 4-GPU allpairs baseline
 ./d2d_peer_bw --devices=4 --pattern=allpairs --size=255M --iters=120 --warmup=10 --no-sync-each-iter
 ​
 # 4-GPU ring baseline
 ./d2d_peer_bw --devices=4 --pattern=ring --size=255M --iters=500 --warmup=20 --no-sync-each-iter
 ​
 # D2H background, run in another terminal
 ./pcie_background_copy --devices=4 --direction=d2h --size=255M
 ​
 # H2D background, run in another terminal
 ./pcie_background_copy --devices=4 --direction=h2d --size=255M

Question

Why does sustained D2H traffic significantly reduce 4-GPU allpairs D2D peer-copy bandwidth, while H2D traffic with the same copy size does not? Also, why are the 4-GPU ring and 2-GPU D2D cases almost unaffected by either D2H or H2D background traffic?

My guess is that allpairs creates enough source-side peer-copy pressure to become sensitive to GPU-local memory/copy/fabric arbitration, and D2H competes with that path more directly than H2D. Is this expected behavior, and which GPU resources are likely being contended?

Source code

d2d_peer_bw.cu

 #include <cuda_runtime.h>
 ​
 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
 #include <string>
 #include <vector>
 ​
 #define CHECK_CUDA(call)                                                     \
   do {                                                                       \
     cudaError_t err__ = (call);                                               \
     if (err__ != cudaSuccess) {                                               \
       std::fprintf(stderr, "CUDA error at %s:%d: %s\n", __FILE__, __LINE__,  \
                    cudaGetErrorString(err__));                                \
       return 1;                                                               \
     }                                                                         \
   } while (0)
 ​
 struct Options {
   int devices = 4;
   size_t bytes = 255ULL * 1024ULL * 1024ULL;
   int warmup = 20;
   int iters = 500;
   std::string pattern = "ring";
   bool sync_each_iter = true;
 };
 ​
 static size_t parse_size(const char* s) {
   char* end = nullptr;
   double value = std::strtod(s, &end);
   if (end == s || value <= 0) return 0;
   double scale = 1.0;
   if (*end == 'k' || *end == 'K') scale = 1024.0;
   if (*end == 'm' || *end == 'M') scale = 1024.0 * 1024.0;
   if (*end == 'g' || *end == 'G') scale = 1024.0 * 1024.0 * 1024.0;
   return static_cast<size_t>(value * scale);
 }
 ​
 static void usage(const char* prog) {
   std::printf(
       "Usage: %s [--devices=N] [--pattern=ring|allpairs]\n"
       "          [--size=255M] [--warmup=20] [--iters=500] [--no-sync-each-iter]\n\n"
       "Patterns:\n"
       "  ring     GPU i -> GPU (i+1) %% devices (default)\n"
       "  allpairs every ordered GPU pair i -> j, i != j\n",
       prog);
 }
 ​
 static int parse_args(int argc, char** argv, Options* opt) {
   for (int i = 1; i < argc; ++i) {
     if (std::strncmp(argv[i], "--devices=", 10) == 0) {
       opt->devices = std::atoi(argv[i] + 10);
     } else if (std::strncmp(argv[i], "--pattern=", 10) == 0) {
       opt->pattern = argv[i] + 10;
     } else if (std::strncmp(argv[i], "--size=", 7) == 0) {
       opt->bytes = parse_size(argv[i] + 7);
     } else if (std::strncmp(argv[i], "--warmup=", 9) == 0) {
       opt->warmup = std::atoi(argv[i] + 9);
     } else if (std::strncmp(argv[i], "--iters=", 8) == 0) {
       opt->iters = std::atoi(argv[i] + 8);
     } else if (std::strcmp(argv[i], "--no-sync-each-iter") == 0) {
       opt->sync_each_iter = false;
     } else if (std::strcmp(argv[i], "-h") == 0 ||
                std::strcmp(argv[i], "--help") == 0) {
       usage(argv[0]);
       std::exit(0);
     } else {
       std::fprintf(stderr, "Unknown option: %s\n", argv[i]);
       usage(argv[0]);
       return 1;
     }
   }
 ​
   if (opt->devices < 2 || opt->bytes == 0 || opt->warmup < 0 ||
       opt->iters <= 0 ||
       (opt->pattern != "ring" && opt->pattern != "allpairs")) {
     usage(argv[0]);
     return 1;
   }
   return 0;
 }
 ​
 static int enable_peer(int dst, int src) {
   int can_access = 0;
   CHECK_CUDA(cudaDeviceCanAccessPeer(&can_access, dst, src));
   if (!can_access) {
     std::fprintf(stderr, "GPU %d cannot directly access GPU %d\n", dst, src);
     return 1;
   }
   CHECK_CUDA(cudaSetDevice(dst));
   cudaError_t err = cudaDeviceEnablePeerAccess(src, 0);
   if (err == cudaErrorPeerAccessAlreadyEnabled) {
     cudaGetLastError();
     return 0;
   }
   if (err != cudaSuccess) {
     std::fprintf(stderr, "cudaDeviceEnablePeerAccess(%d -> %d) failed: %s\n",
                  dst, src, cudaGetErrorString(err));
     return 1;
   }
   return 0;
 }
 ​
 static int issue_copies(const Options& opt, const std::vector<void*>& src,
                         const std::vector<std::vector<void*> >& dst,
                         const std::vector<cudaStream_t>& streams) {
   for (int s = 0; s < opt.devices; ++s) {
     CHECK_CUDA(cudaSetDevice(s));
     if (opt.pattern == "ring") {
       int d = (s + 1) % opt.devices;
       CHECK_CUDA(cudaMemcpyPeerAsync(dst[d][s], d, src[s], s, opt.bytes,
                                      streams[s]));
     } else {
       for (int d = 0; d < opt.devices; ++d) {
         if (d == s) continue;
         CHECK_CUDA(cudaMemcpyPeerAsync(dst[d][s], d, src[s], s, opt.bytes,
                                        streams[s]));
       }
     }
   }
   return 0;
 }
 ​
 int main(int argc, char** argv) {
   Options opt;
   if (parse_args(argc, argv, &opt) != 0) return 1;
 ​
   int device_count = 0;
   CHECK_CUDA(cudaGetDeviceCount(&device_count));
   if (opt.devices > device_count) {
     std::fprintf(stderr, "Requested devices=%d, device_count=%d\n",
                  opt.devices, device_count);
     return 1;
   }
 ​
   std::printf("D2D peer memcpy benchmark\n");
   std::printf("devices=%d pattern=%s size=%zu bytes (%.2f MiB) warmup=%d iters=%d sync_each_iter=%s\n",
               opt.devices, opt.pattern.c_str(), opt.bytes,
               opt.bytes / 1024.0 / 1024.0, opt.warmup, opt.iters,
               opt.sync_each_iter ? "yes" : "no");
   for (int d = 0; d < opt.devices; ++d) {
     cudaDeviceProp prop;
     CHECK_CUDA(cudaGetDeviceProperties(&prop, d));
     std::printf("GPU %d: %s\n", d, prop.name);
   }
 ​
   for (int d = 0; d < opt.devices; ++d) {
     for (int s = 0; s < opt.devices; ++s) {
       if (d != s && enable_peer(d, s) != 0) return 1;
     }
   }
 ​
   std::vector<void*> src(opt.devices, nullptr);
   std::vector<std::vector<void*> > dst(opt.devices,
                                        std::vector<void*>(opt.devices, nullptr));
   std::vector<cudaStream_t> streams(opt.devices, nullptr);
   std::vector<cudaEvent_t> start(opt.devices, nullptr);
   std::vector<cudaEvent_t> stop(opt.devices, nullptr);
 ​
   for (int d = 0; d < opt.devices; ++d) {
     CHECK_CUDA(cudaSetDevice(d));
     CHECK_CUDA(cudaMalloc(&src[d], opt.bytes));
     CHECK_CUDA(cudaMemset(src[d], d & 0xff, opt.bytes));
     for (int s = 0; s < opt.devices; ++s) {
       if (s != d) CHECK_CUDA(cudaMalloc(&dst[d][s], opt.bytes));
     }
     CHECK_CUDA(cudaStreamCreate(&streams[d]));
     CHECK_CUDA(cudaEventCreate(&start[d]));
     CHECK_CUDA(cudaEventCreate(&stop[d]));
   }
 ​
   for (int i = 0; i < opt.warmup; ++i) {
     if (issue_copies(opt, src, dst, streams) != 0) return 1;
     for (int d = 0; d < opt.devices; ++d) {
       CHECK_CUDA(cudaSetDevice(d));
       CHECK_CUDA(cudaStreamSynchronize(streams[d]));
     }
   }
 ​
   for (int d = 0; d < opt.devices; ++d) {
     CHECK_CUDA(cudaSetDevice(d));
     CHECK_CUDA(cudaEventRecord(start[d], streams[d]));
   }
   for (int i = 0; i < opt.iters; ++i) {
     if (issue_copies(opt, src, dst, streams) != 0) return 1;
     if (opt.sync_each_iter) {
       for (int d = 0; d < opt.devices; ++d) {
         CHECK_CUDA(cudaSetDevice(d));
         CHECK_CUDA(cudaStreamSynchronize(streams[d]));
       }
     }
   }
   for (int d = 0; d < opt.devices; ++d) {
     CHECK_CUDA(cudaSetDevice(d));
     CHECK_CUDA(cudaEventRecord(stop[d], streams[d]));
   }
   for (int d = 0; d < opt.devices; ++d) {
     CHECK_CUDA(cudaSetDevice(d));
     CHECK_CUDA(cudaEventSynchronize(stop[d]));
   }
 ​
   double max_sec = 0.0;
   for (int d = 0; d < opt.devices; ++d) {
     float ms = 0.0f;
     CHECK_CUDA(cudaSetDevice(d));
     CHECK_CUDA(cudaEventElapsedTime(&ms, start[d], stop[d]));
     double sec = ms / 1000.0;
     if (sec > max_sec) max_sec = sec;
     std::printf("Elapsed GPU %d source stream: %.3f ms\n", d, ms);
   }
 ​
   int copies_per_iter = opt.devices;
   if (opt.pattern == "allpairs") copies_per_iter = opt.devices * (opt.devices - 1);
   double total_bytes = static_cast<double>(opt.bytes) * opt.iters * copies_per_iter;
   double aggregate_gbs = total_bytes / max_sec / 1e9;
   std::printf("Directed copies per iteration: %d\n", copies_per_iter);
   std::printf("Aggregate D2D bandwidth: %.2f GB/s\n", aggregate_gbs);
   std::printf("Per-copy average bandwidth: %.2f GB/s\n", aggregate_gbs / copies_per_iter);
 ​
   for (int d = 0; d < opt.devices; ++d) {
     CHECK_CUDA(cudaSetDevice(d));
     cudaEventDestroy(start[d]);
     cudaEventDestroy(stop[d]);
     cudaStreamDestroy(streams[d]);
     cudaFree(src[d]);
     for (int s = 0; s < opt.devices; ++s) {
       if (dst[d][s]) cudaFree(dst[d][s]);
     }
   }
   return 0;
 }

pcie_background_copy.cu

 #include <cuda_runtime.h>
 ​
 #include <atomic>
 #include <chrono>
 #include <cstdio>
 #include <cstdlib>
 #include <cstring>
 #include <thread>
 #include <vector>
 ​
 #define CHECK_CUDA_OR_FAIL(call)                                              \
   do {                                                                        \
     cudaError_t err__ = (call);                                                \
     if (err__ != cudaSuccess) {                                                \
       std::fprintf(stderr, "CUDA error at %s:%d: %s\n", __FILE__, __LINE__,   \
                    cudaGetErrorString(err__));                                 \
       failed->store(true, std::memory_order_release);                          \
       return;                                                                  \
     }                                                                          \
   } while (0)
 ​
 struct Options {
   int devices = 2;
   int dev0 = 0;
   int dev1 = 1;
   bool use_pair = false;
   size_t bytes = 255ULL * 1024ULL * 1024ULL;
   int warmup = 5;
   int report_sec = 2;
   const char* direction = "d2h";
 };
 ​
 static size_t parse_size(const char* s) {
   char* end = nullptr;
   double value = std::strtod(s, &end);
   if (end == s || value <= 0) return 0;
   double scale = 1.0;
   if (*end == 'k' || *end == 'K') scale = 1024.0;
   if (*end == 'm' || *end == 'M') scale = 1024.0 * 1024.0;
   if (*end == 'g' || *end == 'G') scale = 1024.0 * 1024.0 * 1024.0;
   return static_cast<size_t>(value * scale);
 }
 ​
 static void usage(const char* prog) {
   std::printf(
       "Usage: %s [--devices=N | --dev0=0 --dev1=1] [--size=255M]\n"
       "          [--direction=d2h|h2d] [--warmup=5] [--report-sec=2]\n",
       prog);
 }
 ​
 static int parse_args(int argc, char** argv, Options* opt) {
   for (int i = 1; i < argc; ++i) {
     if (std::strncmp(argv[i], "--devices=", 10) == 0) {
       opt->devices = std::atoi(argv[i] + 10);
       opt->use_pair = false;
     } else if (std::strncmp(argv[i], "--dev0=", 7) == 0) {
       opt->dev0 = std::atoi(argv[i] + 7);
       opt->use_pair = true;
     } else if (std::strncmp(argv[i], "--dev1=", 7) == 0) {
       opt->dev1 = std::atoi(argv[i] + 7);
       opt->use_pair = true;
     } else if (std::strncmp(argv[i], "--size=", 7) == 0) {
       opt->bytes = parse_size(argv[i] + 7);
     } else if (std::strncmp(argv[i], "--direction=", 12) == 0) {
       opt->direction = argv[i] + 12;
     } else if (std::strncmp(argv[i], "--warmup=", 9) == 0) {
       opt->warmup = std::atoi(argv[i] + 9);
     } else if (std::strncmp(argv[i], "--report-sec=", 13) == 0) {
       opt->report_sec = std::atoi(argv[i] + 13);
     } else if (std::strcmp(argv[i], "-h") == 0 ||
                std::strcmp(argv[i], "--help") == 0) {
       usage(argv[0]);
       std::exit(0);
     } else {
       std::fprintf(stderr, "Unknown option: %s\n", argv[i]);
       usage(argv[0]);
       return 1;
     }
   }
 ​
   if (opt->devices <= 0 || opt->dev0 == opt->dev1 || opt->bytes == 0 ||
       opt->warmup < 0 || opt->report_sec <= 0 ||
       (std::strcmp(opt->direction, "d2h") != 0 &&
        std::strcmp(opt->direction, "h2d") != 0)) {
     usage(argv[0]);
     return 1;
   }
   return 0;
 }
 ​
 static void worker(int device, const Options opt,
                    std::atomic<unsigned long long>* total_bytes,
                    std::atomic<int>* ready, std::atomic<bool>* go,
                    std::atomic<bool>* failed) {
   CHECK_CUDA_OR_FAIL(cudaSetDevice(device));
   void* h = nullptr;
   void* d = nullptr;
   cudaStream_t stream = nullptr;
   CHECK_CUDA_OR_FAIL(cudaMallocHost(&h, opt.bytes));
   CHECK_CUDA_OR_FAIL(cudaMalloc(&d, opt.bytes));
   CHECK_CUDA_OR_FAIL(cudaStreamCreate(&stream));
   CHECK_CUDA_OR_FAIL(cudaMemset(d, device & 0xff, opt.bytes));
   std::memset(h, device & 0xff, opt.bytes);
 ​
   cudaMemcpyKind kind = std::strcmp(opt.direction, "d2h") == 0
                             ? cudaMemcpyDeviceToHost
                             : cudaMemcpyHostToDevice;
   void* dst = kind == cudaMemcpyDeviceToHost ? h : d;
   void* src = kind == cudaMemcpyDeviceToHost ? d : h;
 ​
   for (int i = 0; i < opt.warmup; ++i) {
     CHECK_CUDA_OR_FAIL(cudaMemcpyAsync(dst, src, opt.bytes, kind, stream));
     CHECK_CUDA_OR_FAIL(cudaStreamSynchronize(stream));
   }
 ​
   ready->fetch_add(1, std::memory_order_acq_rel);
   while (!go->load(std::memory_order_acquire) &&
          !failed->load(std::memory_order_acquire)) {
     std::this_thread::yield();
   }
 ​
   while (!failed->load(std::memory_order_acquire)) {
     CHECK_CUDA_OR_FAIL(cudaMemcpyAsync(dst, src, opt.bytes, kind, stream));
     CHECK_CUDA_OR_FAIL(cudaStreamSynchronize(stream));
     total_bytes->fetch_add(opt.bytes, std::memory_order_relaxed);
   }
 }
 ​
 int main(int argc, char** argv) {
   Options opt;
   if (parse_args(argc, argv, &opt) != 0) return 1;
 ​
   int device_count = 0;
   cudaError_t err = cudaGetDeviceCount(&device_count);
   if (err != cudaSuccess) {
     std::fprintf(stderr, "cudaGetDeviceCount failed: %s\n", cudaGetErrorString(err));
     return 1;
   }
 ​
   std::vector<int> devices;
   if (opt.use_pair) {
     devices.push_back(opt.dev0);
     devices.push_back(opt.dev1);
   } else {
     for (int d = 0; d < opt.devices; ++d) devices.push_back(d);
   }
   for (int d : devices) {
     if (d < 0 || d >= device_count) {
       std::fprintf(stderr, "Requested GPU %d but device_count=%d\n", d, device_count);
       return 1;
     }
   }
 ​
   std::printf("PCIe background copy\n");
   std::printf("devices=");
   for (size_t i = 0; i < devices.size(); ++i) std::printf("%s%d", i ? "," : "", devices[i]);
   std::printf(" direction=%s size=%zu bytes (%.2f MiB)\n", opt.direction,
               opt.bytes, opt.bytes / 1024.0 / 1024.0);
   std::printf("Each GPU uses one pinned host buffer, one device buffer, one stream, and repeats memcpy+stream sync forever.\n");
   std::fflush(stdout);
 ​
   std::atomic<unsigned long long> total_bytes{0};
   std::atomic<int> ready{0};
   std::atomic<bool> go{false};
   std::atomic<bool> failed{false};
   std::vector<std::thread> threads;
   for (int d : devices) threads.emplace_back(worker, d, opt, &total_bytes, &ready, &go, &failed);
 ​
   while (ready.load(std::memory_order_acquire) < static_cast<int>(devices.size()) &&
          !failed.load(std::memory_order_acquire)) {
     std::this_thread::sleep_for(std::chrono::milliseconds(10));
   }
   go.store(true, std::memory_order_release);
 ​
   using clock = std::chrono::steady_clock;
   unsigned long long last_bytes = total_bytes.load(std::memory_order_relaxed);
   auto last = clock::now();
   while (!failed.load(std::memory_order_acquire)) {
     std::this_thread::sleep_for(std::chrono::seconds(opt.report_sec));
     auto now = clock::now();
     unsigned long long current = total_bytes.load(std::memory_order_relaxed);
     double seconds = std::chrono::duration<double>(now - last).count();
     double gbps = static_cast<double>(current - last_bytes) / seconds / 1e9;
     std::printf("bytes=%llu window_bytes=%llu time=%.3f s aggregate_bw=%.2f GB/s\n",
                 current, current - last_bytes, seconds, gbps);
     std::fflush(stdout);
     last = now;
     last_bytes = current;
   }
 ​
   for (auto& t : threads) t.join();
   return failed.load() ? 1 : 0;
 }
2

I would like to add an update to clarify the motivation of this issue and the latest results.

The original reason I started looking into this CUDA-only D2D case was that I first observed a performance degradation in nccl-tests under D2H background traffic, across different message sizes. To better understand whether the slowdown was caused by general GPU-to-GPU data movement interference, I then tried to build a CUDA-only D2D memcpy microbenchmark as an ablation experiment.

During that ablation, I found the issue reported in this thread: the D2D peer-copy bandwidth drop only appears in the 4-GPU allpairs cudaMemcpyPeerAsync case, while the 2-GPU case and the 4-GPU ring case are almost unaffected. Based on the current results, this D2D behavior seems more like a bandwidth/arbitration bottleneck that only shows up when the full-mesh traffic pattern creates enough pressure.

I also later tested the NCCL case on an H200 server. The NCCL performance degradation can still be reproduced there under D2H background traffic, but it seems to have a different cause from the CUDA-only D2D allpairs case. In our preliminary analysis of NCCL’s internal synchronization mechanism, we found that some system-scope synchronization operations, especially the sys-scope fence used in the NCCL pipeline, become significantly slower when concurrent D2H traffic is running on the same physical GPU.

Therefore, my current understanding is that the CUDA-only D2D allpairs slowdown in this thread and the NCCL slowdown are related observations triggered by D2H background traffic, but they may not be caused by the exact same mechanism. The D2D issue seems to be exposed only under the allpairs full-mesh traffic pattern, while the NCCL issue appears to be more related to NCCL’s internal kernel synchronization path on newer H-series GPUs as well.