When computing sinks, the cm1 shader was looping r from 0 to Br rather than
to rows_per_thread. I must have copied this from the scalar path (where it is
correct), and somehow it wasn't causing failures on current drivers.
* First attempt
* No permute during convert (fixes qk tensors), proper norm application.
* RoPE = NeoX
* Coherence!
* Migrate xielu params from tensors to hyperparameters
* Simple CUDA kernel
* Revert stupid LLM refactorings
* Chat template support
* configchecker / flake8 errors
* Reorder unary.cu
* I do conclude that LLMs are, in fact, stupid.
* Fix after merge
* Final newline
* Make xIELU an UNARY_OP
* Final newline
* Correctly account for parameter shift
* Argh.
* Update ggml/src/ggml-cpu/unary-ops.cpp
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* Refactor: remove unused methods, inline and factorize softplus, add const modifiers
* Revert CUDA changes, implement xIELU as a separate OP
* Pesky newline
* Add float2half / half2float for F16 inputs/outputs
* CUDA variants, attempt 2
* Actually, attempt 3
* Update ggml/src/ggml-cuda/unary.cu
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* Missing convert header
* Proper formula and reference for xIELU in the comments.
* Modify unary-ops.cpp to add the functor-based logic besides the template system to retain optimizations
* Apply suggestions from code review
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
* Add tensor mappings for Apertus to global list instead
* Fix lazy on scalars
* Update ggml/src/ggml-cuda/unary.cu
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* Add comment about the constraints on positive/negative alpha
* Change `softplus` to `ggml_softplus`
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
Co-authored-by: Sigbjørn Skjæret <sigbjorn.skjaeret@scala.com>
* HIP: Disable ROCWMMA fatt on CDNA when compiled against ROCWMMA 2.0.0
rocwmma 2.0.0 includes a bug in the code fakeing fp16 accumulation on CDNA
* CUDA: Fix volta condition in ggml_cuda_should_use_wmma_fattn
* Work on rope
* Simplify inplace operation generation and combine mul/add generation
* Work on rope variants
* implement neox rope
* rope complete
* Add sub,div,glu operators
* implement scale op
* Update cpy shader to handle cont/more types
* formatting
* Update test vars printing for rope,rms_norm
* Avoid ROPE hardcoded constants
* Add TODO to change ROPE constants to enum
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* fix TODO comment
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* Add support for --carry-initial-prompt
* PR fixes for ruby and go
* Refactoring for readability
* WIP 1
* WIP 2
* PR fixes
* More PR fixes
* PR fix
* Further simplification
* d'oh
* One more logic fix
* Update src/whisper.cpp
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* Truncate prompt_past0 upon initialization
* Slight simplification
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* cli: Fix assignment for vad_min_silence_duration_ms
Found and fixed this simple copy/paste error
* server : fix vad_min_silence_duration_ms assignment
---------
Co-authored-by: Daniel Bevenius <daniel.bevenius@gmail.com>
This commit fixes multiple issues:
* memory leak because vad_segments is never released
* avoid segmentation fault when whisper_vad_segments_from_samples returns nullptr.
* avoid potential segmentation fault when the app fails to allocate memory for filtered samples and the vad context is released but also get released withing state itself when whisper_free_state is called
* vad : fix memory leak by storing ggml_context in vad context struct
This commit addresses a memory leak issue in the voice activity
detection (VAD) where the ggml_context is not stored within the vad
context structure.
The motivation for this change that this is causing the context memory
to stay allocated and the tensor still point to that memory but this
memory is never freed.
* vad : free memory allocated for VAD hparams
This commit frees the model hyperparameters allocated for the VAD
context in the `whisper_vad_free` function. Specifically, it deletes the
`encoder_in_channels`, `encoder_out_channels`, and `kernel_sizes` arrays
allocated with `new[]` in the `whisper_vad_init` function.
The motivation for this is to prevent memory leaks when the VAD.
* vad: free ggml buffer in whisper_vad_free
This commit frees the ggml buffer in the whisper_vad_free function to
prevent memory leaks.
Resolves: https://github.com/ggml-org/whisper.cpp/issues/3452
* Revert "vad : fix memory leak by storing ggml_context in vad context struct"
This reverts commit aeafca437e.
* whisper : free ggml context in whisper_vad_init_context
This commit frees the ggml_context after initializing the VAD context in
the whisper_vad_init_context function.
The motivation for this is to prevent memory leaks.
This commit disables flash-attention for the Java binding test so that
the testFullTranscribe test passes.
Without this change the test was failing because the expected output
mismatches after the flash-attention change:
```console
<And so my fellow Americans ask not what your country can do for you ask what you can do for your country.>
but was:
<and so my fellow Americans ask not what your country can do for you ask what you can do for your country>
```
An alternative would also be to update the expected output but it felt
better to keep the same expected output and disable flash-attention and
not just change the expected output to match the new behavior.
This PR adds additional information to an error message when loading backend library via ld_load_library() fails. This helps spotting why backend library did not load (missing library, missing dependency or unresolved symbol etc.).
* vulkan: 64-bit im2col
Add variants of the im2col shaders that use buffer_device_address/buffer_reference,
and use 64-bit address calculations. This is needed for large convolutions used in
stable-diffusion.cpp.
* fix validation error for large im2col
* metal : support mul_mm with src1->type == GGML_TYPE_F16
* metal : support mul_mm_id with src1->type == GGML_TYPE_F16
[no ci]
* metal : mul_mm support ne00 % 32 != 0
* metal : support mul_mm_id with ne00 % 32 != 0
* cont : remove unnecessary unrolls
* cont : simplify data loading
* metal : optimize mul_mm when output bounds checks are not needed
Jeff Bolz
Reese Levine
Piotr Wilkin (ilintar)
R0CKSTAR
uvos
Eve
lhez
Ruben Ortlam
Andreas Lubbe
Georgi Gerganov
Silviu Caragea
KITAITI Makoto
Daniel Bevenius
anavp-nvidia
Charles Xu
alex-spacemit
Rafal Lewczuk
Sigbjørn Skjæret