Booth硬件验证清单:MI300X/RTX 4060 Ti等真实GPU硅片上它跑通了哪些内核

发布时间:2026/10/11 13:27:24
Booth硬件验证清单:MI300X/RTX 4060 Ti等真实GPU硅片上它跑通了哪些内核
【免费下载链接】BoothOpen-source CUDA, Triton and HIP compiler targeting multiple GPU and CPU architectures.项目地址https://gitcode.com/gh_mirrors/bar/Booth点击查看免费下载Booth 是一个开源的 CUDA、HIP 与 Triton GPU 编译器它把内核源码直接编译成 AMD RDNA/CDNA、NVIDIA PTX、Tenstorrent 以及 x86-64 CPU 机器码全程不依赖 LLVM。这篇文章带你逐条核对它的硬件验证清单在 AMD MI300X 与 NVIDIA RTX 4060 Ti 等真实硅片上Booth 编译出的内核究竟跑通了哪些、结果是否经过数值校验。一分钟了解Booth 编译的是谁的 CUDA 内核Booth 的定位一句话就能说清你平时交给nvcc或 ROCm 的那些.cu/.hip文件直接丢给它就能出机器码。它的测试语料共14 个源文件、35 个内核、约 1,700 条 BIR 指令、约 27,000 字节机器码另有覆盖前端、IR、后端与运行时的 274 个 C 测试用例。验证文档就在 docs/hardware.md完整功能状态见 docs/features.md。四大真实平台验证结果一览硬件架构验证方式结果AMD MI300XCDNA3 / GFX942真实硅片运行✅ 8/8 测试内核通过Monte Carlo 中子输运 k_eff 0.995与参考物理值一致NVIDIA RTX 4060 TiAda真实硅片运行✅ PTX 经 CUDA Driver API 加载、驱动 JIT 编译中子输运基准比单线程 CPU 快 3.8 倍全程无 NVCCAMD RDNA3GFX1100模拟 CI✅ 全套测试经 RDNA3 模拟器通过Tenstorrent BlackholeMetalium编译级验证⏳ 产出有效 Metalium C等待开发套件做硬件验证AMD MI300XCDNA3 硅片上 8/8 内核全绿这是 Booth 最硬的一张成绩单。MI300XGFX942上8/8 测试内核全部通过验证内容包括 MFMA 矩阵指令、Wave64 分支分歧处理与 tinygrad 兼容性详见 CHANGELOG.md。更有说服力的是一个真实物理工作负载一个 654 行的Monte Carlo 中子输运内核跑出了 k_eff 0.995与参考解吻合——这不是能编译而是物理结果正确。同期落地的分歧感知 SSA 寄存器分配--ssa-ra把该内核上 186 处 VGPR 溢出全部消除scratch 流量下降 78%总指令数下降 28%。NVIDIA RTX 4060 Ti没有 NVCC 也能跑还快了 3.8 倍NVIDIA 侧走的是 PTX 路线--nvidia-ptx输出 PTX 文本由 CUDA Driver API 加载、NVIDIA 驱动 JIT 编译整条流水线里没有任何 NVCC。RTX 4060 Ti 上的验证结果是同一个 Monte Carlo 中子输运基准产出正确数值且比单线程 CPU 快3.8 倍。后来这条路线还被进一步扩展llama.cpp 的 ggml-cudaCUDA 半边全部 67 个文件到达 Booth 的 IR其中 47 个内核通过新的原生 SASS 后端直接写出显卡可加载的 cubin。另外用 OCaml 写的 vadd 内核经--bir-in读入后生成的 PTX 也在这张 4060 Ti 上跑通了。RDNA3 模拟器 CI没有显卡也能拦下分歧 bug没有消费级 CDNA 卡Booth 用 tinygrad 的 mockgpu 写了 tests/emu/run_emu.py解析.hsacoELF 中的内核描述符在低地址空间里执行真实机器码逐元素校验输出。它先做三级冒烟检查s_endpgm、VOP1全局存储、SMEM 完整管线再完整执行 256 元素的 vectorAdd任何一位数错都算失败。正是这条模拟路径让 wave 分歧类的 guard 子句 bug 在 CI 里现形而不需要一块真卡。35 个内核都跑通了哪些能力测试语料刻意覆盖了 CUDA 的难啃点对应文件都在 tests/ 下tests/vector_add.cu —— GPU 计算的 hello worldtests/cuda_features.cu —— 原子操作、warp 内禀函数、barrier、goto/switch、短路求值tests/test_tier12.cu —— 向量类型、共享内存、运算符重载tests/notgpt.cu —— 分块 SGEMM、规约、直方图、前缀扫描、模板、半精度、协作组tests/stress.cu —— N-body 模拟、嵌套控制流、位操作、结构体按值传参tests/test_coop_groups.cu、tests/test_launch_bounds.cu —— 协作组降低与__launch_bounds__VGPR 上限强制数值回归 CI每个数字都有 SLATEC 基准值兜底除了能跑Booth 还管算得准。tests/numeric/ 下的数值回归 CI 把 Fortrando concurrent内核经 LFortran 转成 CUDA 再走 Booth 全后端输出对照 SLATEC 已知正确值双轴判定accuracy抓真正算坏的内核regression抓同一后端跨版本的 codegen 漂移基线在 tests/numeric/baseline.json。目前收录了 BLAS-1 元素级内核 saxpy、sscal、scopy、sswap、srot、srotm见 tests/numeric/kernels/并刻意选用 200、255 这类不规则长度来暴露 wave-mask 类 bug。想自己定位答案是数学错还是 codegen 错tests/diff/run_diff.sh 提供差分测试同一内核分别走--cpu和--rv64与主机参考逐元素比对CPU 后端就是天然的 oracle。自己动手复现三步验证你的机器获取二进制后直接运行自检./kath doctor它会报告设备与工具链并执行自检编译经典内核./kath --amdgpu-bin vector_add.cu -o vector_add.hsaco看是否输出约 528 字节的 hsaco用./kath run kernel.cu一条命令编译并跑在本机任意设备上所有后端与参数的完整清单见 docs/usage.md哪些 CUDA 特性已支持、哪些还在路上docs/features.md 写得坦诚直接——目前未覆盖纹理、动态并行、宿主端代码生成等但都不是架构性障碍。结语一张不断变长的跑通清单Booth 的硬件验证逻辑值得借鉴真实硅片MI300X、RTX 4060 Ti验证端到端正确性模拟器 CI 让普通开发者无需显卡也能拦下分歧 bugSLATEC 基准值保证数值可追溯。它的 docs/roadmap.md 里Apple Metal 的硬件验证、确定性模式--deterministic与误差界输出都在路上——对关心编译器承诺能不能兑现的读者来说这份清单本身比任何宣传语都更有说服力。赞分享【免费下载链接】BoothOpen-source CUDA, Triton and HIP compiler targeting multiple GPU and CPU architectures.项目地址https://gitcode.com/gh_mirrors/bar/Booth点击查看免费下载相关推荐4步出片RTX 4060就能跑的AI视频生成模型来了4步出片RTX 4060就能跑的AI视频生成模型来了 导语 Wan2.1 I2V 14B 480P StepDistill CfgDistill Lightx基础模型计算机视觉模型推理服务模型量化4步出片RTX 4060就能跑的AI视频生成模型来了4步出片RTX 4060就能跑的AI视频生成模型来了 导语 Wan2.1 I2V 14B 480P StepDistill CfgDistill Lightx基础模型计算机视觉模型推理服务模型量化4步出片RTX 4060就能跑的AI视频生成模型来了4步出片RTX 4060就能跑的AI视频生成模型来了 ! LightX2V框架标志 https://gitcode ai aigc.obs.cn north基础模型计算机视觉模型推理服务模型量化创作声明:本文部分内容由AI辅助生成(AIGC),仅供参考