Register and indexed tile operations produce wrong results for padded vector and matrix arrays #2282

Bug Description

Register and indexed tile loads and stores produce incorrect results for padded vec3 and mat33 arrays on both CPU and CUDA.

A vec3 holds three float32 values, occupying 12 bytes. In the reproducer, consecutive vectors start 16 bytes apart, leaving four bytes of padding after each vector:

[ x0  y0  z0  padding ][ x1  y1  z1  padding ]
|------ 16 bytes -----|------ 16 bytes -----|

The backing array contains:

[ 0, 1,  2,  3 ]
[ 4, 5,  6,  7 ]
[ 8, 9, 10, 11 ]

Each row represents one vec3. The fourth value is padding and should be skipped, giving (0, 1, 2), (4, 5, 6), and (8, 9, 10).

The following operations fail with both element types on both backends:

• wp.tile_load(..., storage="register")
• wp.tile_store() from a register tile
• wp.tile_load_indexed()
• wp.tile_store_indexed()

Ordinary shared tile loads and stores pass the same padded-layout checks.

Reproducer

import numpy as np
import warp as wp

wp.init()
device = "cuda:0"  # Also reproduces with "cpu".

@wp.kernel(enable_backward=False)
def copy_register(src: wp.array[wp.vec3], dst: wp.array[wp.vec3]):
    t = wp.tile_load(src, shape=3, storage="register")
    wp.tile_store(dst, t)

padded = np.arange(12, dtype=np.float32).reshape(3, 4)
backing = wp.array(padded, dtype=wp.float32, device=device)
src = wp.array(
    ptr=backing.ptr,
    shape=(3,),
    strides=(backing.strides[0],),
    dtype=wp.vec3,
    device=device,
    copy=False,
)
dst = wp.empty(3, dtype=wp.vec3, device=device)
wp.launch_tiled(copy_register, dim=1, inputs=[src, dst], block_dim=32, device=device)
print(dst.numpy())
np.testing.assert_array_equal(dst.numpy(), padded[:, :3])

Expected:

[[ 0.  1.  2.]
 [ 4.  5.  6.]
 [ 8.  9. 10.]]

Actual:

[[0. 1. 2.]
 [3. 4. 5.]
 [6. 7. 8.]]

The assertion fails with six of nine scalar components mismatched.

Application / Workflow Impact

Isolated testing during verification of #2268 exposed this problem. No application failure has been demonstrated, but tiled kernels reading or writing padded vector and matrix arrays can silently produce incorrect results.

The CPU failures also occur on the parent of the tested commit, so they predate #2268’s shared-load optimization.

Warp Version

1.20.0.dev0, built from ae833a3.

System Information

• Linux x86_64, Python 3.12.13.
• CPU: AMD EPYC 9B45.
• GPU: NVIDIA RTX PRO 6000 Blackwell Server Edition, MIG 1g.24gb.
• CUDA Toolkit 13.4, NVIDIA driver 595.58.03.
• CUDA checks use GPU-resident buffers and block_dim=32.