C-Kernel-Engine Parity Testing API Implementation. More...

#include "ck_parity_api.h"
#include "ckernel_quant.h"
#include <math.h>
#include <stdlib.h>
#include <string.h>

Functions
void	attention_forward_causal_head_major_gqa_flash_strided (const float q, const float k, const float v, float output, int num_heads, int num_kv_heads, int num_tokens, int head_dim, int aligned_head_dim, int kv_stride_tokens)

void	attn_gate_sigmoid_mul_forward (const float x, const float gate, float *out, int rows, int num_heads, int state_dim)

int	ck_get_block_q4_k_size (void)
	Get Q4_K block size in bytes.

int	ck_get_block_q6_k_size (void)
	Get Q6_K block size in bytes.

int	ck_get_block_q8_k_size (void)
	Get Q8_K block size in bytes.

int	ck_get_qk_k (void)
	Get QK_K (elements per super-block)

void	ck_test_attention_causal (const float q, const float k, const float v, float out, int num_heads, int num_kv_heads, int tokens, int seq_len, int head_dim)
	Multi-head causal attention for prefill (head-major layout)

void	ck_test_attn_gate_sigmoid_mul (const float x, const float gate, float *out, int rows, int dim)
	Multiply attention output rows by sigmoid(gate) elementwise.

void	ck_test_dequant_q4_0 (const void src, float dst, int n)
	Dequantize Q4_0 data to FP32.

void	ck_test_dequant_q4_k (const void src, float dst, int n)
	Dequantize Q4_K data to FP32.

void	ck_test_dequant_q6_k (const void src, float dst, int n)
	Dequantize Q6_K data to FP32.

void	ck_test_gated_deltanet_autoregressive (const float q, const float k, const float v, const float g, const float beta, const float state_in, float state_out, float out, int num_heads, int state_dim, float norm_eps)
	Gated DeltaNet autoregressive update.

void	ck_test_gemm_q4_k (const void weight_q4k, const float input_f32, float *output, int rows, int cols, int n_tokens)
	Q4_K GEMM - batched matrix multiply with quantized weights.

void	ck_test_gemm_q5_0 (const void weight_q5_0, const float input_f32, float *output, int rows, int cols, int n_tokens)
	Test Q5_0 x Q8_0 GEMM (batch matrix multiply)

void	ck_test_gemm_q6_k (const void weight_q6k, const float input_f32, float *output, int rows, int cols, int n_tokens)
	Test Q6_K x Q8_K GEMM (batch matrix multiply)

void	ck_test_gemm_q8_0 (const void weight_q8_0, const float input_f32, float *output, int rows, int cols, int n_tokens)
	Test Q8_0 x Q8_0 GEMM (batch matrix multiply)

void	ck_test_gemv_q4_k (const void weight_q4k, const float input_f32, float *output, int cols)
	Q4_K GEMV - dot product of quantized weights and FP32 input.

void	ck_test_gemv_q5_0 (const void weight_q5_0, const float input_f32, float *output, int rows, int cols)
	Q5_0 GEMV - matrix-vector multiply with Q5_0 weights.

void	ck_test_gemv_q5_0_q8_0 (const void weight_q5_0, const float input_f32, float *output, int rows, int cols)
	Q5_0 x Q8_0 quantized GEMV - matches llama.cpp's approach.

void	ck_test_gemv_q6_k (const void weight_q6k, const float input_f32, float *output, int cols)
	Q6_K GEMV.

void	ck_test_gemv_q8_0 (const void weight_q8_0, const float input_f32, float *output, int rows, int cols)
	Q8_0 GEMV - matrix-vector multiply with Q8_0 weights.

void	ck_test_gemv_q8_0_q8_0 (const void weight_q8_0, const float input_f32, float *output, int rows, int cols)
	Q8_0 x Q8_0 quantized GEMV - matches llama.cpp's approach.

void	ck_test_outproj_mlp_fused_q5_0 (const float attn_out, const float residual, const float ln2_gamma, const void wo, const void w1, const void w2, float *output, int tokens, int num_heads, int head_dim, int embed_dim, int intermediate, float eps, int w2_is_q6k)
	Test mega-fused OutProj + MLP kernel (Q5_0 weights)

void	ck_test_quantize_q8_k (const float src, void dst, int n)
	Quantize FP32 to Q8_K (for activations)

void	ck_test_recurrent_conv_state_update (const float state_in, const float q, const float k, const float v, float conv_x, float state_out, int history_len, int num_seqs, int num_tokens, int q_dim, int k_dim, int v_dim)
	Build the recurrent convolution input history window.

void	ck_test_recurrent_dt_gate (const float alpha, const float dt_bias, const float a, float gate, int rows, int dim)
	Transform recurrent alpha rows into the DeltaNet gate.

void	ck_test_recurrent_norm_gate (const float x, const float gate, const float weight, float out, int rows, int num_heads, int head_dim, float eps)
	Per-head RMSNorm followed by SiLU(z) gating for recurrent outputs.

void	ck_test_recurrent_qk_l2_norm (float q, float k, int rows, int q_dim, int k_dim, int head_dim, float eps)
	Apply per-head L2 normalization to recurrent Q/K rows in-place.

void	ck_test_recurrent_silu (const float x, float out, int rows, int dim)
	Apply SiLU elementwise to recurrent rows.

void	ck_test_recurrent_split_conv_qkv (const float packed_qkv, float q, float k, float v, int rows, int q_dim, int k_dim, int v_dim)
	Split the post-convolution recurrent packed QKV rows.

void	ck_test_recurrent_split_qkv (const float packed_qkv, float q, float k, float v, int rows, int q_dim, int k_dim, int v_dim)
	Split a packed recurrent QKV matrix into explicit Q, K, and V outputs.

void	ck_test_rmsnorm (const float input, const float weight, float *output, int n_tokens, int dim, float eps)
	RMSNorm.

void	ck_test_rope (float q, float k, int n_tokens, int n_heads, int n_heads_kv, int head_dim, int pos_offset, float theta)
	RoPE (Rotary Position Embedding)

void	ck_test_rope_interleaved (float q, float k, int n_tokens, int n_heads, int n_heads_kv, int head_dim, int pos_offset, float theta)
	RoPE with interleaved format (for llama.cpp compatibility)

void	ck_test_softmax (const float input, float output, int n)
	Softmax (simple, non-causal)

void	ck_test_split_q_gate (const float packed_qg, float q, float *gate, int rows, int q_dim, int gate_dim, int group_dim)
	Split a packed full-attention Q+gate matrix into Q rows and gate rows.

void	ck_test_ssm_conv1d (const float conv_x, const float kernel, float *out, int kernel_size, int num_channels, int num_tokens, int num_seqs)
	qwen3next/Qwen3.5 SSM causal depthwise convolution.

void	ck_test_swiglu (const float gate_up, float output, int n_tokens, int intermediate_dim)
	SwiGLU activation.

void	ck_test_vec_dot_q4_k_q8_k (const void weight_q4_k, const void input_q8_k, float *output, int cols)

void	ck_test_vec_dot_q5_0_q8_0 (const void weight_q5_0, const void input_q8_0, float *output, int cols)
	Direct Q5_0 x Q8_0 dot product test (takes pre-quantized Q8_0 input)

void	ck_test_vec_dot_q6_k_q8_k (const void weight_q6_k, const void input_q8_k, float *output, int cols)

void	ck_test_vec_dot_q8_0_q8_0 (const void weight_q8_0, const void input_q8_0, float *output, int cols)
	Direct Q8_0 x Q8_0 dot product test (takes pre-quantized Q8_0 input)

void	dequant_q4_0_row (const void src, float dst, size_t n_elements)
	Dequantize Q4_0 row (multiple blocks)

void	dequant_q4_k_row (const void src, float dst, size_t n_elements)
	Dequantize Q4_K row (multiple blocks)

void	dequant_q6_k_row (const void src, float dst, size_t n_elements)
	Dequantize Q6_K row (multiple blocks)

void	gated_deltanet_autoregressive_forward (const float q, const float k, const float v, const float g, const float beta, const float state_in, float state_out, float out, int num_heads, int state_dim, float norm_eps)

void	gemm_nt_q4_k_q8_k (const void A_q8, const void B, const float bias, float C, int M, int N, int K)

void	gemm_nt_q5_0_q8_0 (const void A_q8, const void B_q5, const float bias, float C, int M, int N, int K)
	Batch GEMM with Q5_0 weights and Q8_0 activations for prefill.

void	gemm_nt_q6_k_q8_k (const void A_q8, const void B, const float bias, float C, int M, int N, int K)

void	gemm_nt_q8_0_q8_0 (const void A_q8, const void B_q8, const float bias, float C, int M, int N, int K)
	gemm_nt_q8_0_q8_0 with optional bias (matches header signature)

void	gemv_q4_k_q8_k (float y, const void W, const void *x_q8, int M, int K)

void	gemv_q5_0 (float y, const void W, const float *x, int M, int K)
	Auto-dispatch GEMV for Q5_0 weights based on CPU features.

void	gemv_q5_0_q8_0 (float y, const void W, const void *x_q8, int M, int K)
	Matrix-vector multiply with Q5_0 weights and Q8_0 input.

void	gemv_q6_k_q8_k (float y, const void W, const void *x_q8, int M, int K)
	GEMV: y = W @ x where W is Q6_K and x is Q8_K.

void	gemv_q8_0 (float y, const void W, const float *x, int M, int K)
	Auto-dispatch GEMV for Q8_0 weights based on CPU features.

void	gemv_q8_0_q8_0 (float y, const void W, const void *x_q8, int M, int K)
	Matrix-vector multiply with Q8_0 weights and Q8_0 input.

void	mega_fused_outproj_mlp_prefill (float output, const float attn_out, const float residual, const float ln2_gamma, const void wo, const float bo, int wo_dt, const void w1, const float b1, int w1_dt, const void w2, const float b2, int w2_dt, int tokens, int embed_dim, int aligned_embed_dim, int num_heads, int aligned_head_dim, int intermediate_dim, int aligned_intermediate_dim, float eps, void *scratch)

size_t	mega_fused_outproj_mlp_prefill_scratch_size (int tokens, int aligned_embed_dim, int num_heads, int aligned_head_dim, int aligned_intermediate_dim)
	Get scratch buffer size for mega_fused_outproj_mlp_prefill.

void	quantize_row_q8_0 (const float x, void vy, int k)
	Quantize FP32 to Q8_0 format (scalar reference)

void	quantize_row_q8_k (const float x, void vy, int k)

void	recurrent_conv_state_update_forward (const float state_in, const float q, const float k, const float v, float conv_x, float state_out, int history_len, int num_seqs, int num_tokens, int q_dim, int k_dim, int v_dim)

void	recurrent_dt_gate_forward (const float alpha, const float dt_bias, const float a, float gate, int rows, int num_heads, int state_dim)

void	recurrent_norm_gate_forward (const float x, const float gate, const float weight, float out, int rows, int num_heads, int head_dim, float eps)

void	recurrent_qk_l2_norm_forward (float q, float k, int rows, int q_dim, int k_dim, int head_dim, float eps)

void	recurrent_silu_forward (const float x, float out, int rows, int dim)

void	recurrent_split_conv_qkv_forward (const float packed_qkv, float q, float k, float v, int rows, int q_dim, int k_dim, int v_dim)

void	recurrent_split_qkv_forward (const float packed_qkv, float q, float k, float v, int rows, int q_dim, int k_dim, int v_dim)

void	rmsnorm_forward (const float input, const float gamma, float output, float rstd_cache, int tokens, int d_model, int aligned_embed_dim, float eps)

void	rope_forward_qk (float q, float k, const float cos_cache, const float sin_cache, int num_heads, int num_kv_heads, int num_tokens, int head_dim, int aligned_head_dim, int pos_offset)

void	rope_precompute_cache (float cos_cache, float sin_cache, int max_seq_len, int head_dim, float base, int rotary_dim, const char *scaling_type, float scaling_factor)

void	split_q_gate_forward (const float packed_qg, float q, float *gate, int rows, int q_dim, int gate_dim, int group_dim)

void	ssm_conv1d_forward (const float conv_x, const float kernel, float *out, int kernel_size, int num_channels, int num_tokens, int num_seqs)

void	swiglu_forward (const float input, float output, int tokens, int dim)

void	vec_dot_q5_0_q8_0 (int n, float s, const void vx, const void *vy)
	Auto-dispatch quantized dot product Q5_0 x Q8_0.

void	vec_dot_q6_k_q8_k (int n, float s, const void vx, const void *vy)
	Q6_K x Q8_K dot product (single row)

void	vec_dot_q8_0_q8_0 (int n, float s, const void vx, const void *vy)
	Auto-dispatch quantized dot product Q8_0 x Q8_0.

Detailed Description

C-Kernel-Engine Parity Testing API Implementation.

Wraps CK kernels for parity testing against llama.cpp/ggml.

Definition in file ck_parity_api.c.

Function Documentation

◆ attention_forward_causal_head_major_gqa_flash_strided()

void attention_forward_causal_head_major_gqa_flash_strided	(	const float *	q,
		const float *	k,
		const float *	v,
		float *	output,
		int	num_heads,
		int	num_kv_heads,
		int	num_tokens,
		int	head_dim,
		int	aligned_head_dim,
		int	kv_stride_tokens
	)

extern

Flash attention forward with custom KV stride (for KV cache)

Test:

test_flash_attention.py::TestFlashAttention::test_flash_strided

test_kv_cache_attention.py::TestKVCacheAttention::test_flash_attention

Variant with configurable kv_stride_tokens for KV cache layouts where K/V may not be contiguous in memory.

After changes: make test

Definition at line 4496 of file attention_kernels.c.

{
    attention_forward_head_major_gqa_flash_impl(q, k, v, output,
                                                num_heads, num_kv_heads,
                                                num_tokens, head_dim,
                                                aligned_head_dim,
                                                kv_stride_tokens,
                                                /*causal=*/1,
                                                /*round_full_kv_fp16=*/0,
                                                /*output_token_major=*/0,
                                                1.0f / sqrtf((float)head_dim));
}

Referenced by ck_test_attention_causal().

◆ attn_gate_sigmoid_mul_forward()

