SPIR-V on ROCm: A Portable IR for AMD GPUs
12 hours ago
- ROCm adopts SPIR-V as a portable intermediate representation for AMD GPUs, enabling compile-once, specialize-on-device model to replace ahead-of-time per-target compilation.
- SPIR-V addresses scalability issues: eliminates long build times, large fat binaries, and lack of forward compatibility associated with AOT-per-target approach, allowing single binary to run on future GPU architectures.
- Key feature is late-resolved feature predicates (ZCFS), which enable per-arch dispatch without runtime overhead via builtins like __builtin_amdgcn_processor_is and __builtin_amdgcn_is_invocable.
- Benchmark shows SPIR-V provides flat build time and binary size regardless of target count, while fat binaries scale linearly; first-launch JIT latency (~70-100 ms) is mitigated by caching.
- Current status includes shipping amdgcnspirv target and JIT caching; in-flight developments include in-tree LLVM SPIR-V backend and JIT-at-package-install; library coverage remains a limitation.