drop the cuda kernel: metal is the only backend built for now

The cuda variants were shipped from a build nobody was running, and each rebuild had to keep them
alive to avoid committing stale artifacts. `git log -- gguf_cuda` still has the sources, and
`vendor/` keeps ggml's cuda tree, so re-adding the backend is a `build.toml` section again.

README.md

@@ -16,6 +16,7 @@ materializing a dense copy of its weights.
 | op | signature |
 | --- | --- |
 | `dequantize` | `(blocks, ggml_type, rows, cols, dtype) -> (rows, cols)` |
+| `get_rows` | `(blocks, indices, ggml_type, cols, dtype) -> (len(indices), cols)`, unpacks as it gathers |
 | `mul_mat_vec` | `(blocks, x, ggml_type, out_features) -> (rows, out_features)` f32, fused dequant-gemv |
 
 `blocks` is a GGUF weight exactly as stored: `(out_features, bytes_per_row)` uint8. `mul_mat_vec`
@@ -26,15 +27,16 @@ matmul. The quant types each backend implements a gemv for are in `GEMV_TYPES`.
 
 | backend | torch | targets |
 | --- | --- | --- |
-| CUDA | 2.11, 2.12 | x86_64-linux, cu126/128/130/132, sm 7.5–12.0 |
 | Metal | 2.12, 2.13 | aarch64-darwin |
 
+Metal only for now: the cuda kernel is in `git log -- gguf_cuda`, not in this build.
+
 ## Where the kernels come from
 
 [llama.cpp](https://github.com/ggml-org/llama.cpp)'s ggml, vendored rather than reimplemented.
 `vendor/UPSTREAM` records the revision.
 
-CUDA compiles `ggml-cuda`'s sources; Metal compiles `ggml-metal.metal` into the embedded metallib.
+Metal compiles `ggml-metal.metal` into the embedded metallib.
 Only the files listed in `build.toml`'s `src` are built — the rest of the tree rides along so a pin
 bump cannot leave a dangling include.
 
@@ -45,6 +47,6 @@ python vendor.py --rev <llama.cpp commit>   # re-vendor, updates vendor/UPSTREAM
 nix run .#build-and-copy                    # rebuild every variant into build/
 ```
 
-`vendor.py` copies whole trees plus one rename: upstream's `mmvq.cu` lands as `mmvq-impl.cuh`,
-because `gguf_cuda/ggml_dispatch.cu` `#include`s it instead of compiling it as its own translation
-unit.
+`vendor.py` copies whole trees plus one rename: upstream's `mmvq.cu` lands as `mmvq-impl.cuh`, for a
+cuda dispatch that `#include`s it instead of compiling it as its own translation unit. `vendor/` still
+carries ggml's cuda sources, so the rename stays even with the cuda kernel gone.

build.toml

@@ -9,7 +9,9 @@
 name = "ggml-quantization"
 version = 1
 license = "MIT"
-backends = ["cuda", "metal"]
+# Metal only for now. The cuda kernel and its build section were dropped rather than kept building
+# untested -- `git log -- gguf_cuda` has them, and `vendor/` still carries ggml's cuda sources.
+backends = ["metal"]
 
 [general.hub]
 repo-id = "marcsun13/ggml-quantization"
@@ -20,68 +22,11 @@ src = [
     "torch-ext/torch_binding.h",
 ]
 
-[kernel.gguf_cuda]
-backend = "cuda"
-depends = ["torch"]
-cuda-capabilities = ["7.5", "8.0", "8.6", "8.9", "9.0", "10.0", "12.0"]
-include = ["gguf_cuda", "torch-ext", "vendor/include", "vendor/src", "vendor/src/ggml-cuda"]
-src = [
-    "gguf_cuda/gguf_cuda.cu",
-    "gguf_cuda/ggml_dispatch.cu",
-    "gguf_cuda/ggml_stubs.cu",
-    "vendor/src/ggml-cuda/convert.cu",
-    "vendor/src/ggml-cuda/quantize.cu",
-    # headers are listed so the builder stages them; only the files above are compiled
-    "vendor/include/ggml-alloc.h",
-    "vendor/include/ggml-backend.h",
-    "vendor/include/ggml-cuda.h",
-    "vendor/include/ggml.h",
-    "vendor/include/gguf.h",
-    "vendor/src/ggml-common.h",
-    "vendor/src/ggml-cuda/common.cuh",
-    "vendor/src/ggml-cuda/convert.cuh",
-    "vendor/src/ggml-cuda/dequantize.cuh",
-    "vendor/src/ggml-cuda/mma.cuh",
-    "vendor/src/ggml-cuda/mmq-config-ampere.cuh",
-    "vendor/src/ggml-cuda/mmq-config-blackwell.cuh",
-    "vendor/src/ggml-cuda/mmq-config-cdna.cuh",
-    "vendor/src/ggml-cuda/mmq-config-pascal.cuh",
-    "vendor/src/ggml-cuda/mmq-config-rdna2.cuh",
-    "vendor/src/ggml-cuda/mmq-config-rdna3-5.cuh",
-    "vendor/src/ggml-cuda/mmq-config-rdna3.cuh",
-    "vendor/src/ggml-cuda/mmq-config-rdna4.cuh",
-    "vendor/src/ggml-cuda/mmq-load-tiles.cuh",
-    "vendor/src/ggml-cuda/mmq-vec-dot.cuh",
-    "vendor/src/ggml-cuda/mmq.cuh",
-    "vendor/src/ggml-cuda/mmvf.cuh",
-    "vendor/src/ggml-cuda/mmvq.cuh",
-    "vendor/src/ggml-cuda/mmvq-impl.cuh",
-    "vendor/src/ggml-cuda/quantize.cuh",
-    "vendor/src/ggml-cuda/unary.cuh",
-    "vendor/src/ggml-cuda/vecdotq.cuh",
-    "vendor/src/ggml-cuda/vendors/cuda.h",
-    "vendor/src/ggml-impl.h",
-]
-# ggml's half/bfloat16 arithmetic needs the operators torch's build disables
-cuda-flags = [
-    "-DGGML_USE_CUDA",
-    "-DNDEBUG",
-    "-U__CUDA_NO_HALF_OPERATORS__",
-    "-U__CUDA_NO_HALF_CONVERSIONS__",
-    "-U__CUDA_NO_HALF2_OPERATORS__",
-    "-U__CUDA_NO_BFLOAT16_CONVERSIONS__",
-    "-U__CUDA_NO_BFLOAT16_OPERATORS__",
-    "-U__CUDA_NO_BFLOAT162_OPERATORS__",
-    "--expt-relaxed-constexpr",
-    "--expt-extended-lambda",
-]
-cxx-flags = ["-DGGML_USE_CUDA", "-DNDEBUG"]
-
 [kernel.gguf_metal]
 backend = "metal"
 depends = ["torch"]
 # ggml-metal.metal includes "ggml-common.h" from vendor/src, so the shader compile needs it
-# on its include path just as the cuda sources do.
+# on its include path.
 include = ["gguf_metal", "torch-ext", "vendor/include", "vendor/src", "vendor/src/ggml-metal"]
 src = [
     "gguf_metal/gguf_metal.cpp",

build/torch211-cxx11-cu126-x86_64-linux/__init__.py

@@ -1,55 +0,0 @@
-"""Compute directly on the packed blocks of a GGUF checkpoint.
-
-A GGUF weight is stored as blocks of 32 or 256 values sharing a scale. These kernels read that
-layout as it is, so a quantized model can be loaded and run without ever materializing a dense
-copy — which is the whole memory saving.
-
-Ported from llama.cpp's ggml-cuda; `vendor/UPSTREAM` pins the revision. The ops are named for what
-they do rather than for a backend: each backend registers its own implementation of the same
-schema, so calls dispatch on the tensor's device.
-"""
-
-import torch
-
-from ._ops import add_op_namespace_prefix, ops
-
-
-__all__ = ["GEMV_TYPES", "MAX_GEMV_ROWS", "dequantize", "mul_mat_vec"]
-
-# Upstream's MMVQ_MAX_BATCH_SIZE: `mul_mat_vec` has no implementation beyond this many rows, so a
-# caller with more (prefill) dequantizes and uses an ordinary matmul.
-MAX_GEMV_ROWS = 8
-
-# ggml type ids `mul_mat_vec` implements: Q4_0/Q4_1/Q5_0/Q5_1/Q8_0, the K quants, and the IQ quants.
-# `dequantize` covers more, so check this one before choosing the fused path.
-GEMV_TYPES = frozenset({2, 3, 6, 7, 8, 10, 11, 12, 13, 14, 16, 17, 18, 19, 20, 21, 22, 23, 29, 39, 40, 41, 42})
-
-
-def dequantize(
-    blocks: torch.Tensor, ggml_type: int, rows: int, cols: int, dtype: torch.dtype
-) -> torch.Tensor:
-    """`(rows, bytes_per_row)` uint8 blocks -> `(rows, cols)` values of `dtype`."""
-    return ops.dequantize(blocks, ggml_type, rows, cols, dtype)
-
-
-def mul_mat_vec(
-    blocks: torch.Tensor, x: torch.Tensor, ggml_type: int, out_features: int
-) -> torch.Tensor:
-    """Fused dequantize-gemv: `x @ blocks.T` without unpacking `blocks`.
-
-    `x` is `(rows, in_features)` and `rows` must be at most `MAX_GEMV_ROWS`. The result is f32
-    whatever `x`'s dtype was, since the kernel writes an f32 destination; cast it at the call site,
-    where the cast can be fused into whatever consumes it.
-    """
-    return ops.mul_mat_vec(blocks, x, ggml_type, out_features)
-
-
-# Without these, torch.compile cannot trace the ops and breaks the graph at every call.
-@torch.library.register_fake(add_op_namespace_prefix("dequantize"))
-def _dequantize_fake(blocks, ggml_type, rows, cols, dtype):
-    return blocks.new_empty((rows, cols), dtype=dtype)
-
-
-@torch.library.register_fake(add_op_namespace_prefix("mul_mat_vec"))
-def _mul_mat_vec_fake(blocks, x, ggml_type, out_features):
-    return x.new_empty((x.shape[0], out_features), dtype=torch.float32)

build/torch211-cxx11-cu126-x86_64-linux/_gguf_kernels_cuda_e4e2164.abi3.so

@@ -1,3 +0,0 @@
-version https://git-lfs.github.com/spec/v1
-oid sha256:4bc09a45aabd708b54f9ce83b5ec9627ed185826b64aeb2b5673f46cad4c508e
-size 24799632

build/torch211-cxx11-cu126-x86_64-linux/_ops.py

@@ -1,9 +0,0 @@
-import torch
-from . import _gguf_kernels_cuda_e4e2164
-ops = torch.ops._gguf_kernels_cuda_e4e2164
-
-def add_op_namespace_prefix(op_name: str):
-    """
-    Prefix op by namespace.
-    """
-    return f"_gguf_kernels_cuda_e4e2164::{op_name}"

build/torch211-cxx11-cu126-x86_64-linux/gguf_kernels/__init__.py

@@ -1,26 +0,0 @@
-import ctypes
-import importlib.util
-import sys
-from pathlib import Path
-from types import ModuleType
-
-
-def _import_from_path(file_path: Path) -> ModuleType:
-    # We cannot use the module name as-is, after adding it to `sys.modules`,
-    # it would also be used for other imports. So, we make a module name that
-    # depends on the path for it to be unique using the hex-encoded hash of
-    # the path.
-    path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
-    module_name = path_hash
-    spec = importlib.util.spec_from_file_location(module_name, file_path)
-    if spec is None:
-        raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
-    module = importlib.util.module_from_spec(spec)
-    if module is None:
-        raise ImportError(f"Cannot load module {module_name} from spec")
-    sys.modules[module_name] = module
-    spec.loader.exec_module(module)  # type: ignore
-    return module
-
-
-globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))

