diff --git a/Makefile b/Makefile
index bb202379..49560b74 100644
--- a/Makefile
+++ b/Makefile
@@ -17,7 +17,7 @@ LIB_M := h3_metal.m h3_gpu.m h3_tokenizer.m
 LIB_OBJ := $(LIB_C:.c=.o) $(LIB_M:.m=.o)
 CLI_OBJ := main.o h3_cli.o linenoise.o
 
-.PHONY: all test parity real-parity clean
+.PHONY: all test parity real-parity clean h3-cuda cuda-spark cuda-generic cuda
 
 all: h3 libh3.a
 
@@ -196,6 +196,54 @@ real-parity: h3_real_prompt_test h3_real_dit_block_test
 tests/%.o: tests/%.c
 	$(CC) $(CFLAGS) -I. -c $< -o $@
 
+# ---------------------------------------------------------------------------
+# CUDA backend (feat/cuda). The Metal implementation (h3_gpu.m, h3_shaders.metal,
+# h3_metal.m, h3_tokenizer.m) is preserved untouched. This builds the CUDA path
+# on Linux, replacing the Metal GPU layer with h3_cuda.cu and the Foundation
+# tokenizer with h3_tokenizer.c. Usage:
+#   make cuda-spark    DGX Spark / GB10 (omits explicit -arch: fastest on GB10)
+#   make cuda-generic  any local CUDA GPU (nvcc -arch=native)
+#   make cuda CUDA_ARCH=sm_N  explicit arch
+CUDA_HOME ?= /usr/local/cuda-13.0
+NVCC ?= $(CUDA_HOME)/bin/nvcc
+CUDA_ARCH ?= sm_121
+CUDA_CFLAGS := -std=c11 -O3 -MMD -MP -Wall -Wextra -Wpedantic -Wshadow \
+	-Wno-sign-conversion -D_GNU_SOURCE -DH3_CUDA
+CUDA_LDLIBS := -L$(CUDA_HOME)/lib64 -lcudart -lcublas -lm
+
+CUDA_C_SRC := h3.c h3_host.c h3_safetensors.c h3_weights.c h3_text_encoder.c \
+	h3_dit_schedule.c h3_dit.c h3_video_vae.c h3_video_encoder.c h3_audio_vae.c \
+	h3_ffmpeg.c h3_terminal.c h3_vision_encoder.c h3_multimodal.c h3_tokenizer.c
+CUDA_OBJ := $(CUDA_C_SRC:.c=.cuda.o) h3_cuda.cuda.o
+
+CLI_CUDA_OBJ := main.cuda.o h3_cli.cuda.o linenoise.cuda.o
+
+%.cuda.o: %.c
+	$(CC) $(CUDA_CFLAGS) -I. -c $< -o $@
+
+h3_cuda.cuda.o: h3_cuda.cu h3_gpu.h h3_cuda.h
+	$(NVCC) -std=c++17 -arch=$(CUDA_ARCH) -I. -DH3_CUDA -c $< -o $@
+
+linenoise.cuda.o: linenoise.c
+	$(CC) $(CUDA_CFLAGS) -Wno-conversion -Wno-variadic-macro-arguments-omitted -I. -c $< -o $@
+
+h3-cuda: $(CLI_CUDA_OBJ) $(CUDA_OBJ)
+	$(NVCC) -o h3 $^ $(CUDA_LDLIBS)
+
+cuda-spark:
+	$(MAKE) -B h3-cuda CUDA_ARCH=sm_121 CC=gcc
+
+cuda-generic:
+	$(MAKE) -B h3-cuda CUDA_ARCH=native CC=gcc
+
+cuda:
+	@if [ -z "$(strip $(CUDA_ARCH))" ]; then \
+		echo "error: specify CUDA_ARCH, e.g. make cuda CUDA_ARCH=sm_120"; \
+		exit 2; \
+	fi
+	$(MAKE) -B h3-cuda CUDA_ARCH="$(CUDA_ARCH)" CC=gcc
+
+# ---------------------------------------------------------------------------
 # Vendored from Iris. Keep the main project strict without rewriting this small
 # terminal editor for conversion diagnostics unrelated to H3.
 linenoise.o: CFLAGS += -Wno-conversion -Wno-variadic-macro-arguments-omitted
diff --git a/h3.c b/h3.c
index d5dca259..ab4cf7ed 100644
--- a/h3.c
+++ b/h3.c
@@ -4,6 +4,9 @@
 #include "h3_dit.h"
 #include "h3_ffmpeg.h"
 #include "h3_metal.h"
+#ifdef H3_CUDA
+#include "h3_cuda.h"
+#endif
 #include "h3_multimodal.h"
 #include "h3_safetensors.h"
 #include "h3_text_encoder.h"
@@ -132,8 +135,13 @@ static int h3_key_file(h3_key *key, const char *role, const char *path) {
                              strlen(path), path);
     return h3_key_append(key, "|%s=%zu:%s:%lld:%lld:%ld", role, strlen(path),
                          path, (long long)status.st_size,
+#ifdef __APPLE__
                          (long long)status.st_mtimespec.tv_sec,
                          status.st_mtimespec.tv_nsec);
