出现uncorrectable以后GPU Not available

硬件为C500 mxlink，在docker环境下，两个GPU做集合通信以后出现：
[二 9月 8 22:43:34 2026] METAX.B5700.D0.PCI.ERROR metalk4 uncorrectable error, status 0x00004000, count 5
[二 9月 8 22:43:34 2026] METAX.B5700.D0.PCI.ERROR metalk4 uncorrectable AER [14] Completion Timeout, count 5
然后GPU Not available了，这是正常的吗？

尊敬的开发者您好，请参考developer.metax-tech.com/forum/t/fa-tie-qian-bi-kan-jing-xiang-shi-yong-wen-ti-ti-wen-mo-ban/267/ 提供详细信息

你好，我上传详细故障报告，代码文件无法上传：

/* custom_all_reduce.cuh — vLLM's custom all-reduce, ported to MetaX C500 (MACA).
 * ===========================================================================
 * !! DANGEROUS ON C500 — DISABLED BY DEFAULT, DO NOT RUN WITHOUT READING THIS !!
 *
 * Running this on C500 produced UNCORRECTABLE HARDWARE errors in the host driver:
 *     METAX.B5700.D0.PCI.ERROR metalk4 uncorrectable AER [14] Completion Timeout
 * on every GPU it ran on and no other, permanently disabling 4 of 8 GPUs
 * (2,3 then 4,5) — not recoverable by restarting the container.
 *
 * metalk4 is the MetaLink inter-GPU link. This kernel drives that link from the
 * SMs: packed_reduce issues fine-grained 128-bit REMOTE LOADS against every peer's
 * IPC-mapped Finegrained (uncached) buffer, continuously, for the whole run. The
 * vendor collective does not move data that way. The probes passed because they
 * ran for seconds; the link faults need minutes of sustained traffic.
 *
 * The correctness and latency results below remain valid as MEASUREMENTS. What is
 * now known is that the traffic pattern itself is not safe on this hardware, which
 * is a design-level problem, not a bug to patch.
 * ===========================================================================
 * WHAT THIS IS. A NCCL-bypassing one-shot all-reduce: every rank's buffer is
 * IPC-mapped into every other rank, a flag barrier rendezvouses, then each rank
 * BULK-READS all peers' buffers and reduces locally. For the small, latency-bound
 * tensors a TP decode step produces, this beats a ring collective — which is the
 * regime C500 currently has nothing for (vllm_metax ships zero custom-AR ops and
 * its MacaCommunicator never references them; see project_c500_custom_ar).
 *
 * WHY A PORT IS ONLY THREE CHANGES. Upstream is ALREADY two-armed — the barrier
 * primitives are `#if !defined(USE_ROCM)` inline PTX `#else` `__scoped_atomic_*`
 * (upstream custom_all_reduce.cuh:157/243/286), and even the ALLOCATOR is
 * vendor-split (`hipExtMallocWithFlags(hipDeviceMallocUncached)` for ROCm vs plain
 * `cudaMalloc`). Custom AR was never NVIDIA-locked. So MACA is a THIRD arm beside
 * the ROCm one, not a semantics rewrite:
 *
 *   1. the four flag primitives    -> __atomic_{load,store}_n(ACQUIRE/RELEASE)
 *   2. the shared-buffer allocator -> mcExtMallocWithFlags(..., 0x1) Finegrained
 *                                     (lives in custom_all_reduce.cu, not here)
 *   3. the Python-side gates       -> _can_p2p / is_fully_connected / max_size
 *
 * ALL OF THIS IS MEASURED, NOT ASSUMED. src/hetero_tp/probe_maca_custom_ar.cu ran
 * on C500 GPUs 2,3 on 2026-09-08 and gated exactly the parts uccl-c500 does NOT
 * already prove in production:
 *   G1 the release-store/acquire-load flag barrier (NOT atomicAdd_system)  PASS 32/32
 *   G2 the same on DEFAULT (cached) memory                                 FAIL, as predicted
 *   G3 sustained bulk 128-bit reads FROM peer memory, 64 MiB, checksummed  PASS
 *   G4 G1+G3 under heavy background load, 20 reps                          PASS, 0 livelocks
 * G2 is why change (2) is load-bearing: with a plain cudaMalloc the peer's release
 * store never becomes visible and the barrier hangs with a signature that reads
 * like a hardware coherence limit ("arrived but never released"). Same rule, third
 * time on this project — see the Finegrained note in
 * docs/notes/vllm/moe/uccl_ep_vendor_port_debugging.md.
 *
 * NOT A HAZARD HERE: warp-64. The barrier is `threadIdx.x < ngpus` + `__syncthreads()`
 * with no __shfl and no lane mask, unlike UCCL's windowed publish that warp-64 killed.
 *
 * v1 SCOPE (deliberate, so every shipped path is a path the probe covered):
 *   - one-shot only. Upstream's 2-stage (reduce-scatter + allgather) is the
 *     larger-payload path; it is a follow-up, and is marked TODO below.
 *   - eager only: no CUDA-graph buffer registration yet (C500 hetero-TP runs
 *     --enforce-eager today). register_graph_buffers is the other follow-up.
 *   - ngpus 2/4/8.
 * ---------------------------------------------------------------------------
 */
