Core Ultra 9 285K 的 iGPU 能打过 CPU 吗?SYCL 五内存模式实测
EDIT测试代码由AI辅助生成, 文章由AI辅助生成并经过人工修改
测试结果
相同的算法,在 Intel Core Ultra 9 285K 的 CPU 与 核显(iGPU) 上各跑一遍的测试结果:
| 测试项 | CPU (24 Core) | iGPU (64 EU) | 最优设备 | 倍率 |
|---|---|---|---|---|
| GEMM 算力 | 355.4 GFLOP/s | 160.9 GFLOP/s | CPU | 2.2× |
| 压缩吞吐 | 11169.4 MB/s | 448.8 MB/s | CPU | 24.9× |
| 读带宽 | 66.6 GB/s | 61.0 GB/s | CPU | 1.1× |
| FFT 吞吐 | 43.1 MSamples/s | 145.4 MSamples/s | iGPU | 3.4× |
| 查询吞吐 | 122.7 Mqueries/s | 444.8 Mqueries/s | iGPU | 3.6× |
| 写带宽 | 54.0 GB/s | 87.6 GB/s | iGPU | 1.6× |
| 拷贝带宽 | 78.1 GB/s | 77.8 GB/s | 误差范围内表现视为相同 | ≈1.0× |
测试平台与细节
| 项目 | 配置 |
|---|---|
| CPU | Intel Core Ultra 9 285K,D2D/NGU Up to 3200 MHz |
| 内存 | DDR5-6400 CL32 48 GB × 2 |
| iGPU | Arrow Lake-S 核显(Xe 架构,oneAPI 设备名 Intel(R) Graphics),共享系统内存 |
| 编程模型 | SYCL 2020(Intel oneAPI DPC++ Compiler 2026.1) |
| CPU 侧并行 | oneTBB 2023.1,24 线程(tbb_24);另含串行基线(serial) |
| 构建 | Visual Studio 2026,Intel oneAPI Toolkit 2026.1,两侧均 Release |
测试细节
- GPU 内存模式:
malloc_device/malloc_host/malloc_shared/usm_allocator(sycl::usm_allocator+std::vector)/buffer(sycl::buffer + sycl::accessor) - 计时设计:统一使用
std::chrono::steady_clock主机端计时,SYCL 未使用 profiling;每项 warmup 1 次 + 重复 3 次取最优值。 - 测试规模:zip 128 MiB / 2048 块、FFT n = 4,194,304(2²²)复数点、GEMM n = 4096、lookup 1,048,576 次查询、带宽数组 256 MiB、指针追逐 4,194,304 步。
测试项目对比
1. 内存带宽:写入 GPU 强,读取 CPU 略胜,拷贝平手
Workload:256 MiB 数组的 read / write / copy 顺序访问,GB/s = 字节数 / 时间(copy 为 2×字节数)。
| 来源 | 读取 (GB/s) | 写入 (GB/s) | 拷贝 (GB/s) |
|---|---|---|---|
| CPU serial | 7.6 | 18.1 | 47.3 |
| CPU tbb_24 | 66.6 | 54.0 | 78.1 |
| GPU malloc_device | 61.0 | 87.6 | 76.1 |
| GPU malloc_host | 56.2 | 64.1 | 75.8 |
| GPU malloc_shared | 55.7 | 64.0 | 75.3 |
| GPU usm_allocator | 53.9 | 61.5 | 71.1 |
| GPU buffer | 56.5 | 64.0 | 77.8 |
- 为什么拷贝平手?
- 双方处在 71–78.1 GB/s 区间,iGPU 与 CPU 共享同一条 DDR5(DDR5-6400 双通道理论峰值 102.4 GB/s),物理上限本来就是同一个;copy 指标 = 实际总线流量,78.1 GB/s 已是理论峰值的约 76%,两边都触及了代码执行所能到达的最大峰值性能 (跑满带宽需要更大的数组)。
- 为什么 GPU 写入(87.6 GB/s)比 CPU(54.0 GB/s)快?
- 两侧实现都是”单个元素填充”循环,CPU 写路径受微架构中 RFO(Read-For-Ownership,写前读)机制影响:
- CPU 缓存写是 write-allocate 策略:要写入的缓存行不在缓存时,必须先发一次 RFO 读——把整行读回来取得所有权,然后才能标脏、逐级淘汰写回。一次”写”实际混入一次读往返,RFO 读流量还额外占用 DDR5 带宽,流水线被缓存行粒度的分配—标脏—写回拖住;
- GPU 把数千 work-item 的写合并成大段连续写,以大块粒度取得行所有权、避开逐行 RFO 往返,直接灌满内存控制器的写队列——同一条 DDR5,GPU 的写路径少了 RFO 开销,有效吞吐推得更满。
- 为什么 CPU 读(66.6)又能略胜 GPU 读(61.0)?
- 读内核两侧其实不同构:CPU 是一段可被 AVX2 完全向量化的紧致求和循环,硬件预取器把顺序读拉满;GPU 读用的是 SYCL
reduction(树形归约),测的不是纯读带宽——额外的归约组合与同步开销混在指标里。共享 DDR5 下纯读上限本就同量级,这点结构差异足以让 CPU “略胜”。
- 为什么 CPU 串行版低到 7.6 GB/s / 18.1 GB/s?
- 单线程发出的未完成访存请求数有限,打不满多通道内存控制器(读的求和还有累加依赖链);写侧靠合并缓冲吸收,所以串行写(18.1)反而高于串行读(7.6)。这类带宽测试对并发访存请求数极为敏感。
2. GEMM 矩阵乘:CPU 大胜
Workload:分块矩阵乘,n = 4096,算力口径 2·n³ / t;TBB 按行块并行。
| 来源 | GEMM 算力 (GFLOP/s) |
|---|---|
| CPU serial | 23.3 |
| CPU tbb_24 | 355.4 |
| GPU malloc_device | 160.9 |
| GPU malloc_shared | 157.1 |
| GPU usm_allocator | 157.0 |
| GPU buffer | 156.9 |
| GPU malloc_host | 156.2 |
tbb_24 打出 355.4 GFLOP/s,是 GPU 最佳值(160.9)的 2.2 倍。GPU 五种内存模式彼此差距不到 3%——算力瓶颈不在分配方式上。
为什么 CPU 赢? GEMM 是”高频 + SIMD FMA + 缓存分块”的教科书主场:分块 64 让 A/B 小块驻留缓存,内层循环被编译成连续的 AVX2 FMA 指令流,24 个核心的频率与 FMA 火力全开。GPU 侧实现是 16×16 work-group + 双层 local 内存 tile 的朴素 FP32 乘加(未用半精度或矩阵加速指令),而核显 Xe 核在 FP32 算力与频率上本就不是 CPU 大核的对手——输的是硬算力,不是内存模式,五模式挤在 3% 以内也印证了这点。
3. FFT:GPU 完胜
Workload:radix-2 就地 FFT,n = 4,194,304 复数点,算力口径 5·n·log2(n);GPU 侧 22 个蝶形 pass + 1 个 bit-reversal 共 23 次内核提交。
| 来源 | FFT 吞吐 (MSamples/s) | FFT 算力 (GFLOP/s) |
|---|---|---|
| CPU serial | 19.0 | 2.1 |
| CPU tbb_24 | 43.1 | 4.7 |
| GPU malloc_device | 145.4 | 16.0 |
| GPU malloc_shared | 140.1 | 15.4 |
| GPU usm_allocator | 139.8 | 15.4 |
| GPU malloc_host | 139.7 | 15.4 |
| GPU buffer | 125.3 | 13.8 |
GPU 最佳是 tbb_24 的 3.4 倍。buffer 模式在这里掉了约 14%(125.3 vs 145.4),是五模式中差距最大的一项。
为什么 GPU 赢? 蝶形之间零依赖——每个 pass 的 n/2 个蝶形完全独立,GPU 一个 kernel 把它们全部铺开,数千线程轮流吃访存延迟。CPU 侧 TBB 虽然也逐 pass 并行,但每步都吃亏:
- 线程上限 24,可并行的蝶形槽位远少于 GPU;
- 高 pass 的蝶形跨步越来越大(len 从 2 涨到 n),顺序性崩坏、硬件预取失效;
- 每个 pass 结束还有一次 TBB 同步屏障;
- 计时口径里还多含 2n 个 float 的输入拷贝(见测试局限)。
FFT 是”独立小运算 + 规则访存”的教科书场景,正是 GPU 的主场。
4. 批量二分查找:GPU 完胜
Workload:1,048,576 个查询在有序数组上二分查找(命中 524,862 次),每工作项一个查询;两侧校验和一致。
| 来源 | 查询吞吐 (Mqueries/s) |
|---|---|
| CPU serial | 8.8 |
| CPU tbb_24 | 122.7 |
| GPU malloc_device | 444.8 |
| GPU malloc_host | 433.7 |
| GPU malloc_shared | 433.7 |
| GPU usm_allocator | 433.1 |
| GPU buffer | 427.5 |
GPU 最佳是 tbb_24 的 3.6 倍、是串行的 50 倍。这也是五模式差距最小的一项(最差与最佳仅差 4%)。
为什么 GPU 赢? 二分查找每查询最多 log2(1048576) = 20 步串行依赖的随机访问——下一步的 mid 要等上一步结果,CPU 上这是”分支预测失败 + 缓存未命中”的双重打击,串行只有 8.8 Mqueries/s 正是单条依赖链的真实速度。GPU 赢的不是”单次访问更快”(延迟测试里 GPU 单次访问反而慢 3.6 倍),而是查询彼此独立:百万查询摊到数千线程,一个线程等内存时另一个在算,用线程级并行把依赖延迟整个藏掉。五模式仅差 4% 说明此时瓶颈既不在算力也不在带宽,只在延迟隐藏能力。
5. 块压缩(zip):CPU 碾压
Workload:128 MiB 输入 / 2048 块的块压缩(哈希表匹配),仅压缩计时(含哈希表清零),解压回读逐字节校验;两侧压缩比同为 1.30。
| 来源 | 压缩吞吐 (MB/s) |
|---|---|
| CPU serial | 533.8 |
| CPU tbb_24 | 11169.4 |
| GPU buffer | 448.8 |
| GPU malloc_device | 446.9 |
| GPU malloc_host | 443.8 |
| GPU malloc_shared | 443.3 |
| GPU usm_allocator | 442.0 |
这是全场最悬殊的一项:tbb_24 是 GPU 的 24.9 倍,甚至串行 CPU(533.8)都快过所有 GPU 模式(≈447)。
为什么连串行 CPU 都能赢? 三个原因叠加,每条都打在 GPU 的软肋上:
- 块内天生串行:匹配位置只能一步步前推,命中后还要逐字节回填哈希表,块内零并行空间——GPU 的海量线程在块内完全用不上;
- 并行度只有 2048:全场只有 2048 个块(每块一个 GPU 工作项),摊到核显上喂不饱所有执行单元,而 24 个 CPU 线程吃 2048 个块绰绰有余;
- 哈希表随机访问放大缓存差距:每块一张 4096 项(16 KB)的哈希表按块散布,块内全是哈希随机读写。CPU 上 16 KB 轻松驻留 L1/L2,命中即快;核显要在更大的地址空间里踩这些散落的小表。候选字比对这类数据相关分支也一样——CPU 的分支预测与乱序执行远比 SIMT 分支处理高效。
串行 CPU 的哈希表同样在缓存里,块内串行恰好是单核的舒适区——所以 533.8 反超全部 GPU 模式也就顺理成章。
深挖细节
GPU 五种内存模式差异很小,malloc_device 全面最佳
把五模式在各测试项的排名拉通看:malloc_device 几乎全面第一——读/写带宽、FFT、GEMM、lookup、延迟均为五模式最佳。仅两项例外:拷贝带宽与 zip 压缩由 buffer 以不到 3% 的微弱优势领先(77.8 vs 76.1 GB/s、448.8 vs 446.9 MB/s),而 buffer 在 FFT、lookup、延迟上又均居末位。总体差距普遍在 3–14% 以内——选型上可以放心用 malloc_device(USM device 指针),不必为模式选择过度纠结;真正决定性能的是 workload 本身。
延迟:CPU 单线程延迟低得多,GPU 靠聚合吞吐翻盘
指针追逐(4,194,304 步依赖访问,每次访问必须等上一次完成):
| 视角 | CPU | GPU(malloc_device 最佳) |
|---|---|---|
| 单追逐者延迟 | 91.5 ns/access(serial) | 330.4 ns/access |
| 并行聚合吞吐 | 144.2 Maccess/s(24 追逐者) | 417.6 Maccess/s(256 追逐者) |
| 并行摊薄每访问 | 6.9 ns | 2.4 ns |
两个要点:
- 依赖链延迟上 CPU 单线程就是快 3.6 倍(91.5 vs 330.4 ns)——核显核心的依赖访问代价远高于 CPU 核。
- GPU 用 256 个并行追逐者把聚合吞吐做到 417.6 Maccess/s,是 CPU 24 线程的 2.9 倍。表中”2.4 ns”是吞吐摊薄值(1/吞吐)不是真实延迟——并行度把等待藏起来了,单条依赖链的代价依然在那里。这解释了为什么 GPU 赢 FFT/lookup(吞吐型)却输 zip(依赖+分支型)。
主机↔设备传输:256 MiB 约 16.5 ms
| 方向 | 耗时 | 等效带宽 |
|---|---|---|
| H2D 上传 256 MiB | 16.54 ms | 16.2 GB/s |
| D2H 回读 256 MiB | 16.39 ms | 16.4 GB/s |
虽然 iGPU 共享 DDR5、拷贝带宽账面上与 CPU 持平,但 SYCL 的 malloc_device 路径下数据进出”设备”仍要走这笔显式拷贝(read/write 指标不含它)。对一次性小任务,这 16 ms 级固定开销足以吃掉 GPU 的算力优势;malloc_shared / usm_allocator 免显式拷贝,代价是各项带宽略低(差距同样在 5% 内)。
CPU 拷贝阶梯:小尺寸靠缓存,大尺寸回落 DRAM
CPU 侧四档拷贝带宽(copy = 2×字节数 / 时间):
| 档位 | serial (GB/s) | tbb_24 (GB/s) |
|---|---|---|
| 32 KiB | 218.5 | 131.1 |
| 512 KiB | 163.8 | 59.6 |
| 8192 KiB | 31.1 | 224.0 |
| 262144 KiB | 47.1 | 79.1 |
小尺寸完全命中缓存时能冲到 200+ GB/s;256 MiB 落回 DRAM 后就是 47–79 GB/s,与前面的全尺寸带宽一致。tbb 在 32 KiB/512 KiB 反而不如串行(线程分发开销 > 并行收益),8 MiB 档并行收益最大(224.0)。带宽数字不标注尺寸就是耍流氓——同一程序不同档位能差 7 倍。
结论与选型建议
在这台 285K 上,iGPU 不是”小显卡”,而是一个”另一套脾气的协处理器”:
- 吞吐型规则并行优先丢给 GPU——FFT(3.4×)、批量查找(3.6×)、顺序大块写入(1.6×)是它的主场;聚合访存吞吐也能到 CPU 24 线程的近 3 倍。
- 分支密集 / 依赖链重的活留在 CPU——压缩(24.9×,串行都赢)、依赖型延迟(快 3.6 倍)是它的绝对领域;分块 GEMM 这类 SIMD 友好算力也以 2.2× 领先核显。
- 五种 SYCL 内存模式不是性能旋钮——差异 3–14%,
malloc_device几乎全面最佳(仅拷贝带宽与 zip 由buffer微弱领先);按编程便利性选即可。 - 算上传输再做决定——H2D/D2H 每 256 MiB 约 16.5 ms(≈16.3 GB/s),数据已在设备侧的长驻任务才划得来;频繁与主机来回倒数据的短任务,CPU 更省心。
- 共享内存是 iGPU 的结构现实——拷贝带宽与 CPU 打平(≈78 GB/s)说明”显存带宽”这个独立维度在这里不存在,GPU 赢的从来不是带宽,是并发结构。
测试局限
- 单机、单次运行,数值为 warmup 后 3 次重复取最优值,非多次统计均值,绝对值有波动空间;结论看相对量级。
- CPU 侧 FFT 计时含 2n 个 float 输入拷贝,GPU 侧为就地 23 次提交不含同等拷贝——该项 CPU 处于轻微不利口径(量级结论不受影响)。
buffer模式的隐式 H2D/D2H 由运行时管理、不计入指标;malloc_device的 H2D/D2H 单列(上文传输小节)。- 仅测本机 Intel 核显:工具不支持 dGPU 与 AMD iGPU,结论不可外推到独显或其他厂商核显。
- CPU 并行固定 24 线程(
tbb_24);不同线程数、不同 BIOS 功耗墙下数值会变化。
附:复现方式
1 | |
测试工程要求:Intel 6 代及以后带核显的酷睿 CPU、支持 AVX2;不含 dGPU / AMD iGPU 支持。
完整测试工程(源码 + 运行脚本 + 结果转换工具)已开源:github.com/1992724048/sycl-igpu-benchmark。