Slice has no HIR→TIR lowering, so any tensor subscript fails to compile
Affects: main @ e55e4b0
Summary
hir.tensor.Slice has an Op, a typeinfer, an evaluator and an access
relation, but no HIR→TIR lowering. Any kernel containing a tensor subscript
therefore type-checks and evaluates, then dies at lowering with a bare
TypeError.
parser §1.7 documents x[slice0, …] as lifting to a
hir.tensor.Slice call, so this is a documented surface that cannot reach a
compiled kernel.
Reproduction
A static window over a kernel param — no loop, no dynamic bound:
ROWS, K, BK = 64, 512, 256
@func(topologies=(Topology("cta", ROWS),))
def win(a: Tensor[(ROWS, K), "f32"]) -> Tensor[(ROWS, BK), "f32"]:
with Mesh(topology="cta", layout=Layout(shape=(ROWS,), strides=(1,))) as cta:
w = a[:, 0:BK]
r = tf.reshard(w, layout=(ROWS @ cta, BK), storage=rmem)
return tf.reshard(tf.mul(r, r), layout=(ROWS @ cta, BK), storage=gmem)
tilefoundry.compile(win, target="cuda")
Actual
TypeError: hir_to_tir: no lowering registered for Op Slice
src/tilefoundry/passes/transforms/hir_to_tir.py:455
Expected
Either the subscript compiles, or the failure names Slice as an unimplemented
lowering rather than surfacing as a generic registry miss.
Where it sits in the capability matrix
Slice is not alone in being evaluator-only, but it is the one a documented
syntactic form lifts to, so it is reachable without the author opting into
anything. The narrow fix — a diagnostic that says "Slice lowering is not
implemented" — would already be an improvement over the current message.
A trap for whoever implements it
I tried this and got it wrong, so it is worth recording. Lowering the window
with cute::local_tile looks right and is not: local_tile records the tile
offset inside the composed layout, but a ShardTensor built over that window is
addressed as engine.data() + shard_offset (shard_tensor.cuh, local_impl),
which ignores the composed offset entirely. The result is correct at offset 0
and silently wrong everywhere else — in my attempt the static case above still
came out with 16128 of 16384 elements wrong, all rows past the first.
Advancing the pointer explicitly instead of composing the offset into the layout
is the direction that behaves, but I do not have a version I would stand behind,
which is why this is an issue and not a PR.
Related
Two-arg tile() — the other documented way to produce a subscript window —
is separately broken (#73), so that route cannot be used to reach this code
path either.
Slicehas no HIR→TIR lowering, so any tensor subscript fails to compileAffects:
main@e55e4b0Summary
hir.tensor.Slicehas an Op, atypeinfer, an evaluator and an accessrelation, but no HIR→TIR lowering. Any kernel containing a tensor subscript
therefore type-checks and evaluates, then dies at lowering with a bare
TypeError.parser §1.7 documents
x[slice0, …]as lifting to ahir.tensor.Slicecall, so this is a documented surface that cannot reach acompiled kernel.
Reproduction
A static window over a kernel param — no loop, no dynamic bound:
Actual
Expected
Either the subscript compiles, or the failure names
Sliceas an unimplementedlowering rather than surfacing as a generic registry miss.
Where it sits in the capability matrix
Sliceis not alone in being evaluator-only, but it is the one a documentedsyntactic form lifts to, so it is reachable without the author opting into
anything. The narrow fix — a diagnostic that says "
Slicelowering is notimplemented" — would already be an improvement over the current message.
A trap for whoever implements it
I tried this and got it wrong, so it is worth recording. Lowering the window
with
cute::local_tilelooks right and is not:local_tilerecords the tileoffset inside the composed layout, but a
ShardTensorbuilt over that window isaddressed as
engine.data() + shard_offset(shard_tensor.cuh,local_impl),which ignores the composed offset entirely. The result is correct at offset 0
and silently wrong everywhere else — in my attempt the static case above still
came out with 16128 of 16384 elements wrong, all rows past the first.
Advancing the pointer explicitly instead of composing the offset into the layout
is the direction that behaves, but I do not have a version I would stand behind,
which is why this is an issue and not a PR.
Related
Two-arg
tile()— the other documented way to produce a subscript window —is separately broken (#73), so that route cannot be used to reach this code
path either.