void attn_gate_sigmoid_mul_forward	(	const float *	x,
		const float *	gate,
		float *	out,
		int	rows,
		int	num_heads,
		int	state_dim
	)

extern

Definition at line 154 of file hybrid_attention_kernels.c.

                                                  {
    if (!x || !gate || !out || rows <= 0 || num_heads <= 0 ||
        state_dim <= 0 || num_heads > INT_MAX / state_dim) {
        return;
    }
    const int dim = num_heads * state_dim;
    const size_t elements = (size_t)rows * (size_t)dim;
    ck_attn_gate_sigmoid_mul_args_t args = {x, gate, out, dim};
    const size_t minimum_elements_per_thread = 32768u;
 
    const int partial_x_overlap = out != x &&
        hybrid_ranges_overlap(out, x, elements);
    const int gate_overlap = hybrid_ranges_overlap(out, gate, elements);
    if (partial_x_overlap || gate_overlap ||
        elements < 2u * minimum_elements_per_thread) {
        attn_gate_sigmoid_mul_row_range(0, rows, &args);
        return;
    }
 
    ck_threadpool_t *pool = ck_threadpool_global();
    int active = pool ? ck_threadpool_n_threads(pool) : 1;
    size_t useful_threads =
        (elements + minimum_elements_per_thread - 1u) /
        minimum_elements_per_thread;
    if ((size_t)active > useful_threads) {
        active = (int)useful_threads;
    }
    if (active > rows) {
        active = rows;
    }
 
    if (active > 1 && ck_threadpool_thread_id(pool) <= 0) {
        ck_threadpool_parallel_for_n(
            pool, active, 0, rows, 4,
            attn_gate_sigmoid_mul_row_range, &args);
    } else {
        attn_gate_sigmoid_mul_row_range(0, rows, &args);
    }
}

Referenced by ck_test_attn_gate_sigmoid_mul().

◆ ck_get_block_q4_k_size()

int ck_get_block_q4_k_size ( void )

Get Q4_K block size in bytes.

Definition at line 1026 of file ck_parity_api.c.

{
    return sizeof(block_q4_K);
}

◆ ck_get_block_q6_k_size()

int ck_get_block_q6_k_size ( void )

Get Q6_K block size in bytes.

Definition at line 1031 of file ck_parity_api.c.

{
    return sizeof(block_q6_K);
}

◆ ck_get_block_q8_k_size()

int ck_get_block_q8_k_size ( void )

Get Q8_K block size in bytes.

Definition at line 1036 of file ck_parity_api.c.

{
    return sizeof(block_q8_K);
}

◆ ck_get_qk_k()

int ck_get_qk_k ( void )

Get QK_K (elements per super-block)

Definition at line 1041 of file ck_parity_api.c.

{
    return QK_K;
}

References QK_K.

◆ ck_test_attention_causal()

void ck_test_attention_causal	(	const float *	q,
		const float *	k,
		const float *	v,
		float *	out,
		int	num_heads,
		int	num_kv_heads,
		int	tokens,
		int	seq_len,
		int	head_dim
	)

Multi-head causal attention for prefill (head-major layout)

Layout (head-major, matches llama.cpp test): Q: [num_heads, tokens, head_dim] K: [num_kv_heads, seq_len, head_dim] V: [num_kv_heads, seq_len, head_dim] out: [num_heads, tokens, head_dim]

Supports GQA (grouped-query attention) where num_heads > num_kv_heads. Causal masking: token t can only attend to positions 0..t (inclusive).

Parameters

q	Query [num_heads, tokens, head_dim]
k	Key [num_kv_heads, seq_len, head_dim]
v	Value [num_kv_heads, seq_len, head_dim]
out	Output [num_heads, tokens, head_dim]
num_heads	Number of query heads
num_kv_heads	Number of key/value heads (for GQA)
tokens	Number of query tokens
seq_len	Key/value sequence length (for prefill: seq_len == tokens)
head_dim	Dimension per head

Definition at line 884 of file ck_parity_api.c.

{
    /* For prefill, seq_len == tokens, and kv_stride == tokens.
     * The CK kernel expects strided KV layout with kv_stride_tokens parameter.
     * For parity testing with contiguous tensors, kv_stride = seq_len.
     */
    attention_forward_causal_head_major_gqa_flash_strided(
        q, k, v, out,
        num_heads, num_kv_heads, tokens,
        head_dim, head_dim,  /* aligned_head_dim = head_dim for testing */
        seq_len              /* kv_stride_tokens = seq_len for contiguous KV */
    );
}

References attention_forward_causal_head_major_gqa_flash_strided().

◆ ck_test_attn_gate_sigmoid_mul()

void ck_test_attn_gate_sigmoid_mul	(	const float *	x,
		const float *	gate,
		float *	out,
		int	rows,
		int	dim
	)

Multiply attention output rows by sigmoid(gate) elementwise.

Layout: x : [rows, dim] gate : [rows, dim] out : [rows, dim]

Definition at line 859 of file ck_parity_api.c.

{
    attn_gate_sigmoid_mul_forward(x, gate, out, rows, 1, dim);
}

References attn_gate_sigmoid_mul_forward().

◆ ck_test_dequant_q4_0()

void ck_test_dequant_q4_0	(	const void *	src,
		float *	dst,
		int	n
	)

Dequantize Q4_0 data to FP32.

Definition at line 202 of file ck_parity_api.c.

{
    dequant_q4_0_row(src, dst, (size_t)n);
}

References dequant_q4_0_row().

◆ ck_test_dequant_q4_k()

void ck_test_dequant_q4_k	(	const void *	src,
		float *	dst,
		int	n
	)

Dequantize Q4_K data to FP32.

Parameters

src	Input Q4_K blocks
dst	Output FP32 values
n	Number of elements (must be multiple of 256)

Definition at line 192 of file ck_parity_api.c.

{
    dequant_q4_k_row(src, dst, (size_t)n);
}

References dequant_q4_k_row().

◆ ck_test_dequant_q6_k()

void ck_test_dequant_q6_k	(	const void *	src,
		float *	dst,
		int	n
	)

Dequantize Q6_K data to FP32.

Definition at line 197 of file ck_parity_api.c.

{
    dequant_q6_k_row(src, dst, (size_t)n);
}

References dequant_q6_k_row().

◆ ck_test_gated_deltanet_autoregressive()

void ck_test_gated_deltanet_autoregressive	(	const float *	q,
		const float *	k,
		const float *	v,
		const float *	g,
		const float *	beta,
		const float *	state_in,
		float *	state_out,
		float *	out,
		int	num_heads,
		int	state_dim,
		float	norm_eps
	)

Gated DeltaNet autoregressive update.

Layout: q, k, v [num_heads, state_dim] g, beta [num_heads] state_* [num_heads, state_dim, state_dim] row-major per head out [num_heads, state_dim]

This mirrors the single-token recurrent update used by qwen3next in llama.cpp after projections/convolution but before output projection.

Definition at line 741 of file ck_parity_api.c.

{
    gated_deltanet_autoregressive_forward(
        q,
        k,
        v,
        g,
        beta,
        state_in,
        state_out,
        out,
        num_heads,
        state_dim,
        norm_eps);
}

References gated_deltanet_autoregressive_forward().

◆ ck_test_gemm_q4_k()

void ck_test_gemm_q4_k	(	const void *	weight_q4k,
		const float *	input_f32,
		float *	output,
		int	rows,
		int	cols,
		int	n_tokens
	)

Q4_K GEMM - batched matrix multiply with quantized weights.

Computes: output[t,r] = sum_k(weight[r,k] * input[t,k])

Parameters

weight_q4k	Q4_K quantized weights [rows, cols]
input_f32	FP32 input [n_tokens, cols]
output	FP32 output [n_tokens, rows]
rows	Number of output rows
cols	Number of columns (must be multiple of 256)
n_tokens	Batch size

Definition at line 434 of file ck_parity_api.c.

{
    /* Allocate Q8_K buffer for quantized activations */
    int n_blocks_per_row = cols / CK_QK_K;
    block_q8_K *q8_data = (block_q8_K *)malloc(n_tokens * n_blocks_per_row * sizeof(block_q8_K));
    if (!q8_data) {
        memset(output, 0, n_tokens * rows * sizeof(float));
        return;
    }
 
    /* Quantize all input tokens */
    for (int t = 0; t < n_tokens; t++) {
        quantize_row_q8_k(input_f32 + t * cols,
                          q8_data + t * n_blocks_per_row, cols);
    }
 
    /* Use gemm_nt_q4_k_q8_k: C[M,N] = A[M,K] * B[N,K]^T
     * Our layout: output[n_tokens, rows] = input[n_tokens, cols] * weight[rows, cols]^T
     * So: M = n_tokens, N = rows, K = cols
     */
    gemm_nt_q4_k_q8_k(q8_data, weight_q4k, NULL, output, n_tokens, rows, cols);
 
    free(q8_data);
}

References CK_QK_K, gemm_nt_q4_k_q8_k(), and quantize_row_q8_k().

◆ ck_test_gemm_q5_0()

void ck_test_gemm_q5_0	(	const void *	weight_q5_0,
		const float *	input_f32,
		float *	output,
		int	rows,
		int	cols,
		int	n_tokens
	)

Test Q5_0 x Q8_0 GEMM (batch matrix multiply)

Q5_0 GEMM - batched matrix multiply with Q5_0 weights (32-element blocks)

Used for MLP W1 (gate/up projection) and attention Q/K with Q5_0 weights.

Definition at line 533 of file ck_parity_api.c.

{
    /* Allocate Q8_0 buffer for quantized activations */
    int n_blocks_per_row = cols / CK_QK8_0;
    block_q8_0 *q8_data = (block_q8_0 *)malloc(n_tokens * n_blocks_per_row * sizeof(block_q8_0));
    if (!q8_data) {
        memset(output, 0, n_tokens * rows * sizeof(float));
        return;
    }
 
    /* Quantize all input tokens */
    for (int t = 0; t < n_tokens; t++) {
        quantize_row_q8_0(input_f32 + t * cols,
                          q8_data + t * n_blocks_per_row, cols);
    }
 
    /* Use gemm_nt_q5_0_q8_0: C[M,N] = A[M,K] * B[N,K]^T
     * Our layout: output[n_tokens, rows] = input[n_tokens, cols] * weight[rows, cols]^T
     * So: M = n_tokens, N = rows, K = cols
     */
    gemm_nt_q5_0_q8_0(q8_data, weight_q5_0, NULL, output, n_tokens, rows, cols);
 
    free(q8_data);
}

References CK_QK8_0, gemm_nt_q5_0_q8_0(), and quantize_row_q8_0().

◆ ck_test_gemm_q6_k()

void ck_test_gemm_q6_k	(	const void *	weight_q6k,
		const float *	input_f32,
		float *	output,
		int	rows,
		int	cols,
		int	n_tokens
	)

Test Q6_K x Q8_K GEMM (batch matrix multiply)

Q6_K GEMM - batched matrix multiply with Q6_K weights.

Used for MLP W2 (down projection) with Q6_K weights.

Definition at line 467 of file ck_parity_api.c.

{
    /* Allocate Q8_K buffer for quantized activations */
    int n_blocks_per_row = cols / CK_QK_K;
    block_q8_K *q8_data = (block_q8_K *)malloc(n_tokens * n_blocks_per_row * sizeof(block_q8_K));
    if (!q8_data) {
        memset(output, 0, n_tokens * rows * sizeof(float));
        return;
    }
 
    /* Quantize all input tokens */
    for (int t = 0; t < n_tokens; t++) {
        quantize_row_q8_k(input_f32 + t * cols,
                          q8_data + t * n_blocks_per_row, cols);
    }
 
    /* Use gemm_nt_q6_k_q8_k: C[M,N] = A[M,K] * B[N,K]^T
     * Our layout: output[n_tokens, rows] = input[n_tokens, cols] * weight[rows, cols]^T
     * So: M = n_tokens, N = rows, K = cols
     */
    gemm_nt_q6_k_q8_k(q8_data, weight_q6k, NULL, output, n_tokens, rows, cols);
 
    free(q8_data);
}

References CK_QK_K, gemm_nt_q6_k_q8_k(), and quantize_row_q8_k().

◆ ck_test_gemm_q8_0()

void ck_test_gemm_q8_0	(	const void *	weight_q8_0,
		const float *	input_f32,
		float *	output,
		int	rows,
		int	cols,
		int	n_tokens
	)

Test Q8_0 x Q8_0 GEMM (batch matrix multiply)

Q8_0 GEMM - batched matrix multiply with Q8_0 weights (32-element blocks)

Used for attention V projection with Q8_0 weights.

Definition at line 500 of file ck_parity_api.c.

{
    /* Allocate Q8_0 buffer for quantized activations */
    int n_blocks_per_row = cols / CK_QK8_0;
    block_q8_0 *q8_data = (block_q8_0 *)malloc(n_tokens * n_blocks_per_row * sizeof(block_q8_0));
    if (!q8_data) {
        memset(output, 0, n_tokens * rows * sizeof(float));
        return;
    }
 
    /* Quantize all input tokens */
    for (int t = 0; t < n_tokens; t++) {
        quantize_row_q8_0(input_f32 + t * cols,
                          q8_data + t * n_blocks_per_row, cols);
    }
 
    /* Use gemm_nt_q8_0_q8_0: C[M,N] = A[M,K] * B[N,K]^T
     * Our layout: output[n_tokens, rows] = input[n_tokens, cols] * weight[rows, cols]^T
     * So: M = n_tokens, N = rows, K = cols
     */
    gemm_nt_q8_0_q8_0(q8_data, weight_q8_0, NULL, output, n_tokens, rows, cols);
 
    free(q8_data);
}

References CK_QK8_0, gemm_nt_q8_0_q8_0(), and quantize_row_q8_0().

◆ ck_test_gemv_q4_k()

void ck_test_gemv_q4_k	(	const void *	weight_q4k,
		const float *	input_f32,
		float *	output,
		int	cols
	)

Q4_K GEMV - dot product of quantized weights and FP32 input.

Internally quantizes input to Q8_K, then computes dot product.

Parameters

