[bug] gfx/Vulkan: f64 field at a 4 mod 8 offset silently aliases the last element of the previous field

Summary

On the gfx (SPIR-V) backends, root fields are packed with no alignment padding, while buffer accesses derive the element index from a byte offset with a right shift. The shift truncates, so a place SNode whose offset is not a multiple of its primitive size is read and written one slot too low, silently overlapping the last element(s) of the field placed before it.

Reads and writes of the misplaced field agree, so that field looks perfectly correct — the corruption shows up in its neighbour, which makes this very hard to attribute.

Repro (Vulkan)

import taichi as ti

ti.init(arch=ti.vulkan)

N = 1024
s0 = ti.field(ti.u32, shape=())    # 4 B
s1 = ti.field(ti.i32, shape=())    # 4 B
s2 = ti.field(ti.u32, shape=())    # 4 B  -> 12 B of scalars
ints = ti.field(ti.u32, shape=N)   # 4096 B at offset 12, ends at 4108
dbl = ti.field(ti.f64, shape=N)    # placed at 4108, and 4108 % 8 == 4


@ti.kernel
def write_dbl():
    for i in dbl:
        dbl[i] = 1.0


ints.fill(0x0BADF00D)
write_dbl()
print(hex(int(ints.to_numpy()[-1])))   # expected 0xbadf00d
print(dbl.to_numpy()[0])               # 1.0 — the f64 field itself is fine

Expected 0xbadf00d, actual on Vulkan 0x0 — the low 32 bits of f64(1.0). Nothing was written to ints between the two reads. CPU and CUDA are unaffected (the LLVM path builds the layout as a struct type and addresses members by pointer arithmetic).

A variant that needs no data64, for completeness: three 4-byte scalars, then ti.field(ti.u8, shape=1026), then ti.field(ti.u32, shape=512). The u32 field lands at an offset that is 2 (mod 4), and after writing ints[0] = 0x0BADF00D the last two u8 elements read back as 0x0d and 0xf0.

Cause

Two places, both on the gfx path:

1. taichi/codegen/spirv/snode_struct_compiler.cpp, compute_snode_size: children of a container are sorted by ascending size and packed back to back (cell_stride += snode_size). Nothing rounds a child's offset to the alignment of its element type.
2. taichi/codegen/spirv/spirv_codegen.cpp, at_buffer: the root buffer is viewed as an array of the accessed primitive and the element index is computed as byte_offset >> log2(width).

Together these require every place SNode's absolute byte offset to be a multiple of its primitive size, and nothing enforces or checks that.

Why it is easy to hit in practice

Because field sizes determine the order, adding or removing a single 4-byte scalar field shifts every subsequent offset by 4 bytes and can misalign every f64 field in the program.

We hit this in a Monte-Carlo kernel: two variants of the same program differed only in that one declared one extra u32 scalar field. In that variant the last element of a u32 array shared bytes with the low half of the first element of an f64 array, so it read back as the low mantissa word of a double — exactly 0 whenever that double happened to be 0.0, and a large arbitrary integer otherwise. The array held per-path time indices, so the 0 was a valid-looking value and quietly biased results in about half of all runs. Neither fill(), nor from_numpy(), nor kernel writes could keep a correct value in that slot.

Environment

• taichi 1.8.0 (local build), Windows 10, AMD Radeon R7 450 (driver 24.20.11020.3002), Vulkan SDK 1.3.296
• Both code paths above are unchanged on master (ba0e81d) and in v1.7.4, so this is not specific to a release.

Fix

Track the required alignment per SNode (natural alignment of the primitive for a place SNode, strictest child alignment for a container), align each child offset to it, and round the cell stride up so every cell of a container stays aligned. Padding is at most alignment - 1 bytes per field, and layouts that were already aligned come out byte-identical.