build/torch211-cxx11-cu126-x86_64-linux/metadata.json

@@ -1,17 +0,0 @@
-{
-  "name": "gguf-kernels",
-  "id": "_gguf_kernels_cuda_e4e2164",
-  "version": 1,
-  "license": "MIT",
-  "python-depends": [],
-  "backend": {
-    "type": "cuda",
-    "archs": [
-      "7.5",
-      "8.0",
-      "8.6",
-      "8.9",
-      "9.0"
-    ]
-  }
-}

build/torch211-cxx11-cu128-x86_64-linux/__init__.py

@@ -1,55 +0,0 @@
-"""Compute directly on the packed blocks of a GGUF checkpoint.
-
-A GGUF weight is stored as blocks of 32 or 256 values sharing a scale. These kernels read that
-layout as it is, so a quantized model can be loaded and run without ever materializing a dense
-copy — which is the whole memory saving.
-
-Ported from llama.cpp's ggml-cuda; `vendor/UPSTREAM` pins the revision. The ops are named for what
-they do rather than for a backend: each backend registers its own implementation of the same
-schema, so calls dispatch on the tensor's device.
-"""
-
-import torch
-
-from ._ops import add_op_namespace_prefix, ops
-
-
-__all__ = ["GEMV_TYPES", "MAX_GEMV_ROWS", "dequantize", "mul_mat_vec"]
-
-# Upstream's MMVQ_MAX_BATCH_SIZE: `mul_mat_vec` has no implementation beyond this many rows, so a
-# caller with more (prefill) dequantizes and uses an ordinary matmul.
-MAX_GEMV_ROWS = 8
-
-# ggml type ids `mul_mat_vec` implements: Q4_0/Q4_1/Q5_0/Q5_1/Q8_0, the K quants, and the IQ quants.
-# `dequantize` covers more, so check this one before choosing the fused path.
-GEMV_TYPES = frozenset({2, 3, 6, 7, 8, 10, 11, 12, 13, 14, 16, 17, 18, 19, 20, 21, 22, 23, 29, 39, 40, 41, 42})
-
-
-def dequantize(
-    blocks: torch.Tensor, ggml_type: int, rows: int, cols: int, dtype: torch.dtype
-) -> torch.Tensor:
-    """`(rows, bytes_per_row)` uint8 blocks -> `(rows, cols)` values of `dtype`."""
-    return ops.dequantize(blocks, ggml_type, rows, cols, dtype)
-
-
-def mul_mat_vec(
-    blocks: torch.Tensor, x: torch.Tensor, ggml_type: int, out_features: int
-) -> torch.Tensor:
-    """Fused dequantize-gemv: `x @ blocks.T` without unpacking `blocks`.
-
-    `x` is `(rows, in_features)` and `rows` must be at most `MAX_GEMV_ROWS`. The result is f32
-    whatever `x`'s dtype was, since the kernel writes an f32 destination; cast it at the call site,
-    where the cast can be fused into whatever consumes it.
-    """
-    return ops.mul_mat_vec(blocks, x, ggml_type, out_features)
-
-
-# Without these, torch.compile cannot trace the ops and breaks the graph at every call.
-@torch.library.register_fake(add_op_namespace_prefix("dequantize"))
-def _dequantize_fake(blocks, ggml_type, rows, cols, dtype):
-    return blocks.new_empty((rows, cols), dtype=dtype)
-
-
-@torch.library.register_fake(add_op_namespace_prefix("mul_mat_vec"))
-def _mul_mat_vec_fake(blocks, x, ggml_type, out_features):
-    return x.new_empty((x.shape[0], out_features), dtype=torch.float32)

build/torch211-cxx11-cu128-x86_64-linux/_gguf_kernels_cuda_e4e2164.abi3.so

@@ -1,3 +0,0 @@
-version https://git-lfs.github.com/spec/v1
-oid sha256:b71394062f24b9601e17ee704b313494487bbe77b8d20f54adbfd52665d9f300
-size 37124624

build/torch211-cxx11-cu128-x86_64-linux/_ops.py

@@ -1,9 +0,0 @@
-import torch
-from . import _gguf_kernels_cuda_e4e2164
-ops = torch.ops._gguf_kernels_cuda_e4e2164
-
-def add_op_namespace_prefix(op_name: str):
-    """
-    Prefix op by namespace.
-    """
-    return f"_gguf_kernels_cuda_e4e2164::{op_name}"

build/torch211-cxx11-cu128-x86_64-linux/gguf_kernels/__init__.py

@@ -1,26 +0,0 @@
-import ctypes
-import importlib.util
-import sys
-from pathlib import Path
-from types import ModuleType
-
-
-def _import_from_path(file_path: Path) -> ModuleType:
-    # We cannot use the module name as-is, after adding it to `sys.modules`,
-    # it would also be used for other imports. So, we make a module name that
-    # depends on the path for it to be unique using the hex-encoded hash of
-    # the path.
-    path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
-    module_name = path_hash
-    spec = importlib.util.spec_from_file_location(module_name, file_path)
-    if spec is None:
-        raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
-    module = importlib.util.module_from_spec(spec)
-    if module is None:
-        raise ImportError(f"Cannot load module {module_name} from spec")
-    sys.modules[module_name] = module
-    spec.loader.exec_module(module)  # type: ignore
-    return module
-
-
-globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))

build/torch211-cxx11-cu128-x86_64-linux/metadata.json

@@ -1,19 +0,0 @@
-{
-  "name": "gguf-kernels",
-  "id": "_gguf_kernels_cuda_e4e2164",
-  "version": 1,
-  "license": "MIT",
-  "python-depends": [],
-  "backend": {
-    "type": "cuda",
-    "archs": [
-      "10.0",
-      "12.0",
-      "7.5",
-      "8.0",
-      "8.6",
-      "8.9",
-      "9.0"
-    ]
-  }
-}

build/torch211-cxx11-cu130-x86_64-linux/__init__.py

@@ -1,55 +0,0 @@
-"""Compute directly on the packed blocks of a GGUF checkpoint.
-
-A GGUF weight is stored as blocks of 32 or 256 values sharing a scale. These kernels read that
-layout as it is, so a quantized model can be loaded and run without ever materializing a dense
-copy — which is the whole memory saving.
-
-Ported from llama.cpp's ggml-cuda; `vendor/UPSTREAM` pins the revision. The ops are named for what
-they do rather than for a backend: each backend registers its own implementation of the same
-schema, so calls dispatch on the tensor's device.
-"""
-
-import torch
-
-from ._ops import add_op_namespace_prefix, ops
-
-
-__all__ = ["GEMV_TYPES", "MAX_GEMV_ROWS", "dequantize", "mul_mat_vec"]
-
-# Upstream's MMVQ_MAX_BATCH_SIZE: `mul_mat_vec` has no implementation beyond this many rows, so a
-# caller with more (prefill) dequantizes and uses an ordinary matmul.
-MAX_GEMV_ROWS = 8
-
-# ggml type ids `mul_mat_vec` implements: Q4_0/Q4_1/Q5_0/Q5_1/Q8_0, the K quants, and the IQ quants.
-# `dequantize` covers more, so check this one before choosing the fused path.
-GEMV_TYPES = frozenset({2, 3, 6, 7, 8, 10, 11, 12, 13, 14, 16, 17, 18, 19, 20, 21, 22, 23, 29, 39, 40, 41, 42})
-
-
-def dequantize(
-    blocks: torch.Tensor, ggml_type: int, rows: int, cols: int, dtype: torch.dtype
-) -> torch.Tensor:
-    """`(rows, bytes_per_row)` uint8 blocks -> `(rows, cols)` values of `dtype`."""
-    return ops.dequantize(blocks, ggml_type, rows, cols, dtype)
-
-
-def mul_mat_vec(
-    blocks: torch.Tensor, x: torch.Tensor, ggml_type: int, out_features: int
-) -> torch.Tensor:
-    """Fused dequantize-gemv: `x @ blocks.T` without unpacking `blocks`.
-
-    `x` is `(rows, in_features)` and `rows` must be at most `MAX_GEMV_ROWS`. The result is f32
-    whatever `x`'s dtype was, since the kernel writes an f32 destination; cast it at the call site,
-    where the cast can be fused into whatever consumes it.
-    """
-    return ops.mul_mat_vec(blocks, x, ggml_type, out_features)
-
-
-# Without these, torch.compile cannot trace the ops and breaks the graph at every call.
-@torch.library.register_fake(add_op_namespace_prefix("dequantize"))
-def _dequantize_fake(blocks, ggml_type, rows, cols, dtype):
-    return blocks.new_empty((rows, cols), dtype=dtype)
-
-
-@torch.library.register_fake(add_op_namespace_prefix("mul_mat_vec"))
-def _mul_mat_vec_fake(blocks, x, ggml_type, out_features):
-    return x.new_empty((x.shape[0], out_features), dtype=torch.float32)

