Fix kernel OpenCL Mesa/rusticl per llama.cpp #1

Nowe zgłoszenie

Otwarty

otworzone 2026-07-09 21:00:59 +00:00 przez devbadxyz · 1 comment

devbadxyz skomentował(-a) 2026-07-09 21:00:59 +00:00

Członek

Skopiuj link

Copy Source

Questo file è un prompt auto-contenuto da dare a Claude Code direttamente sul
Windows DevKit 2023 (Snapdragon 8cx Gen 3 + GPU Adreno 690, Mesa/rusticl
custom già installata). Non presuppone che chi lo esegue abbia accesso alla
conversazione originale: tutto il contesto necessario è qui sotto.

Obiettivo finale: produrre un file FIX.md con l'elenco esatto e verificato
delle modifiche da riportare in llamacpp-opencl.Containerfile, così da poter
integrare il fix nel container senza dover ripetere il lavoro di scoperta.

---

Prompt da dare a Claude Code

Contesto: sto cercando di far funzionare llama.cpp con il backend OpenCL su
una GPU Qualcomm Adreno 690 (Freedreno via Mesa/rusticl open source, NON il
driver proprietario Qualcomm). L'hardware è un Windows DevKit 2023 (Snapdragon
8cx Gen 3) con Ubuntu 26.04 arm64.

Sul sistema è già installata una build custom di Mesa (da
gitlab.freedesktop.org/mesa/mesa.git, branch main) che include il fix
upstream "freedreno/a6xx: Expose subgroup ops" (commit 185c89084ab7), non
presente in nessuna release stabile. Verifica con:

    clinfo | grep -i "platform name\|device name\|subgroup"

Ti aspetti: Platform "rusticl", Device "FD690", ed estensioni
cl_khr_subgroup_ballot/shuffle/rotate/ecc. Se questo non c'è, FERMATI e
segnalalo: significa che l'installazione Mesa è rotta, non è un problema di
llama.cpp.

## Il problema

`git clone --depth 1 https://github.com/ggml-org/llama.cpp`, poi compila con:

    cmake -B build -G Ninja -DCMAKE_BUILD_TYPE=Release -DGGML_OPENCL=ON \
        -DGGML_OPENCL_EMBED_KERNELS=ON -DGGML_OPENCL_USE_ADRENO_KERNELS=OFF \
        -DBUILD_SHARED_LIBS=ON
    cmake --build build -j$(nproc)

La build C++ compila senza errori. Il problema emerge A RUNTIME (avviando
`llama-server` con un modello reale), quando ggml compila i kernel OpenCL
embedded via `clBuildProgram`. Errore tipico:

    input.cl:N:M: warning: OpenCL extension 'cl_qcom_reqd_sub_group_size'
        unknown or does not require pragma - ignoring
    input.cl:N:M: warning: unknown attribute 'qcom_reqd_sub_group_size' ignored
    input.cl:N:M: error: loading directly from pointer to type '__global half'
        requires cl_khr_fp16. Use vector data load builtin functions instead
    Error executing LLVM compilation action.

## Perché succede (già investigato, non ripartire da zero)

I kernel sorgente sono in `ggml/src/ggml-opencl/kernels/*.cl`. Ne esistono
94 che dichiarano puntatori a `half`/`half2`/`half4`/`half8` (verificato con
`grep -lE '\b(half|half2|half4|half8)\s*\*' ggml/src/ggml-opencl/kernels/*.cl`).
Mesa/rusticl (a differenza del driver Adreno proprietario per cui questi
kernel sono stati scritti) RIFIUTA il dereference diretto di puntatori half
(`*ptr` o `ptr[i]`), richiedendo esplicitamente le funzioni builtin OpenCL
`vload_half()` / `vload_half2()` / `vload_half4()` / `vload_half8()` /
`vstore_half()` ecc., anche con `#pragma OPENCL EXTENSION cl_khr_fp16 :
enable` attivo.

Equivalenza semantica (per lo scalare; esistono le controparti vettoriali
vload_half2/4/8 per i tipi vettore):

    *P        ->  vload_half(0, P)
    P[E]      ->  vload_half(E, P)

Per half2/half4/half8 la firma è identica ma il valore restituito è un
vettore: `vload_half2(offset, p)` legge 2 half consecutivi da `p+offset*2` e
li converte in `float2`. ATTENZIONE alla differenza di offset in elementi
half vs elementi vettore - verificare sempre contro la documentazione OpenCL
ufficiale (khronos.org) prima di applicare, non indovinare.

NON è un problema del fix Mesa dei subgroup, è indipendente: riguarda TUTTI
i kernel che leggono direttamente valori half, a prescindere dal fix
subgroup.

### Cosa è STATO GIÀ provato e scartato

- `-DGGML_OPENCL_USE_ADRENO_KERNELS=OFF`: NON risolve il problema.
  Contrariamente a un vecchio commento di un maintainer nella PR
  https://github.com/ggml-org/llama.cpp/pull/10693 (che diceva "i kernel
  generici dovrebbero funzionare"), nella versione attuale di llama.cpp la
  famiglia di kernel `gemv_noshuffle_*`/`gemm_noshuffle_*` viene comunque
  compilata sempre, indipendentemente da questo flag (verificato nel log di
  build: questi kernel compaiono nella lista "opencl: embedding kernel ..."
  a prescindere dal valore del flag). Quel consiglio è ormai superato.
- Forzare `-D cl_qcom_reqd_sub_group_size` nelle opzioni di compilazione
  kernel (per attivare il ramo `ADRENO_GPU` nei kernel che ne hanno bisogno
  per macro come N_SIMDGROUP/N_DST/N_SIMDWIDTH): questo FUNZIONA e va
  mantenuto, ma da solo non basta - risolve solo l'errore "undeclared
  identifier", non l'errore fp16 successivo.
- Ho già un piccolo script Python che fa la sostituzione automatica
  *P -> vload_half(0,P) / P[E] -> vload_half(E,P) SOLO per i file
  `*flat*.cl` con puntatori `half` scalari (14 file, 27 righe modificate,
  verificate manualmente una per una col diff). Funziona, ma copre solo una
  piccola parte dei 94 file totali e solo il caso scalare.

## Cosa devi fare

1. Estendi (o riscrivi da zero, se più pulito) lo script di patch per
   coprire TUTTI i 94 file, gestendo correttamente half/half2/half4/half8.
   Usa la stessa metodologia rigorosa già validata: per ogni modifica,
   genera il diff e verificalo a mano leggendo il codice circostante prima
   di fidarti — l'obiettivo è zero sostituzioni sbagliate silenziose (offset
   vettoriali errati = pesi del modello letti male = output numerico
   sbagliato ma nessun crash, il peggior tipo di bug qui).

2. Compila e, per OGNI kernel che fallisce, applica la patch, ricompila.
   Hai il vantaggio che noi non avevamo: accesso reale all'hardware. Usa
   `RUSTICL_ENABLE=freedreno ./llama-server -hf <un modello piccolo, es.
   Qwen2.5-0.5B-Instruct-GGUF:Q4_K_M> --no-warmup -ngl 99` per iterare
   velocemente (modello piccolo = build/reload rapidi). Prova anche un
   secondo modello con un formato di quantizzazione diverso (es. uno IQ4_NL
   e uno Q4_K/Q6_K) per essere sicuro di coprire più famiglie di kernel.

3. IMPORTANTE - verifica la CORRETTEZZA NUMERICA, non solo che compili e
   non crashi: fai generare al modello una risposta a un prompt semplice e
   verificabile (es. "What is 2+2? Answer with just the number.") e
   controlla che la risposta sia sensata, non NaN/garbage/ripetizioni
   insensate. Un kernel con offset vload_half sbagliato spesso compila ed
   esegue senza errori, ma produce output numericamente corrotto - l'unico
   modo per scoprirlo è guardare l'output reale, non i log di compilazione.

4. Confronta se possibile l'output con lo stesso identico prompt/seed su
   CPU pura (`-ngl 0`) per lo stesso modello, per avere un termine di
   paragone di sanità (non serve che sia identico token-per-token, ma deve
   essere semanticamente coerente allo stesso modo).

5. Testa anche un modello MoE reale se possibile (es. lo stesso Qwen A3B o
   Gemma usato nei test precedenti) dato che la famiglia di kernel
   `gemm_moe_*`/`gemv_moe_*` è specifica per il routing MoE e potrebbe avere
   pattern di lettura half diversi dai kernel mul_mv standard.

6. Scrivi un file `FIX.md` (nella stessa cartella di questo TODO.md) con:
   - L'elenco esatto dei file modificati e il diff di ciascuno (o un unico
     patch/diff file allegato, se più pratico).
   - Quali modelli/formati di quantizzazione hai testato con successo e
     quali no.
   - Il valore CMake finale raccomandato per
     `GGML_OPENCL_USE_ADRENO_KERNELS` (ON o OFF) alla luce di quanto trovato.
   - Eventuali ulteriori patch necessarie oltre a queste (es. se scopri che
     serve ANCHE differenziare qualcosa per modelli MoE).
   - Se possibile, un confronto tokens/sec CPU (-ngl 0, -t 4) vs GPU
     (-ngl 99) sullo stesso modello/prompt, utile per capire se ne vale la
     pena rispetto a girare su CPU.

Non serve che tu conosca o modifichi il Containerfile del container Podman -
lavora direttamente su un checkout locale di llama.cpp sul DevKit. Il FIX.md
che produci verrà usato da un'altra sessione per applicare le modifiche
corrispondenti al Containerfile.

---

Riferimento: stato attuale del Containerfile (per chi applicherà FIX.md)

llamacpp-opencl.Containerfile contiene già, in ordine:

• PATCH 1: riconosce FD6* come famiglia Adreno (necessaria sempre).
• PATCH 2: accetta cl_khr_subgroup_ballot come prova di supporto
  subgroup, dato che Mesa non espone mai la stringa legacy
  cl_khr_subgroups (necessaria sempre).
• PATCH 3: forza -D cl_qcom_reqd_sub_group_size nelle opzioni di
  compilazione kernel (necessaria se si riattivano kernel che usano
  REQD_SUBGROUP_SIZE_64/ADRENO_GPU).
• PATCH 4: vload_half() per i 14 file *flat*.cl con puntatori half
  scalari (copertura parziale, superata da questo lavoro).
• -DGGML_OPENCL_USE_ADRENO_KERNELS=OFF (da rivalutare in base a FIX.md).

Il pacchetto Mesa custom (.deb in containers/llamacpp/deb/, tracciato via
Git LFS) è generato da build-mesa.sh nella stessa cartella e NON deve essere
toccato per questo lavoro - il problema è tutto lato llama.cpp/kernel OpenCL,
non lato driver.

