/* SPDX-License-Identifier: AGPL-3.0-or-later */ /* Copyright (C) 2026 David (Corey) Herrera */ /* * fc_cuda.h — CUDA GPU compute backend for Foundry * * Module prefix: fc_ (foundry cuda) * Build flag: USE_CUDA=1 * * Provides GPU-accelerated tensor operations via NVIDIA CUDA. * Uses dlopen() for zero compile-time CUDA dependency — falls back * to CPU if CUDA is unavailable at runtime. * * Key difference from Metal: CUDA targets discrete GPUs with separate * VRAM, so buffers require explicit Host-to-Device / Device-to-Host * copies. Weights should be uploaded once and kept resident. * * Target: CUDA 11.4 / sm_37 (K80) — no FP16 math, no tensor cores, * uses __shfl_down() not __shfl_down_sync(). */ #ifndef FOUNDRY_CUDA_H #define FOUNDRY_CUDA_H #include #include #include #ifdef __cplusplus extern "C" { #endif /* ================================================================ * Opaque types * ================================================================ */ typedef struct fc_buffer fc_buffer; /* Per-device CUDA context. Holds one device's context, stream, * cuBLAS handle, PTX modules and kernel function handles. One * instance per K80 die lets a single process drive multiple dies * concurrently. The dlopen'd driver function table is process-wide * and shared by all contexts. Opaque — see fc_cuda.c for fields. */ typedef struct fc_cuda_ctx fc_cuda_ctx; /* ================================================================ * Process-wide driver loading * * The CUDA driver / cuBLAS shared libraries are dlopen'd once per * process; their function pointers are shared by all contexts. * fc_cuda_init_device() loads them lazily on first call, so most * callers never need fc_driver_load() directly. * ================================================================ */ /* dlopen libcuda + libcublas and resolve all driver symbols, then * cuInit(0). Idempotent. Returns 0 on success, -1 on failure * (e.g. no CUDA driver present — the Mac case). */ int fc_driver_load(void); /* Returns true once fc_driver_load() has succeeded. */ bool fc_driver_loaded(void); /* ================================================================ * Per-device context lifecycle (multi-die) * * Each context owns one device's CUcontext, stream, cuBLAS handle, * and a private copy of the PTX modules/kernels. Operations taking * an fc_cuda_ctx* make that context current (push) for the call and * restore the previous context (pop) on return, so contexts for * different dies never bleed into each other. * ================================================================ */ /* Initialize CUDA for a specific device id: dlopen driver (if not * already), cuDeviceGet, cuCtxCreate, stream + cuBLAS create, and * per-context cuModuleLoadData of the embedded PTX. * * @param device_id CUDA ordinal of the die to bind (0-based) * @return Owned context handle, or NULL on failure. * Free with fc_cuda_ctx_free(). * * Thread-safe: no (driver load + cuCtxCreate are not serialized here) * Reentrant: no */ fc_cuda_ctx* fc_cuda_init_device(int device_id); /* Destroy a per-device context: sync, unload modules, destroy * cuBLAS handle, stream and context. Safe to call with NULL. */ void fc_cuda_ctx_free(fc_cuda_ctx *ctx); /* Returns true if the context is initialized and ready. */ bool fc_cuda_ctx_available(const fc_cuda_ctx *ctx); /* Returns the CUDA device ordinal this context is bound to, or -1. */ int fc_cuda_ctx_device_id(const fc_cuda_ctx *ctx); /* ================================================================ * MemoryAware VRAM signals (MemoryAware.md §2.2) * ================================================================ */ /* Query this die's free/total VRAM via cuMemGetInfo, with ctx made current for * the call. Authoritative at the moment of the call but racy under concurrent * alloc, so it is a decision-point cross-check, not a per-token input. Either * output pointer may be NULL. Returns 0 on success, -1 if unavailable. * * Thread-safe: no (makes ctx current; serialize with other ctx ops) * Reentrant: no */ int fc_mem_info_ctx(fc_cuda_ctx *ctx, size_t *free_bytes, size_t *total_bytes); /* Same reading, but for a die this process does NOT own: creates a bare * context (no cuBLAS handle, no PTX modules), reads cuMemGetInfo, destroys it. * The die is left byte-identical. Use this for enumeration/probing; use * fc_mem_info_ctx for a die you already hold a context on. * * Probing with fc_cuda_init_device() instead is what stranded a full ~124 MiB * context on every die a process merely looked at — including dies belonging * to another serve. * * Thread-safe: no (creates/destroys a context on the calling thread) * Reentrant: no */ int fc_cuda_probe_device(int device_id, size_t *free_bytes, size_t *total_bytes); /* Running sum of bytes this context has allocated and not yet freed, tallied at * the buffer choke points. Free + predictive — the number admission schedules * against (MemoryAware.md §6). Returns 0 for a NULL ctx. * * Thread-safe: no (plain read of a counter mutated on the alloc path) * Reentrant: yes */ size_t fc_ctx_alloc_bytes(const fc_cuda_ctx *ctx); /* ================================================================ * Crash-time teardown (MemoryAware.md §8) * ================================================================ */ /* Number of ready contexts currently registered (live dies). */ int fc_cuda_live_ctx_count(void); /* Drain + cuCtxDestroy every live context, then the caller MUST exit so * libcuda closes the /dev/nvidiaN fd — the only step that reclaims Kepler VRAM * (cuCtxDestroy alone does not). Intended for the SIGINT/SIGTERM shutdown path. * * Thread-safe: no * Reentrant: no * Signal-safe: NO — calls into libcuda. Use only from a controlled-shutdown * signal (SIGINT/SIGTERM), never from SIGSEGV/SIGABRT (there, _exit() and * lean on kernel fd-close + startup stranded-VRAM detection). */ void fc_cuda_emergency_teardown(void); /* ================================================================ * Lifecycle (process-default context — back-compat single-die path) * * These keep every existing single-context call site working. The * first fc_init_on() creates a process-default fc_cuda_ctx bound to * the device it names; all the non-ctx functions below (fc_buffer_*, * fc_sgemm, kernels, fc_sync) delegate to it. Multi-die callers should * use the fc_cuda_ctx* variants instead. * ================================================================ */ /* Initialize CUDA on `device_id` into the process-default context: * dlopen libcuda + libcublas, cuInit, cuCtxCreate, load embedded PTX * kernels. Returns 0 on success, -1 on failure. Idempotent — the FIRST * call wins, so a caller that knows its placement must call this before * anything else forces a default context onto a die it does not own. */ int fc_init_on(int device_id); /* fc_init_on(0). Only correct when the process genuinely owns device 0 * (or CUDA_VISIBLE_DEVICES has pinned its die to ordinal 0). Callers that * have a placement MUST use fc_init_on() — this is how a run pinned to * two other dies ended up holding a context on die 0. */ int fc_init(void); /* True once libcuda is loadable and cuInit has succeeded. Unlike * fc_available() this creates NO context, so it is the right gate for * "is there a GPU at all" questions asked before placement is known. */ bool fc_driver_available(void); /* Cleanup the process-default context and unload driver libraries. */ void fc_free(void); /* Returns true if the process-default context is initialized. */ bool fc_available(void); /* Returns the process-default context, or NULL if fc_init() has not * succeeded. Lets a caller bridge from the back-compat path to the * explicit-ctx API without re-initializing. */ fc_cuda_ctx* fc_default_ctx(void); /* ================================================================ * Buffers (discrete GPU memory — explicit H2D/D2H copies) * * Unlike Metal's unified memory, CUDA buffers live in GPU VRAM. * fc_buffer_from_ptr() allocates GPU memory and copies data from * CPU. Use fc_buffer_to_cpu() to read results back. * * For persistent data (model weights), create the buffer once at * load time and reuse it across inference steps. * ================================================================ */ /* --- Per-context variants (multi-die) --- * Each allocation is tagged with its owning context; frees and copies * automatically make that context current. The buffer must only be * used with the context that created it. */ /* Allocate a GPU buffer on ctx's device and copy CPU data into it. */ fc_buffer* fc_buffer_from_ptr_ctx(fc_cuda_ctx *ctx, const void *cpu, size_t bytes); /* Allocate an uninitialized GPU buffer on ctx's device. */ fc_buffer* fc_buffer_alloc_ctx(fc_cuda_ctx *ctx, size_t bytes); /* --- Process-default-context variants (back-compat) --- */ /* Allocate GPU buffer and copy CPU data to it (Host-to-Device). */ fc_buffer* fc_buffer_from_ptr(const void *cpu, size_t bytes); /* Allocate GPU buffer without initializing (for output buffers). */ fc_buffer* fc_buffer_alloc(size_t bytes); /* Free a GPU buffer (uses the buffer's owning context). */ void fc_buffer_free(fc_buffer *buf); /* Non-owning offset view: returns a heap fc_buffer whose device pointer is * base's + byte_off (size shrunk by byte_off), sharing base's owning context. * It does NOT own the device memory — free ONLY with fc_buffer_free_view * (which frees the wrapper struct, never the underlying allocation). Used by * the die-interleave (L5) to launch a slot kernel over a sub-range * [slot_base, slot_base+count) of a persistent per-slot ring while the batch * scratch stays group-local [0,count): the ring buffer is offset by * slot_base*row_bytes so the kernel's slot index 0 lands at slot_base. * Returns NULL on bad args (byte_off past the end). */ fc_buffer* fc_buffer_offset_view(fc_buffer *base, size_t byte_off); /* Free a view created by fc_buffer_offset_view: frees the wrapper struct only, * NOT the aliased device memory. A no-op on NULL. */ void fc_buffer_free_view(fc_buffer *view); /* Copy GPU buffer contents back to CPU (Device-to-Host). */ void fc_buffer_to_cpu(fc_buffer *buf, void *dst); /* Copy CPU data into an existing GPU buffer (Host-to-Device). */ void fc_buffer_from_cpu(fc_buffer *buf, const void *src); /* Get buffer size in bytes. */ size_t fc_buffer_size(fc_buffer *buf); /* Partial copies (first `bytes` of the buffer; clamped to its size). The * batch path sizes persistent buffers for the max batch but moves only the * live n tokens' bytes per call. */ void fc_buffer_to_cpu_n(fc_buffer *buf, void *dst, size_t bytes); void fc_buffer_from_cpu_n(fc_buffer *buf, const void *src, size_t bytes); /* Async device-to-host copy queued on the buffer's owning ctx stream (non- * blocking). Stream-ordered after prior kernel work; the caller MUST * fc_sync_ctx(owner) before reading `dst`. Returns 0 on success, -1 if async * DtoH is unavailable (caller should fall back to the blocking variant). */ int fc_buffer_to_cpu_n_async(fc_buffer *buf, void *dst, size_t bytes); /* Host-to-device copy into a byte offset of the buffer (bounds-checked * no-op on overflow). Used to re-sync a row range of the GPU KV cache * from the host mirror. */ void fc_buffer_from_cpu_at(fc_buffer *buf, size_t dst_off, const void *src, size_t bytes); /* Device-to-host copy from a byte offset of the buffer (bounds-checked * no-op on overflow). Used to refresh a row range of the host KV mirror * from the GPU cache byte-for-byte. */ void fc_buffer_to_cpu_at(fc_buffer *buf, void *dst, size_t src_off, size_t bytes); /* Async device-to-device range copy on ctx's stream (both buffers must * belong to ctx; bounds-checked no-op otherwise). */ void fc_buffer_copy_d2d_ctx(fc_cuda_ctx *ctx, fc_buffer *dst, size_t dst_off, fc_buffer *src, size_t src_off, size_t bytes); /* ================================================================ * Peer (die-to-die) copies — for multi-die tensor/vocab split on cards * whose dies are separate CUDA devices (K80: two GK210 dies behind an * on-card PLX switch, cudaDeviceCanAccessPeer==1, PIX single-bridge path). * ================================================================ */ /* Enable P2P access between two contexts' devices in BOTH directions. * Idempotent: a second call for an already-enabled pair returns 0. Returns * 0 on success (or if the pair is already enabled), -1 if either context is * unavailable, the driver lacks the peer symbols, or the devices cannot * access each other (cuDeviceCanAccessPeer==0). Call once at model load. */ int fc_cuda_enable_peer(fc_cuda_ctx *a, fc_cuda_ctx *b); /* Blocking device-to-device peer copy: copy `bytes` from src (owned by * src's context/device) into dst (owned by dst's context/device), across * the peer link (cuMemcpyPeer). Both buffers may live on different devices. * Synchronous (uses cuMemcpyPeer, which blocks until complete). Bounds- * checked no-op if bytes exceeds either buffer. Requires fc_cuda_enable_peer * to have succeeded for the pair. Returns 0 on success, -1 on error. */ int fc_buffer_copy_peer(fc_buffer *dst, fc_buffer *src, size_t bytes); /* Async (non-blocking) device-to-device peer copy queued on stream_ctx's * stream — stream-ordered after prior work on that stream, does NOT block the * host. stream_ctx must be one of the two buffers' owners (normally dst's, so * the receiving die both issues the copy and can event-sync before reading * dst). Caller must fc_sync_ctx(stream_ctx) or wait on an event recorded after * this before reading dst. Requires fc_cuda_enable_peer for the pair. Returns 0 * on success, -1 if cuMemcpyPeerAsync is unavailable (fall back to the blocking * fc_buffer_copy_peer). */ int fc_buffer_copy_peer_async(fc_buffer *dst, fc_buffer *src, size_t bytes, fc_cuda_ctx *stream_ctx); /* ================================================================ * Dispatch thresholds * * On discrete GPU, even single-token decode can benefit from GPU * since GPU memory bandwidth (240 GB/s on K80) exceeds CPU * (~20 GB/s on server). However, PCIe transfer overhead means * we need sufficient work to amortize the copy cost. * * These thresholds assume weights are already resident in VRAM * (only activations transfer per token). * ================================================================ */ #define FC_GEMM_THRESHOLD (64L * 1024 * 1024) /* ~64M ops */ #define FC_GEMV_THRESHOLD (4096L * 4096) /* ~16M ops */ #define FC_Q8_THRESHOLD (0) /* always use GPU for Q8 matvec */ #define FC_ELEM_THRESHOLD (1024 * 1024) /* 1M elements */ /* ================================================================ * BLAS via cuBLAS (dlopen'd) * * All matrices are row-major float32. cuBLAS is column-major, * so we swap A/B and transpose the problem internally. * ================================================================ */ /* C = alpha * A @ B + beta * C * A: [M, K], B: [K, N], C: [M, N] — all row-major float32. */ void fc_sgemm_ctx(fc_cuda_ctx *ctx, fc_buffer *A, fc_buffer *B, fc_buffer *C, int M, int N, int K, float alpha, float beta); void fc_sgemm(fc_buffer *A, fc_buffer *B, fc_buffer *C, int M, int N, int K, float alpha, float beta); /* y = alpha * A @ x + beta * y * A: [M, K] row-major, x: [K], y: [M] — all float32. */ void fc_sgemv_ctx(fc_cuda_ctx *ctx, fc_buffer *A, fc_buffer *x, fc_buffer *y, int M, int K, float alpha, float beta); void fc_sgemv(fc_buffer *A, fc_buffer *x, fc_buffer *y, int M, int K, float alpha, float beta); /* ================================================================ * Custom kernels (loaded from embedded PTX) * ================================================================ */ /* RMS normalization over last dimension. * x: [n], weight: [n] (may be NULL), out: [n] */ void fc_rms_norm_ctx(fc_cuda_ctx *ctx, fc_buffer *x, fc_buffer *weight, fc_buffer *out, int n, float eps); void fc_rms_norm(fc_buffer *x, fc_buffer *weight, fc_buffer *out, int n, float eps); /* Softmax over last dimension. * x: [rows, cols], out: [rows, cols] */ void fc_softmax_ctx(fc_cuda_ctx *ctx, fc_buffer *x, fc_buffer *out, int rows, int cols); void fc_softmax(fc_buffer *x, fc_buffer *out, int rows, int cols); /* SiLU activation: out = x * sigmoid(x) */ void fc_silu_ctx(fc_cuda_ctx *ctx, fc_buffer *x, fc_buffer *out, int n); void fc_silu(fc_buffer *x, fc_buffer *out, int n); /* Elementwise addition: out = a + b */ void fc_add_ctx(fc_cuda_ctx *ctx, fc_buffer *a, fc_buffer *b, fc_buffer *out, int n); void fc_add(fc_buffer *a, fc_buffer *b, fc_buffer *out, int n); /* Elementwise multiplication: out = a * b */ void fc_mul_ctx(fc_cuda_ctx *ctx, fc_buffer *a, fc_buffer *b, fc_buffer *out, int n); void fc_mul(fc_buffer *a, fc_buffer *b, fc_buffer *out, int n); /* INT8 quantized matrix-vector product. * x_q: [cols] int8, w: [rows, cols] int8, w_scales: [rows] float32, * out: [rows] float32. * Computes: out[i] = x_scale * w_scales[i] * dot(x_q, w[i]) */ void fc_matvec_q8_ctx(fc_cuda_ctx *ctx, fc_buffer *x_q, fc_buffer *w, fc_buffer *w_scales, fc_buffer *out, float x_scale, int rows, int cols); void fc_matvec_q8(fc_buffer *x_q, fc_buffer *w, fc_buffer *w_scales, fc_buffer *out, float x_scale, int rows, int cols); /* ================================================================ * qwen35moe kernels (loaded from embedded PTX) * * The CUDA halves of the qwen35moe forward pass. Each mirrors a CPU * reference (cited in its cuda_kernels .cu source) bit-for-bit and is * validated against the golden vectors in Pantheon/tests/golden/. * ================================================================ */ /* Q3_K dequant + matrix-vector product. * w: raw Q3_K weight bytes (row-major, row_stride bytes/row), * x: [cols] float32 activation, out: [rows] float32. * cols must be a multiple of 256; row_stride = (cols/256) * 110. * Computes: out[r] = dequant_q3k(w[r]) · x. */ void fc_matvec_q3k_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride); void fc_matvec_q3k(fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride); /* ---- Q3_K peer-stream variants (scheme 9a) ---- * BIT-EXACT alternatives to fc_matvec_q3k that differ only in how many weight * loads a warp keeps in flight. They exist for the case where the die running * the kernel reads ANOTHER die's VRAM over the K80's on-card peer link, which * is LATENCY-bound (9.3-12.4 GB/s available, 2.7-4.2 GB/s achieved) — see * plans/results/phase9a_peer_streaming.md. On a die's OWN memory the kernel is * issue-bound and the stock shape is the right one. * r2 = row pair per warp (the stock shape), r1 = one row per warp (2x warps) * pN = N super-blocks' weight loads issued before the first is consumed * All variants also stage x with 128-bit loads. */ enum { FC_Q3K_PS_R2P1 = 0, FC_Q3K_PS_R2P2, FC_Q3K_PS_R2P4, FC_Q3K_PS_R2P8, FC_Q3K_PS_R1P1, FC_Q3K_PS_R1P4, FC_Q3K_PS_R1P8, FC_Q3K_PS_COUNT }; /* Human-readable name of a variant, or NULL if `variant` is out of range. */ const char *fc_matvec_q3k_ps_name(int variant); /* Launch peer-stream variant `variant`. Same contract and same output bytes as * fc_matvec_q3k_ctx. Returns 0 if launched, -1 if the variant is out of range * or absent from this build's PTX — the caller must handle -1 rather than let * a missing kernel silently launch nothing. */ int fc_matvec_q3k_ps_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride, int variant); /* Mixed-quant dequant + matvec family (matvec_kquants.cu). Same contract as * fc_matvec_q3k_ctx for each GGUF quant type the Q3_K_M mini mixes in: * Q4_K (144 B / 256 elems), Q5_K (176 B / 256), Q6_K (210 B / 256), * Q8_0 (34 B / 32), and plain F32 rows (row_stride = cols * 4). */ void fc_matvec_q4k_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride); void fc_matvec_q5k_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride); void fc_matvec_q6k_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride); void fc_matvec_q80_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride); /* Bonsai native ternary (GGUF type 42, Q2_0_g128): conditional add/sub, no * dequant. Repack via fc_repack_weight(42, ...). Warp-per-row like q80. */ void fc_matvec_q2_0_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride); void fc_matvec_f32_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride); /* F16 rows (GGUF type 1): the dense half of an F16 model — attention q/k/v/o, * shared expert, router, LM head. 2-byte load widened in-register (sm_37 has no * fp16 ARITHMETIC; cvt.f32.f16 is a storage conversion and predates it), fp32 * accumulate. row_stride is bytes (cols * 2) and the weight is NOT repacked — * the GGUF bytes are already the device layout, which is what keeps a 70 GB * upload from needing a second host copy of every tensor. * These are OPTIONAL PTX symbols: probe fc_matvec_f16_available_ctx before * routing a model's dense weights here. With the kernel absent the launcher * does NOTHING, so a caller that skipped the probe would write no output at * all instead of falling back. */ int fc_matvec_f16_available_ctx(fc_cuda_ctx *ctx); void fc_matvec_f16_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride); /* BF16 rows (GGUF type 30): the same dense tensors as the F16 pair above, * carrying the values the model actually SHIPS — the checkpoints are bf16 * safetensors, and converting bf16 -> F16 re-rounds every weight (7 stored * mantissa bits against 10, and neither field contains the other), so an F16 * run measures a perturbed model. That fidelity is the only reason to prefer * it: the widen is a shift into the high half rather than a cvt, but ptxas * emits ONE instruction either way and the two decode at the same speed * (measured; see fc_bf16_to_f32 in moe_expert_group.cu). Identical * contract to the F16 launcher otherwise: row_stride in bytes (cols * 2), NOT * repacked, OPTIONAL PTX symbols, silent when absent — probe * fc_matvec_bf16_available_ctx before routing weights here. */ int fc_matvec_bf16_available_ctx(fc_cuda_ctx *ctx); void fc_matvec_bf16_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride); /* ---- Batched (multi-column) matvecs: the prefill/batch GEMM path ---- * out[t][rows] = dequant(w) . x[t] for the n activation columns of a batch: * x is [n][cols] row-major (token-major), out is [n][rows]. Each weight * block is decoded ONCE per 4-column tile and reused across the columns, so * the dequant instruction stream (the single-token kernels' issue bound on * sm_37) amortizes over the batch. BIT-EXACT per column: column t of the * batched output is bit-identical to the corresponding single-column * fc_matvec_* call on x[t] (same per-lane chain order + reduction; verified * by cuda/tests/test_matvec_nx). Same layout/stride contracts as the * single-column variants. */ void fc_matvec_q3k_nx_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride, int n); void fc_matvec_q4k_nx_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride, int n); void fc_matvec_q5k_nx_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride, int n); void fc_matvec_q6k_nx_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride, int n); void fc_matvec_q80_nx_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride, int n); void fc_matvec_f32_nx_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride, int n); void fc_matvec_f16_nx_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride, int n); void fc_matvec_bf16_nx_ctx(fc_cuda_ctx *ctx, fc_buffer *w, fc_buffer *x, fc_buffer *out, int rows, int cols, int row_stride, int n); /* ================================================================ * Quant-block repack (host-side, done once at weight upload) * * The raw GGUF block layouts are unaligned (110/144/176/210/34 B) * and interleave packed scales with the quant bits, which caps the * warp-per-row matvec kernels at 10-40 GB/s (per-lane byte loads + * per-element scale unpack). fc_repack_weight converts a weight's * blocks ONCE into the 4-byte-aligned device layout those kernels * read: numerically identical values, scales pre-unpacked, block * scale d (and dmin) pre-decoded to f32, and the quant bytes * transposed so lane l of a warp issues ONE aligned 32-bit load per * super-block. Per-super-block layouts (offsets in bytes): * * Q3_K 110->148: u32 qw[32] @0 qw[l] = qs[l] | qs[32+l]<<8 | hmask[l]<<16 * s8 sc[16] @128 (6-bit scales, unpacked, -32 applied) * f32 d @144 * Q4_K 144->152: u32 qw[32] @0 qw[l] byte j = qs[32*j + l] * f32 d @128, f32 dmin @132, u8 sc[8] @136, u8 mn[8] @144 * Q5_K 176->184: u32 qw[32] @0 (as Q4_K), u8 qh[32] @128, * f32 d @160, f32 dmin @164, u8 sc[8] @168, u8 mn[8] @176 * Q6_K 210->212: u32 qlw[32] @0 qlw[l] byte t = ql[32*t + l] * u16 qh2[32] @128 qh2[l] = qh[l] | qh[32+l]<<8 * s8 sc[16] @192, f32 d @208 * Q8_0 34-> 36: f32 d @0, s8 qs[32] @4 (per 32-elem block) * * fc_repack_row_bytes returns the repacked bytes-per-row for a GGUF * quant type id (8=Q8_0, 11=Q3_K, 12=Q4_K, 13=Q5_K, 14=Q6_K), or 0 * if the type/cols combination has no repacked form (caller must not * use the quant matvec kernels then). * ================================================================ */ size_t fc_repack_row_bytes(int gguf_quant, int cols); int fc_repack_weight(int gguf_quant, const void *src, void *dst, long rows, int cols); /* Partial NEOX rotary position embedding, in place. * data: [n_heads * head_dim] float32. Only the first rope_dim dims of * each head are rotated (NEOX pairing j vs j+rope_dim/2); dims * [rope_dim, head_dim) pass through unchanged. */ void fc_rope_neox_partial_ctx(fc_cuda_ctx *ctx, fc_buffer *data, int n_heads, int head_dim, int rope_dim, int pos, float theta_base); void fc_rope_neox_partial(fc_buffer *data, int n_heads, int head_dim, int rope_dim, int pos, float theta_base); /* Per-head QK-RMSNorm, in place. * data: [n_heads * head_dim] float32; each head is RMS-normalized over * its own head_dim and scaled by the shared weight [head_dim]. * weight may be NULL (no scaling). */ void fc_qk_rms_norm_ctx(fc_cuda_ctx *ctx, fc_buffer *data, fc_buffer *weight, int n_heads, int head_dim, float eps); void fc_qk_rms_norm(fc_buffer *data, fc_buffer *weight, int n_heads, int head_dim, float eps); /* MoE router: softmax over n_expert logits (in place), top-k select, * norm_topk_prob renormalize, apply expert_weights_scale. * logits: [n_expert] float32 (overwritten with softmax probs), * top_idx: [k] int32 output, top_wt: [k] float32 output. * k must be <= 32 (FC_MOE_MAX_TOPK). */ void fc_moe_router_topk_ctx(fc_cuda_ctx *ctx, fc_buffer *logits, fc_buffer *top_idx, fc_buffer *top_wt, int n_expert, int k, float expert_weights_scale); void fc_moe_router_topk(fc_buffer *logits, fc_buffer *top_idx, fc_buffer *top_wt, int n_expert, int k, float expert_weights_scale); /* Device-side exact top-k over LM-head logits (topk_logits.cu). * logits: [n] float32 device buffer (left intact on return). * out_idx/out_val: HOST arrays of k entries, filled in descending value * order; equal values order by ascending index (the host argmax tie-break, * so the k=1 result IS the host argmax). Values are bit-exact — selection * is compare-only — and only k (idx, val) pairs cross the host link * instead of the full n-float buffer. k must be <= FC_TOPK_MAX. * Synchronous; returns 0 on success, -1 on failure (caller should fall * back to a full-logits D2H copy). */ #define FC_TOPK_MAX 256 int fc_topk_logits_ctx(fc_cuda_ctx *ctx, fc_buffer *logits, int n, int k, int *out_idx, float *out_val); int fc_topk_logits(fc_buffer *logits, int n, int k, int *out_idx, float *out_val); /* Batched (multi-slot) top-k (L8): the same exact selection over every row * of a [nslots][n] logits matrix IN PLACE (no per-row staging copy), one * sync + one nslots*k-pair D2H for the whole batch. out_idx/out_val are * HOST arrays of nslots*k entries (row u's pairs at [u*k .. u*k+k)); each * row's pairs are bit-exact to fc_topk_logits_ctx on that row. Returns 0 * on success, -1 on failure / kernels absent (caller falls back to the * per-slot path). */ int fc_topk_logits_slots_ctx(fc_cuda_ctx *ctx, fc_buffer *logits, int nslots, int n, int k, int *out_idx, float *out_val); /* MoE weighted accumulate: moe_out[i] += weight * expert_out[i]. * Both buffers are [n_embd] float32. */ void fc_moe_weighted_add_ctx(fc_cuda_ctx *ctx, fc_buffer *expert_out, fc_buffer *moe_out, float weight, int n_embd); void fc_moe_weighted_add(fc_buffer *expert_out, fc_buffer *moe_out, float weight, int n_embd); /* MoE shared-expert sigmoid gate: shexp_out[i] *= sigmoid(gate_logit). * shexp_out: [n_embd] float32, gated in place. */ void fc_moe_shared_gate_ctx(fc_cuda_ctx *ctx, fc_buffer *shexp_out, float gate_logit, int n_embd); void fc_moe_shared_gate(fc_buffer *shexp_out, float gate_logit, int n_embd); /* ---- Grouped top-k expert kernels (moe_expert_group.cu) ---- * The whole routed-expert SwiGLU block in two launches, reading the router * selection (top_idx/top_wt) directly on device — no host read-back. */ /* Are the 16-bit expert kernels (F16, BF16) routed to their uint2-load form? * Set by FC_MOE_VEC_LOADS=1 in the environment and read per call. An * unsupported combination FAILS CLOSED (no launch) rather than quietly running * the scalar kernel, so a bench cannot end up timing the wrong thing. * * The vec kernels move the SAME DRAM bytes in HALF the LSU transactions (a * warp-wide 16-bit load fills 64 B of a 128-B transaction; a uint2 load fills * it), and on die 7 that is worth 1.31x on gate+up (0.2418 -> 0.1849 ms, 139 -> * 181 GB/s) and 1.32x on down (0.1304 -> 0.0986 ms, 129 -> 170 GB/s), the same * for both formats. 181 GB/s is 76% of the die's 240 GB/s theoretical, against * 58% for the scalar form. * * It is opt-in and not the default because it is not a pure win: to load four * contiguous weights per lane the kernels must hand lane l the run {4l..4l+3} * instead of the scalar map's stride-32 set {32i+l}. That repartitions the * products across lanes, so the per-lane fmaf chains and the shfl tree * associate differently and the output is NOT bit-identical to the scalar * kernels'. Each set is bit-exact against its own CPU replica and neither is * more accurate; but a run that has to reproduce an earlier run's exact logits * must pin this setting, so flipping the default is a deliberate call, not a * tuning detail. Exposed so a test can pick the matching replica from the same * source of truth the dispatch uses. */ int fc_moe_vec_loads_enabled(void); /* Grouped gate+up SwiGLU for the k selected experts. * gate_all/up_all: merged [n_expert * ff * row_stride] repacked quant bytes, * x: [cols] f32, top_idx: [k] int32 (device), hid: [k*ff] f32 out. * hid[s*ff+row] = silu(gate[e_s,row] . x) * (up[e_s,row] . x). * * gguf_quant selects the decoder (11 = Q3_K, 12 = Q4_K, 1 = F16, 30 = BF16), * exactly as it does * for fc_moe_down_accum_ctx below. FAILS CLOSED: an unsupported quant, or a * supported one whose kernel is missing from the loaded PTX, logs a named * error and launches NOTHING (hid untouched) rather than falling through to * another quant's decode. Returns 0 on launch, -1 otherwise. * * Probe fc_moe_gateup_quant_available_ctx BEFORE committing a model to the * grouped path: the Q4_K kernels are OPTIONAL symbols, so an fc_cuda_ptx.h * regenerated before they existed answers 0 for Q4_K and the caller can take * the per-expert fallback instead of failing on every layer at run time. */ int fc_moe_gateup_quant_available_ctx(fc_cuda_ctx *ctx, int gguf_quant); int fc_moe_gateup_ctx(fc_cuda_ctx *ctx, int gguf_quant, fc_buffer *gate_all, fc_buffer *up_all, fc_buffer *x, fc_buffer *top_idx, fc_buffer *hid, int k, int ff, int cols, int row_stride); /* Q3_K-pinned shorthand for fc_moe_gateup_ctx (benches, older callers). */ void fc_moe_gateup_q3k_ctx(fc_cuda_ctx *ctx, fc_buffer *gate_all, fc_buffer *up_all, fc_buffer *x, fc_buffer *top_idx, fc_buffer *hid, int k, int ff, int cols, int row_stride); /* Grouped down-projection + weighted accumulate: * out[row] = sum_s top_wt[s] * (down[e_s,row] . hid[s]). * gguf_quant selects the decoder (11 = Q3_K, 12 = Q4_K, 14 = Q6_K, 1 = F16, * 30 = BF16), and a * merged ffn_down_exps really is mixed WITHIN one file — Q3_K/Q4_K by layer in * a Q3_K_M, Q4_K/Q6_K by layer in a Q4_K_M. Unsupported quants log a named * error and launch nothing. out is WRITTEN. * * Q6_K is an OPTIONAL PTX symbol (a header regenerated before it existed leaves * it NULL); probe fc_moe_down_quant_available_ctx before committing a model to * the grouped path. */ int fc_moe_down_quant_available_ctx(fc_cuda_ctx *ctx, int gguf_quant); void fc_moe_down_accum_ctx(fc_cuda_ctx *ctx, int gguf_quant, fc_buffer *down_all, fc_buffer *hid, fc_buffer *top_idx, fc_buffer *top_wt, fc_buffer *out, int k, int rows, int ff, int row_stride); /* Warps cooperating per down-projection output row pair (the wpp launch arg * of fc_moe_down_q3k/q4k; FC_MOE_DOWN_WPP env-overridable for tuning, block * size via FC_MOE_DOWN_THREADS). The 512-col down dot is 2 super-blocks, so * one warp per row-pair stalls on latency; wpp warps split the k slots and * shared-mem-reduce. Must divide the block's warp count. wpp=2 with * 128-thread blocks measured best on a K80 (the kernel is * instruction-issue bound at ~4.3 IPC — see moe_expert_group.cu). Mirrored * by the bench_matvec CPU replica. */ #define FC_MOE_DOWN_WPP_DEFAULT 2 #define FC_MOE_DOWN_THREADS_DEFAULT 128 /* MoE expert-union dedup request-tile depth. Must match FC_MOE_DEDUP_RT in * cuda_kernels/moe_expert_group.cu: it sets both the per-block register * footprint and grid.z = ceil(maxr / RT). An expert with R requests has its * rows read ceil(R/RT) times, so RT >= the typical R is what makes the dedup * pay. */ #define FC_MOE_DEDUP_RT 8 /* Fused SwiGLU elementwise: out[i] = silu(a[i]) * b[i] (in-place safe). */ void fc_moe_silu_mul_ctx(fc_cuda_ctx *ctx, fc_buffer *a, fc_buffer *b, fc_buffer *out, int n); /* Shared-expert scalar gate fully on-device: * shexp[i] *= sigmoid(gate_w . x). gate_w, x: [n] f32. */ void fc_moe_shexp_dot_gate_ctx(fc_cuda_ctx *ctx, fc_buffer *gate_w, fc_buffer *x, fc_buffer *shexp, int n); /* ================================================================ * Fused per-layer launches (fused_layer.cu) * * Decode on the K80 is launch-bound (~80 us/launch), so fc_fused_multi * batches up to FC_FUSED_MAX_SEG independent matvec "segments" that read * the SAME x into ONE launch, each with an optional fused epilogue. Every * segment's math is bit-identical to its stand-alone kernel (the kernel * bodies are duplicated verbatim; see fused_layer.cu). * ================================================================ */ #define FC_FUSED_MAX_SEG 8 /* One matvec segment. kind: 0 = quant warp-per-row (qtype 8/11/12/13/14, * repacked layout), 1 = F32 block-per-row (cols >= 256), 2 = quant gate+up * pair (w = gate, w2 = up; writes silu(g.x)*(u.x)). ep (epilogue applied to * the finished row value): 0 none; 1 sigmoid; 2 decay gate out[r] = * aux1[r]*softplus(v+aux0[r]); 3 residual out[r] = v, aux1[r] += v; * 4 moe combine out[r] = aux0[r] + v*aux2[0], aux1[r] += out[r]. */ typedef struct { fc_buffer *w; /* weight bytes (repacked quant, or raw F32 rows) */ fc_buffer *w2; /* kind 2 only: the up weight (same type/stride) */ fc_buffer *out; /* [rows] f32 */ fc_buffer *aux0; /* epilogue operand (see ep) or NULL */ fc_buffer *aux1; fc_buffer *aux2; int rows, cols, row_stride; int kind, qtype, ep; } fc_fused_seg; /* Validate the segments and upload their device descriptor ONCE (all * referenced buffers must be persistent). Returns the descriptor buffer * (free with fc_buffer_free) and the launch shape, or NULL if any segment * is unsupported (caller keeps its unfused path). */ fc_buffer *fc_fused_build_ctx(fc_cuda_ctx *ctx, const fc_fused_seg *segs, int nseg, int *total_blocks, int *shared_bytes); /* Launch the fused multi-matvec: every segment's dot(s) against x in one * kernel. desc/total_blocks/shared_bytes come from fc_fused_build_ctx. */ void fc_fused_multi_ctx(fc_cuda_ctx *ctx, fc_buffer *desc, int nseg, int total_blocks, int shared_bytes, fc_buffer *x); /* Fused per-head QK-RMSNorm + partial NEOX RoPE on Q and K in one launch * (replaces fc_qk_rms_norm x2 + fc_rope_neox_partial x2; per-head math * identical). Returns 0 on launch, -1 if the kernel is unavailable (caller * must fall back to the four stand-alone launches). */ int fc_attn_qknorm_rope_ctx(fc_cuda_ctx *ctx, fc_buffer *q, fc_buffer *k, fc_buffer *q_norm_w, fc_buffer *k_norm_w, int n_head, int n_head_kv, int head_dim, float eps, int rope_dim, int pos, float theta_base); /* ================================================================ * SSM Gated-Delta-Net kernels (loaded from embedded PTX) * * The GPU halves of the single-token Gated-Delta-Net SSM block. Each * mirrors a static helper in Pantheon/src/models/qwen35moe_ssm.c * (cited in ssm_gated_delta.cu) and is validated against the SSM * golden vectors in Pantheon/tests/golden/. * ================================================================ */ /* Per-token delta-rule recurrence (the centerpiece). For one token, per * (head h, row j): decays the transposed state row, computes the rank-1 * delta update in place, and reads y back after the write. * s: [n_v_heads * S * S] transposed state (s[h*S*S+j*S+i]==S_h[i][j]), * updated in place; q,k,v: [n_v_heads * S] (q/k L2-normed + GQA- * repeated on host, v straight); g: [n_v_heads] PRE-exp decay; * beta: [n_v_heads] sigmoid'd; y: [n_v_heads * S] output. * scale = 1/sqrt(state_size). Mirrors qwen35moe_compute_recur_heads. */ void fc_ssm_recurrence_ctx(fc_cuda_ctx *ctx, fc_buffer *s, fc_buffer *q, fc_buffer *k, fc_buffer *v, fc_buffer *g, fc_buffer *beta, fc_buffer *y, int n_v_heads, int state_size, float scale); void fc_ssm_recurrence(fc_buffer *s, fc_buffer *q, fc_buffer *k, fc_buffer *v, fc_buffer *g, fc_buffer *beta, fc_buffer *y, int n_v_heads, int state_size, float scale); /* Causal depthwise conv1d (kernel size conv_kernel) + SiLU over the full * conv_dim stream, priming from conv_state and rolling it forward. * qkv: [conv_dim] raw input, weight: [conv_dim * conv_kernel] row-major, * conv_state: [conv_dim * (conv_kernel-1)] history (updated in place), * out: [conv_dim] SiLU'd output. Mirrors qwen35moe_conv1d_silu. */ void fc_ssm_conv1d_silu_ctx(fc_cuda_ctx *ctx, fc_buffer *qkv, fc_buffer *weight, fc_buffer *conv_state, fc_buffer *out, int conv_dim, int conv_kernel); void fc_ssm_conv1d_silu(fc_buffer *qkv, fc_buffer *weight, fc_buffer *conv_state, fc_buffer *out, int conv_dim, int conv_kernel); /* Gated RMSNorm: gated[h*S+d] = RMSNorm(y_h)[d] * weight[d] * silu(z[h*S+d]), * per V-head over state_size dims (norm weight shared across heads). Omits the * trailing ssm_out projection (reuse fc_matvec_q3k). y,z,gated: [n_v_heads*S], * weight: [state_size]. Mirrors qwen35moe_gated_norm. */ void fc_ssm_gated_norm_ctx(fc_cuda_ctx *ctx, fc_buffer *y, fc_buffer *weight, fc_buffer *z, fc_buffer *gated, int n_v_heads, int state_size, float eps); void fc_ssm_gated_norm(fc_buffer *y, fc_buffer *weight, fc_buffer *z, fc_buffer *gated, int n_v_heads, int state_size, float eps); /* SSM "glue" kernels — replace host round-trips in the qwen35moe GPU forward. */ /* beta = sigmoid(beta), in place [n]. */ void fc_sigmoid_ctx(fc_cuda_ctx *ctx, fc_buffer *x, int n); /* g[h] = ssm_a[h] * softplus(alpha[h] + ssm_dt[h]), per V-head [n]. */ void fc_ssm_decay_gate_ctx(fc_cuda_ctx *ctx, fc_buffer *alpha, fc_buffer *ssm_dt, fc_buffer *ssm_a, fc_buffer *g, int n); /* Per-K-head L2-norm of q,k over state_size dims + GQA interleaved-repeat to * qg/kg [n_v_heads*S]; v (conv_out[2*key_dim..], val_dim) straight-copied to vg. * conv_out layout: [q key_dim | k key_dim | v val_dim]. */ void fc_ssm_qknorm_gqa_ctx(fc_cuda_ctx *ctx, fc_buffer *conv_out, fc_buffer *qg, fc_buffer *kg, fc_buffer *vg, int n_k_heads, int kv_groups, int state_size, int key_dim, int val_dim, float eps); /* ================================================================ * Batched (prefill) kernels — n positions of ONE sequence per launch * * The launch-count killers for batched prefill: each covers all n * positions of a batch in one launch, with per-position math BIT-EXACT * to its single-position counterpart (the stateful SSM kernels scan the * positions inside the kernel, carrying state exactly as n sequential * launches would). Batched activations are token-major: a [n][dim] * buffer holds position t's vector at offset t*dim. * ================================================================ */ /* Batched RMS norm: x/out [rows][n], one block per row. */ void fc_rms_norm_nx_ctx(fc_cuda_ctx *ctx, fc_buffer *x, fc_buffer *weight, fc_buffer *out, int n, float eps, int rows); /* Batched partial NEOX RoPE: data [npos][n_heads*head_dim], row t rotated * at position pos0 + t, in place. */ void fc_rope_neox_partial_nx_ctx(fc_cuda_ctx *ctx, fc_buffer *data, int n_heads, int head_dim, int rope_dim, int pos0, float theta_base, int npos); /* Batched MoE router: logits [ntok][n_expert] (softmax'd in place), * top_idx [ntok][k], top_wt [ntok][k]. */ void fc_moe_router_topk_nx_ctx(fc_cuda_ctx *ctx, fc_buffer *logits, fc_buffer *top_idx, fc_buffer *top_wt, int n_expert, int k, float expert_weights_scale, int ntok); /* Batched grouped gate+up SwiGLU: x [ntok][cols], top_idx [ntok][k], * hid [ntok][k*ff]. gguf_quant and the fail-closed contract are exactly * fc_moe_gateup_ctx's; per token the result is bit-exact with it. */ int fc_moe_gateup_nx_ctx(fc_cuda_ctx *ctx, int gguf_quant, fc_buffer *gate_all, fc_buffer *up_all, fc_buffer *x, fc_buffer *top_idx, fc_buffer *hid, int k, int ff, int cols, int row_stride, int ntok); /* Q3_K-pinned shorthand for fc_moe_gateup_nx_ctx. */ void fc_moe_gateup_q3k_nx_ctx(fc_cuda_ctx *ctx, fc_buffer *gate_all, fc_buffer *up_all, fc_buffer *x, fc_buffer *top_idx, fc_buffer *hid, int k, int ff, int cols, int row_stride, int ntok); /* Expert-union dedup for the batched gate+up. Q3_K GATE/UP ONLY — there is no * Q4_K twin of these kernels (dedup is an opt-in, prefill-only speed lever, not * a format the engine must support), so a caller with Q4_K experts must keep to * fc_moe_gateup_nx_ctx. fc_moe_dedup_build_ctx * turns top_idx [ntok*k] into per-expert request lists (ecount [n_expert], * ereq [n_expert*maxr]); fc_moe_gateup_q3k_dedup_ctx then runs the grouped * gate+up over the union, reading each unique expert's weights once and * applying them to every routing token. hid is bit-exact with the _nx path, * so the existing fc_moe_down_accum_nx_ctx consumes it unchanged. maxr must * be >= the max per-expert request count (== ntok). * * Both dedup calls are no-ops when their kernels are missing (an * fc_cuda_ptx.h predating them), which would leave hid UNWRITTEN rather than * wrong-but-plausible. Callers must gate on fc_moe_dedup_available_ctx and * fall back to fc_moe_gateup_q3k_nx_ctx. * * FC_MOE_DEDUP_SHX=1 (default off) swaps in fc_moe_gateup_q3k_dedup_shx, * which stages x in shared like the _nx kernel and keeps the block-uniform * request list in shared instead of 16 per-thread registers: 62 registers and * 8 blocks/SMX against 95 and 5, and bit-exact with both other paths. * Measured +20% on the kernel at ntok=128 / n_expert=256. */ int fc_moe_dedup_available_ctx(fc_cuda_ctx *ctx); void fc_moe_dedup_build_ctx(fc_cuda_ctx *ctx, fc_buffer *top_idx, fc_buffer *ecount, fc_buffer *ereq, int n_expert, int k, int ntok, int maxr); void fc_moe_gateup_q3k_dedup_ctx(fc_cuda_ctx *ctx, fc_buffer *gate_all, fc_buffer *up_all, fc_buffer *x, fc_buffer *ecount, fc_buffer *ereq, fc_buffer *hid, int k, int ff, int cols, int row_stride, int n_expert, int maxr); /* ---- COMPACTED-GRID dedup gate+up (FC_MOE_DEDUP_COMPACT, default OFF) ---- * * The uncompacted launch above is (ceil(ff/8), n_expert, ceil(maxr/RT)) = * (64, 256, 16) = 262,144 blocks at the shipped prefill shape, of which only * 6.3 % do any work; MEASURED, the empty ones cost 18.6 % of the kernel * (results/moe_dedup_occupancy_probe.md §5.2). The compacted path folds * (expert, request-tile) onto ONE grid axis over a device-built work list and * launches grid.y at fc_moe_dedup_wmax's exact upper bound — 352 instead of * 4096, i.e. 22,528 blocks — with NO device->host sync. Bit-exact: same work * items, same kernel body, one plain store per output, so block order is not * observable (the full argument is in moe_expert_group.cu). * * Sequence per MoE block, replacing the two calls above: * fc_moe_dedup_build_ctx (UNCHANGED — still builds ecount/ereq) * fc_moe_dedup_build_wl_ctx (compacts ecount into wlist) * fc_moe_gateup_q3k_dedup_compact_ctx * * Both new kernels exist only in a * cd cuda_kernels && make NVCCFLAGS_EXTRA=-DFC_MOE_DEDUP_COMPACT=1 * build of the PTX, so with the shipped header fc_moe_dedup_compact_enabled_ctx * is 0 whatever the environment says and callers keep the uncompacted launch. */ /* Exact host-side upper bound on the number of (expert, request-tile) work * items — the grid.y a compacted launch must cover, and the int2 element count * the wlist buffer must hold. W = sum over routed experts of * ceil(R_e / FC_MOE_DEDUP_RT), and each of these three bounds is individually * exact, with T = ntok*k requests and U = min(n_expert, T) routed experts: * 1. W <= (T + (RT-1)*U) / RT — ATTAINED: at T=1024, U=256 a router giving * 160 experts R=1 and 96 experts R=9 produces 160 + 192 = 352 tiles, so * this bound cannot be tightened without knowing the router. * 2. W <= U * ceil(maxr/RT) — an expert's list is at most maxr long * 3. W <= T — every work item holds at least one request * At the shipped prefill shape (ntok=128, k=8, n_expert=256, maxr=128) that is * min(352, 4096, 1024) = 352 against the uncompacted grid.y*grid.z of 4096. * Returns 0 when the shape cannot be compacted (n_expert or maxr outside the * work-item encoding / the builder's scan capacity), which callers must treat as * "no compaction available". Pure function; no ctx, no device. */ int fc_moe_dedup_wmax(int n_expert, int k, int ntok, int maxr); /* Both compaction kernels present in the loaded PTX? */ int fc_moe_dedup_compact_available_ctx(fc_cuda_ctx *ctx); /* ...and FC_MOE_DEDUP_COMPACT=1 in the environment. Default OFF. */ int fc_moe_dedup_compact_enabled_ctx(fc_cuda_ctx *ctx); void fc_moe_dedup_build_wl_ctx(fc_cuda_ctx *ctx, fc_buffer *ecount, fc_buffer *wlist, int n_expert, int wmax); void fc_moe_gateup_q3k_dedup_compact_ctx(fc_cuda_ctx *ctx, fc_buffer *gate_all, fc_buffer *up_all, fc_buffer *x, fc_buffer *ereq, fc_buffer *wlist, fc_buffer *hid, int k, int ff, int cols, int row_stride, int wmax, int maxr); /* Batched grouped down + weighted accumulate: hid [ntok][k*ff], * top_idx/top_wt [ntok][k], out [ntok][rows] (WRITTEN). */ void fc_moe_down_accum_nx_ctx(fc_cuda_ctx *ctx, int gguf_quant, fc_buffer *down_all, fc_buffer *hid, fc_buffer *top_idx, fc_buffer *top_wt, fc_buffer *out, int k, int rows, int ff, int row_stride, int ntok); /* Batched shared-expert scalar gate: x/shexp [ntok][n], per token * shexp[t] *= sigmoid(gate_w . x[t]). */ void fc_moe_shexp_dot_gate_nx_ctx(fc_cuda_ctx *ctx, fc_buffer *gate_w, fc_buffer *x, fc_buffer *shexp, int n, int ntok); /* SSM conv1d + SiLU SCAN over npos positions: qkv/out [npos][conv_dim], * conv_state rolled forward through all npos (updated in place). Requires * conv_kernel <= 4. snaps (optional, NULL = off): per-position conv_state * snapshots [npos][conv_dim*(conv_kernel-1)] — snaps[t] = state after * position t, for speculative-decode partial-accept rollback. */ void fc_ssm_conv1d_silu_seq_ctx(fc_cuda_ctx *ctx, fc_buffer *qkv, fc_buffer *weight, fc_buffer *conv_state, fc_buffer *out, int conv_dim, int conv_kernel, int npos, fc_buffer *snaps); /* SSM delta-rule recurrence SCAN over npos positions: q/k/v/y * [npos][n_v_heads*S], g/beta [npos][n_v_heads], state s updated in place * through all npos. snaps (optional, NULL = off): per-position state * snapshots [npos][n_v_heads*S*S] — snaps[t] = state after position t, for * speculative-decode partial-accept rollback. */ void fc_ssm_recurrence_seq_ctx(fc_cuda_ctx *ctx, fc_buffer *s, fc_buffer *q, fc_buffer *k, fc_buffer *v, fc_buffer *g, fc_buffer *beta, fc_buffer *y, int n_v_heads, int state_size, int npos, float scale, fc_buffer *snaps); /* Batched decay gate: alpha/g [npos][n_heads], dt/a broadcast [n_heads]. */ void fc_ssm_decay_gate_nx_ctx(fc_cuda_ctx *ctx, fc_buffer *alpha, fc_buffer *ssm_dt, fc_buffer *ssm_a, fc_buffer *g, int n_heads, int npos); /* Batched q/k L2-norm + GQA repeat (+ v passthrough): conv_out * [npos][conv_dim], qg/kg/vg [npos][val_dim]. */ void fc_ssm_qknorm_gqa_nx_ctx(fc_cuda_ctx *ctx, fc_buffer *conv_out, fc_buffer *qg, fc_buffer *kg, fc_buffer *vg, int n_k_heads, int kv_groups, int state_size, int key_dim, int val_dim, float eps, int npos); /* ================================================================ * On-die causal GQA flash attention (attention.cu) * * Replaces the host-side O(N^2) causal softmax of the qwen35moe * batched prefill: the KV cache lives on the die ([n_head_kv][max_seq] * [head_dim], the host cache's exact layout) and one launch runs the * tiled QK^T -> online-softmax -> PV pipeline for all n query rows. * Online softmax reduces in tile order (running max + running sum), so * outputs drift ~1e-6 relative vs the host single-pass softmax — same * math, different reduction order, NOT bit-exact by design. * ================================================================ */ /* 1 if this ctx has the full flash-attention kernel set loaded (callers * pick the host-softmax path otherwise). */ int fc_attn_flash_available_ctx(fc_cuda_ctx *ctx); /* Split the joint Q+gate projection: qg [n][n_head][2*head_dim] (per head * [q | z]) into contiguous q and z, each [n][n_head][head_dim]. */ void fc_attn_qz_split_ctx(fc_cuda_ctx *ctx, fc_buffer *qg, fc_buffer *q, fc_buffer *z, int n_head, int head_dim, int npos); /* Append npos roped K/V rows to the GPU ring KV cache: k_new/v_new * [npos][n_head_kv][head_dim], caches [n_head_kv][max_seq][head_dim], * row t landing at (pos0 + t) % max_seq per head. */ void fc_attn_kv_append_ctx(fc_cuda_ctx *ctx, fc_buffer *k_new, fc_buffer *v_new, fc_buffer *k_cache, fc_buffer *v_cache, int n_head_kv, int head_dim, int max_seq, int pos0, int npos); /* Causal GQA flash attention (prefill): q/out [npos][n_head][head_dim], * caches as above with rows [0, pos0+npos) live; query row t attends * cache rows [0, pos0+t]. Requires pos0+npos <= max_seq (no ring wrap), * head_dim a multiple of 16 (<= 1024) and n_head/n_head_kv <= 8. * Returns 0 on launch, -1 if unavailable/unsupported (caller falls back * to its host path). */ /* attn_sink/attn_win drive StreamingLLM sliding-window attention: when * attn_win > 0 each query attends only [0, attn_sink) U the last attn_win keys * (prefill O(N^2) -> O(N*W)); attn_win <= 0 is full attention (byte-identical * to the pre-window kernel). Orthogonal to the Q4 sink/win fp16-shadow. */ int fc_attn_flash_prefill_ctx(fc_cuda_ctx *ctx, fc_buffer *q, fc_buffer *k_cache, fc_buffer *v_cache, fc_buffer *out, int n_head, int n_head_kv, int head_dim, int max_seq, int pos0, int npos, int attn_sink, int attn_win, float scale); /* ---- TurboQuant Q4 KV cache (QWEN35_KV_Q4) ------------------------------ * The cache is stored ggml-Q4_0 (32-element blocks, one fp16 scale each, * ~7x smaller than fp32) and dequantized in-kernel by the flash path — the * lever for the model's full 256K context on one card. Byte layout matches * Pantheon's pn_block_q4_0 so host<->device copies are byte identical. Row * bytes = (head_dim/32)*18; head_dim must be a multiple of 32. */ /* 1 if this ctx has both Q4 KV kernels loaded. */ int fc_attn_flash_q4_available_ctx(fc_cuda_ctx *ctx); /* Append + quantize npos roped K/V rows into the Q4 ring cache, and mirror * each row into the fp16 sink/recent-window shadow so greedy argmax stops * flipping. k_new/v_new fp32 [npos][n_head_kv][head_dim]; caches * [n_head_kv][max_seq][row_bytes]; k/v_sink fp16 [n_head_kv][SINKMAX][head_dim]; * k/v_win fp16 [n_head_kv][WMAX][head_dim] ring. sink/win = live sizes. */ void fc_attn_kv_append_q4_ctx(fc_cuda_ctx *ctx, fc_buffer *k_new, fc_buffer *v_new, fc_buffer *k_cache, fc_buffer *v_cache, fc_buffer *k_sink, fc_buffer *v_sink, fc_buffer *k_win, fc_buffer *v_win, int n_head_kv, int head_dim, int max_seq, int pos0, int npos, int sink, int win); /* Causal GQA flash attention over the mixed Q4 + fp16-shadow cache (fp16 for * sink/recent keys, in-kernel q4 dequant for the middle). Same contract/returns * as fc_attn_flash_prefill_ctx. */ int fc_attn_flash_prefill_q4_ctx(fc_cuda_ctx *ctx, fc_buffer *q, fc_buffer *k_cache, fc_buffer *v_cache, fc_buffer *k_sink, fc_buffer *v_sink, fc_buffer *k_win, fc_buffer *v_win, fc_buffer *out, int n_head, int n_head_kv, int head_dim, int max_seq, int pos0, int npos, int sink, int win, int attn_sink, int attn_win, float scale); /* ================================================================ * Decode-mode split flash attention (QWEN35_FLASH_DECODE) * * A single decode query attends the full cached context [0, seq_len). The KV * sequence is split across n_splits blocks (grid (n_splits, n_head_kv)) so the * die is not left near-idle by the npos=1 prefill launch; each block writes a * partial {m, l, acc} to scratch and a combine pass merges them per head. The * combine is fused into the same ctx-stream launch, so ..._ctx returns once * both kernels are queued. * * Scratch sizing (caller-owned, indexed by global head qh = kvh*kv_groups+g): * m_scr, l_scr : >= n_head * n_splits floats * acc_scr : >= n_head * n_splits * head_dim floats * out: [n_head][head_dim]. Returns 0 on success, -1 on a contract violation / * launch failure. */ int fc_attn_flash_decode_available_ctx(fc_cuda_ctx *ctx); int fc_attn_flash_decode_q4_available_ctx(fc_cuda_ctx *ctx); int fc_attn_flash_decode_ctx(fc_cuda_ctx *ctx, fc_buffer *q, fc_buffer *k_cache, fc_buffer *v_cache, fc_buffer *m_scr, fc_buffer *l_scr, fc_buffer *acc_scr, fc_buffer *out, int n_head, int n_head_kv, int head_dim, int max_seq, int seq_len, int n_splits, int split_len, int attn_sink, int attn_win, float scale); int fc_attn_flash_decode_q4_ctx(fc_cuda_ctx *ctx, fc_buffer *q, fc_buffer *k_cache, fc_buffer *v_cache, fc_buffer *k_sink, fc_buffer *v_sink, fc_buffer *k_win, fc_buffer *v_win, fc_buffer *m_scr, fc_buffer *l_scr, fc_buffer *acc_scr, fc_buffer *out, int n_head, int n_head_kv, int head_dim, int max_seq, int seq_len, int n_splits, int split_len, int sink, int win, int attn_sink, int attn_win, float scale); /* ---- Device-position decode variants (CUDA-graph capture) --------------- * Same math as the host-arg append/decode calls above, but the token position * is read from `pos` (an int32 DEVICE buffer, pos[0] = current position; * seq_len = pos[0] + 1) and the split plan (n_splits/split_len) is recomputed * ON-DEVICE from kps/splits_cap with the identical formula the caller uses — * so every launch argument, buffer address and grid dim is position-invariant * and the whole decode attention block can live inside a captured CUDA graph: * the host updates pos[0] once per token (outside the graph) and replays. * Bit-exact per token to the host-arg variants. The decode launch uses a * fixed grid of splits_cap blocks; dead blocks exit immediately. Scratch must * be sized for splits_cap. Caller guarantees no ring wrap * (pos[0] + 1 <= max_seq). Kernels are OPTIONAL in the PTX blob (older * fc_cuda_ptx.h): probe with fc_attn_decode_graph_available_ctx before use. */ /* 1 if every kernel the graph-captured decode-attention block needs is * loaded (qz split, QK norm, RoPE-from-device-pos, sigmoid/mul glue, the * *_dev append + split + combine for the given cache mode q4 = 0/1). */ int fc_attn_decode_graph_available_ctx(fc_cuda_ctx *ctx, int q4); void fc_attn_kv_append_dev_ctx(fc_cuda_ctx *ctx, fc_buffer *k_new, fc_buffer *v_new, fc_buffer *k_cache, fc_buffer *v_cache, int n_head_kv, int head_dim, int max_seq, fc_buffer *pos); void fc_attn_kv_append_q4_dev_ctx(fc_cuda_ctx *ctx, fc_buffer *k_new, fc_buffer *v_new, fc_buffer *k_cache, fc_buffer *v_cache, fc_buffer *k_sink, fc_buffer *v_sink, fc_buffer *k_win, fc_buffer *v_win, int n_head_kv, int head_dim, int max_seq, fc_buffer *pos, int sink, int win); int fc_attn_flash_decode_dev_ctx(fc_cuda_ctx *ctx, fc_buffer *q, fc_buffer *k_cache, fc_buffer *v_cache, fc_buffer *m_scr, fc_buffer *l_scr, fc_buffer *acc_scr, fc_buffer *out, int n_head, int n_head_kv, int head_dim, int max_seq, fc_buffer *pos, int kps, int splits_cap, int attn_sink, int attn_win, float scale); int fc_attn_flash_decode_q4_dev_ctx(fc_cuda_ctx *ctx, fc_buffer *q, fc_buffer *k_cache, fc_buffer *v_cache, fc_buffer *k_sink, fc_buffer *v_sink, fc_buffer *k_win, fc_buffer *v_win, fc_buffer *m_scr, fc_buffer *l_scr, fc_buffer *acc_scr, fc_buffer *out, int n_head, int n_head_kv, int head_dim, int max_seq, fc_buffer *pos, int kps, int splits_cap, int sink, int win, int attn_sink, int attn_win, float scale); /* ---- Multi-user decode slot attention (L7) ------------------------------ * K users, one decode token each at its OWN position, against PER-SLOT KV * rings laid out [slot][n_head_kv][max_seq][row] (each slot slice * byte-identical to the single-user ring layout). `pos` is the int32 DEVICE * position array (the RoPE position buffer, pos[u] = slot u's position); * every user's split plan is recomputed on-device from its own position with * the single-user formula (kps/splits_cap), so each user's attention output * is bit-identical to fc_attn_flash_decode(_q4) at the same position. * launch_splits must be >= every user's plan (the caller takes the max of * the per-user host plans); partial scratch is strided by launch_splits: * m_scr, l_scr : >= nslots * n_head * launch_splits floats * acc_scr : >= nslots * n_head * launch_splits * head_dim floats * q: [nslots][n_head][head_dim]; out: [nslots][n_head][head_dim]. * Caller guarantees no ring wrap (pos[u] + 1 <= max_seq for every slot). * Kernels are OPTIONAL in the PTX blob (older fc_cuda_ptx.h): probe with the * availability calls below. */ int fc_attn_flash_decode_slots_available_ctx(fc_cuda_ctx *ctx); int fc_attn_flash_decode_slots_q4_available_ctx(fc_cuda_ctx *ctx); /* Append each slot's single roped K/V row ([nslots][n_head_kv][head_dim]) to * its own ring at pos[u] (q4: quantized in-kernel + per-slot fp16 sink / * recent-window shadow mirror, matching the host pn_kvq4_store_row bytes). */ void fc_attn_kv_append_slots_ctx(fc_cuda_ctx *ctx, fc_buffer *k_new, fc_buffer *v_new, fc_buffer *k_cache, fc_buffer *v_cache, int n_head_kv, int head_dim, int max_seq, fc_buffer *pos, int nslots); void fc_attn_kv_append_slots_q4_ctx(fc_cuda_ctx *ctx, fc_buffer *k_new, fc_buffer *v_new, fc_buffer *k_cache, fc_buffer *v_cache, fc_buffer *k_sink, fc_buffer *v_sink, fc_buffer *k_win, fc_buffer *v_win, int n_head_kv, int head_dim, int max_seq, fc_buffer *pos, int nslots, int sink, int win); int fc_attn_flash_decode_slots_ctx(fc_cuda_ctx *ctx, fc_buffer *q, fc_buffer *k_cache, fc_buffer *v_cache, fc_buffer *m_scr, fc_buffer *l_scr, fc_buffer *acc_scr, fc_buffer *out, int n_head, int n_head_kv, int head_dim, int max_seq, fc_buffer *pos, int nslots, int launch_splits, int kps, int splits_cap, int attn_sink, int attn_win, float scale); int fc_attn_flash_decode_slots_q4_ctx(fc_cuda_ctx *ctx, fc_buffer *q, fc_buffer *k_cache, fc_buffer *v_cache, fc_buffer *k_sink, fc_buffer *v_sink, fc_buffer *k_win, fc_buffer *v_win, fc_buffer *m_scr, fc_buffer *l_scr, fc_buffer *acc_scr, fc_buffer *out, int n_head, int n_head_kv, int head_dim, int max_seq, fc_buffer *pos, int nslots, int launch_splits, int kps, int splits_cap, int sink, int win, int attn_sink, int attn_win, float scale); /* ================================================================ * Multi-user decode (slot) kernels — K independent sequences, ONE * token each, per launch * * The batched-decode counterparts of the _seq scans: slot u owns its * OWN state image (slot-major state buffers) and advances exactly one * step, with per-slot math bit-exact to the single-token kernels — * batching K users changes throughput, not any user's output. * Activations are slot-major [K][dim], matching the batch layout. * ================================================================ */ /* Conv1d + SiLU, one step per slot: qkv/out [nslots][conv_dim], * conv_state [nslots][conv_dim*(conv_kernel-1)] (each slot's slice * updated in place). */ void fc_ssm_conv1d_silu_slots_ctx(fc_cuda_ctx *ctx, fc_buffer *qkv, fc_buffer *weight, fc_buffer *conv_state, fc_buffer *out, int conv_dim, int conv_kernel, int nslots); /* Delta-rule recurrence, one step per slot: s [nslots][n_v_heads*S*S] * (each slot's state updated in place), q/k/v/y [nslots][n_v_heads*S], * g/beta [nslots][n_v_heads]. */ void fc_ssm_recurrence_slots_ctx(fc_cuda_ctx *ctx, fc_buffer *s, fc_buffer *q, fc_buffer *k, fc_buffer *v, fc_buffer *g, fc_buffer *beta, fc_buffer *y, int n_v_heads, int state_size, int nslots, float scale); /* Partial NEOX RoPE with a PER-ROW position: data [nrows][n_heads*head_dim] * rotated in place, pos an int32 DEVICE buffer of nrows entries (row t at * position pos[t]) — each user decodes at its own context length. */ void fc_rope_neox_partial_pos_ctx(fc_cuda_ctx *ctx, fc_buffer *data, int n_heads, int head_dim, int rope_dim, fc_buffer *pos, float theta_base, int nrows); /* ================================================================ * CUDA graph capture / replay * * Decode on the K80 is launch-bound: each cuLaunchKernel costs far more * host time than the kernel it starts. A captured graph replays a whole * recorded launch sequence with ONE cuGraphLaunch. Bit-exact by * construction — the graph re-issues the exact kernels, grids and args * that were recorded, so the device does precisely the same work. * * Contract: between begin and end, issue ONLY stream-ordered ops on this * ctx's stream (kernel launches, async copies). Synchronous ops * (fc_sync_ctx, fc_buffer_to_cpu/from_cpu) invalidate the capture. The * recorded launches DO NOT EXECUTE during capture; on a NULL return from * fc_graph_end_ctx the caller must re-issue the sequence directly. * Replay requires every recorded buffer address and grid dim to still be * valid — capture only position-invariant sequences over persistent * buffers. * * Thread scope: capture is THREAD_LOCAL, so the contract above binds the * CAPTURING THREAD only. Other threads may allocate, free, sync and run * their own dies normally while a capture is open — and may hold their own * captures concurrently. (This was not true when capture used GLOBAL mode: * any thread's unsafe call anywhere in the process killed the capture.) * ================================================================ */ typedef struct fc_graph fc_graph; /* Begin capturing ctx's stream. Returns 0 on success, -1 if graphs are * unavailable (old driver / missing symbols) — caller keeps its direct * launch path. */ int fc_graph_begin_ctx(fc_cuda_ctx *ctx); /* End capture and instantiate the executable graph. NULL on failure (the * recorded launches never ran — re-issue them directly). */ fc_graph *fc_graph_end_ctx(fc_cuda_ctx *ctx); /* Replay a captured graph on its owning ctx's stream (async, stream- * ordered like any launch). Returns 0 on success. */ int fc_graph_launch(fc_graph *g); void fc_graph_free(fc_graph *g); /* ================================================================ * CUDA events — non-blocking completion * * fc_sync_ctx blocks the calling host thread (the driver spins or * yields) until the stream drains. The async serve scheduler instead * records an event after launching a card's step and POLLS it, so one * 2-core host keeps 8 cards' streams in flight simultaneously: the * waiter sleeps between polls instead of burning a core per card. * ================================================================ */ typedef struct fc_event fc_event; /* Create a (timing-disabled) event owned by ctx. NULL if events are * unavailable (missing driver symbols / ctx not ready) — caller falls * back to fc_sync_ctx. */ fc_event *fc_event_create_ctx(fc_cuda_ctx *ctx); /* Record ev on ctx's stream: the event completes once all work queued * on the stream BEFORE this call has finished. ev must have been * created on the same ctx. Returns 0 on success. */ int fc_event_record_ctx(fc_cuda_ctx *ctx, fc_event *ev); /* Non-blocking completion probe. Returns 0 = the recorded point has * completed, 1 = still running, -1 = error/unavailable. Never blocks. */ int fc_event_query(fc_event *ev); /* Block the calling host thread until ev's recorded point completes * (cuEventSynchronize). Narrower than fc_sync_ctx, which drains the WHOLE * stream: this waits only for work queued before the record. Returns 0 on * success, -1 if events are unavailable. * * Thread-safe: yes (enters/leaves the owning context) * Reentrant: no * Signal-safe: no */ int fc_event_sync(fc_event *ev); /* CROSS-DIE ORDERING WITHOUT THE HOST. Make `waiter`'s stream block until ev * completes (cuStreamWaitEvent). ev may have been recorded on a DIFFERENT * device's context — that is the point: every other ordering primitive here * (fc_sync_ctx, fc_event_query polling) routes a die-to-die dependency through * a host round-trip, and this one does not. The call itself returns * immediately; it enqueues a wait node, it does not block the host. * * Pair with fc_buffer_copy_peer_async to build a die-to-die producer/consumer * hand-off that costs no host synchronisation: * fc_buffer_copy_peer_async(dst_on_b, src_on_a, n, ctx_b); * fc_event_record_ctx(ctx_b, ev_b); * fc_event_wait_ctx(ctx_a, ev_b); // die A's stream waits on die B * * Returns 0 on success, -1 if cuStreamWaitEvent is unavailable (old driver) or * either handle is unusable. * * Thread-safe: yes (enters/leaves the waiter's context) * Reentrant: no * Signal-safe: no */ int fc_event_wait_ctx(fc_cuda_ctx *waiter, fc_event *ev); void fc_event_free(fc_event *ev); /* ================================================================ * Synchronization * * CUDA operations are asynchronous by default. Call fc_sync() * when CPU needs GPU results (e.g., before fc_buffer_to_cpu). * In practice, cuMemcpyDtoH already synchronizes, so fc_sync() * is mainly for explicit fence points. * ================================================================ */ /* Wait for all pending GPU operations on ctx's stream to complete. */ void fc_sync_ctx(fc_cuda_ctx *ctx); /* Wait for all pending GPU operations on the process-default context. */ void fc_sync(void); #ifdef __cplusplus } #endif #endif /* FOUNDRY_CUDA_H */