当前位置：首页 > news >正文

gpu driven：vello新执行流程

news 2025/10/28 11:48:14

1. 路径标签归约 (pathtag_reduce)：

wgsl
@compute @workgroup_size(256)
fn main(
@builtin(global_invocation_id) global_id: vec3<u32>,
@builtin(local_invocation_id) local_id: vec3<u32>,
) {
let ix = global_id.x;
let tag_word = scene[config.pathtag_base + ix];
var agg = reduce_tag(tag_word);
sh_scratch[local_id.x] = agg;
for (var i = 0u; i < firstTrailingBit(WG_SIZE); i += 1u) {
workgroupBarrier();
if local_id.x + (1u << i) < WG_SIZE {
let other = sh_scratch[local_id.x + (1u << i)];
agg = combine_tag_monoid(agg, other);
}
workgroupBarrier();
sh_scratch[local_id.x] = agg;
}
if local_id.x == 0u {
reduced[ix >> LG_WG_SIZE] = agg;
}
}
代码详细分析：

工作组大小: WG_SIZE = 256u，每个工作组处理 256 个标签字（每个标签字包含 4 个标签）
读取标签: tag_word = scene[config.pathtag_base + ix] 从场景数据中读取一个 32 位的标签字
归约操作: reduce_tag(tag_word) 将这个标签字转换为 TagMonoid 结构
并行归约: 使用并行前缀和算法（Parallel Prefix Sum）合并工作组内所有线程的结果：
使用共享内存 sh_scratch 存储中间结果
通过 firstTrailingBit(WG_SIZE) 次迭代完成归约（对于 256 线程，需要 8 次迭代）
每次迭代中，线程与其相距 1 << i 的线程合并结果
输出结果: 只有线程 0 将归约结果写入 reduced 缓冲区
wgsl
fn reduce_tag(tag_word: u32) -> TagMonoid {
var c: TagMonoid;
let point_count = tag_word & 0x3030303u;
c.pathseg_ix = countOneBits((point_count * 7u) & 0x4040404u);
c.trans_ix = countOneBits(tag_word & (PATH_TAG_TRANSFORM * 0x1010101u));
let n_points = point_count + ((tag_word >> 2u) & 0x1010101u);
var a = n_points + (n_points & (((tag_word >> 3u) & 0x1010101u) * 15u));
a += a >> 8u;
a += a >> 16u;
c.pathseg_offset = a & 0xffu;
// ...
return c;
}
reduce_tag 函数详细分析：

point_count = tag_word & 0x3030303u - 提取每个标签的点计数位（位 0-1）
c.pathseg_ix = countOneBits((point_count * 7u) & 0x4040404u) - 计算路径段数量
(point_count * 7u) & 0x4040404u 会将每个标签的点计数位转换为标志位
countOneBits 计算有多少个非零标签
c.trans_ix = countOneBits(tag_word & (PATH_TAG_TRANSFORM * 0x1010101u)) - 计算变换标签数量
n_points 和后续计算确定路径段数据的偏移量
2. 路径扫描 (pathtag_scan)
路径扫描阶段使用前缀和算法计算每个路径标签在整个数据流中的位置：

wgsl
fn combine_tag_monoid(a: TagMonoid, b: TagMonoid) -> TagMonoid {
var c: TagMonoid;
c.trans_ix = a.trans_ix + b.trans_ix;
c.pathseg_ix = a.pathseg_ix + b.pathseg_ix;
c.pathseg_offset = a.pathseg_offset + b.pathseg_offset;
// ...
return c;
}
这个函数定义了如何合并两个 TagMonoid，简单地将各个计数器相加。

3. 曲线展平 (flatten)
曲线展平将贝塞尔曲线转换为线段：