Questo file è un prompt auto-contenuto da dare a Claude Code **direttamente sul Windows DevKit 2023** (Snapdragon 8cx Gen 3 + GPU Adreno 690, Mesa/rusticl custom già installata). Non presuppone che chi lo esegue abbia accesso alla conversazione originale: tutto il contesto necessario è qui sotto. Obiettivo finale: produrre un file `FIX.md` con l'elenco esatto e verificato delle modifiche da riportare in `llamacpp-opencl.Containerfile`, così da poter integrare il fix nel container senza dover ripetere il lavoro di scoperta. --- ## Prompt da dare a Claude Code ``` Contesto: sto cercando di far funzionare llama.cpp con il backend OpenCL su una GPU Qualcomm Adreno 690 (Freedreno via Mesa/rusticl open source, NON il driver proprietario Qualcomm). L'hardware è un Windows DevKit 2023 (Snapdragon 8cx Gen 3) con Ubuntu 26.04 arm64. Sul sistema è già installata una build custom di Mesa (da gitlab.freedesktop.org/mesa/mesa.git, branch main) che include il fix upstream "freedreno/a6xx: Expose subgroup ops" (commit 185c89084ab7), non presente in nessuna release stabile. Verifica con: clinfo | grep -i "platform name\|device name\|subgroup" Ti aspetti: Platform "rusticl", Device "FD690", ed estensioni cl_khr_subgroup_ballot/shuffle/rotate/ecc. Se questo non c'è, FERMATI e segnalalo: significa che l'installazione Mesa è rotta, non è un problema di llama.cpp. ## Il problema `git clone --depth 1 https://github.com/ggml-org/llama.cpp`, poi compila con: cmake -B build -G Ninja -DCMAKE_BUILD_TYPE=Release -DGGML_OPENCL=ON \ -DGGML_OPENCL_EMBED_KERNELS=ON -DGGML_OPENCL_USE_ADRENO_KERNELS=OFF \ -DBUILD_SHARED_LIBS=ON cmake --build build -j$(nproc) La build C++ compila senza errori. Il problema emerge A RUNTIME (avviando `llama-server` con un modello reale), quando ggml compila i kernel OpenCL embedded via `clBuildProgram`. Errore tipico: input.cl:N:M: warning: OpenCL extension 'cl_qcom_reqd_sub_group_size' unknown or does not require pragma - ignoring input.cl:N:M: warning: unknown attribute 'qcom_reqd_sub_group_size' ignored input.cl:N:M: error: loading directly from pointer to type '__global half' requires cl_khr_fp16. Use vector data load builtin functions instead Error executing LLVM compilation action. ## Perché succede (già investigato, non ripartire da zero) I kernel sorgente sono in `ggml/src/ggml-opencl/kernels/*.cl`. Ne esistono 94 che dichiarano puntatori a `half`/`half2`/`half4`/`half8` (verificato con `grep -lE '\b(half|half2|half4|half8)\s*\*' ggml/src/ggml-opencl/kernels/*.cl`). Mesa/rusticl (a differenza del driver Adreno proprietario per cui questi kernel sono stati scritti) RIFIUTA il dereference diretto di puntatori half (`*ptr` o `ptr[i]`), richiedendo esplicitamente le funzioni builtin OpenCL `vload_half()` / `vload_half2()` / `vload_half4()` / `vload_half8()` / `vstore_half()` ecc., anche con `#pragma OPENCL EXTENSION cl_khr_fp16 : enable` attivo. Equivalenza semantica (per lo scalare; esistono le controparti vettoriali vload_half2/4/8 per i tipi vettore): *P -> vload_half(0, P) P[E] -> vload_half(E, P) Per half2/half4/half8 la firma è identica ma il valore restituito è un vettore: `vload_half2(offset, p)` legge 2 half consecutivi da `p+offset*2` e li converte in `float2`. ATTENZIONE alla differenza di offset in elementi half vs elementi vettore - verificare sempre contro la documentazione OpenCL ufficiale (khronos.org) prima di applicare, non indovinare. NON è un problema del fix Mesa dei subgroup, è indipendente: riguarda TUTTI i kernel che leggono direttamente valori half, a prescindere dal fix subgroup. ### Cosa è STATO GIÀ provato e scartato - `-DGGML_OPENCL_USE_ADRENO_KERNELS=OFF`: NON risolve il problema. Contrariamente a un vecchio commento di un maintainer nella PR https://github.com/ggml-org/llama.cpp/pull/10693 (che diceva "i kernel generici dovrebbero funzionare"), nella versione attuale di llama.cpp la famiglia di kernel `gemv_noshuffle_*`/`gemm_noshuffle_*` viene comunque compilata sempre, indipendentemente da questo flag (verificato nel log di build: questi kernel compaiono nella lista "opencl: embedding kernel ..." a prescindere dal valore del flag). Quel consiglio è ormai superato. - Forzare `-D cl_qcom_reqd_sub_group_size` nelle opzioni di compilazione kernel (per attivare il ramo `ADRENO_GPU` nei kernel che ne hanno bisogno per macro come N_SIMDGROUP/N_DST/N_SIMDWIDTH): questo FUNZIONA e va mantenuto, ma da solo non basta - risolve solo l'errore "undeclared identifier", non l'errore fp16 successivo. - Ho già un piccolo script Python che fa la sostituzione automatica *P -> vload_half(0,P) / P[E] -> vload_half(E,P) SOLO per i file `*flat*.cl` con puntatori `half` scalari (14 file, 27 righe modificate, verificate manualmente una per una col diff). Funziona, ma copre solo una piccola parte dei 94 file totali e solo il caso scalare. ## Cosa devi fare 1. Estendi (o riscrivi da zero, se più pulito) lo script di patch per coprire TUTTI i 94 file, gestendo correttamente half/half2/half4/half8. Usa la stessa metodologia rigorosa già validata: per ogni modifica, genera il diff e verificalo a mano leggendo il codice circostante prima di fidarti — l'obiettivo è zero sostituzioni sbagliate silenziose (offset vettoriali errati = pesi del modello letti male = output numerico sbagliato ma nessun crash, il peggior tipo di bug qui). 2. Compila e, per OGNI kernel che fallisce, applica la patch, ricompila. Hai il vantaggio che noi non avevamo: accesso reale all'hardware. Usa `RUSTICL_ENABLE=freedreno ./llama-server -hf <un modello piccolo, es. Qwen2.5-0.5B-Instruct-GGUF:Q4_K_M> --no-warmup -ngl 99` per iterare velocemente (modello piccolo = build/reload rapidi). Prova anche un secondo modello con un formato di quantizzazione diverso (es. uno IQ4_NL e uno Q4_K/Q6_K) per essere sicuro di coprire più famiglie di kernel. 3. IMPORTANTE - verifica la CORRETTEZZA NUMERICA, non solo che compili e non crashi: fai generare al modello una risposta a un prompt semplice e verificabile (es. "What is 2+2? Answer with just the number.") e controlla che la risposta sia sensata, non NaN/garbage/ripetizioni insensate. Un kernel con offset vload_half sbagliato spesso compila ed esegue senza errori, ma produce output numericamente corrotto - l'unico modo per scoprirlo è guardare l'output reale, non i log di compilazione. 4. Confronta se possibile l'output con lo stesso identico prompt/seed su CPU pura (`-ngl 0`) per lo stesso modello, per avere un termine di paragone di sanità (non serve che sia identico token-per-token, ma deve essere semanticamente coerente allo stesso modo). 5. Testa anche un modello MoE reale se possibile (es. lo stesso Qwen A3B o Gemma usato nei test precedenti) dato che la famiglia di kernel `gemm_moe_*`/`gemv_moe_*` è specifica per il routing MoE e potrebbe avere pattern di lettura half diversi dai kernel mul_mv standard. 6. Scrivi un file `FIX.md` (nella stessa cartella di questo TODO.md) con: - L'elenco esatto dei file modificati e il diff di ciascuno (o un unico patch/diff file allegato, se più pratico). - Quali modelli/formati di quantizzazione hai testato con successo e quali no. - Il valore CMake finale raccomandato per `GGML_OPENCL_USE_ADRENO_KERNELS` (ON o OFF) alla luce di quanto trovato. - Eventuali ulteriori patch necessarie oltre a queste (es. se scopri che serve ANCHE differenziare qualcosa per modelli MoE). - Se possibile, un confronto tokens/sec CPU (-ngl 0, -t 4) vs GPU (-ngl 99) sullo stesso modello/prompt, utile per capire se ne vale la pena rispetto a girare su CPU. Non serve che tu conosca o modifichi il Containerfile del container Podman - lavora direttamente su un checkout locale di llama.cpp sul DevKit. Il FIX.md che produci verrà usato da un'altra sessione per applicare le modifiche corrispondenti al Containerfile. ``` --- ## Riferimento: stato attuale del Containerfile (per chi applicherà FIX.md) `llamacpp-opencl.Containerfile` contiene già, in ordine: - **PATCH 1**: riconosce `FD6*` come famiglia Adreno (necessaria sempre). - **PATCH 2**: accetta `cl_khr_subgroup_ballot` come prova di supporto subgroup, dato che Mesa non espone mai la stringa legacy `cl_khr_subgroups` (necessaria sempre). - **PATCH 3**: forza `-D cl_qcom_reqd_sub_group_size` nelle opzioni di compilazione kernel (necessaria se si riattivano kernel che usano `REQD_SUBGROUP_SIZE_64`/`ADRENO_GPU`). - **PATCH 4**: `vload_half()` per i 14 file `*flat*.cl` con puntatori half scalari (copertura parziale, superata da questo lavoro). - `-DGGML_OPENCL_USE_ADRENO_KERNELS=OFF` (da rivalutare in base a FIX.md). Il pacchetto Mesa custom (`.deb` in `containers/llamacpp/deb/`, tracciato via Git LFS) è generato da `build-mesa.sh` nella stessa cartella e NON deve essere toccato per questo lavoro - il problema è tutto lato llama.cpp/kernel OpenCL, non lato driver.

devbadxyz skomentował(-a) 2026-07-09 21:01:15 +00:00

Autor

Członek

Skopiuj link

Copy Source

FIX.md — llama.cpp OpenCL su Adreno 690 / Mesa rusticl (Windows DevKit 2023)

Risultato del lavoro di debug fatto direttamente sul DevKit reale (Snapdragon
8cx Gen 3, Adreno 690, Mesa/rusticl custom con fix subgroup Freedreno a6xx,
commit 185c89084ab7). Testato su llama.cpp @ 049326a00 (2026-07-09 e
2026-07-10).

🟢 SECONDO AGGIORNAMENTO (2026-07-10): bug reale trovato e fixato — modelli

densi "normali" con n_embd/dimensioni interne > 1024 falliscono sempre

Un secondo giro di debug (partito da un crash reale riscontrato dall'utente
nel container con Qwen3.6-27B-MTP-GGUF:Q2_K_XL) ha portato alla scoperta
di un bug concreto e ben isolato, distinto da tutto quanto scritto sotto:
in 9 punti di ggml-opencl.cpp, il calcolo del local_work_size per un
dispatch dei kernel confrontava nth/cols/ecc. solo contro
CL_KERNEL_WORK_GROUP_SIZE (il limite di prodotto totale del work-group,
2048 su questo device) ma mai contro CL_DEVICE_MAX_WORK_ITEM_SIZES[0]
(il limite per-dimensione, 1024 su questo device — verificato via
clinfo: "Max work item sizes: 1024x1024x64", "Max work group size: 2048").
Questi due limiti sono indipendenti nello standard OpenCL e vanno rispettati
entrambi; su device dove il limite per-dimensione è più stretto del limite
di prodotto (come questo), il codice esistente calcolava local_work_size
fino a 1536+ per modelli con n_embd (o dimensioni analoghe) sopra 1024,
violando il limite per-dimensione con CL_INVALID_WORK_ITEM_SIZE (-55).

Sintomo: qualunque modello con n_embd (o la dimensione rilevante per
quel kernel) superiore a 1024 falliva SEMPRE al caricamento, anche modelli
completamente "normali" senza MoE/SSM/altro — es. Qwen2.5-1.5B-Instruct
(n_embd=1536), fallito in modo riproducibile su kernel_rms_norm_mul con
lws=[1536,1,1]. Spiega perché nella sessione precedente solo il modello
minuscolo Qwen2.5-0.5B (n_embd<1024) risultava funzionante: non era una
questione di "modelli piccoli vs grandi" in generale, ma di questa soglia
precisa.