weight_q4k	Q4_K quantized weights [cols]
input_f32	FP32 input vector [cols]
output	Output scalar [1]
cols	Number of columns (must be multiple of 256)

Definition at line 220 of file ck_parity_api.c.

{
    /* Allocate Q8_K buffer for quantized activations */
    int n_blocks = cols / CK_QK_K;
    block_q8_K *q8_data = (block_q8_K *)malloc(n_blocks * sizeof(block_q8_K));
    if (!q8_data) {
        *output = 0.0f;
        return;
    }
 
    /* Quantize input to Q8_K */
    quantize_row_q8_k(input_f32, q8_data, cols);
 
    /* Compute dot product using GEMV with M=1 */
    gemv_q4_k_q8_k(output, weight_q4k, q8_data, 1, cols);
 
    free(q8_data);
}

References CK_QK_K, gemv_q4_k_q8_k(), and quantize_row_q8_k().

◆ ck_test_gemv_q5_0()

void ck_test_gemv_q5_0	(	const void *	weight_q5_0,
		const float *	input_f32,
		float *	output,
		int	rows,
		int	cols
	)

Q5_0 GEMV - matrix-vector multiply with Q5_0 weights.

Parameters

weight_q5_0	Q5_0 quantized weights [rows * cols]
input_f32	FP32 input vector [cols]
output	FP32 output vector [rows]
rows	Number of output rows
cols	Number of columns (must be multiple of 32)

Definition at line 263 of file ck_parity_api.c.

{
    /* Match llama.cpp's test_gemv_q5_0:
     * 1. Quantize input to Q8_0 format
     * 2. Use quantized dot product (vec_dot_q5_0_q8_0)
     *
     * This ensures parity with llama.cpp which always uses the
     * quantized path, NOT the FP32 dequantization path.
     */
    int n_blocks = cols / CK_QK8_0;
    block_q8_0 *q8_data = (block_q8_0 *)malloc(n_blocks * sizeof(block_q8_0));
    if (!q8_data) {
        for (int r = 0; r < rows; r++) output[r] = 0.0f;
        return;
    }
 
    /* Quantize input to Q8_0 */
    quantize_row_q8_0(input_f32, q8_data, cols);
 
    /* Call the quantized GEMV kernel (same as ck_test_gemv_q5_0_q8_0) */
    gemv_q5_0_q8_0(output, weight_q5_0, q8_data, rows, cols);
 
    free(q8_data);
}

References CK_QK8_0, gemv_q5_0_q8_0(), and quantize_row_q8_0().

◆ ck_test_gemv_q5_0_q8_0()

void ck_test_gemv_q5_0_q8_0	(	const void *	weight_q5_0,
		const float *	input_f32,
		float *	output,
		int	rows,
		int	cols
	)

Q5_0 x Q8_0 quantized GEMV - matches llama.cpp's approach.

This version quantizes the input to Q8_0 first, then uses integer dot products (like llama.cpp does). Use this for parity testing.

Parameters

weight_q5_0	Q5_0 quantized weights [rows * cols]
input_f32	FP32 input vector [cols] - will be quantized to Q8_0
output	FP32 output vector [rows]
rows	Number of output rows
cols	Number of columns (must be multiple of 32)

Definition at line 319 of file ck_parity_api.c.

{
    /* This matches llama.cpp's approach:
     * 1. Quantize input to Q8_0 format
     * 2. Use quantized dot product (integer math)
     * 3. Scale at the end
     */
    int n_blocks = cols / CK_QK8_0;
    block_q8_0 *q8_data = (block_q8_0 *)malloc(n_blocks * sizeof(block_q8_0));
    if (!q8_data) {
        for (int r = 0; r < rows; r++) output[r] = 0.0f;
        return;
    }
 
    /* Quantize input to Q8_0 */
    quantize_row_q8_0(input_f32, q8_data, cols);
 
    /* Call the quantized GEMV kernel */
    gemv_q5_0_q8_0(output, weight_q5_0, q8_data, rows, cols);
 
    free(q8_data);
}

References CK_QK8_0, gemv_q5_0_q8_0(), and quantize_row_q8_0().

◆ ck_test_gemv_q6_k()

void ck_test_gemv_q6_k	(	const void *	weight_q6k,
		const float *	input_f32,
		float *	output,
		int	cols
	)

Q6_K GEMV.

Definition at line 242 of file ck_parity_api.c.

{
    /* Match runtime decode path:
     * 1) quantize FP32 activation to Q8_K
     * 2) run Q6_K x Q8_K GEMV kernel with M=1
     */
    int n_blocks = cols / CK_QK_K;
    block_q8_K *q8_data = (block_q8_K *)malloc(n_blocks * sizeof(block_q8_K));
    if (!q8_data) {
        *output = 0.0f;
        return;
    }
 
    quantize_row_q8_k(input_f32, q8_data, cols);
    gemv_q6_k_q8_k(output, weight_q6k, q8_data, 1, cols);
    free(q8_data);
}

References CK_QK_K, gemv_q6_k_q8_k(), and quantize_row_q8_k().

◆ ck_test_gemv_q8_0()

void ck_test_gemv_q8_0	(	const void *	weight_q8_0,
		const float *	input_f32,
		float *	output,
		int	rows,
		int	cols
	)

Q8_0 GEMV - matrix-vector multiply with Q8_0 weights.

Parameters

weight_q8_0	Q8_0 quantized weights [rows * cols]
input_f32	FP32 input vector [cols]
output	FP32 output vector [rows]
rows	Number of output rows
cols	Number of columns (must be multiple of 32)

Definition at line 291 of file ck_parity_api.c.

{
    /* Match llama.cpp's test_gemv_q8_0:
     * 1. Quantize input to Q8_0 format
     * 2. Use quantized dot product (vec_dot_q8_0_q8_0)
     *
     * This ensures parity with llama.cpp which always uses the
     * quantized path, NOT the FP32 dequantization path.
     */
    int n_blocks = cols / CK_QK8_0;
    block_q8_0 *q8_data = (block_q8_0 *)malloc(n_blocks * sizeof(block_q8_0));
    if (!q8_data) {
        for (int r = 0; r < rows; r++) output[r] = 0.0f;
        return;
    }
 
    /* Quantize input to Q8_0 */
    quantize_row_q8_0(input_f32, q8_data, cols);
 
    /* Call the quantized GEMV kernel (same as ck_test_gemv_q8_0_q8_0) */
    gemv_q8_0_q8_0(output, weight_q8_0, q8_data, rows, cols);
 
    free(q8_data);
}

References CK_QK8_0, gemv_q8_0_q8_0(), and quantize_row_q8_0().

◆ ck_test_gemv_q8_0_q8_0()

void ck_test_gemv_q8_0_q8_0	(	const void *	weight_q8_0,
		const float *	input_f32,
		float *	output,
		int	rows,
		int	cols
	)

Q8_0 x Q8_0 quantized GEMV - matches llama.cpp's approach.

This version quantizes the input to Q8_0 first, then uses integer dot products (like llama.cpp does). Use this for parity testing.

Parameters

weight_q8_0	Q8_0 quantized weights [rows * cols]
input_f32	FP32 input vector [cols] - will be quantized to Q8_0
output	FP32 output vector [rows]
rows	Number of output rows
cols	Number of columns (must be multiple of 32)

Definition at line 345 of file ck_parity_api.c.

{
    /* This matches llama.cpp's approach:
     * 1. Quantize input to Q8_0 format
     * 2. Use quantized dot product (integer math)
     * 3. Scale at the end
     */
    int n_blocks = cols / CK_QK8_0;
    block_q8_0 *q8_data = (block_q8_0 *)malloc(n_blocks * sizeof(block_q8_0));
    if (!q8_data) {
        for (int r = 0; r < rows; r++) output[r] = 0.0f;
        return;
    }
 
    /* Quantize input to Q8_0 */
    quantize_row_q8_0(input_f32, q8_data, cols);
 
    /* Call the quantized GEMV kernel */
    gemv_q8_0_q8_0(output, weight_q8_0, q8_data, rows, cols);
 
    free(q8_data);
}

References CK_QK8_0, gemv_q8_0_q8_0(), and quantize_row_q8_0().

◆ ck_test_outproj_mlp_fused_q5_0()

void ck_test_outproj_mlp_fused_q5_0	(	const float *	attn_out,
		const float *	residual,
		const float *	ln2_gamma,
		const void *	wo,
		const void *	w1,
		const void *	w2,
		float *	output,
		int	tokens,
		int	num_heads,
		int	head_dim,
		int	embed_dim,
		int	intermediate,
		float	eps,
		int	w2_is_q6k
	)

Test mega-fused OutProj + MLP kernel (Q5_0 weights)

This is a simplified wrapper for parity testing that:

Uses Q5_0 for W_o and W1 weights
Uses Q4_K for W2 weights
Allocates scratch internally

Parameters

attn_out	Attention output [num_heads, tokens, head_dim] (FP32, head-major)
residual	Residual input [tokens, embed_dim] (FP32)
ln2_gamma	RMSNorm gamma [embed_dim] (FP32)
wo	OutProj weights [embed_dim, embed_dim] (Q5_0)
w1	MLP W1 weights [2*intermediate, embed_dim] (Q5_0)
w2	MLP W2 weights [embed_dim, intermediate] (Q4_K or Q6_K)
output	Output [tokens, embed_dim] (FP32)
tokens	Number of tokens
num_heads	Number of attention heads
head_dim	Dimension per head
embed_dim	Embedding dimension (= num_heads * head_dim)
intermediate	MLP intermediate dimension
eps	RMSNorm epsilon
w2_is_q6k	If true, W2 is Q6_K; if false, W2 is Q4_K

Definition at line 959 of file ck_parity_api.c.

{
    /* CK uses dtype enum: CK_DT_Q5_0 = 11, CK_DT_Q4_K = 7, CK_DT_Q6_K = 8 */
    const int CK_DT_Q5_0_VAL = 11;
    const int CK_DT_Q4_K_VAL = 7;
    const int CK_DT_Q6_K_VAL = 8;
 
    /* For parity testing, aligned = actual (no padding) */
    int aligned_embed_dim = embed_dim;
    int aligned_head_dim = head_dim;
    int aligned_intermediate = intermediate;
 
    /* Ensure intermediate is multiple of 256 (QK_K) for K-quants */
    if ((intermediate % 256) != 0) {
        aligned_intermediate = ((intermediate + 255) / 256) * 256;
    }
 
    /* Allocate scratch */
    size_t scratch_size = mega_fused_outproj_mlp_prefill_scratch_size(
        tokens, aligned_embed_dim, num_heads, aligned_head_dim, aligned_intermediate);
 
    void *scratch = malloc(scratch_size);
    if (!scratch) {
        return;
    }
 
    /* Call the mega-fused kernel */
    mega_fused_outproj_mlp_prefill(
        output,
        attn_out,
        residual,
        ln2_gamma,
        wo, NULL, CK_DT_Q5_0_VAL,          /* W_o with Q5_0 */
        w1, NULL, CK_DT_Q5_0_VAL,          /* W1 with Q5_0 */
        w2, NULL, w2_is_q6k ? CK_DT_Q6_K_VAL : CK_DT_Q4_K_VAL,  /* W2 with Q4_K or Q6_K */
        tokens,
        embed_dim,
        aligned_embed_dim,
        num_heads,
        aligned_head_dim,
        intermediate,
        aligned_intermediate,
        eps,
        scratch
    );
 
    free(scratch);
}

References mega_fused_outproj_mlp_prefill(), and mega_fused_outproj_mlp_prefill_scratch_size().

◆ ck_test_quantize_q8_k()

void ck_test_quantize_q8_k	(	const float *	src,
		void *	dst,
		int	n
	)

Quantize FP32 to Q8_K (for activations)

Parameters

src	Input FP32 values
dst	Output Q8_K blocks
n	Number of elements (must be multiple of 256)

Definition at line 211 of file ck_parity_api.c.

{
    quantize_row_q8_k(src, dst, n);
}

References quantize_row_q8_k().

◆ ck_test_recurrent_conv_state_update()

void ck_test_recurrent_conv_state_update	(	const float *	state_in,
		const float *	q,
		const float *	k,
		const float *	v,
		float *	conv_x,
		float *	state_out,
		int	history_len,
		int	num_seqs,
		int	num_tokens,
		int	q_dim,
		int	k_dim,
		int	v_dim
	)

Build the recurrent convolution input history window.

Layout: state_in [num_seqs, channels, history_len] q [num_seqs * num_tokens, q_dim] k [num_seqs * num_tokens, k_dim] v [num_seqs * num_tokens, v_dim] conv_x [num_seqs, channels, history_len + num_tokens] state_out [num_seqs, channels, history_len]

Definition at line 811 of file ck_parity_api.c.

{
    recurrent_conv_state_update_forward(
        state_in, q, k, v, conv_x, state_out, history_len, num_seqs, num_tokens, q_dim, k_dim, v_dim);
}

References recurrent_conv_state_update_forward().

◆ ck_test_recurrent_dt_gate()

void ck_test_recurrent_dt_gate	(	const float *	alpha,
		const float *	dt_bias,
		const float *	a,
		float *	gate,
		int	rows,
		int	dim
	)

Transform recurrent alpha rows into the DeltaNet gate.

Layout: alpha : [rows, dim] dt_bias : [dim] a : [dim] gate : [rows, dim]

Definition at line 801 of file ck_parity_api.c.

{
    recurrent_dt_gate_forward(alpha, dt_bias, a, gate, rows, 1, dim);
}

References recurrent_dt_gate_forward().

◆ ck_test_recurrent_norm_gate()

void ck_test_recurrent_norm_gate	(	const float *	x,
		const float *	gate,
		const float *	weight,
		float *	out,
		int	rows,
		int	num_heads,
		int	head_dim,
		float	eps
	)

Per-head RMSNorm followed by SiLU(z) gating for recurrent outputs.

Definition at line 868 of file ck_parity_api.c.

{
    recurrent_norm_gate_forward(x, gate, weight, out, rows, num_heads, head_dim, eps);
}

References recurrent_norm_gate_forward().

◆ ck_test_recurrent_qk_l2_norm()