+#else
+                         (long long)status.st_mtim.tv_sec,
+                         status.st_mtim.tv_nsec);
+#endif
 }
 
 static char *h3_conditioning_key(const char *prompt, const h3_params *params,
@@ -449,6 +457,15 @@ h3_ctx *h3_load_dir(const char *model_dir) {
         h3_free(ctx);
         return NULL;
     }
+#ifdef H3_CUDA
+    char cuda_error[256];
+    if (!h3_cuda_probe(&ctx->device, cuda_error, sizeof(cuda_error))) {
+        h3_set_error(ctx, "%s", cuda_error);
+        snprintf(h3_global_error, sizeof(h3_global_error), "%s", ctx->error);
+        h3_free(ctx);
+        return NULL;
+    }
+#else
     char metal_error[256];
     if (!h3_metal_probe(&ctx->device, metal_error, sizeof(metal_error))) {
         h3_set_error(ctx, "%s", metal_error);
@@ -456,6 +473,7 @@ h3_ctx *h3_load_dir(const char *model_dir) {
         h3_free(ctx);
         return NULL;
     }
+#endif
     return ctx;
 }
 
diff --git a/h3_cli.c b/h3_cli.c
index 79339c82..6ec06745 100644
--- a/h3_cli.c
+++ b/h3_cli.c
@@ -15,6 +15,9 @@
 #include <string.h>
 #include <strings.h>
 #include <sys/stat.h>
+#ifndef __APPLE__
+#include <sys/random.h>
+#endif
 #include <time.h>
 #include <unistd.h>
 
@@ -103,7 +106,12 @@ static int set_directory(char destination[H3_CLI_PATH], const char *path) {
 
 static uint64_t random_seed(void) {
     uint64_t value;
+#ifdef __APPLE__
     arc4random_buf(&value, sizeof(value));
+#else
+    if (getrandom(&value, sizeof(value), 0) != (ssize_t)sizeof(value))
+        return (uint64_t)time(NULL);
+#endif
     return value;
 }
 
diff --git a/h3_cuda.cu b/h3_cuda.cu
new file mode 100644
index 00000000..938ee40e
--- /dev/null
+++ b/h3_cuda.cu
@@ -0,0 +1,1138 @@
+/* h3_cuda.cu - CUDA backend for h3.c (feat/cuda).
+ * Metal backend (h3_gpu.m/h3_shaders.metal) is preserved untouched;
+ * this file implements the same h3_gpu.h C API against CUDA/cuBLAS.
+ * I1 scaffold + I2 tensor layer + I3 elementwise/norm/activation/embedding real;
+ * remaining compute ops are stubs.
+ */
+#include <cuda_runtime.h>
+#include <cublas_v2.h>
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+#include <stdint.h>
+#include <fcntl.h>
+#include <unistd.h>
+#include <errno.h>
+#include <limits.h>
+#include "h3_gpu.h"
+#include "h3.h"
+#include "h3_cuda.h"
+
+#define H3_CUDA_ERR "CUDA backend: op not yet implemented (feat/cuda)"
+#define MIN(a, b) ((a) < (b) ? (a) : (b))
+#define H3_CU_BLOCK 256
+
+struct h3_gpu { void *dev_ctx; h3_gpu_stats stats; char error[512]; };
+struct h3_gpu_tensor { void *device_ptr; h3_gpu_dtype dtype; size_t elements; size_t bytes; h3_gpu *owner; };
+
+static unsigned h3_cu_grid(unsigned n) { return (n + H3_CU_BLOCK - 1) / H3_CU_BLOCK; }
+
+/* BF16 helpers -- match Metal exactly (round-to-nearest-even). __device__ so
+ * kernels can call them; CUDA intrinsics avoid memcpy in device code. */
+__device__ float h3_bf16_to_f32(uint16_t v) {
+    return __int_as_float(((uint32_t)v) << 16);
+}
+__device__ uint16_t h3_f32_to_bf16(float x) {
+    uint32_t bits = __float_as_uint(x);
+    bits += 0x7fffu + ((bits >> 16) & 1u);
+    return (uint16_t)(bits >> 16);
+}
+
+/* Row-major GEMM: C(rows x out) = A(rows x in) @ W^T, W stored (out x in) row-major.
+ * cuBLAS is column-major; compute C^T = W @ A^T via GemmEx(OP_T, OP_T, out, rows, in).
+ * ab = A/B data type, c = C data type, comp = compute type. */
+static cublasStatus_t h3_cu_gemm(h3_gpu *g, cudaDataType ab, cudaDataType c,
+                                 cublasComputeType_t comp,
+                                 const void *w, const void *a, void *out,
+                                 uint32_t rows, uint32_t in, uint32_t out_dim) {
+    cublasHandle_t h = (cublasHandle_t)g->dev_ctx;
+    const float alpha = 1.0f, beta = 0.0f;
+    return cublasGemmEx(h, CUBLAS_OP_T, CUBLAS_OP_T, (int)out_dim, (int)rows, (int)in,
+                        &alpha, w, ab, (int)in, a, ab, (int)rows, &beta, out, c, (int)out_dim,
+                        comp, CUBLAS_GEMM_DEFAULT);
+}
+
+__global__ void h3_cu_linear_bias_f32(float *out, const float *bias,
+                                      uint32_t rows, uint32_t out_dim, int has_bias) {
+    uint32_t i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < rows * out_dim) {
+        uint32_t col = i % out_dim;
+        out[i] += has_bias ? bias[col] : 0.0f;
+    }
+}
+/* BF16 bias: gemm already rounded to bf16; add bias in f32 and round once more
+ * (Metal accumulates bias in f32 before the single final bf16 rounding). */
+__global__ void h3_cu_linear_bias_bf16(uint16_t *out, const uint16_t *bias,
+                                       uint32_t rows, uint32_t out_dim, int has_bias) {
+    uint32_t i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < rows * out_dim) {
+        uint32_t col = i % out_dim;
+        float v = h3_bf16_to_f32(out[i]) + (has_bias ? h3_bf16_to_f32(bias[col]) : 0.0f);
+        out[i] = h3_f32_to_bf16(v);
+    }
+}
+
+static size_t h3_gpu_dtype_size(h3_gpu_dtype dtype) {
+    switch (dtype) {
+    case H3_GPU_F32: return sizeof(float);
+    case H3_GPU_BF16: return sizeof(uint16_t);
+    case H3_GPU_I8: return sizeof(int8_t);
+    case H3_GPU_U32: return sizeof(uint32_t);
+    default: return 0;
+    }
+}
+
+__global__ void h3_cu_silu_f32(const float *in, float *out, unsigned n) {
+    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < n) out[i] = in[i] / (1.0f + expf(-in[i]));
+}
+__global__ void h3_cu_silu_bf16(const uint16_t *in, uint16_t *out, unsigned n) {
+    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < n) { float v = h3_bf16_to_f32(in[i]); out[i] = h3_f32_to_bf16(v / (1.0f + expf(-v))); }
+}
+__global__ void h3_cu_cast_f32_to_bf16(const float *in, uint16_t *out, unsigned n) {
+    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < n) out[i] = h3_f32_to_bf16(in[i]);
+}
+__global__ void h3_cu_cast_bf16_to_f32(const uint16_t *in, float *out, unsigned n) {
+    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < n) out[i] = h3_bf16_to_f32(in[i]);
+}
+__global__ void h3_cu_clip_f32(const float *in, float *out, unsigned n, float mn, float mx) {
+    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < n) { float v = in[i]; out[i] = fminf(fmaxf(v, mn), mx); }
+}
+__global__ void h3_cu_add_bf16(const uint16_t *a, const uint16_t *b, uint16_t *o, unsigned n) {
+    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < n) o[i] = h3_f32_to_bf16(h3_bf16_to_f32(a[i]) + h3_bf16_to_f32(b[i]));
+}
+__global__ void h3_cu_sub_bf16(const uint16_t *a, const uint16_t *b, uint16_t *o, unsigned n) {
+    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < n) o[i] = h3_f32_to_bf16(h3_bf16_to_f32(a[i]) - h3_bf16_to_f32(b[i]));
+}
+__global__ void h3_cu_add_scaled_f32(const float *l, const float *r, float *o, unsigned n, float ls, float rs) {
+    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < n) o[i] = l[i] * ls + r[i] * rs;
+}
+__global__ void h3_cu_geglu_f32(const float *gate, const float *lin, float *o, unsigned n) {
+    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < n) { float x = gate[i]; float c = x*x*x;
+        o[i] = 0.5f*x*(1.0f + tanhf(0.7978845608028654f*(x + 0.044715f*c))) * lin[i]; }
+}
+__global__ void h3_cu_gelu_bf16(const uint16_t *in, uint16_t *o, unsigned n, int approx) {
+    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < n) { float v = h3_bf16_to_f32(in[i]); float act;
+        if (approx) {
+            float inner = 0.7978845608028654f*(v + 0.044715f*v*v*v);
+            act = inner <= -10.0f ? 0.0f : inner >= 10.0f ? v : 0.5f*v*(1.0f + tanhf(inner));
+        } else {
+            act = v <= -10.0f ? 0.0f : v >= 10.0f ? v : 0.5f*v*(1.0f + erf(v*0.7071067811865475f));
+        }
+        o[i] = h3_f32_to_bf16(act); }
+}
+__global__ void h3_cu_swiglu_f32(const float *fused, float *o, unsigned rows, unsigned width) {
+    unsigned col = blockIdx.x * blockDim.x + threadIdx.x; unsigned row = blockIdx.y;
+    if (row < rows && col < width) { unsigned base = row*width*2;
+        float g = fused[base+col]; float u = fused[base+width+col];
+        o[row*width+col] = (g / (1.0f + expf(-g))) * u; }
+}
+__global__ void h3_cu_swiglu_bf16(const uint16_t *fused, uint16_t *o, unsigned rows, unsigned width) {
+    unsigned col = blockIdx.x * blockDim.x + threadIdx.x; unsigned row = blockIdx.y;
+    if (row < rows && col < width) { unsigned base = row*width*2;
+        float g = h3_bf16_to_f32(fused[base+col]); float u = h3_bf16_to_f32(fused[base+width+col]);
+        o[row*width+col] = h3_f32_to_bf16((g / (1.0f + expf(-g))) * u); }
+}
+__global__ void h3_cu_silu_mul_bf16(const uint16_t *gate, const uint16_t *up, uint16_t *o, unsigned n) {
+    unsigned i = blockIdx.x * blockDim.x + threadIdx.x;
+    if (i < n) { float g = h3_bf16_to_f32(gate[i]); float u = h3_bf16_to_f32(up[i]);
+        o[i] = h3_f32_to_bf16((g / (1.0f + expf(-g))) * u); }
+}
+__global__ void h3_cu_scale_add_f32(const float *res, const float *branch, const float *scale, float *o, unsigned rows, unsigned width) {
+    unsigned col = blockIdx.x * blockDim.x + threadIdx.x; unsigned row = blockIdx.y;
+    if (row < rows && col < width) { unsigned idx = row*width+col; o[idx] = res[idx] + branch[idx]*scale[col]; }
+}
+__global__ void h3_cu_embedding_bf16(const uint16_t *w, const unsigned *ids, uint16_t *o, unsigned tokens, unsigned width, unsigned vocab) {
+    unsigned col = blockIdx.x * blockDim.x + threadIdx.x; unsigned token = blockIdx.y;
+    if (token < tokens && col < width) { unsigned id = ids[token];
+        o[token*width+col] = id < vocab ? w[id*width+col] : (uint16_t)0; }
+}
+__global__ void h3_cu_gate_bf16(const uint16_t *res, const uint16_t *branch, const uint16_t *mod, const unsigned *row_map, uint16_t *o, unsigned rows, unsigned width, unsigned slots, unsigned gate_slot) {
+    unsigned col = blockIdx.x * blockDim.x + threadIdx.x; unsigned row = blockIdx.y;
+    if (row < rows && col < width) { unsigned base = row_map[row]*slots*width;
+        float g = h3_bf16_to_f32(mod[base + gate_slot*width + col]); unsigned idx = row*width+col;
+        float v = h3_bf16_to_f32(res[idx]) + h3_bf16_to_f32(branch[idx]) * g;
+        o[idx] = h3_f32_to_bf16(v); }
+}
+__global__ void h3_cu_rms_norm_f32(const float *in, const float *w, float *o, unsigned rows, unsigned width, float eps) {
+    unsigned row = blockIdx.x; unsigned tid = threadIdx.x; unsigned threads = blockDim.x;
+    if (row >= rows) return;
+    extern __shared__ float red[];
+    const float *x = in + row*width;
+    float local = 0.0f;
+    for (unsigned k = tid; k < width; k += threads) { float v = x[k]; local = fmaf(v, v, local); }
+    red[tid] = local; __syncthreads();
+    for (unsigned stride = threads/2; stride; stride >>= 1) { if (tid < stride) red[tid] += red[tid+stride]; __syncthreads(); }
+    float inv = rsqrtf(red[0]/width + eps);
+    for (unsigned col = tid; col < width; col += threads) o[row*width+col] = x[col]*inv*w[col];
+}
+__global__ void h3_cu_rms_norm_bf16(const uint16_t *in, const uint16_t *w, uint16_t *o, unsigned rows, unsigned width, float eps) {
+    unsigned row = blockIdx.x; unsigned tid = threadIdx.x; unsigned threads = blockDim.x;
+    if (row >= rows) return;
+    extern __shared__ float red[];
+    const uint16_t *x = in + row*width;
+    float local = 0.0f;
+    for (unsigned k = tid; k < width; k += threads) { float v = h3_bf16_to_f32(x[k]); local = fmaf(v, v, local); }
+    red[tid] = local; __syncthreads();
+    for (unsigned stride = threads/2; stride; stride >>= 1) { if (tid < stride) red[tid] += red[tid+stride]; __syncthreads(); }
+    float inv = rsqrtf(red[0]/width + eps);
+    for (unsigned col = tid; col < width; col += threads) {
+        float norm = h3_bf16_to_f32(x[col])*inv;
+        o[row*width+col] = h3_f32_to_bf16(norm * h3_bf16_to_f32(w[col])); }
+}
+__global__ void h3_cu_layer_norm_f32(const float *in, const float *w, const float *b, float *o, unsigned rows, unsigned width, float eps) {
+    unsigned row = blockIdx.x; unsigned tid = threadIdx.x; unsigned threads = blockDim.x;
+    if (row >= rows) return;
+    extern __shared__ float red[];
+    const float *x = in + row*width;
+    float local = 0.0f;
+    for (unsigned k = tid; k < width; k += threads) local += x[k];
+    red[tid] = local; __syncthreads();
+    for (unsigned stride = threads/2; stride; stride >>= 1) { if (tid < stride) red[tid] += red[tid+stride]; __syncthreads(); }
+    float mean = red[0]/width; __syncthreads();
+    local = 0.0f;
+    for (unsigned k = tid; k < width; k += threads) { float c = x[k]-mean; local = fmaf(c, c, local); }
+    red[tid] = local; __syncthreads();
+    for (unsigned stride = threads/2; stride; stride >>= 1) { if (tid < stride) red[tid] += red[tid+stride]; __syncthreads(); }
+    float inv = rsqrtf(red[0]/width + eps);
+    for (unsigned col = tid; col < width; col += threads)
+        o[row*width+col] = (x[col]-mean)*inv*w[col] + b[col];
+}
+__global__ void h3_cu_layer_norm_bf16(const uint16_t *in, const uint16_t *w, const uint16_t *b, uint16_t *o, unsigned rows, unsigned width, float eps) {
+    unsigned row = blockIdx.x; unsigned tid = threadIdx.x; unsigned threads = blockDim.x;
+    if (row >= rows) return;
+    extern __shared__ float red[];
+    const uint16_t *x = in + row*width;
+    float local = 0.0f;
+    for (unsigned k = tid; k < width; k += threads) local += h3_bf16_to_f32(x[k]);
+    red[tid] = local; __syncthreads();
+    for (unsigned stride = threads/2; stride; stride >>= 1) { if (tid < stride) red[tid] += red[tid+stride]; __syncthreads(); }
+    float mean = red[0]/width; __syncthreads();
+    local = 0.0f;
+    for (unsigned k = tid; k < width; k += threads) { float c = h3_bf16_to_f32(x[k])-mean; local = fmaf(c, c, local); }
+    red[tid] = local; __syncthreads();
+    for (unsigned stride = threads/2; stride; stride >>= 1) { if (tid < stride) red[tid] += red[tid+stride]; __syncthreads(); }
+    float inv = rsqrtf(red[0]/width + eps);
+    for (unsigned col = tid; col < width; col += threads) {
+        float norm = (h3_bf16_to_f32(x[col])-mean)*inv;
+        o[row*width+col] = h3_f32_to_bf16(fmaf(norm, h3_bf16_to_f32(w[col]), h3_bf16_to_f32(b[col]))); }
+}
+__global__ void h3_cu_head_rms_norm_bf16(uint16_t *t, const uint16_t *w, unsigned seq, unsigned heads, unsigned head_dim, float eps) {
+    unsigned idx = blockIdx.x * blockDim.x + threadIdx.x;
+    unsigned row = idx / heads; unsigned head = idx % heads;
+    if (row >= seq || head >= heads) return;
+    unsigned base = (row*heads + head)*head_dim;
+    float sum = 0.0f;
+    for (unsigned d = 0; d < head_dim; d++) { float v = h3_bf16_to_f32(t[base+d]); sum = fmaf(v, v, sum); }
+    float inv = rsqrtf(sum/head_dim + eps);
+    for (unsigned d = 0; d < head_dim; d++) { float v = h3_bf16_to_f32(t[base+d]);
+        t[base+d] = h3_f32_to_bf16(v*inv*h3_bf16_to_f32(w[d])); }
+}
+__global__ void h3_cu_weight_norm_f32(const float *v, const float *mag, float *o, unsigned outer, unsigned inner) {
+    unsigned row = blockIdx.x * blockDim.x + threadIdx.x;
+    if (row >= outer) return;
+    unsigned base = row*inner;
+    float ss = 0.0f;
+    for (unsigned i = 0; i < inner; i++) { float val = v[base+i]; ss = fmaf(val, val, ss); }
+    float scale = mag[row]*rsqrtf(ss);
+    for (unsigned i = 0; i < inner; i++) o[base+i] = v[base+i]*scale;
+}
+
+/* Allocate device memory (optionally filled from host `values`). */
+static h3_gpu_tensor *h3_gpu_tensor_alloc(h3_gpu *gpu, size_t elements, h3_gpu_dtype dtype, const void *values) {
+    size_t bytes = elements * h3_gpu_dtype_size(dtype);
+    if (bytes == 0) return NULL;
+    void *dptr = NULL;
+    if (cudaMalloc(&dptr, bytes) != cudaSuccess) return NULL;
+    if (values) {
+        if (cudaMemcpy(dptr, values, bytes, cudaMemcpyHostToDevice) != cudaSuccess) {
+            cudaFree(dptr); return NULL;
+        }
+    }
+    h3_gpu_tensor *t = (h3_gpu_tensor *)calloc(1, sizeof(*t));
+    if (!t) { cudaFree(dptr); return NULL; }
+    t->device_ptr = dptr; t->dtype = dtype; t->elements = elements; t->bytes = bytes; t->owner = gpu;
+    if (gpu) {
+        gpu->stats.allocated_bytes += bytes;
+        gpu->stats.live_bytes += bytes;
+        if (gpu->stats.live_bytes > gpu->stats.peak_live_bytes)
+            gpu->stats.peak_live_bytes = gpu->stats.live_bytes;
+        gpu->stats.tensor_allocations++;
+    }
+    return t;
+}
+
+/* Stream a file range into a device tensor via a small host staging buffer. */
+static int h3_gpu_tensor_file_load(h3_gpu *gpu, h3_gpu_tensor *t, const char *path, uint64_t file_offset, char *error, size_t error_size) {
+    if (!t || !path || !*path || file_offset > (uint64_t)INT64_MAX) return 0;
+    size_t bytes = t->bytes;
+    int fd = open(path, O_RDONLY | O_CLOEXEC);
+    if (fd < 0) { if (error && error_size) snprintf(error, error_size, "cannot open %s: %s", path, strerror(errno)); return 0; }
+    size_t chunk = MIN(bytes, (size_t)(1 << 20));
+    void *host = malloc(chunk ? chunk : 1);
+    if (!host) { close(fd); return 0; }
+    size_t completed = 0;
+    while (completed < bytes) {
+        size_t request = MIN(chunk, bytes - completed);
+        size_t got = 0;
+        while (got < request) {
+            ssize_t count = pread(fd, (char *)host + got, request - got, (off_t)(file_offset + completed + got));
+            if (count < 0 && errno == EINTR) continue;
+            if (count <= 0) {
+                if (error && error_size) snprintf(error, error_size, "cannot read %s payload: %s", path, count < 0 ? strerror(errno) : "unexpected end of file");
+                free(host); close(fd); return 0;
+            }
+            got += (size_t)count;
+        }
+        if (cudaMemcpy((char *)t->device_ptr + completed, host, request, cudaMemcpyHostToDevice) != cudaSuccess) {
+            if (error && error_size) snprintf(error, error_size, "cudaMemcpy failed");
+            free(host); close(fd); return 0;
+        }
+        completed += request;
+    }
+    free(host); close(fd);
+    (void)gpu;
+    return 1;
+}
+
+static int h3_gpu_tensor_read_file_bf16_mode(h3_gpu_tensor *tensor, const char *path, uint64_t file_offset, size_t elements, int uncached, char *error, size_t error_size) {
+    (void)uncached;
+    if (error && error_size) error[0] = '\0';
+    if (!tensor || !path || !*path || tensor->dtype != H3_GPU_BF16 || elements != tensor->elements || file_offset > (uint64_t)INT64_MAX) {
+        if (error && error_size) snprintf(error, error_size, "invalid BF16 file read request");
+        return 0;
+    }
+    size_t bytes = elements * sizeof(uint16_t);
+    int fd = open(path, O_RDONLY | O_CLOEXEC);
+    if (fd < 0) { if (error && error_size) snprintf(error, error_size, "cannot open %s: %s", path, strerror(errno)); return 0; }
+    size_t chunk = MIN(bytes, (size_t)(1 << 20));
+    void *host = malloc(chunk ? chunk : 1);
+    if (!host) { close(fd); return 0; }
+    size_t completed = 0;
+    while (completed < bytes) {
+        size_t request = MIN(chunk, bytes - completed);
+        size_t got = 0;
+        while (got < request) {
+            ssize_t count = pread(fd, (char *)host + got, request - got, (off_t)(file_offset + completed + got));
+            if (count < 0 && errno == EINTR) continue;
+            if (count <= 0) {
+                if (error && error_size) snprintf(error, error_size, "cannot read BF16 payload from %s: %s", path, count < 0 ? strerror(errno) : "unexpected end of file");
+                free(host); close(fd); return 0;
+            }
+            got += (size_t)count;
+        }
+        if (cudaMemcpy((char *)tensor->device_ptr + completed, host, request, cudaMemcpyHostToDevice) != cudaSuccess) {
+            if (error && error_size) snprintf(error, error_size, "cudaMemcpy failed");
+            free(host); close(fd); return 0;
+        }
+        completed += request;
+    }
+    free(host); close(fd);
+    return 1;
+}
+
+int h3_cuda_probe(h3_device_info *info, char *error, size_t error_size) {
+    int count = 0;
+    cudaError_t ce = cudaGetDeviceCount(&count);
+    if (ce != cudaSuccess || count < 1) {
+        if (error && error_size)
+            snprintf(error, error_size, "no CUDA device available: %s", cudaGetErrorString(ce));
+        return 0;
+    }
+    if (info) {
+        memset(info, 0, sizeof(*info));
+        cudaDeviceProp prop;
+        cudaGetDeviceProperties(&prop, 0);
+        snprintf(info->name, sizeof(info->name), "%s", prop.name);
+        snprintf(info->architecture, sizeof(info->architecture), "sm_%d", prop.major * 100 + prop.minor * 10);
+        info->physical_memory = (uint64_t)prop.totalGlobalMem;
+        info->unified_memory = 1;
+    }
+    return 1;
+}
+h3_gpu *h3_gpu_create(const char *shader_source_path, char *error, size_t error_size) {
+    (void)shader_source_path;
+    h3_gpu *g = (h3_gpu *)calloc(1, sizeof(*g));
+    if (!g) { if (error && error_size) snprintf(error, error_size, "oom"); return NULL; }
+    cudaError_t ce = cudaSetDevice(0);
+    if (ce != cudaSuccess) {
+        if (error && error_size)
+            snprintf(error, error_size, "cudaSetDevice: %s", cudaGetErrorString(ce));
+        free(g); return NULL;
+    }
+    cublasHandle_t h = NULL;
+    cublasStatus_t st = cublasCreate(&h);
+    if (st != CUBLAS_STATUS_SUCCESS) {
+        if (error && error_size)
+            snprintf(error, error_size, "cublasCreate failed (%d)", (int)st);
+        free(g); return NULL;
+    }
+    g->dev_ctx = (void *)h;
+    return g;
+}
+void h3_gpu_free(h3_gpu *gpu) {
+    if (gpu) {
+        if (gpu->dev_ctx) cublasDestroy((cublasHandle_t)gpu->dev_ctx);
+        free(gpu);
+    }
+}
+const char *h3_gpu_error(const h3_gpu *gpu) {
+    return gpu && gpu->error[0] ? gpu->error : "no error";
+}
+static void h3_cuda_seterr(const h3_gpu *gpu) {
+    if (gpu) snprintf(((h3_gpu *)gpu)->error, sizeof(((h3_gpu *)gpu)->error), "%s", H3_CUDA_ERR);
+}
+
+int h3_gpu_is_m5(const h3_gpu *gpu) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_has_nax_mlp(const h3_gpu *gpu) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_has_int8_mlp(const h3_gpu *gpu) { h3_cuda_seterr(gpu); return (int)0; }
+h3_gpu_tensor * h3_gpu_tensor_new_f32(h3_gpu *gpu, size_t elements) {
+    return h3_gpu_tensor_alloc(gpu, elements, H3_GPU_F32, NULL);
+}
+h3_gpu_tensor * h3_gpu_tensor_new_bf16(h3_gpu *gpu, size_t elements) {
+    return h3_gpu_tensor_alloc(gpu, elements, H3_GPU_BF16, NULL);
+}
+h3_gpu_tensor * h3_gpu_tensor_new_i8(h3_gpu *gpu, size_t elements) {
+    return h3_gpu_tensor_alloc(gpu, elements, H3_GPU_I8, NULL);
+}
+h3_gpu_tensor * h3_gpu_tensor_from_f32(h3_gpu *gpu, const float *values,
+                                      size_t elements) {
+    return h3_gpu_tensor_alloc(gpu, elements, H3_GPU_F32, values);
+}
+h3_gpu_tensor * h3_gpu_tensor_from_bf16(h3_gpu *gpu, const uint16_t *values,
+                                       size_t elements) {
+    return h3_gpu_tensor_alloc(gpu, elements, H3_GPU_BF16, values);
+}
+h3_gpu_tensor * h3_gpu_tensor_from_u32(h3_gpu *gpu, const uint32_t *values,
+                                      size_t elements) {
+    return h3_gpu_tensor_alloc(gpu, elements, H3_GPU_U32, values);
+}
+h3_gpu_tensor * h3_gpu_tensor_load_bf16(h3_gpu *gpu, const char *path,
+                                       uint64_t file_offset, size_t elements) {
+    h3_gpu_tensor *t = h3_gpu_tensor_alloc(gpu, elements, H3_GPU_BF16, NULL);
+    if (!t) return NULL;
+    if (!h3_gpu_tensor_file_load(gpu, t, path, file_offset, NULL, 0)) { h3_gpu_tensor_free(t); return NULL; }
+    return t;
+}
+h3_gpu_tensor * h3_gpu_tensor_load_f32(h3_gpu *gpu, const char *path,
+                                      uint64_t file_offset, size_t elements) {
+    h3_gpu_tensor *t = h3_gpu_tensor_alloc(gpu, elements, H3_GPU_F32, NULL);
+    if (!t) return NULL;
+    if (!h3_gpu_tensor_file_load(gpu, t, path, file_offset, NULL, 0)) { h3_gpu_tensor_free(t); return NULL; }
+    return t;
+}
+int h3_gpu_tensor_read_file_bf16(h3_gpu_tensor *tensor, const char *path,
+                                 uint64_t file_offset, size_t elements,
+                                 char *error, size_t error_size) {
+    return h3_gpu_tensor_read_file_bf16_mode(tensor, path, file_offset, elements, 0, error, error_size);
+}
+int h3_gpu_tensor_stream_file_bf16(h3_gpu_tensor *tensor, const char *path,
+                                   uint64_t file_offset, size_t elements,
+                                   char *error, size_t error_size) {
+    return h3_gpu_tensor_read_file_bf16_mode(tensor, path, file_offset, elements, 1, error, error_size);
+}
+void h3_gpu_tensor_free(h3_gpu_tensor *tensor) {
+    if (!tensor) return;
+    if (tensor->device_ptr) cudaFree(tensor->device_ptr);
+    if (tensor->owner) { tensor->owner->stats.live_bytes = tensor->owner->stats.live_bytes >= tensor->bytes ? tensor->owner->stats.live_bytes - tensor->bytes : 0; }
+    free(tensor);
+}
+size_t h3_gpu_tensor_elements(const h3_gpu_tensor *tensor) {
+    return tensor ? tensor->elements : 0;
+}
+h3_gpu_dtype h3_gpu_tensor_dtype(const h3_gpu_tensor *tensor) {
+    return tensor ? tensor->dtype : H3_GPU_F32;
+}
+int h3_gpu_tensor_read_f32(const h3_gpu_tensor *tensor, float *values,
+                           size_t elements) {
+    return h3_gpu_tensor_read_f32_range(tensor, 0, values, elements);
+}
+int h3_gpu_tensor_read_f32_range(const h3_gpu_tensor *tensor,
+                                 size_t source_offset, float *values,
+                                 size_t elements) {
+    if (!tensor || !values || tensor->dtype != H3_GPU_F32 || source_offset > tensor->elements || elements > tensor->elements - source_offset) return 0;
+    size_t bytes = elements * sizeof(float);
+    return (cudaMemcpy(values, (char *)tensor->device_ptr + source_offset * sizeof(float), bytes, cudaMemcpyDeviceToHost) == cudaSuccess) ? 1 : 0;
+}
+int h3_gpu_tensor_read_bf16(const h3_gpu_tensor *tensor, uint16_t *values,
+                            size_t elements) {
+    if (!tensor || !values || tensor->dtype != H3_GPU_BF16 || elements > tensor->elements) return 0;
+    size_t bytes = elements * sizeof(uint16_t);
+    return (cudaMemcpy(values, tensor->device_ptr, bytes, cudaMemcpyDeviceToHost) == cudaSuccess) ? 1 : 0;
+}
+int h3_gpu_tensor_write_f32(h3_gpu_tensor *tensor, const float *values,
+                            size_t elements) {
+    return h3_gpu_tensor_write_f32_range(tensor, 0, values, elements);
+}
+int h3_gpu_tensor_write_f32_range(h3_gpu_tensor *tensor,
+                                  size_t destination_offset,
+                                  const float *values, size_t elements) {
+    if (!tensor || !values || tensor->dtype != H3_GPU_F32 || destination_offset > tensor->elements || elements > tensor->elements - destination_offset) return 0;
+    size_t bytes = elements * sizeof(float);
+    return (cudaMemcpy((char *)tensor->device_ptr + destination_offset * sizeof(float), values, bytes, cudaMemcpyHostToDevice) == cudaSuccess) ? 1 : 0;
+}
+int h3_gpu_tensor_write_bf16(h3_gpu_tensor *tensor, const uint16_t *values,
+                             size_t elements) {
+    return h3_gpu_tensor_write_bf16_range(tensor, 0, values, elements);
+}
+int h3_gpu_tensor_write_bf16_range(h3_gpu_tensor *tensor,
+                                   size_t destination_offset,
+                                   const uint16_t *values, size_t elements) {
+    if (!tensor || !values || tensor->dtype != H3_GPU_BF16 || destination_offset > tensor->elements || elements > tensor->elements - destination_offset) return 0;
+    size_t bytes = elements * sizeof(uint16_t);
+    return (cudaMemcpy((char *)tensor->device_ptr + destination_offset * sizeof(uint16_t), values, bytes, cudaMemcpyHostToDevice) == cudaSuccess) ? 1 : 0;
+}
+int h3_gpu_begin(h3_gpu *gpu) {
+    (void)gpu;
+    return 1;  /* CUDA has no explicit command buffer */
+}
+int h3_gpu_continue(h3_gpu *gpu) {
+    (void)gpu;
+    return 1;  /* CUDA has no explicit command buffer */
+}
+int h3_gpu_submit(h3_gpu *gpu) {
+    (void)gpu;
+    return 1;  /* CUDA has no explicit command buffer */
+}
+int h3_gpu_get_stats(const h3_gpu *gpu, h3_gpu_stats *stats) {
+    if (!gpu || !stats) return 0;
+    *stats = gpu->stats;
+    return 1;
+}
+void h3_gpu_profile_set_label(h3_gpu *gpu, const char *label) {
+    (void)gpu; (void)label;
+    /* no-op on CUDA */
+}
+void h3_gpu_profile_mark(h3_gpu *gpu, const char *phase) {
+    (void)gpu; (void)phase;
+    /* no-op on CUDA */
+}
+int h3_gpu_linear_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                      const h3_gpu_tensor *input, const h3_gpu_tensor *weight,
+                      const h3_gpu_tensor *bias, uint32_t rows,
+                      uint32_t input_dim, uint32_t output_dim) {
+    if (!gpu || !output || !input || !weight || !output->device_ptr ||
+        !input->device_ptr || !weight->device_ptr) return 0;
+    cublasStatus_t st = h3_cu_gemm(gpu, CUDA_R_32F, CUDA_R_32F, CUBLAS_COMPUTE_32F,
+                                   weight->device_ptr, input->device_ptr, output->device_ptr,
+                                   rows, input_dim, output_dim);
+    if (st != CUBLAS_STATUS_SUCCESS) {
+        snprintf(((h3_gpu *)gpu)->error, sizeof(((h3_gpu *)gpu)->error),
+                 "cublasGemmEx f32 failed (%d)", (int)st);
+        return 0;
+    }
+    if (bias && bias->device_ptr)
+        h3_cu_linear_bias_f32<<<h3_cu_grid(rows * output_dim), H3_CU_BLOCK>>>(
+            (float *)output->device_ptr, (const float *)bias->device_ptr,
+            rows, output_dim, 1);
+    return 1;
+}
+int h3_gpu_patch_linear_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                             const h3_gpu_tensor *input,
+                             const h3_gpu_tensor *weight,
+                             const h3_gpu_tensor *bias, uint32_t rows,
+                             uint32_t input_dim, uint32_t output_dim) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_patch_linear_bf16_offset(
+                             h3_gpu *gpu, h3_gpu_tensor *output,
+                             size_t output_offset,
+                             const h3_gpu_tensor *input, size_t input_offset,
+                             const h3_gpu_tensor *weight,
+                             const h3_gpu_tensor *bias, uint32_t rows,
+                             uint32_t input_dim, uint32_t output_dim) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_patch_linear_bf16_map(
+                             h3_gpu *gpu, h3_gpu_tensor *output,
+                             const h3_gpu_tensor *input,
+                             const h3_gpu_tensor *weight,
+                             const h3_gpu_tensor *bias,
+                             const h3_gpu_tensor *row_map,
+                             uint32_t output_rows, uint32_t rows,
+                             uint32_t input_dim, uint32_t output_dim) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_silu_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                    const h3_gpu_tensor *input, uint32_t elements) {
+    if (!output || !input || output->dtype != H3_GPU_F32 || input->dtype != H3_GPU_F32 || elements > input->elements || elements > output->elements) return 0;
+    h3_cu_silu_f32<<<h3_cu_grid(elements), H3_CU_BLOCK>>>((const float *)input->device_ptr, (float *)output->device_ptr, elements);
+    return 1;
+}
+int h3_gpu_cast_f32_to_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                            const h3_gpu_tensor *input, uint32_t elements) {
+    if (!output || !input || output->dtype != H3_GPU_BF16 || input->dtype != H3_GPU_F32 || elements > input->elements || elements > output->elements) return 0;
+    h3_cu_cast_f32_to_bf16<<<h3_cu_grid(elements), H3_CU_BLOCK>>>((const float *)input->device_ptr, (uint16_t *)output->device_ptr, elements);
+    return 1;
+}
+int h3_gpu_cast_bf16_to_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                            const h3_gpu_tensor *input, uint32_t elements) {
+    if (!output || !input || output->dtype != H3_GPU_F32 || input->dtype != H3_GPU_BF16 || elements > input->elements || elements > output->elements) return 0;
+    h3_cu_cast_bf16_to_f32<<<h3_cu_grid(elements), H3_CU_BLOCK>>>((const uint16_t *)input->device_ptr, (float *)output->device_ptr, elements);
+    return 1;
+}
+int h3_gpu_copy_bf16(h3_gpu *gpu, h3_gpu_tensor *destination,
+                     size_t destination_offset,
+                     const h3_gpu_tensor *source, size_t source_offset,
+                     size_t elements) {
+    if (!destination || !source || destination->dtype != H3_GPU_BF16 || source->dtype != H3_GPU_BF16 || source_offset > source->elements || elements > source->elements - source_offset || destination_offset > destination->elements || elements > destination->elements - destination_offset) return 0;
+    size_t b = elements * sizeof(uint16_t);
+    return (cudaMemcpy((char *)destination->device_ptr + destination_offset*sizeof(uint16_t), (char *)source->device_ptr + source_offset*sizeof(uint16_t), b, cudaMemcpyDeviceToDevice) == cudaSuccess) ? 1 : 0;
+}
+int h3_gpu_copy_f32(h3_gpu *gpu, h3_gpu_tensor *destination,
+                    size_t destination_offset,
+                    const h3_gpu_tensor *source, size_t source_offset,
+                    size_t elements) {
+    if (!destination || !source || destination->dtype != H3_GPU_F32 || source->dtype != H3_GPU_F32 || source_offset > source->elements || elements > source->elements - source_offset || destination_offset > destination->elements || elements > destination->elements - destination_offset) return 0;
+    size_t b = elements * sizeof(float);
+    return (cudaMemcpy((char *)destination->device_ptr + destination_offset*sizeof(float), (char *)source->device_ptr + source_offset*sizeof(float), b, cudaMemcpyDeviceToDevice) == cudaSuccess) ? 1 : 0;
+}
+int h3_gpu_rms_norm_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                        const h3_gpu_tensor *input,
+                        const h3_gpu_tensor *weight, uint32_t rows,
+                        uint32_t width, float epsilon) {
+    if (!output || !input || !weight || output->dtype != H3_GPU_F32 || input->dtype != H3_GPU_F32 || weight->dtype != H3_GPU_F32 || (size_t)rows*width > input->elements || width > weight->elements || (size_t)rows*width > output->elements) return 0;
+    h3_cu_rms_norm_f32<<<rows, H3_CU_BLOCK, H3_CU_BLOCK*sizeof(float)>>>((const float *)input->device_ptr, (const float *)weight->device_ptr, (float *)output->device_ptr, rows, width, epsilon);
+    return 1;
+}
+int h3_gpu_adaln_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                     const h3_gpu_tensor *input,
+                     const h3_gpu_tensor *norm_weight,
+                     const h3_gpu_tensor *modulation,
+                     const h3_gpu_tensor *row_map, uint32_t rows,
+                     uint32_t width, uint32_t slots, uint32_t shift_slot,
+                     uint32_t scale_slot, float epsilon) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_gate_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                    const h3_gpu_tensor *residual,
+                    const h3_gpu_tensor *branch,
+                    const h3_gpu_tensor *modulation,
+                    const h3_gpu_tensor *row_map, uint32_t rows,
+                    uint32_t width, uint32_t slots, uint32_t gate_slot) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_qkv_rope_f32(h3_gpu *gpu, h3_gpu_tensor *query,
+                        h3_gpu_tensor *key, h3_gpu_tensor *value,
+                        const h3_gpu_tensor *qkv,
+                        const h3_gpu_tensor *q_norm,
+                        const h3_gpu_tensor *k_norm,
+                        const h3_gpu_tensor *rope_cos,
+                        const h3_gpu_tensor *rope_sin, uint32_t sequence,
+                        uint32_t heads, uint32_t head_dim,
+                        uint32_t rope_half, float epsilon) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_sdpa_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                    const h3_gpu_tensor *query, const h3_gpu_tensor *key,
+                    const h3_gpu_tensor *value, uint32_t sequence,
+                    uint32_t heads, uint32_t head_dim, float scale) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_swiglu_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                      const h3_gpu_tensor *fused, uint32_t rows,
+                      uint32_t width) {
+    if (!output || !fused || output->dtype != H3_GPU_F32 || fused->dtype != H3_GPU_F32 || (size_t)rows*width*2 > fused->elements || (size_t)rows*width > output->elements) return 0;
+    dim3 g(h3_cu_grid(width), rows);
+    h3_cu_swiglu_f32<<<g, H3_CU_BLOCK>>>((const float *)fused->device_ptr, (float *)output->device_ptr, rows, width);
+    return 1;
+}
+int h3_gpu_scale_add_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                         const h3_gpu_tensor *residual,
+                         const h3_gpu_tensor *branch,
+                         const h3_gpu_tensor *scale, uint32_t rows,
+                         uint32_t width) {
+    if (!output || !residual || !branch || !scale || output->dtype != H3_GPU_F32 || residual->dtype != H3_GPU_F32 || branch->dtype != H3_GPU_F32 || scale->dtype != H3_GPU_F32 || (size_t)rows*width > residual->elements || (size_t)rows*width > branch->elements || width > scale->elements || (size_t)rows*width > output->elements) return 0;
+    dim3 g(h3_cu_grid(width), rows);
+    h3_cu_scale_add_f32<<<g, H3_CU_BLOCK>>>((const float *)residual->device_ptr, (const float *)branch->device_ptr, (const float *)scale->device_ptr, (float *)output->device_ptr, rows, width);
+    return 1;
+}
+int h3_gpu_layer_norm_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                          const h3_gpu_tensor *input,
+                          const h3_gpu_tensor *weight,
+                          const h3_gpu_tensor *bias, uint32_t rows,
+                          uint32_t width, float epsilon) {
+    if (!output || !input || !weight || !bias || output->dtype != H3_GPU_F32 || input->dtype != H3_GPU_F32 || weight->dtype != H3_GPU_F32 || bias->dtype != H3_GPU_F32 || (size_t)rows*width > input->elements || width > weight->elements || width > bias->elements || (size_t)rows*width > output->elements) return 0;
+    h3_cu_layer_norm_f32<<<rows, H3_CU_BLOCK, H3_CU_BLOCK*sizeof(float)>>>((const float *)input->device_ptr, (const float *)weight->device_ptr, (const float *)bias->device_ptr, (float *)output->device_ptr, rows, width, epsilon);
+    return 1;
+}
+int h3_gpu_video_qkv_rope_f32(h3_gpu *gpu, h3_gpu_tensor *query,
+                              h3_gpu_tensor *key, h3_gpu_tensor *value,
+                              const h3_gpu_tensor *qkv,
+                              const h3_gpu_tensor *rope_cos,
+                              const h3_gpu_tensor *rope_sin,
+                              uint32_t sequence, uint32_t heads,
+                              uint32_t head_dim, uint32_t rope_half,
+                              float epsilon) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_conv1d_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                      const h3_gpu_tensor *input,
+                      const h3_gpu_tensor *weight,
+                      const h3_gpu_tensor *bias, uint32_t batch,
+                      uint32_t length, uint32_t input_channels,
+                      uint32_t output_channels, uint32_t kernel,
+                      uint32_t padding, uint32_t dilation) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_conv1d_stride_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                      const h3_gpu_tensor *input,
+                      const h3_gpu_tensor *weight,
+                      const h3_gpu_tensor *bias, uint32_t batch,
+                      uint32_t length, uint32_t input_channels,
+                      uint32_t output_channels, uint32_t kernel,
+                      uint32_t stride, uint32_t padding,
+                      uint32_t dilation) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_conv_transpose1d_f32(
+                      h3_gpu *gpu, h3_gpu_tensor *output,
+                      const h3_gpu_tensor *input,
+                      const h3_gpu_tensor *weight,
+                      const h3_gpu_tensor *bias, uint32_t batch,
+                      uint32_t length, uint32_t input_channels,
+                      uint32_t output_channels, uint32_t kernel,
+                      uint32_t stride, uint32_t padding) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_weight_norm_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                           const h3_gpu_tensor *vector,
+                           const h3_gpu_tensor *magnitude,
+                           uint32_t outer, uint32_t inner) {
+    if (!output || !vector || !magnitude || output->dtype != H3_GPU_F32 || vector->dtype != H3_GPU_F32 || magnitude->dtype != H3_GPU_F32 || (size_t)outer*inner > vector->elements || outer > magnitude->elements || (size_t)outer*inner > output->elements) return 0;
+    h3_cu_weight_norm_f32<<<h3_cu_grid(outer), H3_CU_BLOCK>>>((const float *)vector->device_ptr, (const float *)magnitude->device_ptr, (float *)output->device_ptr, outer, inner);
+    return 1;
+}
+int h3_gpu_add_scaled_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                          const h3_gpu_tensor *left,
+                          const h3_gpu_tensor *right, float left_scale,
+                          float right_scale, uint32_t elements) {
+    if (!output || !left || !right || output->dtype != H3_GPU_F32 || left->dtype != H3_GPU_F32 || right->dtype != H3_GPU_F32 || elements > left->elements || elements > right->elements || elements > output->elements) return 0;
+    h3_cu_add_scaled_f32<<<h3_cu_grid(elements), H3_CU_BLOCK>>>((const float *)left->device_ptr, (const float *)right->device_ptr, (float *)output->device_ptr, elements, left_scale, right_scale);
+    return 1;
+}
+int h3_gpu_alias_free_snake_f32(
+                          h3_gpu *gpu, h3_gpu_tensor *output,
+                          const h3_gpu_tensor *input,
+                          const h3_gpu_tensor *alpha_log,
+                          const h3_gpu_tensor *beta_log,
+                          const h3_gpu_tensor *upsample_filter,
+                          const h3_gpu_tensor *downsample_filter,
+                          uint32_t batch, uint32_t length,
+                          uint32_t channels) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_snake1d_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                       const h3_gpu_tensor *input,
+                       const h3_gpu_tensor *alpha, uint32_t batch,
+                       uint32_t length, uint32_t channels) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_audio_qkv_split_f32(h3_gpu *gpu,
+                       h3_gpu_tensor *query, h3_gpu_tensor *key,
+                       h3_gpu_tensor *value, const h3_gpu_tensor *qkv,
+                       const h3_gpu_tensor *q_bias,
+                       const h3_gpu_tensor *k_bias,
+                       const h3_gpu_tensor *v_bias, uint32_t batch,
+                       uint32_t length, uint32_t heads,
+                       uint32_t head_dim) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_sdpa_causal_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                       const h3_gpu_tensor *query,
+                       const h3_gpu_tensor *key,
+                       const h3_gpu_tensor *value, uint32_t batch,
+                       uint32_t sequence, uint32_t heads,
+                       uint32_t head_dim, float scale) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_audio_attention_pool_f32(h3_gpu *gpu,
+                       h3_gpu_tensor *output,
+                       const h3_gpu_tensor *attended, uint32_t batch,
+                       uint32_t length, uint32_t heads,
+                       uint32_t head_dim, uint32_t output_dim) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_geglu_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                     const h3_gpu_tensor *gate,
+                     const h3_gpu_tensor *linear, uint32_t elements) {
+    if (!output || !gate || !linear || output->dtype != H3_GPU_F32 || gate->dtype != H3_GPU_F32 || linear->dtype != H3_GPU_F32 || elements > gate->elements || elements > linear->elements || elements > output->elements) return 0;
+    h3_cu_geglu_f32<<<h3_cu_grid(elements), H3_CU_BLOCK>>>((const float *)gate->device_ptr, (const float *)linear->device_ptr, (float *)output->device_ptr, elements);
+    return 1;
+}
+int h3_gpu_clip_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                    const h3_gpu_tensor *input, uint32_t elements,
+                    float minimum, float maximum) {
+    if (!output || !input || output->dtype != H3_GPU_F32 || input->dtype != H3_GPU_F32 || elements > input->elements || elements > output->elements) return 0;
+    h3_cu_clip_f32<<<h3_cu_grid(elements), H3_CU_BLOCK>>>((const float *)input->device_ptr, (float *)output->device_ptr, elements, minimum, maximum);
+    return 1;
+}
+int h3_gpu_vae_encoder_pad_f32(
+                    h3_gpu *gpu, h3_gpu_tensor *output,
+                    const h3_gpu_tensor *input, uint32_t batch,
+                    uint32_t depth, uint32_t height, uint32_t width,
+                    uint32_t channels, uint32_t depth_front,
+                    uint32_t height_before, uint32_t height_after,
+                    uint32_t width_before, uint32_t width_after) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_conv3d_f32(h3_gpu *gpu, h3_gpu_tensor *output,
+                      const h3_gpu_tensor *input,
+                      const h3_gpu_tensor *weight,
+                      const h3_gpu_tensor *bias, uint32_t batch,
+                      uint32_t depth, uint32_t height, uint32_t width,
+                      uint32_t input_channels, uint32_t output_channels,
+                      uint32_t kernel_depth, uint32_t kernel_height,
+                      uint32_t kernel_width, uint32_t stride_depth,
+                      uint32_t stride_height, uint32_t stride_width) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_vae_encoder_group_norm_silu_f32(
+                      h3_gpu *gpu, h3_gpu_tensor *output,
+                      const h3_gpu_tensor *input,
+                      const h3_gpu_tensor *weight,
+                      const h3_gpu_tensor *bias, uint32_t batch,
+                      uint32_t depth, uint32_t height, uint32_t width,
+                      uint32_t channels, uint32_t groups, float epsilon) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_linear_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                       const h3_gpu_tensor *input,
+                       const h3_gpu_tensor *weight,
+                       const h3_gpu_tensor *bias, uint32_t rows,
+                       uint32_t input_dim, uint32_t output_dim) {
+    if (!gpu || !output || !input || !weight || !output->device_ptr ||
+        !input->device_ptr || !weight->device_ptr) return 0;
+    cublasStatus_t st = h3_cu_gemm(gpu, CUDA_R_16BF, CUDA_R_16BF, CUBLAS_COMPUTE_32F,
+                                   weight->device_ptr, input->device_ptr, output->device_ptr,
+                                   rows, input_dim, output_dim);
+    if (st != CUBLAS_STATUS_SUCCESS) {
+        snprintf(((h3_gpu *)gpu)->error, sizeof(((h3_gpu *)gpu)->error),
+                 "cublasGemmEx bf16 failed (%d)", (int)st);
+        return 0;
+    }
+    if (bias && bias->device_ptr)
+        h3_cu_linear_bias_bf16<<<h3_cu_grid(rows * output_dim), H3_CU_BLOCK>>>(
+            (uint16_t *)output->device_ptr, (const uint16_t *)bias->device_ptr,
+            rows, output_dim, 1);
+    return 1;
+}
+int h3_gpu_mlp_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                    const h3_gpu_tensor *input,
+                    const h3_gpu_tensor *fc1_weight,
+                    const h3_gpu_tensor *fc2_weight, uint32_t rows,
+                    uint32_t input_dim, uint32_t hidden_dim,
+                    uint32_t output_dim) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_mlp_nax_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                        h3_gpu_tensor *activated,
+                        const h3_gpu_tensor *input,
+                        const h3_gpu_tensor *fc1_weight,
+                        const h3_gpu_tensor *fc2_weight, uint32_t rows,
+                        uint32_t input_dim, uint32_t hidden_dim,
+                        uint32_t output_dim) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_quantize_weight_int8(h3_gpu *gpu, h3_gpu_tensor *output,
+                                h3_gpu_tensor *scales,
+                                const h3_gpu_tensor *input, uint32_t rows,
+                                uint32_t columns) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_linear_int8_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                            h3_gpu_tensor *quantized_input,
+                            h3_gpu_tensor *input_scales,
+                            const h3_gpu_tensor *input,
+                            const h3_gpu_tensor *weight,
+                            const h3_gpu_tensor *weight_scales,
+                            uint32_t rows, uint32_t input_dim,
+                            uint32_t output_dim,
+                            int use_slower_uncached_int8_scales) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_linear_int8_head_major_bf16(
+                            h3_gpu *gpu, h3_gpu_tensor *output,
+                            h3_gpu_tensor *quantized_input,
+                            h3_gpu_tensor *input_scales,
+                            const h3_gpu_tensor *input,
+                            const h3_gpu_tensor *weight,
+                            const h3_gpu_tensor *weight_scales,
+                            uint32_t rows, uint32_t heads,
+                            uint32_t head_dim, uint32_t output_dim) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_mlp_int8_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                         h3_gpu_tensor *activated,
+                         h3_gpu_tensor *quantized_activation,
+                         h3_gpu_tensor *activation_scales,
+                         const h3_gpu_tensor *input,
+                         const h3_gpu_tensor *fc1_weight,
+                         const h3_gpu_tensor *fc1_scales,
+                         const h3_gpu_tensor *fc2_weight,
+                         const h3_gpu_tensor *fc2_scales,
+                         const h3_gpu_tensor *fc1_bf16,
+                         const h3_gpu_tensor *fc2_bf16, uint32_t rows,
+                         uint32_t input_dim, uint32_t hidden_dim,
+                         uint32_t output_dim,
+                         int use_slower_grouped_quantizer,
+                         int use_slower_dynamic_fc1_k,
+                         int use_int8_row_fc2,
+                         int input_is_quantized) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_silu_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                     const h3_gpu_tensor *input, uint32_t elements) {
+    if (!output || !input || output->dtype != H3_GPU_BF16 || input->dtype != H3_GPU_BF16 || elements > input->elements || elements > output->elements) return 0;
+    h3_cu_silu_bf16<<<h3_cu_grid(elements), H3_CU_BLOCK>>>((const uint16_t *)input->device_ptr, (uint16_t *)output->device_ptr, elements);
+    return 1;
+}
+int h3_gpu_rms_norm_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                         const h3_gpu_tensor *input,
+                         const h3_gpu_tensor *weight, uint32_t rows,
+                         uint32_t width, float epsilon) {
+    if (!output || !input || !weight || output->dtype != H3_GPU_BF16 || input->dtype != H3_GPU_BF16 || weight->dtype != H3_GPU_BF16 || (size_t)rows*width > input->elements || width > weight->elements || (size_t)rows*width > output->elements) return 0;
+    h3_cu_rms_norm_bf16<<<rows, H3_CU_BLOCK, H3_CU_BLOCK*sizeof(float)>>>((const uint16_t *)input->device_ptr, (const uint16_t *)weight->device_ptr, (uint16_t *)output->device_ptr, rows, width, epsilon);
+    return 1;
+}
+int h3_gpu_layer_norm_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                           const h3_gpu_tensor *input,
+                           const h3_gpu_tensor *weight,
+                           const h3_gpu_tensor *bias, uint32_t rows,
+                           uint32_t width, float epsilon) {
+    if (!output || !input || !weight || !bias || output->dtype != H3_GPU_BF16 || input->dtype != H3_GPU_BF16 || weight->dtype != H3_GPU_BF16 || bias->dtype != H3_GPU_BF16 || (size_t)rows*width > input->elements || width > weight->elements || width > bias->elements || (size_t)rows*width > output->elements) return 0;
+    h3_cu_layer_norm_bf16<<<rows, H3_CU_BLOCK, H3_CU_BLOCK*sizeof(float)>>>((const uint16_t *)input->device_ptr, (const uint16_t *)weight->device_ptr, (const uint16_t *)bias->device_ptr, (uint16_t *)output->device_ptr, rows, width, epsilon);
+    return 1;
+}
+int h3_gpu_gelu_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                     const h3_gpu_tensor *input, uint32_t elements,
+                     int approximate) {
+    if (!output || !input || output->dtype != H3_GPU_BF16 || input->dtype != H3_GPU_BF16 || elements > input->elements || elements > output->elements) return 0;
+    h3_cu_gelu_bf16<<<h3_cu_grid(elements), H3_CU_BLOCK>>>((const uint16_t *)input->device_ptr, (uint16_t *)output->device_ptr, elements, approximate);
+    return 1;
+}
+int h3_gpu_vision_qkv_rope_bf16(
+                     h3_gpu *gpu, h3_gpu_tensor *query,
+                     h3_gpu_tensor *key, h3_gpu_tensor *value,
+                     const h3_gpu_tensor *qkv,
+                     const h3_gpu_tensor *rope_cos,
+                     const h3_gpu_tensor *rope_sin, uint32_t sequence,
+                     uint32_t heads, uint32_t head_dim,
+                     uint32_t rope_half) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_adaln_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                      const h3_gpu_tensor *input,
+                      const h3_gpu_tensor *norm_weight,
+                      const h3_gpu_tensor *modulation,
+                      const h3_gpu_tensor *row_map, uint32_t rows,
+                      uint32_t width, uint32_t slots, uint32_t shift_slot,
+                      uint32_t scale_slot, float epsilon) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_adaln_bf16_offset(h3_gpu *gpu, h3_gpu_tensor *output,
+                      const h3_gpu_tensor *input, size_t input_offset,
+                      const h3_gpu_tensor *norm_weight,
+                      const h3_gpu_tensor *modulation,
+                      const h3_gpu_tensor *row_map, uint32_t rows,
+                      uint32_t width, uint32_t slots, uint32_t shift_slot,
+                      uint32_t scale_slot, float epsilon) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_adaln_linear_bf16(
+                      h3_gpu *gpu, h3_gpu_tensor *output,
+                      h3_gpu_tensor *inverse,
+                      const h3_gpu_tensor *input, size_t input_offset,
+                      const h3_gpu_tensor *norm_weight,
+                      const h3_gpu_tensor *modulation,
+                      const h3_gpu_tensor *row_map,
+                      const h3_gpu_tensor *weight,
+                      const h3_gpu_tensor *bias, uint32_t rows,
+                      uint32_t width, uint32_t output_dim, uint32_t slots,
+                      uint32_t shift_slot, uint32_t scale_slot,
+                      float epsilon) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_gate_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                     const h3_gpu_tensor *residual,
+                     const h3_gpu_tensor *branch,
+                     const h3_gpu_tensor *modulation,
+                     const h3_gpu_tensor *row_map, uint32_t rows,
+                     uint32_t width, uint32_t slots, uint32_t gate_slot) {
+    if (!output || !residual || !branch || !modulation || !row_map || output->dtype != H3_GPU_BF16 || residual->dtype != H3_GPU_BF16 || branch->dtype != H3_GPU_BF16 || modulation->dtype != H3_GPU_BF16 || row_map->dtype != H3_GPU_U32 || (size_t)rows*width > residual->elements || (size_t)rows*width > branch->elements || (size_t)rows*width > output->elements || rows > row_map->elements) return 0;
+    dim3 g(h3_cu_grid(width), rows);
+    h3_cu_gate_bf16<<<g, H3_CU_BLOCK>>>((const uint16_t *)residual->device_ptr, (const uint16_t *)branch->device_ptr, (const uint16_t *)modulation->device_ptr, (const unsigned *)row_map->device_ptr, (uint16_t *)output->device_ptr, rows, width, slots, gate_slot);
+    return 1;
+}
+int h3_gpu_gate_adaln_bf16(
+                     h3_gpu *gpu, h3_gpu_tensor *gated_residual,
+                     h3_gpu_tensor *output,
+                     const h3_gpu_tensor *residual,
+                     const h3_gpu_tensor *branch,
+                     const h3_gpu_tensor *norm_weight,
+                     const h3_gpu_tensor *gate_modulation,
+                     const h3_gpu_tensor *norm_modulation,
+                     const h3_gpu_tensor *row_map, uint32_t rows,
+                     uint32_t width, uint32_t slots, uint32_t gate_slot,
+                     uint32_t shift_slot, uint32_t scale_slot,
+                     float epsilon) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_gate_adaln_quantize_int8(
+                     h3_gpu *gpu, h3_gpu_tensor *gated_residual,
+                     h3_gpu_tensor *quantized_output,
+                     h3_gpu_tensor *quantized_scales,
+                     const h3_gpu_tensor *residual,
+                     const h3_gpu_tensor *branch,
+                     const h3_gpu_tensor *norm_weight,
+                     const h3_gpu_tensor *gate_modulation,
+                     const h3_gpu_tensor *norm_modulation,
+                     const h3_gpu_tensor *row_map, uint32_t rows,
+                     uint32_t padded_rows, uint32_t width, uint32_t slots,
+                     uint32_t gate_slot, uint32_t shift_slot,
+                     uint32_t scale_slot, float epsilon) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_qkv_rope_bf16(h3_gpu *gpu, h3_gpu_tensor *query,
+                         h3_gpu_tensor *key, h3_gpu_tensor *value,
+                         const h3_gpu_tensor *qkv,
+                         const h3_gpu_tensor *q_norm,
+                         const h3_gpu_tensor *k_norm,
+                         const h3_gpu_tensor *rope_cos,
+                         const h3_gpu_tensor *rope_sin, uint32_t sequence,
+                         uint32_t heads, uint32_t head_dim,
+                         uint32_t rope_half, float epsilon) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_grouped_qkv_rope_bf16(h3_gpu *gpu, h3_gpu_tensor *query,
+                                 h3_gpu_tensor *key, h3_gpu_tensor *value,
+                                 const h3_gpu_tensor *qkv,
+                                 const h3_gpu_tensor *q_norm,
+                                 const h3_gpu_tensor *k_norm,
+                                 const h3_gpu_tensor *rope_cos,
+                                 const h3_gpu_tensor *rope_sin,
+                                 uint32_t sequence, uint32_t heads,
+                                 uint32_t head_dim, uint32_t rope_half,
+                                 float epsilon) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_grouped_qkv_linear_rope_bf16(
+                                 h3_gpu *gpu,
+                                 h3_gpu_tensor *query,
+                                 h3_gpu_tensor *key,
+                                 h3_gpu_tensor *value,
+                                 h3_gpu_tensor *qkv,
+                                 const h3_gpu_tensor *input,
+                                 const h3_gpu_tensor *weight,
+                                 const h3_gpu_tensor *q_norm,
+                                 const h3_gpu_tensor *k_norm,
+                                 const h3_gpu_tensor *rope_cos,
+                                 const h3_gpu_tensor *rope_sin,
+                                 uint32_t rows, uint32_t input_dim,
+                                 uint32_t heads, uint32_t head_dim,
+                                 uint32_t rope_half, float epsilon) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_grouped_qkv_linear_rope_int8(
+                                 h3_gpu *gpu,
+                                 h3_gpu_tensor *query,
+                                 h3_gpu_tensor *key,
+                                 h3_gpu_tensor *value,
+                                 h3_gpu_tensor *quantized_input,
+                                 h3_gpu_tensor *input_scales,
+                                 const h3_gpu_tensor *input,
+                                 const h3_gpu_tensor *weight,
+                                 const h3_gpu_tensor *weight_scales,
+                                 const h3_gpu_tensor *q_norm,
+                                 const h3_gpu_tensor *k_norm,
+                                 const h3_gpu_tensor *rope_cos,
+                                 const h3_gpu_tensor *rope_sin,
+                                 uint32_t rows, uint32_t input_dim,
+                                 uint32_t heads, uint32_t head_dim,
+                                 uint32_t rope_half, float epsilon,
+                                 int input_is_quantized,
+                                 int use_slower_unfused_qkv_rope,
+                                 int use_slower_scalar_qkv_rms,
+                                 int use_slower_uncached_int8_scales) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_sdpa_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                     const h3_gpu_tensor *query, const h3_gpu_tensor *key,
+                     const h3_gpu_tensor *value, uint32_t sequence,
+                     uint32_t heads, uint32_t head_dim, float scale) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_sdpa_bf16_head_major_output(
+                     h3_gpu *gpu, h3_gpu_tensor *output,
+                     const h3_gpu_tensor *query, const h3_gpu_tensor *key,
+                     const h3_gpu_tensor *value, uint32_t sequence,
+                     uint32_t heads, uint32_t head_dim, float scale) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_swiglu_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                       const h3_gpu_tensor *fused, uint32_t rows,
+                       uint32_t width) {
+    if (!output || !fused || output->dtype != H3_GPU_BF16 || fused->dtype != H3_GPU_BF16 || (size_t)rows*width*2 > fused->elements || (size_t)rows*width > output->elements) return 0;
+    dim3 g(h3_cu_grid(width), rows);
+    h3_cu_swiglu_bf16<<<g, H3_CU_BLOCK>>>((const uint16_t *)fused->device_ptr, (uint16_t *)output->device_ptr, rows, width);
+    return 1;
+}
+int h3_gpu_embedding_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                          const h3_gpu_tensor *weight,
+                          const h3_gpu_tensor *token_ids, uint32_t tokens,
+                          uint32_t vocab_size, uint32_t width) {
+    if (!output || !weight || !token_ids || output->dtype != H3_GPU_BF16 || weight->dtype != H3_GPU_BF16 || token_ids->dtype != H3_GPU_U32 || (size_t)tokens*width > output->elements || (size_t)vocab_size*width > weight->elements || tokens > token_ids->elements) return 0;
+    dim3 g(h3_cu_grid(width), tokens);
+    h3_cu_embedding_bf16<<<g, H3_CU_BLOCK>>>((const uint16_t *)weight->device_ptr, (const unsigned *)token_ids->device_ptr, (uint16_t *)output->device_ptr, tokens, width, vocab_size);
+    return 1;
+}
+int h3_gpu_text_qk_rope_bf16(h3_gpu *gpu,
+                             h3_gpu_tensor *query_output,
+                             h3_gpu_tensor *key_output,
+                             const h3_gpu_tensor *query_input,
+                             const h3_gpu_tensor *key_input,
+                             const h3_gpu_tensor *q_norm,
+                             const h3_gpu_tensor *k_norm,
+                             const h3_gpu_tensor *rope_cos,
+                             const h3_gpu_tensor *rope_sin,
+                             uint32_t sequence, uint32_t query_heads,
+                             uint32_t kv_heads, uint32_t head_dim,
+                             float epsilon) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_head_rms_norm_bf16(h3_gpu *gpu, h3_gpu_tensor *tensor,
+                              const h3_gpu_tensor *weight,
+                              uint32_t sequence, uint32_t heads,
+                              uint32_t head_dim, float epsilon) {
+    if (!tensor || !weight || tensor->dtype != H3_GPU_BF16 || weight->dtype != H3_GPU_BF16 || (size_t)sequence*heads*head_dim > tensor->elements || head_dim > weight->elements) return 0;
+    h3_cu_head_rms_norm_bf16<<<h3_cu_grid((unsigned)sequence*heads), H3_CU_BLOCK>>>((uint16_t *)tensor->device_ptr, (const uint16_t *)weight->device_ptr, sequence, heads, head_dim, epsilon);
+    return 1;
+}
+int h3_gpu_rope_text_bf16(h3_gpu *gpu, h3_gpu_tensor *query,
+                          h3_gpu_tensor *key,
+                          const h3_gpu_tensor *rope_cos_f32,
+                          const h3_gpu_tensor *rope_sin_f32,
+                          uint32_t sequence, uint32_t query_heads,
+                          uint32_t kv_heads, uint32_t head_dim) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_gqa_causal_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                           const h3_gpu_tensor *query,
+                           const h3_gpu_tensor *key,
+                           const h3_gpu_tensor *value,
+                           uint32_t sequence, uint32_t query_heads,
+                           uint32_t kv_heads, uint32_t head_dim,
+                           float scale) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_add_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                    const h3_gpu_tensor *left, const h3_gpu_tensor *right,
+                    uint32_t elements) {
+    if (!output || !left || !right || output->dtype != H3_GPU_BF16 || left->dtype != H3_GPU_BF16 || right->dtype != H3_GPU_BF16 || elements > left->elements || elements > right->elements || elements > output->elements) return 0;
+    h3_cu_add_bf16<<<h3_cu_grid(elements), H3_CU_BLOCK>>>((const uint16_t *)left->device_ptr, (const uint16_t *)right->device_ptr, (uint16_t *)output->device_ptr, elements);
+    return 1;
+}
+int h3_gpu_sub_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                    const h3_gpu_tensor *left, const h3_gpu_tensor *right,
+                    uint32_t elements) {
+    if (!output || !left || !right || output->dtype != H3_GPU_BF16 || left->dtype != H3_GPU_BF16 || right->dtype != H3_GPU_BF16 || elements > left->elements || elements > right->elements || elements > output->elements) return 0;
+    h3_cu_sub_bf16<<<h3_cu_grid(elements), H3_CU_BLOCK>>>((const uint16_t *)left->device_ptr, (const uint16_t *)right->device_ptr, (uint16_t *)output->device_ptr, elements);
+    return 1;
+}
+int h3_gpu_token_pool_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                           const h3_gpu_tensor *input,
+                           size_t input_offset,
+                           h3_gpu_tensor *original,
+                           size_t original_offset,
+                           h3_gpu_tensor *baseline,
+                           size_t baseline_offset,
+                           const h3_gpu_tensor *baseline_indices,
+                           const h3_gpu_tensor *pairs, uint32_t input_rows,
+                           uint32_t rows, uint32_t baseline_rows,
+                           uint32_t width) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_token_pool_adaln_bf16(
+                           h3_gpu *gpu, h3_gpu_tensor *residual,
+                           h3_gpu_tensor *output,
+                           const h3_gpu_tensor *input, size_t input_offset,
+                           h3_gpu_tensor *original, size_t original_offset,
+                           h3_gpu_tensor *baseline, size_t baseline_offset,
+                           const h3_gpu_tensor *baseline_indices,
+                           const h3_gpu_tensor *pairs,
+                           const h3_gpu_tensor *norm_weight,
+                           const h3_gpu_tensor *modulation,
+                           const h3_gpu_tensor *row_map,
+                           uint32_t input_rows, uint32_t rows,
+                           uint32_t baseline_rows, uint32_t width,
+                           uint32_t slots, uint32_t shift_slot,
+                           uint32_t scale_slot, float epsilon) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_token_expand_delta_bf16(
+                           h3_gpu *gpu, h3_gpu_tensor *output,
+                           const h3_gpu_tensor *original,
+                           size_t original_offset,
+                           const h3_gpu_tensor *reduced,
+                           const h3_gpu_tensor *baseline,
+                           size_t baseline_offset,
+                           const h3_gpu_tensor *baseline_indices,
+                           const h3_gpu_tensor *parents, uint32_t rows,
+                           uint32_t reduced_rows, uint32_t baseline_rows,
+                           uint32_t width,
+                           uint32_t exact_prefix_rows,
+                           float update_scale) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_token_expand_adaln_bf16(
+                           h3_gpu *gpu, h3_gpu_tensor *residual,
+                           h3_gpu_tensor *output,
+                           const h3_gpu_tensor *original,
+                           size_t original_offset,
+                           const h3_gpu_tensor *reduced,
+                           const h3_gpu_tensor *baseline,
+                           size_t baseline_offset,
+                           const h3_gpu_tensor *baseline_indices,
+                           const h3_gpu_tensor *parents,
+                           const h3_gpu_tensor *norm_weight,
+                           const h3_gpu_tensor *modulation,
+                           const h3_gpu_tensor *row_map,
+                           uint32_t rows, uint32_t reduced_rows,
+                           uint32_t baseline_rows, uint32_t width,
+                           uint32_t exact_prefix_rows, float update_scale,
+                           uint32_t slots, uint32_t shift_slot,
+                           uint32_t scale_slot, float epsilon) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_euler_bf16(h3_gpu *gpu, h3_gpu_tensor *sample,
+                      size_t sample_offset, const h3_gpu_tensor *last,
+                      const h3_gpu_tensor *previous, uint32_t elements,
+                      float delta, float ratio) { h3_cuda_seterr(gpu); return (int)0; }
+int h3_gpu_silu_mul_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
+                         const h3_gpu_tensor *gate,
+                         const h3_gpu_tensor *up, uint32_t elements) {
+    if (!output || !gate || !up || output->dtype != H3_GPU_BF16 || gate->dtype != H3_GPU_BF16 || up->dtype != H3_GPU_BF16 || elements > gate->elements || elements > up->elements || elements > output->elements) return 0;
+    h3_cu_silu_mul_bf16<<<h3_cu_grid(elements), H3_CU_BLOCK>>>((const uint16_t *)gate->device_ptr, (const uint16_t *)up->device_ptr, (uint16_t *)output->device_ptr, elements);
+    return 1;
+}
diff --git a/h3_cuda.h b/h3_cuda.h
new file mode 100644
index 00000000..955171cf
--- /dev/null
+++ b/h3_cuda.h
@@ -0,0 +1,16 @@
+#ifndef H3_CUDA_H
+#define H3_CUDA_H
+
+#include "h3.h"
+
+#ifdef __cplusplus
+extern "C" {
+#endif
+
+int h3_cuda_probe(h3_device_info *info, char *error, size_t error_size);
+
+#ifdef __cplusplus
+}
+#endif
+
+#endif
diff --git a/h3_ffmpeg.c b/h3_ffmpeg.c
index 66762425..cb730f3e 100644
--- a/h3_ffmpeg.c
+++ b/h3_ffmpeg.c
@@ -12,6 +12,10 @@
 #include <sys/wait.h>
 #include <unistd.h>
 