wgsl
fn flatten_cubic(cubic: CubicPoints, path_ix: u32, local_to_device: Transform, offset: f32) {
// ...
let err_v = 3.0 * (p2 - p1) + p0 - p3;
let err = dot(err_v, err_v);
let ACCURACY = 0.25;
let Q_ACCURACY = ACCURACY * 0.1;
let REM_ACCURACY = ACCURACY - Q_ACCURACY;
let MAX_HYPOT2 = 432.0 * Q_ACCURACY * Q_ACCURACY;
let scaled_sqrt_tol = sqrt(REM_ACCURACY / scale);
var n_quads = max(u32(ceil(pow(err * (1.0 / MAX_HYPOT2), 1.0 / 6.0)) * scale), 1u);
n_quads = min(n_quads, MAX_QUADS);
// ...
}
详细分析：

误差计算: 通过计算 3.0 * (p2 - p1) + p0 - p3 来估计三次贝塞尔曲线与二次贝塞尔曲线的近似误差
细分计算: 根据误差计算需要将曲线细分为多少个二次贝塞尔曲线段
进一步细分: 对每个二次贝塞尔曲线段进行进一步细分以满足精度要求
输出线段: output_line_with_transform 将最终的线段写入 lines 缓冲区
wgsl
fn output_line_with_transform(path_ix: u32, p0: vec2f, p1: vec2f, transform: Transform) {
let line_ix = atomicAdd(&bump.lines, 1u);
write_line_with_transform(line_ix, path_ix, p0, p1, transform);
}
使用原子操作 atomicAdd 从凸起分配器中分配内存，确保线程安全。

4. 绘制对象处理 (draw_leaf)
绘制对象处理阶段将绘制标签转换为实际的绘制命令：

wgsl
if tag_word == DRAWTAG_FILL_COLOR || tag_word == DRAWTAG_FILL_LIN_GRADIENT ||
tag_word == DRAWTAG_FILL_RAD_GRADIENT || tag_word == DRAWTAG_FILL_IMAGE
{
// ...
switch tag_word {
// DRAWTAG_FILL_COLOR
case 0x44u: {
info[di] = draw_flags;
}
// DRAWTAG_FILL_LIN_GRADIENT
case 0x114u: {
info[di] = draw_flags;
var p0 = bitcast<vec2<f32>>(vec2(scene[dd + 1u], scene[dd + 2u]));
var p1 = bitcast<vec2<f32>>(vec2(scene[dd + 3u], scene[dd + 4u]));
p0 = transform_apply(transform, p0);
p1 = transform_apply(transform, p1);
let dxy = p1 - p0;
let scale = 1.0 / dot(dxy, dxy);
let line_xy = dxy * scale;
let line_c = -dot(p0, line_xy);
info[di + 1u] = bitcast<u32>(line_xy.x);
info[di + 2u] = bitcast<u32>(line_xy.y);
info[di + 3u] = bitcast<u32>(line_c);
}
// ...
}
}
详细分析：

颜色填充: 简单地将绘制标志存储到 info 缓冲区
线性渐变: 计算渐变线的参数，包括方向向量和偏移
径向渐变: 实现复杂的径向渐变算法，包括焦点处理
5. 分箱 (binning)
分箱阶段将绘制对象分配到对应的瓦片：

wgsl
// 将边界框转换为瓦片坐标
x0 = i32(floor(bbox.x * SX));
y0 = i32(floor(bbox.y * SY));
x1 = i32(ceil(bbox.z * SX));
y1 = i32(ceil(bbox.w * SY));

// 使用位图跟踪每个瓦片中的对象
let my_slice = local_id.x / 32u;
let my_mask = 1u << (local_id.x & 31u);
while y < y1 {
atomicOr(&sh_bitmaps[my_slice][y * width_in_bins + x], my_mask);
// ...
}
详细分析：

坐标转换: 将边界框坐标转换为瓦片索引
位图标记: 使用位图（bitmaps）标记哪些瓦片被当前绘制对象覆盖
内存分配: 使用原子操作 atomicAdd(&bump.binning, element_count) 从凸起分配器分配内存
数据写入: 将绘制对象索引写入对应的分箱数据缓冲区
6. 瓦片分配 (tile_alloc)
为每个路径分配瓦片存储：

