Radxa平台OpenCL调用计时波动大的原因？是否使用共享内存？

测试环境 #

• 硬件：Radxa Dragon Q6A（搭载Adreno 643，Qualcomm QCS6490 SoC）
• 系统：Ubuntu 24.04.4 aarch64
• 软件：Python 3.14.5、pyopencl、clinfo

可用OpenCL平台 #

(.venv) rock@radxa-dragon-q6a:~/bench$ clinfo -l
clinfo: /lib/aarch64-linux-gnu/libOpenCL.so.1: no version information available (required by clinfo)
Platform #0: QUALCOMM Snapdragon(TM)
 `-- Device #0: QUALCOMM Adreno(TM) 635
Platform #1: rusticl
 `-- Device #0: FD643

Python测试脚本 #

import numpy as np, pyopencl as cl, os, argparse
from time import perf_counter

parser = argparse.ArgumentParser()
parser.add_argument('-p', '--platform', default='0')
os.environ['PYOPENCL_CTX'] = parser.parse_args().platform

rng = np.random.default_rng()
ctx = cl.create_some_context()
queue = cl.CommandQueue(ctx)
mf = cl.mem_flags

prg = cl.Program(ctx, """
__kernel void sum(
    __global const float *a_g, __global const float *b_g, __global float *res_g)
{
  int gid = get_global_id(0);
  res_g[gid] = a_g[gid] + b_g[gid];
}
""").build()
knl = prg.sum  

for i in range(10):
  a_np = rng.random(50000, dtype=np.float32)
  b_np = rng.random(50000, dtype=np.float32)
  res_np = np.empty_like(a_np)
  t0 = perf_counter()
  a_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=a_np)
  b_g = cl.Buffer(ctx, mf.READ_ONLY | mf.COPY_HOST_PTR, hostbuf=b_np)
  res_g = cl.Buffer(ctx, mf.WRITE_ONLY, a_np.nbytes)
  t1 = perf_counter()
  knl(queue, a_np.shape, None, a_g, b_g, res_g)
  t2 = perf_counter()
  cl.enqueue_copy(queue, res_np, res_g)
  t3 = perf_counter()

  print(f'{1000*(t1-t0):.3f}  {1000*(t2-t1):.3f}  {1000*(t3-t2):.3f} [ms]  {np.allclose(res_np, a_np + b_np)}')

测试结果 #

Adreno 635平台结果 #

Choosing only available device: <pyopencl.Device 'QUALCOMM Adreno(TM) 635' on 'QUALCOMM Snapdragon(TM)' at 0xffff8e60e900>
0.667  0.625  2.956 [ms]  True
1.102  0.096  1.042 [ms]  True
0.706  0.079  1.022 [ms]  True
0.775  0.115  1.553 [ms]  True
0.870  0.081  1.202 [ms]  True
1.649  0.277  2.219 [ms]  True
1.132  0.100  1.155 [ms]  True
0.782  0.083  0.950 [ms]  True
0.788  0.073  0.996 [ms]  True
0.806  0.075  1.441 [ms]  True

rusticl平台结果 #

Choosing only available device: <pyopencl.Device 'FD643' on 'rusticl' at 0x2d14ae10>
3.128  0.092  4.975 [ms]  True
1.537  0.082  3.084 [ms]  True
0.957  0.079  4.040 [ms]  True
1.750  0.116  3.285 [ms]  True
2.099  0.166  4.157 [ms]  True
0.752  0.087  3.471 [ms]  True
0.588  0.078  4.282 [ms]  True
0.940  0.076  2.782 [ms]  True
1.433  0.120  3.222 [ms]  True
0.903  0.091  2.811 [ms]  True

疑问 #

1. 为什么两个平台的计时结果均存在明显波动，尤其是缓冲阶段？
2. 本测试是否用到了Qualcomm QCS6490 SoC应具备的CPU/GPU共享内存？

一、计时波动的原因 #

• 系统后台干扰：即使仅打开终端，Ubuntu仍有后台守护进程、内存回收、磁盘I/O等活动，会抢占CPU或总线资源，导致内存分配、数据拷贝的延迟不稳定。
• 动态内存分配：每次循环都重新创建cl.Buffer，系统需要动态分配设备内存，这一过程受当前内存碎片、可用内存量影响，缓冲阶段波动因此被放大。
• OpenCL运行时调度：第一次执行时运行时会完成内核缓存、设备初始化等操作，后续循环虽会复用，但运行时的动态调度仍可能带来延迟波动。
• 小数据量放大误差：测试仅处理约200KB数据，操作本身耗时极短，微小的系统干扰就会导致相对明显的波动。

二、是否用到CPU/GPU共享内存？ #

你的测试没有用到QCS6490的CPU/GPU共享内存，原因如下：

• 当前代码使用mf.COPY_HOST_PTR创建Buffer，会将主机内存的数据拷贝到设备专属内存，而非直接映射共享内存区域。
• 要利用共享内存，需创建共享内存对象：
  1. 先确认设备支持CL_MEM_ALLOC_HOST_PTR或CL_MEM_USE_HOST_PTR（对应统一内存架构特性）。
  2. 使用mf.READ_ONLY | mf.USE_HOST_PTR等共享内存标志创建Buffer，避免数据拷贝，让GPU直接访问CPU内存区域。
• 针对Qualcomm Adreno平台，通常还需启用统一内存相关的OpenCL扩展，或使用特定内存分配方式才能激活共享内存特性。

内容的提问来源于stack exchange，提问作者Paul Jurczak