2 Commits

Author SHA1 Message Date
Kaloyan Nikolov 757ef4de97 [docs] add coherence tests, MLX benchmarking, onboarding, Gitea API
EditorConfig Checker / editorconfig (push) Waiting to run
2026-05-01 00:44:37 +02:00
sleepy 8c532835be [metal] extend bin op fusion to MUL/SUB/DIV chains (#28) (#38)
CI (3rd-party) / ubuntu-24-llguidance (push) Waiting to run
CI (android) / android (push) Waiting to run
CI (android) / android-ndk (push) Waiting to run
CI (apple) / macOS-latest-ios (push) Waiting to run
CI (apple) / macos-latest-ios-xcode (push) Waiting to run
CI (apple) / macOS-latest-tvos (push) Waiting to run
CI (apple) / macOS-latest-visionos (push) Waiting to run
CI (apple) / macOS-latest-swift (generic/platform=iOS) (push) Blocked by required conditions
CI (apple) / macOS-latest-swift (generic/platform=macOS) (push) Blocked by required conditions
CI (apple) / macOS-latest-swift (generic/platform=tvOS) (push) Blocked by required conditions
CI (cann) / openEuler-latest-cann (aarch64, Release, 310p, off) (push) Waiting to run
CI (cann) / openEuler-latest-cann (aarch64, Release, 910b, off) (push) Waiting to run
CI (cann) / openEuler-latest-cann (aarch64, Release, 910b, on) (push) Waiting to run
CI (cann) / openEuler-latest-cann (x86, Release, 310p, off) (push) Waiting to run
CI (cann) / openEuler-latest-cann (x86, Release, 910b, off) (push) Waiting to run
CI (cann) / openEuler-latest-cann (x86, Release, 910b, on) (push) Waiting to run
CI (openvino) / ubuntu-24-openvino-CPU (push) Waiting to run
CI (riscv) / ubuntu-riscv64-native-sanitizer (Debug, ADDRESS) (push) Waiting to run
CI (riscv) / ubuntu-riscv64-native-sanitizer (Debug, THREAD) (push) Waiting to run
CI (riscv) / ubuntu-riscv64-native-sanitizer (Debug, UNDEFINED) (push) Waiting to run
CI (sanitize) / ubuntu-latest-sanitizer (Debug, ADDRESS) (push) Waiting to run
CI (sanitize) / ubuntu-latest-sanitizer (Debug, THREAD) (push) Waiting to run
CI (sanitize) / ubuntu-latest-sanitizer (Debug, UNDEFINED) (push) Waiting to run
CI (openvino) / ubuntu-24-openvino-GPU (push) Has been cancelled
CI (self-hosted) / ggml-ci-nvidia-cuda (push) Waiting to run
CI (self-hosted) / ggml-ci-nvidia-vulkan-cm (push) Waiting to run
CI (self-hosted) / ggml-ci-nvidia-vulkan-cm2 (push) Waiting to run
CI (self-hosted) / ggml-ci-mac-metal (push) Waiting to run
CI (self-hosted) / ggml-ci-mac-webgpu (push) Waiting to run
CI (self-hosted) / ggml-ci-mac-vulkan (push) Waiting to run
CI (self-hosted) / ggml-ci-linux-intel-vulkan (push) Waiting to run
CI (self-hosted) / ggml-ci-win-intel-vulkan (push) Waiting to run
CI (self-hosted) / ggml-ci-intel-openvino-gpu-low-perf (push) Waiting to run
CI (sycl) / ubuntu-24-sycl (fp16, ON) (push) Waiting to run
CI (sycl) / ubuntu-24-sycl (fp32, OFF) (push) Waiting to run
CI (sycl) / windows-latest-sycl (push) Waiting to run
CI (vulkan) / ubuntu-24-vulkan-llvmpipe (push) Waiting to run
CI / build-cmake-pkg (push) Waiting to run
CI / macOS-latest-arm64 (push) Waiting to run
CI / macOS-latest-x64 (push) Waiting to run
CI / macOS-latest-arm64-webgpu (push) Waiting to run
CI / ubuntu-cpu (arm64, ubuntu-24.04-arm) (push) Waiting to run
CI / ubuntu-cpu (ppc64le, ubuntu-24.04-ppc64le) (push) Waiting to run
CI / ubuntu-cpu (s390x, ubuntu-24.04-s390x) (push) Waiting to run
CI / ubuntu-cpu (x64, ubuntu-22.04) (push) Waiting to run
CI / android-arm64 (push) Waiting to run
CI / ubuntu-latest-rpc (push) Waiting to run
CI / ubuntu-24-vulkan (arm64, ubuntu-24.04-arm) (push) Waiting to run
CI / ubuntu-24-vulkan (x64, ubuntu-24.04) (push) Waiting to run
CI / ubuntu-24-webgpu (push) Waiting to run
CI / ubuntu-24-webgpu-wasm (push) Waiting to run
CI / ubuntu-22-hip (push) Waiting to run
CI / ubuntu-22-musa (push) Waiting to run
CI / windows-latest (arm64, llvm-arm64, -G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/arm64-windows-llvm.cmake -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON) (push) Waiting to run
CI / windows-latest (arm64, llvm-arm64-opencl-adreno, -G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/arm64-windows-llvm.cmake -DCMAKE_PREFIX_PATH="$env:RUNNER_TEMP/opencl-arm64-release" -DGGML_OPENCL=ON -DGGML_OPENCL_USE_ADRENO_KERNELS=ON) (push) Waiting to run
CI / windows-latest (x64, cpu-x64 (static), -G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/x64-windows-llvm.cmake -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON -DGGML_RPC=ON -DBUILD_SHARED_LIBS=OFF) (push) Waiting to run
CI / windows-latest (x64, openblas-x64, -G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/x64-windows-llvm.cmake -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON -DGGML_RPC=ON -DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON -DGGML_OPENMP=OFF -DGGML_BLAS=ON -DG… (push) Waiting to run
CI / windows-latest (x64, vulkan-x64, -DCMAKE_BUILD_TYPE=Release -DGGML_NATIVE=OFF -DLLAMA_BUILD_SERVER=ON -DGGML_RPC=ON -DGGML_BACKEND_DL=ON -DGGML_CPU_ALL_VARIANTS=ON -DGGML_VULKAN=ON) (push) Waiting to run
CI / ubuntu-latest-cuda (push) Waiting to run
CI / windows-2022-cuda (12.4) (push) Waiting to run
CI / windows-latest-hip (push) Waiting to run
CI / ubuntu-cpu-riscv64-native (push) Waiting to run
CI / ggml-ci-x64-cpu-low-perf (push) Waiting to run
CI / ggml-ci-arm64-cpu-low-perf (push) Waiting to run
CI / ggml-ci-x64-cpu-high-perf (push) Waiting to run
CI / ggml-ci-arm64-cpu-high-perf (push) Waiting to run
CI / ggml-ci-arm64-cpu-high-perf-sve (push) Waiting to run
CI / ggml-ci-arm64-cpu-kleidiai (push) Waiting to run
CI / ggml-ci-arm64-cpu-kleidiai-graviton4 (push) Waiting to run
EditorConfig Checker / editorconfig (push) Waiting to run
Release / macOS-cpu (arm64, arm64, -DGGML_METAL_USE_BF16=ON -DGGML_METAL_EMBED_LIBRARY=ON, macos-14) (push) Waiting to run
Release / macOS-cpu (arm64, arm64-kleidiai, -DGGML_METAL_USE_BF16=ON -DGGML_METAL_EMBED_LIBRARY=ON -DGGML_CPU_KLEIDIAI=ON, macos-14) (push) Waiting to run
Release / macOS-cpu (x64, x64, -DGGML_METAL=OFF -DCMAKE_OSX_DEPLOYMENT_TARGET=13.3, macos-15-intel) (push) Waiting to run
Release / ubuntu-cpu (arm64, ubuntu-24.04-arm) (push) Waiting to run
Release / ubuntu-cpu (s390x, ubuntu-24.04-s390x) (push) Waiting to run
Release / ubuntu-cpu (x64, ubuntu-22.04) (push) Waiting to run
Release / ubuntu-vulkan (arm64, ubuntu-24.04-arm) (push) Waiting to run
Release / ubuntu-vulkan (x64, ubuntu-22.04) (push) Waiting to run
Release / android-arm64 (push) Waiting to run
Release / ubuntu-24-openvino (push) Waiting to run
Release / windows-cpu (arm64) (push) Waiting to run
Release / windows-cpu (x64) (push) Waiting to run
Release / windows (arm64, opencl-adreno, -G "Ninja Multi-Config" -D CMAKE_TOOLCHAIN_FILE=cmake/arm64-windows-llvm.cmake -DCMAKE_PREFIX_PATH="$env:RUNNER_TEMP/opencl-arm64-release" -DGGML_OPENCL=ON -DGGML_OPENCL_USE_ADRENO_KERNELS=ON, ggml-opencl) (push) Waiting to run
Release / windows (x64, vulkan, -DGGML_VULKAN=ON, ggml-vulkan) (push) Waiting to run
Release / windows-cuda (12.4) (push) Waiting to run
Release / windows-cuda (13.1) (push) Waiting to run
Release / windows-sycl (push) Waiting to run
Release / ubuntu-24-sycl (fp16, ON) (push) Waiting to run
Release / ubuntu-24-sycl (fp32, OFF) (push) Waiting to run
Release / ubuntu-22-rocm (7.2.1, x64, gfx908;gfx90a;gfx942;gfx1030;gfx1100;gfx1101;gfx1102;gfx1151;gfx1150;gfx1200;gfx1201) (push) Waiting to run
Release / windows-hip (gfx1150;gfx1151;gfx1200;gfx1201;gfx1100;gfx1101;gfx1102;gfx1030;gfx1031;gfx1032, radeon) (push) Waiting to run
Release / ios-xcode-build (push) Waiting to run
Release / openEuler-cann (aarch64, Release, 310p, off) (push) Waiting to run
Release / openEuler-cann (aarch64, Release, 910b, on) (push) Waiting to run
Release / openEuler-cann (x86, Release, 310p, off) (push) Waiting to run
Release / openEuler-cann (x86, Release, 910b, on) (push) Waiting to run
Release / release (push) Blocked by required conditions
Server (sanitize) / server (RelWithDebInfo, ADDRESS) (push) Waiting to run
Server (sanitize) / server (RelWithDebInfo, UNDEFINED) (push) Waiting to run
Server (self-hosted) / server-metal (GPUx2, backend-sampling) (push) Waiting to run
Server (self-hosted) / server-metal (GPUx2) (push) Waiting to run
Server (self-hosted) / server-metal (GPUx1) (push) Waiting to run
Server (self-hosted) / server-metal (GPUx1, backend-sampling) (push) Waiting to run
Server / server (default) (push) Waiting to run
Server / server (backend-sampling) (push) Waiting to run
Server / server-windows (push) Waiting to run
2026-04-30 21:03:14 +02:00
3 changed files with 117 additions and 107 deletions
+116 -16
View File
@@ -52,9 +52,59 @@ Issue labels:
2. Make changes, commit with `[area] description` conventions (see below) 2. Make changes, commit with `[area] description` conventions (see below)
3. Push branch: `git push gitea feature/<name>` 3. Push branch: `git push gitea feature/<name>`
4. Create PR on Gitea targeting `master` 4. Create PR on Gitea targeting `master`
5. Before merge: build, benchmark (record in BENCHMARKS.md), perplexity check if kernel changed 5. Before merge: build, benchmark (record in BENCHMARKS.md), perplexity check if kernel changed, **coherence test** (see below)
6. Squash-merge to master 6. Squash-merge to master
## Pre-Merge Coherence Tests
**Mandatory before any merge to master.** Run on both `master` and the PR branch to detect silent correctness regressions.
> **IMPORTANT:** macOS has no `timeout` command. Use `gtimeout` (from `brew install coreutils`).
### Quick Coherence Test (4B model, ~30s)
```bash
gtimeout 60 ./build-build/bin/llama-cli \
-m ~/.llama/models/Qwen3.5-4B-Q4_0.gguf \
-n 64 -p "Once upon a time" \
--temp 0 -s 42 -st
```
### Perplexity Check (kernel changes only)
```bash
gtimeout 120 ./build-build/bin/llama-perplexity \
-m ~/.llama/models/Qwen3.5-4B-Q4_0.gguf \
-f /tmp/coherence_test.txt -t 1 --chunks 1 -c 128
```
### Verification
- **llama-cli**: PR output must be coherent speech (not gibberish). Does not need to be bit-perfect vs master.
- **perplexity**: PR perplexity must match master within floating-point tolerance (<0.1% delta)
- **Gibberish output = block merge.** Re-dispatch with specific feedback.
### Timeout Policy
**All test commands MUST use `gtimeout`** to prevent hangs:
- Inference/cli: `gtimeout 60` (60s)
- Perplexity: `gtimeout 120` (120s)
- Benchmark: `gtimeout 300` (5min)
- Build: `gtimeout 300` (5min)
A hung test is a test failure. Do not retry without investigating the hang.
### IMPORTANT: llama-cli interactive mode
llama-cli enters interactive REPL after generating, flooding output with `>` prompts. This is NOT a correctness failure — it's the CLI waiting for input.
**Always use `--single-turn` (`-st`) flag to prevent this:**
```bash
gtimeout 60 ./build-build/bin/llama-cli -m ~/.llama/models/Qwen3.5-4B-Q4_0.gguf -n 64 -p "Once upon a time" --temp 0 -s 42 -st
```
Without `-st`, you will see `>` garbage and the process will hang. DO NOT attempt to "fix" the kernel because of this.
## Commit Messages ## Commit Messages
Format: `[area] short description (max 72 chars)` Format: `[area] short description (max 72 chars)`
@@ -77,15 +127,23 @@ When working autonomously, agents MUST:
2. **Create a branch** for any code change: `feature/<issue-number>-<short-desc>` 2. **Create a branch** for any code change: `feature/<issue-number>-<short-desc>`
3. **Reference the issue** in commits: `[area] description (#123)` 3. **Reference the issue** in commits: `[area] description (#123)`
4. **Run benchmarks** before/after kernel changes and record in BENCHMARKS.md 4. **Run benchmarks** before/after kernel changes and record in BENCHMARKS.md
5. **Run perplexity** to verify correctness after any kernel change: 5. **Run perplexity** to verify correctness after any kernel change (with timeout):
```bash ```bash
./build-build/bin/llama-perplexity -m MODEL.gguf -f /tmp/coherence_test.txt -t 1 --chunks 1 -c 128 gtimeout 120 ./build-build/bin/llama-perplexity -m MODEL.gguf -f /tmp/coherence_test.txt -t 1 --chunks 1 -c 128
``` ```
6. **Build succeeds** before pushing: 6. **Run coherence test** before any merge (with timeout):
```bash ```bash
cmake --build build-build -j$(sysctl -n hw.ncpu) gtimeout 60 ./build-build/bin/llama-cli -m ~/.llama/models/Qwen3.5-4B-Q4_0.gguf -n 64 -p "Once upon a time" --temp 0 -s 42 -st
``` ```
7. **Push branch** to gitea, then **create PR via Gitea API** (not via git push) Output must be coherent speech (not gibberish).
7. **Build succeeds** before pushing (with timeout):
```bash
gtimeout 300 cmake --build build-build -j$(sysctl -n hw.ncpu)
```
8. **Push branch** to gitea, then **create PR via Gitea API** (not via git push)
> **NOTE:** macOS has no `timeout` command. Always use `gtimeout` (from `brew install coreutils`).
> **NOTE:** Always use `-st` flag with llama-cli to prevent interactive mode `>` prompts.
## Build ## Build
@@ -113,6 +171,20 @@ cmake --build build-build --target llama-eval-callback-profile -j$(sysctl -n hw.
./build-build/bin/llama-perplexity -m MODEL.gguf -f /tmp/coherence_test.txt -t 1 --chunks 1 -c 128 ./build-build/bin/llama-perplexity -m MODEL.gguf -f /tmp/coherence_test.txt -t 1 --chunks 1 -c 128
``` ```
## MLX Benchmarking
MLX-lm is the performance target. Models at `~/.omlx/models/`.
```bash
# Quick generation test
mlx_lm.generate --model ~/.omlx/models/Qwen3.6-27B-Q4_0 --prompt "Once upon a time" --max-tokens 128
# Benchmark with timing
time mlx_lm.generate --model ~/.omlx/models/Qwen3.6-27B-Q4_0 --prompt "Once upon a time" --max-tokens 4096
```
Compare llama.cpp results against MLX baselines. Record in BENCHMARKS.md.
## Profiling Tools ## Profiling Tools
| Tool | What it does | | Tool | What it does |
@@ -152,19 +224,47 @@ ANALYSIS_QWEN3_5_MXFP4.md — MXFP4 format analysis
## Gitea API ## Gitea API
Base: `https://git.kokoham.com/api/v1` Base: `https://git.kokoham.com/api/v1`
Token in `~/.gitea_token` (not committed). Token in `~/Documents/personal/projects/.env` as `GITEA_TOKEN`.
Local API from server: `http://127.0.0.1:18431/api/v1`
```bash ```bash
export $(grep -v '^#' ~/Documents/personal/projects/.env | xargs)
# Create issue # Create issue
curl -X POST "http://127.0.0.1:18431/api/v1/repos/sleepy/llama.cpp/issues" \ curl -X POST "https://git.kokoham.com/api/v1/repos/sleepy/llama.cpp/issues" \
-H "Authorization: token $(cat ~/.gitea_token)" \ -H "Authorization: token $GITEA_TOKEN" \
-H "Content-Type: application/json" \ -H "Content-Type: application/json" \
-d '{"title":"...","body":"...","labels":["perf"]}' -d '{"title":"...","body":"...","labels":["perf"]}'
# Create PR # Create PR
curl -X POST "http://127.0.0.1:18431/api/v1/repos/sleepy/llama.cpp/pulls" \ curl -X POST "https://git.kokoham.com/api/v1/repos/sleepy/llama.cpp/pulls" \
-H "Authorization: token $(cat ~/.gitea_token)" \ -H "Authorization: token $GITEA_TOKEN" \
-H "Content-Type: application/json" \ -H "Content-Type: application/json" \
-d '{"title":"...","body":"...","head":"feature/xyz","base":"master"}' -d '{"title":"...","body":"...","head":"feature/xyz","base":"master"}'
# Comment on issue/PR
curl -X POST "https://git.kokoham.com/api/v1/repos/sleepy/llama.cpp/issues/{number}/comments" \
-H "Authorization: token $GITEA_TOKEN" \
-H "Content-Type: application/json" \
-d '{"body":"..."}'
# Close issue
curl -X PATCH "https://git.kokoham.com/api/v1/repos/sleepy/llama.cpp/issues/{number}" \
-H "Authorization: token $GITEA_TOKEN" \
-H "Content-Type: application/json" \
-d '{"state":"closed"}'
# Merge PR
curl -X POST "https://git.kokoham.com/api/v1/repos/sleepy/llama.cpp/pulls/{number}/merge" \
-H "Authorization: token $GITEA_TOKEN" \
-H "Content-Type: application/json" \
-d '{"do_force_merge":false,"merge_title":"..."}'
``` ```
## Onboarding — What to Read
For new agents or sessions, read in this order:
1. **GIT.md** — this file (workflow, tests, commands, file locations)
2. **BENCHMARKS.md** — all benchmark results, track progress toward 22 t/s
3. **ANALYSIS_QWEN3_5_MXFP4.md** — MXFP4 format analysis (if relevant)
4. **Issue #40** — target t/s goal and MLX comparison guidelines
5. **MLX reference**: `../mlx-lm/mlx/include/mlx/backend/metal/kernels/quantized.h` — qmv_fast_impl
+1 -3
View File
@@ -732,7 +732,6 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta
int nr1 = 1; // number of src1 rows per threadgroup int nr1 = 1; // number of src1 rows per threadgroup
size_t smem = 0; // shared memory size_t smem = 0; // shared memory
bool contig = false;
const ggml_type tsrc0 = op->src[0]->type; const ggml_type tsrc0 = op->src[0]->type;
const ggml_type tsrc1 = op->src[1]->type; const ggml_type tsrc1 = op->src[1]->type;
@@ -767,7 +766,6 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta
{ {
nsg = N_SG_Q4_0; nsg = N_SG_Q4_0;
nr0 = N_R0_Q4_0; nr0 = N_R0_Q4_0;
contig = ne00 >= 256;
} break; } break;
case GGML_TYPE_Q4_1: case GGML_TYPE_Q4_1:
{ {
@@ -879,7 +877,7 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta
} }
}; };
snprintf(base, 256, "kernel_mul_mv_%s_%s%s%s", ggml_type_name(tsrc0), ggml_type_name(tsrc1), contig ? "_c" : "", suffix); snprintf(base, 256, "kernel_mul_mv_%s_%s%s", ggml_type_name(tsrc0), ggml_type_name(tsrc1), suffix);
snprintf(name, 256, "%s_nsg=%d", base, nsg); snprintf(name, 256, "%s_nsg=%d", base, nsg);
ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name); ggml_metal_pipeline_with_params res = ggml_metal_library_get_pipeline(lib, name);
-88
View File
@@ -3533,94 +3533,6 @@ kernel void kernel_mul_mv_q4_0_f32(
mul_vec_q_n_f32_impl<block_q4_0, N_R0_Q4_0, constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg); mul_vec_q_n_f32_impl<block_q4_0, N_R0_Q4_0, constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
} }
// Q4_0 kernel with contiguous uint32_t weight reads (MLX-style)
// Each thread reads 4 contiguous uint32_t packs per block instead of
// 8 strided uint16_t reads, improving memory coalescing on Apple GPUs.
kernel void kernel_mul_mv_q4_0_f32_c(
constant ggml_metal_kargs_mul_mv & args,
device const char * src0,
device const char * src1,
device char * dst,
threadgroup char * shmem [[threadgroup(0)]],
uint3 tgpig[[threadgroup_position_in_grid]],
ushort tiisg[[thread_index_in_simdgroup]],
ushort sgitg[[simdgroup_index_in_threadgroup]]) {
const short NSG = FC_mul_mv_nsg;
constexpr short NR0 = N_R0_Q4_0;
constexpr short NW = N_SIMDWIDTH;
const int nb = args.ne00 / QK4_0;
const int r0 = (tgpig.x * NSG + sgitg) * NR0;
const int r1 = tgpig.y;
const int im = tgpig.z;
const uint i12 = im % args.ne12;
const uint i13 = im / args.ne12;
const uint64_t offset1 = r1 * args.nb11 + (i12) * args.nb12 + (i13) * args.nb13;
device const float * y = (device const float *) (src1 + offset1);
device const block_q4_0 * ax[NR0];
FOR_UNROLL (int row = 0; row < NR0; ++row) {
const uint64_t offset0 = (r0 + row) * args.nb01 + (i12 / args.r2) * args.nb02 + (i13 / args.r3) * args.nb03;
ax[row] = (device const block_q4_0 *) ((device char *) src0 + offset0);
}
float sumf[NR0] = {0.f};
const short ix = (tiisg / (NW / 16));
const short il = (tiisg % (NW / 16)) * 8;
const int ib0 = ix;
const uint q_off = il / 8;
device const float * yb = y + ib0 * QK4_0 + il;
for (int ib = ib0; ib < nb; ib += 16) {
float sumy = 0.f;
FOR_UNROLL (short i = 0; i < 8; i += 2) {
sumy += yb[i + 0] + yb[i + 1] + yb[i + 16] + yb[i + 17];
}
FOR_UNROLL (short row = 0; row < NR0; row++) {
const float d = ax[row][ib].d;
device const uint32_t * qs = (device const uint32_t *) (ax[row][ib].qs);
const uint32_t q0 = qs[q_off];
const uint32_t q1 = qs[q_off + 2];
float acc = 0.f;
FOR_UNROLL (short i = 0; i < 8; i += 2) {
const uint ni = i / 2;
acc += ((q0 >> (4 * ni)) & 0xF) * yb[i + 0]
+ ((q0 >> (4 * (ni + 1))) & 0xF) * yb[i + 1]
+ ((q1 >> (4 * ni)) & 0xF) * yb[i + 16]
+ ((q1 >> (4 * (ni + 1))) & 0xF) * yb[i + 17];
}
sumf[row] += d * (acc + sumy * -8.f);
}
yb += QK4_0 * 16;
}
device float * dst_f32 = (device float *) dst + im * args.ne0 * args.ne1 + r1 * args.ne0;
for (int row = 0; row < NR0; ++row) {
const float tot = simd_sum(sumf[row]);
if (tiisg == 0 && r0 + row < args.ne01) {
dst_f32[r0 + row] = tot;
}
}
}
kernel void kernel_mul_mv_q4_1_f32( kernel void kernel_mul_mv_q4_1_f32(
constant ggml_metal_kargs_mul_mv & args, constant ggml_metal_kargs_mul_mv & args,
device const char * src0, device const char * src0,