feat: HIGGS quantization kernels for HF Kernel Hub

Dequantize (codebook lookup), quantize fp16, quantize bf16.
Ported from github.com/galqiwi/higgs-kernels.
Source only - build/ to be generated with nix run .#build-and-copy

.gitattributes

@@ -1,35 +1,2 @@
-*.7z filter=lfs diff=lfs merge=lfs -text
-*.arrow filter=lfs diff=lfs merge=lfs -text
-*.bin filter=lfs diff=lfs merge=lfs -text
-*.bz2 filter=lfs diff=lfs merge=lfs -text
-*.ckpt filter=lfs diff=lfs merge=lfs -text
-*.ftz filter=lfs diff=lfs merge=lfs -text
-*.gz filter=lfs diff=lfs merge=lfs -text
-*.h5 filter=lfs diff=lfs merge=lfs -text
-*.joblib filter=lfs diff=lfs merge=lfs -text
-*.lfs.* filter=lfs diff=lfs merge=lfs -text
-*.mlmodel filter=lfs diff=lfs merge=lfs -text
-*.model filter=lfs diff=lfs merge=lfs -text
-*.msgpack filter=lfs diff=lfs merge=lfs -text
-*.npy filter=lfs diff=lfs merge=lfs -text
-*.npz filter=lfs diff=lfs merge=lfs -text
-*.onnx filter=lfs diff=lfs merge=lfs -text
-*.ot filter=lfs diff=lfs merge=lfs -text
-*.parquet filter=lfs diff=lfs merge=lfs -text
-*.pb filter=lfs diff=lfs merge=lfs -text
-*.pickle filter=lfs diff=lfs merge=lfs -text
-*.pkl filter=lfs diff=lfs merge=lfs -text
-*.pt filter=lfs diff=lfs merge=lfs -text
-*.pth filter=lfs diff=lfs merge=lfs -text
-*.rar filter=lfs diff=lfs merge=lfs -text
 *.safetensors filter=lfs diff=lfs merge=lfs -text
-saved_model/**/* filter=lfs diff=lfs merge=lfs -text
-*.tar.* filter=lfs diff=lfs merge=lfs -text
-*.tar filter=lfs diff=lfs merge=lfs -text
-*.tflite filter=lfs diff=lfs merge=lfs -text
-*.tgz filter=lfs diff=lfs merge=lfs -text
-*.wasm filter=lfs diff=lfs merge=lfs -text
-*.xz filter=lfs diff=lfs merge=lfs -text
-*.zip filter=lfs diff=lfs merge=lfs -text
-*.zst filter=lfs diff=lfs merge=lfs -text
-*tfevents* filter=lfs diff=lfs merge=lfs -text
+*.so filter=lfs diff=lfs merge=lfs -text

LICENSE

