Repository: pytorch
Issue: CUDA Invalid Memory Access caused by torch.distributed.barrier() with NCCL
Labels: oncall: distributed,triaged

### 🐛 Describe the bug

```
TORCH_NCCL_AVOID_RECORD_STREAMS=1 PYTORCH_NO_CUDA_MEMORY_CACHING=1 torchrun --standalone --nnodes=1 --nproc_per_node=2 --no-python compute-sanitizer python -c 'import torch; torch.distributed.init_process_group(backend="nccl"); torch.distributed.barrier(device_ids=[torch.distributed.get_rank()])'
```
causes
```
========= COMPUTE-SANITIZER
========= Invalid __global__ read of size 4 bytes
=========     at ncclDevKernel_AllReduce_Sum_u8_RING_LL(ncclDevComm *, unsigned long, ncclWork *)+0x1ad0
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0x7f4ac0200000 is out of bounds
=========     and is inside the nearest allocation at 0x7f4ac0200000 of size 1 bytes
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame: [0x332833]
=========                in /lib/x86_64-linux-gnu/libcuda.so.1
=========     Host Frame: [0x14cb8]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/../../../../libcudart.so.12
=========     Host Frame:cudaLaunchKernelExC [0x6c163]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/../../../../libcudart.so.12
=========     Host Frame:ncclLaunchKernel(ncclComm*, ncclKernelPlan*) in /opt/conda/conda-bld/pytorch_1721288503779/work/third_party/nccl/nccl/src/enqueue.cc:1402 [0x3964581]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so
=========     Host Frame:groupLaunch(ncclAsyncJob*) in /opt/conda/conda-bld/pytorch_1721288503779/work/third_party/nccl/nccl/src/group.cc:339 [0x396bb42]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so
=========     Host Frame:ncclGroupEndInternal() in /opt/conda/conda-bld/pytorch_1721288503779/work/third_party/nccl/nccl/src/group.cc:368 [0x396c707]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so
=========     Host Frame:ncclEnqueueCheck(ncclInfo*) in /opt/conda/conda-bld/pytorch_1721288503779/work/third_party/nccl/nccl/src/enqueue.cc:2032 [0x3964dae]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so
=========     Host Frame:ncclAllReduce in /opt/conda/conda-bld/pytorch_1721288503779/work/third_party/nccl/nccl/src/collectives.cc:50 [0x39566ce]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so
=========     Host Frame:c10::detail::intrusive_target_default_null_type<c10::intrusive_ptr> c10d::ProcessGroupNCCL::collective<c10d::ProcessGroupNCCL::allreduce_impl(at::Tensor&, c10d::AllreduceOptions const&)::{lambda(at::Tensor&, at::Tensor&, ncclComm*, c10::cuda::CUDAStream&)#1}, c10d::ProcessGroupNCCL::collective<{lambda(at::Tensor&, at::Tensor&, ncclComm*, c10::cuda::CUDAStream&)#1}>(at::Tensor&, at::Tensor&, {lambda(at::Tensor&, at::Tensor&, ncclComm*, c10::cuda::CUDAStream&)#1}, c10d::OpType, char const*, bool)::{lambda(c10::cuda::CUDAStream&, c10d::ProcessGroupNCCL::collective<c10d::ProcessGroupNCCL::WorkNCCL, c10::detail<char const*> >&)#1}, c10d::ProcessGroupNCCL::allreduce_impl(at::Tensor&, c10d::AllreduceOptions const&)::{lambda(at::Tensor&, at::Tensor&, ncclComm*, c10::cuda::CUDAStream&)#1}<{lambda(at::Tensor&, at::Tensor&, ncclComm*, c10::cuda::CUDAStream&)#1}>(at::Tensor&, at::Tensor&, c10::intrusive_ptr<c10d::Work, c10::detail::intrusive_target_default_null_type<c10::intrusive_ptr> >, {lambda(at::Tensor&, at::Tensor&, ncclComm*, c10::cuda::CUDAStream&)#1}, char const, bool)::{lambda(c10::cuda::CUDAStream&, c10d::ProcessGroupNCCL::collective<c10d::ProcessGroupNCCL::WorkNCCL, c10::detail<char const*> >)#2}>(at::Tensor&, at::Tensor&, c10::intrusive_ptr<c10d::Work, c10::detail::intrusive_target_default_null_type<c10::intrusive_ptr> >, c10d::ProcessGroupNCCL::collective<{lambda(at::Tensor&, at::Tensor&, ncclComm*, c10::cuda::CUDAStream&)#1}>(at::Tensor&, at::Tensor&, {lambda(at::Tensor&, at::Tensor&, ncclComm*, c10::cuda::CUDAStream&)#1}, c10d::OpType, char const*, bool)::{lambda(c10::cuda::CUDAStream&, c10d::ProcessGroupNCCL::collective<c10d::ProcessGroupNCCL::WorkNCCL, c10::detail<char const*> >&)#1}, c10d::ProcessGroupNCCL::allreduce_impl(at::Tensor&, c10d::AllreduceOptions const&)::{lambda(at::Tensor&, at::Tensor&, ncclComm*, c10::cuda::CUDAStream&)#1}<{lambda(at::Tensor&, at::Tensor&, ncclComm*, c10::cuda::CUDAStream&)#1}>(at::Tensor&, at::Tensor&, c10::intrusive_ptr<c10d::Work, c10::detail::intrusive_target_default_null_type<c10::intrusive_ptr> >, {lambda(at::Tensor&, at::Tensor&, ncclComm*, c10::cuda::CUDAStream&)#1}, char const, bool)::{lambda(c10::cuda::CUDAStream&, c10d::ProcessGroupNCCL::collective<c10d::ProcessGroupNCCL::WorkNCCL, c10::detail<char const*> >)#2}, c10d::ProcessGroupNCCL::allreduce_impl(at::Tensor&, c10d::AllreduceOptions const&)::{lambda(at::Tensor&, at::Tensor&, ncclComm*, c10::cuda::CUDAStream&)#1}, char const, bool) [clone .constprop.0] [0x121bd58]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so
=========     Host Frame:c10d::ProcessGroupNCCL::allreduce_impl(at::Tensor&, c10d::AllreduceOptions const&) [0x121ce3f]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so
=========     Host Frame:c10d::ProcessGroupNCCL::barrier(c10d::BarrierOptions const&) [0x1229d8b]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/libtorch_cuda.so
=========     Host Frame:c10d::ops::(anonymous namespace)::barrierCUDA(at::Tensor, c10::intrusive_ptr<c10d::ProcessGroup, c10::detail::intrusive_target_default_null_type<c10d::ProcessGroup> > const&, std::vector<long, std::allocator<long> > const&, long) [0x597a4c1]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/libtorch_cpu.so
=========     Host Frame:c10::impl::make_boxed_from_unboxed_functor<c10::impl::detail::WrapFunctionIntoRuntimeFunctor_<c10::intrusive_ptr<c10d::Work, c10::detail::intrusive_target_default_null_type<c10d::Work> > (*)(at::Tensor, c10::intrusive_ptr<c10d::ProcessGroup, c10::detail::intrusive_target_default_null_type<c10d::ProcessGroup> > const&, std::vector<long, std::allocator<long> > const&, long), c10::intrusive_ptr<c10d::Work, c10::detail::intrusive_target_default_null_type<c10d::Work> >, c10::guts::typelist::typelist<at::Tensor, c10::intrusive_ptr<c10d::ProcessGroup, c10::detail::intrusive_target_default_null_type<c10d::ProcessGroup> > const&, std::vector<long, std::allocator<long> > const&, long> >, false>::call(c10::OperatorKernel*, c10::OperatorHandle const&, c10::DispatchKeySet, std::vector<c10::IValue, std::allocator<c10::IValue> >*) [0x5984c54]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/libtorch_cpu.so
=========     Host Frame:c10::OperatorHandle::redispatchBoxed(c10::DispatchKeySet, std::vector<c10::IValue, std::allocator<c10::IValue> >*) const [0x4fad03a]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/libtorch_cpu.so
=========     Host Frame:torch::autograd::basicAutogradNotImplementedFallbackImpl(c10::OperatorHandle const&, c10::DispatchKeySet, std::vector<c10::IValue, std::allocator<c10::IValue> >*) [0x4faa8b3]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/libtorch_cpu.so
=========     Host Frame:void c10::BoxedKernel::make_boxed_function<&(anonymous namespace)::autograd_fallback>(c10::OperatorKernel*, c10::OperatorHandle const&, c10::DispatchKeySet, std::vector<c10::IValue, std::allocator<c10::IValue> >*) [0x177b367]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/libtorch_cpu.so
=========     Host Frame:c10::impl::BoxedKernelWrapper<c10::intrusive_ptr<c10d::Work, c10::detail::intrusive_target_default_null_type<c10d::Work> > (at::Tensor, c10::intrusive_ptr<c10d::ProcessGroup, c10::detail::intrusive_target_default_null_type<c10d::ProcessGroup> > const&, std::vector<long, std::allocator<long> > const&, long), void>::call(c10::BoxedKernel const&, c10::OperatorHandle const&, c10::DispatchKeySet, at::Tensor, c10::intrusive_ptr<c10d::ProcessGroup, c10::detail::intrusive_target_default_null_type<c10d::ProcessGroup> > const&, std::vector<long, std::allocator<long> > const&, long) [0x598ebc3]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/libtorch_cpu.so
=========     Host Frame:c10d::ProcessGroup::barrier(c10d::BarrierOptions const&) [0x598f954]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/libtorch_cpu.so
=========     Host Frame:pybind11::cpp_function::initialize<pybind11::cpp_function::initialize<c10::intrusive_ptr<c10d::Work, c10::detail::intrusive_target_default_null_type<c10d::Work> >, c10d::ProcessGroup, c10d::BarrierOptions const&, pybind11::name, pybind11::is_method, pybind11::sibling, pybind11::arg_v, pybind11::call_guard<pybind11::gil_scoped_release> >(c10::intrusive_ptr<c10d::Work, c10::detail::intrusive_target_default_null_type<c10d::Work> > (c10d::ProcessGroup::*)(c10d::BarrierOptions const&), pybind11::name const&, pybind11::is_method const&, pybind11::sibling const&, pybind11::arg_v const&, pybind11::call_guard<pybind11::gil_scoped_release> const&)::{lambda(c10d::ProcessGroup*, c10d::BarrierOptions const&)#1}, c10::intrusive_ptr<c10d::Work, c10::detail::intrusive_target_default_null_type<c10d::Work> >, c10d::ProcessGroup*, c10d::BarrierOptions const&, pybind11::name, pybind11::is_method, pybind11::sibling, pybind11::arg_v, pybind11::call_guard<pybind11::gil_scoped_release> >(pybind11::cpp_function::initialize<c10::intrusive_ptr<c10d::Work, c10::detail::intrusive_target_default_null_type<c10d::Work> >, c10d::ProcessGroup, c10d::BarrierOptions const&, pybind11::name, pybind11::is_method, pybind11::sibling, pybind11::arg_v, pybind11::call_guard<pybind11::gil_scoped_release> >(c10::intrusive_ptr<c10d::Work, c10::detail::intrusive_target_default_null_type<c10d::Work> > (c10d::ProcessGroup::*)(c10d::BarrierOptions const&), pybind11::name const&, pybind11::is_method const&, pybind11::sibling const&, pybind11::arg_v const&, pybind11::call_guard<pybind11::gil_scoped_release> const&)::{lambda(c10d::ProcessGroup*, c10d::BarrierOptions const&)#1}&&, c10::intrusive_ptr<c10d::Work, c10::detail::intrusive_target_default_null_type<c10d::Work> > (*)(c10d::ProcessGroup*, c10d::BarrierOptions const&), pybind11::name const&, pybind11::is_method const&, pybind11::sibling const&, pybind11::arg_v const&, pybind11::call_guard<pybind11::gil_scoped_release> const&)::{lambda(pybind11::detail::function_call&)#3}::_FUN(pybind11::detail::function_call) [0xde7da7]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/libtorch_python.so
=========     Host Frame:pybind11::cpp_function::dispatcher(_object*, _object*, _object*) [0x4c7cc3]
=========                in /home/lcw/conda_env/lib/python3.10/site-packages/torch/lib/libtorch_python.so
=========     Host Frame:cfunction_call in /usr/local/src/conda/python-3.10.14/Objects/methodobject.c:543 [0x1445a5]
=========                in /home/lcw/conda_env/bin/python
=========     Host Frame:_PyObject_MakeTpCall in /usr/local/src/conda/python-3.10.14/Objects/call.c:215 [0x13da6a]
=========                in /home/lcw/conda_env/bin/python
=========     Host Frame:method_vectorcall in /usr/local/src/conda/python-3.10.14/Objects/classobject.c:53 [0x150865]
=========                in /home/lcw/conda_env/bin/python
=========     Host Frame:_PyEval_EvalFrameDefault in /usr/local/src/conda/python-3.10.14/Python/ceval.c:4231 [0x1358f9]
=========                in /home/lcw/conda_env/bin/python
=========     Host Frame:_PyFunction_Vectorcall in /usr/local/src/conda/python-3.10.14/Objects/call.c:342 [0x144a2b]
=========                in /home/lcw/conda_env/bin/python
=========     Host Frame:PyObject_Call in /usr/local/src/conda/python-3.10.14/Objects/call.c:317 [0x150f1b]
=========                in /home/lcw/conda_env/bin/python
=========     Host Frame:_PyEval_EvalFrameDefault in /usr/local/src/conda/python-3.10.14/Python/ceval.c:4277 [0x1372b2]
=========                in /home/lcw/conda_env/bin/python
=========     Host Frame:_PyFunction_Vectorcall in /usr/local/src/conda/python-3.10.14/Objects/call.c:342 [0x144a2b]
=========                in /home/lcw/conda_env/bin/python
=========     Host Frame:_PyEval_EvalFrameDefault in /usr/local/src/conda/python-3.10.14/Python/ceval.c:4231 [0x1358f9]
=========                in /home/lcw/conda_env/bin/python
=========     Host Frame:_PyEval_Vector in /usr/local/src/conda/python-3.10.14/Python/ceval.c:5067 [0x1d7c5f]
=========                in /home/lcw/conda_env/bin/python
=========     Host Frame:PyEval_EvalCode in /usr/local/src/conda/python-3.10.14/Python/ceval.c:1134 [0x1d7ba6]
=========                in /home/lcw/conda_env/bin/python
=========     Host Frame:run_eval_code_obj in /usr/local/src/conda/python-3.10.14/Python/pythonrun.c:1291 [0x208129]
=========                in /home/lcw/conda_env/bin/python
=========     Host Frame:run_mod in /usr/local/src/conda/python-3.10.14/Python/pythonrun.c:1312 [0x203522]
=========                in /home/lcw/conda_env/bin/python
=========     Host Frame:PyRun_StringFlags in /usr/local/src/conda/python-3.10.14/Python/pythonrun.c:1183 [0x1fb91c]
=========                in /home/lcw/conda_env/bin/python
=========     Host Frame:PyRun_SimpleStringFlags in /usr/local/src/conda/python-3.10.14/Python/pythonrun.c:503 [0x1fb75b]
=========                in /home/lcw/conda_env/bin/python
=========     Host Frame:Py_RunMain in /usr/local/src/conda/python-3.10.14/Modules/main.c:670 [0x1fa66a]
=========                in /home/lcw/conda_env/bin/python
=========     Host Frame:Py_BytesMain in /usr/local/src/conda/python-3.10.14/Modules/main.c:1090 [0x1cb1f6]
=========                in /home/lcw/conda_env/bin/python
=========     Host Frame:__libc_start_main [0x24082]
=========                in /lib/x86_64-linux-gnu/libc.so.6
=========     Host Frame: [0x1cb0f0]
=========                in /home/lcw/conda_env/bin/python
=========
```
plus a lot of other downstream errors.

