Chapter 109: OpenCL Fundamentals

OpenCL 解决的是跨设备通用计算的工程问题：同一段宿主程序需要发现可用计算设备，把 kernel 编译成设备可执行对象，分配 buffer 或 image，把命令排入队列，再把输出交回图形管线、离线工具或 CPU 侧数据流程。读完本章，读者应能追踪一次 OpenCL 任务从 platform 选择到 kernel enqueue 的完整路径，并能判断一个图像处理或粒子更新任务卡在设备选择、内存传输、队列依赖、kernel 编译还是 work-group 参数上。

本章的贯穿材料是一帧后处理图像：渲染器已经产生一张 1920 × 1080 的线性颜色图，工具链希望用 OpenCL 执行一次亮度阈值筛选，把高亮像素写入输出 buffer，后续可能继续做 bloom、粒子生成或离线分析。这个任务足够小，可以用一个 kernel 解释 OpenCL 的基础对象；它也足够接近图形工作流，可以暴露图像资源、同步、拷贝和调试证据之间的关系。

OpenCL 的核心判断顺序可以压缩成一句话：先确认目标平台和设备能承担这类并行任务，再确认 context、command queue、program、kernel、memory object 和 event 是否形成同一条有效数据路径。Khronos 在 OpenCL 主页 中把 OpenCL 描述为面向异构系统的开放并行编程标准；OpenCL 统一 API 规范 中的基础模型也围绕 host、device、command queue、memory object 和 kernel command 展开。本章只使用这些稳定对象建立工程理解，版本相关能力通过运行时 query 处理。

109.1 OpenCL 架构与跨平台特性

OpenCL 的架构起点是 host 与 device 的分工。Host 是运行普通 C、C++、Python binding 或其它语言绑定的 CPU 侧程序，它负责发现平台、选择设备、创建 context、编译 program、配置 kernel 参数、提交命令和读取状态。Device 是执行 kernel 的计算设备，可以是 GPU、CPU、DSP、FPGA 或其它加速器。这个分工让 OpenCL 能覆盖多种硬件，也把大量硬件差异暴露给应用层处理。

跨平台特性来自一组稳定抽象，而非来自完全相同的硬件行为。Platform 表示某个 OpenCL 实现入口，通常由驱动或 runtime 提供；device 表示该实现下可用的计算设备；context 表示一组 device、memory object、program object 和 queue 的共享边界；command queue 表示 host 向某个 device 提交命令的顺序入口。只要程序通过 query 读取能力、限制和扩展，就能用同一套 API 组织不同设备上的计算路径。

在本章的后处理图像任务中，跨平台目标不是让每个设备给出完全相同的性能曲线，而是让程序能用同一套步骤表达“读取输入像素、执行阈值 kernel、写出结果”的计算意图。CPU device 可能适合调试和离线校验，集成 GPU 可能适合小批量图像处理，独立 GPU 可能适合高分辨率图像和粒子更新。OpenCL 提供的是可移植执行层，性能判断仍要回到设备能力、内存路径、queue 同步和 kernel 访问模式。

OpenCL 的可移植性有三个边界。第一，OpenCL 3.x 采用能力查询和扩展模型，旧设备可能只提供 OpenCL 1.2 基线能力，某些 2.x 或 3.x 功能需要逐项查询。第二，kernel 语言、SPIR-V ingestion、image 支持、共享虚拟内存、外部内存互操作等能力由设备和驱动声明。第三，图形互操作通常依赖具体图形 API、操作系统和扩展组合，工程上要准备 shared resource path 与 staging copy path 两条路线。

这条边界会直接影响后处理图像的设计。如果 OpenCL 和图形 API 之间有可用互操作，渲染输出可以通过共享 buffer、共享 image 或外部内存进入 OpenCL；如果没有这类能力，程序就要把图像数据拷贝到 OpenCL memory object，再把输出交回渲染端或离线端。前者主要检查资源所有权和同步对象，后者主要检查带宽、拷贝次数和队列等待。