Fix: aggiunto un campo backend_ctx->max_work_item_size0 (popolato una
volta all'init del device via clGetDeviceInfo(..., CL_DEVICE_MAX_WORK_ITEM_SIZES, ...))
e usato per clampare max_workgroup_size con MIN(...) in tutti i 9 punti
di dispatch interessati: kernel_rms_norm_mul, kernel_norm_mul_add,
kernel_l2_norm_f32, kernel_get_rows_*, kernel_set_rows/ggml_cl_set,
ggml_cl_add_id, kernel_argsort_f32_i32 (nel supports_op), kernel_group_norm_mul_add,
kernel_cumsum_blk. Altri 3 punti che usano lo stesso helper
(get_kernel_workgroup_size) sono stati verificati come già sicuri (limitano
esplicitamente a 64/256 indipendentemente dal device, quindi non toccano mai
il limite di 1024) e lasciati invariati.

Verificato funzionante (boot pulito, testato due volte separatamente):
Qwen2.5-0.5B e Qwen2.5-1.5B-Instruct caricano ed eseguono correttamente
sulla GPU dopo il fix (il secondo falliva sempre prima). Diff completo nel
file opencl-mesa-rusticl.patch aggiornato accanto a questo file.

⚠️ Questo fix è indipendente e complementare ai problemi ancora aperti
descritti sotto (MoE, SSM/MTP) — risolve una classe di crash diversa e più
generale (colpisce anche modelli densi "normali" abbastanza grandi), ma
NON risolve i problemi MoE/SSM descritti più sotto, che restano aperti.

🔴 AGGIORNAMENTO CRITICO (2026-07-10): i modelli MoE mandano in hang la GPU

Dopo aver scritto la prima versione di questo file (sezioni sotto, ancora
valide per i modelli densi), ho testato dei modelli MoE reali
(granite-3.0-1b-a400m, 84MB di soli pesi offloadati, e
Qwen3.6-35B-A3B-UD-IQ4_NL, 18GB) e in ENTRAMBI i casi il processo
llama-server crasha con CL_OUT_OF_RESOURCES (-5) sulla primissima
scrittura OpenCL della sessione (anche un tensore F32 di poche KB tipo
output_norm.weight), indipendentemente da -ngl, --ctx-size, --fit,
presenza del componente vision, o dimensione del modello.

Non è un limite di risorse software: ho verificato nel sorgente Mesa
(src/gallium/frontends/rusticl/core/queue.rs) che CL_OUT_OF_RESOURCES in
questo punto corrisponde esattamente a
pipe.device_reset_status() != PIPE_NO_RESET — cioè rusticl sta segnalando
che la GPU ha subito un vero reset hardware. Confermato anche via
dmesg (accesso root):

msm_dpu ae01000.display-controller: [drm:hangcheck_handler [msm]] *ERROR* hangcheck detected gpu lockup rb 0!
msm_dpu ae01000.display-controller: [drm:recover_worker [msm]] *ERROR* hangcheck recover!
msm_dpu ae01000.display-controller: [drm:recover_worker [msm]] *ERROR* offending task: ... (comando llama-server con il modello MoE in uso)

19 hang consecutivi registrati durante questa sessione, uno per ogni processo
llama-server lanciato con un modello MoE. Un devcoredump completo del
crash (/sys/class/devcoredump/) è stato salvato in
.crash_dumps/devcd6_moe_crash_*.data nel checkout locale usato per questo
lavoro (non incluso qui, contiene lo stream di comandi GPU codificato
ascii85 — servirebbe crashdec di Mesa, non buildato per mancanza di tempo,
per decodificarlo fino al comando esatto che ha causato il fault).

Più grave: la recovery del driver (hangcheck recover!) NON restituisce
sempre un stato pulito. Dopo alcuni hang causati da modelli MoE, ANCHE il
modello denso Qwen2.5-0.5B (che aveva funzionato perfettamente in precedenza
nella stessa sessione) ha iniziato a fallire allo stesso modo — la GPU resta
degradata anche per carichi di lavoro diversi da quello che ha causato il
primo hang. Questo probabilmente spiega anche i riavvii completi del
sistema (non solo del processo) osservati nella prima parte di questa
sessione, prima che i 3 fix qui sotto fossero applicati: se la recovery
del driver a volte fallisce del tutto invece di limitarsi a un hangcheck
recover, il risultato può essere un crash dell'intero kernel/sistema.

Root cause esatta non confermata (richiederebbe crashdec per decodificare
il command stream GPU), ma l'ipotesi più probabile: con
GGML_OPENCL_USE_ADRENO_KERNELS=OFF, i tensori MoE ("*_exps.weight",
dimensione 3D con l'asse esperti) passano dal path di conversione SOA_Q
generico (non quello Adreno-specifico _trans4_ns, escluso dal flag) —
questo path potrebbe non gestire correttamente la dimensione "esperti"
nell'indicizzazione del kernel di conversione, causando un accesso GPU fuori
dai limiti che il driver Mesa/Freedreno non degrada in modo pulito.

Raccomandazione pratica:

1. Riavviare il DevKit prima di qualunque altro test — la GPU risulta
   attualmente degradata da questa sessione di debug.
2. Non usare modelli MoE con questo backend OpenCL per ora — solo i
   modelli densi (testati: Qwen2.5-0.5B Q4_K_M) sono confermati sicuri e
   funzionanti. Il container aggiornato con le patch sotto va bene per
   modelli densi; per MoE serve indagine ulteriore (idealmente con
   crashdec per identificare il kernel esatto, o testando
   ADRENO_KERNELS=ON nonostante i suoi problemi noti di compilazione sui
   kernel _ns, per vedere se il path Adreno-specifico per i tensori MoE
   evita il fault).
3. I fix in questo file restano validi e consigliati indipendentemente da
   questo problema (sono richiesti anche solo per far funzionare i modelli
   densi, e non sono la causa degli hang MoE - il codice coinvolto negli
   hang è diverso, nel path di conversione SOA_Q per tensori quantizzati).

🔴 TERZO PROBLEMA APERTO (2026-07-10): modelli ibridi SSM/MTP (famiglia

Qwen3.5/Qwen3.6) mandano in hang la GPU, indipendentemente dalla dimensione

Dopo aver applicato il fix del work-item-size sopra, ho testato
Qwen3.5-2B-UD-Q2_K_XL (1GB, --no-mmproj, niente MTP attivo) aspettandomi
che funzionasse dato che è piccolo — invece ha causato un hang GPU reale
identico a quello dei modelli MoE (clEnqueueWriteBuffer → CL_OUT_OF_RESOURCES
sulla primissima scrittura, output_norm.weight, prima ancora che qualunque
kernel venga lanciato — confermato con instrumentazione che stampa il nome di
ogni tensore appena prima della set_tensor: una sola riga stampata,
poi il crash). Stesso identico comportamento per Qwen3.6-27B-UD-Q2_K_XL
(12GB), sia con -ngl 0 che con -ngl 8.

Conferma che non è una questione di dimensione: 1GB e 12GB falliscono
allo stesso identico modo, nello stesso punto. Deve essere qualcosa di
specifico all'architettura.

Ipotesi (non confermata): entrambi i modelli usano un'architettura
ibrida SSM (State Space Model, tipo Mamba) + attention — il file GGUF
contiene tensori blk.N.ssm_a, ssm_alpha.weight, ssm_beta.weight,
ssm_conv1d.weight, ssm_dt.bias, ssm_norm.weight, ssm_out.weight per
ogni layer, oltre a una testa MTP (blk.N.nextn.* — multi-token
prediction/speculative decoding integrata nel modello). Nessuno degli altri
modelli testati con successo (Qwen2.5, granite pre-fix) ha tensori SSM.
Il kernel ssm_conv (e il correlato gated_delta_net, altra architettura
SSM-like) sono nella lista dei kernel sempre compilati
(ggml/src/ggml-opencl/CMakeLists.txt) — non ancora verificato se il bug è
lì dentro o altrove nella gestione C++ di questi tensori.

Non ancora investigato: quale specifica chiamata OpenCL (probabilmente
async, dato che il fallimento emerge solo alla primissima scrittura
bloccante successiva) avvelena la coda GPU. Servirebbe la stessa
metodologia usata per il bug MoE (instrumentazione mirata + eventualmente
crashdec per il devcoredump) ma applicata a questo caso specifico — non
fatto per esaurimento tempo in questa sessione.

Nota operativa importante: dopo un hang GPU, il sistema resta
degradato — ho riconfermato che ANCHE configurazioni prima funzionanti
(Qwen2.5-1.5B con il fix già applicato) hanno ripreso a fallire subito dopo
aver innescato questi hang SSM/MTP, fino al riavvio del DevKit. Se si
riproduce questo problema, riavviare prima di trarre conclusioni su
qualunque altro test.

Non è un problema generale di "modelli grandi" o di altre famiglie:
testato anche gemma-4-E2B-it (Gemma 3n, 2.3GB, nessun tensore SSM/MTP,
n_embd=1536 — stessa taglia di dimensione che falliva per il bug
work-item-size) su un boot pulito subito dopo i due hang Qwen3.5/3.6: ha
dato un errore completamente diverso, pulito (non un hang), vedi sezione
dedicata subito sotto. Questo conferma che l'hang SSM/MTP è specifico alla
famiglia Qwen3.5/3.6 (o più precisamente ai tensori SSM/nextn), non un
problema architetturale generico che colpisce ogni modello "insolito".

🟡 QUARTO PROBLEMA (2026-07-10, RISOLTO CON WORKAROUND): Gemma 3n e

flash-attention — CL_INVALID_KERNEL_ARGS, non un hang GPU

Testando gemma-4-E2B-it-UD-Q2_K_XL (Gemma 3n, 2.3GB, nessun tensore
SSM/MTP) con il fix work-item-size già applicato, il caricamento falliva
con un errore diverso da tutti i precedenti:

DEBUG enqueue failed err=-54 kernel=flash_attn_f32_f16 tensor=node_235 op=FLASH_ATTN_EXT

-54 è CL_INVALID_KERNEL_ARGS (non tutti gli argomenti del kernel sono
stati impostati correttamente, o uno è invalido) — un rifiuto pulito del
driver, confermato NON essere un hang GPU (contatore hangcheck rimasto
a 0 sia prima che dopo). Causa non ancora identificata nel dettaglio (il
kernel flash-attention ha moltissime varianti compilate a runtime in base
a DK/DV/BLOCK_M/BLOCK_N, plausibile un bug di clSetKernelArg
mancante o non allineato per la specifica combinazione di head-dimension
usata da Gemma 3n).

Workaround verificato: avviando con --flash-attn off (o
LLAMA_ARG_FLASH_ATTN=off), Gemma carica e genera correttamente
("2+2" → "4"), nessun hang, nessun crash. Root cause nel kernel FA non
ancora investigata a fondo (fuori scope per questa sessione), ma il
workaround è sufficiente per usare il modello subito.

TL;DR (modelli densi, sezione originale)

• -DGGML_OPENCL_USE_ADRENO_KERNELS=OFF confermato corretto — non è
  superato, resta la scelta giusta (motivo diverso da quello originariamente
  ipotizzato nella PR #10693, vedi sotto).
• La PATCH 4 esistente nel Containerfile (vload_half sui soli file
  *flat*.cl) è superata: quei file non vengono nemmeno compilati con
  ADRENO_KERNELS=OFF (vedi analisi CMakeLists sotto), quindi la patch era
  innocua ma inutile. Va rimossa e sostituita dalle 3 patch mirate qui
  sotto, che coprono i file realmente compilati ed eseguiti.
• 3 bug reali trovati e fixati, verificati sia a compile-time
  (clBuildProgram su tutti i 150 kernel) sia a runtime (inferenza reale su
  GPU, output numericamente corretto) — ma solo per modelli densi, vedi
  aggiornamento critico sopra per i modelli MoE.
• ⚠️ ATTENZIONE HARDWARE: eseguire kernel OpenCL non fixati (che
  compilano ma leggono memoria male, o che il driver rifiuta a metà) su
  questo Mesa/rusticl ancora sperimentale può causare riavvio dell'intero
  DevKit, non solo crash del processo — successo più volte durante questo
  lavoro. Vedi sezione "Nota sicurezza hardware" più sotto E l'aggiornamento
  critico in cima al file.

Perché GGML_OPENCL_USE_ADRENO_KERNELS=OFF è la scelta giusta

Analisi di ggml/src/ggml-opencl/CMakeLists.txt: la flag NON controlla
quali kernel .cl vengono compilati ed embeddati (con
GGML_OPENCL_EMBED_KERNELS=ON) — la lista GGML_OPENCL_KERNELS include
SEMPRE tutti i kernel _flat, _noshuffle, _moe/_ns, indipendentemente
dal flag. L'unica eccezione è gemm_xmem_f16_f32_os8, aggiunto alla lista
solo se ADRENO_KERNELS=ON (CMakeLists.txt righe 213-215).

Quello che il flag controlla davvero è quale codice C++ in
ggml-opencl.cpp viene compilato (#ifdef GGML_OPENCL_USE_ADRENO_KERNELS):
un blocco enorme (righe ~3199-4123) che copre TUTTA la build/dispatch di
gemv_noshuffle_*, gemm_noshuffle_*, gemm_moe_*_ns, gemv_moe_*_ns,
gemm_xmem_f16_f32_os8. Con ADRENO_KERNELS=OFF questo blocco è escluso
dalla compilazione: i file .cl corrispondenti restano innocuamente
embeddati come stringhe header ma non vengono mai passati a
clBuildProgram, quindi non vengono mai eseguiti sul device.

Questo è rilevante perché 7 kernel gemm_moe_*_ns.cl + 1
gemm_xmem_f16_f32_os8.cl non compilano affatto su Mesa/rusticl (vedi
sotto) — con ADRENO_KERNELS=ON questo causerebbe un crash fatale
all'avvio del server (o al primo uso, per i moe) nel momento in cui
clBuildProgram fallisce con fatal=true. Con OFF il problema
semplicemente non si presenta, perché quel codice non viene mai eseguito.

Le famiglie di kernel "standard" (mul_mv_id_* per il routing MoE,
mul_mm_*_l4_lm, mul_mv_*) restano invece SEMPRE compilate/usate
(non gated dal flag), e sono quelle effettivamente coinvolte nell'inferenza
reale — inclusi i modelli MoE, che usano mul_mv_id_* per il routing, non
i kernel _ns esclusivi Adreno.

I 3 bug di compilazione fixati (sessione 2026-07-09; il 4° bug, quello

sul work-item-size, è descritto sopra ed è di natura diversa - runtime,

non compilazione)

Metodologia: harness standalone in C (clCreateProgramWithSource +
clBuildProgram diretto su ogni singolo file .cl, con le stesse opzioni
di compilazione usate realmente da ggml-opencl.cpp, incluse le macro -D
dinamiche iniettate per-kernel dove applicabile — es. SIMDGROUP_WIDTH,
DK/DV/BLOCK_M/BLOCK_N per flash-attention) per isolare rapidamente
quali dei 150 kernel falliscono davvero, senza dover riavviare
llama-server (lento, e rischioso vedi sotto) per ogni iterazione.

Su 150 file, 22 fallivano la compilazione. Di questi, 19 erano falsi
positivi del test harness (mancavano le macro -D dinamiche che il C++
reale inietta a runtime — verificato ricompilando con le macro corrette
prese da ggml-opencl.cpp e dalla tabella di tuning fa_tune.h): tutti i 6
flash_attn_*.cl e tutti i 4 gemv_noshuffle_*.cl compilano puliti con le
macro reali.

3 bug erano reali, tutti nel set di kernel sempre compilato
(indipendente da ADRENO_KERNELS):

1. kernels/mul_mv_q4_k_f32.cl — dereference diretto di puntatore half

-            float dall = dh[0];
-            float dmin = dh[1];
+            float dall = vload_half(0, dh);
+            float dmin = vload_half(1, dh);

dh è global half *. Stesso pattern già noto e documentato nella PATCH 4
esistente (Mesa/rusticl richiede vload_half() esplicito, il driver
proprietario Adreno permette il dereference diretto). Semanticamente
equivalente (vload_half(offset, p) == *(p+offset) convertito a float).

2. kernels/mul_mm_q4_k_f32_l4_lm.cl — address space implicito non valido

-                char * scales = src0_s + ib * 12;
+                global uchar * scales = src0_s + ib * 12;

src0_s è dichiarato global uchar *. Il codice assegnava il puntatore a
una variabile locale char * senza qualificatore di address space (quindi
default private), cambiando implicitamente lo address space da global a
private — cosa che OpenCL C standard vieta e che Mesa/rusticl segnala come
errore ("changes address space of pointer"), mentre il compiler Adreno
proprietario evidentemente lo tollera. Non è un bug indotto dalle patch
Mesa: è un'inconsistenza nel sorgente llama.cpp stesso — il file gemello
mul_mm_q5_k_f32_l4_lm.cl ha lo stesso identico pattern ma con il
qualificatore corretto (global uchar * scales = ...), confermando che si
tratta di una svista upstream, non di un problema di portabilità.

3. kernels/mul_mv_f16_f32_l4.cl — builtin proprietario Qualcomm applicato per errore

-#ifdef ADRENO_GPU
+#if defined(ADRENO_GPU) && defined(cl_qcom_subgroup_shuffle)
 #pragma OPENCL EXTENSION cl_qcom_subgroup_shuffle : enable
 #define sub_group_shuffle_xor(val, mask) qcom_sub_group_shuffle_xor((val), (mask), CLK_SUB_GROUP_SHUFFLE_WIDTH_WAVE_SIZE_QCOM, 0.0f)
+#endif
 
+#ifdef ADRENO_GPU
 REQD_SUBGROUP_SIZE_64

Il file macro-ridefinisce sub_group_shuffle_xor (la funzione builtin
standard di cl_khr_subgroup_shuffle, già usata correttamente altrove nel
file) con la variante proprietaria qcom_sub_group_shuffle_xor, ma lo fa
sotto #ifdef ADRENO_GPU — che noi definiamo forzatamente con la PATCH 3
esistente (-D cl_qcom_reqd_sub_group_size) per attivare i rami di codice
che impostano N_SIMDGROUP/N_DST/ecc. Su hardware Adreno reale con
driver proprietario, ADRENO_GPU definito implica sempre anche il supporto
di cl_qcom_subgroup_shuffle; noi però forziamo solo la prima macro, non
la seconda, che Mesa infatti non implementa (verificato via clinfo: solo
cl_khr_subgroup_shuffle standard è esposto, non l'estensione proprietaria
cl_qcom_subgroup_shuffle). Fix: la macro-redefinizione ora richiede
esplicitamente che l'estensione proprietaria sia davvero disponibile, così
su Mesa il codice ricade sulla funzione standard sub_group_shuffle_xor
già corretta e già usata nel resto del file.

Questo file è nel set di kernel sempre compilato, indipendentemente da
ADRENO_KERNELS — senza questo fix, llama-server crasha già in fase di
caricamento kernel (load_cl_kernels()), prima ancora di caricare un
modello.

Bug noti ma NON fixati (fuori scope, irrilevanti con ADRENO_KERNELS=OFF)

Questi kernel falliscono la compilazione su Mesa/rusticl per motivi reali,
non falsi positivi — ma dato che ADRENO_KERNELS=OFF, il codice C++ che li
compilerebbe/userebbe è escluso a compile-time, quindi restano innocui
(file .cl embeddato ma mai passato a clBuildProgram):

• gemm_moe_{q4_0,q4_1,q5_0,q5_1,q4_k,q5_k,q6_k}_f32_ns.cl (7 file):
  usano float32 come tipo di variabile (__private float32 reg_c = ...)
  e swizzle vettoriali estesi fino a .sv (32 componenti). Questo non è
  OpenCL C standard (il massimo vettore standard è float16, swizzle fino
  a .sf) — è quasi certamente un'estensione proprietaria del compiler
  Qualcomm per vettori più larghi. Un fix richiederebbe riscrivere ogni
  kernel per usare due registri float16 (.lo/.hi) al posto di un
  singolo float32, con ~40 righe da toccare per file e un rimappaggio
  attento degli indici s0..sv → lo.s0..sf / hi.s0..sf. Non tentato:
  rischio alto di bug numerici silenziosi per il beneficio marginale (con
  ADRENO_KERNELS=OFF il routing MoE passa comunque da mul_mv_id_*, che
  compila ed esegue correttamente, vedi test granite sotto).
• gemm_moe_mxfp4_f32_ns.cl: stesso problema float32 più un mismatch
  di rank tra half8 e float scalare in 4 punti — probabilmente
  correlato allo stesso pattern di refactor necessario sopra.
• gemm_xmem_f16_f32_os8.cl: usa builtin Qualcomm genuinamente
  proprietari senza equivalente standard (qcom_get_physical_sub_group_id,
  qcom_sub_group_constant_load8, qcom_sub_group_sync,
  QCOM_CLK_CONST_LOAD_SYNC). Non fixabile senza il driver proprietario;
  questo è l'unico file la cui embedding dipende davvero da
  ADRENO_KERNELS (vedi sopra), quindi con OFF è semplicemente escluso.

Se in futuro si vuole tentare ADRENO_KERNELS=ON per inseguire le
performance dei kernel "_ns"/flat ottimizzati, questi bug vanno risolti
prima — altrimenti il server crasha al boot (kernel sempre compilati
eagerly in load_cl_kernels_adreno).

Nota sicurezza hardware (importante per chi ripete il lavoro)

Durante questo lavoro il DevKit si è riavviato più volte (crash-loop con
boot da pochi minuti l'uno) esattamente nel momento in cui venivano
eseguiti kernel OpenCL (non solo compilati) contenenti i bug sopra — in
particolare il bug #1 (vload_half mancante) causa una lettura di memoria
non valida quando il kernel viene davvero eseguito sull'hardware, non
solo compilato. Su un driver Mesa/Freedreno ancora sperimentale come
questo, un accesso di memoria non valido da parte della GPU può bloccare
il driver kernel-space abbastanza da far scattare un watchdog reset
dell'intero sistema, non solo un crash del processo llama-server.

Compilare (clBuildProgram) senza mai eseguire (clEnqueueNDRangeKernel)
è invece risultato sempre sicuro nei test — utile saperlo se si vuole
verificare rapidamente nuovi kernel senza rischiare un riavvio.
Consiglio: dopo aver applicato queste patch e PRIMA di lanciare
llama-server -ngl 99 con un modello nuovo/quantizzazione nuova, verificare
prima con un harness di solo-compile (vedi .clcheck.c non incluso qui ma
banale da riscrivere: clCreateProgramWithSource + clBuildProgram con le
stesse opzioni di ggml-opencl.cpp) che tutti i kernel coinvolti
compilino, prima di eseguire davvero.

Modelli testati

Modello	Quant	Architettura	GPU (-ngl 99)	Note
Qwen2.5-0.5B-Instruct	Q4_K_M	densa, n_embd<1024	✅ output corretto ("2+2"→"4")	mul_mv_q4_k_f32 + mul_mm_q4_k_f32_l4_lm + mul_mv_f16_f32_l4 tutti esercitati
Qwen2.5-1.5B-Instruct	Q4_0	densa, n_embd=1536	✅ output corretto ("2+2"→"4"), dopo fix work-item-size	falliva sempre prima del fix (CL_INVALID_WORK_ITEM_SIZE)
granite-3.0-1b-a400m	Q4_K_M	MoE piccola	❌ hang GPU reale	vedi sezione MoE sopra, causa non identificata
Qwen3.6-35B-A3B	UD-IQ4_NL	MoE grande + VL	❌ hang GPU reale	vedi sezione MoE sopra
Qwen3.5-2B-MTP	UD-Q2_K_XL	SSM+MTP piccola	❌ hang GPU reale	vedi sezione SSM/MTP sopra, causa non identificata
Qwen3.6-27B-MTP	UD-Q2_K_XL	SSM+MTP grande	❌ hang GPU reale	stesso problema di Qwen3.5-2B-MTP, confermato non è questione di dimensione
gemma-4-E2B-it	UD-Q2_K_XL	densa Gemma 3n, n_embd=1536	✅ output corretto ("2+2"→"4") con --flash-attn off	senza il flag: CL_INVALID_KERNEL_ARGS pulito su flash_attn_f32_f16, NON un hang GPU — problema diverso e minore dal bug SSM/MTP

Performance (Qwen2.5-0.5B Q4_K_M, stesso prompt "What is 2+2?", CPU vs GPU):

• GPU: prompt 152.6 tok/s, generazione 20.3 tok/s
• CPU (4 thread): prompt 140.6 tok/s, generazione 73.4 tok/s
• (su un modello così piccolo la CPU vince in generazione: overhead di
  dispatch OpenCL dominante — non rappresentativo per modelli più grandi,
  ma non è stato possibile completare un confronto su un modello denso
  medio/grande dato che gli unici modelli grandi disponibili nella cache
  locale sono tutti MoE o SSM/MTP, entrambi ancora non funzionanti su GPU)

⚠️ Su questo modello minuscolo (0.5B) la CPU genera più veloce della
GPU — atteso: l'overhead di dispatch OpenCL domina su un modello così
piccolo. Non è indicativo delle prestazioni relative su modelli più grandi
(dove la GPU dovrebbe vincere nettamente), ma non sono riuscito a
completare il test con un modello più grande/MoE/IQ4_NL in questa sessione
prima di dover consegnare questo file — vedi "Cosa manca" sotto.

Cosa manca (da fare nella prossima sessione)

• Un secondo formato di quantizzazione denso (Q6_K, IQ4_NL) — non ho
  trovato rapidamente un modello piccolo con quant IQ4_NL disponibile su HF
  (bartowski et al. tendono a pubblicare IQ4_XS invece); i repo con
  IQ4_NL trovati erano quasi tutti modelli 7-8B, più lenti da iterare.
• Modello MoE: testato (granite-3.0-1b-a400m e Qwen3.6-35B-A3B),
  risultato: hang GPU reale, vedi "AGGIORNAMENTO CRITICO" in cima al file.
  Serve crashdec (tool Mesa in src/freedreno/decode/, non buildato per
  tempo) per identificare il comando GPU esatto che causa il fault, poi
  capire se è nel path di conversione SOA_Q generico per tensori 3D
  (ipotesi principale) o altrove.
• Confronto tokens/sec CPU vs GPU su un modello di dimensione più
  realistica (es. 3-8B denso), dove il vantaggio della GPU dovrebbe essere
  visibile (sul modello 0.5B testato la CPU vince, ma non è rappresentativo).
  Non fatto per esaurimento tempo dopo l'indagine sugli hang MoE.

Modifiche da riportare nel Containerfile

Sostituire la PATCH 4 esistente (script Python vload_half sui file
*flat*.cl) con il diff completo qui sotto (identico al contenuto di
opencl-mesa-rusticl.patch accanto a questo file, applicabile con
git apply direttamente sul checkout in /build/llama.cpp). Include sia
i 3 fix di compilazione della sessione 2026-07-09 sia il fix del
work-item-size della sessione 2026-07-10. Le PATCH 1/2/3 esistenti nel
Containerfile restano necessarie e corrette così come sono (il diff qui
sotto le ri-applica comunque, essendo un diff completo su ggml-opencl.cpp).

diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp
index 5c96b9a9f..6e1b19c89 100644
--- a/ggml/src/ggml-opencl/ggml-opencl.cpp
+++ b/ggml/src/ggml-opencl/ggml-opencl.cpp
@@ -529,6 +529,7 @@ struct ggml_backend_opencl_context {
     size_t global_mem_size;
     size_t max_alloc_size;
     size_t max_workgroup_size;
+    size_t max_work_item_size0;
     bool fp16_support;
     bool has_vector_subgroup_broadcast;
     bool has_subgroup_shuffle = false;       // cl_khr_subgroup_shuffle or cl_qcom_subgroup_shuffle
@@ -1160,7 +1161,8 @@ static void load_cl_kernels_argsort(ggml_backend_opencl_context *backend_ctx) {
         std::string("CL") + std::to_string(backend_ctx->opencl_c_version.major) + "." + std::to_string(backend_ctx->opencl_c_version.minor);
     std::string compile_opts = std::string("-cl-std=") + opencl_c_std +
                                " -cl-mad-enable -cl-unsafe-math-optimizations"
-                               " -cl-finite-math-only -cl-fast-relaxed-math";
+                               " -cl-finite-math-only -cl-fast-relaxed-math"
+                               " -D cl_qcom_reqd_sub_group_size";
 
     // argsort
     if (!backend_ctx->kernels_loaded_argsort) {
@@ -1203,7 +1205,8 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) {
         std::string("CL") + std::to_string(backend_ctx->opencl_c_version.major) + "." + std::to_string(backend_ctx->opencl_c_version.minor);
     std::string compile_opts = std::string("-cl-std=") + opencl_c_std +
                                " -cl-mad-enable -cl-unsafe-math-optimizations"
-                               " -cl-finite-math-only -cl-fast-relaxed-math";
+                               " -cl-finite-math-only -cl-fast-relaxed-math"
+                               " -D cl_qcom_reqd_sub_group_size";
 
     if (backend_ctx->adreno_use_large_buffer) {
         compile_opts += " -qcom-enable-large-buffer ";
@@ -5343,6 +5346,7 @@ static bool ggml_opencl_is_device_supported(ggml_backend_dev_t dev) {
     GGML_ASSERT(dev_ctx->device);
 
     if (strstr(dev_ctx->device_name.c_str(), "Adreno") ||
+        strstr(dev_ctx->device_name.c_str(), "FD6") ||
         strstr(dev_ctx->device_name.c_str(), "Qualcomm") ||
         strstr(dev_ctx->device_version.c_str(), "Adreno")) {
         dev_ctx->gpu_family = GPU_FAMILY::ADRENO;
@@ -5394,7 +5398,8 @@ static bool ggml_opencl_is_device_supported(ggml_backend_dev_t dev) {
     // If OpenCL 3.0 is supported, then check for cl_khr_subgroups, which becomes
     // optional in OpenCL 3.0 (cl_khr_subgroup is mandatory in OpenCL 2.x)
     if (opencl_c_version.major == 3 && strstr(ext_buffer, "cl_khr_subgroups") == NULL &&
-        strstr(ext_buffer, "cl_intel_subgroups") == NULL) {
+        strstr(ext_buffer, "cl_intel_subgroups") == NULL &&
+        strstr(ext_buffer, "cl_khr_subgroup_ballot") == NULL) {
         GGML_LOG_WARN("ggml_opencl: device does not support subgroups (cl_khr_subgroups or cl_intel_subgroups) "
             "(note that subgroups is an optional feature in OpenCL 3.0)\n");
         return false;
@@ -5502,6 +5507,15 @@ static ggml_backend_opencl_context * ggml_cl_init(ggml_backend_dev_t dev) {
     CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof(size_t), &backend_ctx->image2d_max_width, NULL));
     CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof(size_t), &backend_ctx->image2d_max_height, NULL));
     CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), &backend_ctx->max_workgroup_size, NULL));
+    {
+        // CL_DEVICE_MAX_WORK_ITEM_SIZES is a per-dimension cap that can be
+        // stricter than CL_DEVICE_MAX_WORK_GROUP_SIZE (the overall product
+        // cap) - e.g. this Mesa/rusticl Adreno device reports 1024 here but
+        // 2048 for MAX_WORK_GROUP_SIZE. A 1-D dispatch must respect both.
+        size_t max_work_item_sizes[3] = {0, 0, 0};
+        CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(max_work_item_sizes), max_work_item_sizes, NULL));
+        backend_ctx->max_work_item_size0 = max_work_item_sizes[0];
+    }
     CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_SVM_CAPABILITIES, sizeof(cl_device_svm_capabilities), &backend_ctx->svm_caps, 0));
 
     if (opencl_c_version.major >= 3) {
@@ -6748,7 +6762,7 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te
             load_cl_kernels_argsort(backend_ctx);
 
             cl_kernel kernel = backend_ctx->kernel_argsort_f32_i32;
-            int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel);
+            int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0);
 
             int cols = 1;
             while (cols < op->ne[0]) {
@@ -10491,7 +10505,7 @@ static void ggml_cl_get_rows(ggml_backend_t backend, const ggml_tensor * src0, c
     CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_ulong), &nb2));
     CL_CHECK(clSetKernelArg(kernel, 16, sizeof(cl_ulong), &nb3));
 
-    int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel);
+    int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0);
     int nth = 1;
     while (nth < ne00 && 2*nth <= max_workgroup_size) {
         nth *= 2;
@@ -10685,7 +10699,7 @@ static void ggml_cl_set_rows(ggml_backend_t backend, const ggml_tensor * src0, c
         nth0 = 64;
     }
 
-    int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel);
+    int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0);
     while (nth0 < nblk0 && nth0 < max_workgroup_size) {
         nth0 *= 2;
     }
