Implementation:Deepspeedai DeepSpeed Dequantization Utils
| Knowledge Sources | |
|---|---|
| Domains | Quantization, Model_Compression, CUDA_Kernels |
| Last Updated | 2026-02-09 00:00 GMT |
Overview
Device-side dequantization framework for converting INT4/INT8 quantized tensors back to FP16 with group-wise scaling and offset application.
Description
The dequantization_utils.h header provides the inverse operations to quantization_utils.h, implementing efficient unpacking and reconstruction of quantized values. It reuses the Params and Type definitions from quantization for consistency, providing chunk-based processing that loads quantized INT4/INT8 data, applies group-specific dequantization parameters (scale and optional offset), and stores results as FP16 in global memory. The implementation handles 4-bit packed representations where two values share a byte using the PackedInt4 structure for bitwise extraction. The main to_global function orchestrates parallel dequantization across thread blocks with configurable unrolling for memory bandwidth optimization, supporting both symmetric (scale-only) and asymmetric (scale+offset) modes that match the original quantization configuration.
Usage
Use these utilities when implementing weight or activation dequantization for inference, or when decompressing gradients in distributed training. The group-wise dequantization preserves the accuracy characteristics of the original quantization while providing efficient memory-to-compute transformation.
Code Reference
Source Location
- Repository: DeepSpeed
- File: csrc/includes/dequantization_utils.h
Signature
namespace dequantize {
using Type = quantize::Type;
template <Type qType, int numBits>
using Params = quantize::Params<qType, numBits>;
// Dequantize chunk (16 bytes of quantized data)
template <typename T, int numBits, Type qType>
DS_D_INLINE void chunk(T* local_output, const int8_t* data,
Params<qType, numBits> q_params);
// Full dequantization to global memory
template <typename T, int numBits, Type qType, int unroll, int threads>
DS_D_INLINE void to_global(T* global_output,
const int8_t* data,
const float* global_params,
const int elems_per_group,
const int total_elems);
}
Import
#include "csrc/includes/dequantization_utils.h"
I/O Contract
| Input | Type | Description |
|---|---|---|
| data | const int8_t* | Packed quantized values (INT4/INT8) |
| global_params | const float* | Quantization parameters per group |
| elems_per_group | int | Group size used during quantization |
| total_elems | int | Total number of elements to dequantize |
| numBits | int (template) | 4 or 8 bits per quantized value |
| qType | Type (template) | Symmetric or Asymmetric mode |
| Output | Type | Description |
|---|---|---|
| global_output | T* | Dequantized FP16 values |
Usage Examples
8-bit Symmetric Dequantization:
__global__ void dequantize_int8_kernel(const int8_t* quantized,
const float* scales,
__half* output,
int total_elems,
int group_size) {
constexpr int unroll = 4;
constexpr int threads = 256;
dequantize::to_global<__half, 8, dequantize::Type::Symmetric, unroll, threads>(
output, quantized, scales, group_size, total_elems);
}
// Launch
int threads = 256;
int groups = (total_elems + group_size - 1) / group_size;
int elems_per_thread = unroll * 16; // unroll × 16 bytes
int blocks = (total_elems + threads * elems_per_thread - 1) /
(threads * elems_per_thread);
dequantize_int8_kernel<<<blocks, threads>>>(
quantized_data, scale_params, fp16_output, total_elems, group_size);
4-bit Asymmetric Dequantization:
__global__ void dequantize_int4_asym_kernel(const int8_t* packed_data,
const float* params,
__half* output,
int n, int group_size) {
// Each group has 2 params: scale and offset
constexpr int unroll = 8;
constexpr int threads = 256;
dequantize::to_global<__half, 4, dequantize::Type::Asymmetric, unroll, threads>(
output, packed_data, params, group_size, n);
}
// For 4-bit: 2 values per byte
int compressed_size = (total_elems + 1) / 2;
int param_size = (total_elems / group_size) * 2; // scale + offset per group
dequantize_int4_asym_kernel<<<blocks, threads>>>(
compressed_data, quantization_params, fp16_weights, total_elems, group_size);
Manual Chunk Processing:
__device__ void dequantize_chunk_example(const int8_t* quant_data,
const float* params,
__half* output,
int chunk_id, int group_id) {
// Load quantization parameters for this group
dequantize::Params<dequantize::Type::Symmetric, 8> q_params(params, group_id);
// Dequantize one chunk (16 bytes → 16 halfs)
__half local_buffer[16];
dequantize::chunk<__half, 8, dequantize::Type::Symmetric>(
local_buffer, quant_data + chunk_id * 16, q_params);
// Store to global memory
for (int i = 0; i < 16; i++) {
output[chunk_id * 16 + i] = local_buffer[i];
}
}
Weight Dequantization for Inference:
class QuantizedLinear {
int8_t* quantized_weight; // [out_features × in_features] quantized
float* weight_params; // Scales per group
__half* dequant_buffer; // Temporary FP16 buffer
int in_features, out_features, group_size;
public:
void dequantize_weights(cudaStream_t stream) {
int total_elems = in_features * out_features;
int groups = total_elems / group_size;
dim3 block(256);
dim3 grid((total_elems + 256 * 4 * 16 - 1) / (256 * 4 * 16));
dequantize_kernel<8><<<grid, block, 0, stream>>>(
quantized_weight, weight_params, dequant_buffer,
total_elems, group_size);
}
template <int numBits>
__global__ void dequantize_kernel(const int8_t* q_data,
const float* params,
__half* output,
int n, int gs) {
dequantize::to_global<__half, numBits,
dequantize::Type::Symmetric, 4, 256>(
output, q_data, params, gs, n);
}
};
Gradient Decompression:
void decompress_gradients_after_allreduce(
const int8_t* compressed_grads,
const float* grad_scales,
__half* decompressed_grads,
int total_params, int group_size,
cudaStream_t stream) {
constexpr int num_bits = 8;
constexpr int unroll = 4;
int threads = 256;
int blocks = (total_params + threads * unroll * 16 - 1) /
(threads * unroll * 16);
auto dequant_kernel = [] __global__ (
const int8_t* q_data, const float* scales,
__half* output, int n, int gs) {
dequantize::to_global<__half, num_bits,
dequantize::Type::Symmetric, unroll, 256>(
output, q_data, scales, gs, n);
};
dequant_kernel<<<blocks, threads, 0, stream>>>(
compressed_grads, grad_scales, decompressed_grads,
total_params, group_size);
}
Batched Dequantization:
__global__ void batch_dequantize_kernel(const int8_t** quantized_ptrs,
const float** param_ptrs,
__half** output_ptrs,
const int* elem_counts,
const int* group_sizes,
int num_tensors) {
int tensor_id = blockIdx.y;
if (tensor_id >= num_tensors) return;
int total = elem_counts[tensor_id];
int gs = group_sizes[tensor_id];
dequantize::to_global<__half, 8, dequantize::Type::Symmetric, 4, 256>(
output_ptrs[tensor_id],
quantized_ptrs[tensor_id],
param_ptrs[tensor_id],
gs, total);
}
// Launch with 2D grid: X for parallelism, Y for tensor selection
dim3 grid(compute_blocks(max_elems), num_tensors);
dim3 block(256);
batch_dequantize_kernel<<<grid, block>>>(q_ptrs, p_ptrs, o_ptrs,
counts, groups, num_tensors);
Related Pages
- Quantization Utils - Forward quantization operation
- Quantization Binding - PyTorch interface
- Conversion Utils - Type conversions used internally