Initialize proxy global-exit request state

[zianglih]

Summary

@HumansAnd

Initialize the proxy's global-exit request state and exit code immediately after allocating their pinned host storage. This prevents allocation contents from being interpreted as an application request to exit.

• Bug: calloc() clears proxy_state_t, including its pointer fields, but does not initialize the two subsequent cudaMallocHost() allocations. Both the minimal and full proxy progress paths read global_exit_request_state; a value equal to PROXY_GLOBAL_EXIT_REQUESTED (2) causes nvshmem_global_exit(*global_exit_code).
• Fix: Store PROXY_GLOBAL_EXIT_NOT_REQUESTED and 0 before nvshmemi_proxy_prep_minimal_state(), device-state publication, and either progress-thread creation. This follows the existing initialization of the proxy channel counters. No API, transport selection, device request protocol, or runtime option changes.
• Related application integration: sgl-project/sglang#39210. This fixes the allocation initialization in NVSHMEM itself; it does not disable the proxy.

Evidence and validation

The missing initialization is still present on devel at b0d9d3dc08fc3ee0840fdb6f3a2c11932d85a2e2. This PR changes one file by two lines. git diff --check passes. No behavioral regression test is added.

The following observations were collected separately on NVSHMEM 3.4.5, before porting the same two stores to this branch:

• Environment: One eight-GPU NVIDIA B300 node, four application ranks, driver 590.48.01; lmsysorg/sglang:nightly-dev-cu13-20260918-20518d85 (sha256:b518f4f8cd15664cf0f733e9bf4fd2105c9c994882a369b247364394db99f596), PyTorch 2.13.0+cu130, CUDA 13.0, NVSHMEM4Py 0.3.1, SGLang 6d8a58f177a1488e10e8567a3fa4b4fccd0a26e2, FlashInfer ad0a5e5e78e57070ec7c582efe733cb55cd8839f.
• Application trigger: Real-weight GLM-5.2 NVFP4 startup with FlashInfer MegaMoE, TP/DP/EP 4, max running requests 8, static memory fraction 0.80, and EAGLE draft settings 3/1/4. This is allocation-history sensitive; a standalone deterministic reproducer is not yet available.
• Native observation: Hardware breakpoints sampled the request and code immediately after the successful allocation returns, before that proxy's progress thread was created. All four ranks returned request value 2 and code 0. The same proxy pointers later reached nvshmem_global_exit@plt from the full progress path with exit argument 0. No inferior code, data, or computational register values were changed; debugger stops did perturb timing. The allocator's prior writer was not observed.

rank                 0    1    2    3
request after alloc  2    2    2    2
code after alloc     0    0    0    0
proxy global_exit    0    0    0    0

A separate matched source-build comparison used the public v3.4.5-0 source (131da55f643ac87c810ba0bc51d359258bf433a1) with the two stores as the only original-source difference. Both builds used CUDA compiler 13.0.88, GCC 13.3, sm_103, and identical build settings. UCX was disabled because its headers were unavailable; tests, examples, Python, packaging, and bitcode builds were disabled. These are not claimed to be binary-equivalent to the distributed wheel.

Startup observation	Pristine build	Two-store build
Natural host-init wrapper returned	4 ranks	4 ranks
First symmetric allocation entered / returned	4 / 0	4 / 4
nvshmem API called before nvshmem_init messages	4	0
HTTP readiness	Not reached	Reached

Limits: The current devel patch has not been built or run on a GPU. The matched 3.4.5 builds completed with existing CMake/PMIx warnings. The startup comparison sent no performance requests and establishes no numerical or performance result. Legitimate device global_exit propagation, other transports, and multi-node behavior were not tested. The pristine run's initially incomplete cleanup check was retained; a separate later idle check passed. The patched server was stopped deliberately after readiness with no recorded remaining owners or cleanup errors.

[benhg]

Thank you for your contribution. We are somewhat busy finishing up the 3.8 release but will take a look. From first glance I don't see any problems but will go through init in more detail before committing ot that stance.

[benhg]

This has been merged internally and should come back out to GitHub in the next push from our internal mirror.