void ck_test_recurrent_qk_l2_norm	(	float *	q,
		float *	k,
		int	rows,
		int	q_dim,
		int	k_dim,
		int	head_dim,
		float	eps
	)

Apply per-head L2 normalization to recurrent Q/K rows in-place.

Layout: q : [rows, q_dim] k : [rows, k_dim] where q_dim and k_dim are multiples of head_dim.

Definition at line 848 of file ck_parity_api.c.

{
    recurrent_qk_l2_norm_forward(q, k, rows, q_dim, k_dim, head_dim, eps);
}

References recurrent_qk_l2_norm_forward().

◆ ck_test_recurrent_silu()

void ck_test_recurrent_silu	(	const float *	x,
		float *	out,
		int	rows,
		int	dim
	)

Apply SiLU elementwise to recurrent rows.

Definition at line 828 of file ck_parity_api.c.

{
    recurrent_silu_forward(x, out, rows, dim);
}

References recurrent_silu_forward().

◆ ck_test_recurrent_split_conv_qkv()

void ck_test_recurrent_split_conv_qkv	(	const float *	packed_qkv,
		float *	q,
		float *	k,
		float *	v,
		int	rows,
		int	q_dim,
		int	k_dim,
		int	v_dim
	)

Split the post-convolution recurrent packed QKV rows.

Definition at line 836 of file ck_parity_api.c.

{
    recurrent_split_conv_qkv_forward(packed_qkv, q, k, v, rows, q_dim, k_dim, v_dim);
}

References recurrent_split_conv_qkv_forward().

◆ ck_test_recurrent_split_qkv()

void ck_test_recurrent_split_qkv	(	const float *	packed_qkv,
		float *	q,
		float *	k,
		float *	v,
		int	rows,
		int	q_dim,
		int	k_dim,
		int	v_dim
	)

Split a packed recurrent QKV matrix into explicit Q, K, and V outputs.

Layout: packed_qkv : [rows, q_dim + k_dim + v_dim] q : [rows, q_dim] k : [rows, k_dim] v : [rows, v_dim]

Definition at line 789 of file ck_parity_api.c.

{
    recurrent_split_qkv_forward(packed_qkv, q, k, v, rows, q_dim, k_dim, v_dim);
}

References recurrent_split_qkv_forward().

◆ ck_test_rmsnorm()

void ck_test_rmsnorm	(	const float *	input,
		const float *	weight,
		float *	output,
		int	n_tokens,
		int	dim,
		float	eps
	)

RMSNorm.

Computes: output = (input / rms(input)) * weight where rms(x) = sqrt(mean(x^2) + eps)

Parameters

input	Input tensor [n_tokens, dim]
weight	Normalization weights [dim]
output	Output tensor [n_tokens, dim]
n_tokens	Number of tokens
dim	Hidden dimension
eps	Epsilon for numerical stability

Definition at line 565 of file ck_parity_api.c.

{
    /* CK rmsnorm_forward has aligned_embed_dim parameter
     * For testing, use dim as aligned_embed_dim (no padding) */
    rmsnorm_forward(input, weight, output, NULL, n_tokens, dim, dim, eps);
}

References rmsnorm_forward().

◆ ck_test_rope()

void ck_test_rope	(	float *	q,
		float *	k,
		int	n_tokens,
		int	n_heads,
		int	n_heads_kv,
		int	head_dim,
		int	pos_offset,
		float	theta
	)

RoPE (Rotary Position Embedding)

Applies rotary position embeddings to Q and K tensors.

NOTE: CK uses rotate-half format (split first/second halves) while some implementations use interleaved format. The test harness should account for this.

Parameters

q	Query tensor [n_tokens, n_heads * head_dim], modified in-place
k	Key tensor [n_tokens, n_heads_kv * head_dim], modified in-place
n_tokens	Number of tokens
n_heads	Number of query heads
n_heads_kv	Number of key/value heads
head_dim	Dimension per head
pos_offset	Starting position for RoPE
theta	RoPE base frequency (typically 10000.0)

Definition at line 575 of file ck_parity_api.c.

{
    /* Precompute cos/sin cache */
    int half_dim = head_dim / 2;
    int max_seq = pos_offset + n_tokens;
 
    float *cos_cache = (float *)malloc(max_seq * half_dim * sizeof(float));
    float *sin_cache = (float *)malloc(max_seq * half_dim * sizeof(float));
    if (!cos_cache || !sin_cache) {
        free(cos_cache);
        free(sin_cache);
        return;
    }
 
    rope_precompute_cache(cos_cache, sin_cache, max_seq, head_dim, theta,
                          head_dim, "none", 1.0f);
 
    /* CK RoPE expects layout [num_heads, num_tokens, head_dim]
     * Reshape from [n_tokens, n_heads * head_dim] to [n_heads, n_tokens, head_dim]
     */
    float *q_reorder = (float *)malloc(n_heads * n_tokens * head_dim * sizeof(float));
    float *k_reorder = (float *)malloc(n_heads_kv * n_tokens * head_dim * sizeof(float));
 
    if (q_reorder && k_reorder) {
        /* Reorder Q: [T, H*D] -> [H, T, D] */
        for (int t = 0; t < n_tokens; t++) {
            for (int h = 0; h < n_heads; h++) {
                for (int d = 0; d < head_dim; d++) {
                    q_reorder[h * n_tokens * head_dim + t * head_dim + d] =
                        q[t * n_heads * head_dim + h * head_dim + d];
                }
            }
        }
 
        /* Reorder K: [T, H_kv*D] -> [H_kv, T, D] */
        for (int t = 0; t < n_tokens; t++) {
            for (int h = 0; h < n_heads_kv; h++) {
                for (int d = 0; d < head_dim; d++) {
                    k_reorder[h * n_tokens * head_dim + t * head_dim + d] =
                        k[t * n_heads_kv * head_dim + h * head_dim + d];
                }
            }
        }
 
        /* Apply RoPE */
        rope_forward_qk(q_reorder, k_reorder,
                        cos_cache, sin_cache,
                        n_heads, n_heads_kv, n_tokens,
                        head_dim, head_dim, pos_offset);
 
        /* Reorder back: [H, T, D] -> [T, H*D] */
        for (int t = 0; t < n_tokens; t++) {
            for (int h = 0; h < n_heads; h++) {
                for (int d = 0; d < head_dim; d++) {
                    q[t * n_heads * head_dim + h * head_dim + d] =
                        q_reorder[h * n_tokens * head_dim + t * head_dim + d];
                }
            }
        }
 
        for (int t = 0; t < n_tokens; t++) {
            for (int h = 0; h < n_heads_kv; h++) {
                for (int d = 0; d < head_dim; d++) {
                    k[t * n_heads_kv * head_dim + h * head_dim + d] =
                        k_reorder[h * n_tokens * head_dim + t * head_dim + d];
                }
            }
        }
    }
 
    free(q_reorder);
    free(k_reorder);
    free(cos_cache);
    free(sin_cache);
}

References rope_forward_qk(), and rope_precompute_cache().

◆ ck_test_rope_interleaved()

void ck_test_rope_interleaved	(	float *	q,
		float *	k,
		int	n_tokens,
		int	n_heads,
		int	n_heads_kv,
		int	head_dim,
		int	pos_offset,
		float	theta
	)

RoPE with interleaved format (for llama.cpp compatibility)

Uses interleaved format: (x0, x1) -> (x0*cos - x1*sin, x0*sin + x1*cos)

Definition at line 653 of file ck_parity_api.c.

{
    /* Interleaved RoPE format (matches llama.cpp):
     * (x0, x1) -> (x0*cos - x1*sin, x0*sin + x1*cos)
     * Applied to consecutive pairs of elements
     */
 
    /* Precompute inverse frequencies */
    float *inv_freq = (float *)malloc((head_dim / 2) * sizeof(float));
    if (!inv_freq) return;
 
    for (int i = 0; i < head_dim / 2; i++) {
        inv_freq[i] = 1.0f / powf(theta, (float)(2 * i) / head_dim);
    }
 
    /* Apply RoPE to Q */
    for (int t = 0; t < n_tokens; t++) {
        int pos = pos_offset + t;
        for (int h = 0; h < n_heads; h++) {
            float *qh = q + t * n_heads * head_dim + h * head_dim;
 
            for (int i = 0; i < head_dim / 2; i++) {
                float freq = pos * inv_freq[i];
                float cos_val = cosf(freq);
                float sin_val = sinf(freq);
 
                /* Interleaved format */
                float x0 = qh[i * 2];
                float x1 = qh[i * 2 + 1];
                qh[i * 2]     = x0 * cos_val - x1 * sin_val;
                qh[i * 2 + 1] = x0 * sin_val + x1 * cos_val;
            }
        }
    }
 
    /* Apply RoPE to K */
    for (int t = 0; t < n_tokens; t++) {
        int pos = pos_offset + t;
        for (int h = 0; h < n_heads_kv; h++) {
            float *kh = k + t * n_heads_kv * head_dim + h * head_dim;
 
            for (int i = 0; i < head_dim / 2; i++) {
                float freq = pos * inv_freq[i];
                float cos_val = cosf(freq);
                float sin_val = sinf(freq);
 
                float x0 = kh[i * 2];
                float x1 = kh[i * 2 + 1];
                kh[i * 2]     = x0 * cos_val - x1 * sin_val;
                kh[i * 2 + 1] = x0 * sin_val + x1 * cos_val;
            }
        }
    }
 
    free(inv_freq);
}

◆ ck_test_softmax()

void ck_test_softmax	(	const float *	input,
		float *	output,
		int	n
	)

Softmax (simple, non-causal)

Computes: output[i] = exp(input[i]) / sum(exp(input))

Parameters

input	Input tensor [n]
output	Output tensor [n]
n	Number of elements

Definition at line 719 of file ck_parity_api.c.

{
    /* Find max for numerical stability */
    float max_val = input[0];
    for (int i = 1; i < n; i++) {
        if (input[i] > max_val) max_val = input[i];
    }
 
    /* Compute exp and sum */
    float sum = 0.0f;
    for (int i = 0; i < n; i++) {
        output[i] = expf(input[i] - max_val);
        sum += output[i];
    }
 
    /* Normalize */
    float inv_sum = 1.0f / sum;
    for (int i = 0; i < n; i++) {
        output[i] *= inv_sum;
    }
}

◆ ck_test_split_q_gate()

void ck_test_split_q_gate	(	const float *	packed_qg,
		float *	q,
		float *	gate,
		int	rows,
		int	q_dim,
		int	gate_dim,
		int	group_dim
	)

Split a packed full-attention Q+gate matrix into Q rows and gate rows.

Layout: packed_qg : [rows, q_dim + gate_dim] q : [rows, q_dim] gate : [rows, gate_dim]

Definition at line 778 of file ck_parity_api.c.

{
    split_q_gate_forward(packed_qg, q, gate, rows, q_dim, gate_dim, group_dim);
}

References split_q_gate_forward().

◆ ck_test_ssm_conv1d()

void ck_test_ssm_conv1d	(	const float *	conv_x,
		const float *	kernel,
		float *	out,
		int	kernel_size,
		int	num_channels,
		int	num_tokens,
		int	num_seqs
	)

qwen3next/Qwen3.5 SSM causal depthwise convolution.

Layout: conv_x [num_seqs, num_channels, kernel_size - 1 + num_tokens] kernel [num_channels, kernel_size] out [num_seqs, num_tokens, num_channels]

This mirrors ggml's GGML_OP_SSM_CONV used immediately before the DeltaNet recurrent update in qwen3next/Qwen3.5.

Definition at line 767 of file ck_parity_api.c.

{
    ssm_conv1d_forward(conv_x, kernel, out, kernel_size, num_channels, num_tokens, num_seqs);
}

References ssm_conv1d_forward().

◆ ck_test_swiglu()

void ck_test_swiglu	(	const float *	gate_up,
		float *	output,
		int	n_tokens,
		int	intermediate_dim
	)

SwiGLU activation.

Computes: output = SiLU(gate) * up where SiLU(x) = x * sigmoid(x)

Parameters

gate_up	Input tensor [n_tokens, 2 * intermediate_dim] Layout: [gate_0..gate_D-1, up_0..up_D-1] per token
output	Output tensor [n_tokens, intermediate_dim]
n_tokens	Number of tokens
intermediate_dim	Intermediate dimension

Definition at line 712 of file ck_parity_api.c.

{
    swiglu_forward(gate_up, output, n_tokens, intermediate_dim);
}

References swiglu_forward().

◆ ck_test_vec_dot_q4_k_q8_k()

void ck_test_vec_dot_q4_k_q8_k	(	const void *	weight_q4_k,
		const void *	input_q8_k,
		float *	output,
		int	cols
	)

Direct Q4_K x Q8_K dot product using identical pre-quantized bytes.

Definition at line 375 of file ck_parity_api.c.

{
    /* M=1 reaches the production Q4 dispatch without requantizing input. */
    gemv_q4_k_q8_k(output, weight_q4_k, input_q8_k, 1, cols);
}

References gemv_q4_k_q8_k().

◆ ck_test_vec_dot_q5_0_q8_0()

void ck_test_vec_dot_q5_0_q8_0	(	const void *	weight_q5_0,
		const void *	input_q8_0,
		float *	output,
		int	cols
	)

Direct Q5_0 x Q8_0 dot product test (takes pre-quantized Q8_0 input)

Direct Q5_0 x Q8_0 dot product (takes pre-quantized Q8_0 input)

This is a "direct" test that bypasses FP32-to-Q8_0 conversion. Useful for isolating kernel bugs from quantization bugs.

Parameters

weight_q5_0	Q5_0 quantized weights [cols]
input_q8_0	Q8_0 quantized input [cols] (pre-quantized!)
output	Output scalar [1]
cols	Number of elements (must be multiple of 32)

Definition at line 406 of file ck_parity_api.c.

{
    vec_dot_q5_0_q8_0(cols, output, weight_q5_0, input_q8_0);
}

References vec_dot_q5_0_q8_0().

◆ ck_test_vec_dot_q6_k_q8_k()

void ck_test_vec_dot_q6_k_q8_k	(	const void *	weight_q6_k,
		const void *	input_q8_k,
		float *	output,
		int	cols
	)

Direct Q6_K x Q8_K dot product using identical pre-quantized bytes.

