471 lines
13 KiB
Common Lisp
471 lines
13 KiB
Common Lisp
|
||
/**
|
||
* Zigzag 映射：将有符号字节映射为非负字节
|
||
* 0→0, -1→1, 1→2, -2→3, 2→4, ...
|
||
*/
|
||
uchar zigzag_map(char x) {
|
||
int v = (int)x; // char 转为 int（保留符号）
|
||
return (uchar)((v >= 0) ? (v << 1) : ((-v << 1) - 1));
|
||
}
|
||
|
||
/**
|
||
* Zigzag 逆映射：将非负字节还原为有符号字节
|
||
* 0→0, 1→-1, 2→1, 3→-2, 4→2, ...
|
||
*/
|
||
char zigzag_unmap(uchar y) {
|
||
int v = (int)y; // 转为无符号 0~255
|
||
return (char)((v & 1) == 0 ? (v >> 1) : (-((v + 1) >> 1)));
|
||
}
|
||
|
||
|
||
|
||
|
||
/**
|
||
* 像素安全加法：RGBA 四个通道分别相加，结果裁剪到 0~255
|
||
*/
|
||
int pixelAdd(int a, int b) {
|
||
int ar = (a >> 16) & 0xFF;
|
||
int ag = (a >> 8) & 0xFF;
|
||
int ab = a & 0xFF;
|
||
int aa = (a >> 24) & 0xFF;
|
||
|
||
int br = (b >> 16) & 0xFF;
|
||
int bg = (b >> 8) & 0xFF;
|
||
int bb = b & 0xFF;
|
||
int ba = (b >> 24) & 0xFF;
|
||
|
||
int r = ar + br;
|
||
int g = ag + bg;
|
||
int bv = ab + bb;
|
||
int av = aa + ba;
|
||
|
||
// 用 &0xFF 替代 min/max，保留低 8 位（自动溢出，等同于裁剪）
|
||
return ((av & 0xFF) << 24) |
|
||
((r & 0xFF) << 16) |
|
||
((g & 0xFF) << 8) |
|
||
(bv & 0xFF);
|
||
}
|
||
/**
|
||
* 像素安全减法（等效于 Java 的 pixelSub）
|
||
* 逐通道相减，结果裁剪到 0~255，组装回 ARGB
|
||
*/
|
||
int pixelSub(int a, int b) {
|
||
int ar = (a >> 16) & 0xFF;
|
||
int ag = (a >> 8) & 0xFF;
|
||
int ab = a & 0xFF;
|
||
int aa = (a >> 24) & 0xFF;
|
||
|
||
int br = (b >> 16) & 0xFF;
|
||
int bg = (b >> 8) & 0xFF;
|
||
int bb = b & 0xFF;
|
||
int ba = (b >> 24) & 0xFF;
|
||
|
||
int r = ar - br;
|
||
int g = ag - bg;
|
||
int bv = ab - bb;
|
||
int av = aa - ba;
|
||
|
||
return ((av & 0xFF) << 24) |
|
||
((r & 0xFF) << 16) |
|
||
((g & 0xFF) << 8) |
|
||
(bv & 0xFF);
|
||
}
|
||
|
||
|
||
/**
|
||
* 位平面编码（矩阵转置风格）- 等效于 Java 版本
|
||
*
|
||
* 核心逻辑：
|
||
* 对于输出数组的每个字节，依次从输入数组的8个位平面各取1个bit
|
||
* 即：输出字节的 bit i = 输入数组第 i 个位平面的当前位
|
||
*
|
||
* 输出数组的每个字节，由8个位平面的各1个bit组成
|
||
* 等价于：把 8×len 的 bit 矩阵转置为 len×8 的 bit 矩阵
|
||
*/
|
||
__kernel void encodePacked(
|
||
__global const uchar* input, // 输入字节数组
|
||
__global uchar* output, // 输出字节数组（长度与输入相同）
|
||
int len // 输入/输出长度
|
||
) {
|
||
int outIdx = get_global_id(0);
|
||
if (outIdx >= len) return;
|
||
|
||
int outByte = 0;
|
||
|
||
// 从8个位平面各取1个bit，组装成一个字节
|
||
for (int plane = 0; plane < 8; plane++) {
|
||
int bitpos = (outIdx << 3) + plane;
|
||
int inputPos = bitpos % len;
|
||
int bitShift = bitpos / len;
|
||
int bit = (input[inputPos] >> (7 - bitShift)) & 1;
|
||
outByte |= (bit << plane);
|
||
}
|
||
|
||
output[outIdx] = (uchar)outByte;
|
||
}
|
||
|
||
__kernel void decodePacked(__global const unsigned char* datain,
|
||
__global unsigned char* dataout,
|
||
int size) {
|
||
int outIdx = get_global_id(0);
|
||
|
||
// 边界检查
|
||
if (outIdx >= size) {
|
||
return;
|
||
}
|
||
|
||
int outByte = 0;
|
||
|
||
// 从 8 个位平面各取 1 个 bit，组装成一个字节
|
||
// plane 0 → bit7, plane 1 → bit6, ..., plane 7 → bit0
|
||
for (int plane = 0; plane < 8; plane++) {
|
||
int bitpos = plane * size + outIdx;
|
||
int inputpos = bitpos >> 3;
|
||
int inputshift = bitpos & 0b111;
|
||
int bit = (datain[inputpos] >> inputshift) & 1;
|
||
outByte |= (bit << (7 - plane));
|
||
}
|
||
|
||
dataout[outIdx] = (unsigned char)outByte;
|
||
}
|
||
|
||
/**
|
||
* 序列化 + Zigzag 映射（RGBA 交错）
|
||
*/
|
||
__kernel void serializeRGBAWithZigzag(
|
||
__global int* pixels,
|
||
__global uchar* output,
|
||
int size
|
||
) {
|
||
int i = get_global_id(0);
|
||
if (i >= size) return;
|
||
|
||
int pixel = pixels[i];
|
||
|
||
char r = (char)((pixel >> 16) & 0xFF);
|
||
char g = (char)((pixel >> 8) & 0xFF);
|
||
char b = (char)(pixel & 0xFF);
|
||
char a = (char)((pixel >> 24) & 0xFF);
|
||
|
||
int idx = i * 4;
|
||
output[idx + 0] = zigzag_map(r);
|
||
output[idx + 1] = zigzag_map(g);
|
||
output[idx + 2] = zigzag_map(b);
|
||
output[idx + 3] = zigzag_map(a);
|
||
}
|
||
// 对应 deserializeRGBAWithZigzag
|
||
__kernel void deserializeRGBAWithZigzag(
|
||
__global const uchar* data, // 输入：Zigzag 映射后的字节流（RGBA 交错）
|
||
__global int* pixels, // 输出：像素数组（ARGB）
|
||
int size // 像素数量 (w * h)
|
||
) {
|
||
int i = get_global_id(0);
|
||
if (i >= size) return;
|
||
|
||
int idx = i * 4;
|
||
|
||
// 读取并逆映射四个通道
|
||
uchar r_raw = data[idx];
|
||
uchar g_raw = data[idx + 1];
|
||
uchar b_raw = data[idx + 2];
|
||
uchar a_raw = data[idx + 3];
|
||
|
||
char r = zigzag_unmap(r_raw);
|
||
char g = zigzag_unmap(g_raw);
|
||
char b = zigzag_unmap(b_raw);
|
||
char a = zigzag_unmap(a_raw);
|
||
|
||
// 组装成 ARGB (Java 的 int 格式)
|
||
// 注意：OpenCL 的 int 是 32 位有符号，和 Java 一致
|
||
pixels[i] = ((int)(a & 0xFF) << 24) |
|
||
((int)(r & 0xFF) << 16) |
|
||
((int)(g & 0xFF) << 8) |
|
||
(int)(b & 0xFF);
|
||
}
|
||
/**
|
||
* 序列化 + Zigzag 映射（平面 RGBA）
|
||
*/
|
||
__kernel void serializePlannarRGBAWithZigzag(
|
||
__global int* pixels,
|
||
__global uchar* output,
|
||
int size
|
||
) {
|
||
int i = get_global_id(0);
|
||
if (i >= size) return;
|
||
|
||
int pixel = pixels[i];
|
||
|
||
char r = (char)((pixel >> 16) & 0xFF);
|
||
char g = (char)((pixel >> 8) & 0xFF);
|
||
char b = (char)(pixel & 0xFF);
|
||
char a = (char)((pixel >> 24) & 0xFF);
|
||
|
||
int rOffset = 0;
|
||
int gOffset = size;
|
||
int bOffset = size * 2;
|
||
int aOffset = size * 3;
|
||
|
||
output[rOffset + i] = zigzag_map(r);
|
||
output[gOffset + i] = zigzag_map(g);
|
||
output[bOffset + i] = zigzag_map(b);
|
||
output[aOffset + i] = zigzag_map(a);
|
||
}
|
||
// 对应 deserializePlannarRGBAWithZigzag
|
||
__kernel void deserializePlannarRGBAWithZigzag(
|
||
__global const uchar* data, // 输入：Zigzag 映射后的字节流（RRRR...GGGG...BBBB...AAAA...）
|
||
__global int* pixels, // 输出：像素数组（ARGB）
|
||
int size // 像素数量 (w * h)
|
||
) {
|
||
int i = get_global_id(0);
|
||
if (i >= size) return;
|
||
|
||
int rOffset = 0;
|
||
int gOffset = size;
|
||
int bOffset = size * 2;
|
||
int aOffset = size * 3;
|
||
|
||
// 从四个平面分别读取并逆映射
|
||
uchar r_raw = data[rOffset + i];
|
||
uchar g_raw = data[gOffset + i];
|
||
uchar b_raw = data[bOffset + i];
|
||
uchar a_raw = data[aOffset + i];
|
||
|
||
char r = zigzag_unmap(r_raw);
|
||
char g = zigzag_unmap(g_raw);
|
||
char b = zigzag_unmap(b_raw);
|
||
char a = zigzag_unmap(a_raw);
|
||
|
||
// 组装成 ARGB
|
||
pixels[i] = ((int)(a & 0xFF) << 24) |
|
||
((int)(r & 0xFF) << 16) |
|
||
((int)(g & 0xFF) << 8) |
|
||
(int)(b & 0xFF);
|
||
}
|
||
|
||
/**
|
||
* 水平预测（非原地，完美并行）
|
||
* residual = pixelSub(current, left)
|
||
* 每个线程独立处理一个像素
|
||
*/
|
||
__kernel void horizontalPredictParallel(
|
||
__global const int* input, // 原始像素
|
||
__global int* output, // 残差输出
|
||
int w,
|
||
int h
|
||
) {
|
||
int idx = get_global_id(0);
|
||
int size = w * h;
|
||
if (idx >= size) return;
|
||
|
||
int x = idx % w;
|
||
|
||
if (x == 0) {
|
||
// 第一列：残差 = 原值（没有左邻居）
|
||
output[idx] = input[idx];
|
||
} else {
|
||
int left = input[idx - 1];
|
||
output[idx] = pixelSub(input[idx], left);
|
||
}
|
||
}
|
||
|
||
/**
|
||
* 垂直预测（非原地，完美并行）
|
||
* residual = pixelSub(current, above)
|
||
* 每个线程独立处理一个像素
|
||
*/
|
||
__kernel void verticalPredictParallel(
|
||
__global const int* input, // 原始像素（或水平预测后的残差）
|
||
__global int* output, // 残差输出
|
||
int w,
|
||
int h
|
||
) {
|
||
int idx = get_global_id(0);
|
||
int size = w * h;
|
||
if (idx >= size) return;
|
||
|
||
int y = idx / w;
|
||
|
||
if (y == 0) {
|
||
// 第一行：残差 = 原值（没有上邻居）
|
||
output[idx] = input[idx];
|
||
} else {
|
||
int above = input[idx - w];
|
||
output[idx] = pixelSub(input[idx], above);
|
||
}
|
||
}
|
||
|
||
/**
|
||
* 水平预测（使用 pixelSub）
|
||
* residual = pixelSub(current, left)
|
||
* 从右向左，保证 left 是原始值
|
||
*/
|
||
__kernel void horizontalPredictInPlace(
|
||
__global int* data,
|
||
int w,
|
||
int h
|
||
) {
|
||
int row = get_global_id(0);
|
||
if (row >= h) return;
|
||
|
||
int base = row * w;
|
||
|
||
// 从右向左（第一列不变）
|
||
for (int x = w - 1; x >= 1; x--) {
|
||
int idx = base + x;
|
||
int left = data[idx - 1];
|
||
data[idx] = pixelSub(data[idx], left);
|
||
}
|
||
}
|
||
/**
|
||
* 水平逆预测（原地操作）
|
||
* 每个工作组处理一行，行内从左到右串行
|
||
*
|
||
* 数据依赖：左侧像素 (x-1, y) 必须先被还原
|
||
*//**
|
||
* 水平逆预测（寄存器优化版）
|
||
*
|
||
* 核心优化：预测值在寄存器中传递，减少显存读取
|
||
*/
|
||
__kernel void horizontalInverseInPlace(
|
||
__global int* data,
|
||
int w,
|
||
int h
|
||
) {
|
||
int row = get_global_id(0);
|
||
if (row >= h) return;
|
||
|
||
int base = row * w;
|
||
|
||
// 第一列保持不变
|
||
// 直接用 data[base] 作为初始预测值
|
||
int pred = data[base]; // ✅ 只读一次显存
|
||
|
||
// 行内从左到右串行，但 pred 在寄存器中传递
|
||
for (int x = 1; x < w; x++) {
|
||
int idx = base + x;
|
||
int residual = data[idx]; // ✅ 只读残差
|
||
int result = pixelAdd(pred, residual);
|
||
data[idx] = result; // ✅ 只写一次显存
|
||
pred = result; // ✅ 寄存器传递（下次循环直接使用）
|
||
}
|
||
|
||
}
|
||
/**
|
||
* 垂直预测（使用 pixelSub）
|
||
* residual = pixelSub(current, above)
|
||
* 从下向上，保证 above 是原始值
|
||
*/
|
||
__kernel void verticalPredictInPlace(
|
||
__global int* data,
|
||
int w,
|
||
int h
|
||
) {
|
||
int col = get_global_id(0);
|
||
if (col >= w) return;
|
||
|
||
// 从下向上（第一行不变）
|
||
for (int y = h - 1; y >= 1; y--) {
|
||
int idx = y * w + col;
|
||
int above = data[(y - 1) * w + col];
|
||
data[idx] = pixelSub(data[idx], above);
|
||
}
|
||
}
|
||
/**
|
||
* 垂直逆预测（原地操作）
|
||
* 每个工作组处理一列，列内从上到下串行
|
||
*
|
||
* 数据依赖：上方像素 (x, y-1) 必须先被还原
|
||
*//**
|
||
* 垂直逆预测（寄存器优化版）
|
||
*/
|
||
__kernel void verticalInverseInPlace(
|
||
__global int* data,
|
||
int w,
|
||
int h
|
||
) {
|
||
int col = get_global_id(0);
|
||
if (col >= w) return;
|
||
|
||
// 第一行：用 data[col] 作为初始预测值
|
||
int pred = data[col]; // ✅ 只读一次显存
|
||
|
||
for (int y = 1; y < h; y++) {
|
||
int idx = y * w + col;
|
||
int residual = data[idx]; // ✅ 只读残差
|
||
int result = pixelAdd(pred, residual);
|
||
data[idx] = result;
|
||
pred = result; // ✅ 寄存器传递
|
||
}
|
||
}
|
||
/**
|
||
* RGBA 颜色变换（原地版本）- 等效于 Java 的 colorTransformRGBAInPlace
|
||
*
|
||
* 编码：R' = R - G, G' = G, B' = B - G, A' = A
|
||
* 结果存储为：A | R' | G | B'（与 Java 版本完全一致）
|
||
*/
|
||
__kernel void colorTransformRGBAInPlace(
|
||
__global int* data, // 输入/输出像素数组
|
||
int size // 像素总数 (w * h)
|
||
) {
|
||
int idx = get_global_id(0);
|
||
if (idx >= size) return;
|
||
|
||
int pixel = data[idx];
|
||
|
||
int r = (pixel >> 16) & 0xFF;
|
||
int g = (pixel >> 8) & 0xFF;
|
||
int b = pixel & 0xFF;
|
||
int a = (pixel >> 24) & 0xFF;
|
||
|
||
int rg = r - g; // 范围 -255 ~ 255
|
||
int bg = b - g; // 范围 -255 ~ 255
|
||
|
||
// 与 Java 版本完全一致：
|
||
// ((a & 0xFF) << 24) | ((rg & 0xFF) << 16) | ((g & 0xFF) << 8) | (bg & 0xFF)
|
||
data[idx] = ((a & 0xFF) << 24) |
|
||
((rg & 0xFF) << 16) |
|
||
((g & 0xFF) << 8) |
|
||
(bg & 0xFF);
|
||
}
|
||
|
||
/**
|
||
* 颜色逆变换（RGBA，逐像素并行）- 与 colorTransformRGBAInPlace 完全对应
|
||
*
|
||
* Java 版本：
|
||
* a = (packed >> 24) & 0xFF
|
||
* rg = (packed >> 16) & 0xFF
|
||
* g = (packed >> 8) & 0xFF
|
||
* bg = packed & 0xFF
|
||
* r = g + rg
|
||
* b = g + bg
|
||
* result = (a << 24) | (r << 16) | (g << 8) | b
|
||
*/
|
||
__kernel void colorInverseRGBAInPlace(
|
||
__global int* data, // 输入：颜色差分数据，输出：还原后的 RGBA
|
||
int size // 像素总数 (w * h)
|
||
) {
|
||
int idx = get_global_id(0);
|
||
if (idx >= size) return;
|
||
|
||
int packed = data[idx];
|
||
|
||
// 提取各通道
|
||
int a = (packed >> 24) & 0xFF;
|
||
int rg = (packed >> 16) & 0xFF;
|
||
int g = (packed >> 8) & 0xFF;
|
||
int bg = packed & 0xFF;
|
||
|
||
// 还原 R 和 B
|
||
int r = g + rg;
|
||
int b = g + bg;
|
||
|
||
// 裁剪到 0~255（用 &0xFF 保留低 8 位）
|
||
r = r & 0xFF;
|
||
b = b & 0xFF;
|
||
|
||
// 组装回 ARGB
|
||
data[idx] = ((a & 0xFF) << 24) |
|
||
((r & 0xFF) << 16) |
|
||
((g & 0xFF) << 8) |
|
||
(b & 0xFF);
|
||
} |