The GitHub Actions job "Lint" on tvm.git/main has succeeded. Run started by GitHub user spectrometerHBH (triggered by spectrometerHBH).
Head commit for run: af52b2f015e1b58c0d3dd6325da803540f6ef136 / spectrometerHBH <[email protected]> fix(lower-tirx): keep buffer identity coherent across buffer rebuilds After the typed-buffer-variable migration, a buffer's identity is the variable itself. A pass that rewrites the program while rebuilding buffers is applying a substitution over buffer identities, and that substitution must reach every expression position — including the ones tucked inside buffer types and pass-internal folds. Several spots missed this and broke kernels using shared-memory pools, views sized/offset by local scalars, or host-side tensormap encodes of device-declared views. FlattenBuffer is restructured around its invariant instead of patched: each n-d buffer flattens to a 1-d storage husk buf' (same data origin, dtype, alignment and scope; no layout, no elem_offset), and every access buf[x] rewrites to buf'[f(x)] with f(x) = layout.apply(x, shape) + elem_offset. Since the folded indices live in the rewritten program, f's coefficients (runtime elem_offset, symbolic shapes/strides, layout iters) must be rewritten too. The pass now walks the AST top-down and derives, exactly once at each buffer's definition point (AllocBuffer/DeclBuffer/params), the pair {fold view = rewritten geometry, buf'}; use sites only look the pair up. The substitution is applied exactly once per subexpression by construction, and a use before its definition is a hard error at the fault instead of a stale reference that surfaces later in MakePackedAPI. A buffer_data projection is only resolvable where the buffer's definition is visible, and TMA host-init statements violated that: they are emitted inside the kernel region but hoisted to host scope, carrying projections of device-local views. TilePrimitiveDispatch now tracks each buffer's storage root as it walks definitions and rewrites those projections onto the root — a PrimFunc parameter, visible on the host — at the moment the statements are hoisted. With host statements self-contained, LowerTIRxCleanup's alias lookup no longer has to tolerate forward references, so a projection with no visible definition is an error there instead of silently resolving to itself. Two more substitution obligations were unmet: - LowerTIRxCleanup rebuilt buffers without rewriting the buffer-type shape/stride fields, so a view sized by a local scalar kept pointing at the pre-rebuild variable. - LowerTIRxOpaque's unit-loop Var visitor returned non-unit-loop variables verbatim, shadowing the base visitor's buffer remapping. The shared-memory pool no longer needs buffer-referencing metadata at all. Its allocation is an extern placeholder whose extent never reaches the generated code, so the dynamic-shared-memory launch size is declared directly instead of being patched into the allocation: - SMEMPool.commit() appends a leaf AttrStmt (node=0, key="tirx.dyn_smem_bytes", IntImm value) to the kernel region; hand-written kernels declare theirs with the existing T.attr({"tirx.dyn_smem_bytes": n}) sugar. - SplitHostDevice reads the declaration as the dynamic-shared-memory launch argument and strips it when finalizing the kernel; allocation extents are no longer consulted, and a shared.dyn allocation without a declaration is a hard error. - LowerTIRxOpaque's pool-size collection and extent patching are removed. Fork tests updated for the typed-buffer surface; test_tcgen05_mma_ss_no_tma builds its smem descriptors with first-class ptr_to() instead of a reinterpreting access_ptr over the raw arena, satisfying tvm_access_ptr's element-type check; and test_transform_flatten_buffer.py pins the FlattenBuffer invariant (rewritten references in view shapes and folded elem_offsets; identity preservation for already-flat buffers); all three of its tests fail on the pre-fix pass. Report URL: https://github.com/apache/tvm/actions/runs/30672675385 With regards, GitHub Actions via GitBox --------------------------------------------------------------------- To unsubscribe, e-mail: [email protected] For additional commands, e-mail: [email protected]