Definition at line 384 of file ck_parity_api.c.

{
    /* Exercise the production M=1 provider. The scalar vec_dot helper is an
     * internal architecture-neutral oracle and does not preserve the x86
     * provider's declared lane reduction order. */
    gemv_q6_k_q8_k(output, weight_q6_k, input_q8_k, 1, cols);
}

References gemv_q6_k_q8_k().

◆ ck_test_vec_dot_q8_0_q8_0()

void ck_test_vec_dot_q8_0_q8_0	(	const void *	weight_q8_0,
		const void *	input_q8_0,
		float *	output,
		int	cols
	)

Direct Q8_0 x Q8_0 dot product test (takes pre-quantized Q8_0 input)

Direct Q8_0 x Q8_0 dot product (takes pre-quantized Q8_0 input)

Parameters

weight_q8_0	Q8_0 quantized weights [cols]
input_q8_0	Q8_0 quantized input [cols] (pre-quantized!)
output	Output scalar [1]
cols	Number of elements (must be multiple of 32)

Definition at line 422 of file ck_parity_api.c.

{
    vec_dot_q8_0_q8_0(cols, output, weight_q8_0, input_q8_0);
}

References vec_dot_q8_0_q8_0().

◆ dequant_q4_0_row()

void dequant_q4_0_row	(	const void *	src,
		float *	dst,
		size_t	n_elements
	)

extern

Dequantize Q4_0 row (multiple blocks)

Parameters

src	Q4_0 data
dst	FP32 output
n_elements	Number of elements to dequantize

Definition at line 63 of file dequant_kernels.c.

{
    const block_q4_0 *blocks = (const block_q4_0 *)src;
    const size_t n_blocks = n_elements / QK4_0;
 
    for (size_t b = 0; b < n_blocks; b++) {
        dequant_q4_0_block(&blocks[b], &dst[b * QK4_0]);
    }
}

Referenced by ck_test_dequant_q4_0(), and dequant_row().

◆ dequant_q4_k_row()

void dequant_q4_k_row	(	const void *	src,
		float *	dst,
		size_t	n_elements
	)

extern

Dequantize Q4_K row (multiple blocks)

Definition at line 372 of file dequant_kernels.c.

{
    const block_q4_K *blocks = (const block_q4_K *)src;
    const size_t n_blocks = n_elements / QK_K;
 
    for (size_t b = 0; b < n_blocks; b++) {
        dequant_q4_k_block(&blocks[b], &dst[b * QK_K]);
    }
}

Referenced by ck_test_dequant_q4_k(), and dequant_row().

◆ dequant_q6_k_row()

void dequant_q6_k_row	(	const void *	src,
		float *	dst,
		size_t	n_elements
	)

extern

Dequantize Q6_K row (multiple blocks)

Definition at line 422 of file dequant_kernels.c.

{
    const block_q6_K *blocks = (const block_q6_K *)src;
    const size_t n_blocks = n_elements / QK_K;
 
    for (size_t b = 0; b < n_blocks; b++) {
        dequant_q6_k_block(&blocks[b], &dst[b * QK_K]);
    }
}

Referenced by ck_test_dequant_q6_k(), and dequant_row().

◆ gated_deltanet_autoregressive_forward()

void gated_deltanet_autoregressive_forward	(	const float *	q,
		const float *	k,
		const float *	v,
		const float *	g,
		const float *	beta,
		const float *	state_in,
		float *	state_out,
		float *	out,
		int	num_heads,
		int	state_dim,
		float	norm_eps
	)

extern

Definition at line 1956 of file deltanet_kernels.c.

{
    if (!q || !k || !v || !g || !beta || !state_in || !state_out || !out) {
        return;
    }
    if (num_heads <= 0 || state_dim <= 0) {
        return;
    }
 
    /*
     * q and k arrive pre-normalized by recurrent_qk_l2_norm, so the
     * ISA-specialized kernels can follow the same contract as the scalar ref.
     */
    if (ck_strict_parity_enabled() || ck_deltanet_force_ref()) {
        gated_deltanet_autoregressive_forward_ref(
            q, k, v, g, beta, state_in, state_out, out, num_heads, state_dim, norm_eps);
        return;
    }
#if defined(__AVX512F__)
    gated_deltanet_autoregressive_forward_avx512(
        q, k, v, g, beta, state_in, state_out, out, num_heads, state_dim, norm_eps);
#elif defined(__AVX2__)
    gated_deltanet_autoregressive_forward_avx2(
        q, k, v, g, beta, state_in, state_out, out, num_heads, state_dim, norm_eps);
#elif defined(__AVX__)
    gated_deltanet_autoregressive_forward_avx(
        q, k, v, g, beta, state_in, state_out, out, num_heads, state_dim, norm_eps);
#else
    gated_deltanet_autoregressive_forward_ref(
        q, k, v, g, beta, state_in, state_out, out, num_heads, state_dim, norm_eps);
#endif
}

Referenced by ck_test_gated_deltanet_autoregressive(), and gated_deltanet_prefill_forward().

◆ gemm_nt_q4_k_q8_k()

void gemm_nt_q4_k_q8_k	(	const void *	A_q8,
		const void *	B,
		const float *	bias,
		float *	C,
		int	M,
		int	N,
		int	K
	)

extern

Definition at line 648 of file gemm_kernels_q4k_q8k.c.

{
    if (!A_q8 || !B || !C) {
        return;
    }
    if (M <= 0 || N <= 0 || K <= 0) {
        return;
    }
 
    gemm_q4_k_q8_k(C, B, A_q8, /*M_out=*/N, /*N_batch=*/M, K);
 
    if (!bias) {
        return;
    }
 
    for (int i = 0; i < M; ++i) {
        float *row = C + (size_t)i * (size_t)N;
        for (int j = 0; j < N; ++j) {
            row[j] += bias[j];
        }
    }
}

Referenced by ck_test_gemm_q4_k().

◆ gemm_nt_q5_0_q8_0()

void gemm_nt_q5_0_q8_0	(	const void *	A_q8,
		const void *	B_q5,
		const float *	bias,
		float *	C,
		int	M,
		int	N,
		int	K
	)

extern

Batch GEMM with Q5_0 weights and Q8_0 activations for prefill.

Computes C = A @ B^T + bias where: A: [M x K] Q8_0 quantized activations (M tokens, K features) B: [N x K] Q5_0 quantized weights (N outputs, K features) C: [M x N] FP32 output

This is the INT8 batch kernel for prefill, using pre-quantized activations to avoid FP32->Q8_0 conversion overhead per operation.

Parameters

A_q8	Input activations in Q8_0 format [M rows of K/32 blocks each]
B_q5	Weights in Q5_0 format [N rows of K/32 blocks each]
bias	Optional bias vector [N], NULL if not used
C	Output matrix [M x N], row-major FP32
M	Batch size (number of tokens)
N	Output dimension (number of output features)
K	Input dimension (must be multiple of 32)

Definition at line 1728 of file gemm_kernels_q5_0.c.

{
    const block_q5_0 *weights = (const block_q5_0 *)B_q5;
    const block_q8_0 *inputs = (const block_q8_0 *)A_q8;
    const int blocks_per_row = K / QK5_0;
 
    for (int m = 0; m < M; m++) {
        const block_q8_0 *input_row = &inputs[m * blocks_per_row];
 
        for (int n = 0; n < N; n++) {
            const block_q5_0 *weight_row = &weights[n * blocks_per_row];
            float *out = &C[m * N + n];
 
            /* Dispatches to vec_dot_q5_0_q8_0_avx (2x block unrolled) on AVX */
            vec_dot_q5_0_q8_0(K, out, weight_row, input_row);
 
            if (bias) {
                *out += bias[n];
            }
        }
    }
}

Referenced by ck_test_gemm_q5_0(), and gemm_nt_q5_0_q8_0_m2n4_tile().

◆ gemm_nt_q6_k_q8_k()

void gemm_nt_q6_k_q8_k	(	const void *	A_q8,
		const void *	B,
		const float *	bias,
		float *	C,
		int	M,
		int	N,
		int	K
	)

extern

Definition at line 1548 of file gemm_kernels_q6k_q8k.c.

{
    if (!A_q8 || !B || !C) {
        return;
    }
    if (M <= 0 || N <= 0 || K <= 0) {
        return;
    }
#if defined(__AVX2__) && defined(__FMA__)
    if (M >= 8 && ck_strict_parity_enabled()) {
        gemm_nt_q6_k_q8_k_tiled_parity(A_q8, B, bias, C, M, N, K);
        return;
    }
#endif
 
    /* Prefill GEMM uses the pinned x86 leaf reduction for short row groups
     * and the tiled reduction for strict batches of at least eight rows.
     * CK_DEBUG_Q6K_Q8K_REF=1 still forces the scalar reference leaf. */
    const block_q8_K *A = (const block_q8_K *)A_q8;
    const block_q6_K *W = (const block_q6_K *)B;
    const int blocks_per_vec = K / QK_K;
    const int blocks_per_row = K / QK_K;
 
    for (int m = 0; m < M; ++m) {
        const block_q8_K *a_row = A + (size_t)m * (size_t)blocks_per_vec;
        float *c_row = C + (size_t)m * (size_t)N;
        for (int n = 0; n < N; ++n) {
            const block_q6_K *w_row = W + (size_t)n * (size_t)blocks_per_row;
            const float b = bias ? bias[n] : 0.0f;
            c_row[n] = ck_dot_q6_k_q8_k_fast_or_ref(w_row, a_row, K) + b;
        }
    }
}

Referenced by ck_test_gemm_q6_k().

◆ gemm_nt_q8_0_q8_0()

void gemm_nt_q8_0_q8_0	(	const void *	A,
		const void *	B,
		const float *	bias,
		float *	C,
		int	M,
		int	N,
		int	K
	)

extern

gemm_nt_q8_0_q8_0 with optional bias (matches header signature)

C[m,n] = A[m,K] @ B[n,K]^T + bias[n]

Definition at line 552 of file gemm_batch_int8.c.

{
    /* First compute GEMM */
#if defined(__AVX2__)
    /*
     * The production contract is bit-exact with llama.cpp's eight-lane
     * FP32 accumulation tree.  AVX-512/VNNI changes the integer dot, but it
     * must not silently replace that FP32 reduction with the scalar-per-block
     * candidate above.  The AVX2-named entry point delegates each activation
     * row to the certified x4 provider, which also uses VNNI instructions when
     * they are available while preserving the declared reduction order.
     */
    gemm_nt_q8_0_q8_0_avx2(A, B, C, M, N, K);
#elif defined(__AVX__)
    gemm_nt_q8_0_q8_0_avx(A, B, C, M, N, K);
#else
    gemm_nt_q8_0_q8_0_ref(A, B, C, M, N, K);
#endif
 
    /* Add bias if provided */
    if (bias != NULL) {
        for (int m = 0; m < M; m++) {
            for (int n = 0; n < N; n++) {
                C[(size_t)m * N + n] += bias[n];
            }
        }
    }
}

Referenced by ck_test_gemm_q8_0().

◆ gemv_q4_k_q8_k()

void gemv_q4_k_q8_k	(	float *	y,
		const void *	W,
		const void *	x_q8,
		int	M,
		int	K
	)

extern

Definition at line 497 of file gemm_kernels_q4k_q8k.c.

{
    if (ck_q4k_q8k_force_ref()) {
        gemv_q4_k_q8_k_ref(y, W, x_q8, M, K);
        return;
    }
#if defined(__AVX512VNNI__) && defined(__AVX512VL__) && !defined(CK_NO_AVX512_VNNI)
    /* VNNI: Best for decode (single token) - INT8 dot product acceleration */
    gemv_q4_k_q8_k_vnni(y, W, x_q8, M, K);
#elif defined(__AVX2__)
    gemv_q4_k_q8_k_avx2(y, W, x_q8, M, K);
#elif defined(__AVX__)
    /* AVX version uses maddubs_epi16 (more efficient than SSE) */
    gemv_q4_k_q8_k_avx(y, W, x_q8, M, K);
#elif defined(__SSE4_1__)
    gemv_q4_k_q8_k_sse(y, W, x_q8, M, K);
#elif defined(__ARM_FEATURE_DOTPROD)
    /* SDOT NEON is integer-exact against the scalar reference (identical
     * block arithmetic and FP32 epilogue); CK_DEBUG_Q4K_Q8_REF=1 still
     * forces the reference above. Baseline armv8-a builds (no dot-product
     * extension) keep the reference path, which the compiler already
     * autovectorizes well. */
    gemv_q4_k_q8_k_neon(y, W, x_q8, M, K);
#else
    gemv_q4_k_q8_k_ref(y, W, x_q8, M, K);
#endif
}

Referenced by ck_test_gemv_q4_k(), and ck_test_vec_dot_q4_k_q8_k().

◆ gemv_q5_0()

void gemv_q5_0	(	float *	y,
		const void *	W,
		const float *	x,
		int	M,
		int	K
	)

extern

Auto-dispatch GEMV for Q5_0 weights based on CPU features.

Dispatch priority (best available):

AVX-512 (512-bit vectors) - Intel Skylake-X+
AVX2+FMA (256-bit vectors) - Intel Haswell+
AVX (256-bit vectors) - Intel Sandy Bridge+
SSE4.1 (128-bit vectors) - Intel Nehalem+
Reference (scalar) - Fallback

Uses ck_features.h for standardized feature detection.

Parameters

y	Output vector [M]
W	Weight matrix in Q5_0 format [M x K]
x	Input vector [K]
M	Number of output rows
K	Number of input columns (hidden dimension)

Definition at line 550 of file gemm_kernels_q5_0.c.

{
// Dispatch order: AVX512 > AVX2 > AVX > SSE > ref
#if defined(__AVX512F__)
    gemv_q5_0_avx512(y, W, x, M, K);
#elif defined(__AVX2__)
    gemv_q5_0_avx2(y, W, x, M, K);
#elif defined(__AVX__)
    gemv_q5_0_avx(y, W, x, M, K);
#elif defined(__SSE4_1__)
    gemv_q5_0_ref(y, W, x, M, K);
#else
    gemv_q5_0_ref(y, W, x, M, K);
#endif
}

