get_nccl_comm_handle reuses a null _comm_ptr() when the group's NCCL comm is lazily uninitialized — ncclTeamWorld null-deref at ElasticBuffer construction (vLLM serve, 16/16 ranks)

## Summary

`get_nccl_comm_handle` (`deep_ep/utils/comm.py`) reuses PyTorch's NCCL communicator when `EP_REUSE_NCCL_COMM=1` (the default since #637):

```python
backend = group._get_backend(torch.device('cuda'))
if not force_new_comm and hasattr(backend, '_comm_ptr') and int(os.getenv('EP_REUSE_NCCL_COMM', '1')):
_storage[group] = NCCLCommHandle(backend._comm_ptr(), False)
return _storage[group]
```

`backend._comm_ptr()` is a **passive accessor**: PyTorch creates NCCL communicators **lazily**, and `ProcessGroupNCCL::getCommPtr()` returns **0** when the group+device has no communicator yet ("ncclComm is a nullptr if the communicator does not exist", `ProcessGroupNCCL.cpp`, torch 2.11). Comms materialize at the first collective on that group+device — or eagerly only when `init_process_group(device_id=...)` was used.

So any consumer that constructs an `ElasticBuffer` **before any collective has run on the EP group** gets `NCCLCommHandle(0)`, and the constructor immediately calls `calculate_elastic_buffer_size(handle.get(), ...)` → `deep_ep::nccl::get_physical_domain_size` (`csrc/kernels/backend/nccl.cu`) → `ncclTeamWorld(comm)` which dereferences `comm->nRanks` (NCCL `src/nccl_device/core.cc`) → segfault.

## Live reproduction (not hypothetical)

vLLM expert-parallel serving (Qwen3-30B-A3B-FP8, DP16/EP16, 2× p5en over EFA) hits this deterministically on **all 16 ranks** at serve init:

```
ncclTeamWorld (nccl_device/core.cc:16)
<- deep_ep::nccl::get_physical_domain_size
<- ElasticBuffer::calculate_buffer_size (csrc/elastic/buffer.hpp)
<- pybind
```

vLLM's default path is exactly the cold-group shape: `init_process_group` without `device_id` (all groups lazy), TP=1 (no TP collectives), DP-sync NCCL disabled under async scheduling, gloo-only barriers — so the EP `device_group` has had zero collectives when `ElasticBuffer` is constructed from `process_weights_after_loading`. `_comm_ptr()` returns 0 and every rank crashes identically.

Setting `EP_REUSE_NCCL_COMM=0` (create-own-comm path) fixes the serve on the same substrate — confirmed E2E with sustained load (c=1→64 sweep, 0 errors).

## Why DeepEP's own tests don't catch it

`tests/utils` `init_dist` passes `device_id=torch.device(f'cuda:{local_rank}')` to `init_process_group`, which makes torch **eager-create** communicators — `_comm_ptr()` is always valid in the test suite. The reuse branch's null case is only reachable from lazy-init consumers.

## Proposed fix

Only take the reuse branch when the pointer is non-null; otherwise fall through to the existing create-own-comm path (whose `all_gather_object` is itself a collective, so it warms torch's comm as a side effect and all ranks take the same branch on a cold boot):

```python
comm_ptr = backend._comm_ptr()
if comm_ptr != 0:
_storage[group] = NCCLCommHandle(comm_ptr, False)
return _storage[group]
# fall through: create our own comm
```

PR incoming with this guard.