llama.cpp with the Entrit d2 kernel (D2NL)

LICENSE

@@ -0,0 +1,22 @@
+MIT License
+
+Copyright (c) 2026 Entrit Systems
+Copyright (c) 2023-2026 The ggml authors
+
+Permission is hereby granted, free of charge, to any person obtaining a copy
+of this software and associated documentation files (the "Software"), to deal
+in the Software without restriction, including without limitation the rights
+to use, copy, modify, merge, publish, distribute, sublicense, and/or sell
+copies of the Software, and to permit persons to whom the Software is
+furnished to do so, subject to the following conditions:
+
+The above copyright notice and this permission notice shall be included in all
+copies or substantial portions of the Software.
+
+THE SOFTWARE IS PROVIDED "AS IS", WITHOUT WARRANTY OF ANY KIND, EXPRESS OR
+IMPLIED, INCLUDING BUT NOT LIMITED TO THE WARRANTIES OF MERCHANTABILITY,
+FITNESS FOR A PARTICULAR PURPOSE AND NONINFRINGEMENT. IN NO EVENT SHALL THE
+AUTHORS OR COPYRIGHT HOLDERS BE LIABLE FOR ANY CLAIM, DAMAGES OR OTHER
+LIABILITY, WHETHER IN AN ACTION OF CONTRACT, TORT OR OTHERWISE, ARISING FROM,
+OUT OF OR IN CONNECTION WITH THE SOFTWARE OR THE USE OR OTHER DEALINGS IN THE
+SOFTWARE.

LLAMA_CPP_VERSION

@@ -0,0 +1 @@
+62061f91088281e65071cc38c5f69ee95c39f14e

README.md

