Taichi GPU 内核如何用 ti.sync() 正确测量执行时间?
【免费下载链接】taichiProductive, portable, and performant GPU programming in Python.项目地址: https://gitcode.com/GitHub_Trending/ta/taichi
在 Taichi 中给 GPU 内核计时是一个容易踩坑的操作:使用 GPU 后端时,内核被编译后是发送到 GPU 队列执行的,Python 程序不会等待内核执行完就继续执行下一条语句。如果你直接在内核调用前后用time.time()掐表,测出来的时间并不包含内核在 GPU 上的实际执行时间。本文基于 kernel_sync 文档,说明如何在内核调用后显式调用ti.sync()等待内核完成,从而得到正确的执行时间。
适用前提:程序运行在 GPU 后端(ti.init(arch=ti.gpu)),内核本身没有返回值、不访问 field、也没有print等需要同步数据的操作。
为什么直接计时会漏掉 GPU 执行时间
下面的代码来自官方文档,定义了一个计算量很大的benchmark内核并尝试测量其耗时:
import time import taichi as ti ti.init(arch=ti.gpu) @ti.kernel def benchmark(): x = 1.0 for i in range(1000): for j in range(10000): x += ti.sin(float(i + j)) start = time.time() benchmark() end = time.time() print(end - start)文档指出:虽然benchmark是一个计算密集的内核,但程序不会等待它执行完才执行后续语句,benchmark的执行时间被排除在计时之外。也就是说,这里打印出来的end - start只覆盖了内核编译与提交到 GPU 队列的开销,远小于 GPU 上真正的计算时间。
在内核调用后加 ti.sync() 再计时
修正方法是在内核调用之后、读取结束时间之前,显式调用ti.sync():
import time import taichi as ti ti.init(arch=ti.gpu) @ti.kernel def benchmark(): x = 1.0 for i in range(10000): for j in range(100000): x += ti.sin(float(i + j)) start = time.time() benchmark() ti.sync() end = time.time() print(end - start)ti.sync()会阻塞程序,等待 GPU 队列中排在前面的所有内核任务执行完毕后才继续执行后续语句,因此end - start此时包含了内核在 GPU 上的执行时间。FAQ 中也提到,ti.sync()的作用类似 CUDA 的cudaStreamSynchronize()。
验证方式同样是行为层面的:对比两段代码,第一段(无ti.sync())的计时结果不包含内核执行时间,第二段在benchmark()与end = time.time()之间加入ti.sync()后,计时结果才会反映 GPU 上的实际执行耗时。文档没有给出固定的预期数值,实际输出取决于你的硬件。
哪些情况下不需要手动 ti.sync()
大多数时候 Taichi 会自动处理数据同步,以下情形无需手动调用ti.sync():
- 内核有返回值:Taichi 会等待内核完成后才把返回值交还给 Python 作用域;
- 在 Python 作用域调用
x.to_numpy():如果 fieldx正被其他内核使用,Taichi 会确保使用该 field 的内核先完成; - 在 Python 作用域访问或修改 Taichi field:同样会先确保使用 field 的内核结束。
文档示例之所以必须显式调用ti.sync(),正是因为这个内核没有返回值、不含 print 语句、也不涉及任何 field 操作,属于典型的"发射后不管"内核。
另外注意 CPU 后端的差异:在 CPU 后端上所有内核调用都是阻塞的,程序会等待当前内核结束再执行下一条语句,因此CPU 后端既不用担心数据依赖,也不需要调用ti.sync()。ti.sync()的计时修正只针对 GPU 后端。
可选:用 KernelProfiler 获取内核级性能数据
如果目标不止是测一次耗时,而是想看每个内核的统计信息,可以在ti.init()时启用kernel_profiler=True,再用ti.profiler.print_kernel_profiler_info()输出结果(支持 "count" 与 "trace" 两种模式)。Profiler 文档 特别强调了两点:
KernelProfiler目前仅支持 CPU 和 CUDA 后端;- 程序运行在 GPU 上时,执行性能分析前要先调用
ti.sync(),与本文的计时要求一致。
文档还建议在多次运行性能分析后观察最小或平均执行时间,而不是只看单次结果。
【免费下载链接】taichiProductive, portable, and performant GPU programming in Python.项目地址: https://gitcode.com/GitHub_Trending/ta/taichi
创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考