+#ifndef SSIZE_MAX
+#define SSIZE_MAX ((ssize_t)(~((size_t)0) >> 1))
+#endif
+
 extern char **environ;
 
 static const char *ffmpeg_program(void) {
diff --git a/h3_gpu.h b/h3_gpu.h
index 3a47cc35..2f1db0f0 100644
--- a/h3_gpu.h
+++ b/h3_gpu.h
@@ -1,6 +1,10 @@
 #ifndef H3_GPU_H
 #define H3_GPU_H
 
+#ifdef __cplusplus
+extern "C" {
+#endif
+
 #include <stddef.h>
 #include <stdint.h>
 
@@ -610,4 +614,8 @@ int h3_gpu_silu_mul_bf16(h3_gpu *gpu, h3_gpu_tensor *output,
                          const h3_gpu_tensor *gate,
                          const h3_gpu_tensor *up, uint32_t elements);
 
+#ifdef __cplusplus
+}
+#endif
+
 #endif
diff --git a/h3_host.c b/h3_host.c
index a04a2a0a..d637d92a 100644
--- a/h3_host.c
+++ b/h3_host.c
@@ -1,6 +1,8 @@
 #include "h3_host.h"
 
+#ifndef H3_CUDA
 #include <Accelerate/Accelerate.h>
+#endif
 
 #include <float.h>
 #include <limits.h>
@@ -555,14 +557,15 @@ int h3_resize_rgb24_high_quality(const uint8_t *input, int frames,
         free(pixels);
         return 0;
     }
+    size_t input_frame_bytes = input_area * 3;
+    size_t output_frame_bytes = output_area * 3;
+#ifndef H3_CUDA
     uint8_t *source_argb = malloc(input_area * 4);
     uint8_t *output_argb = malloc(output_area * 4);
     if (!source_argb || !output_argb) {
         free(source_argb); free(output_argb); free(pixels);
         return 0;
     }
-    size_t input_frame_bytes = input_area * 3;
-    size_t output_frame_bytes = output_area * 3;
     vImage_Buffer source_buffer = {
         source_argb, (vImagePixelCount)input_height,
         (vImagePixelCount)input_width, (size_t)input_width * 4
@@ -594,6 +597,39 @@ int h3_resize_rgb24_high_quality(const uint8_t *input, int frames,
         }
     }
     free(source_argb); free(output_argb);
+#else
+    /* Portable bilinear RGB24 resize (CUDA build; Metal keeps vImage). */
+    for (int frame = 0; frame < frames; frame++) {
+        const uint8_t *src = input + (size_t)frame * input_frame_bytes;
+        uint8_t *dst = pixels + (size_t)frame * output_frame_bytes;
+        for (size_t y = 0; y < output_height; y++) {
+            float gy = (output_height == 1) ? 0.0f :
+                       (float)y * (input_height - 1) / (output_height - 1);
+            int y0 = (int)gy;
+            if (y0 > input_height - 2) y0 = input_height - 2;
+            float fy = gy - (float)y0;
+            for (size_t x = 0; x < output_width; x++) {
+                float gx = (output_width == 1) ? 0.0f :
+                           (float)x * (input_width - 1) / (output_width - 1);
+                int x0 = (int)gx;
+                if (x0 > input_width - 2) x0 = input_width - 2;
+                float fx = gx - (float)x0;
+                const uint8_t *p00 = src + ((size_t)y0 * input_width + x0) * 3;
+                const uint8_t *p10 = p00 + 3;
+                const uint8_t *p01 = src + ((size_t)(y0 + 1) * input_width + x0) * 3;
+                const uint8_t *p11 = p01 + 3;
+                uint8_t *o = dst + (y * output_width + x) * 3;
+                for (int c = 0; c < 3; c++) {
+                    float v = (1.0f - fy) * (1.0f - fx) * p00[c] +
+                              (1.0f - fy) * fx * p10[c] +
+                              fy * (1.0f - fx) * p01[c] +
+                              fy * fx * p11[c];
+                    o[c] = (uint8_t)(v + 0.5f);
+                }
+            }
+        }
+    }
+#endif
     *output = pixels;
     return 1;
 }
