Core Ultra 9 285K 的 iGPU 能打过 CPU 吗?SYCL 五内存模式实测

发布于:
更新于:
文章字数:
3621 字
阅读时间:
15 分钟
浏览量: 加载中...

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_allocatorsycl::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
  1. 为什么拷贝平手?
  • 双方处在 71–78.1 GB/s 区间,iGPU 与 CPU 共享同一条 DDR5(DDR5-6400 双通道理论峰值 102.4 GB/s),物理上限本来就是同一个;copy 指标 = 实际总线流量,78.1 GB/s 已是理论峰值的约 76%,两边都触及了代码执行所能到达的最大峰值性能 (跑满带宽需要更大的数组)。
  1. 为什么 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 开销,有效吞吐推得更满。
  1. 为什么 CPU 读(66.6)又能略胜 GPU 读(61.0)?
  • 读内核两侧其实不同构:CPU 是一段可被 AVX2 完全向量化的紧致求和循环,硬件预取器把顺序读拉满;GPU 读用的是 SYCL reduction(树形归约),测的不是纯读带宽——额外的归约组合与同步开销混在指标里。共享 DDR5 下纯读上限本就同量级,这点结构差异足以让 CPU “略胜”。
  1. 为什么 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 并行,但每步都吃亏:

  1. 线程上限 24,可并行的蝶形槽位远少于 GPU;
  2. 高 pass 的蝶形跨步越来越大(len 从 2 涨到 n),顺序性崩坏、硬件预取失效
  3. 每个 pass 结束还有一次 TBB 同步屏障;
  4. 计时口径里还多含 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 的软肋上:

  1. 块内天生串行:匹配位置只能一步步前推,命中后还要逐字节回填哈希表,块内零并行空间——GPU 的海量线程在块内完全用不上;
  2. 并行度只有 2048:全场只有 2048 个块(每块一个 GPU 工作项),摊到核显上喂不饱所有执行单元,而 24 个 CPU 线程吃 2048 个块绰绰有余;
  3. 哈希表随机访问放大缓存差距:每块一张 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

两个要点:

  1. 依赖链延迟上 CPU 单线程就是快 3.6 倍(91.5 vs 330.4 ns)——核显核心的依赖访问代价远高于 CPU 核。
  2. 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 不是”小显卡”,而是一个”另一套脾气的协处理器”:

  1. 吞吐型规则并行优先丢给 GPU——FFT(3.4×)、批量查找(3.6×)、顺序大块写入(1.6×)是它的主场;聚合访存吞吐也能到 CPU 24 线程的近 3 倍。
  2. 分支密集 / 依赖链重的活留在 CPU——压缩(24.9×,串行都赢)、依赖型延迟(快 3.6 倍)是它的绝对领域;分块 GEMM 这类 SIMD 友好算力也以 2.2× 领先核显。
  3. 五种 SYCL 内存模式不是性能旋钮——差异 3–14%,malloc_device 几乎全面最佳(仅拷贝带宽与 zip 由 buffer 微弱领先);按编程便利性选即可。
  4. 算上传输再做决定——H2D/D2H 每 256 MiB 约 16.5 ms(≈16.3 GB/s),数据已在设备侧的长驻任务才划得来;频繁与主机来回倒数据的短任务,CPU 更省心。
  5. 共享内存是 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
2
3
./run_bench.cmd dpcpp                    # iGPU(SYCL)侧
./run_bench.cmd cpu --threads 24 # CPU 侧
python ./tools/bench_to_labplot.py # 汇总为宽表 CSV + LabPlot 工程

测试工程要求:Intel 6 代及以后带核显的酷睿 CPU、支持 AVX2;不含 dGPU / AMD iGPU 支持。

完整测试工程(源码 + 运行脚本 + 结果转换工具)已开源:github.com/1992724048/sycl-igpu-benchmark


术语表
  1. Visual Studio
  2. DPC++