build/torch211-cxx11-cu130-x86_64-linux/_gguf_kernels_cuda_e4e2164.abi3.so

@@ -1,3 +0,0 @@
-version https://git-lfs.github.com/spec/v1
-oid sha256:5b19aec224f81561100e3133d414393779bacfabf2c703040e64b95fe0b69cef
-size 37378504

build/torch211-cxx11-cu130-x86_64-linux/_ops.py

@@ -1,9 +0,0 @@
-import torch
-from . import _gguf_kernels_cuda_e4e2164
-ops = torch.ops._gguf_kernels_cuda_e4e2164
-
-def add_op_namespace_prefix(op_name: str):
-    """
-    Prefix op by namespace.
-    """
-    return f"_gguf_kernels_cuda_e4e2164::{op_name}"

build/torch211-cxx11-cu130-x86_64-linux/gguf_kernels/__init__.py

@@ -1,26 +0,0 @@
-import ctypes
-import importlib.util
-import sys
-from pathlib import Path
-from types import ModuleType
-
-
-def _import_from_path(file_path: Path) -> ModuleType:
-    # We cannot use the module name as-is, after adding it to `sys.modules`,
-    # it would also be used for other imports. So, we make a module name that
-    # depends on the path for it to be unique using the hex-encoded hash of
-    # the path.
-    path_hash = "{:x}".format(ctypes.c_size_t(hash(file_path.absolute())).value)
-    module_name = path_hash
-    spec = importlib.util.spec_from_file_location(module_name, file_path)
-    if spec is None:
-        raise ImportError(f"Cannot load spec for {module_name} from {file_path}")
-    module = importlib.util.module_from_spec(spec)
-    if module is None:
-        raise ImportError(f"Cannot load module {module_name} from spec")
-    sys.modules[module_name] = module
-    spec.loader.exec_module(module)  # type: ignore
-    return module
-
-
-globals().update(vars(_import_from_path(Path(__file__).parent.parent / "__init__.py")))

build/torch211-cxx11-cu130-x86_64-linux/metadata.json

@@ -1,19 +0,0 @@
-{
-  "name": "gguf-kernels",
-  "id": "_gguf_kernels_cuda_e4e2164",
-  "version": 1,
-  "license": "MIT",
-  "python-depends": [],
-  "backend": {
-    "type": "cuda",
-    "archs": [
-      "10.0",
-      "12.0",
-      "7.5",
-      "8.0",
-      "8.6",
-      "8.9",
-      "9.0"
-    ]
-  }
-}

build/torch212-cxx11-cu126-x86_64-linux/__init__.py

@@ -1,62 +0,0 @@
-"""Compute directly on the packed blocks of a GGUF checkpoint.
-
-A GGUF weight is stored as blocks of 32 or 256 values sharing a scale. These kernels read that
-layout as it is, so a quantized model can be loaded and run without ever materializing a dense
-copy — which is the whole memory saving.
-
-Ported from llama.cpp's ggml-cuda; `vendor/UPSTREAM` pins the revision. The ops are named for what
-they do rather than for a backend: each backend registers its own implementation of the same
-schema, so calls dispatch on the tensor's device.
-"""
-
-import torch
-
-from ._ops import add_op_namespace_prefix, ops
-
-
-__all__ = ["GEMV_TYPES", "MAX_GEMV_ROWS", "dequantize", "mul_mat_vec"]
-
-# Upstream's MMVQ_MAX_BATCH_SIZE: `mul_mat_vec` has no implementation beyond this many rows, so a
-# caller with more (prefill) dequantizes and uses an ordinary matmul.
-MAX_GEMV_ROWS = 8
-
-# ggml type ids `mul_mat_vec` implements, as this build reports them -- it is backend-specific, and
-# a type routed into a gemv that has no kernel for it is a fault, not a fallback. CUDA covers
-# Q4_0/Q4_1/Q5_0/Q5_1/Q8_0, the K quants, the IQ quants, MXFP4/NVFP4 and Q1_0/Q2_0; Metal is the
-# same minus NVFP4/Q1_0/Q2_0, which its metallib has no kernels for. `dequantize` covers more, so
-# check this before choosing the fused path.
-try:
-    GEMV_TYPES = frozenset(ops.gemv_types())
-except AttributeError:
-    # A build published before `gemv_types` existed; those are CUDA-only, so its list is theirs.
-    GEMV_TYPES = frozenset({2, 3, 6, 7, 8, 10, 11, 12, 13, 14, 16, 17, 18, 19, 20, 21, 22, 23, 29, 39, 40, 41, 42})
-
-
-def dequantize(
-    blocks: torch.Tensor, ggml_type: int, rows: int, cols: int, dtype: torch.dtype
-) -> torch.Tensor:
-    """`(rows, bytes_per_row)` uint8 blocks -> `(rows, cols)` values of `dtype`."""
-    return ops.dequantize(blocks, ggml_type, rows, cols, dtype)
-
-
-def mul_mat_vec(
-    blocks: torch.Tensor, x: torch.Tensor, ggml_type: int, out_features: int
-) -> torch.Tensor:
-    """Fused dequantize-gemv: `x @ blocks.T` without unpacking `blocks`.
-
-    `x` is `(rows, in_features)` and `rows` must be at most `MAX_GEMV_ROWS`. The result is f32
-    whatever `x`'s dtype was, since the kernel writes an f32 destination; cast it at the call site,
-    where the cast can be fused into whatever consumes it.
-    """
-    return ops.mul_mat_vec(blocks, x, ggml_type, out_features)
-
-
-# Without these, torch.compile cannot trace the ops and breaks the graph at every call.
-@torch.library.register_fake(add_op_namespace_prefix("dequantize"))
-def _dequantize_fake(blocks, ggml_type, rows, cols, dtype):
-    return blocks.new_empty((rows, cols), dtype=dtype)
-
-
-@torch.library.register_fake(add_op_namespace_prefix("mul_mat_vec"))
-def _mul_mat_vec_fake(blocks, x, ggml_type, out_features):
-    return x.new_empty((x.shape[0], out_features), dtype=torch.float32)

build/torch212-cxx11-cu126-x86_64-linux/_gguf_kernels_cuda_1df3b99.abi3.so

@@ -1,3 +0,0 @@
-version https://git-lfs.github.com/spec/v1
-oid sha256:a0dcdaf3e54554158476036dad0d2e6412aa42d4d1bfefa8d68e019e841283da
-size 24822056

build/torch212-cxx11-cu126-x86_64-linux/_ops.py

@@ -1,9 +0,0 @@
-import torch
-from . import _gguf_kernels_cuda_1df3b99
-ops = torch.ops._gguf_kernels_cuda_1df3b99
-
-def add_op_namespace_prefix(op_name: str):
-    """
-    Prefix op by namespace.
-    """
-    return f"_gguf_kernels_cuda_1df3b99::{op_name}"

build/torch212-cxx11-cu126-x86_64-linux/metadata.json

@@ -1,36 +0,0 @@
-{
-  "name": "gguf-kernels",
-  "id": "_gguf_kernels_cuda_1df3b99",
-  "version": 1,
-  "license": "MIT",
-  "python-depends": [],
-  "backend": {
-    "type": "cuda",
-    "archs": [
-      "7.5",
-      "8.0",
-      "8.6",
-      "8.9",
-      "9.0"
-    ]
-  },
-  "digest": {
-    "algorithm": "sha256",
-    "files": {
-      "__init__.py": "GoFcHJbbJ9C8b40XNznMxsTcC+CV9kaFE8a1DP763kU=",
-      "_gguf_kernels_cuda_1df3b99.abi3.so": "oNza8+VFVBWEdgNtrQ0uZBKqQtTRv++o1o4BnoQSg9o=",
-      "_ops.py": "dRIZkil9x8IS6KhIwXlRzO6qzQoGYUQkGL87Z3gpEzY="
-    }
-  },
-  "provenance": {
-    "kernel-builder": {
-      "version": "0.17.0-dev0",
-      "sha": "81f55ea30fd8f819dcf93a3c934dd584c895bd2f",
-      "dirty": false
-    },
-    "kernel": {
-      "sha": "1df3b999686b77ef298ae13d066a39958503164c",
-      "dirty": false
-    }
-  }
-}

build/torch212-cxx11-cu130-x86_64-linux/__init__.py