下面的图把 OpenCL 在图形工作流中的位置压缩成一条资源路径。它只描述本章使用的最小路径，不覆盖所有高级互操作或多队列调度细节。

图中的关键关系是 context 把 queue、memory object、program 和 kernel 放进同一生命周期边界。Command queue 负责把写 buffer、执行 kernel、读 buffer 或同步 event 排成设备命令。Kernel 本身只看到参数和地址空间，它并不知道 host 如何选择 platform，也不知道输出最终进入 bloom pass、粒子系统还是离线工具。

因此，OpenCL 的基础理解要从“对象归属”开始。一个 buffer 归属某个 context，一个 queue 绑定某个 context 和 device，一个 kernel 来自同一 context 中构建出的 program。对象归属不一致时，错误往往出现在 enqueue 或 set arg 阶段；对象归属一致时，下一步才进入 work size、内存带宽和同步依赖的性能判断。

109.2 设备/平台/队列模型解析

Platform、device 和 queue 是 OpenCL 运行模型的三层入口。Platform 回答“当前机器有哪些 OpenCL 实现”；device 回答“这个实现暴露哪些可执行计算设备”；queue 回答“host 把命令提交给哪个 device，以及命令按什么顺序完成”。这三层的选择决定后续 context、memory object 和 kernel command 的合法边界。

Platform 选择的常见误区是只取第一个 platform。工程程序应读取 platform name、vendor、version、extension，再列出每个 platform 下的 device。多平台机器并不少见，例如一个系统可能同时有 Intel CPU runtime、NVIDIA GPU runtime、AMD GPU runtime 或 Portable Computing Language 实现。只取第一个 platform 会把设备选择交给驱动枚举顺序，后续性能和功能差异就会变成偶发现象。

Device 选择要把任务类型和能力限制一起看。后处理阈值 kernel 主要是线性读写，适合优先检查 CL_DEVICE_TYPE_GPU、global memory size、maximum work-group size、image support、preferred vector width、OpenCL C version 和扩展列表。体数据处理可能更关注 3D image、地址空间限制和缓存行为；粒子模拟可能更关注全局内存带宽、原子操作和 subgroup 相关能力；离线工具可能更关注双精度、长时间稳定性和错误日志。

Queue 是命令进入 device 的提交入口。In-order queue 会让同一 queue 中的命令按提交顺序建立默认依赖，适合初学路径和确定性调试。Out-of-order queue 允许命令在满足 event 依赖后重排执行，适合独立图像块、分阶段 pipeline 或多 stream 工作，但它要求应用显式维护 event wait list。Profiling queue 会记录 queued、submit、start、end 等时间戳，适合把问题定位到 host 提交、设备排队或 kernel 执行。

对于一帧后处理图像，最小可解释队列路径是 write input → enqueue threshold kernel → read output。如果使用 in-order queue，三个命令在同一 queue 中天然形成顺序。这个路径的可观察结果很直接：输入图像中亮度超过阈值的像素保留或标记，输出 buffer 中其余像素变暗或清零。若画面结果正确但耗时异常，queue profiling 能把时间拆成传输阶段和 kernel 阶段。

下面这段 host 侧骨架展示对象关系。它是简化代码，省略了完整错误处理和释放路径，只用于说明 platform、device、context、queue、program、kernel 和 buffer 的创建顺序。

// Simplified OpenCL host path for one image-processing kernel.
cl_platform_id platform = pick_platform_by_vendor_or_feature();
cl_device_id device = pick_device(platform, CL_DEVICE_TYPE_GPU);

cl_int err = CL_SUCCESS;
cl_context context = clCreateContext(NULL, 1, &device, NULL, NULL, &err);