Referenced by dot_q5_0(), gemm_nt_q5_0(), and gemm_q5_0().

◆ gemv_q5_0_q8_0()

void gemv_q5_0_q8_0	(	float *	y,
		const void *	W,
		const void *	x_q8,
		int	M,
		int	K
	)

extern

Matrix-vector multiply with Q5_0 weights and Q8_0 input.

Parameters

y	Output vector [M]
W	Weight matrix in Q5_0 format [M x K]
x_q8	Input vector in Q8_0 format [K]
M	Number of output rows
K	Number of columns (must be multiple of 32)

Definition at line 1640 of file gemm_kernels_q5_0.c.

{
    const block_q5_0 *w_blocks = (const block_q5_0 *)W;
    const block_q8_0 *x_blocks = (const block_q8_0 *)x_q8;
    const int blocks_per_row = K / QK5_0;
 
    for (int row = 0; row < M; row++) {
        vec_dot_q5_0_q8_0(K, &y[row],
                          &w_blocks[row * blocks_per_row],
                          x_blocks);
    }
}

Referenced by ck_test_gemv_q5_0(), and ck_test_gemv_q5_0_q8_0().

◆ gemv_q6_k_q8_k()

void gemv_q6_k_q8_k	(	float *	y,
		const void *	W,
		const void *	x_q8,
		int	M,
		int	K
	)

extern

GEMV: y = W @ x where W is Q6_K and x is Q8_K.

Definition at line 1356 of file gemm_kernels_q6k_q8k.c.

{
    if (ck_strict_parity_enabled() || ck_q6k_q8k_force_ref()) {
        gemv_q6_k_q8_k_ref(y, W, x_q8, M, K);
        return;
    }
 
#if defined(__AVX2__)
    /* llama.cpp's x86 Q6_K production graph keeps the AVX2 reduction order
     * even when AVX-512 is available. Wider ISA availability is not a license
     * to change this numerical contract. */
    gemv_q6_k_q8_k_avx2(y, W, x_q8, M, K);
    return;
#elif defined(__AVX__)
    gemv_q6_k_q8_k_avx(y, W, x_q8, M, K);
    return;
#elif defined(__SSE4_1__)
    gemv_q6_k_q8_k_sse(y, W, x_q8, M, K);
    return;
#elif defined(__ARM_NEON) || defined(__aarch64__)
    /* NEON keeps the same integer-exact per-block dot as the scalar
     * reference but accumulates the FP32 block contributions sequentially
     * instead of through eight partial lanes, so strict parity mode (handled
     * above) still uses the reference reduction order. */
    gemv_q6_k_q8_k_neon(y, W, x_q8, M, K);
    return;
#endif
    gemv_q6_k_q8_k_ref(y, W, x_q8, M, K);
}

Referenced by ck_test_gemv_q6_k(), ck_test_vec_dot_q6_k_q8_k(), and gemm_q6_k_q8_k().

◆ gemv_q8_0()

void gemv_q8_0	(	float *	y,
		const void *	W,
		const float *	x,
		int	M,
		int	K
	)

extern

Auto-dispatch GEMV for Q8_0 weights based on CPU features.

Dispatch priority (best available):

AVX-512 (512-bit vectors) - Intel Skylake-X+
AVX2+FMA (256-bit vectors) - Intel Haswell+
AVX (256-bit vectors) - Intel Sandy Bridge+
SSE4.1 (128-bit vectors) - Intel Nehalem+
Reference (scalar) - Fallback

Uses ck_features.h for standardized feature detection.

Parameters

y	Output vector [M]
W	Weight matrix in Q8_0 format [M x K]
x	Input vector [K]
M	Number of output rows
K	Number of input columns (hidden dimension)

Definition at line 694 of file gemm_kernels_q8_0.c.

{
    if (ck_q8_0_debug_ref()) {
        gemv_q8_0_ref(y, W, x, M, K);
        return;
    }
 
// Dispatch order: AVX512 > AVX2 > AVX > SSE > ref
#if defined(__AVX512F__)
    gemv_q8_0_avx512(y, W, x, M, K);
#elif defined(__AVX2__)
    gemv_q8_0_avx2(y, W, x, M, K);
#elif defined(__AVX__)
    gemv_q8_0_avx(y, W, x, M, K);
#elif defined(__SSE4_1__)
    gemv_q8_0_sse(y, W, x, M, K);
#else
    gemv_q8_0_ref(y, W, x, M, K);
#endif
}

Referenced by dot_q8_0(), gemm_nt_q8_0_rowloop(), and gemm_q8_0().

◆ gemv_q8_0_q8_0()

void gemv_q8_0_q8_0	(	float *	y,
		const void *	W,
		const void *	x_q8,
		int	M,
		int	K
	)

extern

Matrix-vector multiply with Q8_0 weights and Q8_0 input.

Parameters

y	Output vector [M]
W	Weight matrix in Q8_0 format [M x K]
x_q8	Input vector in Q8_0 format [K]
M	Number of output rows
K	Number of columns (must be multiple of 32)

Definition at line 1405 of file gemm_kernels_q8_0.c.

{
    const block_q8_0 *w_blocks = (const block_q8_0 *)W;
    const block_q8_0 *x_blocks = (const block_q8_0 *)x_q8;
    const int blocks_per_row = K / QK8_0;
 
    for (int row = 0; row < M; row++) {
        vec_dot_q8_0_q8_0(K, &y[row],
                          &w_blocks[row * blocks_per_row],
                          x_blocks);
    }
}

Referenced by ck_test_gemv_q8_0(), ck_test_gemv_q8_0_q8_0(), gemm_q8_0_q8_0_m2n4_strided(), and gemv_q8_0_q8_0_x4().

◆ mega_fused_outproj_mlp_prefill()

void mega_fused_outproj_mlp_prefill	(	float *	output,
		const float *	attn_out,
		const float *	residual,
		const float *	ln2_gamma,
		const void *	wo,
		const float *	bo,
		int	wo_dt,
		const void *	w1,
		const float *	b1,
		int	w1_dt,
		const void *	w2,
		const float *	b2,
		int	w2_dt,
		int	tokens,
		int	embed_dim,
		int	aligned_embed_dim,
		int	num_heads,
		int	aligned_head_dim,
		int	intermediate_dim,
		int	aligned_intermediate_dim,
		float	eps,
		void *	scratch
	)

extern

Referenced by ck_test_outproj_mlp_fused_q5_0().

◆ mega_fused_outproj_mlp_prefill_scratch_size()

size_t mega_fused_outproj_mlp_prefill_scratch_size	(	int	tokens,
		int	aligned_embed_dim,
		int	num_heads,
		int	aligned_head_dim,
		int	aligned_intermediate_dim
	)

extern

Get scratch buffer size for mega_fused_outproj_mlp_prefill.

Definition at line 159 of file mega_fused_outproj_mlp_prefill.c.

{
    if (tokens <= 0 || aligned_embed_dim <= 0 || num_heads <= 0 ||
        aligned_head_dim <= 0 || aligned_intermediate_dim <= 0) {
        return 0;
    }
 
    const size_t q8_row_bytes = ck_dtype_row_bytes(CK_DT_Q8_0,
                                                   (size_t)aligned_head_dim);
    const size_t attn_q8_bytes = (size_t)num_heads * (size_t)tokens * q8_row_bytes;
    const size_t h1_bytes = (size_t)tokens * (size_t)aligned_embed_dim * sizeof(float);
    const size_t ln2_bytes = h1_bytes;
    const size_t mlp_scratch = fused_mlp_swiglu_prefill_w1w2_quant_scratch_size(
        aligned_embed_dim, aligned_intermediate_dim);
 
    return align_up_size(attn_q8_bytes, 64) +
           align_up_size(h1_bytes, 64) +
           align_up_size(ln2_bytes, 64) +
           align_up_size(mlp_scratch, 64);
}

Referenced by ck_test_outproj_mlp_fused_q5_0().

◆ quantize_row_q8_0()

void quantize_row_q8_0	(	const float *	x,
		void *	vy,
		int	k
	)

extern

Quantize FP32 to Q8_0 format (scalar reference)

Parameters

x	Input FP32 values
vy	Output Q8_0 blocks
k	Number of elements (must be multiple of 32)

Definition at line 125 of file gemm_kernels_q8_0.c.

{
    block_q8_0 *y = (block_q8_0 *)vy;
    const int nb = k / QK8_0;  /* QK8_0 = 32 */
 
#if defined(__AVX__)
    const __m256 sign_bit = _mm256_set1_ps(-0.0f);
 
    for (int i = 0; i < nb; i++) {
        __m256 v0 = _mm256_loadu_ps(x + 0);
        __m256 v1 = _mm256_loadu_ps(x + 8);
        __m256 v2 = _mm256_loadu_ps(x + 16);
        __m256 v3 = _mm256_loadu_ps(x + 24);
        x += QK8_0;
 
        __m256 max_abs = _mm256_andnot_ps(sign_bit, v0);
        max_abs = _mm256_max_ps(max_abs, _mm256_andnot_ps(sign_bit, v1));
        max_abs = _mm256_max_ps(max_abs, _mm256_andnot_ps(sign_bit, v2));
        max_abs = _mm256_max_ps(max_abs, _mm256_andnot_ps(sign_bit, v3));
 
        __m128 max4 = _mm_max_ps(_mm256_extractf128_ps(max_abs, 1),
                                 _mm256_castps256_ps128(max_abs));
        max4 = _mm_max_ps(max4, _mm_movehl_ps(max4, max4));
        max4 = _mm_max_ss(max4, _mm_movehdup_ps(max4));
        const float max_scalar = _mm_cvtss_f32(max4);
 
#if defined(__INTEL_LLVM_COMPILER)
        const float d = ck_q8_0_div_rounded_f32(max_scalar, 127.0f);
        const float id = max_scalar != 0.0f
            ? ck_q8_0_div_rounded_f32(127.0f, max_scalar)
            : 0.0f;
#else
        const float d = max_scalar / 127.0f;
        const float id = max_scalar != 0.0f ? 127.0f / max_scalar : 0.0f;
#endif
        y[i].d = CK_FP32_TO_FP16(d);
 
        const __m256 mul = _mm256_set1_ps(id);
        v0 = _mm256_mul_ps(v0, mul);
        v1 = _mm256_mul_ps(v1, mul);
        v2 = _mm256_mul_ps(v2, mul);
        v3 = _mm256_mul_ps(v3, mul);
 
        /* Match llama.cpp x86 Q8 quantization: nearest-even rounding. */
        v0 = _mm256_round_ps(v0, _MM_ROUND_NEAREST);
        v1 = _mm256_round_ps(v1, _MM_ROUND_NEAREST);
        v2 = _mm256_round_ps(v2, _MM_ROUND_NEAREST);
        v3 = _mm256_round_ps(v3, _MM_ROUND_NEAREST);
 
        __m256i i0 = _mm256_cvtps_epi32(v0);
        __m256i i1 = _mm256_cvtps_epi32(v1);
        __m256i i2 = _mm256_cvtps_epi32(v2);
        __m256i i3 = _mm256_cvtps_epi32(v3);
 
#if defined(__AVX2__)
        i0 = _mm256_packs_epi32(i0, i1);
        i2 = _mm256_packs_epi32(i2, i3);
        i0 = _mm256_packs_epi16(i0, i2);
 
        const __m256i perm = _mm256_setr_epi32(0, 4, 1, 5, 2, 6, 3, 7);
        i0 = _mm256_permutevar8x32_epi32(i0, perm);
        _mm256_storeu_si256((__m256i *)y[i].qs, i0);
#else
        __m128i ni0 = _mm256_castsi256_si128(i0);
        __m128i ni1 = _mm256_extractf128_si256(i0, 1);
        __m128i ni2 = _mm256_castsi256_si128(i1);
        __m128i ni3 = _mm256_extractf128_si256(i1, 1);
        __m128i ni4 = _mm256_castsi256_si128(i2);
        __m128i ni5 = _mm256_extractf128_si256(i2, 1);
        __m128i ni6 = _mm256_castsi256_si128(i3);
        __m128i ni7 = _mm256_extractf128_si256(i3, 1);
 
        ni0 = _mm_packs_epi32(ni0, ni1);
        ni2 = _mm_packs_epi32(ni2, ni3);
        ni4 = _mm_packs_epi32(ni4, ni5);
        ni6 = _mm_packs_epi32(ni6, ni7);
 
        ni0 = _mm_packs_epi16(ni0, ni2);
        ni4 = _mm_packs_epi16(ni4, ni6);
 
        _mm_storeu_si128((__m128i *)(y[i].qs + 0), ni0);
        _mm_storeu_si128((__m128i *)(y[i].qs + 16), ni4);
#endif
    }
#else
    for (int i = 0; i < nb; i++) {
        const float *xb = x + i * QK8_0;
 
        /* Find max absolute value in block */
        float amax = 0.0f;
        for (int j = 0; j < QK8_0; j++) {
            float av = xb[j] >= 0 ? xb[j] : -xb[j];
            if (av > amax) amax = av;
        }
 
        /* Compute scale: d = max / 127 */
        float d = amax / 127.0f;
        float id = d != 0.0f ? 127.0f / amax : 0.0f;
 
        /* Store scale as FP16 */
        y[i].d = CK_FP32_TO_FP16(d);
 
        /* Quantize values */
        for (int j = 0; j < QK8_0; j++) {
            float v = xb[j] * id;
            int q = ck_nearest_int_q8_0(v);
            if (q > 127) q = 127;
            if (q < -127) q = -127;
            y[i].qs[j] = (int8_t)q;
        }
    }
#endif
}

Referenced by ck_test_gemm_q5_0(), ck_test_gemm_q8_0(), ck_test_gemv_q5_0(), ck_test_gemv_q5_0_q8_0(), ck_test_gemv_q8_0(), and ck_test_gemv_q8_0_q8_0().

◆ quantize_row_q8_k()

void quantize_row_q8_k	(	const float *	x,
		void *	vy,
		int	k
	)

extern

