Files
derek.holloway 7502bf64a4
CodeQL / Analyze (actions) (push) Canceled after 0s
CodeQL / Analyze (python) (push) Canceled after 0s
pre-commit / pre-commit (push) Canceled after 0s
Multi-Arch CI / setup (push) Canceled after 0s
Publish build_manylinux_x86_64 images / publish_build_manylinux_x86_64 (push) Canceled after 0s
Setup (Multi-Arch) / setup (push) Canceled after 0s
Multi-Arch CI / Linux::${{ fromJSON(needs.setup.outputs.linux_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubi10 image / publish_no_rocm_image_ubi10 (push) Canceled after 0s
Multi-Arch CI / Windows::${{ fromJSON(needs.setup.outputs.windows_build_config || '{}').build_variant_label || 'skip' }} (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Multi-Arch CI / Manifest Diff (push) Canceled after 0s
TheRock Manifest Diff Report / Generate Manifest Diff Report (push) Canceled after 0s
Multi-Arch CI / CI Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_ocl_rt images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04_rocgdb images / publish_no_rocm_image_ubuntu24_04_rocgdb (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_ocl_rt (push) Canceled after 0s
Publish TheRock Docker image / build-and-push-image (push) Canceled after 0s
Unit Tests / Unit Tests :: ubuntu-24.04 (push) Canceled after 0s
Unit Tests / Unit Tests :: windows-2022 (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_media (push) Canceled after 0s
Unit Tests / Unit Tests Summary (push) Canceled after 0s
Publish no_rocm_image_ubuntu24_04 images / publish_no_rocm_image_ubuntu24_04_openmpi (push) Canceled after 0s
Flatten subprojects into a single base monorepo
2026-09-02 16:34:47 -07:00

2.4 KiB

GEMM with Quantization

Theory

This example demonstrates GEMM (General Matrix Multiplication) with quantized inputs or weights. Quantization is a technique to reduce memory and computation by representing values with lower-precision integer types (e.g., int8), commonly used for efficient inference in deep learning.

Mathematical Formulation:

  • Quantized GEMM: C = \text{dequant}(A_q) \times \text{dequant}(B_q)
  • A_q, B_q: quantized matrices (e.g., int8)
  • \text{dequant}(x_q) = (x_q - z) \cdot s (scale s, zero-point z)
  • C: output matrix (often in higher precision, e.g., float32 or float16)

Algorithmic Background:

  • Quantized values are dequantized on-the-fly during GEMM computation.
  • Accumulation is performed in higher precision for accuracy.
  • Supports symmetric and asymmetric quantization.

How to Run

Prerequisites

Please follow the instructions in the main Build Guide section as a prerequisite to building and running this example.

Build and run

cd composable_kernel/example/14_gemm_quantization
mkdir build && cd build
cmake -DCMAKE_CXX_COMPILER=/opt/rocm/bin/hipcc ..
make -j

# Example run
./gemm_quantization_xdl --verify=1 --time=1

Source Code Structure

Directory Layout

example/14_gemm_quantization/
├── gemm_quantization_xdl.cpp         # Main example: sets up, runs, and verifies quantized GEMM
include/ck/tensor_operation/gpu/device/
│   └── device_gemm_quantized.hpp       # Device-level quantized GEMM API
include/ck/tensor_operation/gpu/device/impl/
│   └── device_gemm_quantized_impl.hpp  # Implementation
include/ck/tensor_operation/gpu/grid/
│   └── gridwise_gemm_quantized.hpp     # Grid-level quantized GEMM kernel
include/ck/tensor_operation/gpu/element/
    └── quantization_operations.hpp     # Quantization/dequantization utilities

Key Classes and Functions

  • DeviceGemmQuantized (in device_gemm_quantized.hpp):
    Device API for quantized GEMM.
  • gridwise_gemm_quantized (in gridwise_gemm_quantized.hpp):
    Implements the tiled/blocking quantized GEMM kernel.
  • quantization_operations (in quantization_operations.hpp):
    Defines quantization and dequantization functions.

This example demonstrates how Composable Kernel supports efficient quantized matrix multiplication for deep learning inference.