From 9e37a1b4f1b07e088de95f003280e21005b63bd7 Mon Sep 17 00:00:00 2001 From: 0xShug0 <231717474+0xShug0@users.noreply.github.com> Date: Sat, 12 Sep 2026 16:42:00 -0400 Subject: [PATCH 1/3] Release v0.7.4 --- CONTRIBUTING.md | 73 +++++++++++++++++++++++++++++++------------------ README.md | 6 ++-- 2 files changed, 49 insertions(+), 30 deletions(-) diff --git a/CONTRIBUTING.md b/CONTRIBUTING.md index b75869fad..81cad80f6 100644 --- a/CONTRIBUTING.md +++ b/CONTRIBUTING.md @@ -9,30 +9,46 @@ High-impact areas right now: - UI and app-facing workflows - API server behavior, especially OpenAI-compatible serving - Pipeline and workflow subsystem improvements +- Community model ports +- Extending existing model families with additional variants - Model documentation, examples, and validation reports - Cross-platform build and packaging polish -These areas help many model families at once. If you are unsure where to start, improving one of these shared surfaces is usually more valuable than adding another copy of an already-supported model. +> [!TIP] +> **High-impact optimization work:** as the framework evolves, newer models often gain better performance and memory-management patterns that older integrations have not yet adopted. Backporting those patterns is very useful, but it must be done model by model with careful regression tests for output quality, performance, and memory use. -## Before Adding a Model +These areas either improve shared surfaces used by many model families or make existing model support more complete. For already-supported model families, please add new variants to the existing implementation instead of creating a duplicate port. + +## Pull Request Requirements + +Keep PRs focused. A model port, a server change, a pipeline change, and a broad refactor are easier to review and validate when they are separate. + +If a model port needs framework or runtime changes, split them when possible: -Please check the supported model table in [README.md](README.md) before starting a new model port. Some model families are already released, and others are implemented but still marked as testing while they are validated, polished, or promoted into the broader released surface. +- Additive framework/runtime support can be included if it is clearly scoped, opt-in, and genuinely reusable by other models. Do not promote model-specific logic into the framework just because it is convenient; prefer a model-local implementation unless the shared abstraction is clearly useful beyond one model. +- Non-additive framework/runtime behavior changes should be a separate PR before the model port depends on them. +- Backend or GGML changes should be separate PRs so regressions (if any) are easier to trace. -If you want to add support for a model family that is already listed, please focus on improving the existing implementation instead of opening a duplicate port. +If a PR changes runtime behavior, model implementation, backend logic, performance paths, or shared framework code, include actual runtime validation. Code inspection, generated tests, or "this should work" are not enough by themselves. + +For optimization PRs, include before/after measurements. For behavior changes, explain whether outputs, performance, memory, or only docs/build wiring are affected, and include regression checks for the affected route. + +**Please keep no more than 3 PRs open at the same time, including draft PRs.** + +## Before Adding a Model -When a loader is registered (or parked), keep the **package catalog** in sync. Installable `ModelPackage` entries must not advertise families that `audiocpp_cli --list-loaders` does not expose. Follow the checklist in [docs/maintainers/loader_and_catalog.md](docs/maintainers/loader_and_catalog.md) and run: +Please check the supported model table in [README.md](README.md), the community model docs, and open PRs/issues before starting a new model port. Many requests are better handled by extending an existing family with a new variant, route, GGUF package, backend fix, or UI surface. -```bash -python3 tools/check_loader_catalog_sync.py --self-test -python3 tools/check_loader_catalog_sync.py -``` +If the model family is already listed, build on the existing implementation instead of opening a duplicate port. Add the new variant to the same family/spec/package layout unless there is a real architecture or task boundary that requires a separate family. + +For new or updated model packages, keep the loader, model spec, docs, and UI/model-manager package entries aligned. A package should not advertise a model that the CLI/server cannot actually load, and a spec should not claim options or routes the implementation does not support. -Do not leave a live Hugging Face `SnapshotSource` for a loader that is commented out of `registry.cpp` — mark it `UnsupportedSource` (or remove it) and update the README package table. Good follow-up work for existing model families includes: - Better CLI or server examples - More complete path tests - Clearer model-manager package entries +- Additional GGUF variants or model-size variants - Backend coverage improvements - Memory, latency, or portability improvements - Documentation for real user workflows @@ -41,6 +57,10 @@ Good follow-up work for existing model families includes: New standalone model ports should normally start under `community_models/`. This keeps ownership clear and lets useful model ports land with a lighter review bar than core framework models. Models can graduate into the core model tree later after they are validated, polished, and maintained as part of the main release surface. +Please read [Welcoming Community Models & Using Model Spec v1 for New Models](https://github.com/0xShug0/audio.cpp/issues/54) before starting a new model PR. + +New models should use model spec v1, follow existing model/package naming conventions, and reuse framework modules, runtimes, helpers, and ownership patterns when they fit. + Even for community models, PRs should include enough evidence for maintainers and users to understand exactly what was tested. Follow the validation style shown in [PR #19](https://github.com/0xShug0/audio.cpp/pull/19) and [PR #63](https://github.com/0xShug0/audio.cpp/pull/63). Please include: @@ -66,14 +86,6 @@ Framework modules are a high-impact but higher-risk contribution area because th If your contribution is a variant of an existing module, add it as a separate experimental module, for example `xxxExp`, instead of branching inside or rewriting the existing module. The new module can replace the existing one later, after it has shown no regressions across all models that rely on the current implementation. -## Pull Request Notes - -Keep PRs focused. A model port, a server change, a pipeline change, and a broad refactor are easier to review and validate when they are separate. - -When changing shared framework code, explain which model families or routes were checked. When changing model behavior, explain whether the change affects outputs, performance, memory, or only docs/build wiring. - -If a PR intentionally leaves a model under testing, say what remains before it should be marked released. - ## Acknowledgments audio.cpp is moving faster because people keep showing up with real fixes, careful testing, and useful pressure on the parts that matter. Thank you to: @@ -92,8 +104,8 @@ audio.cpp is moving faster because people keep showing up with real fixes, caref - [@xashr](https://github.com/xashr) for Dockerfiles, Docker examples, Docker documentation, portable CPU Docker builds, published-image docs, GGUF/Qwen3-TTS examples, CI workflow polish, CUDA probe cleanup, and portable CUDA architecture selection in [#30](https://github.com/0xShug0/audio.cpp/pull/30), [#51](https://github.com/0xShug0/audio.cpp/pull/51), [#81](https://github.com/0xShug0/audio.cpp/pull/81), [#84](https://github.com/0xShug0/audio.cpp/pull/84), [#107](https://github.com/0xShug0/audio.cpp/pull/107), [#109](https://github.com/0xShug0/audio.cpp/pull/109), and [#280](https://github.com/0xShug0/audio.cpp/pull/280). - [@5uck1ess](https://github.com/5uck1ess) for improving Citrinet CTC decoding through the SentencePiece model, hardening PocketTTS FlowLM step allocation, and adding live PCM transcription ingest to the server in [#49](https://github.com/0xShug0/audio.cpp/pull/49), [#59](https://github.com/0xShug0/audio.cpp/pull/59), and [#144](https://github.com/0xShug0/audio.cpp/pull/144). - [@dkruyt](https://github.com/dkruyt) for the first multipart transcription upload support in [#25](https://github.com/0xShug0/audio.cpp/pull/25). -- [@CaptainArni](https://github.com/CaptainArni) for fixing PocketTTS empty output when switching cached voices, keeping the Windows CUDA build path healthy, and adding ACE-Step 1.5 XL DiT variants in [#22](https://github.com/0xShug0/audio.cpp/pull/22), [#93](https://github.com/0xShug0/audio.cpp/pull/93), and [#235](https://github.com/0xShug0/audio.cpp/pull/235). -- [@IIIIIllllIIIIIlllll](https://github.com/IIIIIllllIIIIIlllll) for the experimental ROCm/HIP backend, Linux HIP build path, Windows HIP/ROCm distribution preparation, HIP build documentation, VibeVoice HIP enablement, backend device listing, Vulkan AMD integer-dot guard, IndexTTS-2.5 support, base64 voice references in server requests, IndexTTS2 speech-rate controls, IndexTTS2 HIP/text-normalization fixes, and BreezeTTS 2 performance and VRAM improvements in [#48](https://github.com/0xShug0/audio.cpp/pull/48), [#148](https://github.com/0xShug0/audio.cpp/pull/148), [#153](https://github.com/0xShug0/audio.cpp/pull/153), [#159](https://github.com/0xShug0/audio.cpp/pull/159), [#164](https://github.com/0xShug0/audio.cpp/pull/164), [#168](https://github.com/0xShug0/audio.cpp/pull/168), [#171](https://github.com/0xShug0/audio.cpp/pull/171), [#193](https://github.com/0xShug0/audio.cpp/pull/193), [#210](https://github.com/0xShug0/audio.cpp/pull/210), [#226](https://github.com/0xShug0/audio.cpp/pull/226), [#239](https://github.com/0xShug0/audio.cpp/pull/239), [#247](https://github.com/0xShug0/audio.cpp/pull/247), [#259](https://github.com/0xShug0/audio.cpp/pull/259), [#299](https://github.com/0xShug0/audio.cpp/pull/299), [#305](https://github.com/0xShug0/audio.cpp/pull/305), [#393](https://github.com/0xShug0/audio.cpp/pull/393), and [#431](https://github.com/0xShug0/audio.cpp/pull/431). +- [@CaptainArni](https://github.com/CaptainArni) for fixing PocketTTS empty output when switching cached voices, keeping the Windows CUDA build path healthy, adding ACE-Step 1.5 XL DiT variants, allowing finished GGUFs to be re-converted, and fixing Windows CUDA setup when `PATH` contains stray quotes in [#22](https://github.com/0xShug0/audio.cpp/pull/22), [#93](https://github.com/0xShug0/audio.cpp/pull/93), [#235](https://github.com/0xShug0/audio.cpp/pull/235), [#464](https://github.com/0xShug0/audio.cpp/pull/464), and [#511](https://github.com/0xShug0/audio.cpp/pull/511). +- [@IIIIIllllIIIIIlllll](https://github.com/IIIIIllllIIIIIlllll) for the experimental ROCm/HIP backend, Linux HIP build path, Windows HIP/ROCm distribution preparation, HIP build documentation, VibeVoice HIP enablement, backend device listing, Vulkan AMD integer-dot guard, IndexTTS-2.5 support, base64 voice references in server requests, IndexTTS2 speech-rate controls, IndexTTS2 HIP/text-normalization fixes, BreezeTTS 2 performance and VRAM improvements, ModelScope downloads, HIP build guards, and the Fish Audio HIP Fast-AR CMake target in [#48](https://github.com/0xShug0/audio.cpp/pull/48), [#148](https://github.com/0xShug0/audio.cpp/pull/148), [#153](https://github.com/0xShug0/audio.cpp/pull/153), [#159](https://github.com/0xShug0/audio.cpp/pull/159), [#164](https://github.com/0xShug0/audio.cpp/pull/164), [#168](https://github.com/0xShug0/audio.cpp/pull/168), [#171](https://github.com/0xShug0/audio.cpp/pull/171), [#193](https://github.com/0xShug0/audio.cpp/pull/193), [#210](https://github.com/0xShug0/audio.cpp/pull/210), [#226](https://github.com/0xShug0/audio.cpp/pull/226), [#239](https://github.com/0xShug0/audio.cpp/pull/239), [#247](https://github.com/0xShug0/audio.cpp/pull/247), [#259](https://github.com/0xShug0/audio.cpp/pull/259), [#299](https://github.com/0xShug0/audio.cpp/pull/299), [#305](https://github.com/0xShug0/audio.cpp/pull/305), [#393](https://github.com/0xShug0/audio.cpp/pull/393), [#431](https://github.com/0xShug0/audio.cpp/pull/431), [#450](https://github.com/0xShug0/audio.cpp/pull/450), [#480](https://github.com/0xShug0/audio.cpp/pull/480), and [#490](https://github.com/0xShug0/audio.cpp/pull/490). - [@francescobozzo](https://github.com/francescobozzo) for Nix ROCm/HIP backend support, selectable model targets, and Nix CI/package fixes in [#162](https://github.com/0xShug0/audio.cpp/pull/162), [#163](https://github.com/0xShug0/audio.cpp/pull/163), and [#172](https://github.com/0xShug0/audio.cpp/pull/172). - [@patrickvonplaten](https://github.com/patrickvonplaten) for Metal backend fixes, Voxtral Realtime streaming speedups, live audio streaming from stdin, and incremental transcript deltas in [#102](https://github.com/0xShug0/audio.cpp/pull/102), [#116](https://github.com/0xShug0/audio.cpp/pull/116), [#118](https://github.com/0xShug0/audio.cpp/pull/118), and [#127](https://github.com/0xShug0/audio.cpp/pull/127). - [@dleiferives](https://github.com/dleiferives) for Parakeet-TDT 0.6B v3 ASR support and follow-up standalone GGUF validation/docs in [#111](https://github.com/0xShug0/audio.cpp/pull/111) and [#139](https://github.com/0xShug0/audio.cpp/pull/139). @@ -103,7 +115,7 @@ audio.cpp is moving faster because people keep showing up with real fixes, caref - [@liuzl](https://github.com/liuzl) for speeding up Metal `conv_transpose_1d` dispatch in [#149](https://github.com/0xShug0/audio.cpp/pull/149). - [@JayDataEngineer](https://github.com/JayDataEngineer) for fixing PocketTTS `clone_audio_path` option typing in [#147](https://github.com/0xShug0/audio.cpp/pull/147). - [@jasonchen31](https://github.com/jasonchen31) for adding server-side voice library folder support for name-based voice cloning, helping land SenseVoice-Small, fixing crashes on binaries built for older GPUs, adding Audio8 TTS, and fixing VoxCPM1 WebUI download/Yue language handling in [#191](https://github.com/0xShug0/audio.cpp/pull/191), [#219](https://github.com/0xShug0/audio.cpp/pull/219), [#240](https://github.com/0xShug0/audio.cpp/pull/240), [#333](https://github.com/0xShug0/audio.cpp/pull/333), and [#424](https://github.com/0xShug0/audio.cpp/pull/424). -- [@mirek190](https://github.com/mirek190) for the native WebUI follow-up work around model management, package handling, request controls, package labels, new GGUF package surfacing, CUDA BF16 cuBLAS output on Ampere, MiniMax-H3 WebUI polish, retiring the legacy Python WebUI, adding the reusable native model package manager, clearing stale WebUI service workers, clarifying ACE-Step GGUF package labels, documenting server host/port options, and fixing Qwen compact-logits graph reuse in [#199](https://github.com/0xShug0/audio.cpp/pull/199), [#206](https://github.com/0xShug0/audio.cpp/pull/206), [#207](https://github.com/0xShug0/audio.cpp/pull/207), [#208](https://github.com/0xShug0/audio.cpp/pull/208), [#211](https://github.com/0xShug0/audio.cpp/pull/211), [#213](https://github.com/0xShug0/audio.cpp/pull/213), [#229](https://github.com/0xShug0/audio.cpp/pull/229), [#230](https://github.com/0xShug0/audio.cpp/pull/230), [#257](https://github.com/0xShug0/audio.cpp/pull/257), [#258](https://github.com/0xShug0/audio.cpp/pull/258), [#381](https://github.com/0xShug0/audio.cpp/pull/381), and [#426](https://github.com/0xShug0/audio.cpp/pull/426). +- [@mirek190](https://github.com/mirek190) for the native WebUI follow-up work around model management, package handling, request controls, package labels, new GGUF package surfacing, CUDA BF16 cuBLAS output on Ampere, MiniMax-H3 WebUI polish, retiring the legacy Python WebUI, adding the reusable native model package manager, clearing stale WebUI service workers, clarifying ACE-Step GGUF package labels, documenting server host/port options, fixing Qwen compact-logits graph reuse, accelerating MiraTTS on CPU/Vulkan, adding Kokoro multilingual GGUF synthesis and CPU optimization, and sharing the eSpeak-ng phonemizer across model frontends in [#199](https://github.com/0xShug0/audio.cpp/pull/199), [#206](https://github.com/0xShug0/audio.cpp/pull/206), [#207](https://github.com/0xShug0/audio.cpp/pull/207), [#208](https://github.com/0xShug0/audio.cpp/pull/208), [#211](https://github.com/0xShug0/audio.cpp/pull/211), [#213](https://github.com/0xShug0/audio.cpp/pull/213), [#229](https://github.com/0xShug0/audio.cpp/pull/229), [#230](https://github.com/0xShug0/audio.cpp/pull/230), [#257](https://github.com/0xShug0/audio.cpp/pull/257), [#258](https://github.com/0xShug0/audio.cpp/pull/258), [#381](https://github.com/0xShug0/audio.cpp/pull/381), [#426](https://github.com/0xShug0/audio.cpp/pull/426), [#491](https://github.com/0xShug0/audio.cpp/pull/491), [#496](https://github.com/0xShug0/audio.cpp/pull/496), and [#502](https://github.com/0xShug0/audio.cpp/pull/502). - [@nikich340](https://github.com/nikich340) for adding explicit server model unloading support in [#197](https://github.com/0xShug0/audio.cpp/pull/197). - [@utsl42](https://github.com/utsl42) for fixing CUDA linking on NixOS in [#214](https://github.com/0xShug0/audio.cpp/pull/214). - [@yegorius](https://github.com/yegorius) for fixing PocketTTS handling in the model manager in [#205](https://github.com/0xShug0/audio.cpp/pull/205). @@ -114,22 +126,31 @@ audio.cpp is moving faster because people keep showing up with real fixes, caref - [@Blakeolson21](https://github.com/Blakeolson21) for documenting macOS CPU-only builds and fixing HIP backend classification in tests in [#237](https://github.com/0xShug0/audio.cpp/pull/237) and [#238](https://github.com/0xShug0/audio.cpp/pull/238). - [@odest](https://github.com/odest) for adding Vulkan backend support to Windows builds in [#242](https://github.com/0xShug0/audio.cpp/pull/242). - [@JoeMattie](https://github.com/JoeMattie) for bringing up MiniMax Music3, fixing BF16 convolution issues, repairing conversion, wiring native WebUI support, and honoring required request options in Studio in [#243](https://github.com/0xShug0/audio.cpp/pull/243) and [#245](https://github.com/0xShug0/audio.cpp/pull/245). -- [@Orion-zhen](https://github.com/Orion-zhen) for adding `--voice-dir` parsing to the server CLI in [#246](https://github.com/0xShug0/audio.cpp/pull/246). +- [@Orion-zhen](https://github.com/Orion-zhen) for adding `--voice-dir` parsing to the server CLI and fixing a HIP build error in [#246](https://github.com/0xShug0/audio.cpp/pull/246) and [#486](https://github.com/0xShug0/audio.cpp/pull/486). - [@derekja](https://github.com/derekja) for multipart ASR recognition prompts, clearer large-request errors, ASR graph allocation fixes, and CUDA graph-cache cleanup for concurrent ASR pressure in [#262](https://github.com/0xShug0/audio.cpp/pull/262), [#264](https://github.com/0xShug0/audio.cpp/pull/264), [#276](https://github.com/0xShug0/audio.cpp/pull/276), and [#293](https://github.com/0xShug0/audio.cpp/pull/293). - [@kawshikbuet17](https://github.com/kawshikbuet17) for adding OmniVoice weight-type benchmark coverage, validation reports, and Python-to-C++ tests in [#269](https://github.com/0xShug0/audio.cpp/pull/269). - [@tareko](https://github.com/tareko) for adding F5-TTS community scaffolding with Habibi Arabic aliases in [#275](https://github.com/0xShug0/audio.cpp/pull/275). - [@jrohde](https://github.com/jrohde) for adding the MOSS-VoiceGenerator community model in [#278](https://github.com/0xShug0/audio.cpp/pull/278). -- [@LysanderdeJong](https://github.com/LysanderdeJong) for adding the MMS-300M-1130 forced aligner community model in [#279](https://github.com/0xShug0/audio.cpp/pull/279). -- [@drzsdrtfg](https://github.com/drzsdrtfg) for the tag-driven prebuilt release pipeline, Supertonic voice-preset request-option fix, Soprano TTS community model, and Qwen cached-graph prefill input fix in [#286](https://github.com/0xShug0/audio.cpp/pull/286), [#302](https://github.com/0xShug0/audio.cpp/pull/302), [#323](https://github.com/0xShug0/audio.cpp/pull/323), and [#331](https://github.com/0xShug0/audio.cpp/pull/331). +- [@LysanderdeJong](https://github.com/LysanderdeJong) for adding the MMS-300M-1130 forced aligner community model, Sortformer v2.1 streaming diarization, the Orukeet r3 Parakeet-TDT weight variant, and VibeVoice ASR Streaming 7B performance work in [#279](https://github.com/0xShug0/audio.cpp/pull/279), [#493](https://github.com/0xShug0/audio.cpp/pull/493), [#523](https://github.com/0xShug0/audio.cpp/pull/523), and [#526](https://github.com/0xShug0/audio.cpp/pull/526). +- [@drzsdrtfg](https://github.com/drzsdrtfg) for the tag-driven prebuilt release pipeline, Supertonic voice-preset request-option fix, Soprano TTS community model, Soprano maintainer metadata correction, and Qwen cached-graph prefill input fix in [#286](https://github.com/0xShug0/audio.cpp/pull/286), [#302](https://github.com/0xShug0/audio.cpp/pull/302), [#323](https://github.com/0xShug0/audio.cpp/pull/323), [#331](https://github.com/0xShug0/audio.cpp/pull/331), and [#451](https://github.com/0xShug0/audio.cpp/pull/451). - [@Hi5808](https://github.com/Hi5808) for documenting Jetson Orin bring-up and correcting native architecture wording in [#288](https://github.com/0xShug0/audio.cpp/pull/288). - [@SelfRef](https://github.com/SelfRef) for making embedded WebUI work behind a path-prefix reverse proxy and adding server model LRU controls in [#297](https://github.com/0xShug0/audio.cpp/pull/297) and [#298](https://github.com/0xShug0/audio.cpp/pull/298). - [@bjhengen](https://github.com/bjhengen) for dropping out-of-span chunk speech metadata instead of aborting the whole run in [#301](https://github.com/0xShug0/audio.cpp/pull/301). -- [@gqf2008](https://github.com/gqf2008) for adding server idle unload and pre-load memory guard behavior, tightening indeterminate-footprint handling, adding Audio8 ASR, and improving default CJK text chunking in [#306](https://github.com/0xShug0/audio.cpp/pull/306), [#308](https://github.com/0xShug0/audio.cpp/pull/308), [#337](https://github.com/0xShug0/audio.cpp/pull/337), and [#441](https://github.com/0xShug0/audio.cpp/pull/441). +- [@gqf2008](https://github.com/gqf2008) for adding server idle unload and pre-load memory guard behavior, tightening indeterminate-footprint handling, adding Audio8 ASR, improving default CJK text chunking, and bringing the unit test suite into Linux/macOS/Windows CI in [#306](https://github.com/0xShug0/audio.cpp/pull/306), [#308](https://github.com/0xShug0/audio.cpp/pull/308), [#337](https://github.com/0xShug0/audio.cpp/pull/337), [#441](https://github.com/0xShug0/audio.cpp/pull/441), and [#442](https://github.com/0xShug0/audio.cpp/pull/442). - [@ampersandru](https://github.com/ampersandru) for adding IBM Granite Speech 5.0 470M TurboCTC ASR as a community model in [#311](https://github.com/0xShug0/audio.cpp/pull/311). - [@iamwavecut](https://github.com/iamwavecut) for MiniMax Music3 performance work around native RoPE/SwiGLU lowering, opt-in CFG reuse, chunk hop support, batched ensemble takes, and Q4_K depth decoder support in [#321](https://github.com/0xShug0/audio.cpp/pull/321). - [@XythQ](https://github.com/XythQ) for resolving model contracts once per loaded model instead of per request in [#328](https://github.com/0xShug0/audio.cpp/pull/328). - [@pannagaps](https://github.com/pannagaps) for adding Chatterbox Turbo TTS as a community model in [#394](https://github.com/0xShug0/audio.cpp/pull/394). - [@gsaon](https://github.com/gsaon) for fixing Granite Speech ASR mel filterbank construction on the continuous frequency axis in [#384](https://github.com/0xShug0/audio.cpp/pull/384). -- [@reezex0-ux](https://github.com/reezex0-ux) for adding the HIP Fast-AR top-k sampler path for Fish Audio in [#386](https://github.com/0xShug0/audio.cpp/pull/386). +- [@reezex0-ux](https://github.com/reezex0-ux) for adding the HIP Fast-AR top-k sampler path for Fish Audio and keeping Fast-AR inputs on the HIP stream in [#386](https://github.com/0xShug0/audio.cpp/pull/386) and [#500](https://github.com/0xShug0/audio.cpp/pull/500). - [@feng19](https://github.com/feng19) for adding `HF_ENDPOINT` mirror support to model downloads in [#397](https://github.com/0xShug0/audio.cpp/pull/397). - [@CryptVenture](https://github.com/CryptVenture) for adding richer WAV output options, correcting the shared text chunk mode preset, accepting the MOSS VoiceGenerator instruction spelling used by the speech route, and forwarding `language` only when a model contract accepts it in [#358](https://github.com/0xShug0/audio.cpp/pull/358), [#370](https://github.com/0xShug0/audio.cpp/pull/370), [#371](https://github.com/0xShug0/audio.cpp/pull/371), and [#400](https://github.com/0xShug0/audio.cpp/pull/400). +- [@voidash](https://github.com/voidash) for adding the sanoTTS voice family and expanding it with eleven more language voices plus acoustic calibration adapter support in [#449](https://github.com/0xShug0/audio.cpp/pull/449) and [#495](https://github.com/0xShug0/audio.cpp/pull/495). +- [@XsquirrelC](https://github.com/XsquirrelC) for adding additive INT8/ternary GGML types, fused CPU ops, and the VibeASR INT8/ternary ASR pipeline in [#447](https://github.com/0xShug0/audio.cpp/pull/447) and [#448](https://github.com/0xShug0/audio.cpp/pull/448). +- [@niksedk](https://github.com/niksedk) for fixing Higgs Audio TTS codec decode seams by giving the codec decoder right-hand context at the end of the stream in [#454](https://github.com/0xShug0/audio.cpp/pull/454). +- [@Th-Underscore](https://github.com/Th-Underscore) for adding the `Ctrl` variant to the BreezeTTS model entry in [#463](https://github.com/0xShug0/audio.cpp/pull/463). +- [@DrewThomasson](https://github.com/DrewThomasson) for adding the one-click Colab WebUI notebook and CUDA T4 release build path in [#501](https://github.com/0xShug0/audio.cpp/pull/501). +- [@Anai-Guo](https://github.com/Anai-Guo) for fixing the Voxtral Realtime warmbench path to use the warmup case streaming flag in [#506](https://github.com/0xShug0/audio.cpp/pull/506). +- [@dpmm99](https://github.com/dpmm99) for fixing the configured-model Yue2 UI parameter surface in [#509](https://github.com/0xShug0/audio.cpp/pull/509). +- [@weidu-tpvision](https://github.com/weidu-tpvision) for preventing 32-bit `size_t` overflow in HiFT vocoder and model configuration allocation paths in [#512](https://github.com/0xShug0/audio.cpp/pull/512) and [#513](https://github.com/0xShug0/audio.cpp/pull/513). +- [@SmilingWolf](https://github.com/SmilingWolf) for making PocketTTS reference audio fields visible again in [#528](https://github.com/0xShug0/audio.cpp/pull/528). diff --git a/README.md b/README.md index 5e9438965..28ac6a274 100644 --- a/README.md +++ b/README.md @@ -12,9 +12,7 @@ ModelScope repo mirror: https://www.modelscope.cn/models/HereIsMark/audio.cpp-gg > [!IMPORTANT] > -> **2026-09-10 - Dev testing: Yue2 3B:** Yue2 native song generation is available on the [dev branch](https://github.com/0xShug0/audio.cpp/tree/dev) for community testing and optimization. -> -> **2026-09-09 - VibeVoice ASR Streaming 7B and Irodori-TTS v4.1 Anime:** New GGUF packages are available for streaming VibeVoice ASR 7B and the anime fine-tuned Irodori-TTS v4.1 Small variant. +> **2026-09-10 - Dev testing: Yue2 3B:** Yue2 native song generation is available on the [dev branch](https://github.com/0xShug0/audio.cpp/tree/dev) for community testing and optimization. Dev binaries are available from [Actions -> Release -> latest dev artifacts](https://github.com/0xShug0/audio.cpp/actions/workflows/release.yml?query=branch%3Adev). > > **Arena UI:** The new Arena tab makes it easier to compare local models side by side for TTS, voice conversion, and ASR. Use one shared input, queue multiple models or GGUF variants, then review outputs with metrics! > @@ -53,7 +51,7 @@ audio.cpp would not be moving this quickly without generous contributors bringin ## News > [!IMPORTANT] -> **2026-09-04 - Release 0.7.2:** This release adds BreezeTTS 2, CosyVoice3, Chatterbox Turbo TTS, Audio8 TTS, and Audio8 ASR, plus the new multipart audio alignment endpoint. +> **2026-09-12 - Release 0.7.4:** This release adds VibeVoice ASR Streaming 7B, Irodori-TTS v4.1 Anime, Moonshine Streaming ASR, and Kokoro 82M, plus GGUF/package and UI updates for the latest community models. Thanks [@DrewThomasson](https://github.com/DrewThomasson) for the Colab UI! > > **2026-08-26 - Release 0.7:** This release adds MiniMax Music 3, MagpieTTS, PersonaPlex, MeanVC2, AudioSR, ControlFoley, FireRedTTS3, FireRedAudio, MiDashengLM-Gen, F5-TTS/Habibi, Granite Speech 5.0 TurboCTC, MMS Forced Aligner, and MOSS-VoiceGenerator, plus DotTTS Edit and ACE-Step 1.5 XL variants, bringing audio.cpp to **62** total model families and **85+** model variants! It also introduces the new Arena UI for side-by-side TTS, voice-conversion, and ASR comparison with shared inputs, queued runs, metrics, and result sorting. > From 78458ef0cf04f2e7635c51fd502d0918f9837145 Mon Sep 17 00:00:00 2001 From: 0xShug0 <231717474+0xShug0@users.noreply.github.com> Date: Sat, 12 Sep 2026 17:27:42 -0400 Subject: [PATCH 2/3] Add explicit ggml lowering APIs --- external/ggml/include/ggml.h | 117 +++++++++++++++++++- external/ggml/src/ggml.c | 206 ++++++++++++++++++++++++++++++++++- 2 files changed, 316 insertions(+), 7 deletions(-) diff --git a/external/ggml/include/ggml.h b/external/ggml/include/ggml.h index 12ad92b45..275e577be 100644 --- a/external/ggml/include/ggml.h +++ b/external/ggml/include/ggml.h @@ -448,8 +448,14 @@ extern "C" { // op hint enum ggml_op_hint { - GGML_HINT_NONE = 0, - GGML_HINT_SRC0_IS_HADAMARD = 1, + GGML_HINT_NONE = 0, + GGML_HINT_SRC0_IS_HADAMARD = 1, + }; + + enum ggml_mul_mat_lowering { + GGML_MUL_MAT_LOWERING_DEFAULT = 0, + GGML_MUL_MAT_LOWERING_CUDA_NVFP4_F16_ACTIVATION = 2, + GGML_MUL_MAT_LOWERING_CUDA_TILE_F16_ACCUM_OUTPUT = 3, }; // model file types @@ -604,6 +610,11 @@ extern "C" { GGML_OP_MUL_MAT_ADD, GGML_OP_MUL_MAT_ADD_RELU, GGML_OP_IM2COL_ASYM, + GGML_OP_CONV_3D_CONCAT_PAD_SPATIAL_GEMM, + GGML_OP_RMS_NORM_CHANNELS, + GGML_OP_RMS_NORM_CHANNELS_SILU, + GGML_OP_RMS_NORM_CHANNELS_ADD_BIAS_SILU, + GGML_OP_ROPE_INTERLEAVED_PAIRS, GGML_OP_COUNT, }; @@ -678,6 +689,33 @@ extern "C" { GGML_TRI_TYPE_LOWER = 3 }; + enum ggml_concat_lowering { + GGML_CONCAT_LOWERING_DEFAULT = 0, + GGML_CONCAT_LOWERING_CUDA_CONTIGUOUS_4D = 1, + }; + + enum ggml_im2col_2d_lowering { + GGML_IM2COL_2D_LOWERING_DEFAULT = 0, + GGML_IM2COL_2D_LOWERING_CUDA_N_K3_PAD1_X8 = 1, + GGML_IM2COL_2D_LOWERING_CUDA_N_K3_NOPAD_X8 = 2, + }; + + enum ggml_im2col_3d_lowering { + GGML_IM2COL_3D_LOWERING_DEFAULT = 0, + GGML_IM2COL_3D_LOWERING_CUDA_N1_K3_NOPAD_X8 = 1, + }; + + enum ggml_rms_norm_channels_lowering { + GGML_RMS_NORM_CHANNELS_LOWERING_DEFAULT = 0, + GGML_RMS_NORM_CHANNELS_LOWERING_CUDA_COALESCED = 1, + }; + + enum ggml_conv_3d_concat_pad_spatial_gemm_lowering { + GGML_CONV_3D_CONCAT_PAD_SPATIAL_GEMM_LOWERING_DEFAULT = 0, + GGML_CONV_3D_CONCAT_PAD_SPATIAL_GEMM_LOWERING_CUDA_C48 = 1, + GGML_CONV_3D_CONCAT_PAD_SPATIAL_GEMM_LOWERING_CUDA_TILED_C48 = 2, + }; + struct ggml_init_params { // memory pool size_t mem_size; // bytes @@ -1110,6 +1148,46 @@ extern "C" { struct ggml_tensor * b, int dim); + GGML_API struct ggml_tensor * ggml_rope_interleaved_pairs( + struct ggml_context * ctx, + struct ggml_tensor * even, + struct ggml_tensor * odd, + struct ggml_tensor * cos, + struct ggml_tensor * sin); + + GGML_API void ggml_concat_set_lowering( + struct ggml_tensor * tensor, + enum ggml_concat_lowering lowering); + + GGML_API struct ggml_tensor * ggml_conv_3d_concat_pad_spatial_gemm( + struct ggml_context * ctx, + struct ggml_tensor * a, + struct ggml_tensor * b, + struct ggml_tensor * w, + int lp0, + int rp0, + int lp1, + int rp1, + int lp2, + int rp2); + + GGML_API struct ggml_tensor * ggml_conv_3d_concat_pad_spatial_gemm_ex( + struct ggml_context * ctx, + struct ggml_tensor * a, + struct ggml_tensor * b, + struct ggml_tensor * w, + int lp0, + int rp0, + int lp1, + int rp1, + int lp2, + int rp2, + enum ggml_type dst_type); + + GGML_API void ggml_conv_3d_concat_pad_spatial_gemm_set_lowering( + struct ggml_tensor * tensor, + enum ggml_conv_3d_concat_pad_spatial_gemm_lowering lowering); + GGML_API struct ggml_tensor * ggml_abs( struct ggml_context * ctx, struct ggml_tensor * a); @@ -1406,6 +1484,29 @@ extern "C" { struct ggml_tensor * a, float eps); + GGML_API struct ggml_tensor * ggml_rms_norm_channels( + struct ggml_context * ctx, + struct ggml_tensor * a, + struct ggml_tensor * gamma, + float eps); + + GGML_API struct ggml_tensor * ggml_rms_norm_channels_silu( + struct ggml_context * ctx, + struct ggml_tensor * a, + struct ggml_tensor * gamma, + float eps); + + GGML_API struct ggml_tensor * ggml_rms_norm_channels_add_bias_silu( + struct ggml_context * ctx, + struct ggml_tensor * a, + struct ggml_tensor * bias, + struct ggml_tensor * gamma, + float eps); + + GGML_API void ggml_rms_norm_channels_set_lowering( + struct ggml_tensor * tensor, + enum ggml_rms_norm_channels_lowering lowering); + // group normalize along ne0*ne1*n_groups // used in stable-diffusion GGML_API struct ggml_tensor * ggml_group_norm( @@ -1464,6 +1565,10 @@ extern "C" { struct ggml_tensor * a, enum ggml_op_hint hint); + GGML_API void ggml_mul_mat_set_lowering( + struct ggml_tensor * a, + enum ggml_mul_mat_lowering lowering); + // indirect matrix multiplication GGML_API struct ggml_tensor * ggml_mul_mat_id( struct ggml_context * ctx, @@ -2028,6 +2133,10 @@ extern "C" { bool is_2D, enum ggml_type dst_type); + GGML_API void ggml_im2col_2d_set_lowering( + struct ggml_tensor * tensor, + enum ggml_im2col_2d_lowering lowering); + GGML_API struct ggml_tensor * ggml_im2col_back( struct ggml_context * ctx, struct ggml_tensor * a, // convolution kernel @@ -2125,6 +2234,10 @@ extern "C" { int d2, // dilation depth enum ggml_type dst_type); + GGML_API void ggml_im2col_3d_set_lowering( + struct ggml_tensor * tensor, + enum ggml_im2col_3d_lowering lowering); + // a: [OC*IC, KD, KH, KW] // b: [N*IC, ID, IH, IW] // result: [N*OC, OD, OH, OW] diff --git a/external/ggml/src/ggml.c b/external/ggml/src/ggml.c index 39319bfc7..431d7f987 100644 --- a/external/ggml/src/ggml.c +++ b/external/ggml/src/ggml.c @@ -1131,9 +1131,14 @@ static const char * GGML_OP_NAME[GGML_OP_COUNT] = { "MUL_MAT_ADD", "MUL_MAT_ADD_RELU", "IM2COL_ASYM", + "CONV_3D_CONCAT_PAD_SPATIAL_GEMM", + "RMS_NORM_CHANNELS", + "RMS_NORM_CHANNELS_SILU", + "RMS_NORM_CHANNELS_ADD_BIAS_SILU", + "ROPE_INTERLEAVED_PAIRS", }; -static_assert(GGML_OP_COUNT == 107, "GGML_OP_COUNT != 107"); +static_assert(GGML_OP_COUNT == 112, "GGML_OP_COUNT != 112"); static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "none", @@ -1253,9 +1258,14 @@ static const char * GGML_OP_SYMBOL[GGML_OP_COUNT] = { "a*b+bias", "relu(a*b+bias)", "im2col_asym(x)", + "conv_3d_concat_pad_spatial_gemm(x)", + "rms_norm_channels(x)", + "rms_norm_channels_silu(x)", + "rms_norm_channels_add_bias_silu(x)", + "rope_interleaved_pairs(even, odd, cos, sin)", }; -static_assert(GGML_OP_COUNT == 107, "GGML_OP_COUNT != 107"); +static_assert(GGML_OP_COUNT == 112, "GGML_OP_COUNT != 112"); static_assert(GGML_OP_POOL_COUNT == 2, "GGML_OP_POOL_COUNT != 2"); @@ -2120,9 +2130,7 @@ static struct ggml_tensor * ggml_add_cast_impl( struct ggml_tensor * a, struct ggml_tensor * b, enum ggml_type type) { - // TODO: support less-strict constraint - // GGML_ASSERT(ggml_can_repeat(b, a)); - GGML_ASSERT(ggml_can_repeat_rows(b, a)); + GGML_ASSERT(ggml_can_repeat(b, a)); // currently only supported for quantized input and f16 GGML_ASSERT(ggml_is_quantized(a->type) || @@ -2676,6 +2684,105 @@ struct ggml_tensor * ggml_concat( return result; } +struct ggml_tensor * ggml_rope_interleaved_pairs( + struct ggml_context * ctx, + struct ggml_tensor * even, + struct ggml_tensor * odd, + struct ggml_tensor * cos, + struct ggml_tensor * sin) { + GGML_ASSERT(even->type == odd->type); + GGML_ASSERT(cos->type == GGML_TYPE_F32 && sin->type == GGML_TYPE_F32); + GGML_ASSERT(ggml_nelements(even) == ggml_nelements(odd)); + GGML_ASSERT(ggml_nelements(even) == ggml_nelements(cos)); + GGML_ASSERT(ggml_nelements(even) == ggml_nelements(sin)); + + int64_t ne[GGML_MAX_DIMS] = { even->ne[0], even->ne[1], even->ne[2], even->ne[3] }; + ne[0] = 2; + + struct ggml_tensor * result = ggml_new_tensor(ctx, even->type, GGML_MAX_DIMS, ne); + + result->op = GGML_OP_ROPE_INTERLEAVED_PAIRS; + result->src[0] = even; + result->src[1] = odd; + result->src[2] = cos; + result->src[3] = sin; + + return result; +} + +void ggml_concat_set_lowering(struct ggml_tensor * tensor, enum ggml_concat_lowering lowering) { + GGML_ASSERT(tensor->op == GGML_OP_CONCAT); + ggml_set_op_params_i32(tensor, 1, (int32_t) lowering); +} + +struct ggml_tensor * ggml_conv_3d_concat_pad_spatial_gemm_ex( + struct ggml_context * ctx, + struct ggml_tensor * a, + struct ggml_tensor * b, + struct ggml_tensor * w, + int lp0, + int rp0, + int lp1, + int rp1, + int lp2, + int rp2, + enum ggml_type dst_type) { + GGML_ASSERT(a->type == GGML_TYPE_F32 || a->type == GGML_TYPE_F16); + GGML_ASSERT(b->type == GGML_TYPE_F32 || b->type == GGML_TYPE_F16); + GGML_ASSERT(w->type == GGML_TYPE_F32); + GGML_ASSERT(dst_type == GGML_TYPE_F16 || dst_type == GGML_TYPE_F32); + GGML_ASSERT(a->ne[0] == b->ne[0]); + GGML_ASSERT(a->ne[1] == b->ne[1]); + GGML_ASSERT(a->ne[3] == b->ne[3]); + GGML_ASSERT(w->ne[0] == 3 && w->ne[1] == 3 && w->ne[2] == 3); + GGML_ASSERT(lp0 >= 0 && rp0 >= 0 && lp1 >= 0 && rp1 >= 0 && lp2 >= 0 && rp2 >= 0); + + const int64_t ic = a->ne[3]; + GGML_ASSERT(w->ne[3] % ic == 0); + const int64_t oc = w->ne[3] / ic; + const int64_t ow = a->ne[0] + lp0 + rp0 - 2; + const int64_t oh = a->ne[1] + lp1 + rp1 - 2; + const int64_t od = a->ne[2] + b->ne[2] + lp2 + rp2 - 2; + GGML_ASSERT(ow > 0 && oh > 0 && od > 0); + + struct ggml_tensor * result = ggml_new_tensor_4d(ctx, dst_type, ow, oh, od, oc); + + ggml_set_op_params_i32(result, 0, lp0); + ggml_set_op_params_i32(result, 1, rp0); + ggml_set_op_params_i32(result, 2, lp1); + ggml_set_op_params_i32(result, 3, rp1); + ggml_set_op_params_i32(result, 4, lp2); + ggml_set_op_params_i32(result, 5, rp2); + + result->op = GGML_OP_CONV_3D_CONCAT_PAD_SPATIAL_GEMM; + result->src[0] = a; + result->src[1] = b; + result->src[2] = w; + + return result; +} + +void ggml_conv_3d_concat_pad_spatial_gemm_set_lowering( + struct ggml_tensor * tensor, + enum ggml_conv_3d_concat_pad_spatial_gemm_lowering lowering) { + GGML_ASSERT(tensor->op == GGML_OP_CONV_3D_CONCAT_PAD_SPATIAL_GEMM); + ggml_set_op_params_i32(tensor, 6, (int32_t) lowering); +} + +struct ggml_tensor * ggml_conv_3d_concat_pad_spatial_gemm( + struct ggml_context * ctx, + struct ggml_tensor * a, + struct ggml_tensor * b, + struct ggml_tensor * w, + int lp0, + int rp0, + int lp1, + int rp1, + int lp2, + int rp2) { + return ggml_conv_3d_concat_pad_spatial_gemm_ex(ctx, a, b, w, lp0, rp0, lp1, rp1, lp2, rp2, GGML_TYPE_F32); +} + // ggml_abs struct ggml_tensor * ggml_abs( @@ -3230,6 +3337,70 @@ struct ggml_tensor * ggml_rms_norm_inplace( return ggml_rms_norm_impl(ctx, a, eps, true); } +struct ggml_tensor * ggml_rms_norm_channels( + struct ggml_context * ctx, + struct ggml_tensor * a, + struct ggml_tensor * gamma, + float eps) { + GGML_ASSERT(gamma->ne[0] == a->ne[3]); + struct ggml_tensor * result = ggml_new_tensor(ctx, GGML_TYPE_F32, GGML_MAX_DIMS, a->ne); + + ggml_set_op_params(result, &eps, sizeof(eps)); + + result->op = GGML_OP_RMS_NORM_CHANNELS; + result->src[0] = a; + result->src[1] = gamma; + + return result; +} + +struct ggml_tensor * ggml_rms_norm_channels_silu( + struct ggml_context * ctx, + struct ggml_tensor * a, + struct ggml_tensor * gamma, + float eps) { + GGML_ASSERT(gamma->ne[0] == a->ne[3]); + struct ggml_tensor * result = ggml_new_tensor(ctx, GGML_TYPE_F32, GGML_MAX_DIMS, a->ne); + + ggml_set_op_params(result, &eps, sizeof(eps)); + + result->op = GGML_OP_RMS_NORM_CHANNELS_SILU; + result->src[0] = a; + result->src[1] = gamma; + + return result; +} + +struct ggml_tensor * ggml_rms_norm_channels_add_bias_silu( + struct ggml_context * ctx, + struct ggml_tensor * a, + struct ggml_tensor * bias, + struct ggml_tensor * gamma, + float eps) { + GGML_ASSERT(bias->ne[0] == a->ne[3]); + GGML_ASSERT(gamma->ne[0] == a->ne[3]); + struct ggml_tensor * result = ggml_new_tensor(ctx, GGML_TYPE_F32, GGML_MAX_DIMS, a->ne); + + ggml_set_op_params(result, &eps, sizeof(eps)); + + result->op = GGML_OP_RMS_NORM_CHANNELS_ADD_BIAS_SILU; + result->src[0] = a; + result->src[1] = bias; + result->src[2] = gamma; + + return result; +} + +void ggml_rms_norm_channels_set_lowering( + struct ggml_tensor * tensor, + enum ggml_rms_norm_channels_lowering lowering) { + GGML_ASSERT( + tensor->op == GGML_OP_RMS_NORM_CHANNELS || + tensor->op == GGML_OP_RMS_NORM_CHANNELS_SILU || + tensor->op == GGML_OP_RMS_NORM_CHANNELS_ADD_BIAS_SILU); + ggml_set_op_params_i32(tensor, 1, (int32_t) lowering); +} + // ggml_rms_norm_back struct ggml_tensor * ggml_rms_norm_back( @@ -3379,6 +3550,16 @@ void ggml_mul_mat_set_hint( ggml_set_op_params_i32(a, 1, hint_i32); } +void ggml_mul_mat_set_lowering( + struct ggml_tensor * a, + enum ggml_mul_mat_lowering lowering) { + GGML_ASSERT(a->op == GGML_OP_MUL_MAT || a->op == GGML_OP_MUL_MAT_PACK4); + + const int32_t lowering_i32 = (int32_t) lowering; + + ggml_set_op_params_i32(a, 1, lowering_i32); +} + // ggml_mul_mat_id /* @@ -4560,6 +4741,14 @@ struct ggml_tensor * ggml_im2col( return result; } +void ggml_im2col_2d_set_lowering( + struct ggml_tensor * tensor, + enum ggml_im2col_2d_lowering lowering) { + GGML_ASSERT(tensor->op == GGML_OP_IM2COL || tensor->op == GGML_OP_IM2COL_FAST_1D); + GGML_ASSERT(ggml_get_op_params_i32(tensor, 6) == 1); + ggml_set_op_params_i32(tensor, 7, (int32_t) lowering); +} + struct ggml_tensor * ggml_im2col_back( struct ggml_context * ctx, struct ggml_tensor * a, @@ -4827,6 +5016,13 @@ struct ggml_tensor * ggml_im2col_3d( return result; } +void ggml_im2col_3d_set_lowering( + struct ggml_tensor * tensor, + enum ggml_im2col_3d_lowering lowering) { + GGML_ASSERT(tensor->op == GGML_OP_IM2COL_3D); + ggml_set_op_params_i32(tensor, 10, (int32_t) lowering); +} + // a: [OC*IC, KD, KH, KW] // b: [N*IC, ID, IH, IW] // result: [N*OC, OD, OH, OW] From 01842a66e2ee5a448465220a86129e8887ef9856 Mon Sep 17 00:00:00 2001 From: 0xShug0 <231717474+0xShug0@users.noreply.github.com> Date: Sat, 12 Sep 2026 17:28:01 -0400 Subject: [PATCH 3/3] Add CUDA lowering implementations --- external/ggml/src/ggml-cuda/binbcast.cu | 1078 ++++++++++++---------- external/ggml/src/ggml-cuda/binbcast.cuh | 1 + external/ggml/src/ggml-cuda/concat.cu | 89 ++ external/ggml/src/ggml-cuda/conv3d.cu | 352 +++++++ external/ggml/src/ggml-cuda/conv3d.cuh | 5 + external/ggml/src/ggml-cuda/ggml-cuda.cu | 105 ++- external/ggml/src/ggml-cuda/im2col.cu | 320 ++++++- external/ggml/src/ggml-cuda/mmq.cu | 31 +- external/ggml/src/ggml-cuda/norm.cu | 240 +++++ external/ggml/src/ggml-cuda/norm.cuh | 6 + external/ggml/src/ggml-cuda/quantize.cu | 25 +- external/ggml/src/ggml-cuda/quantize.cuh | 5 + 12 files changed, 1709 insertions(+), 548 deletions(-) create mode 100644 external/ggml/src/ggml-cuda/conv3d.cu create mode 100644 external/ggml/src/ggml-cuda/conv3d.cuh diff --git a/external/ggml/src/ggml-cuda/binbcast.cu b/external/ggml/src/ggml-cuda/binbcast.cu index b23dc43e2..1a66e3579 100644 --- a/external/ggml/src/ggml-cuda/binbcast.cu +++ b/external/ggml/src/ggml-cuda/binbcast.cu @@ -1,153 +1,153 @@ -#include "binbcast.cuh" -#include -#include - -static __device__ __forceinline__ float op_repeat(const float a, const float b) { - return b; - GGML_UNUSED(a); -} - -static __device__ __forceinline__ float op_add(const float a, const float b) { - return a + b; -} - -static __device__ __forceinline__ float op_sub(const float a, const float b) { - return a - b; -} - -static __device__ __forceinline__ float op_mul(const float a, const float b) { - return a * b; -} - -static __device__ __forceinline__ float op_div(const float a, const float b) { - return a / b; -} - -template -static __global__ void k_bin_bcast(const src0_t * src0, - const src1_t * src1, - dst_t * dst, - const int ne0, - const int ne1, - const int ne2, - const uint3 ne3, - const uint3 ne10, - const uint3 ne11, - const uint3 ne12, - const uint3 ne13, - /*const int s0,*/ - const int s1, - const int s2, - const int s3, - const int s00, - const int s01, - const int s02, - const int s03, - const int s10, - const int s11, - const int s12, - const int s13, - src1_ptrs... src1s) { - const uint32_t i0s = blockDim.x * blockIdx.x + threadIdx.x; - const uint32_t i1 = (blockDim.y * blockIdx.y + threadIdx.y); - const uint32_t i2 = fastdiv((blockDim.z * blockIdx.z + threadIdx.z), ne3); - const uint32_t i3 = (blockDim.z * blockIdx.z + threadIdx.z) - (i2 * ne3.z); - - if (i0s >= (uint32_t)ne0 || i1 >= (uint32_t)ne1 || i2 >= (uint32_t)ne2 || i3 >= ne3.z) { - return; - } - - const uint32_t i11 = fastmodulo(i1, ne11); - const uint32_t i12 = fastmodulo(i2, ne12); - const uint32_t i13 = fastmodulo(i3, ne13); - - const size_t i_src0 = i3*s03 + i2*s02 + i1*s01; - const size_t i_src1 = i13*s13 + i12*s12 + i11*s11; - const size_t i_dst = i3*s3 + i2*s2 + i1*s1; - - const src0_t * src0_row = src0 ? (src0 + i_src0) : nullptr; - dst_t * dst_row = dst + i_dst; - - for (int i0 = i0s; i0 < ne0; i0 += blockDim.x * gridDim.x) { - const uint32_t i10 = fastmodulo(i0, ne10); - - float result = src0_row ? (float) src0_row[i0*s00] : 0.0f; - if constexpr (sizeof...(src1_ptrs) > 0) { - result = (..., (result = bin_op(result, (float)src1s[i_src1 + i10*s10]))); - } else { - result = bin_op(result, (float)src1[i_src1 + i10*s10]); - } - - dst_row[i0] = (dst_t) result; - } -} - -template +#include "binbcast.cuh" +#include +#include + +static __device__ __forceinline__ float op_repeat(const float a, const float b) { + return b; + GGML_UNUSED(a); +} + +static __device__ __forceinline__ float op_add(const float a, const float b) { + return a + b; +} + +static __device__ __forceinline__ float op_sub(const float a, const float b) { + return a - b; +} + +static __device__ __forceinline__ float op_mul(const float a, const float b) { + return a * b; +} + +static __device__ __forceinline__ float op_div(const float a, const float b) { + return a / b; +} + +template +static __global__ void k_bin_bcast(const src0_t * src0, + const src1_t * src1, + dst_t * dst, + const int ne0, + const int ne1, + const int ne2, + const uint3 ne3, + const uint3 ne10, + const uint3 ne11, + const uint3 ne12, + const uint3 ne13, + /*const int s0,*/ + const int s1, + const int s2, + const int s3, + const int s00, + const int s01, + const int s02, + const int s03, + const int s10, + const int s11, + const int s12, + const int s13, + src1_ptrs... src1s) { + const uint32_t i0s = blockDim.x * blockIdx.x + threadIdx.x; + const uint32_t i1 = (blockDim.y * blockIdx.y + threadIdx.y); + const uint32_t i2 = fastdiv((blockDim.z * blockIdx.z + threadIdx.z), ne3); + const uint32_t i3 = (blockDim.z * blockIdx.z + threadIdx.z) - (i2 * ne3.z); + + if (i0s >= (uint32_t)ne0 || i1 >= (uint32_t)ne1 || i2 >= (uint32_t)ne2 || i3 >= ne3.z) { + return; + } + + const uint32_t i11 = fastmodulo(i1, ne11); + const uint32_t i12 = fastmodulo(i2, ne12); + const uint32_t i13 = fastmodulo(i3, ne13); + + const size_t i_src0 = i3*s03 + i2*s02 + i1*s01; + const size_t i_src1 = i13*s13 + i12*s12 + i11*s11; + const size_t i_dst = i3*s3 + i2*s2 + i1*s1; + + const src0_t * src0_row = src0 ? (src0 + i_src0) : nullptr; + dst_t * dst_row = dst + i_dst; + + for (int i0 = i0s; i0 < ne0; i0 += blockDim.x * gridDim.x) { + const uint32_t i10 = fastmodulo(i0, ne10); + + float result = src0_row ? (float) src0_row[i0*s00] : 0.0f; + if constexpr (sizeof...(src1_ptrs) > 0) { + result = (..., (result = bin_op(result, (float)src1s[i_src1 + i10*s10]))); + } else { + result = bin_op(result, (float)src1[i_src1 + i10*s10]); + } + + dst_row[i0] = (dst_t) result; + } +} + +template static __global__ void k_bin_bcast_unravel(const src0_t * src0, - const src1_t * src1, - dst_t * dst, - const uint3 ne0, - const uint3 ne1, - const uint3 ne2, - const uint32_t ne3, - const uint3 prod_012, - const uint3 prod_01, - const uint3 ne10, - const uint3 ne11, - const uint3 ne12, - const uint3 ne13, - /*const int s0,*/ - const int s1, - const int s2, - const int s3, - const int s00, - const int s01, - const int s02, - const int s03, - const int s10, - const int s11, - const int s12, - const int s13, - src1_ptrs... src1s) { - const int i = blockDim.x*blockIdx.x + threadIdx.x; - - const uint32_t i3 = fastdiv(i, prod_012); - const uint32_t i2 = fastdiv(i - i3 * prod_012.z, prod_01); - const uint32_t i1 = fastdiv(i - i3 * prod_012.z - i2 * prod_01.z, ne0); - const uint32_t i0 = i - i3 * prod_012.z - i2 * prod_01.z - i1 * ne0.z; - - if (i0 >= ne0.z || i1 >= ne1.z || i2 >= ne2.z || i3 >= ne3) { - return; - } - - const int i11 = fastmodulo(i1, ne11); - const int i12 = fastmodulo(i2, ne12); - const int i13 = fastmodulo(i3, ne13); - - const size_t i_src0 = i3*s03 + i2*s02 + i1*s01; - const size_t i_src1 = i13*s13 + i12*s12 + i11*s11; - const size_t i_dst = i3*s3 + i2*s2 + i1*s1; - - const src0_t * src0_row = src0 ? (src0 + i_src0) : nullptr; - dst_t * dst_row = dst + i_dst; - - const int i10 = fastmodulo(i0, ne10); - - float result = src0_row ? (float) src0_row[i0*s00] : 0.0f; - if constexpr (sizeof...(src1_ptrs) > 0) { - result = (..., (result = bin_op(result, (float)src1s[i_src1 + i10*s10]))); - } else { - result = bin_op(result, (float)src1[i_src1 + i10*s10]); - } - + const src1_t * src1, + dst_t * dst, + const uint3 ne0, + const uint3 ne1, + const uint3 ne2, + const uint32_t ne3, + const uint3 prod_012, + const uint3 prod_01, + const uint3 ne10, + const uint3 ne11, + const uint3 ne12, + const uint3 ne13, + /*const int s0,*/ + const int s1, + const int s2, + const int s3, + const int s00, + const int s01, + const int s02, + const int s03, + const int s10, + const int s11, + const int s12, + const int s13, + src1_ptrs... src1s) { + const int i = blockDim.x*blockIdx.x + threadIdx.x; + + const uint32_t i3 = fastdiv(i, prod_012); + const uint32_t i2 = fastdiv(i - i3 * prod_012.z, prod_01); + const uint32_t i1 = fastdiv(i - i3 * prod_012.z - i2 * prod_01.z, ne0); + const uint32_t i0 = i - i3 * prod_012.z - i2 * prod_01.z - i1 * ne0.z; + + if (i0 >= ne0.z || i1 >= ne1.z || i2 >= ne2.z || i3 >= ne3) { + return; + } + + const int i11 = fastmodulo(i1, ne11); + const int i12 = fastmodulo(i2, ne12); + const int i13 = fastmodulo(i3, ne13); + + const size_t i_src0 = i3*s03 + i2*s02 + i1*s01; + const size_t i_src1 = i13*s13 + i12*s12 + i11*s11; + const size_t i_dst = i3*s3 + i2*s2 + i1*s1; + + const src0_t * src0_row = src0 ? (src0 + i_src0) : nullptr; + dst_t * dst_row = dst + i_dst; + + const int i10 = fastmodulo(i0, ne10); + + float result = src0_row ? (float) src0_row[i0*s00] : 0.0f; + if constexpr (sizeof...(src1_ptrs) > 0) { + result = (..., (result = bin_op(result, (float)src1s[i_src1 + i10*s10]))); + } else { + result = bin_op(result, (float)src1[i_src1 + i10*s10]); + } + dst_row[i0] = (dst_t) result; } @@ -204,166 +204,166 @@ static __global__ void k_bin_bcast_two_sided(const src0_t * src0, template static void launch_bin_bcast_pack(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, - const src0_t * src0_dd, const src1_t * src1_dd, dst_t * dst_dd, - cudaStream_t stream, std::index_sequence) { - GGML_TENSOR_BINARY_OP_LOCALS - - int nr0 = ne10 / ne0; - int nr1 = ne11 / ne1; - int nr2 = ne12 / ne2; - int nr3 = ne13 / ne3; - - int nr[4] = { nr0, nr1, nr2, nr3 }; - - int64_t cne[] = { ne0, ne1, ne2, ne3 }; - int64_t cne0[] = { ne00, ne01, ne02, ne03 }; - int64_t cne1[] = { ne10, ne11, ne12, ne13 }; - - size_t cnb[] = { nb0, nb1, nb2, nb3 }; - size_t cnb0[] = { nb00, nb01, nb02, nb03 }; - size_t cnb1[] = { nb10, nb11, nb12, nb13 }; - - auto collapse = [](int64_t cne[]) { - cne[0] *= cne[1]; - cne[1] = cne[2]; - cne[2] = cne[3]; - cne[3] = 1; - }; - - auto collapse_nb = [](size_t cnb[], const int64_t cne[]) { - cnb[1] *= cne[1]; - cnb[2] *= cne[2]; - cnb[3] *= cne[3]; - }; - - if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1) && !ggml_is_permuted(src0) && !ggml_is_permuted(src1)) { - for (int i = 0; i < 4; i++) { - if (nr[i] != 1) { - break; - } - if (i > 0) { - collapse_nb(cnb, cne); - collapse_nb(cnb0, cne0); - collapse_nb(cnb1, cne1); - collapse(cne); - collapse(cne0); - collapse(cne1); - } - } - } - - { - int64_t ne0 = cne[0]; - int64_t ne1 = cne[1]; - int64_t ne2 = cne[2]; - int64_t ne3 = cne[3]; - - //int64_t ne00 = cne0[0]; GGML_UNUSED(ne00); - //int64_t ne01 = cne0[1]; GGML_UNUSED(ne01); - //int64_t ne02 = cne0[2]; GGML_UNUSED(ne02); - //int64_t ne03 = cne0[3]; GGML_UNUSED(ne03); - - size_t nb0 = cnb[0]; - size_t nb1 = cnb[1]; - size_t nb2 = cnb[2]; - size_t nb3 = cnb[3]; - - size_t nb00 = cnb0[0]; - size_t nb01 = cnb0[1]; - size_t nb02 = cnb0[2]; - size_t nb03 = cnb0[3]; - - size_t nb10 = cnb1[0]; - size_t nb11 = cnb1[1]; - size_t nb12 = cnb1[2]; - size_t nb13 = cnb1[3]; - - //size_t s0 = nb0 / sizeof(dst_t); - size_t s1 = nb1 / sizeof(dst_t); - size_t s2 = nb2 / sizeof(dst_t); - size_t s3 = nb3 / sizeof(dst_t); - - size_t s10 = nb10 / sizeof(src1_t); - size_t s11 = nb11 / sizeof(src1_t); - size_t s12 = nb12 / sizeof(src1_t); - size_t s13 = nb13 / sizeof(src1_t); - - size_t s00 = nb00 / sizeof(src0_t); - size_t s01 = nb01 / sizeof(src0_t); - size_t s02 = nb02 / sizeof(src0_t); - size_t s03 = nb03 / sizeof(src0_t); - - GGML_ASSERT(nb0 % sizeof(dst_t) == 0); - GGML_ASSERT(nb1 % sizeof(dst_t) == 0); - GGML_ASSERT(nb2 % sizeof(dst_t) == 0); - GGML_ASSERT(nb3 % sizeof(dst_t) == 0); - - GGML_ASSERT(nb00 % sizeof(src0_t) == 0); - GGML_ASSERT(nb01 % sizeof(src0_t) == 0); - GGML_ASSERT(nb02 % sizeof(src0_t) == 0); - GGML_ASSERT(nb03 % sizeof(src0_t) == 0); - - GGML_ASSERT(nb10 % sizeof(src1_t) == 0); - GGML_ASSERT(nb11 % sizeof(src1_t) == 0); - GGML_ASSERT(nb12 % sizeof(src1_t) == 0); - GGML_ASSERT(nb13 % sizeof(src1_t) == 0); - - const int block_size = 128; - - int64_t hne0 = std::max(ne0 / 2LL, 1LL); - - dim3 block_dims; - block_dims.x = std::min(hne0, block_size); - block_dims.y = std::min(ne1, block_size / block_dims.x); - block_dims.z = std::min(std::min(ne2 * ne3, block_size / block_dims.x / block_dims.y), 64U); - - dim3 block_nums((hne0 + block_dims.x - 1) / block_dims.x, (ne1 + block_dims.y - 1) / block_dims.y, - (ne2 * ne3 + block_dims.z - 1) / block_dims.z); - - const uint3 ne10 = init_fastdiv_values((uint32_t) cne1[0]); - const uint3 ne11 = init_fastdiv_values((uint32_t) cne1[1]); - const uint3 ne12 = init_fastdiv_values((uint32_t) cne1[2]); - const uint3 ne13 = init_fastdiv_values((uint32_t) cne1[3]); - - if (block_nums.z > 65535 || block_nums.y > 65535) { - int block_num = (ne0 * ne1 * ne2 * ne3 + block_size - 1) / block_size; - const uint3 prod_012 = init_fastdiv_values((uint32_t) (ne0 * ne1 * ne2)); - const uint3 prod_01 = init_fastdiv_values((uint32_t) (ne0 * ne1)); - const uint3 ne0_fastdiv = init_fastdiv_values((uint32_t) ne0); - const uint3 ne1_fastdiv = init_fastdiv_values((uint32_t) ne1); - const uint3 ne2_fastdiv = init_fastdiv_values((uint32_t) ne2); - - if constexpr (sizeof...(I) > 0) { - k_bin_bcast_unravel<<>>( - src0_dd, src1_dd, dst_dd, ne0_fastdiv, ne1_fastdiv, ne2_fastdiv, ne3, prod_012, prod_01, ne10, ne11, - ne12, ne13, - /*s0,*/ s1, s2, s3, - s00, s01, s02, s03, - s10, s11, s12, s13, (const src1_t *) dst->src[I + 1]->data...); - } else { - k_bin_bcast_unravel - <<>>(src0_dd, src1_dd, dst_dd, ne0_fastdiv, ne1_fastdiv, - ne2_fastdiv, ne3, prod_012, prod_01, ne10, ne11, ne12, ne13, - /*s0,*/ s1, s2, s3, - s00, s01, s02, s03, - s10, s11, s12, s13); - } - } else { - const uint3 ne3_fastdiv = init_fastdiv_values((uint32_t) ne3); - if constexpr (sizeof...(I) > 0) { - k_bin_bcast<<>>( - src0_dd, src1_dd, dst_dd, ne0, ne1, ne2, ne3_fastdiv, ne10, ne11, ne12, ne13, - /*s0,*/ s1, s2, s3, - s00 ,s01, s02, s03, - s10, s11, s12, s13, (const src1_t *) dst->src[I + 1]->data...); - } else { - k_bin_bcast<<>>( - src0_dd, src1_dd, dst_dd, ne0, ne1, ne2, ne3_fastdiv, ne10, ne11, ne12, ne13, - /*s0,*/ s1, s2, s3, - s00, s01, s02, s03, - s10, s11, s12, s13); - } - } + const src0_t * src0_dd, const src1_t * src1_dd, dst_t * dst_dd, + cudaStream_t stream, std::index_sequence) { + GGML_TENSOR_BINARY_OP_LOCALS + + int nr0 = ne10 / ne0; + int nr1 = ne11 / ne1; + int nr2 = ne12 / ne2; + int nr3 = ne13 / ne3; + + int nr[4] = { nr0, nr1, nr2, nr3 }; + + int64_t cne[] = { ne0, ne1, ne2, ne3 }; + int64_t cne0[] = { ne00, ne01, ne02, ne03 }; + int64_t cne1[] = { ne10, ne11, ne12, ne13 }; + + size_t cnb[] = { nb0, nb1, nb2, nb3 }; + size_t cnb0[] = { nb00, nb01, nb02, nb03 }; + size_t cnb1[] = { nb10, nb11, nb12, nb13 }; + + auto collapse = [](int64_t cne[]) { + cne[0] *= cne[1]; + cne[1] = cne[2]; + cne[2] = cne[3]; + cne[3] = 1; + }; + + auto collapse_nb = [](size_t cnb[], const int64_t cne[]) { + cnb[1] *= cne[1]; + cnb[2] *= cne[2]; + cnb[3] *= cne[3]; + }; + + if (ggml_is_contiguous(src0) && ggml_is_contiguous(src1) && !ggml_is_permuted(src0) && !ggml_is_permuted(src1)) { + for (int i = 0; i < 4; i++) { + if (nr[i] != 1) { + break; + } + if (i > 0) { + collapse_nb(cnb, cne); + collapse_nb(cnb0, cne0); + collapse_nb(cnb1, cne1); + collapse(cne); + collapse(cne0); + collapse(cne1); + } + } + } + + { + int64_t ne0 = cne[0]; + int64_t ne1 = cne[1]; + int64_t ne2 = cne[2]; + int64_t ne3 = cne[3]; + + //int64_t ne00 = cne0[0]; GGML_UNUSED(ne00); + //int64_t ne01 = cne0[1]; GGML_UNUSED(ne01); + //int64_t ne02 = cne0[2]; GGML_UNUSED(ne02); + //int64_t ne03 = cne0[3]; GGML_UNUSED(ne03); + + size_t nb0 = cnb[0]; + size_t nb1 = cnb[1]; + size_t nb2 = cnb[2]; + size_t nb3 = cnb[3]; + + size_t nb00 = cnb0[0]; + size_t nb01 = cnb0[1]; + size_t nb02 = cnb0[2]; + size_t nb03 = cnb0[3]; + + size_t nb10 = cnb1[0]; + size_t nb11 = cnb1[1]; + size_t nb12 = cnb1[2]; + size_t nb13 = cnb1[3]; + + //size_t s0 = nb0 / sizeof(dst_t); + size_t s1 = nb1 / sizeof(dst_t); + size_t s2 = nb2 / sizeof(dst_t); + size_t s3 = nb3 / sizeof(dst_t); + + size_t s10 = nb10 / sizeof(src1_t); + size_t s11 = nb11 / sizeof(src1_t); + size_t s12 = nb12 / sizeof(src1_t); + size_t s13 = nb13 / sizeof(src1_t); + + size_t s00 = nb00 / sizeof(src0_t); + size_t s01 = nb01 / sizeof(src0_t); + size_t s02 = nb02 / sizeof(src0_t); + size_t s03 = nb03 / sizeof(src0_t); + + GGML_ASSERT(nb0 % sizeof(dst_t) == 0); + GGML_ASSERT(nb1 % sizeof(dst_t) == 0); + GGML_ASSERT(nb2 % sizeof(dst_t) == 0); + GGML_ASSERT(nb3 % sizeof(dst_t) == 0); + + GGML_ASSERT(nb00 % sizeof(src0_t) == 0); + GGML_ASSERT(nb01 % sizeof(src0_t) == 0); + GGML_ASSERT(nb02 % sizeof(src0_t) == 0); + GGML_ASSERT(nb03 % sizeof(src0_t) == 0); + + GGML_ASSERT(nb10 % sizeof(src1_t) == 0); + GGML_ASSERT(nb11 % sizeof(src1_t) == 0); + GGML_ASSERT(nb12 % sizeof(src1_t) == 0); + GGML_ASSERT(nb13 % sizeof(src1_t) == 0); + + const int block_size = 128; + + int64_t hne0 = std::max(ne0 / 2LL, 1LL); + + dim3 block_dims; + block_dims.x = std::min(hne0, block_size); + block_dims.y = std::min(ne1, block_size / block_dims.x); + block_dims.z = std::min(std::min(ne2 * ne3, block_size / block_dims.x / block_dims.y), 64U); + + dim3 block_nums((hne0 + block_dims.x - 1) / block_dims.x, (ne1 + block_dims.y - 1) / block_dims.y, + (ne2 * ne3 + block_dims.z - 1) / block_dims.z); + + const uint3 ne10 = init_fastdiv_values((uint32_t) cne1[0]); + const uint3 ne11 = init_fastdiv_values((uint32_t) cne1[1]); + const uint3 ne12 = init_fastdiv_values((uint32_t) cne1[2]); + const uint3 ne13 = init_fastdiv_values((uint32_t) cne1[3]); + + if (block_nums.z > 65535 || block_nums.y > 65535) { + int block_num = (ne0 * ne1 * ne2 * ne3 + block_size - 1) / block_size; + const uint3 prod_012 = init_fastdiv_values((uint32_t) (ne0 * ne1 * ne2)); + const uint3 prod_01 = init_fastdiv_values((uint32_t) (ne0 * ne1)); + const uint3 ne0_fastdiv = init_fastdiv_values((uint32_t) ne0); + const uint3 ne1_fastdiv = init_fastdiv_values((uint32_t) ne1); + const uint3 ne2_fastdiv = init_fastdiv_values((uint32_t) ne2); + + if constexpr (sizeof...(I) > 0) { + k_bin_bcast_unravel<<>>( + src0_dd, src1_dd, dst_dd, ne0_fastdiv, ne1_fastdiv, ne2_fastdiv, ne3, prod_012, prod_01, ne10, ne11, + ne12, ne13, + /*s0,*/ s1, s2, s3, + s00, s01, s02, s03, + s10, s11, s12, s13, (const src1_t *) dst->src[I + 1]->data...); + } else { + k_bin_bcast_unravel + <<>>(src0_dd, src1_dd, dst_dd, ne0_fastdiv, ne1_fastdiv, + ne2_fastdiv, ne3, prod_012, prod_01, ne10, ne11, ne12, ne13, + /*s0,*/ s1, s2, s3, + s00, s01, s02, s03, + s10, s11, s12, s13); + } + } else { + const uint3 ne3_fastdiv = init_fastdiv_values((uint32_t) ne3); + if constexpr (sizeof...(I) > 0) { + k_bin_bcast<<>>( + src0_dd, src1_dd, dst_dd, ne0, ne1, ne2, ne3_fastdiv, ne10, ne11, ne12, ne13, + /*s0,*/ s1, s2, s3, + s00 ,s01, s02, s03, + s10, s11, s12, s13, (const src1_t *) dst->src[I + 1]->data...); + } else { + k_bin_bcast<<>>( + src0_dd, src1_dd, dst_dd, ne0, ne1, ne2, ne3_fastdiv, ne10, ne11, ne12, ne13, + /*s0,*/ s1, s2, s3, + s00, s01, s02, s03, + s10, s11, s12, s13); + } + } } } @@ -417,76 +417,76 @@ static void launch_bin_bcast_two_sided(const ggml_tensor * src0, const ggml_tens template static __global__ void k_repeat_back( - const T * __restrict__ src, T * __restrict__ dst, const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03, - const size_t s00, const size_t s01, const size_t s02, const size_t s03, - const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3) { - - const int64_t tid0 = int64_t(blockIdx.x)*blockDim.x + threadIdx.x; - const int64_t tid1 = int64_t(blockIdx.y)*blockDim.y + threadIdx.y; - const int64_t tid23 = int64_t(blockIdx.z)*blockDim.z + threadIdx.z; - const int64_t tid2 = tid23 % ne2; - const int64_t tid3 = tid23 / ne2; - - if (tid0 >= ne0) { - return; - } - - T sum = 0; - for (int64_t i3 = tid3; i3 < ne03; i3 += ne3) { - for (int64_t i2 = tid2; i2 < ne02; i2 += ne2) { - for (int64_t i1 = tid1; i1 < ne01; i1 += ne1) { - for (int64_t i0 = tid0; i0 < ne00; i0 += ne0) { - sum += src[i3*s03 + i2*s02 + i1*s01 + i0*s00]; - } - } - } - } - dst[tid3*ne2*ne1*ne0 + tid2*ne1*ne0 + tid1*ne0 + tid0] = sum; -} - -template -struct bin_bcast_cuda { - template - void operator()(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst, - const src0_t * src0_dd, const src1_t * src1_dd, dst_t * dst_dd, - cudaStream_t stream) { - launch_bin_bcast_pack( - src0, src1, dst, src0_dd, src1_dd, dst_dd, stream, std::make_index_sequence{}); - } -}; - -template -static void repeat_back_cuda( - const T * src, T * dst, const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03, - const size_t s00, const size_t s01, const size_t s02, const size_t s03, - const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3, cudaStream_t stream) { - - const dim3 block_dims(WARP_SIZE, 1, 1); - const dim3 block_nums((ne0 + WARP_SIZE - 1) / WARP_SIZE, ne1, ne2*ne3); - k_repeat_back<<>> - (src, dst, ne00, ne01, ne02, ne03, s00, s01, s02, s03, ne0, ne1, ne2, ne3); -} - -template -static void ggml_cuda_op_bin_bcast( - const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, - const void * src0_dd, const void * src1_dd, void * dst_dd, cudaStream_t stream) { - - GGML_ASSERT(src1->type == GGML_TYPE_F32 || src1->type == GGML_TYPE_F16); - - if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) { - op()(src0, src1, dst, (const float *)src0_dd, (const float *)src1_dd, (float *)dst_dd, stream); - } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F16) { - op()(src0, src1, dst, (const half *) src0_dd, (const half *)src1_dd, (half *) dst_dd, stream); - } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F16) { - op()(src0, src1, dst, (const half *) src0_dd, (const float *)src1_dd, (half *) dst_dd, stream); - } else if (src0->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F32) { - op()(src0, src1, dst, (const half *) src0_dd, (const float *)src1_dd, (float *)dst_dd, stream); - } else { - fprintf(stderr, "%s: unsupported types: dst: %s, src0: %s, src1: %s\n", __func__, - ggml_type_name(dst->type), ggml_type_name(src0->type), ggml_type_name(src1->type)); - GGML_ABORT("fatal error"); - } + const T * __restrict__ src, T * __restrict__ dst, const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03, + const size_t s00, const size_t s01, const size_t s02, const size_t s03, + const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3) { + + const int64_t tid0 = int64_t(blockIdx.x)*blockDim.x + threadIdx.x; + const int64_t tid1 = int64_t(blockIdx.y)*blockDim.y + threadIdx.y; + const int64_t tid23 = int64_t(blockIdx.z)*blockDim.z + threadIdx.z; + const int64_t tid2 = tid23 % ne2; + const int64_t tid3 = tid23 / ne2; + + if (tid0 >= ne0) { + return; + } + + T sum = 0; + for (int64_t i3 = tid3; i3 < ne03; i3 += ne3) { + for (int64_t i2 = tid2; i2 < ne02; i2 += ne2) { + for (int64_t i1 = tid1; i1 < ne01; i1 += ne1) { + for (int64_t i0 = tid0; i0 < ne00; i0 += ne0) { + sum += src[i3*s03 + i2*s02 + i1*s01 + i0*s00]; + } + } + } + } + dst[tid3*ne2*ne1*ne0 + tid2*ne1*ne0 + tid1*ne0 + tid0] = sum; +} + +template +struct bin_bcast_cuda { + template + void operator()(const struct ggml_tensor * src0, const struct ggml_tensor * src1, struct ggml_tensor * dst, + const src0_t * src0_dd, const src1_t * src1_dd, dst_t * dst_dd, + cudaStream_t stream) { + launch_bin_bcast_pack( + src0, src1, dst, src0_dd, src1_dd, dst_dd, stream, std::make_index_sequence{}); + } +}; + +template +static void repeat_back_cuda( + const T * src, T * dst, const int64_t ne00, const int64_t ne01, const int64_t ne02, const int64_t ne03, + const size_t s00, const size_t s01, const size_t s02, const size_t s03, + const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3, cudaStream_t stream) { + + const dim3 block_dims(WARP_SIZE, 1, 1); + const dim3 block_nums((ne0 + WARP_SIZE - 1) / WARP_SIZE, ne1, ne2*ne3); + k_repeat_back<<>> + (src, dst, ne00, ne01, ne02, ne03, s00, s01, s02, s03, ne0, ne1, ne2, ne3); +} + +template +static void ggml_cuda_op_bin_bcast( + const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst, + const void * src0_dd, const void * src1_dd, void * dst_dd, cudaStream_t stream) { + + GGML_ASSERT(src1->type == GGML_TYPE_F32 || src1->type == GGML_TYPE_F16); + + if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) { + op()(src0, src1, dst, (const float *)src0_dd, (const float *)src1_dd, (float *)dst_dd, stream); + } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F16) { + op()(src0, src1, dst, (const half *) src0_dd, (const half *)src1_dd, (half *) dst_dd, stream); + } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F16) { + op()(src0, src1, dst, (const half *) src0_dd, (const float *)src1_dd, (half *) dst_dd, stream); + } else if (src0->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F32) { + op()(src0, src1, dst, (const half *) src0_dd, (const float *)src1_dd, (float *)dst_dd, stream); + } else { + fprintf(stderr, "%s: unsupported types: dst: %s, src0: %s, src1: %s\n", __func__, + ggml_type_name(dst->type), ggml_type_name(src0->type), ggml_type_name(src1->type)); + GGML_ABORT("fatal error"); + } } template @@ -548,6 +548,80 @@ void ggml_cuda_op_mul(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { ggml_cuda_op_bin_bcast>(dst->src[0], dst->src[1], dst, dst->src[0]->data, dst->src[1]->data, dst->data, ctx.stream()); } +template +static __global__ void rope_interleaved_pairs_kernel( + const T * __restrict__ even, + const T * __restrict__ odd, + const float * __restrict__ cos, + const float * __restrict__ sin, + T * __restrict__ dst, + int64_t n, + int64_t even_s, + int64_t odd_s, + int64_t cos_s, + int64_t sin_s, + int64_t dst_s) { + for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x; i < n; i += int64_t(blockDim.x) * gridDim.x) { + const float x0 = float(even[i * even_s]); + const float x1 = float(odd[i * odd_s]); + const float c = cos[i * cos_s]; + const float s = sin[i * sin_s]; + dst[i * dst_s + 0] = T(x0 * c - x1 * s); + dst[i * dst_s + 1] = T(x1 * c + x0 * s); + } +} + +template +static void rope_interleaved_pairs_cuda( + const ggml_tensor * even, + const ggml_tensor * odd, + const ggml_tensor * cos, + const ggml_tensor * sin, + ggml_tensor * dst, + cudaStream_t stream) { + GGML_ASSERT(even->ne[0] == 1 && odd->ne[0] == 1 && cos->ne[0] == 1 && sin->ne[0] == 1); + GGML_ASSERT(dst->ne[0] == 2); + const int64_t n = ggml_nelements(even); + constexpr int block_size = 256; + const int64_t blocks = std::min((n + block_size - 1) / block_size, 65535); + rope_interleaved_pairs_kernel<<>>( + (const T *) even->data, + (const T *) odd->data, + (const float *) cos->data, + (const float *) sin->data, + (T *) dst->data, + n, + even->nb[1] / sizeof(T), + odd->nb[1] / sizeof(T), + cos->nb[1] / sizeof(float), + sin->nb[1] / sizeof(float), + dst->nb[1] / sizeof(T)); +} + +void ggml_cuda_op_rope_interleaved_pairs(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + GGML_ASSERT(dst->op == GGML_OP_ROPE_INTERLEAVED_PAIRS); + const ggml_tensor * even = dst->src[0]; + const ggml_tensor * odd = dst->src[1]; + const ggml_tensor * cos = dst->src[2]; + const ggml_tensor * sin = dst->src[3]; + GGML_ASSERT(even != nullptr && odd != nullptr && cos != nullptr && sin != nullptr); + GGML_ASSERT(cos->type == GGML_TYPE_F32 && sin->type == GGML_TYPE_F32); + GGML_ASSERT(dst->type == even->type && dst->type == odd->type); + + ggml_cuda_pool_alloc tmp(ctx.pool(), ggml_nbytes(dst)); + ggml_tensor tmp_concat = *dst; + tmp_concat.data = tmp.get(); + + if (dst->type == GGML_TYPE_F32) { + rope_interleaved_pairs_cuda(even, odd, cos, sin, &tmp_concat, ctx.stream()); + } else if (dst->type == GGML_TYPE_F16) { + rope_interleaved_pairs_cuda(even, odd, cos, sin, &tmp_concat, ctx.stream()); + } else { + GGML_ABORT("fatal error"); + } + CUDA_CHECK(cudaMemcpyAsync(dst->data, tmp.get(), ggml_nbytes(dst), cudaMemcpyDeviceToDevice, ctx.stream())); +} + void ggml_cuda_op_div(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { if (!ggml_are_same_shape(dst->src[0], dst)) { ggml_cuda_op_bin_bcast_two_sided( @@ -556,126 +630,126 @@ void ggml_cuda_op_div(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { } ggml_cuda_op_bin_bcast>(dst->src[0], dst->src[1], dst, dst->src[0]->data, dst->src[1]->data, dst->data, ctx.stream()); } - -template -static void ggml_cuda_op_fused_binbcast_impl(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { - cudaStream_t stream = ctx.stream(); - - const ggml_tensor * src0 = dst->src[0]; - const ggml_tensor * src1 = dst->src[1]; - - if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) { - launch_bin_bcast_pack(src0, src1, dst, - (const float *) src0->data, (const float *) src1->data, (float *) dst->data, - stream, std::make_index_sequence{}); - } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F16) { - launch_bin_bcast_pack(src0, src1, dst, - (const half *) src0->data, (const half *) src1->data, (half *) dst->data, - stream, std::make_index_sequence{}); - } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F16) { - launch_bin_bcast_pack(src0, src1, dst, - (const half *) src0->data, (const float *) src1->data, (half *) dst->data, - stream, std::make_index_sequence{}); - } else if (src0->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F32) { - launch_bin_bcast_pack(src0, src1, dst, - (const half *) src0->data, (const float *) src1->data, (float *) dst->data, - stream, std::make_index_sequence{}); - } else { - fprintf(stderr, - "%s: unsupported types for fusion: dst: %s, src0: %s, src1: %s\n", - __func__, ggml_type_name(dst->type), ggml_type_name(src0->type), ggml_type_name(src1->type)); - GGML_ABORT("fatal error"); - } -} - - -void ggml_cuda_op_fused_add(ggml_backend_cuda_context & ctx, ggml_tensor * dst, int n_fuse) { - GGML_ASSERT(2 <= n_fuse && n_fuse <= 8); - - switch (n_fuse) { - case 2: - ggml_cuda_op_fused_binbcast_impl(ctx, dst); - break; - case 3: - ggml_cuda_op_fused_binbcast_impl(ctx, dst); - break; - case 4: - ggml_cuda_op_fused_binbcast_impl(ctx, dst); - break; - case 5: - ggml_cuda_op_fused_binbcast_impl(ctx, dst); - break; - case 6: - ggml_cuda_op_fused_binbcast_impl(ctx, dst); - break; - case 7: - ggml_cuda_op_fused_binbcast_impl(ctx, dst); - break; - case 8: - ggml_cuda_op_fused_binbcast_impl(ctx, dst); - break; - default: - GGML_ASSERT(false && "Unsupported n_fuse value"); - } -} - -void ggml_cuda_op_fused_mul(ggml_backend_cuda_context & ctx, ggml_tensor * dst, int n_fuse) { - GGML_ASSERT(2 <= n_fuse && n_fuse <= 8); - - switch (n_fuse) { - case 2: - ggml_cuda_op_fused_binbcast_impl(ctx, dst); - break; - case 3: - ggml_cuda_op_fused_binbcast_impl(ctx, dst); - break; - case 4: - ggml_cuda_op_fused_binbcast_impl(ctx, dst); - break; - case 5: - ggml_cuda_op_fused_binbcast_impl(ctx, dst); - break; - case 6: - ggml_cuda_op_fused_binbcast_impl(ctx, dst); - break; - case 7: - ggml_cuda_op_fused_binbcast_impl(ctx, dst); - break; - case 8: - ggml_cuda_op_fused_binbcast_impl(ctx, dst); - break; - default: - GGML_ASSERT(false && "Unsupported n_fuse value"); - } -} - -void ggml_cuda_op_repeat_back(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { - const ggml_tensor * src0 = dst->src[0]; - - GGML_ASSERT(src0->type == dst->type); - GGML_ASSERT(ggml_is_contiguous(dst)); - GGML_ASSERT(ggml_can_repeat(dst, src0)); - - cudaStream_t stream = ctx.stream(); - - GGML_TENSOR_UNARY_OP_LOCALS; - - GGML_ASSERT(ne2*ne3 <= (1 << 15)); - - const size_t ts = ggml_type_size(src0->type); - const size_t s00 = nb00 / ts; - const size_t s01 = nb01 / ts; - const size_t s02 = nb02 / ts; - const size_t s03 = nb03 / ts; - - switch (dst->type) { - case GGML_TYPE_F32: { - const float * src0_d = (const float *) src0->data; - float * dst_d = (float *) dst->data; - repeat_back_cuda(src0_d, dst_d, ne00, ne01, ne02, ne03, s00, s01, s02, s03, ne0, ne1, ne2, ne3, stream); - } break; - default: { - GGML_ASSERT(false); - } break; - } -} + +template +static void ggml_cuda_op_fused_binbcast_impl(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + cudaStream_t stream = ctx.stream(); + + const ggml_tensor * src0 = dst->src[0]; + const ggml_tensor * src1 = dst->src[1]; + + if (src0->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32) { + launch_bin_bcast_pack(src0, src1, dst, + (const float *) src0->data, (const float *) src1->data, (float *) dst->data, + stream, std::make_index_sequence{}); + } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F16) { + launch_bin_bcast_pack(src0, src1, dst, + (const half *) src0->data, (const half *) src1->data, (half *) dst->data, + stream, std::make_index_sequence{}); + } else if (src0->type == GGML_TYPE_F16 && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F16) { + launch_bin_bcast_pack(src0, src1, dst, + (const half *) src0->data, (const float *) src1->data, (half *) dst->data, + stream, std::make_index_sequence{}); + } else if (src0->type == GGML_TYPE_F16 && dst->type == GGML_TYPE_F32) { + launch_bin_bcast_pack(src0, src1, dst, + (const half *) src0->data, (const float *) src1->data, (float *) dst->data, + stream, std::make_index_sequence{}); + } else { + fprintf(stderr, + "%s: unsupported types for fusion: dst: %s, src0: %s, src1: %s\n", + __func__, ggml_type_name(dst->type), ggml_type_name(src0->type), ggml_type_name(src1->type)); + GGML_ABORT("fatal error"); + } +} + + +void ggml_cuda_op_fused_add(ggml_backend_cuda_context & ctx, ggml_tensor * dst, int n_fuse) { + GGML_ASSERT(2 <= n_fuse && n_fuse <= 8); + + switch (n_fuse) { + case 2: + ggml_cuda_op_fused_binbcast_impl(ctx, dst); + break; + case 3: + ggml_cuda_op_fused_binbcast_impl(ctx, dst); + break; + case 4: + ggml_cuda_op_fused_binbcast_impl(ctx, dst); + break; + case 5: + ggml_cuda_op_fused_binbcast_impl(ctx, dst); + break; + case 6: + ggml_cuda_op_fused_binbcast_impl(ctx, dst); + break; + case 7: + ggml_cuda_op_fused_binbcast_impl(ctx, dst); + break; + case 8: + ggml_cuda_op_fused_binbcast_impl(ctx, dst); + break; + default: + GGML_ASSERT(false && "Unsupported n_fuse value"); + } +} + +void ggml_cuda_op_fused_mul(ggml_backend_cuda_context & ctx, ggml_tensor * dst, int n_fuse) { + GGML_ASSERT(2 <= n_fuse && n_fuse <= 8); + + switch (n_fuse) { + case 2: + ggml_cuda_op_fused_binbcast_impl(ctx, dst); + break; + case 3: + ggml_cuda_op_fused_binbcast_impl(ctx, dst); + break; + case 4: + ggml_cuda_op_fused_binbcast_impl(ctx, dst); + break; + case 5: + ggml_cuda_op_fused_binbcast_impl(ctx, dst); + break; + case 6: + ggml_cuda_op_fused_binbcast_impl(ctx, dst); + break; + case 7: + ggml_cuda_op_fused_binbcast_impl(ctx, dst); + break; + case 8: + ggml_cuda_op_fused_binbcast_impl(ctx, dst); + break; + default: + GGML_ASSERT(false && "Unsupported n_fuse value"); + } +} + +void ggml_cuda_op_repeat_back(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + const ggml_tensor * src0 = dst->src[0]; + + GGML_ASSERT(src0->type == dst->type); + GGML_ASSERT(ggml_is_contiguous(dst)); + GGML_ASSERT(ggml_can_repeat(dst, src0)); + + cudaStream_t stream = ctx.stream(); + + GGML_TENSOR_UNARY_OP_LOCALS; + + GGML_ASSERT(ne2*ne3 <= (1 << 15)); + + const size_t ts = ggml_type_size(src0->type); + const size_t s00 = nb00 / ts; + const size_t s01 = nb01 / ts; + const size_t s02 = nb02 / ts; + const size_t s03 = nb03 / ts; + + switch (dst->type) { + case GGML_TYPE_F32: { + const float * src0_d = (const float *) src0->data; + float * dst_d = (float *) dst->data; + repeat_back_cuda(src0_d, dst_d, ne00, ne01, ne02, ne03, s00, s01, s02, s03, ne0, ne1, ne2, ne3, stream); + } break; + default: { + GGML_ASSERT(false); + } break; + } +} diff --git a/external/ggml/src/ggml-cuda/binbcast.cuh b/external/ggml/src/ggml-cuda/binbcast.cuh index 2bd0f0b72..00d8d59d4 100644 --- a/external/ggml/src/ggml-cuda/binbcast.cuh +++ b/external/ggml/src/ggml-cuda/binbcast.cuh @@ -4,6 +4,7 @@ void ggml_cuda_op_repeat(ggml_backend_cuda_context & ctx, ggml_tensor * dst); void ggml_cuda_op_add(ggml_backend_cuda_context & ctx, ggml_tensor * dst); void ggml_cuda_op_sub(ggml_backend_cuda_context & ctx, ggml_tensor * dst); void ggml_cuda_op_mul(ggml_backend_cuda_context & ctx, ggml_tensor * dst); +void ggml_cuda_op_rope_interleaved_pairs(ggml_backend_cuda_context & ctx, ggml_tensor * dst); void ggml_cuda_op_div(ggml_backend_cuda_context & ctx, ggml_tensor * dst); void ggml_cuda_op_repeat_back(ggml_backend_cuda_context & ctx, ggml_tensor * dst); diff --git a/external/ggml/src/ggml-cuda/concat.cu b/external/ggml/src/ggml-cuda/concat.cu index 9a4537b37..109a2689a 100644 --- a/external/ggml/src/ggml-cuda/concat.cu +++ b/external/ggml/src/ggml-cuda/concat.cu @@ -77,6 +77,87 @@ static void concat_cuda( concat_cont<<>>(x, y, dst, ne00, ne01, ne02, ne0, ne1, ne2); } +template +static __global__ void __launch_bounds__(CUDA_CONCAT_BLOCK_SIZE) concat_cont_4d( + const T * x, + const T * y, + T * dst, + int64_t ne00, + int64_t ne01, + int64_t ne02, + int64_t ne0, + int64_t ne1, + int64_t ne2, + int64_t ne3) { + static_assert(dim >= 0 && dim <= 2, "dim must be in [0, 2]"); + + const int64_t n = ne0 * ne1 * ne2 * ne3; + for (int64_t i = (int64_t) blockIdx.x * blockDim.x + threadIdx.x; i < n; i += (int64_t) blockDim.x * gridDim.x) { + if constexpr (dim == 0) { + const int64_t row = i / ne0; + const int64_t i0 = i - row * ne0; + if (i0 < ne00) { + dst[i] = x[row * ne00 + i0]; + } else { + dst[i] = y[row * (ne0 - ne00) + (i0 - ne00)]; + } + } else if constexpr (dim == 1) { + const int64_t dst_vol = ne0 * ne1 * ne2; + const int64_t src0_plane = ne0 * ne01; + const int64_t src1_plane = ne0 * (ne1 - ne01); + const int64_t src0_vol = src0_plane * ne2; + const int64_t src1_vol = src1_plane * ne2; + const int64_t i3 = i / dst_vol; + const int64_t within = i - i3 * dst_vol; + const int64_t i2 = within / (ne0 * ne1); + const int64_t i01 = within - i2 * ne0 * ne1; + if (i01 < src0_plane) { + dst[i] = x[i3 * src0_vol + i2 * src0_plane + i01]; + } else { + dst[i] = y[i3 * src1_vol + i2 * src1_plane + (i01 - src0_plane)]; + } + } else { + const int64_t dst_vol = ne0 * ne1 * ne2; + const int64_t src0_size = ne0 * ne1 * ne02; + const int64_t src1_size = dst_vol - src0_size; + const int64_t i3 = i / dst_vol; + const int64_t within = i - i3 * dst_vol; + if (within < src0_size) { + dst[i] = x[i3 * src0_size + within]; + } else { + dst[i] = y[i3 * src1_size + (within - src0_size)]; + } + } + } +} + +template +static void concat_cuda_4d( + const T * x, + const T * y, + T * dst, + int64_t ne00, + int64_t ne01, + int64_t ne02, + int64_t ne0, + int64_t ne1, + int64_t ne2, + int64_t ne3, + int dim, + cudaStream_t stream) { + const int64_t n = ne0 * ne1 * ne2 * ne3; + const int num_blocks = (n + CUDA_CONCAT_BLOCK_SIZE - 1) / CUDA_CONCAT_BLOCK_SIZE; + if (dim == 0) { + concat_cont_4d<<>>(x, y, dst, ne00, ne01, ne02, ne0, ne1, ne2, ne3); + return; + } + if (dim == 1) { + concat_cont_4d<<>>(x, y, dst, ne00, ne01, ne02, ne0, ne1, ne2, ne3); + return; + } + concat_cont_4d<<>>(x, y, dst, ne00, ne01, ne02, ne0, ne1, ne2, ne3); +} + // non-contiguous kernel (slow) template static __global__ void __launch_bounds__(CUDA_CONCAT_BLOCK_SIZE) concat_non_cont( @@ -146,6 +227,14 @@ static void concat_cuda_typed(ggml_tensor * dst, cudaStream_t stream, int32_t di T * dst_d = (T *) dst->data; if (dim != 3) { + if (ggml_get_op_params_i32(dst, 1) == GGML_CONCAT_LOWERING_CUDA_CONTIGUOUS_4D) { + concat_cuda_4d( + src0_d, src1_d, dst_d, + src0->ne[0], src0->ne[1], src0->ne[2], + dst->ne[0], dst->ne[1], dst->ne[2], dst->ne[3], + dim, stream); + return; + } for (int i3 = 0; i3 < dst->ne[3]; i3++) { concat_cuda( src0_d + i3 * (src0->nb[3] / sizeof(T)), diff --git a/external/ggml/src/ggml-cuda/conv3d.cu b/external/ggml/src/ggml-cuda/conv3d.cu new file mode 100644 index 000000000..ce86b27a3 --- /dev/null +++ b/external/ggml/src/ggml-cuda/conv3d.cu @@ -0,0 +1,352 @@ +#include "conv3d.cuh" +#include "convert.cuh" + +#include +#include + +#define CUDA_CONV3D_BLOCK_SIZE 256 +#define CUDA_CONV3D_SPATIAL_IM2COL_BLOCK_SIZE 256 + +template +static __device__ float conv3d_load_as_float(const T * ptr) { + return static_cast(*ptr); +} + +template <> +__device__ float conv3d_load_as_float(const half * ptr) { + return __half2float(*ptr); +} + +template +static __global__ void conv3d_concat_pad_spatial_im2col_kernel( + const TA * __restrict__ a, + const TB * __restrict__ b, + half * __restrict__ dst, + int IC, int AT, int BT, int H, int W, int OD, int OH, int OW, + int lp0, int lp1, int lp2, + int src_t_offset, + int spatial_y_offset, + int tile_oh, + int a_stride_q, int a_stride_z, int a_stride_y, int a_stride_x, + int b_stride_q, int b_stride_z, int b_stride_y, int b_stride_x) { + const int local = threadIdx.x; + if (local >= C_TILE * 9) { + return; + } + + const int groups = (IC + C_TILE - 1) / C_TILE; + const int group = blockIdx.x % groups; + const int local_src_t = blockIdx.x / groups; + const int src_t_pad = local_src_t + src_t_offset; + const int base_iow = blockIdx.y * X_TILE; + const int local_ioh = blockIdx.z; + const int ioh = spatial_y_offset + local_ioh; + + const int local_channel = local / 9; + const int iic = group * C_TILE + local_channel; + if (iic >= IC) { + return; + } + + const int rem0 = local - local_channel * 9; + const int ikh = rem0 / 3; + const int ikw = rem0 - ikh * 3; + const int src_y_pad = ioh + ikh; + const bool y_ok = src_y_pad >= lp1 && src_y_pad < lp1 + H; + const int src_y = src_y_pad - lp1; + + const bool t_in_a = src_t_pad >= lp2 && src_t_pad < lp2 + AT; + const bool t_in_b = src_t_pad >= lp2 + AT && src_t_pad < lp2 + AT + BT; + const int src_t_a = src_t_pad - lp2; + const int src_t_b = src_t_pad - lp2 - AT; + + const int K2 = IC * 9; + const int M = tile_oh * OW; + const int64_t dst_base = + (static_cast(local_src_t) * M + static_cast(local_ioh) * OW + base_iow) * K2 + + static_cast(iic) * 9 + rem0; + +#pragma unroll + for (int dx = 0; dx < X_TILE; ++dx) { + const int iow = base_iow + dx; + if (iow >= OW) { + return; + } + + const int src_x_pad = iow + ikw; + float value = 0.0f; + if (y_ok && src_x_pad >= lp0 && src_x_pad < lp0 + W) { + const int src_x = src_x_pad - lp0; + if (t_in_a) { + value = conv3d_load_as_float(a + iic * a_stride_q + src_t_a * a_stride_z + src_y * a_stride_y + src_x * a_stride_x); + } else if (t_in_b) { + value = conv3d_load_as_float(b + iic * b_stride_q + src_t_b * b_stride_z + src_y * b_stride_y + src_x * b_stride_x); + } + } + dst[dst_base + dx * K2] = __float2half(value); + } + + GGML_UNUSED(OD); + GGML_UNUSED(OH); +} + +template +static void launch_conv3d_concat_pad_spatial_im2col( + cudaStream_t stream, + dim3 grid, + int threads, + const ggml_tensor * a, + const ggml_tensor * b, + half * im2col, + int IC, int AT, int BT, int H, int W, int OD, int OH, int OW, + int lp0, int lp1, int lp2, int src_t_offset, int spatial_y_offset, int tile_oh, + int a_type_size, int b_type_size) { + conv3d_concat_pad_spatial_im2col_kernel<<>>( + reinterpret_cast(a->data), + reinterpret_cast(b->data), + im2col, + IC, AT, BT, H, W, OD, OH, OW, + lp0, lp1, lp2, src_t_offset, spatial_y_offset, tile_oh, + static_cast(a->nb[3] / a_type_size), static_cast(a->nb[2] / a_type_size), + static_cast(a->nb[1] / a_type_size), static_cast(a->nb[0] / a_type_size), + static_cast(b->nb[3] / b_type_size), static_cast(b->nb[2] / b_type_size), + static_cast(b->nb[1] / b_type_size), static_cast(b->nb[0] / b_type_size)); +} + +template +static void dispatch_conv3d_concat_pad_spatial_im2col( + cudaStream_t stream, + dim3 grid, + int threads, + const ggml_tensor * a, + const ggml_tensor * b, + half * im2col, + int IC, int AT, int BT, int H, int W, int OD, int OH, int OW, + int lp0, int lp1, int lp2, int src_t_offset, int spatial_y_offset, int tile_oh) { + const int a_type_size = ggml_type_size(a->type); + const int b_type_size = ggml_type_size(b->type); + if (a->type == GGML_TYPE_F16 && b->type == GGML_TYPE_F16) { + launch_conv3d_concat_pad_spatial_im2col( + stream, grid, threads, a, b, im2col, IC, AT, BT, H, W, OD, OH, OW, lp0, lp1, lp2, src_t_offset, + spatial_y_offset, tile_oh, a_type_size, b_type_size); + } else if (a->type == GGML_TYPE_F16) { + launch_conv3d_concat_pad_spatial_im2col( + stream, grid, threads, a, b, im2col, IC, AT, BT, H, W, OD, OH, OW, lp0, lp1, lp2, src_t_offset, + spatial_y_offset, tile_oh, a_type_size, b_type_size); + } else if (b->type == GGML_TYPE_F16) { + launch_conv3d_concat_pad_spatial_im2col( + stream, grid, threads, a, b, im2col, IC, AT, BT, H, W, OD, OH, OW, lp0, lp1, lp2, src_t_offset, + spatial_y_offset, tile_oh, a_type_size, b_type_size); + } else { + launch_conv3d_concat_pad_spatial_im2col( + stream, grid, threads, a, b, im2col, IC, AT, BT, H, W, OD, OH, OW, lp0, lp1, lp2, src_t_offset, + spatial_y_offset, tile_oh, a_type_size, b_type_size); + } +} + +static __global__ void conv3d_zero_spatial_prefix_kernel( + half * __restrict__ dst, + int rows, + int ld, + int OC) { + const int64_t n = int64_t(rows) * OC; + for (int64_t i = int64_t(blockIdx.x) * blockDim.x + threadIdx.x; i < n; i += int64_t(blockDim.x) * gridDim.x) { + const int row = i % rows; + const int oc = i / rows; + dst[row + oc * ld] = __float2half(0.0f); + } +} + +static __global__ void conv3d_k3_weight_f32_to_f16_spatial_kernel( + const float * __restrict__ w, + half * __restrict__ dst, + int IC, int OC) { + const int K2 = IC * 9; + const int total = 3 * K2 * OC; + const int i = blockIdx.x * blockDim.x + threadIdx.x; + if (i >= total) { + return; + } + + const int rem_oc = i % (K2 * OC); + const int kd = i / (K2 * OC); + const int oc = rem_oc / K2; + const int k2 = rem_oc - oc * K2; + const int ic = k2 / 9; + const int rem = k2 - ic * 9; + const int kh = rem / 3; + const int kw = rem - kh * 3; + const int src = kw + kh * 3 + kd * 9 + (oc * IC + ic) * 27; + dst[i] = __float2half(w[src]); +} + +void ggml_cuda_op_conv3d_concat_pad_spatial_gemm(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + const ggml_tensor * a = dst->src[0]; + const ggml_tensor * b = dst->src[1]; + const ggml_tensor * w = dst->src[2]; + + GGML_ASSERT(a->type == GGML_TYPE_F16 || a->type == GGML_TYPE_F32); + GGML_ASSERT(b->type == GGML_TYPE_F16 || b->type == GGML_TYPE_F32); + GGML_ASSERT(w->type == GGML_TYPE_F32); + GGML_ASSERT(dst->type == GGML_TYPE_F16 || dst->type == GGML_TYPE_F32); + GGML_ASSERT(w->ne[0] == 3 && w->ne[1] == 3 && w->ne[2] == 3); + + const int32_t * params = reinterpret_cast(dst->op_params); + const int lp0 = params[0]; + const int lp1 = params[2]; + const int lp2 = params[4]; + + const int64_t IC = a->ne[3]; + const int64_t AT = a->ne[2]; + const int64_t BT = b->ne[2]; + const int64_t H = a->ne[1]; + const int64_t W = a->ne[0]; + const int64_t OD = dst->ne[2]; + const int64_t OH = dst->ne[1]; + const int64_t OW = dst->ne[0]; + const int64_t OC = dst->ne[3]; + const int64_t K2 = IC * 9; + const int64_t M = OH * OW; + const int64_t source_frames = OD + 2; + + GGML_ASSERT(w->ne[3] == IC * OC); + GGML_ASSERT(IC <= INT_MAX && AT <= INT_MAX && BT <= INT_MAX && H <= INT_MAX && W <= INT_MAX); + GGML_ASSERT(OD <= INT_MAX && OH <= INT_MAX && OW <= INT_MAX && OC <= INT_MAX); + const int a_type_size = ggml_type_size(a->type); + const int b_type_size = ggml_type_size(b->type); + GGML_ASSERT(a->nb[0] / a_type_size <= INT_MAX && a->nb[1] / a_type_size <= INT_MAX && + a->nb[2] / a_type_size <= INT_MAX && a->nb[3] / a_type_size <= INT_MAX); + GGML_ASSERT(b->nb[0] / b_type_size <= INT_MAX && b->nb[1] / b_type_size <= INT_MAX && + b->nb[2] / b_type_size <= INT_MAX && b->nb[3] / b_type_size <= INT_MAX); + + cudaStream_t stream = ctx.stream(); + ggml_cuda_pool_alloc weight(ctx.pool(), static_cast(3 * K2 * OC)); + ggml_cuda_pool_alloc out; + half * out_data = nullptr; + if (dst->type == GGML_TYPE_F16) { + out_data = reinterpret_cast(dst->data); + } else { + out.alloc(ctx.pool(), static_cast(OD * M * OC)); + out_data = out.get(); + } + + const bool skip_front_zero = false; + const bool use_skip_front_zero = skip_front_zero && lp2 == 1 && source_frames == OD + 2 && OD > 1; + + const int groups = (static_cast(IC) + 27) / 28; + const int src_t_offset = use_skip_front_zero ? 1 : 0; + const int weight_blocks = (static_cast(3 * K2 * OC) + CUDA_CONV3D_BLOCK_SIZE - 1) / CUDA_CONV3D_BLOCK_SIZE; + conv3d_k3_weight_f32_to_f16_spatial_kernel<<>>( + reinterpret_cast(w->data), weight.get(), static_cast(IC), static_cast(OC)); + + CUBLAS_CHECK(cublasSetStream(ctx.cublas_handle(), stream)); + + const half alpha = __float2half(1.0f); + const half beta_zero = __float2half(0.0f); + const half beta_one = __float2half(1.0f); + if (use_skip_front_zero) { + const int64_t zero_rows = M * OC; + const int zero_blocks = (static_cast(zero_rows) + CUDA_CONV3D_BLOCK_SIZE - 1) / CUDA_CONV3D_BLOCK_SIZE; + conv3d_zero_spatial_prefix_kernel<<>>( + out_data, static_cast(M), static_cast(OD * M), static_cast(OC)); + } + + const int32_t conv_lowering = ggml_get_op_params_i32(dst, 6); + const bool use_tiled_f16_path = + conv_lowering == GGML_CONV_3D_CONCAT_PAD_SPATIAL_GEMM_LOWERING_CUDA_TILED_C48; + if (!use_tiled_f16_path) { + ggml_cuda_pool_alloc im2col(ctx.pool(), static_cast(source_frames * M * K2)); + const int active_source_frames = static_cast(source_frames) - src_t_offset; + if (conv_lowering == GGML_CONV_3D_CONCAT_PAD_SPATIAL_GEMM_LOWERING_CUDA_C48) { + const int groups_c48 = (static_cast(IC) + 47) / 48; + dim3 im2col_grid(groups_c48 * active_source_frames, (OW + 7) / 8, OH); + dispatch_conv3d_concat_pad_spatial_im2col<8, 48>( + stream, im2col_grid, 48 * 9, a, b, im2col.get(), + static_cast(IC), static_cast(AT), static_cast(BT), + static_cast(H), static_cast(W), static_cast(OD), static_cast(OH), static_cast(OW), + lp0, lp1, lp2, src_t_offset, 0, static_cast(OH)); + } else { + dim3 im2col_grid(groups * active_source_frames, (OW + 7) / 8, OH); + dispatch_conv3d_concat_pad_spatial_im2col<8, 28>( + stream, im2col_grid, CUDA_CONV3D_SPATIAL_IM2COL_BLOCK_SIZE, + a, b, im2col.get(), static_cast(IC), static_cast(AT), static_cast(BT), + static_cast(H), static_cast(W), static_cast(OD), static_cast(OH), static_cast(OW), + lp0, lp1, lp2, src_t_offset, 0, static_cast(OH)); + } + for (int kd = 0; kd < 3; ++kd) { + const int row_offset = use_skip_front_zero && kd == 0 ? static_cast(M) : 0; + const int gemm_rows = static_cast(OD * M) - row_offset; + const half * a_ptr = im2col.get() + (static_cast(kd) * M + row_offset) * K2; + const half * b_ptr = weight.get() + static_cast(kd) * K2 * OC; + half * c_ptr = out_data + row_offset; + const half * beta = kd == 0 ? &beta_zero : &beta_one; + CUBLAS_CHECK(cublasGemmEx( + ctx.cublas_handle(), CUBLAS_OP_T, CUBLAS_OP_N, + gemm_rows, static_cast(OC), static_cast(K2), + &alpha, + a_ptr, CUDA_R_16F, static_cast(K2), + b_ptr, CUDA_R_16F, static_cast(K2), + beta, + c_ptr, CUDA_R_16F, static_cast(OD * M), + CUBLAS_COMPUTE_16F, + CUBLAS_GEMM_DEFAULT_TENSOR_OP)); + } + } else { + constexpr size_t kMaxIm2ColBytes = 128ull * 1024ull * 1024ull; + const int64_t max_source_frames = + std::max(3, static_cast(kMaxIm2ColBytes / (sizeof(half) * static_cast(M) * static_cast(K2)))); + const int64_t max_tile_frames = std::max(1, max_source_frames - 2); + const int64_t tile_frames = std::min(OD, max_tile_frames); + for (int64_t tile_start = 0; tile_start < OD; tile_start += tile_frames) { + const int64_t current_frames = std::min(tile_frames, OD - tile_start); + const int current_source_frames = static_cast(current_frames + 2); + const int64_t max_tile_rows = + std::max(1, static_cast(kMaxIm2ColBytes / (sizeof(half) * static_cast(current_source_frames) * static_cast(OW) * static_cast(K2)))); + const int64_t tile_oh = std::min(OH, max_tile_rows); + for (int64_t y_start = 0; y_start < OH; y_start += tile_oh) { + const int current_tile_oh = static_cast(std::min(tile_oh, OH - y_start)); + const int64_t tile_m = static_cast(current_tile_oh) * OW; + ggml_cuda_pool_alloc im2col(ctx.pool(), static_cast(current_source_frames * tile_m * K2)); + if (conv_lowering == GGML_CONV_3D_CONCAT_PAD_SPATIAL_GEMM_LOWERING_CUDA_TILED_C48) { + const int groups_c48 = (static_cast(IC) + 47) / 48; + dim3 im2col_grid(groups_c48 * current_source_frames, (OW + 7) / 8, current_tile_oh); + dispatch_conv3d_concat_pad_spatial_im2col<8, 48>( + stream, im2col_grid, 48 * 9, a, b, im2col.get(), + static_cast(IC), static_cast(AT), static_cast(BT), + static_cast(H), static_cast(W), static_cast(current_frames), static_cast(OH), static_cast(OW), + lp0, lp1, lp2, src_t_offset + static_cast(tile_start), static_cast(y_start), current_tile_oh); + } else { + dim3 im2col_grid(groups * current_source_frames, (OW + 7) / 8, current_tile_oh); + dispatch_conv3d_concat_pad_spatial_im2col<8, 28>( + stream, im2col_grid, CUDA_CONV3D_SPATIAL_IM2COL_BLOCK_SIZE, + a, b, im2col.get(), static_cast(IC), static_cast(AT), static_cast(BT), + static_cast(H), static_cast(W), static_cast(current_frames), static_cast(OH), static_cast(OW), + lp0, lp1, lp2, src_t_offset + static_cast(tile_start), static_cast(y_start), current_tile_oh); + } + for (int64_t frame = 0; frame < current_frames; ++frame) { + for (int kd = 0; kd < 3; ++kd) { + const half * a_ptr = im2col.get() + (static_cast(kd) + frame) * tile_m * K2; + const half * b_ptr = weight.get() + static_cast(kd) * K2 * OC; + half * c_ptr = out_data + (tile_start + frame) * M + y_start * OW; + const half * beta = kd == 0 ? &beta_zero : &beta_one; + CUBLAS_CHECK(cublasGemmEx( + ctx.cublas_handle(), CUBLAS_OP_T, CUBLAS_OP_N, + static_cast(tile_m), static_cast(OC), static_cast(K2), + &alpha, + a_ptr, CUDA_R_16F, static_cast(K2), + b_ptr, CUDA_R_16F, static_cast(K2), + beta, + c_ptr, CUDA_R_16F, static_cast(OD * M), + CUBLAS_COMPUTE_16F, + CUBLAS_GEMM_DEFAULT_TENSOR_OP)); + } + } + } + } + } + + if (dst->type == GGML_TYPE_F32) { + const to_fp32_cuda_t to_fp32_cuda = ggml_get_to_fp32_cuda(GGML_TYPE_F16); + to_fp32_cuda(out_data, reinterpret_cast(dst->data), OD * M * OC, stream); + } +} diff --git a/external/ggml/src/ggml-cuda/conv3d.cuh b/external/ggml/src/ggml-cuda/conv3d.cuh new file mode 100644 index 000000000..0696e61ee --- /dev/null +++ b/external/ggml/src/ggml-cuda/conv3d.cuh @@ -0,0 +1,5 @@ +#pragma once + +#include "common.cuh" + +void ggml_cuda_op_conv3d_concat_pad_spatial_gemm(ggml_backend_cuda_context & ctx, ggml_tensor * dst); diff --git a/external/ggml/src/ggml-cuda/ggml-cuda.cu b/external/ggml/src/ggml-cuda/ggml-cuda.cu index 278a49c43..b8ffb35bf 100644 --- a/external/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/external/ggml/src/ggml-cuda/ggml-cuda.cu @@ -18,6 +18,7 @@ #include "ggml-cuda/conv2d.cuh" #include "ggml-cuda/conv2d-dw.cuh" #include "ggml-cuda/conv2d-transpose.cuh" +#include "ggml-cuda/conv3d.cuh" #include "ggml-cuda/convert.cuh" #include "ggml-cuda/count-equal.cuh" #include "ggml-cuda/cpy.cuh" @@ -1842,32 +1843,65 @@ static void ggml_cuda_op_mul_mat_cublas( 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 bool tile_f16_output = + ggml_get_op_params_i32(dst, 1) == GGML_MUL_MAT_LOWERING_CUDA_TILE_F16_ACCUM_OUTPUT && + ldc == row_diff; + ggml_cuda_pool_alloc dst_f16(ctx.pool(id)); const half alpha_f16 = 1.0f; const half beta_f16 = 0.0f; + if (tile_f16_output) { + constexpr int64_t max_scratch_elements = 64ll * 1024ll * 1024ll; + const int64_t col_tile = std::max(1, std::min(src1_ncols, max_scratch_elements / row_diff)); + dst_f16.alloc(row_diff * col_tile); + const to_fp32_cuda_t to_fp32_cuda = ggml_get_to_fp32_cuda(GGML_TYPE_F16); + for (int64_t col = 0; col < src1_ncols; col += col_tile) { + const int64_t current_cols = std::min(col_tile, src1_ncols - col); #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); + GGML_UNUSED_VARS(alpha_f16, beta_f16); + ggml_hipblaslt_gemm(ctx, stream, + row_diff, current_cols, ne10, + src0_ptr, CUDA_R_16F, ne00, 0, + src1_ptr + col * ne10, CUDA_R_16F, ne10, 0, + dst_f16.get(), CUDA_R_16F, row_diff, 0, + 1); #else - CUBLAS_CHECK( - cublasGemmEx(ctx.cublas_handle(id), CUBLAS_OP_T, CUBLAS_OP_N, + CUBLAS_CHECK( + cublasGemmEx(ctx.cublas_handle(id), CUBLAS_OP_T, CUBLAS_OP_N, + row_diff, current_cols, ne10, + &alpha_f16, src0_ptr, CUDA_R_16F, ne00, + src1_ptr + col * ne10, CUDA_R_16F, ne10, + &beta_f16, dst_f16.get(), CUDA_R_16F, row_diff, + CUBLAS_COMPUTE_16F, + CUBLAS_GEMM_DEFAULT_TENSOR_OP)); +#endif // defined(GGML_USE_HIP) && defined(GGML_HIP_USE_HIPBLASLT) + to_fp32_cuda(dst_f16.get(), dst_dd_i + col * ldc, row_diff * current_cols, stream); + } + } else { + dst_f16.alloc(row_diff*src1_ncols); +#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, - &alpha_f16, src0_ptr, CUDA_R_16F, ne00, - src1_ptr, CUDA_R_16F, ne10, - &beta_f16, dst_f16.get(), CUDA_R_16F, ldc, - CUBLAS_COMPUTE_16F, - CUBLAS_GEMM_DEFAULT_TENSOR_OP)); + 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, + &alpha_f16, src0_ptr, CUDA_R_16F, ne00, + src1_ptr, CUDA_R_16F, ne10, + &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); + 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); + } } } else { ggml_cuda_pool_alloc src0_ddq_as_f32(ctx.pool(id)); @@ -2719,11 +2753,15 @@ static void ggml_cuda_mul_mat(ggml_backend_cuda_context & ctx, const ggml_tensor && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32; bool use_mul_mat_f = !ggml_is_quantized(src0->type) && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32; + const int device_cc = ggml_cuda_info().devices[ctx.device].cc; + const bool use_nvfp4_f16_mmq = !split && src0->type == GGML_TYPE_NVFP4 && src1->type == GGML_TYPE_F16 && + dst->type == GGML_TYPE_F32 && blackwell_mma_available(device_cc) && + ggml_get_op_params_i32(dst, 1) == GGML_MUL_MAT_LOWERING_CUDA_NVFP4_F16_ACTIVATION; bool use_mul_mat_vec_q = ggml_is_quantized(src0->type) && !bad_padding_clear && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32 && src1->ne[1] <= MMVQ_MAX_BATCH_SIZE; bool use_mul_mat_q = ggml_is_quantized(src0->type) && !bad_padding_clear - && src1->type == GGML_TYPE_F32 && dst->type == GGML_TYPE_F32; + && (src1->type == GGML_TYPE_F32 || use_nvfp4_f16_mmq) && dst->type == GGML_TYPE_F32; bool any_gpus_with_slow_fp16 = false; @@ -3114,6 +3152,9 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg case GGML_OP_CONCAT: ggml_cuda_op_concat(ctx, dst); break; + case GGML_OP_ROPE_INTERLEAVED_PAIRS: + ggml_cuda_op_rope_interleaved_pairs(ctx, dst); + break; case GGML_OP_UPSCALE: ggml_cuda_op_upscale(ctx, dst); break; @@ -3138,6 +3179,15 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg case GGML_OP_RMS_NORM: ggml_cuda_op_rms_norm(ctx, dst); break; + case GGML_OP_RMS_NORM_CHANNELS: + ggml_cuda_op_rms_norm_channels(ctx, dst); + break; + case GGML_OP_RMS_NORM_CHANNELS_SILU: + ggml_cuda_op_rms_norm_channels_silu(ctx, dst); + break; + case GGML_OP_RMS_NORM_CHANNELS_ADD_BIAS_SILU: + ggml_cuda_op_rms_norm_channels_add_bias_silu(ctx, dst); + break; case GGML_OP_RMS_NORM_BACK: ggml_cuda_op_rms_norm_back(ctx, dst); break; @@ -3212,6 +3262,9 @@ static bool ggml_cuda_compute_forward(ggml_backend_cuda_context & ctx, struct gg case GGML_OP_CONV_2D: ggml_cuda_op_conv2d(ctx, dst); break; + case GGML_OP_CONV_3D_CONCAT_PAD_SPATIAL_GEMM: + ggml_cuda_op_conv3d_concat_pad_spatial_gemm(ctx, dst); + break; case GGML_OP_CONV_2D_DW: ggml_cuda_op_conv2d_dw(ctx, dst); break; @@ -5586,6 +5639,9 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g break; case GGML_OP_NORM: case GGML_OP_RMS_NORM: + case GGML_OP_RMS_NORM_CHANNELS: + case GGML_OP_RMS_NORM_CHANNELS_SILU: + case GGML_OP_RMS_NORM_CHANNELS_ADD_BIAS_SILU: case GGML_OP_L2_NORM: return true; case GGML_OP_RMS_NORM_BACK: @@ -5648,6 +5704,11 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_OP_ROPE_BACK: { return op->src[0]->nb[0] == ggml_type_size(op->src[0]->type) && ggml_is_contiguous_2(op->src[0]); } + case GGML_OP_ROPE_INTERLEAVED_PAIRS: + return op->src[0]->type == op->type && + op->src[1]->type == op->type && + op->src[2]->type == GGML_TYPE_F32 && + op->src[3]->type == GGML_TYPE_F32; case GGML_OP_IM2COL: case GGML_OP_IM2COL_FAST_1D: case GGML_OP_IM2COL_3D: @@ -5656,6 +5717,14 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_OP_CONV_TRANSPOSE_2D: case GGML_OP_POOL_2D: return true; + case GGML_OP_CONV_3D_CONCAT_PAD_SPATIAL_GEMM: + return op->src[0]->type == GGML_TYPE_F32 && + op->src[1]->type == GGML_TYPE_F32 && + op->src[2]->type == GGML_TYPE_F32 && + (op->type == GGML_TYPE_F16 || op->type == GGML_TYPE_F32) && + ggml_is_contiguous(op->src[0]) && + ggml_is_contiguous(op->src[1]) && + ggml_is_contiguous(op->src[2]); case GGML_OP_COL2IM_1D: return ggml_is_contiguous(op->src[0]) && (op->src[0]->type == GGML_TYPE_F32 || diff --git a/external/ggml/src/ggml-cuda/im2col.cu b/external/ggml/src/ggml-cuda/im2col.cu index 40700ab6d..2ef3c9d36 100644 --- a/external/ggml/src/ggml-cuda/im2col.cu +++ b/external/ggml/src/ggml-cuda/im2col.cu @@ -43,16 +43,184 @@ static __global__ void im2col_kernel( GGML_UNUSED(KH); } +template +static __global__ void im2col_n_k3_pad1_kernel( + const float * __restrict__ x, T * __restrict__ dst, + int64_t IC, int64_t IW, int64_t IH, int64_t OH, int64_t OW, + int64_t groups, int64_t IC_KH_KW, + int64_t IC_IH_IW, int64_t IH_IW) { + const int64_t local = threadIdx.x; + if (local >= 28 * 9) { + return; + } + const int64_t group = blockIdx.x % groups; + const int64_t in = blockIdx.x / groups; + const int64_t iow = blockIdx.y; + const int64_t ioh = blockIdx.z; + + const int64_t local_channel = local / 9; + const int64_t iic = group * 28 + local_channel; + if (iic >= IC) { + return; + } + const int64_t rem0 = local - local_channel * 9; + const int64_t ikh = rem0 / 3; + const int64_t ikw = rem0 - ikh * 3; + const int64_t iih = ioh + ikh - 1; + const int64_t iiw = iow + ikw - 1; + + const int64_t offset_dst = + ((in * OH + ioh) * OW + iow) * IC_KH_KW + + iic * 9 + + rem0; + if (iih < 0 || iih >= IH || iiw < 0 || iiw >= IW) { + dst[offset_dst] = T(0.0f); + } else { + dst[offset_dst] = T(x[iic * IC_IH_IW + in * IH_IW + iih * IW + iiw]); + } +} + +template +static __global__ void im2col_n_k3_pad1_xtile_kernel( + const float * __restrict__ x, T * __restrict__ dst, + int IC, int IW, int IH, int OH, int OW, + int groups, int IC_KH_KW, + int IC_IH_IW, int IH_IW) { + const int local = threadIdx.x; + if (local >= 28 * 9) { + return; + } + const int group = blockIdx.x % groups; + const int in = blockIdx.x / groups; + const int base_iow = blockIdx.y * X_TILE; + const int ioh = blockIdx.z; + + const int local_channel = local / 9; + const int iic = group * 28 + local_channel; + if (iic >= IC) { + return; + } + const int rem0 = local - local_channel * 9; + const int ikh = rem0 / 3; + const int ikw = rem0 - ikh * 3; + const int iih = ioh + ikh - 1; + +#pragma unroll + for (int dx = 0; dx < X_TILE; ++dx) { + const int iow = base_iow + dx; + if (iow >= OW) { + return; + } + const int iiw = iow + ikw - 1; + const int offset_dst = + ((in * OH + ioh) * OW + iow) * IC_KH_KW + + iic * 9 + + rem0; + if (iih < 0 || iih >= IH || iiw < 0 || iiw >= IW) { + dst[offset_dst] = T(0.0f); + } else { + dst[offset_dst] = T(x[iic * IC_IH_IW + in * IH_IW + iih * IW + iiw]); + } + } +} + +template +static __global__ void im2col_n_k3_nopad_xtile_kernel( + const float * __restrict__ x, T * __restrict__ dst, + int IC, int IW, int IH, int OH, int OW, + int groups, int IC_KH_KW, + int IC_IH_IW, int IH_IW) { + const int local = threadIdx.x; + if (local >= 28 * 9) { + return; + } + const int group = blockIdx.x % groups; + const int in = blockIdx.x / groups; + const int base_iow = blockIdx.y * X_TILE; + const int ioh = blockIdx.z; + + const int local_channel = local / 9; + const int iic = group * 28 + local_channel; + if (iic >= IC) { + return; + } + const int rem0 = local - local_channel * 9; + const int ikh = rem0 / 3; + const int ikw = rem0 - ikh * 3; + const int iih = ioh + ikh; + +#pragma unroll + for (int dx = 0; dx < X_TILE; ++dx) { + const int iow = base_iow + dx; + if (iow >= OW) { + return; + } + const int iiw = iow + ikw; + const int offset_dst = + ((in * OH + ioh) * OW + iow) * IC_KH_KW + + iic * 9 + + rem0; + dst[offset_dst] = T(x[iic * IC_IH_IW + in * IH_IW + iih * IW + iiw]); + } + + GGML_UNUSED(IH); +} + // im2col: [N, IC, IH, IW] => [N, OH, OW, IC*KH*KW] template static void im2col_cuda(const float * x, T* dst, int64_t IW, int64_t IH, int64_t OW, int64_t OH, int64_t KW, int64_t KH, int64_t IC, int64_t N, int64_t IC_IH_IW, int64_t IH_IW, - int s0,int s1,int p0,int p1,int d0,int d1, cudaStream_t stream) { + int s0,int s1,int p0,int p1,int d0,int d1, int32_t lowering, cudaStream_t stream) { const int64_t IC_KH_KW = IC * KH * KW; const int64_t num_blocks = (IC_KH_KW + CUDA_IM2COL_BLOCK_SIZE - 1) / CUDA_IM2COL_BLOCK_SIZE; const int64_t N_OH = N * OH; const int64_t KH_KW = KW*KH; + const bool use_n_k3_pad1_kernel = lowering == GGML_IM2COL_2D_LOWERING_CUDA_N_K3_PAD1_X8; + const bool use_n_k3_pad1_x8_kernel = lowering == GGML_IM2COL_2D_LOWERING_CUDA_N_K3_PAD1_X8; + const bool use_n_k3_nopad_x8_kernel = lowering == GGML_IM2COL_2D_LOWERING_CUDA_N_K3_NOPAD_X8; + if (use_n_k3_nopad_x8_kernel && + KW == 3 && KH == 3 && + s0 == 1 && s1 == 1 && + p0 == 0 && p1 == 0 && + d0 == 1 && d1 == 1 && + IC <= INT_MAX && IW <= INT_MAX && IH <= INT_MAX && + OH <= INT_MAX && OW <= INT_MAX && + IC_KH_KW <= INT_MAX && IC_IH_IW <= INT_MAX && IH_IW <= INT_MAX) { + const int64_t channel_groups = (IC + 27) / 28; + if (channel_groups <= INT_MAX) { + dim3 block_nums(channel_groups * N, MIN((OW + 7) / 8, MAX_GRIDDIM_Y), MIN(OH, MAX_GRIDDIM_Z)); + im2col_n_k3_nopad_xtile_kernel<<>>( + x, dst, + static_cast(IC), static_cast(IW), static_cast(IH), + static_cast(OH), static_cast(OW), static_cast(channel_groups), + static_cast(IC_KH_KW), static_cast(IC_IH_IW), static_cast(IH_IW)); + return; + } + } + if (use_n_k3_pad1_kernel && + KW == 3 && KH == 3 && + s0 == 1 && s1 == 1 && + p0 == 1 && p1 == 1 && + d0 == 1 && d1 == 1) { + const int64_t channel_groups = (IC + 27) / 28; + if (use_n_k3_pad1_x8_kernel && + IC <= INT_MAX && IW <= INT_MAX && IH <= INT_MAX && + OH <= INT_MAX && OW <= INT_MAX && channel_groups <= INT_MAX && + IC_KH_KW <= INT_MAX && IC_IH_IW <= INT_MAX && IH_IW <= INT_MAX) { + dim3 block_nums(channel_groups * N, MIN((OW + 7) / 8, MAX_GRIDDIM_Y), MIN(OH, MAX_GRIDDIM_Z)); + im2col_n_k3_pad1_xtile_kernel<<>>( + x, dst, + static_cast(IC), static_cast(IW), static_cast(IH), + static_cast(OH), static_cast(OW), static_cast(channel_groups), + static_cast(IC_KH_KW), static_cast(IC_IH_IW), static_cast(IH_IW)); + } else { + dim3 block_nums(channel_groups * N, MIN(OW, MAX_GRIDDIM_Y), MIN(OH, MAX_GRIDDIM_Z)); + im2col_n_k3_pad1_kernel<<>>( + x, dst, IC, IW, IH, OH, OW, channel_groups, IC_KH_KW, IC_IH_IW, IH_IW); + } + return; + } dim3 block_nums(num_blocks, MIN(OW, MAX_GRIDDIM_Y), MIN(N_OH, MAX_GRIDDIM_Z)); im2col_kernel<<>>(x, dst, IC, IW, IH, OH, OW, KW, KH, IC_IH_IW, IH_IW, N_OH, KH_KW, IC_KH_KW, @@ -62,17 +230,17 @@ static void im2col_cuda(const float * x, T* dst, static void im2col_cuda_f16(const float * x, half * dst, int64_t IW, int64_t IH, int64_t OW, int64_t OH, int64_t KW, int64_t KH, int64_t IC, int64_t N, int64_t IC_IH_IW, int64_t IH_IW, - int s0,int s1,int p0,int p1,int d0,int d1, cudaStream_t stream) { + int s0,int s1,int p0,int p1,int d0,int d1, int32_t lowering, cudaStream_t stream) { - im2col_cuda(x, dst, IW, IH, OW, OH, KW, KH, IC, N, IC_IH_IW, IH_IW, s0, s1, p0, p1, d0, d1, stream); + im2col_cuda(x, dst, IW, IH, OW, OH, KW, KH, IC, N, IC_IH_IW, IH_IW, s0, s1, p0, p1, d0, d1, lowering, stream); } static void im2col_cuda_f32(const float * x, float * dst, int64_t IW, int64_t IH, int64_t OW, int64_t OH, int64_t KW, int64_t KH, int64_t IC, int64_t N, int64_t IC_IH_IW, int64_t IH_IW, - int s0,int s1,int p0,int p1,int d0,int d1, cudaStream_t stream) { + int s0,int s1,int p0,int p1,int d0,int d1, int32_t lowering, cudaStream_t stream) { - im2col_cuda(x, dst, IW, IH, OW, OH, KW, KH, IC, N, IC_IH_IW, IH_IW, s0, s1, p0, p1, d0, d1, stream); + im2col_cuda(x, dst, IW, IH, OW, OH, KW, KH, IC, N, IC_IH_IW, IH_IW, s0, s1, p0, p1, d0, d1, lowering, stream); } void ggml_cuda_op_im2col(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { @@ -107,11 +275,12 @@ void ggml_cuda_op_im2col(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { const int64_t IC_IH_IW = src1->nb[is_2D ? 2 : 1] / 4; // nb is byte offset, src is type float32 const int64_t N = src1->ne[is_2D ? 3 : 2]; const int64_t IH_IW = src1->nb[is_2D ? 3 : 2] / 4; // nb is byte offset, src is type float32 + const int32_t lowering = ggml_get_op_params_i32(dst, 7); if(dst->type == GGML_TYPE_F16) { - im2col_cuda_f16(src1_d, (half *) dst_d, IW, IH, OW, OH, KW, KH, IC, N, IC_IH_IW, IH_IW, s0, s1, p0, p1, d0, d1, stream); + im2col_cuda_f16(src1_d, (half *) dst_d, IW, IH, OW, OH, KW, KH, IC, N, IC_IH_IW, IH_IW, s0, s1, p0, p1, d0, d1, lowering, stream); } else { - im2col_cuda_f32(src1_d, (float *) dst_d, IW, IH, OW, OH, KW, KH, IC, N, IC_IH_IW, IH_IW, s0, s1, p0, p1, d0, d1, stream); + im2col_cuda_f32(src1_d, (float *) dst_d, IW, IH, OW, OH, KW, KH, IC, N, IC_IH_IW, IH_IW, s0, s1, p0, p1, d0, d1, lowering, stream); } } @@ -160,13 +329,105 @@ static __global__ void im2col_3d_kernel( } } +template +static __global__ void im2col_3d_n1_k3_nopad_kernel( + const float * __restrict__ src, T * __restrict__ dst, + int64_t IC, int64_t IH, int64_t IW, int64_t OH, int64_t OW, + int64_t groups, int64_t IC_KD_KH_KW, + int64_t OH_OW_IC_KD_KH_KW, int64_t OW_IC_KD_KH_KW, + int64_t stride_q, int64_t stride_z, int64_t stride_y, int64_t stride_x) { + const int64_t local = threadIdx.x; + if (local >= 9 * 27) { + return; + } + const int64_t group = blockIdx.x % groups; + const int64_t iod = blockIdx.x / groups; + const int64_t iow = blockIdx.y; + const int64_t ioh = blockIdx.z; + + const int64_t local_channel = local / 27; + const int64_t iic = group * 9 + local_channel; + if (iic >= IC) { + return; + } + const int64_t rem0 = local - local_channel * 27; + const int64_t ikd = rem0 / 9; + const int64_t rem1 = rem0 - ikd * 9; + const int64_t ikh = rem1 / 3; + const int64_t ikw = rem1 - ikh * 3; + + const int64_t offset_dst = + iod * OH_OW_IC_KD_KH_KW + + ioh * OW_IC_KD_KH_KW + + iow * IC_KD_KH_KW + + iic * 27 + + rem0; + const int64_t offset_src = + iic * stride_q + + (iod + ikd) * stride_z + + (ioh + ikh) * stride_y + + (iow + ikw) * stride_x; + dst[offset_dst] = T(src[offset_src]); + + GGML_UNUSED(IH); + GGML_UNUSED(IW); +} + +template +static __global__ void im2col_3d_n1_k3_nopad_xtile_kernel( + const float * __restrict__ src, T * __restrict__ dst, + int IC, int OH, int OW, + int groups, int IC_KD_KH_KW, + int OH_OW_IC_KD_KH_KW, int OW_IC_KD_KH_KW, + int stride_q, int stride_z, int stride_y, int stride_x) { + const int local = threadIdx.x; + if (local >= 9 * 27) { + return; + } + const int group = blockIdx.x % groups; + const int iod = blockIdx.x / groups; + const int base_iow = blockIdx.y * X_TILE; + const int ioh = blockIdx.z; + + const int local_channel = local / 27; + const int iic = group * 9 + local_channel; + if (iic >= IC) { + return; + } + const int rem0 = local - local_channel * 27; + const int ikd = rem0 / 9; + const int rem1 = rem0 - ikd * 9; + const int ikh = rem1 / 3; + const int ikw = rem1 - ikh * 3; + +#pragma unroll + for (int dx = 0; dx < X_TILE; ++dx) { + const int iow = base_iow + dx; + if (iow >= OW) { + return; + } + const int offset_dst = + iod * OH_OW_IC_KD_KH_KW + + ioh * OW_IC_KD_KH_KW + + iow * IC_KD_KH_KW + + iic * 27 + + rem0; + const int offset_src = + iic * stride_q + + (iod + ikd) * stride_z + + (ioh + ikh) * stride_y + + (iow + ikw) * stride_x; + dst[offset_dst] = T(src[offset_src]); + } +} + // [N*IC, ID, IH, IW] => [N*OD, OH, OW, IC * KD * KH * KW] template static void im2col_3d_cuda(const float * src, T* dst, int64_t N, int64_t IC, int64_t ID, int64_t IH, int64_t IW, int64_t OC, int64_t KD, int64_t KH, int64_t KW, int64_t OD, int64_t OH, int64_t OW, int64_t stride_q, int64_t stride_z, int64_t stride_y, int64_t stride_x, - int s0, int s1, int s2, int p0, int p1, int p2, int d0, int d1, int d2, cudaStream_t stream) { + int s0, int s1, int s2, int p0, int p1, int p2, int d0, int d1, int d2, int32_t lowering, cudaStream_t stream) { const int64_t OH_OW = OH*OW; const int64_t KD_KH_KW = KD*KH*KW; const int64_t ID_IH_IW = ID*IH*IW; @@ -181,6 +442,36 @@ static void im2col_3d_cuda(const float * src, T* dst, const int64_t OH_OW_IC_KD_KH_KW = OH*OW*IC*KD*KH*KW; const int64_t OW_IC_KD_KH_KW = OW*IC*KD*KH*KW; const int64_t num_blocks = (IC_KD_KH_KW + CUDA_IM2COL_BLOCK_SIZE - 1) / CUDA_IM2COL_BLOCK_SIZE; + const int64_t N_OD_OH_OW_IC_KD_KH_KW = N*OD*OH*OW*IC*KD*KH*KW; + const bool use_n1_k3_nopad_kernel = lowering == GGML_IM2COL_3D_LOWERING_CUDA_N1_K3_NOPAD_X8; + if (use_n1_k3_nopad_kernel && + N == 1 && KD == 3 && KH == 3 && KW == 3 && + s0 == 1 && s1 == 1 && s2 == 1 && + p0 == 0 && p1 == 0 && p2 == 0 && + d0 == 1 && d1 == 1 && d2 == 1) { + const bool dst_offsets_fit_i32 = N_OD_OH_OW_IC_KD_KH_KW <= INT_MAX; + const int64_t channel_groups = (IC + 8) / 9; + if (IC <= INT_MAX && OH <= INT_MAX && OW <= INT_MAX && + channel_groups <= INT_MAX && IC_KD_KH_KW <= INT_MAX && + dst_offsets_fit_i32 && + OH_OW_IC_KD_KH_KW <= INT_MAX && OW_IC_KD_KH_KW <= INT_MAX && + stride_q <= INT_MAX && stride_z <= INT_MAX && stride_y <= INT_MAX && stride_x <= INT_MAX) { + dim3 block_nums(channel_groups * OD, MIN((OW + 7) / 8, MAX_GRIDDIM_Y), MIN(OH, MAX_GRIDDIM_Z)); + im2col_3d_n1_k3_nopad_xtile_kernel<<>>( + src, dst, + static_cast(IC), static_cast(OH), static_cast(OW), + static_cast(channel_groups), static_cast(IC_KD_KH_KW), + static_cast(OH_OW_IC_KD_KH_KW), static_cast(OW_IC_KD_KH_KW), + static_cast(stride_q), static_cast(stride_z), static_cast(stride_y), static_cast(stride_x)); + return; + } + dim3 block_nums(channel_groups * OD, MIN(OW, MAX_GRIDDIM_Y), MIN(OH, MAX_GRIDDIM_Z)); + im2col_3d_n1_k3_nopad_kernel<<>>( + src, dst, IC, IH, IW, OH, OW, channel_groups, IC_KD_KH_KW, + OH_OW_IC_KD_KH_KW, OW_IC_KD_KH_KW, + stride_q, stride_z, stride_y, stride_x); + return; + } dim3 block_nums(num_blocks, MIN(OW, MAX_GRIDDIM_Y), MIN(N_OD_OH, MAX_GRIDDIM_Z)); im2col_3d_kernel<<>>(src, dst, N, IC, ID, IH, IW, OC, KD, KH, KW, OD, OH, OW, OH_OW, KD_KH_KW, ID_IH_IW, KH_KW, IH_IW, IC_ID_IH_IW, @@ -194,22 +485,22 @@ static void im2col_3d_cuda_f16(const float * src, half * dst, int64_t N, int64_t IC, int64_t ID, int64_t IH, int64_t IW, int64_t OC, int64_t KD, int64_t KH, int64_t KW, int64_t OD, int64_t OH, int64_t OW, int64_t stride_q, int64_t stride_z, int64_t stride_y, int64_t stride_x, - int s0, int s1, int s2, int p0, int p1, int p2, int d0, int d1, int d2, cudaStream_t stream) { + int s0, int s1, int s2, int p0, int p1, int p2, int d0, int d1, int d2, int32_t lowering, cudaStream_t stream) { im2col_3d_cuda(src, dst, N, IC, ID, IH, IW, OC, KD, KH, KW, OD, OH, OW, stride_q, stride_z, stride_y, stride_x, - s0, s1, s2, p0, p1, p2, d0, d1, d2, stream); + s0, s1, s2, p0, p1, p2, d0, d1, d2, lowering, stream); } static void im2col_3d_cuda_f32(const float * src, float * dst, int64_t N, int64_t IC, int64_t ID, int64_t IH, int64_t IW, int64_t OC, int64_t KD, int64_t KH, int64_t KW, int64_t OD, int64_t OH, int64_t OW, int64_t stride_q, int64_t stride_z, int64_t stride_y, int64_t stride_x, - int s0, int s1, int s2, int p0, int p1, int p2, int d0, int d1, int d2, cudaStream_t stream) { + int s0, int s1, int s2, int p0, int p1, int p2, int d0, int d1, int d2, int32_t lowering, cudaStream_t stream) { im2col_3d_cuda(src, dst, N, IC, ID, IH, IW, OC, KD, KH, KW, OD, OH, OW, stride_q, stride_z, stride_y, stride_x, - s0, s1, s2, p0, p1, p2, d0, d1, d2, stream); + s0, s1, s2, p0, p1, p2, d0, d1, d2, lowering, stream); } void ggml_cuda_op_im2col_3d(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { @@ -254,14 +545,15 @@ void ggml_cuda_op_im2col_3d(ggml_backend_cuda_context & ctx, ggml_tensor * dst) const int64_t stride_y = src1->nb[1] / es; const int64_t stride_z = src1->nb[2] / es; const int64_t stride_q = src1->nb[3] / es; + const int32_t lowering = ggml_get_op_params_i32(dst, 10); if(dst->type == GGML_TYPE_F16) { im2col_3d_cuda_f16(src1_d, (half *) dst_d, N, IC, ID, IH, IW, OC, KD, KH, KW, OD, OH, OW, stride_q, stride_z, stride_y, stride_x, - s0, s1, s2, p0, p1, p2, d0, d1, d2, stream); + s0, s1, s2, p0, p1, p2, d0, d1, d2, lowering, stream); } else { im2col_3d_cuda_f32(src1_d, (float *) dst_d, N, IC, ID, IH, IW, OC, KD, KH, KW, OD, OH, OW, stride_q, stride_z, stride_y, stride_x, - s0, s1, s2, p0, p1, p2, d0, d1, d2, stream); + s0, s1, s2, p0, p1, p2, d0, d1, d2, lowering, stream); } } diff --git a/external/ggml/src/ggml-cuda/mmq.cu b/external/ggml/src/ggml-cuda/mmq.cu index e8b265ad6..315302cb1 100644 --- a/external/ggml/src/ggml-cuda/mmq.cu +++ b/external/ggml/src/ggml-cuda/mmq.cu @@ -76,7 +76,6 @@ static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, con void ggml_cuda_mul_mat_q( ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * ids, ggml_tensor * dst) { - GGML_ASSERT( src1->type == GGML_TYPE_F32); GGML_ASSERT( dst->type == GGML_TYPE_F32); GGML_ASSERT(!ids || ids->type == GGML_TYPE_I32); // Optional, used for batched GGML_MUL_MAT_ID. @@ -84,6 +83,11 @@ void ggml_cuda_mul_mat_q( cudaStream_t stream = ctx.stream(); const int cc = ggml_cuda_info().devices[ggml_cuda_get_device()].cc; + const bool use_nvfp4_f16_activation = + blackwell_mma_available(cc) && + src0->type == GGML_TYPE_NVFP4 && + src1->type == GGML_TYPE_F16 && + ggml_get_op_params_i32(dst, 1) == GGML_MUL_MAT_LOWERING_CUDA_NVFP4_F16_ACTIVATION; const size_t ts_src0 = ggml_type_size(src0->type); const size_t ts_src1 = ggml_type_size(src1->type); @@ -93,10 +97,11 @@ void ggml_cuda_mul_mat_q( GGML_ASSERT( nb10 == ts_src1); GGML_ASSERT( nb0 == ts_dst); GGML_ASSERT(!ids || ids->nb[0] == ggml_type_size(ids->type)); + GGML_ASSERT(src1->type == GGML_TYPE_F32 || use_nvfp4_f16_activation); - const char * src0_d = (const char *) src0->data; - const float * src1_d = (const float *) src1->data; - float * dst_d = (float *) dst->data; + const char * src0_d = (const char *) src0->data; + const void * src1_d = (const void *) src1->data; + float * dst_d = (float *) dst->data; // If src0 is a temporary compute buffer, clear any potential padding. if (ggml_backend_buffer_get_usage(src0->buffer) == GGML_BACKEND_BUFFER_USAGE_COMPUTE) { @@ -133,13 +138,16 @@ void ggml_cuda_mul_mat_q( const int64_t s11 = src1->nb[1] / ts_src1; const int64_t s12 = src1->nb[2] / ts_src1; const int64_t s13 = src1->nb[3] / ts_src1; - if (use_native_fp4) { + if (use_nvfp4_f16_activation) { + quantize_mmq_nvfp4_f16_cuda((const half *) src1_d, nullptr, src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, + ne11, ne12, ne13, stream); + } else if (use_native_fp4) { static_assert(sizeof(block_fp4_mmq) == 4 * sizeof(block_q8_1)); - quantize_mmq_fp4_cuda(src1_d, nullptr, src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, + quantize_mmq_fp4_cuda((const float *) src1_d, nullptr, src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11, ne12, ne13, stream); } else { - quantize_mmq_q8_1_cuda(src1_d, nullptr, src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, + quantize_mmq_q8_1_cuda((const float *) src1_d, nullptr, src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11, ne12, ne13, stream); } CUDA_CHECK(cudaGetLastError()); @@ -196,11 +204,14 @@ void ggml_cuda_mul_mat_q( const int64_t s12 = src1->nb[2] / ts_src1; const int64_t s13 = src1->nb[3] / ts_src1; - if (use_native_fp4) { - quantize_mmq_fp4_cuda(src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, s11, s12, s13, + if (use_nvfp4_f16_activation) { + quantize_mmq_nvfp4_f16_cuda((const half *) src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, s11, s12, s13, + ne10_padded, ne11_flat, ne12_flat, ne13_flat, stream); + } else if (use_native_fp4) { + quantize_mmq_fp4_cuda((const float *) src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11_flat, ne12_flat, ne13_flat, stream); } else { - quantize_mmq_q8_1_cuda(src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, s11, s12, s13, + quantize_mmq_q8_1_cuda((const float *) src1_d, ids_src1.get(), src1_q8_1.get(), src0->type, ne10, s11, s12, s13, ne10_padded, ne11_flat, ne12_flat, ne13_flat, stream); } CUDA_CHECK(cudaGetLastError()); diff --git a/external/ggml/src/ggml-cuda/norm.cu b/external/ggml/src/ggml-cuda/norm.cu index e943cf414..e2387af59 100644 --- a/external/ggml/src/ggml-cuda/norm.cu +++ b/external/ggml/src/ggml-cuda/norm.cu @@ -150,6 +150,117 @@ static __global__ void rms_norm_f32(const float * x, } } +template +static __global__ void rms_norm_channels_f32( + const T * x, + const float * bias, + const float * gamma, + float * dst, + int64_t ne0, + int64_t ne1, + int64_t ne2, + int64_t ne3, + int64_t nb0, + int64_t nb1, + int64_t nb2, + int64_t nb3, + int64_t dnb0, + int64_t dnb1, + int64_t dnb2, + int64_t dnb3, + int64_t gnb0, + const float eps) { + const int64_t row = int64_t(blockIdx.x); + const int tid = threadIdx.x; + if (row >= ne0 * ne1 * ne2) { + return; + } + + const int64_t i0 = row % ne0; + const int64_t tmp0 = row / ne0; + const int64_t i1 = tmp0 % ne1; + const int64_t i2 = tmp0 / ne1; + const char * x_base = reinterpret_cast(x) + i0 * nb0 + i1 * nb1 + i2 * nb2; + char * dst_base = reinterpret_cast(dst) + i0 * dnb0 + i1 * dnb1 + i2 * dnb2; + const char * bias_base = reinterpret_cast(bias); + const char * gamma_base = reinterpret_cast(gamma); + + float sum = 0.0f; + for (int64_t c = tid; c < ne3; c += block_size) { + const float v = float(*reinterpret_cast(x_base + c * nb3)); + const float shifted = add_bias ? v + *reinterpret_cast(bias_base + c * gnb0) : v; + sum += shifted * shifted; + } + + extern __shared__ float s_sum[]; + sum = block_reduce(sum, s_sum); + const float scale = rsqrtf(sum / ne3 + eps); + + for (int64_t c = tid; c < ne3; c += block_size) { + const float v = float(*reinterpret_cast(x_base + c * nb3)); + const float g = *reinterpret_cast(gamma_base + c * gnb0); + const float shifted = add_bias ? v + *reinterpret_cast(bias_base + c * gnb0) : v; + const float normalized = shifted * scale * g; + if constexpr (apply_silu) { + *reinterpret_cast(dst_base + c * dnb3) = normalized / (1.0f + expf(-normalized)); + } else { + *reinterpret_cast(dst_base + c * dnb3) = normalized; + } + } +} + +template +static __global__ void rms_norm_channels_coalesced_f32( + const T * x, + const float * bias, + const float * gamma, + float * dst, + int64_t rows, + int64_t channels, + const float eps) { + const int lane = threadIdx.x & (WARP_SIZE - 1); + const int warp = threadIdx.x / WARP_SIZE; + const int64_t row = int64_t(blockIdx.x) * rows_per_block + lane; + const bool row_valid = row < rows; + + __shared__ float partial[channels_per_block][rows_per_block]; + __shared__ float scale_s[rows_per_block]; + + float sum = 0.0f; + for (int64_t c = warp; c < channels; c += channels_per_block) { + const float v = row_valid ? float(x[c * rows + row]) : 0.0f; + const float shifted = add_bias ? v + bias[c] : v; + sum += shifted * shifted; + } + partial[warp][lane] = sum; + __syncthreads(); + + if (warp == 0) { + float row_sum = 0.0f; +#pragma unroll + for (int c = 0; c < channels_per_block; ++c) { + row_sum += partial[c][lane]; + } + scale_s[lane] = rsqrtf(row_sum / channels + eps); + } + __syncthreads(); + + if (!row_valid) { + return; + } + const float scale = scale_s[lane]; + for (int64_t c = warp; c < channels; c += channels_per_block) { + const float v = float(x[c * rows + row]); + const float shifted = add_bias ? v + bias[c] : v; + const float normalized = shifted * scale * gamma[c]; + if constexpr (apply_silu) { + dst[c * rows + row] = normalized / (1.0f + expf(-normalized)); + } else { + dst[c * rows + row] = normalized; + } + } +} + template static __global__ void rms_norm_back_f32( const float * grad, const float * xf, float * dst, const int ncols, const float eps) { @@ -473,6 +584,135 @@ void ggml_cuda_op_rms_norm(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { rms_norm_f32_cuda(src0_d, dst_d, ne00, ne01, ne02, ne03, s01, s02, s03, eps, stream); } +template +static void rms_norm_channels_coalesced_launch( + const T * src0_d, + const float * bias_d, + const float * gamma_d, + float * dst_d, + int64_t rows, + int64_t channels, + float eps, + bool apply_silu, + bool add_bias, + cudaStream_t stream) { + const int64_t blocks = (rows + rows_per_block - 1) / rows_per_block; + const dim3 block_dims(rows_per_block * channels_per_block, 1, 1); + if (apply_silu && add_bias) { + rms_norm_channels_coalesced_f32<<>>( + src0_d, bias_d, gamma_d, dst_d, rows, channels, eps); + } else if (apply_silu) { + rms_norm_channels_coalesced_f32<<>>( + src0_d, nullptr, gamma_d, dst_d, rows, channels, eps); + } else { + rms_norm_channels_coalesced_f32<<>>( + src0_d, nullptr, gamma_d, dst_d, rows, channels, eps); + } +} + +template +static void rms_norm_channels_block_launch( + const T * src0_d, + const float * bias_d, + const float * gamma_d, + float * dst_d, + const ggml_tensor * src0, + const ggml_tensor * gamma, + const ggml_tensor * dst, + int64_t rows, + float eps, + bool apply_silu, + bool add_bias, + cudaStream_t stream) { + if (apply_silu && add_bias) { + rms_norm_channels_f32<<>>( + src0_d, bias_d, gamma_d, dst_d, + src0->ne[0], src0->ne[1], src0->ne[2], src0->ne[3], + src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3], + dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], + gamma->nb[0], + eps); + } else if (apply_silu) { + rms_norm_channels_f32<<>>( + src0_d, nullptr, gamma_d, dst_d, + src0->ne[0], src0->ne[1], src0->ne[2], src0->ne[3], + src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3], + dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], + gamma->nb[0], + eps); + } else { + rms_norm_channels_f32<<>>( + src0_d, nullptr, gamma_d, dst_d, + src0->ne[0], src0->ne[1], src0->ne[2], src0->ne[3], + src0->nb[0], src0->nb[1], src0->nb[2], src0->nb[3], + dst->nb[0], dst->nb[1], dst->nb[2], dst->nb[3], + gamma->nb[0], + eps); + } +} + +static void ggml_cuda_op_rms_norm_channels_impl(ggml_backend_cuda_context & ctx, ggml_tensor * dst, bool apply_silu, bool add_bias) { + const ggml_tensor * src0 = dst->src[0]; + const ggml_tensor * bias = add_bias ? dst->src[1] : nullptr; + const ggml_tensor * gamma = add_bias ? dst->src[2] : dst->src[1]; + const float * bias_d = add_bias ? (const float *) bias->data : nullptr; + const float * gamma_d = (const float *) gamma->data; + cudaStream_t stream = ctx.stream(); + + GGML_ASSERT(src0->type == GGML_TYPE_F16 || src0->type == GGML_TYPE_F32); + GGML_ASSERT(!add_bias || bias->type == GGML_TYPE_F32); + GGML_ASSERT(gamma->type == GGML_TYPE_F32); + GGML_ASSERT(dst->type == GGML_TYPE_F32); + GGML_ASSERT(!add_bias || bias->ne[0] == src0->ne[3]); + GGML_ASSERT(gamma->ne[0] == src0->ne[3]); + GGML_ASSERT(ggml_is_contiguous(dst)); + + float eps; + memcpy(&eps, dst->op_params, sizeof(float)); + GGML_ASSERT(eps >= 0.0f); + + const int64_t rows = src0->ne[0] * src0->ne[1] * src0->ne[2]; + const bool use_coalesced_kernel = + ggml_get_op_params_i32(dst, 1) == GGML_RMS_NORM_CHANNELS_LOWERING_CUDA_COALESCED; + if (use_coalesced_kernel && + ggml_is_contiguous(src0) && + gamma->nb[0] == (int64_t) sizeof(float) && + (!add_bias || bias->nb[0] == (int64_t) sizeof(float)) && + src0->nb[0] == (int64_t) ggml_type_size(src0->type) && + dst->nb[0] == (int64_t) sizeof(float) && + src0->ne[3] <= 512) { + constexpr int rows_per_block = 32; + constexpr int channels_per_block = 8; + if (src0->type == GGML_TYPE_F16) { + rms_norm_channels_coalesced_launch( + (const half *) src0->data, bias_d, gamma_d, (float *) dst->data, rows, src0->ne[3], eps, apply_silu, add_bias, stream); + } else { + rms_norm_channels_coalesced_launch( + (const float *) src0->data, bias_d, gamma_d, (float *) dst->data, rows, src0->ne[3], eps, apply_silu, add_bias, stream); + } + return; + } + if (src0->type == GGML_TYPE_F16) { + rms_norm_channels_block_launch( + (const half *) src0->data, bias_d, gamma_d, (float *) dst->data, src0, gamma, dst, rows, eps, apply_silu, add_bias, stream); + } else { + rms_norm_channels_block_launch( + (const float *) src0->data, bias_d, gamma_d, (float *) dst->data, src0, gamma, dst, rows, eps, apply_silu, add_bias, stream); + } +} + +void ggml_cuda_op_rms_norm_channels(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + ggml_cuda_op_rms_norm_channels_impl(ctx, dst, false, false); +} + +void ggml_cuda_op_rms_norm_channels_silu(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + ggml_cuda_op_rms_norm_channels_impl(ctx, dst, true, false); +} + +void ggml_cuda_op_rms_norm_channels_add_bias_silu(ggml_backend_cuda_context & ctx, ggml_tensor * dst) { + ggml_cuda_op_rms_norm_channels_impl(ctx, dst, true, true); +} + void ggml_cuda_op_rms_norm_fused(ggml_backend_cuda_context & ctx, ggml_tensor * dst, ggml_tensor * mul_tensor) { const ggml_tensor * rms_norm_src = (ggml_tensor *) dst->src[0]; float eps = 0.0f; diff --git a/external/ggml/src/ggml-cuda/norm.cuh b/external/ggml/src/ggml-cuda/norm.cuh index 120920160..2f997d05a 100644 --- a/external/ggml/src/ggml-cuda/norm.cuh +++ b/external/ggml/src/ggml-cuda/norm.cuh @@ -6,6 +6,12 @@ void ggml_cuda_op_group_norm(ggml_backend_cuda_context & ctx, ggml_tensor * dst) void ggml_cuda_op_rms_norm(ggml_backend_cuda_context & ctx, ggml_tensor * dst); +void ggml_cuda_op_rms_norm_channels(ggml_backend_cuda_context & ctx, ggml_tensor * dst); + +void ggml_cuda_op_rms_norm_channels_silu(ggml_backend_cuda_context & ctx, ggml_tensor * dst); + +void ggml_cuda_op_rms_norm_channels_add_bias_silu(ggml_backend_cuda_context & ctx, ggml_tensor * dst); + void ggml_cuda_op_rms_norm_fused(ggml_backend_cuda_context & ctx, ggml_tensor * dst, ggml_tensor * mul_tensor); void ggml_cuda_op_rms_norm_fused_add(ggml_backend_cuda_context & ctx, diff --git a/external/ggml/src/ggml-cuda/quantize.cu b/external/ggml/src/ggml-cuda/quantize.cu index 8d824264c..b4efb1add 100644 --- a/external/ggml/src/ggml-cuda/quantize.cu +++ b/external/ggml/src/ggml-cuda/quantize.cu @@ -71,8 +71,9 @@ __device__ __forceinline__ uint8_t compute_e8m0_scale(float amax) { } +template static __global__ void quantize_mmq_nvfp4( - const float * __restrict__ x, const int32_t * __restrict__ ids, void * __restrict__ vy, + const T * __restrict__ x, const int32_t * __restrict__ ids, void * __restrict__ vy, const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03, const int64_t ne0, const int64_t ne1, const int64_t ne2) { #if defined(BLACKWELL_MMA_AVAILABLE) @@ -105,7 +106,7 @@ static __global__ void quantize_mmq_nvfp4( for (int k = 0; k < QK_NVFP4_SUB; k++) { const int64_t i00 = i0_base + k; if (i00 < ne00) { - const float v = x[base_idx + i00]; + const float v = float(x[base_idx + i00]); vals_raw[k] = v; amax_raw = fmaxf(amax_raw, fabsf(v)); } else { @@ -219,7 +220,7 @@ static __global__ void quantize_mmq_mxfp4(const float * __restrict__ x, #pragma unroll for (int b = 0; b < 2; ++b) { const int64_t i0 = warp_start_offset + b * vals_per_scale + lane_id_32; - const float xi = (i0 < ne00) ? x[base_pos + i0] : 0.0f; + const float xi = (i0 < ne00) ? float(x[base_pos + i0]) : 0.0f; float amax = fabsf(xi); #pragma unroll @@ -425,7 +426,7 @@ void quantize_mmq_fp4_cuda( const int64_t block_num_y = (ne0 + QK_NVFP4_SUB * nvfp4_block_size - 1) / (QK_NVFP4_SUB * nvfp4_block_size); const dim3 block_size(nvfp4_block_size, 1, 1); const dim3 num_blocks(ne1, block_num_y, ne2 * ne3); - quantize_mmq_nvfp4<<>>( + quantize_mmq_nvfp4<<>>( x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); } else { GGML_ASSERT(ne0 % (2 * QK_MXFP4) == 0); @@ -441,3 +442,19 @@ void quantize_mmq_fp4_cuda( quantize_mmq_mxfp4<<>>(x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); } } + +void quantize_mmq_nvfp4_f16_cuda( + const half * x, const int32_t * ids, void * vy, const ggml_type type_src0, + const int64_t ne00, const int64_t s01, const int64_t s02, const int64_t s03, + const int64_t ne0, const int64_t ne1, const int64_t ne2, const int64_t ne3, cudaStream_t stream) { + GGML_ASSERT(type_src0 == GGML_TYPE_NVFP4); + GGML_ASSERT(ne0 > 0); + GGML_ASSERT(ne00 % QK_NVFP4 == 0); + + constexpr int nvfp4_block_size = 128; + const int64_t block_num_y = (ne0 + QK_NVFP4_SUB * nvfp4_block_size - 1) / (QK_NVFP4_SUB * nvfp4_block_size); + const dim3 block_size(nvfp4_block_size, 1, 1); + const dim3 num_blocks(ne1, block_num_y, ne2 * ne3); + quantize_mmq_nvfp4<<>>( + x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2); +} diff --git a/external/ggml/src/ggml-cuda/quantize.cuh b/external/ggml/src/ggml-cuda/quantize.cuh index 2d684d385..f56645a74 100644 --- a/external/ggml/src/ggml-cuda/quantize.cuh +++ b/external/ggml/src/ggml-cuda/quantize.cuh @@ -39,3 +39,8 @@ void quantize_mmq_fp4_cuda(const float * x, int64_t ne2, int64_t ne3, cudaStream_t stream); + +void quantize_mmq_nvfp4_f16_cuda( + const half * x, const int32_t * ids, void * vy, ggml_type type_src0, + int64_t ne00, int64_t s01, int64_t s02, int64_t s03, + int64_t ne0, int64_t ne1, int64_t ne2, int64_t ne3, cudaStream_t stream);