wgsl
let tile_count = (ux1 - ux0) * (uy1 - uy0);
var total_tile_count = tile_count;
sh_tile_count[local_id.x] = tile_count;
for (var i = 0u; i < firstTrailingBit(WG_SIZE); i += 1u) {
workgroupBarrier();
if local_id.x >= (1u << i) {
total_tile_count += sh_tile_count[local_id.x - (1u << i)];
}
workgroupBarrier();
sh_tile_count[local_id.x] = total_tile_count;
}
if local_id.x == WG_SIZE - 1u {
let count = sh_tile_count[WG_SIZE - 1u];
var offset = atomicAdd(&bump.tile, count);
if offset + count > config.tiles_size {
offset = 0u;
atomicOr(&bump.failed, STAGE_TILE_ALLOC);
}
paths[drawobj_ix].tiles = offset;
}
详细分析：

瓦片计数: 计算每个绘制对象覆盖的瓦片数量
并行归约: 使用并行前缀和算法计算总瓦片需求
内存分配: 使用原子操作 atomicAdd(&bump.tile, count) 从凸起分配器分配瓦片内存
失败处理: 如果分配失败，标记失败标志以便后续重新分配
7. 粗栅格化 (coarse)
粗栅格化阶段生成每瓦片的命令列表：

wgsl
fn write_path(tile: Tile, tile_ix: u32, draw_flags: u32) {
// 分配线段内存
let n_segs = tile.segment_count_or_ix;
if n_segs != 0u {
var seg_ix = atomicAdd(&bump.segments, n_segs);
tiles[tile_ix].segment_count_or_ix = ~seg_ix;
// 写入 PTCL（Per-Tile Command List）命令
alloc_cmd(4u);
ptcl[cmd_offset] = CMD_FILL;
let even_odd = (draw_flags & DRAW_INFO_FLAGS_FILL_RULE_BIT) != 0u;
let size_and_rule = (n_segs << 1u) | u32(even_odd);
let fill = CmdFill(size_and_rule, seg_ix, tile.backdrop);
ptcl[cmd_offset + 1u] = fill.size_and_rule;
ptcl[cmd_offset + 2u] = fill.seg_data;
ptcl[cmd_offset + 3u] = u32(fill.backdrop);
cmd_offset += 4u;
}
}
详细分析：

线段分配: 为每个瓦片分配线段内存
命令生成: 生成填充命令（CMD_FILL）并写入 PTCL 缓冲区
参数编码: 将线段数量、奇偶填充规则等信息编码到命令中
8. 细栅格化 (fine)
最终的像素渲染阶段：

wgsl
fn fill_path_ms(fill: CmdFill, local_id: vec2<u32>, result: ptr<function, array<f32, PIXELS_PER_THREAD>>) {
let even_odd = (fill.size_and_rule & 1u) != 0u;
if even_odd {
fill_path_ms_evenodd(fill, local_id, result);
return;
}
let n_segs = fill.size_and_rule >> 1u;
// 初始化缠绕数数组
if th_ix < 64u {
if th_ix < 4u {
atomicStore(&sh_winding_y[th_ix], 0x80808080u);
}
atomicStore(&sh_winding[th_ix], 0x80808080u);
}
// 处理每个线段
for (var batch = 0u; batch < n_batch; batch++) {
let seg_ix = batch * WG_SIZE + th_ix;
let seg_off = fill.seg_data + seg_ix;
var count = 0u;
if th_ix < slice_size {
let segment = segments[seg_off];
let xy0 = segment.point0;
let xy1 = segment.point1;
var y_edge_f = f32(TILE_HEIGHT);
var delta = select(-1, 1, xy1.x <= xy0.x);
// 计算跨越的像素数
if !(xy0.y == xy1.y && xy0.y == floor(xy0.y)) {
count = span(xy0.x, xy1.x) + span(xy0.y, xy1.y) - 1u;
}
// 更新缠绕数
let y_edge = u32(ceil(y_edge_f));
if y_edge < TILE_HEIGHT {
atomicAdd(&sh_winding_y[y_edge >> 2u], u32(delta) << ((y_edge & 3u) << 3u));
}
}
// ...
}
}
详细分析：