const cl_queue_properties queue_props[] = {
    CL_QUEUE_PROPERTIES, CL_QUEUE_PROFILING_ENABLE,
    0
};
cl_command_queue queue = clCreateCommandQueueWithProperties(
    context, device, queue_props, &err);

cl_program program = clCreateProgramWithSource(
    context, 1, &kernel_source, NULL, &err);
err = clBuildProgram(program, 1, &device, build_options, NULL, NULL);

cl_kernel kernel = clCreateKernel(program, "luminance_threshold", &err);
cl_mem input_buffer = clCreateBuffer(context, CL_MEM_READ_ONLY, input_bytes, NULL, &err);
cl_mem output_buffer = clCreateBuffer(context, CL_MEM_WRITE_ONLY, output_bytes, NULL, &err);

这段代码的关键点不是函数数量，而是对象生命周期。context 是后续对象的共同父边界；queue 绑定到同一个 context 和一个具体 device；program 要针对同一个 device build；kernel 来自已构建的 program；input_buffer 和 output_buffer 属于同一个 context。这些对象形成闭合路径后，host 才能把图像数据、kernel 参数和 ND-range dispatch 串成一次可执行命令序列。

队列模型还决定同步语义。clEnqueueWriteBuffer、clEnqueueNDRangeKernel、clEnqueueReadBuffer 都可以返回 event。Event 既是依赖句柄，也是 profiling 证据。若多个 queue 共享同一个 output buffer，应用要让消费命令等待生产 event；若 host 要读取结果，应用要等待 read event 或调用 clFinish。在图形工作流中，event 语义要进一步对齐图形 API 的 fence、semaphore 或 command buffer 完成状态。

109.3 OpenCL Platform Device Queue Kernel Buffer Path

一条可维护的 OpenCL 路径应从 platform 和 device 选择开始，以 enqueue 后的 event 或输出资源结束。把这条路径拆清楚，能让图像处理、粒子模拟、体数据转换和离线烘焙共用同一套排查顺序。本节继续使用亮度阈值后处理，把数据路径固定为 host input buffer → OpenCL input buffer → kernel → OpenCL output buffer → host or render consume。

Kernel 负责把输入像素映射到输出像素。每个 work-item 处理一个像素时，global ID 就对应线性像素索引；global size 应覆盖像素总数；local size 决定 work-group 的分块方式。这个 kernel 使用全局内存 buffer，能在支持 OpenCL C 的设备上表达最小计算路径。

__kernel void luminance_threshold(
    __global const uchar4* input_pixels,
    __global uchar4* output_pixels,
    const uint pixel_count,
    const float threshold)
{
    const uint gid = (uint)get_global_id(0);
    if (gid >= pixel_count) {
        return;
    }

    const uchar4 src = input_pixels[gid];
    const float3 rgb = (float3)(src.x, src.y, src.z) / 255.0f;
    const float luminance = dot(rgb, (float3)(0.2126f, 0.7152f, 0.0722f));
    const uchar keep = (uchar)(luminance >= threshold ? 255 : 0);

    output_pixels[gid] = (uchar4)(src.x & keep, src.y & keep, src.z & keep, src.w);
}

这段 kernel 的输入是 input_pixels、pixel_count 和 threshold，输出是 output_pixels。执行粒度是 work-item，每个 work-item 读取一个 uchar4 并写回一个 uchar4。性能上，它主要受全局内存读写带宽影响，ALU 成本很小。视觉结果上，阈值越高，输出中保留的高亮像素越少；阈值越低，更多像素保留。这种“输入像素 → work-item → 输出像素”的一对一映射，是 OpenCL 入门时最适合验证对象路径的形态。

Host enqueue 路径要把 kernel 参数、内存传输和 dispatch 组合起来。下面的骨架保留关键 API 调用，并把 event 作为后续 profiling 与同步证据保存下来。

size_t pixel_count = width * height;
size_t bytes = pixel_count * sizeof(cl_uchar4);