@@ -1,62 +0,0 @@
-"""Compute directly on the packed blocks of a GGUF checkpoint.
-
-A GGUF weight is stored as blocks of 32 or 256 values sharing a scale. These kernels read that
-layout as it is, so a quantized model can be loaded and run without ever materializing a dense
-copy — which is the whole memory saving.
-
-Ported from llama.cpp's ggml-cuda; `vendor/UPSTREAM` pins the revision. The ops are named for what
-they do rather than for a backend: each backend registers its own implementation of the same
-schema, so calls dispatch on the tensor's device.
-"""
-
-import torch
-
-from ._ops import add_op_namespace_prefix, ops
-
-
-__all__ = ["GEMV_TYPES", "MAX_GEMV_ROWS", "dequantize", "mul_mat_vec"]
-
-# Upstream's MMVQ_MAX_BATCH_SIZE: `mul_mat_vec` has no implementation beyond this many rows, so a
-# caller with more (prefill) dequantizes and uses an ordinary matmul.
-MAX_GEMV_ROWS = 8
-
-# ggml type ids `mul_mat_vec` implements, as this build reports them -- it is backend-specific, and
-# a type routed into a gemv that has no kernel for it is a fault, not a fallback. CUDA covers
-# Q4_0/Q4_1/Q5_0/Q5_1/Q8_0, the K quants, the IQ quants, MXFP4/NVFP4 and Q1_0/Q2_0; Metal is the
-# same minus NVFP4/Q1_0/Q2_0, which its metallib has no kernels for. `dequantize` covers more, so
-# check this before choosing the fused path.
-try:
-    GEMV_TYPES = frozenset(ops.gemv_types())
-except AttributeError:
-    # A build published before `gemv_types` existed; those are CUDA-only, so its list is theirs.
-    GEMV_TYPES = frozenset({2, 3, 6, 7, 8, 10, 11, 12, 13, 14, 16, 17, 18, 19, 20, 21, 22, 23, 29, 39, 40, 41, 42})
-
-
-def dequantize(
-    blocks: torch.Tensor, ggml_type: int, rows: int, cols: int, dtype: torch.dtype
-) -> torch.Tensor:
-    """`(rows, bytes_per_row)` uint8 blocks -> `(rows, cols)` values of `dtype`."""
-    return ops.dequantize(blocks, ggml_type, rows, cols, dtype)
-
-
-def mul_mat_vec(
-    blocks: torch.Tensor, x: torch.Tensor, ggml_type: int, out_features: int
-) -> torch.Tensor:
-    """Fused dequantize-gemv: `x @ blocks.T` without unpacking `blocks`.
-
-    `x` is `(rows, in_features)` and `rows` must be at most `MAX_GEMV_ROWS`. The result is f32
-    whatever `x`'s dtype was, since the kernel writes an f32 destination; cast it at the call site,
-    where the cast can be fused into whatever consumes it.
-    """
-    return ops.mul_mat_vec(blocks, x, ggml_type, out_features)
-
-
-# Without these, torch.compile cannot trace the ops and breaks the graph at every call.
-@torch.library.register_fake(add_op_namespace_prefix("dequantize"))
-def _dequantize_fake(blocks, ggml_type, rows, cols, dtype):
-    return blocks.new_empty((rows, cols), dtype=dtype)
-
-
-@torch.library.register_fake(add_op_namespace_prefix("mul_mat_vec"))
-def _mul_mat_vec_fake(blocks, x, ggml_type, out_features):
-    return x.new_empty((x.shape[0], out_features), dtype=torch.float32)

build/torch212-cxx11-cu130-x86_64-linux/_gguf_kernels_cuda_1df3b99.abi3.so

@@ -1,3 +0,0 @@
-version https://git-lfs.github.com/spec/v1
-oid sha256:7355d69df74e88b5163248a7ada198f95c5310b2b182268e951ee7da67f4a537
-size 37396976

build/torch212-cxx11-cu130-x86_64-linux/_ops.py

@@ -1,9 +0,0 @@
-import torch
-from . import _gguf_kernels_cuda_1df3b99
-ops = torch.ops._gguf_kernels_cuda_1df3b99
-
-def add_op_namespace_prefix(op_name: str):
-    """
-    Prefix op by namespace.
-    """
-    return f"_gguf_kernels_cuda_1df3b99::{op_name}"

build/torch212-cxx11-cu130-x86_64-linux/metadata.json

@@ -1,38 +0,0 @@
-{
-  "name": "gguf-kernels",
-  "id": "_gguf_kernels_cuda_1df3b99",
-  "version": 1,
-  "license": "MIT",
-  "python-depends": [],
-  "backend": {
-    "type": "cuda",
-    "archs": [
-      "10.0",
-      "12.0",
-      "7.5",
-      "8.0",
-      "8.6",
-      "8.9",
-      "9.0"
-    ]
-  },
-  "digest": {
-    "algorithm": "sha256",
-    "files": {
-      "__init__.py": "GoFcHJbbJ9C8b40XNznMxsTcC+CV9kaFE8a1DP763kU=",
-      "_gguf_kernels_cuda_1df3b99.abi3.so": "c1XWnfdOiLUWMkinraGY+VxTELKxgiaOlR7n2mf0pTc=",
-      "_ops.py": "dRIZkil9x8IS6KhIwXlRzO6qzQoGYUQkGL87Z3gpEzY="
-    }
-  },
-  "provenance": {
-    "kernel-builder": {
-      "version": "0.17.0-dev0",
-      "sha": "81f55ea30fd8f819dcf93a3c934dd584c895bd2f",
-      "dirty": false
-    },
-    "kernel": {
-      "sha": "1df3b999686b77ef298ae13d066a39958503164c",
-      "dirty": false
-    }
-  }
-}

build/torch212-cxx11-cu132-x86_64-linux/__init__.py

@@ -1,62 +0,0 @@
-"""Compute directly on the packed blocks of a GGUF checkpoint.
-
-A GGUF weight is stored as blocks of 32 or 256 values sharing a scale. These kernels read that
-layout as it is, so a quantized model can be loaded and run without ever materializing a dense
-copy — which is the whole memory saving.
-
-Ported from llama.cpp's ggml-cuda; `vendor/UPSTREAM` pins the revision. The ops are named for what
-they do rather than for a backend: each backend registers its own implementation of the same
-schema, so calls dispatch on the tensor's device.
-"""
-
-import torch
-
-from ._ops import add_op_namespace_prefix, ops
-
-
-__all__ = ["GEMV_TYPES", "MAX_GEMV_ROWS", "dequantize", "mul_mat_vec"]
-
-# Upstream's MMVQ_MAX_BATCH_SIZE: `mul_mat_vec` has no implementation beyond this many rows, so a
-# caller with more (prefill) dequantizes and uses an ordinary matmul.
-MAX_GEMV_ROWS = 8
-
-# ggml type ids `mul_mat_vec` implements, as this build reports them -- it is backend-specific, and
-# a type routed into a gemv that has no kernel for it is a fault, not a fallback. CUDA covers
-# Q4_0/Q4_1/Q5_0/Q5_1/Q8_0, the K quants, the IQ quants, MXFP4/NVFP4 and Q1_0/Q2_0; Metal is the
-# same minus NVFP4/Q1_0/Q2_0, which its metallib has no kernels for. `dequantize` covers more, so
-# check this before choosing the fused path.
-try:
-    GEMV_TYPES = frozenset(ops.gemv_types())
-except AttributeError:
-    # A build published before `gemv_types` existed; those are CUDA-only, so its list is theirs.
-    GEMV_TYPES = frozenset({2, 3, 6, 7, 8, 10, 11, 12, 13, 14, 16, 17, 18, 19, 20, 21, 22, 23, 29, 39, 40, 41, 42})
-
-
-def dequantize(
-    blocks: torch.Tensor, ggml_type: int, rows: int, cols: int, dtype: torch.dtype
-) -> torch.Tensor:
-    """`(rows, bytes_per_row)` uint8 blocks -> `(rows, cols)` values of `dtype`."""
-    return ops.dequantize(blocks, ggml_type, rows, cols, dtype)
-
-
-def mul_mat_vec(
-    blocks: torch.Tensor, x: torch.Tensor, ggml_type: int, out_features: int
-) -> torch.Tensor:
-    """Fused dequantize-gemv: `x @ blocks.T` without unpacking `blocks`.
-
-    `x` is `(rows, in_features)` and `rows` must be at most `MAX_GEMV_ROWS`. The result is f32
-    whatever `x`'s dtype was, since the kernel writes an f32 destination; cast it at the call site,
-    where the cast can be fused into whatever consumes it.
-    """
-    return ops.mul_mat_vec(blocks, x, ggml_type, out_features)
-
-
-# Without these, torch.compile cannot trace the ops and breaks the graph at every call.
-@torch.library.register_fake(add_op_namespace_prefix("dequantize"))
-def _dequantize_fake(blocks, ggml_type, rows, cols, dtype):
-    return blocks.new_empty((rows, cols), dtype=dtype)
-
-
-@torch.library.register_fake(add_op_namespace_prefix("mul_mat_vec"))
-def _mul_mat_vec_fake(blocks, x, ggml_type, out_features):
-    return x.new_empty((x.shape[0], out_features), dtype=torch.float32)

build/torch212-cxx11-cu132-x86_64-linux/_gguf_kernels_cuda_1df3b99.abi3.so

@@ -1,3 +0,0 @@
-version https://git-lfs.github.com/spec/v1
-oid sha256:50a408d91abc14355e1b1666ac6aa591f61271f8333dc863f1e9395822f0b1a1
-size 37415488

build/torch212-cxx11-cu132-x86_64-linux/_ops.py

@@ -1,9 +0,0 @@
-import torch
-from . import _gguf_kernels_cuda_1df3b99
-ops = torch.ops._gguf_kernels_cuda_1df3b99
-
-def add_op_namespace_prefix(op_name: str):
-    """
-    Prefix op by namespace.
-    """
-    return f"_gguf_kernels_cuda_1df3b99::{op_name}"

build/torch212-cxx11-cu132-x86_64-linux/metadata.json

@@ -1,38 +0,0 @@
-{
-  "name": "gguf-kernels",
-  "id": "_gguf_kernels_cuda_1df3b99",
-  "version": 1,
-  "license": "MIT",
-  "python-depends": [],
-  "backend": {
-    "type": "cuda",
-    "archs": [
-      "10.0",
-      "12.0",
-      "7.5",
-      "8.0",
-      "8.6",
-      "8.9",
-      "9.0"
-    ]
-  },
-  "digest": {
-    "algorithm": "sha256",
-    "files": {
-      "__init__.py": "GoFcHJbbJ9C8b40XNznMxsTcC+CV9kaFE8a1DP763kU=",
-      "_gguf_kernels_cuda_1df3b99.abi3.so": "UKQI2Rq8FDVeGxZmrGqlkfYScfgzPchj8ek5WCLwsaE=",
-      "_ops.py": "dRIZkil9x8IS6KhIwXlRzO6qzQoGYUQkGL87Z3gpEzY="
-    }
-  },
-  "provenance": {
-    "kernel-builder": {
-      "version": "0.17.0-dev0",
-      "sha": "81f55ea30fd8f819dcf93a3c934dd584c895bd2f",
-      "dirty": false
-    },
-    "kernel": {
-      "sha": "1df3b999686b77ef298ae13d066a39958503164c",
-      "dirty": false
-    }
-  }
-}

build/torch213-cxx11-cu126-x86_64-linux/__init__.py

