Yunxiao

最近在用C500+RDMA+GPUDirect，驱动版本:
+---------------------------------------------------------------------------------+
| MX-SMI 2.3.4 Kernel Mode Driver Version: 3.9.25 |
| MACA Version: 3.8.0.23 BIOS Version: 1.35.4.0 |
遇到以下问题，咨询一下沐曦官方的老师：
1. L2 的位置：网卡经 PCIe P2P 写入 C500 显存（比如 ibv_reg_dmabuf_mr 注册的 buffer）时，这笔写入是否经过 GPU 的 L2？L2 是在内存一侧（所有访问方共享），还是在 SM 一侧？
2. 用默认分配时会不会读到旧数据：如果用 mcExtMallocWithFlags(…, 0x0)（默认、可缓存）分配接收缓冲区，SM 之前读过同一地址，之后网卡又写入新数据，SM 再读时有没有可能拿到 L2 里的旧行？如果有，推荐的做法是什么，比如某个失效接口或者 acquire 语义的 fence？
3. Finegrained 的缓存语义：mcDeviceMallocFinegrained（0x1）分配的内存，SM 读取时会不会被缓存在 L2？我们测到，同一个 GEMM 如果 A 放在 0x1 内存里，重复读取的成本高约 57%。
4. 写入长期留在 L2：C500 上 SM 写入后留在 L2 的脏行，会不会一直驻留，直到被写入挤出？只读访问能不能把它们挤出？我们观察到的现象是：写入约 1–5 MB 后，之后几毫秒内访存密集的 kernel 都会慢 5–20%；各种存储缓存提示（.cs / .wt / .cg / evict_first）以及后续 32 MB 的只读访问都消除不了。
mcDeviceCacheEarlyWriteBack 的作用：这个设备配置具体改变什么？进程退出后会不会保留？

需要vllm-metax:0.26.0-maca.ai3.8.2.208-torch2.10-py312-ubuntu22.04-amd64镜像包离线下载。
谢谢！

你好，跟文档描述不同：
developer.metax-tech.com/api/client/document/preview/1475/split_files/memory%E9%AA%8C%E6%94%B6%E6%B5%8B%E8%AF%95%E5%B7%A5%E5%85%B7.html#h3hdfawbw5w51
文档中描述--kernel-copy是通过读写方式计算带宽，默认为只读。我的理解是不指定该参数的时候是只读，指定以后为读写混合。
所以读写混合带宽比读带宽低一半是正常的吗？还是说我的服务器有问题需要维修呢？

你好，已通过mxvs测试，结果如下：
dpu4@dpu4:~$ sudo /opt/maca/bin/mxvs memory benchmark --devices 0 --data-size 1GB --detail
MetaX GPU Information
──────────────────────────────────────────────────────────────────────────────
MACA 3.8.2.6
TIMESTAMP Sun Sep 20 17:51:00 2026
──────────────────────────────────────────────────────────────────────────────
GPU ID 0
MODEL MetaX C500
BDF 0000:4f:00.0
KMD 3.9.25
VBIOS 1.35.4.0
NUMA NODE 0
CURRENT PCIE speed: 16.0 GT/s width: x16
MAXIMUM PCIE speed: 16.0 GT/s width: x16
METAXLINK PORTS 4,5,6
──────────────────────────────────────────────────────────────────────────────
HBM BANDWIDTH BENCHMARK TEST

