[BUG] Weird memory-related errors when dealing with nested array of structs

Bug Description

I'm currently experiencing issues where the buffers I've copied to the GPU contain garbage values that weren't initially there before. For context, I was creating a half-edge data structure on the CPU and uploading it to the GPU. Once the copying has finished I've printed out the values with a test kernel (via wp.print) and confirmed it was fine. However, after doing some unrelated stuff, I then found out that some (not all) of the values in the buffers were overriden to zero. To summarize the current code:

class CollisionDetector:
    def __init__(self, model: Model, ...):
        self.model = model

        # the first kernel launch prints out the values in the array normally
        wp.launch(kernel=print_mesh, dim=1024, inputs=[model])

        # unrelated stuff that doesn't modify the model at all
        self.data = make_kindata(model, nenv, rg=requires_grad)

        # the second kernel launch prints out some garbage values in the middle
         wp.launch(kernel=print_mesh, dim=1024, inputs=[model])

        .... // omitted

builder = ModelBuilder()
model = builder.finalize(device)
collision_detector = CollisionDetector(model, ...)

I was suspecting that this may be some kind of memory-related error, and one that especially has to do with dealing with complex nested structures (wp.struct containing wp.array of wp.struct). To roughly describe the current data structure:

@wp.struct
class Model:
    ... // omitted
    half_edge_meshes: wp.array(dtype=HalfEdgeMesh, ndim=1)

@wp.struct
class HalfEdgeMesh:
    num_verts: int
    num_faces: int
    num_halfedges: int
    center_coord: wp.vec3
    vert_coords: wp.array(ndim=1, dtype=wp.vec3)
    face_halfedge: wp.array(ndim=1, dtype=wp.int32) // this buffer seems to be corrupted for no reason...
    ... // omitted

I've tried my best to create a minimal reproducible example, and although the situation might not be exactly the same, I think it's related enough to my current problem. Here's the code:

import warp as wp

wp.config.mode = "debug"
wp.config.verify_cuda = True
wp.config.lineinfo = True
wp.config.enable_backward = False

@wp.struct
class Item:
    count: int
    data: wp.array(ndim=1, dtype=wp.int32)

@wp.struct
class ItemList:
    items: wp.array(dtype=Item, ndim=1)

def make_item(data_count: int) -> ItemList:
    mesh = Item()
    mesh.count = data_count
    mesh.data = wp.array([i for i in range(data_count)], dtype=wp.int32)
    return mesh

@wp.kernel
def print_data(
        item_list: ItemList,
):
    item = item_list.items[0]
    for i in range(item.count):
        wp.printf("i = %d, item.data[i] = %d\n", i, item.data[i])

num_threads = 1

device = wp.get_device()

class ItemListBuilder:
    def make_item_list(self) -> ItemList:
        item_list = ItemList()
        items = []
        items.append(make_item(200))
        items.append(make_item(200))
        item_list.items = wp.array(items, dtype=Item)
        return item_list

builder = ItemListBuilder()
item_list = builder.make_item_list()

wp.launch(
    kernel=print_data,
    dim=num_threads,
    inputs=[item_list]
)

If I run this code directly the values are printed without any problem, but when running this through compute-sanitizer I get the following error:

========= Invalid __global__ read of size 4 bytes
=========     at T1 wp::load<int>(T1 *)+0x56f0 in array.h:795
=========     by thread (0,0,0) in block (0,0,0)
=========     Address 0xed2000200 is out of bounds
=========     and is 385 bytes after the nearest allocation at 0xed2000000 of size 128 bytes
=========         Device Frame: print_data_f21a1d80_cuda_kernel_forward+0x56f0 in item_list.py:29
=========     Saved host backtrace up to driver entry point at kernel launch time
=========         Host Frame: cuda_launch_kernel [0x6f4abe] in warp.so
=========         Host Frame: launch in context.py:5842
=========         Host Frame: <module> in item_list.py:53
=========         Host Frame: _run_code in <frozen runpy>:88
=========         Host Frame: _run_module_as_main in <frozen runpy>:198

The weirder thing is that if I change the code to directly creating the ItemList instead of going through the ItemListBuilder class, then I don't get any errors with compute-sanitizer at all. For a code snippet:

... // same as before
item_list = ItemList()
items = []
items.append(make_item(200))
items.append(make_item(200))
item_list.items = wp.array(items, dtype=Item)

wp.launch(
    kernel=print_data,
    dim=num_threads,
    inputs=[item_list]
)

System Information

• GTX 4090 with driver version 550.120 and CUDA version 12.4
• Python 3.12.3, Warp 1.7.0
• Ubuntu 24.04

Reproduce the nested ItemList case in item_list.py, especially the print_data kernel at line 29 and the launch at line 53, under compute-sanitizer. Compare construction through ItemListBuilder with direct construction; done means the invalid global read is gone and the nested data prints correctly.