@@ -1,62 +0,0 @@
-"""Compute directly on the packed blocks of a GGUF checkpoint.
-
-A GGUF weight is stored as blocks of 32 or 256 values sharing a scale. These kernels read that
-layout as it is, so a quantized model can be loaded and run without ever materializing a dense
-copy — which is the whole memory saving.
-
-Ported from llama.cpp's ggml-cuda; `vendor/UPSTREAM` pins the revision. The ops are named for what
-they do rather than for a backend: each backend registers its own implementation of the same
-schema, so calls dispatch on the tensor's device.
-"""
-
-import torch
-
-from ._ops import add_op_namespace_prefix, ops
-
-
-__all__ = ["GEMV_TYPES", "MAX_GEMV_ROWS", "dequantize", "mul_mat_vec"]
-
-# Upstream's MMVQ_MAX_BATCH_SIZE: `mul_mat_vec` has no implementation beyond this many rows, so a
-# caller with more (prefill) dequantizes and uses an ordinary matmul.
-MAX_GEMV_ROWS = 8
-
-# ggml type ids `mul_mat_vec` implements, as this build reports them -- it is backend-specific, and
-# a type routed into a gemv that has no kernel for it is a fault, not a fallback. CUDA covers
-# Q4_0/Q4_1/Q5_0/Q5_1/Q8_0, the K quants, the IQ quants, MXFP4/NVFP4 and Q1_0/Q2_0; Metal is the
-# same minus NVFP4/Q1_0/Q2_0, which its metallib has no kernels for. `dequantize` covers more, so
-# check this before choosing the fused path.
-try:
-    GEMV_TYPES = frozenset(ops.gemv_types())
-except AttributeError:
-    # A build published before `gemv_types` existed; those are CUDA-only, so its list is theirs.
-    GEMV_TYPES = frozenset({2, 3, 6, 7, 8, 10, 11, 12, 13, 14, 16, 17, 18, 19, 20, 21, 22, 23, 29, 39, 40, 41, 42})
-
-
-def dequantize(
-    blocks: torch.Tensor, ggml_type: int, rows: int, cols: int, dtype: torch.dtype
-) -> torch.Tensor:
-    """`(rows, bytes_per_row)` uint8 blocks -> `(rows, cols)` values of `dtype`."""
-    return ops.dequantize(blocks, ggml_type, rows, cols, dtype)
-
-
-def mul_mat_vec(
-    blocks: torch.Tensor, x: torch.Tensor, ggml_type: int, out_features: int
-) -> torch.Tensor:
-    """Fused dequantize-gemv: `x @ blocks.T` without unpacking `blocks`.
-
-    `x` is `(rows, in_features)` and `rows` must be at most `MAX_GEMV_ROWS`. The result is f32
-    whatever `x`'s dtype was, since the kernel writes an f32 destination; cast it at the call site,
-    where the cast can be fused into whatever consumes it.
-    """
-    return ops.mul_mat_vec(blocks, x, ggml_type, out_features)
-
-
-# Without these, torch.compile cannot trace the ops and breaks the graph at every call.
-@torch.library.register_fake(add_op_namespace_prefix("dequantize"))
-def _dequantize_fake(blocks, ggml_type, rows, cols, dtype):
-    return blocks.new_empty((rows, cols), dtype=dtype)
-
-
-@torch.library.register_fake(add_op_namespace_prefix("mul_mat_vec"))
-def _mul_mat_vec_fake(blocks, x, ggml_type, out_features):
-    return x.new_empty((x.shape[0], out_features), dtype=torch.float32)

build/torch213-cxx11-cu126-x86_64-linux/_gguf_kernels_cuda_1df3b99.abi3.so

@@ -1,3 +0,0 @@
-version https://git-lfs.github.com/spec/v1
-oid sha256:9a7b5651497660db8a6b4cbe857610fc5ac448d52c36ad44b16c6e199d0d09e6
-size 24821896

build/torch213-cxx11-cu126-x86_64-linux/_ops.py

@@ -1,9 +0,0 @@
-import torch
-from . import _gguf_kernels_cuda_1df3b99
-ops = torch.ops._gguf_kernels_cuda_1df3b99
-
-def add_op_namespace_prefix(op_name: str):
-    """
-    Prefix op by namespace.
-    """
-    return f"_gguf_kernels_cuda_1df3b99::{op_name}"

build/torch213-cxx11-cu126-x86_64-linux/metadata.json

@@ -1,36 +0,0 @@
-{
-  "name": "gguf-kernels",
-  "id": "_gguf_kernels_cuda_1df3b99",
-  "version": 1,
-  "license": "MIT",
-  "python-depends": [],
-  "backend": {
-    "type": "cuda",
-    "archs": [
-      "7.5",
-      "8.0",
-      "8.6",
-      "8.9",
-      "9.0"
-    ]
-  },
-  "digest": {
-    "algorithm": "sha256",
-    "files": {
-      "__init__.py": "GoFcHJbbJ9C8b40XNznMxsTcC+CV9kaFE8a1DP763kU=",
-      "_gguf_kernels_cuda_1df3b99.abi3.so": "mntWUUl2YNuKa0y+hXYQ/FrESNUsNq1EsWxuGZ0NCeY=",
-      "_ops.py": "dRIZkil9x8IS6KhIwXlRzO6qzQoGYUQkGL87Z3gpEzY="
-    }
-  },
-  "provenance": {
-    "kernel-builder": {
-      "version": "0.17.0-dev0",
-      "sha": "81f55ea30fd8f819dcf93a3c934dd584c895bd2f",
-      "dirty": false
-    },
-    "kernel": {
-      "sha": "1df3b999686b77ef298ae13d066a39958503164c",
-      "dirty": false
-    }
-  }
-}

build/torch213-cxx11-cu130-x86_64-linux/__init__.py

@@ -1,62 +0,0 @@
-"""Compute directly on the packed blocks of a GGUF checkpoint.
-
-A GGUF weight is stored as blocks of 32 or 256 values sharing a scale. These kernels read that
-layout as it is, so a quantized model can be loaded and run without ever materializing a dense
-copy — which is the whole memory saving.
-
-Ported from llama.cpp's ggml-cuda; `vendor/UPSTREAM` pins the revision. The ops are named for what
-they do rather than for a backend: each backend registers its own implementation of the same
-schema, so calls dispatch on the tensor's device.
-"""
-
-import torch
-
-from ._ops import add_op_namespace_prefix, ops
-
-
-__all__ = ["GEMV_TYPES", "MAX_GEMV_ROWS", "dequantize", "mul_mat_vec"]
-
-# Upstream's MMVQ_MAX_BATCH_SIZE: `mul_mat_vec` has no implementation beyond this many rows, so a
-# caller with more (prefill) dequantizes and uses an ordinary matmul.
-MAX_GEMV_ROWS = 8
-
-# ggml type ids `mul_mat_vec` implements, as this build reports them -- it is backend-specific, and
-# a type routed into a gemv that has no kernel for it is a fault, not a fallback. CUDA covers
-# Q4_0/Q4_1/Q5_0/Q5_1/Q8_0, the K quants, the IQ quants, MXFP4/NVFP4 and Q1_0/Q2_0; Metal is the
-# same minus NVFP4/Q1_0/Q2_0, which its metallib has no kernels for. `dequantize` covers more, so
-# check this before choosing the fused path.
-try:
-    GEMV_TYPES = frozenset(ops.gemv_types())
-except AttributeError:
-    # A build published before `gemv_types` existed; those are CUDA-only, so its list is theirs.
-    GEMV_TYPES = frozenset({2, 3, 6, 7, 8, 10, 11, 12, 13, 14, 16, 17, 18, 19, 20, 21, 22, 23, 29, 39, 40, 41, 42})
-
-
-def dequantize(
-    blocks: torch.Tensor, ggml_type: int, rows: int, cols: int, dtype: torch.dtype
-) -> torch.Tensor:
-    """`(rows, bytes_per_row)` uint8 blocks -> `(rows, cols)` values of `dtype`."""
-    return ops.dequantize(blocks, ggml_type, rows, cols, dtype)
-
-
-def mul_mat_vec(
-    blocks: torch.Tensor, x: torch.Tensor, ggml_type: int, out_features: int
-) -> torch.Tensor:
-    """Fused dequantize-gemv: `x @ blocks.T` without unpacking `blocks`.
-
-    `x` is `(rows, in_features)` and `rows` must be at most `MAX_GEMV_ROWS`. The result is f32
-    whatever `x`'s dtype was, since the kernel writes an f32 destination; cast it at the call site,
-    where the cast can be fused into whatever consumes it.
-    """
-    return ops.mul_mat_vec(blocks, x, ggml_type, out_features)
-
-
-# Without these, torch.compile cannot trace the ops and breaks the graph at every call.
-@torch.library.register_fake(add_op_namespace_prefix("dequantize"))
-def _dequantize_fake(blocks, ggml_type, rows, cols, dtype):
-    return blocks.new_empty((rows, cols), dtype=dtype)
-
-
-@torch.library.register_fake(add_op_namespace_prefix("mul_mat_vec"))
-def _mul_mat_vec_fake(blocks, x, ggml_type, out_features):
-    return x.new_empty((x.shape[0], out_features), dtype=torch.float32)

build/torch213-cxx11-cu130-x86_64-linux/_gguf_kernels_cuda_1df3b99.abi3.so

@@ -1,3 +0,0 @@
-version https://git-lfs.github.com/spec/v1
-oid sha256:aa7b02e5b0e1976c3d244e697710bf692943a598e9a56fee5c475cbd0d63015c
-size 37396824

build/torch213-cxx11-cu130-x86_64-linux/_ops.py

@@ -1,9 +0,0 @@
-import torch
-from . import _gguf_kernels_cuda_1df3b99
-ops = torch.ops._gguf_kernels_cuda_1df3b99
-
-def add_op_namespace_prefix(op_name: str):
-    """
-    Prefix op by namespace.
-    """
-    return f"_gguf_kernels_cuda_1df3b99::{op_name}"

build/torch213-cxx11-cu130-x86_64-linux/metadata.json