BOARD GPU DIE BDF SIZE(B) BANDWIDTH
──────────────────────────────────────────────────────────────────────────────
0 0 0 0000:4f:00.0 1073741824 1633.39 GB/s
dpu4@dpu4:~$ sudo /opt/maca/bin/mxvs memory benchmark --devices 0 --data-size 1GB --detail --kernel-co
py
MetaX GPU Information
──────────────────────────────────────────────────────────────────────────────
MACA 3.8.2.6
TIMESTAMP Sun Sep 20 17:51:21 2026
──────────────────────────────────────────────────────────────────────────────
GPU ID 0
MODEL MetaX C500
BDF 0000:4f:00.0
KMD 3.9.25
VBIOS 1.35.4.0
NUMA NODE 0
CURRENT PCIE speed: 16.0 GT/s width: x16
MAXIMUM PCIE speed: 16.0 GT/s width: x16
METAXLINK PORTS 4,5,6
──────────────────────────────────────────────────────────────────────────────
HBM BANDWIDTH BENCHMARK TEST

BOARD GPU DIE BDF SIZE(B) BANDWIDTH
──────────────────────────────────────────────────────────────────────────────
0 0 0 0000:4f:00.0 1073741824 727.69 GB/s
kernel-copy的带宽只有727GB/s，正常的memory benchmark是1633，请问这是正常的吗？kernel-copy是读+写操作吗？

测试脚本已上传，镜像为vllm-metax-0.23.0-maca.ai3.8.0.103-torch2.10-py312-ubuntu22.04-amd64，宿主机的驱动版本
| MX-SMI 2.3.4 Kernel Mode Driver Version: 3.9.25 |
| MACA Version: 3.8.2.6 BIOS Version: 1.35.4.0 |
经过测试，除了一张GPU读写方向带宽基本一致以外，其余GPU读带宽均大于写带宽，读方向带宽为1600GB/s左右：
➜ ✗ python3 benchmark/probe_rw_direction.py --out /tmp/rw.json 2>&1 | tail -22
buffer 1024 MiB fp32, per-GPU allocate/free, min-of-20 cuda-event timing

gpu name w4 read w4 write w8 read w8 write rd/wr copy free GiB
0 MetaX C500 1603 412 1602 415 3.89 575 62.10
1 MetaX C500 1605 413 1595 416 3.88 574 62.10
2 MetaX C500 1605 411 1578 414 3.90 573 62.10
3 MetaX C500 1611 413 1608 415 3.90 574 62.10
4 MetaX C500 1457 1365 1214 1154 1.07 1189 31.67
5 MetaX C500 1605 412 1606 415 3.87 576 62.10
6 MetaX C500 1600 412 1610 415 3.88 574 62.09
执行脚本方法：
python3 probe_rw_direction.py --out /tmp/rw.json 2>&1
请问这是正常现象吗？

import argparse, json, os, statistics as st

import torch, triton, triton.language as tl


@triton.jit
def k_read(x, out, N, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)
    v = tl.load(x + offs, mask=offs < N, other=0.0)
    tl.store(out + pid, tl.sum(v))


@triton.jit
def k_write(y, N, BLOCK: tl.constexpr):
    pid = tl.program_id(0)
    offs = pid * BLOCK + tl.arange(0, BLOCK)
    tl.store(y + offs, 1.0, mask=offs < N)


def t(fn, iters=20, warm=6):
    for _ in range(warm):
        fn()
    torch.cuda.synchronize()
    ts = []
    for _ in range(iters):
        a = torch.cuda.Event(True)
        b = torch.cuda.Event(True)
        a.record()
        fn()
        b.record()
        b.synchronize()
        ts.append(a.elapsed_time(b) / 1e3)
    return min(ts)


