[BUG] Vectors do not change even though writing to them in kernel

[JonathanKuelz]

Bug Description

So I don't know if I'm not seeing an obvious mistake I made or if this is a bug, but I can't wrap my hand around this. I wrote this very simple kernel to get the aabb of a point cloud:

@wp.kernel
def get_aabb(points: wp.array(dtype=wp.vec3), lower: wp.vec3, upper: wp.vec3):

    tid = wp.tid()

    for i in range(points.shape[0]):
        lower[tid] = wp.min(lower[tid], points[i][tid])
        upper[tid] = wp.max(upper[tid], points[i][tid])
    print(lower[tid])
    points[0][tid] = 10.

It computes the axis-wise minimum of the points array. As you can see, there are two print statements in there for debugging. Also, setting the first points to [10, 10, 10] is for debugging purposes only.

I call it like this:

wp.init()
aabb_lower = wp.vec3(wp.INF, wp.INF, wp.INF)
aabb_upper = wp.vec3(-wp.INF, -wp.INF, -wp.INF)
points = wp.array([[-1., 1., -1.], [1., -2., 0.]], dtype=wp.vec3)
wp.launch(get_mesh_aabb, dim=3, inputs=[points], outputs=[aabb_lower, aabb_upper])
print(aabb_lower)
print(aabb_upper)
print(points)

I would expect aabb_lower and aabb_upper to contain minimums/maximums of the points, but instead I get the following:

[inf, inf, inf]
[-inf, -inf, -inf]
[[10. 10. 10.]
 [ 1. -2.  0.]]
-1
-2
-1

So inside the kernel, aabb_lower is computed correctly. However, this doesn't "propagate" to the outside of the kernel, my variable remains unchanged. Note that this is not the case for the points, where the print outside the kernel reflects the changes, so this seems to happen only for vectors not wrapped in arrays. Is there a logical explanation for this behavior?

System Information

No response

[shi-eric]

Things like vectors and scalars are passed by value, so you'll need to wrap those vectors in an array if you want to "persist" the state of those variables after running the kernel.