@@ -1,38 +0,0 @@
-{
-  "name": "gguf-kernels",
-  "id": "_gguf_kernels_cuda_1df3b99",
-  "version": 1,
-  "license": "MIT",
-  "python-depends": [],
-  "backend": {
-    "type": "cuda",
-    "archs": [
-      "10.0",
-      "12.0",
-      "7.5",
-      "8.0",
-      "8.6",
-      "8.9",
-      "9.0"
-    ]
-  },
-  "digest": {
-    "algorithm": "sha256",
-    "files": {
-      "__init__.py": "GoFcHJbbJ9C8b40XNznMxsTcC+CV9kaFE8a1DP763kU=",
-      "_gguf_kernels_cuda_1df3b99.abi3.so": "qnsC5bDhl2w9JE5pdxC/aSlDpZjppW/uXEdcvQ1jAVw=",
-      "_ops.py": "dRIZkil9x8IS6KhIwXlRzO6qzQoGYUQkGL87Z3gpEzY="
-    }
-  },
-  "provenance": {
-    "kernel-builder": {
-      "version": "0.17.0-dev0",
-      "sha": "81f55ea30fd8f819dcf93a3c934dd584c895bd2f",
-      "dirty": false
-    },
-    "kernel": {
-      "sha": "1df3b999686b77ef298ae13d066a39958503164c",
-      "dirty": false
-    }
-  }
-}

build/torch213-cxx11-cu132-x86_64-linux/__init__.py

@@ -1,62 +0,0 @@
-"""Compute directly on the packed blocks of a GGUF checkpoint.
-
-A GGUF weight is stored as blocks of 32 or 256 values sharing a scale. These kernels read that
-layout as it is, so a quantized model can be loaded and run without ever materializing a dense
-copy — which is the whole memory saving.
-
-Ported from llama.cpp's ggml-cuda; `vendor/UPSTREAM` pins the revision. The ops are named for what
-they do rather than for a backend: each backend registers its own implementation of the same
-schema, so calls dispatch on the tensor's device.
-"""
-
-import torch
-
-from ._ops import add_op_namespace_prefix, ops
-
-
-__all__ = ["GEMV_TYPES", "MAX_GEMV_ROWS", "dequantize", "mul_mat_vec"]
-
-# Upstream's MMVQ_MAX_BATCH_SIZE: `mul_mat_vec` has no implementation beyond this many rows, so a
-# caller with more (prefill) dequantizes and uses an ordinary matmul.
-MAX_GEMV_ROWS = 8
-
-# ggml type ids `mul_mat_vec` implements, as this build reports them -- it is backend-specific, and
-# a type routed into a gemv that has no kernel for it is a fault, not a fallback. CUDA covers
-# Q4_0/Q4_1/Q5_0/Q5_1/Q8_0, the K quants, the IQ quants, MXFP4/NVFP4 and Q1_0/Q2_0; Metal is the
-# same minus NVFP4/Q1_0/Q2_0, which its metallib has no kernels for. `dequantize` covers more, so
-# check this before choosing the fused path.
-try:
-    GEMV_TYPES = frozenset(ops.gemv_types())
-except AttributeError:
-    # A build published before `gemv_types` existed; those are CUDA-only, so its list is theirs.
-    GEMV_TYPES = frozenset({2, 3, 6, 7, 8, 10, 11, 12, 13, 14, 16, 17, 18, 19, 20, 21, 22, 23, 29, 39, 40, 41, 42})
-
-
-def dequantize(
-    blocks: torch.Tensor, ggml_type: int, rows: int, cols: int, dtype: torch.dtype
-) -> torch.Tensor:
-    """`(rows, bytes_per_row)` uint8 blocks -> `(rows, cols)` values of `dtype`."""
-    return ops.dequantize(blocks, ggml_type, rows, cols, dtype)
-
-
-def mul_mat_vec(
-    blocks: torch.Tensor, x: torch.Tensor, ggml_type: int, out_features: int
-) -> torch.Tensor:
-    """Fused dequantize-gemv: `x @ blocks.T` without unpacking `blocks`.
-
-    `x` is `(rows, in_features)` and `rows` must be at most `MAX_GEMV_ROWS`. The result is f32
-    whatever `x`'s dtype was, since the kernel writes an f32 destination; cast it at the call site,
-    where the cast can be fused into whatever consumes it.
-    """
-    return ops.mul_mat_vec(blocks, x, ggml_type, out_features)
-
-
-# Without these, torch.compile cannot trace the ops and breaks the graph at every call.
-@torch.library.register_fake(add_op_namespace_prefix("dequantize"))
-def _dequantize_fake(blocks, ggml_type, rows, cols, dtype):
-    return blocks.new_empty((rows, cols), dtype=dtype)
-
-
-@torch.library.register_fake(add_op_namespace_prefix("mul_mat_vec"))
-def _mul_mat_vec_fake(blocks, x, ggml_type, out_features):
-    return x.new_empty((x.shape[0], out_features), dtype=torch.float32)

build/torch213-cxx11-cu132-x86_64-linux/_gguf_kernels_cuda_1df3b99.abi3.so

@@ -1,3 +0,0 @@
-version https://git-lfs.github.com/spec/v1
-oid sha256:732504ad19d3eb0692fca0a1085cd27eeb838c068c6f27b657d9b8224504f629
-size 37415328

build/torch213-cxx11-cu132-x86_64-linux/_ops.py

@@ -1,9 +0,0 @@
-import torch
-from . import _gguf_kernels_cuda_1df3b99
-ops = torch.ops._gguf_kernels_cuda_1df3b99
-
-def add_op_namespace_prefix(op_name: str):
-    """
-    Prefix op by namespace.
-    """
-    return f"_gguf_kernels_cuda_1df3b99::{op_name}"

build/torch213-cxx11-cu132-x86_64-linux/metadata.json

@@ -1,38 +0,0 @@
-{
-  "name": "gguf-kernels",
-  "id": "_gguf_kernels_cuda_1df3b99",
-  "version": 1,
-  "license": "MIT",
-  "python-depends": [],
-  "backend": {
-    "type": "cuda",
-    "archs": [
-      "10.0",
-      "12.0",
-      "7.5",
-      "8.0",
-      "8.6",
-      "8.9",
-      "9.0"
-    ]
-  },
-  "digest": {
-    "algorithm": "sha256",
-    "files": {
-      "__init__.py": "GoFcHJbbJ9C8b40XNznMxsTcC+CV9kaFE8a1DP763kU=",
-      "_gguf_kernels_cuda_1df3b99.abi3.so": "cyUErRnT6waS/KChCFzSfuuDjAaMbye2V9m4IkUE9ik=",
-      "_ops.py": "dRIZkil9x8IS6KhIwXlRzO6qzQoGYUQkGL87Z3gpEzY="
-    }
-  },
-  "provenance": {
-    "kernel-builder": {
-      "version": "0.17.0-dev0",
-      "sha": "81f55ea30fd8f819dcf93a3c934dd584c895bd2f",
-      "dirty": false
-    },
-    "kernel": {
-      "sha": "1df3b999686b77ef298ae13d066a39958503164c",
-      "dirty": false
-    }
-  }
-}

gguf_cuda/ggml_dispatch.cu