Definition at line 229 of file gemm_kernels_q4k_q8k.c.

                                                        {
    const char *ref_env = getenv("CK_DEBUG_Q8K_REF");
    if (ref_env && atoi(ref_env) != 0) {
        quantize_row_q8_k_ref(x, vy, k);
        return;
    }
#if defined(__AVX512F__) && defined(__AVX512BW__)
    quantize_row_q8_k_avx512(x, vy, k);
#elif defined(__AVX2__)
    quantize_row_q8_k_avx2(x, vy, k);
#elif defined(__AVX__)
    quantize_row_q8_k_avx(x, vy, k);
#elif defined(__SSE4_1__)
    quantize_row_q8_k_sse(x, vy, k);
#elif defined(__ARM_NEON) || defined(__aarch64__)
    quantize_row_q8_k_neon(x, vy, k);
#else
    quantize_row_q8_k_ref(x, vy, k);
#endif
}

Referenced by ck_test_gemm_q4_k(), ck_test_gemm_q6_k(), ck_test_gemv_q4_k(), ck_test_gemv_q6_k(), and ck_test_quantize_q8_k().

◆ recurrent_conv_state_update_forward()

void recurrent_conv_state_update_forward	(	const float *	state_in,
		const float *	q,
		const float *	k,
		const float *	v,
		float *	conv_x,
		float *	state_out,
		int	history_len,
		int	num_seqs,
		int	num_tokens,
		int	q_dim,
		int	k_dim,
		int	v_dim
	)

extern

Definition at line 100 of file recurrent_state_kernels.c.

                                                    {
    if (!state_in || !q || !k || !v || !conv_x || !state_out ||
        history_len < 0 || num_seqs <= 0 || num_tokens < 0 ||
        q_dim < 0 || k_dim < 0 || v_dim < 0 ||
        q_dim > INT_MAX - k_dim || q_dim + k_dim > INT_MAX - v_dim ||
        history_len > INT_MAX - num_tokens) {
        return;
    }
    const int channels = q_dim + k_dim + v_dim;
    const int total_len = history_len + num_tokens;
    if (channels <= 0 || total_len <= 0 || num_seqs > INT_MAX / channels ||
        (num_tokens > 0 && num_seqs > INT_MAX / num_tokens)) {
        return;
    }
 
    ck_recurrent_conv_state_args_t args = {
        state_in, q, k, v, conv_x, state_out,
        history_len, num_tokens, q_dim, k_dim, v_dim, channels, total_len,
    };
    const int jobs = num_seqs * channels;
    const size_t elements = (size_t)jobs * (size_t)num_tokens;
    const size_t minimum_elements_per_thread = 32768u;
    if (num_tokens <= 1 || elements < 2u * minimum_elements_per_thread) {
        ck_recurrent_conv_state_channel_range(0, jobs, &args);
        return;
    }
 
    ck_threadpool_t *pool = ck_threadpool_global();
    int active = pool ? ck_threadpool_n_threads(pool) : 1;
    size_t useful_threads =
        (elements + minimum_elements_per_thread - 1u) /
        minimum_elements_per_thread;
    if ((size_t)active > useful_threads) active = (int)useful_threads;
    if (active > jobs) active = jobs;
 
    if (active > 1 && ck_threadpool_thread_id(pool) <= 0) {
        ck_threadpool_parallel_for_n(
            pool, active, 0, jobs, 32,
            ck_recurrent_conv_state_channel_range, &args);
    } else {
        ck_recurrent_conv_state_channel_range(0, jobs, &args);
    }
}

Referenced by ck_test_recurrent_conv_state_update().

◆ recurrent_dt_gate_forward()

void recurrent_dt_gate_forward	(	const float *	alpha,
		const float *	dt_bias,
		const float *	a,
		float *	gate,
		int	rows,
		int	num_heads,
		int	state_dim
	)

extern

Definition at line 66 of file recurrent_gate_kernels.c.

                                              {
    const int dim = num_heads * state_dim;
    for (int row = 0; row < rows; ++row) {
        const float *alpha_row = alpha + (size_t) row * (size_t) dim;
        float *gate_row = gate + (size_t) row * (size_t) dim;
        for (int col = 0; col < dim; ++col) {
            const float x = alpha_row[col] + dt_bias[col];
            gate_row[col] = recurrent_softplus(x) * a[col];
        }
    }
}

Referenced by ck_test_recurrent_dt_gate().

◆ recurrent_norm_gate_forward()

void recurrent_norm_gate_forward	(	const float *	x,
		const float *	gate,
		const float *	weight,
		float *	out,
		int	rows,
		int	num_heads,
		int	head_dim,
		float	eps
	)

extern

Definition at line 18 of file recurrent_norm_kernels.c.

                                            {
    const int inner_dim = num_heads * head_dim;
    for (int row = 0; row < rows; ++row) {
        const float *x_row = x + (size_t) row * (size_t) inner_dim;
        const float *gate_row = gate + (size_t) row * (size_t) inner_dim;
        float *out_row = out + (size_t) row * (size_t) inner_dim;
 
        for (int head = 0; head < num_heads; ++head) {
            const float *x_head = x_row + (size_t) head * (size_t) head_dim;
            const float *gate_head = gate_row + (size_t) head * (size_t) head_dim;
            float *out_head = out_row + (size_t) head * (size_t) head_dim;
 
            float ms = 0.0f;
            for (int col = 0; col < head_dim; ++col) {
                ms += x_head[col] * x_head[col];
            }
            ms /= (float) head_dim;
            const float inv_rms = 1.0f / sqrtf(ms + eps);
 
            for (int col = 0; col < head_dim; ++col) {
                const float g = gate_head[col];
                const float silu = g * recurrent_sigmoid_local(g);
                out_head[col] = x_head[col] * inv_rms * weight[col] * silu;
            }
        }
    }
}

Referenced by ck_test_recurrent_norm_gate().

◆ recurrent_qk_l2_norm_forward()

void recurrent_qk_l2_norm_forward	(	float *	q,
		float *	k,
		int	rows,
		int	q_dim,
		int	k_dim,
		int	head_dim,
		float	eps
	)

extern

Definition at line 124 of file recurrent_qk_norm_kernels.c.

                                             {
    if (!q || !k || rows <= 0 || q_dim <= 0 || k_dim <= 0 ||
        head_dim <= 0 || q_dim % head_dim != 0 || k_dim % head_dim != 0) {
        return;
    }
    const int q_heads = q_dim / head_dim;
    const int k_heads = k_dim / head_dim;
    if (q_heads > INT_MAX - k_heads) {
        return;
    }
    const int heads_per_row = q_heads + k_heads;
    if (rows > INT_MAX / heads_per_row) {
        return;
    }
 
    ck_recurrent_qk_l2_args_t args = {
        q, k, q_heads, k_heads, head_dim, eps,
    };
    const int jobs = rows * heads_per_row;
    const size_t elements = (size_t)jobs * (size_t)head_dim;
    const size_t minimum_elements_per_thread = 32768u;
    const uintptr_t q_begin = (uintptr_t)q;
    const uintptr_t q_end = q_begin +
        (size_t)rows * (size_t)q_dim * sizeof(float);
    const uintptr_t k_begin = (uintptr_t)k;
    const uintptr_t k_end = k_begin +
        (size_t)rows * (size_t)k_dim * sizeof(float);
    const int ranges_overlap = q_begin < k_end && k_begin < q_end;
    if (ranges_overlap || elements < 2u * minimum_elements_per_thread) {
        recurrent_l2_norm_rows_forward_one(q, rows, q_dim, head_dim, eps);
        recurrent_l2_norm_rows_forward_one(k, rows, k_dim, head_dim, eps);
        return;
    }
 
    ck_threadpool_t *pool = ck_threadpool_global();
    int active = pool ? ck_threadpool_n_threads(pool) : 1;
    size_t useful_threads =
        (elements + minimum_elements_per_thread - 1u) /
        minimum_elements_per_thread;
    if ((size_t)active > useful_threads) {
        active = (int)useful_threads;
    }
    if (active > jobs) {
        active = jobs;
    }
 
    if (active > 1 && ck_threadpool_thread_id(pool) <= 0) {
        ck_threadpool_parallel_for_n(
            pool, active, 0, jobs, 8,
            recurrent_qk_l2_norm_head_range, &args);
    } else {
        recurrent_l2_norm_rows_forward_one(q, rows, q_dim, head_dim, eps);
        recurrent_l2_norm_rows_forward_one(k, rows, k_dim, head_dim, eps);
    }
}

Referenced by ck_test_recurrent_qk_l2_norm().

◆ recurrent_silu_forward()

void recurrent_silu_forward	(	const float *	x,
		float *	out,
		int	rows,
		int	dim
	)

extern

Definition at line 138 of file recurrent_gate_kernels.c.

                                     {
    for (int row = 0; row < rows; ++row) {
        const float *x_row = x + (size_t) row * (size_t) dim;
        float *out_row = out + (size_t) row * (size_t) dim;
        for (int col = 0; col < dim; ++col) {
            const float xv = x_row[col];
            out_row[col] = xv * recurrent_sigmoid(xv);
        }
    }
}

Referenced by ck_test_recurrent_silu().

◆ recurrent_split_conv_qkv_forward()

void recurrent_split_conv_qkv_forward	(	const float *	packed_qkv,
		float *	q,
		float *	k,
		float *	v,
		int	rows,
		int	q_dim,
		int	k_dim,
		int	v_dim
	)

extern

Definition at line 97 of file recurrent_split_kernels.c.

                                                 {
    recurrent_split_qkv_forward(packed_qkv, q, k, v, rows, q_dim, k_dim, v_dim);
}

Referenced by ck_test_recurrent_split_conv_qkv().

◆ recurrent_split_qkv_forward()

void recurrent_split_qkv_forward	(	const float *	packed_qkv,
		float *	q,
		float *	k,
		float *	v,
		int	rows,
		int	q_dim,
		int	k_dim,
		int	v_dim
	)

extern

Definition at line 5 of file recurrent_split_kernels.c.

                                            {
    const int packed_dim = q_dim + k_dim + v_dim;
    for (int row = 0; row < rows; ++row) {
        const float *src = packed_qkv + (size_t) row * (size_t) packed_dim;
        float *q_dst = q + (size_t) row * (size_t) q_dim;
        float *k_dst = k + (size_t) row * (size_t) k_dim;
        float *v_dst = v + (size_t) row * (size_t) v_dim;
        memcpy(q_dst, src, (size_t) q_dim * sizeof(float));
        memcpy(k_dst, src + q_dim, (size_t) k_dim * sizeof(float));
        memcpy(v_dst, src + q_dim + k_dim, (size_t) v_dim * sizeof(float));
    }
}

Referenced by ck_test_recurrent_split_qkv(), and recurrent_split_conv_qkv_forward().

◆ rmsnorm_forward()

void rmsnorm_forward	(	const float *	input,
		const float *	gamma,
		float *	output,
		float *	rstd_cache,
		int	tokens,
		int	d_model,
		int	aligned_embed_dim,
		float	eps
	)

extern

Definition at line 735 of file rmsnorm_kernels.c.

{
    rmsnorm_forward_strided_f32(
        input,
        gamma,
        output,
        rstd_cache,
        tokens,
        d_model,
        aligned_embed_dim,
        aligned_embed_dim,
        eps
    );
}

Referenced by ck_test_rmsnorm().

◆ rope_forward_qk()

void rope_forward_qk	(	float *	q,
		float *	k,
		const float *	cos_cache,
		const float *	sin_cache,
		int	num_heads,
		int	num_kv_heads,
		int	num_tokens,
		int	head_dim,
		int	aligned_head_dim,
		int	pos_offset
	)

extern

RoPE forward for both Q and K (common inference pattern)

Test:

test_rope.py::TestRoPEForward::test_rope_forward_qk

test_fused_attention_decode.py::TestFusedAttentionDecode::test_qk_rope

test_parity.py::test_rope_qk_parity

Combined RoPE forward for both Q and K in one call. q: [num_heads, num_tokens, head_dim] k: [num_kv_heads, num_tokens, head_dim]

After changes: make test && make llamacpp-parity-full

Definition at line 1094 of file rope_kernels.c.

{
    rope_forward_qk_with_rotary_dim(q, k, cos_cache, sin_cache, num_heads, num_kv_heads,
                                    num_tokens, head_dim, aligned_head_dim, pos_offset, head_dim);
}

Referenced by ck_test_rope().

◆ rope_precompute_cache()

void rope_precompute_cache	(	float *	cos_cache,
		float *	sin_cache,
		int	max_seq_len,
		int	head_dim,
		float	base,
		int	rotary_dim,
		const char *	scaling_type,
		float	scaling_factor
	)

extern

Precompute RoPE cos/sin cache with rotary_dim and scaling support

Test:

test_rope.py::TestRoPECache::test_cache_computation

test_rope.py::TestRoPECache::test_cache_values

Precomputes cos(m * theta_i) and sin(m * theta_i) for positions 0..max_seq_len-1. Only computes for first rotary_dim channels; remaining head_dim - rotary_dim channels are NOT rotated (pass through unchanged).

Scaling types:

"none": Standard RoPE
"linear": Scale positions by 1/scaling_factor
"dynamic": NTK-aware dynamic scaling
"yarn": YaRN scaling (beta-based)

Parameters

cos_cache	Output: [max_seq_len, rotary_dim/2] cos values
sin_cache	Output: [max_seq_len, rotary_dim/2] sin values
max_seq_len	Maximum sequence length for cache
head_dim	Full head dimension (for frequency computation)
base	RoPE base frequency (theta)
rotary_dim	Number of dimensions to rotate (0 = use head_dim)
scaling_type	Scaling type string: "none", "linear", "dynamic", "yarn"
scaling_factor	Scaling factor (1.0 = no scaling)

After changes: make test

Definition at line 384 of file rope_kernels.c.

