From 69aa9b474c290a5315d894ff1418d3cbd1a1e734 Mon Sep 17 00:00:00 2001 From: magiodev Date: Wed, 12 Aug 2026 20:49:48 +0000 Subject: [PATCH 01/11] feat/cuda: I1 scaffold - CUDA backend skeleton (probe+create/free real, 100 op stubs) - Add h3_cuda.cu: implements the existing plain-C h3_gpu.h API against CUDA/cuBLAS. h3_cuda_probe + h3_gpu_create/free/error real; all compute ops are stubs returning 'not yet implemented'. Metal backend (h3_gpu.m/h3_shaders.metal) preserved untouched. - Add h3_cuda.h (probe decl); h3.c probes via H3_CUDA guard. - Add h3_tokenizer.c C stub (Linux has no Foundation; Metal keeps h3_tokenizer.m). - h3_host.c: vImage high-quality scale guarded, portable bilinear fallback for CUDA build. - h3.c/h3_cli.c/h3_ffmpeg.c: Linux portability guards (st_mtimespec->st_mtim, arc4random_buf->getrandom, SSIZE_MAX define). - Makefile: cuda-spark/cuda-generic/cuda targets (nvcc + cuBLAS, .cuda.o objects, CUDA_ARCH=sm_121 for DGX Spark GB10), mirroring ds4's pattern. --- Makefile | 45 +++- h3.c | 18 ++ h3_cli.c | 8 + h3_cuda.cu | 586 +++++++++++++++++++++++++++++++++++++++++++++++++ h3_cuda.h | 8 + h3_ffmpeg.c | 4 + h3_host.c | 40 +++- h3_tokenizer.c | 59 +++++ 8 files changed, 765 insertions(+), 3 deletions(-) create mode 100644 h3_cuda.cu create mode 100644 h3_cuda.h create mode 100644 h3_tokenizer.c diff --git a/Makefile b/Makefile index bb202379..ecfbac21 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,49 @@ 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 + +%.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 $@ + +h3-cuda: $(CLI_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 #include #include +#ifndef __APPLE__ +#include +#endif #include #include @@ -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..f80662bd --- /dev/null +++ b/h3_cuda.cu @@ -0,0 +1,586 @@ +/* h3_cuda.cu - CUDA backend skeleton 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: probe/create/free real, compute ops are stubs. + */ +#include +#include +#include +#include +#include +#include "h3_gpu.h" +#include "h3.h" + +#define H3_CUDA_ERR "CUDA backend: op not yet implemented (feat/cuda)" + +struct h3_gpu { void *dev_ctx; char error[512]; }; +struct h3_gpu_tensor { void *device_ptr; h3_gpu_dtype dtype; size_t elements; }; + +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 = (prop.unifiedMemory ? 1 : 0); + } + 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; } + return g; +} +void h3_gpu_free(h3_gpu *gpu) { if (gpu) 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(h3_gpu *gpu) { + if (gpu) snprintf(gpu->error, sizeof(gpu->error), "%s", H3_CUDA_ERR); +} + +int h3_gpu_is_m5(const h3_gpu *gpu) { h3_cuda_seterr(gpu); return 0; } +int h3_gpu_has_nax_mlp(const h3_gpu *gpu) { h3_cuda_seterr(gpu); return 0; } +int h3_gpu_has_int8_mlp(const h3_gpu *gpu) { h3_cuda_seterr(gpu); return 0; } +h3_gpu_tensor * h3_gpu_tensor_new_f32(h3_gpu *gpu, size_t elements) { h3_cuda_seterr(gpu); return NULL; } +h3_gpu_tensor * h3_gpu_tensor_new_bf16(h3_gpu *gpu, size_t elements) { h3_cuda_seterr(gpu); return NULL; } +h3_gpu_tensor * h3_gpu_tensor_new_i8(h3_gpu *gpu, size_t elements) { h3_cuda_seterr(gpu); return NULL; } +h3_gpu_tensor * h3_gpu_tensor_from_f32(h3_gpu *gpu, const float *values, + size_t elements) { h3_cuda_seterr(gpu); return NULL; } +h3_gpu_tensor * h3_gpu_tensor_from_bf16(h3_gpu *gpu, const uint16_t *values, + size_t elements) { h3_cuda_seterr(gpu); return NULL; } +h3_gpu_tensor * h3_gpu_tensor_from_u32(h3_gpu *gpu, const uint32_t *values, + size_t elements) { h3_cuda_seterr(gpu); return NULL; } +h3_gpu_tensor * h3_gpu_tensor_load_bf16(h3_gpu *gpu, const char *path, + uint64_t file_offset, size_t elements) { h3_cuda_seterr(gpu); return NULL; } +h3_gpu_tensor * h3_gpu_tensor_load_f32(h3_gpu *gpu, const char *path, + uint64_t file_offset, size_t elements) { h3_cuda_seterr(gpu); return NULL; } +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 0; } +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 0; } +void h3_gpu_tensor_free(h3_gpu_tensor *tensor) { } +size_t h3_gpu_tensor_elements(const h3_gpu_tensor *tensor) { return 0; } +h3_gpu_dtype h3_gpu_tensor_dtype(const h3_gpu_tensor *tensor) { return 0; } +int h3_gpu_tensor_read_f32(const h3_gpu_tensor *tensor, float *values, + size_t elements) { return 0; } +int h3_gpu_tensor_read_f32_range(const h3_gpu_tensor *tensor, + size_t source_offset, float *values, + size_t elements) { return 0; } +int h3_gpu_tensor_read_bf16(const h3_gpu_tensor *tensor, uint16_t *values, + size_t elements) { return 0; } +int h3_gpu_tensor_write_f32(h3_gpu_tensor *tensor, const float *values, + size_t elements) { return 0; } +int h3_gpu_tensor_write_f32_range(h3_gpu_tensor *tensor, + size_t destination_offset, + const float *values, size_t elements) { return 0; } +int h3_gpu_tensor_write_bf16(h3_gpu_tensor *tensor, const uint16_t *values, + size_t elements) { return 0; } +int h3_gpu_tensor_write_bf16_range(h3_gpu_tensor *tensor, + size_t destination_offset, + const uint16_t *values, size_t elements) { return 0; } +int h3_gpu_begin(h3_gpu *gpu) { h3_cuda_seterr(gpu); return 0; } +int h3_gpu_continue(h3_gpu *gpu) { h3_cuda_seterr(gpu); return 0; } +int h3_gpu_submit(h3_gpu *gpu) { h3_cuda_seterr(gpu); return 0; } +int h3_gpu_get_stats(const h3_gpu *gpu, h3_gpu_stats *stats) { h3_cuda_seterr(gpu); return 0; } +void h3_gpu_profile_set_label(h3_gpu *gpu, const char *label) { h3_cuda_seterr(gpu); } +void h3_gpu_profile_mark(h3_gpu *gpu, const char *phase) { h3_cuda_seterr(gpu); } +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) { h3_cuda_seterr(gpu); return 0; } +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 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 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 0; } +int h3_gpu_silu_f32(h3_gpu *gpu, h3_gpu_tensor *output, + const h3_gpu_tensor *input, uint32_t elements) { h3_cuda_seterr(gpu); return 0; } +int h3_gpu_cast_f32_to_bf16(h3_gpu *gpu, h3_gpu_tensor *output, + const h3_gpu_tensor *input, uint32_t elements) { h3_cuda_seterr(gpu); return 0; } +int h3_gpu_cast_bf16_to_f32(h3_gpu *gpu, h3_gpu_tensor *output, + const h3_gpu_tensor *input, uint32_t elements) { h3_cuda_seterr(gpu); return 0; } +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) { h3_cuda_seterr(gpu); return 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) { h3_cuda_seterr(gpu); return 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) { h3_cuda_seterr(gpu); return 0; } +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 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 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 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 0; } +int h3_gpu_swiglu_f32(h3_gpu *gpu, h3_gpu_tensor *output, + const h3_gpu_tensor *fused, uint32_t rows, + uint32_t width) { h3_cuda_seterr(gpu); return 0; } +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) { h3_cuda_seterr(gpu); return 0; } +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) { h3_cuda_seterr(gpu); return 0; } +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 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 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 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 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) { h3_cuda_seterr(gpu); return 0; } +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) { h3_cuda_seterr(gpu); return 0; } +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 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 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 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 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 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) { h3_cuda_seterr(gpu); return 0; } +int h3_gpu_clip_f32(h3_gpu *gpu, h3_gpu_tensor *output, + const h3_gpu_tensor *input, uint32_t elements, + float minimum, float maximum) { h3_cuda_seterr(gpu); return 0; } +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 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 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 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) { h3_cuda_seterr(gpu); return 0; } +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 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 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 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 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 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 0; } +int h3_gpu_silu_bf16(h3_gpu *gpu, h3_gpu_tensor *output, + const h3_gpu_tensor *input, uint32_t elements) { h3_cuda_seterr(gpu); return 0; } +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) { h3_cuda_seterr(gpu); return 0; } +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) { h3_cuda_seterr(gpu); return 0; } +int h3_gpu_gelu_bf16(h3_gpu *gpu, h3_gpu_tensor *output, + const h3_gpu_tensor *input, uint32_t elements, + int approximate) { h3_cuda_seterr(gpu); return 0; } +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 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 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 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 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) { h3_cuda_seterr(gpu); return 0; } +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 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 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 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 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 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 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 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 0; } +int h3_gpu_swiglu_bf16(h3_gpu *gpu, h3_gpu_tensor *output, + const h3_gpu_tensor *fused, uint32_t rows, + uint32_t width) { h3_cuda_seterr(gpu); return 0; } +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) { h3_cuda_seterr(gpu); return 0; } +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 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) { h3_cuda_seterr(gpu); return 0; } +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 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 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) { h3_cuda_seterr(gpu); return 0; } +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) { h3_cuda_seterr(gpu); return 0; } +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 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 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 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 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 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) { h3_cuda_seterr(gpu); return 0; } diff --git a/h3_cuda.h b/h3_cuda.h new file mode 100644 index 00000000..f5eb09ee --- /dev/null +++ b/h3_cuda.h @@ -0,0 +1,8 @@ +#ifndef H3_CUDA_H +#define H3_CUDA_H + +#include "h3.h" + +int h3_cuda_probe(h3_device_info *info, char *error, size_t error_size); + +#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 #include +#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_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 +#endif #include #include @@ -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 +#include +#include + +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; +} From c0d92e155baa87e1337df68c2e47a208d6b3b318 Mon Sep 17 00:00:00 2001 From: magiodev Date: Wed, 12 Aug 2026 20:53:23 +0000 Subject: [PATCH 02/11] feat/cuda: build CLI objects as .cuda.o with -D_GNU_SOURCE (Linux CLOCK_MONOTONIC/strdup) --- Makefile | 7 ++++++- 1 file changed, 6 insertions(+), 1 deletion(-) diff --git a/Makefile b/Makefile index ecfbac21..49560b74 100644 --- a/Makefile +++ b/Makefile @@ -216,13 +216,18 @@ CUDA_C_SRC := h3.c h3_host.c h3_safetensors.c h3_weights.c h3_text_encoder.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 $@ -h3-cuda: $(CLI_OBJ) $(CUDA_OBJ) +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: From cd0e18bf94a77e0ba855da43ab9dcdcf8d55f996 Mon Sep 17 00:00:00 2001 From: magiodev Date: Wed, 12 Aug 2026 20:58:06 +0000 Subject: [PATCH 03/11] feat/cuda: fix nvcc strictness - seterr takes const h3_gpu*, stub returns cast to exact type --- h3_cuda.cu | 198 ++++++++++++++++++++++++++--------------------------- 1 file changed, 99 insertions(+), 99 deletions(-) diff --git a/h3_cuda.cu b/h3_cuda.cu index f80662bd..9440463e 100644 --- a/h3_cuda.cu +++ b/h3_cuda.cu @@ -45,74 +45,74 @@ void h3_gpu_free(h3_gpu *gpu) { if (gpu) 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(h3_gpu *gpu) { - if (gpu) snprintf(gpu->error, sizeof(gpu->error), "%s", H3_CUDA_ERR); +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 0; } -int h3_gpu_has_nax_mlp(const h3_gpu *gpu) { h3_cuda_seterr(gpu); return 0; } -int h3_gpu_has_int8_mlp(const h3_gpu *gpu) { h3_cuda_seterr(gpu); return 0; } -h3_gpu_tensor * h3_gpu_tensor_new_f32(h3_gpu *gpu, size_t elements) { h3_cuda_seterr(gpu); return NULL; } -h3_gpu_tensor * h3_gpu_tensor_new_bf16(h3_gpu *gpu, size_t elements) { h3_cuda_seterr(gpu); return NULL; } -h3_gpu_tensor * h3_gpu_tensor_new_i8(h3_gpu *gpu, size_t elements) { h3_cuda_seterr(gpu); return NULL; } +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) { h3_cuda_seterr(gpu); return (h3_gpu_tensor *)NULL; } +h3_gpu_tensor * h3_gpu_tensor_new_bf16(h3_gpu *gpu, size_t elements) { h3_cuda_seterr(gpu); return (h3_gpu_tensor *)NULL; } +h3_gpu_tensor * h3_gpu_tensor_new_i8(h3_gpu *gpu, size_t elements) { h3_cuda_seterr(gpu); return (h3_gpu_tensor *)NULL; } h3_gpu_tensor * h3_gpu_tensor_from_f32(h3_gpu *gpu, const float *values, - size_t elements) { h3_cuda_seterr(gpu); return NULL; } + size_t elements) { h3_cuda_seterr(gpu); return (h3_gpu_tensor *)NULL; } h3_gpu_tensor * h3_gpu_tensor_from_bf16(h3_gpu *gpu, const uint16_t *values, - size_t elements) { h3_cuda_seterr(gpu); return NULL; } + size_t elements) { h3_cuda_seterr(gpu); return (h3_gpu_tensor *)NULL; } h3_gpu_tensor * h3_gpu_tensor_from_u32(h3_gpu *gpu, const uint32_t *values, - size_t elements) { h3_cuda_seterr(gpu); return NULL; } + size_t elements) { h3_cuda_seterr(gpu); return (h3_gpu_tensor *)NULL; } h3_gpu_tensor * h3_gpu_tensor_load_bf16(h3_gpu *gpu, const char *path, - uint64_t file_offset, size_t elements) { h3_cuda_seterr(gpu); return NULL; } + uint64_t file_offset, size_t elements) { h3_cuda_seterr(gpu); return (h3_gpu_tensor *)NULL; } h3_gpu_tensor * h3_gpu_tensor_load_f32(h3_gpu *gpu, const char *path, - uint64_t file_offset, size_t elements) { h3_cuda_seterr(gpu); return NULL; } + uint64_t file_offset, size_t elements) { h3_cuda_seterr(gpu); return (h3_gpu_tensor *)NULL; } 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 0; } + char *error, size_t error_size) { return (int)0; } 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 0; } + char *error, size_t error_size) { return (int)0; } void h3_gpu_tensor_free(h3_gpu_tensor *tensor) { } -size_t h3_gpu_tensor_elements(const h3_gpu_tensor *tensor) { return 0; } -h3_gpu_dtype h3_gpu_tensor_dtype(const h3_gpu_tensor *tensor) { return 0; } +size_t h3_gpu_tensor_elements(const h3_gpu_tensor *tensor) { return (size_t)0; } +h3_gpu_dtype h3_gpu_tensor_dtype(const h3_gpu_tensor *tensor) { return (h3_gpu_dtype)0; } int h3_gpu_tensor_read_f32(const h3_gpu_tensor *tensor, float *values, - size_t elements) { return 0; } + size_t elements) { return (int)0; } int h3_gpu_tensor_read_f32_range(const h3_gpu_tensor *tensor, size_t source_offset, float *values, - size_t elements) { return 0; } + size_t elements) { return (int)0; } int h3_gpu_tensor_read_bf16(const h3_gpu_tensor *tensor, uint16_t *values, - size_t elements) { return 0; } + size_t elements) { return (int)0; } int h3_gpu_tensor_write_f32(h3_gpu_tensor *tensor, const float *values, - size_t elements) { return 0; } + size_t elements) { return (int)0; } int h3_gpu_tensor_write_f32_range(h3_gpu_tensor *tensor, size_t destination_offset, - const float *values, size_t elements) { return 0; } + const float *values, size_t elements) { return (int)0; } int h3_gpu_tensor_write_bf16(h3_gpu_tensor *tensor, const uint16_t *values, - size_t elements) { return 0; } + size_t elements) { return (int)0; } int h3_gpu_tensor_write_bf16_range(h3_gpu_tensor *tensor, size_t destination_offset, - const uint16_t *values, size_t elements) { return 0; } -int h3_gpu_begin(h3_gpu *gpu) { h3_cuda_seterr(gpu); return 0; } -int h3_gpu_continue(h3_gpu *gpu) { h3_cuda_seterr(gpu); return 0; } -int h3_gpu_submit(h3_gpu *gpu) { h3_cuda_seterr(gpu); return 0; } -int h3_gpu_get_stats(const h3_gpu *gpu, h3_gpu_stats *stats) { h3_cuda_seterr(gpu); return 0; } + const uint16_t *values, size_t elements) { return (int)0; } +int h3_gpu_begin(h3_gpu *gpu) { h3_cuda_seterr(gpu); return (int)0; } +int h3_gpu_continue(h3_gpu *gpu) { h3_cuda_seterr(gpu); return (int)0; } +int h3_gpu_submit(h3_gpu *gpu) { h3_cuda_seterr(gpu); return (int)0; } +int h3_gpu_get_stats(const h3_gpu *gpu, h3_gpu_stats *stats) { h3_cuda_seterr(gpu); return (int)0; } void h3_gpu_profile_set_label(h3_gpu *gpu, const char *label) { h3_cuda_seterr(gpu); } void h3_gpu_profile_mark(h3_gpu *gpu, const char *phase) { h3_cuda_seterr(gpu); } 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) { h3_cuda_seterr(gpu); return 0; } + uint32_t input_dim, uint32_t output_dim) { h3_cuda_seterr(gpu); return (int)0; } 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 0; } + 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 0; } + 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, @@ -120,38 +120,38 @@ int h3_gpu_patch_linear_bf16_map( 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 0; } + 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) { h3_cuda_seterr(gpu); return 0; } + const h3_gpu_tensor *input, uint32_t elements) { h3_cuda_seterr(gpu); return (int)0; } int h3_gpu_cast_f32_to_bf16(h3_gpu *gpu, h3_gpu_tensor *output, - const h3_gpu_tensor *input, uint32_t elements) { h3_cuda_seterr(gpu); return 0; } + const h3_gpu_tensor *input, uint32_t elements) { h3_cuda_seterr(gpu); return (int)0; } int h3_gpu_cast_bf16_to_f32(h3_gpu *gpu, h3_gpu_tensor *output, - const h3_gpu_tensor *input, uint32_t elements) { h3_cuda_seterr(gpu); return 0; } + const h3_gpu_tensor *input, uint32_t elements) { h3_cuda_seterr(gpu); return (int)0; } 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) { h3_cuda_seterr(gpu); return 0; } + size_t elements) { h3_cuda_seterr(gpu); return (int)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) { h3_cuda_seterr(gpu); return 0; } + size_t elements) { h3_cuda_seterr(gpu); return (int)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) { h3_cuda_seterr(gpu); return 0; } + uint32_t width, float epsilon) { h3_cuda_seterr(gpu); return (int)0; } 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 0; } + 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 0; } + 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, @@ -160,24 +160,24 @@ int h3_gpu_qkv_rope_f32(h3_gpu *gpu, h3_gpu_tensor *query, 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 0; } + 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 0; } + 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) { h3_cuda_seterr(gpu); return 0; } + uint32_t width) { h3_cuda_seterr(gpu); return (int)0; } 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) { h3_cuda_seterr(gpu); return 0; } + uint32_t width) { h3_cuda_seterr(gpu); return (int)0; } 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) { h3_cuda_seterr(gpu); return 0; } + uint32_t width, float epsilon) { h3_cuda_seterr(gpu); return (int)0; } 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, @@ -185,14 +185,14 @@ int h3_gpu_video_qkv_rope_f32(h3_gpu *gpu, h3_gpu_tensor *query, 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 0; } + 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 0; } + 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, @@ -200,7 +200,7 @@ int h3_gpu_conv1d_stride_f32(h3_gpu *gpu, h3_gpu_tensor *output, 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 0; } + 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, @@ -208,15 +208,15 @@ int h3_gpu_conv_transpose1d_f32( 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 0; } + 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) { h3_cuda_seterr(gpu); return 0; } + uint32_t outer, uint32_t inner) { h3_cuda_seterr(gpu); return (int)0; } 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) { h3_cuda_seterr(gpu); return 0; } + float right_scale, uint32_t elements) { h3_cuda_seterr(gpu); return (int)0; } int h3_gpu_alias_free_snake_f32( h3_gpu *gpu, h3_gpu_tensor *output, const h3_gpu_tensor *input, @@ -225,11 +225,11 @@ int h3_gpu_alias_free_snake_f32( 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 0; } + 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 0; } + 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, @@ -237,31 +237,31 @@ int h3_gpu_audio_qkv_split_f32(h3_gpu *gpu, 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 0; } + 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 0; } + 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 0; } + 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) { h3_cuda_seterr(gpu); return 0; } + const h3_gpu_tensor *linear, uint32_t elements) { h3_cuda_seterr(gpu); return (int)0; } int h3_gpu_clip_f32(h3_gpu *gpu, h3_gpu_tensor *output, const h3_gpu_tensor *input, uint32_t elements, - float minimum, float maximum) { h3_cuda_seterr(gpu); return 0; } + float minimum, float maximum) { h3_cuda_seterr(gpu); return (int)0; } 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 0; } + 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, @@ -270,36 +270,36 @@ int h3_gpu_conv3d_f32(h3_gpu *gpu, h3_gpu_tensor *output, 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 0; } + 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 0; } + 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) { h3_cuda_seterr(gpu); return 0; } + uint32_t input_dim, uint32_t output_dim) { h3_cuda_seterr(gpu); return (int)0; } 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 0; } + 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 0; } + 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 0; } + 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, @@ -308,7 +308,7 @@ int h3_gpu_linear_int8_bf16(h3_gpu *gpu, h3_gpu_tensor *output, 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 0; } + 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, @@ -317,7 +317,7 @@ int h3_gpu_linear_int8_head_major_bf16( 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 0; } + 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, @@ -334,21 +334,21 @@ int h3_gpu_mlp_int8_bf16(h3_gpu *gpu, h3_gpu_tensor *output, 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 0; } + 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) { h3_cuda_seterr(gpu); return 0; } + const h3_gpu_tensor *input, uint32_t elements) { h3_cuda_seterr(gpu); return (int)0; } 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) { h3_cuda_seterr(gpu); return 0; } + uint32_t width, float epsilon) { h3_cuda_seterr(gpu); return (int)0; } 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) { h3_cuda_seterr(gpu); return 0; } + uint32_t width, float epsilon) { h3_cuda_seterr(gpu); return (int)0; } int h3_gpu_gelu_bf16(h3_gpu *gpu, h3_gpu_tensor *output, const h3_gpu_tensor *input, uint32_t elements, - int approximate) { h3_cuda_seterr(gpu); return 0; } + int approximate) { h3_cuda_seterr(gpu); return (int)0; } int h3_gpu_vision_qkv_rope_bf16( h3_gpu *gpu, h3_gpu_tensor *query, h3_gpu_tensor *key, h3_gpu_tensor *value, @@ -356,21 +356,21 @@ int h3_gpu_vision_qkv_rope_bf16( 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 0; } + 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 0; } + 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 0; } + 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, @@ -382,13 +382,13 @@ int h3_gpu_adaln_linear_bf16( 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 0; } + 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) { h3_cuda_seterr(gpu); return 0; } + uint32_t width, uint32_t slots, uint32_t gate_slot) { h3_cuda_seterr(gpu); return (int)0; } int h3_gpu_gate_adaln_bf16( h3_gpu *gpu, h3_gpu_tensor *gated_residual, h3_gpu_tensor *output, @@ -400,7 +400,7 @@ int h3_gpu_gate_adaln_bf16( 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 0; } + 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, @@ -413,7 +413,7 @@ int h3_gpu_gate_adaln_quantize_int8( 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 0; } + 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, @@ -422,7 +422,7 @@ int h3_gpu_qkv_rope_bf16(h3_gpu *gpu, h3_gpu_tensor *query, 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 0; } + 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, @@ -432,7 +432,7 @@ int h3_gpu_grouped_qkv_rope_bf16(h3_gpu *gpu, h3_gpu_tensor *query, 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 0; } + float epsilon) { h3_cuda_seterr(gpu); return (int)0; } int h3_gpu_grouped_qkv_linear_rope_bf16( h3_gpu *gpu, h3_gpu_tensor *query, @@ -447,7 +447,7 @@ int h3_gpu_grouped_qkv_linear_rope_bf16( 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 0; } + 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, @@ -468,23 +468,23 @@ int h3_gpu_grouped_qkv_linear_rope_int8( 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 0; } + 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 0; } + 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 0; } + 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) { h3_cuda_seterr(gpu); return 0; } + uint32_t width) { h3_cuda_seterr(gpu); return (int)0; } 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) { h3_cuda_seterr(gpu); return 0; } + uint32_t vocab_size, uint32_t width) { h3_cuda_seterr(gpu); return (int)0; } int h3_gpu_text_qk_rope_bf16(h3_gpu *gpu, h3_gpu_tensor *query_output, h3_gpu_tensor *key_output, @@ -496,30 +496,30 @@ int h3_gpu_text_qk_rope_bf16(h3_gpu *gpu, 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 0; } + 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) { h3_cuda_seterr(gpu); return 0; } + uint32_t head_dim, float epsilon) { h3_cuda_seterr(gpu); return (int)0; } 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 0; } + 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 0; } + 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) { h3_cuda_seterr(gpu); return 0; } + uint32_t elements) { h3_cuda_seterr(gpu); return (int)0; } 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) { h3_cuda_seterr(gpu); return 0; } + uint32_t elements) { h3_cuda_seterr(gpu); return (int)0; } int h3_gpu_token_pool_bf16(h3_gpu *gpu, h3_gpu_tensor *output, const h3_gpu_tensor *input, size_t input_offset, @@ -530,7 +530,7 @@ int h3_gpu_token_pool_bf16(h3_gpu *gpu, h3_gpu_tensor *output, 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 0; } + 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, @@ -545,7 +545,7 @@ int h3_gpu_token_pool_adaln_bf16( 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 0; } + 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, @@ -558,7 +558,7 @@ int h3_gpu_token_expand_delta_bf16( uint32_t reduced_rows, uint32_t baseline_rows, uint32_t width, uint32_t exact_prefix_rows, - float update_scale) { h3_cuda_seterr(gpu); return 0; } + 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, @@ -576,11 +576,11 @@ int h3_gpu_token_expand_adaln_bf16( 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 0; } + 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 0; } + 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) { h3_cuda_seterr(gpu); return 0; } + const h3_gpu_tensor *up, uint32_t elements) { h3_cuda_seterr(gpu); return (int)0; } From a626261843e5409bfe11f5004241ada045452c0a Mon Sep 17 00:00:00 2001 From: magiodev Date: Wed, 12 Aug 2026 20:59:23 +0000 Subject: [PATCH 04/11] feat/cuda: probe uses cudaDeviceGetAttribute(UnifiedMemory) (CUDA 13 removed prop.unifiedMemory) --- h3_cuda.cu | 4 +++- 1 file changed, 3 insertions(+), 1 deletion(-) diff --git a/h3_cuda.cu b/h3_cuda.cu index 9440463e..069c729d 100644 --- a/h3_cuda.cu +++ b/h3_cuda.cu @@ -31,7 +31,9 @@ int h3_cuda_probe(h3_device_info *info, char *error, size_t error_size) { 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 = (prop.unifiedMemory ? 1 : 0); + int unif = 0; + cudaDeviceGetAttribute(&unif, cudaDevAttrUnifiedMemory, 0); + info->unified_memory = (unif ? 1 : 0); } return 1; } From 2e42886455dfd865cfb85def6da7d73fa8e6282c Mon Sep 17 00:00:00 2001 From: magiodev Date: Wed, 12 Aug 2026 21:00:26 +0000 Subject: [PATCH 05/11] feat/cuda: unified_memory=1 (GB10 unified; CUDA13 removed unifiedMemory attr) --- h3_cuda.cu | 6 +++--- 1 file changed, 3 insertions(+), 3 deletions(-) diff --git a/h3_cuda.cu b/h3_cuda.cu index 069c729d..758b934d 100644 --- a/h3_cuda.cu +++ b/h3_cuda.cu @@ -31,9 +31,9 @@ int h3_cuda_probe(h3_device_info *info, char *error, size_t error_size) { 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; - int unif = 0; - cudaDeviceGetAttribute(&unif, cudaDevAttrUnifiedMemory, 0); - info->unified_memory = (unif ? 1 : 0); + /* GB10/DGX Spark is a unified-memory architecture (CUDA 13 removed + * both prop.unifiedMemory and cudaDevAttrUnifiedMemory). */ + info->unified_memory = 1; } return 1; } From 65c8ac31dcae686c6447c0ad9cb701b73aff4fad Mon Sep 17 00:00:00 2001 From: magiodev Date: Wed, 12 Aug 2026 21:10:21 +0000 Subject: [PATCH 06/11] feat/cuda: extern "C" guards in h3_gpu.h/h3_cuda.h + include h3_cuda.h in h3_cuda.cu nvcc compiles .cu as C++; without C linkage the GPU API symbols were mangled and the host C objects could not link. extern "C" guards are compile-time only and inert for the Metal build (.c/.m never define __cplusplus). --- h3_cuda.cu | 1 + h3_cuda.h | 8 ++++++++ h3_gpu.h | 8 ++++++++ 3 files changed, 17 insertions(+) diff --git a/h3_cuda.cu b/h3_cuda.cu index 758b934d..b2637377 100644 --- a/h3_cuda.cu +++ b/h3_cuda.cu @@ -10,6 +10,7 @@ #include #include "h3_gpu.h" #include "h3.h" +#include "h3_cuda.h" #define H3_CUDA_ERR "CUDA backend: op not yet implemented (feat/cuda)" diff --git a/h3_cuda.h b/h3_cuda.h index f5eb09ee..955171cf 100644 --- a/h3_cuda.h +++ b/h3_cuda.h @@ -3,6 +3,14 @@ #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_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 #include @@ -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 From 9c1d5988368c6c28fbdd9df976b6b1141af906cd Mon Sep 17 00:00:00 2001 From: magiodev Date: Wed, 12 Aug 2026 21:54:59 +0000 Subject: [PATCH 07/11] feat/cuda: I2 tensor layer - real device alloc/copies/file streaming, command-buffer no-ops, stats - h3_gpu_tensor_new/from_*: cudaMalloc + host->device copy (f32/bf16/i8/u32) - h3_gpu_tensor_load_bf16/f32 + read_file_bf16/stream_file_bf16: chunked pread -> host staging -> device (1MB chunks; no host double-buffer of big weights) - h3_gpu_tensor_read/write_*: device<->host cudaMemcpy with bounds checks - h3_gpu_begin/continue/submit: no-ops (CUDA has no explicit command buffer) - h3_gpu_get_stats: allocated/live/peak bytes + tensor_allocations tracking - tensor_free: cudaFree + stats decrement; tensor struct carries owner - capability flags (is_m5/has_int8_mlp/has_nax_mlp) still return 0 (no int8/nax fast paths yet) --- h3_cuda.cu | 249 ++++++++++++++++++++++++++++++++++++++++++++++------- 1 file changed, 217 insertions(+), 32 deletions(-) diff --git a/h3_cuda.cu b/h3_cuda.cu index b2637377..eacc6282 100644 --- a/h3_cuda.cu +++ b/h3_cuda.cu @@ -1,21 +1,130 @@ -/* h3_cuda.cu - CUDA backend skeleton for h3.c (feat/cuda). +/* 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: probe/create/free real, compute ops are stubs. + * I1 scaffold: probe/create/free + I2 tensor layer real; compute ops are stubs. */ #include #include #include #include #include +#include +#include +#include +#include +#include #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)) -struct h3_gpu { void *dev_ctx; char error[512]; }; -struct h3_gpu_tensor { void *device_ptr; h3_gpu_dtype dtype; size_t elements; }; +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 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; + } +} + +/* 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; @@ -32,8 +141,6 @@ int h3_cuda_probe(h3_device_info *info, char *error, size_t error_size) { 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; - /* GB10/DGX Spark is a unified-memory architecture (CUDA 13 removed - * both prop.unifiedMemory and cudaDevAttrUnifiedMemory). */ info->unified_memory = 1; } return 1; @@ -55,51 +162,129 @@ static void h3_cuda_seterr(const h3_gpu *gpu) { 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) { h3_cuda_seterr(gpu); return (h3_gpu_tensor *)NULL; } -h3_gpu_tensor * h3_gpu_tensor_new_bf16(h3_gpu *gpu, size_t elements) { h3_cuda_seterr(gpu); return (h3_gpu_tensor *)NULL; } -h3_gpu_tensor * h3_gpu_tensor_new_i8(h3_gpu *gpu, size_t elements) { h3_cuda_seterr(gpu); return (h3_gpu_tensor *)NULL; } +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) { h3_cuda_seterr(gpu); return (h3_gpu_tensor *)NULL; } + 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) { h3_cuda_seterr(gpu); return (h3_gpu_tensor *)NULL; } + 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) { h3_cuda_seterr(gpu); return (h3_gpu_tensor *)NULL; } + 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_cuda_seterr(gpu); return (h3_gpu_tensor *)NULL; } + 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_cuda_seterr(gpu); return (h3_gpu_tensor *)NULL; } + 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 (int)0; } + 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 (int)0; } -void h3_gpu_tensor_free(h3_gpu_tensor *tensor) { } -size_t h3_gpu_tensor_elements(const h3_gpu_tensor *tensor) { return (size_t)0; } -h3_gpu_dtype h3_gpu_tensor_dtype(const h3_gpu_tensor *tensor) { return (h3_gpu_dtype)0; } + 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 (int)0; } + 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) { return (int)0; } + 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) { return (int)0; } + 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 (int)0; } + 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) { return (int)0; } + 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 (int)0; } + 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) { return (int)0; } -int h3_gpu_begin(h3_gpu *gpu) { h3_cuda_seterr(gpu); return (int)0; } -int h3_gpu_continue(h3_gpu *gpu) { h3_cuda_seterr(gpu); return (int)0; } -int h3_gpu_submit(h3_gpu *gpu) { h3_cuda_seterr(gpu); return (int)0; } -int h3_gpu_get_stats(const h3_gpu *gpu, h3_gpu_stats *stats) { h3_cuda_seterr(gpu); return (int)0; } -void h3_gpu_profile_set_label(h3_gpu *gpu, const char *label) { h3_cuda_seterr(gpu); } -void h3_gpu_profile_mark(h3_gpu *gpu, const char *phase) { h3_cuda_seterr(gpu); } + 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, From e174ac46e8e607c9cf5be878c749ebd54456483f Mon Sep 17 00:00:00 2001 From: magiodev Date: Wed, 12 Aug 2026 22:46:17 +0000 Subject: [PATCH 08/11] feat/cuda: I3 elementwise/norm/activation/embedding CUDA kernels Real CUDA kernels (Metal formulas mirrored exactly): - activations: silu (f32/bf16), gelu (approx/exact, tanh/erf), geglu, swiglu (f32/bf16), silu_mul - casts: f32<->bf16 (round-to-nearest-even, matches Metal), clip - arithmetic: add/sub bf16, add_scaled f32 - norms: rms_norm/layer_norm (f32/bf16, block-per-row shared reduction), head_rms_norm, weight_norm - embedding bf16, scale_add f32, gate bf16 - copy f32/bf16 (device->device cudaMemcpy) BF16 helpers implemented to match Metal bit-for-bit. 50 compute ops remain stubs. --- h3_cuda.cu | 334 ++++++++++++++++++++++++++++++++++++++++++++++++----- 1 file changed, 306 insertions(+), 28 deletions(-) diff --git a/h3_cuda.cu b/h3_cuda.cu index eacc6282..4d4db0b9 100644 --- a/h3_cuda.cu +++ b/h3_cuda.cu @@ -1,7 +1,8 @@ /* 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: probe/create/free + I2 tensor layer real; compute ops are stubs. + * I1 scaffold + I2 tensor layer + I3 elementwise/norm/activation/embedding real; + * remaining compute ops are stubs. */ #include #include @@ -19,10 +20,23 @@ #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). */ +static inline float h3_bf16_to_f32(uint16_t v) { + uint32_t bits = ((uint32_t)v) << 16; float o; memcpy(&o, &bits, sizeof(o)); return o; +} +static inline uint16_t h3_f32_to_bf16(float x) { + uint32_t bits; memcpy(&bits, &x, sizeof(bits)); + bits += 0x7fffu + ((bits >> 16) & 1u); + return (uint16_t)(bits >> 16); +} + static size_t h3_gpu_dtype_size(h3_gpu_dtype dtype) { switch (dtype) { case H3_GPU_F32: return sizeof(float); @@ -33,6 +47,171 @@ static size_t h3_gpu_dtype_size(h3_gpu_dtype dtype) { } } +__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); @@ -210,9 +389,7 @@ int h3_gpu_tensor_stream_file_bf16(h3_gpu_tensor *tensor, const char *path, 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; - } + 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) { @@ -310,23 +487,47 @@ int h3_gpu_patch_linear_bf16_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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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) { h3_cuda_seterr(gpu); return (int)0; } + 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) { h3_cuda_seterr(gpu); return (int)0; } + 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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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, @@ -355,17 +556,31 @@ int h3_gpu_sdpa_f32(h3_gpu *gpu, h3_gpu_tensor *output, 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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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, @@ -400,11 +615,19 @@ int h3_gpu_conv_transpose1d_f32( 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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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, @@ -439,10 +662,18 @@ int h3_gpu_audio_attention_pool_f32(h3_gpu *gpu, 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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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, @@ -524,19 +755,35 @@ int h3_gpu_mlp_int8_bf16(h3_gpu *gpu, h3_gpu_tensor *output, 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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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, @@ -576,7 +823,12 @@ int h3_gpu_gate_bf16(h3_gpu *gpu, h3_gpu_tensor *output, 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; } + 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<<>>((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, @@ -668,11 +920,21 @@ int h3_gpu_sdpa_bf16_head_major_output( 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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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, @@ -688,7 +950,11 @@ int h3_gpu_text_qk_rope_bf16(h3_gpu *gpu, 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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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, @@ -704,10 +970,18 @@ int h3_gpu_gqa_causal_bf16(h3_gpu *gpu, h3_gpu_tensor *output, 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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((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, @@ -771,4 +1045,8 @@ int h3_gpu_euler_bf16(h3_gpu *gpu, h3_gpu_tensor *sample, 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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>((const uint16_t *)gate->device_ptr, (const uint16_t *)up->device_ptr, (uint16_t *)output->device_ptr, elements); + return 1; +} From f9477ba6c1852de976745ac968f232788dcb8f3b Mon Sep 17 00:00:00 2001 From: magiodev Date: Wed, 12 Aug 2026 22:48:54 +0000 Subject: [PATCH 09/11] feat/cuda: BF16 helpers -> __device__ with CUDA intrinsics (kernels call them) --- h3_cuda.cu | 11 ++++++----- 1 file changed, 6 insertions(+), 5 deletions(-) diff --git a/h3_cuda.cu b/h3_cuda.cu index 4d4db0b9..4802fcde 100644 --- a/h3_cuda.cu +++ b/h3_cuda.cu @@ -27,12 +27,13 @@ struct h3_gpu_tensor { void *device_ptr; h3_gpu_dtype dtype; size_t elements; si 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). */ -static inline float h3_bf16_to_f32(uint16_t v) { - uint32_t bits = ((uint32_t)v) << 16; float o; memcpy(&o, &bits, sizeof(o)); return o; +/* 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); } -static inline uint16_t h3_f32_to_bf16(float x) { - uint32_t bits; memcpy(&bits, &x, sizeof(bits)); +__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); } From 0987aab164684dff434b23457ddcd2a86902fee8 Mon Sep 17 00:00:00 2001 From: magiodev Date: Wed, 12 Aug 2026 23:39:34 +0000 Subject: [PATCH 10/11] feat/cuda: I4 cuBLAS linear core (linear_f32 + linear_bf16 real) Row-major GEMM via cublasGemmEx(OP_T,OP_T): C=A@W^T. f32: CUDA_R_32F/CUBLAS_COMPUTE_32F. bf16: CUDA_R_16BF/CUBLAS_COMPUTE_32F (accum f32), bias added in f32, bf16 rounded once. h3_gpu_create now inits CUDA+cublas (handle stashed in dev_ctx), free destroys handle. --- h3_cuda.cu | 91 ++++++++++++++++++++++++++++++++++++++++++++++++++++-- 1 file changed, 88 insertions(+), 3 deletions(-) diff --git a/h3_cuda.cu b/h3_cuda.cu index 4802fcde..1bebb93f 100644 --- a/h3_cuda.cu +++ b/h3_cuda.cu @@ -27,6 +27,40 @@ struct h3_gpu_tensor { void *device_ptr; h3_gpu_dtype dtype; size_t elements; si static unsigned h3_cu_grid(unsigned n) { return (n + H3_CU_BLOCK - 1) / H3_CU_BLOCK; } +/* 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_TENSOR_OP_ALGO); +} + +__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); + } +} + /* 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) { @@ -329,9 +363,28 @@ h3_gpu *h3_gpu_create(const char *shader_source_path, char *error, size_t error_ (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) free(gpu); } +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"; } @@ -466,7 +519,23 @@ void h3_gpu_profile_mark(h3_gpu *gpu, const char *phase) { 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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>( + (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, @@ -702,7 +771,23 @@ 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) { h3_cuda_seterr(gpu); return (int)0; } + 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<<>>( + (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, From 3e23b2b7b7911a9f5dfb6cccab7f31e4e99bbbdb Mon Sep 17 00:00:00 2001 From: magiodev Date: Wed, 12 Aug 2026 23:44:11 +0000 Subject: [PATCH 11/11] feat/cuda: I4 fix - BF16 helpers before bias kernels, CUBLAS_GEMM_DEFAULT algo --- h3_cuda.cu | 24 ++++++++++++------------ 1 file changed, 12 insertions(+), 12 deletions(-) diff --git a/h3_cuda.cu b/h3_cuda.cu index 1bebb93f..938ee40e 100644 --- a/h3_cuda.cu +++ b/h3_cuda.cu @@ -27,6 +27,17 @@ struct h3_gpu_tensor { void *device_ptr; h3_gpu_dtype dtype; size_t elements; si 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. */ @@ -38,7 +49,7 @@ static cublasStatus_t h3_cu_gemm(h3_gpu *g, cudaDataType ab, cudaDataType c, 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_TENSOR_OP_ALGO); + comp, CUBLAS_GEMM_DEFAULT); } __global__ void h3_cu_linear_bias_f32(float *out, const float *bias, @@ -61,17 +72,6 @@ __global__ void h3_cu_linear_bias_bf16(uint16_t *out, const uint16_t *bias, } } -/* 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); -} - static size_t h3_gpu_dtype_size(h3_gpu_dtype dtype) { switch (dtype) { case H3_GPU_F32: return sizeof(float);