GPU Vendor Porting Guide (Tutti cuda_like Framework)

This document describes the three-layer GPU-vendor abstraction in Tutti and the integration steps a new vendor (MUSA / MACA / future) must follow to plug their device compiler + runtime SDK into the project.

Audience

  • Tutti maintainers — own the framework contracts.

  • Vendor port engineers (e.g. Metax 沐曦) — fill in the per-vendor shim and compile their toolchain against Tutti’s kernels.

The three layers

┌─────────────────────────────────────────────────────────────┐
│ Layer A: Vendor selection framework                          │
│   cmake/accelerators/<VENDOR>.cmake     ← toolchain + libs   │
│   include/tutti/gpu_vendor/<vendor>.h   ← runtime API shim   │
│   include/tutti/cuda_like.h             ← profile selector  │
├─────────────────────────────────────────────────────────────┤
│ Layer B: Kernel primitives macro layer                       │
│   data_paths/.../io/tutti_gpu_primitives.cuh                │
│   ↑ device-only macros (TUTTI_DEVICE, TUTTI_THREADFENCE_*,   │
│     TUTTI_ATOMIC_ADD, ...) with semantic contracts           │
│   vendor fills: include/tutti/gpu_vendor/<vendor>_primitives.h│
├─────────────────────────────────────────────────────────────┤
│ Layer C: Direct includes → cuda_like                         │
│   All host/.cu/.cuh files include <tutti/cuda_like.h>       │
│   instead of <cuda_runtime.h> directly                       │
└─────────────────────────────────────────────────────────────┘

A single profile is selected at CMake configure time (-DTUTTI_ACCELERATOR=CUDA|HOST|MUSA|MACA). The build is single-vendor (no runtime coexistence). Each profile wires its own:

  • Runtime API surface (cudaMalloc / cudaMemcpy / cudaStream_t / …)

  • Device compiler (nvcc / mcc / …)

  • Architecture flags

  • Link libraries

Layer A — Vendor selection framework

cmake/accelerators/<VENDOR>.cmake

Each profile provides a tutti_configure_cuda_like(<target_name>) function that:

  1. Defines TUTTI_USE_<VENDOR>=1 on the target.

  2. Adds include/ to the target’s include path.

  3. Links the vendor’s runtime + driver libraries.

Each profile also sets two CMake variables that downstream CMakeLists use to link the runtime/driver without naming the vendor-specific target:

set(TUTTI_ACCEL_RUNTIME_LIBS "CUDA::cudart" CACHE INTERNAL "")
set(TUTTI_ACCEL_DRIVER_LIBS  "CUDA::cuda_driver" CACHE INTERNAL "")

Downstream targets link ${TUTTI_ACCEL_RUNTIME_LIBS} / ${TUTTI_ACCEL_DRIVER_LIBS} (never CUDA::cudart directly) so vendor ports don’t touch every CMakeLists.txt.

Vendor step A.1: copy cmake/accelerators/MUSA.cmake (or MACA.cmake) and replace the TODO markers:

# TODO(Metax): locate MUSA SDK
find_path(MUSA_INCLUDE_DIR musa_runtime.h PATHS /opt/musa/include)
find_library(MUSA_RUNTIME_LIB  NAMES musart  PATHS /opt/musa/lib)
find_library(MUSA_DRIVER_LIB   NAME  musa     PATHS /opt/musa/lib)
target_include_directories(${target_name} INTERFACE ${MUSA_INCLUDE_DIR})
target_link_libraries(${target_name} INTERFACE ${MUSA_RUNTIME_LIB} ${MUSA_DRIVER_LIB})
set(TUTTI_ACCEL_RUNTIME_LIBS "${MUSA_RUNTIME_LIB}" CACHE INTERNAL "")
set(TUTTI_ACCEL_DRIVER_LIBS  "${MUSA_DRIVER_LIB}"  CACHE INTERNAL "")

If the vendor compiler is NOT nvcc, enable the device language in the profile module:

enable_language(MUSA)  # or whatever the toolchain exposes

… and adjust project(... LANGUAGES ...) in the root CMakeLists.txt accordingly.

include/tutti/gpu_vendor/<vendor>.h

A header that supplies the runtime API surface Tutti’s host code calls. This is the host-side shim — it does NOT define kernel-side primitives (those live in Layer B).

