diff --git a/.github/workflows/release-prism.yml b/.github/workflows/release-prism.yml new file mode 100644 index 000000000000..43fa9e72330c --- /dev/null +++ b/.github/workflows/release-prism.yml @@ -0,0 +1,255 @@ +name: Release (Prism) + +on: + workflow_dispatch: + inputs: + create_release: + description: 'Create new release' + required: true + type: boolean + +concurrency: + group: ${{ github.workflow }}-${{ github.head_ref && github.ref || github.run_id }} + cancel-in-progress: true + +env: + BRANCH_NAME: ${{ github.head_ref || github.ref_name }} + CMAKE_ARGS: "-DLLAMA_BUILD_EXAMPLES=OFF -DLLAMA_BUILD_TESTS=OFF -DLLAMA_BUILD_TOOLS=ON -DLLAMA_BUILD_SERVER=ON -DGGML_RPC=ON" + +jobs: + macOS-arm64: + runs-on: macos-14 + + steps: + - name: Clone + uses: actions/checkout@v6 + with: + fetch-depth: 0 + + - name: ccache + uses: ggml-org/ccache-action@v1.2.16 + with: + key: macOS-latest-cmake-arm64 + evict-old-files: 1d + + - name: Build + run: | + cmake -B build \ + -DCMAKE_INSTALL_RPATH='@loader_path' \ + -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \ + -DLLAMA_FATAL_WARNINGS=ON \ + -DGGML_METAL_USE_BF16=ON \ + -DGGML_METAL_EMBED_LIBRARY=ON \ + -DGGML_RPC=ON \ + ${{ env.CMAKE_ARGS }} + cmake --build build --config Release -j $(sysctl -n hw.logicalcpu) + + - name: Determine tag name + id: tag + uses: ./.github/actions/get-tag-name + + - name: Pack artifacts + run: | + cp LICENSE ./build/bin/ + tar -czvf llama-${{ steps.tag.outputs.name }}-bin-macos-arm64.tar.gz -s ",./,llama-${{ steps.tag.outputs.name }}/," -C ./build/bin . + + - name: Upload artifacts + uses: actions/upload-artifact@v6 + with: + path: llama-${{ steps.tag.outputs.name }}-bin-macos-arm64.tar.gz + name: llama-bin-macos-arm64.tar.gz + + linux-cuda: + runs-on: ubuntu-22.04 + + strategy: + matrix: + include: + - cuda: '12.4' + cuda_pkg: '12-4' + - cuda: '12.8' + cuda_pkg: '12-8' + - cuda: '13.1' + cuda_pkg: '13-1' + + steps: + - name: Clone + uses: actions/checkout@v6 + with: + fetch-depth: 0 + + - name: ccache + uses: ggml-org/ccache-action@v1.2.16 + with: + key: ubuntu-22-cmake-cuda-${{ matrix.cuda }} + evict-old-files: 1d + + - name: Install CUDA toolkit + run: | + wget -q https://developer.download.nvidia.com/compute/cuda/repos/ubuntu2204/x86_64/cuda-keyring_1.1-1_all.deb + sudo dpkg -i cuda-keyring_1.1-1_all.deb + sudo apt-get update + sudo apt-get -y install cuda-toolkit-${{ matrix.cuda_pkg }} + echo "/usr/local/cuda-${{ matrix.cuda }}/bin" >> $GITHUB_PATH + echo "CUDA_PATH=/usr/local/cuda-${{ matrix.cuda }}" >> $GITHUB_ENV + echo "LD_LIBRARY_PATH=/usr/local/cuda-${{ matrix.cuda }}/lib64:$LD_LIBRARY_PATH" >> $GITHUB_ENV + + - name: Build + run: | + cmake -B build \ + -DCMAKE_INSTALL_RPATH='$ORIGIN' \ + -DCMAKE_BUILD_WITH_INSTALL_RPATH=ON \ + -DGGML_NATIVE=OFF \ + -DGGML_CUDA=ON \ + ${{ env.CMAKE_ARGS }} + cmake --build build --config Release -j $(nproc) 2>&1 | grep -v "^nvcc warning" + + - name: Determine tag name + id: tag + uses: ./.github/actions/get-tag-name + + - name: Pack artifacts + run: | + cp LICENSE ./build/bin/ + tar -czvf llama-${{ steps.tag.outputs.name }}-bin-linux-cuda-${{ matrix.cuda }}-x64.tar.gz --transform "s,./,llama-${{ steps.tag.outputs.name }}/," -C ./build/bin . + + - name: Upload artifacts + uses: actions/upload-artifact@v6 + with: + path: llama-${{ steps.tag.outputs.name }}-bin-linux-cuda-${{ matrix.cuda }}-x64.tar.gz + name: llama-bin-linux-cuda-${{ matrix.cuda }}-x64.tar.gz + + windows-cuda: + runs-on: windows-2022 + + strategy: + matrix: + cuda: ['12.4', '13.1'] + + steps: + - name: Clone + uses: actions/checkout@v6 + + - name: Install ccache + uses: ggml-org/ccache-action@v1.2.16 + with: + key: windows-cuda-${{ matrix.cuda }} + variant: ccache + evict-old-files: 1d + + - name: Install Cuda Toolkit + uses: ./.github/actions/windows-setup-cuda + with: + cuda_version: ${{ matrix.cuda }} + + - name: Install Ninja + run: choco install ninja + + - name: Build + shell: cmd + run: | + call "C:\Program Files\Microsoft Visual Studio\2022\Enterprise\VC\Auxiliary\Build\vcvarsall.bat" x64 + cmake -S . -B build -G "Ninja Multi-Config" ^ + -DGGML_NATIVE=OFF ^ + -DGGML_CUDA=ON ^ + -DLLAMA_BUILD_BORINGSSL=ON ^ + -DCMAKE_CUDA_FLAGS="-diag-suppress=221" ^ + ${{ env.CMAKE_ARGS }} + set /A NINJA_JOBS=%NUMBER_OF_PROCESSORS%-1 + cmake --build build --config Release -j %NINJA_JOBS% + + - name: Determine tag name + id: tag + uses: ./.github/actions/get-tag-name + + - name: Pack artifacts + run: | + 7z a -snl llama-${{ steps.tag.outputs.name }}-bin-win-cuda-${{ matrix.cuda }}-x64.zip .\build\bin\Release\* + + - name: Upload artifacts + uses: actions/upload-artifact@v6 + with: + path: llama-${{ steps.tag.outputs.name }}-bin-win-cuda-${{ matrix.cuda }}-x64.zip + name: llama-bin-win-cuda-${{ matrix.cuda }}-x64.zip + + - name: Copy and pack Cuda runtime + run: | + echo "Cuda install location: ${{ env.CUDA_PATH }}" + $dst='.\build\bin\cudart\' + robocopy "${{env.CUDA_PATH}}\bin" $dst cudart64_*.dll cublas64_*.dll cublasLt64_*.dll + robocopy "${{env.CUDA_PATH}}\lib" $dst cudart64_*.dll cublas64_*.dll cublasLt64_*.dll + robocopy "${{env.CUDA_PATH}}\bin\x64" $dst cudart64_*.dll cublas64_*.dll cublasLt64_*.dll + 7z a cudart-llama-bin-win-cuda-${{ matrix.cuda }}-x64.zip $dst\* + + - name: Upload Cuda runtime + uses: actions/upload-artifact@v6 + with: + path: cudart-llama-bin-win-cuda-${{ matrix.cuda }}-x64.zip + name: cudart-llama-bin-win-cuda-${{ matrix.cuda }}-x64.zip + + release: + if: ${{ github.event.inputs.create_release == 'true' }} + + permissions: + contents: write + + runs-on: ubuntu-latest + + needs: + - macOS-arm64 + - linux-cuda + - windows-cuda + + steps: + - name: Clone + uses: actions/checkout@v6 + with: + fetch-depth: 0 + + - name: Determine tag name + id: tag + uses: ./.github/actions/get-tag-name + + - name: Download artifacts + uses: actions/download-artifact@v7 + with: + path: ./artifact + merge-multiple: true + + - name: Move artifacts + run: | + mkdir -p release + mv -v artifact/*.tar.gz release/ 2>/dev/null || true + mv -v artifact/*.zip release/ 2>/dev/null || true + ls -lh release/ + + - name: Create release + id: create_release + uses: ggml-org/action-create-release@v1 + env: + GITHUB_TOKEN: ${{ secrets.GITHUB_TOKEN }} + with: + tag_name: ${{ steps.tag.outputs.name }} + body: | + Pre-built binaries (PrismML fork with Q1_0 1-bit quantization support). + + **macOS:** + - [macOS Apple Silicon (arm64)](https://github.com/${{ github.repository }}/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-macos-arm64.tar.gz) + + **Linux:** + - [Linux x64 (CUDA 12.4)](https://github.com/${{ github.repository }}/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-linux-cuda-12.4-x64.tar.gz) + - [Linux x64 (CUDA 12.8)](https://github.com/${{ github.repository }}/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-linux-cuda-12.8-x64.tar.gz) + - [Linux x64 (CUDA 13.1)](https://github.com/${{ github.repository }}/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-linux-cuda-13.1-x64.tar.gz) + + **Windows:** + - [Windows x64 (CUDA 12.4)](https://github.com/${{ github.repository }}/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-win-cuda-12.4-x64.zip) - [CUDA 12.4 DLLs](https://github.com/${{ github.repository }}/releases/download/${{ steps.tag.outputs.name }}/cudart-llama-bin-win-cuda-12.4-x64.zip) + - [Windows x64 (CUDA 13.1)](https://github.com/${{ github.repository }}/releases/download/${{ steps.tag.outputs.name }}/llama-${{ steps.tag.outputs.name }}-bin-win-cuda-13.1-x64.zip) - [CUDA 13.1 DLLs](https://github.com/${{ github.repository }}/releases/download/${{ steps.tag.outputs.name }}/cudart-llama-bin-win-cuda-13.1-x64.zip) + + - name: Upload release + env: + GH_TOKEN: ${{ secrets.GITHUB_TOKEN }} + run: | + for file in release/*; do + echo "Uploading $(basename $file)..." + gh release upload ${{ steps.tag.outputs.name }} "$file" --clobber + done diff --git a/README.md b/README.md index be23abcea67f..28faf73c2542 100644 --- a/README.md +++ b/README.md @@ -1,3 +1,49 @@ +> [!IMPORTANT] +> This is a fork of the [Prism-ML fork of llama.cpp](https://github.com/PrismML-Eng/llama.cpp) that is synced to the main llama.cpp repo +> It is not yet ready for production use and should be considered experimental +> +> The primary benefit of this fork is an up-to-date version of llama.cpp that also merges the capabiilities of the Prism-ML fork +> to support the [Bonsai 1-bit models](https://huggingface.co/prism-ml/Bonsai-8B-gguf) +> +> Note: this is not an official fork and is not supported by the Prism-ML team - this is just a personal fork to demo Bonsai until official support is added + + +## How to use this fork + +```bash +# On MacOS +git clone https://github.com/Mintplex-Labs/prism-ml-llama.cpp +cd prism-ml-llama.cpp +cmake -B build && cmake --build build -j +``` + +### You __must__ recode the public Bonsai 1-bit models to work with this fork +```bash +# Download the public Bonsai 1-bit models +wget https://huggingface.co/prism-ml/Bonsai-8B-gguf/resolve/main/Bonsai-8B.gguf -O Bonsai-8B.gguf +``` + +### Run the model +```bash +# Llama cli +./build/bin/llama-cli \ + -m Bonsai-8B.gguf \ + -p "Explain quantum computing in simple terms." \ + -n 256 \ + --temp 0.5 \ + --top-p 0.85 \ + --top-k 20 \ + -ngl 99 + +# llama server +./build/bin/llama-server \ + -m Bonsai-8B.gguf \ + --host 0.0.0.0 \ + --port 8080 \ + -ngl 99 + --ctx-size 65536 +``` + # llama.cpp ![llama](https://user-images.githubusercontent.com/1991296/230134379-7181e485-c521-4d23-a0d6-f7b3b61ba524.png) diff --git a/bench_attn_rot.py b/bench_attn_rot.py new file mode 100755 index 000000000000..a00b6a6c5aff --- /dev/null +++ b/bench_attn_rot.py @@ -0,0 +1,523 @@ +#!/usr/bin/env python3 +""" +Benchmark script for testing KV cache rotation (ATTN_ROT) with different context sizes. +Compares performance and quality with rotation enabled vs disabled. +""" + +import subprocess +import time +import os +import pty +import select +import argparse +import json +import re +from dataclasses import dataclass +from typing import Optional + +@dataclass +class BenchResult: + context_size: int + kv_type: str + attn_rot_enabled: bool + prompt_tokens: int + generation_tokens: int + prompt_time_ms: float + generation_time_ms: float + prompt_tps: float + tokens_per_sec: float + output_text: str + success: bool + context_memory_mib: float = 0.0 # KV cache memory + total_memory_mib: float = 0.0 # Total GPU memory used + error: Optional[str] = None + +def run_benchmark( + model_path: str, + context_size: int, + kv_type: str, + attn_rot_enabled: bool, + prompt: str, + n_predict: int = 128, + n_gpu_layers: int = 99, + llama_cli_path: str = "./build/bin/llama-cli", + temp: float = 0.5, + top_p: float = 0.85, + top_k: int = 20, + debug: bool = False, +) -> BenchResult: + """Run a single benchmark with specified parameters.""" + + env = os.environ.copy() + if not attn_rot_enabled: + env["LLAMA_ATTN_ROT_DISABLE"] = "1" + elif "LLAMA_ATTN_ROT_DISABLE" in env: + del env["LLAMA_ATTN_ROT_DISABLE"] + + cmd = [ + llama_cli_path, + "-m", model_path, + "-c", str(context_size), + "-ctk", kv_type, + "-ctv", kv_type, + "-p", prompt, + "-n", str(n_predict), + "-ngl", str(n_gpu_layers), + "--temp", str(temp), + "--top-p", str(top_p), + "--top-k", str(top_k), + "--perf", # enable performance timing output + ] + + print(f"\n{'='*60}") + print(f"Context: {context_size}, KV: {kv_type}, ATTN_ROT: {'ON' if attn_rot_enabled else 'OFF'}") + print(f"{'='*60}") + + try: + start = time.time() + + # Use pty to capture all terminal output (including TTY-only output) + output_chunks = [] + + def read_output(fd): + """Read from file descriptor with timeout.""" + output = b"" + while True: + ready, _, _ = select.select([fd], [], [], 0.1) + if ready: + try: + chunk = os.read(fd, 4096) + if not chunk: + break + output += chunk + except OSError: + break + else: + # Check if process is still running + break + return output + + master_fd, slave_fd = pty.openpty() + + process = subprocess.Popen( + cmd, + env=env, + stdin=subprocess.DEVNULL, + stdout=slave_fd, + stderr=slave_fd, + close_fds=True, + ) + + os.close(slave_fd) + + # Read output until process completes + output = b"" + while True: + ready, _, _ = select.select([master_fd], [], [], 1.0) + if ready: + try: + chunk = os.read(master_fd, 4096) + if chunk: + output += chunk + else: + break + except OSError: + break + + # Check if process has finished + if process.poll() is not None: + # Read any remaining output + while True: + ready, _, _ = select.select([master_fd], [], [], 0.1) + if ready: + try: + chunk = os.read(master_fd, 4096) + if chunk: + output += chunk + else: + break + except OSError: + break + else: + break + break + + # Timeout check + if time.time() - start > 300: + process.kill() + raise subprocess.TimeoutExpired(cmd, 300) + + os.close(master_fd) + process.wait() + + output = output.decode('utf-8', errors='replace') + elapsed = time.time() - start + + # Create a mock result object for compatibility + class Result: + def __init__(self, returncode, stdout): + self.returncode = returncode + self.stdout = stdout + self.stderr = "" + + result = Result(process.returncode, output) + + if debug: + print(f"\n--- DEBUG: Captured output ({len(output)} chars) ---") + print(output[-2000:] if len(output) > 2000 else output) + print("--- END DEBUG ---\n") + + # Parse timing from llama.cpp output + prompt_tps = 0.0 + gen_tps = 0.0 + prompt_time = 0.0 + gen_time = 0.0 + prompt_tokens = 0 + gen_tokens = 0 + + # Search entire output for the timing line (handles multiline) + # New format: [ Prompt: 414.1 t/s | Generation: 115.4 t/s ] + prompt_match = re.search(r'Prompt:\s*([\d.]+)\s*t/s', output) + gen_match = re.search(r'Generation:\s*([\d.]+)\s*t/s', output) + + if prompt_match: + prompt_tps = float(prompt_match.group(1)) + if gen_match: + gen_tps = float(gen_match.group(1)) + + # Fallback: Old format parsing + if prompt_tps == 0 or gen_tps == 0: + for line in output.split('\n'): + # Old format: llama_print_timings: prompt eval time = 123.45 ms / 100 tokens + if 'prompt eval time' in line.lower(): + try: + parts = line.split('=')[1].split('/') + prompt_time = float(parts[0].strip().replace('ms', '').strip()) + prompt_tokens = int(parts[1].strip().split()[0]) + except: + pass + elif 'eval time' in line.lower() and 'prompt' not in line.lower(): + try: + parts = line.split('=')[1].split('/') + gen_time = float(parts[0].strip().replace('ms', '').strip()) + gen_tokens = int(parts[1].strip().split()[0]) + except: + pass + + # Use new format values if available, otherwise calculate from old format + if gen_tps == 0 and gen_time > 0 and gen_tokens > 0: + gen_tps = gen_tokens / gen_time * 1000 + if prompt_tps == 0 and prompt_time > 0 and prompt_tokens > 0: + prompt_tps = prompt_tokens / prompt_time * 1000 + + tps = gen_tps + + # Parse memory breakdown + # Format: | - MTL0 (Apple M4 Max) | 36864 = 32867 + (3995 = 1099 + 2592 + 304) + 0 | + # Values: total = free + (used = self + model + context + compute) + unaccounted + context_memory = 0.0 + total_memory = 0.0 + + mem_match = re.search( + r'llama_memory_breakdown_print:.*?\|\s*(\d+)\s*=\s*(\d+)\s*\+\s*\((\d+)\s*=\s*(\d+)\s*\+\s*(\d+)\s*\+\s*(\d+)', + output + ) + if mem_match: + total = float(mem_match.group(1)) + free = float(mem_match.group(2)) + # used = self + model + context + compute + # self = group(4), model = group(5), context = group(6) + context_memory = float(mem_match.group(6)) + total_memory = total - free + + # Extract generated text (everything before timing output) + gen_text = result.stdout.split('\nllama_print_timings')[0] if result.stdout else "" + + return BenchResult( + context_size=context_size, + kv_type=kv_type, + attn_rot_enabled=attn_rot_enabled, + prompt_tokens=prompt_tokens, + generation_tokens=gen_tokens, + prompt_time_ms=prompt_time, + generation_time_ms=gen_time, + prompt_tps=prompt_tps, + tokens_per_sec=tps, + output_text=gen_text.strip()[:500], # Truncate for display + success=result.returncode == 0, + context_memory_mib=context_memory, + total_memory_mib=total_memory, + error=result.stderr if result.returncode != 0 else None, + ) + + except subprocess.TimeoutExpired: + return BenchResult( + context_size=context_size, + kv_type=kv_type, + attn_rot_enabled=attn_rot_enabled, + prompt_tokens=0, + generation_tokens=0, + prompt_time_ms=0, + generation_time_ms=0, + prompt_tps=0, + tokens_per_sec=0, + output_text="", + success=False, + error="Timeout", + ) + except Exception as e: + return BenchResult( + context_size=context_size, + kv_type=kv_type, + attn_rot_enabled=attn_rot_enabled, + prompt_tokens=0, + generation_tokens=0, + prompt_time_ms=0, + generation_time_ms=0, + prompt_tps=0, + tokens_per_sec=0, + output_text="", + success=False, + error=str(e), + ) + + +def generate_long_prompt(target_tokens: int) -> str: + """Generate a prompt that will fill context to approximately target_tokens.""" + base = "The following is a detailed analysis of complex systems. " + filler = "Consider the intricate relationships between variables in dynamic environments. " + + # Rough estimate: 1 token ≈ 4 chars + target_chars = target_tokens * 4 + prompt = base + while len(prompt) < target_chars: + prompt += filler + + prompt += "\n\nBased on the above, provide a brief summary: " + return prompt + + +def run_comparison( + model_path: str, + kv_type: str, + context_size: int, + prompt: str, + n_predict: int, + n_gpu_layers: int, + llama_cli_path: str, + debug: bool = False, +) -> tuple[BenchResult, BenchResult]: + """Run benchmark with ATTN_ROT on and off, return both results.""" + + # With rotation + result_on = run_benchmark( + model_path=model_path, + context_size=context_size, + kv_type=kv_type, + attn_rot_enabled=True, + prompt=prompt, + n_predict=n_predict, + n_gpu_layers=n_gpu_layers, + llama_cli_path=llama_cli_path, + debug=debug, + ) + + # Without rotation + result_off = run_benchmark( + model_path=model_path, + context_size=context_size, + kv_type=kv_type, + attn_rot_enabled=False, + prompt=prompt, + n_predict=n_predict, + n_gpu_layers=n_gpu_layers, + llama_cli_path=llama_cli_path, + debug=debug, + ) + + return result_on, result_off + + +def print_result(r: BenchResult): + """Print a single benchmark result.""" + status = "✓" if r.success else "✗" + rot = "ON" if r.attn_rot_enabled else "OFF" + mem_str = f", KV={r.context_memory_mib:.0f}MiB" if r.context_memory_mib > 0 else "" + print(f" [{status}] ATTN_ROT={rot}: prompt={r.prompt_tps:.1f} t/s, gen={r.tokens_per_sec:.1f} t/s{mem_str}") + if not r.success and r.error: + print(f" Error: {r.error[:100]}") + + +def print_comparison(on: BenchResult, off: BenchResult): + """Print comparison between two results.""" + if on.success and off.success and off.tokens_per_sec > 0: + speedup = on.tokens_per_sec / off.tokens_per_sec + print(f" → Speed ratio (ON/OFF): {speedup:.2f}x") + print() + + +def main(): + parser = argparse.ArgumentParser( + description="Benchmark KV cache rotation (ATTN_ROT) feature" + ) + parser.add_argument( + "-m", "--model", + default="Bonsai-8B_patched.gguf", + help="Path to model file" + ) + parser.add_argument( + "--llama-cli", + default="./build/bin/llama-completion", + help="Path to llama-completion binary" + ) + parser.add_argument( + "-c", "--contexts", + type=int, + nargs="+", + default=[64], + help="Context sizes to test" + ) + parser.add_argument( + "--kv-types", + nargs="+", + default=["q4_0", "q8_0"], + help="KV cache quantization types to test" + ) + parser.add_argument( + "-n", "--n-predict", + type=int, + default=128, + help="Number of tokens to generate" + ) + parser.add_argument( + "--ngl", + type=int, + default=99, + help="Number of GPU layers" + ) + parser.add_argument( + "-p", "--prompt", + default="Explain quantum computing in simple terms.", + help="Prompt to use (or 'fill' to auto-generate long prompts)" + ) + parser.add_argument( + "--fill-ratio", + type=float, + default=0.5, + help="When using 'fill' prompt, what ratio of context to fill (0.0-0.9)" + ) + parser.add_argument( + "--quick", + action="store_true", + help="Quick test with minimal contexts" + ) + parser.add_argument( + "-o", "--output", + help="Save results to JSON file" + ) + parser.add_argument( + "--debug", + action="store_true", + help="Print captured output for debugging" + ) + parser.add_argument( + "--include-f16", + action="store_true", + help="Include f16 KV cache baseline for memory comparison" + ) + + args = parser.parse_args() + + if args.quick: + args.contexts = [2048, 8192] + args.kv_types = ["q4_0"] + args.n_predict = 64 + + # Add f16 baseline if requested + kv_types_to_test = args.kv_types.copy() + if args.include_f16 and "f16" not in kv_types_to_test: + kv_types_to_test.insert(0, "f16") + + print("=" * 70) + print("ATTN_ROT Benchmark") + print("=" * 70) + print(f"Model: {args.model}") + print(f"Contexts: {args.contexts}") + print(f"KV Types: {kv_types_to_test}") + print(f"Generate: {args.n_predict} tokens") + if args.include_f16: + print("(Including f16 baseline for memory comparison)") + print("=" * 70) + + all_results = [] + + for kv_type in kv_types_to_test: + print(f"\n{'#'*70}") + print(f"# KV Cache Type: {kv_type}") + print(f"{'#'*70}") + + for ctx in args.contexts: + # Generate prompt + if args.prompt == "fill": + target_tokens = int(ctx * args.fill_ratio) + prompt = generate_long_prompt(target_tokens) + print(f"\nContext {ctx} (filling ~{target_tokens} tokens):") + else: + prompt = args.prompt + print(f"\nContext {ctx}:") + + result_on, result_off = run_comparison( + model_path=args.model, + kv_type=kv_type, + context_size=ctx, + prompt=prompt, + n_predict=args.n_predict, + n_gpu_layers=args.ngl, + llama_cli_path=args.llama_cli, + debug=args.debug, + ) + + print_result(result_on) + print_result(result_off) + print_comparison(result_on, result_off) + + all_results.append({ + "context": ctx, + "kv_type": kv_type, + "attn_rot_on": { + "success": result_on.success, + "prompt_tps": result_on.prompt_tps, + "gen_tps": result_on.tokens_per_sec, + "kv_memory_mib": result_on.context_memory_mib, + }, + "attn_rot_off": { + "success": result_off.success, + "prompt_tps": result_off.prompt_tps, + "gen_tps": result_off.tokens_per_sec, + "kv_memory_mib": result_off.context_memory_mib, + }, + }) + + # Summary + print("\n" + "=" * 100) + print("SUMMARY - Generation Speed (t/s) and KV Cache Memory (MiB)") + print("=" * 100) + print(f"{'Context':<10} {'KV':<6} {'ON gen':<10} {'OFF gen':<10} {'Ratio':<8} {'KV Mem':<10}") + print("-" * 100) + + for r in all_results: + on_g = r["attn_rot_on"]["gen_tps"] + off_g = r["attn_rot_off"]["gen_tps"] + kv_mem = r["attn_rot_on"]["kv_memory_mib"] # Same for ON/OFF + ratio = on_g / off_g if off_g > 0 else 0 + mem_str = f"{kv_mem:.0f}" if kv_mem > 0 else "N/A" + print(f"{r['context']:<10} {r['kv_type']:<6} {on_g:<10.1f} {off_g:<10.1f} {ratio:<8.2f}x {mem_str:<10}") + + if args.output: + with open(args.output, 'w') as f: + json.dump(all_results, f, indent=2) + print(f"\nResults saved to: {args.output}") + + +if __name__ == "__main__": + main() diff --git a/ggml/include/ggml.h b/ggml/include/ggml.h index 669f66b650fb..ebab07cf3645 100644 --- a/ggml/include/ggml.h +++ b/ggml/include/ggml.h @@ -428,7 +428,9 @@ extern "C" { // GGML_TYPE_IQ4_NL_8_8 = 38, GGML_TYPE_MXFP4 = 39, // MXFP4 (1 block) GGML_TYPE_NVFP4 = 40, // NVFP4 (4 blocks, E4M3 scale) - GGML_TYPE_COUNT = 41, + GGML_TYPE_Q1_0_g128 = 41, + GGML_TYPE_Q1_0 = 42, + GGML_TYPE_COUNT = 43, }; // precision @@ -465,6 +467,8 @@ extern "C" { GGML_FTYPE_MOSTLY_BF16 = 24, // except 1d tensors GGML_FTYPE_MOSTLY_MXFP4 = 25, // except 1d tensors GGML_FTYPE_MOSTLY_NVFP4 = 26, // except 1d tensors + GGML_FTYPE_MOSTLY_Q1_0_g128 = 27, // except 1d tensors + GGML_FTYPE_MOSTLY_Q1_0 = 28, // except 1d tensors }; // available tensor operations: diff --git a/ggml/src/ggml-common.h b/ggml/src/ggml-common.h index 92cf739e7a7b..e5d0aedd5b7a 100644 --- a/ggml/src/ggml-common.h +++ b/ggml/src/ggml-common.h @@ -93,6 +93,13 @@ typedef sycl::half2 ggml_half2; // QR = QK / number of values before dequantization // QI = number of 32 bit integers before dequantization +#define QI1_0 (QK1_0 / 32) // Number of int32s needed for QK1_0 bits (QK1_0/32) +#define QR1_0 1 // 1 bit per quantized element (matches the 1-bit nature of Q1_0) + +#define QI1_0_g128 (QK1_0_g128 / 32) // Number of int32s needed for QK1_0_g128 bits (QK1_0_g128/32) +#define QR1_0_g128 1 // 1 bit per quantized element (matches the 1-bit nature of Q1_0_g128) + + #define QI4_0 (QK4_0 / (4 * QR4_0)) #define QR4_0 2 @@ -170,6 +177,20 @@ typedef sycl::half2 ggml_half2; #define GGML_EXTENSION __extension__ #endif // _MSC_VER +#define QK1_0 32 // MUST match QK8_0 for vec_dot computation! TODO see if we can do larger blocks later +typedef struct { + ggml_half d; // delta + uint8_t qs[QK1_0 / 8]; // bits / quants +} block_q1_0; +static_assert(sizeof(block_q1_0) == sizeof(ggml_half) + QK1_0 / 8, "wrong q1_0 block size/padding"); + +#define QK1_0_g128 128 +typedef struct { + ggml_half d; // delta + uint8_t qs[QK1_0_g128 / 8]; // bits / quants +} block_q1_0_g128; +static_assert(sizeof(block_q1_0_g128) == sizeof(ggml_half) + QK1_0_g128 / 8, "wrong q1_0_g128 block size/padding"); + #define QK4_0 32 typedef struct { ggml_half d; // delta diff --git a/ggml/src/ggml-cpu/arch/arm/quants.c b/ggml/src/ggml-cpu/arch/arm/quants.c index 82b048bb3ae4..3cf935a38edf 100644 --- a/ggml/src/ggml-cpu/arch/arm/quants.c +++ b/ggml/src/ggml-cpu/arch/arm/quants.c @@ -137,6 +137,186 @@ void quantize_row_q8_K(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, in //===================================== Dot products ================================= +void ggml_vec_dot_q1_0_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) { + const int qk = QK1_0; + const int nb = n / qk; + + assert(n % qk == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const block_q1_0 * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + float sumf = 0.0f; + +#if defined(__ARM_NEON) + float32x4_t sumv = vdupq_n_f32(0.0f); + + for (int i = 0; i < nb; i++) { + const float d0 = GGML_CPU_FP16_TO_FP32(x[i].d); + const float d1 = GGML_CPU_FP16_TO_FP32(y[i].d); + + const uint8_t * bits = x[i].qs; + + const int8x16_t y0 = vld1q_s8(y[i].qs); + const int8x16_t y1 = vld1q_s8(y[i].qs + 16); + + const uint64_t expand0 = table_b2b_0[bits[0]]; + const uint64_t expand1 = table_b2b_0[bits[1]]; + const uint64_t expand2 = table_b2b_0[bits[2]]; + const uint64_t expand3 = table_b2b_0[bits[3]]; + + uint8x8_t e0 = vcreate_u8(expand0); + uint8x8_t e1 = vcreate_u8(expand1); + uint8x8_t e2 = vcreate_u8(expand2); + uint8x8_t e3 = vcreate_u8(expand3); + + int8x8_t s0 = vreinterpret_s8_u8(vshr_n_u8(e0, 4)); + int8x8_t s1 = vreinterpret_s8_u8(vshr_n_u8(e1, 4)); + int8x8_t s2 = vreinterpret_s8_u8(vshr_n_u8(e2, 4)); + int8x8_t s3 = vreinterpret_s8_u8(vshr_n_u8(e3, 4)); + + int8x8_t one = vdup_n_s8(1); + s0 = vsub_s8(vadd_s8(s0, s0), one); + s1 = vsub_s8(vadd_s8(s1, s1), one); + s2 = vsub_s8(vadd_s8(s2, s2), one); + s3 = vsub_s8(vadd_s8(s3, s3), one); + + int8x16_t signs0 = vcombine_s8(s0, s1); + int8x16_t signs1 = vcombine_s8(s2, s3); + + int32x4_t p0 = ggml_vdotq_s32(vdupq_n_s32(0), signs0, y0); + int32x4_t p1 = ggml_vdotq_s32(p0, signs1, y1); + + sumv = vmlaq_n_f32(sumv, vcvtq_f32_s32(p1), d0 * d1); + } + + sumf = vaddvq_f32(sumv); +#else + ggml_vec_dot_q1_0_q8_0_generic(n, &sumf, bs, vx, bx, vy, by, 1); +#endif + + *s = sumf; +} + +void ggml_vec_dot_q1_0_g128_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) { + const int qk = QK1_0_g128; // 128 + const int nb = n / qk; + + assert(n % qk == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const block_q1_0_g128 * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + float sumf = 0.0f; + +#if defined(__ARM_NEON) + // Process one Q1_0_g128 block at a time + // Each block has 128 1-bit values and needs 4 Q8_0 blocks (4 * 32 = 128) + // + // Strategy: For 1-bit quants, bit=1 means +1, bit=0 means -1 + // dot_product = sum(xi * yi) where xi is +1 or -1 + // = sum_where_bit_1(yi) - sum_where_bit_0(yi) + // = 2 * sum_where_bit_1(yi) - sum_all(yi) + // + // We use the lookup table approach: expand each byte of bits to 8 bytes + // where each byte is either 0x00 (bit=0) or 0x10 (bit=1), then use as mask + + float32x4_t sumv = vdupq_n_f32(0.0f); + + for (int i = 0; i < nb; i++) { + const float d0 = GGML_CPU_FP16_TO_FP32(x[i].d); + + // Process 4 Q8_0 blocks (each has 32 elements) + for (int k = 0; k < 4; k++) { + const block_q8_0 * GGML_RESTRICT yb = &y[i * 4 + k]; + const float d1 = GGML_CPU_FP16_TO_FP32(yb->d); + + // Get the 4 bytes of bits for this Q8_0 block (32 bits = 4 bytes) + // Bits are at offset k*4 bytes in x[i].qs + const uint8_t * bits = &x[i].qs[k * 4]; + + // Load 32 int8 values from y + const int8x16_t y0 = vld1q_s8(yb->qs); + const int8x16_t y1 = vld1q_s8(yb->qs + 16); + + // Byte 0-1: bits for y0[0..15] + const uint64_t expand0 = table_b2b_0[bits[0]]; + const uint64_t expand1 = table_b2b_0[bits[1]]; + // Byte 2-3: bits for y1[0..15] + const uint64_t expand2 = table_b2b_0[bits[2]]; + const uint64_t expand3 = table_b2b_0[bits[3]]; + + // Build the sign vectors by reinterpreting the table values + uint8x8_t e0 = vcreate_u8(expand0); + uint8x8_t e1 = vcreate_u8(expand1); + uint8x8_t e2 = vcreate_u8(expand2); + uint8x8_t e3 = vcreate_u8(expand3); + + // Shift right by 4 to get 0 or 1 + int8x8_t s0 = vreinterpret_s8_u8(vshr_n_u8(e0, 4)); + int8x8_t s1 = vreinterpret_s8_u8(vshr_n_u8(e1, 4)); + int8x8_t s2 = vreinterpret_s8_u8(vshr_n_u8(e2, 4)); + int8x8_t s3 = vreinterpret_s8_u8(vshr_n_u8(e3, 4)); + + // Convert 0/1 to -1/+1: sign = 2*val - 1 + int8x8_t one = vdup_n_s8(1); + s0 = vsub_s8(vadd_s8(s0, s0), one); // 2*s0 - 1 + s1 = vsub_s8(vadd_s8(s1, s1), one); + s2 = vsub_s8(vadd_s8(s2, s2), one); + s3 = vsub_s8(vadd_s8(s3, s3), one); + + // Combine into 16-element vectors + int8x16_t signs0 = vcombine_s8(s0, s1); + int8x16_t signs1 = vcombine_s8(s2, s3); + + // Multiply signs with y values and accumulate + // dot(signs, y) where signs are +1/-1 + int32x4_t p0 = ggml_vdotq_s32(vdupq_n_s32(0), signs0, y0); + int32x4_t p1 = ggml_vdotq_s32(p0, signs1, y1); + + // Scale by d1 and accumulate + sumv = vmlaq_n_f32(sumv, vcvtq_f32_s32(p1), d0 * d1); + } + } + + sumf = vaddvq_f32(sumv); +#else + // Scalar fallback + for (int i = 0; i < nb; i++) { + const float d0 = GGML_FP16_TO_FP32(x[i].d); + + // Process 4 Q8_0 blocks + for (int k = 0; k < 4; k++) { + const float d1 = GGML_FP16_TO_FP32(y[i*4 + k].d); + + int sumi = 0; + for (int j = 0; j < QK8_0; j++) { + const int bit_index = k * QK8_0 + j; + const int byte_index = bit_index / 8; + const int bit_offset = bit_index % 8; + + const int xi = ((x[i].qs[byte_index] >> bit_offset) & 1) ? 1 : -1; + sumi += xi * y[i*4 + k].qs[j]; + } + sumf += d0 * d1 * sumi; + } + } +#endif + + *s = sumf; +} + + void ggml_vec_dot_q4_0_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) { const int qk = QK8_0; const int nb = n / qk; diff --git a/ggml/src/ggml-cpu/arch/x86/quants.c b/ggml/src/ggml-cpu/arch/x86/quants.c index 74d699f633d3..e4130ef22f93 100644 --- a/ggml/src/ggml-cpu/arch/x86/quants.c +++ b/ggml/src/ggml-cpu/arch/x86/quants.c @@ -540,6 +540,14 @@ static inline __m128i get_scale_shuffle(int i) { } #endif +void ggml_vec_dot_q1_0_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) { + ggml_vec_dot_q1_0_q8_0_generic(n, s, bs, vx, bx, vy, by, nrc); +} + +void ggml_vec_dot_q1_0_g128_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) { + ggml_vec_dot_q1_0_g128_q8_0_generic(n, s, bs, vx, bx, vy, by, nrc); +} + void ggml_vec_dot_q4_0_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) { const int qk = QK8_0; const int nb = n / qk; diff --git a/ggml/src/ggml-cpu/ggml-cpu.c b/ggml/src/ggml-cpu/ggml-cpu.c index 7486acc2b5d6..becaa3aff5ea 100644 --- a/ggml/src/ggml-cpu/ggml-cpu.c +++ b/ggml/src/ggml-cpu/ggml-cpu.c @@ -217,6 +217,18 @@ static const struct ggml_type_traits_cpu type_traits_cpu[GGML_TYPE_COUNT] = { .vec_dot_type = GGML_TYPE_F16, .nrows = 1, }, + [GGML_TYPE_Q1_0] = { + .from_float = quantize_row_q1_0, + .vec_dot = ggml_vec_dot_q1_0_q8_0, + .vec_dot_type = GGML_TYPE_Q8_0, + .nrows = 1, + }, + [GGML_TYPE_Q1_0_g128] = { + .from_float = quantize_row_q1_0_g128, + .vec_dot = ggml_vec_dot_q1_0_g128_q8_0, + .vec_dot_type = GGML_TYPE_Q8_0, + .nrows = 1, + }, [GGML_TYPE_Q4_0] = { .from_float = quantize_row_q4_0, .vec_dot = ggml_vec_dot_q4_0_q8_0, diff --git a/ggml/src/ggml-cpu/ops.cpp b/ggml/src/ggml-cpu/ops.cpp index 765ce07f06c2..330c31e11e39 100644 --- a/ggml/src/ggml-cpu/ops.cpp +++ b/ggml/src/ggml-cpu/ops.cpp @@ -4829,6 +4829,8 @@ void ggml_compute_forward_get_rows( const ggml_tensor * src0 = dst->src[0]; switch (src0->type) { + case GGML_TYPE_Q1_0: + case GGML_TYPE_Q1_0_g128: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -5554,6 +5556,8 @@ void ggml_compute_forward_clamp( ggml_compute_forward_clamp_f16(params, dst); } break; case GGML_TYPE_BF16: + case GGML_TYPE_Q1_0: + case GGML_TYPE_Q1_0_g128: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: diff --git a/ggml/src/ggml-cpu/quants.c b/ggml/src/ggml-cpu/quants.c index 7ebbb9c6f15b..37e3c84d5d49 100644 --- a/ggml/src/ggml-cpu/quants.c +++ b/ggml/src/ggml-cpu/quants.c @@ -22,6 +22,14 @@ #define UNUSED GGML_UNUSED +void quantize_row_q1_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { + quantize_row_q1_0_ref(x, y, k); +} + +void quantize_row_q1_0_g128(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { + quantize_row_q1_0_g128_ref(x, y, k); +} + void quantize_row_q4_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k) { quantize_row_q4_0_ref(x, y, k); } @@ -116,6 +124,95 @@ void quantize_row_q8_K_generic(const float * GGML_RESTRICT x, void * GGML_RESTRI //===================================== Dot products ================================= +void ggml_vec_dot_q1_0_q8_0_generic(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) { + const int qk = QK8_0; + const int nb = n / qk; + + assert(n % qk == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const block_q1_0 * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + + float sumf = 0.0; + + for (int i = 0; i < nb; i++) { + const float d0 = GGML_FP16_TO_FP32(x[i].d); + const float d1 = GGML_FP16_TO_FP32(y[i].d); + + int sumi = 0; + + for (int j = 0; j < QK1_0; j++) { + const int bit_index = j; + const int byte_index = bit_index / 8; + const int bit_offset = bit_index % 8; + + // Extract bit: 1 = +1, 0 = -1 + const int xi = ((x[i].qs[byte_index] >> bit_offset) & 1) ? 1 : -1; + const int yi = y[i].qs[j]; + + sumi += xi * yi; + } + + sumf += d0 * d1 * sumi; + } + + *s = sumf; +} + +void ggml_vec_dot_q1_0_g128_q8_0_generic(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) { + const int qk = QK1_0_g128; + const int nb = n / qk; + + assert(n % qk == 0); + assert(nrc == 1); + UNUSED(nrc); + UNUSED(bx); + UNUSED(by); + UNUSED(bs); + + const block_q1_0_g128 * GGML_RESTRICT x = vx; + const block_q8_0 * GGML_RESTRICT y = vy; + + + float sumf = 0.0; + + // Each Q1_0_g128 block has 128 elements, each Q8_0 block has 32 elements + // So we need 4 Q8_0 blocks per Q1_0_g128 block + for (int i = 0; i < nb; i++) { + const float d0 = GGML_FP16_TO_FP32(x[i].d); + + float sumi = 0.0f; + + for (int k = 0; k < 4; k++) { + const float d1 = GGML_FP16_TO_FP32(y[i*4 + k].d); + + int sumi_block = 0; + + for (int j = 0; j < QK8_0; j++) { + const int bit_index = k * QK8_0 + j; + const int byte_index = bit_index / 8; + const int bit_offset = bit_index % 8; + + const int xi = ((x[i].qs[byte_index] >> bit_offset) & 1) ? 1 : -1; + sumi_block += xi * y[i*4 + k].qs[j]; + } + + sumi += d1 * sumi_block; + } + + sumf += d0 * sumi; + } + + *s = sumf; +} + + void ggml_vec_dot_q4_0_q8_0_generic(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) { const int qk = QK8_0; const int nb = n / qk; diff --git a/ggml/src/ggml-cpu/quants.h b/ggml/src/ggml-cpu/quants.h index 3584aaa43e8c..fb74c3c76ae4 100644 --- a/ggml/src/ggml-cpu/quants.h +++ b/ggml/src/ggml-cpu/quants.h @@ -12,6 +12,8 @@ extern "C" { #endif // Quantization +void quantize_row_q1_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void quantize_row_q1_0_g128(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q4_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q4_1(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); void quantize_row_q5_0(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); @@ -36,6 +38,8 @@ void quantize_row_iq4_nl (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, void quantize_row_iq4_xs (const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); // Dot product +void ggml_vec_dot_q1_0_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 ggml_vec_dot_q1_0_g128_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 ggml_vec_dot_q4_0_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 ggml_vec_dot_q4_1_q8_1(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 ggml_vec_dot_q5_0_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); @@ -68,6 +72,8 @@ void ggml_vec_dot_iq3_s_q8_K (int n, float * GGML_RESTRICT s, size_t bs, const void quantize_row_q8_0_generic(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k); void quantize_row_q8_1_generic(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k); void quantize_row_q8_K_generic(const float * GGML_RESTRICT x, void * GGML_RESTRICT y, int64_t k); +void ggml_vec_dot_q1_0_q8_0_generic(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 ggml_vec_dot_q1_0_g128_q8_0_generic(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 ggml_vec_dot_q4_0_q8_0_generic(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 ggml_vec_dot_q4_1_q8_1_generic(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 ggml_vec_dot_q5_0_q8_0_generic(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); diff --git a/ggml/src/ggml-cuda/argsort.cu b/ggml/src/ggml-cuda/argsort.cu index 38fdf3678c11..676830eaaf02 100644 --- a/ggml/src/ggml-cuda/argsort.cu +++ b/ggml/src/ggml-cuda/argsort.cu @@ -2,7 +2,7 @@ #ifdef GGML_CUDA_USE_CUB # include -# if (CCCL_MAJOR_VERSION >= 3 && CCCL_MINOR_VERSION >= 1) +# if (CCCL_MAJOR_VERSION == 3 && CCCL_MINOR_VERSION == 1) # define STRIDED_ITERATOR_AVAILABLE # endif using namespace cub; diff --git a/ggml/src/ggml-cuda/common.cuh b/ggml/src/ggml-cuda/common.cuh index 9affe023403e..a09435922d79 100644 --- a/ggml/src/ggml-cuda/common.cuh +++ b/ggml/src/ggml-cuda/common.cuh @@ -953,6 +953,20 @@ struct ggml_cuda_type_traits { static constexpr int qi = QI8_0; }; +template<> +struct ggml_cuda_type_traits { + static constexpr int qk = QK1_0; + static constexpr int qr = QR1_0; + static constexpr int qi = QI1_0; +}; + +template<> +struct ggml_cuda_type_traits { + static constexpr int qk = QK1_0_g128; + static constexpr int qr = QR1_0_g128; + static constexpr int qi = QI1_0_g128; +}; + template<> struct ggml_cuda_type_traits { static constexpr int qk = QK_MXFP4; diff --git a/ggml/src/ggml-cuda/convert.cu b/ggml/src/ggml-cuda/convert.cu index 79ccfe568a23..7e5fb3efe314 100644 --- a/ggml/src/ggml-cuda/convert.cu +++ b/ggml/src/ggml-cuda/convert.cu @@ -711,6 +711,10 @@ to_bf16_cuda_t ggml_get_to_bf16_cuda(ggml_type type) { to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) { switch (type) { + case GGML_TYPE_Q1_0: + return dequantize_block_cont_cuda; + case GGML_TYPE_Q1_0_g128: + return dequantize_block_cont_cuda; case GGML_TYPE_Q4_0: return dequantize_row_q4_0_cuda; case GGML_TYPE_Q4_1: @@ -767,6 +771,10 @@ to_fp16_cuda_t ggml_get_to_fp16_cuda(ggml_type type) { to_fp32_cuda_t ggml_get_to_fp32_cuda(ggml_type type) { switch (type) { + case GGML_TYPE_Q1_0: + return dequantize_block_cont_cuda; + case GGML_TYPE_Q1_0_g128: + return dequantize_block_cont_cuda; case GGML_TYPE_Q4_0: return dequantize_row_q4_0_cuda; case GGML_TYPE_Q4_1: @@ -822,6 +830,10 @@ to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type) { switch (type) { case GGML_TYPE_F32: return convert_unary_cuda; + case GGML_TYPE_Q1_0: + return dequantize_block_cuda; + case GGML_TYPE_Q1_0_g128: + return dequantize_block_cuda; case GGML_TYPE_Q4_0: return dequantize_block_cuda; case GGML_TYPE_Q4_1: @@ -843,6 +855,10 @@ to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type) { switch (type) { case GGML_TYPE_F32: return convert_unary_cuda; + case GGML_TYPE_Q1_0: + return dequantize_block_cuda; + case GGML_TYPE_Q1_0_g128: + return dequantize_block_cuda; case GGML_TYPE_Q4_0: return dequantize_block_cuda; case GGML_TYPE_Q4_1: @@ -864,6 +880,10 @@ to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type) { switch (type) { case GGML_TYPE_F16: return convert_unary_cuda; + case GGML_TYPE_Q1_0: + return dequantize_block_cuda; + case GGML_TYPE_Q1_0_g128: + return dequantize_block_cuda; case GGML_TYPE_Q4_0: return dequantize_block_cuda; case GGML_TYPE_Q4_1: diff --git a/ggml/src/ggml-cuda/dequantize.cuh b/ggml/src/ggml-cuda/dequantize.cuh index e060fb29fdc0..3f4215170629 100644 --- a/ggml/src/ggml-cuda/dequantize.cuh +++ b/ggml/src/ggml-cuda/dequantize.cuh @@ -1,5 +1,51 @@ #include "common.cuh" +static __device__ __forceinline__ void dequantize_q1_0(const void * vx, const int64_t ib, const int iqs, float2 & v){ + const block_q1_0 * x = (const block_q1_0 *) vx; + + const float d = x[ib].d; + const float neg_d = -d; + + const int bit_index_0 = iqs; + const int bit_index_1 = iqs + 1; + + const int byte_index_0 = bit_index_0 / 8; + const int bit_offset_0 = bit_index_0 % 8; + + const int byte_index_1 = bit_index_1 / 8; + const int bit_offset_1 = bit_index_1 % 8; + + // Extract bits: 1 = +d, 0 = -d + const uint8_t bit_0 = (x[ib].qs[byte_index_0] >> bit_offset_0) & 1; + const uint8_t bit_1 = (x[ib].qs[byte_index_1] >> bit_offset_1) & 1; + + v.x = bit_0 ? d : neg_d; + v.y = bit_1 ? d : neg_d; +} + +static __device__ __forceinline__ void dequantize_q1_0_g128(const void * vx, const int64_t ib, const int iqs, float2 & v){ + const block_q1_0_g128 * x = (const block_q1_0_g128 *) vx; + + const float d = x[ib].d; + const float neg_d = -d; + + const int bit_index_0 = iqs; + const int bit_index_1 = iqs + 1; + + const int byte_index_0 = bit_index_0 / 8; + const int bit_offset_0 = bit_index_0 % 8; + + const int byte_index_1 = bit_index_1 / 8; + const int bit_offset_1 = bit_index_1 % 8; + + // Extract bits: 1 = +d, 0 = -d + const uint8_t bit_0 = (x[ib].qs[byte_index_0] >> bit_offset_0) & 1; + const uint8_t bit_1 = (x[ib].qs[byte_index_1] >> bit_offset_1) & 1; + + v.x = bit_0 ? d : neg_d; + v.y = bit_1 ? d : neg_d; +} + static __device__ __forceinline__ void dequantize_q4_0(const void * vx, const int64_t ib, const int iqs, float2 & v){ const block_q4_0 * x = (const block_q4_0 *) vx; diff --git a/ggml/src/ggml-cuda/ggml-cuda.cu b/ggml/src/ggml-cuda/ggml-cuda.cu index 75b62129adea..4a4025b94430 100644 --- a/ggml/src/ggml-cuda/ggml-cuda.cu +++ b/ggml/src/ggml-cuda/ggml-cuda.cu @@ -4785,6 +4785,8 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g switch (a->type) { case GGML_TYPE_F32: case GGML_TYPE_F16: + case GGML_TYPE_Q1_0: + case GGML_TYPE_Q1_0_g128: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -4822,6 +4824,8 @@ static bool ggml_backend_cuda_device_supports_op(ggml_backend_dev_t dev, const g case GGML_TYPE_F32: case GGML_TYPE_BF16: case GGML_TYPE_I32: + case GGML_TYPE_Q1_0: + case GGML_TYPE_Q1_0_g128: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: diff --git a/ggml/src/ggml-cuda/mmq.cu b/ggml/src/ggml-cuda/mmq.cu index 27b4145ac9ab..85c5d9ad4eb8 100644 --- a/ggml/src/ggml-cuda/mmq.cu +++ b/ggml/src/ggml-cuda/mmq.cu @@ -5,6 +5,13 @@ static void ggml_cuda_mul_mat_q_switch_type(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) { switch (args.type_x) { + // TODO: Q1_0/Q1_0_g128 MMQ disabled due to accuracy issues; for now commenting these to use cuBLAS fallback + case GGML_TYPE_Q1_0: + mul_mat_q_case(ctx, args, stream); + break; + case GGML_TYPE_Q1_0_g128: + mul_mat_q_case(ctx, args, stream); + break; case GGML_TYPE_Q4_0: mul_mat_q_case(ctx, args, stream); break; @@ -270,6 +277,9 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t bool mmq_supported; switch (type) { + // TODO: Q1_0 and Q1_0_g128 MMQ implementation exists but is currently disabled due to accuracy issues + case GGML_TYPE_Q1_0: + case GGML_TYPE_Q1_0_g128: case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: @@ -301,6 +311,10 @@ bool ggml_cuda_should_use_mmq(enum ggml_type type, int cc, int64_t ne11, int64_t return false; } + if ((type == GGML_TYPE_Q1_0 || type == GGML_TYPE_Q1_0_g128) && !turing_mma_available(cc)) { + return false; + } + if (turing_mma_available(cc)) { return true; } diff --git a/ggml/src/ggml-cuda/mmq.cuh b/ggml/src/ggml-cuda/mmq.cuh index 51e8dad4ce7b..a2fadaae786a 100644 --- a/ggml/src/ggml-cuda/mmq.cuh +++ b/ggml/src/ggml-cuda/mmq.cuh @@ -11,6 +11,7 @@ using namespace ggml_cuda_mma; #define MMQ_DP4A_MAX_BATCH_SIZE 64 // Max. batch size to use for dp4a MMQ kernels when FP16 tensor cores are available. #define MMQ_ITER_K 256 +#define MMQ_ITER_K_Q1_0 128 // For Q1_0: 32 blocks per row, QI1_0=1, so threads_per_row = 128/(4*1) = 32 #define MMQ_ITER_K_MXFP4_FP4 512 #define MMQ_NWARPS 8 @@ -57,6 +58,9 @@ static_assert(sizeof(block_fp4_mmq) == sizeof(block_q8_1_mmq), "Unexpected b static mmq_q8_1_ds_layout mmq_get_q8_1_ds_layout(const ggml_type type_x) { switch (type_x) { + case GGML_TYPE_Q1_0: + case GGML_TYPE_Q1_0_g128: + return MMQ_Q8_1_DS_LAYOUT_D4; case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: return MMQ_Q8_1_DS_LAYOUT_DS4; @@ -209,15 +213,17 @@ static constexpr __host__ __device__ tile_x_sizes mmq_get_dp4a_tile_x_sizes(ggml } } -#define MMQ_MMA_TILE_X_K_Q8_0 (2*MMQ_TILE_NE_K + 2*MMQ_TILE_NE_K/QI8_0 + 4) -#define MMQ_MMA_TILE_X_K_FP4 (2*MMQ_TILE_NE_K + 8 + 4) // MXFP4 -#define MMQ_MMA_TILE_X_K_NVFP4 (2*MMQ_TILE_NE_K + MMQ_TILE_NE_K/2 + 4) // NVFP4 -#define MMQ_MMA_TILE_X_K_Q8_1 (2*MMQ_TILE_NE_K + 2*MMQ_TILE_NE_K/QI8_0 + 4) -#define MMQ_MMA_TILE_X_K_Q2_K (2*MMQ_TILE_NE_K + MMQ_TILE_NE_K + 4) -#define MMQ_MMA_TILE_X_K_Q3_K (2*MMQ_TILE_NE_K + MMQ_TILE_NE_K/2 + 4) -#define MMQ_MMA_TILE_X_K_Q6_K (2*MMQ_TILE_NE_K + MMQ_TILE_NE_K/QI6_K + MMQ_TILE_NE_K/8 + 7) +#define MMQ_MMA_TILE_X_K_Q8_0 (2*MMQ_TILE_NE_K + 2*MMQ_TILE_NE_K/QI8_0 + 4) +#define MMQ_MMA_TILE_X_K_Q8_0_g128 (8*MMQ_TILE_NE_K + 8*MMQ_TILE_NE_K/QI8_0 + 4) +#define MMQ_MMA_TILE_X_K_FP4 (2*MMQ_TILE_NE_K + 8 + 4) // MXFP4 +#define MMQ_MMA_TILE_X_K_NVFP4 (2*MMQ_TILE_NE_K + MMQ_TILE_NE_K/2 + 4) // NVFP4 +#define MMQ_MMA_TILE_X_K_Q8_1 (2*MMQ_TILE_NE_K + 2*MMQ_TILE_NE_K/QI8_0 + 4) +#define MMQ_MMA_TILE_X_K_Q2_K (2*MMQ_TILE_NE_K + MMQ_TILE_NE_K + 4) +#define MMQ_MMA_TILE_X_K_Q3_K (2*MMQ_TILE_NE_K + MMQ_TILE_NE_K/2 + 4) +#define MMQ_MMA_TILE_X_K_Q6_K (2*MMQ_TILE_NE_K + MMQ_TILE_NE_K/QI6_K + MMQ_TILE_NE_K/8 + 7) static_assert(MMQ_MMA_TILE_X_K_Q8_0 % 8 == 4, "Wrong padding."); +static_assert(MMQ_MMA_TILE_X_K_Q8_0_g128 % 8 == 4, "Wrong padding."); static_assert(MMQ_MMA_TILE_X_K_Q8_1 % 8 == 4, "Wrong padding."); static_assert(MMQ_MMA_TILE_X_K_Q2_K % 8 == 4, "Wrong padding."); static_assert(MMQ_MMA_TILE_X_K_Q3_K % 8 == 4, "Wrong padding."); @@ -229,6 +235,8 @@ static_assert(MMQ_MMA_TILE_X_K_NVFP4 % 8 == 4, "Wrong padding."); static constexpr __host__ __device__ int mmq_get_mma_tile_x_k(ggml_type type) { switch (type) { + case GGML_TYPE_Q1_0: return MMQ_MMA_TILE_X_K_Q8_0; + case GGML_TYPE_Q1_0_g128: return MMQ_MMA_TILE_X_K_Q8_0; case GGML_TYPE_Q4_0: return MMQ_MMA_TILE_X_K_Q8_0; case GGML_TYPE_Q4_1: return MMQ_MMA_TILE_X_K_Q8_1; case GGML_TYPE_Q5_0: return MMQ_MMA_TILE_X_K_Q8_0; @@ -302,6 +310,149 @@ static constexpr __device__ int mmq_get_nwarps_device() { // ------------------------------------------------------------ +template static __device__ __forceinline__ void load_tiles_q1_0( + const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { +#if !(defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE)) + GGML_UNUSED_VARS(x, x_tile, kbx0, i_max, stride, mmq_y, need_check); + NO_DEVICE_CODE; +#else + constexpr int nwarps = mmq_get_nwarps_device(); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); + + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + 2*MMQ_TILE_NE_K); + constexpr int blocks_per_iter = MMQ_ITER_K / QK1_0; + constexpr int threads_per_row = blocks_per_iter * QI1_0; + constexpr int nrows = warp_size / threads_per_row; + constexpr int scale_entries_per_row = blocks_per_iter * (QK1_0 / QK8_1); + const int txi = threadIdx.x % threads_per_row; + const int kbx = txi / QI1_0; + + +#pragma unroll + for (int i0 = 0; i0 < mmq_y; i0 += nrows*nwarps) { + int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row; + + if (need_check) { + i = min(i, i_max); + } + + const block_q1_0 * bxi = (const block_q1_0 *) x + kbx0 + i*stride + kbx; + + // Q1_0 has 32 bits (4 bytes) for 32 elements at 1 bit each + // Read all 4 bytes safely to avoid alignment issues + const int qs0 = bxi->qs[0] | (bxi->qs[1] << 8) | (bxi->qs[2] << 16) | (bxi->qs[3] << 24); + + // For MMA: unpack 1-bit values to signed bytes (-1 or +1) + // Process all 32 bits, 4 at a time + int unpacked_bytes[8]; +#pragma unroll + for (int j = 0; j < 8; ++j) { + const int shift = j * 4; + const int bits4 = (qs0 >> shift) & 0x0F; + const int b0 = (bits4 & 0x01) ? 1 : -1; + const int b1 = (bits4 & 0x02) ? 1 : -1; + const int b2 = (bits4 & 0x04) ? 1 : -1; + const int b3 = (bits4 & 0x08) ? 1 : -1; + unpacked_bytes[j] = (b0 & 0xFF) | ((b1 & 0xFF) << 8) | ((b2 & 0xFF) << 16) | ((b3 & 0xFF) << 24); + } + // Store unpacked values +#pragma unroll + for (int j = 0; j < 8; ++j) { + x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + kbx*QI8_0 + j] = unpacked_bytes[j]; + } + } + + constexpr int rows_per_warp = warp_size / scale_entries_per_row; + const int kbxd = threadIdx.x % scale_entries_per_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 / scale_entries_per_row; + + if (need_check) { + i = min(i, i_max); + } + + const block_q1_0 * bxi = (const block_q1_0 *) x + kbx0 + i*stride + kbxd; + + x_df[i*MMQ_MMA_TILE_X_K_Q8_0 + kbxd] = bxi->d; + } +#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) +} + +template static __device__ __forceinline__ void load_tiles_q1_0_g128( + const char * __restrict__ x, int * __restrict__ x_tile, const int kbx0, const int i_max, const int stride) { +#if !(defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE)) + GGML_UNUSED_VARS(x, x_tile, kbx0, i_max, stride, mmq_y, need_check); + NO_DEVICE_CODE; +#else + constexpr int nwarps = mmq_get_nwarps_device(); + constexpr int warp_size = ggml_cuda_get_physical_warp_size(); + + int * x_qs = (int *) x_tile; + float * x_df = (float *) (x_qs + 2*MMQ_TILE_NE_K); + + constexpr int blocks_per_iter = MMQ_ITER_K / QK1_0_g128; + constexpr int threads_per_row = blocks_per_iter * QI1_0_g128; + constexpr int nrows = warp_size / threads_per_row; + constexpr int scale_entries_per_block = QK1_0_g128 / QK8_1; + constexpr int scale_entries_per_row = blocks_per_iter * scale_entries_per_block; + + const int txi = threadIdx.x % threads_per_row; + const int kbx = txi / QI1_0_g128; + const int kqsx = txi % QI1_0_g128; + +#pragma unroll + for (int i0 = 0; i0 < mmq_y; i0 += nrows*nwarps) { + int i = i0 + threadIdx.y*nrows + threadIdx.x/threads_per_row; + + if (need_check) { + i = min(i, i_max); + } + + const block_q1_0_g128 * bxi = (const block_q1_0_g128 *) x + kbx0 + i*stride + kbx; + const int qs_offset = 4*kqsx; + const int qs0 = bxi->qs[qs_offset + 0] | (bxi->qs[qs_offset + 1] << 8) | + (bxi->qs[qs_offset + 2] << 16) | (bxi->qs[qs_offset + 3] << 24); + + int unpacked_bytes[8]; +#pragma unroll + for (int j = 0; j < 8; ++j) { + const int shift = j * 4; + const int bits4 = (qs0 >> shift) & 0x0F; + const int b0 = (bits4 & 0x01) ? 1 : -1; + const int b1 = (bits4 & 0x02) ? 1 : -1; + const int b2 = (bits4 & 0x04) ? 1 : -1; + const int b3 = (bits4 & 0x08) ? 1 : -1; + unpacked_bytes[j] = (b0 & 0xFF) | ((b1 & 0xFF) << 8) | ((b2 & 0xFF) << 16) | ((b3 & 0xFF) << 24); + } + + const int dst_offset = kbx*(scale_entries_per_block*QI8_0) + kqsx*QI8_0; +#pragma unroll + for (int j = 0; j < 8; ++j) { + x_qs[i*MMQ_MMA_TILE_X_K_Q8_0 + dst_offset + j] = unpacked_bytes[j]; + } + } + + const int ksx = threadIdx.x % scale_entries_per_row; + const int scale_block = ksx / scale_entries_per_block; + +#pragma unroll + for (int i0 = 0; i0 < mmq_y; i0 += nwarps) { + int i = i0 + threadIdx.y; + + if (need_check) { + i = min(i, i_max); + } + + const block_q1_0_g128 * bxi = (const block_q1_0_g128 *) x + kbx0 + i*stride + scale_block; + + x_df[i*MMQ_MMA_TILE_X_K_Q8_0 + ksx] = bxi->d; + } +#endif // defined(AMD_MFMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) +} + template static __device__ __forceinline__ void load_tiles_q4_0( 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(); @@ -363,6 +514,15 @@ template static __device__ __forceinline__ void loa } } +template +static __device__ __forceinline__ void vec_dot_q1_mmq_dp4a_disabled( + const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { + // Q1_0 and Q1_0_g128 intentionally target the MMA path only on this branch. + // If DP4A support is needed later for older GPUs, it should be reintroduced and validated separately. + GGML_UNUSED_VARS(x, y, sum, k00, mmq_x, mmq_y); + NO_DEVICE_CODE; +} + template static __device__ __forceinline__ void vec_dot_q4_0_q8_1_dp4a( const int * __restrict__ x, const int * __restrict__ y, float * __restrict__ sum, const int k00) { @@ -3274,6 +3434,23 @@ static __device__ __forceinline__ void mmq_write_back_mma( template struct mmq_type_traits; +template +struct mmq_type_traits { + static constexpr int vdr = VDR_Q1_0_Q8_1_MMQ; + static constexpr load_tiles_mmq_t load_tiles = load_tiles_q1_0; + static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_q8_0_q8_1_mma; + static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q1_mmq_dp4a_disabled; +}; + +template +struct mmq_type_traits { + static constexpr int vdr = VDR_Q1_0_g128_Q8_1_MMQ; + static constexpr load_tiles_mmq_t load_tiles = load_tiles_q1_0_g128; + static constexpr vec_dot_mmq_t vec_dot_mma = vec_dot_q8_0_q8_1_mma; + // The DP4A path is intentionally disabled; keep the MMA path as the validated route. + static constexpr vec_dot_mmq_t vec_dot_dp4a = vec_dot_q1_mmq_dp4a_disabled; +}; + template struct mmq_type_traits { static constexpr int vdr = VDR_Q4_0_Q8_1_MMQ; @@ -4137,6 +4314,8 @@ void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cuda #define DECL_MMQ_CASE(type) \ template void mul_mat_q_case(ggml_backend_cuda_context & ctx, const mmq_args & args, cudaStream_t stream) \ +extern DECL_MMQ_CASE(GGML_TYPE_Q1_0); +extern DECL_MMQ_CASE(GGML_TYPE_Q1_0_g128); extern DECL_MMQ_CASE(GGML_TYPE_Q4_0); extern DECL_MMQ_CASE(GGML_TYPE_Q4_1); extern DECL_MMQ_CASE(GGML_TYPE_Q5_0); diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu index 07b10167bc46..acf25d1023fe 100644 --- a/ggml/src/ggml-cuda/mmvq.cu +++ b/ggml/src/ggml-cuda/mmvq.cu @@ -9,6 +9,8 @@ typedef float (*vec_dot_q_cuda_t)(const void * __restrict__ vbq, const block_q8_ static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type) { switch (type) { + case GGML_TYPE_Q1_0: return vec_dot_q1_0_q8_1; + case GGML_TYPE_Q1_0_g128: return vec_dot_q1_0_g128_q8_1; case GGML_TYPE_Q4_0: return vec_dot_q4_0_q8_1; case GGML_TYPE_Q4_1: return vec_dot_q4_1_q8_1; case GGML_TYPE_Q5_0: return vec_dot_q5_0_q8_1; @@ -36,6 +38,8 @@ static constexpr __device__ vec_dot_q_cuda_t get_vec_dot_q_cuda(ggml_type type) static constexpr __host__ __device__ int get_vdr_mmvq(ggml_type type) { switch (type) { + case GGML_TYPE_Q1_0: return VDR_Q1_0_Q8_1_MMVQ; + case GGML_TYPE_Q1_0_g128: return VDR_Q1_0_g128_Q8_1_MMVQ; case GGML_TYPE_Q4_0: return VDR_Q4_0_Q8_1_MMVQ; case GGML_TYPE_Q4_1: return VDR_Q4_1_Q8_1_MMVQ; case GGML_TYPE_Q5_0: return VDR_Q5_0_Q8_1_MMVQ; @@ -886,6 +890,18 @@ static void mul_mat_vec_q_switch_type( const int nsamples_x, const int nsamples_dst, const int stride_sample_x, const int stride_sample_y, const int stride_sample_dst, const int ids_stride, cudaStream_t stream) { switch (type_x) { + case GGML_TYPE_Q1_0: + mul_mat_vec_q_switch_ncols_dst + (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_Q1_0_g128: + mul_mat_vec_q_switch_ncols_dst + (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_Q4_0: mul_mat_vec_q_switch_ncols_dst (vx, vy, ids, fusion, dst, ncols_x, nrows_x, ncols_dst, stride_row_x, stride_col_y, stride_col_dst, diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu index 4300ffc148cf..28eebe3f07e2 100644 --- a/ggml/src/ggml-cuda/quantize.cu +++ b/ggml/src/ggml-cuda/quantize.cu @@ -297,6 +297,7 @@ void quantize_mmq_q8_1_cuda( const int64_t block_num_y = (ne0 + 4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ - 1) / (4*CUDA_QUANTIZE_BLOCK_SIZE_MMQ); const dim3 num_blocks(ne1, block_num_y, ne2*ne3); const dim3 block_size(CUDA_QUANTIZE_BLOCK_SIZE_MMQ, 1, 1); + switch (mmq_get_q8_1_ds_layout(type_src0)) { case MMQ_Q8_1_DS_LAYOUT_D4: quantize_mmq_q8_1 diff --git a/ggml/src/ggml-cuda/template-instances/generate_cu_files.py b/ggml/src/ggml-cuda/template-instances/generate_cu_files.py index 40d51f93fa4d..a7146e71d482 100755 --- a/ggml/src/ggml-cuda/template-instances/generate_cu_files.py +++ b/ggml/src/ggml-cuda/template-instances/generate_cu_files.py @@ -32,6 +32,7 @@ SOURCE_FATTN_MMA_CASE = "DECL_FATTN_MMA_F16_CASE({head_size_kq}, {head_size_v}, {ncols1}, {ncols2});\n" TYPES_MMQ = [ + "GGML_TYPE_Q1_0", "GGML_TYPE_Q1_0_g128", "GGML_TYPE_Q4_0", "GGML_TYPE_Q4_1", "GGML_TYPE_Q5_0", "GGML_TYPE_Q5_1", "GGML_TYPE_Q8_0", "GGML_TYPE_Q2_K", "GGML_TYPE_Q3_K", "GGML_TYPE_Q4_K", "GGML_TYPE_Q5_K", "GGML_TYPE_Q6_K", "GGML_TYPE_IQ2_XXS", "GGML_TYPE_IQ2_XS", "GGML_TYPE_IQ2_S", "GGML_TYPE_IQ3_XXS", "GGML_TYPE_IQ3_S", diff --git a/ggml/src/ggml-cuda/template-instances/mmq-instance-q1_0.cu b/ggml/src/ggml-cuda/template-instances/mmq-instance-q1_0.cu new file mode 100644 index 000000000000..f0686b0d0d85 --- /dev/null +++ b/ggml/src/ggml-cuda/template-instances/mmq-instance-q1_0.cu @@ -0,0 +1,5 @@ +// This file has been autogenerated by generate_cu_files.py, do not edit manually. + +#include "../mmq.cuh" + +DECL_MMQ_CASE(GGML_TYPE_Q1_0); diff --git a/ggml/src/ggml-cuda/template-instances/mmq-instance-q1_0_g128.cu b/ggml/src/ggml-cuda/template-instances/mmq-instance-q1_0_g128.cu new file mode 100644 index 000000000000..3283041beca6 --- /dev/null +++ b/ggml/src/ggml-cuda/template-instances/mmq-instance-q1_0_g128.cu @@ -0,0 +1,5 @@ +// This file has been autogenerated by generate_cu_files.py, do not edit manually. + +#include "../mmq.cuh" + +DECL_MMQ_CASE(GGML_TYPE_Q1_0_g128); diff --git a/ggml/src/ggml-cuda/top-k.cu b/ggml/src/ggml-cuda/top-k.cu index 785a18389f29..f4c3fd538fa2 100644 --- a/ggml/src/ggml-cuda/top-k.cu +++ b/ggml/src/ggml-cuda/top-k.cu @@ -3,7 +3,7 @@ #ifdef GGML_CUDA_USE_CUB # include -# if (CCCL_MAJOR_VERSION >= 3 && CCCL_MINOR_VERSION >= 2) +# if (CCCL_MAJOR_VERSION == 3 && CCCL_MINOR_VERSION == 2) # define CUB_TOP_K_AVAILABLE using namespace cub; # endif // CCCL_MAJOR_VERSION >= 3 && CCCL_MINOR_VERSION >= 2 diff --git a/ggml/src/ggml-cuda/vecdotq.cuh b/ggml/src/ggml-cuda/vecdotq.cuh index 40b2b41e7e82..13d43105db68 100644 --- a/ggml/src/ggml-cuda/vecdotq.cuh +++ b/ggml/src/ggml-cuda/vecdotq.cuh @@ -106,6 +106,57 @@ static __device__ __forceinline__ uint32_t unpack_ksigns(const uint8_t v) { // VDR = vec dot ratio, how many contiguous integers each thread processes when the vec dot kernel is called // MMVQ = mul_mat_vec_q, MMQ = mul_mat_q +#define VDR_Q1_0_Q8_1_MMVQ 1 +#define VDR_Q1_0_Q8_1_MMQ 1 // Changed from 2 to 1: Q1_0 has only 32 bits (1 int) per block +#define VDR_Q1_0_g128_Q8_1_MMVQ 1 // Process one 32-element chunk at a time for parallelism +#define VDR_Q1_0_g128_Q8_1_MMQ 4 // Q1_0_g128 has 128 bits (4 ints) per block + +template static __device__ __forceinline__ float vec_dot_q1_0_q8_1_impl( + const int * v, const int * u, const float & d1, const half2 & ds8) { + + int sumi = 0; + +#pragma unroll + for (int i = 0; i < vdr; ++i) { + const int vi = v[i]; + + // Unpack 32 bits into 32 signed values (-1 or +1) + // Each bit: 0 -> -1, 1 -> +1 + // Process all 32 bits, converting each to a signed byte + + int vi_bytes[8]; + +#pragma unroll + for (int j = 0; j < 8; ++j) { + // Extract 4 bits and convert each to -1 or +1 + const int shift = j * 4; + const int bits4 = (vi >> shift) & 0x0F; + + // Convert each of the 4 bits to a signed byte, then pack into int + // bit=1 -> +1, bit=0 -> -1 + const int b0 = (bits4 & 0x01) ? 1 : -1; + const int b1 = (bits4 & 0x02) ? 1 : -1; + const int b2 = (bits4 & 0x04) ? 1 : -1; + const int b3 = (bits4 & 0x08) ? 1 : -1; + + // Pack 4 signed bytes into a single int for dp4a + vi_bytes[j] = (b0 & 0xFF) | ((b1 & 0xFF) << 8) | ((b2 & 0xFF) << 16) | ((b3 & 0xFF) << 24); + } + + // Perform dot product using dp4a (4-way int8 dot product) +#pragma unroll + for (int j = 0; j < 8; ++j) { + sumi = ggml_cuda_dp4a(vi_bytes[j], u[8*i + j], sumi); + } + } + + const float2 ds8f = __half22float2(ds8); + + // Q1_0 is symmetric (no offset), so we just multiply by scales + // ds8f.x is the scale from Q8_1, ds8f.y is the precomputed sum (not needed for symmetric quant) + return d1 * ds8f.x * sumi; +} + #define VDR_Q4_0_Q8_1_MMVQ 2 #define VDR_Q4_0_Q8_1_MMQ 4 @@ -669,6 +720,72 @@ static __device__ __forceinline__ float vec_dot_q6_K_q8_1_impl_mmq( return d6 * sumf_d; } +static __device__ __forceinline__ float vec_dot_q1_0_q8_1( + const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { + + const block_q1_0 * bq1_0 = (const block_q1_0 *) vbq + kbx; + + int v[VDR_Q1_0_Q8_1_MMVQ]; + int u[8*VDR_Q1_0_Q8_1_MMVQ]; + + // Q1_0 has 32 bits per block, stored in 4 bytes + // Read all 4 bytes and pack into a single int32 + v[0] = bq1_0->qs[0] | (bq1_0->qs[1] << 8) | (bq1_0->qs[2] << 16) | (bq1_0->qs[3] << 24); + + // Load 8 int32s (each containing 4 int8 values) for all 32 Q8_1 values +#pragma unroll + for (int j = 0; j < 8; ++j) { + u[j] = get_int_b4(bq8_1->qs, j); + } + + return vec_dot_q1_0_q8_1_impl(v, u, bq1_0->d, bq8_1->ds); +} + +static __device__ __forceinline__ float vec_dot_q1_0_g128_q8_1( + const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { + + const block_q1_0_g128 * bq1_0_g128 = (const block_q1_0_g128 *) vbq + kbx; + + // Q1_0_g128: 128 elements with ONE scale + // Q8_1: 32 elements per block with individual scales + // iqs selects which of the 4 chunks of 32 elements to process (0-3) + + const float d1 = bq1_0_g128->d; + + // Process only the chunk specified by iqs + const block_q8_1 * bq8_1_chunk = bq8_1 + iqs; + + // Load 32 bits (4 bytes) for this chunk from Q1_0_g128 + const int offset = iqs * 4; + const int v = bq1_0_g128->qs[offset + 0] | (bq1_0_g128->qs[offset + 1] << 8) | + (bq1_0_g128->qs[offset + 2] << 16) | (bq1_0_g128->qs[offset + 3] << 24); + + // Unpack 32 bits into 32 signed values (-1 or +1) + int vi_bytes[8]; +#pragma unroll + for (int j = 0; j < 8; ++j) { + const int shift = j * 4; + const int bits4 = (v >> shift) & 0x0F; + const int b0 = (bits4 & 0x01) ? 1 : -1; + const int b1 = (bits4 & 0x02) ? 1 : -1; + const int b2 = (bits4 & 0x04) ? 1 : -1; + const int b3 = (bits4 & 0x08) ? 1 : -1; + vi_bytes[j] = (b0 & 0xFF) | ((b1 & 0xFF) << 8) | ((b2 & 0xFF) << 16) | ((b3 & 0xFF) << 24); + } + + // Compute dot product for this 32-element chunk + int sumi = 0; +#pragma unroll + for (int j = 0; j < 8; ++j) { + const int u = get_int_b4(bq8_1_chunk->qs, j); + sumi = ggml_cuda_dp4a(vi_bytes[j], u, sumi); + } + + // Apply Q1_0_g128's single scale and this chunk's Q8_1 scale + const float2 ds8f = __half22float2(bq8_1_chunk->ds); + return d1 * ds8f.x * sumi; +} + static __device__ __forceinline__ float vec_dot_q4_0_q8_1( const void * __restrict__ vbq, const block_q8_1 * __restrict__ bq8_1, const int & kbx, const int & iqs) { diff --git a/ggml/src/ggml-metal/ggml-metal-device.cpp b/ggml/src/ggml-metal/ggml-metal-device.cpp index 89539bd7615c..79ab67bced85 100644 --- a/ggml/src/ggml-metal/ggml-metal-device.cpp +++ b/ggml/src/ggml-metal/ggml-metal-device.cpp @@ -736,6 +736,16 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta suffix = ne00 % 4 == 0 ? "_4" : ""; } } break; + case GGML_TYPE_Q1_0: + { + nsg = N_SG_Q1_0; + nr0 = N_R0_Q1_0; + } break; + case GGML_TYPE_Q1_0_g128: + { + nsg = N_SG_Q1_0_g128; + nr0 = N_R0_Q1_0_g128; + } break; case GGML_TYPE_Q4_0: { nsg = N_SG_Q4_0; @@ -948,6 +958,16 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id(ggml_m smem = 32*sizeof(float)*nr0; suffix = ne00 % 4 == 0 ? "_4" : ""; } break; + case GGML_TYPE_Q1_0: + { + nsg = N_SG_Q1_0; + nr0 = N_R0_Q1_0; + } break; + case GGML_TYPE_Q1_0_g128: + { + nsg = N_SG_Q1_0_g128; + nr0 = N_R0_Q1_0_g128; + } break; case GGML_TYPE_Q4_0: { nsg = N_SG_Q4_0; diff --git a/ggml/src/ggml-metal/ggml-metal-impl.h b/ggml/src/ggml-metal/ggml-metal-impl.h index eb2253e029ad..a5ff876cbf75 100644 --- a/ggml/src/ggml-metal/ggml-metal-impl.h +++ b/ggml/src/ggml-metal/ggml-metal-impl.h @@ -8,6 +8,12 @@ // // TODO: for optimal performance, become function of the device and work size +#define N_R0_Q1_0 4 +#define N_SG_Q1_0 2 + +#define N_R0_Q1_0_g128 4 +#define N_SG_Q1_0_g128 2 + #define N_R0_Q4_0 4 #define N_SG_Q4_0 2 diff --git a/ggml/src/ggml-metal/ggml-metal-ops.cpp b/ggml/src/ggml-metal/ggml-metal-ops.cpp index 3cda21be43e6..919a4ea394a5 100644 --- a/ggml/src/ggml-metal/ggml-metal-ops.cpp +++ b/ggml/src/ggml-metal/ggml-metal-ops.cpp @@ -2047,6 +2047,8 @@ int ggml_metal_op_mul_mat(ggml_metal_op_t ctx, int idx) { op->src[0]->type == GGML_TYPE_F32 || // TODO: helper function op->src[0]->type == GGML_TYPE_F16 || op->src[0]->type == GGML_TYPE_BF16 || + op->src[0]->type == GGML_TYPE_Q1_0 || + op->src[0]->type == GGML_TYPE_Q1_0_g128 || op->src[0]->type == GGML_TYPE_Q4_0 || op->src[0]->type == GGML_TYPE_Q4_1 || op->src[0]->type == GGML_TYPE_Q5_0 || diff --git a/ggml/src/ggml-metal/ggml-metal.metal b/ggml/src/ggml-metal/ggml-metal.metal index 2074211594ce..b6afb013ad44 100644 --- a/ggml/src/ggml-metal/ggml-metal.metal +++ b/ggml/src/ggml-metal/ggml-metal.metal @@ -118,6 +118,112 @@ void dequantize_bf16_t4(device const bfloat4 * src, short il, thread type4 & reg } #endif +template +void dequantize_q1_0(device const block_q1_0 * xb, short il, thread type4x4 & reg) { + device const uint8_t * qs = xb->qs; + const float d = xb->d; + + float4x4 reg_f; + + // Process 16 bits (2 bytes) for each call, since we have il=0,1 + const int offset = il * 16; + + for (int i = 0; i < 16; i++) { + const int bit_idx = offset + i; + const int byte_idx = bit_idx / 8; + const int bit_offset = bit_idx % 8; + + const bool bit_val = (qs[byte_idx] >> bit_offset) & 1; + const float val = bit_val ? d : -d; + + reg_f[i/4][i%4] = val; + } + + reg = (type4x4) reg_f; +} + +template +void dequantize_q1_0_t4(device const block_q1_0 * xb, short il, thread type4 & reg) { + device const uint8_t * qs = xb->qs; + const float d = xb->d; + + float4 reg_f; + + // Process 4 bits for each call + const int offset = il * 4; + + for (int i = 0; i < 4; i++) { + const int bit_idx = offset + i; + const int byte_idx = bit_idx / 8; + const int bit_offset = bit_idx % 8; + + const bool bit_val = (qs[byte_idx] >> bit_offset) & 1; + reg_f[i] = bit_val ? d : -d; + } + + reg = (type4) reg_f; +} + +template +void dequantize_q1_0_g128(device const block_q1_0_g128 * xb, short il, thread type4x4 & reg) { + device const uint8_t * qs = xb->qs; + const float d = xb->d; + const float neg_d = -d; + + // Process 16 bits starting at offset il*16 + // Optimization: process 2 bytes (16 bits) at once for better memory access + const int byte_offset = il * 2; // il*16 bits = il*2 bytes + const uint8_t b0 = qs[byte_offset]; + const uint8_t b1 = qs[byte_offset + 1]; + + float4x4 reg_f; + + // Unroll completely for better ILP + // First byte (bits 0-7) + reg_f[0][0] = (b0 & 0x01) ? d : neg_d; + reg_f[0][1] = (b0 & 0x02) ? d : neg_d; + reg_f[0][2] = (b0 & 0x04) ? d : neg_d; + reg_f[0][3] = (b0 & 0x08) ? d : neg_d; + reg_f[1][0] = (b0 & 0x10) ? d : neg_d; + reg_f[1][1] = (b0 & 0x20) ? d : neg_d; + reg_f[1][2] = (b0 & 0x40) ? d : neg_d; + reg_f[1][3] = (b0 & 0x80) ? d : neg_d; + + // Second byte (bits 8-15) + reg_f[2][0] = (b1 & 0x01) ? d : neg_d; + reg_f[2][1] = (b1 & 0x02) ? d : neg_d; + reg_f[2][2] = (b1 & 0x04) ? d : neg_d; + reg_f[2][3] = (b1 & 0x08) ? d : neg_d; + reg_f[3][0] = (b1 & 0x10) ? d : neg_d; + reg_f[3][1] = (b1 & 0x20) ? d : neg_d; + reg_f[3][2] = (b1 & 0x40) ? d : neg_d; + reg_f[3][3] = (b1 & 0x80) ? d : neg_d; + + reg = (type4x4) reg_f; +} + +template +void dequantize_q1_0_g128_t4(device const block_q1_0_g128 * xb, short il, thread type4 & reg) { + device const uint8_t * qs = xb->qs; + const float d = xb->d; + + float4 reg_f; + + // Process 4 bits for each call + const int offset = il * 4; + + for (int i = 0; i < 4; i++) { + const int bit_idx = offset + i; + const int byte_idx = bit_idx / 8; + const int bit_offset = bit_idx % 8; + + const bool bit_val = (qs[byte_idx] >> bit_offset) & 1; + reg_f[i] = bit_val ? d : -d; + } + + reg = (type4) reg_f; +} + template void dequantize_q4_0(device const block_q4_0 * xb, short il, thread type4x4 & reg) { device const uint16_t * qs = ((device const uint16_t *)xb + 1); @@ -3116,6 +3222,54 @@ kernel void kernel_group_norm_f32( } } +// function for calculate inner product between half a q1_0 block and 16 floats (yl), sumy is SUM(yl[i]) +// il indicates where the q1 quants begin (0 or QK1_0/4) +// we assume that the yl's have been multiplied with the appropriate scale factor +inline float block_q_n_dot_y(device const block_q1_0 * qb_curr, float sumy, thread float * yl, int il) { + float d = qb_curr->d; + + float acc = 0.0f; + + // il represents which half of the block (0 or 16) + // 16 weights = 16 bits = 2 bytes + // il=0 → bytes 0-1 (bits 0-15), il=16 → bytes 2-3 (bits 16-31) + // TODO: if we increase Q1_0 block size this might need to change + const int byte_offset = il / 8; // 0 or 2 + device const uint8_t * qs = qb_curr->qs + byte_offset; + + for (int i = 0; i < 16; i++) { + const uint8_t byte_idx = i / 8; + const uint8_t bit_idx = i % 8; + const int8_t qval = ((qs[byte_idx] >> bit_idx) & 1) ? 1 : -1; + acc += yl[i] * qval; + } + + return d * acc; +} + +// function for calculate inner product between part of a q1_0_g128 block and 16 floats (yl), sumy is SUM(yl[i]) +// il indicates where the q1 quants begin (0, 16, 32, ..., 112 for 128-element block) +// we assume that the yl's have been multiplied with the appropriate scale factor +inline float block_q_n_dot_y(device const block_q1_0_g128 * qb_curr, float sumy, thread float * yl, int il) { + float d = qb_curr->d; + + float acc = 0.0f; + + // il represents which 16-element chunk of the 128-element block (0, 16, 32, ..., 112) + // 16 weights = 16 bits = 2 bytes + const int byte_offset = il / 8; + device const uint8_t * qs = qb_curr->qs + byte_offset; + + for (int i = 0; i < 16; i++) { + const uint8_t byte_idx = i / 8; + const uint8_t bit_idx = i % 8; + const int8_t qval = ((qs[byte_idx] >> bit_idx) & 1) ? 1 : -1; + acc += yl[i] * qval; + } + + return d * acc; +} + // function for calculate inner product between half a q4_0 block and 16 floats (yl), sumy is SUM(yl[i]) // il indicates where the q4 quants begin (0 or QK4_0/4) // we assume that the yl's have been multiplied with the appropriate scale factor @@ -3337,6 +3491,148 @@ void mul_vec_q_n_f32_impl( } } +kernel void kernel_mul_mv_q1_0_f32( + constant ggml_metal_kargs_mul_mv & args, + device const char * src0, + device const char * src1, + device char * dst, + uint3 tgpig[[threadgroup_position_in_grid]], + ushort tiisg[[thread_index_in_simdgroup]], + ushort sgitg[[simdgroup_index_in_threadgroup]]) { + // Q1_0-specific implementation + const int nb = args.ne00/QK1_0; + + const int r0 = tgpig.x; + const int r1 = tgpig.y; + const int im = tgpig.z; + + const int first_row = (r0 * N_SG_Q1_0 + sgitg) * N_R0_Q1_0; + + const uint i12 = im%args.ne12; + const uint i13 = im/args.ne12; + + const uint64_t offset1 = r1*args.nb11 + (i12)*args.nb12 + (i13)*args.nb13; + + device const float * y = (device const float *) (src1 + offset1); + + // pointers to src0 rows + device const block_q1_0 * ax[N_R0_Q1_0]; + for (int row = 0; row < N_R0_Q1_0; ++row) { + const uint64_t offset0 = (first_row + row)*args.nb01 + (i12/args.r2)*args.nb02 + (i13/args.r3)*args.nb03; + + ax[row] = (device const block_q1_0 *) ((device char *) src0 + offset0); + } + + float yl[16]; // src1 vector cache + float sumf[N_R0_Q1_0] = {0.f}; + + const short ix = (tiisg/2); + const short il = (tiisg%2)*16; // 0 or 16 - which half of the 32-element block + + device const float * yb = y + ix*QK1_0 + il; + + // each thread in a SIMD group deals with half a block. + for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/2) { + float sumy = 0.f; + + // Q1_0: simple copy, no fancy scaling (unlike Q4_0) +#pragma unroll + for (short i = 0; i < 16; i++) { + yl[i] = yb[i]; + sumy += yb[i]; + } + +#pragma unroll + for (short row = 0; row < N_R0_Q1_0; row++) { + sumf[row] += block_q_n_dot_y(ax[row] + ib, sumy, yl, il); + } + + yb += QK1_0 * 16; + } + + device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0; + + for (int row = 0; row < N_R0_Q1_0; ++row) { + const float tot = simd_sum(sumf[row]); + + if (tiisg == 0 && first_row + row < args.ne01) { + dst_f32[first_row + row] = tot; + } + } +} + +kernel void kernel_mul_mv_q1_0_g128_f32( + constant ggml_metal_kargs_mul_mv & args, + device const char * src0, + device const char * src1, + device char * dst, + uint3 tgpig[[threadgroup_position_in_grid]], + ushort tiisg[[thread_index_in_simdgroup]], + ushort sgitg[[simdgroup_index_in_threadgroup]]) { + // Q1_0_g128-specific implementation with 128-element blocks + const int nb = args.ne00/QK1_0_g128; + + const int r0 = tgpig.x; + const int r1 = tgpig.y; + const int im = tgpig.z; + + const int first_row = (r0 * N_SG_Q1_0_g128 + sgitg) * N_R0_Q1_0_g128; + + const uint i12 = im%args.ne12; + const uint i13 = im/args.ne12; + + const uint64_t offset1 = r1*args.nb11 + (i12)*args.nb12 + (i13)*args.nb13; + + device const float * y = (device const float *) (src1 + offset1); + + // pointers to src0 rows + device const block_q1_0_g128 * ax[N_R0_Q1_0_g128]; + for (int row = 0; row < N_R0_Q1_0_g128; ++row) { + const uint64_t offset0 = (first_row + row)*args.nb01 + (i12/args.r2)*args.nb02 + (i13/args.r3)*args.nb03; + + ax[row] = (device const block_q1_0_g128 *) ((device char *) src0 + offset0); + } + + float yl[16]; // src1 vector cache + float sumf[N_R0_Q1_0_g128] = {0.f}; + + // For 128-element blocks, we need 8 passes of 16 elements each + // Each thread processes a different 16-element chunk + const short ix = (tiisg/8); // which block (0 to 3 for 32 threads / 8) + const short il = (tiisg%8)*16; // which 16-element chunk within the 128-element block (0, 16, 32, ..., 112) + + device const float * yb = y + ix*QK1_0_g128 + il; + + // each thread in a SIMD group deals with 1/8 of a block (16 elements out of 128) + for (int ib = ix; ib < nb; ib += N_SIMDWIDTH/8) { + float sumy = 0.f; + + // Q1_0_g128: simple copy +#pragma unroll + for (short i = 0; i < 16; i++) { + yl[i] = yb[i]; + sumy += yb[i]; + } + +#pragma unroll + for (short row = 0; row < N_R0_Q1_0_g128; row++) { + sumf[row] += block_q_n_dot_y(ax[row] + ib, sumy, yl, il); + } + + yb += QK1_0_g128 * (N_SIMDWIDTH/8); + } + + device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0; + + for (int row = 0; row < N_R0_Q1_0_g128; ++row) { + const float tot = simd_sum(sumf[row]); + + if (tiisg == 0 && first_row + row < args.ne01) { + dst_f32[first_row + row] = tot; + } + } +} + kernel void kernel_mul_mv_q4_0_f32( constant ggml_metal_kargs_mul_mv & args, device const char * src0, @@ -3729,6 +4025,16 @@ template [[host_name("kernel_mul_mv_ext_bf16_f32_r1_4")]] kernel mul_mv_ext_q4 template [[host_name("kernel_mul_mv_ext_bf16_f32_r1_5")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<5, bfloat4, 4, dequantize_bf16_t4>; #endif +template [[host_name("kernel_mul_mv_ext_q1_0_f32_r1_2")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<2, block_q1_0, 32, dequantize_q1_0_t4>; +template [[host_name("kernel_mul_mv_ext_q1_0_f32_r1_3")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<3, block_q1_0, 32, dequantize_q1_0_t4>; +template [[host_name("kernel_mul_mv_ext_q1_0_f32_r1_4")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<4, block_q1_0, 32, dequantize_q1_0_t4>; +template [[host_name("kernel_mul_mv_ext_q1_0_f32_r1_5")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<5, block_q1_0, 32, dequantize_q1_0_t4>; + +template [[host_name("kernel_mul_mv_ext_q1_0_g128_f32_r1_2")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<2, block_q1_0_g128, 128, dequantize_q1_0_g128_t4>; +template [[host_name("kernel_mul_mv_ext_q1_0_g128_f32_r1_3")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<3, block_q1_0_g128, 128, dequantize_q1_0_g128_t4>; +template [[host_name("kernel_mul_mv_ext_q1_0_g128_f32_r1_4")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<4, block_q1_0_g128, 128, dequantize_q1_0_g128_t4>; +template [[host_name("kernel_mul_mv_ext_q1_0_g128_f32_r1_5")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<5, block_q1_0_g128, 128, dequantize_q1_0_g128_t4>; + template [[host_name("kernel_mul_mv_ext_q4_0_f32_r1_2")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<2, block_q4_0, 32, dequantize_q4_0_t4>; template [[host_name("kernel_mul_mv_ext_q4_0_f32_r1_3")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<3, block_q4_0, 32, dequantize_q4_0_t4>; template [[host_name("kernel_mul_mv_ext_q4_0_f32_r1_4")]] kernel mul_mv_ext_q4_f32_t kernel_mul_mv_ext_q4_f32_disp<4, block_q4_0, 32, dequantize_q4_0_t4>; @@ -9838,6 +10144,8 @@ template [[host_name("kernel_mul_mm_f16_f32")]] kernel mul_mm_t kernel_mul_m #if defined(GGML_METAL_HAS_BF16) template [[host_name("kernel_mul_mm_bf16_f32")]] kernel mul_mm_t kernel_mul_mm; #endif +template [[host_name("kernel_mul_mm_q1_0_f32")]] kernel mul_mm_t kernel_mul_mm; +template [[host_name("kernel_mul_mm_q1_0_g128_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q4_0_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q4_1_f32")]] kernel mul_mm_t kernel_mul_mm; template [[host_name("kernel_mul_mm_q5_0_f32")]] kernel mul_mm_t kernel_mul_mm; diff --git a/ggml/src/ggml-quants.c b/ggml/src/ggml-quants.c index 48695a61ea33..8f0a89935f05 100644 --- a/ggml/src/ggml-quants.c +++ b/ggml/src/ggml-quants.c @@ -32,6 +32,75 @@ static inline int best_index_int8(int n, const int8_t * val, float x) { return x - val[mu-1] < val[mu] - x ? mu-1 : mu; } +// reference implementation for deterministic creation of model files +void quantize_row_q1_0_ref(const float * GGML_RESTRICT x, block_q1_0 * GGML_RESTRICT y, int64_t k) { + static const int qk = QK1_0; + + assert(k % qk == 0); + + const int nb = k / qk; + + for (int i = 0; i < nb; i++) { + float sum_abs = 0.0f; + for (int j = 0; j < qk; j++) { + sum_abs += fabsf(x[i*qk + j]); + } + const float d = sum_abs / qk; + + y[i].d = GGML_FP32_TO_FP16(d); + + // Clear all bits first + for (int j = 0; j < qk / 8; ++j) { + y[i].qs[j] = 0; + } + + // Just store sign of each weight directly (no normalization) + for (int j = 0; j < qk; ++j) { + const int bit_index = j; + const int byte_index = bit_index / 8; + const int bit_offset = bit_index % 8; + + if (x[i*qk + j] >= 0.0f) { + y[i].qs[byte_index] |= (1 << bit_offset); + } + } + } +} + +void quantize_row_q1_0_g128_ref(const float * GGML_RESTRICT x, block_q1_0_g128 * GGML_RESTRICT y, int64_t k) { + static const int qk = QK1_0_g128; + + assert(k % qk == 0); + + const int nb = k / qk; + + for (int i = 0; i < nb; i++) { + float sum_abs = 0.0f; + for (int j = 0; j < qk; j++) { + sum_abs += fabsf(x[i*qk + j]); + } + const float d = sum_abs / qk; + + y[i].d = GGML_FP32_TO_FP16(d); + + // Clear all bits first + for (int j = 0; j < qk / 8; ++j) { + y[i].qs[j] = 0; + } + + // Just store sign of each weight directly (no normalization) + for (int j = 0; j < qk; ++j) { + const int bit_index = j; + const int byte_index = bit_index / 8; + const int bit_offset = bit_index % 8; + + if (x[i*qk + j] >= 0.0f) { + y[i].qs[byte_index] |= (1 << bit_offset); + } + } + } +} + // reference implementation for deterministic creation of model files void quantize_row_q4_0_ref(const float * GGML_RESTRICT x, block_q4_0 * GGML_RESTRICT y, int64_t k) { static const int qk = QK4_0; @@ -339,6 +408,47 @@ void quantize_row_nvfp4_ref(const float * GGML_RESTRICT x, block_nvfp4 * GGML_RE } } +void dequantize_row_q1_0(const block_q1_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + static const int qk = QK1_0; + + assert(k % qk == 0); + + const int nb = k / qk; + + for (int i = 0; i < nb; i++) { + const float d = GGML_FP16_TO_FP32(x[i].d); + const float neg_d = -d; + + // Simple bit unpacking + for (int j = 0; j < qk; ++j) { + const int byte_index = j / 8; + const int bit_offset = j % 8; + const uint8_t bit = (x[i].qs[byte_index] >> bit_offset) & 1; + y[i*qk + j] = bit ? d : neg_d; + } + } +} + +void dequantize_row_q1_0_g128(const block_q1_0_g128 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { + static const int qk = QK1_0_g128; + + assert(k % qk == 0); + + const int nb = k / qk; + + for (int i = 0; i < nb; i++) { + const float d = GGML_FP16_TO_FP32(x[i].d); + const float neg_d = -d; + + for (int j = 0; j < qk; ++j) { + const int byte_index = j / 8; + const int bit_offset = j % 8; + const uint8_t bit = (x[i].qs[byte_index] >> bit_offset) & 1; + y[i*qk + j] = bit ? d : neg_d; + } + } +} + void dequantize_row_q4_0(const block_q4_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k) { static const int qk = QK4_0; @@ -1978,6 +2088,37 @@ static void quantize_row_q4_0_impl(const float * GGML_RESTRICT x, block_q4_0 * G } } +size_t quantize_q1_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { + if (!quant_weights) { + quantize_row_q1_0_ref(src, dst, (int64_t)nrow*n_per_row); + return nrow * ggml_row_size(GGML_TYPE_Q1_0, n_per_row); + } + size_t row_size = ggml_row_size(GGML_TYPE_Q1_0, n_per_row); + char * qrow = (char *)dst; + for (int64_t row = 0; row < nrow; ++row) { + quantize_row_q1_0_ref(src, (block_q1_0*)qrow, n_per_row); + src += n_per_row; + qrow += row_size; + } + return nrow * row_size; +} + +size_t quantize_q1_0_g128(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { + if (!quant_weights) { + quantize_row_q1_0_g128_ref(src, dst, (int64_t)nrow*n_per_row); + return nrow * ggml_row_size(GGML_TYPE_Q1_0_g128, n_per_row); + } + size_t row_size = ggml_row_size(GGML_TYPE_Q1_0_g128, n_per_row); + char * qrow = (char *)dst; + for (int64_t row = 0; row < nrow; ++row) { + quantize_row_q1_0_g128_ref(src, (block_q1_0_g128*)qrow, n_per_row); + src += n_per_row; + qrow += row_size; + } + return nrow * row_size; +} + + size_t quantize_q4_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrow, int64_t n_per_row, const float * quant_weights) { if (!quant_weights) { quantize_row_q4_0_ref(src, dst, (int64_t)nrow*n_per_row); @@ -5286,6 +5427,14 @@ bool ggml_validate_row_data(enum ggml_type type, const void * data, size_t nbyte } } } break; + case GGML_TYPE_Q1_0: + { + VALIDATE_ROW_DATA_D_F16_IMPL(block_q1_0, data, nb); + } break; + case GGML_TYPE_Q1_0_g128: + { + VALIDATE_ROW_DATA_D_F16_IMPL(block_q1_0_g128, data, nb); + } break; case GGML_TYPE_Q4_0: { VALIDATE_ROW_DATA_D_F16_IMPL(block_q4_0, data, nb); diff --git a/ggml/src/ggml-quants.h b/ggml/src/ggml-quants.h index 00604f75c0e9..c64029263933 100644 --- a/ggml/src/ggml-quants.h +++ b/ggml/src/ggml-quants.h @@ -14,6 +14,8 @@ extern "C" { // NOTE: these functions are defined as GGML_API because they used by the CPU backend // Quantization +GGML_API void quantize_row_q1_0_ref(const float * GGML_RESTRICT x, block_q1_0 * GGML_RESTRICT y, int64_t k); +GGML_API void quantize_row_q1_0_g128_ref(const float * GGML_RESTRICT x, block_q1_0_g128 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q4_0_ref(const float * GGML_RESTRICT x, block_q4_0 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q4_1_ref(const float * GGML_RESTRICT x, block_q4_1 * GGML_RESTRICT y, int64_t k); GGML_API void quantize_row_q5_0_ref(const float * GGML_RESTRICT x, block_q5_0 * GGML_RESTRICT y, int64_t k); @@ -41,6 +43,8 @@ GGML_API void quantize_row_iq3_s_ref (const float * GGML_RESTRICT x, block_iq3_ GGML_API void quantize_row_iq2_s_ref (const float * GGML_RESTRICT x, block_iq2_s * GGML_RESTRICT y, int64_t k); // Dequantization +GGML_API void dequantize_row_q1_0(const block_q1_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); +GGML_API void dequantize_row_q1_0_g128(const block_q1_0_g128 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_q4_0(const block_q4_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_q4_1(const block_q4_1 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); GGML_API void dequantize_row_q5_0(const block_q5_0 * GGML_RESTRICT x, float * GGML_RESTRICT y, int64_t k); @@ -90,6 +94,8 @@ GGML_API size_t quantize_q3_K(const float * GGML_RESTRICT src, void * GGML_RESTR GGML_API size_t quantize_q4_K(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_q5_K(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_q6_K(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_q1_0(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_q1_0_g128(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_q4_0(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_q4_1(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_q5_0(const float * GGML_RESTRICT src, void * GGML_RESTRICT dst, int64_t nrows, int64_t n_per_row, const float * imatrix); diff --git a/ggml/src/ggml.c b/ggml/src/ggml.c index e9b6720c0afb..30aaff2e6052 100644 --- a/ggml/src/ggml.c +++ b/ggml/src/ggml.c @@ -651,6 +651,22 @@ static const struct ggml_type_traits type_traits[GGML_TYPE_COUNT] = { .to_float = (ggml_to_float_t) ggml_fp16_to_fp32_row, .from_float_ref = (ggml_from_float_t) ggml_fp32_to_fp16_row, }, + [GGML_TYPE_Q1_0] = { + .type_name = "q1_0", + .blck_size = QK1_0, + .type_size = sizeof(block_q1_0), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_q1_0, + .from_float_ref = (ggml_from_float_t) quantize_row_q1_0_ref, + }, + [GGML_TYPE_Q1_0_g128] = { + .type_name = "q1_0_g128", + .blck_size = QK1_0_g128, + .type_size = sizeof(block_q1_0_g128), + .is_quantized = true, + .to_float = (ggml_to_float_t) dequantize_row_q1_0_g128, + .from_float_ref = (ggml_from_float_t) quantize_row_q1_0_g128_ref, + }, [GGML_TYPE_Q4_0] = { .type_name = "q4_0", .blck_size = QK4_0, @@ -1384,6 +1400,8 @@ enum ggml_type ggml_ftype_to_ggml_type(enum ggml_ftype ftype) { case GGML_FTYPE_MOSTLY_BF16: wtype = GGML_TYPE_BF16; break; case GGML_FTYPE_MOSTLY_Q4_0: wtype = GGML_TYPE_Q4_0; break; case GGML_FTYPE_MOSTLY_Q4_1: wtype = GGML_TYPE_Q4_1; break; + case GGML_FTYPE_MOSTLY_Q1_0: wtype = GGML_TYPE_Q1_0; break; + case GGML_FTYPE_MOSTLY_Q1_0_g128: wtype = GGML_TYPE_Q1_0_g128; break; case GGML_FTYPE_MOSTLY_Q5_0: wtype = GGML_TYPE_Q5_0; break; case GGML_FTYPE_MOSTLY_Q5_1: wtype = GGML_TYPE_Q5_1; break; case GGML_FTYPE_MOSTLY_Q8_0: wtype = GGML_TYPE_Q8_0; break; @@ -7652,6 +7670,8 @@ size_t ggml_quantize_chunk( size_t result = 0; switch (type) { + case GGML_TYPE_Q1_0: result = quantize_q1_0(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; + case GGML_TYPE_Q1_0_g128: result = quantize_q1_0_g128(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q4_0: result = quantize_q4_0(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q4_1: result = quantize_q4_1(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; case GGML_TYPE_Q5_0: result = quantize_q5_0(src + start, (char *) dst + start_row * row_size, nrows, n_per_row, imatrix); break; diff --git a/gguf-py/gguf/constants.py b/gguf-py/gguf/constants.py index 3ebd9de5f6ed..18d0732c0709 100644 --- a/gguf-py/gguf/constants.py +++ b/gguf-py/gguf/constants.py @@ -3987,6 +3987,8 @@ class GGMLQuantizationType(IntEnum): TQ2_0 = 35 MXFP4 = 39 NVFP4 = 40 + Q1_0_g128 = 41 + Q1_0 = 42 class ExpertGatingFuncType(IntEnum): @@ -4040,6 +4042,8 @@ class LlamaFileType(IntEnum): MOSTLY_TQ2_0 = 37 # except 1d tensors MOSTLY_MXFP4_MOE = 38 # except 1d tensors MOSTLY_NVFP4 = 39 # except 1d tensors + MOSTLY_Q1_0_g128 = 40 # except 1d tensors + MOSTLY_Q1_0 = 41 # except 1d tensors GUESSED = 1024 # not specified in the model file @@ -4150,7 +4154,9 @@ class VisionProjectorType: GGMLQuantizationType.TQ1_0: (256, 2 + 4 * 13), GGMLQuantizationType.TQ2_0: (256, 2 + 64), GGMLQuantizationType.MXFP4: (32, 1 + 16), - GGMLQuantizationType.NVFP4: (64, 4 + 32), + GGMLQuantizationType.NVFP4: (64, 4 + 32), + GGMLQuantizationType.Q1_0: (32, 2 + 4), # 2 bytes fp16 scale + 4 bytes (32 bits) + GGMLQuantizationType.Q1_0_g128: (128, 2 + 16), # 2 bytes fp16 scale + 16 bytes (128 bits) } diff --git a/include/llama.h b/include/llama.h index a940f9d648a0..1836bc5c8eb9 100644 --- a/include/llama.h +++ b/include/llama.h @@ -154,6 +154,8 @@ extern "C" { LLAMA_FTYPE_MOSTLY_TQ2_0 = 37, // except 1d tensors LLAMA_FTYPE_MOSTLY_MXFP4_MOE = 38, // except 1d tensors LLAMA_FTYPE_MOSTLY_NVFP4 = 39, // except 1d tensors + LLAMA_FTYPE_MOSTLY_Q1_0_g128 = 40, // except 1d tensors + LLAMA_FTYPE_MOSTLY_Q1_0 = 41, // except 1d tensors LLAMA_FTYPE_GUESSED = 1024, // not specified in the model file }; diff --git a/src/llama-arch.cpp b/src/llama-arch.cpp index e210dcdae21f..aae6a0abc078 100644 --- a/src/llama-arch.cpp +++ b/src/llama-arch.cpp @@ -2800,7 +2800,7 @@ LLM_TN_IMPL::LLM_TN_IMPL(llm_arch arch, llm_tensor tensor, const char * suffix, std::string LLM_TN_IMPL::str() const { if (LLM_TENSOR_NAMES.find(tensor) == LLM_TENSOR_NAMES.end()) { - GGML_ABORT("unknown tensor name for tensor id %d", static_cast(tensor)); + return "unknown"; } if (model_tensors.find(tensor) == model_tensors.end()) { diff --git a/src/llama-model-loader.cpp b/src/llama-model-loader.cpp index 3d549cae5b6b..20194745bd3c 100644 --- a/src/llama-model-loader.cpp +++ b/src/llama-model-loader.cpp @@ -36,6 +36,8 @@ static std::string llama_model_ftype_name(llama_ftype ftype) { case LLAMA_FTYPE_ALL_F32: return "all F32"; case LLAMA_FTYPE_MOSTLY_F16: return "F16"; case LLAMA_FTYPE_MOSTLY_BF16: return "BF16"; + case LLAMA_FTYPE_MOSTLY_Q1_0: return "Q1_0"; + case LLAMA_FTYPE_MOSTLY_Q1_0_g128: return "Q1_0_g128"; case LLAMA_FTYPE_MOSTLY_Q4_0: return "Q4_0"; case LLAMA_FTYPE_MOSTLY_Q4_1: return "Q4_1"; case LLAMA_FTYPE_MOSTLY_Q5_0: return "Q5_0"; @@ -757,6 +759,8 @@ llama_model_loader::llama_model_loader( 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; + case GGML_TYPE_Q1_0: ftype = LLAMA_FTYPE_MOSTLY_Q1_0; break; + case GGML_TYPE_Q1_0_g128: ftype = LLAMA_FTYPE_MOSTLY_Q1_0_g128; break; default: { LLAMA_LOG_WARN("%s: unknown type %s\n", __func__, ggml_type_name(type_max)); diff --git a/src/llama-quant.cpp b/src/llama-quant.cpp index 322cb313f1c1..5753b41502c7 100644 --- a/src/llama-quant.cpp +++ b/src/llama-quant.cpp @@ -801,6 +801,9 @@ ggml_type llama_ftype_get_default_type(llama_ftype ftype) { case LLAMA_FTYPE_ALL_F32: return GGML_TYPE_F32; case LLAMA_FTYPE_MOSTLY_MXFP4_MOE: return GGML_TYPE_MXFP4; + case LLAMA_FTYPE_MOSTLY_NVFP4: return GGML_TYPE_NVFP4; + case LLAMA_FTYPE_MOSTLY_Q1_0: return GGML_TYPE_Q1_0; + case LLAMA_FTYPE_MOSTLY_Q1_0_g128: return GGML_TYPE_Q1_0_g128; // K-quants case LLAMA_FTYPE_MOSTLY_Q2_K_S: diff --git a/tests/test-quantize-fns.cpp b/tests/test-quantize-fns.cpp index a8fb1926231f..1e36efc84171 100644 --- a/tests/test-quantize-fns.cpp +++ b/tests/test-quantize-fns.cpp @@ -16,6 +16,7 @@ constexpr float MAX_QUANTIZATION_REFERENCE_ERROR = 0.0001f; constexpr float MAX_QUANTIZATION_TOTAL_ERROR = 0.002f; +constexpr float MAX_QUANTIZATION_TOTAL_ERROR_BINARY = 0.025f; constexpr float MAX_QUANTIZATION_TOTAL_ERROR_TERNARY = 0.01f; constexpr float MAX_QUANTIZATION_TOTAL_ERROR_2BITS = 0.0075f; constexpr float MAX_QUANTIZATION_TOTAL_ERROR_3BITS = 0.0040f; @@ -24,6 +25,7 @@ constexpr float MAX_QUANTIZATION_TOTAL_ERROR_FP4 = 0.0030f; constexpr float MAX_DOT_PRODUCT_ERROR = 0.02f; constexpr float MAX_DOT_PRODUCT_ERROR_LOWBIT = 0.04f; constexpr float MAX_DOT_PRODUCT_ERROR_FP4 = 0.03f; +constexpr float MAX_DOT_PRODUCT_ERROR_BINARY = 0.40f; constexpr float MAX_DOT_PRODUCT_ERROR_TERNARY = 0.15f; static const char* RESULT_STR[] = {"ok", "FAILED"}; @@ -145,6 +147,7 @@ int main(int argc, char * argv[]) { if (qfns_cpu->from_float && qfns->to_float) { const float total_error = total_quantization_error(qfns, qfns_cpu, test_size, test_data.data()); const float max_quantization_error = + (type == GGML_TYPE_Q1_0 || type == GGML_TYPE_Q1_0_g128) ? MAX_QUANTIZATION_TOTAL_ERROR_BINARY : type == GGML_TYPE_TQ1_0 ? MAX_QUANTIZATION_TOTAL_ERROR_TERNARY : type == GGML_TYPE_TQ2_0 ? MAX_QUANTIZATION_TOTAL_ERROR_TERNARY : type == GGML_TYPE_Q2_K ? MAX_QUANTIZATION_TOTAL_ERROR_2BITS : @@ -170,6 +173,8 @@ int main(int argc, char * argv[]) { const float max_allowed_error = type == GGML_TYPE_Q2_K || type == GGML_TYPE_IQ2_XS || type == GGML_TYPE_IQ2_XXS || type == GGML_TYPE_IQ3_XXS || type == GGML_TYPE_IQ3_S || type == GGML_TYPE_IQ2_S ? MAX_DOT_PRODUCT_ERROR_LOWBIT + : (type == GGML_TYPE_Q1_0 || type == GGML_TYPE_Q1_0_g128) + ? MAX_DOT_PRODUCT_ERROR_BINARY : type == GGML_TYPE_TQ1_0 || type == GGML_TYPE_TQ2_0 ? MAX_DOT_PRODUCT_ERROR_TERNARY : type == GGML_TYPE_NVFP4 diff --git a/tools/quantize/quantize.cpp b/tools/quantize/quantize.cpp index b727c9dd39f3..970e7f32a854 100644 --- a/tools/quantize/quantize.cpp +++ b/tools/quantize/quantize.cpp @@ -29,6 +29,8 @@ struct quant_option { }; static const std::vector QUANT_OPTIONS = { + { "Q1_0", LLAMA_FTYPE_MOSTLY_Q1_0, " ~1.5 bpw quantization", }, + { "Q1_0_g128",LLAMA_FTYPE_MOSTLY_Q1_0_g128," 1.125 bpw quantization (group size 128)", }, { "Q4_0", LLAMA_FTYPE_MOSTLY_Q4_0, " 4.34G, +0.4685 ppl @ Llama-3-8B", }, { "Q4_1", LLAMA_FTYPE_MOSTLY_Q4_1, " 4.78G, +0.4511 ppl @ Llama-3-8B", }, { "MXFP4_MOE",LLAMA_FTYPE_MOSTLY_MXFP4_MOE," MXFP4 MoE", }, diff --git a/vendor/cpp-httplib/httplib.cpp b/vendor/cpp-httplib/httplib.cpp index 8ff1da57bb52..9e81d3d6c5ed 100644 --- a/vendor/cpp-httplib/httplib.cpp +++ b/vendor/cpp-httplib/httplib.cpp @@ -1460,8 +1460,8 @@ bool mmap::open(const char *path) { auto wpath = u8string_to_wstring(path); if (wpath.empty()) { return false; } - hFile_ = ::CreateFile2(wpath.c_str(), GENERIC_READ, FILE_SHARE_READ, - OPEN_EXISTING, NULL); + hFile_ = ::CreateFileW(wpath.c_str(), GENERIC_READ, FILE_SHARE_READ, + NULL, OPEN_EXISTING, FILE_ATTRIBUTE_NORMAL, NULL); if (hFile_ == INVALID_HANDLE_VALUE) { return false; }