Named kernels — `gpu fn` and `.launch`

A forall is an anonymous kernel with inferred launch geometry. When you need a reusable kernel, or explicit control over the grid and block shape (for shared-memory tiling or warp operations), declare a gpu fn and launch it yourself:

gpu fn scale(src Array<f32, 1024>, dst out Array<f32, 1024>)
    let i = kernel.global_idx.x
    if i < src.length()
        dst[i] = src[i] * 2.0

fn main()
    gpu let src = Array<f32, 1024>()
    gpu var dst = Array<f32, 1024>()
    scale(src, dst).launch(Dim3(4, 1, 1), Dim3(256, 1, 1))   // grid, block

Buffer parameters must be gpu-resident at the launch site (gpu let/gpu var). An out parameter is writable, everything else is read-only. .launch(grid, block) takes two Dim3s: the number of workgroups per axis and the threads per workgroup. The block shape must be a compile-time literal — it is stamped into the generated shader as @workgroup_size. Dispatch is synchronous.

Named kernels — gpu fn and .launch

Named kernels — `gpu fn` and `.launch`