diff --git a/CMakeLists.txt b/CMakeLists.txt index dd84de80..0e46556f 100644 --- a/CMakeLists.txt +++ b/CMakeLists.txt @@ -49,6 +49,7 @@ if (NOT MSVC) endif() set(ENGINE_DEFAULT_ENABLE_CUDA OFF) +set(ENGINE_DEFAULT_ENABLE_HIP OFF) set(ENGINE_DEFAULT_ENABLE_VULKAN OFF) set(ENGINE_DEFAULT_ENABLE_METAL OFF) set(ENGINE_DEFAULT_ENABLE_LLAMAFILE ON) @@ -60,6 +61,9 @@ endif() if (DEFINED GGML_CUDA) set(ENGINE_DEFAULT_ENABLE_CUDA ${GGML_CUDA}) endif() +if (DEFINED GGML_HIP) + set(ENGINE_DEFAULT_ENABLE_HIP ${GGML_HIP}) +endif() if (DEFINED GGML_VULKAN) set(ENGINE_DEFAULT_ENABLE_VULKAN ${GGML_VULKAN}) endif() @@ -74,6 +78,7 @@ if (DEFINED GGML_CUDA_GRAPHS) endif() option(ENGINE_ENABLE_CUDA "Build ggml with CUDA backend support" ${ENGINE_DEFAULT_ENABLE_CUDA}) +option(ENGINE_ENABLE_HIP "Build ggml with HIP/ROCm backend support" ${ENGINE_DEFAULT_ENABLE_HIP}) option(ENGINE_ENABLE_VULKAN "Build ggml with Vulkan backend support" ${ENGINE_DEFAULT_ENABLE_VULKAN}) option(ENGINE_ENABLE_METAL "Build ggml with Metal backend support" ${ENGINE_DEFAULT_ENABLE_METAL}) option(ENGINE_ENABLE_LLAMAFILE "Build ggml with llamafile SGEMM support" ${ENGINE_DEFAULT_ENABLE_LLAMAFILE}) @@ -84,7 +89,7 @@ option(ENGINE_BUILD_EXAMPLES "Build example binaries" OFF) option(ENGINE_BUILD_TESTS "Build framework unit tests" OFF) option(ENGINE_BUILD_WARMBENCH "Build warmbench helper binaries" OFF) -if (ENGINE_ENABLE_CUDA) +if (ENGINE_ENABLE_CUDA AND NOT ENGINE_ENABLE_HIP) enable_language(CUDA) find_package(CUDAToolkit REQUIRED) if (NOT MSVC) @@ -101,10 +106,14 @@ set(GGML_BUILD_TESTS OFF CACHE BOOL "Do not build upstream tests" FORCE) set(GGML_ALL_WARNINGS OFF CACHE BOOL "Keep upstream warnings quiet in this sample" FORCE) set(GGML_CCACHE OFF CACHE BOOL "Do not probe for ccache in this local setup" FORCE) set(GGML_CUDA ${ENGINE_ENABLE_CUDA} CACHE BOOL "Build ggml with CUDA backend support" FORCE) +set(GGML_HIP ${ENGINE_ENABLE_HIP} CACHE BOOL "Build ggml with HIP backend support" FORCE) set(GGML_VULKAN ${ENGINE_ENABLE_VULKAN} CACHE BOOL "Build ggml with Vulkan backend support" FORCE) set(GGML_METAL ${ENGINE_ENABLE_METAL} CACHE BOOL "Build ggml with Metal backend support" FORCE) set(GGML_LLAMAFILE ${ENGINE_ENABLE_LLAMAFILE} CACHE BOOL "Build ggml with llamafile SGEMM support" FORCE) set(GGML_CUDA_GRAPHS ${ENGINE_ENABLE_CUDA_GRAPHS} CACHE BOOL "Enable ggml CUDA graphs support" FORCE) +# HIP graphs are controlled by a separate upstream option that defaults to ON; +# keep it in sync so ENGINE_ENABLE_CUDA_GRAPHS=OFF also disables graphs on HIP builds +set(GGML_HIP_GRAPHS ${ENGINE_ENABLE_CUDA_GRAPHS} CACHE BOOL "Enable ggml HIP graphs support" FORCE) set(GGML_NATIVE ${ENGINE_ENABLE_NATIVE_CPU} CACHE BOOL "Build ggml CPU kernels with native host ISA flags" FORCE) set(SPM_ENABLE_SHARED OFF CACHE BOOL "Build sentencepiece shared libs" FORCE) @@ -529,7 +538,11 @@ if (ENGINE_ENABLE_OPENMP) target_link_libraries(engine_runtime PRIVATE OpenMP::OpenMP_CXX) endif() -if (ENGINE_ENABLE_CUDA) +if (MSVC) + target_compile_options(engine_runtime PRIVATE /utf-8) +endif() + +if (ENGINE_ENABLE_CUDA AND NOT ENGINE_ENABLE_HIP) target_sources(engine_runtime PRIVATE src/framework/audio/istft_cuda_runtime.cu src/framework/sampling/torch_random_cuda_runtime.cu @@ -849,6 +862,16 @@ if (ENGINE_BUILD_TESTS) COMMAND depthwise_conv1d_lowering_test ) + add_engine_unittest(conv_lowering_matrix_test tests/unittests/test_conv_lowering_matrix.cpp) + if (ENGINE_ENABLE_HIP) + target_compile_definitions(conv_lowering_matrix_test PRIVATE ENGINE_TEST_ENABLE_HIP=1) + endif() + + add_test( + NAME conv_lowering_matrix_test + COMMAND conv_lowering_matrix_test + ) + add_engine_unittest(gguf_tensor_source_test tests/unittests/test_gguf_tensor_source.cpp) target_include_directories(gguf_tensor_source_test PRIVATE ${CMAKE_CURRENT_SOURCE_DIR}/tests/unittests) diff --git a/app/cli/args.cpp b/app/cli/args.cpp index bdc32d81..dbd4c01d 100644 --- a/app/cli/args.cpp +++ b/app/cli/args.cpp @@ -96,6 +96,9 @@ engine::core::BackendType parse_backend(const std::string & value) { if (value == "cuda") { return engine::core::BackendType::Cuda; } + if (value == "hip" || value == "rocm") { + return engine::core::BackendType::Hip; + } if (value == "vulkan") { return engine::core::BackendType::Vulkan; } diff --git a/app/cli/main.cpp b/app/cli/main.cpp index d378b24d..c812032a 100644 --- a/app/cli/main.cpp +++ b/app/cli/main.cpp @@ -45,7 +45,7 @@ void print_task_list_help() { << " --task vad|asr|diar|sep|gen|tts|clon|vc|s2s|align|vdes|spk|svc\n" << " --family \n" << " --model \n" - << " --backend cpu|cuda|vulkan|metal|best\n" + << " --backend cpu|cuda|hip|rocm|vulkan|metal|best (rocm is an alias for hip)\n" << " --mode offline|streaming default offline\n" << " --device \n" << " --threads Backend and OpenMP worker threads, default 4\n" diff --git a/app/server/runtime.cpp b/app/server/runtime.cpp index 442cd975..d1a9c0a9 100644 --- a/app/server/runtime.cpp +++ b/app/server/runtime.cpp @@ -64,6 +64,8 @@ const char * backend_name(engine::core::BackendType type) { return "cpu"; case engine::core::BackendType::Cuda: return "cuda"; + case engine::core::BackendType::Hip: + return "hip"; case engine::core::BackendType::Vulkan: return "vulkan"; case engine::core::BackendType::Metal: diff --git a/docs/HIP.md b/docs/HIP.md new file mode 100644 index 00000000..0ad823e1 --- /dev/null +++ b/docs/HIP.md @@ -0,0 +1,373 @@ +# HIP/ROCm Backend Support + +## Overview + +This document describes the HIP/ROCm backend support added to audio.cpp for running models on AMD GPUs. + +The GGML library vendored at `external/ggml/` already contains a HIP backend that compiles CUDA kernels (`ggml-cuda/*.cu`) as HIP code via a vendor header mapping (`ggml-cuda/vendors/hip.h`). The changes in this PR expose that functionality through audio.cpp's build system, backend abstraction layer, and CLI. + +### How It Works + +``` +audio.cpp ENGINE_ENABLE_HIP=ON + -> CMakeLists.txt sets GGML_HIP=ON + -> external/ggml/src/ggml-hip/CMakeLists.txt + -> compiles ../ggml-cuda/*.cu with GGML_USE_HIP + -> vendor header maps cudaMalloc->hipMalloc, cublasCreate->hipblasCreate, etc. + -> sets GGML_USE_CUDA on ggml target (same register path as native CUDA) + -> audio.cpp backend.cpp sees #ifdef GGML_USE_CUDA == true + -> ggml_backend_cuda_init() / ggml_backend_cuda_reg() work for HIP + -> device name "ROCm0" distinguishes HIP from "CUDA0" +``` + +Key insight: HIP and CUDA share **the same GGML backend code**. The `ggml-hip/` directory contains only a `CMakeLists.txt` -- zero lines of unique kernel code. The runtime distinction is made via device name prefix (`"ROCm"` vs `"CUDA"`). + +--- + +## File-by-File Changes + +### 1. `CMakeLists.txt` (root) + +**L24** -- New option default: +```cmake +set(ENGINE_DEFAULT_ENABLE_HIP OFF) +``` + +**L36-L38** -- Forward pre-existing GGML_HIP variable (allows `-DGGML_HIP=ON` to auto-enable): +```cmake +if (DEFINED GGML_HIP) + set(ENGINE_DEFAULT_ENABLE_HIP ${GGML_HIP}) +endif() +``` + +**L53** -- New CMake option: +```cmake +option(ENGINE_ENABLE_HIP "Build ggml with HIP/ROCm backend support" ${ENGINE_DEFAULT_ENABLE_HIP}) +``` + +**L64** -- Guard CUDA language/Toolkit to CUDA-only builds: +```cmake +if (ENGINE_ENABLE_CUDA AND NOT ENGINE_ENABLE_HIP) + enable_language(CUDA) + find_package(CUDAToolkit REQUIRED) +``` + +**L81** -- Forward to vendored GGML: +```cmake +set(GGML_HIP ${ENGINE_ENABLE_HIP} CACHE BOOL "Build ggml with HIP backend support" FORCE) +``` + +**L497** -- Guard audio.cpp's own `.cu` files to CUDA-only builds: +```cmake +if (ENGINE_ENABLE_CUDA AND NOT ENGINE_ENABLE_HIP) + target_sources(engine_runtime PRIVATE + src/framework/audio/istft_cuda_runtime.cu + src/framework/sampling/torch_random_cuda_runtime.cu + ) +``` + +> **Rationale:** `istft_cuda_runtime.cu` and `torch_random_cuda_runtime.cu` call CUDA APIs directly (`cuda_runtime.h`, `cufft.h`). They are optional GPU-accelerated paths with CPU fallbacks (gated by `ENGINE_HAS_CUDA_ISTFT` and `ENGINE_HAS_CUDA_TORCH_RANDOM`). Skipping them on HIP builds is safe and avoids adding a HIP port of these audio.cpp-specific routines. + +--- + +### 2. `include/engine/framework/core/module.h` + +**L16** -- New enum value: +```cpp +enum class BackendType { + Cpu, + Cuda, + Hip, // <-- added + Vulkan, + Metal, + BestAvailable, +}; +``` + +--- + +### 3. `src/framework/core/backend.cpp` + +All HIP code paths reuse `#ifdef GGML_USE_CUDA` because `ggml-hip/CMakeLists.txt` defines `GGML_USE_CUDA` on the `ggml` target for static builds (line 92 of `external/ggml/src/ggml-hip/CMakeLists.txt`). + +**L107-L117** -- `init_backend()` Hip case: +```cpp +case BackendType::Hip: { +#ifdef GGML_USE_CUDA + ggml_backend_t backend = ggml_backend_cuda_init(config.device); + // ... +#endif +} +``` + +**L183-L190** -- `backend_type()` CUDA/HIP distinction: +```cpp +if (is_cuda_backend_handle(backend)) { +#ifdef GGML_USE_CUDA + if (backend_name_has_prefix(backend, "ROCm")) { + return BackendType::Hip; + } +#endif + return BackendType::Cuda; +} +``` + +`is_cuda_backend_handle()` compares device registration against `ggml_backend_cuda_reg()`, which returns the same registry for both CUDA and HIP. We disambiguate by checking the backend name prefix: HIP devices are named `"ROCm0"`, `"ROCm1"`, etc. + +**L217** -- `release_backend_graph_resources()` by backend handle: +```cpp +if (backend_name_has_prefix(backend, "CUDA") || backend_name_has_prefix(backend, "ROCm")) { +``` + +**L228** -- `release_backend_graph_resources()` by BackendType: +```cpp +if (backend_type == BackendType::Cuda || backend_type == BackendType::Hip) { +``` + +**L318** -- `query_backend_memory()` by BackendConfig: +```cpp +case BackendType::Cuda: +case BackendType::Hip: // <-- fall-through to CUDA path +``` + +--- + +### 4. `app/cli/args.cpp` + +CLI string to BackendType mapping (`rocm` accepted as an alias): +```cpp +if (value == "hip" || value == "rocm") { + return engine::core::BackendType::Hip; +} +``` + +### 5. `app/cli/main.cpp` + +Help text: +``` +--backend cpu|cuda|hip|rocm|vulkan|metal|best (rocm is an alias for hip) +``` + +--- + +### 6. `external/sentencepiece/src/CMakeLists.txt` (build compatibility fix) +**L266** -- Disable `-fPIC` on Windows (clang Windows target does not support it): +```cmake +# Original: +if (NOT MSVC) +# Fixed: +if (NOT MSVC AND NOT WIN32) +``` + +> This is not HIP-specific, but is required when compiling with clang (the HIP compiler) on Windows. + +--- + +### 7. `external/ggml` -- hipBLASLt GEMM path + +rocBLAS does not ship Tensile kernels for every AMD GPU arch (e.g. gfx1103 on Windows), while hipBLASLt covers more arches. HIP builds therefore route every cuBLAS-equivalent GEMM through hipBLASLt by default. + +**`external/ggml/src/ggml-hip/CMakeLists.txt`** -- New option `GGML_HIP_HIPBLASLT` (default `ON`). Runs `find_package(hipblaslt)`; when found, defines `GGML_HIP_USE_HIPBLASLT` and links `roc::hipblaslt`. When not found (e.g. older ROCm without the hipblaslt dev package), it warns and falls back to the original hipBLAS (rocBLAS) path, so Linux builds are unaffected. + +**`external/ggml/src/ggml-cuda/common.cuh`** -- Includes `` when `GGML_HIP_USE_HIPBLASLT` is defined; adds a `HIPBLASLT_CHECK` error macro; adds lazily-created per-device `hipblasLtHandle_t` handles and a 32 MiB per-device workspace to `ggml_backend_cuda_context` (freed in the context destructor in `ggml-cuda.cu`). + +**`external/ggml/src/ggml-cuda/ggml-cuda.cu`** -- New `ggml_hipblaslt_gemm()` helper mirroring the cublas call semantics (`OP_T`/`OP_N`, column-major, strided-batch support), plus a `hipblasDatatype_t` -> `hipDataType` conversion (the legacy hipBLAS enum values, 150+, differ from hipDataType, 0-based, on ROCm < 6.5). All GEMM call sites are switched to it when `GGML_HIP_USE_HIPBLASLT` is defined, with the original cublas code kept as `#else` fallback: + +- `ggml_cuda_op_mul_mat_cublas()`: BF16, FP16->FP32, FP16->FP16, and FP32 GEMM paths +- `ggml_cuda_mul_mat_batched_cublas_impl()`: strided-batched path (native hipBLASLt batched layouts) and pointer-array batched path (emulated with a per-batch-element GEMM loop; hipBLASLt has no pointer-array API) + +All Lt GEMMs use `HIPBLAS_COMPUTE_32F` with FP32 scale for accuracy. The algorithm heuristic is queried per call -- caching heuristics per shape is future work. + +### 8. `scripts/build_windows_hip.ps1` + +Dedicated Windows HIP build script. Auto-detects ROCm (`HIP_PATH` or `C:\Program Files\AMD\ROCm\*`), GPU targets (`amdgpu-arch`), cmake, and ninja (PATH or the VS-bundled copy), then configures and builds with `ENGINE_ENABLE_HIP=ON`. See the Windows build section below for options. + +--- + +## Build Instructions + +### Linux + +Prerequisites: +- ROCm 6.1+ installed (`/opt/rocm` or custom path) +- hipBLAS, rocBLAS, hipBLASLt packages (hipBLASLt is used for GEMM by default; without it the build falls back to hipBLAS/rocBLAS) +- Find your GPU target: `rocminfo | grep gfx | head -1 | awk '{print $2}'` + (e.g. `gfx1100` for RX 7900 XTX, `gfx1151` for Strix Halo iGPU) + +```bash +# Configure +cmake -S . -B build_hip \ + -DENGINE_ENABLE_HIP=ON \ + -DGPU_TARGETS=gfx1151 \ + -DCMAKE_C_COMPILER="$(hipconfig -l)/clang" \ + -DCMAKE_CXX_COMPILER="$(hipconfig -l)/clang++" \ + -DCMAKE_BUILD_TYPE=Release + +# Build +cmake --build build_hip -j$(nproc) + +# Run +./build_hip/bin/audiocpp_server --config server.json +./build_hip/bin/audiocpp_cli --task tts --family index_tts2 --model ./model --backend hip --device 0 +``` + +Multiple GPU targets (for distribution): +```bash +-DGPU_TARGETS="gfx1030;gfx1100;gfx1101;gfx1102;gfx1150;gfx1151" +``` + +**iGPU vs discrete GPU (Linux):** + +- **Discrete GPU (gfx1100/1101/1102, gfx1200/1201):** the defaults are fine. `ENGINE_ENABLE_CUDA_GRAPHS` is ON by default and dGPUs have the VRAM headroom for it. Optionally add `-DGGML_HIP_NO_VMM=OFF` — HIP VMM works on Linux dGPUs and improves ggml's memory-pool reuse under varying shapes. +- **iGPU (gfx1103 780M, gfx1150, gfx1151):** add `-DENGINE_ENABLE_CUDA_GRAPHS=OFF` if you see `out of memory` during graph warmup — each cached graph reserves its own buffers out of shared system memory (see Known Limitations). Keep `GGML_HIP_NO_VMM=ON` (the default). +- **gfx1103 note:** with the default hipBLASLt GEMM path, no `HSA_OVERRIDE_GFX_VERSION=11.0.0` is needed. That override is only required if you force the legacy rocBLAS path with `-DGGML_HIP_HIPBLASLT=OFF`. +- **rocWMMA fattn** (`GGML_HIP_ROCWMMA_FATTN`): keep OFF on both — the default `fattn-tile` kernels are faster on RDNA3/RDNA4. + +### Windows + +> **hipBLASLt GEMM (default):** HIP builds use hipBLASLt instead of hipBLAS (rocBLAS) for all cuBLAS-equivalent GEMM calls. hipBLASLt ships Tensile kernels for more GPU architectures than rocBLAS on Windows — notably **gfx1103 (Radeon 780M) works**, even though rocBLAS has no gfx1103 library. Disable with `-DGGML_HIP_HIPBLASLT=OFF` to fall back to hipBLAS. + +**Build script (recommended):** + +```powershell +# Auto-detects ROCm (HIP_PATH), GPU targets (amdgpu-arch), cmake, and ninja +powershell -ExecutionPolicy Bypass -File scripts\build_windows_hip.ps1 + +# Useful options: +# -GpuTargets gfx1103 override target arch +# -NoHipblasLt use hipBLAS (rocBLAS) instead of hipBLASLt +# -ForceMmq route quantized matmul through GGML MMQ kernels +# -WithVmm enable HIP virtual memory management +# -ConfigureOnly / -Clean / -Target audiocpp_cli / -Jobs 8 +``` + +**Manual build:** + +```cmd +# Set ROCm environment +set PATH=C:\Program Files\AMD\ROCm\6.4\bin;%PATH% +set HIP_PATH=C:\Program Files\AMD\ROCm\6.4 +set ROCM_PATH=C:\Program Files\AMD\ROCm\6.4 + +# Configure (ninja: bundled with Visual Studio CMake tools or standalone) +cmake -S . -B build/hip -G Ninja ^ + -DENGINE_ENABLE_HIP=ON ^ + -DENGINE_ENABLE_OPENMP=OFF ^ + -DGPU_TARGETS=gfx1103 ^ + -DCMAKE_C_COMPILER="%HIP_PATH%\bin\clang.exe" ^ + -DCMAKE_CXX_COMPILER="%HIP_PATH%\bin\clang++.exe" ^ + -DCMAKE_BUILD_TYPE=Release + +# Build +cmake --build build/hip +``` + +The resulting binaries need the ROCm `bin` directory on `PATH` at runtime (for `amdhip64_6.dll`, `hipblas.dll`, `hipblaslt.dll`, ...). + +**iGPU vs discrete GPU:** + +The script defaults are tuned for memory-constrained iGPUs (780M, Strix Point/Halo). On a discrete GPU (gfx1100/1101/1102, gfx1200/1201) adjust the following: + +```powershell +powershell -ExecutionPolicy Bypass -File scripts\build_windows_hip.ps1 ` + -GpuTargets gfx1100 ` # your dGPU arch; see the support matrix below + -Graphs ` # re-enable CUDA graphs: dGPUs have VRAM headroom, avoids per-request graph rebuild overhead + -WithVmm # optional: HIP VMM improves ggml memory-pool reuse; keep OFF on iGPUs +``` + +- **CUDA graphs** (`-Graphs`): each cached graph reserves its own VRAM buffers. Fine on a dGPU with 8+ GB; on UMA iGPUs it can exhaust shared memory during warmup (see Known Limitations). +- **VMM** (`-WithVmm`): with `GGML_HIP_NO_VMM=ON` (the default) ggml's memory pool reuses less, which costs performance under varying shapes. It is required on Windows iGPUs but generally works on dGPUs. +- **hipBLASLt**: keep the default ON. Unlike gfx1103, rocBLAS does ship Tensile kernels for the dGPU arches on Windows, so `-NoHipblasLt` is viable there — but there is no performance reason to prefer it. +- **rocWMMA fattn**: stays OFF on dGPUs too; the default `fattn-tile` kernels are the faster path on RDNA3/RDNA4. + +> **Legacy workaround (pre-hipBLASLt):** `-DGGML_CUDA_FORCE_MMQ=ON` bypasses BLAS for *quantized* matmul only; FP16/FP32 GEMM still required rocBLAS and failed on gfx1103. Renaming rocBLAS `gfx1100` Tensile files to `gfx1103` segfaults with ROCm 6.4 on Windows — do not use. hipBLASLt is the supported path. + +--- + +## Architecture Support Matrix + +| GPU Target | Type | rocBLAS Tensile (Linux) | rocBLAS Tensile (Windows) | hipBLASLt (Linux + Windows) | +|---|---|---|---|---| +| gfx1100 | RDNA3 discrete (7900 XTX/XT) | Yes | Yes | Yes | +| gfx1101 | RDNA3 discrete (7900 GRE) | Yes | Yes | Yes | +| gfx1102 | RDNA3 discrete (7600 XT) | Yes | Yes | Yes | +| **gfx1103** | **RDNA3 iGPU (780M)** | Yes 1 | **No** | **Yes** | +| gfx1150 | RDNA3.5 iGPU (Strix Point) | Yes | Yes | Yes | +| gfx1151 | RDNA3.5 iGPU (Strix Halo) | Yes | Yes | Yes | +| gfx1200 | RDNA4 discrete | Yes | Yes | Yes | +| gfx1201 | RDNA4 discrete | Yes | Yes | Yes | + +> 1 gfx1103 on Linux requires `HSA_OVERRIDE_GFX_VERSION=11.0.0` for rocBLAS. This environment variable is [not supported on Windows](https://github.com/ROCm/ROCm/issues/2654). With the hipBLASLt GEMM path (default), gfx1103 works on both Linux and Windows without any override. + +--- + +## Known Limitations & Future Work + +### Model-level GPU optimizations not yet enabled for HIP + +The following locations check `BackendType::Cuda` specifically and will not apply their CUDA optimizations for HIP backends. They fall back to generic CPU-equivalent paths, which are functionally correct but may have lower performance. + +| File | Line | Optimization | +|---|---|---| +| `src/framework/sampling/torch_random.cpp` | 311 | TorchCUDA random sampling policy probe | +| `src/framework/modules/conv_modules.cpp` | 254 | CUDA depthwise conv1d fast path | +| `src/models/irodori_tts/rf_dit.cpp` | 141, 535 | CUDA-specific attention path | +| `src/models/demucs/pipeline.cpp` | 474, 561 | CUDA-specific tensor pipeline | +| `src/models/demucs/session.cpp` | 154 | CUDA-specific tensor storage | +| `src/framework/modules/optimizations/fast_projection_modules.cpp` | 62 | CUDA projection acceleration | +| `src/models/miocodec/audio_pipeline.cpp` | 294 | CUDA audio pipeline | +| `src/models/index_tts2/gpt.cpp` | 83 | CUDA-required GPT module | +| `src/models/vibevoice/session.cpp` | 164 | CUDA-specific max seconds | +| `src/models/vibevoice_asr/session.cpp` | 543, 545 | CUDA-specific weight storage | +| `src/models/vibevoice_asr/speech_encoder.cpp` | 54 | CUDA-specific speech encoder | + +**To enable:** Change `== BackendType::Cuda` to `== BackendType::Cuda || == BackendType::Hip` on a per-file basis after verifying each optimization works correctly on AMD hardware. + +### Switch statements missing `Hip` case (compiler warnings) + +- `app/server/runtime.cpp` (`backend_name()`) and `src/models/ace_step/planner.cpp` (`planner_prefill_uses_host_backend()`) now handle `Hip` explicitly. +- `src/models/moss/moss_tts_local/session.cpp` (`resolve_auto_weight_type()`) intentionally keeps HIP on the `Native` dtype default; enabling the CUDA-style BF16 path for HIP needs validation with real models first. + +### hipBLASLt GEMM path notes + +- Heuristics are queried per GEMM call (shape-keyed algo caching is future work). +- The pointer-array batched GEMM (`cublasGemmBatchedEx` equivalent) is emulated with a per-batch-element loop. +- Verified on gfx1103 / ROCm 6.4 / Windows: F32 and F16 2D GEMM, strided-batched, and broadcast-batched all match CPU reference within 1e-4. + +### CUDA graphs on iGPUs / limited VRAM + +`ENGINE_ENABLE_CUDA_GRAPHS` defaults to ON (inherited from the CUDA build), but every cached graph reserves its own VRAM buffers. On memory-constrained GPUs (notably UMA iGPUs like the 780M, where `GGML_HIP_NO_VMM` also reduces ggml's memory-pool reuse), a burst of new shapes (e.g. the second inference request) can exhaust VRAM with `cudaMalloc failed: out of memory` during graph warmup. `scripts/build_windows_hip.ps1` therefore builds with `-DENGINE_ENABLE_CUDA_GRAPHS=OFF` by default; pass `-Graphs` to re-enable on discrete GPUs with headroom (recommended there — it removes per-request graph rebuild overhead). + +### audio.cpp CUDA-specific `.cu` files + +- `src/framework/audio/istft_cuda_runtime.cu` -- uses `cufft` +- `src/framework/sampling/torch_random_cuda_runtime.cu` -- uses CUDA Driver API + +These are skipped entirely on HIP builds. To enable GPU acceleration for ISTFT and TorchRandom on AMD GPUs, they would need to be ported (`cufft` -> `hipfft`, `cuInit`/`cuDeviceGet` -> `hipInit`/`hipDeviceGet`). + +--- + +## Verification Checklist + +- [ ] Linux: `cmake -DENGINE_ENABLE_HIP=ON ...` configures without errors +- [ ] Linux: `cmake --build build_hip` completes without errors +- [ ] Linux: `--backend hip --device 0` initializes and detects AMD GPU +- [x] Windows: HIP build configures and compiles (ROCm 6.4, gfx1103) +- [x] Windows: `--backend hip` appears in `--help` output +- [x] Windows: HIP backend initializes and detects the GPU (`ROCm0`, gfx1103) +- [x] Windows: F32/F16/BF16 GEMM via hipBLASLt matches CPU reference (2D, strided-batched, broadcast-batched) +- [ ] `backend_type()` returns `BackendType::Hip` for ROCm-initialized backends +- [ ] `query_backend_memory()` works for HIP +- [ ] `release_backend_graph_resources()` works for HIP +- [ ] Model inference runs on HIP without crashes +- [ ] Model inference produces numerically correct results +- [ ] Server JSON config accepts `"backend": "hip"` + +--- + +## References + +- GGML HIP vendor header: `external/ggml/src/ggml-cuda/vendors/hip.h` (~300 lines, ~150 CUDA-to-HIP mappings) +- GGML HIP backend CMake: `external/ggml/src/ggml-hip/CMakeLists.txt` +- GGML backend registry: `external/ggml/src/ggml-backend-reg.cpp` (L116: HIP registers via `GGML_USE_CUDA`) +- llama.cpp HIP build docs: `docs/build.md` (upstream GGML documentation) diff --git a/external/ggml/src/ggml-cuda/common.cuh b/external/ggml/src/ggml-cuda/common.cuh index 0e523239..c0964f2a 100644 --- a/external/ggml/src/ggml-cuda/common.cuh +++ b/external/ggml/src/ggml-cuda/common.cuh @@ -31,6 +31,9 @@ #if defined(GGML_USE_HIP) #include "vendors/hip.h" +#if defined(GGML_HIP_USE_HIPBLASLT) +#include +#endif // defined(GGML_HIP_USE_HIPBLASLT) #elif defined(GGML_USE_MUSA) #include "vendors/musa.h" #else @@ -188,6 +191,15 @@ void ggml_cuda_error(const char * stmt, const char * func, const char * file, in #define CUBLAS_CHECK(err) CUDA_CHECK_GEN(err, CUBLAS_STATUS_SUCCESS, cublas_get_error_str) +#if defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) + static const char * hipblaslt_get_error_str(const hipblasStatus_t err) { + return hipblasStatusToString(err); + } +#define HIPBLASLT_CHECK(err) CUDA_CHECK_GEN(err, HIPBLAS_STATUS_SUCCESS, hipblaslt_get_error_str) + +#define HIPBLASLT_WORKSPACE_SIZE ((size_t) 32 * 1024 * 1024) +#endif // defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) + #ifdef GGML_USE_NCCL #define NCCL_CHECK(err) CUDA_CHECK_GEN(err, ncclSuccess, ncclGetErrorString) #endif // GGML_USE_NCCL @@ -1458,6 +1470,32 @@ struct ggml_backend_cuda_context { return cublas_handle(device); } +#if defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) + hipblasLtHandle_t hipblaslt_handles[GGML_CUDA_MAX_DEVICES] = {nullptr}; + void * hipblaslt_workspaces[GGML_CUDA_MAX_DEVICES] = {nullptr}; + + hipblasLtHandle_t hipblaslt_handle(int device) { + if (hipblaslt_handles[device] == nullptr) { + ggml_cuda_set_device(device); + HIPBLASLT_CHECK(hipblasLtCreate(&hipblaslt_handles[device])); + } + return hipblaslt_handles[device]; + } + + hipblasLtHandle_t hipblaslt_handle() { + return hipblaslt_handle(device); + } + + // hipBLASLt requires a user-provided workspace; allocate it once per device + void * hipblaslt_workspace(int device) { + if (hipblaslt_workspaces[device] == nullptr) { + ggml_cuda_set_device(device); + CUDA_CHECK(cudaMalloc(&hipblaslt_workspaces[device], HIPBLASLT_WORKSPACE_SIZE)); + } + return hipblaslt_workspaces[device]; + } +#endif // defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) + // pool std::unique_ptr pools[GGML_CUDA_MAX_DEVICES][GGML_CUDA_MAX_STREAMS]; diff --git a/external/ggml/src/ggml-cuda/ggml-cuda.cu b/external/ggml/src/ggml-cuda/ggml-cuda.cu index c354877e..ecfdc939 100644 --- a/external/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/external/ggml/src/ggml-cuda/ggml-cuda.cu @@ -616,6 +616,14 @@ ggml_backend_cuda_context::~ggml_backend_cuda_context() { if (cublas_handles[i] != nullptr) { CUBLAS_CHECK(cublasDestroy(cublas_handles[i])); } +#if defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) + if (hipblaslt_handles[i] != nullptr) { + HIPBLASLT_CHECK(hipblasLtDestroy(hipblaslt_handles[i])); + } + if (hipblaslt_workspaces[i] != nullptr) { + CUDA_CHECK(cudaFree(hipblaslt_workspaces[i])); + } +#endif // defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) } } @@ -1622,6 +1630,88 @@ static const cublas_force_compute_type & ggml_cuda_cublas_get_force_compute_type return compute_type; } +#if defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) +// hipBLASLt equivalent of the cublasGemm* calls used below. +// rocBLAS does not ship Tensile kernels for every AMD GPU arch (e.g. gfx1103 on Windows), +// while hipBLASLt covers them, so HIP builds route GEMM through hipBLASLt when available. +// Computes C = op(A) * op(B) with op(A) = A^T, op(B) = B (column-major, same as the cublas calls). +// hipBLASLt only accepts hipDataType. ROCm < 6.5 routes cudaDataType_t to the legacy +// hipblasDatatype_t enum (150/151/168), while ROCm >= 6.5 uses hipDataType (0/2/14) directly. +// Accept the raw integer value and map both numbering schemes, so this compiles on all ROCm versions. +static hipDataType ggml_hipblaslt_convert_type(int type) { + switch (type) { + case 150: return HIP_R_16F; // legacy HIPBLAS_R_16F + case 151: return HIP_R_32F; // legacy HIPBLAS_R_32F + case 168: return HIP_R_16BF; // legacy HIPBLAS_R_16B + default: + GGML_ASSERT(type == HIP_R_16F || type == HIP_R_32F || type == HIP_R_16BF); + return (hipDataType) type; + } +} + +static void ggml_hipblaslt_gemm( + ggml_backend_cuda_context & ctx, cudaStream_t stream, + int64_t m, int64_t n, int64_t k, + const void * A, int type_a, int64_t lda, int64_t stride_a, + const void * B, int type_b, int64_t ldb, int64_t stride_b, + void * C, int type_c, int64_t ldc, int64_t stride_c, + int64_t batch_count) { + + const hipblasOperation_t trans_a = HIPBLAS_OP_T; + const hipblasOperation_t trans_b = HIPBLAS_OP_N; + + const float alpha = 1.0f; + const float beta = 0.0f; + + hipblasLtHandle_t lt = ctx.hipblaslt_handle(); + void * workspace = ctx.hipblaslt_workspace(ctx.device); + + hipblasLtMatmulDesc_t matmul_desc; + hipblasLtMatrixLayout_t layout_a, layout_b, layout_c; + hipblasLtMatmulPreference_t pref; + + HIPBLASLT_CHECK(hipblasLtMatmulDescCreate(&matmul_desc, HIPBLAS_COMPUTE_32F, HIP_R_32F)); + HIPBLASLT_CHECK(hipblasLtMatmulDescSetAttribute(matmul_desc, HIPBLASLT_MATMUL_DESC_TRANSA, &trans_a, sizeof(trans_a))); + HIPBLASLT_CHECK(hipblasLtMatmulDescSetAttribute(matmul_desc, HIPBLASLT_MATMUL_DESC_TRANSB, &trans_b, sizeof(trans_b))); + + // layout dims describe the stored (pre-op) matrix: A is stored [k, m], B is stored [k, n], C is [m, n] + HIPBLASLT_CHECK(hipblasLtMatrixLayoutCreate(&layout_a, ggml_hipblaslt_convert_type(type_a), k, m, lda)); + HIPBLASLT_CHECK(hipblasLtMatrixLayoutCreate(&layout_b, ggml_hipblaslt_convert_type(type_b), k, n, ldb)); + HIPBLASLT_CHECK(hipblasLtMatrixLayoutCreate(&layout_c, ggml_hipblaslt_convert_type(type_c), m, n, ldc)); + + if (batch_count > 1) { + int batch_count_i32 = (int) batch_count; + HIPBLASLT_CHECK(hipblasLtMatrixLayoutSetAttribute(layout_a, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &batch_count_i32, sizeof(batch_count_i32))); + HIPBLASLT_CHECK(hipblasLtMatrixLayoutSetAttribute(layout_a, HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET, &stride_a, sizeof(stride_a))); + HIPBLASLT_CHECK(hipblasLtMatrixLayoutSetAttribute(layout_b, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &batch_count_i32, sizeof(batch_count_i32))); + HIPBLASLT_CHECK(hipblasLtMatrixLayoutSetAttribute(layout_b, HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET, &stride_b, sizeof(stride_b))); + HIPBLASLT_CHECK(hipblasLtMatrixLayoutSetAttribute(layout_c, HIPBLASLT_MATRIX_LAYOUT_BATCH_COUNT, &batch_count_i32, sizeof(batch_count_i32))); + HIPBLASLT_CHECK(hipblasLtMatrixLayoutSetAttribute(layout_c, HIPBLASLT_MATRIX_LAYOUT_STRIDED_BATCH_OFFSET, &stride_c, sizeof(stride_c))); + } + + HIPBLASLT_CHECK(hipblasLtMatmulPreferenceCreate(&pref)); + size_t max_workspace = HIPBLASLT_WORKSPACE_SIZE; + HIPBLASLT_CHECK(hipblasLtMatmulPreferenceSetAttribute(pref, HIPBLASLT_MATMUL_PREF_MAX_WORKSPACE_BYTES, &max_workspace, sizeof(max_workspace))); + + hipblasLtMatmulHeuristicResult_t heuristic; + int algo_count = 0; + HIPBLASLT_CHECK(hipblasLtMatmulAlgoGetHeuristic(lt, matmul_desc, layout_a, layout_b, layout_c, layout_c, + pref, 1, &heuristic, &algo_count)); + GGML_ASSERT(algo_count > 0); + + HIPBLASLT_CHECK(hipblasLtMatmul(lt, matmul_desc, + &alpha, A, layout_a, B, layout_b, + &beta, C, layout_c, C, layout_c, + &heuristic.algo, workspace, max_workspace, stream)); + + HIPBLASLT_CHECK(hipblasLtMatmulPreferenceDestroy(pref)); + HIPBLASLT_CHECK(hipblasLtMatrixLayoutDestroy(layout_a)); + HIPBLASLT_CHECK(hipblasLtMatrixLayoutDestroy(layout_b)); + HIPBLASLT_CHECK(hipblasLtMatrixLayoutDestroy(layout_c)); + HIPBLASLT_CHECK(hipblasLtMatmulDescDestroy(matmul_desc)); +} +#endif // defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) + static void ggml_cuda_op_mul_mat_cublas( ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, const char * src0_dd_i, const float * src1_ddf_i, @@ -1673,6 +1763,15 @@ static void ggml_cuda_op_mul_mat_cublas( const float alpha_f32 = 1.0f; const float beta_f32 = 0.0f; +#if defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) + ggml_hipblaslt_gemm(ctx, stream, + row_diff, src1_ncols, ne10, + src0_ptr, CUDA_R_16BF, ne00, 0, + src1_ptr, CUDA_R_16BF, ne10, 0, + dst_bf16.get(), CUDA_R_16BF, ldc, 0, + 1); + GGML_UNUSED_VARS(alpha_f32, beta_f32); +#else CUBLAS_CHECK(cublasSetStream(ctx.cublas_handle(id), stream)); CUBLAS_CHECK( cublasGemmEx(ctx.cublas_handle(id), CUBLAS_OP_T, CUBLAS_OP_N, @@ -1682,6 +1781,7 @@ static void ggml_cuda_op_mul_mat_cublas( &beta_f32, dst_bf16.get(), CUDA_R_16BF, ldc, CUBLAS_COMPUTE_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP)); +#endif // defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) const to_fp32_cuda_t to_fp32_cuda = ggml_get_to_fp32_cuda(GGML_TYPE_BF16); to_fp32_cuda(dst_bf16.get(), dst_dd_i, row_diff*src1_ncols, stream); @@ -1718,6 +1818,15 @@ static void ggml_cuda_op_mul_mat_cublas( { const float alpha = 1.0f; const float beta = 0.0f; +#if defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) + GGML_UNUSED_VARS(alpha, beta); + ggml_hipblaslt_gemm(ctx, stream, + row_diff, src1_ncols, ne10, + src0_ptr, CUDA_R_16F, ne00, 0, + src1_ptr, CUDA_R_16F, ne10, 0, + dst_dd_i, CUDA_R_32F, ldc, 0, + 1); +#else CUBLAS_CHECK( cublasGemmEx(ctx.cublas_handle(id), CUBLAS_OP_T, CUBLAS_OP_N, row_diff, src1_ncols, ne10, @@ -1726,12 +1835,22 @@ static void ggml_cuda_op_mul_mat_cublas( &beta, dst_dd_i, CUDA_R_32F, ldc, CUBLAS_COMPUTE_32F, CUBLAS_GEMM_DEFAULT_TENSOR_OP)); +#endif // defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) } else { ggml_cuda_pool_alloc dst_f16(ctx.pool(id), row_diff*src1_ncols); const half alpha_f16 = 1.0f; const half beta_f16 = 0.0f; +#if defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) + GGML_UNUSED_VARS(alpha_f16, beta_f16); + ggml_hipblaslt_gemm(ctx, stream, + row_diff, src1_ncols, ne10, + src0_ptr, CUDA_R_16F, ne00, 0, + src1_ptr, CUDA_R_16F, ne10, 0, + dst_f16.get(), CUDA_R_16F, ldc, 0, + 1); +#else CUBLAS_CHECK( cublasGemmEx(ctx.cublas_handle(id), CUBLAS_OP_T, CUBLAS_OP_N, row_diff, src1_ncols, ne10, @@ -1740,6 +1859,7 @@ static void ggml_cuda_op_mul_mat_cublas( &beta_f16, dst_f16.get(), CUDA_R_16F, ldc, CUBLAS_COMPUTE_16F, CUBLAS_GEMM_DEFAULT_TENSOR_OP)); +#endif // defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) const to_fp32_cuda_t to_fp32_cuda = ggml_get_to_fp32_cuda(GGML_TYPE_F16); to_fp32_cuda(dst_f16.get(), dst_dd_i, row_diff*src1_ncols, stream); @@ -1767,6 +1887,15 @@ static void ggml_cuda_op_mul_mat_cublas( const float alpha = 1.0f; const float beta = 0.0f; +#if defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) + GGML_UNUSED_VARS(alpha, beta); + ggml_hipblaslt_gemm(ctx, stream, + row_diff, src1_ncols, ne10, + src0_ddf_i, CUDA_R_32F, ne00, 0, + src1_ddf1_i, CUDA_R_32F, ne10, 0, + dst_dd_i, CUDA_R_32F, ldc, 0, + 1); +#else CUBLAS_CHECK(cublasSetStream(ctx.cublas_handle(id), stream)); CUBLAS_CHECK( cublasSgemm(ctx.cublas_handle(id), CUBLAS_OP_T, CUBLAS_OP_N, @@ -1774,6 +1903,7 @@ static void ggml_cuda_op_mul_mat_cublas( &alpha, src0_ddf_i, ne00, src1_ddf1_i, ne10, &beta, dst_dd_i, ldc)); +#endif // defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) } GGML_UNUSED_VARS(dst, src1_ddq_i, src1_padded_row_size); @@ -2251,6 +2381,9 @@ static void ggml_cuda_mul_mat_batched_cublas_impl(ggml_backend_cuda_context & ct size_t nbd3 = dst->nb[3]; cublasComputeType_t cu_compute_type = traits::compute_type; +#if defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) + GGML_UNUSED(cu_compute_type); // only referenced by the cublas fallback paths +#endif // defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) cudaDataType_t cu_data_type = traits::data_type; cudaDataType_t cu_data_type_a = traits::data_type; cudaDataType_t cu_data_type_b = traits::data_type; @@ -2302,6 +2435,15 @@ static void ggml_cuda_mul_mat_batched_cublas_impl(ggml_backend_cuda_context & ct // there is no broadcast and src0, src1 are contiguous across dims 2, 3 // use cublasGemmStridedBatchedEx +#if defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) + GGML_UNUSED_VARS(alpha, beta); + ggml_hipblaslt_gemm(ctx, main_stream, + ne01, ne11, ne10, + src0_ptr, cu_data_type_a, nb01/nb00, sma, + src1_ptr, cu_data_type_b, s11, smb, + dst_t, cu_data_type, ne0, ne1*ne0, + ne12*ne13); +#else CUBLAS_CHECK( cublasGemmStridedBatchedEx(ctx.cublas_handle(), CUBLAS_OP_T, CUBLAS_OP_N, ne01, ne11, ne10, @@ -2311,7 +2453,27 @@ static void ggml_cuda_mul_mat_batched_cublas_impl(ggml_backend_cuda_context & ct ne12*ne13, cu_compute_type, CUBLAS_GEMM_DEFAULT_TENSOR_OP)); +#endif // defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) } else { +#if defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) + // hipBLASLt has no pointer-array batched GEMM; issue one GEMM per batch element instead. + GGML_UNUSED_VARS(alpha, beta); + const size_t src1_nb2 = (src1->type == src0_type) ? nb12 : s12*sizeof(cuda_t); + const size_t src1_nb3 = (src1->type == src0_type) ? nb13 : s13*sizeof(cuda_t); + for (int64_t i13 = 0; i13 < ne13; i13++) { + for (int64_t i12 = 0; i12 < ne12; i12++) { + const char * ptr_a = (const char *) src0_ptr + (i12/r2)*nb02 + (i13/r3)*nb03; + const char * ptr_b = (const char *) src1_ptr + i12*src1_nb2 + i13*src1_nb3; + char * ptr_c = ( char *) dst_t + i12*nbd2 + i13*nbd3; + ggml_hipblaslt_gemm(ctx, main_stream, + ne01, ne11, ne10, + ptr_a, cu_data_type_a, nb01/nb00, 0, + ptr_b, cu_data_type_b, s11, 0, + ptr_c, cu_data_type, ne0, 0, + 1); + } + } +#else // use cublasGemmBatchedEx const int64_t ne23 = ne12*ne13; @@ -2350,6 +2512,7 @@ static void ggml_cuda_mul_mat_batched_cublas_impl(ggml_backend_cuda_context & ct ne23, cu_compute_type, CUBLAS_GEMM_DEFAULT_TENSOR_OP)); +#endif // defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) } // Convert output back to F32 if needed diff --git a/external/ggml/src/ggml-hip/CMakeLists.txt b/external/ggml/src/ggml-hip/CMakeLists.txt index e0b04bfa..19e77896 100644 --- a/external/ggml/src/ggml-hip/CMakeLists.txt +++ b/external/ggml/src/ggml-hip/CMakeLists.txt @@ -47,6 +47,18 @@ find_package(hip REQUIRED) find_package(hipblas REQUIRED) find_package(rocblas REQUIRED) +option(GGML_HIP_HIPBLASLT "ggml: use hipBLASLt instead of hipBLAS (rocBLAS) for GEMM" ON) + +if (GGML_HIP_HIPBLASLT) + find_package(hipblaslt) + if (hipblaslt_FOUND) + set(GGML_HIP_USE_HIPBLASLT ON) + message(STATUS "HIP: using hipBLASLt for GEMM") + else() + message(WARNING "GGML_HIP_HIPBLASLT is ON but hipblaslt was not found; falling back to hipBLAS (rocBLAS)") + endif() +endif() + if (GGML_HIP_RCCL) find_package(rccl REQUIRED) endif() @@ -114,6 +126,10 @@ if (GGML_HIP_NO_VMM) add_compile_definitions(GGML_HIP_NO_VMM) endif() +if (GGML_HIP_USE_HIPBLASLT) + add_compile_definitions(GGML_HIP_USE_HIPBLASLT) +endif() + if (GGML_HIP_ROCWMMA_FATTN) add_compile_definitions(GGML_HIP_ROCWMMA_FATTN) endif() @@ -155,3 +171,7 @@ if (GGML_HIP_RCCL) endif() target_link_libraries(ggml-hip PRIVATE ggml-base hip::host roc::rocblas roc::hipblas) + +if (GGML_HIP_USE_HIPBLASLT) + target_link_libraries(ggml-hip PRIVATE roc::hipblaslt) +endif() diff --git a/external/sentencepiece/src/CMakeLists.txt b/external/sentencepiece/src/CMakeLists.txt index b4637a33..45074155 100644 --- a/external/sentencepiece/src/CMakeLists.txt +++ b/external/sentencepiece/src/CMakeLists.txt @@ -263,7 +263,7 @@ endif() set_target_properties(sentencepiece-static PROPERTIES OUTPUT_NAME "sentencepiece") set_target_properties(sentencepiece_train-static PROPERTIES OUTPUT_NAME "sentencepiece_train") -if (NOT MSVC) +if (NOT MSVC AND NOT WIN32) set(CMAKE_CXX_FLAGS "-O3 -Wall -fPIC ${CMAKE_CXX_FLAGS}") if (SPM_NO_THREADLOCAL) add_definitions(-DSPM_NO_THREADLOCAL=1) diff --git a/include/engine/framework/core/module.h b/include/engine/framework/core/module.h index 4182f4c0..19949b84 100644 --- a/include/engine/framework/core/module.h +++ b/include/engine/framework/core/module.h @@ -13,6 +13,7 @@ namespace engine::core { enum class BackendType { Cpu, Cuda, + Hip, Vulkan, Metal, BestAvailable, diff --git a/scripts/build_windows_hip.ps1 b/scripts/build_windows_hip.ps1 new file mode 100644 index 00000000..dd1902b0 --- /dev/null +++ b/scripts/build_windows_hip.ps1 @@ -0,0 +1,155 @@ +[CmdletBinding()] +param( + [string]$RocmPath = "", + [string]$GpuTargets = "", + [string]$BuildType = "Release", + [string]$Target = "", + [int]$Jobs = 0, + [switch]$ConfigureOnly, + [switch]$Clean, + [switch]$NoHipblasLt, # fall back to hipBLAS (rocBLAS) GEMM + [switch]$ForceMmq, # route quantized matmul through GGML MMQ kernels instead of BLAS + [switch]$NoVmm = $true, # disable HIP virtual memory management (required on Windows iGPUs) + [switch]$WithVmm, # explicitly re-enable VMM + [switch]$Graphs # enable CUDA graphs (memory-hungry on iGPUs: each cached graph reserves its own VRAM) +) + +Set-StrictMode -Version Latest +$ErrorActionPreference = "Stop" + +function Invoke-Checked { + param( + [Parameter(Mandatory = $true)][string]$FilePath, + [Parameter()][string[]]$Arguments = @() + ) + Write-Host "> $FilePath $($Arguments -join ' ')" + & $FilePath @Arguments + if ($LASTEXITCODE -ne 0) { + throw "Command failed with exit code $LASTEXITCODE`: $FilePath $($Arguments -join ' ')" + } +} + +function Convert-ToCMakePath { + param([Parameter(Mandatory = $true)][string]$Path) + return ($Path -replace "\\", "/") +} + +function Find-FirstFile { + param([Parameter(Mandatory = $true)][string[]]$Patterns) + foreach ($pattern in $Patterns) { + $found = Get-ChildItem -Path $pattern -ErrorAction SilentlyContinue | Sort-Object FullName -Descending | Select-Object -First 1 + if ($null -ne $found) { + return $found.FullName + } + } + return "" +} + +# --- ROCm --- +if ($RocmPath -eq "") { + if ($env:HIP_PATH -and (Test-Path (Join-Path $env:HIP_PATH "bin\clang++.exe"))) { + $RocmPath = (Resolve-Path $env:HIP_PATH).Path + } else { + $clang = Find-FirstFile @("C:\Program Files\AMD\ROCm\*\bin\clang++.exe") + if ($clang -eq "") { + throw "ROCm was not found. Install the AMD HIP SDK or pass -RocmPath." + } + $RocmPath = (Resolve-Path (Join-Path (Split-Path $clang -Parent) "..")).Path + } +} +$clangxx = Join-Path $RocmPath "bin\clang++.exe" +$clangc = Join-Path $RocmPath "bin\clang.exe" +if (-not (Test-Path $clangxx) -or -not (Test-Path $clangc)) { + throw "ROCm clang not found under $RocmPath\bin" +} + +# --- GPU targets --- +if ($GpuTargets -eq "") { + $amdgpuArch = Join-Path $RocmPath "bin\amdgpu-arch.exe" + if (Test-Path $amdgpuArch) { + $detected = & $amdgpuArch 2>$null | Where-Object { $_ -match "^gfx" } | Select-Object -Unique + if ($detected) { + $GpuTargets = ($detected | Sort-Object -Unique) -join ";" + } + } + if ($GpuTargets -eq "") { + throw "Could not detect GPU targets with amdgpu-arch. Pass -GpuTargets gfxXXXX." + } +} + +# --- tools --- +$cmake = Get-Command "cmake.exe" -ErrorAction SilentlyContinue +if ($null -eq $cmake) { + throw "cmake.exe was not found on PATH" +} +$cmake = $cmake.Source + +$ninjaCmd = Get-Command "ninja.exe" -ErrorAction SilentlyContinue +if ($null -ne $ninjaCmd) { + $ninja = $ninjaCmd.Source +} else { + $ninja = Find-FirstFile @("C:\Program Files\Microsoft Visual Studio\*\*\Common7\IDE\CommonExtensions\Microsoft\CMake\Ninja\ninja.exe") + if ($ninja -eq "") { + throw "ninja.exe was not found. Install ninja or Visual Studio with the CMake tools component." + } +} + +$sourceDir = Split-Path $PSScriptRoot -Parent +$buildDir = Join-Path (Join-Path $sourceDir "build") "hip" + +$hipblasLtValue = if ($NoHipblasLt) { "OFF" } else { "ON" } +$forceMmqValue = if ($ForceMmq) { "ON" } else { "OFF" } +$noVmmValue = if ($WithVmm) { "OFF" } elseif ($NoVmm) { "ON" } else { "OFF" } +$graphsValue = if ($Graphs) { "ON" } else { "OFF" } + +Write-Host "ROCm: $RocmPath" +Write-Host "GPU targets: $GpuTargets" +Write-Host "CMake: $cmake" +Write-Host "Ninja: $ninja" +Write-Host "Build dir: $buildDir" +Write-Host "hipBLASLt GEMM: $hipblasLtValue, FORCE_MMQ: $forceMmqValue, NO_VMM: $noVmmValue, CUDA graphs: $graphsValue" + +if ($Clean) { + Invoke-Checked $cmake @("--build", $buildDir, "--target", "clean") + exit 0 +} + +$env:HIP_PATH = $RocmPath +$env:ROCM_PATH = $RocmPath +$env:PATH = @((Join-Path $RocmPath "bin"), $env:PATH) -join [IO.Path]::PathSeparator + +$configureArgs = @( + "-S", $sourceDir, + "-B", $buildDir, + "-G", "Ninja", + "-DCMAKE_MAKE_PROGRAM=$(Convert-ToCMakePath $ninja)", + "-DCMAKE_BUILD_TYPE=$BuildType", + "-DCMAKE_C_COMPILER=$(Convert-ToCMakePath $clangc)", + "-DCMAKE_CXX_COMPILER=$(Convert-ToCMakePath $clangxx)", + "-DENGINE_ENABLE_HIP=ON", + "-DENGINE_ENABLE_OPENMP=OFF", + "-DGPU_TARGETS=$GpuTargets", + "-DGGML_HIP_HIPBLASLT=$hipblasLtValue", + "-DGGML_CUDA_FORCE_MMQ=$forceMmqValue", + "-DGGML_HIP_NO_VMM=$noVmmValue", + "-DENGINE_ENABLE_CUDA_GRAPHS=$graphsValue" +) + +Invoke-Checked $cmake $configureArgs + +if ($ConfigureOnly) { + exit 0 +} + +$effectiveJobs = if ($Jobs -gt 0) { $Jobs } else { [Math]::Max(2, [Environment]::ProcessorCount) } +$buildArgs = @("--build", $buildDir, "-j", $effectiveJobs.ToString()) +if ($Target -ne "") { + $buildArgs += @("--target", $Target) +} + +Write-Host "Build jobs: $effectiveJobs" +Invoke-Checked $cmake $buildArgs + +Write-Host "" +Write-Host "Binaries: $buildDir\bin" +Write-Host "Note: run with '$RocmPath\bin' on PATH so the ROCm DLLs (amdhip64, hipblas, hipblaslt) are found." diff --git a/src/framework/core/backend.cpp b/src/framework/core/backend.cpp index b34cdea3..3845862d 100644 --- a/src/framework/core/backend.cpp +++ b/src/framework/core/backend.cpp @@ -93,17 +93,28 @@ ggml_backend_t init_backend(const BackendConfig & config) { } return backend; } - case BackendType::Cuda: { -#ifdef GGML_USE_CUDA - ggml_backend_t backend = ggml_backend_cuda_init(config.device); - if (backend == nullptr) { - throw std::runtime_error("Failed to initialize CUDA backend"); + case BackendType::Cuda: { +#ifdef GGML_USE_CUDA + ggml_backend_t backend = ggml_backend_cuda_init(config.device); + if (backend == nullptr) { + throw std::runtime_error("Failed to initialize CUDA backend"); } return backend; #else - throw std::runtime_error("CUDA backend requested but this build does not include GGML_USE_CUDA"); -#endif - } + throw std::runtime_error("CUDA backend requested but this build does not include GGML_USE_CUDA"); +#endif + } + case BackendType::Hip: { +#ifdef GGML_USE_CUDA + ggml_backend_t backend = ggml_backend_cuda_init(config.device); + if (backend == nullptr) { + throw std::runtime_error("Failed to initialize HIP backend"); + } + return backend; +#else + throw std::runtime_error("HIP backend requested but this build does not include GGML_USE_CUDA"); +#endif + } case BackendType::Vulkan: { #ifdef GGML_USE_VULKAN if (config.device < 0) { @@ -169,8 +180,13 @@ BackendType backend_type(ggml_backend_t backend) { if (is_host_backend(backend)) { return BackendType::Cpu; } - if (is_cuda_backend_handle(backend)) { - return BackendType::Cuda; + if (is_cuda_backend_handle(backend)) { +#ifdef GGML_USE_CUDA + if (backend_name_has_prefix(backend, "ROCm")) { + return BackendType::Hip; + } +#endif + return BackendType::Cuda; } if (is_vulkan_backend_handle(backend)) { return BackendType::Vulkan; @@ -197,19 +213,19 @@ void release_backend_graph_resources(ggml_backend_t backend, ggml_cgraph * graph if (backend == nullptr || graph == nullptr) { return; } -#ifdef GGML_USE_CUDA - if (backend_name_has_prefix(backend, "CUDA")) { - ggml_backend_cuda_clear_graph(backend, graph); - } -#endif -} - -void release_backend_graph_resources(BackendType backend_type, ggml_backend_t backend, ggml_cgraph * graph) { - if (backend == nullptr || graph == nullptr) { - return; - } -#ifdef GGML_USE_CUDA - if (backend_type == BackendType::Cuda) { +#ifdef GGML_USE_CUDA + if (backend_name_has_prefix(backend, "CUDA") || backend_name_has_prefix(backend, "ROCm")) { + ggml_backend_cuda_clear_graph(backend, graph); + } +#endif +} + +void release_backend_graph_resources(BackendType backend_type, ggml_backend_t backend, ggml_cgraph * graph) { + if (backend == nullptr || graph == nullptr) { + return; + } +#ifdef GGML_USE_CUDA + if (backend_type == BackendType::Cuda || backend_type == BackendType::Hip) { ggml_backend_cuda_clear_graph(backend, graph); } #else @@ -298,7 +314,8 @@ BackendMemorySnapshot query_backend_memory(ggml_backend_t backend, int device_hi BackendMemorySnapshot query_backend_memory(const BackendConfig & config) { BackendMemorySnapshot snapshot; switch (config.type) { - case BackendType::Cuda: + case BackendType::Cuda: + case BackendType::Hip: #ifdef GGML_USE_CUDA { size_t free_bytes = 0; diff --git a/src/models/ace_step/planner.cpp b/src/models/ace_step/planner.cpp index cf7395c5..62a58e59 100644 --- a/src/models/ace_step/planner.cpp +++ b/src/models/ace_step/planner.cpp @@ -158,6 +158,7 @@ bool planner_prefill_uses_host_backend(core::BackendType backend_type) { case core::BackendType::Metal: case core::BackendType::Cpu: case core::BackendType::Cuda: + case core::BackendType::Hip: case core::BackendType::BestAvailable: return false; } diff --git a/tests/moss_tts_local/codec_decode_parity.cpp b/tests/moss_tts_local/codec_decode_parity.cpp index 93eef290..24c5a3cf 100644 --- a/tests/moss_tts_local/codec_decode_parity.cpp +++ b/tests/moss_tts_local/codec_decode_parity.cpp @@ -5,6 +5,7 @@ #include "engine/framework/core/backend.h" #include "engine/framework/core/execution_context.h" +#include "engine/framework/assets/tensor_source.h" #include "engine/models/moss/shared/audio_tokenizer_decoder.h" #include @@ -101,8 +102,9 @@ int main(int argc, char ** argv) { std::cout << "codec=" << codec_dir.string() << "\n"; std::cout << "loading decoder weights...\n" << std::flush; + const auto source = engine::assets::open_tensor_source(codec_dir); engine::models::moss::MossAudioTokenizerDecoder decoder( - codec_dir, execution_context, num_quantizers, kWeightContextBytes, kGraphArenaBytes); + *source, execution_context, num_quantizers, kWeightContextBytes, kGraphArenaBytes); std::cout << "decoding " << frames << " frames...\n" << std::flush; const auto stereo = decoder.decode(codes); diff --git a/tests/moss_tts_local/codec_dequant_parity.cpp b/tests/moss_tts_local/codec_dequant_parity.cpp index 04925bd6..0591788b 100644 --- a/tests/moss_tts_local/codec_dequant_parity.cpp +++ b/tests/moss_tts_local/codec_dequant_parity.cpp @@ -4,6 +4,8 @@ #include "engine/models/moss/shared/audio_tokenizer_quantizer.h" +#include "engine/framework/assets/tensor_source.h" + #include #include #include @@ -33,7 +35,8 @@ int main(int argc, char ** argv) { } try { - engine::models::moss::MossAudioTokenizerQuantizer dequantizer(codec_dir, kNumQuantizers); + const auto source = engine::assets::open_tensor_source(codec_dir); + engine::models::moss::MossAudioTokenizerQuantizer dequantizer(*source, kNumQuantizers); const std::vector latent = dequantizer.decode(codes); // [code_dim, steps] const int64_t code_dim = dequantizer.code_dim(); diff --git a/tests/moss_tts_local/codec_encode_parity.cpp b/tests/moss_tts_local/codec_encode_parity.cpp index d021bdbf..9af3c259 100644 --- a/tests/moss_tts_local/codec_encode_parity.cpp +++ b/tests/moss_tts_local/codec_encode_parity.cpp @@ -8,6 +8,7 @@ #include "engine/framework/core/backend.h" #include "engine/framework/core/execution_context.h" +#include "engine/framework/assets/tensor_source.h" #include "engine/models/moss/shared/audio_tokenizer_encoder.h" #include @@ -115,8 +116,9 @@ int main(int argc, char ** argv) { engine::core::ExecutionContext execution_context(backend_config); std::cout << "loading codec encoder weights...\n" << std::flush; + const auto source = engine::assets::open_tensor_source(codec_dir); engine::models::moss::MossAudioTokenizerEncoder encoder( - codec_dir, execution_context, num_quantizers, kWeightContextBytes, kGraphArenaBytes); + *source, execution_context, num_quantizers, kWeightContextBytes, kGraphArenaBytes); std::cout << "encoding...\n" << std::flush; const auto codes = encoder.encode(stereo); diff --git a/tests/perf/model_perf.cpp b/tests/perf/model_perf.cpp index 70e4b6bf..820036b0 100644 --- a/tests/perf/model_perf.cpp +++ b/tests/perf/model_perf.cpp @@ -121,6 +121,9 @@ engine::core::BackendType parse_backend(const std::string & value) { if (lowered == "cuda") { return engine::core::BackendType::Cuda; } + if (lowered == "hip" || lowered == "rocm") { + return engine::core::BackendType::Hip; + } if (lowered == "vulkan") { return engine::core::BackendType::Vulkan; } diff --git a/tests/perf/model_perf_cases.json b/tests/perf/model_perf_cases.json index 130711dc..26d22342 100644 --- a/tests/perf/model_perf_cases.json +++ b/tests/perf/model_perf_cases.json @@ -25,6 +25,12 @@ "seed" ] }, + { + "id": "index_tts2", + "family": "index_tts2", + "case_id": "index_tts2_voice_clone", + "request_id": "clone" + }, { "id": "chatterbox", "family": "chatterbox", diff --git a/tests/perf/model_perf_request_cases.json b/tests/perf/model_perf_request_cases.json index 9c694149..441b68fe 100644 --- a/tests/perf/model_perf_request_cases.json +++ b/tests/perf/model_perf_request_cases.json @@ -354,6 +354,27 @@ } ] }, + { + "id": "index_tts2_voice_clone", + "coverage": "IndexTTS2 English voice clone, HIP/ROCm perf case", + "family": "index_tts2", + "model": "models/IndexTTS-2", + "task": "clon", + "mode": "offline", + "outputs": [ + "audio" + ], + "requests": [ + { + "id": "clone", + "text": "Hey, you. You're finally awake. You were trying to cross the border, right? Walked right into that Imperial ambush, same as us, and that thief over there.", + "language": "English", + "voice_ref": "assets/resources/a.wav", + "seed": 1234, + "do_sample": false + } + ] + }, { "id": "chatterbox_en_long_cfg", "coverage": "Chatterbox English voice clone, T3, S3 tokenizer, CFM/HiFT, CFG two-batch path, long text", diff --git a/tests/unittests/test_conv_lowering_matrix.cpp b/tests/unittests/test_conv_lowering_matrix.cpp new file mode 100644 index 00000000..c68b8644 --- /dev/null +++ b/tests/unittests/test_conv_lowering_matrix.cpp @@ -0,0 +1,844 @@ +#include "engine/framework/core/backend.h" +#include "engine/framework/modules/conv_modules.h" +#include "engine/framework/modules/streaming_conv_modules.h" +#include "engine/framework/modules/structural_modules.h" + +#include +#include + +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include +#include + +namespace { + +constexpr size_t kGraphBytes = 512 * 1024 * 1024; +constexpr size_t kGraphNodes = 16384; +constexpr int kWarmupRounds = 1; +constexpr int kMeasureRounds = 5; + +struct DiffStats { + float max_abs = 0.0f; + double mean_abs = 0.0; + double cosine = 1.0; +}; + +struct RunResult { + bool supported = false; + std::string error; + engine::core::TensorShape shape = {}; + std::vector values; + double avg_ms = 0.0; +}; + +std::vector make_patterned_f32(size_t count, float phase, float scale) { + std::vector values(count, 0.0f); + for (size_t i = 0; i < count; ++i) { + const float x = static_cast(i); + values[i] = scale * ( + std::sin(phase + 0.113f * x) + + 0.5f * std::cos(phase * 0.7f + 0.071f * x)); + } + return values; +} + +int64_t conv_out(int64_t input, int64_t kernel, int stride, int padding, int dilation) { + return (input + 2 * padding - dilation * (kernel - 1) - 1) / stride + 1; +} + +int64_t conv_transpose_out(int64_t input, int64_t kernel, int stride, int padding, int dilation) { + return (input - 1) * stride - 2 * padding + dilation * (kernel - 1) + 1; +} + +const char * backend_name(engine::core::BackendType backend_type) { + switch (backend_type) { + case engine::core::BackendType::Cpu: return "cpu"; + case engine::core::BackendType::Cuda: return "cuda"; +#ifdef ENGINE_TEST_ENABLE_HIP + case engine::core::BackendType::Hip: return "hip"; +#endif + case engine::core::BackendType::Vulkan: return "vulkan"; + case engine::core::BackendType::Metal: return "metal"; + default: return "unknown"; + } +} + +bool same_shape(const engine::core::TensorShape & lhs, const engine::core::TensorShape & rhs) { + if (lhs.rank != rhs.rank) { + return false; + } + for (size_t i = 0; i < lhs.rank; ++i) { + if (lhs.dims[i] != rhs.dims[i]) { + return false; + } + } + return true; +} + +DiffStats diff_values(const std::vector & reference, const std::vector & actual) { + if (reference.size() != actual.size()) { + throw std::runtime_error("value count mismatch"); + } + DiffStats stats; + double dot = 0.0; + double ref_norm = 0.0; + double actual_norm = 0.0; + for (size_t i = 0; i < reference.size(); ++i) { + const float diff = std::fabs(reference[i] - actual[i]); + stats.max_abs = std::max(stats.max_abs, diff); + stats.mean_abs += diff; + dot += static_cast(reference[i]) * static_cast(actual[i]); + ref_norm += static_cast(reference[i]) * static_cast(reference[i]); + actual_norm += static_cast(actual[i]) * static_cast(actual[i]); + } + stats.mean_abs /= static_cast(reference.size()); + if (ref_norm > 0.0 && actual_norm > 0.0) { + stats.cosine = dot / (std::sqrt(ref_norm) * std::sqrt(actual_norm)); + } + return stats; +} + +engine::core::TensorValue add_bias_3d( + engine::core::ModuleBuildContext & ctx, + const engine::core::TensorValue & output, + int64_t channels, + const std::optional & bias) { + if (!bias.has_value()) { + return output; + } + auto output_contiguous = engine::core::ensure_backend_addressable_layout(ctx, output); + auto bias_view = engine::core::reshape_tensor(ctx, *bias, engine::core::TensorShape::from_dims({1, channels, 1})); + auto repeated = engine::core::wrap_tensor( + ggml_repeat(ctx.ggml, bias_view.tensor, output_contiguous.tensor), + output.shape, + GGML_TYPE_F32); + return engine::core::wrap_tensor(ggml_add(ctx.ggml, output_contiguous.tensor, repeated.tensor), output.shape, GGML_TYPE_F32); +} + +engine::core::TensorValue add_bias_4d( + engine::core::ModuleBuildContext & ctx, + const engine::core::TensorValue & output, + int64_t channels, + const std::optional & bias) { + if (!bias.has_value()) { + return output; + } + auto output_contiguous = engine::core::ensure_backend_addressable_layout(ctx, output); + auto bias_view = engine::core::reshape_tensor(ctx, *bias, engine::core::TensorShape::from_dims({1, channels, 1, 1})); + auto repeated = engine::core::wrap_tensor( + ggml_repeat(ctx.ggml, bias_view.tensor, output_contiguous.tensor), + output.shape, + GGML_TYPE_F32); + return engine::core::wrap_tensor(ggml_add(ctx.ggml, output_contiguous.tensor, repeated.tensor), output.shape, GGML_TYPE_F32); +} + +engine::core::TensorValue view_batch_matrix( + engine::core::ModuleBuildContext & ctx, + const engine::core::TensorValue & input, + int64_t batch_index, + int64_t channels, + int64_t frames) { + auto * view = ggml_view_2d( + ctx.ggml, + input.tensor, + frames, + channels, + input.tensor->nb[1], + static_cast(batch_index) * input.tensor->nb[2]); + return engine::core::wrap_tensor(view, engine::core::TensorShape::from_dims({channels, frames}), input.type); +} + +class GraphRunner { +public: + GraphRunner(const char * name, engine::core::BackendType backend_type) : backend_type_(backend_type) { + backend_ = engine::core::init_backend({backend_type, 0, 8}); + engine::core::set_backend_threads(backend_, 8); + ggml_init_params params{}; + params.mem_size = kGraphBytes; + params.mem_buffer = nullptr; + params.no_alloc = true; + ggml_ = ggml_init(params); + if (ggml_ == nullptr) { + throw std::runtime_error("failed to initialize ggml test context"); + } + ctx_.ggml = ggml_; + ctx_.module_instance_name = name; + ctx_.backend_type = backend_type; + } + + ~GraphRunner() { + if (buffer_ != nullptr) { + ggml_backend_buffer_free(buffer_); + } + if (ggml_ != nullptr) { + ggml_free(ggml_); + } + if (backend_ != nullptr) { + ggml_backend_free(backend_); + } + } + + engine::core::TensorValue make_f32(const engine::core::TensorShape & shape) { + return engine::core::make_tensor(ctx_, GGML_TYPE_F32, shape); + } + + engine::core::ModuleBuildContext & ctx() noexcept { return ctx_; } + + RunResult run( + const engine::core::TensorValue & output, + const std::vector>> & writes) { + ggml_cgraph * graph = ggml_new_graph_custom(ggml_, kGraphNodes, false); + ggml_build_forward_expand(graph, output.tensor); + engine::core::validate_backend_graph_supported(backend_, graph, "conv_lowering_matrix"); + buffer_ = ggml_backend_alloc_ctx_tensors(ggml_, backend_); + if (buffer_ == nullptr) { + throw std::runtime_error("failed to allocate backend tensors"); + } + for (const auto & write : writes) { + engine::core::write_tensor_f32(write.first, write.second); + } + for (int i = 0; i < kWarmupRounds; ++i) { + if (ggml_backend_graph_compute(backend_, graph) != GGML_STATUS_SUCCESS) { + throw std::runtime_error("warmup graph compute failed"); + } + } + double total_ms = 0.0; + for (int i = 0; i < kMeasureRounds; ++i) { + const auto start = std::chrono::steady_clock::now(); + if (ggml_backend_graph_compute(backend_, graph) != GGML_STATUS_SUCCESS) { + throw std::runtime_error("graph compute failed"); + } + const auto end = std::chrono::steady_clock::now(); + total_ms += std::chrono::duration(end - start).count(); + } + RunResult result; + result.supported = true; + result.shape = output.shape; + result.avg_ms = total_ms / static_cast(kMeasureRounds); + engine::core::read_tensor_f32_into(output.tensor, result.values); + return result; + } + +private: + engine::core::BackendType backend_type_; + ggml_backend_t backend_ = nullptr; + ggml_backend_buffer_t buffer_ = nullptr; + ggml_context * ggml_ = nullptr; + engine::core::ModuleBuildContext ctx_{}; +}; + +template +RunResult run_guarded( + const char * label, + engine::core::BackendType backend_type, + Fn && fn) { + try { + GraphRunner runner(label, backend_type); + return fn(runner); + } catch (const std::exception & ex) { + RunResult result; + result.supported = false; + result.error = ex.what(); + return result; + } +} + +struct Conv1dCase { + const char * name; + int64_t batch; + int64_t in_channels; + int64_t out_channels; + int64_t frames; + int64_t kernel; + int stride; + int padding; + int dilation; + bool bias; +}; + +RunResult run_conv1d(const Conv1dCase & c, const char * candidate, engine::core::BackendType backend_type) { + return run_guarded(candidate, backend_type, [&](GraphRunner & runner) { + const auto input_shape = engine::core::TensorShape::from_dims({c.batch, c.in_channels, c.frames}); + const auto weight_shape = engine::core::TensorShape::from_dims({c.out_channels, c.in_channels, c.kernel}); + const auto bias_shape = engine::core::TensorShape::from_dims({c.out_channels}); + auto input = runner.make_f32(input_shape); + auto weight = runner.make_f32(weight_shape); + std::optional bias = c.bias ? std::optional(runner.make_f32(bias_shape)) : std::nullopt; + + engine::core::TensorValue output; + if (std::string(candidate) == "native") { + const auto output_shape = engine::core::TensorShape::from_dims( + {c.batch, c.out_channels, conv_out(c.frames, c.kernel, c.stride, c.padding, c.dilation)}); + if (c.batch == 1) { + output = engine::core::wrap_tensor( + ggml_conv_1d(runner.ctx().ggml, weight.tensor, input.tensor, c.stride, c.padding, c.dilation), + output_shape, + GGML_TYPE_F32); + } else { + for (int64_t batch = 0; batch < c.batch; ++batch) { + auto batch_input = view_batch_matrix(runner.ctx(), input, batch, c.in_channels, c.frames); + auto batch_output = engine::core::wrap_tensor( + ggml_conv_1d(runner.ctx().ggml, weight.tensor, batch_input.tensor, c.stride, c.padding, c.dilation), + engine::core::TensorShape::from_dims({1, c.out_channels, output_shape.dims[2]}), + GGML_TYPE_F32); + output = output.valid() ? engine::modules::ConcatModule({0}).build(runner.ctx(), output, batch_output) : batch_output; + } + } + output = add_bias_3d(runner.ctx(), output, c.out_channels, bias); + } else if (std::string(candidate) == "conv2d_normal") { + auto x4 = engine::core::reshape_tensor(runner.ctx(), input, engine::core::TensorShape::from_dims({c.batch, c.in_channels, 1, c.frames})); + auto w4 = engine::core::reshape_tensor(runner.ctx(), weight, engine::core::TensorShape::from_dims({c.out_channels, c.in_channels, 1, c.kernel})); + auto y4 = engine::core::wrap_tensor( + ggml_conv_2d(runner.ctx().ggml, w4.tensor, x4.tensor, c.stride, 1, c.padding, 0, c.dilation, 1), + engine::core::TensorShape::from_dims({c.batch, c.out_channels, 1, conv_out(c.frames, c.kernel, c.stride, c.padding, c.dilation)}), + GGML_TYPE_F32); + y4 = add_bias_4d(runner.ctx(), y4, c.out_channels, bias); + output = engine::core::reshape_tensor(runner.ctx(), y4, engine::core::TensorShape::from_dims({c.batch, c.out_channels, y4.shape.dims[3]})); + } else if (std::string(candidate) == "conv2d_direct") { + auto x4 = engine::core::reshape_tensor(runner.ctx(), input, engine::core::TensorShape::from_dims({c.batch, c.in_channels, 1, c.frames})); + auto w4 = engine::core::reshape_tensor(runner.ctx(), weight, engine::core::TensorShape::from_dims({c.out_channels, c.in_channels, 1, c.kernel})); + auto y4 = engine::core::wrap_tensor( + ggml_conv_2d_direct(runner.ctx().ggml, w4.tensor, x4.tensor, c.stride, 1, c.padding, 0, c.dilation, 1), + engine::core::TensorShape::from_dims({c.batch, c.out_channels, 1, conv_out(c.frames, c.kernel, c.stride, c.padding, c.dilation)}), + GGML_TYPE_F32); + y4 = add_bias_4d(runner.ctx(), y4, c.out_channels, bias); + output = engine::core::reshape_tensor(runner.ctx(), y4, engine::core::TensorShape::from_dims({c.batch, c.out_channels, y4.shape.dims[3]})); + } else { + throw std::runtime_error("unknown conv1d candidate"); + } + + std::vector>> writes; + writes.push_back({input, make_patterned_f32(static_cast(input_shape.num_elements()), 0.19f, 0.031f)}); + writes.push_back({weight, make_patterned_f32(static_cast(weight_shape.num_elements()), 0.47f, 0.017f)}); + if (bias) { + writes.push_back({*bias, make_patterned_f32(static_cast(bias_shape.num_elements()), 0.83f, 0.011f)}); + } + return runner.run(output, writes); + }); +} + +struct Conv2dCase { + const char * name; + int64_t batch; + int64_t in_channels; + int64_t out_channels; + int64_t height; + int64_t width; + int64_t kernel_h; + int64_t kernel_w; + int stride_h; + int stride_w; + int padding_h; + int padding_w; + int dilation_h; + int dilation_w; + bool bias; +}; + +RunResult run_conv2d(const Conv2dCase & c, const char * candidate, engine::core::BackendType backend_type) { + return run_guarded(candidate, backend_type, [&](GraphRunner & runner) { + const auto input_shape = engine::core::TensorShape::from_dims({c.batch, c.in_channels, c.height, c.width}); + const auto weight_shape = engine::core::TensorShape::from_dims({c.out_channels, c.in_channels, c.kernel_h, c.kernel_w}); + const auto bias_shape = engine::core::TensorShape::from_dims({c.out_channels}); + auto input = runner.make_f32(input_shape); + auto weight = runner.make_f32(weight_shape); + std::optional bias = c.bias ? std::optional(runner.make_f32(bias_shape)) : std::nullopt; + + engine::core::TensorValue output; + const auto output_shape = engine::core::TensorShape::from_dims({ + c.batch, + c.out_channels, + conv_out(c.height, c.kernel_h, c.stride_h, c.padding_h, c.dilation_h), + conv_out(c.width, c.kernel_w, c.stride_w, c.padding_w, c.dilation_w), + }); + if (std::string(candidate) == "im2col_matmul") { + output = engine::core::wrap_tensor( + ggml_conv_2d(runner.ctx().ggml, weight.tensor, input.tensor, c.stride_w, c.stride_h, c.padding_w, c.padding_h, c.dilation_w, c.dilation_h), + output_shape, + GGML_TYPE_F32); + output = add_bias_4d(runner.ctx(), output, c.out_channels, bias); + } else if (std::string(candidate) == "direct") { + output = engine::core::wrap_tensor( + ggml_conv_2d_direct(runner.ctx().ggml, weight.tensor, input.tensor, c.stride_w, c.stride_h, c.padding_w, c.padding_h, c.dilation_w, c.dilation_h), + output_shape, + GGML_TYPE_F32); + output = add_bias_4d(runner.ctx(), output, c.out_channels, bias); + } else { + throw std::runtime_error("unknown conv2d candidate"); + } + std::vector>> writes; + writes.push_back({input, make_patterned_f32(static_cast(input_shape.num_elements()), 0.21f, 0.021f)}); + writes.push_back({weight, make_patterned_f32(static_cast(weight_shape.num_elements()), 0.51f, 0.013f)}); + if (bias) { + writes.push_back({*bias, make_patterned_f32(static_cast(bias_shape.num_elements()), 0.91f, 0.009f)}); + } + return runner.run(output, writes); + }); +} + +struct Depthwise1dCase { + const char * name; + int64_t batch; + int64_t channels; + int64_t frames; + int64_t kernel; + int stride; + int padding; + int dilation; + bool bias; +}; + +RunResult run_depthwise1d(const Depthwise1dCase & c, const char * candidate, engine::core::BackendType backend_type) { + return run_guarded(candidate, backend_type, [&](GraphRunner & runner) { + const auto input_shape = engine::core::TensorShape::from_dims({c.batch, c.channels, c.frames}); + const auto weight_shape = engine::core::TensorShape::from_dims({c.channels, 1, c.kernel}); + const auto bias_shape = engine::core::TensorShape::from_dims({c.channels}); + auto input = runner.make_f32(input_shape); + auto weight = runner.make_f32(weight_shape); + std::optional bias = c.bias ? std::optional(runner.make_f32(bias_shape)) : std::nullopt; + engine::core::TensorValue output; + if (std::string(candidate) == "dw2d_direct") { + auto x4 = engine::core::reshape_tensor(runner.ctx(), input, engine::core::TensorShape::from_dims({c.batch, c.channels, 1, c.frames})); + auto w4 = engine::core::reshape_tensor(runner.ctx(), weight, engine::core::TensorShape::from_dims({c.channels, 1, 1, c.kernel})); + auto y4 = engine::core::wrap_tensor( + ggml_conv_2d_dw_direct(runner.ctx().ggml, w4.tensor, x4.tensor, c.stride, 1, c.padding, 0, c.dilation, 1), + engine::core::TensorShape::from_dims({c.batch, c.channels, 1, conv_out(c.frames, c.kernel, c.stride, c.padding, c.dilation)}), + GGML_TYPE_F32); + y4 = add_bias_4d(runner.ctx(), y4, c.channels, bias); + output = engine::core::reshape_tensor(runner.ctx(), y4, engine::core::TensorShape::from_dims({c.batch, c.channels, y4.shape.dims[3]})); + } else if (std::string(candidate) == "native_1d_dw") { + if (c.batch != 1) { + throw std::runtime_error("native ggml_conv_1d_dw asserts for batched rank-3 input; slice batch first"); + } + output = engine::core::wrap_tensor( + ggml_conv_1d_dw(runner.ctx().ggml, weight.tensor, input.tensor, c.stride, c.padding, c.dilation), + engine::core::TensorShape::from_dims({c.batch, c.channels, conv_out(c.frames, c.kernel, c.stride, c.padding, c.dilation)}), + GGML_TYPE_F32); + output = add_bias_3d(runner.ctx(), output, c.channels, bias); + } else { + throw std::runtime_error("unknown depthwise1d candidate"); + } + std::vector>> writes; + writes.push_back({input, make_patterned_f32(static_cast(input_shape.num_elements()), 0.23f, 0.025f)}); + writes.push_back({weight, make_patterned_f32(static_cast(weight_shape.num_elements()), 0.53f, 0.015f)}); + if (bias) { + writes.push_back({*bias, make_patterned_f32(static_cast(bias_shape.num_elements()), 0.93f, 0.007f)}); + } + return runner.run(output, writes); + }); +} + +struct Pointwise1dCase { + const char * name; + int64_t batch; + int64_t in_channels; + int64_t out_channels; + int64_t frames; + bool bias; +}; + +RunResult run_pointwise1d(const Pointwise1dCase & c, const char * candidate, engine::core::BackendType backend_type) { + return run_guarded(candidate, backend_type, [&](GraphRunner & runner) { + const auto input_shape = engine::core::TensorShape::from_dims({c.batch, c.in_channels, c.frames}); + const auto weight_shape = engine::core::TensorShape::from_dims({c.out_channels, c.in_channels, 1}); + const auto bias_shape = engine::core::TensorShape::from_dims({c.out_channels}); + auto input = runner.make_f32(input_shape); + auto weight = runner.make_f32(weight_shape); + std::optional bias = c.bias ? std::optional(runner.make_f32(bias_shape)) : std::nullopt; + + engine::core::TensorValue output; + if (std::string(candidate) == "conv1d_kernel1") { + const auto output_shape = engine::core::TensorShape::from_dims({c.batch, c.out_channels, c.frames}); + if (c.batch == 1) { + output = engine::core::wrap_tensor( + ggml_conv_1d(runner.ctx().ggml, weight.tensor, input.tensor, 1, 0, 1), + output_shape, + GGML_TYPE_F32); + } else { + for (int64_t batch = 0; batch < c.batch; ++batch) { + auto batch_input = view_batch_matrix(runner.ctx(), input, batch, c.in_channels, c.frames); + auto batch_output = engine::core::wrap_tensor( + ggml_conv_1d(runner.ctx().ggml, weight.tensor, batch_input.tensor, 1, 0, 1), + engine::core::TensorShape::from_dims({1, c.out_channels, c.frames}), + GGML_TYPE_F32); + output = output.valid() ? engine::modules::ConcatModule({0}).build(runner.ctx(), output, batch_output) : batch_output; + } + } + output = add_bias_3d(runner.ctx(), output, c.out_channels, bias); + } else if (std::string(candidate) == "linear_matmul") { + auto x = engine::modules::TransposeModule({{0, 2, 1, 3}, 3}).build(runner.ctx(), input); + x = engine::core::ensure_backend_addressable_layout(runner.ctx(), x); + auto matrix = engine::core::reshape_tensor(runner.ctx(), x, engine::core::TensorShape::from_dims({c.batch * c.frames, c.in_channels})); + auto w2 = engine::core::reshape_tensor(runner.ctx(), weight, engine::core::TensorShape::from_dims({c.out_channels, c.in_channels})); + auto projected = engine::core::wrap_tensor( + ggml_mul_mat(runner.ctx().ggml, w2.tensor, matrix.tensor), + engine::core::TensorShape::from_dims({c.batch * c.frames, c.out_channels}), + GGML_TYPE_F32); + if (bias) { + projected = engine::core::wrap_tensor(ggml_add(runner.ctx().ggml, projected.tensor, bias->tensor), projected.shape, GGML_TYPE_F32); + } + auto y = engine::core::reshape_tensor(runner.ctx(), projected, engine::core::TensorShape::from_dims({c.batch, c.frames, c.out_channels})); + output = engine::modules::TransposeModule({{0, 2, 1, 3}, 3}).build(runner.ctx(), y); + } else { + throw std::runtime_error("unknown pointwise1d candidate"); + } + std::vector>> writes; + writes.push_back({input, make_patterned_f32(static_cast(input_shape.num_elements()), 0.24f, 0.027f)}); + writes.push_back({weight, make_patterned_f32(static_cast(weight_shape.num_elements()), 0.54f, 0.014f)}); + if (bias) { + writes.push_back({*bias, make_patterned_f32(static_cast(bias_shape.num_elements()), 0.94f, 0.007f)}); + } + return runner.run(output, writes); + }); +} + +RunResult run_depthwise2d(const Conv2dCase & c, const char * candidate, engine::core::BackendType backend_type) { + return run_guarded(candidate, backend_type, [&](GraphRunner & runner) { + const auto input_shape = engine::core::TensorShape::from_dims({c.batch, c.in_channels, c.height, c.width}); + const auto weight_shape = engine::core::TensorShape::from_dims({c.in_channels, 1, c.kernel_h, c.kernel_w}); + const auto bias_shape = engine::core::TensorShape::from_dims({c.in_channels}); + auto input = runner.make_f32(input_shape); + auto weight = runner.make_f32(weight_shape); + std::optional bias = c.bias ? std::optional(runner.make_f32(bias_shape)) : std::nullopt; + const auto output_shape = engine::core::TensorShape::from_dims({ + c.batch, + c.in_channels, + conv_out(c.height, c.kernel_h, c.stride_h, c.padding_h, c.dilation_h), + conv_out(c.width, c.kernel_w, c.stride_w, c.padding_w, c.dilation_w), + }); + engine::core::TensorValue output; + if (std::string(candidate) == "direct") { + output = engine::core::wrap_tensor( + ggml_conv_2d_dw_direct(runner.ctx().ggml, weight.tensor, input.tensor, c.stride_w, c.stride_h, c.padding_w, c.padding_h, c.dilation_w, c.dilation_h), + output_shape, + GGML_TYPE_F32); + output = add_bias_4d(runner.ctx(), output, c.in_channels, bias); + } else if (std::string(candidate) == "im2col_matmul") { + output = engine::core::wrap_tensor( + ggml_conv_2d_dw(runner.ctx().ggml, weight.tensor, input.tensor, c.stride_w, c.stride_h, c.padding_w, c.padding_h, c.dilation_w, c.dilation_h), + output_shape, + GGML_TYPE_F32); + output = add_bias_4d(runner.ctx(), output, c.in_channels, bias); + } else { + throw std::runtime_error("unknown depthwise2d candidate"); + } + std::vector>> writes; + writes.push_back({input, make_patterned_f32(static_cast(input_shape.num_elements()), 0.25f, 0.023f)}); + writes.push_back({weight, make_patterned_f32(static_cast(weight_shape.num_elements()), 0.55f, 0.012f)}); + if (bias) { + writes.push_back({*bias, make_patterned_f32(static_cast(bias_shape.num_elements()), 0.95f, 0.008f)}); + } + return runner.run(output, writes); + }); +} + +struct ConvTranspose1dCase { + const char * name; + int64_t batch; + int64_t in_channels; + int64_t out_channels; + int64_t frames; + int64_t kernel; + int stride; + int padding; + int dilation; + bool bias; +}; + +engine::core::TensorValue build_conv_transpose_native( + engine::core::ModuleBuildContext & ctx, + const ConvTranspose1dCase & c, + const engine::core::TensorValue & input, + const engine::core::TensorValue & weight, + const std::optional & bias) { + if (c.padding != 0 || c.dilation != 1) { + throw std::runtime_error("native ggml_conv_transpose_1d supports only padding=0 and dilation=1"); + } + engine::core::TensorValue output; + for (int64_t batch = 0; batch < c.batch; ++batch) { + auto matrix = view_batch_matrix(ctx, input, batch, c.in_channels, c.frames); + auto batch_out = engine::core::wrap_tensor( + ggml_conv_transpose_1d(ctx.ggml, weight.tensor, matrix.tensor, c.stride, c.padding, c.dilation), + engine::core::TensorShape::from_dims({1, c.out_channels, conv_transpose_out(c.frames, c.kernel, c.stride, c.padding, c.dilation)}), + GGML_TYPE_F32); + output = output.valid() ? engine::modules::ConcatModule({0}).build(ctx, output, batch_out) : batch_out; + } + return add_bias_3d(ctx, output, c.out_channels, bias); +} + +engine::core::TensorValue build_conv_transpose_col2im( + engine::core::ModuleBuildContext & ctx, + const ConvTranspose1dCase & c, + const engine::core::TensorValue & input, + const engine::core::TensorValue & weight, + const std::optional & bias) { + if (c.dilation != 1) { + throw std::runtime_error("col2im lowering currently supports only dilation=1"); + } + auto * weight_perm = ggml_reshape_2d( + ctx.ggml, + ggml_cont(ctx.ggml, ggml_permute(ctx.ggml, weight.tensor, 1, 2, 0, 3)), + c.in_channels, + c.kernel * c.out_channels); + ggml_tensor * bias_matrix = nullptr; + if (c.bias) { + if (!bias.has_value()) { + throw std::runtime_error("missing bias"); + } + bias_matrix = ggml_reshape_2d(ctx.ggml, bias->tensor, 1, c.out_channels); + } + engine::core::TensorValue output; + for (int64_t batch = 0; batch < c.batch; ++batch) { + auto * batch_input = ggml_view_2d( + ctx.ggml, + input.tensor, + input.tensor->ne[0], + input.tensor->ne[1], + input.tensor->nb[1], + static_cast(batch) * input.tensor->nb[2]); + auto * transposed_input = ggml_cont(ctx.ggml, ggml_transpose(ctx.ggml, batch_input)); + auto * columns = ggml_mul_mat(ctx.ggml, weight_perm, transposed_input); + auto * batch_output = ggml_col2im_1d(ctx.ggml, columns, c.stride, static_cast(c.out_channels), c.padding); + if (bias_matrix != nullptr) { + batch_output = ggml_add(ctx.ggml, batch_output, bias_matrix); + } + auto batch_value = engine::core::wrap_tensor( + ggml_reshape_3d(ctx.ggml, batch_output, batch_output->ne[0], batch_output->ne[1], 1), + engine::core::TensorShape::from_dims({1, c.out_channels, batch_output->ne[0]}), + GGML_TYPE_F32); + output = output.valid() ? engine::modules::ConcatModule({0}).build(ctx, output, batch_value) : batch_value; + } + return output; +} + +RunResult run_conv_transpose1d(const ConvTranspose1dCase & c, const char * candidate, engine::core::BackendType backend_type) { + if (backend_type == engine::core::BackendType::Cpu && std::string(candidate) == "matmul_col2im") { + RunResult result; + result.supported = false; + result.error = "current ggml CPU backend aborts for COL2IM_1D"; + return result; + } + return run_guarded(candidate, backend_type, [&](GraphRunner & runner) { + const auto input_shape = engine::core::TensorShape::from_dims({c.batch, c.in_channels, c.frames}); + const auto weight_shape = engine::core::TensorShape::from_dims({c.in_channels, c.out_channels, c.kernel}); + const auto bias_shape = engine::core::TensorShape::from_dims({c.out_channels}); + auto input = runner.make_f32(input_shape); + auto weight = runner.make_f32(weight_shape); + std::optional bias = c.bias ? std::optional(runner.make_f32(bias_shape)) : std::nullopt; + engine::core::TensorValue output; + if (std::string(candidate) == "native_direct") { + output = build_conv_transpose_native(runner.ctx(), c, input, weight, bias); + } else if (std::string(candidate) == "matmul_col2im") { + output = build_conv_transpose_col2im(runner.ctx(), c, input, weight, bias); + } else { + throw std::runtime_error("unknown conv_transpose1d candidate"); + } + std::vector>> writes; + writes.push_back({input, make_patterned_f32(static_cast(input_shape.num_elements()), 0.27f, 0.019f)}); + writes.push_back({weight, make_patterned_f32(static_cast(weight_shape.num_elements()), 0.57f, 0.011f)}); + if (bias) { + writes.push_back({*bias, make_patterned_f32(static_cast(bias_shape.num_elements()), 0.97f, 0.006f)}); + } + return runner.run(output, writes); + }); +} + +struct MatrixRow { + std::string module; + std::string case_name; + std::string candidate; + engine::core::BackendType backend; + RunResult result; + std::optional diff; +}; + +void print_row(const MatrixRow & row) { + std::cout << "| " << row.module + << " | " << row.case_name + << " | " << row.candidate + << " | " << backend_name(row.backend) + << " | "; + if (!row.result.supported) { + std::string error = row.result.error; + std::replace(error.begin(), error.end(), '|', '/'); + std::cout << "unsupported | - | - | - | - | " << error << " |\n"; + return; + } + std::cout << "ok | " << row.result.shape.to_string() + << " | " << std::fixed << std::setprecision(4) << row.result.avg_ms + << " | "; + if (row.diff.has_value()) { + std::cout << std::scientific << std::setprecision(3) << row.diff->max_abs + << " | " << row.diff->mean_abs + << " | " << std::fixed << std::setprecision(9) << row.diff->cosine << " |\n"; + } else { + std::cout << "- | - | - |\n"; + } +} + +void add_result( + std::vector & rows, + const std::string & module, + const std::string & case_name, + const std::string & candidate, + engine::core::BackendType backend, + const RunResult & result, + const RunResult & reference) { + MatrixRow row{module, case_name, candidate, backend, result, std::nullopt}; + if (result.supported && reference.supported) { + if (!same_shape(reference.shape, result.shape)) { + row.result.supported = false; + row.result.error = "shape mismatch vs reference " + reference.shape.to_string(); + } else { + row.diff = diff_values(reference.values, result.values); + } + } + rows.push_back(std::move(row)); +} + +bool backend_available(engine::core::BackendType backend_type) { + try { + GraphRunner runner("conv_lowering_matrix.probe", backend_type); + return true; + } catch (...) { + return false; + } +} + +} // namespace + +int main() { + try { + std::vector backends = {engine::core::BackendType::Cpu}; + if (backend_available(engine::core::BackendType::Cuda)) { + backends.push_back(engine::core::BackendType::Cuda); + } else { + std::cout << "[SKIP] cuda backend unavailable\n"; + } +#ifdef ENGINE_TEST_ENABLE_HIP + if (backend_available(engine::core::BackendType::Hip)) { + backends.push_back(engine::core::BackendType::Hip); + } else { + std::cout << "[SKIP] hip backend unavailable\n"; + } +#endif + if (backend_available(engine::core::BackendType::Vulkan)) { + backends.push_back(engine::core::BackendType::Vulkan); + } else { + std::cout << "[SKIP] vulkan backend unavailable\n"; + } + + std::vector rows; + + const std::vector conv1d_cases = { + {"citrinet_like_large_regular", 1, 80, 256, 256, 11, 1, 5, 1, true}, + {"bigvgan_like_resblock", 1, 192, 192, 384, 7, 1, 3, 1, true}, + {"batched_stride_regular", 2, 64, 128, 160, 5, 2, 2, 1, true}, + {"dilated_regular", 1, 128, 128, 192, 3, 1, 2, 2, false}, + }; + const std::vector conv1d_candidates = {"native", "conv2d_normal", "conv2d_direct"}; + for (const auto & c : conv1d_cases) { + const auto reference = run_conv1d(c, "native", engine::core::BackendType::Cpu); + for (const auto backend : backends) { + for (const auto & candidate : conv1d_candidates) { + add_result(rows, "Conv1dModule", c.name, candidate, backend, run_conv1d(c, candidate.c_str(), backend), reference); + } + } + } + + const std::vector conv2d_cases = { + {"spectrogram_small_kernel", 1, 64, 128, 20, 160, 3, 3, 1, 1, 1, 1, 1, 1, true}, + {"conv1d_lowered_shape", 1, 256, 256, 1, 384, 1, 7, 1, 1, 0, 3, 1, 1, true}, + {"batched_feature_map", 2, 32, 64, 12, 96, 3, 5, 1, 2, 1, 2, 1, 1, false}, + }; + const std::vector conv2d_candidates = {"im2col_matmul", "direct"}; + for (const auto & c : conv2d_cases) { + const auto reference = run_conv2d(c, "im2col_matmul", engine::core::BackendType::Cpu); + for (const auto backend : backends) { + for (const auto & candidate : conv2d_candidates) { + add_result(rows, "Conv2dModule", c.name, candidate, backend, run_conv2d(c, candidate.c_str(), backend), reference); + } + } + } + + const std::vector depthwise1d_cases = { + {"conformer_like_depthwise", 1, 256, 192, 31, 1, 15, 1, true}, + {"tokenizer_stride_depthwise", 2, 96, 256, 7, 2, 3, 1, true}, + {"dilated_depthwise", 1, 128, 160, 5, 1, 4, 2, false}, + }; + const std::vector depthwise1d_candidates = {"dw2d_direct", "native_1d_dw"}; + for (const auto & c : depthwise1d_cases) { + const auto reference = run_depthwise1d(c, "dw2d_direct", engine::core::BackendType::Cpu); + for (const auto backend : backends) { + for (const auto & candidate : depthwise1d_candidates) { + add_result(rows, "DepthwiseConv1dModule", c.name, candidate, backend, run_depthwise1d(c, candidate.c_str(), backend), reference); + } + } + } + + const std::vector pointwise1d_cases = { + {"conformer_projection", 1, 256, 512, 192, true}, + {"batched_token_projection", 2, 192, 384, 160, true}, + {"vocoder_channel_mix", 1, 192, 192, 384, false}, + }; + const std::vector pointwise1d_candidates = {"conv1d_kernel1", "linear_matmul"}; + for (const auto & c : pointwise1d_cases) { + const auto reference = run_pointwise1d(c, "conv1d_kernel1", engine::core::BackendType::Cpu); + for (const auto backend : backends) { + for (const auto & candidate : pointwise1d_candidates) { + add_result(rows, "PointwiseConv1dModule", c.name, candidate, backend, run_pointwise1d(c, candidate.c_str(), backend), reference); + } + } + } + + const std::vector depthwise2d_cases = { + {"depthwise_1d_lowered_shape", 1, 192, 192, 1, 384, 1, 7, 1, 1, 0, 3, 1, 1, true}, + {"image_depthwise_small", 1, 64, 64, 24, 80, 3, 3, 1, 1, 1, 1, 1, 1, true}, + }; + const std::vector depthwise2d_candidates = {"direct", "im2col_matmul"}; + for (const auto & c : depthwise2d_cases) { + const auto reference = run_depthwise2d(c, "direct", engine::core::BackendType::Cpu); + for (const auto backend : backends) { + for (const auto & candidate : depthwise2d_candidates) { + add_result(rows, "DepthwiseConv2dModule", c.name, candidate, backend, run_depthwise2d(c, candidate.c_str(), backend), reference); + } + } + } + + const std::vector conv_transpose_cases = { + {"qwen3_like_stride5_padding0", 1, 256, 128, 96, 10, 5, 0, 1, true}, + {"vocoder_stride2_padding1", 1, 192, 96, 192, 4, 2, 1, 1, true}, + {"batched_stride2_no_bias", 2, 128, 128, 96, 2, 2, 0, 1, false}, + {"dilated_unsupported_probe", 1, 64, 64, 96, 3, 2, 0, 2, true}, + }; + const std::vector conv_transpose_candidates = {"native_direct", "matmul_col2im"}; + for (const auto & c : conv_transpose_cases) { + const auto reference = run_conv_transpose1d(c, "native_direct", engine::core::BackendType::Cpu); + for (const auto backend : backends) { + for (const auto & candidate : conv_transpose_candidates) { + add_result(rows, "ConvTranspose1dModule", c.name, candidate, backend, run_conv_transpose1d(c, candidate.c_str(), backend), reference); + } + } + } + + std::cout << "| module | case | candidate | backend | status | shape | avg_ms | max_abs_vs_cpu_ref | mean_abs_vs_cpu_ref | cosine_vs_cpu_ref |\n"; + std::cout << "| --- | --- | --- | --- | --- | --- | ---: | ---: | ---: | ---: |\n"; + for (const auto & row : rows) { + print_row(row); + } + } catch (const std::exception & ex) { + std::cerr << "[FAIL] " << ex.what() << '\n'; + return 1; + } + return 0; +}