diff --git a/h3_tokenizer.c b/h3_tokenizer.c
new file mode 100644
index 00000000..3993390e
--- /dev/null
+++ b/h3_tokenizer.c
@@ -0,0 +1,59 @@
+/* h3_tokenizer.c - C stub for the CUDA/Linux build (feat/cuda).
+ *
+ * The Metal build uses h3_tokenizer.m (Objective-C/Foundation). On Linux there
+ * is no Foundation, so the CUDA build compiles this C file against the same
+ * h3_tokenizer.h API. It is a scaffold stub for I1: it satisfies the link so
+ * the host binary builds and `--info` runs. Real tokenizer port is a later
+ * iteration (I20). Generation paths that need tokenization return an error.
+ */
+#include "h3_tokenizer.h"
+
+#include <stdio.h>
+#include <stdlib.h>
+#include <string.h>
+
+struct h3_tokenizer {
+    int unused;
+};
+
+h3_tokenizer *h3_tokenizer_load(const char *tokenizer_json,
+                                char *error, size_t error_size) {
+    (void)tokenizer_json;
+    if (error && error_size) {
+        snprintf(error, error_size,
+                 "CUDA build: tokenizer not yet ported (feat/cuda)");
+    }
+    return NULL;
+}
+
+void h3_tokenizer_free(h3_tokenizer *tokenizer) {
+    (void)tokenizer;
+}
+
+int h3_tokenizer_encode(const h3_tokenizer *tokenizer, const char *utf8,
+                        int pad_empty, uint32_t **ids, size_t *count,
+                        char *error, size_t error_size) {
+    (void)tokenizer; (void)utf8; (void)pad_empty;
+    if (ids) *ids = NULL;
+    if (count) *count = 0;
+    if (error && error_size) {
+        snprintf(error, error_size,
+                 "CUDA build: tokenizer not yet ported (feat/cuda)");
+    }
+    return 0;
+}
+
+void h3_tokenizer_ids_free(uint32_t *ids) {
+    free(ids);
+}
+
+char *h3_tokenizer_decode(const h3_tokenizer *tokenizer,
+                          const uint32_t *ids, size_t count,
+                          char *error, size_t error_size) {
+    (void)tokenizer; (void)ids; (void)count;
+    if (error && error_size) {
+        snprintf(error, error_size,
+                 "CUDA build: tokenizer not yet ported (feat/cuda)");
+    }
+    return NULL;
+}