cl_event write_event = NULL;
clEnqueueWriteBuffer(
    queue, input_buffer, CL_FALSE,
    0, bytes, host_pixels,
    0, NULL, &write_event);

clSetKernelArg(kernel, 0, sizeof(cl_mem), &input_buffer);
clSetKernelArg(kernel, 1, sizeof(cl_mem), &output_buffer);
clSetKernelArg(kernel, 2, sizeof(cl_uint), &pixel_count);
clSetKernelArg(kernel, 3, sizeof(cl_float), &threshold);

const size_t local_size = 256;
const size_t global_size = ((pixel_count + local_size - 1) / local_size) * local_size;

cl_event kernel_event = NULL;
clEnqueueNDRangeKernel(
    queue, kernel, 1,
    NULL, &global_size, &local_size,
    1, &write_event, &kernel_event);

cl_event read_event = NULL;
clEnqueueReadBuffer(
    queue, output_buffer, CL_TRUE,
    0, bytes, host_output,
    1, &kernel_event, &read_event);

这条路径体现了三类依赖。第一，数据依赖：kernel 读取 input_buffer 前，写入命令要完成。第二，执行依赖：read output 前，kernel 要完成。第三，范围依赖：global size 向上对齐到 local size 后，kernel 内部必须用 gid >= pixel_count 保护尾部 work-item。尾部保护把 dispatch 对齐和真实像素数量分开，便于保留稳定 local size。

Build 阶段是 OpenCL 初学路径中最容易被低估的阶段。clBuildProgram 失败时，host 侧只能得到错误码；真正的语法错误、地址空间错误、类型错误和目标设备编译限制，要通过 clGetProgramBuildInfo 读取 build log。工程代码应把 build log 与 platform name、device name、driver version、build options 一起输出。这样才能区分 kernel 源码错误、设备语言版本限制和编译选项不匹配。

Buffer 路径与 image 路径要按访问模式选择。线性后处理、粒子数组、体素列表、prefix sum 和 compacted list 更适合 buffer，因为它们使用线性索引和结构化数据。二维采样、硬件过滤、边界地址模式和图像格式转换更适合 OpenCL image object。图形工程中常见做法是先用 buffer 路径验证计算正确性，再根据采样需求和互操作能力切换 image 路径。

Platform、device、queue、kernel 和 buffer path 的可复用检查顺序如下：确认 platform 枚举结果；确认目标 device 类型和能力；创建 context 并记录设备归属；创建 profiling queue；构建 program 并保存 build log；创建 memory object 并检查大小和 flag；设置 kernel arg 并记录每个参数的语义；enqueue write、kernel、read 或 interop acquire/release；用 event 验证依赖和耗时。这个顺序让错误能够停在最早可解释的位置。

109.4 OpenCL Kernel Launch Memory and Synchronization Pitfalls

OpenCL 的常见故障通常集中在四类位置：work-group 参数、内存对象和传输、barrier 语义、queue dependency 和 kernel build。它们都能在本章的后处理图像任务中复现，也都能用可观察证据定位。

Work-group size 的问题先看合法性，再看性能。合法性由 CL_DEVICE_MAX_WORK_GROUP_SIZE、kernel work-group info、local memory 使用量、维度限制和 global size 整除条件共同决定。OpenCL 2.0 以后支持非均匀 work-group 的设备可以处理某些尾部 group，但为了兼容旧设备和简化调试，本章示例把 global size 向上对齐，并在 kernel 中用边界判断保护尾部 work-item。性能上，local size 过小会增加调度开销，过大会降低 occupancy 或触发寄存器、local memory 压力。

内存传输问题要分清 host-device copy、device memory access 和图形互操作。clEnqueueWriteBuffer 和 clEnqueueReadBuffer 的时间主要反映传输路径和同步等待；kernel event 的 start-end 时间主要反映设备执行。若输出画面正确但 read 阶段耗时高，瓶颈在回读路径；若 kernel 时间高且内存访问是连续读写，瓶颈可能在 global memory bandwidth；若图形端和 OpenCL 端之间频繁 acquire/release 或 copy，瓶颈可能来自跨 API 所有权转换。

