feat: add kernel-based copy mode to nvlink_bw_test

[staryxchen]

Summary

Add a __global__ copy kernel as an alternative to cudaMemcpyDeviceToDevice in example/nvlink_bw_test, selectable via a new -m/--mode memcpy|kernel flag (default: memcpy). The kernel uses vectorized uint4 (16-byte) load/stores and launches on the source GPU, writing to the destination GPU over P2P/NVLink.

This produces NVLink traffic that nvlink_monitor can observe, mirroring the memcpy path but exercising the SMs instead of the copy engine — useful for validating that NVML counters pick up SM-initiated traffic (they do).

Changes

New files

• example/copy_kernel.h — declares launchCopyKernel()
• example/copy_kernel.cu — __global__ copyKernel (uint4, 256 threads/block) + launch wrapper. The only file compiled by nvcc; g++ compiles everything else and does the final link.

Modified

• example/arg_parser.{h,cpp} — CopyMode enum + mode field + -m/--mode parsing. Stays CUDA-free (arg_parser.h is included by the GPU-less test binary).
• example/nvlink_bw_test.cpp — measureCopyTimeKernel() mirrors measureCopyTime()'s event timing; testCopyPerformance() dispatches on mode; mode printed in Configuration; -m in printUsage().
• Makefile — NVCC/NVCC_ARCH (sm_70/sm_80/sm_90 + compute_90 PTX for forward-compat), build/copy_kernel.o rule, linked into nvlink_bw_test. make test/check-format stay nvcc-free.
• test/test_bw_test_args.cpp — 6 new --mode test cases.
• README.md — documents -m/--mode and the nvcc requirement.

Verification

• make check-format — clean (clang-format 14)
• make test — 58/58 pass (52 existing + 6 new mode tests), no GPU/nvcc needed
• make example — builds with nvcc 12.8 + g++ link
• End-to-end on 8x H20 (sm_90): kernel mode ~266 GiB/s avg, memcpy ~285 GiB/s avg (100MB D2D). Kernel slightly slower as expected (SM load/store vs copy-engine DMA).
• Cross-checked against nvlink_monitor: with a 3.5s memcpy test (500×2GB), bw_test reported 369.45 GiB/s vs nvlink_monitor steady-state 368.4 GiB/s — within 0.3%, confirming both tools agree when the test spans multiple monitor sampling windows.

Notes

• CopyMode enum deliberately lives in arg_parser.h (CUDA-free) so CI's make test (no CUDA toolkit) still builds.
• nvcc multi-arch covers V100/A100/H100 + PTX for future GPUs.

🤖 Generated with CodeBuddy Code