def probe_one(g, mib, warps):
    """Everything for one GPU. Allocates and frees inside, so a card already
    holding a server's weights is fine as long as the buffers fit."""
    torch.cuda.set_device(g)
    dev = f"cuda:{g}"
    N = mib * 1024 * 1024 // 4                        # fp32 elements
    x = torch.randn(N, dtype=torch.float32, device=dev)
    y = torch.empty_like(x)
    BLOCK = 4096
    G = triton.cdiv(N, BLOCK)
    red = torch.empty(G, dtype=torch.float32, device=dev)
    B = N * 4

    r = {"triton": {}, "torch": {}}
    for nw in warps:
        tr = t(lambda: k_read[(G,)](x, red, N, BLOCK=BLOCK, num_warps=nw))
        tw = t(lambda: k_write[(G,)](y, N, BLOCK=BLOCK, num_warps=nw))
        r["triton"][str(nw)] = {"read": B / 1e9 / tr, "write": B / 1e9 / tw}
    for nm, fn, by in (("sum_read", lambda: x.sum(), B),
                       ("fill_write", lambda: y.fill_(1.0), B),
                       ("copy_rdwr", lambda: y.copy_(x), 2 * B),
                       ("mul_rmw", lambda: x.mul_(1.0001), 2 * B)):
        r["torch"][nm] = by / 1e9 / t(fn)
    del x, y, red
    torch.cuda.empty_cache()
    return r


def main():
    p = argparse.ArgumentParser()
    p.add_argument("--gpus", default="",
                   help="comma list of torch device indices; default every visible GPU")
    p.add_argument("--mib", type=int, default=1024, help="buffer size, well past L2")
    p.add_argument("--warps", default="4,8", help="comma list of num_warps to try")
    p.add_argument("--out", default="", help="write the full result as JSON here")
    a = p.parse_args()

    warps = [int(w) for w in a.warps.split(",") if w != ""]
    gpus = ([int(x) for x in a.gpus.split(",") if x != ""]
            or list(range(torch.cuda.device_count())))

    cvd = os.environ.get("CUDA_VISIBLE_DEVICES") or os.environ.get("MACA_VISIBLE_DEVICES")
    if cvd:
        print(f"!! CUDA_VISIBLE_DEVICES={cvd} is set -- the 'gpu' column below is the\n"
              f"!! TORCH index, which is a position within that list, NOT the physical\n"
              f"!! index mx-smi/nvidia-smi prints. Unset it before producing a report\n"
              f"!! that names a card.")
    print(f"buffer {a.mib} MiB fp32, per-GPU allocate/free, min-of-20 cuda-event timing\n")

    out = {"mib": a.mib, "visible_devices": cvd, "gpus": {}}
    hdr = "".join(f"{f'w{w} read':>11}{f'w{w} write':>11}" for w in warps)
    print(f"  {'gpu':<4}{'name':<24}{hdr}{'rd/wr':>8}{'copy':>9}{'free GiB':>10}")
    for g in gpus:
        torch.cuda.set_device(g)
        free, _ = torch.cuda.mem_get_info(g)
        name = torch.cuda.get_device_name(g)
        if free < a.mib * 1024 * 1024 * 3:
            print(f"  {g:<4}{name[:23]:<24}   skipped: needs 3x the buffer free "
                  f"({free / 1024 ** 3:.1f} GiB)")
            continue
        r = probe_one(g, a.mib, warps)
        r["name"] = name
        r["free_gib"] = round(free / 1024 ** 3, 2)
        best = max(warps, key=lambda w: r["triton"][str(w)]["read"])
        r["ratio_read_over_write"] = (r["triton"][str(best)]["read"]
                                      / r["triton"][str(best)]["write"])
        out["gpus"][str(g)] = r
        cells = "".join(f"{r['triton'][str(w)]['read']:>11.0f}"
                        f"{r['triton'][str(w)]['write']:>11.0f}" for w in warps)
        print(f"  {g:<4}{name[:23]:<24}{cells}{r['ratio_read_over_write']:>8.2f}"
              f"{r['torch']['copy_rdwr']:>9.0f}{r['free_gib']:>10.2f}")

    got = out["gpus"]
    if len(got) > 1:
        for k, get in (("read", lambda r: r["triton"][str(warps[0])]["read"]),
                       ("write", lambda r: r["triton"][str(warps[0])]["write"]),
                       ("copy", lambda r: r["torch"]["copy_rdwr"])):
            v = {g: get(r) for g, r in got.items()}
            lo, hi = min(v, key=v.get), max(v, key=v.get)
            out[f"{k}_spread"] = round(v[hi] / v[lo], 3)
            print(f"\n  {k:<6} spread {v[hi] / v[lo]:.2f}x across GPUs "
                  f"(min {v[lo]:.0f} on gpu {lo}, max {v[hi]:.0f} on gpu {hi})")
        print("\n  A spread well above 1.1x on identical parts under no load is the\n"
              "  finding. Report the whole table, not one card.")
    else:
        print("\n  ONE GPU ONLY. This number has no reference. Nothing here can\n"
              "  distinguish a slow card from a slow card model -- run every card.")

    if a.out:
        json.dump(out, open(a.out, "w"), indent=1)
        print(f"\n  -> {a.out}")