Barrier 语义只在同一个 work-group 内提供同步。Kernel 中的 barrier(CLK_LOCAL_MEM_FENCE) 可以约束同组 work-item 对 local memory 的访问顺序；它不提供跨 work-group 的全局同步。跨 work-group 的阶段分隔应拆成多个 kernel enqueue，并通过 event 或同一 in-order queue 建立顺序。这个边界会影响粒子模拟、prefix sum、tile reduction 和体数据统计等任务。

Queue dependency 的问题通常表现为偶发输出错误。In-order queue 能让同一 queue 中的命令按提交顺序依赖，多个 queue 或 out-of-order queue 则要求显式 event wait list。若一个 queue 写入 output_buffer，另一个 queue 读取同一 buffer，没有 event 依赖时就可能读到旧数据或部分写入结果。图形互操作路径还要把 OpenCL event 对齐到图形 queue 的 fence 或 semaphore，资源所有权转换完成后再让下游 pass 消费。

Kernel build 错误要用 build log 定位。OpenCL C 的地址空间限定符、向量类型、函数支持、编译选项和扩展开关都可能导致不同设备上的构建差异。工程上应把 build 作为可诊断阶段处理：失败时输出源码版本、build options、device name、OpenCL C version 和 log；成功时也记录 program binary cache key，便于后续减少重复 build 成本。

下面这组症状可以作为排查入口。它们按“结果可见性 → API 错误 → 计时证据 → 资源依赖”的顺序组织，适合在没有大型 profiler 的情况下使用。

• 输出全黑：先检查 kernel 是否真的 enqueue 成功，再检查阈值参数、输入 buffer 写入、global size 和 gid 边界。若 build 失败或 kernel name 拼错，输出 buffer 可能保持初始值。
• 输出随机闪烁：优先检查 queue 之间的 event 依赖、图形互操作的 acquire/release 顺序、host 是否过早复用输入内存。闪烁往往来自命令顺序不稳定。
• kernel 很快但整帧很慢：检查 write/read buffer 的 event 时间、blocking read、host 侧 clFinish 调用位置和图形端同步等待。这个症状说明计算阶段可能已经完成，数据移动或等待成为主要开销。
• 只在某些设备失败：读取 device limits、OpenCL C version、extension list、maximum work-group size 和 build log。跨设备失败通常来自能力假设、编译器差异或 work-group 参数越界。
• 调大 local size 后变慢：检查 occupancy、寄存器压力、local memory 使用和内存合并访问。local size 提升只改变分块方式，不自动提升内存带宽。

对于本章的后处理 kernel，可复用的排查顺序是：先用小图像和 CPU device 校验输出，再切到目标 GPU；先用 in-order profiling queue 固定命令顺序，再尝试多 queue；先用 buffer 路径建立正确性，再引入 image object 或图形互操作；先记录 event 时间，再用 RenderDoc、Nsight、Xcode GPU tools 或 OpenCL intercept layer 观察更细的资源和命令证据。工具只回答“命令、资源和时间发生了什么”，最终判断仍要回到数据路径。

109.5 在图形工作流中的应用场景

OpenCL 在图形工作流中的价值，来自它把通用并行计算放在渲染管线旁边。它适合处理与 draw call 关系松散、数据并行明显、可用 buffer 或 image 表达的任务。它的输出可以进入下一帧渲染、离线资产管线、科学可视化工具或机器视觉预处理路径。