@@ -1,157 +0,0 @@
-/* Thin raw-pointer wrappers around llama.cpp's ggml-cuda kernels.
- *
- * This file is the only place that sees ggml headers; the torch bindings call these
- * extern "C" entry points, so torch and ggml headers never meet. Upstream sources are
- * used unmodified: mmvq.cu is #include'd because its raw-pointer dispatch
- * (mul_mat_vec_q_switch_type) is static, and patching upstream would make syncing harder.
- *
- * Scratch buffers (q8_1 activations) are allocated by the caller so they come from torch's
- * caching allocator and stay capturable in a CUDA graph.
- */
-
-#include "common.cuh"
-#include "convert.cuh"
-#include "quantize.cuh"
-
-#include "mmvq-impl.cuh"  // upstream mmvq.cu, renamed by vendor.py: #include’d, never its own TU
-
-#include <cstdint>
-
-enum gguf_dtype { GGUF_DTYPE_F32 = 0, GGUF_DTYPE_F16 = 1, GGUF_DTYPE_BF16 = 2 };
-
-// ---------------------------------------------------------------------------
-// q8_1 activation quantization, templated on the input dtype.
-//
-// Upstream's quantize_row_q8_1_cuda only accepts f32 (llama.cpp's graph always feeds f32
-// activations). Ours are bf16, and at M=1 a separate .to(float) pass costs more than the
-// matmul itself — everything here is launch-bound, so the win is in doing fewer kernels.
-// Same math and same block_q8_1 layout as upstream, restricted to the contiguous 2D case.
-// ---------------------------------------------------------------------------
-
-template <typename src_t>
-static __global__ void quantize_q8_1_typed(const src_t * __restrict__ x, void * __restrict__ vy,
-                                           const int64_t ne00, const int64_t ne0) {
-    const int64_t i0 = (int64_t) blockDim.x * blockIdx.x + threadIdx.x;
-    if (i0 >= ne0) {
-        return;
-    }
-    const int64_t i1     = blockIdx.y;  // row
-    const int64_t i_cont = i1 * ne0 + i0;
-
-    block_q8_1 * y = (block_q8_1 *) vy;
-
-    const int64_t ib  = i_cont / QK8_1;
-    const int64_t iqs = i_cont % QK8_1;
-
-    const float xi = i0 < ne00 ? ggml_cuda_cast<float>(x[i1 * ne00 + i0]) : 0.0f;
-    float amax = fabsf(xi);
-    float sum  = xi;
-
-    amax = warp_reduce_max<QK8_1>(amax);
-    sum  = warp_reduce_sum<QK8_1>(sum);
-
-    const float  d = amax / 127.0f;
-    const int8_t q = amax == 0.0f ? 0 : roundf(xi / d);
-
-    y[ib].qs[iqs] = q;
-
-    if (iqs > 0) {
-        return;
-    }
-    y[ib].ds = make_half2(d, sum);
-}
-
-template <typename src_t>
-static void launch_quantize_q8_1(const void * x, void * vy, int64_t K, int64_t k_padded, int64_t M,
-                                 cudaStream_t stream) {
-    const dim3 num_blocks((k_padded + CUDA_QUANTIZE_BLOCK_SIZE - 1) / CUDA_QUANTIZE_BLOCK_SIZE, M, 1);
-    const dim3 block_size(CUDA_QUANTIZE_BLOCK_SIZE, 1, 1);
-    quantize_q8_1_typed<src_t><<<num_blocks, block_size, 0, stream>>>((const src_t *) x, vy, K, k_padded);
-}
-
-extern "C" {
-
-// ---------------------------------------------------------------------------
-// dequantize: packed blocks -> dense f32/f16/bf16
-// ---------------------------------------------------------------------------
-
-int gguf_dequantize_cuda(const void * w, void * out, int type, int64_t n_elems, int out_dtype,
-                         cudaStream_t stream) {
-    const ggml_type t = (ggml_type) type;
-    switch (out_dtype) {
-        case GGUF_DTYPE_F32: {
-            to_fp32_cuda_t f = ggml_get_to_fp32_cuda(t);
-            if (!f) return -1;
-            f(w, (float *) out, n_elems, stream);
-            return 0;
-        }
-        case GGUF_DTYPE_F16: {
-            to_fp16_cuda_t f = ggml_get_to_fp16_cuda(t);
-            if (!f) return -1;
-            f(w, (half *) out, n_elems, stream);
-            return 0;
-        }
-        case GGUF_DTYPE_BF16: {
-            to_bf16_cuda_t f = ggml_get_to_bf16_cuda(t);
-            if (!f) return -1;
-            f(w, (nv_bfloat16 *) out, n_elems, stream);
-            return 0;
-        }
-        default:
-            return -1;
-    }
-}
-
-// ---------------------------------------------------------------------------
-// mmvq: fused dequant-gemv over packed blocks. M (ncols_dst) must be <= 8.
-// W is (N rows, K cols) quantized row-major; x is (M, K) f32/f16/bf16; dst is (M, N) f32
-// (the upstream kernel hard-codes a float dst; casting is left to the caller, where
-// torch.compile can fuse it into the consumer).
-// ---------------------------------------------------------------------------
-
-// bytes of q8_1 scratch needed for the quantized activations
-size_t gguf_q8_1_scratch_bytes(int64_t K, int64_t M) {
-    const int64_t k_padded = GGML_PAD(K, MATRIX_ROW_PADDING);
-    return (size_t) (M * k_padded) * sizeof(block_q8_1) / QK8_1;
-}
-
-int gguf_mul_mat_vec_q_cuda(const void * w, const void * x, int x_dtype, float * dst,
-                            void * q8_scratch, int type, int64_t K, int64_t N, int64_t M,
-                            cudaStream_t stream) {
-    const ggml_type t = (ggml_type) type;
-    if (M > MMVQ_MAX_BATCH_SIZE) return -1;
-    const int64_t blck = ggml_blck_size(t);
-    if (blck <= 0 || K % blck != 0) return -1;
-
-    const int64_t k_padded = GGML_PAD(K, MATRIX_ROW_PADDING);
-
-    switch (x_dtype) {
-        case GGUF_DTYPE_F32:  launch_quantize_q8_1<float>       (x, q8_scratch, K, k_padded, M, stream); break;
-        case GGUF_DTYPE_F16:  launch_quantize_q8_1<half>        (x, q8_scratch, K, k_padded, M, stream); break;
-        case GGUF_DTYPE_BF16: launch_quantize_q8_1<nv_bfloat16> (x, q8_scratch, K, k_padded, M, stream); break;
-        default: return -1;
-    }
-
-    const int64_t s01 = K / blck;          // src0 row stride, in blocks
-    const int64_t s11 = k_padded / QK8_1;  // q8_1 row stride, in blocks
-    const int64_t s1  = N;                 // dst row stride, in floats
-
-    ggml_cuda_mm_fusion_args_device fusion{};
-
-    mul_mat_vec_q_switch_type(
-        w, t, q8_scratch, /*ids=*/nullptr, fusion, dst,
-        /*ncols_x=*/K, /*nrows_x=*/N, /*ncols_dst=*/M,
-        /*stride_row_x=*/s01, /*stride_col_y=*/s11, /*stride_col_dst=*/s1,
-        /*nchannels_x=*/1, /*nchannels_y=*/1, /*nchannels_dst=*/1,
-        /*stride_channel_x=*/N * s01, /*stride_channel_y=*/M * s11, /*stride_channel_dst=*/M * s1,
-        /*nsamples_x=*/1, /*nsamples_dst=*/1,
-        /*stride_sample_x=*/N * s01, /*stride_sample_y=*/M * s11, /*stride_sample_dst=*/M * s1,
-        /*ids_stride=*/0, stream);
-    return 0;
-}
-
-// NOTE: mmq (quantized gemm) and mmf (dense gemm) are deliberately not ported: both take a
-// ggml_backend_cuda_context for their pool allocator, which would pull in the whole ggml backend.
-// Above MMVQ_MAX_BATCH_SIZE rows the caller dequantizes and uses an ordinary matmul.
-
-}  // extern "C"

gguf_cuda/ggml_stubs.cu

@@ -1,192 +0,0 @@
-/* Minimal replacements for the ggml symbols that llama.cpp's ggml-cuda kernels reference,
- * so those kernels can be built as a standalone torch extension with no libggml and no
- * ggml backend/scheduler runtime.
- *
- * The kernels only need:
- *   - ggml_cuda_info() / ggml_cuda_get_device() / ggml_cuda_set_device()  (declared in common.cuh,
- *     defined in ggml-cuda.cu, which would drag in the entire backend)
- *   - a few ggml core type traits from ggml.c
- *
- * Everything else in common.cuh is inline/template.
- */
-
-#include "common.cuh"
-
-#include <cstdio>
-#include <cstdlib>
-#include <mutex>
-
-// ---------------------------------------------------------------------------
-// device info
-// ---------------------------------------------------------------------------
-
-static ggml_cuda_device_info build_device_info() {
-    ggml_cuda_device_info info = {};
-
-    int device_count = 0;
-    if (cudaGetDeviceCount(&device_count) != cudaSuccess) {
-        fprintf(stderr, "%s: cudaGetDeviceCount failed\n", __func__);
-        return info;
-    }
-    info.device_count          = device_count;
-    info.physical_device_count = device_count;
-
-    for (int id = 0; id < device_count; ++id) {
-        cudaDeviceProp prop;
-        CUDA_CHECK(cudaGetDeviceProperties(&prop, id));
-
-        info.default_tensor_split[id] = float(id) / float(device_count);
-
-        info.devices[id].nsm                         = prop.multiProcessorCount;
-        info.devices[id].smpb                        = prop.sharedMemPerBlock;
-        info.devices[id].smpbo                       = prop.sharedMemPerBlockOptin;
-        info.devices[id].warp_size                   = prop.warpSize;
-        info.devices[id].integrated                  = prop.integrated;
-        info.devices[id].vmm                         = false;  // no virtual memory pool here
-        info.devices[id].vmm_granularity             = 0;
-        info.devices[id].total_vram                  = prop.totalGlobalMem;
-        info.devices[id].supports_cooperative_launch = prop.cooperativeLaunch;
-        info.devices[id].physical_device             = id;
-        info.devices[id].physical_share_count        = 1;
-        info.devices[id].virtual_index               = 0;
-        info.devices[id].cc = 100 * prop.major + 10 * prop.minor;
-    }
-    return info;
-}
-
-const ggml_cuda_device_info & ggml_cuda_info() {
-    static ggml_cuda_device_info info = build_device_info();
-    return info;
-}
-
-int ggml_cuda_get_device() {
-    int id = 0;
-    CUDA_CHECK(cudaGetDevice(&id));
-    return id;
-}
-
-void ggml_cuda_set_device(int device) {
-    int current = 0;
-    CUDA_CHECK(cudaGetDevice(&current));
-    if (current == device) {
-        return;
-    }
-    CUDA_CHECK(cudaSetDevice(device));
-}
-
-// ---------------------------------------------------------------------------
-// ggml core type traits (subset of ggml.c). Values come from the block structs and QK_*
-// macros in ggml-common.h, so they cannot drift from the kernels that consume them.
-// ---------------------------------------------------------------------------
-
-#define GGML_TRAIT(t, blck_, size_, quant_) \
-    case GGML_TYPE_##t: return trait{ blck_, size_, quant_ };
-
-namespace {
-struct trait {
-    int64_t blck;
-    size_t  size;
-    bool    quantized;
-};
-
-trait type_trait(ggml_type type) {
-    switch (type) {
-        GGML_TRAIT(F32,     1,      sizeof(float),          false)
-        GGML_TRAIT(F16,     1,      sizeof(half),           false)
-        GGML_TRAIT(BF16,    1,      sizeof(nv_bfloat16),    false)
-        GGML_TRAIT(I8,      1,      sizeof(int8_t),         false)
-        GGML_TRAIT(I16,     1,      sizeof(int16_t),        false)
-        GGML_TRAIT(I32,     1,      sizeof(int32_t),        false)
-        GGML_TRAIT(I64,     1,      sizeof(int64_t),        false)
-        GGML_TRAIT(F64,     1,      sizeof(double),         false)
-        GGML_TRAIT(Q1_0,    QK1_0,  sizeof(block_q1_0),     true)
-        GGML_TRAIT(Q2_0,    QK2_0,  sizeof(block_q2_0),     true)
-        GGML_TRAIT(Q4_0,    QK4_0,  sizeof(block_q4_0),     true)
-        GGML_TRAIT(Q4_1,    QK4_1,  sizeof(block_q4_1),     true)
-        GGML_TRAIT(Q5_0,    QK5_0,  sizeof(block_q5_0),     true)
-        GGML_TRAIT(Q5_1,    QK5_1,  sizeof(block_q5_1),     true)
-        GGML_TRAIT(Q8_0,    QK8_0,  sizeof(block_q8_0),     true)
-        GGML_TRAIT(Q8_1,    QK8_1,  sizeof(block_q8_1),     true)
-        GGML_TRAIT(Q2_K,    QK_K,   sizeof(block_q2_K),     true)
-        GGML_TRAIT(Q3_K,    QK_K,   sizeof(block_q3_K),     true)
-        GGML_TRAIT(Q4_K,    QK_K,   sizeof(block_q4_K),     true)
-        GGML_TRAIT(Q5_K,    QK_K,   sizeof(block_q5_K),     true)
-        GGML_TRAIT(Q6_K,    QK_K,   sizeof(block_q6_K),     true)
-        GGML_TRAIT(Q8_K,    QK_K,   sizeof(block_q8_K),     true)
-        GGML_TRAIT(IQ1_S,   QK_K,   sizeof(block_iq1_s),    true)
-        GGML_TRAIT(IQ1_M,   QK_K,   sizeof(block_iq1_m),    true)
-        GGML_TRAIT(IQ2_XXS, QK_K,   sizeof(block_iq2_xxs),  true)
-        GGML_TRAIT(IQ2_XS,  QK_K,   sizeof(block_iq2_xs),   true)
-        GGML_TRAIT(IQ2_S,   QK_K,   sizeof(block_iq2_s),    true)
-        GGML_TRAIT(IQ3_XXS, QK_K,   sizeof(block_iq3_xxs),  true)
-        GGML_TRAIT(IQ3_S,   QK_K,   sizeof(block_iq3_s),    true)
-        GGML_TRAIT(IQ4_NL,  QK4_NL, sizeof(block_iq4_nl),   true)
-        GGML_TRAIT(IQ4_XS,  QK_K,   sizeof(block_iq4_xs),   true)
-        GGML_TRAIT(TQ1_0,   QK_K,   sizeof(block_tq1_0),    true)
-        GGML_TRAIT(TQ2_0,   QK_K,   sizeof(block_tq2_0),    true)
-        GGML_TRAIT(MXFP4,   QK_MXFP4, sizeof(block_mxfp4),  true)
-        GGML_TRAIT(NVFP4,   QK_NVFP4, sizeof(block_nvfp4),  true)
-        default: return trait{ 0, 0, false };  // callers must reject blck == 0
-    }
-}
-}  // namespace
-
-extern "C" {
-
-int64_t ggml_blck_size(enum ggml_type type)  { return type_trait(type).blck; }
-size_t  ggml_type_size(enum ggml_type type)  { return type_trait(type).size; }
-bool    ggml_is_quantized(enum ggml_type type) { return type_trait(type).quantized; }
-
-// Referenced only by ggml's own error/abort messages.
-const char * ggml_type_name(enum ggml_type type) {
-    switch (type) {
-        case GGML_TYPE_F32:  return "f32";
-        case GGML_TYPE_F16:  return "f16";
-        case GGML_TYPE_BF16: return "bf16";
-        case GGML_TYPE_Q4_K: return "q4_K";
-        case GGML_TYPE_Q5_K: return "q5_K";
-        case GGML_TYPE_Q6_K: return "q6_K";
-        case GGML_TYPE_Q8_0: return "q8_0";
-        default: return "unknown";
-    }
-}
-
-void ggml_abort(const char * file, int line, const char * fmt, ...) {
-    fprintf(stderr, "ggml abort at %s:%d: %s\n", file, line, fmt ? fmt : "");
-    abort();
-}
-
-}  // extern "C"
-
-[[noreturn]] void ggml_cuda_error(const char * stmt, const char * func, const char * file, int line,
-                                  const char * msg) {
-    fprintf(stderr, "CUDA error: %s\n  %s at %s:%d\n  %s\n", msg, func, file, line, stmt);
-    abort();
-}
-
-// ---------------------------------------------------------------------------
-// Stubs. These are referenced only by the ggml_tensor-based entry points in mmvq.cu
-// (ggml_cuda_mul_mat_vec_q / ggml_cuda_op_mul_mat_vec_q), which get compiled but are never
-// called — we go straight to the raw-pointer dispatch. If one of them ever runs, abort loudly
-// rather than silently misbehave.
-// ---------------------------------------------------------------------------
-
-#define GGML_SHIM_UNREACHABLE(name)                                                  \
-    do {                                                                             \
-        fprintf(stderr, "ggml-quantization: %s is a stub and must not be called\n", name); \
-        abort();                                                                     \
-    } while (0)
-
-extern "C" {
-int64_t ggml_nelements(const ggml_tensor *)                    { GGML_SHIM_UNREACHABLE("ggml_nelements"); }
-size_t  ggml_nbytes(const ggml_tensor *)                       { GGML_SHIM_UNREACHABLE("ggml_nbytes"); }
-bool    ggml_is_contiguous(const ggml_tensor *)                { GGML_SHIM_UNREACHABLE("ggml_is_contiguous"); }
-bool    ggml_is_contiguously_allocated(const ggml_tensor *)    { GGML_SHIM_UNREACHABLE("ggml_is_contiguously_allocated"); }
-bool    ggml_are_same_stride(const ggml_tensor *, const ggml_tensor *) { GGML_SHIM_UNREACHABLE("ggml_are_same_stride"); }
-enum ggml_backend_buffer_usage ggml_backend_buffer_get_usage(ggml_backend_buffer_t) { GGML_SHIM_UNREACHABLE("ggml_backend_buffer_get_usage"); }
-size_t  ggml_backend_buffer_get_alloc_size(ggml_backend_buffer_t, const ggml_tensor *) { GGML_SHIM_UNREACHABLE("ggml_backend_buffer_get_alloc_size"); }
-}
-
-std::unique_ptr<ggml_cuda_pool> ggml_backend_cuda_context::new_pool_for_device(int, int) {
-    GGML_SHIM_UNREACHABLE("ggml_backend_cuda_context::new_pool_for_device");
-}