@@ -10945,7 +10959,7 @@ static void ggml_cl_add_id(ggml_backend_t backend, const ggml_tensor * src0, con
     CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int),      &ne0));
     CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int),      &ne1));
 
-    int nth = MIN(ne00, (int) backend_ctx->get_kernel_workgroup_size(kernel));
+    int nth = MIN(ne00, MIN((int) backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0));
     size_t global_work_size[] = { (size_t)ne01*nth, (size_t)ne02, 1 };
     size_t local_work_size[] = { (size_t)nth, 1, 1 };
 
@@ -12095,7 +12109,7 @@ static void ggml_opencl_op_rms_norm_fused(ggml_backend_t backend, ggml_tensor *
     cl_kernel kernel = backend_ctx->kernel_rms_norm_mul;
 
     int nth = sgs;
-    int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel);
+    int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0);
     while (nth < ne00 && nth < max_workgroup_size) {
         nth *= 2;
     }
@@ -12173,7 +12187,7 @@ static void ggml_opencl_op_norm_fused(ggml_backend_t backend, ggml_tensor * norm
     cl_kernel kernel = backend_ctx->kernel_norm_mul_add;
 
     int nth = sgs;
-    int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel);
+    int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0);
     while (nth < ne00/4 && nth < max_workgroup_size) nth *= 2;
     nth = MIN(nth, max_workgroup_size);
     nth = MIN(nth, ne00/4);
