AMD backend doesn't lower string literals yet. E110 refusal fires on any device-side "..." until this lands.
Best guess iiiiis: ELF emitter grows a .rodata section, the bytes go there, isel_global_ref for byte-init globals does a PC-relative address load instead of the hidden-kernarg dance it does for __device__ globals. Helper to check: bir_global_is_bytes(M, gi).
Hardest of the three Phase 2 backends, because of the section bookkeeping. Soft-float runtime doesn't need this, so no pressure on the Tenstorrent side from it.
AMD backend doesn't lower string literals yet. E110 refusal fires on any device-side
"..."until this lands.Best guess iiiiis: ELF emitter grows a
.rodatasection, the bytes go there,isel_global_reffor byte-init globals does a PC-relative address load instead of the hidden-kernarg dance it does for__device__globals. Helper to check:bir_global_is_bytes(M, gi).Hardest of the three Phase 2 backends, because of the section bookkeeping. Soft-float runtime doesn't need this, so no pressure on the Tenstorrent side from it.