69 lines
2.0 KiB
Plaintext
69 lines
2.0 KiB
Plaintext
/* SPDX-License-Identifier: AGPL-3.0-or-later */
|
|
/* Copyright (C) 2026 David (Corey) Herrera */
|
|
/*
|
|
* matvec_q8.cu — INT8 quantized matrix-vector product for Foundry
|
|
*
|
|
* Target: sm_37 (K80, Kepler)
|
|
* Compile: nvcc -ptx -arch=sm_37 matvec_q8.cu -o matvec_q8.ptx
|
|
*
|
|
* Computes: out[row] = x_scale * w_scales[row] * dot(x_q[cols], w[row, cols])
|
|
* where x_q and w are int8, out is float32.
|
|
*
|
|
* One block per output row. Threads cooperatively compute the dot product
|
|
* then reduce via shared memory.
|
|
*/
|
|
|
|
extern "C" {
|
|
|
|
/*
|
|
* INT8 quantized matvec.
|
|
*
|
|
* x_q: [cols] int8 — quantized input vector
|
|
* w: [rows, cols] int8 — quantized weight matrix (row-major)
|
|
* w_scales: [rows] float32 — per-row weight scales
|
|
* out: [rows] float32 — output vector
|
|
* x_scale: float32 — input vector scale
|
|
* rows: number of output elements
|
|
* cols: inner dimension (input size)
|
|
*
|
|
* Launched with rows blocks, blockDim.x threads.
|
|
* Shared memory: blockDim.x * sizeof(float)
|
|
*/
|
|
__global__ void fc_matvec_q8(const signed char *x_q, const signed char *w,
|
|
const float *w_scales, float *out,
|
|
float x_scale, int rows, int cols) {
|
|
extern __shared__ float sdata[];
|
|
|
|
int row = blockIdx.x;
|
|
if (row >= rows) return;
|
|
|
|
int tid = threadIdx.x;
|
|
int threads = blockDim.x;
|
|
|
|
const signed char *w_row = w + row * cols;
|
|
|
|
/* Each thread computes a partial dot product over its stride */
|
|
int partial_sum = 0;
|
|
for (int i = tid; i < cols; i += threads) {
|
|
partial_sum += (int)x_q[i] * (int)w_row[i];
|
|
}
|
|
|
|
sdata[tid] = (float)partial_sum;
|
|
__syncthreads();
|
|
|
|
/* Shared memory reduction */
|
|
for (int s = threads / 2; s > 0; s >>= 1) {
|
|
if (tid < s) {
|
|
sdata[tid] += sdata[tid + s];
|
|
}
|
|
__syncthreads();
|
|
}
|
|
|
|
/* Thread 0 writes the final result */
|
|
if (tid == 0) {
|
|
out[row] = x_scale * w_scales[row] * sdata[0];
|
|
}
|
|
}
|
|
|
|
} /* extern "C" */
|