if __name__ == "__main__":
    main()

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

/* 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

硬件为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/api/client/document/preview/1289/split_files/%E7%BC%96%E8%AF%91%E5%92%8C%E8%B0%83%E8%AF%95.html#yjwrqerouotc1的文档描述，MACA_EXT_DIRECT_DISPATCH设置为1时MACA_EXT_SYNC_POLICY 自动设置成3。但是在/opt/maca/include/mcctrl/feature.h里面没有找到相关的自动设置代码，请问是以文档为准还是以代码为准？

收到，谢谢管理员

请问管理员：本地镜像maca中/opt/maca/include/mcctrl/feature.h里面的配置参数有没有相关的说明文档？具体是哪一个文档？
例如MACA_EXT_CPU_THREAD_POLICY和MACA_EXT_NUMA_MEMORY_POLICY，我在实验中发现MACA_EXT_NUMA_MEMORY_POLICY设置为0会忽略外部numactl绑定，MACA_EXT_CPU_THREAD_POLICY默认值0（balance policy）的确切语义是什么？
maca版本为：3.8.0.103
谢谢

8专家并行，docker是vllm-metax-0.23.0-maca.ai3.8.0.103-torch2.10-py312-ubuntu22.04-amd64.tar.xz，先是带CUDA graph的启动失败，报驱动层面的trace，index_elementwise_kernel越界；换eager启动以后，发一个请求出现：
(Worker_DP4_EP4 pid=237836) WARNING 08-19 10:28:15 [fused_moe.py:1379] Using default MoE config. Performance might be sub-optimal! Config file not found at /opt/conda/lib/python3.12/site-packages/vllm_metax/model_executor/layers/fused_moe/configs/H=2048/H=2048,E=32,N=512,device_name=MXC500.json, /opt/conda/lib/python3.12/site-packages/vllm_metax/model_executor/layers/fused_moe/configs/E=32,N=512,device_name=MXC500.json
[10:28:15.415][MXC][E]xnack(0x8): kernel causes atu address translation error
[10:28:15.415][MCR][E]mx_trapProcess.cpp :1405: call get queue logic id fail
[10:28:15.415][MCR][E]mx_trapProcess.cpp :130 : trapping pipeID=0,queueID=0!
[10:28:15.415][MCR][E]mx_trapProcess.cpp :131 : trapping virtualDevice=0x7f907c000c70!
[10:28:15.416][MCR][E]mx_trapProcess.cpp :998 : Node ID is 3, Wave ID is 32646
[10:28:15.416][MCR][E]mx_trapProcess.cpp :1021: debug info: 0x7c,0x00,0x00,0x10,0x16,0x2c,0x7f,0x8c,0x7c,0x00,0x00,0x10,0x14,0x24,0x7f,0x8c,
[10:28:15.416][MCR][E]mx_device.cpp :10571: trap:precise positioning.
[10:28:15.416][MCR][E]mx_device.cpp :10576: trapping: kernelName: _ZN7mctlass6KernelINS_4gemm6kernel11MacaGemmMoeINS1_11threadblock20MacaMoeMmaMultistageINS1_9GemmShapeILi16ELi128ELi128EEENS6_ILi16ELi32ELi128EEE15__maca_bfloat16NS_6layout8RowMajorES9_NSA_11ColumnMajorEiSB_Li2ELi1ELb1ELb0ENS_4arch4Sm80ELb0EbEENS_8epilogue11threadblock30MacaMoeGemmEpilogueDirectStoreIS7_S9_SB_S9_SC_Li1ENSG_6thread24MacaLinearCombinationMoeIfLi2EfS9_LNSJ_9ScaleType13ScaleBiasKindE11ELb0ELNS_15FloatRoundStyleE2EEES9_SB_iLi2ELi1ELb1ELb0EbEENS4_37GemmBatchedIdentityThreadblockSwizzleEbEEEEvNT_6ParamsE ,commandIndex: 6573 , trapType: Xnack Error/ATU Fault(0x8)
[10:28:15.416][MCR][E]mx_trapProcess.cpp :918 : Xnack(0x8) exception happened in the shader, the mcruntime api will be disabled
[10:28:15.428][MCR][E]mc_runtime_api.cpp :215 : 237839: [7f9368735740] mcGetDevice: Returned mcErrorIllegalAddress
terminate called after throwing an instance of 'c10::AcceleratorError'
what(): CUDA error: an illegal memory access was encountered
Search for cudaErrorIllegalAddress' in https://docs.nvidia.com/cuda/cuda-runtime-api/group__CUDART__TYPES.html for more information. CUDA kernel errors might be asynchronously reported at some other API call, so the stacktrace below might be incorrect. For debugging consider passing CUDA_LAUNCH_BLOCKING=1 Compile withTORCH_USE_CUDA_DSA` to enable device-side assertions.

问题：C500 显存经ibv_reg_dmabuf_mr注册给mlx5网卡。GPU kernel写这块显存后，GPU的写停在L2域，网卡通过dma-buf P2P读HBM读到的是旧值。
提问：
1. C500 GPUDirect 是否支持网卡从显存 READ(从显存发送),还是只支持网卡写显存(接收)?
2. 若支持，有没有让GPU kernel的显存写对网卡P2P读可见的 API（L2→HBM flush/system-scope store/dma-buf 注册前的正确内存属性)?mcDeviceFlushGPUDirectRDMAWrites是接收侧,有无发送侧对应？

我现在有一个.cu文件，如何把他转换成MACA源代码比如.maca或者直接编译成可执行文件

需要vllm-metax:0.21.0-maca.ai3.7.1.106-torch2.8-py310-ubuntu22.04-amd64离线镜像包导出

目前想把C500的RDMA传输控制面从CPU改为由GPU kernel来做，需要GPU和MACA支持GDAKI / GPUDirect Async，请问目前是否支持，有相关文档吗？

谢谢，已发帖申请

需要vllm-metax:0.18.0-maca.ai3.5.3.405-torch2.8-py310-ubuntu22.04-amd64

能否优化一下呢？我重新复制docker命令以后并没有触发断点续传，是完全重新下载了，在时间范围内下载不完，这个错误就循环出现了，这是我刚刚复制命令后下载的：
Login Succeeded
0.18.0-maca.ai3.5.3.405-torch2.8-py310-ubuntu22.04-amd64: Pulling from public-ai-release/maca/vllm-metax
43f89b94cd7d: Already exists
f2b7c1644807: Pull complete
077977940afd: Pull complete
b1b8f989434a: Pull complete
84398257dc67: Downloading 6.488MB/94.22MB
443c93c96edb: Downloading 4.866MB/3.572GB
2f427e59f857: Downloading 1.229MB/3.485MB
b0c2e660298b: Waiting
0c1d591115bb: Waiting
30e855b71d68: Waiting
06830532ea96: Pulling fs layer
4f4fb700ef54: Waiting
afe16f59bac8: Waiting
1d0f707f6f52: Waiting
96b54463a433: Waiting
37bec76d27ba: Waiting