{
    // Use rotary_dim = head_dim if not specified
    if (rotary_dim <= 0 || rotary_dim > head_dim) {
        rotary_dim = head_dim;
    }
 
    // Use no scaling if not specified
    int is_linear_scaling = 0;
    if (scaling_type != NULL && strcmp(scaling_type, "linear") == 0 && scaling_factor > 0.0f && scaling_factor != 1.0f) {
        is_linear_scaling = 1;
    }
 
    int rotary_half = rotary_dim / 2;
 
    for (int pos = 0; pos < max_seq_len; ++pos) {
        // Apply linear scaling to position if needed
        float effective_pos = (float)pos;
        if (is_linear_scaling) {
            effective_pos = (float)pos / scaling_factor;
        }
 
        for (int i = 0; i < rotary_half; ++i) {
            // Match the FP32 reference contract directly. Computing this via
            // long-double log/exp makes the final float depend on the host
            // libm and long-double ABI.
            const float exponent = ((float)(2 * i)) / (float)rotary_dim;
            const float freq_f = 1.0f / powf(base, exponent);
            float angle_f = effective_pos * freq_f;
            cos_cache[pos * rotary_half + i] = cosf(angle_f);
            sin_cache[pos * rotary_half + i] = sinf(angle_f);
        }
    }
}

Referenced by ck_test_rope().

◆ split_q_gate_forward()

void split_q_gate_forward	(	const float *	packed_qg,
		float *	q,
		float *	gate,
		int	rows,
		int	q_dim,
		int	gate_dim,
		int	group_dim
	)

extern

Definition at line 41 of file hybrid_attention_kernels.c.

                                         {
    const int packed_dim = q_dim + gate_dim;
    if (!packed_qg || !q || !gate || rows <= 0 || q_dim <= 0 || gate_dim <= 0) {
        return;
    }
    if (group_dim <= 0) {
        group_dim = q_dim;
    }
    const int q_groups = q_dim / group_dim;
    const int gate_group_dim = (q_groups > 0 && gate_dim % q_groups == 0) ? (gate_dim / q_groups) : gate_dim;
    for (int row = 0; row < rows; ++row) {
        const float *src = packed_qg + (size_t) row * (size_t) packed_dim;
        float *q_dst = q + (size_t) row * (size_t) q_dim;
        float *gate_dst = gate + (size_t) row * (size_t) gate_dim;
        if (q_groups > 0 && q_groups * group_dim == q_dim && q_groups * gate_group_dim == gate_dim) {
            for (int group = 0; group < q_groups; ++group) {
                const size_t src_group_off = (size_t) group * (size_t) (group_dim + gate_group_dim);
                memcpy(
                    q_dst + (size_t) group * (size_t) group_dim,
                    src + src_group_off,
                    (size_t) group_dim * sizeof(float));
                memcpy(
                    gate_dst + (size_t) group * (size_t) gate_group_dim,
                    src + src_group_off + (size_t) group_dim,
                    (size_t) gate_group_dim * sizeof(float));
            }
        } else {
            memcpy(q_dst, src, (size_t) q_dim * sizeof(float));
            memcpy(gate_dst, src + q_dim, (size_t) gate_dim * sizeof(float));
        }
    }
}

Referenced by ck_test_split_q_gate().

◆ ssm_conv1d_forward()

void ssm_conv1d_forward	(	const float *	conv_x,
		const float *	kernel,
		float *	out,
		int	kernel_size,
		int	num_channels,
		int	num_tokens,
		int	num_seqs
	)

extern

Definition at line 128 of file ssm_kernels.c.

{
    ssm_conv1d_forward_ref(conv_x, kernel, out, kernel_size, num_channels, num_tokens, num_seqs);
}

Referenced by ck_test_ssm_conv1d().

◆ swiglu_forward()

void swiglu_forward	(	const float *	input,
		float *	output,
		int	tokens,
		int	dim
	)

extern

SwiGLU forward pass

Test:

test_swiglu.py::TestSwiGLUForward::test_forward_tokens

test_swiglu.py::TestSwiGLUForward::test_forward_single

test_mlp.py::TestMLPForward::test_swiglu_mlp

test_fused_swiglu_decode.py::TestFusedSwiGLUDecode::test_fused_swiglu_decode

test_parity.py::test_swiglu_parity

SwiGLU: y = silu(gate) * up where silu(x) = x * sigmoid(x)

After changes: make test && make llamacpp-parity-full

Definition at line 283 of file swiglu_kernels.c.

{
    const char *fast_env = getenv("CK_SWIGLU_FAST");
    const char *exact_env = getenv("CK_SWIGLU_EXACT");
    if (ck_strict_parity_enabled() ||
        !(fast_env && atoi(fast_env) != 0) ||
        (exact_env && atoi(exact_env) != 0)) {
        swiglu_forward_exact(input, output, tokens, dim);
        return;
    }
 
    int T = tokens;
    int D = dim;
 
    for (int t = 0; t < T; ++t) {
        const float *row = input + (size_t)t * (2 * D);
        float *out_row = output + (size_t)t * D;
        int d = 0;
 
#if defined(__AVX512F__)
        // AVX-512: Process 16 floats at a time
        for (; d + 16 <= D; d += 16) {
            __m512 a = _mm512_loadu_ps(&row[d]);         // gate
            __m512 b = _mm512_loadu_ps(&row[D + d]);     // value
 
            __m512 s = sigmoid512_fast(a);              // sigmoid(a)
            __m512 silu = _mm512_mul_ps(a, s);          // silu(a) = a * sigmoid(a)
            __m512 y = _mm512_mul_ps(silu, b);          // y = silu(a) * b
 
            _mm512_storeu_ps(&out_row[d], y);
        }
#elif defined(__AVX2__)
        // AVX2: Process 8 floats at a time
        for (; d + 8 <= D; d += 8) {
            __m256 a = _mm256_loadu_ps(&row[d]);         // gate
            __m256 b = _mm256_loadu_ps(&row[D + d]);     // value
 
            __m256 s = sigmoid256_fast(a);              // sigmoid(a)
            __m256 silu = _mm256_mul_ps(a, s);          // silu(a) = a * sigmoid(a)
            __m256 y = _mm256_mul_ps(silu, b);          // y = silu(a) * b
 
            _mm256_storeu_ps(&out_row[d], y);
        }
#elif defined(__AVX__)
        // AVX1: Vectorize arithmetic, use scalar sigmoid
        float a_arr[8] __attribute__((aligned(32)));
        float s_arr[8] __attribute__((aligned(32)));
 
        for (; d + 8 <= D; d += 8) {
            __m256 a = _mm256_loadu_ps(&row[d]);         // gate
            __m256 b = _mm256_loadu_ps(&row[D + d]);     // value
 
            // Compute sigmoid scalarly
            _mm256_store_ps(a_arr, a);
            for (int j = 0; j < 8; ++j) {
                s_arr[j] = sigmoid_scalar(a_arr[j]);
            }
            __m256 s = _mm256_load_ps(s_arr);
 
            __m256 silu = _mm256_mul_ps(a, s);          // silu(a) = a * sigmoid(a)
            __m256 y = _mm256_mul_ps(silu, b);          // y = silu(a) * b
 
            _mm256_storeu_ps(&out_row[d], y);
        }
#endif
 
        // Scalar fallback for remaining elements
        for (; d < D; ++d) {
            float a = row[d];       // gate
            float b = row[D + d];   // value
 
            float s = sigmoid_scalar(a);         // sigmoid(a)
            float silu = a * s;                  // silu(a) = a * sigmoid(a)
 
            out_row[d] = silu * b;
        }
    }
}

Referenced by ck_test_swiglu().

◆ vec_dot_q5_0_q8_0()

void vec_dot_q5_0_q8_0	(	int	n,
		float *	s,
		const void *	vx,
		const void *	vy
	)

extern

Auto-dispatch quantized dot product Q5_0 x Q8_0.

Dispatch priority:

AVX512 (best performance on modern Intel/AMD)
AVX (256-bit float ops, works on Sandy/Ivy Bridge and newer)
SSSE3 (128-bit fallback)
Reference scalar (last resort)

Definition at line 1602 of file gemm_kernels_q5_0.c.

{
#if defined(__AVX2__)
    /* llama.cpp uses the packed AVX2 dot on AVX-512 hosts as well. It keeps
     * Q5/Q8 data in byte lanes and avoids the per-block 32-bit lane expansion
     * overhead of the baseline AVX-512 path. */
    vec_dot_q5_0_q8_0_avx2(n, s, vx, vy);
#elif defined(__AVX512F__)
    vec_dot_q5_0_q8_0_avx512(n, s, vx, vy);
#elif defined(__ARM_NEON) || defined(__aarch64__)
    vec_dot_q5_0_q8_0_neon(n, s, vx, vy);
#elif defined(__AVX__)
    /* AVX for 256-bit float ops (works on Ivy Bridge and newer) */
    vec_dot_q5_0_q8_0_avx(n, s, vx, vy);
#elif defined(__SSSE3__)
    /* SSSE3 - most efficient on older CPUs */
    vec_dot_q5_0_q8_0_sse(n, s, vx, vy);
#else
    vec_dot_q5_0_q8_0_ref(n, s, vx, vy);
#endif
}

Referenced by ck_test_vec_dot_q5_0_q8_0().

◆ vec_dot_q6_k_q8_k()

void vec_dot_q6_k_q8_k	(	int	n,
		float *	s,
		const void *	vx,
		const void *	vy
	)

extern

Q6_K x Q8_K dot product (single row)

Definition at line 1339 of file gemm_kernels_q6k_q8k.c.

{
    if (!s || !vx || !vy || n <= 0) {
        return;
    }
 
    const block_q6_K *x = (const block_q6_K *)vx;
    const block_q8_K *y = (const block_q8_K *)vy;
 
    /* This is the architecture-neutral scalar oracle. Production x86 dispatch
     * uses the separately parity-tested SIMD reduction tree. */
    *s = dot_q6_k_q8_k_ref(x, y, n);
}

◆ vec_dot_q8_0_q8_0()

void vec_dot_q8_0_q8_0	(	int	n,
		float *	s,
		const void *	vx,
		const void *	vy
	)

extern

Auto-dispatch quantized dot product Q8_0 x Q8_0.

Definition at line 1368 of file gemm_kernels_q8_0.c.

{
    if (ck_q8_0_q8_0_debug_ref()) {
        vec_dot_q8_0_q8_0_ref(n, s, vx, vy);
        return;
    }
#ifdef __AVX512F__
    vec_dot_q8_0_q8_0_avx512(n, s, vx, vy);
#elif defined(__AVX2__)
    vec_dot_q8_0_q8_0_avx2(n, s, vx, vy);
#elif defined(__ARM_NEON) || defined(__aarch64__)
    vec_dot_q8_0_q8_0_neon(n, s, vx, vy);
#elif defined(__AVX__)
    vec_dot_q8_0_q8_0_avx(n, s, vx, vy);
#elif defined(__SSE4_1__)
    vec_dot_q8_0_q8_0_sse(n, s, vx, vy);
#else
    vec_dot_q8_0_q8_0_ref(n, s, vx, vy);
#endif
}

Referenced by ck_test_vec_dot_q8_0_q8_0().

Functions

Detailed Description

Function Documentation

◆ attention_forward_causal_head_major_gqa_flash_strided()

◆ attn_gate_sigmoid_mul_forward()

◆ ck_get_block_q4_k_size()

◆ ck_get_block_q6_k_size()

◆ ck_get_block_q8_k_size()

◆ ck_get_qk_k()

◆ ck_test_attention_causal()

◆ ck_test_attn_gate_sigmoid_mul()

◆ ck_test_dequant_q4_0()

◆ ck_test_dequant_q4_k()

◆ ck_test_dequant_q6_k()

◆ ck_test_gated_deltanet_autoregressive()

◆ ck_test_gemm_q4_k()

◆ ck_test_gemm_q5_0()

◆ ck_test_gemm_q6_k()

◆ ck_test_gemm_q8_0()

◆ ck_test_gemv_q4_k()

◆ ck_test_gemv_q5_0()

◆ ck_test_gemv_q5_0_q8_0()

◆ ck_test_gemv_q6_k()

◆ ck_test_gemv_q8_0()

◆ ck_test_gemv_q8_0_q8_0()

◆ ck_test_outproj_mlp_fused_q5_0()

◆ ck_test_quantize_q8_k()

◆ ck_test_recurrent_conv_state_update()

◆ ck_test_recurrent_dt_gate()

◆ ck_test_recurrent_norm_gate()

◆ ck_test_recurrent_qk_l2_norm()

◆ ck_test_recurrent_silu()

◆ ck_test_recurrent_split_conv_qkv()

◆ ck_test_recurrent_split_qkv()

◆ ck_test_rmsnorm()

◆ ck_test_rope()

◆ ck_test_rope_interleaved()

◆ ck_test_softmax()

◆ ck_test_split_q_gate()

◆ ck_test_ssm_conv1d()

◆ ck_test_swiglu()

◆ ck_test_vec_dot_q4_k_q8_k()

◆ ck_test_vec_dot_q5_0_q8_0()

◆ ck_test_vec_dot_q6_k_q8_k()

◆ ck_test_vec_dot_q8_0_q8_0()

◆ dequant_q4_0_row()

◆ dequant_q4_k_row()

◆ dequant_q6_k_row()

◆ gated_deltanet_autoregressive_forward()

◆ gemm_nt_q4_k_q8_k()

◆ gemm_nt_q5_0_q8_0()

◆ gemm_nt_q6_k_q8_k()

◆ gemm_nt_q8_0_q8_0()

◆ gemv_q4_k_q8_k()

◆ gemv_q5_0()

◆ gemv_q5_0_q8_0()

◆ gemv_q6_k_q8_k()

◆ gemv_q8_0()

◆ gemv_q8_0_q8_0()

◆ mega_fused_outproj_mlp_prefill()

◆ mega_fused_outproj_mlp_prefill_scratch_size()

◆ quantize_row_q8_0()

◆ quantize_row_q8_k()

◆ recurrent_conv_state_update_forward()

◆ recurrent_dt_gate_forward()

◆ recurrent_norm_gate_forward()

◆ recurrent_qk_l2_norm_forward()

◆ recurrent_silu_forward()

◆ recurrent_split_conv_qkv_forward()

◆ recurrent_split_qkv_forward()

◆ rmsnorm_forward()

◆ rope_forward_qk()

◆ rope_precompute_cache()

◆ split_q_gate_forward()

◆ ssm_conv1d_forward()

◆ swiglu_forward()

◆ vec_dot_q5_0_q8_0()

◆ vec_dot_q6_k_q8_k()

◆ vec_dot_q8_0_q8_0()