@@ -0,0 +1,201 @@
+                                 Apache License
+                           Version 2.0, January 2004
+                        http://www.apache.org/licenses/
+
+   TERMS AND CONDITIONS FOR USE, REPRODUCTION, AND DISTRIBUTION
+
+   1. Definitions.
+
+      "License" shall mean the terms and conditions for use, reproduction,
+      and distribution as defined by Sections 1 through 9 of this document.
+
+      "Licensor" shall mean the copyright owner or entity authorized by
+      the copyright owner that is granting the License.
+
+      "Legal Entity" shall mean the union of the acting entity and all
+      other entities that control, are controlled by, or are under common
+      control with that entity. For the purposes of this definition,
+      "control" means (i) the power, direct or indirect, to cause the
+      direction or management of such entity, whether by contract or
+      otherwise, or (ii) ownership of fifty percent (50%) or more of the
+      outstanding shares, or (iii) beneficial ownership of such entity.
+
+      "You" (or "Your") shall mean an individual or Legal Entity
+      exercising permissions granted by this License.
+
+      "Source" form shall mean the preferred form for making modifications,
+      including but not limited to software source code, documentation
+      source, and configuration files.
+
+      "Object" form shall mean any form resulting from mechanical
+      transformation or translation of a Source form, including but
+      not limited to compiled object code, generated documentation,
+      and conversions to other media types.
+
+      "Work" shall mean the work of authorship, whether in Source or
+      Object form, made available under the License, as indicated by a
+      copyright notice that is included in or attached to the work
+      (an example is provided in the Appendix below).
+
+      "Derivative Works" shall mean any work, whether in Source or Object
+      form, that is based on (or derived from) the Work and for which the
+      editorial revisions, annotations, elaborations, or other modifications
+      represent, as a whole, an original work of authorship. For the purposes
+      of this License, Derivative Works shall not include works that remain
+      separable from, or merely link (or bind by name) to the interfaces of,
+      the Work and Derivative Works thereof.
+
+      "Contribution" shall mean any work of authorship, including
+      the original version of the Work and any modifications or additions
+      to that Work or Derivative Works thereof, that is intentionally
+      submitted to Licensor for inclusion in the Work by the copyright owner
+      or by an individual or Legal Entity authorized to submit on behalf of
+      the copyright owner. For the purposes of this definition, "submitted"
+      means any form of electronic, verbal, or written communication sent
+      to the Licensor or its representatives, including but not limited to
+      communication on electronic mailing lists, source code control systems,
+      and issue tracking systems that are managed by, or on behalf of, the
+      Licensor for the purpose of discussing and improving the Work, but
+      excluding communication that is conspicuously marked or otherwise
+      designated in writing by the copyright owner as "Not a Contribution."
+
+      "Contributor" shall mean Licensor and any individual or Legal Entity
+      on behalf of whom a Contribution has been received by Licensor and
+      subsequently incorporated within the Work.
+
+   2. Grant of Copyright License. Subject to the terms and conditions of
+      this License, each Contributor hereby grants to You a perpetual,
+      worldwide, non-exclusive, no-charge, royalty-free, irrevocable
+      copyright license to reproduce, prepare Derivative Works of,
+      publicly display, publicly perform, sublicense, and distribute the
+      Work and such Derivative Works in Source or Object form.
+
+   3. Grant of Patent License. Subject to the terms and conditions of
+      this License, each Contributor hereby grants to You a perpetual,
+      worldwide, non-exclusive, no-charge, royalty-free, irrevocable
+      (except as stated in this section) patent license to make, have made,
+      use, offer to sell, sell, import, and otherwise transfer the Work,
+      where such license applies only to those patent claims licensable
+      by such Contributor that are necessarily infringed by their
+      Contribution(s) alone or by combination of their Contribution(s)
+      with the Work to which such Contribution(s) was submitted. If You
+      institute patent litigation against any entity (including a
+      cross-claim or counterclaim in a lawsuit) alleging that the Work
+      or a Contribution incorporated within the Work constitutes direct
+      or contributory patent infringement, then any patent licenses
+      granted to You under this License for that Work shall terminate
+      as of the date such litigation is filed.
+
+   4. Redistribution. You may reproduce and distribute copies of the
+      Work or Derivative Works thereof in any medium, with or without
+      modifications, and in Source or Object form, provided that You
+      meet the following conditions:
+
+      (a) You must give any other recipients of the Work or
+          Derivative Works a copy of this License; and
+
+      (b) You must cause any modified files to carry prominent notices
+          stating that You changed the files; and
+
+      (c) You must retain, in the Source form of any Derivative Works
+          that You distribute, all copyright, patent, trademark, and
+          attribution notices from the Source form of the Work,
+          excluding those notices that do not pertain to any part of
+          the Derivative Works; and
+
+      (d) If the Work includes a "NOTICE" text file as part of its
+          distribution, then any Derivative Works that You distribute must
+          include a readable copy of the attribution notices contained
+          within such NOTICE file, excluding those notices that do not
+          pertain to any part of the Derivative Works, in at least one
+          of the following places: within a NOTICE text file distributed
+          as part of the Derivative Works; within the Source form or
+          documentation, if provided along with the Derivative Works; or,
+          within a display generated by the Derivative Works, if and
+          wherever such third-party notices normally appear. The contents
+          of the NOTICE file are for informational purposes only and
+          do not modify the License. You may add Your own attribution
+          notices within Derivative Works that You distribute, alongside
+          or as an addendum to the NOTICE text from the Work, provided
+          that such additional attribution notices cannot be construed
+          as modifying the License.
+
+      You may add Your own copyright statement to Your modifications and
+      may provide additional or different license terms and conditions
+      for use, reproduction, or distribution of Your modifications, or
+      for any such Derivative Works as a whole, provided Your use,
+      reproduction, and distribution of the Work otherwise complies with
+      the conditions stated in this License.
+
+   5. Submission of Contributions. Unless You explicitly state otherwise,
+      any Contribution intentionally submitted for inclusion in the Work
+      by You to the Licensor shall be under the terms and conditions of
+      this License, without any additional terms or conditions.
+      Notwithstanding the above, nothing herein shall supersede or modify
+      the terms of any separate license agreement you may have executed
+      with Licensor regarding such Contributions.
+
+   6. Trademarks. This License does not grant permission to use the trade
+      names, trademarks, service marks, or product names of the Licensor,
+      except as required for reasonable and customary use in describing the
+      origin of the Work and reproducing the content of the NOTICE file.
+
+   7. Disclaimer of Warranty. Unless required by applicable law or
+      agreed to in writing, Licensor provides the Work (and each
+      Contributor provides its Contributions) on an "AS IS" BASIS,
+      WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or
+      implied, including, without limitation, any warranties or conditions
+      of TITLE, NON-INFRINGEMENT, MERCHANTABILITY, or FITNESS FOR A
+      PARTICULAR PURPOSE. You are solely responsible for determining the
+      appropriateness of using or redistributing the Work and assume any
+      risks associated with Your exercise of permissions under this License.
+
+   8. Limitation of Liability. In no event and under no legal theory,
+      whether in tort (including negligence), contract, or otherwise,
+      unless required by applicable law (such as deliberate and grossly
+      negligent acts) or agreed to in writing, shall any Contributor be
+      liable to You for damages, including any direct, indirect, special,
+      incidental, or consequential damages of any character arising as a
+      result of this License or out of the use or inability to use the
+      Work (including but not limited to damages for loss of goodwill,
+      work stoppage, computer failure or malfunction, or any and all
+      other commercial damages or losses), even if such Contributor
+      has been advised of the possibility of such damages.
+
+   9. Accepting Warranty or Additional Liability. While redistributing
+      the Work or Derivative Works thereof, You may choose to offer,
+      and charge a fee for, acceptance of support, warranty, indemnity,
+      or other liability obligations and/or rights consistent with this
+      License. However, in accepting such obligations, You may act only
+      on Your own behalf and on Your sole responsibility, not on behalf
+      of any other Contributor, and only if You agree to indemnify,
+      defend, and hold each Contributor harmless for any liability
+      incurred by, or claims asserted against, such Contributor by reason
+      of your accepting any such warranty or additional liability.
+
+   END OF TERMS AND CONDITIONS
+
+   APPENDIX: How to apply the Apache License to your work.
+
+      To apply the Apache License to your work, attach the following
+      boilerplate notice, with the fields enclosed by brackets "[]"
+      replaced with your own identifying information. (Don't include
+      the brackets!)  The text should be enclosed in the appropriate
+      comment syntax for the file format. We also recommend that a
+      file or class name and description of purpose be included on the
+      same "printed page" as the copyright notice for easier
+      identification within third-party archives.
+
+   Copyright [yyyy] [name of copyright owner]
+
+   Licensed under the Apache License, Version 2.0 (the "License");
+   you may not use this file except in compliance with the License.
+   You may obtain a copy of the License at
+
+       http://www.apache.org/licenses/LICENSE-2.0
+
+   Unless required by applicable law or agreed to in writing, software
+   distributed under the License is distributed on an "AS IS" BASIS,
+   WITHOUT WARRANTIES OR CONDITIONS OF ANY KIND, either express or implied.
+   See the License for the specific language governing permissions and
+   limitations under the License.

