QVAC-21582 — Pull latest from upstream whisper.cpp (v1.9.1)#73
Open
freddy311082 wants to merge 290 commits into
Open
QVAC-21582 — Pull latest from upstream whisper.cpp (v1.9.1)#73freddy311082 wants to merge 290 commits into
freddy311082 wants to merge 290 commits into
Conversation
This reverts commit a4d9176.
This reverts commit be867ea. Sorry about this noise, I thought it was worth a try.
This commit adds explicit casts to float for -INFINITY. The motivation for this is that in CUDA 11.8.0, the -INFINITY macro is defined as a double (a header provided NVCC). This triggers a warning and hence causes a CI failure in whisper.cpp. I belive that this header might have been updated in CUDA 12 which is why we don't see this warning. Refs: https://github.com/ggml-org/whisper.cpp/actions/runs/25713948217/job/75500081939?pr=3803 Refs: ggml-org/llama.cpp#22824
…f/bfloat16 in CUDA 11.8" This reverts commit 5cd2284. Reverting in favor of: ggml-org/llama.cpp#22994
…f2 types" This reverts commit a2839b4. Reverting this as after closer inspection these only warnings and not errors.
…/1477) For a given output position j on the time axis, only input positions i such that i*s0 <= j < i*s0 + K contribute -- i.e. i in [ceil((j - K + 1)/s0), floor(j/s0)] intersected with [0, IL-1]. That's at most ceil(K/s0) values (typically 2 for stride==K/2 transposed convs). The current kernel iterates the full IL range and filters with an `if`, amplifying per-thread work by IL/ceil(K/s0) (~160x for IL=320, K=10, s0=5 -- a representative codec-decoder shape). On Apple M1 the wasted work trips the macOS GPU watchdog (kIOGPUCommandBufferCallbackErrorImpactingInteractivity) on long graphs. Compute i_min, i_max analytically before the inner loop and iterate only [i_min, i_max]. Output is bit-identical (same multiplies and adds in the same order); loop bound shrinks by IL/ceil(K/s0). Tested on M1 with a downstream consumer running a TTS codec at full T_codec; end-to-end codec decode ~3-4x faster, zero watchdog hits across long synthesis runs vs ~30% pre-patch.
Add missing `#include <mutex>` in ggml-backend-device.cpp. Fixes: #22809 Signed-off-by: Oliver Walsh <owalsh@redhat.com>
* add im2col_3d * format code * update the ops.md
Before, we relied on a transient import from `cub/cub.cuh`, which is bad practice to do as cub may not always expose cuda/iterator
* cuda: tighten snake fusion type checks for all operands (defensive, sync vulkan) * cuda: reject snake fusion when ne[2] or ne[3] > 1 (mirror vulkan PR review) * cuda: merge type_ok and types_ok into a single types_ok (address am17an review) * cuda: filter ADD/SUB/MUL/DIV in supports_op to F32/F16 bin_bcast only dispatches F32/F16 type triplets, mirror the vulkan filter so unsupported types fall back through cpy instead of aborting. * test-backend-ops: extend snake_fuse to rank-4 with ne[2]/ne[3] > 1 cases
`im2col_cuda` and `im2col_3d_cuda` both dispatch with `block_nums.y = OW`. CUDA caps grid Y at 65535. Conv1d encoders on raw 16 kHz audio with T > 65535 (~ 4 s) trip the limit -- e.g. SEANet at 11 s lands at OW = 176000 -- and the launch returns `invalid configuration argument`. Clamp `block_nums.y` to `MIN(OW, MAX_GRIDDIM_Y)` and loop inside the kernel with stride `MAX_GRIDDIM_Y`. Same in-kernel stride pattern already used for the z axis (`MAX_GRIDDIM_Z`). Both 2D `im2col_kernel` and 3D `im2col_3d_kernel` need the same fix. Bit-identical for OW <= 65535 (single iteration of the new outer loop). Tested on T4 / Jetson Orin with a SEANet encoder running on 11 s / 16 kHz audio (im2col reaching OW ~ 176000); pre-fix launch returns `invalid configuration argument`, post-fix runs to completion. Existing test-backend-ops im2col cases unchanged.
* Q4_1 MoE CLC pass sanity check * remove unnecessary code * opencl: remove unnecessary asserts and reformat * opencl: fix supports_op for q4_1 moe * q4_1 moe is supported by Adreno with certain shapes --------- Co-authored-by: Li He <lih@qti.qualcomm.com>
…lama/22711) * metal : promote mul_mv/mul_mm batch divisors to function constants * metal : take op directly in get_pipeline_mul_mv_ext
…s for Xe2 and newer (llama/22461) * refactor * Use l_warptile only when coopamt is available for BF16
* fix(mixed-types): use f32 for precision and update the shared memory calculation logic for f32 * fix(unary): correct the gelu, gelu quick and gelu erf functions * fix(flash-attn-tile): fix the hardcode v type * fix(flash_attn): fix tile path * fix: pass editorconfig and address the type conflicts * fix: remove reduant pipeline keys * fix: remove inline min/max group size functions and revert the flash attn path order * fix: use clamp to avoid NaN for GELU * fix: use the right range for exp, 80 is safer for f32 exp
* Enable to run gpt-oss-20b and refactor mulmat-q * disable test-backend-ops in ubuntu-24-webgpu
* ggml-opencl: add Adreno xmem F16xF32 GEMM for prefill * ggml-opencl: address Adreno xmem review comments * ggml-opencl: align xmem gemm kernel naming --------- Co-authored-by: Your Name <your@email.com>
* hexagon: add hvx_vec_repl helpers and use those for splat-from-vtcm usecase * hmx-mm: optimize per-group scale handling * hmx-fa: optimize slope load from vtcm * hmx-fa: use aligned access where possible in hmx-utils * hexagon: add hvx_vec_repl_2x_f16 helper and consolidate repl helpers --------- Co-authored-by: Max Krasnyansky <maxk@qti.qualcomm.com>
…and Flash Attention (llama/24123) * vulkan: add support for valve fp16 dot2 extension * use macro for dot2 path choice * properly check for the feature * add dot_product abstraction to reduce preprocessor branching
* Add missing syncthreads before resuing cub_temp_storage __syncthreads() is required before being allowed to resue TempStorage smem: https://nvidia.github.io/cccl/unstable/cub/api/classcub_1_1BlockLoad.html#_CPPv4I0EN3cub9BlockLoad4LoadEv20RandomAccessIteratorRA14ItemsPerThread_1Ti * Add one more missing __syncthreads Could also double-buffer, but alternative is to simply ensure all threads have read smem* before writing to it again in the next loop iteration * Remove unused smem from ssm_scan_f32
* Make ggml_gated_delta_net take only the initial recurrent state (D, 1, n_seqs) and passes the snapshot count K as an op parameter instead of inferring it from state->ne[1]. Remove the padding hack and copy all emitted snapshots into the recurrent cache with a single strided ggml_cpy * Make GDN changes in all backends. Address review comments. * Fix CI build errors
* vulkan: use medium matmul tile on Asahi Linux * vulkan: switch Apple detection to Honeykrisp driver id
Fixes build/CI after #24306.
* opencl: add q5_0 adreno support * opencl: add q5_1 adreno support * opencl: cosmetic fix --------- Co-authored-by: Li He <lih@qti.qualcomm.com>
* cuda: support concat for scalar types * Update concat.cu * fix metal ci issue
Adds a `--version` option to whisper-cli that prints the library version via `whisper_version()` and exits, plus a corresponding entry in the help output. Mirrors the existing `-h`/`--help` handling. Closes ggml-org#608
* ci : only trigger release jobs for tags This commit removes the building of the release jobs on pushed to master. The motivation for this is that it can be confusing at the momement when releasing that the push to master also triggers the release jobs but the actual release will be skipped. With this change the release job is only run when a tag is pushed which should result in a single Release github actions job and make it easier to follow. * ci : add GGML_NATIVE=OFF for ubuntu-22-gcc
* parakeet : add support for NVIDIA Parakeet Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* Add Whisper::Parakeet::Params * Add tests for Parakeet::Params * Remove unused variabel * Add callbacks to Parakeet::Params * Group callback and user_data params * Undefine local macros * Define GetParakeetParams * Remove unused variable * Use ITERATE_CALLBACK_PARAMS * Use ITERATE_CALLBACK_PARAMS instead of ITERATE_USER_DATA_PARAMS * Fix memsize * Remove unnecessary macros * Simplify params registration * Define Parakeet * Add hook methods to Parakeet::Params * Fix typo * Check callback container in GetParakeetParams * Reduce if * Free parakeet_full_params * Implement Parakeet::Context#initialize * Add TestParakeetContext * Add Parakeet::Segment * Prevent double-free * Add Parakeet::Context#transcribe * Add Parakeet::Context#each_segment * Define Parakeet::Segment attributes * Define Parakeet::Segment#deconstruct_keys * Add tests for Parakeet::Segment#deconstruct_keys * Run Parakeet::Context#transcribe without GVL * Make it to abort for Parakeet * Add Parakeet.log_set * Define Parakeet::Token * Define Parakeet::Segment#each_token * Implement some hooks of Parakeet::Params * Convert int to VALUE * Implement hooks for Parakeet * Implement Parakeet::Context#full * Add tests for Parakeet::Context#full * Add Parakeet to RBS * Fix ruby_whisper_parakeet_params_free * Free ruby_whisper_parakeet_context * Add tests for hooks * Add Parakeet section to README * Add more attributes of Parakeet::Context * Add tests for Parakeet::Context's attributes * Update RBS * Register parakeet-tdt-0.6b-v3 * Narrow scope of log constants * Extract activate and deactivate of log_queue * Make start_log_callback_thread private * Don't call start_log_callback_thread unncecessarilly * Early return from log_queue_enqueue when not active * Gropu log_queue members * is_active -> is_open * Fix English * Share parakeet full body function * ruby_whisper_parakeet_abort_callback_user_data -> ruby_whisper_abort_callback_user_data * NULL check for callback containers * Fix Parakeet.log_set * Omit Parakeet tests on CI * Extract Whisper::LogSettable * Join log callback thread in a log queue function * Revert Join log callback thread in a log queue function * Extract output methods to modules * Move Parakeet init functions into init_parakeet() * Add output methods to Parakeet classes * Add Parakeet's output methods to RBS * Use Whisper::Output in RBS * Add LogSettable to RBS * Fix module Token -> class Token * Add Parakeet::Model * Add test for Parakeet::Model * Add Parakeet::Model to RBS * Move position of Parakeet::Model in RBS * Parakeet -> TestBase::Parakeet * Add Parakeet::Context#model in RBS * Add Whisper::Output * Fix nil check * Define ruby_whisper_parakeet_model_memsize * Fix order of declaration in ruby_whisper_parakeet_model_get_xxx * Define Parakeet.system_info_str * Add test for Parakeet.system_info_str * Add signature of Parakeet.system_info_str * Define Parakeet::VERSION * Add test for Parakeet::VERSION * Add signature of Parakeet::VERSION * Add Parakeet::Context::Params * Make Parakeet::Context.new accept Context::Params * Add test for Parakeet::Context.new with Context::Params * Update RBS * Remove params from Parakeet::Params which are moved from whisper_parakeet_full_params * Remove tests for removed params * Make Parakeet tests follow original behavior changes * Add Parakeet model shortcuts * Alloc token data in factory instead of alloc func * Fix variable name * Update RBS * Refactor log settable module * Use log settable for Whisper * Address deadlock * Make test follow change of log queue implementation * Refactor to make abort callback use the same way to parakeet's way * Remove redundant structs * Fix test name * Fix README * Add missing parallel transcription * Fix test for parakeet info * Remove removed params * Wait for logs dequeued * Fix instance variable name * Load etc feature * Remove unnecessary comment * Remove unnecessary thread safety check * Remove outdated comment * Skip downloading model if cache exists * Change Hugging Face URI for Parakeet models * Bump required Ruby version to 3.3 * Fix English
…3891) * ci : add GGML_NATIVE=OFF and build all cpu-variants This commit adds -DGGML_BACKEND_DL=ON, -DGGML_NATIVE=OFF, and -DGGML_CPU_ALL_VARIANTS=ON to the releases. The motivation for this is that currently the Windows BLAS build uses the native CPU instructions and if target systems do not support these instructions, the build will fail like the linked issue reports. Resolves: ggml-org#3889 * ci : update ubuntu-cpu release job for all variants [no ci] This commit enables the ubuntu-cpu job to include all cpu variants and ensures that the ggml backend libraries are built into the bin directory similar to how llama.cpp does it. The following is a build on my fork with this change: https://github.com/danbev/whisper.cpp/releases/tag/untagged-fc3c71f0bf0f7bf19d19
…am v1.9.1 Merge upstream tag v1.9.1 (ggml-org/whisper.cpp) into the Tether fork, bumping the version string from 1.8.5 to 1.9.1 (ggml 0.15.1). Conflict resolution: - talk-llama, bindings/ruby, bindings/javascript, examples/server, examples/common-whisper, cmake/whisper.pc.in, scripts/sync-ggml.last, and 49 non-Tether ggml files: take upstream. - .github/workflows/build.yml: keep ours (Tether CI). - CMakeLists.txt / src/CMakeLists.txt: keep the `whisper::whisper` EXPORT (whisper-targets) packaging contract the vcpkg consumer needs, while taking upstream's version string and parakeet target props. - ggml/CMakeLists.txt: take upstream ggml version (0.15.1). - ggml-vulkan.cpp: take upstream (bf16 KV flash-attn, multithreaded shader compile refactor, Honeykrisp, ssm-scan K) and keep the Tether Adreno/Qualcomm matmul _l->_m fallback. Preserved Tether work: Adreno/Qualcomm Vulkan fixes (QVAC-19213), Android per-arch dlopen fallback (QVAC-18993), BCI windowed attention. Gained from upstream: exception-safe whisper_init (no throw across the C ABI), max_tokens timestamp fix (PR ggml-org#3798), and native NVIDIA Parakeet support (PR ggml-org#3735, built but not yet wired into the package). Local validation (macOS/Metal): full build (89 targets) clean; 13/14 ctest pass incl. all test-whisper-cli-*, test-vad, test-vad-streaming, test-parakeet. test-vad-full fails only due to a missing real ggml-base.en.bin model (pre-existing env issue, unrelated to the merge). Co-authored-by: Cursor <cursoragent@cursor.com>
Review StatusCurrent Status: ❌ PENDING Pending reviews: Needs 1 Management or Team Lead, and 1 more from Management, Team Lead, or Member. |
![github-advanced-security[bot]](https://avatars.githubusercontent.com/in/57789?s=60&v=4){kind=small}
There was a problem hiding this comment.
CodeQL found more than 20 potential problems in the proposed changes. Check the Files changed tab for more details.
…m merge The v1.9.1 merge pulled in upstream's GitHub Actions workflows (build-android/clang/coreml/cpu/freebsd/gcc/macos/quantize/sanitize/ self-hosted/sycl/vad/wasm/windows, release.yml, ccache-clear action) and modified bindings-*/docker/examples + renamed examples-wasm.yml. None of these are used by the Tether fork. Restore .github to exactly match master so the PR introduces no CI-workflow changes. Co-authored-by: Cursor <cursoragent@cursor.com>
| GGML_ASSERT(n_tokens*n_embd <= (int64_t) embd_pre_norm.size); | ||
| ggml_backend_tensor_get_async(backend_h, t_h_pre_norm, embd_pre_norm.data, 0, n_tokens*n_embd*sizeof(float)); | ||
| const uint32_t n_embd = hparams.n_embd_out(); | ||
| GGML_ASSERT(n_tokens*n_embd <= (int64_t) embd_nextn.size); |
| ggml_backend_tensor_get_async(backend_h, t_h_pre_norm, embd_pre_norm.data, 0, n_tokens*n_embd*sizeof(float)); | ||
| const uint32_t n_embd = hparams.n_embd_out(); | ||
| GGML_ASSERT(n_tokens*n_embd <= (int64_t) embd_nextn.size); | ||
| ggml_backend_tensor_get_async(backend_h, t_h_nextn, embd_nextn.data, 0, n_tokens*n_embd*sizeof(float)); |
| if (it == ctx_map.end()) { | ||
| ggml_init_params params = { | ||
| /*.mem_size =*/ size_t(2u*(1 + n_stream)*n_layer_kv*ggml_tensor_overhead()), | ||
| /*.mem_size =*/ size_t(2u*(1 + n_stream)*n_layer*ggml_tensor_overhead()), |
| layer.indexer_k_norm_b = create_tensor(tn(LLM_TENSOR_INDEXER_K_NORM, "bias", i), {hparams.indexer_head_size}, flags); | ||
| layer.indexer_proj = create_tensor(tn(LLM_TENSOR_INDEXER_PROJ, "weight", i), {n_embd, hparams.indexer_n_head}, flags); | ||
| layer.indexer_attn_k = create_tensor(tn(LLM_TENSOR_INDEXER_ATTN_K, "weight", i), {n_embd, hparams.indexer_head_size}, flags); | ||
| layer.indexer_attn_q_b = create_tensor(tn(LLM_TENSOR_INDEXER_ATTN_Q_B, "weight", i), {q_lora_rank, hparams.indexer_n_head * hparams.indexer_head_size}, flags); |
| break; | ||
| default: | ||
| GGML_ABORT("Unsupported K type for flash attention shader"); | ||
| } |
| GGML_ABORT("Unsupported K type for flash attention shader"); | ||
| } | ||
| variant += std::string("_k") + ggml_type_name(key.k_type); | ||
|
|
| variant += std::string("_k") + ggml_type_name(key.k_type); | ||
|
|
||
| switch (key.v_type) { | ||
| case GGML_TYPE_F32: |
| switch (key.v_type) { | ||
| case GGML_TYPE_F32: | ||
| defines.push_back("V_F32"); | ||
| break; |
| defines.push_back("V_F32"); | ||
| break; | ||
| case GGML_TYPE_F16: | ||
| defines.push_back("V_F16"); |
This file contains hidden or bidirectional Unicode text that may be interpreted or compiled differently than what appears below. To review, open the file in an editor that reveals hidden Unicode characters.
Learn more about bidirectional Unicode characters
Sign up for free
to join this conversation on GitHub.
Already have an account?
Sign in to comment
20 participants
Add this suggestion to a batch that can be applied as a single commit.This suggestion is invalid because no changes were made to the code.Suggestions cannot be applied while the pull request is closed.Suggestions cannot be applied while viewing a subset of changes.Only one suggestion per line can be applied in a batch.Add this suggestion to a batch that can be applied as a single commit.Applying suggestions on deleted lines is not supported.You must change the existing code in this line in order to create a valid suggestion.Outdated suggestions cannot be applied.This suggestion has been applied or marked resolved.Suggestions cannot be applied from pending reviews.Suggestions cannot be applied on multi-line comments.Suggestions cannot be applied while the pull request is queued to merge.Suggestion cannot be applied right now. Please check back later.
What
Merges upstream tag v1.9.1 from
ggml-org/whisper.cppinto the Tether fork,bumping the version string
1.8.5 → 1.9.1(bundled ggml0.15.1), whilepreserving all Tether-specific work.
Why
Periodic upstream sync (ticket QVAC-21582): pull the latest whisper.cpp, keep our
fork current, and pick up upstream fixes/features relevant to our packages.
Conflict resolution
talk-llama,bindings/ruby,bindings/javascript,examples/server,examples/common-whisper,cmake/whisper.pc.in,scripts/sync-ggml.last, and the ~49 non-Tetherggml/*files.
.github/workflows/build.yml(Tether CI).CMakeLists.txt/src/CMakeLists.txt: keep thewhisper::whisperEXPORT(whisper-targets) packaging contract our vcpkgconsumers rely on, while adopting upstream's version props and the new
parakeettarget.ggml/CMakeLists.txt: take upstream ggml version (0.15.1).ggml/src/ggml-vulkan/ggml-vulkan.cpp: take upstream (bf16 KV flash-attn,multithreaded shader-compile refactor, Honeykrisp driver handling, ssm-scan K)
and keep the Tether Adreno/Qualcomm matmul
_l → _mfallback.Tether work preserved
dlopenfallback (QVAC-18993)Gained from upstream
whisper_init(no throw across the C ABI)max_tokenstimestamp fix (PR Fix #984: About max_tokens skipping remaining audio issue ggml-org/whisper.cpp#3798)Testing
(incl.
test-whisper-cli-*,test-vad,test-vad-streaming,test-parakeet)after downloading the real
ggml-base.en.binmodel required bytest-vad-full.tetherto/qvac): consumed this commit via a vcpkg registrybranch and ran the addon pipelines against
whisper-cpp 1.9.1:transcription-whispercpp: prebuilds on all 10 platforms + desktop & mobile(Android/iOS) integration tests green.
bci-whispercpp: validated against 1.9.1 as well.(Only the
verified-label merge-gate is red; no build/test failures.)