Feature/repeat penalty by beiller · Pull Request #20 · ggml-org/llama.cpp

beiller · 2023-03-11T19:25:03Z

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.

Should probably still scale by temp even if penalized

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

ggerganov · 2023-03-12T09:30:54Z

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.

fix: Mention of incorrect filename for MacOS cmake build artifact

Logging changes

* 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 ---------

QVAC-6114: Corrected build interface for libmtmd

docs: Add initial design for reload feature

…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>

…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>

…lt-model-VRqa7 Update agent default model to GLM-4.7-Flash

…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>

* 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>

…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>

For LLaMA-3.1 models: * It is better to quantize all of attn_v with iq3_k instead of half of attn_v with iq4_k * Quantizing attn_output with iq3_k results in a larger PPL decrease compared to what one expects from the added bpw. Co-authored-by: Iwan Kawrakow <iwan.kawrakow@gmail.com>

…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>

ROCm 7.2 / RDNA3 build support + gated_delta_net perf fix

…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>

… 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>

…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>

…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>

… 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>

…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>

…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>

… 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>

…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>

beiller added 4 commits March 11, 2023 14:23

Adding repeat penalization

d2b96b5

Update utils.h

3f6a118

Update utils.cpp

78651d5

Numeric fix

c90e78e

Should probably still scale by temp even if penalized

G2G2G2G mentioned this pull request Mar 12, 2023

Inference on the model shawwn/llama-dl#1

Open

Update comments, more proper application

340bff0

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

ggerganov reviewed Mar 12, 2023

View reviewed changes

Comment thread utils.cpp

Minor formatting

ebb357f

ggerganov merged commit 129c7d1 into ggml-org:master Mar 12, 2023

Piezoid mentioned this pull request Mar 20, 2023

Improving the repetition penalty #331

Closed

abetlen pushed a commit to abetlen/llama.cpp that referenced this pull request Apr 10, 2023

Merge pull request ggml-org#20 from BrutalCoding/patch-1

e84c446

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

Merge pull request ggml-org#20 from SlyEcho/server_refactor

eee8b28

Logging changes

phymbert referenced this pull request Mar 23, 2024

split: add gguf-split in the make build target (#6262)

21cad01

mofosyne mentioned this pull request May 29, 2024

Bug: Inconsistent ggml-4-x86-cuda-v100 ci failures on master #7613

Closed

Bearsaerker mentioned this pull request Mar 12, 2025

Eval bug: Gemma 3 extremly slow prompt processing when using quantized kv cache. #12352

Closed

jesusmb1995 pushed a commit to jesusmb1995/llama.cpp that referenced this pull request Sep 29, 2025

Merge pull request ggml-org#20 from jpgaribotti/QVAC-6114

285713d

QVAC-6114: Corrected build interface for libmtmd

YANGPuxu pushed a commit to YANGPuxu/llama.cpp that referenced this pull request Oct 30, 2025

Merge pull request ggml-org#20 from YANGPuxu/ypx/feat/reload

65cdb63

docs: Add initial design for reload feature

uttampc1 mentioned this pull request Nov 18, 2025

Throughput improvement for small batch sizes #17342

Open

sainnhe mentioned this pull request Jan 25, 2026

Eval bug: coredump due to ops of discontinuous tensor memory #19078

Closed

feyleth mentioned this pull request Mar 11, 2026

Eval bug: when image is to big the vision model crash #20418

Closed

julien-c pushed a commit to julien-c/llama.cpp that referenced this pull request Mar 17, 2026

Merge pull request ggml-org#20 from gary149/claude/update-agent-defau…

5637533

…lt-model-VRqa7 Update agent default model to GLM-4.7-Flash

rubin55 mentioned this pull request Mar 26, 2026

Eval bug: Unresolved Symbol <__memcpy_chk> when running (any?) model #21041

Closed

Cyrille37 mentioned this pull request Apr 13, 2026

Misc. bug: new EuroLLM version "EuroLLM-9B-Instruct-2512" failed after sereral requests #21856

Open

heislera763 mentioned this pull request May 2, 2026

TP: fix ggml context size calculation #22616

Merged

freelinuxer-cn mentioned this pull request May 9, 2026

Misc. bug: Meta backend ggml_context pool exhaustion with --split-mode tensor #22404

Closed

meh pushed a commit to meh/llama.cpp that referenced this pull request May 19, 2026

Merge pull request ggml-org#20 from flamme-demon/buun-rocm

330664a

ROCm 7.2 / RDNA3 build support + gated_delta_net perf fix

savaszamanianimations-alt mentioned this pull request May 20, 2026

Eval bug: Core dumped when running qwen 3.6 27b on vulkan. No-mmproj works #23430

Open

kenneth55555 mentioned this pull request May 23, 2026

Eval bug: MTP 崩溃，看不懂，我会同步反馈模型厂商 #23551

Closed

Provide feedback

Saved searches

Use saved searches to filter your results more quickly

Feature/repeat penalty#20

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

beiller commented Mar 11, 2023 •

edited

Loading

Uh oh!

Uh oh!

ggerganov commented Mar 12, 2023

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

Conversation

beiller commented Mar 11, 2023 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Uh oh!

ggerganov commented Mar 12, 2023

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

2 participants

beiller commented Mar 11, 2023 •

edited

Loading