### Versions

```
PyTorch version: 2.5.0.dev20240718
Is debug build: False
CUDA used to build PyTorch: 12.1
ROCM used to build PyTorch: N/A

OS: Ubuntu 20.04.6 LTS (x86_64)
GCC version: (Ubuntu 9.4.0-1ubuntu1~20.04.2) 9.4.0
Clang version: Could not collect
CMake version: version 3.29.5
Libc version: glibc-2.31

Python version: 3.10.14 | packaged by conda-forge | (main, Mar 20 2024, 12:45:18) [GCC 12.3.0] (64-bit runtime)
Python platform: Linux-5.15.0-1063-aws-x86_64-with-glibc2.31
Is CUDA available: True
CUDA runtime version: 12.1.105
CUDA_MODULE_LOADING set to: LAZY
GPU models and configuration:
GPU 0: NVIDIA H100 80GB HBM3
GPU 1: NVIDIA H100 80GB HBM3
GPU 2: NVIDIA H100 80GB HBM3
GPU 3: NVIDIA H100 80GB HBM3
GPU 4: NVIDIA H100 80GB HBM3
GPU 5: NVIDIA H100 80GB HBM3
GPU 6: NVIDIA H100 80GB HBM3
GPU 7: NVIDIA H100 80GB HBM3

Nvidia driver version: 535.183.01
cuDNN version: Could not collect
HIP runtime version: N/A
MIOpen runtime version: N/A
Is XNNPACK available: True

CPU:
Architecture:                       x86_64
CPU op-mode(s):                     32-bit, 64-bit
Byte Order:                         Little Endian
Address sizes:                      48 bits physical, 48 bits virtual
CPU(s):                             192
On-line CPU(s) list:                0-191
Thread(s) per core:                 2
Core(s) per socket:                 48
Socket(s):                          2
NUMA node(s):                       2
Vendor ID:                          AuthenticAMD
CPU family:                         25
Model:                              1
Model name:                         AMD EPYC 7R13 Processor
Stepping:                           1
CPU MHz:                            2649.998
BogoMIPS:                           5299.99
Hypervisor vendor:                  KVM
Virtualization type:                full
L1d cache:                          3 MiB
L1i cache:                          3 MiB
L2 cache:                           48 MiB
L3 cache:                           384 MiB
NUMA node0 CPU(s):                  0-47,96-143
NUMA node1 CPU(s):                  48-95,144-191
Vulnerability Gather data sampling: Not affected
Vulnerability Itlb multihit:        Not affected
Vulnerability L1tf:                 Not affected
Vulnerability Mds:                  Not affected
Vulnerability Meltdown:             Not affected
Vulnerability Mmio stale data:      Not affected
Vulnerability Retbleed:             Not affected
Vulnerability Spec rstack overflow: Mitigation; safe RET, no microcode
Vulnerability Spec store bypass:    Mitigation; Speculative Store Bypass disabled via prctl and seccomp
Vulnerability Spectre v1:           Mitigation; usercopy/swapgs barriers and __user pointer sanitization
Vulnerability Spectre v2:           Mitigation; Retpolines; IBPB conditional; IBRS_FW; STIBP always-on; RSB filling; PBRSB-eIBRS Not affected; BHI Not affected
Vulnerability Srbds:                Not affected
Vulnerability Tsx async abort:      Not affected
Flags:                              fpu vme de pse tsc msr pae mce cx8 apic sep mtrr pge mca cmov pat pse36 clflush mmx fxsr sse sse2 ht syscall nx mmxext fxsr_opt pdpe1gb rdtscp lm constant_tsc rep_good nopl nonstop_tsc cpuid extd_apicid aperfmperf tsc_known_freq pni pclmulqdq monitor ssse3 fma cx16 pcid sse4_1 sse4_2 x2apic movbe popcnt aes xsave avx f16c rdrand hypervisor lahf_lm cmp_legacy cr8_legacy abm sse4a misalignsse 3dnowprefetch topoext perfctr_core invpcid_single ssbd ibrs ibpb stibp vmmcall fsgsbase bmi1 avx2 smep bmi2 invpcid rdseed adx smap clflushopt clwb sha_ni xsaveopt xsavec xgetbv1 clzero xsaveerptr rdpru wbnoinvd arat npt nrip_save vaes vpclmulqdq rdpid

Versions of relevant libraries:
[pip3] flake8==7.0.0
[pip3] lovely-numpy==0.2.11
[pip3] mypy-extensions==1.0.0
[pip3] numpy==1.24.4
[pip3] torch==2.5.0.dev20240718
[pip3] torchmetrics==0.10.3
[pip3] torchvision==0.20.0.dev20240718
[pip3] triton==3.0.0
[conda] Could not collect
```

cc @XilunWu @H-Huang @awgu @kwen2501 @wanchaol @fegin @fduwjj @wz337 @wconstab @d4l3k @c-p-i-o