Skip to content

Feature/repeat penalty - #20

Merged
ggerganov merged 6 commits into
ggml-org:masterfrom
beiller:feature/repeat_penalty
Mar 12, 2023
Merged

Feature/repeat penalty#20
ggerganov merged 6 commits into
ggml-org:masterfrom
beiller:feature/repeat_penalty

Conversation

@beiller

@beiller beiller commented Mar 11, 2023

Copy link
Copy Markdown
Contributor

Closes #4

I tried hacking together penalizing repeated tokens for n parameters. I'm actually not sure of the correct approach and I am not a great C++ programmer! But it appears to work somewhat. I think this can be improved much more with better sampling code.

I see that numbers can go negative so a fix from a referenced commit
Comment thread utils.cpp
@ggerganov
ggerganov merged commit 129c7d1 into ggml-org:master Mar 12, 2023
@ggerganov

Copy link
Copy Markdown
Member

Maybe should reconsider bringing back the top_k + top_p sampling:

https://github.com/ggerganov/llama.cpp/blob/129c7d1ea886e52ac1b87ff6184310bab3158806/utils.h#L75-L89

I decided to use top_p only because this is what is in the original Python repo, but I think I see comments that top_k helps.

abetlen pushed a commit to abetlen/llama.cpp that referenced this pull request Apr 10, 2023
fix: Mention of incorrect filename for MacOS cmake build artifact
SlyEcho pushed a commit to SlyEcho/llama.cpp that referenced this pull request Jun 12, 2023
flowgrad pushed a commit to flowgrad/llama.cpp that referenced this pull request Jun 27, 2023
* Cuda:
1) changed tensor split calculation to work on free vram instead of total vram (the main GPU often consumes a few GB from desktop work)

Changed the loading progress callback:
1) Added a numerical percentage to it
2) Added a status message to display what it is doing right now ("loading tensors: CPU)
3) normalized the length of the bar and designed it from dots to an actual progress bar
Example :
[============================================>-----]  88%  Loading tensor (GPU-Split)
Todo:
The progressbar can be changed into 3-4 lines, so it displays important flags and changes alongside the progress.
This would replace the list of "log entries" before the progress starts counting and preceding log-lines are only printed for errors and warnings.

* 1) Added a tensor split prepare wrapper, this allows to take the tensor split given by -ts immediately

2) Created a CUDA status struct with getter and update functions.
Moved g_main_device and g_num_devices into the new struct.
typedef struct {
    int num_devices;
    int main_device_id;
    size_t total_vram;
    size_t total_free_vram;
    struct cudaDeviceProp device_props[GGML_CUDA_MAX_DEVICES];
    size_t device_vram_free[GGML_CUDA_MAX_DEVICES];
    size_t device_vram_total[GGML_CUDA_MAX_DEVICES];
} GPUStatus;

3) Replaced the previous log output with a print function that gives better CUDA information on init
Example:
CUDA Device Summary - 2 devices found
+------------------------------------+------------+-----------+-----------+-----------+-----------+
| Device                             | VRAM Total | VRAM Free | VRAM Used |   Split % | Device ID |
+------------------------------------+------------+-----------+-----------+-----------+-----------+
| NVIDIA GeForce RTX 4090            |   24563 MB |  23006 MB |   1557 MB |      0.0% |  0 (Main) |
+------------------------------------+------------+-----------+-----------+-----------+-----------+
| NVIDIA GeForce RTX 3090            |   24575 MB |  23318 MB |   1257 MB |     67.2% |  1        |
+------------------------------------+------------+-----------+-----------+-----------+-----------+
Total VRAM: 47.99 GB, Free VRAM: 45.24 GB
--------------------

4) cleaned up libfalcon.cpp to use the new information instead of making cuda requests and device changes

TODO: The VRAM reported by CUDA is not the same as reported by CPU-Z which also is different from HWINFO64 (this one is totally off).

* CUDA:
1) Tensor split now allows to disable either device (usage proportion of 0)
- this currently still occupies minimal VRAM and if main-device it will still be used for non split tensors
2) Corrected (reduced) total free VRAM calculation if a device proportion is disabled (does not take main device into account)

* vram_total bugfix

* status table alignment

* bugfixes

* added perf counters also into makefile