Look at include/tutti/gpu_vendor/host.h for the full surface (expected symbols: cudaMalloc, cudaFree, cudaMemcpy, cudaStreamCreate, cudaEventCreate, cudaGetDeviceProperties, cudaPointerGetAttributes, …).

Vendor step A.2: replace the #error at the top of include/tutti/gpu_vendor/musa.h with the real SDK #include and define / typedef every symbol listed in the TODO skeleton.

include/tutti/cuda_like.h

The profile selector — single header that downstream files include. It dispatches based on TUTTI_USE_<VENDOR> and includes the matching vendor header. Already wired; vendor does not modify this file.

Layer B — Kernel primitives macro layer

File: tutti/data_paths/local_nvme/io/tutti_gpu_primitives.cuh

This file provides macros for the CUDA C++ language extensions used by Tutti’s IO kernels. Each macro has a semantic contract documented in the file. The macros fall into groups:

Group

Macros

Notes

Function attrs

TUTTI_DEVICE, TUTTI_HOST, TUTTI_GLOBAL, TUTTI_FORCEINLINE

Mark function entry / linkage

Variable attrs

TUTTI_SHARED

Per-block shared memory

Built-ins

TUTTI_THREAD_IDX_X, TUTTI_BLOCK_IDX_X, TUTTI_BLOCK_DIM_X

Thread / block indices

Fences

TUTTI_THREADFENCE_SYSTEM (CRITICAL), TUTTI_THREADFENCE, TUTTI_SYNC_THREADS

Memory / control barriers — see contract

Atomics

TUTTI_ATOMIC_ADD, TUTTI_ATOMIC_CAS

Device-global scope

Warp

TUTTI_SHFL_SYNC, TUTTI_BALLOT_SYNC

Full-warp mask required

Timing

TUTTI_NANOSLEEP, TUTTI_CLOCK, TUTTI_CLOCK64

Best-effort

Why TUTTI_THREADFENCE_SYSTEM is the命根

NVMe doorbell writes must observe a precise memory-ordering contract:

  1. CPU / GPU writes the SQE (submission queue entry) into the SQ ring.

  2. The thread issues TUTTI_THREADFENCE_SYSTEM to flush all outstanding writes — including the SQE — to a point of system-wide visibility (GPU + host + peer PCIe / NVMe controller).

  3. Then the thread writes the doorbell register (sq.db).

If the fence is missing or has weaker scope (device-only), the NVMe controller may observe a stale doorbell value (wrote SQ tail = N, but the SQE for slot N is not yet visible to the controller) → silent IO corruption / hang.

The contract is documented in the header — every vendor port MUST preserve it.

Vendor step B.1: provide <vendor>_primitives.h

Create tutti/include/tutti/gpu_vendor/musa_primitives.h (or maca_primitives.h) with the macro definitions for the vendor’s device compiler. Two cases:

Case 1 — vendor compiler accepts CUDA-style keywords verbatim (the most common case for CUDA-like SDKs). Copy the CUDA block from tutti_gpu_primitives.cuh:

#pragma once
#define TUTTI_DEVICE              __device__
#define TUTTI_GLOBAL              __global__
#define TUTTI_FORCEINLINE         __forceinline__
#define TUTTI_THREAD_IDX_X        threadIdx.x
// ... (copy from tutti_gpu_primitives.cuh CUDA branch)

Case 2 — vendor compiler uses different syntax. Provide the MUSA / MACA equivalents per the contract table in tutti_gpu_primitives.cuh.

tutti_gpu_primitives.cuh already uses __has_include(...) to pick up your vendor header, so once the file exists the build will work.

Kernel launch syntax (<<<>>>)

Tutti’s host launchers (submit_one.cu, fused_submit_kernel.cu) use CUDA’s kernel<<<blocks, threads, smem, stream>>>(args...) triple-bracket syntax. Every CUDA-like compiler we have surveyed (nvcc, mcc, …) accepts this syntax. If the vendor’s compiler does NOT, raise it on the Tutti issue tracker — the launchers will be updated to use a TUTTI_LAUNCH_KERNEL macro.

Layer C — Direct includes

Every host / .cu / .cuh file in tutti/ includes <tutti/cuda_like.h> instead of <cuda_runtime.h> directly. The single source of truth for “which vendor are we building against” is the TUTTI_USE_<VENDOR> compile definition propagated by the tutti_cuda_like interface target.

Files NOT modified (libnvm upstream headers)

