Zaneham/Booth

Tensix backend doesn't lower string literals yet. E110 refusal fires on any device-side "..." until this lands. Two paths to do, in increasing pain order: Metalium (--tensix): emit each byte-init global as static const char .strN[len] = { ... }; at file scope in the generated compute C++, lower BIR_GLOBAL_REF to a reference. Should be a Sunday-afternoon job. Native RV32IM (--rv-elf): same shape as AMD #93, the bytes go into the ELF read-only segment and BIR_GLOBAL_REF becomes a lui+addi of the symbol's PC-relative address. More involved. bir_global_is_bytes(M, gi) tells you which globals to materialise; bytes are at M->strings[c->d.bytes.off].

PTX backend doesn't lower string literals yet. E110 refusal fires on any device-side "..." until this lands. Easiest of the three. For each byte-init global, emit .const .align 1 .b8 .strN[len] = { 'h', 'e', ... }; at module scope, and lower BIR_GLOBAL_REF to mov.u64 %rd, .strN. The existing PTX BIR_GLOBAL_REF handler is a no-op stub today, so it's basically a new code path rather than touching anything. Bytes live at M->strings[c->d.bytes.off], length c->d.bytes.len. bir_global_is_bytes(M, gi) tells you which globals.

Triton exposes a family of atomic operations (tl.atomic_add, tl.atomic_xchg, tl.atomic_cas, tl.atomic_min, tl.atomic_max, tl.atomic_and, tl.atomic_or, tl.atomic_xor) used in reductions, histograms, and sparse update patterns. None of them are in the sema intrinsic table yet, and the lowerer has no handler. The BIR side already covers atomics through BIR_ATOMIC_ADD and friends (subop carries the memory ordering), so this is mostly: add entries to the tl.* intrinsic table in src/triton/sema.c, and add the corresponding switch arms in l_intrinsic_call in src/triton/lower.c that route to the right BIR_ATOMIC_* opcode. Honour the mask= keyword the same way once mask honouring lands (issue tracked separately).

Triton kernels use tuple-unpacking assignment forms like pid_m, pid_n = tl.program_id(0), tl.program_id(1) and the matmul-style a_ptrs += BLOCK_K * stride_ak with multi-component pointer arithmetic. The parser builds the Tuple AST node correctly on either side, but the lowering currently only handles Name targets in Assign and explicitly rejects Tuple targets through s_bind_assign_target. The fix is to walk the Tuple's children and bind each independently; if the RHS is also a Tuple we pair element-wise, otherwise we treat the RHS as iterable and project. For sitting-one purposes the element-wise pairing case is enough.

Sema currently rejects tl.math.exp(x) with the diagnostic unknown intrinsic: tl.math because the submodule chain is not part of the recognised symbol set. The same gap applies to tl.extra.libdevice.* and any other nested namespace inside the tl module. Two-part fix in src/triton/sema.c: (a) when the Attr resolver sees a Name resolved to TN_SYM_MODULE with module id TN_MOD_TL, check whether the attribute name is one of the known submodules (math, extra, etc.) and annotate the Attr as a sub-module rather than as an unknown intrinsic; (b) extend the intrinsic table so the full math.exp / math.log / extra.libdevice.fast_expf etc. style names resolve correctly when they chain off a sub-module. Cheap fix to a real gap; affects any kernel that wants the higher-precision math routines or the libdevice ones.