README.md

@@ -0,0 +1,36 @@
+# higgs-kernels
+
+CUDA kernels for [HIGGS](https://arxiv.org/abs/2410.20939) quantization, packaged for the [Hugging Face Kernel Hub](https://huggingface.co/docs/kernels).
+
+Extracted from [galqiwi/higgs-kernels](https://github.com/galqiwi/higgs-kernels).
+
+## Kernels
+
+- `higgs_dequantize_2_256` - codebook lookup: uint8 indices -> 2D fp16/bf16 vectors
+- `higgs_quantize_2_256_f16` - nearest codebook entry search (fp16)
+- `higgs_quantize_2_256_bf16` - nearest codebook entry search (bf16)
+
+## Usage
+
+```python
+from kernels import get_kernel
+
+higgs = get_kernel("galqiwi/higgs-kernels")
+
+out = higgs.higgs_dequantize_2_256_kernel(x_uint8, grid_256x2)
+indices = higgs.higgs_quantize_2_256_kernel(x_fp16_Nx2, grid_256x2, grid_norms_256)
+```
+
+## Grid data
+
+Pre-trained codebook included in `grids.safetensors` (256x2, key `"2_256"`).
+
+```python
+from kernels import get_kernel
+higgs = get_kernel("galqiwi/higgs-kernels")
+grid = higgs.load_optimal_grid_2_256(device="cuda", dtype=torch.float16)
+```
+
+## License
+
+Apache-2.0

build.toml

@@ -0,0 +1,24 @@
+[general]
+name = "higgs-kernels"
+universal = false
+
+[torch]
+src = [
+  "torch-ext/torch_binding.cpp",
+  "torch-ext/torch_binding.h",
+]
+
+[kernel.dequant]
+backend = "cuda"
+depends = ["torch"]
+src = ["csrc/dequant.cu"]
+
+[kernel.quant_f16]
+backend = "cuda"
+depends = ["torch"]
+src = ["csrc/quant_f16.cu"]
+
+[kernel.quant_bf16]
+backend = "cuda"
+depends = ["torch"]
+src = ["csrc/quant_bf16.cu"]

csrc/dequant.cu

@@ -0,0 +1,75 @@
+#include <cstdint>
+#include <cuda.h>
+#include <cuda_runtime.h>
+#include <cuda_fp16.h>
+#include <ATen/cuda/CUDAContext.h>
+#include <c10/cuda/CUDAException.h>
+
+// Each thread processes 4 indices: loads 4 bytes (uint32), stores 16 bytes (uint4)
+static constexpr int ELEMENTS_PER_THREAD = 4;
+
+static __global__ void higgs_dequantize_2_256_ptr_cuda_portable_kernel(
+	const uint8_t* __restrict__ x,
+	const uint32_t* __restrict__ grid_packed,
+	uint32_t* __restrict__ out_packed,
+	long long out_dim) {
+	__shared__ uint32_t s_grid[256];
+
+	// Load codebook to shared memory
+	for (int idx = threadIdx.x; idx < 256; idx += blockDim.x) {
+		s_grid[idx] = grid_packed[idx];
+	}
+	__syncthreads();
+
+	// Each thread processes 4 indices
+	long long base_idx = (static_cast<long long>(blockIdx.x) * blockDim.x + threadIdx.x) * ELEMENTS_PER_THREAD;
+
+	if (base_idx >= out_dim) return;
+
+	// Check if we have a full 4 elements to process
+	if (base_idx + ELEMENTS_PER_THREAD <= out_dim) {
+		// Vectorized load: load 4 uint8 indices as uint32 (4 bytes)
+		uint32_t indices_packed = *reinterpret_cast<const uint32_t*>(&x[base_idx]);
+
+		// Extract individual bytes
+		uint8_t idx0 = indices_packed & 0xFF;
+		uint8_t idx1 = (indices_packed >> 8) & 0xFF;
+		uint8_t idx2 = (indices_packed >> 16) & 0xFF;
+		uint8_t idx3 = (indices_packed >> 24) & 0xFF;
+
+		// Lookup all 4 values
+		uint32_t val0 = s_grid[idx0];
+		uint32_t val1 = s_grid[idx1];
+		uint32_t val2 = s_grid[idx2];
+		uint32_t val3 = s_grid[idx3];
+
+		// Vectorized store: write 4 uint32 values as uint4 (16 bytes)
+		uint4 result = make_uint4(val0, val1, val2, val3);
+		*reinterpret_cast<uint4*>(&out_packed[base_idx]) = result;
+	} else {
+		// Handle remainder (less than 4 elements at the end)
+		for (long long i = base_idx; i < out_dim; i++) {
+			out_packed[i] = s_grid[x[i]];
+		}
+	}
+}
+
+extern "C" void higgs_dequantize_2_256_ptr_cuda_portable(
+	uint64_t x_ptr,
+	uint64_t grid_ptr,
+	uint64_t out_ptr,
+	int64_t out_dim) {
+	const uint8_t* x = reinterpret_cast<const uint8_t*>(x_ptr);
+	const uint32_t* grid_packed = reinterpret_cast<const uint32_t*>(grid_ptr);
+	uint32_t* out_packed = reinterpret_cast<uint32_t*>(out_ptr);
+
+	constexpr int threads_per_block = 256;
+	constexpr int elements_per_block = threads_per_block * ELEMENTS_PER_THREAD;
+	int blocks = static_cast<int>((out_dim + elements_per_block - 1) / elements_per_block);
+
+	auto stream = at::cuda::getCurrentCUDAStream();
+	higgs_dequantize_2_256_ptr_cuda_portable_kernel<<<blocks, threads_per_block, 0, stream.stream()>>>(
+		x, grid_packed, out_packed, static_cast<long long>(out_dim));
+
+	C10_CUDA_KERNEL_LAUNCH_CHECK();
+}

csrc/quant_bf16.cu

@@ -0,0 +1,103 @@
+#include <cuda_runtime.h>
+#include <math_constants.h>
+#include <cstdint>
+#include <limits>
+#include <ATen/cuda/CUDAContext.h>
+#include <c10/cuda/CUDAException.h>
+
+static constexpr int kCodebookSize = 256;
+static constexpr int kInDim = 2;
+
+__device__ __forceinline__ float bf16_to_fp32(uint16_t v) {
+	uint32_t u = static_cast<uint32_t>(v) << 16;
+	return __uint_as_float(u);
+}
+__device__ __forceinline__ uint16_t fp32_to_bf16_rne(float f) {
+	uint32_t x = __float_as_uint(f);
+	uint32_t lsb = (x >> 16) & 1U;
+	uint32_t rounding_bias = 0x00007FFFU + lsb;
+	x += rounding_bias;
+	return static_cast<uint16_t>(x >> 16);
+}
+
+__device__ __forceinline__ uint16_t bf16_mul(uint16_t a_bf16, uint16_t b_bf16) {
+	float a = bf16_to_fp32(a_bf16);
+	float b = bf16_to_fp32(b_bf16);
+	return fp32_to_bf16_rne(a * b);
+}
+
+__global__ void higgs_quantize_2_256_ptr_bf16_cuda_portable_kernel(
+	const uint16_t* __restrict__ x,
+	const uint16_t* __restrict__ grid,
+	const uint16_t* __restrict__ grid_norms,
+	unsigned char* __restrict__ out,
+	int64_t out_dim)
+{
+	__shared__ float s_grid[kCodebookSize][kInDim];
+	__shared__ uint16_t s_norms_bf16[kCodebookSize];
+
+	for (int idx = threadIdx.x; idx < kCodebookSize * kInDim; idx += blockDim.x) {
+		int r = idx / kInDim;
+		int c = idx % kInDim;
+		s_grid[r][c] = bf16_to_fp32(grid[r * kInDim + c]);
+	}
+	for (int idx = threadIdx.x; idx < kCodebookSize; idx += blockDim.x) {
+		s_norms_bf16[idx] = grid_norms[idx];
+	}
+	__syncthreads();
+
+	int64_t row = blockIdx.x * blockDim.x + threadIdx.x;
+	if (row >= out_dim) return;
+
+	const uint16_t two_bf16 = fp32_to_bf16_rne(2.0f);
+
+	uint16_t x0_b = x[row * kInDim + 0];
+	uint16_t x1_b = x[row * kInDim + 1];
+	float x0 = bf16_to_fp32(x0_b);
+	float x1 = bf16_to_fp32(x1_b);
+
+	float best_score = -CUDART_INF_F;
+	unsigned int best_index = 0u;
+
+	#pragma unroll 8
+	for (int c = 0; c < kCodebookSize; ++c) {
+		float g0 = s_grid[c][0];
+		float g1 = s_grid[c][1];
+		float dot_fp32 = x0 * g0 + x1 * g1;
+		uint16_t dot_bf16 = fp32_to_bf16_rne(dot_fp32);
+
+		uint16_t twice_dot_bf16 = bf16_mul(dot_bf16, two_bf16);
+		float twice_dot = bf16_to_fp32(twice_dot_bf16);
+
+		uint16_t grid_norm_bf16 = s_norms_bf16[c];
+		float score = bf16_to_fp32(fp32_to_bf16_rne(twice_dot - bf16_to_fp32(grid_norm_bf16)));
+
+		if (score > best_score) {
+			best_score = score;
+			best_index = static_cast<unsigned int>(c);
+		}
+	}
+
+	out[row] = static_cast<unsigned char>(best_index);
+}
+
+extern "C" void higgs_quantize_2_256_ptr_bf16_cuda_portable(
+	uint64_t x_ptr,
+	uint64_t grid_ptr,
+	uint64_t grid_norms_ptr,
+	uint64_t out_ptr,
+	int64_t out_dim)
+{
+	const uint16_t* x = reinterpret_cast<const uint16_t*>(x_ptr);
+	const uint16_t* grid = reinterpret_cast<const uint16_t*>(grid_ptr);
+	const uint16_t* grid_norms = reinterpret_cast<const uint16_t*>(grid_norms_ptr);
+	unsigned char* out = reinterpret_cast<unsigned char*>(out_ptr);
+
+	int threads = 256;
+	int blocks = static_cast<int>((out_dim + threads - 1) / threads);
+
+	auto stream = at::cuda::getCurrentCUDAStream();
+	higgs_quantize_2_256_ptr_bf16_cuda_portable_kernel<<<blocks, threads, 0, stream>>>(x, grid, grid_norms, out, out_dim);
+
+	C10_CUDA_KERNEL_LAUNCH_CHECK();
+}

csrc/quant_f16.cu

@@ -0,0 +1,78 @@
+#include <cuda_fp16.h>
+#include <stdint.h>
+#include <float.h>
+#include <ATen/cuda/CUDAContext.h>
+#include <c10/cuda/CUDAException.h>
+
+static constexpr int codebookSize = 256;
+static constexpr int codebookDim = 2;
+
+static __global__ void higgs_quantize_2_256_ptr_f16_cuda_portable_kernel(
+	const __half* __restrict__ x,
+	const __half* __restrict__ grid,
+	const __half* __restrict__ grid_norms,
+	uint8_t* __restrict__ out,
+	int64_t out_dim)
+{
+	__shared__ float s_grid[codebookSize][codebookDim];
+	__shared__ __half s_norms[codebookSize];
+
+	for (int idx = threadIdx.x; idx < codebookSize * codebookDim; idx += blockDim.x) {
+		int r = idx / codebookDim;
+		int c = idx % codebookDim;
+		s_grid[r][c] = __half2float(grid[r * codebookDim + c]);
+	}
+	for (int idx = threadIdx.x; idx < codebookSize; idx += blockDim.x) {
+		s_norms[idx] = grid_norms[idx];
+	}
+	__syncthreads();
+
+	int64_t row = static_cast<int64_t>(blockIdx.x) * blockDim.x + threadIdx.x;
+	if (row >= out_dim) return;
+
+	const __half two_h = __float2half(2.0f);
+
+	const __half x0_h = x[row * codebookDim + 0];
+	const __half x1_h = x[row * codebookDim + 1];
+
+	float best_score = -FLT_MAX;
+	uint8_t best_index = 0;
+
+	#pragma unroll 8
+	for (int c = 0; c < codebookSize; ++c) {
+		const float g0 = s_grid[c][0];
+		const float g1 = s_grid[c][1];
+		const float dot_f = __half2float(x0_h) * g0 + __half2float(x1_h) * g1;
+		const __half dot_h = __float2half_rn(dot_f);
+		const __half twice_dot_h = __hmul(dot_h, two_h);
+		const __half score_h = __hsub(twice_dot_h, s_norms[c]);
+		const float score_f = __half2float(score_h);
+		if (score_f > best_score) {
+			best_score = score_f;
+			best_index = static_cast<uint8_t>(c);
+		}
+	}
+
+	out[row] = best_index;
+}
+
+extern "C" void higgs_quantize_2_256_ptr_f16_cuda_portable(
+	uint64_t x_ptr,
+	uint64_t grid_ptr,
+	uint64_t grid_norms_ptr,
+	uint64_t out_ptr,
+	int64_t out_dim)
+{
+	const __half* x = reinterpret_cast<const __half*>(x_ptr);
+	const __half* grid = reinterpret_cast<const __half*>(grid_ptr);
+	const __half* grid_norms = reinterpret_cast<const __half*>(grid_norms_ptr);
+	uint8_t* out = reinterpret_cast<uint8_t*>(out_ptr);
+
+	const int threads = 256;
+	const int blocks = static_cast<int>((out_dim + threads - 1) / threads);
+
+	auto stream = at::cuda::getCurrentCUDAStream();
+	higgs_quantize_2_256_ptr_f16_cuda_portable_kernel<<<blocks, threads, 0, stream>>>(x, grid, grid_norms, out, out_dim);
+
+    C10_CUDA_KERNEL_LAUNCH_CHECK();
+}

flake.nix

@@ -0,0 +1,17 @@
+{
+  description = "CUDA kernels for HIGGS quantization";
+
+  inputs = {
+    kernel-builder.url = "github:huggingface/kernels";
+  };
+
+  outputs =
+    {
+      self,
+      kernel-builder,
+    }:
+    kernel-builder.lib.genKernelFlakeOutputs {
+      inherit self;
+      path = ./.;
+    };
+}

torch-ext/higgs_kernels/__init__.py

@@ -0,0 +1,65 @@
+import os
+import functools
+
+import torch
+from ._ops import ops
+
+import safetensors.torch
+
+PKG_PATH = os.path.dirname(os.path.realpath(__file__))
+
+
+@functools.cache
+def load_optimal_grid_2_256(device="cpu", dtype=torch.float16):
+    return safetensors.torch.load_file(
+        os.path.join(PKG_PATH, "grids.safetensors"), device=device
+    )["2_256"].to(dtype)
+
+
+def higgs_dequantize_2_256_kernel(x, grid):
+    x = x.contiguous()
+    grid = grid.contiguous()
+
+    assert grid.device == x.device
+    assert "cuda" in str(x.device)
+
+    assert grid.dtype in (torch.float16, torch.bfloat16)
+    assert grid.shape == (256, 2)
+
+    assert x.dtype == torch.uint8
+    (out_dim,) = x.shape
+    assert out_dim > 0
+
+    assert x.data_ptr() % 4 == 0, "Input tensor must be 4-byte aligned for vectorized loads"
+
+    out = torch.zeros((out_dim, 2), dtype=grid.dtype, device=grid.device)
+
+    assert out.data_ptr() % 16 == 0, "Output tensor must be 16-byte aligned for vectorized stores"
+
+    ops.higgs_dequantize_2_256(x, grid, out)
+
+    return out
+
+
+def higgs_quantize_2_256_kernel(x, grid, grid_norms):
+    assert x.dtype == grid.dtype == grid_norms.dtype
+    assert x.device == grid.device == grid_norms.device
+    assert "cuda" in str(x.device)
+
+    assert x.dtype in (torch.bfloat16, torch.float16)
+
+    assert grid.shape == (256, 2), grid.shape
+    assert grid_norms.shape == (256,), grid_norms.shape
+
+    out_dim, in_dim = x.shape
+    assert in_dim == 2
+    assert out_dim > 0
+
+    out = torch.empty((out_dim,), dtype=torch.uint8, device=x.device)
+
+    if x.dtype == torch.bfloat16:
+        ops.higgs_quantize_2_256_bf16(x, grid, grid_norms, out)
+    else:
+        ops.higgs_quantize_2_256_f16(x, grid, grid_norms, out)
+
+    return out

torch-ext/higgs_kernels/grids.safetensors

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

torch-ext/torch_binding.cpp

@@ -0,0 +1,64 @@
+#include <torch/library.h>
+
+#include "registration.h"
+#include "torch_binding.h"
+
+extern "C" void higgs_dequantize_2_256_ptr_cuda_portable(uint64_t x_ptr,
+                                                         uint64_t grid_ptr,
+                                                         uint64_t out_ptr,
+                                                         int64_t out_dim);
+
+extern "C" void higgs_quantize_2_256_ptr_f16_cuda_portable(
+    uint64_t x_ptr, uint64_t grid_ptr, uint64_t grid_norms_ptr,
+    uint64_t out_ptr, int64_t out_dim);
+
+extern "C" void higgs_quantize_2_256_ptr_bf16_cuda_portable(
+    uint64_t x_ptr, uint64_t grid_ptr, uint64_t grid_norms_ptr,
+    uint64_t out_ptr, int64_t out_dim);
+
+void higgs_dequantize_2_256(torch::Tensor x, torch::Tensor grid,
+                            torch::Tensor out) {
+  int64_t out_dim = x.size(0);
+  higgs_dequantize_2_256_ptr_cuda_portable(
+      reinterpret_cast<uint64_t>(x.data_ptr()),
+      reinterpret_cast<uint64_t>(grid.data_ptr()),
+      reinterpret_cast<uint64_t>(out.data_ptr()), out_dim);
+}
+
+void higgs_quantize_2_256_f16(torch::Tensor x, torch::Tensor grid,
+                              torch::Tensor grid_norms, torch::Tensor out) {
+  int64_t out_dim = x.size(0);
+  higgs_quantize_2_256_ptr_f16_cuda_portable(
+      reinterpret_cast<uint64_t>(x.data_ptr()),
+      reinterpret_cast<uint64_t>(grid.data_ptr()),
+      reinterpret_cast<uint64_t>(grid_norms.data_ptr()),
+      reinterpret_cast<uint64_t>(out.data_ptr()), out_dim);
+}
+
+void higgs_quantize_2_256_bf16(torch::Tensor x, torch::Tensor grid,
+                               torch::Tensor grid_norms, torch::Tensor out) {
+  int64_t out_dim = x.size(0);
+  higgs_quantize_2_256_ptr_bf16_cuda_portable(
+      reinterpret_cast<uint64_t>(x.data_ptr()),
+      reinterpret_cast<uint64_t>(grid.data_ptr()),
+      reinterpret_cast<uint64_t>(grid_norms.data_ptr()),
+      reinterpret_cast<uint64_t>(out.data_ptr()), out_dim);
+}
+
+TORCH_LIBRARY_EXPAND(TORCH_EXTENSION_NAME, ops) {
+  ops.def(
+      "higgs_dequantize_2_256(Tensor x, Tensor grid, Tensor! out) -> ()");
+  ops.impl("higgs_dequantize_2_256", torch::kCUDA, &higgs_dequantize_2_256);
+
+  ops.def("higgs_quantize_2_256_f16(Tensor x, Tensor grid, Tensor "
+          "grid_norms, Tensor! out) -> ()");
+  ops.impl("higgs_quantize_2_256_f16", torch::kCUDA,
+           &higgs_quantize_2_256_f16);
+
+  ops.def("higgs_quantize_2_256_bf16(Tensor x, Tensor grid, Tensor "
+          "grid_norms, Tensor! out) -> ()");
+  ops.impl("higgs_quantize_2_256_bf16", torch::kCUDA,
+           &higgs_quantize_2_256_bf16);
+}
+
+REGISTER_EXTENSION(TORCH_EXTENSION_NAME)

torch-ext/torch_binding.h

@@ -0,0 +1,12 @@
+#pragma once
+
+#include <torch/types.h>
+
+void higgs_dequantize_2_256(torch::Tensor x, torch::Tensor grid,
+                            torch::Tensor out);
+
+void higgs_quantize_2_256_f16(torch::Tensor x, torch::Tensor grid,
+                              torch::Tensor grid_norms, torch::Tensor out);
+
+void higgs_quantize_2_256_bf16(torch::Tensor x, torch::Tensor grid,
+                               torch::Tensor grid_norms, torch::Tensor out);