---------
jesusmb1995 pushed a commit to jesusmb1995/llama.cpp that referenced this pull request Sep 29, 2025
QVAC-6114: Corrected build interface for libmtmd
YANGPuxu pushed a commit to YANGPuxu/llama.cpp that referenced this pull request Oct 30, 2025
docs: Add initial design for reload feature
reeselevine added a commit that referenced this pull request Mar 10, 2026
…better shader parameter handling (#20173)

* K quant speedup (#20)

* Basic JIT compilation for mul_mat, get_rows, and scale (#17)

* scale jit working

* preliminary working jit for getrows and mulmat, needs refining

* simplified mul_mat preprocessing switch statement

* get_rows fixes, mul_mat refinement

* formatted + last edits

* removed some extraneous prints

* fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish

* small fix

* some changes, working

* get_rows and mul_mat jit fixed and working

* Update formatting

* formatting

* Add header

---------

Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local>
Co-authored-by: Reese Levine <reeselevine1@gmail.com>

* Start work on all-encompassing shader library

* refactor argmax, set_rows

* Refactor all but flashattention, mat mul

* no gibberish, all k quants added, merged

* vec memory fix

* q6_k matching metal on my machine, tests passing

* Set tile size for q6_k separately

* Separate out fast shaders

---------

Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>

* Move towards writeBuffer for params

* Move away from multiple buffers for set_rows errors, remove host buffer for parameter buffers, minor cleanups

* Remove extra file

* Formatting

---------

Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
julien-c pushed a commit to julien-c/llama.cpp that referenced this pull request Mar 17, 2026
…lt-model-VRqa7

Update agent default model to GLM-4.7-Flash
Seunghhon pushed a commit to Seunghhon/llama.cpp that referenced this pull request Apr 26, 2026
* Adding repeat penalization

* Update utils.h

* Update utils.cpp

* Numeric fix

Should probably still scale by temp even if penalized

* Update comments, more proper application

I see that numbers can go negative so a fix from a referenced commit

* Minor formatting

---------

Co-authored-by: Georgi Gerganov <ggerganov@gmail.com>
Seunghhon pushed a commit to Seunghhon/llama.cpp that referenced this pull request Apr 26, 2026
…better shader parameter handling (ggml-org#20173)

* K quant speedup (ggml-org#20)

* Basic JIT compilation for mul_mat, get_rows, and scale (ggml-org#17)

* scale jit working

* preliminary working jit for getrows and mulmat, needs refining

* simplified mul_mat preprocessing switch statement

* get_rows fixes, mul_mat refinement

* formatted + last edits

* removed some extraneous prints

* fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish

* small fix

* some changes, working

* get_rows and mul_mat jit fixed and working

* Update formatting

* formatting

* Add header

---------

Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local>
Co-authored-by: Reese Levine <reeselevine1@gmail.com>

* Start work on all-encompassing shader library

* refactor argmax, set_rows

* Refactor all but flashattention, mat mul

* no gibberish, all k quants added, merged

* vec memory fix

* q6_k matching metal on my machine, tests passing

* Set tile size for q6_k separately

* Separate out fast shaders

---------

Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>

* Move towards writeBuffer for params

* Move away from multiple buffers for set_rows errors, remove host buffer for parameter buffers, minor cleanups

* Remove extra file

* Formatting

---------

Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
my-other-github-account pushed a commit to my-other-github-account/llama.cpp that referenced this pull request May 15, 2026
…better shader parameter handling (ggml-org#20173)

* K quant speedup (ggml-org#20)

* Basic JIT compilation for mul_mat, get_rows, and scale (ggml-org#17)

* scale jit working

* preliminary working jit for getrows and mulmat, needs refining

* simplified mul_mat preprocessing switch statement

* get_rows fixes, mul_mat refinement

* formatted + last edits

* removed some extraneous prints

* fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish

* small fix

* some changes, working

* get_rows and mul_mat jit fixed and working

* Update formatting

* formatting

* Add header

---------

Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local>
Co-authored-by: Reese Levine <reeselevine1@gmail.com>

* Start work on all-encompassing shader library

* refactor argmax, set_rows

* Refactor all but flashattention, mat mul

* no gibberish, all k quants added, merged

* vec memory fix

* q6_k matching metal on my machine, tests passing

* Set tile size for q6_k separately

* Separate out fast shaders

---------

Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>

* Move towards writeBuffer for params

* Move away from multiple buffers for set_rows errors, remove host buffer for parameter buffers, minor cleanups

* Remove extra file

* Formatting

---------

Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
meh pushed a commit to meh/llama.cpp that referenced this pull request May 19, 2026
ROCm 7.2 / RDNA3 build support + gated_delta_net perf fix
jimbothigpen added a commit to jimbothigpen/llama.cpp that referenced this pull request May 21, 2026
…buun 018092c)

Lifts buun 018092c "cuda: optimize turbo3 tcq set_rows" into yggdrasil's
k_set_rows_turboq3_tcq kernel and its dispatch in set-rows.cu.

Algorithmic / structural changes (bit-identical math vs prior implementation):

1. FWHT rewrite — the first 5 of 7 butterfly stages now use intra-warp
   __shfl_xor_sync (1 warp == 32 lanes covers them), so the loop drops
   from ~7 __syncthreads() to 3 (one per cross-warp stage + final).

2. Viterbi forward pass — predecessor minimum is independent of sid[8:6]
   output bits, so the 8-way scan only runs on 64 threads (sid<64) into
   pred_min_cost[64]/pred_min_p[64]; the 512-thread cost write then reads
   from pred_min_cost. Backtrace shrinks 8x: bt[t*64 + (sid & 0x3F)]
   instead of bt[t*512 + sid].

3. Backtrace storage — opt-in shared-memory backtrace when the device
   exposes >= 128*64=8192 bytes of opt-in shared/block (CUDA only;
   HIP/MUSA always use global bt_buf, gated by a one-shot device probe
   plus TURBO_TCQ_SHARED_BT=0 escape). New use_shared_bt kernel param
   and extern __shared__ bt_shared[]; global path's per-block stride
   shrinks 512 -> 64 bytes.

4. Parallel bitpack — 49 threads each pack one byte of dst_blk->qs
   (one byte per sid in [0,49)), replacing the serial OR loop on sid==0.

Yggdrasil-vs-buun naming adaptations (mainline-forward-sync convention):

- buun's diff lives in turbo-quant-cuda.cuh + set-rows.cu; yggdrasil
  consolidates the kernel + dispatch inside set-rows.cu, so all changes
  land in the single file.
- d_turbo3_tcq_codebook -> d_turboq3_tcq_codebook
- d_turbo_wht_signs{1,2}  -> TURBO_WHT_SIGNS{1,2}  (yggdrasil macro names)
- iq_is_k                  -> innerq_is_k          (yggdrasil arg name)
- block_turbo3_tcq         -> block_turboq3_tcq
- k_set_rows_turbo3_tcq    -> k_set_rows_turboq3_tcq
- __shfl_xor_sync gains a 4th WARP_SIZE arg per yggdrasil's HIP shim
  convention (see commit cdaeb0a).

Divergence vs buun:

- buun keeps per-device tcq3_use_shared_bt[GGML_CUDA_MAX_DEVICES] /
  tcq3_bt_checked[GGML_CUDA_MAX_DEVICES] arrays and a per-device
  ensure_tcq_bt_buf(device, bytes). Yggdrasil's tcq_bt_buf is single-
  device (matches the rest of yggdrasil's TCQ surface); this port keeps
  the single-device shape and uses plain static bool / static int for
  the one-shot probe rather than introducing the per-device array.
  Per [[port-fidelity-to-mainline-llamacpp]], the per-device upgrade is
  a separate scope and is not included here.

Untouched:

- k_set_rows_turboq2_tcq and set_rows_cuda_turboq2_tcq (2-bit TCQ) —
  buun's commit does not modify them, so neither does this port.
- InnerQ scaling / d_innerq_* paths and d_tcq_dump_* error-dump paths
  in the 3-bit kernel — yggdrasil-original code preserved verbatim.

The [[tcq-vulkan-port]] s6 SET_ROWS dispatch-sizing bug flag attached to
this row in recon/09 is a Vulkan concern; yggdrasil's Vulkan TCQ is
Phase 3c (net-new, not yet implemented) so the CUDA-only changes here
do not interact with it.

PPL target: bit-identical to s60/s61/s62 anchor PPL=6.9020 +- 0.05337
(Qwen3.5-9B-Q4_K_M turboq3_tcq KV, ai00 ROCm, n_seq=1, -c 4096 -ub 512
-b 512, GGML_CUDA_DISABLE_GRAPHS=1, 50 chunks).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
jimbothigpen added a commit to jimbothigpen/llama.cpp that referenced this pull request May 21, 2026
… bitpack

Replaces `uint8_t * outputs = (uint8_t *)x;` (alias onto __shared__ float
x[128]) with a dedicated `__shared__ uint8_t s_outputs[128]` array.

Root cause (session-65-resume-cell-c-ppl bisect, 2026-05-17): the buun ggml-org#20
parallel 49-thread bitpack reads outputs[sym_idx] from sids 1..48 (sid==0
writes the canonical winning-path bytes during backtrack, the __syncthreads()
following the backtrack is meant to publish those writes to all sids). Under
HIP/ROCm, however, the uint8_t-into-float[] alias is a strict-aliasing
violation, and the compiler can hoist or cache cross-thread reads of
outputs[] above the __syncthreads(), so sids 1..48 observed stale
(non-winning) values. Bytes 1..48 of dst_blk->qs[] were therefore packed
from arbitrary leftover bits, corrupting the symbol bitstream past byte 0
and producing the +12.7% PPL regression session-64 measured.

The fix is purely a storage-type cleanup: declare s_outputs[] as a typed
uint8_t shared array (128 bytes) and point the local `outputs` pointer at
it. All backtrack writes, recon_norm reads, and parallel-bitpack reads now
land in a type-consistent shared buffer, eliminating the alias hazard. The
__syncthreads() between the backtrack section and the bitpack section is
unchanged — the publish-to-all-sids semantics are now actually honored.

No algorithmic / arithmetic change. Buun's perf-gain intent (parallel
bitpack across 49 threads) is retained.

Expected PPL: bit-identical to anchor `[[phase-3-anchor-post-s60]]` =
6.9020 +/- 0.05337 at chunks 1-4 (Qwen3.5-9B-Q4_K_M turboq3_tcq KV, ai00
ROCm, n_seq=1, -c 4096 -ub 512 -b 512, GGML_CUDA_DISABLE_GRAPHS=1).

Untouched: k_set_rows_turboq2_tcq retains its (uint8_t *)x alias because
that kernel still uses the serial sid==0-only bitpack (writer and reader
are the same thread, so the strict-aliasing UB doesn't manifest as a
cross-thread visibility bug). If turboq2 ever gains a parallel bitpack,
it will need the same fix.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
jimbothigpen added a commit to jimbothigpen/llama.cpp that referenced this pull request May 21, 2026
…ive s_outputs[] fix

Source: buun 2e239fb "perf: port turbo3_tcq optimizations to turbo2_tcq encoder"

Ports six optimizations to k_set_rows_turboq2_tcq:
  - 128x64 backtrace (was 128x256); predecessor depends only on sid's low 6 bits
  - 64-group predecessor-minima precomputation (replaces 256 4-way scans/step)
  - shared-memory backtrace opt-in (CUDA only; HIP/MUSA stay on global bt)
  - FWHT warp shuffles (5 stages warp + 2 stages shared)
  - parallel 33-thread bitpack (one byte per thread, no atomics)
  - warp-shuffle final-state min reduction (replaces serial 8-warp scan)

Additionally folds in the strict-aliasing fix from Phase 3a ggml-org#20 (commit
70b3dd5) to k_set_rows_turboq2_tcq's outputs[] handling. Reason:
buun's port introduces a parallel-bitpack section that reads outputs[]
cross-thread, which under HIP/ROCm would have triggered the same
hoist-past-__syncthreads symbol corruption that produced the +12.7%
PPL regression session-64 measured on turboq3. Declare __shared__
uint8_t s_outputs[128] and point the local outputs pointer at it
instead of the (uint8_t *)x alias over __shared__ float x[128].

Per the 70b3dd5 commit-message forecast: "If turboq2 ever gains a
parallel bitpack, it will need the same fix."

Expected PPL chunks 1-4: bit-identical to the pre-port TURBOQ2_TCQ
baseline measured in session-phase-3a-23-turbo2-encoder-opt §3.3
(turboq2-baseline.log).

Mirrors ygg's ggml-org#21 dead-store removal (no pred_min_p[64] declared —
the same value is already written to bt[t*64+sid] for backtrack).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
jimbothigpen added a commit to jimbothigpen/llama.cpp that referenced this pull request May 25, 2026
…buun 018092c)

Lifts buun 018092c "cuda: optimize turbo3 tcq set_rows" into yggdrasil's
k_set_rows_turboq3_tcq kernel and its dispatch in set-rows.cu.

Algorithmic / structural changes (bit-identical math vs prior implementation):

1. FWHT rewrite — the first 5 of 7 butterfly stages now use intra-warp
   __shfl_xor_sync (1 warp == 32 lanes covers them), so the loop drops
   from ~7 __syncthreads() to 3 (one per cross-warp stage + final).

2. Viterbi forward pass — predecessor minimum is independent of sid[8:6]
   output bits, so the 8-way scan only runs on 64 threads (sid<64) into
   pred_min_cost[64]/pred_min_p[64]; the 512-thread cost write then reads
   from pred_min_cost. Backtrace shrinks 8x: bt[t*64 + (sid & 0x3F)]
   instead of bt[t*512 + sid].

3. Backtrace storage — opt-in shared-memory backtrace when the device
   exposes >= 128*64=8192 bytes of opt-in shared/block (CUDA only;
   HIP/MUSA always use global bt_buf, gated by a one-shot device probe
   plus TURBO_TCQ_SHARED_BT=0 escape). New use_shared_bt kernel param
   and extern __shared__ bt_shared[]; global path's per-block stride
   shrinks 512 -> 64 bytes.

4. Parallel bitpack — 49 threads each pack one byte of dst_blk->qs
   (one byte per sid in [0,49)), replacing the serial OR loop on sid==0.

Yggdrasil-vs-buun naming adaptations (mainline-forward-sync convention):

- buun's diff lives in turbo-quant-cuda.cuh + set-rows.cu; yggdrasil
  consolidates the kernel + dispatch inside set-rows.cu, so all changes
  land in the single file.
- d_turbo3_tcq_codebook -> d_turboq3_tcq_codebook
- d_turbo_wht_signs{1,2}  -> TURBO_WHT_SIGNS{1,2}  (yggdrasil macro names)
- iq_is_k                  -> innerq_is_k          (yggdrasil arg name)
- block_turbo3_tcq         -> block_turboq3_tcq
- k_set_rows_turbo3_tcq    -> k_set_rows_turboq3_tcq
- __shfl_xor_sync gains a 4th WARP_SIZE arg per yggdrasil's HIP shim
  convention (see commit cdaeb0a).

Divergence vs buun:

- buun keeps per-device tcq3_use_shared_bt[GGML_CUDA_MAX_DEVICES] /
  tcq3_bt_checked[GGML_CUDA_MAX_DEVICES] arrays and a per-device
  ensure_tcq_bt_buf(device, bytes). Yggdrasil's tcq_bt_buf is single-
  device (matches the rest of yggdrasil's TCQ surface); this port keeps
  the single-device shape and uses plain static bool / static int for
  the one-shot probe rather than introducing the per-device array.
  Per [[port-fidelity-to-mainline-llamacpp]], the per-device upgrade is
  a separate scope and is not included here.

Untouched:

- k_set_rows_turboq2_tcq and set_rows_cuda_turboq2_tcq (2-bit TCQ) —
  buun's commit does not modify them, so neither does this port.
- InnerQ scaling / d_innerq_* paths and d_tcq_dump_* error-dump paths
  in the 3-bit kernel — yggdrasil-original code preserved verbatim.

The [[tcq-vulkan-port]] s6 SET_ROWS dispatch-sizing bug flag attached to
this row in recon/09 is a Vulkan concern; yggdrasil's Vulkan TCQ is
Phase 3c (net-new, not yet implemented) so the CUDA-only changes here
do not interact with it.

PPL target: bit-identical to s60/s61/s62 anchor PPL=6.9020 +- 0.05337
(Qwen3.5-9B-Q4_K_M turboq3_tcq KV, ai00 ROCm, n_seq=1, -c 4096 -ub 512
-b 512, GGML_CUDA_DISABLE_GRAPHS=1, 50 chunks).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
jimbothigpen added a commit to jimbothigpen/llama.cpp that referenced this pull request May 25, 2026
… bitpack

Replaces `uint8_t * outputs = (uint8_t *)x;` (alias onto __shared__ float
x[128]) with a dedicated `__shared__ uint8_t s_outputs[128]` array.

Root cause (session-65-resume-cell-c-ppl bisect, 2026-05-17): the buun ggml-org#20
parallel 49-thread bitpack reads outputs[sym_idx] from sids 1..48 (sid==0
writes the canonical winning-path bytes during backtrack, the __syncthreads()
following the backtrack is meant to publish those writes to all sids). Under
HIP/ROCm, however, the uint8_t-into-float[] alias is a strict-aliasing
violation, and the compiler can hoist or cache cross-thread reads of
outputs[] above the __syncthreads(), so sids 1..48 observed stale
(non-winning) values. Bytes 1..48 of dst_blk->qs[] were therefore packed
from arbitrary leftover bits, corrupting the symbol bitstream past byte 0
and producing the +12.7% PPL regression session-64 measured.

The fix is purely a storage-type cleanup: declare s_outputs[] as a typed
uint8_t shared array (128 bytes) and point the local `outputs` pointer at
it. All backtrack writes, recon_norm reads, and parallel-bitpack reads now
land in a type-consistent shared buffer, eliminating the alias hazard. The
__syncthreads() between the backtrack section and the bitpack section is
unchanged — the publish-to-all-sids semantics are now actually honored.

No algorithmic / arithmetic change. Buun's perf-gain intent (parallel
bitpack across 49 threads) is retained.

Expected PPL: bit-identical to anchor `[[phase-3-anchor-post-s60]]` =
6.9020 +/- 0.05337 at chunks 1-4 (Qwen3.5-9B-Q4_K_M turboq3_tcq KV, ai00
ROCm, n_seq=1, -c 4096 -ub 512 -b 512, GGML_CUDA_DISABLE_GRAPHS=1).

Untouched: k_set_rows_turboq2_tcq retains its (uint8_t *)x alias because
that kernel still uses the serial sid==0-only bitpack (writer and reader
are the same thread, so the strict-aliasing UB doesn't manifest as a
cross-thread visibility bug). If turboq2 ever gains a parallel bitpack,
it will need the same fix.

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
jimbothigpen added a commit to jimbothigpen/llama.cpp that referenced this pull request May 25, 2026
…ive s_outputs[] fix

Source: buun 2e239fb "perf: port turbo3_tcq optimizations to turbo2_tcq encoder"

Ports six optimizations to k_set_rows_turboq2_tcq:
  - 128x64 backtrace (was 128x256); predecessor depends only on sid's low 6 bits
  - 64-group predecessor-minima precomputation (replaces 256 4-way scans/step)
  - shared-memory backtrace opt-in (CUDA only; HIP/MUSA stay on global bt)
  - FWHT warp shuffles (5 stages warp + 2 stages shared)
  - parallel 33-thread bitpack (one byte per thread, no atomics)
  - warp-shuffle final-state min reduction (replaces serial 8-warp scan)

Additionally folds in the strict-aliasing fix from Phase 3a ggml-org#20 (commit
70b3dd5) to k_set_rows_turboq2_tcq's outputs[] handling. Reason:
buun's port introduces a parallel-bitpack section that reads outputs[]
cross-thread, which under HIP/ROCm would have triggered the same
hoist-past-__syncthreads symbol corruption that produced the +12.7%
PPL regression session-64 measured on turboq3. Declare __shared__
uint8_t s_outputs[128] and point the local outputs pointer at it
instead of the (uint8_t *)x alias over __shared__ float x[128].

Per the 70b3dd5 commit-message forecast: "If turboq2 ever gains a
parallel bitpack, it will need the same fix."

Expected PPL chunks 1-4: bit-identical to the pre-port TURBOQ2_TCQ
baseline measured in session-phase-3a-23-turbo2-encoder-opt §3.3
(turboq2-baseline.log).

Mirrors ygg's ggml-org#21 dead-store removal (no pred_min_p[64] declared —
the same value is already written to bt[t*64+sid] for backtrack).

Co-Authored-By: Claude Opus 4.7 (1M context) <noreply@anthropic.com>
fewtarius pushed a commit to fewtarius/CachyLLama that referenced this pull request May 30, 2026
…better shader parameter handling (ggml-org#20173)

* K quant speedup (ggml-org#20)

* Basic JIT compilation for mul_mat, get_rows, and scale (ggml-org#17)

* scale jit working

* preliminary working jit for getrows and mulmat, needs refining

* simplified mul_mat preprocessing switch statement

* get_rows fixes, mul_mat refinement

* formatted + last edits

* removed some extraneous prints

* fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish

* small fix

* some changes, working

* get_rows and mul_mat jit fixed and working

* Update formatting

* formatting

* Add header

---------

Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local>
Co-authored-by: Reese Levine <reeselevine1@gmail.com>

* Start work on all-encompassing shader library

* refactor argmax, set_rows

* Refactor all but flashattention, mat mul

* no gibberish, all k quants added, merged

* vec memory fix

* q6_k matching metal on my machine, tests passing

* Set tile size for q6_k separately

* Separate out fast shaders

---------

Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>

* Move towards writeBuffer for params

* Move away from multiple buffers for set_rows errors, remove host buffer for parameter buffers, minor cleanups

* Remove extra file

* Formatting

---------

Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
MrLordCat referenced this pull request in MrLordCat/llama.cpp-rdna-lab Jul 16, 2026
…better shader parameter handling (#20173)

* K quant speedup (#20)

* Basic JIT compilation for mul_mat, get_rows, and scale (#17)

* scale jit working

* preliminary working jit for getrows and mulmat, needs refining

* simplified mul_mat preprocessing switch statement

* get_rows fixes, mul_mat refinement

* formatted + last edits

* removed some extraneous prints

* fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish

* small fix

* some changes, working

* get_rows and mul_mat jit fixed and working

* Update formatting

* formatting

* Add header

---------

Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local>
Co-authored-by: Reese Levine <reeselevine1@gmail.com>

* Start work on all-encompassing shader library

* refactor argmax, set_rows

* Refactor all but flashattention, mat mul

* no gibberish, all k quants added, merged

* vec memory fix

* q6_k matching metal on my machine, tests passing

* Set tile size for q6_k separately

* Separate out fast shaders

---------

Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>

* Move towards writeBuffer for params

* Move away from multiple buffers for set_rows errors, remove host buffer for parameter buffers, minor cleanups

* Remove extra file

* Formatting

---------

Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
SolshineCode pushed a commit to SolshineCode/llama.cpp that referenced this pull request Jul 28, 2026
…rser

chat : add Kimi K3 chat format (reasoning, content, typed tool calls)
kainlan added a commit to kainlan/llama.cpp-intel-optimizations that referenced this pull request Aug 1, 2026
…gap at its site

Both findings from the final cross-task integration review over 0ab1b1a..f8906c5.

Test ggml-org#20 (sycl-reset-model-weight-lease-preserve) carried no
ONEAPI_DEVICE_SELECTOR. Its own source falls back to level_zero:0 -- the B70,
the benchmark card -- when nothing is set, so a bare `ctest -R` or CLAUDE.md's
throttled full sweep would land it there and perturb a run in progress. Its
sibling test-sycl-zone-reset-live-refusal, written in the same session, pins
level_zero:1 with an explicit rationale, as do unpin-event and env-report.

This is a seam no per-task review could see: the placement in
ggml-sycl/CMakeLists.txt was principled (it matches the RESIDENCY_TEST link-only
precedent for build reasons), and it silently inherited that file's convention --
none of its ~30 GPU tests pin a selector at all -- rather than the stronger one
established minutes earlier in tests/CMakeLists.txt. Both registrations were
individually correct; the gap existed only between them. The comment now says
explicitly that the surrounding tests' silence is not the local rule.

Also drops the "ljb9" task-id label suffix, which no other test in that file
carries and its own sibling does not, and records the measurement that justified
dropping the hazard labels (free -g flat at 210 GB, peak RSS 317 MB, against a
family that reaches ~227 GB) so a future reader does not have to re-derive it.

Second finding: reset_model_weight_entries() explained preserve-and-continue but
carried no pointer to llama.cpp-0qlw, which establishes that the fix is
incomplete -- in_use_count == 0 means "nobody is resolving this at this instant",
not "no live model owns it", so a live model's idle weights are still freed. The
caveat was accurate but recorded only in the regression test's header and in
f8906c5's message; a reader opening unified-cache.cpp, the natural place to
look, would not have seen it. Not a false claim anywhere -- an incompleteness in
where the true claim lives.

Verified: build exit 0, 0 errors, GGML_SYCL:BOOL=ON, ldd 2. Registry regenerated
and confirmed fresh; ctest -N shows the pin and LABELS as intended; the test still
survives -LE 'residency|mem-handle|cache' (returns 1, so the safe sweep runs it);
Test ggml-org#20 passes in 1.14s under the pin, and the other three new gates stay green.

Co-Authored-By: Claude Opus 5 (1M context) <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01EHYpYxXNwtkx2t8dQJNumx
juliantorr-es pushed a commit to Tribunus-dev/tessera that referenced this pull request Aug 14, 2026
Pipeline refactor phases 1-4 (expert-eval seam, eval cache, KV-joint
measurement plumbing, dispatch decomposition + CLI unification), the
real per-tensor activation-capture feature for GA/AWQ fitness, two
argument-parsing fixes (champq/w4a4 --model requirement, value-taking
flags silently swallowing the next flag), and KV-joint reconstruction
item 5 (scale migration: fold K/V's migrated scale into W_q/W_o).

See docs/tessera-refactor-implementation-spec.md and
docs/tessera-kv-joint-reconstruction-design.md for the design context.
Phase gates (ggml-org#10, ggml-org#20, ggml-org#23, ggml-org#26) remain outstanding pending a real
model run, not available in this environment.
Madreag added a commit to Madreag/turbo3-cuda that referenced this pull request Aug 15, 2026
proxy v6.3: slot restore always attempted (A1 slots-dir bug), truncated
streams now deliver error frame + [DONE]/error event to CLIENT on both paths
(A3), passthrough allowlisted — /slots and raw completion endpoints blocked
(A4), swap shielded + owner=None during transition (TheTom#11), robust tripwire
classifier (TheTom#12), literal think-tags after content start stay content +
finish() flushes tag residue (TheTom#10), Anthropic images translated not dropped
(TheTom#14), Anthropic heartbeat pings (TheTom#4), mid-stream upstream death caught (TheTom#9),
retry guarded post-prepare (TheTom#8), partial-artifact persist only on abort (TheTom#6),
persistent trace handle + capture freshness/prune (TheTom#7,TheTom#17), generic blob
stripper (TheTom#16), hop-by-hop headers stripped (ggml-org#20), non-dict bodies 400 (TheTom#15),
read_task awaited (TheTom#19).

scripts: stop.sh verified kills (poll→SIGKILL, net only for survivors, port
check); start-long-38.sh double-start guard, log rotation, health-timeout
aborts before proxy launch, --api-key-file (key off cmdline), --metrics,
--swa-full dropped (dead flag), corrected comments+rollback; status.sh
probes ports+health. keys.json 600; stale systemd units + qwen36-era slot
sentinel archived to attic/.

test_proxy.py: runner moved to EOF (was silently skipping 5 tests — A6),
script-mode fixture shim, +7 regression tests → 49/49.

Claude-Session: https://claude.ai/code/session_01XTXkWYoiB8SvYyQ8VduRat
zommiommy pushed a commit to zommiommy/llama.cpp that referenced this pull request Aug 18, 2026
…better shader parameter handling (ggml-org#20173)

* K quant speedup (ggml-org#20)

* Basic JIT compilation for mul_mat, get_rows, and scale (ggml-org#17)

* scale jit working

* preliminary working jit for getrows and mulmat, needs refining

* simplified mul_mat preprocessing switch statement

* get_rows fixes, mul_mat refinement

* formatted + last edits

* removed some extraneous prints

* fixed get_rows, fixed workgroup dispatch in mul_mat. no gibberish

* small fix

* some changes, working

* get_rows and mul_mat jit fixed and working

* Update formatting

* formatting

* Add header

---------

Co-authored-by: Neha Abbas <nehaabbas@ReeseLevines-MacBook-Pro.local>
Co-authored-by: Reese Levine <reeselevine1@gmail.com>

* Start work on all-encompassing shader library

* refactor argmax, set_rows

* Refactor all but flashattention, mat mul

* no gibberish, all k quants added, merged

* vec memory fix

* q6_k matching metal on my machine, tests passing

* Set tile size for q6_k separately

* Separate out fast shaders

---------

Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>

* Move towards writeBuffer for params

* Move away from multiple buffers for set_rows errors, remove host buffer for parameter buffers, minor cleanups

* Remove extra file

* Formatting

---------

Co-authored-by: neha-ha <137219201+neha-ha@users.noreply.github.com>
platima added a commit to platima/llama.cpp-SpacemiT-K3 that referenced this pull request Sep 1, 2026
… path (+35.6% tg)

Cherry-picks only the ggml/src/ggml-cpu/spacemit/ portion of SpacemiT 0.1.7
(commits 5a23f07 ggml-org#25 and 6ad6d85 ggml-org#27). The remainder of those commits is
the ONNX media / smt-mtmd / Qwen3-TTS wrapper path, which this build does not
enable, so taking them whole would add a large untested surface for no gain.

Corrects an earlier assumption in MODELS.md that the newer spacemit commits
were ONNX-only. That holds for ggml-org#19/ggml-org#20/ggml-org#21 but not 0.1.7, which adds a
gemm_m == 1 (decode) fast path for q4_0/q8_0 that streams wide Q8 output heads
straight through TCM, replaces the even/odd pair-barrier lockstep, and adds RVV
GELU/TANH/GEGLU kernels wired in as GGML_OP_UNARY/GGML_OP_GLU.

Applied with git apply -3, no conflicts. Verified our backend patches survived:
patch 20's GGML_OP_CONT bypass in get_tensor_traits, the patch 14 dispatch and
supports log hooks, and patch 28's max_perfer_threads clamp.

Same-session A/B (llama-bench -t 8 -p 128 -n 128 -mmp 0 -fa 1 -ub 128 -r 3),
stashing and rebuilding between measurements so there is no cross-session drift:

  gemma-4-E2B_q4_0   pp 116.39 -> 136.55 (+17.3%)   tg 11.42 -> 15.49 (+35.6%)
  Qwen3-4B-Q4_K_M    pp  57.63 ->  57.58            tg  9.20 ->  9.20

q4_K is unchanged within noise, as expected: the fast path is if constexpr
gated to q4_0/q8_0. This matters because it is a decode win, which patch 17
found the IME2 HP tile could not deliver (M=1 is bandwidth-bound). Output
coherence spot-checked on both quants at --temp 0.

Co-Authored-By: Claude Opus 5 <noreply@anthropic.com>
Sign up for free to join this conversation on GitHub. Already have an account? Sign in to comment

Labels

None yet

Projects

None yet

Development

Successfully merging this pull request may close these issues.

Repetition penalty

2 participants