图像处理是最直接的应用。阈值、模糊、边缘检测、色彩转换、直方图、mipmap 预处理和离线降噪都能映射到 work-item 处理像素或 tile 的模式。图形工程中要判断它是否适合 OpenCL，先看输入输出是否已经在设备侧，再看是否存在高频回读。若每帧都把 render target 从 GPU 拷贝到 host 再传回 OpenCL，传输成本可能超过 kernel 成本；若资源能通过互操作或同设备内存路径共享，OpenCL 更容易发挥作用。

粒子模拟适合用 OpenCL 表达更新阶段。每个粒子可以作为一个结构体或 SoA buffer 元素，kernel 根据速度、力场、寿命和碰撞代理更新状态，渲染端再用输出 buffer 作为 instance data 或 vertex stream。这里的关键依赖是 compute update 与 render consume 的顺序。OpenCL 端写完粒子 buffer 后，图形 queue 要等到更新完成；渲染端读取完成后，下一轮模拟才能安全覆盖或双缓冲写入。

体数据处理适合 OpenCL 的原因是数据规模大、访问模式可分块、算法阶段清晰。医学体数据、流体体素、SDF 生成、体纹理预处理和切片统计都可以用 OpenCL 做离线或准实时处理。工程上要关注 3D image 支持、内存容量、分块 streaming、边界采样和中间结果格式。若任务最终进入 volume rendering，输出格式还要对齐后续采样方式和 transfer function 需求。

离线工具是 OpenCL 的稳定使用场景。纹理批处理、mesh attribute bake、光照图预计算、SDF bake、视频帧预处理和数据集转换都能接受较长启动时间，也能容忍不同设备的性能差异。离线场景的优先目标是吞吐和可复现结果，因此 build log、设备记录、输入 hash、输出校验和错误重试比单帧延迟更关键。

跨设备计算是 OpenCL 的经典优势，也是图形系统中的复杂点。一个工具可以在没有高端独立 GPU 的机器上使用 CPU device 完成同一算法，也可以在工作站上选择 GPU device 提升吞吐。跨设备策略要把“能力可用”和“性能合适”分开判断：CPU device 便于调试和精确日志，GPU device 便于大规模并行，FPGA 或专用加速器可能适合固定工作负载。统一 API 降低了移植成本，性能模型仍要按设备重建。

在现代渲染引擎里，OpenCL 需要和 compute shader、CUDA、Vulkan compute、Metal compute、Direct3D compute 放在同一组维度比较。OpenCL 的优势是开放标准、异构设备覆盖和离线/工具链友好；compute shader 的优势是与渲染 API 的资源、pipeline state 和同步模型高度一致；CUDA 的优势是 NVIDIA 生态内的工具、库和硬件特性；Metal 和 Direct3D compute 的优势是平台集成。选择 OpenCL 时，最重要的判断不是 API 名称，而是资源是否能低成本进入 OpenCL，输出是否能低成本回到消费端。

回到本章贯穿材料，亮度阈值后处理可以有三种落地方式。工具型应用可以把输入图像读入 OpenCL buffer，执行 kernel 后写出文件；实时渲染可以通过互操作共享资源，减少 host 回读；跨平台编辑器可以保留 CPU fallback，在 OpenCL GPU device 可用时启用加速路径。三种方式共用同一套对象模型，但同步证据、性能瓶颈和失败处理不同。

本章最终建立的理解是：OpenCL 程序不是一个孤立 kernel，而是一条由 platform、device、context、queue、program、kernel、memory object 和 event 连接起来的数据路径。读者面对一个 OpenCL 图形计算任务时，应先把这条路径画出来，再检查对象归属、能力查询、内存移动、命令依赖和输出消费。只有这条路径闭合，kernel 优化和高级互操作才有可靠前提。

最小自检任务

给定一个 1024 × 1024 的 RGBA 输入图像，你要用 OpenCL 执行一次亮度阈值筛选，并把结果交给后续 bloom 预处理。请写出最小对象路径和排查顺序：需要哪些 OpenCL 对象，kernel 的 global size 和 local size 如何设定，哪些 event 需要保留，若输出全黑或整帧耗时过高应先检查哪些证据。任务不要求运行完整工程，也不要求安装特定 GPU profiler。