#pragma once

#include <cuda_runtime.h>
#include <cuda_fp16.h>
#include <cuda_bf16.h>
#include <cstdint>
#include <cstdio>
#include <limits>
#include <map>
#include <unordered_map>
#include <vector>

namespace hetero_car {

constexpr int kMaxBlocks = 36;

// Counter/flag state shared across ranks. `start`/`end` are per-block, per-rank
// slots: block b on rank r writes peer[b][r] and spins on self[b][peer]. Per-block
// (not grid-wide) is sufficient because every rank walks the payload with the SAME
// grid stride, so thread tid only ever reads what the peer's thread tid wrote —
// upstream states the same invariant ("visibility across devices is only guaranteed
// between threads that have the same tid").
struct Signal {
  alignas(128) uint32_t start[kMaxBlocks][8];
  alignas(128) uint32_t end[kMaxBlocks][8];
  alignas(128) uint32_t _flag[kMaxBlocks];  // incremented per barrier, per block
};

struct __align__(16) RankData {
  const void* ptrs[8];
};

struct __align__(16) RankSignals {
  Signal* signals[8];
};

template <typename T, int sz>
struct __align__(alignof(T) * sz) array_t {
  T data[sz];
  using type = T;
  static constexpr int size = sz;
};

// packed type maximises memory efficiency; goal is 128-bit loads/stores — the
// exact access shape probe G3 validated against peer memory on C500.
template <typename T>
struct packed_t {
  using P = array_t<T, 16 / sizeof(T)>;
  using A = array_t<float, 16 / sizeof(T)>;
};

#define DINLINE __device__ __forceinline__

/* ── THE THREE ARMS ─────────────────────────────────────────────────────────
 * MACA note: `mxcc` on its own defines neither __XCORE_WN__ nor __NVCC__ and would
 * silently take the NVIDIA PTX arm below. The compiler that selects this arm is
 * $CUDA_PATH/bin/nvcc -> /opt/maca/tools/cu-bridge/bin/cucc. build.sh enforces it,
 * and the runtime prints which arm compiled so a wrong build cannot pass silently.
 */
DINLINE void st_flag_release(uint32_t* flag_addr, uint32_t flag) {
#if defined(__XCORE_WN__)
  __atomic_store_n(flag_addr, flag, __ATOMIC_RELEASE);
#elif defined(USE_ROCM)
  __scoped_atomic_store_n(flag_addr, flag, __ATOMIC_RELEASE,
                          __MEMORY_SCOPE_SYSTEM);
#else
  asm volatile("st.release.sys.global.u32 [%1], %0;" ::"r"(flag), "l"(flag_addr)
               : "memory");
#endif
}

DINLINE uint32_t ld_flag_acquire(uint32_t* flag_addr) {
#if defined(__XCORE_WN__)
  return __atomic_load_n(flag_addr, __ATOMIC_ACQUIRE);
#elif defined(USE_ROCM)
  return __scoped_atomic_load_n(flag_addr, __ATOMIC_ACQUIRE,
                                __MEMORY_SCOPE_SYSTEM);
#else
  uint32_t flag;
  asm volatile("ld.acquire.sys.global.u32 %0, [%1];" : "=r"(flag) : "l"(flag_addr));
  return flag;
#endif
}

DINLINE void st_flag_volatile(uint32_t* flag_addr, uint32_t flag) {
#if defined(__XCORE_WN__)
  __atomic_store_n(flag_addr, flag, __ATOMIC_RELAXED);
#elif defined(USE_ROCM)
  __scoped_atomic_store_n(flag_addr, flag, __ATOMIC_RELAXED,
                          __MEMORY_SCOPE_SYSTEM);
#else
  asm volatile("st.volatile.global.u32 [%1], %0;" ::"r"(flag), "l"(flag_addr));
#endif
}

DINLINE uint32_t ld_flag_volatile(uint32_t* flag_addr) {
#if defined(__XCORE_WN__)
  return __atomic_load_n(flag_addr, __ATOMIC_RELAXED);
#elif defined(USE_ROCM)
  return __scoped_atomic_load_n(flag_addr, __ATOMIC_RELAXED,
                                __MEMORY_SCOPE_SYSTEM);
#else
  uint32_t flag;
  asm volatile("ld.volatile.global.u32 %0, [%1];" : "=r"(flag) : "l"(flag_addr));
  return flag;
#endif
}

/* ── FAILURE REPORTING ──────────────────────────────────────────────────────
 * Why this exists: the first serving run wedged BOTH GPUs of the pair into a
 * driver-level "Not Available" state. The spins below used to be UNBOUNDED
 * (upstream's shape, faithfully copied) — and an unbounded spin in a kernel that
 * also WRITES INTO A PEER'S memory is the worst combination available:
 *   - a kernel that never returns cannot be killed from the host; SIGKILL removes
 *     the process but leaves the kernel resident, so context teardown fails and
 *     the device is left unusable;
 *   - and because the barrier stores target the PEER's buffer, a survivor still
 *     spinning keeps writing into a mapping that is being torn down, which is how
 *     one dead rank takes its neighbour's GPU down with it.
 * The probe (probe_maca_custom_ar.cu) always had a bound and a trap; dropping it
 * in this port was the mistake. Bounded now, with the FIRST failure recorded
 * (read the first timeout, not the last) so a desync is diagnosable instead of
 * being a silent hang.
 *
 * A timeout does NOT __trap(): trap tears down the context and can itself leave
 * the device dirty. It records, returns, and lets the host raise — the process
 * then exits cleanly and the GPU stays recoverable.
 */
struct CarStatus {
  int code;       // 0 = clean; 1 = start barrier; 2 = end barrier; 3 = final barrier
  int block;      // blockIdx.x that timed out
  int peer;       // which peer's slot never advanced
  unsigned expected;
  unsigned observed;
  int _pad;
};

// Device-memory status (NOT the IPC-shared region: a peer dying must not affect
// our ability to report). First writer wins.
DINLINE void car_report(CarStatus* st, int code, int peer, unsigned expected,
                        unsigned observed) {
  if (st == nullptr) return;
  if (atomicCAS(&st->code, 0, code) == 0) {
    st->block = blockIdx.x;
    st->peer = peer;
    st->expected = expected;
    st->observed = observed;
    __threadfence();
  }
}

/* Barrier entering the reduction: publish my counter into every peer, then wait
 * for every peer's counter in my own slot. Release/acquire so the payload stores
 * that precede it are visible to the peer that reads them afterwards.
 *
 * Returns false on timeout. NOTE the shape: a timing-out thread must NOT return
 * early — its block-mates are still heading for __syncthreads(), and a divergent
 * __syncthreads() is itself a hang. The whole block leaves together, via a shared
 * flag, exactly as the probe does.
 */
template <int ngpus>
DINLINE bool barrier_at_start(const RankSignals& sg, Signal* self_sg, int rank,
                              CarStatus* st, long long max_spins) {
  __shared__ int blk_bad;
  if (threadIdx.x == 0) blk_bad = 0;
  __syncthreads();
  uint32_t flag = self_sg->_flag[blockIdx.x] + 1;
  if (threadIdx.x < ngpus) {
    // The payload writes must land before the peer observes my flag.
    __threadfence_system();
    st_flag_release(&sg.signals[threadIdx.x]->start[blockIdx.x][rank], flag);
    long long spins = 0;
    uint32_t seen;
    while ((seen = ld_flag_acquire(&self_sg->start[blockIdx.x][threadIdx.x])) !=
           flag) {
      if (++spins > max_spins) {
        car_report(st, 1, threadIdx.x, flag, seen);
        blk_bad = 1;
        break;
      }
    }
  }
  __syncthreads();
  if (blk_bad) return false;          // every thread of the block returns here
  if (threadIdx.x == 0) self_sg->_flag[blockIdx.x] = flag;
  return true;
}

// Barrier leaving the reduction. final_sync=true skips the visibility guarantee
// (nothing after it reads peer memory), matching upstream. Same bounded shape.
template <int ngpus, bool final_sync = false>
DINLINE bool barrier_at_end(const RankSignals& sg, Signal* self_sg, int rank,
                            CarStatus* st, long long max_spins) {
  __shared__ int blk_bad_e;
  if (threadIdx.x == 0) blk_bad_e = 0;
  __syncthreads();
  uint32_t flag = self_sg->_flag[blockIdx.x] + 1;
  if (threadIdx.x < ngpus) {
    long long spins = 0;
    uint32_t seen;
    if constexpr (!final_sync) {
      __threadfence_system();
      st_flag_release(&sg.signals[threadIdx.x]->end[blockIdx.x][rank], flag);
      while ((seen = ld_flag_acquire(&self_sg->end[blockIdx.x][threadIdx.x])) !=
             flag) {
        if (++spins > max_spins) {
          car_report(st, 2, threadIdx.x, flag, seen);
          blk_bad_e = 1;
          break;
        }
      }
    } else {
      st_flag_volatile(&sg.signals[threadIdx.x]->end[blockIdx.x][rank], flag);
      while ((seen = ld_flag_volatile(&self_sg->end[blockIdx.x][threadIdx.x])) !=
             flag) {
        if (++spins > max_spins) {
          car_report(st, 3, threadIdx.x, flag, seen);
          blk_bad_e = 1;
          break;
        }
      }
    }
  }
  __syncthreads();                    // unconditional: see the note above
  if (blk_bad_e) return false;
  if (threadIdx.x == 0) self_sg->_flag[blockIdx.x] = flag;
  return true;
}

// ── scalar/packed arithmetic ──
DINLINE float upcast_s(half val) { return __half2float(val); }
template <typename T>
DINLINE T downcast_s(float val);
template <>
DINLINE half downcast_s(float val) {
  return __float2half(val);
}
DINLINE half& assign_add(half& a, half b) {
  a = __hadd(a, b);
  return a;
}
DINLINE float& assign_add(float& a, float b) { return a += b; }
DINLINE float upcast_s(float val) { return val; }
template <>
DINLINE float downcast_s(float val) {
  return val;
}

// bf16. __CUDA_ARCH__ IS defined in cu-bridge's device pass (the "it is undefined
// under cucc" theory was tested with an #error and refuted), and C500 reports
// capability 9.0, so this arm compiles on MACA.
#if (__CUDA_ARCH__ >= 800 || !defined(__CUDA_ARCH__))
DINLINE float upcast_s(__nv_bfloat16 val) { return __bfloat162float(val); }
template <>
DINLINE __nv_bfloat16 downcast_s(float val) {
  return __float2bfloat16(val);
}
DINLINE __nv_bfloat16& assign_add(__nv_bfloat16& a, __nv_bfloat16 b) {
  a = __hadd(a, b);
  return a;
}
#endif

template <typename T, int N>
DINLINE array_t<T, N>& packed_assign_add(array_t<T, N>& a, array_t<T, N> b) {
#pragma unroll
  for (int i = 0; i < N; i++) assign_add(a.data[i], b.data[i]);
  return a;
}

template <typename T, int N>
DINLINE array_t<float, N> upcast(array_t<T, N> val) {
  array_t<float, N> out;
#pragma unroll
  for (int i = 0; i < N; i++) out.data[i] = upcast_s(val.data[i]);
  return out;
}

template <typename O, typename A>
DINLINE O downcast(A val) {
  O out;
#pragma unroll
  for (int i = 0; i < O::size; i++)
    out.data[i] = downcast_s<typename O::type>(val.data[i]);
  return out;
}

// Accumulation order is fixed (ptrs are NOT rotated per rank) so every rank
// produces a bitwise-identical result. Do not "optimise" this into a rotation.
template <typename P, int ngpus, typename A>
DINLINE P packed_reduce(const P* ptrs[], int idx) {
  A tmp = upcast(ptrs[0][idx]);
#pragma unroll
  for (int i = 1; i < ngpus; i++) packed_assign_add(tmp, upcast(ptrs[i][idx]));
  return downcast<P>(tmp);
}

// One-shot: barrier, then every rank reads every peer's whole buffer and reduces.
// The peer-read inner loop is exactly what probe G3 checksummed over 64 MiB.
// TODO(v2): cross_device_reduce_2stage (reduce-scatter + allgather) for payloads
// where one-shot's ngpus-fold read amplification stops paying.
template <typename T, int ngpus>
__global__ void __launch_bounds__(512, 1)
    cross_device_reduce_1stage(RankData* _dp, RankSignals sg, Signal* self_sg,
                               T* __restrict__ result, int rank, int size,
                               CarStatus* st, long long max_spins) {
  using P = typename packed_t<T>::P;
  using A = typename packed_t<T>::A;
  auto dp = *_dp;
  // Bail on a timed-out rendezvous rather than reducing against a peer that never
  // arrived: the host sees st->code != 0 and raises, so no wrong result is served.
  if (!barrier_at_start<ngpus>(sg, self_sg, rank, st, max_spins)) return;
  for (int idx = blockIdx.x * blockDim.x + threadIdx.x; idx < size;
       idx += gridDim.x * blockDim.x) {
    ((P*)result)[idx] = packed_reduce<P, ngpus, A>((const P**)&dp.ptrs[0], idx);
  }
  barrier_at_end<ngpus, true>(sg, self_sg, rank, st, max_spins);
}

}  // namespace hetero_car

尊敬的开发者您好，请宿主机执行bash /opt/maca/samples/mccl_tests/perf/mccl.sh 8