Vulkan 不认识 ONNX，也不知道什么是卷积、MatMul、Attention。
那么，一个 ONNX OCR 模型究竟怎样变成几百次甚至上千次 GPU Compute Dispatch？这篇文章结合开源项目 lw.PPOCR.Vulkan 的真实代码，从 Vulkan 最基础的对象讲起，一直拆到一张图片完成 DET、CLS、REC 和 CTC 的全过程。
项目地址：
本文分析基于仓库当前 main 分支，主要涉及：
src/vulkan_context.cpp
src/graph.cpp
src/onnx_import.cpp
src/graph_optimizer.cpp
src/ocr_pipeline.cpp
src/ocr_host.cpp
src/rec_batch.cpp
src/shaders/
include/lw_ppocr_vulkan.h
当前正式维护版本为 v1.0.1。
项目使用 C++17 实现 PP-OCRv6 Tiny / Small / Medium 的 Vulkan GPU 推理，默认采用 FP32 Vulkan Compute，运行时不依赖 CUDA、ONNX Runtime 或 OpenCV DNN，同时提供：
- C ABI；
- C# WinForms Demo；
- Python 示例；
- HTTP / Web 服务；
- Windows x64 / Linux x64 部署包。
这篇文章不重点讲“怎么点 Demo”，而是回答一个更底层的问题：
Vulkan 到底是怎样把一个神经网络跑起来的？
一、先说结论：Vulkan 本身不会推理 ONNX
这是理解整个项目最重要的一句话。
Vulkan 是一个 GPU 图形与计算 API。
它能帮我们做这些事情：
选择 GPU
分配显存
创建 Buffer
加载 GPU Shader
创建 Compute Pipeline
绑定输入/输出 Buffer
向 GPU 提交计算任务
等待 GPU 完成
把结果读回来
但 Vulkan 完全不知道：
Conv 是什么
DepthwiseConv 是什么
MatMul 是什么
LayerNorm 是什么
Attention 是什么
Softmax 是什么
PP-OCR 是什么
ONNX 又是什么
所以：
ONNX
↓
Vulkan
中间必须存在一个推理 Runtime。
lw.PPOCR.Vulkan 做的其实就是这层工作。
整个关系可以画成：
PP-OCR ONNX
│
▼
┌───────────────────┐
│ ONNX Parser │
│ 模型解析 │
└─────────┬─────────┘
│
▼
┌───────────────────┐
│ Graph Optimizer │
│ 图转换 / 算子融合 │
└─────────┬─────────┘
│
▼
┌───────────────────┐
│ Workspace Planner │
│ Tensor 显存规划 │
└─────────┬─────────┘
│
▼
┌───────────────────┐
│ Kernel Selection │
│ 选择 Compute Shader│
└─────────┬─────────┘
│
▼
┌───────────────────┐
│ Vulkan Runtime │
│ Buffer / Pipeline │
│ Descriptor / Cmd │
└─────────┬─────────┘
│
▼
GPU
所以项目真正有意思的地方，不只是调用 Vulkan API。
而是：
自己实现了从模型图到 GPU Compute 的完整执行链。
二、Vulkan 为什么能跑神经网络？
很多开发者第一次接触 Vulkan，是在游戏渲染领域。
例如：
Vertex Shader
Fragment Shader
Texture
Framebuffer
SwapChain
但 Vulkan 除了 Graphics Pipeline，还有：
Compute Pipeline
Compute Pipeline 不需要窗口，也不需要画三角形。
它的工作就是：
给 GPU 一批数据
+
给 GPU 一个计算程序
=
让大量 GPU 线程并行执行
例如我们有两个数组：
C[i] = A[i] + B[i];
可以写成一个 Compute Shader：
#version 450
layout(local_size_x = 256) in;
layout(binding = 0, std430) readonly buffer A {
float a[];
};
layout(binding = 1, std430) readonly buffer B {
float b[];
};
layout(binding = 2, std430) writeonly buffer C {
float c[];
};
layout(push_constant) uniform Params {
uint count;
} p;
void main()
{
uint i = gl_GlobalInvocationID.x;
if (i < p.count)
c[i] = a[i] + b[i];
}
然后 CPU 调：
vkCmdDispatch(cmd, groupX, 1, 1);
GPU 就会启动大量线程并行完成这些加法。
卷积、矩阵乘法、激活函数，本质上也是把大量数学运算拆成这种并行计算。
区别只是实际 Shader 会复杂很多。
三、CUDA 和 Vulkan 可以怎么对应理解？
熟悉 CUDA 的开发者可以先看这张表。
CUDA Vulkan CUDA Device VkPhysicalDevice CUDA Context VkDevice CUDA Stream VkQueue cudaMalloc vkCreateBuffer + vkAllocateMemory cudaMemcpy vkCmdCopyBuffer / mapped memory CUDA Kernel Compute Shader <<<grid, block>>> vkCmdDispatch Kernel Parameters Descriptor + Push Constants Shared Memory GLSL shared threadIdx.x gl_LocalInvocationID.x blockIdx.x gl_WorkGroupID.x Stream Submit vkQueueSubmit CUDA Event / Sync Fence / Query
所以 Vulkan Compute 并没有神秘到完全无法理解。
例如 Shader：
layout(local_size_x = 128) in;
基本可以理解为：
一个 WorkGroup 有 128 个 GPU Invocation
如果：
vkCmdDispatch(cmd, 100, 20, 1);
那么 GPU 将调度：
100 × 20 × 128
=
256000 个 Invocation
当然实际 GPU 硬件会按照自己的 Wave / Warp / Subgroup 进一步执行。
四、一次 Vulkan 计算，需要哪些核心对象？
我们先把 Vulkan 的核心对象全部串起来。
VkInstance
│
▼
VkPhysicalDevice
│
▼
VkDevice
│
├──────── VkQueue
│
├──────── VkBuffer
│ │
│ ▼
│ VkDeviceMemory
│
├──────── VkShaderModule
│ │
│ ▼
│ VkPipeline
│
├──────── VkDescriptorSet
│
├──────── VkCommandBuffer
│
└──────── VkFence
对应神经网络，可以这样理解：
Vulkan 对象 推理 Runtime 中的意义 VkInstance Vulkan Runtime 入口 VkPhysicalDevice RTX 4060、AMD iGPU 等真实 GPU VkDevice 应用创建的逻辑 GPU VkQueue GPU 任务提交队列 VkBuffer Tensor、权重、输入、输出 VkDeviceMemory Buffer 对应实际显存 VkShaderModule 编译后的 GPU 算子 VkPipeline 可执行的 Compute Kernel VkDescriptorSet Kernel 的 Tensor 参数 VkCommandBuffer 一串神经网络计算命令 VkFence 判断本轮 GPU 是否执行完成
下面结合项目实际代码继续往下看。
五、第一步：创建 Vulkan Instance
代码在：
src/vulkan_context.cpp
项目使用 Vulkan 1.1：
VkApplicationInfo app{
VK_STRUCTURE_TYPE_APPLICATION_INFO
};
app.pApplicationName = "lw.PPOCR.Vulkan";
app.apiVersion = VK_API_VERSION_1_1;
VkInstanceCreateInfo ci{
VK_STRUCTURE_TYPE_INSTANCE_CREATE_INFO
};
ci.pApplicationInfo = &app;
VkInstance instance{};
vkCreateInstance(
&ci,
nullptr,
&instance
);
VkInstance 可以理解为 Vulkan Runtime 的入口。
没有 Instance，后面连 GPU 都无法枚举。
六、第二步：枚举显卡
接下来调用：
uint32_t count = 0;
vkEnumeratePhysicalDevices(
instance,
&count,
nullptr
);
得到 GPU 数量。
然后：
std::vector<VkPhysicalDevice> devices(count);
vkEnumeratePhysicalDevices(
instance,
&count,
devices.data()
);
这样就得到：
GPU 0
GPU 1
GPU 2
...
项目中的设备编号就是这个 Vulkan 枚举顺序。
需要特别注意：
Vulkan Device 0
不一定等于：
CUDA Device 0
多 GPU 机器上尤其不要直接假设两者编号相同。
项目因此专门提供：
lwvk_device_count()
lwvk_device_get()
让应用先枚举，再明确指定设备。
七、第三步：找到 Compute Queue
GPU 内部可能有不同 Queue Family：
Graphics
Compute
Transfer
Video
...
OCR 推理至少需要：
VK_QUEUE_COMPUTE_BIT
项目调用：
vkGetPhysicalDeviceQueueFamilyProperties(...)
然后查找：
families[i].queueFlags & VK_QUEUE_COMPUTE_BIT
lw.PPOCR.Vulkan 当前策略是：
优先选择 dedicated compute queue，没有的话再选择 graphics + compute queue。
找到后创建逻辑设备：
vkCreateDevice(
physical,
&deviceCreateInfo,
nullptr,
&device
);
然后拿到 Queue：
vkGetDeviceQueue(
device,
queueFamily,
0,
&queue
);
现在我们就有：
Physical GPU
│
▼
VkDevice
│
▼
VkQueue
后面的神经网络任务最终都会提交到这个 Queue。
八、第四步：Tensor 在 Vulkan 中是什么？
在 PyTorch 中：
tensor = torch.randn(...)
在 Vulkan Runtime 中没有这种高级 Tensor 对象。
GPU 真正看到的主要还是：
VkBuffer
+
VkDeviceMemory
+
offset
+
size
例如：
Input Tensor
Weight Tensor
Bias Tensor
Intermediate Tensor
Output Tensor
最终都可以映射成 Buffer 中的一段区域。
创建 Buffer：
VkBufferCreateInfo ci{
VK_STRUCTURE_TYPE_BUFFER_CREATE_INFO
};
ci.size = bytes;
ci.usage =
VK_BUFFER_USAGE_STORAGE_BUFFER_BIT |
VK_BUFFER_USAGE_TRANSFER_SRC_BIT |
VK_BUFFER_USAGE_TRANSFER_DST_BIT;
vkCreateBuffer(
device,
&ci,
nullptr,
&buffer
);
接下来查询显存需求：
vkGetBufferMemoryRequirements(
device,
buffer,
&requirements
);
分配 Memory：
vkAllocateMemory(
device,
&allocateInfo,
nullptr,
&memory
);
最后绑定：
vkBindBufferMemory(
device,
buffer,
memory,
0
);
于是：
VkBuffer
│
▼
VkDeviceMemory
建立起来了。
九、Host Buffer 和 Device Local Buffer
做 GPU 推理时，经常存在两类 Buffer。
一种 CPU 可以直接访问，例如：
HOST_VISIBLE
主要负责：
CPU → GPU 输入上传
GPU → CPU 输出回读
另一种放在 GPU 本地显存：
DEVICE_LOCAL
主要负责：
权重
中间 Tensor
计算输出
CPU 可访问的 Memory 可以：
vkMapMemory(...)
拿到地址。
然后：
memcpy(...)
即可写入。
项目中的 Buffer 类将这些细节封装起来：
Buffer::write()
Buffer::read()
并且对于非 HOST_COHERENT Memory，还显式处理：
vkFlushMappedMemoryRanges()
vkInvalidateMappedMemoryRanges()
这非常重要。
Map 成功，不等于 CPU Cache 和 GPU Cache 自动保持一致。
十、权重不能每次推理都上传
如果每处理一张图片都：
加载 ONNX
解析权重
创建 Buffer
上传 Weight
创建 Pipeline
执行
销毁
GPU 推理性能会非常差。
正确方式应该是：
初始化阶段
↓
读取 ONNX
↓
上传权重
↓
权重常驻 GPU
↓
创建 Pipeline
↓
录制 Command
↓
-------------------
之后每张图
-------------------
只更新 Input
↓
Submit
↓
Readback
lw.PPOCR.Vulkan 正是这么做的。
这也是：
lwvk_ocr_create()
和：
lwvk_ocr_run_bgr()
必须分开的原因。
十一、ONNX 怎样变成 Vulkan 计算？
这是整个推理 Runtime 最关键的地方。
项目没有在运行时把 ONNX 请求转给 ONNX Runtime。
而是自己解析模型。
代码：
src/onnx_import.cpp
项目实现了一个受限的 ONNX Protobuf Parser。
注意这里的“受限”。
它不是准备实现：
任意 ONNX 模型
任意算子
任意 Dynamic Shape
当前目标明确限定在：
PP-OCRv6 Tiny
PP-OCRv6 Small
PP-OCRv6 Medium
这种设计非常重要。
通用 Runtime 和专用 Runtime 是两个完全不同的工程方向。
专用 Runtime 可以针对：
固定模型结构
固定算子集合
固定 Shape 规律
做更激进的优化。
十二、模型解析以后，还不能马上执行
ONNX 原始图通常不是最适合 GPU 执行的形式。
所以项目还有：
src/graph_optimizer.cpp
用于做图优化。
例如：
Conv
↓
Bias
↓
HardSwish
如果条件满足，可以把它变成：
Conv + Bias + HardSwish
一个 Shader Epilogue 直接完成。
这样可以少一次：
Tensor 写显存
↓
Tensor 再读显存
↓
额外 Dispatch
尤其大特征图情况下：
少一次完整显存读写，往往比少几个浮点运算更值钱。
项目当前已经包含多种经过验证的融合，例如：
Conv + Bias
GELU
SiLU
HardSwish
部分 Affine / ReLU
当然并不是看到连续节点就随便融合。
还必须确认：
中间 Tensor 是否还有其他消费者
是否是 Graph Output
参数是否符合融合形式
计算顺序是否允许保持
数值结果是否通过回归测试
否则很容易出现：
速度快了
结果错了
十三、Shader 是怎样进入 Vulkan 的？
开发时我们写的是：
*.comp
例如：
conv_pointwise_wide.comp
这是 GLSL Compute Shader。
但 Vulkan GPU 最终执行的不是 GLSL 文本，而是：
SPIR-V
所以构建阶段：
GLSL
│
│ glslc
▼
SPIR-V
│
▼
嵌入程序
运行时读取对应 SPIR-V，然后：
VkShaderModuleCreateInfo sm{
VK_STRUCTURE_TYPE_SHADER_MODULE_CREATE_INFO
};
sm.codeSize = code.bytes;
sm.pCode = code.data;
vkCreateShaderModule(
device,
&sm,
nullptr,
&shaderModule
);
接下来再创建 Compute Pipeline。
十四、Compute Pipeline 是什么？
可以粗略理解为：
“这个 GPU 算子以后应该怎么执行”的可复用对象。
创建过程包括：
ShaderModule
Descriptor Layout
Push Constant Layout
Pipeline Layout
Compute Pipeline
项目中封装成：
Context::pipeline()
并且使用：
pipeline name
进行缓存。
比如：
conv_dw
conv_dw4
conv_gemm
conv_gemm_tiled
conv_pointwise
conv_pointwise_vector
conv_pointwise_wide
layernorm
attn
softmax
softmax_parallel
第一次遇到：
conv_pointwise_wide
创建 Pipeline。
后面再次使用：
直接复用
千万不要每一个神经网络节点都重新：
vkCreateShaderModule
vkCreateComputePipelines
这是非常昂贵的。
十五、DescriptorSet：Shader 怎么知道哪个 Buffer 是输入？
假设 Compute Shader：
layout(binding = 0, std430)
readonly buffer X {
vec4 x[];
};
layout(binding = 1, std430)
readonly buffer W {
vec4 w[];
};
layout(binding = 2, std430)
readonly buffer B {
vec4 b[];
};
layout(binding = 3, std430)
writeonly buffer O {
vec4 o[];
};
那么 GPU 需要知道：
binding 0 → 哪个 Buffer
binding 1 → 哪个 Buffer
binding 2 → 哪个 Buffer
binding 3 → 哪个 Buffer
这就是 Descriptor 的作用。
大致过程：
VkDescriptorSetLayout
↓
VkDescriptorPool
↓
VkDescriptorSet
↓
vkUpdateDescriptorSets
例如：
VkDescriptorBufferInfo info{};
info.buffer = inputBuffer;
info.offset = inputOffset;
info.range = inputBytes;
然后绑定到：
binding = 0
Shader 就知道 Input 在哪里。
因此可以把 Descriptor 理解为：
GPU Kernel 的“大参数表”
十六、Push Constants：小参数不要再建 Buffer
有些参数非常小。
比如 GEMM：
M
N
K
activation flags
如果每次都专门创建 Buffer 太麻烦。
Vulkan 提供：
Push Constants
例如项目中的 Shader：
layout(push_constant) uniform PC
{
uint M;
uint N;
uint K;
uint flags;
} p;
CPU 侧：
vkCmdPushConstants(
commandBuffer,
pipelineLayout,
VK_SHADER_STAGE_COMPUTE_BIT,
0,
sizeof(params),
¶ms
);
这非常适合：
Tensor Shape
Stride
Padding
Activation Flag
Width / Height
等小规模动态参数。
十七、真正让 GPU 计算的核心：vkCmdDispatch
到这里终于来到 Vulkan Compute 最关键的函数：
vkCmdDispatch(...)
例如项目中的一个 Pointwise Shader：
layout(local_size_x = 128) in;
表示每个 WorkGroup 包含 128 个 Invocation。
假设：
vkCmdDispatch(
command,
8,
12,
1
);
那么：
8 × 12
个 WorkGroup 会被调度。
每组：
128 Invocation
注意：
vkCmdDispatch() 的参数不是：
元素数量
而是：
WorkGroup 数量
十八、看一个项目里的真实 FP32 GEMM Shader
项目为了优化 Medium REC 的大通道 Projection，实现了：
src/shaders/conv_pointwise_wide.comp
它采用：
M32 / N128 / K32
Tile。
工作组：
layout(local_size_x = 128) in;
并使用共享内存：
shared float aa[1024];
shared vec4 bb[1024];
核心思想可以简化成：
Global Memory
│
├── Input Tile
│
└── Weight Tile
│
▼
Shared Memory
│
▼
128 GPU Invocation
│
▼
FP32 FMA
│
▼
Register
│
▼
Output Buffer
为什么使用 Shared Memory？
因为如果 128 个线程都反复从全局显存读取相同 Input / Weight：
显存带宽浪费
把 Tile 先搬到：
shared
之后同一个 WorkGroup 内重复利用，就能减少 Global Memory 访问。
这和 CUDA 的：
__shared__
是一个概念。
十九、为什么需要 barrier()？
Shader 中可以看到：
barrier();
原因很简单。
假设：
Thread 0 写 shared
Thread 1 写 shared
Thread 2 已经开始读 shared
如果没有同步：
Thread 2 可能读到还没有写完的数据
所以必须：
加载 Tile
↓
barrier()
↓
所有线程开始计算
↓
barrier()
↓
覆盖 Shared Memory，加载下一 Tile
这属于 WorkGroup 内部同步。
但神经网络算子之间还需要另一层 Vulkan Pipeline Barrier。
二十、神经网络为什么需要 vkCmdPipelineBarrier？
例如：
Conv A
│
│ write
▼
Tensor X
│
│ read
▼
Conv B
GPU 是异步执行系统。
不能简单认为：
vkCmdDispatch(A);
vkCmdDispatch(B);
就一定代表：
A 所有写操作对 B 都已经可见
所以 Runtime 需要建立正确的：
Execution Dependency
+
Memory Dependency
项目里会使用：
vkCmdPipelineBarrier(...)
处理类似：
Transfer → Compute
Compute → Compute
Compute → Transfer
Host → Compute
这些依赖。
这也是 Vulkan 相比很多高级推理框架“麻烦”的地方：
同步由你自己负责。
但反过来：
你也因此获得了完整控制权。
二十一、CommandBuffer：不是调用一个算子就等一次 GPU
理解 Vulkan 性能，必须理解 Command Buffer。
CPU 可以先：
vkBeginCommandBuffer(...)
然后连续记录：
Bind Conv Pipeline
Bind Descriptor
Push Params
Dispatch Conv
Barrier
Bind Depthwise Pipeline
Bind Descriptor
Push Params
Dispatch Depthwise
Barrier
Bind Activation Pipeline
Dispatch
Barrier
Bind MatMul Pipeline
Dispatch
...
全部录完：
vkEndCommandBuffer(...)
此时：
GPU 还不一定已经执行。
我们只是准备好了一张“计算任务清单”。
这对神经网络特别重要。
如果一个模型有几百个算子，不是必须：
CPU 提交一次
等 GPU
CPU 再提交一次
等 GPU
...
完全可以把一大串 Dispatch 录制好，然后一次提交。
二十二、真正开始执行：vkQueueSubmit
真正让 GPU 干活的是：
vkQueueSubmit(...)
项目中的 Plan::submit_readback() 最核心的逻辑可以抽象成：
vkResetFences(
device,
1,
&fence
);
VkSubmitInfo submit{
VK_STRUCTURE_TYPE_SUBMIT_INFO
};
submit.commandBufferCount = 1;
submit.pCommandBuffers = &command;
VkResult result = vkQueueSubmit(
queue,
1,
&submit,
fence
);
此时：
CPU
│
│ vkQueueSubmit
▼
GPU Queue
│
├─ Preprocess
├─ Conv
├─ Depthwise
├─ MatMul
├─ LayerNorm
├─ Attention
├─ Softmax
└─ Output Copy
GPU 开始真正执行。
二十三、CPU 怎么知道 GPU 执行完了？
项目使用 Fence。
提交：
vkQueueSubmit(
queue,
1,
&submit,
fence
);
等待：
vkWaitForFences(
device,
1,
&fence,
VK_TRUE,
timeout
);
于是：
CPU
│
│ Submit
▼
GPU 正在计算
│
│
│
▼
Fence Signaled
│
▼
CPU 继续
项目还特别处理了一个重要问题：
Fence Timeout 不是 GPU Cancel
也就是说：
vkWaitForFences 超时
并不代表之前提交的 GPU 工作立即停止。
因此如果出现超时或 Device Failure，项目会把对应 Plan 标记为：
poisoned
避免继续复用可能仍被 GPU 引用的资源。
这是比较重要的工程细节。
二十四、GPU 结果怎么回到 CPU？
最终输出如果 CPU 需要使用，就要回读。
典型方式：
Device Local Output
│
│ vkCmdCopyBuffer
▼
Host Visible Readback Buffer
│
▼
Fence
│
▼
CPU read()
比如 DET。
GPU 最终生成：
文字概率图
但是后续 DB 后处理仍然在 CPU。
所以需要把 Probability Map 读回来。
REC 又不完全一样。
项目为了减少巨大的分类概率回读，对 REC 可以直接在 GPU 进行：
Softmax / Argmax
然后只回读紧凑标签结果。
最后 CPU 再做：
CTC 去重
blank 删除
字典映射
UTF-8 拼接
这样可以明显减少：
GPU → CPU 数据量
二十五、一次单网络 Vulkan 推理究竟发生了什么？
现在可以把整个流程简化成：
CPU Input
│
▼
Host Visible Buffer
│
▼
Input Copy
│
▼
Device Local Arena
│
▼
┌────────────────────────┐
│ Compute Pipeline │
│ │
│ Conv │
│ Depthwise │
│ Pointwise │
│ LayerNorm │
│ Attention │
│ MatMul │
│ Softmax │
└──────────┬─────────────┘
│
▼
GPU Output
│
▼
Readback Buffer
│
▼
Fence
│
▼
CPU Result
项目中真正执行一次 Plan 的代码逻辑就是：
upload_->write(input, inputBytes);
submit_readback(
command_,
output,
ctc
);
而 submit_readback() 中：
vkResetFences
↓
vkQueueSubmit
↓
vkWaitForFences
↓
Buffer::read
这就是最核心的一次 Vulkan 神经网络执行。
二十六、但是完整 OCR 不是只有一个神经网络
PP-OCR 完整流程是：
DET
↓
CLS
↓
REC
实际上还夹杂大量图像处理和后处理。
lw.PPOCR.Vulkan 的默认完整流程更加接近：
BGR 原图
│
▼
GPU 上传原图
│
▼
GPU DET Resize / Normalize
│
▼
GPU DET Network
│
▼
Probability Map
│
▼
CPU DB PostProcess
│
├─ Threshold
├─ Contour
├─ Box Score
├─ Unclip
└─ Reading Order
│
▼
Text Boxes
│
▼
GPU Perspective Crop
│
▼
GPU CLS Preprocess
│
▼
GPU CLS
│
▼
CPU 判断 0° / 180°
│
▼
GPU REC Resize / Padding / Normalize
│
▼
GPU REC Network
│
▼
GPU Greedy Label
│
▼
CPU CTC Decode
│
▼
Dictionary UTF-8
│
▼
最终 OCR JSON
这才是：
一次完整 Vulkan OCR
二十七、第一阶段：原始 BGR 图片进入 GPU
项目的 C ABI 接受：
BGR8
width
height
stride
buffer bytes
而不是要求调用方：
JPEG → Base64
这样可以直接对接：
OpenCV Mat
工业相机 Buffer
视频帧
共享内存
完整 OCR API：
lwvk_ocr_run_bgr(...)
图片进入后，如果 GPU 能力满足，项目默认可以启用：
GPU DET Preprocess
GPU TEXT Preprocess
GPU Crop Preprocess
原图只上传一次。
这是非常重要的优化。
如果流程是：
CPU 原图
↓
GPU DET
↓
CPU Crop
↓
GPU CLS
↓
CPU
↓
GPU REC
中间会不断：
Host ↔ Device
搬数据。
项目当前共享原图路径则尽量做到：
原图只上传一次
DET / Crop / CLS / REC
共享同一个 Vulkan Context
CPU 主要传：
Box 坐标
方向信息
控制参数
二十八、第二阶段：DET
DET 的任务是：
图片
↓
文字区域概率图
模型输入通常先经过：
Resize
Normalize
Layout
GPU 前处理会直接把 BGR 图转换到网络输入 Tensor。
之后运行 DET 网络。
GraphEngine 会根据每个节点选择 Shader。
例如：
普通卷积
Depthwise Conv
1×1 Pointwise
GEMM
Pooling
Resize
Concat
LayerNorm
Attention
Softmax
都会变成不同：
dispatch(...)
最终内部还是：
vkCmdBindPipeline
vkCmdBindDescriptorSets
vkCmdPushConstants
vkCmdDispatch
二十九、第三阶段：DB 后处理为什么还留在 CPU？
DET GPU 输出只是概率图。
OCR 还需要：
阈值
连通区域
轮廓
Box Score
Unclip
四点框
阅读顺序
这些属于 DB 后处理。
目前项目把这部分放在 CPU。
原因并不难理解。
网络中：
Conv
GEMM
MatMul
高度规则、计算密集，非常适合 GPU。
而：
Contour
动态候选框
几何排序
条件分支
这种操作未必马上搬到 GPU 就会更快。
所以当前架构遵循一个很实用的原则：
适合 GPU 的交给 GPU，不适合的暂时保留 CPU。
“GPU OCR”并不意味着：
所有代码都必须跑 GPU。
三十、第四阶段：GPU 透视裁剪
CPU 得到四点框之后，要把文字区域裁出来。
传统方法：
GPU DET
↓
概率图回 CPU
↓
CPU Perspective Crop
↓
每个 Crop 再传 GPU
项目目前在支持条件下，可以：
原图保持 GPU
↓
CPU 只传四点框
↓
GPU Perspective Crop
得到：
BGR8 Crop
然后继续用于：
CLS
REC
这里项目还刻意保留：
Perspective Crop
和：
Network Resize / Normalize
两个阶段。
没有为了少一次操作强行融合成一次采样。
原因是要保留原有：
插值
BGR8 舍入
语义，避免为了性能改变 OCR 结果。
三十一、第五阶段：CLS
CLS 主要负责判断文字方向，例如：
0°
180°
CLS 输入经过：
Resize
Normalize
之后执行小型分类网络。
如果判定为 180°：
REC 前重新按照旋转方向采样
项目中 CLS / REC 都可以使用 GPU Text Preprocess。
三十二、第六阶段：REC
真正的识别通常是 OCR 中非常重要的一部分。
REC 输入通常是：
高度固定
宽度动态
例如：
[1, 3, 48, W]
其中：
W = 32 ... 960
并按照 8 对齐。
宽度不同意味着：
Tensor Shape 不同
Workspace 不同
Dispatch 参数不同
Command Binding 可能不同
这就是为什么 REC 比固定 Shape 模型复杂。
三十三、为什么项目需要 REC Plan Cache？
假设连续识别：
宽 64
宽 128
宽 320
宽 96
宽 512
...
如果每次宽度变化都：
重新计算 Tensor Layout
重新分配 Buffer
重新更新 Descriptor
重新录制 CommandBuffer
CPU 和 Vulkan Driver 开销会非常明显。
所以项目会缓存执行 Plan。
v1.0.1 对这一块进行了专门优化：
8 个 REC Slot
×
每个最多 16 个 Width Plan
=
最多 128 个缓存 Plan
但要注意：
这不是：
128 份完整神经网络显存
核心 Arena 仍然共享。
主要缓存的是：
执行计划
Descriptor
Command
相关 Metadata
这样大小图交替时，就不会频繁互相淘汰。
这是典型的：
用有限内存换稳定延迟。
三十四、REC 为什么成为 Vulkan 优化重点？
Tiny、Small、Medium 的 REC 计算规模并不一样。
尤其 Medium 中存在较大的：
1×1 Projection
GEMM
大 Channel MatMul
所以项目最近大量优化都集中在：
Pointwise
Wide Tile
Vocabulary Projection
REC Batch
例如 NVIDIA 路径针对某些大 Channel Projection 会选择：
conv_pointwise_wide
而不是普通：
conv_pointwise
为什么不所有 GPU 都使用 Wide Kernel？
因为项目实测发现：
RTX 4060
上可能加速，但同样策略在某些 AMD 集显上会退化。
所以代码里会根据：
vendorID
shared memory
shape
task
做 Kernel Selection。
这也是自己写 Runtime 最大的价值之一：
你可以针对真实模型和真实 GPU 做调度。
三十五、GPU Greedy CTC
REC 网络最后通常会得到类似：
[T, Classes]
的输出。
如果 Classes 很大：
Small / Medium = 18710
直接把：
T × 18710 FP32
全部传回 CPU，会产生明显 PCIe / Memory Copy 开销。
所以项目对 REC 尾部进行了：
Softmax + Argmax
GPU 化。
最终 GPU 只需要给 CPU：
每个时间步的 label
+
confidence
然后 CPU 做：
去 blank
去连续重复
字典映射
UTF-8 拼接
这正是专用 Runtime 可以做的优化。
并不是：
GPU 算完模型
所有 Logit 全部搬 CPU
再慢慢处理
三十六、外部程序到底怎么调用？
说完内部原理，再看应用层。
项目冻结了 C ABI v1。
完整 OCR 最主要的接口是：
lwvk_ocr_config_default()
lwvk_ocr_create()
lwvk_ocr_run_bgr()
lwvk_ocr_result_json()
lwvk_ocr_result_destroy()
lwvk_ocr_destroy()
调用过程非常清晰：
枚举 GPU
↓
初始化配置
↓
加载模型
↓
创建 OCR Engine
↓
传入 BGR 图像
↓
Vulkan OCR
↓
获得 Result Handle
↓
查询 JSON 长度
↓
复制 JSON
↓
释放 Result
↓
复用 OCR Engine 处理下一张图
三十七、一个完整 C API 调用示例
下面给一个接近真实业务的调用结构：
#include "lw_ppocr_vulkan.h"
lwvk_ocr_config config{};
if (lwvk_ocr_config_default(&config) != LWVK_OK)
{
return -1;
}
config.device_index = 0;
config.enable_classifier = 1;
lwvk_ocr_handle engine = nullptr;
if (lwvk_ocr_create(
"models/onnx/ppocrv6-tiny",
&config,
&engine) != LWVK_OK)
{
printf(
"create failed: %s\n",
lwvk_last_error()
);
return -1;
}
假设我们已经有：
uint8_t* bgr
width
height
stride
bufferBytes
执行：
lwvk_ocr_result_handle result = nullptr;
lwvk_status status = lwvk_ocr_run_bgr(
engine,
bgr,
bufferBytes,
width,
height,
stride,
&result
);
得到结果以后先查询长度：
uint64_t required = 0;
lwvk_ocr_result_json(
result,
nullptr,
0,
&required
);
申请：
std::vector<char> json(required);
复制：
lwvk_ocr_result_json(
result,
json.data(),
json.size(),
&required
);
现在：
printf("%s\n", json.data());
即可得到完整 OCR JSON。
最后：
lwvk_ocr_result_destroy(result);
lwvk_ocr_destroy(engine);
实际批量业务不要每张图片：
Create → Run → Destroy
应该：
Create 一次
↓
Run
Run
Run
Run
Run
↓
Destroy
这样才能复用：
权重
Pipeline
CommandBuffer
Workspace
Plan Cache
三十八、为什么 API 要传 buffer_bytes 和 stride？
项目的 BGR API 不是只收：
width
height
pointer
还要求：
buffer_bytes
stride
原因是实际图像内存不一定：
width × 3
连续紧密排列。
例如相机和 Bitmap 经常存在：
Padding
Stride
安全的 Buffer 长度至少应该满足：
(height - 1) × stride
+
width × 3
通过显式传递：
stride
buffer_bytes
Runtime 可以在 C ABI 边界进行越界检查。
这个设计对于：
工业相机
C#
C++
OpenCV
共享内存
接入都更加稳妥。
三十九、C# 为什么也能调用？
Vulkan Runtime 本身是：
C++
但项目提供稳定 C ABI。
因此 C# 只需要：
P/Invoke
即可调用。
大概结构：
[DllImport(
"lw.PPOCR.Vulkan.dll",
CallingConvention = CallingConvention.Cdecl)]
static extern int lwvk_ocr_run_bgr(
IntPtr engine,
IntPtr pixels,
ulong bytes,
uint width,
uint height,
uint stride,
out IntPtr result);
C# 不需要知道：
VkInstance
VkDevice
VkDescriptorSet
VkCommandBuffer
这些都封装在原生 DLL 内。
因此业务层看到的是：
Bitmap / BGR
↓
C ABI
↓
Vulkan Runtime
↓
OCR JSON
这也是项目提供 WinForms Demo 的原因。
四十、Python 也可以直接调用
项目同样提供 Python 封装。
例如完整图片 OCR：
python examples/python/ocr_image.py ^
--library build/local/Release/lw.PPOCR.Vulkan.dll ^
--models models/onnx/ppocrv6-tiny ^
--image test-images/sample.jpg ^
--device 1 ^
--draw build/ocr-boxes.png
Python 只是负责：
图片读取
C ABI 调用
结果显示
真正的神经网络还是：
C++ Vulkan
四十一、一次完整调用，从 API 一路追到 GPU
现在把代码调用链完整串起来。
应用调用：
lwvk_ocr_run_bgr()
进入：
src/api.cpp
内部调用：
OcrEngine::run()
代码：
src/ocr_pipeline.cpp
这里先建立：
DET Graph
CLS Graph
REC Graph
如果条件满足，三者还共享：
Context / Device
然后进入：
run_ocr_host()
负责整体 OCR 流程。
某个网络真正执行时：
GraphEngine
↓
Plan
↓
Plan::run()
或
Plan::run_bgr()
↓
Plan::submit_readback()
最终来到：
vkResetFences
↓
vkQueueSubmit
↓
vkWaitForFences
↓
Buffer::read
因此整个完整调用链就是：
lwvk_ocr_run_bgr
│
▼
OcrEngine
│
▼
OCR Host
│
├──── DET GraphEngine
│ │
│ ▼
│ Plan
│ │
│ ▼
│ vkQueueSubmit
│
├──── CPU DB
│
├──── GPU Crop
│
├──── CLS GraphEngine
│ │
│ ▼
│ vkQueueSubmit
│
└──── REC GraphEngine
│
▼
Plan
│
▼
vkQueueSubmit
│
▼
CTC
│
▼
OCR JSON
这就是从业务 API 一直追到底层 Vulkan 的完整路径。
四十二、为什么 Vulkan 推理性能不能只看 Shader？
做这种 Runtime，很容易一开始只关注：
Conv 快不快
GEMM 快不快
实际上端到端耗时还包括：
图片前处理
数据上传
Plan 创建
Descriptor 更新
Command Recording
Queue Submit
GPU Execute
Fence Wait
GPU → CPU Readback
DB
Crop
CTC
JSON
所以：
某 Kernel 快 30%
并不代表：
完整 OCR 快 30%
lw.PPOCR.Vulkan 最近的很多实验也证明了这一点。
一些短 REC Kernel：
单算子 Benchmark 有提升
但放回：
完整 OCR
变尺寸流
收益不稳定，所以没有默认开启。
这个原则非常重要：
微基准快，不等于业务快。
四十三、当前项目做了哪些性能优化？
现在回头看，就很容易理解项目里的优化了。
主要包括：
1. 权重常驻 GPU
2. Pipeline Cache
3. CommandBuffer 复用
4. Tensor 生命周期显存复用
5. 有界 Plan Cache
6. FP32 Tile GEMM
7. vec4 访问
8. Shared Memory Tile
9. Depthwise Vectorization
10. GELU / SiLU / HardSwish Fusion
11. GPU DET Preprocess
12. GPU TEXT Preprocess
13. GPU Perspective Crop
14. GPU Greedy CTC
15. REC Width Cache
16. Medium REC Wide Projection
而且项目没有把：
FP16
Cooperative Matrix
Multiple REC Lane
直接作为正式默认路径。
它们目前属于：
实验路径
这是因为性能优化除了“更快”，还有两个门槛：
结果正确
+
不同设备稳定
四十四、为什么不是直接把 FP32 改 FP16？
很多人会说：
GPU 推理想快，全部改 FP16 不就行了吗？
事情没那么简单。
例如 OCR REC 最终输出有：
18710 类
一个很小的 Logit 变化，都可能改变：
blank
repeat
character
选择。
所以：
FP32 → FP16
不能只看：
最大绝对误差
还要看：
最终文字
检测框
置信度
CER
业务准确率
因此当前默认路径保持：
FP32
FP16 / Cooperative Matrix 则单独研究和验证。
这是比“跑起来”更重要的工程问题。
四十五、Vulkan 相比 TensorRT 到底怎么样？
这个问题也需要比较客观地看。
当前项目自己的测试明确指出：
不能宣称 Vulkan 已经全面超过 TensorRT。
在 RTX 4060 Laptop GPU 的部分测试中：
Tiny / Small
差距已经比较有限
但：
Medium
尤其 REC
仍明显落后于已有 FP16-enabled TensorRT 路径。
这很正常。
TensorRT 背后有 NVIDIA 多年的：
Kernel Library
Graph Optimizer
Tensor Core
Tactic Search
Kernel Auto-Tuning
自己写 Vulkan Runtime，短时间内全面超过 TensorRT 并不现实。
但 Vulkan Runtime 的价值并不只在：
RTX 4060 上能不能比 TRT 快 2 ms
而在：
Runtime 完全可控
Shader 完全可改
无 CUDA Runtime
跨 GPU 厂商基础能力
可以围绕固定模型深度优化
部署依赖可控
整个执行链透明
这对于：
边缘部署
工业软件
OCR SDK
自己的推理框架研究
都很有价值。
四十六、如果自己从零写一个 Vulkan 推理 Runtime，最小流程是什么？
把这篇文章压缩成代码，其实就是：
// 1. Vulkan Runtime
create_instance();
select_physical_device();
create_device();
get_compute_queue();
// 2. Model
parse_onnx();
optimize_graph();
// 3. Memory
create_weight_buffer();
upload_weights();
plan_tensor_lifetime();
create_workspace();
// 4. Kernels
for (auto& kernel : required_kernels)
{
create_shader_module(kernel);
create_compute_pipeline(kernel);
}
// 5. Execution Plan
for (auto& node : graph)
{
bind_pipeline(node);
bind_descriptors(node);
push_constants(node);
dispatch(node);
insert_barrier(node);
}
// 6. Inference
upload_input();
vkQueueSubmit(...);
vkWaitForFences(...);
read_output();
// 7. Postprocess
decode_result();
看上去只有几十行。
但真正复杂的是：
ONNX Parser
Graph Optimizer
Shape Inference
Tensor Lifetime
Workspace Planner
Kernel Library
Kernel Selection
Synchronization
Plan Cache
错误恢复
资源上限
跨 GPU 测试
精度回归
所以一个真正可用的 Vulkan Runtime，工作量远远大于：
写几个 vkCmdDispatch
四十七、我认为这个项目最值得看的，不只是 OCR
lw.PPOCR.Vulkan 当前确实是一个 PP-OCR Runtime。
但它背后的几个模块实际上已经具有更通用的价值：
ONNX Parser
│
▼
Graph Optimizer
│
▼
Workspace Planner
│
▼
Vulkan Buffer
│
▼
Compute Pipeline
│
▼
Descriptor Binding
│
▼
Command Recording
│
▼
Queue Submit
│
▼
GPU Profiling
也就是说：
如果以后想继续尝试：
目标检测
图像分割
超分辨率
分类
其他固定 ONNX 模型
真正能够复用的，正是这一层基础设施。
当然当前项目仍然明确：
它不是通用 ONNX Runtime。
支持一个新的模型，仍需要检查：
算子
Shape
Layout
Preprocess
Postprocess
Kernel
是否已经覆盖。
四十八、最后再看一次完整 Vulkan OCR
现在理解所有细节以后，再看这张流程就很清楚了：
BGR Image
│
▼
Host Upload Buffer
│
▼
vkCmdCopyBuffer
│
▼
GPU DET Preprocess
│
▼
DET Network
│
多次 vkCmdDispatch
│
▼
Probability Map
│
▼
CPU DB
│
▼
Text Boxes
│
▼
GPU Perspective Crop
│
▼
CLS
│
▼
Rotation Decision
│
▼
GPU REC Preprocess
│
▼
REC
│
多次 vkCmdDispatch
│
▼
GPU Softmax / Argmax
│
▼
Compact Labels
│
▼
CPU CTC
│
▼
Dictionary
│
▼
UTF-8 Text
│
▼
OCR JSON
而网络内部每一个算子，本质上又落到：
vkCmdBindPipeline
↓
vkCmdBindDescriptorSets
↓
vkCmdPushConstants
↓
vkCmdDispatch
↓
vkCmdPipelineBarrier
整个图录制完成后：
vkQueueSubmit
↓
GPU Execute
↓
vkWaitForFences
↓
Readback
这就是一次真正的：
Vulkan GPU 推理。
结语
以前使用 ONNX Runtime 或 TensorRT 时，我们看到的是：
session.Run();
一行代码。
而在这一行代码下面，实际上隐藏着：
模型解析
图优化
Tensor 生命周期
显存规划
Kernel 选择
Shader
Pipeline
Descriptor
Command Buffer
GPU Dispatch
同步
数据回读
lw.PPOCR.Vulkan 做的事情，就是把这些原来隐藏在推理框架内部的部分，一层一层重新实现出来。
所以这个项目对我来说，价值不只是：
“又做了一个 OCR Demo。”
更有意思的是：
我们开始真正控制一个神经网络从 ONNX 图，一直到 GPU 指令调度之间的整条路径。
当性能有问题的时候，可以继续往下追：
是 CPU Preprocess？
是 Host → Device？
是 Conv？
是 Depthwise？
是 GEMM？
是 Attention？
是 Barrier 太多？
是 Dispatch 太碎？
是 REC Width Cache？
还是 GPU 空闲降频？
这也是自己写 Runtime 最有意思的地方。
所有东西：
都能继续往下挖。
项目地址
lw.PPOCR.Vulkan
项目当前提供：
- PP-OCRv6 Tiny / Small / Medium；
- Vulkan FP32 GPU 推理；
- 原生 ONNX 解析；
- 完整 DET → CLS → REC；
- GPU 图像前处理与透视裁剪；
- C ABI；
- C# WinForms；
- Python 示例；
- HTTP / Web；
- Windows / Linux 构建与部署。
如果你正在研究：
Vulkan Compute
GPU 推理
ONNX Runtime
OCR 部署
C++ 推理引擎
自定义 GPU Kernel
欢迎一起交流。
天天代码码天天
GitHub：
github.com/lxw112190