答案要点

最小对象路径应包含 platform、device、context、command queue、program、kernel、input buffer、output buffer 和 event。Platform 和 device 通过 query 选择，context 作为 queue、program、kernel 和 buffer 的共同归属边界。Command queue 可以先使用 in-order profiling queue，便于把 write、kernel 和 read 的时间拆开观察。

Global size 应覆盖像素总数，1024 × 1024 对应 1,048,576 个 work-item。Local size 可以先选 128 或 256，并检查 device 和 kernel 的 work-group 限制。若 global size按 local size 对齐，kernel 内部要用 gid >= pixel_count 保护尾部 work-item；本题中像素总数能被 256 整除，边界判断仍应保留，便于迁移到其它尺寸。

Event 至少保留 input write event、kernel event 和 output read event。Kernel enqueue 应等待 input write event，output read 或后续消费命令应等待 kernel event。若输出直接交给图形 queue，OpenCL 完成事件还要转换为图形端可等待的 fence、semaphore 或对应互操作同步流程。

输出全黑时，先查 API 返回值和 clBuildProgram 的 build log，再查 kernel name、kernel 参数顺序、threshold 值、input buffer 写入范围、global size 和 gid 边界。整帧耗时过高时，先用 profiling event 区分 write、kernel、read 三段时间。若 write/read 高，重点检查拷贝和同步等待；若 kernel 高，检查内存访问模式、local size、occupancy 和设备带宽；若图形端等待高，检查资源互操作和队列依赖。

本章知识点总结

• OpenCL 定位：OpenCL 是面向异构设备的并行计算 API，工程价值来自 host 侧统一组织 device kernel、memory object 和 command queue。
• Host 角色：Host 负责平台枚举、设备选择、context 创建、program build、kernel 参数设置、命令提交和结果同步。
• Device 角色：Device 执行 kernel，并把 work-item 映射到 compute unit 和 processing element，具体映射由实现和硬件决定。
• Context 边界：Context 是 queue、program、kernel 和 memory object 的共同归属边界，对象归属一致是 enqueue 成功的前提。
• Queue 语义：Command queue 负责提交写入、kernel 执行、读取和同步命令，profiling queue 能提供阶段耗时证据。
• Kernel 路径：一次 OpenCL kernel launch 需要 program build、kernel object、buffer 或 image、kernel arg、global size、local size 和 event 依赖共同成立。
• Work-group 判断：Local size 要同时满足设备限制、kernel 资源压力和访问模式，global size 对齐时要用边界判断保护尾部 work-item。
• 内存成本：图形工作流中的 OpenCL 性能经常受拷贝、回读、互操作同步和 global memory bandwidth 限制。
• 同步边界：Kernel 内 barrier 只同步同一 work-group，跨 queue、跨 kernel 或跨图形 API 的顺序要通过 event、fence 或互操作同步建立。
• Build 日志：Kernel 编译失败要读取 build log，并记录 device name、OpenCL C version、build options 和源码版本。
• 图像处理场景：OpenCL 适合阈值、模糊、边缘检测、色彩转换和离线降噪等数据并行图像任务。
• 粒子场景：OpenCL 可更新粒子 buffer，渲染端消费前必须等待 compute update 完成，并常用双缓冲管理读写冲突。
• 体数据场景：OpenCL 适合体素预处理、SDF 生成和切片统计，输出格式要对齐后续 volume rendering 的采样需求。
• 选择依据：OpenCL、compute shader、CUDA、Metal compute 和 Direct3D compute 应按资源进入成本、同步模型、平台覆盖和工具证据比较。
• 排查顺序：先画出 platform、device、context、queue、kernel、memory 和 event 路径，再定位能力、传输、编译、同步或执行瓶颈。