@@ -0,0 +1,162 @@
+---
+license: mit
+tags:
+  - entrit
+  - d2
+  - llama.cpp
+  - gguf
+  - quantization
+---
+
+# Run Entrit d2 models with llama.cpp
+
+This installs **llama.cpp**, the free program that runs AI models on your own computer, with
+Entrit's **d2 kernel** added. Our d2 models store each weight as one of 9 levels, and stock
+llama.cpp can't read that format. The install script downloads llama.cpp, adds our kernel, and
+builds it for your graphics card. With no graphics card, it builds a slower version that runs on
+the processor.
+
+You'll type a few commands into a terminal. You don't need to know how to program.
+
+## Which models it runs
+
+| Model | File to download | Runs here? |
+|---|---|---|
+| [Qwen3.8-27B d2](https://huggingface.co/Entrit/Qwen3.8-27B-d2-GGUF) | `Qwen3.8-27B-d2-D2NL.gguf` (15.3 GB) | yes |
+| [Muse-Glimmer-30B d2](https://huggingface.co/Entrit/Muse-Glimmer-30B-d2) | — | no: llama.cpp doesn't support this architecture yet. Use `transformers`. |
+| [DeepSeek-V4-Flash d2](https://huggingface.co/Entrit/DeepSeek-V4-Flash-d2swap-mp2) | — | no: needs DeepSeek's own inference code. |
+
+## What you need
+
+- A Mac with Apple Silicon (M1 or newer), or a Linux PC.
+  Windows works through WSL (see the end of this page).
+- About 5 GB of free disk space for the program, plus room for the model.
+- Enough memory for the model. Qwen3.8-27B needs a graphics card with **24 GB** (RTX 3090, 4090,
+  or similar), or a Mac with at least **32 GB**.
+- Internet, for the first install.
+
+## Step 1: Open a terminal
+
+- **Mac:** press Cmd-Space, type `Terminal`, press Enter.
+- **Linux:** press Ctrl-Alt-T.
+
+Every grey box below is a command. Copy it, paste it into the terminal, and press Enter.
+
+## Step 2: Install the build tools (once)
+
+**Mac:**
+```bash
+xcode-select --install
+```
+A window pops up. Click Install and wait for it to finish. Then install Homebrew (a tool installer)
+if you don't have it, from https://brew.sh, and run:
+```bash
+brew install cmake
+```
+
+**Linux (Ubuntu or Debian):**
+```bash
+sudo apt update && sudo apt install -y build-essential cmake curl patch
+```
+Then add the part for your graphics card:
+
+| Your graphics card | Extra command |
+|---|---|
+| NVIDIA | Install the CUDA Toolkit from https://developer.nvidia.com/cuda-downloads |
+| AMD or Intel | `sudo apt install -y libvulkan-dev glslc` |
+| None | nothing; you'll get the processor version |
+
+## Step 3: Download and install
+
+```bash
+mkdir ~/entrit-llama && cd ~/entrit-llama
+curl -fsSLO https://huggingface.co/Entrit/llama.cpp-d2nl/resolve/main/install.sh
+bash install.sh
+```
+The first two lines make a folder and download the installer into it. The installer then downloads
+the rest of our files and llama.cpp, finds your graphics card, and builds for it. The first build
+takes 5 to 20 minutes. At the end you should see:
+```
+==> Self-test: Entrit kernel on ... vs the CPU
+PASS
+==> Done
+```
+**PASS** means our kernel gives the same answers on your graphics card as on the processor.
+
+## Step 4: Download a model
+
+```bash
+./get-model.sh Entrit/Qwen3.8-27B-d2-GGUF Qwen3.8-27B-d2-D2NL.gguf
+```
+It's 15 GB, so give it a while. If the download stops, run the same command again and it picks up
+where it left off.
+
+## Step 5: Chat with it
+
+```bash
+./run.sh models/Qwen3.8-27B-d2-D2NL.gguf
+```
+Wait until the terminal says `server is listening`. Then open **http://localhost:8080** in your web
+browser and type a message. To stop, go back to the terminal and press Ctrl-C.
+
+Other programs can talk to it too. It answers the OpenAI chat API at `http://localhost:8080/v1`.
+
+## Updating
+
+When we release a new version of the kernel:
+```bash
+cd ~/entrit-llama
+./install.sh --update
+```
+It downloads the newest files and rebuilds. Your downloaded models stay where they are.
+
+## If something goes wrong
+
+| What you see | What to do |
+|---|---|
+| `'cmake' is missing` (or c++, curl, patch) | Do Step 2 again. |
+| `could not download ...` | Check your internet connection and try again. |
+| `Vulkan needs glslc` | Run `sudo apt install -y libvulkan-dev glslc`, or install the Vulkan SDK from https://vulkan.lunarg.com |
+| `configure failed` or `build failed` | Open `build-config.log` or `build.log`. The last 30 lines usually say what is missing. |
+| `the kernel patch did not apply` | Run `./install.sh --update` to get matching files. |
+| The self-test doesn't say PASS | Open an issue on this page's Community tab with `build.log` and the name of your graphics card. |
+| `out of memory`, or the model loads very slowly | The model is too big for your card. Add `-ngl 30` to put only part of it on the graphics card: `./run.sh models/<file>.gguf -ngl 30` |
+| It picked the wrong backend | Choose one yourself: `./install.sh --backend vulkan` (or cuda, metal, hip, cpu) |
+
+## Windows
+
+Install WSL, which gives you Linux inside Windows. Open PowerShell as administrator and run
+`wsl --install`, restart, then open "Ubuntu" from the Start menu. From there, follow the Linux steps.
+NVIDIA cards work inside WSL with NVIDIA's normal Windows driver.
+
+## How it works
+
+d2 stores each weight as one of 9 levels, `{0, ±1, ±2.83, ±5.20, ±8}`, times one scale per group of
+64 weights. The kernel adds a GGUF type, **D2NL**, that holds the level index in the same 4-bit block
+llama.cpp uses for its IQ4_NL type, and expands it through a 16-entry int8 table:
+
+```
+level   0    ±1    ±2.828   ±5.196   ±8
+int8    0   ±11    ±31      ±57      ±88      (10.981 × level, within 0.19% of each level)
+```
+
+So every IQ4_NL code path serves D2NL with the table swapped in: dp4a and int8 tensor cores on
+NVIDIA, Vulkan shaders on AMD and Intel, Metal on Apple, and plain C on the processor. The AMD ROCm
+route (`--backend hip`) builds from the CUDA code but hasn't been tested; use `--backend vulkan` if it
+gives you trouble.
+
+| File | What it does |
+|---|---|
+| `install.sh` | downloads llama.cpp, adds our kernel, builds, runs the self-test |
+| `get-model.sh` | downloads a model from Hugging Face into `models/` |
+| `run.sh` | starts the chat server |
+| `patches/entrit-d2nl.diff` | our changes to llama.cpp |
+| `LLAMA_CPP_VERSION` | the exact llama.cpp version our changes apply to (from [PrismML-Eng/llama.cpp](https://github.com/PrismML-Eng/llama.cpp)) |
+
+After installing, the folder also holds `llama.cpp/` (the source), `build/` (the program) and
+`models/`. You can delete `llama.cpp/` and `build/` at any time; `./install.sh` makes them again.
+
+## License and notice
+
+The patch and scripts are released under the MIT License, the same as llama.cpp. The d2 codec is the
+subject of pending patent applications by Entrit Systems. Contact: eric@entrit.io

get-model.sh

@@ -0,0 +1,32 @@
+#!/usr/bin/env bash
+# Download a model file from Hugging Face into models/.
+#
+#   ./get-model.sh <repo>            list the .gguf files in the repo
+#   ./get-model.sh <repo> <file>     download one of them
+#
+# Private repos need a token: export HF_TOKEN=hf_...   (huggingface.co -> Settings -> Access Tokens)
+set -euo pipefail
+
+HERE=$(cd "$(dirname "$0")" && pwd)
+REPO=${1:?usage: ./get-model.sh <huggingface-repo> [file.gguf]}
+FILE=${2:-}
+AUTH=()
+[[ -n "${HF_TOKEN:-}" ]] && AUTH=(-H "Authorization: Bearer $HF_TOKEN")
+
+if [[ -z "$FILE" ]]; then
+    echo "GGUF files in $REPO:"
+    curl -fsSL ${AUTH[@]+"${AUTH[@]}"} "https://huggingface.co/api/models/$REPO" \
+        | grep -o '"rfilename":"[^"]*\.gguf"' | sed 's/"rfilename":"\(.*\)"/  \1/' \
+        || { echo "Could not read $REPO. Check the name, and set HF_TOKEN if the repo is private."; exit 1; }
+    echo
+    echo "Download one with:  ./get-model.sh $REPO <file>"
+    exit 0
+fi
+
+mkdir -p "$HERE/models"
+OUT="$HERE/models/$(basename "$FILE")"
+echo "Downloading $FILE to models/ (it resumes if interrupted)"
+curl -fL -C - --progress-bar ${AUTH[@]+"${AUTH[@]}"} -o "$OUT" "https://huggingface.co/$REPO/resolve/main/$FILE" \
+    || { echo "Download failed. Check the file name, and set HF_TOKEN if the repo is private."; exit 1; }
+echo "Saved $OUT"
+echo "Run it with:  ./run.sh models/$(basename "$FILE")"

install.sh

@@ -0,0 +1,122 @@
+#!/usr/bin/env bash
+# Install llama.cpp with the Entrit d2 kernel (D2NL).
+#
+#   bash install.sh                   first install: fetches the other files, then builds
+#   ./install.sh --update             download the newest kernel, then rebuild
+#   ./install.sh --backend vulkan     force one backend: cuda | vulkan | metal | hip | cpu
+#
+# Private repo? Set your Hugging Face token first:  export HF_TOKEN=hf_...
+set -euo pipefail
+
+HERE=$(cd "$(dirname "$0")" && pwd)
+REPO_URL=${ENTRIT_REPO:-https://huggingface.co/Entrit/llama.cpp-d2nl/resolve/main}
+FILES=(install.sh get-model.sh run.sh README.md LLAMA_CPP_VERSION patches/entrit-d2nl.diff)
+BACKEND=auto
+UPDATE=0
+
+while [[ $# -gt 0 ]]; do
+    case "$1" in
+        --backend)  BACKEND=${2:?--backend needs a value}; shift 2 ;;
+        --update)   UPDATE=1; shift ;;
+        --fetched)  UPDATE=0; FETCHED=1; shift ;;
+        -h|--help)  sed -n '2,9p' "$0"; exit 0 ;;
+        *) echo "unknown option: $1 (try --help)"; exit 1 ;;
+    esac
+done
+
+say()  { printf '\n\033[1m==> %s\033[0m\n' "$*"; }
+fail() { printf '\n\033[31mERROR: %s\033[0m\n' "$*"; exit 1; }
+OS=$(uname -s)
+
+need() {
+    command -v "$1" >/dev/null 2>&1 && return
+    case "$OS" in
+        Darwin) fail "'$1' is missing. Install the build tools:  xcode-select --install   and then:  brew install cmake" ;;
+        *)      fail "'$1' is missing. On Ubuntu/Debian:  sudo apt install -y build-essential cmake curl patch" ;;
+    esac
+}
+need curl
+
+# ------------------------------------------------------------------ 1. our files
+AUTH=()
+[[ -n "${HF_TOKEN:-}" ]] && AUTH=(-H "Authorization: Bearer $HF_TOKEN")
+missing=0
+for f in "${FILES[@]}"; do [[ -f "$HERE/$f" ]] || missing=1; done
+if [[ "$UPDATE" == 1 || ( "$missing" == 1 && -z "${FETCHED:-}" ) ]]; then
+    say "Downloading the Entrit files"
+    mkdir -p "$HERE/patches"
+    for f in "${FILES[@]}"; do
+        curl -fsSL ${AUTH[@]+"${AUTH[@]}"} -o "$HERE/$f.new" "$REPO_URL/$f" \
+            || fail "could not download $f. If the repository is private, set HF_TOKEN (see the top of this file)."
+        mv "$HERE/$f.new" "$HERE/$f"
+    done
+    chmod +x "$HERE/install.sh" "$HERE/get-model.sh" "$HERE/run.sh"
+    exec bash "$HERE/install.sh" --fetched --backend "$BACKEND"    # run the version we just downloaded
+fi
+
+# ------------------------------------------------------------------ 2. tools
+say "Checking tools"
+need cmake; need c++; need tar; need patch
+echo "cmake, a C++ compiler, tar and patch are present"
+
+# ------------------------------------------------------------------ 3. backend
+if [[ "$BACKEND" == auto ]]; then
+    if   [[ "$OS" == Darwin ]];                                               then BACKEND=metal
+    elif command -v nvcc >/dev/null 2>&1 || [[ -x /usr/local/cuda/bin/nvcc ]]; then BACKEND=cuda
+    elif command -v hipcc >/dev/null 2>&1;                                     then BACKEND=hip
+    elif command -v glslc >/dev/null 2>&1;                                     then BACKEND=vulkan
+    else                                                                            BACKEND=cpu
+    fi
+fi
+case "$BACKEND" in
+    metal)  FLAGS=(-DGGML_METAL=ON) ;;
+    cuda)   [[ -x /usr/local/cuda/bin/nvcc ]] && export PATH="/usr/local/cuda/bin:$PATH"
+            FLAGS=(-DGGML_CUDA=ON -DCMAKE_CUDA_ARCHITECTURES=native) ;;   # build for the GPU in this machine only
+    hip)    FLAGS=(-DGGML_HIP=ON) ;;
+    vulkan) command -v glslc >/dev/null 2>&1 || fail "Vulkan needs glslc. Ubuntu: sudo apt install -y libvulkan-dev glslc   Other: install the Vulkan SDK from https://vulkan.lunarg.com"
+            FLAGS=(-DGGML_VULKAN=ON) ;;
+    cpu)    FLAGS=(-DGGML_CUDA=OFF -DGGML_VULKAN=OFF -DGGML_METAL=OFF) ;;
+    *)      fail "unknown backend '$BACKEND' (use cuda, vulkan, metal, hip or cpu)" ;;
+esac
+say "Backend: $BACKEND"
+[[ "$BACKEND" == cpu ]] && echo "No GPU toolkit found, so this build runs on the CPU only. It works, but it is slow."
+
+# ------------------------------------------------------------------ 4. source
+VERSION=$(tr -d '[:space:]' < "$HERE/LLAMA_CPP_VERSION")
+SRC="$HERE/llama.cpp"
+say "Downloading llama.cpp (version ${VERSION:0:7})"
+rm -rf "$SRC" && mkdir -p "$SRC"
+curl -fsSL "https://github.com/PrismML-Eng/llama.cpp/archive/$VERSION.tar.gz" | tar -xz -C "$SRC" --strip-components 1 \
+    || fail "could not download llama.cpp from GitHub. Check your internet connection."
+
+say "Adding the Entrit kernel"
+POPTS=(-p1 -s)
+patch --version 2>/dev/null | grep -q GNU && POPTS+=(--no-backup-if-mismatch)   # GNU patch (Linux); macOS patch has no .orig files to avoid
+(cd "$SRC" && patch "${POPTS[@]}" < "$HERE/patches/entrit-d2nl.diff") \
+    || fail "the kernel patch did not apply. Run ./install.sh --update and try again."
+chmod +x "$SRC"/entrit/*.sh "$SRC"/entrit/*.py 2>/dev/null || true
+
+# ------------------------------------------------------------------ 5. build
+say "Building (this takes 5-20 minutes the first time)"
+BUILD="$HERE/build"
+JOBS=$( (command -v nproc >/dev/null && nproc) || sysctl -n hw.ncpu 2>/dev/null || echo 4)
+cmake -S "$SRC" -B "$BUILD" -DCMAKE_BUILD_TYPE=Release "${FLAGS[@]}" > "$HERE/build-config.log" 2>&1 \
+    || fail "configure failed; the details are in build-config.log"
+cmake --build "$BUILD" -j "$JOBS" --target llama-server llama-cli llama-completion llama-bench test-backend-ops \
+    > "$HERE/build.log" 2>&1 || fail "build failed; the details are in build.log"
+
+# ------------------------------------------------------------------ 6. self-test
+DEV=""
+case "$BACKEND" in cuda|hip) DEV=CUDA0 ;; vulkan) DEV=Vulkan0 ;; metal) DEV=MTL0 ;; esac
+if [[ -n "$DEV" ]]; then
+    say "Self-test: Entrit kernel on $DEV vs the CPU"
+    if "$BUILD/bin/test-backend-ops" test -o MUL_MAT -b "$DEV" -p d2nl 2>&1 | grep -q "tests passed"; then
+        echo "PASS"
+    else
+        echo "The self-test did not pass. The build may still run, but please report this, with build.log."
+    fi
+fi
+
+say "Done"
+echo "Next: download a model with   ./get-model.sh <huggingface-repo>"
+echo "Then run it with              ./run.sh models/<file>.gguf"

patches/entrit-d2nl.diff

@@ -0,0 +1,1352 @@
+--- a/ggml/include/ggml.h
++++ b/ggml/include/ggml.h
+@@ -430,7 +430,9 @@ extern "C" {
+         GGML_TYPE_NVFP4   = 40, // NVFP4 (4 blocks, E4M3 scale)
+         GGML_TYPE_Q1_0    = 41,
+         GGML_TYPE_Q2_0    = 42,
+-        GGML_TYPE_COUNT   = 43,
++        // GGML_TYPE_43 is reserved
++        GGML_TYPE_D2NL    = 44, // Entrit d2: 9 power-mapped levels through an int8 table, IQ4_NL layout
++        GGML_TYPE_COUNT   = 45,
+     };
+ 
+     // precision
+@@ -475,6 +477,7 @@ extern "C" {
+         GGML_FTYPE_MOSTLY_NVFP4   = 26, // except 1d tensors
+         GGML_FTYPE_MOSTLY_Q1_0    = 27, // except 1d tensors
+         GGML_FTYPE_MOSTLY_Q2_0    = 28, // except 1d tensors
++        GGML_FTYPE_MOSTLY_D2NL    = 129, // except 1d tensors
+     };
+ 
+     // available tensor operations:
+--- a/ggml/src/ggml-common.h
++++ b/ggml/src/ggml-common.h
+@@ -451,6 +451,22 @@ typedef struct {
+ } block_iq4_nl;
+ static_assert(sizeof(block_iq4_nl) == sizeof(ggml_half) + QK4_NL/2, "wrong iq4_nl block size/padding");
+ 
++// d2nl: Entrit d2 weights. Each weight is one of 9 levels {0, +-1, +-2.828, +-5.196, +-8} (k^1.5 for
++// k = 0..4) times one scale per group of 64 weights. Stored in IQ4_NL's block so IQ4_NL's kernels can
++// run it with kvalues_d2nl in place of kvalues_iq4nl: 32 weights = fp16 scale + 16 bytes of 4-bit level
++// indices (4.50 bpw). qs[j] low nibble = weight j, high nibble = weight j+16. Index 4 is zero.
++// The stored scale is the group scale divided by GGML_D2NL_K.
++#define QK_D2NL 32
++#define QR_D2NL 2
++#define QI_D2NL (QK_D2NL / (4 * QR_D2NL))
++typedef struct {
++    ggml_half d;
++    uint8_t   qs[QK_D2NL/2];
++} block_d2nl;
++static_assert(sizeof(block_d2nl) == sizeof(ggml_half) + QK_D2NL/2, "wrong d2nl block size/padding");
++// kvalues_d2nl = round(GGML_D2NL_K * level). This K keeps every level within 0.19% of its exact value.
++#define GGML_D2NL_K 10.981f
++
+ typedef struct {
+     ggml_half d;
+     uint16_t scales_h;
+@@ -1121,6 +1137,11 @@ GGML_TABLE_BEGIN(int8_t, kvalues_iq4nl, 16)
+     -127, -104, -83, -65, -49, -35, -22, -10, 1, 13, 25, 38, 53, 69, 89, 113,
+ GGML_TABLE_END()
+ 
++// Entrit d2 levels {0, +-1, +-2.828, +-5.196, +-8} x GGML_D2NL_K, as int8. Index 4 is zero; 9..15 are unused.
++GGML_TABLE_BEGIN(int8_t, kvalues_d2nl, 16)
++    -88, -57, -31, -11, 0, 11, 31, 57, 88, 0, 0, 0, 0, 0, 0, 0,
++GGML_TABLE_END()
++
+ // e2m1 values (doubled)
+ // ref: https://www.opencompute.org/documents/ocp-microscaling-formats-mx-v1-0-spec-final-pdf
+ GGML_TABLE_BEGIN(int8_t, kvalues_mxfp4, 16)
+--- a/ggml/src/ggml-cpu/ggml-cpu.c
++++ b/ggml/src/ggml-cpu/ggml-cpu.c
+@@ -383,6 +383,12 @@ static const struct ggml_type_traits_cpu type_traits_cpu[GGML_TYPE_COUNT] = {
+         .vec_dot_type             = GGML_TYPE_Q8_0,
+         .nrows                    = 1,
+     },
++    [GGML_TYPE_D2NL] = {
++        .from_float               = quantize_row_d2nl,
++        .vec_dot                  = ggml_vec_dot_d2nl_q8_0,
++        .vec_dot_type             = GGML_TYPE_Q8_0,
++        .nrows                    = 1,
++    },
+     [GGML_TYPE_IQ4_XS] = {
+         .from_float               = quantize_row_iq4_xs,
+         .vec_dot                  = ggml_vec_dot_iq4_xs_q8_K,
+--- a/ggml/src/ggml-cpu/ops.cpp
++++ b/ggml/src/ggml-cpu/ops.cpp
+@@ -685,6 +685,7 @@ void ggml_compute_forward_add(
+         case GGML_TYPE_IQ3_XXS:
+         case GGML_TYPE_IQ1_S:
+         case GGML_TYPE_IQ1_M:
++        case GGML_TYPE_D2NL:
+         case GGML_TYPE_IQ4_NL:
+         case GGML_TYPE_IQ4_XS:
+         case GGML_TYPE_IQ3_S:
+@@ -1137,6 +1138,7 @@ void ggml_compute_forward_add1(
+         case GGML_TYPE_IQ3_XXS:
+         case GGML_TYPE_IQ1_S:
+         case GGML_TYPE_IQ1_M:
++        case GGML_TYPE_D2NL:
+         case GGML_TYPE_IQ4_NL:
+         case GGML_TYPE_IQ4_XS:
+         case GGML_TYPE_IQ3_S:
+@@ -1268,6 +1270,7 @@ void ggml_compute_forward_acc(
+         case GGML_TYPE_IQ3_XXS:
+         case GGML_TYPE_IQ1_S:
+         case GGML_TYPE_IQ1_M:
++        case GGML_TYPE_D2NL:
+         case GGML_TYPE_IQ4_NL:
+         case GGML_TYPE_IQ4_XS:
+         case GGML_TYPE_IQ3_S:
+@@ -4438,6 +4441,7 @@ void ggml_compute_forward_out_prod(
+         case GGML_TYPE_IQ3_XXS:
+         case GGML_TYPE_IQ1_S:
+         case GGML_TYPE_IQ1_M:
++        case GGML_TYPE_D2NL:
+         case GGML_TYPE_IQ4_NL:
+         case GGML_TYPE_IQ4_XS:
+         case GGML_TYPE_IQ3_S:
+@@ -4716,6 +4720,7 @@ void ggml_compute_forward_set(
+         case GGML_TYPE_IQ3_XXS:
+         case GGML_TYPE_IQ1_S:
+         case GGML_TYPE_IQ1_M:
++        case GGML_TYPE_D2NL:
+         case GGML_TYPE_IQ4_NL:
+         case GGML_TYPE_IQ4_XS:
+         case GGML_TYPE_IQ3_S:
+@@ -4941,6 +4946,7 @@ void ggml_compute_forward_get_rows(
+         case GGML_TYPE_IQ3_XXS:
+         case GGML_TYPE_IQ1_S:
+         case GGML_TYPE_IQ1_M:
++        case GGML_TYPE_D2NL:
+         case GGML_TYPE_IQ4_NL:
+         case GGML_TYPE_IQ4_XS:
+         case GGML_TYPE_IQ3_S:
+@@ -5668,6 +5674,7 @@ void ggml_compute_forward_clamp(
+         case GGML_TYPE_IQ3_XXS:
+         case GGML_TYPE_IQ1_S:
+         case GGML_TYPE_IQ1_M:
++        case GGML_TYPE_D2NL:
+         case GGML_TYPE_IQ4_NL:
+         case GGML_TYPE_IQ4_XS:
+         case GGML_TYPE_IQ3_S:
+--- a/ggml/src/ggml-cpu/quants.c
++++ b/ggml/src/ggml-cpu/quants.c
+@@ -1332,6 +1332,31 @@ void quantize_row_iq4_nl(const float * GGML_RESTRICT x, void * GGML_RESTRICT y,
+     quantize_row_iq4_nl_ref(x, y, k);
+ }
+ 
++// d2nl: portable C. GPU backends are the fast path; this keeps CPU-only and offloaded runs correct.
++void quantize_row_d2nl(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) {
++    quantize_row_d2nl_ref(x, (block_d2nl *) y, k);
++}
++
++void ggml_vec_dot_d2nl_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) {
++    assert(nrc == 1);
++    UNUSED(nrc); UNUSED(bs); UNUSED(bx); UNUSED(by);
++    assert(n % QK_D2NL == 0);
++    static_assert(QK_D2NL == QK8_0, "QK_D2NL and QK8_0 must be the same");
++    const block_d2nl * GGML_RESTRICT x = vx;
++    const block_q8_0 * GGML_RESTRICT y = vy;
++    const int nb = n / QK_D2NL;
++    float sumf = 0.0f;
++    for (int ib = 0; ib < nb; ++ib) {
++        int sumi1 = 0, sumi2 = 0;
++        for (int j = 0; j < QK_D2NL/2; ++j) {
++            sumi1 += y[ib].qs[j]             * kvalues_d2nl[x[ib].qs[j] & 0xf];
++            sumi2 += y[ib].qs[j + QK_D2NL/2] * kvalues_d2nl[x[ib].qs[j] >>  4];
++        }
++        sumf += GGML_CPU_FP16_TO_FP32(x[ib].d) * GGML_CPU_FP16_TO_FP32(y[ib].d) * (sumi1 + sumi2);
++    }
++    *s = sumf;
++}
++
+ void quantize_row_iq4_xs(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) {
+     assert(k % QK_K == 0);
+     quantize_iq4_xs(x, y, 1, k, NULL);
+--- a/ggml/src/ggml-cpu/quants.h
++++ b/ggml/src/ggml-cpu/quants.h
+@@ -35,6 +35,8 @@ void quantize_row_tq1_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, i
+ void quantize_row_tq2_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+ 
+ void quantize_row_iq4_nl (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
++void quantize_row_d2nl(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
++void ggml_vec_dot_d2nl_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc);
+ void quantize_row_iq4_xs (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k);
+ 
+ // Dot product
+--- a/ggml/src/ggml-cuda/common.cuh
++++ b/ggml/src/ggml-cuda/common.cuh
+@@ -1105,6 +1105,13 @@ struct ggml_cuda_type_traits<GGML_TYPE_IQ4_NL> {
+     static constexpr int qi = QI4_NL;
+ };
+ 
++template<>
++struct ggml_cuda_type_traits<GGML_TYPE_D2NL> {
++    static constexpr int qk = QK_D2NL;
++    static constexpr int qr = QR_D2NL;
++    static constexpr int qi = QI_D2NL;
++};
++
+ template<>
+ struct ggml_cuda_type_traits<GGML_TYPE_IQ4_XS> {
+     static constexpr int qk = QK_K;
+--- a/ggml/src/ggml-cuda/convert.cu
++++ b/ggml/src/ggml-cuda/convert.cu
+@@ -451,6 +451,24 @@ static __global__ void dequantize_block_iq4_nl(const void * __restrict__ vx, dst
+     }
+ }
+ 
++template<typename dst_t>
++static __global__ void dequantize_block_d2nl(const void * __restrict__ vx, dst_t * __restrict__ yy) {
++
++    const int64_t i   = blockIdx.x;
++    const block_d2nl * x = (const block_d2nl *) vx + i*(QK_K/QK_D2NL);
++
++    const int64_t tid = threadIdx.x;
++    const int64_t il = tid/8; // 0...3
++    const int64_t ib = tid%8; // 0...7
++    dst_t * y = yy + i*QK_K + 32*ib + 4*il;
++    const uint8_t  * q4 = x[ib].qs + 4*il;
++    const float d = (float)x[ib].d;
++    for (int j = 0; j < 4; ++j) {
++        y[j+ 0] = d * kvalues_d2nl[q4[j] & 0xf];
++        y[j+16] = d * kvalues_d2nl[q4[j] >>  4];
++    }
++}
++
+ template<typename dst_t>
+ static __global__ void dequantize_block_iq4_xs(const void * __restrict__ vx, dst_t * __restrict__ yy) {
+     const int64_t i   = blockIdx.x;
+@@ -599,6 +617,12 @@ static void dequantize_row_iq4_nl_cuda(const void * vx, dst_t * y, const int64_t
+     dequantize_block_iq4_nl<<<nb, 32, 0, stream>>>(vx, y);
+ }
+ 
++template<typename dst_t>
++static void dequantize_row_d2nl_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
++    const int nb = (k + QK_K - 1) / QK_K;
++    dequantize_block_d2nl<<<nb, 32, 0, stream>>>(vx, y);
++}
++
+ template<typename dst_t>
+ static void dequantize_row_iq1_m_cuda(const void * vx, dst_t * y, const int64_t k, cudaStream_t stream) {
+     const int nb = k / QK_K;
+@@ -752,6 +776,8 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) {
+             return dequantize_row_iq1_m_cuda;
+         case GGML_TYPE_IQ4_NL:
+             return dequantize_row_iq4_nl_cuda;
++        case GGML_TYPE_D2NL:
++            return dequantize_row_d2nl_cuda;
+         case GGML_TYPE_IQ4_XS:
+             return dequantize_row_iq4_xs_cuda;
+         case GGML_TYPE_IQ3_S:
+@@ -809,6 +835,8 @@ to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) {
+             return dequantize_row_iq1_m_cuda;
+         case GGML_TYPE_IQ4_NL:
+             return dequantize_row_iq4_nl_cuda;
++        case GGML_TYPE_D2NL:
++            return dequantize_row_d2nl_cuda;
+         case GGML_TYPE_IQ4_XS:
+             return dequantize_row_iq4_xs_cuda;
+         case GGML_TYPE_IQ3_S:
+--- a/ggml/src/ggml-cuda/dequantize.cuh
++++ b/ggml/src/ggml-cuda/dequantize.cuh
+@@ -56,6 +56,17 @@ static __device__ __forceinline__ void dequantize_q4_0(const void * vx, const in
+     v.y = (v.y - 8.0f) * d;
+ }
+ 
++static __device__ __forceinline__ void dequantize_d2nl(const void * vx, const int64_t ib, const int iqs, float2 & v){
++    const block_d2nl * x = (const block_d2nl *) vx;
++
++    const float d = x[ib].d;
++
++    const int vui = x[ib].qs[iqs];
++
++    v.x = d * kvalues_d2nl[vui & 0xF];
++    v.y = d * kvalues_d2nl[vui >> 4];
++}
++
+ static __device__ __forceinline__ void dequantize_q4_1(const void * vx, const int64_t ib, const int iqs, float2 & v){
+     const block_q4_1 * x = (const block_q4_1 *) vx;
+ 
+--- a/ggml/src/ggml-cuda/getrows.cu
++++ b/ggml/src/ggml-cuda/getrows.cu
+@@ -205,6 +205,10 @@ static void ggml_cuda_get_rows_switch_src0_type(
+             get_rows_cuda_q<QK2_0, QR2_0, dequantize_q2_0>(src0_d, src1_d, dst_d,
+                 ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
+             break;
++        case GGML_TYPE_D2NL:
++            get_rows_cuda_q<QK_D2NL, QR_D2NL, dequantize_d2nl>(src0_d, src1_d, dst_d,
++                ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
++            break;
+         case GGML_TYPE_Q4_0:
+             get_rows_cuda_q<QK4_0, QR4_0, dequantize_q4_0>(src0_d, src1_d, dst_d,
+                 ne00, nb01, nb02, nb03, ne10, ne11, ne12, nb10, nb11, nb12, nb1, nb2, nb3, stream);
+--- a/ggml/src/ggml-cuda/ggml-cuda.cu
++++ b/ggml/src/ggml-cuda/ggml-cuda.cu
+@@ -5179,6 +5179,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
+                     case GGML_TYPE_IQ2_XXS:
+                     case GGML_TYPE_IQ3_S:
+                     case GGML_TYPE_IQ3_XXS:
++                    case GGML_TYPE_D2NL:
+                     case GGML_TYPE_IQ4_NL:
+                     case GGML_TYPE_IQ4_XS:
+                     case GGML_TYPE_BF16:
+@@ -5198,6 +5199,7 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g
+                     case GGML_TYPE_I32:
+                     case GGML_TYPE_Q1_0:
+                     case GGML_TYPE_Q2_0:
++                    case GGML_TYPE_D2NL:
+                     case GGML_TYPE_Q4_0:
+                     case GGML_TYPE_Q4_1:
+                     case GGML_TYPE_Q5_0:
+--- a/ggml/src/ggml-cuda/mmq.cu
++++ b/ggml/src/ggml-cuda/mmq.cu
+@@ -71,6 +71,9 @@ static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, con
+         case GGML_TYPE_IQ4_NL:
+             mul_mat_q_case<GGML_TYPE_IQ4_NL>(ctx, args, stream);
+             break;
++        case GGML_TYPE_D2NL:
++            mul_mat_q_case<GGML_TYPE_D2NL>(ctx, args, stream);
++            break;
+         default:
+             GGML_ABORT("fatal error");
+             break;
+@@ -299,6 +302,9 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t
+         case GGML_TYPE_IQ4_NL:
+             mmq_supported = true;
+             break;
++        case GGML_TYPE_D2NL:
++            mmq_supported = true;
++            break;
+         default:
+             mmq_supported = false;
+             break;
+--- a/ggml/src/ggml-cuda/mmq.cuh
++++ b/ggml/src/ggml-cuda/mmq.cuh
+@@ -95,6 +95,8 @@ static mmq_q8_1_ds_layout mmq_get_q8_1_ds_layout(const ggml_type type_x) {
+         case GGML_TYPE_IQ4_XS:
+         case GGML_TYPE_IQ4_NL:
+             return MMQ_Q8_1_DS_LAYOUT_D4;
++        case GGML_TYPE_D2NL:
++            return MMQ_Q8_1_DS_LAYOUT_D4;
+         default:
+             GGML_ABORT("fatal error");
+             break;
+@@ -214,6 +216,7 @@ static constexpr __host__ __device__ tile_x_sizes mmq_get_dp4a_tile_x_sizes(ggml
+         case GGML_TYPE_IQ1_S:   return MMQ_DP4A_TXS_Q8_0;
+         case GGML_TYPE_IQ4_XS:  return MMQ_DP4A_TXS_Q8_0;
+         case GGML_TYPE_IQ4_NL:  return MMQ_DP4A_TXS_Q8_0;
++        case GGML_TYPE_D2NL:  return MMQ_DP4A_TXS_Q8_0;
+         default:                return tile_x_sizes{0, 0, 0};
+     }
+ }
+@@ -265,6 +268,7 @@ static constexpr __host__ __device__ int mmq_get_mma_tile_x_k(ggml_type type) {
+         case GGML_TYPE_IQ1_S:   return MMQ_MMA_TILE_X_K_Q8_0;
+         case GGML_TYPE_IQ4_XS:  return MMQ_MMA_TILE_X_K_Q8_0;
+         case GGML_TYPE_IQ4_NL:  return MMQ_MMA_TILE_X_K_Q8_0;
++        case GGML_TYPE_D2NL:  return MMQ_MMA_TILE_X_K_Q8_0;
+         default:                return 0;
+     }
+ }
+@@ -2836,6 +2840,71 @@ template <int mmq_y, bool need_check> static __device__ __forceinline__ void loa
+     }
+ }
+ 
++template <int mmq_y, bool need_check> static __device__ __forceinline__ void load_tiles_d2nl(
++    const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) {
++    constexpr int nwarps = mmq_get_nwarps_device();
++    constexpr int warp_size = ggml_cuda_get_physical_warp_size();
++
++#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
++    int   * x_qs = (int   *)  x_tile;
++    float * x_df = (float *) (x_qs + MMQ_TILE_NE_K*2);
++#else
++    constexpr tile_x_sizes txs = mmq_get_dp4a_tile_x_sizes(GGML_TYPE_D2NL, mmq_y);
++    int   * x_qs = (int   *)  x_tile;
++    float * x_df = (float *) (x_qs + txs.qs);
++#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
++
++    constexpr int threads_per_row = MMQ_ITER_K / (4 * QR_D2NL);
++    constexpr int nrows = warp_size / threads_per_row;
++    const int txi = warp_size > threads_per_row ? threadIdx.x % threads_per_row : threadIdx.x;
++    const int kbx  = txi / QI_D2NL;
++    const int kqsx = txi % QI_D2NL;
++
++#pragma unroll
++    for (int i0 = 0; i0 < mmq_y; i0 += nrows*nwarps) {
++        int i = i0 + (nrows == 1 ? threadIdx.y : threadIdx.y*nrows + threadIdx.x/threads_per_row);
++
++        if (need_check) {
++            i = min(i, i_max);
++        }
++
++        const block_d2nl * bxi = (const block_d2nl *) x + kbx0 + i*stride + kbx;
++
++        const int aux_q4 = get_int_b2(bxi->qs, kqsx);
++        const int2 v = get_int_from_table_16(aux_q4, kvalues_d2nl);
++        const int k0 = kbx * (2 * QI_D2NL) + kqsx;
++
++#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
++        x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + k0 + 0]      = v.x;
++        x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + k0 + QI_D2NL] = v.y;
++#else
++        x_qs[i*(2*MMQ_TILE_NE_K + 1) + k0 + 0]      = v.x;
++        x_qs[i*(2*MMQ_TILE_NE_K + 1) + k0 + QI_D2NL] = v.y;
++#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
++    }
++
++    constexpr int blocks_per_tile_x_row = MMQ_TILE_NE_K / QI_D2NL;
++    constexpr int rows_per_warp = warp_size / blocks_per_tile_x_row;
++    const int kbxd = threadIdx.x % blocks_per_tile_x_row;
++
++#pragma unroll
++    for (int i0 = 0; i0 < mmq_y; i0 += nwarps * rows_per_warp) {
++        int i = i0 + threadIdx.y * rows_per_warp + threadIdx.x / blocks_per_tile_x_row;
++
++        if (need_check) {
++            i = min(i, i_max);
++        }
++
++        const block_d2nl * bxi = (const block_d2nl *) x + kbx0 + i*stride + kbxd;
++
++#if defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
++        x_df[i*MMQ_MMA_TILE_X_K_Q8_0             + kbxd] = __half2float(bxi->d);
++#else
++        x_df[i*(MMQ_TILE_NE_K/QI_D2NL) + i/QI_D2NL + kbxd] = __half2float(bxi->d);
++#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE)
++    }
++}
++
+ template <int mmq_y, bool need_check> static __device__ __forceinline__ void load_tiles_iq2_xxs(
+     const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) {
+     constexpr int nwarps = mmq_get_nwarps_device();
+@@ -3541,6 +3610,14 @@ struct mmq_type_traits<mmq_x, mmq_y, need_check, GGML_TYPE_IQ4_NL> {
+     static constexpr vec_dot_mmq_t    vec_dot_dp4a = vec_dot_q8_0_q8_1_dp4a<mmq_x, mmq_y>;
+ };
+ 
++template <int mmq_x, int mmq_y, bool need_check>
++struct mmq_type_traits<mmq_x, mmq_y, need_check, GGML_TYPE_D2NL> {
++    static constexpr int              vdr          = VDR_D2NL_Q8_1_MMQ;
++    static constexpr load_tiles_mmq_t load_tiles   = load_tiles_d2nl<mmq_y, need_check>;
++    static constexpr vec_dot_mmq_t    vec_dot_mma  = vec_dot_q8_0_q8_1_mma<mmq_x, mmq_y, MMQ_Q8_1_DS_LAYOUT_D4>;
++    static constexpr vec_dot_mmq_t    vec_dot_dp4a = vec_dot_q8_0_q8_1_dp4a<mmq_x, mmq_y>;
++};
++
+ template <int mmq_x, int mmq_y, bool need_check>
+ struct mmq_type_traits<mmq_x, mmq_y, need_check, GGML_TYPE_IQ4_XS> {
+     static constexpr int              vdr          = VDR_IQ4_XS_Q8_1_MMQ;
+@@ -4265,6 +4342,7 @@ extern DECL_MMQ_CASE(GGML_TYPE_IQ3_XXS);
+ extern DECL_MMQ_CASE(GGML_TYPE_IQ3_S);
+ extern DECL_MMQ_CASE(GGML_TYPE_IQ1_S);
+ extern DECL_MMQ_CASE(GGML_TYPE_IQ4_NL);
++extern DECL_MMQ_CASE(GGML_TYPE_D2NL);
+ extern DECL_MMQ_CASE(GGML_TYPE_IQ4_XS);
+ 
+ // -------------------------------------------------------------------------------------------------------------------------
+--- a/ggml/src/ggml-cuda/mmvq.cu
++++ b/ggml/src/ggml-cuda/mmvq.cu
+@@ -30,6 +30,7 @@ static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type)
+         case GGML_TYPE_IQ1_S:   return vec_dot_iq1_s_q8_1;
+         case GGML_TYPE_IQ1_M:   return vec_dot_iq1_m_q8_1;
+         case GGML_TYPE_IQ4_NL:  return vec_dot_iq4_nl_q8_1;
++        case GGML_TYPE_D2NL:  return vec_dot_d2nl_q8_1;
+         case GGML_TYPE_IQ4_XS:  return vec_dot_iq4_xs_q8_1;
+         case GGML_TYPE_IQ3_S:   return vec_dot_iq3_s_q8_1;
+         default:                return nullptr;
+@@ -58,6 +59,7 @@ static constexpr __host__ __device__ int get_vdr_mmvq(ggml_type type) {
+         case GGML_TYPE_IQ3_XXS: return VDR_IQ3_XXS_Q8_1_MMVQ;
+         case GGML_TYPE_IQ3_S:   return VDR_IQ3_S_Q8_1_MMVQ;
+         case GGML_TYPE_IQ4_NL:  return VDR_IQ4_NL_Q8_1_MMVQ;
++        case GGML_TYPE_D2NL:  return VDR_D2NL_Q8_1_MMVQ;
+         case GGML_TYPE_IQ4_XS:  return VDR_IQ4_XS_Q8_1_MMVQ;
+         default:                return 1;
+     }
+@@ -121,6 +123,7 @@ static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_pascal_older(gg
+         case GGML_TYPE_IQ3_S:   return 4;
+         case GGML_TYPE_IQ3_XXS: return 4;
+         case GGML_TYPE_IQ4_NL:  return 6;
++        case GGML_TYPE_D2NL:  return 6;
+         case GGML_TYPE_IQ4_XS:  return 5;
+         case GGML_TYPE_MXFP4:   return 4;
+         case GGML_TYPE_NVFP4:   return 4;
+@@ -161,6 +164,7 @@ static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_gcn(ggml_type t
+         case GGML_TYPE_IQ3_S:   return 4;
+         case GGML_TYPE_IQ3_XXS: return 4;
+         case GGML_TYPE_IQ4_NL:  return 6;
++        case GGML_TYPE_D2NL:  return 6;
+         case GGML_TYPE_IQ4_XS:  return 4;
+         case GGML_TYPE_Q2_K:    return 4;
+         case GGML_TYPE_Q3_K:    return 4;
+@@ -211,6 +215,7 @@ static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_rdna3(ggml_type
+         case GGML_TYPE_IQ3_S:   return 4;
+         case GGML_TYPE_IQ3_XXS: return 4;
+         case GGML_TYPE_IQ4_NL:  return 6;
++        case GGML_TYPE_D2NL:  return 6;
+         case GGML_TYPE_IQ4_XS:  return 6;
+         case GGML_TYPE_Q4_K:    return 4;
+         case GGML_TYPE_Q5_K:    return 4;
+@@ -229,6 +234,7 @@ static constexpr __host__ __device__ int get_mmvq_mmid_max_batch_rdna4(ggml_type
+         case GGML_TYPE_IQ3_S:   return 4;
+         case GGML_TYPE_IQ3_XXS: return 4;
+         case GGML_TYPE_IQ4_NL:  return 7;
++        case GGML_TYPE_D2NL:  return 7;
+         case GGML_TYPE_IQ4_XS:  return 5;
+         case GGML_TYPE_MXFP4:   return 5;
+         case GGML_TYPE_NVFP4:   return 5;
+@@ -394,6 +400,7 @@ static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_d
+                 case GGML_TYPE_Q4_K:
+                 case GGML_TYPE_Q5_K:
+                 case GGML_TYPE_Q6_K:
++                case GGML_TYPE_D2NL:
+                 case GGML_TYPE_IQ4_NL:
+                 case GGML_TYPE_IQ4_XS:
+                     return 8;
+@@ -419,6 +426,8 @@ static constexpr __host__ __device__ int calc_nwarps(ggml_type type, int ncols_d
+                     return 2;
+                 case GGML_TYPE_IQ4_NL:
+                     return 8;
++                case GGML_TYPE_D2NL:
++                    return 8;
+                 default:
+                     return 1;
+             }
+@@ -1111,6 +1120,12 @@ static void mul_mat_vec_q_switch_type(
+                  nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst,
+                  nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream);
+             break;
++        case GGML_TYPE_D2NL:
++            mul_mat_vec_q_switch_ncols_dst<GGML_TYPE_D2NL>
++                (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst,
++                 nchannels_x, nchannels_y, nchannels_dst, stride_channel_x, stride_channel_y, stride_channel_dst,
++                 nsamples_x, nsamples_dst, stride_sample_x, stride_sample_y, stride_sample_dst, ids_stride, stream);
++            break;
+         case GGML_TYPE_IQ4_XS:
+             mul_mat_vec_q_switch_ncols_dst<GGML_TYPE_IQ4_XS>
+                 (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst,
+--- /dev/null
++++ b/ggml/src/ggml-cuda/template-instances/mmq-instance-d2nl.cu
+@@ -0,0 +1,5 @@
++// Entrit d2nl: the IQ4_NL MMQ path with the d2 table.
++
++#include "../mmq.cuh"
++
++DECL_MMQ_CASE(GGML_TYPE_D2NL);
+--- a/ggml/src/ggml-cuda/vecdotq.cuh
++++ b/ggml/src/ggml-cuda/vecdotq.cuh
+@@ -1336,6 +1336,8 @@ static __device__ __forceinline__ float vec_dot_iq1_m_q8_1(
+ 
+ #define VDR_IQ4_NL_Q8_1_MMVQ 2
+ #define VDR_IQ4_NL_Q8_1_MMQ  4
++#define VDR_D2NL_Q8_1_MMVQ 2
++#define VDR_D2NL_Q8_1_MMQ  4
+ 
+ static __device__ __forceinline__ float vec_dot_iq4_nl_q8_1(
+     const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {
+@@ -1358,6 +1360,27 @@ static __device__ __forceinline__ float vec_dot_iq4_nl_q8_1(
+     return d * sumi;
+ }
+ 
++static __device__ __forceinline__ float vec_dot_d2nl_q8_1(
++    const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) {
++
++    const block_d2nl * bq4 = (const block_d2nl *) vbq + kbx;
++
++    const int * q8 = (const int *) bq8_1->qs + iqs;
++
++    int sumi = 0;
++#pragma unroll
++    for (int l = 0; l < VDR_Q4_0_Q8_1_MMVQ; ++l) {
++        const int aux_q4 = get_int_b2(bq4->qs, iqs + l);
++        const int2 v = get_int_from_table_16(aux_q4, kvalues_d2nl);
++
++        sumi = ggml_cuda_dp4a(v.x, q8[l + 0], sumi);
++        sumi = ggml_cuda_dp4a(v.y, q8[l + 4], sumi);
++    }
++
++    const float d = __half2float(bq4->d) * __low2float(bq8_1->ds);
++    return d * sumi;
++}
++
+ #define VDR_IQ4_XS_Q8_1_MMVQ 4
+ #define VDR_IQ4_XS_Q8_1_MMQ  4
+ 
+--- a/ggml/src/ggml-metal/ggml-metal-device.cpp
++++ b/ggml/src/ggml-metal/ggml-metal-device.cpp
+@@ -993,6 +993,12 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta
+                 nr0 = N_R0_IQ4_NL;
+                 smem = 32*sizeof(float);
+             } break;
++        case GGML_TYPE_D2NL:
++            {
++                nsg = N_SG_D2NL;
++                nr0 = N_R0_D2NL;
++                smem = 32*sizeof(float);
++            } break;
+         case GGML_TYPE_IQ4_XS:
+             {
+                 nsg = N_SG_IQ4_XS;
+@@ -1222,6 +1228,12 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id(ggml_m
+                 nr0 = N_R0_IQ4_NL;
+                 smem = 32*sizeof(float);
+             } break;
++        case GGML_TYPE_D2NL:
++            {
++                nsg = N_SG_D2NL;
++                nr0 = N_R0_D2NL;
++                smem = 32*sizeof(float);
++            } break;
+         case GGML_TYPE_IQ4_XS:
+             {
+                 nsg = N_SG_IQ4_XS;
+--- a/ggml/src/ggml-metal/ggml-metal-impl.h
++++ b/ggml/src/ggml-metal/ggml-metal-impl.h
+@@ -84,6 +84,9 @@
+ #define N_R0_IQ4_NL 2
+ #define N_SG_IQ4_NL 2
+ 
++#define N_R0_D2NL 2
++#define N_SG_D2NL 2
++
+ #define N_R0_IQ4_XS 2
+ #define N_SG_IQ4_XS 2
+ 
+--- a/ggml/src/ggml-metal/ggml-metal-ops.cpp
++++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp
+@@ -2247,6 +2247,7 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) {
+            op->src[0]->type == GGML_TYPE_Q8_0 ||
+            op->src[0]->type == GGML_TYPE_MXFP4 ||
+            op->src[0]->type == GGML_TYPE_IQ4_NL ||
++           op->src[0]->type == GGML_TYPE_D2NL ||
+            false) && (ne11 >= 2 && ne11 <= ne11_ext_max)
+          ) ||
+          (
+--- a/ggml/src/ggml-metal/ggml-metal.metal
++++ b/ggml/src/ggml-metal/ggml-metal.metal
+@@ -46,6 +46,11 @@ constexpr constant static float kvalues_iq4nl_f[16] = {
+     -127.f, -104.f, -83.f, -65.f, -49.f, -35.f, -22.f, -10.f, 1.f, 13.f, 25.f, 38.f, 53.f, 69.f, 89.f, 113.f
+ };
+ 
++// Entrit d2nl: d2 levels {0, +-1, +-2.828, +-5.196, +-8} x 10.981; index 4 is zero (kvalues_d2nl in ggml-common.h)
++constexpr constant static float kvalues_d2nl_f[16] = {
++    -88.f, -57.f, -31.f, -11.f, 0.f, 11.f, 31.f, 57.f, 88.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f, 0.f
++};
++
+ constexpr constant static float kvalues_mxfp4_f[16] = {
+     0, .5f, 1.f, 1.5f, 2.f, 3.f, 4.f, 6.f, -0, -.5f, -1.f, -1.5f, -2.f, -3.f, -4.f, -6.f
+ };
+@@ -986,6 +991,21 @@ void dequantize_iq4_nl(device const block_iq4_nl * xb, short il, thread type4x4
+     }
+ }
+ 
++template <typename type4x4>
++void dequantize_d2nl(device const block_d2nl * xb, short il, thread type4x4 & reg) {
++    device const uint16_t * q4 = (device const uint16_t *)xb->qs;
++    const float d = xb->d;
++    uint32_t aux32;
++    thread const uint8_t * q8 = (thread const uint8_t *)&aux32;
++    for (int i = 0; i < 4; ++i) {
++        aux32 = ((q4[2*i] | (q4[2*i+1] << 16)) >> 4*il) & 0x0f0f0f0f;
++        reg[i][0] = d * kvalues_d2nl_f[q8[0]];
++        reg[i][1] = d * kvalues_d2nl_f[q8[1]];
++        reg[i][2] = d * kvalues_d2nl_f[q8[2]];
++        reg[i][3] = d * kvalues_d2nl_f[q8[3]];
++    }
++}
++
+ template <typename type4>
+ void dequantize_iq4_nl_t4(device const block_iq4_nl * xb, short il, thread type4 & reg) {
+     device const uint16_t * q4 = (device const uint16_t *)xb->qs;
+@@ -999,6 +1019,19 @@ void dequantize_iq4_nl_t4(device const block_iq4_nl * xb, short il, thread type4
+     reg[3] = d * kvalues_iq4nl_f[q8[3]];
+ }
+ 
++template <typename type4>
++void dequantize_d2nl_t4(device const block_d2nl * xb, short il, thread type4 & reg) {
++    device const uint16_t * q4 = (device const uint16_t *)xb->qs;
++    const float d = xb->d;
++    uint32_t aux32;
++    thread const uint8_t * q8 = (thread const uint8_t *)&aux32;
++    aux32 = ((q4[2*(il%4)] | (q4[2*(il%4)+1] << 16)) >> 4*(il/4)) & 0x0f0f0f0f;
++    reg[0] = d * kvalues_d2nl_f[q8[0]];
++    reg[1] = d * kvalues_d2nl_f[q8[1]];
++    reg[2] = d * kvalues_d2nl_f[q8[2]];
++    reg[3] = d * kvalues_d2nl_f[q8[3]];
++}
++
+ template <typename type4x4>
+ void dequantize_iq4_xs(device const block_iq4_xs * xb, short il, thread type4x4 & reg) {
+     // il is 0...15 for QK_K = 256 => index of block of 32 is il/2
+@@ -4451,9 +4484,13 @@ template [[host_name("kernel_mul_mv_ext_mxfp4_f32_r1_4")]]  kernel mul_mv_ext_q4
+ template [[host_name("kernel_mul_mv_ext_mxfp4_f32_r1_5")]]  kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<5, block_mxfp4,  32, dequantize_mxfp4_t4>;
+ 
+ template [[host_name("kernel_mul_mv_ext_iq4_nl_f32_r1_2")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<2, block_iq4_nl, 32, dequantize_iq4_nl_t4>;
++template [[host_name("kernel_mul_mv_ext_d2nl_f32_r1_2")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<2, block_d2nl, 32, dequantize_d2nl_t4>;
+ template [[host_name("kernel_mul_mv_ext_iq4_nl_f32_r1_3")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<3, block_iq4_nl, 32, dequantize_iq4_nl_t4>;
++template [[host_name("kernel_mul_mv_ext_d2nl_f32_r1_3")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<3, block_d2nl, 32, dequantize_d2nl_t4>;
+ template [[host_name("kernel_mul_mv_ext_iq4_nl_f32_r1_4")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<4, block_iq4_nl, 32, dequantize_iq4_nl_t4>;
++template [[host_name("kernel_mul_mv_ext_d2nl_f32_r1_4")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<4, block_d2nl, 32, dequantize_d2nl_t4>;
+ template [[host_name("kernel_mul_mv_ext_iq4_nl_f32_r1_5")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<5, block_iq4_nl, 32, dequantize_iq4_nl_t4>;
++template [[host_name("kernel_mul_mv_ext_d2nl_f32_r1_5")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<5, block_d2nl, 32, dequantize_d2nl_t4>;
+ 
+ template [[host_name("kernel_mul_mv_ext_q4_K_f32_r1_2")]] kernel mul_mv_ext_q4x4_f32_t kernel_mul_mv_ext_q4x4_f32_disp<2, block_q4_K, 256, dequantize_q4_K>;
+ template [[host_name("kernel_mul_mv_ext_q4_K_f32_r1_3")]] kernel mul_mv_ext_q4x4_f32_t kernel_mul_mv_ext_q4x4_f32_disp<3, block_q4_K, 256, dequantize_q4_K>;
+@@ -9484,6 +9521,102 @@ void kernel_mul_mv_iq4_nl_f32_impl(
+     }
+ }
+ 
++template<int NR0, typename args_t>
++void kernel_mul_mv_d2nl_f32_impl(
++        args_t args,
++        device const char * src0,
++        device const char * src1,
++        device       char * dst,
++        threadgroup  char * shmem,
++        uint3  tgpig,
++        ushort tiisg,
++        ushort sgitg) {
++    const short NSG = FC_mul_mv_nsg;
++
++    threadgroup float * shmem_f32 = (threadgroup float *) shmem;
++
++    const int r0 = tgpig.x;
++    const int r1 = tgpig.y;
++    const int im = tgpig.z;
++
++    const int first_row = (r0 * NSG + sgitg) * NR0;
++
++    const uint i12 = im%FC_mul_mv_ne12;
++    const uint i13 = im/FC_mul_mv_ne12;
++
++    const uint64_t offset0 = first_row*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03;
++    const uint64_t offset1 =        r1*args.nb11 + (i12        )*args.nb12 + (i13        )*args.nb13;
++
++    device const block_d2nl * x = (device const block_d2nl *) (src0 + offset0);
++    device const float        * y = (device const float        *) (src1 + offset1);
++
++    const int nb   = args.ne00/QK_D2NL;
++    const int ns01 = args.nb01/args.nb00;
++
++    const short ix = tiisg/2;  // 0...15
++    const short it = tiisg%2;  // 0 or 1
++
++    shmem_f32[tiisg] = kvalues_d2nl_f[tiisg%16];
++    threadgroup_barrier(mem_flags::mem_threadgroup);
++
++    float4 yl[4];
++    float sumf[NR0]={0.f};
++
++    device const float * yb = y + ix*QK_D2NL + it*8;
++
++    uint32_t aux32[2];
++    thread const uint8_t * q8 = (thread const uint8_t *)aux32;
++
++    float4 qf1, qf2;
++
++    // [TAG_MUL_MV_WEIRD]
++    for (int ib = ix; ib < nb && ib < ns01; ib += 16) {
++        device const float4 * y4 = (device const float4 *)yb;
++        yl[0] = y4[0];
++        yl[1] = y4[4];
++        yl[2] = y4[1];
++        yl[3] = y4[5];
++
++        for (short row = 0; row < NR0; row++) {
++            device const block_d2nl & xb = x[row*ns01 + ib];
++            device const uint16_t * q4 = (device const uint16_t *)(xb.qs + 8*it);
++
++            float4 acc1 = {0.f}, acc2 = {0.f};
++
++            aux32[0] = q4[0] | (q4[1] << 16);
++            aux32[1] = (aux32[0] >> 4) & 0x0f0f0f0f;
++            aux32[0] &= 0x0f0f0f0f;
++            qf1 = {shmem_f32[q8[0]], shmem_f32[q8[1]], shmem_f32[q8[2]], shmem_f32[q8[3]]};
++            qf2 = {shmem_f32[q8[4]], shmem_f32[q8[5]], shmem_f32[q8[6]], shmem_f32[q8[7]]};
++            acc1 += yl[0] * qf1;
++            acc2 += yl[1] * qf2;
++
++            aux32[0] = q4[2] | (q4[3] << 16);
++            aux32[1] = (aux32[0] >> 4) & 0x0f0f0f0f;
++            aux32[0] &= 0x0f0f0f0f;
++            qf1 = {shmem_f32[q8[0]], shmem_f32[q8[1]], shmem_f32[q8[2]], shmem_f32[q8[3]]};
++            qf2 = {shmem_f32[q8[4]], shmem_f32[q8[5]], shmem_f32[q8[6]], shmem_f32[q8[7]]};
++            acc1 += yl[2] * qf1;
++            acc2 += yl[3] * qf2;
++
++            acc1 += acc2;
++
++            sumf[row] += (float)xb.d * (acc1[0] + acc1[1] + acc1[2] + acc1[3]);
++        }
++
++        yb += 16 * QK_D2NL;
++    }
++
++    device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
++
++    for (int row = 0; row < NR0 && first_row + row < args.ne0; ++row) {
++        float sum_all = simd_sum(sumf[row]);
++        if (tiisg == 0) {
++            dst_f32[first_row + row] = sum_all;
++        }
++    }
++}
++
+ [[host_name("kernel_mul_mv_iq4_nl_f32")]]
+ kernel void kernel_mul_mv_iq4_nl_f32(
+         constant ggml_metal_kargs_mul_mv & args,
+@@ -9498,6 +9631,20 @@ kernel void kernel_mul_mv_iq4_nl_f32(
+     kernel_mul_mv_iq4_nl_f32_impl<N_R0_IQ4_NL, constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
+ }
+ 
++[[host_name("kernel_mul_mv_d2nl_f32")]]
++kernel void kernel_mul_mv_d2nl_f32(
++        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]]) {
++
++    kernel_mul_mv_d2nl_f32_impl<N_R0_D2NL, constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
++}
++
+ template<int NR0, typename args_t>
+ void kernel_mul_mv_iq4_xs_f32_impl(
+         args_t args,
+@@ -10711,6 +10858,7 @@ template [[host_name("kernel_get_rows_iq2_s")]]   kernel get_rows_q_t kernel_get
+ template [[host_name("kernel_get_rows_iq1_s")]]   kernel get_rows_q_t kernel_get_rows_q<block_iq1_s,   QK_NL, dequantize_iq1_s>;
+ template [[host_name("kernel_get_rows_iq1_m")]]   kernel get_rows_q_t kernel_get_rows_q<block_iq1_m,   QK_NL, dequantize_iq1_m>;
+ template [[host_name("kernel_get_rows_iq4_nl")]]  kernel get_rows_q_t kernel_get_rows_q<block_iq4_nl,  2,     dequantize_iq4_nl>;
++template [[host_name("kernel_get_rows_d2nl")]]  kernel get_rows_q_t kernel_get_rows_q<block_d2nl,  2,     dequantize_d2nl>;
+ template [[host_name("kernel_get_rows_iq4_xs")]]  kernel get_rows_q_t kernel_get_rows_q<block_iq4_xs,  QK_NL, dequantize_iq4_xs>;
+ 
+ //
+@@ -10775,6 +10923,7 @@ template [[host_name("kernel_mul_mm_iq2_s_f32")]]   kernel mul_mm_t kernel_mul_m
+ template [[host_name("kernel_mul_mm_iq1_s_f32")]]   kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq1_s,   QK_NL, dequantize_iq1_s,   float,  float4x4,  float, float2x4>;
+ template [[host_name("kernel_mul_mm_iq1_m_f32")]]   kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq1_m,   QK_NL, dequantize_iq1_m,   float,  float4x4,  float, float2x4>;
+ template [[host_name("kernel_mul_mm_iq4_nl_f32")]]  kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq4_nl,  2,     dequantize_iq4_nl,  float,  float4x4,  float, float2x4>;
++template [[host_name("kernel_mul_mm_d2nl_f32")]]  kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_d2nl,  2,     dequantize_d2nl,  float,  float4x4,  float, float2x4>;
+ template [[host_name("kernel_mul_mm_iq4_xs_f32")]]  kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq4_xs,  QK_NL, dequantize_iq4_xs,  float,  float4x4,  float, float2x4>;
+ 
+ template [[host_name("kernel_mul_mm_f32_f16")]]     kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   float4x4,      1,     dequantize_f32,     float,  float4x4,  half, half2x4>;
+@@ -10800,6 +10949,7 @@ template [[host_name("kernel_mul_mm_iq2_s_f16")]]   kernel mul_mm_t kernel_mul_m
+ template [[host_name("kernel_mul_mm_iq1_s_f16")]]   kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq1_s,   QK_NL, dequantize_iq1_s,   float,  float4x4,  half, half2x4>;
+ template [[host_name("kernel_mul_mm_iq1_m_f16")]]   kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq1_m,   QK_NL, dequantize_iq1_m,   float,  float4x4,  half, half2x4>;
+ template [[host_name("kernel_mul_mm_iq4_nl_f16")]]  kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq4_nl,  2,     dequantize_iq4_nl,  float,  float4x4,  half, half2x4>;
++template [[host_name("kernel_mul_mm_d2nl_f16")]]  kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_d2nl,  2,     dequantize_d2nl,  float,  float4x4,  half, half2x4>;
+ template [[host_name("kernel_mul_mm_iq4_xs_f16")]]  kernel mul_mm_t kernel_mul_mm<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq4_xs,  QK_NL, dequantize_iq4_xs,  float,  float4x4,  half, half2x4>;
+ 
+ //
+@@ -10834,6 +10984,7 @@ template [[host_name("kernel_mul_mm_id_iq2_s_f32")]]   kernel mul_mm_id kernel_m
+ template [[host_name("kernel_mul_mm_id_iq1_s_f32")]]   kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq1_s,   QK_NL, dequantize_iq1_s,   float,  float4x4,  float, float2x4>;
+ template [[host_name("kernel_mul_mm_id_iq1_m_f32")]]   kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq1_m,   QK_NL, dequantize_iq1_m,   float,  float4x4,  float, float2x4>;
+ template [[host_name("kernel_mul_mm_id_iq4_nl_f32")]]  kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq4_nl,  2,     dequantize_iq4_nl,  float,  float4x4,  float, float2x4>;
++template [[host_name("kernel_mul_mm_id_d2nl_f32")]]  kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_d2nl,  2,     dequantize_d2nl,  float,  float4x4,  float, float2x4>;
+ template [[host_name("kernel_mul_mm_id_iq4_xs_f32")]]  kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq4_xs,  QK_NL, dequantize_iq4_xs,  float,  float4x4,  float, float2x4>;
+ 
+ template [[host_name("kernel_mul_mm_id_f32_f16")]]     kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   float4x4,      1,     dequantize_f32,     float,  float4x4,  half, half2x4>;
+@@ -10859,6 +11010,7 @@ template [[host_name("kernel_mul_mm_id_iq2_s_f16")]]   kernel mul_mm_id kernel_m
+ template [[host_name("kernel_mul_mm_id_iq1_s_f16")]]   kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq1_s,   QK_NL, dequantize_iq1_s,   float,  float4x4,  half, half2x4>;
+ template [[host_name("kernel_mul_mm_id_iq1_m_f16")]]   kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq1_m,   QK_NL, dequantize_iq1_m,   float,  float4x4,  half, half2x4>;
+ template [[host_name("kernel_mul_mm_id_iq4_nl_f16")]]  kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq4_nl,  2,     dequantize_iq4_nl,  float,  float4x4,  half, half2x4>;
++template [[host_name("kernel_mul_mm_id_d2nl_f16")]]  kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_d2nl,  2,     dequantize_d2nl,  float,  float4x4,  half, half2x4>;
+ template [[host_name("kernel_mul_mm_id_iq4_xs_f16")]]  kernel mul_mm_id kernel_mul_mm_id<half,   half4x4,   simdgroup_half8x8,   half,   half2x4,   simdgroup_half8x8,   block_iq4_xs,  QK_NL, dequantize_iq4_xs,  float,  float4x4,  half, half2x4>;
+ 
+ //
+@@ -11016,6 +11168,7 @@ template [[host_name("kernel_mul_mv_id_iq3_xxs_f32")]] kernel kernel_mul_mv_id_t
+ template [[host_name("kernel_mul_mv_id_iq3_s_f32")]]   kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq3_s_f32_impl  <N_R0_IQ3_S>>>;
+ template [[host_name("kernel_mul_mv_id_iq2_s_f32")]]   kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq2_s_f32_impl  <N_R0_IQ2_S>>>;
+ template [[host_name("kernel_mul_mv_id_iq4_nl_f32")]]  kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq4_nl_f32_impl <N_R0_IQ4_NL>>>;
++template [[host_name("kernel_mul_mv_id_d2nl_f32")]]  kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_d2nl_f32_impl <N_R0_D2NL>>>;
+ template [[host_name("kernel_mul_mv_id_iq4_xs_f32")]]  kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq4_xs_f32_impl <N_R0_IQ4_XS>>>;
+ 
+ kernel void kernel_pool_2d_max_f32(
+--- a/ggml/src/ggml-quants.c
++++ b/ggml/src/ggml-quants.c
+@@ -5096,6 +5096,80 @@ size_t quantize_iq4_nl(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst
+     return nrow * nblock * sizeof(block_iq4_nl);
+ }
+ 
++// ============================ d2nl (Entrit d2) ============================
++// Converts weights that already lie on the d2 grid. Each 64-group's scale is recovered by trying the
++// four possible top levels and keeping the one that reconstructs the group best; it is written into
++// both 32-weight blocks of the group.
++static const float ggml_d2nl_levels[9] = { -8.0f, -5.19615242f, -2.82842712f, -1.0f, 0.0f,
++                                            1.0f,  2.82842712f,  5.19615242f,  8.0f };
++
++void quantize_row_d2nl_ref(const float * GGML_RESTRICT x, block_d2nl * GGML_RESTRICT y, int64_t k) {
++    assert(k % QK_D2NL == 0);
++    const int64_t nb = k / QK_D2NL;
++    static const float tops[4] = { 8.0f, 5.19615242f, 2.82842712f, 1.0f };
++    for (int64_t ib = 0; ib < nb; ib += 2) {
++        const int n = (ib + 1 < nb) ? 2*QK_D2NL : QK_D2NL;
++        const float * xg = x + ib*QK_D2NL;
++        float amax = 0.0f;
++        for (int j = 0; j < n; j++) {
++            const float a = fabsf(xg[j]);
++            if (a > amax) amax = a;
++        }
++        uint8_t idx[2*QK_D2NL];
++        float   d = 0.0f;
++        float   bestres = INFINITY;
++        for (int t = 0; t < 4; t++) {
++            const float dt = amax / tops[t];
++            const float it = dt ? 1.0f/dt : 0.0f;
++            uint8_t cand[2*QK_D2NL];
++            float   res = 0.0f;
++            for (int j = 0; j < n; j++) {
++                const float v = xg[j] * it;
++                int   b  = 4;
++                float bd = fabsf(v);
++                for (int c = 0; c < 9; c++) {
++                    const float e = fabsf(v - ggml_d2nl_levels[c]);
++                    if (e < bd) { bd = e; b = c; }
++                }
++                cand[j] = (uint8_t) b;
++                res += bd;
++            }
++            if (res < bestres) { bestres = res; d = dt; memcpy(idx, cand, n); }
++        }
++        const ggml_half dh = GGML_FP32_TO_FP16(d / GGML_D2NL_K);
++        for (int b = 0; b < n/QK_D2NL; b++) {
++            y[ib + b].d = dh;
++            for (int j = 0; j < QK_D2NL/2; j++) {
++                y[ib + b].qs[j] = idx[b*QK_D2NL + j] | (idx[b*QK_D2NL + j + QK_D2NL/2] << 4);
++            }
++        }
++    }
++}
++
++void dequantize_row_d2nl(const block_d2nl * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) {
++    assert(k % QK_D2NL == 0);
++    const int64_t nb = k / QK_D2NL;
++    for (int64_t i = 0; i < nb; i++) {
++        const float d = GGML_FP16_TO_FP32(x[i].d);
++        for (int j = 0; j < QK_D2NL/2; j++) {
++            y[i*QK_D2NL + j]             = d * kvalues_d2nl[x[i].qs[j] & 0xf];
++            y[i*QK_D2NL + j + QK_D2NL/2] = d * kvalues_d2nl[x[i].qs[j] >>  4];
++        }
++    }
++}
++
++size_t quantize_d2nl(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * quant_weights) {
++    (void) quant_weights;
++    const size_t row_size = ggml_row_size(GGML_TYPE_D2NL, n_per_row);
++    char * qrow = (char *) dst;
++    for (int64_t row = 0; row < nrows; ++row) {
++        quantize_row_d2nl_ref(src, (block_d2nl *) qrow, n_per_row);
++        src  += n_per_row;
++        qrow += row_size;
++    }
++    return nrows * row_size;
++}
++
+ //void quantize_row_iq4_nl_ref(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) {
+ void quantize_row_iq4_nl_ref(const float * GGML_RESTRICT x, block_iq4_nl * GGML_RESTRICT y, int64_t k) {
+     GGML_ASSERT(k%QK4_NL == 0);
+@@ -5649,6 +5723,21 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte
+             {
+                 VALIDATE_ROW_DATA_D_F16_IMPL(block_iq4_nl, data, nb);
+             } break;
++        case GGML_TYPE_D2NL:
++            {
++                const block_d2nl * q = (const block_d2nl *) data;
++                for (size_t i = 0; i < nb; ++i) {
++                    if (!validate_fp16(q[i].d, i)) {
++                        return false;
++                    }
++                    for (int j = 0; j < QK_D2NL/2; ++j) {
++                        if ((q[i].qs[j] & 0xf) > 8 || (q[i].qs[j] >> 4) > 8) {
++                            fprintf(stderr, "ggml_validate_row_data: d2nl block %zu byte %d has a level index > 8\n", i, j);
++                            return false;
++                        }
++                    }
++                }
++            } break;
+ 
+         case GGML_TYPE_I8:
+         case GGML_TYPE_I16:
+--- a/ggml/src/ggml-quants.h
++++ b/ggml/src/ggml-quants.h
+@@ -83,6 +83,9 @@ GGML_API size_t quantize_iq3_xxs(const float * GGML_RESTRICT src, void * GGML_RE
+ GGML_API size_t quantize_iq1_s  (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+ GGML_API size_t quantize_iq1_m  (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+ GGML_API size_t quantize_iq4_nl (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
++GGML_API void   quantize_row_d2nl_ref(const float * GGML_RESTRICT x, block_d2nl * GGML_RESTRICT y, int64_t k);
++GGML_API void   dequantize_row_d2nl(const block_d2nl * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k);
++GGML_API size_t quantize_d2nl(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+ GGML_API size_t quantize_iq4_xs (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+ GGML_API size_t quantize_iq3_s  (const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix);
+ 
+--- a/ggml/src/ggml-vulkan/ggml-vulkan.cpp
++++ b/ggml/src/ggml-vulkan/ggml-vulkan.cpp
+@@ -3391,6 +3391,7 @@ static bool ggml_vk_matmul_shmem_support(const vk_device& device, const std::vec
+     case GGML_TYPE_IQ3_S:
+         lut_size = 4*512;
+         break;
++    case GGML_TYPE_D2NL:
+     case GGML_TYPE_IQ4_NL:
+     case GGML_TYPE_IQ4_XS:
+     case GGML_TYPE_MXFP4:
+@@ -4050,6 +4051,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
+         CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_IQ3_S],   matmul_iq3_s_f16,   mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3)
+         CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_IQ4_XS],  matmul_iq4_xs_f16,  mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3)
+         CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_IQ4_NL],  matmul_iq4_nl_f16,  mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3)
++        CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_D2NL],  matmul_d2nl_f16,  mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3)
+         CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_MXFP4],   matmul_mxfp4_f16,   mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3)
+         CREATE_MM2(pipeline_dequant_mul_mat_mat_f16[GGML_TYPE_NVFP4],   matmul_nvfp4_f16,   mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3)
+ 
+@@ -4082,6 +4084,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
+         CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ3_S],   matmul_id_subgroup_iq3_s_f16,   mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
+         CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ4_XS],  matmul_id_subgroup_iq4_xs_f16,  mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
+         CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ4_NL],  matmul_id_subgroup_iq4_nl_f16,  mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
++        CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_D2NL],  matmul_id_subgroup_d2nl_f16,  mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
+         CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_MXFP4],   matmul_id_subgroup_mxfp4_f16,   mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
+         CREATE_MM2(pipeline_dequant_mul_mat_mat_id[GGML_TYPE_NVFP4],   matmul_id_subgroup_nvfp4_f16,   mmqid_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, 5)
+ #undef CREATE_MM
+@@ -4147,6 +4150,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
+             CREATE_MM2(GGML_TYPE_IQ3_S,   pipeline_dequant_mul_mat_mat[GGML_TYPE_IQ3_S],   matmul_iq3_s_f32,   mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
+             CREATE_MM2(GGML_TYPE_IQ4_XS,  pipeline_dequant_mul_mat_mat[GGML_TYPE_IQ4_XS],  matmul_iq4_xs_f32,  mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
+             CREATE_MM2(GGML_TYPE_IQ4_NL,  pipeline_dequant_mul_mat_mat[GGML_TYPE_IQ4_NL],  matmul_iq4_nl_f32,  mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
++            CREATE_MM2(GGML_TYPE_D2NL,  pipeline_dequant_mul_mat_mat[GGML_TYPE_D2NL],  matmul_d2nl_f32,  mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
+             CREATE_MM2(GGML_TYPE_MXFP4,   pipeline_dequant_mul_mat_mat[GGML_TYPE_MXFP4],   matmul_mxfp4_f32,   mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
+             CREATE_MM2(GGML_TYPE_NVFP4,   pipeline_dequant_mul_mat_mat[GGML_TYPE_NVFP4],   matmul_nvfp4_f32,   mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
+         } else {
+@@ -4172,6 +4176,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
+             CREATE_MM(GGML_TYPE_IQ3_S,   pipeline_dequant_mul_mat_mat[GGML_TYPE_IQ3_S].f32acc,   matmul_iq3_s_f32,   , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
+             CREATE_MM(GGML_TYPE_IQ4_XS,  pipeline_dequant_mul_mat_mat[GGML_TYPE_IQ4_XS].f32acc,  matmul_iq4_xs_f32,  , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
+             CREATE_MM(GGML_TYPE_IQ4_NL,  pipeline_dequant_mul_mat_mat[GGML_TYPE_IQ4_NL].f32acc,  matmul_iq4_nl_f32,  , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
++            CREATE_MM(GGML_TYPE_D2NL,  pipeline_dequant_mul_mat_mat[GGML_TYPE_D2NL].f32acc,  matmul_d2nl_f32,  , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
+             CREATE_MM(GGML_TYPE_MXFP4,   pipeline_dequant_mul_mat_mat[GGML_TYPE_MXFP4].f32acc,   matmul_mxfp4_f32,   , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
+             CREATE_MM(GGML_TYPE_NVFP4,   pipeline_dequant_mul_mat_mat[GGML_TYPE_NVFP4].f32acc,   matmul_nvfp4_f32,   , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, );
+         }
+@@ -4208,6 +4213,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
+         CREATE_MM2(GGML_TYPE_IQ3_S,   pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ3_S],   matmul_id_subgroup_iq3_s_f32,   mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
+         CREATE_MM2(GGML_TYPE_IQ4_XS,  pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ4_XS],  matmul_id_subgroup_iq4_xs_f32,  mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
+         CREATE_MM2(GGML_TYPE_IQ4_NL,  pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ4_NL],  matmul_id_subgroup_iq4_nl_f32,  mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
++        CREATE_MM2(GGML_TYPE_D2NL,  pipeline_dequant_mul_mat_mat_id[GGML_TYPE_D2NL],  matmul_id_subgroup_d2nl_f32,  mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
+         CREATE_MM2(GGML_TYPE_MXFP4,   pipeline_dequant_mul_mat_mat_id[GGML_TYPE_MXFP4],   matmul_id_subgroup_mxfp4_f32,   mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
+         CREATE_MM2(GGML_TYPE_NVFP4,   pipeline_dequant_mul_mat_mat_id[GGML_TYPE_NVFP4],   matmul_id_subgroup_nvfp4_f32,   mmq_wg_denoms, warptile_mmq, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id);
+ #undef CREATE_MM2
+@@ -4275,6 +4281,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
+         CREATE_MM2(GGML_TYPE_IQ3_S,   pipeline_dequant_mul_mat_mat[GGML_TYPE_IQ3_S],   matmul_iq3_s_f32,   mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
+         CREATE_MM2(GGML_TYPE_IQ4_XS,  pipeline_dequant_mul_mat_mat[GGML_TYPE_IQ4_XS],  matmul_iq4_xs_f32,  mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
+         CREATE_MM2(GGML_TYPE_IQ4_NL,  pipeline_dequant_mul_mat_mat[GGML_TYPE_IQ4_NL],  matmul_iq4_nl_f32,  mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
++        CREATE_MM2(GGML_TYPE_D2NL,  pipeline_dequant_mul_mat_mat[GGML_TYPE_D2NL],  matmul_d2nl_f32,  mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
+         CREATE_MM2(GGML_TYPE_MXFP4,   pipeline_dequant_mul_mat_mat[GGML_TYPE_MXFP4],   matmul_mxfp4_f32,   mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
+         CREATE_MM2(GGML_TYPE_NVFP4,   pipeline_dequant_mul_mat_mat[GGML_TYPE_NVFP4],   matmul_nvfp4_f32,   mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
+ 
+@@ -4323,6 +4330,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
+             CREATE_MM2(GGML_TYPE_IQ3_S,   pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ3_S],   matmul_id_subgroup_iq3_s_f32,   mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
+             CREATE_MM2(GGML_TYPE_IQ4_XS,  pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ4_XS],  matmul_id_subgroup_iq4_xs_f32,  mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
+             CREATE_MM2(GGML_TYPE_IQ4_NL,  pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ4_NL],  matmul_id_subgroup_iq4_nl_f32,  mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
++            CREATE_MM2(GGML_TYPE_D2NL,  pipeline_dequant_mul_mat_mat_id[GGML_TYPE_D2NL],  matmul_id_subgroup_d2nl_f32,  mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
+             CREATE_MM2(GGML_TYPE_MXFP4,   pipeline_dequant_mul_mat_mat_id[GGML_TYPE_MXFP4],   matmul_id_subgroup_mxfp4_f32,   mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
+             CREATE_MM2(GGML_TYPE_NVFP4,   pipeline_dequant_mul_mat_mat_id[GGML_TYPE_NVFP4],   matmul_id_subgroup_nvfp4_f32,   mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
+ 
+@@ -4370,6 +4378,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
+             CREATE_MM2(GGML_TYPE_IQ3_S,   pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ3_S],   matmul_id_iq3_s_f32,   mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
+             CREATE_MM2(GGML_TYPE_IQ4_XS,  pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ4_XS],  matmul_id_iq4_xs_f32,  mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
+             CREATE_MM2(GGML_TYPE_IQ4_NL,  pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ4_NL],  matmul_id_iq4_nl_f32,  mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
++            CREATE_MM2(GGML_TYPE_D2NL,  pipeline_dequant_mul_mat_mat_id[GGML_TYPE_D2NL],  matmul_id_d2nl_f32,  mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
+             CREATE_MM2(GGML_TYPE_MXFP4,   pipeline_dequant_mul_mat_mat_id[GGML_TYPE_MXFP4],   matmul_id_mxfp4_f32,   mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
+             CREATE_MM2(GGML_TYPE_NVFP4,   pipeline_dequant_mul_mat_mat_id[GGML_TYPE_NVFP4],   matmul_id_nvfp4_f32,   mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
+ 
+@@ -4447,6 +4456,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
+         CREATE_MM(GGML_TYPE_IQ3_S,   pipeline_dequant_mul_mat_mat[GGML_TYPE_IQ3_S].f32acc,   matmul_iq3_s_f32,   , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
+         CREATE_MM(GGML_TYPE_IQ4_XS,  pipeline_dequant_mul_mat_mat[GGML_TYPE_IQ4_XS].f32acc,  matmul_iq4_xs_f32,  , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
+         CREATE_MM(GGML_TYPE_IQ4_NL,  pipeline_dequant_mul_mat_mat[GGML_TYPE_IQ4_NL].f32acc,  matmul_iq4_nl_f32,  , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
++        CREATE_MM(GGML_TYPE_D2NL,  pipeline_dequant_mul_mat_mat[GGML_TYPE_D2NL].f32acc,  matmul_d2nl_f32,  , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
+         CREATE_MM(GGML_TYPE_MXFP4,   pipeline_dequant_mul_mat_mat[GGML_TYPE_MXFP4].f32acc,   matmul_mxfp4_f32,   , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
+         CREATE_MM(GGML_TYPE_NVFP4,   pipeline_dequant_mul_mat_mat[GGML_TYPE_NVFP4].f32acc,   matmul_nvfp4_f32,   , mmq_wg_denoms, warptile_mmq, vk_mat_mat_push_constants, 3, , 0);
+ 
+@@ -4493,6 +4503,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
+             CREATE_MM(GGML_TYPE_IQ3_S,   pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ3_S].f32acc,   matmul_id_subgroup_iq3_s_f32,   , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
+             CREATE_MM(GGML_TYPE_IQ4_XS,  pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ4_XS].f32acc,  matmul_id_subgroup_iq4_xs_f32,  , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
+             CREATE_MM(GGML_TYPE_IQ4_NL,  pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ4_NL].f32acc,  matmul_id_subgroup_iq4_nl_f32,  , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
++            CREATE_MM(GGML_TYPE_D2NL,  pipeline_dequant_mul_mat_mat_id[GGML_TYPE_D2NL].f32acc,  matmul_id_subgroup_d2nl_f32,  , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
+             CREATE_MM(GGML_TYPE_MXFP4,   pipeline_dequant_mul_mat_mat_id[GGML_TYPE_MXFP4].f32acc,   matmul_id_subgroup_mxfp4_f32,   , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
+             CREATE_MM(GGML_TYPE_NVFP4,   pipeline_dequant_mul_mat_mat_id[GGML_TYPE_NVFP4].f32acc,   matmul_id_subgroup_nvfp4_f32,   , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, mul_mat_subgroup_size);
+         } else {
+@@ -4522,6 +4533,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
+             CREATE_MM(GGML_TYPE_IQ3_S,   pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ3_S].f32acc,   matmul_id_iq3_s_f32,   , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
+             CREATE_MM(GGML_TYPE_IQ4_XS,  pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ4_XS].f32acc,  matmul_id_iq4_xs_f32,  , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
+             CREATE_MM(GGML_TYPE_IQ4_NL,  pipeline_dequant_mul_mat_mat_id[GGML_TYPE_IQ4_NL].f32acc,  matmul_id_iq4_nl_f32,  , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
++            CREATE_MM(GGML_TYPE_D2NL,  pipeline_dequant_mul_mat_mat_id[GGML_TYPE_D2NL].f32acc,  matmul_id_d2nl_f32,  , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
+             CREATE_MM(GGML_TYPE_MXFP4,   pipeline_dequant_mul_mat_mat_id[GGML_TYPE_MXFP4].f32acc,   matmul_id_mxfp4_f32,   , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
+             CREATE_MM(GGML_TYPE_NVFP4,   pipeline_dequant_mul_mat_mat_id[GGML_TYPE_NVFP4].f32acc,   matmul_id_nvfp4_f32,   , mmq_wg_denoms, warptile_mmqid, vk_mat_mat_id_push_constants, mul_mat_id_param_count, _id, 0);
+         }
+@@ -4617,6 +4629,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
+             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_IQ3_S][i],   "mul_mat_vec_iq3_s_f32_f32",   arr_dmmv_iq3_s_f32_f32_len[reduc16],   arr_dmmv_iq3_s_f32_f32_data[reduc16],   "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
+             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_IQ4_XS][i],  "mul_mat_vec_iq4_xs_f32_f32",  arr_dmmv_iq4_xs_f32_f32_len[reduc16],  arr_dmmv_iq4_xs_f32_f32_data[reduc16],  "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
+             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_IQ4_NL][i],  "mul_mat_vec_iq4_nl_f32_f32",  arr_dmmv_iq4_nl_f32_f32_len[reduc16],  arr_dmmv_iq4_nl_f32_f32_data[reduc16],  "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
++            ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_D2NL][i],  "mul_mat_vec_d2nl_f32_f32",  arr_dmmv_d2nl_f32_f32_len[reduc16],  arr_dmmv_d2nl_f32_f32_data[reduc16],  "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
+             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_MXFP4][i],   "mul_mat_vec_mxfp4_f32_f32",   arr_dmmv_mxfp4_f32_f32_len[reduc16],   arr_dmmv_mxfp4_f32_f32_data[reduc16],   "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
+             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f32_f32[w][GGML_TYPE_NVFP4][i],   "mul_mat_vec_nvfp4_f32_f32",   arr_dmmv_nvfp4_f32_f32_len[reduc16],   arr_dmmv_nvfp4_f32_f32_data[reduc16],   "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
+ 
+@@ -4644,6 +4657,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
+             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_IQ3_S][i],   "mul_mat_vec_iq3_s_f16_f32",   arr_dmmv_iq3_s_f16_f32_len[reduc16],   arr_dmmv_iq3_s_f16_f32_data[reduc16],   "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
+             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_IQ4_XS][i],  "mul_mat_vec_iq4_xs_f16_f32",  arr_dmmv_iq4_xs_f16_f32_len[reduc16],  arr_dmmv_iq4_xs_f16_f32_data[reduc16],  "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
+             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_IQ4_NL][i],  "mul_mat_vec_iq4_nl_f16_f32",  arr_dmmv_iq4_nl_f16_f32_len[reduc16],  arr_dmmv_iq4_nl_f16_f32_data[reduc16],  "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
++            ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_D2NL][i],  "mul_mat_vec_d2nl_f16_f32",  arr_dmmv_d2nl_f16_f32_len[reduc16],  arr_dmmv_d2nl_f16_f32_data[reduc16],  "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
+             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_MXFP4][i],   "mul_mat_vec_mxfp4_f16_f32",   arr_dmmv_mxfp4_f16_f32_len[reduc16],   arr_dmmv_mxfp4_f16_f32_data[reduc16],   "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
+             ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_f16_f32[w][GGML_TYPE_NVFP4][i],   "mul_mat_vec_nvfp4_f16_f32",   arr_dmmv_nvfp4_f16_f32_len[reduc16],   arr_dmmv_nvfp4_f16_f32_data[reduc16],   "main", mul_mat_vec_num_bindings, sizeof(vk_mat_vec_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq, i+1}, 1, true, use_subgroups16, force_subgroup_size16);
+ 
+@@ -4697,6 +4711,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
+         ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_IQ3_S],   "mul_mat_vec_id_iq3_s_f32",   arr_dmmv_id_iq3_s_f32_f32_len[reduc16],   arr_dmmv_id_iq3_s_f32_f32_data[reduc16],   "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq}, 1, true, use_subgroups16, force_subgroup_size16);
+         ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_IQ4_XS],  "mul_mat_vec_id_iq4_xs_f32",  arr_dmmv_id_iq4_xs_f32_f32_len[reduc16],  arr_dmmv_id_iq4_xs_f32_f32_data[reduc16],  "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq}, 1, true, use_subgroups16, force_subgroup_size16);
+         ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_IQ4_NL],  "mul_mat_vec_id_iq4_nl_f32",  arr_dmmv_id_iq4_nl_f32_f32_len[reduc16],  arr_dmmv_id_iq4_nl_f32_f32_data[reduc16],  "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq}, 1, true, use_subgroups16, force_subgroup_size16);
++        ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_D2NL],  "mul_mat_vec_id_d2nl_f32",  arr_dmmv_id_d2nl_f32_f32_len[reduc16],  arr_dmmv_id_d2nl_f32_f32_data[reduc16],  "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq}, 1, true, use_subgroups16, force_subgroup_size16);
+         ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_MXFP4],   "mul_mat_vec_id_mxfp4_f32",   arr_dmmv_id_mxfp4_f32_f32_len[reduc16],   arr_dmmv_id_mxfp4_f32_f32_data[reduc16],   "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq}, 1, true, use_subgroups16, force_subgroup_size16);
+         ggml_vk_create_pipeline(device, device->pipeline_dequant_mul_mat_vec_id_f32[w][GGML_TYPE_NVFP4],   "mul_mat_vec_id_nvfp4_f32",   arr_dmmv_id_nvfp4_f32_f32_len[reduc16],   arr_dmmv_id_nvfp4_f32_f32_data[reduc16],   "main", mul_mat_vec_id_num_bindings, sizeof(vk_mat_vec_id_push_constants), {rm_iq, 1, 1}, {wg_size_subgroup16, rm_iq}, 1, true, use_subgroups16, force_subgroup_size16);
+ 
+@@ -4754,6 +4769,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
+     ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_IQ3_S],   "dequant_iq3_s",   dequant_iq3_s_len,   dequant_iq3_s_data,   "main", 2, 5 * sizeof(uint32_t), {256 * 32, 1, 1}, {}, 1);
+     ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_IQ4_XS],  "dequant_iq4_xs",  dequant_iq4_xs_len,  dequant_iq4_xs_data,  "main", 2, 5 * sizeof(uint32_t), {256 * 32, 1, 1}, {}, 1);
+     ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_IQ4_NL],  "dequant_iq4_nl",  dequant_iq4_nl_len,  dequant_iq4_nl_data,  "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1);
++    ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_D2NL],  "dequant_d2nl",  dequant_d2nl_len,  dequant_d2nl_data,  "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1);
+     ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_MXFP4],   "dequant_mxfp4",   dequant_mxfp4_len,   dequant_mxfp4_data,   "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1);
+     ggml_vk_create_pipeline(device, device->pipeline_dequant[GGML_TYPE_NVFP4],   "dequant_nvfp4",   dequant_nvfp4_len,   dequant_nvfp4_data,   "main", 2, 5 * sizeof(uint32_t), {256 * 16, 1, 1}, {}, 1);
+ 
+@@ -4782,6 +4798,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
+     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_IQ3_S],   "get_rows_iq3_s",   get_rows_iq3_s_len,   get_rows_iq3_s_data,   "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
+     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_IQ4_XS],  "get_rows_iq4_xs",  get_rows_iq4_xs_len,  get_rows_iq4_xs_data,  "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
+     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_IQ4_NL],  "get_rows_iq4_nl",  get_rows_iq4_nl_len,  get_rows_iq4_nl_data,  "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
++    ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_D2NL],  "get_rows_d2nl",  get_rows_d2nl_len,  get_rows_d2nl_data,  "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
+     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_MXFP4],   "get_rows_mxfp4",   get_rows_mxfp4_len,   get_rows_mxfp4_data,   "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
+     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_NVFP4],   "get_rows_nvfp4",   get_rows_nvfp4_len,   get_rows_nvfp4_data,   "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
+     ggml_vk_create_pipeline(device, device->pipeline_get_rows[GGML_TYPE_I32],     "get_rows_i32",     get_rows_i32_len,     get_rows_i32_data,     "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
+@@ -4810,6 +4827,7 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
+     ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_IQ3_S],   "get_rows_iq3_s_f32",   get_rows_iq3_s_f32_len,   get_rows_iq3_s_f32_data,   "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
+     ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_IQ4_XS],  "get_rows_iq4_xs_f32",  get_rows_iq4_xs_f32_len,  get_rows_iq4_xs_f32_data,  "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
+     ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_IQ4_NL],  "get_rows_iq4_nl_f32",  get_rows_iq4_nl_f32_len,  get_rows_iq4_nl_f32_data,  "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
++    ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_D2NL],  "get_rows_d2nl_f32",  get_rows_d2nl_f32_len,  get_rows_d2nl_f32_data,  "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
+     ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_MXFP4],   "get_rows_mxfp4_f32",   get_rows_mxfp4_f32_len,   get_rows_mxfp4_f32_data,   "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
+     ggml_vk_create_pipeline(device, device->pipeline_get_rows_f32[GGML_TYPE_NVFP4],   "get_rows_nvfp4_f32",   get_rows_nvfp4_f32_len,   get_rows_nvfp4_f32_data,   "main", 3, sizeof(vk_op_binary_push_constants), {1024, 1, 1}, {}, 1);
+ 
+@@ -6769,6 +6787,7 @@ static vk_pipeline ggml_vk_get_to_fp16(ggml_backend_vk_context * ctx, ggml_type
+         case GGML_TYPE_IQ3_XXS:
+         case GGML_TYPE_IQ3_S:
+         case GGML_TYPE_IQ4_XS:
++        case GGML_TYPE_D2NL:
+         case GGML_TYPE_IQ4_NL:
+         case GGML_TYPE_MXFP4:
+         case GGML_TYPE_NVFP4:
+@@ -6843,6 +6862,7 @@ static vk_matmul_pipeline ggml_vk_get_mul_mat_mat_pipeline(ggml_backend_vk_conte
+         case GGML_TYPE_IQ3_XXS:
+         case GGML_TYPE_IQ3_S:
+         case GGML_TYPE_IQ4_XS:
++        case GGML_TYPE_D2NL:
+         case GGML_TYPE_IQ4_NL:
+         case GGML_TYPE_MXFP4:
+         case GGML_TYPE_NVFP4:
+@@ -6911,6 +6931,7 @@ static vk_pipeline ggml_vk_get_dequantize_mul_mat_vec(ggml_backend_vk_context *
+         case GGML_TYPE_IQ3_XXS:
+         case GGML_TYPE_IQ3_S:
+         case GGML_TYPE_IQ4_XS:
++        case GGML_TYPE_D2NL:
+         case GGML_TYPE_IQ4_NL:
+         case GGML_TYPE_MXFP4:
+         case GGML_TYPE_NVFP4:
+@@ -7004,6 +7025,7 @@ static vk_matmul_pipeline ggml_vk_get_mul_mat_mat_id_pipeline(ggml_backend_vk_co
+         case GGML_TYPE_IQ3_XXS:
+         case GGML_TYPE_IQ3_S:
+         case GGML_TYPE_IQ4_XS:
++        case GGML_TYPE_D2NL:
+         case GGML_TYPE_IQ4_NL:
+         case GGML_TYPE_MXFP4:
+         case GGML_TYPE_NVFP4:
+@@ -7075,6 +7097,7 @@ static vk_pipeline ggml_vk_get_dequantize_mul_mat_vec_id(ggml_backend_vk_context
+         case GGML_TYPE_IQ3_XXS:
+         case GGML_TYPE_IQ3_S:
+         case GGML_TYPE_IQ4_XS:
++        case GGML_TYPE_D2NL:
+         case GGML_TYPE_IQ4_NL:
+         case GGML_TYPE_MXFP4:
+         case GGML_TYPE_NVFP4:
+@@ -16528,6 +16551,7 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
+                     case GGML_TYPE_IQ3_XXS:
+                     case GGML_TYPE_IQ3_S:
+                     case GGML_TYPE_IQ4_XS:
++                    case GGML_TYPE_D2NL:
+                     case GGML_TYPE_IQ4_NL:
+                     case GGML_TYPE_MXFP4:
+                     case GGML_TYPE_NVFP4:
+@@ -16634,6 +16658,7 @@ static bool ggml_backend_vk_device_supports_op(ggml_backend_dev_t dev, const ggm
+                     case GGML_TYPE_IQ3_XXS:
+                     case GGML_TYPE_IQ3_S:
+                     case GGML_TYPE_IQ4_XS:
++                    case GGML_TYPE_D2NL:
+                     case GGML_TYPE_IQ4_NL:
+                     case GGML_TYPE_MXFP4:
+                     case GGML_TYPE_NVFP4:
+--- a/ggml/src/ggml-vulkan/vulkan-shaders/types.glsl
++++ b/ggml/src/ggml-vulkan/vulkan-shaders/types.glsl
+@@ -1757,8 +1757,14 @@ struct block_nvfp4_packed32
+ 
+ #if defined(DATA_A_IQ4_NL) || defined(DATA_A_IQ4_XS)
+ const int8_t kvalues_iq4nl_const[16] = {
++#if defined(D2NL_TABLE)
++    // Entrit d2nl: d2 levels {0, +-1, +-2.828, +-5.196, +-8} x 10.981; index 4 is zero (kvalues_d2nl in ggml-common.h)
++    int8_t(-88), int8_t(-57), int8_t(-31), int8_t(-11), int8_t(0), int8_t(11), int8_t(31), int8_t(57),
++    int8_t(88), int8_t(0), int8_t(0), int8_t(0), int8_t(0), int8_t(0), int8_t(0), int8_t(0)
++#else
+     int8_t(-127), int8_t(-104), int8_t(-83), int8_t(-65), int8_t(-49), int8_t(-35), int8_t(-22), int8_t(-10),
+     int8_t(1), int8_t(13), int8_t(25), int8_t(38), int8_t(53), int8_t(69), int8_t(89), int8_t(113)
++#endif
+ };
+ 
+ shared FLOAT_TYPE kvalues_iq4nl[16];
+--- a/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp
++++ b/ggml/src/ggml-vulkan/vulkan-shaders/vulkan-shaders-gen.cpp
+@@ -66,6 +66,7 @@ const std::vector<std::string> type_names = {
+     "iq3_s",
+     "iq4_xs",
+     "iq4_nl",
++    "d2nl",
+     "mxfp4",
+     "nvfp4",
+     "bf16",
+@@ -408,7 +409,18 @@ std::map<std::string, std::string> merge_maps(const std::map<std::string, std::s
+ }
+ 
+ static std::vector<std::future<void>> compiles;
+-void string_to_spv(std::string name, const std::string& source, const std::map<std::string, std::string>& defines, bool fp16 = true, bool coopmat = false, bool coopmat2 = false, bool f16acc = false, const std::string& suffix = "") {
++void string_to_spv(std::string name, const std::string& source_in, const std::map<std::string, std::string>& defines_in, bool fp16 = true, bool coopmat = false, bool coopmat2 = false, bool f16acc = false, const std::string& suffix = "") {
++    // Entrit d2nl shaders are the IQ4_NL shaders with the d2 table (types.glsl, D2NL_TABLE).
++    std::string source = source_in;
++    std::map<std::string, std::string> defines = defines_in;
++    if (defines.erase("DATA_A_D2NL")) {
++        defines["DATA_A_IQ4_NL"] = "1";
++        defines["D2NL_TABLE"] = "1";
++    }
++    const size_t d2nl_pos = source.find("d2nl");
++    if (d2nl_pos != std::string::npos) {
++        source.replace(d2nl_pos, 4, "iq4_nl");
++    }
+     name = name + (f16acc ? "_f16acc" : "") + (coopmat ? "_cm1" : "") + (coopmat2 ? "_cm2" : (fp16 ? "" : "_fp32")) + suffix;
+     std::string out_path = join_paths(output_dir, name + ".spv");
+ 
+@@ -558,7 +570,7 @@ void matmul_shaders(bool fp16, MatMulIdType matmul_id_type, bool coopmat, bool c
+         std::string load_vec_quant = "2";
+         if ((tname == "q1_0") || (tname == "q4_0") || (tname == "q4_1") || (tname == "q5_1") || (tname == "iq1_s") || (tname == "iq1_m") || (tname == "iq2_xxs") || (tname == "iq2_xs") || (tname == "iq2_s"))
+             load_vec_quant = "8";
+-        else if ((tname == "q2_0") || (tname == "q5_0") || (tname == "q8_0") || (tname == "q2_k") || (tname == "q4_k") || (tname == "q5_k") || (tname == "iq3_xxs") || (tname == "iq3_s") || (tname == "iq4_xs") || (tname == "iq4_nl") || (tname == "mxfp4") || (tname == "nvfp4"))
++        else if ((tname == "q2_0") || (tname == "q5_0") || (tname == "q8_0") || (tname == "q2_k") || (tname == "q4_k") || (tname == "q5_k") || (tname == "iq3_xxs") || (tname == "iq3_s") || (tname == "iq4_xs") || (tname == "iq4_nl") || (tname == "d2nl") || (tname == "mxfp4") || (tname == "nvfp4"))
+             load_vec_quant = "4";
+ 
+         if (tname == "bf16") {
+--- a/ggml/src/ggml.c
++++ b/ggml/src/ggml.c
+@@ -861,6 +861,20 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = {
+         .to_float                 = (ggml_to_float_t) dequantize_row_iq4_nl,
+         .from_float_ref           = (ggml_from_float_t)quantize_row_iq4_nl_ref,
+     },
++    [43] = { // reserved
++        .type_name                = "TYPE_43 RESERVED",
++        .blck_size                = 0,
++        .type_size                = 0,
++        .is_quantized             = false,
++    },
++    [GGML_TYPE_D2NL] = {
++        .type_name                = "d2nl",
++        .blck_size                = QK_D2NL,
++        .type_size                = sizeof(block_d2nl),
++        .is_quantized             = true,
++        .to_float                 = (ggml_to_float_t) dequantize_row_d2nl,
++        .from_float_ref           = (ggml_from_float_t) quantize_row_d2nl_ref,
++    },
+     [GGML_TYPE_IQ4_XS] = {
+         .type_name                = "iq4_xs",
+         .blck_size                = QK_K,
+@@ -1433,6 +1447,7 @@ enum ggml_type ggml_ftype_to_ggml_type(enum ggml_ftype ftype) {
+         case GGML_FTYPE_MOSTLY_IQ1_S:         wtype = GGML_TYPE_IQ1_S;    break;
+         case GGML_FTYPE_MOSTLY_IQ1_M:         wtype = GGML_TYPE_IQ1_M;    break;
+         case GGML_FTYPE_MOSTLY_IQ4_NL:        wtype = GGML_TYPE_IQ4_NL;   break;
++        case GGML_FTYPE_MOSTLY_D2NL:          wtype = GGML_TYPE_D2NL;     break;
+         case GGML_FTYPE_MOSTLY_IQ4_XS:        wtype = GGML_TYPE_IQ4_XS;   break;
+         case GGML_FTYPE_MOSTLY_IQ3_S:         wtype = GGML_TYPE_IQ3_S;    break;
+         case GGML_FTYPE_MOSTLY_IQ2_S:         wtype = GGML_TYPE_IQ2_S;    break;
+@@ -7786,6 +7801,7 @@ size_t ggml_quantize_chunk(
+         case GGML_TYPE_IQ1_S:   result = quantize_iq1_s  (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+         case GGML_TYPE_IQ1_M:   result = quantize_iq1_m  (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+         case GGML_TYPE_IQ4_NL:  result = quantize_iq4_nl (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
++        case GGML_TYPE_D2NL:    result = quantize_d2nl   (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+         case GGML_TYPE_IQ4_XS:  result = quantize_iq4_xs (src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break;
+         case GGML_TYPE_F16:
+             {
+--- a/gguf-py/gguf/constants.py
++++ b/gguf-py/gguf/constants.py
+@@ -4406,6 +4406,7 @@ class GGMLQuantizationType(IntEnum):
+     NVFP4   = 40
+     Q1_0    = 41
+     Q2_0    = 42
++    D2NL    = 44
+ 
+ 
+ class ExpertGatingFuncType(IntEnum):
+@@ -4461,6 +4462,7 @@ class LlamaFileType(IntEnum):
+     MOSTLY_NVFP4         = 39  # except 1d tensors
+     MOSTLY_Q1_0          = 40  # except 1d tensors
+     MOSTLY_Q2_0          = 41  # except 1d tensors
++    MOSTLY_D2NL          = 43  # except 1d tensors
+ 
+     GUESSED              = 1024  # not specified in the model file
+ 
+@@ -4587,6 +4589,7 @@ GGML_QUANT_SIZES: dict[GGMLQuantizationType, tuple[int, int]] = {
+     GGMLQuantizationType.NVFP4:   (64, 4 + 32),
+     GGMLQuantizationType.Q1_0:    (128, 2 + 16),
+     GGMLQuantizationType.Q2_0:    (128, 2 + 32),
++    GGMLQuantizationType.D2NL:    (32,  2 + 16),
+ }
+ 
+ 
+--- a/include/llama.h
++++ b/include/llama.h
+@@ -156,6 +156,7 @@ extern "C" {
+         LLAMA_FTYPE_MOSTLY_NVFP4         = 39, // except 1d tensors
+         LLAMA_FTYPE_MOSTLY_Q1_0          = 40, // except 1d tensors
+         LLAMA_FTYPE_MOSTLY_Q2_0          = 41, // except 1d tensors
++        LLAMA_FTYPE_MOSTLY_D2NL          = 43, // except 1d tensors
+ 
+         LLAMA_FTYPE_GUESSED = 1024, // not specified in the model file
+     };
+--- a/src/llama-model-loader.cpp
++++ b/src/llama-model-loader.cpp
+@@ -66,6 +66,7 @@ static std::string llama_model_ftype_name(llama_ftype ftype) {
+         case LLAMA_FTYPE_MOSTLY_IQ1_S:    return "IQ1_S - 1.5625 bpw";
+         case LLAMA_FTYPE_MOSTLY_IQ1_M:    return "IQ1_M - 1.75 bpw";
+         case LLAMA_FTYPE_MOSTLY_IQ4_NL:   return "IQ4_NL - 4.5 bpw";
++        case LLAMA_FTYPE_MOSTLY_D2NL:     return "D2NL - 4.5 bpw (Entrit d2)";
+         case LLAMA_FTYPE_MOSTLY_IQ4_XS:   return "IQ4_XS - 4.25 bpw";
+         case LLAMA_FTYPE_MOSTLY_IQ3_S:    return "IQ3_S - 3.4375 bpw";
+         case LLAMA_FTYPE_MOSTLY_IQ3_M:    return "IQ3_S mix - 3.66 bpw";
+@@ -758,6 +759,7 @@ llama_model_loader::llama_model_loader(
+             case GGML_TYPE_IQ1_S:   ftype = LLAMA_FTYPE_MOSTLY_IQ1_S;   break;
+             case GGML_TYPE_IQ1_M:   ftype = LLAMA_FTYPE_MOSTLY_IQ1_M;   break;
+             case GGML_TYPE_IQ4_NL:  ftype = LLAMA_FTYPE_MOSTLY_IQ4_NL;  break;
++            case GGML_TYPE_D2NL:    ftype = LLAMA_FTYPE_MOSTLY_D2NL;    break;
+             case GGML_TYPE_IQ4_XS:  ftype = LLAMA_FTYPE_MOSTLY_IQ4_XS;  break;
+             case GGML_TYPE_IQ3_S:   ftype = LLAMA_FTYPE_MOSTLY_IQ3_S;   break;
+             case GGML_TYPE_NVFP4:   ftype = LLAMA_FTYPE_MOSTLY_NVFP4;   break;
+--- a/tests/test-backend-ops.cpp
++++ b/tests/test-backend-ops.cpp
+@@ -7596,7 +7596,7 @@ static const ggml_type all_types[] = {
+     // GGML_TYPE_TQ1_0, GGML_TYPE_TQ2_0, // TODO: implement for all backends
+     GGML_TYPE_IQ2_XXS, GGML_TYPE_IQ2_XS, GGML_TYPE_IQ2_S,
+     GGML_TYPE_IQ3_XXS, GGML_TYPE_IQ1_S, GGML_TYPE_IQ1_M,
+-    GGML_TYPE_IQ4_NL, GGML_TYPE_IQ3_S, GGML_TYPE_IQ4_XS,
++    GGML_TYPE_IQ4_NL, GGML_TYPE_IQ3_S, GGML_TYPE_IQ4_XS, GGML_TYPE_D2NL,
+ };
+ 
+ static const ggml_type base_types[] = {
+@@ -7623,7 +7623,7 @@ static const ggml_type other_types[] = {
+     // GGML_TYPE_TQ1_0, GGML_TYPE_TQ2_0, // TODO: implement for all backends
+     GGML_TYPE_IQ2_XS, GGML_TYPE_IQ2_S,
+     GGML_TYPE_IQ3_XXS, GGML_TYPE_IQ1_S, GGML_TYPE_IQ1_M,
+-    GGML_TYPE_IQ4_NL, GGML_TYPE_IQ3_S, GGML_TYPE_IQ4_XS,
++    GGML_TYPE_IQ4_NL, GGML_TYPE_IQ3_S, GGML_TYPE_IQ4_XS, GGML_TYPE_D2NL,
+     GGML_TYPE_BF16,
+ };
+ 

run.sh

@@ -0,0 +1,22 @@
+#!/usr/bin/env bash
+# Start a chat server for a model, then open http://localhost:8080 in your browser.
+#
+#   ./run.sh models/<file>.gguf [extra llama-server options]
+#
+# Stop it with Ctrl-C.
+set -euo pipefail
+
+HERE=$(cd "$(dirname "$0")" && pwd)
+MODEL=${1:?usage: ./run.sh models/<file>.gguf}
+shift
+SERVER="$HERE/build/bin/llama-server"
+[[ -x "$SERVER" ]] || { echo "llama.cpp is not built yet. Run ./install.sh first."; exit 1; }
+[[ -f "$MODEL" ]]  || { echo "No model at $MODEL. Download one with ./get-model.sh first."; exit 1; }
+
+OPTS=(-m "$MODEL" -ngl 99 -c 8192 --port 8080)
+# On Linux, read the whole model into memory instead of paging it from disk. A model paged
+# from a slow disk can make every token wait on the drive. Macs page from fast SSD, so skip it.
+[[ "$(uname -s)" != Darwin ]] && OPTS+=(--no-mmap)
+
+echo "Starting. When you see 'server is listening', open http://localhost:8080"
+exec "$SERVER" "${OPTS[@]}" "$@"