The following headers under tutti/device_manager/nvme/libnvm/include/ are upstream libnvm sources and continue to #include <cuda.h> / #include <cuda_runtime.h> / #include <cuda/atomic> directly:

  • buffer.h, event.h, queue.h, util.h — #include "cuda.h"

  • nvm_util.h — #include <cuda.h>, #include <cuda_runtime.h>

  • nvm_types.h — #include <cuda/atomic>, #include <cuda_runtime.h> (under #ifdef NVM_CUDA_ENABLED)

Modifying these creates upstream divergence. Two options for the vendor:

  1. Patch the headers locally (fork of libnvm) — straightforward, maintain a small diff.

  2. Provide a vendor shim header that libnvm’s #include "cuda.h" picks up via -I — e.g. create <vendor_sdk>/include/cuda.h that #includes the MUSA / MACA equivalent. Less invasive but requires the include path ordering to land the shim ahead of the system CUDA path.

The choice is the vendor’s; document the decision in the vendor’s internal Tutti fork.

Kernel module P2P adaptation (Metax)

The snvme kernel module (tutti/device_manager/nvme/kernel_modules/) uses the peer_memory_ops function-pointer table in peer_memory/peer_memory.h as the single point of contact with the GPU driver’s P2P API.

Current state

  • NVIDIA path (CUDA default): peer_memory/nvidia.c resolves nvidia_p2p_* symbols at module-load time via __symbol_get and routes calls through the peer_memory_ops table. The module body (map.c, pci.c) never touches NVIDIA types directly.

  • Metax path (MUSA/MACA default): peer_memory/metax.c includes metax_p2p.h (ported from ljye2023/Tutti PR #1) and resolves metax_p2p_* symbols at module-load time via __symbol_get (symmetric to the NVIDIA path). The peer_memory_ops implementation is complete; it works as soon as the Metax driver is loaded and exports the symbols.

metax_p2p.h API (ported from Mooncake/Metax PR)

The vendor-supplied metax_p2p.h defines these functions. Tutti does not currently vendor this header; provide its directory through SNVME_P2P_INCLUDE_DIR.

Function

Purpose

NVIDIA equivalent

metax_p2p_acquire_mem

Pin a GPU VA range

nvidia_p2p_get_pages

metax_p2p_get_mem

Get bus addresses as sg_table

(part of dma_map_pages)

metax_p2p_put_mem

Unpin mem

nvidia_p2p_dma_unmap_pages

metax_p2p_release_mem

Release handle

nvidia_p2p_put_pages

metax_p2p_get_pci_dev

Get GPU’s PCI device

(from page_table)

metax_p2p_get_bus_offset

Bus→CPU address offset

(implicit in dma_mapping)

metax_p2p_get_page_size

Page size of pinned region

(GPU_PAGE_SIZE)

Metax integration steps

  1. Configure with -DTUTTI_BUILD_KERNEL_MODULE=ON, -DTUTTI_P2P_BACKEND=metax, and -DSNVME_P2P_INCLUDE_DIR=<directory-containing-metax_p2p.h>.

  2. Load the Metax GPU driver (must export metax_p2p_* symbols to the kernel symbol namespace).

  3. insmod snvme.ko — peer_init() resolves the symbols via __symbol_get; if the Metax driver is not loaded, peer_init() fails gracefully (module load fails with an error message, same as the NVIDIA path when nvidia_p2p_* symbols are missing).

  4. No SNVMe common-source changes are needed — the backend implementation lives in peer_memory/metax.c.

The NVIDIA→Metax API mapping table is documented inline in peer_memory/metax.c.

Note on map.c signature differences

The PR (#1) also modifies map.c’s call to nvfs_nvidia_p2p_dma_map_pages to add page_size and n_addrs parameters, and changes pages->entries to pages->virtual_entries. These are Metax driver API differences — in Tutti’s architecture, map.c calls through the peer_memory_ops function-pointer table, so the signature is fixed by peer_memory.h. The Metax implementation of peer_memory_ops handles the API differences internally; map.c is not modified.

Verification

A vendor port is complete when:

  1. CUDA profile: zero behavioral change — ctest -LE "hardware" passes 15/15 (the hardware-free contract suite).

  2. HOST profile: ctest -L host passes (excluding the mount_manager test which depends on nvmeservice and is host-broken by design).

  3. Vendor profile: cmake -S . -B build -DTUTTI_ACCELERATOR=<VENDOR> configures successfully with the vendor’s toolchain, and the stub-removal #errors in <vendor>.h no longer fire.

  4. Vendor’s device compiler compiles the __CUDACC__-guarded section of submit_one.cuh / fused_submit_kernel.cuh / nvme_submit_primitives.cuh (these are the device code entry points).

  5. Hardware smoke: run tutti_layerwise_kv_overlap (KV-cache pipeline simulator) end-to-end. If the port is sound, this exercises the kernel primitives + launch path.

Reference: host.h as the canonical shim

The HOST profile (include/tutti/gpu_vendor/host.h) is the reference implementation of the Layer A surface — it implements the entire cuda_like API on top of std::malloc / std::memcpy / std::memset with no GPU at all. Vendors should mirror its symbol layout (same names, same signatures) so call sites compile unchanged.

Quick start for Metax (MUSA / MACA)

大部分工作已经从 Mooncake 搬过来并验证。 Metax 只需补齐 libnvm upstream 头的兼容层。

已从 Mooncake 搬过来的确定性实现

文件

来源

状态

tutti/include/tutti/gpu_vendor/musa.h

Mooncake gpu_vendor/musa.h（Apache 2.0）

✅ 完整 runtime API 宏映射（cuda* → musa*），去掉 Mooncake 特有的 GPU_PREFIX 和 IBGDA driver API

tutti/include/tutti/gpu_vendor/maca.h

Mooncake gpu_vendor/maca.h（Apache 2.0）

✅ 完整 runtime API 宏映射（cuda* → mc*），含模板函数 cudaHostGetDevicePointer / cudaFuncGetAttributes

tutti/include/tutti/gpu_vendor/musa_primitives.h

Mooncake musa_ops.cuh 验证

✅ Case 1（mcc 接受 CUDA-style 关键字），含 Mooncake 验证过的已知 SDK 4.3.3 bug 注释（atomicAdd_system SelectionDAG bug、named barriers 不可用、grid sync 不可用——Tutti 的 kernel 都不使用这些有 bug 的特性）

tutti/include/tutti/gpu_vendor/maca_primitives.h

Mooncake maca_ops.cuh 验证

✅ Case 1（cu-bridge 接受 CUDA-like intrinsics），Mooncake 确认 __threadfence_system / __ldg / __syncthreads 可用

tutti/cmake/accelerators/MUSA.cmake

Mooncake common.cmake 验证

✅ SDK 路径 /usr/local/musa/include + /usr/local/musa/lib，库名 musa musart

tutti/cmake/accelerators/MACA.cmake

Mooncake common.cmake 验证

✅ SDK root $MACA_HOME 或 /opt/maca，include ${MACA_ROOT}/include，lib ${MACA_ROOT}/lib64 或 ${MACA_ROOT}/lib

Metax 剩余工作（仅 libnvm upstream 头兼容层）

libnvm upstream 头（tutti/device_manager/nvme/libnvm/include/）仍直引 cuda.h / cuda_runtime.h / cuda/atomic：

  • buffer.h, event.h, queue.h, util.h — #include "cuda.h"

  • nvm_util.h — #include <cuda.h>, #include <cuda_runtime.h>

  • nvm_types.h — #include <cuda/atomic>, #include <cuda_runtime.h>（__CUDACC__ guard 内）

Metax 选一种：

  1. Vendor shim cuda.h（推荐，零 fork）：在 MUSA SDK 的 include 路径下放一个 cuda.h → #include <musa.h> + 宏映射。cuda_runtime.h 同理。cuda/atomic → 创建 cuda/atomic 头文件转发到 musa/atomic。这样 libnvm 头不改一行。

  2. Fork patch libnvm：把 libnvm 头里的 cuda.h / cuda_runtime.h 改为 #include <tutti/cuda_like.h>，cuda/atomic 改为 vendor 等价物。维护一个小 diff。

方案 1 更轻量——Metax 的 SDK 已经提供了 musa.h / musa_runtime.h（Mooncake 验证），只需要补一个 cuda/atomic 转发头即可。

验证步骤

  1. cmake -S . -B build -DTUTTI_ACCELERATOR=MUSA（configure 应该成功，无 WARNING）

  2. cmake --build --preset cuda --target tutti_layerwise_kv_overlap（如果链接成功且跑通，框架级通过）

  3. 跑 CUDA profile 回归确认零影响：cmake --preset cuda --fresh && ctest --preset cuda -LE "hardware|mount_manager"