* rpc : report actual free memory
Start reporting the free memory on every device instead of using
fixed values. Now llama-cli users can get a nice memory breakdown
when using RPC devices.
* drop --mem in rpc-server
* vulkan: implement SSM scan operation
Add State Space Model scan operation to the Vulkan backend.
Signed-off-by: Giuseppe Scrivano <gscrivan@redhat.com>
* vulkan: implement SSM conv operation
Add State Space Model conv operation to the Vulkan backend.
Signed-off-by: Giuseppe Scrivano <gscrivan@redhat.com>
---------
Signed-off-by: Giuseppe Scrivano <gscrivan@redhat.com>
Fix incorrect task-to-batch index calculation in the quantization phase.
The bug caused out-of-bounds access to qnbitgemm_args array when
compute_idx exceeded per_gemm_block_count_m, leading to invalid
pointer dereferences and SIGBUS errors.
Correctly map tasks to batches by dividing compute_idx by
per_gemm_block_count_m instead of block_size_m.
Example:
batch_feature=1, gemm_m=30, block_size_m=4
per_gemm_block_count_m = 8, task_count = 8
Old: gemm_idx = 4/4 = 1 (out of bounds New: gemm_idx = 4/8 = 0 (correct)
Tested on SpaceMit K1 RISC-V64 with qwen2.5:0.5b model.
Co-authored-by: muggle <mingjun.rong@spacemit.com>
* Add power management utilities to NPU device context and update DCVS settings
* Update DCVS settings in power_utils to use v3 API and enhance power management
* wip
* Enhance dequantization functions by adding load_dequant_table support and updating signatures for improved performance
* use lut
* wip
* fix test failure
* wip
* Refactor load_qual_block_generic to improve block handling and optimize vector operations
* Enhance load_dual_block_generic and load_qual_block_generic to accept a mask parameter for improved block handling
* Refactor flash_attn_impl to optimize mask l2 prefetch
* wip
* wip
* wip
* wip
* add log
* link against shared libraries instead of static ones
* fix swiglu
* wip
* refactor expf_fix to handle overflow for different data types
* enhance is_glu_op_supported to validate shapes for multiple sources
* wip
* refactor logging macros to use hexagon namespace and improve formatting
* fix printf format error
* wip
* refactor: update static_assert messages for block size validation and add HVX_VectorPred_x3 type alias
* rename
* feat: enhance fa with mask
* wip
* wip
* refactor: replace instances of Q6_V_vzero() with kZeroV for consistency
* wip
* wip
* wip
* fix: improve address alignment check in HVX_Vector handling
* refactor: streamline vector dot product implementations for improved readability
* refactor: q4k add hvx intrinsic impl
* refactor: enhance dequantize_row_q4_K for clarity and performance
* refactor: optimize scale mask usage in dequantization functions for improved performance
* refactor: optimize dequantize_row_q4_K for intrinsic usage and performance improvements
* refactor: move GLU operation implementation into separated file
* sync after swiglu
* wip
* wip
* wip
* feat: increase prc main thread stack size
* fix: replace hardcoded stack size with NPU_THREAD_STACK_SIZE constant
* wip
* feat: add optimized vector operations for exponential and division with overflow handling
* wip
* feat: refactor exponential function to handle overflow and underflow with improved logic
* wip
* wip
* feat: add vector loading and scaling functions for improved performance in block processing
* wip
* feat: optimize block loading by refactoring scale index handling for improved performance
* use Q6_Vb_vlut32_VbVbR_nomatch instead
* feat: enhance scale loading by adding static assertion and restructuring block handling
* wip
* feat: refactor vec_dot_product_mixed_impl for improved clarity and performance
* wip
* feat: simplify vector loading functions and improve alignment handling
* wip
* feat: enhance scale loading mask with quantization block size validation
* wip
* feat: implement make_scale_load_mask function and refactor vector handling in vec_ops
* feat: enhance load_dual_block_generic to include scale indices for improved vector loading
* revert q8 dequant
* wip
* feat: optimize dequantization functions by removing unnecessary masking and updating lookup methods
* wip
* wip
* add qurt_mutex
* Add DMA transfer class and integrate into thread pool
* Enhance DMA transfer functionality by adding support for multiple descriptors and initiating transfers in parallel
* fix dma crash
* fix failed unit tests
* wip
* use alignas
* Improve DMA transfer error handling and update descriptor completion check
* Fix VTCM cache size calculation in element-wise operations
* Add cache clean operations before DMA transfers in element-wise operations
* reduce cache clean operations
* Refactor DMA transfer functions to support 1D operations and rename for clarity
* Enhance DMA transfer functionality by adding 2D submission support and improving descriptor initialization
* Update read buffer method to support forced invalidation and remove unnecessary invalidation calls in element-wise operations
* wip
* Improve DMA transfer handling in mul_mat_gemv_impl by replacing memcpy with initiate_dma_row_transfer and adding wait_for_dma logic
* fix 2d dma
* feat: add DMA plane cache
* rename
* wip
* use memcpy for debug
* fix cache plane calc
* refactor: remove debug logging from mul_mat_impl and optimize cache handling
* rename
* fix 2d dma type
* refactor: enhance DMA transfer handling in mul_mat_gemv_impl and wait functions
* refactor: optimize DMA transfer handling in mul_mat_gemv_impl and wait functions
* wip
* wip
* move op impl into sub dir
* add log
* fix: correct pointer usage in mul_mat_gemv_impl for next plane access
* fix: improve DMA transfer error handling in mul_mat_impl and mul_mat_gemv_impl
* fix: fix crash by using the entire row bytes
* wip
* wip
* fix: prevent parallelization for scalar src1 in is_mul_mat_supported
* fix: add dimension checks for 2D DMA transfers and fallback to 1D if necessary
* wip
* fix: enable thread barrier for mul multiplication operations
* feat: add synchronization checks for tensor operations and update related functions
* wip
* fix: remove invalidation flag from get_read_buffer calls in element-wise and matrix multiplication operations
* Revert "fix: remove invalidation flag from get_read_buffer calls in element-wise and matrix multiplication operations"
This reverts commit af3441e67e706b2e5122369dc160353796867dd3.
* wip
* wip
* add comment
* fix: improve DMA transfer handling in mul_mat_gemv_impl for quantized source tensors
* add log
* try fix mulmat gemv
* wip
* fix: enhance DMA transfer handling in mul_mat_gemv_impl for quantized source tensors
* fix: optimize cache offset calculation and remove redundant swap in mul_mat_gemv_impl
* fix: refactor DMA transfer handling in mul_mat_gemv_impl for improved clarity and maintainability
* wip
* wip
* wip
* fix: enhance mul_mat_impl for improved cache handling and clarity
* fix: refactor tensor unflattening and DMA transfer initialization for improved clarity and type safety
* fix: improve cache handling of quant
* wip
* fix: improve cache handling in mul_mat_impl and mul_mat_gemv_impl for better memory efficiency
* rename
* add load_hexa_block_generic
* wip
* extract dequant block into separated function
* refactor: enhance dequantization functions with table parameter
* fix load_dual_block_generic
* refactor: rename dequantization functions for clarity and enhance block handling
* refactor: simplify dequantization logic by consolidating block handling and removing unused parameters
* wip
* wip
* feat: add make_qs_load_mask function and update load_dual_block_generic to use qs_indices
* fix load_dual_block_generic
* refactor: update load functions to use qs_indices for improved block loading
* wip
* fix: update loop indices and boundary checks to use size_t for better efficiency
* wip
* update make_scale_load_mask, to make it available for q8
* feat: add vec_dot_product_quant_impl for quantized dot product computation
* refactoring: move come quant func to dedicated file
* refactor: rename dequantization functions for clarity and consistency
* wip
* feat: enhance vec_dot_product_quant_impl with dual dequantization and improved assertions
* add vec_dot_product_vqf32_q40_f32
* wip
* wip
* wip
* wip
* implement vec_mpy_qf32_qf32_qf32 function and update vec_dot_product_vqf32_q40_f32 to use it
* wip
* add src0_plane_write_cache_offset
* wip
* enhance mul_mat_f32 to handle NPU_DATA_TYPE_Q4_0 for quantized matrix multiplication
* wip
* wip
* update test func
* refactor mul_mat_gemv_quant_impl to use get_nb for row stride and remove unused test function in init_f16_f32_table
* wip
* Add support for 4-block dequantization in vec_quant and update dot product implementation
* Refactor vec_dot_product_quant_impl to improve variable handling and enhance readability
* Refactor vec_dot_product_quant_impl to replace template function with inline vector operations
* use Q6_Vqf32_vmpy_VsfVsf instead of Q6_Vqf32_vmpy_Vqf32Vqf32
* Revert "use Q6_Vqf32_vmpy_VsfVsf instead of Q6_Vqf32_vmpy_Vqf32Vqf32"
This reverts commit 54839166fddbe40a0392adee5863c59070ccdbe4.
* wip
* improve log print in graph
* Refactor batched_row_dot to accept additional arguments and remove batched_row_dot_with_table
* Refactor synchronization functions to include previous operation and NE type parameters
* Refactor synchronization checks in several operations
* Update synchronization checks to include NPU_OP_COUNT in required conditions
* Add performance tracking to buffer management functions
* add memset
* add log
* fix: update backend device type from ACCEL to IGPU
* fix comment
This commit applies .clang-format rules to all source files under the
ggml-cann directory to ensure consistent coding style and readability.
The .clang-format option `SortIncludes: false` has been set to disable
automatic reordering of include directives.
No functional changes are introduced.
Co-authored-by: hipudding <huafengchun@gmail.com>
## Why it failed
When compiling with strict compiler flags (-Wwrite-strings -Werror=discarded-qualifiers),
the build fails with the following error:
```
cmake \
-S . \
-B ../llama.cpp.build \
--preset=x64-linux-gcc-debug \
-DCMAKE_INSTALL_PREFIX=/tmp/local \
-DCMAKE_C_FLAGS="-Wwrite-strings -Werror=discarded-qualifiers" && \
cmake --build ../llama.cpp.build/
...
/home/otegami/work/cpp/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c: In function ‘ggml_cpu_init’:
/home/otegami/work/cpp/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c:3572:24: error: passing argument 1 of ‘putenv’ discards ‘const’ qualifier from pointer target type [-Werror=discarded-qualifiers]
3572 | putenv("KMP_BLOCKTIME=200"); // 200ms
| ^~~~~~~~~~~~~~~~~~~
In file included from /home/otegami/work/cpp/llama.cpp/ggml/src/./ggml-impl.h:10,
from /home/otegami/work/cpp/llama.cpp/ggml/src/ggml-cpu/ggml-cpu-impl.h:6,
from /home/otegami/work/cpp/llama.cpp/ggml/src/ggml-cpu/traits.h:3,
from /home/otegami/work/cpp/llama.cpp/ggml/src/ggml-cpu/ggml-cpu.c:6:
/usr/include/stdlib.h:786:26: note: expected ‘char *’ but argument is of type ‘const char *’
786 | extern int putenv (char *__string) __THROW __nonnull ((1));
| ~~~~~~^~~~~~~~
cc1: some warnings being treated as errors
ninja: build stopped: subcommand failed.
```
The issue is that putenv() expects a non-const char * but receives a string literal (const char *).
## How to fix
This PR replaces putenv("KMP_BLOCKTIME=200") with setenv("KMP_BLOCKTIME", "200", 0).
Benefits of setenv():
- Accepts const char * parameters (no qualifier warnings)
- Makes copies of the strings (safer memory handling)
- The third parameter (0) ensures we don't overwrite if already set
* CPU: Add support for FLOOR,CEIL,ROUND and TRUNC unary operators
- Added the operators to unary op enum
- Implemented API functions
- Implemented forward and unary-op logic in CPU backend
- Updated ggml_get_n_tasks
- Updated operators names array and static_assert
- Updated docs and enabled automatic tests
* docs: add documentation for ggml_trunc and ggml_trunc_inplace in ggml.h
* chore: remove trailing whitespace from ggml.h
* Remove unresolved merge markers
* Apply review suggestions: cleanup formatting, enum order and leftover artifacts
* Regenerate ops.md using create_ops_docs.py
* opencl: add mm_q8_0_f32
* opencl: fix data loading for incomplete tile
* opencl: use q8_0 mm for larger matrix
* opencl: add some tests to cover the path
* optimise GGML_OP_SUM
* add non-contiguous tests by permuting the input
* change tests to require full contiguity of OP_SUM
* cuda : add check GGML_OP_SUM
---------
Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
* CUDA set scheduling strategy to spinning for cc121
* Using prop.major and prop.minor, include HIP and MUSA
* Exclude HIP and MUSA
* Remove trailing whitespace
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* Remove empty line
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
---------
Co-authored-by: Johannes Gäßler <johannesg@5d6.de>
* ggml : fix build broken with -march=armv9-a on MacOS
Signed-off-by: Jie Fu <jiefu@tencent.com>
* Add #pragma message
Signed-off-by: Jie Fu <jiefu@tencent.com>
* Address review comment.
Signed-off-by: Jie Fu <jiefu@tencent.com>
* Update ggml/src/ggml-cpu/ggml-cpu.c
---------
Signed-off-by: Jie Fu <jiefu@tencent.com>
Co-authored-by: Diego Devesa <slarengh@gmail.com>
This commit fixes a CPU-side memory leak issue in the CANN backend,
which occurred when intermediate aclTensorList objects were not properly
released after operator execution. The leak happened during repeated
invocations of CANN ops (e.g., FlashAttention), leading to increasing
host memory usage over time.
Proper resource cleanup (aclDestroyTensorList and related release logic)
has been added to ensure that all temporary tensors are correctly freed.
Many Ascend operators internally use FP16 precision for computation.
If input data is in FP32, it must first be cast to FP16 before
computation, and then cast back to FP32 after computation, which
introduces unnecessary cast operations. Moreover, FP16 computation
requires significantly less workload compared to FP32, leading to
noticeable efficiency improvements.
In this change, `get_rows`, `rms_norm`, and `flash_attn_ext` are extended
to support multiple data types. Validation on the Qwen2 0.5b model shows
correct accuracy and about 10% performance gain in concurrent scenarios.
Co-authored-by: noemotiovon <757486878@qq.com>
* scaffold to support opt step adamw on metal (not written so far)
* add opt-step-adamw kernel for metal
* pass op->src[4] as a separate buffer to the pipeline
* add bounds check to opt-step-adamw kernel
* complete scaffold for GGML_OP_SUM
* naive GGML_OP_SUM kernel
* remove unwanted comment
* change OP_SUM capability gate
* Add has_simdgroup_reduction to both ops to pass CI
* fix/refactor OP argsort, pad
* fix count-equal op
* update SYCL OP list
* fix format issue
---------
Co-authored-by: Zhang Jianyu <zhang.jianyu@outlook.com>
The previous SVE implementation for `ggml_vec_dot_f16_unroll` contained a bug due to a copy-paste error. The wrong variable was used in an FMA instruction, leading to incorrect results. This commit corrects the variable usage and improves the clarity of the code by renaming variables to avoid confusion.
Co-authored-by: Aaron <shelhamer.aaron@gmail.com>
* CANN: improve ACL graph matching
Record `ne` and `nb` information for src tensors and include them in the
graph matching check. This enhances the robustness of ACL graph matching
by preventing incorrect matches when src tensors share the same data
address but differ in shape or stride.
* CANN: add op_params match
* refactor to support soft_max_ext
* fix error and support soft_max_back
* rm unused functions
* fix format issue
---------
Co-authored-by: Zhang Jianyu <zhang.jianyu@outlook.com>
* metal : better unroll in the FA kernels
* metal : index FA blocks
* tests : restore [no ci]
* metal : prevent division by zero in FA kernels
* metal : fix -INF detection logic
* Add profiling
* More detailed profiling
* Rework command submission to avoid global locks
* Update wait handling
* try new method of waiting on futures
* Add serializing of command submission in some cases
* Add new pool for timestamp queries and clean up logging
* Serialize command submission in CI and leave a TODO note
* Update webgpu CI
* Add myself as WebGPU codeowner
* Deadlock avoidance
* Leave WebGPU/Vulkan CI serialized
* Fix divide by 0
* Fix logic in division by inflight_threads
* Update CODEOWNERS and remove serialize submit option
* metal : pad K, V and Mask when needed
* cont : simplify
* cuda : add TODO about KV padding requirement
* metal : add comments
* metal : remove mask padding requirement
* tests : add -INF blocks to the KQ mask in the FA tests
* cont : bump -INF block size to 64
Co-authored-by: Jeff Bolz <jbolz@nvidia.com>
* ggml : prevent division by zero in FA CPU op
---------
Co-authored-by: Jeff Bolz <jbolz@nvidia.com>
* metal : ssm_scan minor opts
* metal : get_rows optimize
* metal : cpy optimize
* metal : ssm_conv opt
* metal : ssm_scan simplify
* metal : ssm_Scan opt
This commit updates the leftover handling in ggml_vec_scale_f32.
The motivation for this is that the code currently incorrectly assumes
there would be fewer than ggml_f32_epr leftover elements. However,
since the main loop processes 2*ggml_f32_epr elements per iteration
, there can be up to (2*ggml_f32_epr - 1) leftover elements.
The original single-pass leftover code could only process ggml_f32_epr
elements, leaving some elements unscaled.
Example scenario with 256-bit SVE:
```
ggml_f32_epr = 8 (elements per register)
ggml_f32_step = 16 (two registers per iteration)
n = 25
np = 16
leftovers = 9 elements (16-24)
Original : processes only elements 16-23, misses element 24
This commit : loop processes elements 16-23, then element 24
```
Refs: https://github.com/ggml-org/llama.cpp/actions/runs/18070620247/job/51419855630
* Add power management utilities to NPU device context and update DCVS settings
* Update DCVS settings in power_utils to use v3 API and enhance power management
* wip
* Enhance dequantization functions by adding load_dequant_table support and updating signatures for improved performance
* use lut
* wip
* fix test failure
* wip
* Refactor load_qual_block_generic to improve block handling and optimize vector operations
* Enhance load_dual_block_generic and load_qual_block_generic to accept a mask parameter for improved block handling
* Refactor flash_attn_impl to optimize mask l2 prefetch
* wip
* wip
* wip
* wip
* add log
* link against shared libraries instead of static ones
* fix swiglu
* wip
* refactor expf_fix to handle overflow for different data types
* enhance is_glu_op_supported to validate shapes for multiple sources
* wip
* refactor logging macros to use hexagon namespace and improve formatting
* fix printf format error
* wip
* refactor: update static_assert messages for block size validation and add HVX_VectorPred_x3 type alias
* rename
* feat: enhance fa with mask
* wip
* wip
* refactor: replace instances of Q6_V_vzero() with kZeroV for consistency
* wip
* wip
* wip
* fix: improve address alignment check in HVX_Vector handling
* refactor: streamline vector dot product implementations for improved readability
* refactor: q4k add hvx intrinsic impl
* refactor: enhance dequantize_row_q4_K for clarity and performance
* refactor: optimize scale mask usage in dequantization functions for improved performance
* refactor: optimize dequantize_row_q4_K for intrinsic usage and performance improvements
* refactor: move GLU operation implementation into separated file
* sync after swiglu
* wip
* wip
* wip
* feat: increase prc main thread stack size
* fix: replace hardcoded stack size with NPU_THREAD_STACK_SIZE constant
* wip
* feat: add optimized vector operations for exponential and division with overflow handling
* wip
* feat: refactor exponential function to handle overflow and underflow with improved logic
* wip
* wip
* feat: add vector loading and scaling functions for improved performance in block processing
* wip
* feat: optimize block loading by refactoring scale index handling for improved performance
* use Q6_Vb_vlut32_VbVbR_nomatch instead
* feat: enhance scale loading by adding static assertion and restructuring block handling
* wip
* feat: refactor vec_dot_product_mixed_impl for improved clarity and performance
* wip
* feat: simplify vector loading functions and improve alignment handling
* wip
* feat: enhance scale loading mask with quantization block size validation
* wip
* feat: implement make_scale_load_mask function and refactor vector handling in vec_ops
* feat: enhance load_dual_block_generic to include scale indices for improved vector loading
* revert q8 dequant
* wip
* feat: optimize dequantization functions by removing unnecessary masking and updating lookup methods
* wip
* wip
* add qurt_mutex
* Add DMA transfer class and integrate into thread pool
* Enhance DMA transfer functionality by adding support for multiple descriptors and initiating transfers in parallel
* fix dma crash
* fix failed unit tests
* wip
* use alignas
* Improve DMA transfer error handling and update descriptor completion check
* Fix VTCM cache size calculation in element-wise operations
* Add cache clean operations before DMA transfers in element-wise operations
* reduce cache clean operations
* Refactor DMA transfer functions to support 1D operations and rename for clarity
* Enhance DMA transfer functionality by adding 2D submission support and improving descriptor initialization
* Update read buffer method to support forced invalidation and remove unnecessary invalidation calls in element-wise operations
* wip
* Improve DMA transfer handling in mul_mat_gemv_impl by replacing memcpy with initiate_dma_row_transfer and adding wait_for_dma logic
* fix 2d dma
* feat: add DMA plane cache
* rename
* wip
* use memcpy for debug
* fix cache plane calc
* refactor: remove debug logging from mul_mat_impl and optimize cache handling
* rename
* fix 2d dma type
* refactor: enhance DMA transfer handling in mul_mat_gemv_impl and wait functions
* refactor: optimize DMA transfer handling in mul_mat_gemv_impl and wait functions
* wip
* wip
* move op impl into sub dir
* add log
* fix: correct pointer usage in mul_mat_gemv_impl for next plane access
* fix: improve DMA transfer error handling in mul_mat_impl and mul_mat_gemv_impl
* fix: fix crash by using the entire row bytes
* wip
* wip
* fix: prevent parallelization for scalar src1 in is_mul_mat_supported
* fix: add dimension checks for 2D DMA transfers and fallback to 1D if necessary
* wip
* fix: enable thread barrier for mul multiplication operations
* feat: add synchronization checks for tensor operations and update related functions
* wip
* fix: remove invalidation flag from get_read_buffer calls in element-wise and matrix multiplication operations
* Revert "fix: remove invalidation flag from get_read_buffer calls in element-wise and matrix multiplication operations"
This reverts commit af3441e67e706b2e5122369dc160353796867dd3.
* wip
* wip
* add comment
* fix: improve DMA transfer handling in mul_mat_gemv_impl for quantized source tensors
* add log
* try fix mulmat gemv
* wip
* fix: enhance DMA transfer handling in mul_mat_gemv_impl for quantized source tensors
* fix: optimize cache offset calculation and remove redundant swap in mul_mat_gemv_impl
* fix: refactor DMA transfer handling in mul_mat_gemv_impl for improved clarity and maintainability
* wip
* wip
* wip
* fix: enhance mul_mat_impl for improved cache handling and clarity
* fix: refactor tensor unflattening and DMA transfer initialization for improved clarity and type safety
* fix: improve cache handling of quant
* wip
* fix: improve cache handling in mul_mat_impl and mul_mat_gemv_impl for better memory efficiency
* rename
* add load_hexa_block_generic
* wip
* extract dequant block into separated function
* refactor: enhance dequantization functions with table parameter
* fix load_dual_block_generic
* refactor: rename dequantization functions for clarity and enhance block handling
* refactor: simplify dequantization logic by consolidating block handling and removing unused parameters
* wip
* wip
* rpc : add support for multiple devices
Allow rpc-server to expose multiple devices from a single endpoint.
Change RPC protocol to include device identifier where needed.
closes: #15210
* fixes
* use ggml_backend_reg_t
* address review comments
* fix llama-bench backend report
* address review comments, change device naming
* fix cmd order
* vulkan (DRAFT): split shader generation by GLSL source file, to improve incremental build times
* support dep-files so shaders are recompiled if their included files change
* rename shader files which are used as "headers" to use .glsl extension
* move glslc extension detection shaders to separate folders
* the above is to prevent them from getting glob'd with the actual compute shaders that need to be compiled
* vulkan : only write embedded shader .hpp/.cpp when they change
* avoid recompiling ggml-vulkan.cpp when editing shaders
* pass single --source argument instead of --input-dir & --filter to shader gen
* check for source file match earlier
* fix hang in vulkan-shaders-gen when there are compilation errors
* early out did not decrement compile_count
* clean up
* fix glslc integer dot product test
* unconditionally write the embedded shader cpp output
* replace output filepath in generated dep-files to match output in CMakeLists
---------
Co-authored-by: Jeff Bolz <jbolz@nvidia.com>
* vulkan: Replace uses of maxMemoryAllocationSize and VK_WHOLE_SIZE
Replace maxMemoryAllocationSize check with maxBufferSize when creating buffers.
The maxMemoryAllocationSize limit is a "soft" limit and allocations can succeed
beyond that limit. This allows > 4GB buffers to be allocated on some
implementations (e.g. NVIDIA) and tensors this large can be used for im2col
and mul_mat.
For temporary buffers (prealloc_x/y/etc) check against maxStorageBufferRange.
I'm not sure this check is ideal, but we always use these buffers as a single
full size binding and the limit may be smaller than maxMemoryAllocationSize
or maxBufferSize, so I think this is reasonable.
Replace descriptor range uses of VK_WHOLE_SIZE with a manually computed range.
The maxStorageBufferRange may be smaller than the maxBufferSize or
maxMemoryAllocationSize (and the Vulkan spec warns about this in a note) and
it's invalid usage if VK_WHOLE_SIZE computes a range larger than
maxStorageBufferRange.
With this change, it should be possible to generate videos using wan networks
in stable-diffusion.cpp.
* vulkan: Add env var GGML_VK_FORCE_MAX_BUFFER_SIZE and use stoull
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>
This commit removes the `-dev` suffix from the version string in
CMakeLists.txt and the release script. The version will now be
just be formatted as `MAJOR.MINOR.PATCH`.
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
* vulkan: handle mat_mul with A matrix > 4GB
This change splits mat_mul operations with huge A matrix into chunks in the M
dimension. This works well for stable-diffusion use cases where the im2col
matrix has very large M.
Fix the order of setting the stride in mul_mm_cm2 - setting the dimension
clobbers the stride, so stride should be set after.
* build fixes
The "Clamp" spec constant is already based on whether KV is a multiple of Bc,
so use that to control whether bounds checking is performed. Add bounds checking
to the scalar and coopmat1 paths. Coopmat2 didn't need any changes (the K/V
tensors are already optionally clamped, nothing else needed to be changed).
Aaron Teo
Johannes Gäßler
Jeff Bolz
Aman Gupta
Shawn Gu
Radoslav Gerganov
Giuseppe Scrivano
muggle-stack
Ilia Ilmer
GittyBurstein
nullname
Chenguang Li
takuya kodama
yael-works
safranowith
lhez
Sam/Samuel
Julius Tischbein
Georgi Gerganov
SavicStefan
Anav Prasad
Jie Fu (傅杰)
hipudding
Neo Zhang Jianyu
sirus20x6
Diego Devesa
Prajwal B Mehendarkar
duduta
Charles Xu
ai-fonsi
Reese Levine
Daniel Bevenius
Eve
Acly
Piotr Wilkin (ilintar)
R0CKSTAR
uvos
alex-spacemit
Rafal Lewczuk
Sigbjørn Skjæret