@@ -12246,7 +12260,7 @@ static void ggml_opencl_op_group_norm_fused(ggml_backend_t backend, ggml_tensor
     memcpy(&eps, (char *)gn_tensor->op_params + sizeof(int), sizeof(float));
 
     cl_kernel kernel = backend_ctx->kernel_group_norm_mul_add;
-    int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel);
+    int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0);
     int ne = ggml_nelements(src0);
     int group_size = ne / groups;
 
@@ -12352,7 +12366,8 @@ static void ggml_cl_l2_norm(ggml_backend_t backend, const ggml_tensor * src0, co
     cl_kernel kernel = backend_ctx->kernel_l2_norm_f32;
 
     int nth = sgs;
-    while (nth < ne00 && nth < (int)backend_ctx->get_kernel_workgroup_size(kernel)) {
+    int max_workgroup_size_l2n = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0);
+    while (nth < ne00 && nth < max_workgroup_size_l2n) {
         nth *= 2;
     }
 
@@ -21401,7 +21416,7 @@ static void ggml_cl_set(ggml_backend_t backend, const ggml_tensor * src0, const
     CL_CHECK(clSetKernelArg(kernel, 18, sizeof(cl_ulong), &pnb2));
     CL_CHECK(clSetKernelArg(kernel, 19, sizeof(cl_ulong), &pnb3));
 
-    int max_local_size = backend_ctx->get_kernel_workgroup_size(kernel);
+    int max_local_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0);
 
     const int nth = MIN(max_local_size, ne00);
 
@@ -22112,7 +22127,7 @@ static void ggml_cl_cumsum(ggml_backend_t backend, const ggml_tensor * src0, con
 
     cl_kernel kernel = backend_ctx->kernel_cumsum_blk;
 
-    int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel);
+    int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0);
     int nth = 1;
     while (nth < ne00 && 2*nth <= max_workgroup_size) {
         nth *= 2;
diff --git a/ggml/src/ggml-opencl/kernels/mul_mm_q4_k_f32_l4_lm.cl b/ggml/src/ggml-opencl/kernels/mul_mm_q4_k_f32_l4_lm.cl
index 2235b1ae8..dc191b1d2 100644
--- a/ggml/src/ggml-opencl/kernels/mul_mm_q4_k_f32_l4_lm.cl
+++ b/ggml/src/ggml-opencl/kernels/mul_mm_q4_k_f32_l4_lm.cl
@@ -90,7 +90,7 @@ kernel void kernel_mul_mm_q4_k_f32_l4_lm(
                 int is = 2 * n + b;
                 int qsi = n * 32 + (iqs % 16) * 2;
 
-                char * scales = src0_s + ib * 12;
+                global uchar * scales = src0_s + ib * 12;
 
                 int scidx0 = (is < 4) ? is : (is + 4);
                 int scidx1 = (is < 4) ? is : (is - 4);
diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl b/ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl
index da2e14ae9..b293c73d6 100644
--- a/ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl
+++ b/ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl
@@ -180,10 +180,12 @@ kernel void kernel_mul_mat_f16_f32_l4_dr(
 // Kernels for decoding, Adreno only for now
 #define MUL_MAT_F16_F32_L4_DR_LS_R2_MAX 8
 
-#ifdef ADRENO_GPU
+#if defined(ADRENO_GPU) && defined(cl_qcom_subgroup_shuffle)
 #pragma OPENCL EXTENSION cl_qcom_subgroup_shuffle : enable
 #define sub_group_shuffle_xor(val, mask) qcom_sub_group_shuffle_xor((val), (mask), CLK_SUB_GROUP_SHUFFLE_WIDTH_WAVE_SIZE_QCOM, 0.0f)
+#endif
 
+#ifdef ADRENO_GPU
 REQD_SUBGROUP_SIZE_64
 kernel void kernel_mul_mat_f16_f32_l4_dr_ls(
         global char * src0,
diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_q4_k_f32.cl b/ggml/src/ggml-opencl/kernels/mul_mv_q4_k_f32.cl
index 71ab98982..64cff077f 100644
--- a/ggml/src/ggml-opencl/kernels/mul_mv_q4_k_f32.cl
+++ b/ggml/src/ggml-opencl/kernels/mul_mv_q4_k_f32.cl
@@ -151,8 +151,8 @@ kernel void kernel_mul_mv_q4_K_f32(
                 acc2.s3 += yh[i+9] * (q2[i/2] & 0xF000);
             }
 
-            float dall = dh[0];
-            float dmin = dh[1];
+            float dall = vload_half(0, dh);
+            float dmin = vload_half(1, dh);
             sumf[row] += dall * ((acc1.s0 + 1.f/256.f * acc1.s1) * sc8[0] +
                                  (acc1.s2 + 1.f/256.f * acc1.s3) * sc8[1] * 1.f/16.f +
                                  (acc2.s0 + 1.f/256.f * acc2.s1) * sc8[4] +

# FIX.md — llama.cpp OpenCL su Adreno 690 / Mesa rusticl (Windows DevKit 2023) Risultato del lavoro di debug fatto direttamente sul DevKit reale (Snapdragon 8cx Gen 3, Adreno 690, Mesa/rusticl custom con fix subgroup Freedreno a6xx, commit 185c89084ab7). Testato su `llama.cpp` @ `049326a00` (2026-07-09 e 2026-07-10). ## 🟢 SECONDO AGGIORNAMENTO (2026-07-10): bug reale trovato e fixato — modelli ## densi "normali" con `n_embd`/dimensioni interne > 1024 falliscono sempre Un secondo giro di debug (partito da un crash reale riscontrato dall'utente nel container con `Qwen3.6-27B-MTP-GGUF:Q2_K_XL`) ha portato alla scoperta di un **bug concreto e ben isolato**, distinto da tutto quanto scritto sotto: in 9 punti di `ggml-opencl.cpp`, il calcolo del `local_work_size` per un dispatch dei kernel confrontava `nth`/`cols`/ecc. solo contro `CL_KERNEL_WORK_GROUP_SIZE` (il limite di prodotto totale del work-group, 2048 su questo device) ma **mai contro `CL_DEVICE_MAX_WORK_ITEM_SIZES[0]`** (il limite per-dimensione, **1024** su questo device — verificato via `clinfo`: "Max work item sizes: 1024x1024x64", "Max work group size: 2048"). Questi due limiti sono indipendenti nello standard OpenCL e vanno rispettati entrambi; su device dove il limite per-dimensione è più stretto del limite di prodotto (come questo), il codice esistente calcolava local_work_size fino a 1536+ per modelli con `n_embd` (o dimensioni analoghe) sopra 1024, violando il limite per-dimensione con `CL_INVALID_WORK_ITEM_SIZE` (-55). **Sintomo**: qualunque modello con `n_embd` (o la dimensione rilevante per quel kernel) superiore a 1024 falliva SEMPRE al caricamento, anche modelli completamente "normali" senza MoE/SSM/altro — es. **Qwen2.5-1.5B-Instruct** (`n_embd`=1536), fallito in modo riproducibile su `kernel_rms_norm_mul` con `lws=[1536,1,1]`. Spiega perché nella sessione precedente solo il modello minuscolo Qwen2.5-0.5B (`n_embd`<1024) risultava funzionante: non era una questione di "modelli piccoli vs grandi" in generale, ma di questa soglia precisa. **Fix**: aggiunto un campo `backend_ctx->max_work_item_size0` (popolato una volta all'init del device via `clGetDeviceInfo(..., CL_DEVICE_MAX_WORK_ITEM_SIZES, ...)`) e usato per clampare `max_workgroup_size` con `MIN(...)` in tutti i 9 punti di dispatch interessati: `kernel_rms_norm_mul`, `kernel_norm_mul_add`, `kernel_l2_norm_f32`, `kernel_get_rows_*`, `kernel_set_rows`/`ggml_cl_set`, `ggml_cl_add_id`, `kernel_argsort_f32_i32` (nel supports_op), `kernel_group_norm_mul_add`, `kernel_cumsum_blk`. Altri 3 punti che usano lo stesso helper (`get_kernel_workgroup_size`) sono stati verificati come già sicuri (limitano esplicitamente a 64/256 indipendentemente dal device, quindi non toccano mai il limite di 1024) e lasciati invariati. **Verificato funzionante** (boot pulito, testato due volte separatamente): `Qwen2.5-0.5B` e `Qwen2.5-1.5B-Instruct` caricano ed eseguono correttamente sulla GPU dopo il fix (il secondo falliva sempre prima). Diff completo nel file `opencl-mesa-rusticl.patch` aggiornato accanto a questo file. ⚠️ Questo fix è indipendente e complementare ai problemi ancora aperti descritti sotto (MoE, SSM/MTP) — risolve una classe di crash diversa e più generale (colpisce anche modelli densi "normali" abbastanza grandi), ma NON risolve i problemi MoE/SSM descritti più sotto, che restano aperti. ## 🔴 AGGIORNAMENTO CRITICO (2026-07-10): i modelli MoE mandano in hang la GPU Dopo aver scritto la prima versione di questo file (sezioni sotto, ancora valide per i modelli densi), ho testato dei modelli MoE reali (`granite-3.0-1b-a400m`, 84MB di soli pesi offloadati, e `Qwen3.6-35B-A3B-UD-IQ4_NL`, 18GB) e in ENTRAMBI i casi il processo `llama-server` crasha con `CL_OUT_OF_RESOURCES` (-5) sulla primissima scrittura OpenCL della sessione (anche un tensore F32 di poche KB tipo `output_norm.weight`), indipendentemente da `-ngl`, `--ctx-size`, `--fit`, presenza del componente vision, o dimensione del modello. **Non è un limite di risorse software**: ho verificato nel sorgente Mesa (`src/gallium/frontends/rusticl/core/queue.rs`) che `CL_OUT_OF_RESOURCES` in questo punto corrisponde esattamente a `pipe.device_reset_status() != PIPE_NO_RESET` — cioè rusticl sta segnalando che **la GPU ha subito un vero reset hardware**. Confermato anche via `dmesg` (accesso root): ``` msm_dpu ae01000.display-controller: [drm:hangcheck_handler [msm]] *ERROR* hangcheck detected gpu lockup rb 0! msm_dpu ae01000.display-controller: [drm:recover_worker [msm]] *ERROR* hangcheck recover! msm_dpu ae01000.display-controller: [drm:recover_worker [msm]] *ERROR* offending task: ... (comando llama-server con il modello MoE in uso) ``` 19 hang consecutivi registrati durante questa sessione, uno per ogni processo `llama-server` lanciato con un modello MoE. Un devcoredump completo del crash (`/sys/class/devcoredump/`) è stato salvato in `.crash_dumps/devcd6_moe_crash_*.data` nel checkout locale usato per questo lavoro (non incluso qui, contiene lo stream di comandi GPU codificato ascii85 — servirebbe `crashdec` di Mesa, non buildato per mancanza di tempo, per decodificarlo fino al comando esatto che ha causato il fault). **Più grave**: la recovery del driver (`hangcheck recover!`) NON restituisce sempre un stato pulito. Dopo alcuni hang causati da modelli MoE, ANCHE il modello denso Qwen2.5-0.5B (che aveva funzionato perfettamente in precedenza nella stessa sessione) ha iniziato a fallire allo stesso modo — la GPU resta degradata anche per carichi di lavoro diversi da quello che ha causato il primo hang. **Questo probabilmente spiega anche i riavvii completi del sistema (non solo del processo) osservati nella prima parte di questa sessione**, prima che i 3 fix qui sotto fossero applicati: se la recovery del driver a volte fallisce del tutto invece di limitarsi a un hangcheck recover, il risultato può essere un crash dell'intero kernel/sistema. Root cause esatta non confermata (richiederebbe `crashdec` per decodificare il command stream GPU), ma l'ipotesi più probabile: con `GGML_OPENCL_USE_ADRENO_KERNELS=OFF`, i tensori MoE ("`*_exps.weight`", dimensione 3D con l'asse esperti) passano dal path di conversione SOA_Q *generico* (non quello Adreno-specifico `_trans4_ns`, escluso dal flag) — questo path potrebbe non gestire correttamente la dimensione "esperti" nell'indicizzazione del kernel di conversione, causando un accesso GPU fuori dai limiti che il driver Mesa/Freedreno non degrada in modo pulito. **Raccomandazione pratica**: 1. **Riavviare il DevKit prima di qualunque altro test** — la GPU risulta attualmente degradata da questa sessione di debug. 2. **Non usare modelli MoE con questo backend OpenCL per ora** — solo i modelli densi (testati: Qwen2.5-0.5B Q4_K_M) sono confermati sicuri e funzionanti. Il container aggiornato con le patch sotto va bene per modelli densi; per MoE serve indagine ulteriore (idealmente con `crashdec` per identificare il kernel esatto, o testando `ADRENO_KERNELS=ON` nonostante i suoi problemi noti di compilazione sui kernel `_ns`, per vedere se il path Adreno-specifico per i tensori MoE evita il fault). 3. I fix in questo file restano validi e consigliati **indipendentemente** da questo problema (sono richiesti anche solo per far funzionare i modelli densi, e non sono la causa degli hang MoE - il codice coinvolto negli hang è diverso, nel path di conversione SOA_Q per tensori quantizzati). ## 🔴 TERZO PROBLEMA APERTO (2026-07-10): modelli ibridi SSM/MTP (famiglia ## Qwen3.5/Qwen3.6) mandano in hang la GPU, indipendentemente dalla dimensione Dopo aver applicato il fix del work-item-size sopra, ho testato `Qwen3.5-2B-UD-Q2_K_XL` (1GB, `--no-mmproj`, niente MTP attivo) aspettandomi che funzionasse dato che è piccolo — invece ha causato un hang GPU reale identico a quello dei modelli MoE (`clEnqueueWriteBuffer` → `CL_OUT_OF_RESOURCES` sulla primissima scrittura, `output_norm.weight`, prima ancora che qualunque kernel venga lanciato — confermato con instrumentazione che stampa il nome di ogni tensore appena prima della `set_tensor`: **una sola riga stampata**, poi il crash). Stesso identico comportamento per `Qwen3.6-27B-UD-Q2_K_XL` (12GB), sia con `-ngl 0` che con `-ngl 8`. **Conferma che non è una questione di dimensione**: 1GB e 12GB falliscono allo stesso identico modo, nello stesso punto. Deve essere qualcosa di specifico all'architettura. **Ipotesi (non confermata)**: entrambi i modelli usano un'architettura ibrida SSM (State Space Model, tipo Mamba) + attention — il file GGUF contiene tensori `blk.N.ssm_a`, `ssm_alpha.weight`, `ssm_beta.weight`, `ssm_conv1d.weight`, `ssm_dt.bias`, `ssm_norm.weight`, `ssm_out.weight` per ogni layer, oltre a una testa MTP (`blk.N.nextn.*` — multi-token prediction/speculative decoding integrata nel modello). Nessuno degli altri modelli testati con successo (Qwen2.5, granite pre-fix) ha tensori SSM. Il kernel `ssm_conv` (e il correlato `gated_delta_net`, altra architettura SSM-like) sono nella lista dei kernel sempre compilati (`ggml/src/ggml-opencl/CMakeLists.txt`) — non ancora verificato se il bug è lì dentro o altrove nella gestione C++ di questi tensori. **Non ancora investigato**: quale specifica chiamata OpenCL (probabilmente async, dato che il fallimento emerge solo alla primissima scrittura bloccante successiva) avvelena la coda GPU. Servirebbe la stessa metodologia usata per il bug MoE (instrumentazione mirata + eventualmente `crashdec` per il devcoredump) ma applicata a questo caso specifico — non fatto per esaurimento tempo in questa sessione. **Nota operativa importante**: dopo un hang GPU, il sistema resta degradato — ho riconfermato che ANCHE configurazioni prima funzionanti (Qwen2.5-1.5B con il fix già applicato) hanno ripreso a fallire subito dopo aver innescato questi hang SSM/MTP, fino al riavvio del DevKit. Se si riproduce questo problema, riavviare prima di trarre conclusioni su qualunque altro test. **Non è un problema generale di "modelli grandi" o di altre famiglie**: testato anche `gemma-4-E2B-it` (Gemma 3n, 2.3GB, nessun tensore SSM/MTP, `n_embd`=1536 — stessa taglia di dimensione che falliva per il bug work-item-size) su un boot pulito subito dopo i due hang Qwen3.5/3.6: ha dato un errore completamente diverso, pulito (non un hang), vedi sezione dedicata subito sotto. Questo conferma che l'hang SSM/MTP è specifico alla famiglia Qwen3.5/3.6 (o più precisamente ai tensori SSM/nextn), non un problema architetturale generico che colpisce ogni modello "insolito". ## 🟡 QUARTO PROBLEMA (2026-07-10, RISOLTO CON WORKAROUND): Gemma 3n e ## flash-attention — `CL_INVALID_KERNEL_ARGS`, non un hang GPU Testando `gemma-4-E2B-it-UD-Q2_K_XL` (Gemma 3n, 2.3GB, nessun tensore SSM/MTP) con il fix work-item-size già applicato, il caricamento falliva con un errore diverso da tutti i precedenti: ``` DEBUG enqueue failed err=-54 kernel=flash_attn_f32_f16 tensor=node_235 op=FLASH_ATTN_EXT ``` `-54` è `CL_INVALID_KERNEL_ARGS` (non tutti gli argomenti del kernel sono stati impostati correttamente, o uno è invalido) — un rifiuto pulito del driver, **confermato NON essere un hang GPU** (contatore hangcheck rimasto a 0 sia prima che dopo). Causa non ancora identificata nel dettaglio (il kernel flash-attention ha moltissime varianti compilate a runtime in base a `DK`/`DV`/`BLOCK_M`/`BLOCK_N`, plausibile un bug di `clSetKernelArg` mancante o non allineato per la specifica combinazione di head-dimension usata da Gemma 3n). **Workaround verificato**: avviando con `--flash-attn off` (o `LLAMA_ARG_FLASH_ATTN=off`), Gemma carica e genera correttamente ("2+2" → "4"), nessun hang, nessun crash. Root cause nel kernel FA non ancora investigata a fondo (fuori scope per questa sessione), ma il workaround è sufficiente per usare il modello subito. ## TL;DR (modelli densi, sezione originale) - `-DGGML_OPENCL_USE_ADRENO_KERNELS=OFF` **confermato corretto** — non è superato, resta la scelta giusta (motivo diverso da quello originariamente ipotizzato nella PR #10693, vedi sotto). - La PATCH 4 esistente nel Containerfile (`vload_half` sui soli file `*flat*.cl`) è **superata**: quei file non vengono nemmeno compilati con `ADRENO_KERNELS=OFF` (vedi analisi CMakeLists sotto), quindi la patch era innocua ma inutile. Va **rimossa e sostituita** dalle 3 patch mirate qui sotto, che coprono i file realmente compilati ed eseguiti. - **3 bug reali trovati e fixati**, verificati sia a compile-time (`clBuildProgram` su tutti i 150 kernel) sia a runtime (inferenza reale su GPU, output numericamente corretto) — **ma solo per modelli densi, vedi aggiornamento critico sopra per i modelli MoE**. - ⚠️ **ATTENZIONE HARDWARE**: eseguire kernel OpenCL non fixati (che compilano ma leggono memoria male, o che il driver rifiuta a metà) su questo Mesa/rusticl ancora sperimentale può causare **riavvio dell'intero DevKit**, non solo crash del processo — successo più volte durante questo lavoro. Vedi sezione "Nota sicurezza hardware" più sotto E l'aggiornamento critico in cima al file. ## Perché `GGML_OPENCL_USE_ADRENO_KERNELS=OFF` è la scelta giusta Analisi di `ggml/src/ggml-opencl/CMakeLists.txt`: la flag **NON** controlla quali kernel `.cl` vengono compilati ed embeddati (con `GGML_OPENCL_EMBED_KERNELS=ON`) — la lista `GGML_OPENCL_KERNELS` include SEMPRE tutti i kernel `_flat`, `_noshuffle`, `_moe`/`_ns`, indipendentemente dal flag. L'unica eccezione è `gemm_xmem_f16_f32_os8`, aggiunto alla lista solo se `ADRENO_KERNELS=ON` (CMakeLists.txt righe 213-215). Quello che il flag controlla davvero è quale **codice C++** in `ggml-opencl.cpp` viene compilato (`#ifdef GGML_OPENCL_USE_ADRENO_KERNELS`): un blocco enorme (righe ~3199-4123) che copre TUTTA la build/dispatch di `gemv_noshuffle_*`, `gemm_noshuffle_*`, `gemm_moe_*_ns`, `gemv_moe_*_ns`, `gemm_xmem_f16_f32_os8`. Con `ADRENO_KERNELS=OFF` questo blocco è escluso dalla compilazione: i file `.cl` corrispondenti restano innocuamente embeddati come stringhe header ma **non vengono mai passati a `clBuildProgram`**, quindi non vengono mai eseguiti sul device. Questo è rilevante perché **7 kernel `gemm_moe_*_ns.cl` + 1 `gemm_xmem_f16_f32_os8.cl` non compilano affatto su Mesa/rusticl** (vedi sotto) — con `ADRENO_KERNELS=ON` questo causerebbe un crash fatale all'avvio del server (o al primo uso, per i moe) nel momento in cui `clBuildProgram` fallisce con `fatal=true`. Con `OFF` il problema semplicemente non si presenta, perché quel codice non viene mai eseguito. Le famiglie di kernel "standard" (`mul_mv_id_*` per il routing MoE, `mul_mm_*_l4_lm`, `mul_mv_*`) restano invece SEMPRE compilate/usate (non gated dal flag), e sono quelle effettivamente coinvolte nell'inferenza reale — inclusi i modelli MoE, che usano `mul_mv_id_*` per il routing, non i kernel `_ns` esclusivi Adreno. ## I 3 bug di compilazione fixati (sessione 2026-07-09; il 4° bug, quello ## sul work-item-size, è descritto sopra ed è di natura diversa - runtime, ## non compilazione) Metodologia: harness standalone in C (`clCreateProgramWithSource` + `clBuildProgram` diretto su ogni singolo file `.cl`, con le stesse opzioni di compilazione usate realmente da `ggml-opencl.cpp`, incluse le macro `-D` dinamiche iniettate per-kernel dove applicabile — es. `SIMDGROUP_WIDTH`, `DK`/`DV`/`BLOCK_M`/`BLOCK_N` per flash-attention) per isolare rapidamente quali dei 150 kernel falliscono davvero, senza dover riavviare `llama-server` (lento, e rischioso vedi sotto) per ogni iterazione. Su 150 file, 22 fallivano la compilazione. Di questi, **19 erano falsi positivi** del test harness (mancavano le macro `-D` dinamiche che il C++ reale inietta a runtime — verificato ricompilando con le macro corrette prese da `ggml-opencl.cpp` e dalla tabella di tuning `fa_tune.h`): tutti i 6 `flash_attn_*.cl` e tutti i 4 `gemv_noshuffle_*.cl` compilano puliti con le macro reali. **3 bug erano reali**, tutti nel set di kernel sempre compilato (indipendente da `ADRENO_KERNELS`): ### 1. `kernels/mul_mv_q4_k_f32.cl` — dereference diretto di puntatore half ```diff - float dall = dh[0]; - float dmin = dh[1]; + float dall = vload_half(0, dh); + float dmin = vload_half(1, dh); ``` `dh` è `global half *`. Stesso pattern già noto e documentato nella PATCH 4 esistente (Mesa/rusticl richiede `vload_half()` esplicito, il driver proprietario Adreno permette il dereference diretto). Semanticamente equivalente (`vload_half(offset, p) == *(p+offset)` convertito a float). ### 2. `kernels/mul_mm_q4_k_f32_l4_lm.cl` — address space implicito non valido ```diff - char * scales = src0_s + ib * 12; + global uchar * scales = src0_s + ib * 12; ``` `src0_s` è dichiarato `global uchar *`. Il codice assegnava il puntatore a una variabile locale `char *` senza qualificatore di address space (quindi default `private`), cambiando implicitamente lo address space da `global` a `private` — cosa che OpenCL C standard vieta e che Mesa/rusticl segnala come errore ("changes address space of pointer"), mentre il compiler Adreno proprietario evidentemente lo tollera. **Non è un bug indotto dalle patch Mesa**: è un'inconsistenza nel sorgente llama.cpp stesso — il file gemello `mul_mm_q5_k_f32_l4_lm.cl` ha lo stesso identico pattern ma con il qualificatore corretto (`global uchar * scales = ...`), confermando che si tratta di una svista upstream, non di un problema di portabilità. ### 3. `kernels/mul_mv_f16_f32_l4.cl` — builtin proprietario Qualcomm applicato per errore ```diff -#ifdef ADRENO_GPU +#if defined(ADRENO_GPU) && defined(cl_qcom_subgroup_shuffle) #pragma OPENCL EXTENSION cl_qcom_subgroup_shuffle : enable #define sub_group_shuffle_xor(val, mask) qcom_sub_group_shuffle_xor((val), (mask), CLK_SUB_GROUP_SHUFFLE_WIDTH_WAVE_SIZE_QCOM, 0.0f) +#endif +#ifdef ADRENO_GPU REQD_SUBGROUP_SIZE_64 ``` Il file macro-ridefinisce `sub_group_shuffle_xor` (la funzione builtin *standard* di `cl_khr_subgroup_shuffle`, già usata correttamente altrove nel file) con la variante proprietaria `qcom_sub_group_shuffle_xor`, ma lo fa sotto `#ifdef ADRENO_GPU` — che noi definiamo forzatamente con la PATCH 3 esistente (`-D cl_qcom_reqd_sub_group_size`) per attivare i rami di codice che impostano `N_SIMDGROUP`/`N_DST`/ecc. Su hardware Adreno reale con driver proprietario, `ADRENO_GPU` definito implica sempre anche il supporto di `cl_qcom_subgroup_shuffle`; noi però forziamo solo la prima macro, non la seconda, che Mesa infatti non implementa (verificato via `clinfo`: solo `cl_khr_subgroup_shuffle` standard è esposto, non l'estensione proprietaria `cl_qcom_subgroup_shuffle`). Fix: la macro-redefinizione ora richiede esplicitamente che l'estensione proprietaria sia davvero disponibile, così su Mesa il codice ricade sulla funzione standard `sub_group_shuffle_xor` già corretta e già usata nel resto del file. Questo file è nel set di kernel **sempre compilato**, indipendentemente da `ADRENO_KERNELS` — senza questo fix, `llama-server` crasha già in fase di caricamento kernel (`load_cl_kernels()`), prima ancora di caricare un modello. ## Bug noti ma NON fixati (fuori scope, irrilevanti con `ADRENO_KERNELS=OFF`) Questi kernel falliscono la compilazione su Mesa/rusticl per motivi reali, non falsi positivi — ma dato che `ADRENO_KERNELS=OFF`, il codice C++ che li compilerebbe/userebbe è escluso a compile-time, quindi restano innocui (file `.cl` embeddato ma mai passato a `clBuildProgram`): - **`gemm_moe_{q4_0,q4_1,q5_0,q5_1,q4_k,q5_k,q6_k}_f32_ns.cl`** (7 file): usano `float32` come tipo di variabile (`__private float32 reg_c = ...`) e swizzle vettoriali estesi fino a `.sv` (32 componenti). Questo non è OpenCL C standard (il massimo vettore standard è `float16`, swizzle fino a `.sf`) — è quasi certamente un'estensione proprietaria del compiler Qualcomm per vettori più larghi. Un fix richiederebbe riscrivere ogni kernel per usare due registri `float16` (`.lo`/`.hi`) al posto di un singolo `float32`, con ~40 righe da toccare per file e un rimappaggio attento degli indici `s0..sv` → `lo.s0..sf` / `hi.s0..sf`. Non tentato: rischio alto di bug numerici silenziosi per il beneficio marginale (con `ADRENO_KERNELS=OFF` il routing MoE passa comunque da `mul_mv_id_*`, che compila ed esegue correttamente, vedi test granite sotto). - **`gemm_moe_mxfp4_f32_ns.cl`**: stesso problema `float32` più un mismatch di rank tra `half8` e `float` scalare in 4 punti — probabilmente correlato allo stesso pattern di refactor necessario sopra. - **`gemm_xmem_f16_f32_os8.cl`**: usa builtin Qualcomm genuinamente proprietari senza equivalente standard (`qcom_get_physical_sub_group_id`, `qcom_sub_group_constant_load8`, `qcom_sub_group_sync`, `QCOM_CLK_CONST_LOAD_SYNC`). Non fixabile senza il driver proprietario; questo è l'unico file la cui *embedding* dipende davvero da `ADRENO_KERNELS` (vedi sopra), quindi con `OFF` è semplicemente escluso. Se in futuro si vuole tentare `ADRENO_KERNELS=ON` per inseguire le performance dei kernel "_ns"/flat ottimizzati, questi bug vanno risolti prima — altrimenti il server crasha al boot (kernel sempre compilati eagerly in `load_cl_kernels_adreno`). ## Nota sicurezza hardware (importante per chi ripete il lavoro) Durante questo lavoro il DevKit si è riavviato più volte (crash-loop con boot da pochi minuti l'uno) esattamente nel momento in cui venivano eseguiti kernel OpenCL (non solo compilati) contenenti i bug sopra — in particolare il bug #1 (`vload_half` mancante) causa una lettura di memoria non valida quando il kernel viene davvero *eseguito* sull'hardware, non solo compilato. Su un driver Mesa/Freedreno ancora sperimentale come questo, un accesso di memoria non valido da parte della GPU può bloccare il driver kernel-space abbastanza da far scattare un watchdog reset dell'intero sistema, non solo un crash del processo `llama-server`. Compilare (`clBuildProgram`) senza mai eseguire (`clEnqueueNDRangeKernel`) è invece risultato sempre sicuro nei test — utile saperlo se si vuole verificare rapidamente nuovi kernel senza rischiare un riavvio. **Consiglio**: dopo aver applicato queste patch e PRIMA di lanciare `llama-server -ngl 99` con un modello nuovo/quantizzazione nuova, verificare prima con un harness di solo-compile (vedi `.clcheck.c` non incluso qui ma banale da riscrivere: `clCreateProgramWithSource` + `clBuildProgram` con le stesse opzioni di `ggml-opencl.cpp`) che tutti i kernel coinvolti compilino, prima di eseguire davvero. ## Modelli testati | Modello | Quant | Architettura | GPU (-ngl 99) | Note | |---|---|---|---|---| | Qwen2.5-0.5B-Instruct | Q4_K_M | densa, n_embd<1024 | ✅ output corretto ("2+2"→"4") | mul_mv_q4_k_f32 + mul_mm_q4_k_f32_l4_lm + mul_mv_f16_f32_l4 tutti esercitati | | Qwen2.5-1.5B-Instruct | Q4_0 | densa, n_embd=1536 | ✅ output corretto ("2+2"→"4"), dopo fix work-item-size | falliva sempre prima del fix (CL_INVALID_WORK_ITEM_SIZE) | | granite-3.0-1b-a400m | Q4_K_M | MoE piccola | ❌ hang GPU reale | vedi sezione MoE sopra, causa non identificata | | Qwen3.6-35B-A3B | UD-IQ4_NL | MoE grande + VL | ❌ hang GPU reale | vedi sezione MoE sopra | | Qwen3.5-2B-MTP | UD-Q2_K_XL | SSM+MTP piccola | ❌ hang GPU reale | vedi sezione SSM/MTP sopra, causa non identificata | | Qwen3.6-27B-MTP | UD-Q2_K_XL | SSM+MTP grande | ❌ hang GPU reale | stesso problema di Qwen3.5-2B-MTP, confermato non è questione di dimensione | | gemma-4-E2B-it | UD-Q2_K_XL | densa Gemma 3n, n_embd=1536 | ✅ output corretto ("2+2"→"4") con `--flash-attn off` | senza il flag: CL_INVALID_KERNEL_ARGS pulito su flash_attn_f32_f16, NON un hang GPU — problema diverso e minore dal bug SSM/MTP | **Performance** (Qwen2.5-0.5B Q4_K_M, stesso prompt "What is 2+2?", CPU vs GPU): - GPU: prompt 152.6 tok/s, generazione 20.3 tok/s - CPU (4 thread): prompt 140.6 tok/s, generazione 73.4 tok/s - (su un modello così piccolo la CPU vince in generazione: overhead di dispatch OpenCL dominante — non rappresentativo per modelli più grandi, ma non è stato possibile completare un confronto su un modello denso medio/grande dato che gli unici modelli grandi disponibili nella cache locale sono tutti MoE o SSM/MTP, entrambi ancora non funzionanti su GPU) ⚠️ Su questo modello minuscolo (0.5B) la CPU genera **più veloce** della GPU — atteso: l'overhead di dispatch OpenCL domina su un modello così piccolo. Non è indicativo delle prestazioni relative su modelli più grandi (dove la GPU dovrebbe vincere nettamente), ma **non sono riuscito a completare il test con un modello più grande/MoE/IQ4_NL in questa sessione** prima di dover consegnare questo file — vedi "Cosa manca" sotto. ## Cosa manca (da fare nella prossima sessione) - Un secondo formato di quantizzazione denso (Q6_K, IQ4_NL) — non ho trovato rapidamente un modello piccolo con quant IQ4_NL disponibile su HF (bartowski et al. tendono a pubblicare IQ4_XS invece); i repo con IQ4_NL trovati erano quasi tutti modelli 7-8B, più lenti da iterare. - Modello MoE: **testato** (`granite-3.0-1b-a400m` e `Qwen3.6-35B-A3B`), risultato: hang GPU reale, vedi "AGGIORNAMENTO CRITICO" in cima al file. Serve `crashdec` (tool Mesa in `src/freedreno/decode/`, non buildato per tempo) per identificare il comando GPU esatto che causa il fault, poi capire se è nel path di conversione SOA_Q generico per tensori 3D (ipotesi principale) o altrove. - Confronto tokens/sec CPU vs GPU su un modello di dimensione più realistica (es. 3-8B denso), dove il vantaggio della GPU dovrebbe essere visibile (sul modello 0.5B testato la CPU vince, ma non è rappresentativo). Non fatto per esaurimento tempo dopo l'indagine sugli hang MoE. ## Modifiche da riportare nel Containerfile Sostituire la PATCH 4 esistente (script Python `vload_half` sui file `*flat*.cl`) con il diff completo qui sotto (identico al contenuto di `opencl-mesa-rusticl.patch` accanto a questo file, applicabile con `git apply` direttamente sul checkout in `/build/llama.cpp`). Include sia i 3 fix di compilazione della sessione 2026-07-09 sia il fix del work-item-size della sessione 2026-07-10. Le PATCH 1/2/3 esistenti nel Containerfile restano necessarie e corrette così come sono (il diff qui sotto le ri-applica comunque, essendo un diff completo su `ggml-opencl.cpp`). ```diff diff --git a/ggml/src/ggml-opencl/ggml-opencl.cpp b/ggml/src/ggml-opencl/ggml-opencl.cpp index 5c96b9a9f..6e1b19c89 100644 --- a/ggml/src/ggml-opencl/ggml-opencl.cpp +++ b/ggml/src/ggml-opencl/ggml-opencl.cpp @@ -529,6 +529,7 @@ struct ggml_backend_opencl_context { size_t global_mem_size; size_t max_alloc_size; size_t max_workgroup_size; + size_t max_work_item_size0; bool fp16_support; bool has_vector_subgroup_broadcast; bool has_subgroup_shuffle = false; // cl_khr_subgroup_shuffle or cl_qcom_subgroup_shuffle @@ -1160,7 +1161,8 @@ static void load_cl_kernels_argsort(ggml_backend_opencl_context *backend_ctx) { std::string("CL") + std::to_string(backend_ctx->opencl_c_version.major) + "." + std::to_string(backend_ctx->opencl_c_version.minor); std::string compile_opts = std::string("-cl-std=") + opencl_c_std + " -cl-mad-enable -cl-unsafe-math-optimizations" - " -cl-finite-math-only -cl-fast-relaxed-math"; + " -cl-finite-math-only -cl-fast-relaxed-math" + " -D cl_qcom_reqd_sub_group_size"; // argsort if (!backend_ctx->kernels_loaded_argsort) { @@ -1203,7 +1205,8 @@ static void load_cl_kernels(ggml_backend_opencl_context *backend_ctx) { std::string("CL") + std::to_string(backend_ctx->opencl_c_version.major) + "." + std::to_string(backend_ctx->opencl_c_version.minor); std::string compile_opts = std::string("-cl-std=") + opencl_c_std + " -cl-mad-enable -cl-unsafe-math-optimizations" - " -cl-finite-math-only -cl-fast-relaxed-math"; + " -cl-finite-math-only -cl-fast-relaxed-math" + " -D cl_qcom_reqd_sub_group_size"; if (backend_ctx->adreno_use_large_buffer) { compile_opts += " -qcom-enable-large-buffer "; @@ -5343,6 +5346,7 @@ static bool ggml_opencl_is_device_supported(ggml_backend_dev_t dev) { GGML_ASSERT(dev_ctx->device); if (strstr(dev_ctx->device_name.c_str(), "Adreno") || + strstr(dev_ctx->device_name.c_str(), "FD6") || strstr(dev_ctx->device_name.c_str(), "Qualcomm") || strstr(dev_ctx->device_version.c_str(), "Adreno")) { dev_ctx->gpu_family = GPU_FAMILY::ADRENO; @@ -5394,7 +5398,8 @@ static bool ggml_opencl_is_device_supported(ggml_backend_dev_t dev) { // If OpenCL 3.0 is supported, then check for cl_khr_subgroups, which becomes // optional in OpenCL 3.0 (cl_khr_subgroup is mandatory in OpenCL 2.x) if (opencl_c_version.major == 3 && strstr(ext_buffer, "cl_khr_subgroups") == NULL && - strstr(ext_buffer, "cl_intel_subgroups") == NULL) { + strstr(ext_buffer, "cl_intel_subgroups") == NULL && + strstr(ext_buffer, "cl_khr_subgroup_ballot") == NULL) { GGML_LOG_WARN("ggml_opencl: device does not support subgroups (cl_khr_subgroups or cl_intel_subgroups) " "(note that subgroups is an optional feature in OpenCL 3.0)\n"); return false; @@ -5502,6 +5507,15 @@ static ggml_backend_opencl_context * ggml_cl_init(ggml_backend_dev_t dev) { CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_WIDTH, sizeof(size_t), &backend_ctx->image2d_max_width, NULL)); CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_IMAGE2D_MAX_HEIGHT, sizeof(size_t), &backend_ctx->image2d_max_height, NULL)); CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(size_t), &backend_ctx->max_workgroup_size, NULL)); + { + // CL_DEVICE_MAX_WORK_ITEM_SIZES is a per-dimension cap that can be + // stricter than CL_DEVICE_MAX_WORK_GROUP_SIZE (the overall product + // cap) - e.g. this Mesa/rusticl Adreno device reports 1024 here but + // 2048 for MAX_WORK_GROUP_SIZE. A 1-D dispatch must respect both. + size_t max_work_item_sizes[3] = {0, 0, 0}; + CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(max_work_item_sizes), max_work_item_sizes, NULL)); + backend_ctx->max_work_item_size0 = max_work_item_sizes[0]; + } CL_CHECK(clGetDeviceInfo(device, CL_DEVICE_SVM_CAPABILITIES, sizeof(cl_device_svm_capabilities), &backend_ctx->svm_caps, 0)); if (opencl_c_version.major >= 3) { @@ -6748,7 +6762,7 @@ static bool ggml_opencl_supports_op(ggml_backend_dev_t dev, const struct ggml_te load_cl_kernels_argsort(backend_ctx); cl_kernel kernel = backend_ctx->kernel_argsort_f32_i32; - int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel); + int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); int cols = 1; while (cols < op->ne[0]) { @@ -10491,7 +10505,7 @@ static void ggml_cl_get_rows(ggml_backend_t backend, const ggml_tensor * src0, c CL_CHECK(clSetKernelArg(kernel, 15, sizeof(cl_ulong), &nb2)); CL_CHECK(clSetKernelArg(kernel, 16, sizeof(cl_ulong), &nb3)); - int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel); + int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); int nth = 1; while (nth < ne00 && 2*nth <= max_workgroup_size) { nth *= 2; @@ -10685,7 +10699,7 @@ static void ggml_cl_set_rows(ggml_backend_t backend, const ggml_tensor * src0, c nth0 = 64; } - int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel); + int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); while (nth0 < nblk0 && nth0 < max_workgroup_size) { nth0 *= 2; } @@ -10945,7 +10959,7 @@ static void ggml_cl_add_id(ggml_backend_t backend, const ggml_tensor * src0, con CL_CHECK(clSetKernelArg(kernel, 12, sizeof(int), &ne0)); CL_CHECK(clSetKernelArg(kernel, 13, sizeof(int), &ne1)); - int nth = MIN(ne00, (int) backend_ctx->get_kernel_workgroup_size(kernel)); + int nth = MIN(ne00, MIN((int) backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0)); size_t global_work_size[] = { (size_t)ne01*nth, (size_t)ne02, 1 }; size_t local_work_size[] = { (size_t)nth, 1, 1 }; @@ -12095,7 +12109,7 @@ static void ggml_opencl_op_rms_norm_fused(ggml_backend_t backend, ggml_tensor * cl_kernel kernel = backend_ctx->kernel_rms_norm_mul; int nth = sgs; - int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel); + int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); while (nth < ne00 && nth < max_workgroup_size) { nth *= 2; } @@ -12173,7 +12187,7 @@ static void ggml_opencl_op_norm_fused(ggml_backend_t backend, ggml_tensor * norm cl_kernel kernel = backend_ctx->kernel_norm_mul_add; int nth = sgs; - int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel); + int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); while (nth < ne00/4 && nth < max_workgroup_size) nth *= 2; nth = MIN(nth, max_workgroup_size); nth = MIN(nth, ne00/4); @@ -12246,7 +12260,7 @@ static void ggml_opencl_op_group_norm_fused(ggml_backend_t backend, ggml_tensor memcpy(&eps, (char *)gn_tensor->op_params + sizeof(int), sizeof(float)); cl_kernel kernel = backend_ctx->kernel_group_norm_mul_add; - int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel); + int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); int ne = ggml_nelements(src0); int group_size = ne / groups; @@ -12352,7 +12366,8 @@ static void ggml_cl_l2_norm(ggml_backend_t backend, const ggml_tensor * src0, co cl_kernel kernel = backend_ctx->kernel_l2_norm_f32; int nth = sgs; - while (nth < ne00 && nth < (int)backend_ctx->get_kernel_workgroup_size(kernel)) { + int max_workgroup_size_l2n = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); + while (nth < ne00 && nth < max_workgroup_size_l2n) { nth *= 2; } @@ -21401,7 +21416,7 @@ static void ggml_cl_set(ggml_backend_t backend, const ggml_tensor * src0, const CL_CHECK(clSetKernelArg(kernel, 18, sizeof(cl_ulong), &pnb2)); CL_CHECK(clSetKernelArg(kernel, 19, sizeof(cl_ulong), &pnb3)); - int max_local_size = backend_ctx->get_kernel_workgroup_size(kernel); + int max_local_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); const int nth = MIN(max_local_size, ne00); @@ -22112,7 +22127,7 @@ static void ggml_cl_cumsum(ggml_backend_t backend, const ggml_tensor * src0, con cl_kernel kernel = backend_ctx->kernel_cumsum_blk; - int max_workgroup_size = backend_ctx->get_kernel_workgroup_size(kernel); + int max_workgroup_size = MIN(backend_ctx->get_kernel_workgroup_size(kernel), (int)backend_ctx->max_work_item_size0); int nth = 1; while (nth < ne00 && 2*nth <= max_workgroup_size) { nth *= 2; diff --git a/ggml/src/ggml-opencl/kernels/mul_mm_q4_k_f32_l4_lm.cl b/ggml/src/ggml-opencl/kernels/mul_mm_q4_k_f32_l4_lm.cl index 2235b1ae8..dc191b1d2 100644 --- a/ggml/src/ggml-opencl/kernels/mul_mm_q4_k_f32_l4_lm.cl +++ b/ggml/src/ggml-opencl/kernels/mul_mm_q4_k_f32_l4_lm.cl @@ -90,7 +90,7 @@ kernel void kernel_mul_mm_q4_k_f32_l4_lm( int is = 2 * n + b; int qsi = n * 32 + (iqs % 16) * 2; - char * scales = src0_s + ib * 12; + global uchar * scales = src0_s + ib * 12; int scidx0 = (is < 4) ? is : (is + 4); int scidx1 = (is < 4) ? is : (is - 4); diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl b/ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl index da2e14ae9..b293c73d6 100644 --- a/ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl +++ b/ggml/src/ggml-opencl/kernels/mul_mv_f16_f32_l4.cl @@ -180,10 +180,12 @@ kernel void kernel_mul_mat_f16_f32_l4_dr( // Kernels for decoding, Adreno only for now #define MUL_MAT_F16_F32_L4_DR_LS_R2_MAX 8 -#ifdef ADRENO_GPU +#if defined(ADRENO_GPU) && defined(cl_qcom_subgroup_shuffle) #pragma OPENCL EXTENSION cl_qcom_subgroup_shuffle : enable #define sub_group_shuffle_xor(val, mask) qcom_sub_group_shuffle_xor((val), (mask), CLK_SUB_GROUP_SHUFFLE_WIDTH_WAVE_SIZE_QCOM, 0.0f) +#endif +#ifdef ADRENO_GPU REQD_SUBGROUP_SIZE_64 kernel void kernel_mul_mat_f16_f32_l4_dr_ls( global char * src0, diff --git a/ggml/src/ggml-opencl/kernels/mul_mv_q4_k_f32.cl b/ggml/src/ggml-opencl/kernels/mul_mv_q4_k_f32.cl index 71ab98982..64cff077f 100644 --- a/ggml/src/ggml-opencl/kernels/mul_mv_q4_k_f32.cl +++ b/ggml/src/ggml-opencl/kernels/mul_mv_q4_k_f32.cl @@ -151,8 +151,8 @@ kernel void kernel_mul_mv_q4_K_f32( acc2.s3 += yh[i+9] * (q2[i/2] & 0xF000); } - float dall = dh[0]; - float dmin = dh[1]; + float dall = vload_half(0, dh); + float dmin = vload_half(1, dh); sumf[row] += dall * ((acc1.s0 + 1.f/256.f * acc1.s1) * sc8[0] + (acc1.s2 + 1.f/256.f * acc1.s3) * sc8[1] * 1.f/16.f + (acc2.s0 + 1.f/256.f * acc2.s1) * sc8[4] + ```

Zaloguj się, aby dołączyć do tej rozmowy.

Nie określono gałęzi/etykiety

Gałęzie Tagi

main

No results found.

Etykiety

Brak elementów

Brak etykiety

Kamień milowy

Brak elementów

Brak kamienia milowego

Projekty

Wyczyść projekty

Brak projektu

Przypisani

Usuń przypisanych

badstorm (Marco Olimpi) devbadxyz (Developer BadStorm)

Brak przypisanych

Uczestnicy 1

Powiadomienia

Subskrybuj

Termin realizacji

Brak ustawionego terminu realizacji.

Zależności

No dependencies set.

Reference: SRV/bdi_podman_serverconf#1