diff --git a/.vscode/settings.json b/.vscode/settings.json new file mode 100644 index 0000000..d5ee39f --- /dev/null +++ b/.vscode/settings.json @@ -0,0 +1,56 @@ +{ + "files.associations": { + "*.json": "jsonc", + "array": "cpp", + "atomic": "cpp", + "bit": "cpp", + "cctype": "cpp", + "cfenv": "cpp", + "clocale": "cpp", + "cmath": "cpp", + "compare": "cpp", + "concepts": "cpp", + "cstdarg": "cpp", + "cstddef": "cpp", + "cstdint": "cpp", + "cstdio": "cpp", + "cstdlib": "cpp", + "cstring": "cpp", + "ctime": "cpp", + "cwchar": "cpp", + "cwctype": "cpp", + "deque": "cpp", + "string": "cpp", + "unordered_map": "cpp", + "vector": "cpp", + "exception": "cpp", + "algorithm": "cpp", + "functional": "cpp", + "iterator": "cpp", + "memory": "cpp", + "memory_resource": "cpp", + "numeric": "cpp", + "optional": "cpp", + "random": "cpp", + "string_view": "cpp", + "system_error": "cpp", + "tuple": "cpp", + "type_traits": "cpp", + "utility": "cpp", + "fstream": "cpp", + "initializer_list": "cpp", + "iosfwd": "cpp", + "iostream": "cpp", + "istream": "cpp", + "limits": "cpp", + "new": "cpp", + "numbers": "cpp", + "ostream": "cpp", + "sstream": "cpp", + "stdexcept": "cpp", + "streambuf": "cpp", + "typeinfo": "cpp", + + }, + +} \ No newline at end of file diff --git a/build.sh b/build.sh new file mode 100755 index 0000000..4bdeef0 --- /dev/null +++ b/build.sh @@ -0,0 +1,36 @@ +#!/usr/bin/env bash +# Simple Linux build script for gemv_quantv1.cpp using OpenCL +# Usage: ./build.sh [output_name] +set -euo pipefail + +CXX=${CXX:-g++} +SRC=gemv_quantv1.cpp +OUT=${1:-gemv_quantv1} +CXXFLAGS_DEFAULT="-O3 -std=c++17 -DNDEBUG" +CXXFLAGS=${CXXFLAGS:-$CXXFLAGS_DEFAULT} + +# Gather OpenCL flags via pkg-config when available; otherwise fall back to -lOpenCL +CL_CFLAGS="" +CL_LDFLAGS="" +if command -v pkg-config >/dev/null 2>&1 && pkg-config --exists OpenCL; then + CL_CFLAGS="$(pkg-config --cflags OpenCL)" + CL_LDFLAGS="$(pkg-config --libs OpenCL)" +else + # Fallback: common system defaults + CL_LDFLAGS="-lOpenCL" +fi + +# Optional manual overrides similar to the Windows .bat +# e.g. OPENCL_HEADERS=/usr/include OPENCL_LIB=/usr/lib/x86_64-linux-gnu ./build.sh +if [[ -n "${OPENCL_HEADERS:-}" ]]; then + CL_CFLAGS+=" -I${OPENCL_HEADERS}" +fi +if [[ -n "${OPENCL_LIB:-}" ]]; then + CL_LDFLAGS+=" -L${OPENCL_LIB}" +fi + +set -x +"${CXX}" ${CXXFLAGS} ${CL_CFLAGS} -o "${OUT}" "${SRC}" ${CL_LDFLAGS} +set +x + +echo "Build finished: ${OUT}" diff --git a/gemv_quantv1 b/gemv_quantv1 new file mode 100755 index 0000000..856e43a Binary files /dev/null and b/gemv_quantv1 differ diff --git a/gemv_quantv1.cpp b/gemv_quantv1.cpp index bf5ff41..01bda7e 100644 --- a/gemv_quantv1.cpp +++ b/gemv_quantv1.cpp @@ -7,7 +7,7 @@ 4. 优化内存访问，权重拆分 5. 使用image图像内存存储权重； 6. ... -*/ +*/ /* 量化作业示例-权重A：BlockQ8_0，激活值B：half 内核测试： @@ -26,6 +26,8 @@ #include #include #include "half.hpp" +#include +#include using namespace std; using half_float::half; @@ -135,6 +137,28 @@ inline int8_t clamp_int8(int v) { return static_cast(std::max(-128, std::min(127, v))); } + +// 激活 B 的对称 per-32 量化：输出 qB(int8) 与 sB(half per-block) +void quantB_q8_per32(const vector &B, vector &qB, vector &sB, int k) +{ + int blocks = k / Block_size; + qB.resize(k); + sB.resize(blocks); + for (int j = 0, jj = 0; j < k; j += Block_size, ++jj) + { + float max_abs = 0.f; + for (int l = j; l < j + Block_size; ++l) + max_abs = max(max_abs, fabsf((float)B[l])); + half sd = (half)((max_abs > 0.f) ? (max_abs / QMAX) : EPS); + sB[jj] = sd; + for (int l = j; l < j + Block_size; ++l) + { + float v = (float)B[l] / (float)sd; + int iv = (int)roundf(v); + qB[l] = clamp_int8(iv); + } + } +} // 量化最后一维 量化块存储 // M: m x k ; produce BlockQ8_0: m x (k / 32) void quantv1(const vector &M, vector &q_M, int m, int k, int blocks) @@ -192,42 +216,20 @@ void quantv2(const vector &M, vector &q_M, vector &q_d_M, int } } -// kernel 示例 -const char *kernelSource = R"CLC( - #define CL_TARGET_OPENCL_VERSION 200 - #pragma OPENCL EXTENSION cl_khr_fp16 : enable - typedef struct { - half d; - char qs[32]; - } BlockQ8_0; - __kernel void gemv_q8_base(__global const BlockQ8_0 *A, __global const half *B, __global half *C, - int as, int ars, int acs, int bs, int brs, int bcs, - int cs, int crs, int ccs, - int M, int N, int K, float alpha, float beta) { - - int row_id = get_global_id(0); - BlockQ8_0 valueA; - half valueB = 0.0h; - half sum = 0.0h; - - for (int i = 0; i < K / 32; i++) { - valueA = *(A + row_id * ars + i * acs); - half value = 0.0h; - for (int j = 0; j < 32; j++) { - valueB = *(B + (i * 32 + j) * brs); - value += (half)(int)valueA.qs[j] * valueB; - } - sum += value * valueA.d; - } - - __global half *p = C + row_id * crs; - if (beta != 0) - *p = (half)(beta * (*p) + alpha * (float)sum); - else - *p = (half)(alpha * (float)sum); +static std::string loadTextFile(const std::string &path) +{ + std::ifstream ifs(path, std::ios::in | std::ios::binary); + if (!ifs) + { + throw std::runtime_error("Failed to open file: " + path); } + std::ostringstream oss; + oss << ifs.rdbuf(); + return oss.str(); +} -)CLC"; +const char *kernelPath = "./kernel.cl"; +// kernel 示例 // ---------- Host 辅助：编译、运行、衡量函数 ---------- void printDeviceInfo(cl_device_id device) @@ -251,19 +253,23 @@ void printDeviceInfo(cl_device_id device) } double KernelTest(const string &kernelName, - const char *kernelSrc, + // const char *kernelSrc, cl_context context, cl_device_id device, cl_command_queue queue, const vector &BlockA, // m x (k / 32) - const vector &B, // n * k + const vector &qB, // n * k (int8) + const vector &sB, // k/32 per-block scale for B (n==1) vector &C_gpu, const vector &C_ref, int m, int n, int k, float alpha, float beta) { + auto src = loadTextFile(kernelPath); + const char *kernelSource = src.c_str(); + const size_t src_len = src.size(); cl_int err; - cl_program program = clCreateProgramWithSource(context, 1, &kernelSrc, NULL, &err); + cl_program program = clCreateProgramWithSource(context, 1, &kernelSource, &src_len, &err); checkErr(err, "clCreateProgramWithSource"); const char *buildOptions = "-cl-std=CL2.0"; err = clBuildProgram(program, 1, &device, buildOptions, NULL, NULL); @@ -283,25 +289,30 @@ double KernelTest(const string &kernelName, int blocks = k / Block_size; size_t sizeBlockA = (size_t)(m * blocks) * sizeof(BlockQ8_0); - size_t sizeB = (size_t)n * (size_t)k * sizeof(half); + size_t sizeQB = (size_t)n * (size_t)k * sizeof(int8_t); + size_t sizeSB = (size_t)(k / Block_size) * sizeof(half); size_t sizeC = (size_t)n * (size_t)m * sizeof(half); cl_mem bufA = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeBlockA, (void *)BlockA.data(), &err); checkErr(err, "clCreateBuffer A"); - cl_mem bufB = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeB, (void *)B.data(), &err); - checkErr(err, "clCreateBuffer B"); + cl_mem bufQB = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeQB, (void *)qB.data(), &err); + checkErr(err, "clCreateBuffer qB"); + cl_mem bufSB = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeSB, (void *)sB.data(), &err); + checkErr(err, "clCreateBuffer sB"); cl_mem bufC = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeC, (void *)C_gpu.data(), &err); checkErr(err, "clCreateBuffer C"); // set kernel args int argIdx = 0; clSetKernelArg(kernel, argIdx++, sizeof(cl_mem), &bufA); - clSetKernelArg(kernel, argIdx++, sizeof(cl_mem), &bufB); + clSetKernelArg(kernel, argIdx++, sizeof(cl_mem), &bufQB); + clSetKernelArg(kernel, argIdx++, sizeof(cl_mem), &bufSB); clSetKernelArg(kernel, argIdx++, sizeof(cl_mem), &bufC); // strides // A:m*k B:n*k C:n*m // A * B^T = C^T int as = m * k / Block_size, ars = k / Block_size, acs = 1; + // qB layout matches B: n*k contiguous; for n==1, row-stride 1, col-stride k int bs = n * k, brs = 1, bcs = k; int cs = n * m, crs = 1, ccs = n; clSetKernelArg(kernel, argIdx++, sizeof(int), &as); @@ -319,9 +330,24 @@ double KernelTest(const string &kernelName, clSetKernelArg(kernel, argIdx++, sizeof(float), &alpha); clSetKernelArg(kernel, argIdx++, sizeof(float), &beta); - // NDRange: 1D: (m) - size_t gws[1] = {(size_t)m}; - size_t lws[1] = {256}; // 可以动态调整 + // NDRange: 1D: (m) — 选择安全的 LWS，并将 GWS 向上取整到 LWS 的倍数 + size_t devMaxWG = 0; + clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(devMaxWG), &devMaxWG, NULL); + size_t devMaxItems[3] = {0, 0, 0}; + clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(devMaxItems), &devMaxItems, NULL); + size_t kernelMaxWG = 0; + clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(kernelMaxWG), &kernelMaxWG, NULL); + + size_t lwsCandidate = 256; // 目标 LWS + size_t lws0 = lwsCandidate; + lws0 = std::min(lws0, devMaxWG); + lws0 = std::min(lws0, devMaxItems[0]); + lws0 = std::min(lws0, kernelMaxWG); + if (lws0 == 0) + lws0 = 1; // 兜底 + + size_t lws[1] = {lws0}; + size_t gws[1] = {((size_t)m + lws[0] - 1) / lws[0] * lws[0]}; // warmup for (int i = 0; i < NUM_WARMUP; ++i) @@ -357,12 +383,13 @@ double KernelTest(const string &kernelName, half max_abs = computeAbsoluteError(C_ref, C_gpu); - printf("Kernel: %s | avg_time = %.3f ms | max_abs_err = %f", + printf("Kernel: %s | avg_time = %.3f ms | max_abs_err = %f\n", kernelName.c_str(), avg_ms, (float)max_abs); // cleanup clReleaseMemObject(bufA); - clReleaseMemObject(bufB); + clReleaseMemObject(bufQB); + clReleaseMemObject(bufSB); clReleaseMemObject(bufC); clReleaseKernel(kernel); clReleaseProgram(program); @@ -370,12 +397,117 @@ double KernelTest(const string &kernelName, return avg_ms; } +// 将 Aq(int8, M*K) 打包成 RGBA 像素，宽=K/4，高=M +static void packAqToRGBA(const vector& Aq, int m, int k, vector& rgba) +{ + rgba.resize((size_t)m * (size_t)(k/4) * 4); + size_t idx = 0; + for (int r = 0; r < m; ++r) + { + const char* row = Aq.data() + (size_t)r * (size_t)k; + for (int x = 0; x < k; x += 4) + { + rgba[idx++] = row[x+0]; + rgba[idx++] = row[x+1]; + rgba[idx++] = row[x+2]; + rgba[idx++] = row[x+3]; + } + } +} + +double KernelTestImage(const string &kernelName, + cl_context context, + cl_device_id device, + cl_command_queue queue, + const vector& Aq, + const vector& Ad, + const vector& qB, + const vector& sB, + vector& C_gpu, + const vector& C_ref, + int m, int k, float alpha, float beta) +{ + auto src = loadTextFile(kernelPath); + const char *kernelSource = src.c_str(); + const size_t src_len = src.size(); + cl_int err; + cl_program program = clCreateProgramWithSource(context, 1, &kernelSource, &src_len, &err); + checkErr(err, "clCreateProgramWithSource"); + const char *buildOptions = "-cl-std=CL2.0"; + err = clBuildProgram(program, 1, &device, buildOptions, NULL, NULL); + if (err != CL_SUCCESS) { + size_t log_size = 0; clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, 0, NULL, &log_size); + string log(log_size, '\0'); clGetProgramBuildInfo(program, device, CL_PROGRAM_BUILD_LOG, log_size, &log[0], NULL); + cerr << "Build failed:\n" << log << endl; exit(1); + } + cl_kernel kernel = clCreateKernel(program, kernelName.c_str(), &err); + checkErr(err, "clCreateKernel img"); + + // 创建 image2D — 使用 CL_R/CL_RGBA 与 INT 格式 + int width = k / 4, height = m; + vector rgba; packAqToRGBA(Aq, m, k, rgba); + cl_image_format fmt; fmt.image_channel_order = CL_RGBA; fmt.image_channel_data_type = CL_SIGNED_INT8; + cl_image_desc desc = {}; + desc.image_type = CL_MEM_OBJECT_IMAGE2D; + desc.image_width = width; + desc.image_height = height; + desc.image_row_pitch = 0; // let runtime choose + cl_mem imgA = clCreateImage(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, &fmt, &desc, (void*)rgba.data(), &err); + checkErr(err, "clCreateImage A"); + + size_t sizeAd = (size_t)m * (size_t)(k/Block_size) * sizeof(half); + size_t sizeQB = (size_t)k * sizeof(int8_t); + size_t sizeSB = (size_t)(k/Block_size) * sizeof(half); + size_t sizeC = (size_t)m * sizeof(half); + cl_mem bufAd = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeAd, (void*)Ad.data(), &err); + checkErr(err, "clCreateBuffer Ad"); + cl_mem bufQB = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeQB, (void*)qB.data(), &err); + checkErr(err, "clCreateBuffer qB"); + cl_mem bufSB = clCreateBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR, sizeSB, (void*)sB.data(), &err); + checkErr(err, "clCreateBuffer sB"); + cl_mem bufC = clCreateBuffer(context, CL_MEM_READ_WRITE | CL_MEM_COPY_HOST_PTR, sizeC, (void*)C_gpu.data(), &err); + checkErr(err, "clCreateBuffer C"); + + int arg = 0; + clSetKernelArg(kernel, arg++, sizeof(cl_mem), &imgA); + clSetKernelArg(kernel, arg++, sizeof(cl_mem), &bufAd); + clSetKernelArg(kernel, arg++, sizeof(cl_mem), &bufQB); + clSetKernelArg(kernel, arg++, sizeof(cl_mem), &bufSB); + clSetKernelArg(kernel, arg++, sizeof(cl_mem), &bufC); + clSetKernelArg(kernel, arg++, sizeof(int), &m); + clSetKernelArg(kernel, arg++, sizeof(int), &k); + clSetKernelArg(kernel, arg++, sizeof(float), &alpha); + clSetKernelArg(kernel, arg++, sizeof(float), &beta); + + size_t devMaxWG = 0, kernelMaxWG = 0; size_t devMaxItems[3] = {0,0,0}; + clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_GROUP_SIZE, sizeof(devMaxWG), &devMaxWG, NULL); + clGetDeviceInfo(device, CL_DEVICE_MAX_WORK_ITEM_SIZES, sizeof(devMaxItems), &devMaxItems, NULL); + clGetKernelWorkGroupInfo(kernel, device, CL_KERNEL_WORK_GROUP_SIZE, sizeof(kernelMaxWG), &kernelMaxWG, NULL); + size_t lws0 = std::min({(size_t)256, devMaxWG, devMaxItems[0], kernelMaxWG}); if (lws0==0) lws0=1; + size_t lws[1] = {lws0}; size_t gws[1] = {((size_t)m + lws[0]-1)/lws[0]*lws[0]}; + + for (int i=0;i evs(NUM_ITER); + for (int i=0;i C_gpu(n * m); vector C_ref(n * m); - // 量化 A -> qchars + scales + // 量化 A -> qchars + scales（v1:AoS） int blocks = k / Block_size; vector BlockA(m * blocks); quantv1(A, BlockA, m, k, blocks); + // 量化 A -> SoA（v2） + vector Aq; vector Ad; + quantv2(A, Aq, Ad, m, k); + + // 量化激活 B -> qB(int8) + sB(half per-32) + vector qB(k); + vector sB(blocks); + quantB_q8_per32(B, qB, sB, k); // CPU 计算 gemv(C_ref, A, B, m, n, k, alpha, beta); @@ -425,9 +565,13 @@ int main() cl_command_queue queue = clCreateCommandQueueWithProperties(context, device, props, &err); checkErr(err, "clCreateCommandQueueWithProperties"); - // Kernel test - KernelTest("gemv_q8_base", kernelSource, context, device, queue, - BlockA, B, C_gpu, C_ref, m, n, k, alpha, beta); + // Kernel test: AoS buffer + KernelTest("gemv_q8_base", context, device, queue, + BlockA, qB, sB, C_gpu, C_ref, m, n, k, alpha, beta); + // Kernel test: image2D weights (SoA) + std::fill(C_gpu.begin(), C_gpu.end(), half(0)); + KernelTestImage("gemv_q8_soa_img", context, device, queue, + Aq, Ad, qB, sB, C_gpu, C_ref, m, k, alpha, beta); // cleanup clReleaseCommandQueue(queue); clReleaseContext(context); diff --git a/kernel.cl b/kernel.cl new file mode 100644 index 0000000..3c590d7 --- /dev/null +++ b/kernel.cl @@ -0,0 +1,119 @@ +#define CL_TARGET_OPENCL_VERSION 200 +#pragma OPENCL EXTENSION cl_khr_fp16 : enable +typedef struct +{ + half d; + char qs[32]; +} BlockQ8_0; +// 激活值量化到低位：B 拆分为 qB(int8) + sB(half, 每32元素一块) +__kernel void gemv_q8_base(__global const BlockQ8_0 *A, + __global const char *qB, + __global const half *sB, + __global half *C, + int as, int ars, int acs, int bs, int brs, int bcs, + int cs, int crs, int ccs, + int M, int N, int K, float alpha, float beta) +{ + + int row_id = get_global_id(0); + // 允许 GWS 向上取整：越界线程直接返回 + if (row_id >= M) + return; + + // 本地缓存当前块的 qB[32]，供同一工作组复用 + __local char l_qB[32]; + + BlockQ8_0 valueA; + float sum = 0.0f; + + for (int i = 0; i < K / 32; i++) + { + // 组内协作：加载 qB 的第 i 个 32 元素块到本地内存 + int lid = get_local_id(0); + if (lid < 32) + { + l_qB[lid] = qB[(i * 32 + lid) * brs]; + } + barrier(CLK_LOCAL_MEM_FENCE); + + // 读取 A 的对应块 + valueA = *(A + row_id * ars + i * acs); + + // 使用 char4 向量分段做乘加，减少循环开销 + int acci = 0; + #pragma unroll + for (int t = 0; t < 8; ++t) + { + char4 aa = vload4(t, valueA.qs); + char4 bb = vload4(t, l_qB); + acci += (int)aa.s0 * (int)bb.s0; + acci += (int)aa.s1 * (int)bb.s1; + acci += (int)aa.s2 * (int)bb.s2; + acci += (int)aa.s3 * (int)bb.s3; + } + + // 缩放：权重块 d 与激活块 sB[i] + half sb = sB[i]; + sum += (float)acci * (float)valueA.d * (float)sb; + + barrier(CLK_LOCAL_MEM_FENCE); + } + + __global half *p = C + row_id * crs; + if (beta != 0) + *p = (half)(beta * (*p) + alpha * sum); + else + *p = (half)(alpha * sum); +} + +// 使用 image2D 存储权重的 SoA 版：A 的 int8 权重放入 RGBA 像素（每像素4个有符号int8） +__kernel void gemv_q8_soa_img(read_only image2d_t AqImg, + __global const half *Ad, + __global const char *qB, + __global const half *sB, + __global half *C, + int M, int K, float alpha, float beta) +{ + int row = get_global_id(0); + if (row >= M) return; + + int blocks = K >> 5; // K/32 + int baseX = 0; // 每个块占 8 个像素（8*4 = 32） + __local char l_qB[32]; + float sum = 0.0f; + + for (int bi = 0; bi < blocks; ++bi) + { + int lid = get_local_id(0); + if (lid < 32) + l_qB[lid] = qB[bi * 32 + lid]; + barrier(CLK_LOCAL_MEM_FENCE); + + int x0 = baseX + bi * 8; + int acci = 0; + #pragma unroll + for (int t = 0; t < 8; ++t) + { + int2 coord = (int2)(x0 + t, row); + int4 px = read_imagei(AqImg, coord); + char qb0 = l_qB[4 * t + 0]; + char qb1 = l_qB[4 * t + 1]; + char qb2 = l_qB[4 * t + 2]; + char qb3 = l_qB[4 * t + 3]; + acci += px.x * (int)qb0; + acci += px.y * (int)qb1; + acci += px.z * (int)qb2; + acci += px.w * (int)qb3; + } + + float d = (float)Ad[row * blocks + bi]; + float sb = (float)sB[bi]; + sum += (float)acci * d * sb; + + barrier(CLK_LOCAL_MEM_FENCE); + } + + half oldv = C[row]; + C[row] = (beta != 0.0f) ? (half)(beta * (float)oldv + alpha * sum) + : (half)(alpha * sum); +} diff --git a/quant_gemv_test_origin b/quant_gemv_test_origin new file mode 160000 index 0000000..d043d9f --- /dev/null +++ b/quant_gemv_test_origin @@ -0,0 +1 @@ +Subproject commit d043d9fded6701759727e2f3c6e79aeb5a1f5b77 diff --git a/readme.md b/readme.md new file mode 100644 index 0000000..f069f9d --- /dev/null +++ b/readme.md @@ -0,0 +1,2 @@ +* build.sh:linux 环境编译文件 +* kernel.cl:整理代码结构，内核代码单独放置