gguf_cuda/gguf_cuda.cu

@@ -1,84 +0,0 @@
-/* CUDA implementation of the two entry points.
- *
- * Deliberately includes no ggml header: all ggml contact is behind the extern "C" boundary in
- * ggml_dispatch.cu, so torch and ggml headers never meet. Outputs and scratch come from torch's
- * caching allocator and everything runs on torch's current stream, which keeps the ops
- * CUDA-graph capturable.
- */
-
-#include <ATen/cuda/CUDAContext.h>
-#include <c10/cuda/CUDAGuard.h>
-#include <torch/torch.h>
-
-#include "torch_binding.h"
-
-extern "C" {
-int    gguf_dequantize_cuda(const void *w, void *out, int type, int64_t n_elems, int out_dtype,
-                            cudaStream_t stream);
-size_t gguf_q8_1_scratch_bytes(int64_t K, int64_t M);
-int    gguf_mul_mat_vec_q_cuda(const void *w, const void *x, int x_dtype, float *dst,
-                               void *q8_scratch, int type, int64_t K, int64_t N, int64_t M,
-                               cudaStream_t stream);
-}
-
-namespace {
-
-// mirrors gguf_dtype in ggml_dispatch.cu
-int dtype_code(at::ScalarType t) {
-  switch (t) {
-    case at::kFloat:    return 0;
-    case at::kHalf:     return 1;
-    case at::kBFloat16: return 2;
-    default: TORCH_CHECK(false, "ggml-quantization: unsupported dtype ", t);
-  }
-}
-
-}  // namespace
-
-// The cases upstream's `mul_mat_vec_q_switch_type` dispatches (mmvq.cu): Q4_0/Q4_1/Q5_0/Q5_1/Q8_0,
-// the K quants, the IQ quants, MXFP4/NVFP4 and Q1_0/Q2_0. Kept as a literal because that switch is
-// static and exposes no way to enumerate it; a type added upstream shows up here when the pin moves.
-std::vector<int64_t> gemv_types() {
-  return {2, 3, 6, 7, 8, 10, 11, 12, 13, 14, 16, 17, 18, 19, 20, 21, 22, 23, 29, 39, 40, 41, 42};
-}
-
-at::Tensor dequantize(const at::Tensor &blocks, int64_t ggml_type, int64_t rows, int64_t cols,
-                      at::ScalarType dtype) {
-  TORCH_CHECK(blocks.is_cuda() && blocks.scalar_type() == at::kByte, "blocks must be cuda uint8");
-  TORCH_CHECK(blocks.is_contiguous(), "blocks must be contiguous");
-  const at::cuda::OptionalCUDAGuard guard(at::device_of(blocks));
-
-  auto out = at::empty({rows, cols}, blocks.options().dtype(dtype));
-  const int rc = gguf_dequantize_cuda(blocks.data_ptr(), out.data_ptr(), (int)ggml_type, rows * cols,
-                                      dtype_code(dtype), at::cuda::getCurrentCUDAStream());
-  TORCH_CHECK(rc == 0, "ggml-quantization: dequantize has no implementation for ggml type ", ggml_type);
-  return out;
-}
-
-at::Tensor get_rows(const at::Tensor &blocks, const at::Tensor &indices, int64_t ggml_type,
-                    int64_t cols, at::ScalarType dtype) {
-  // Upstream's CUDA dequantize walks a whole tensor and takes no indices, so the rows are gathered
-  // first and unpacked after -- correct, but without Metal's saving of never touching the rest.
-  return dequantize(blocks.index_select(0, indices), ggml_type, indices.numel(), cols, dtype);
-}
-
-at::Tensor mul_mat_vec(const at::Tensor &blocks, const at::Tensor &x, int64_t ggml_type,
-                       int64_t out_features) {
-  TORCH_CHECK(blocks.is_cuda() && blocks.scalar_type() == at::kByte, "blocks must be cuda uint8");
-  TORCH_CHECK(x.is_cuda() && x.dim() == 2, "x must be a 2D cuda tensor");
-  const at::cuda::OptionalCUDAGuard guard(at::device_of(blocks));
-
-  const int64_t rows = x.size(0), in_features = x.size(1);
-  const auto xc = x.contiguous();
-  // The activations are quantized to q8_1 first; that scratch is a tensor so it comes from torch's
-  // allocator rather than a pool of ggml's own.
-  auto scratch = at::empty({(int64_t)gguf_q8_1_scratch_bytes(in_features, rows)}, blocks.options());
-  auto out = at::empty({rows, out_features}, x.options().dtype(at::kFloat));
-
-  const int rc = gguf_mul_mat_vec_q_cuda(blocks.data_ptr(), xc.data_ptr(),
-                                         dtype_code(xc.scalar_type()), out.data_ptr<float>(),
-                                         scratch.data_ptr(), (int)ggml_type, in_features,
-                                         out_features, rows, at::cuda::getCurrentCUDAStream());
-  TORCH_CHECK(rc == 0, "ggml-quantization: no gemv for ggml type ", ggml_type, " at ", rows, " rows");
-  return out;
-}