[BUG] Partially compiled kernel appears in module code

Bug Description

When running test_mat_constructors.py (to be added to main, split off from test_mat.py), we see a strange recompilation:

wp___main___1f923c6: Why is a GPU module being compiled with #define WP_TILE_BLOCK_DIM 1?

wp___main___72d9243: Why are test_anon_constructor_error_type_mismatch__locals__kernel_a00b1546_cuda_kernel_forward and test_anon_constructor_error_type_mismatch__locals__kernel_a00b1546_cuda_kernel_backward ending up in the generated module? The kernel shouldn't have compiled:

    @wp.kernel
    def kernel():
        wp.matrix(1.0, shape=(3, 2), dtype=wp.float16)

But part of it ended up in the module:

extern "C" __global__ void test_anon_constructor_error_type_mismatch__locals__kernel_a00b1546_cuda_kernel_forward(
    wp::launch_bounds_t dim)
{
    for (size_t _idx = static_cast<size_t>(blockDim.x) * static_cast<size_t>(blockIdx.x) + static_cast<size_t>(threadIdx.x);
         _idx < dim.size;
         _idx += static_cast<size_t>(blockDim.x) * static_cast<size_t>(gridDim.x))
    {
        // reset shared memory allocator
        wp::tile_alloc_shared(0, true);

        //---------
        // primal vars
        const wp::float32 var_0 = 1.0;
        const wp::int32 var_1 = 3;
        const wp::int32 var_2 = 2;
        wp::tuple_t<wp::int32, wp::int32> var_3;
        //---------
        // forward
        // def kernel():                                                                          <L 68>
        // wp.matrix(1.0, shape=(3, 2), dtype=wp.float16)                                         <L 69>
        var_3 = wp::tuple(var_1, var_2);
    }
}



extern "C" __global__ void test_anon_constructor_error_type_mismatch__locals__kernel_a00b1546_cuda_kernel_backward(
    wp::launch_bounds_t dim)
{
    for (size_t _idx = static_cast<size_t>(blockDim.x) * static_cast<size_t>(blockIdx.x) + static_cast<size_t>(threadIdx.x);
         _idx < dim.size;
         _idx += static_cast<size_t>(blockDim.x) * static_cast<size_t>(gridDim.x))
    {
        // reset shared memory allocator
        wp::tile_alloc_shared(0, true);

        //---------
        // primal vars
        const wp::float32 var_0 = 1.0;
        const wp::int32 var_1 = 3;
        const wp::int32 var_2 = 2;
        wp::tuple_t<wp::int32, wp::int32> var_3;
        //---------
        // dual vars
        wp::float32 adj_0 = {};
        wp::int32 adj_1 = {};
        wp::int32 adj_2 = {};
        wp::tuple_t<wp::int32, wp::int32> adj_3 = {};
        //---------
        // forward
        // def kernel():                                                                          <L 68>
        // wp.matrix(1.0, shape=(3, 2), dtype=wp.float16)                                         <L 69>
        var_3 = wp::tuple(var_1, var_2);
        //---------
        // reverse
        // adj: wp.matrix(1.0, shape=(3, 2), dtype=wp.float16)                                    <L 69>
        // adj: def kernel():                                                                     <L 68>
        continue;
    }
}

System Information

No response

Start with the planned test_mat_constructors.py, comparing it with test_mat.py and the generated modules named in the report. Reproduce the matrix constructor type-mismatch case and inspect why code from the failed kernel appears in the generated forward and backward kernels. Done means the invalid kernel no longer contributes partially compiled code or unexpected modules.