Test the device heap

Cover allocations that reach the heap (a `Ref` handed to a `@noinline`
callee, a boxed `Any` field), independence of the per-work-item heaps,
objects staying valid across later allocations, the reported out-of-memory
condition on exhaustion, and that only allocating kernels carry an arena.

Co-Authored-By: Claude Fable 5.1 <noreply@anthropic.com>
Claude-Session: https://claude.ai/code/session_01Fy9hCZkRFuf9fbRQNagCxf

Author: maleadt

test/execution.jl

@@ -742,3 +742,88 @@ end
     end for _ in 1:2])
     @test all(results)
 end
+
+############################################################################################
+
+# Device-side allocation: Julia objects that survive optimization are served by a
+# per-work-item heap in private memory (src/device/runtime.jl). Top-level definitions, as a
+# closure capturing test state would not be a bitstype.
+
+@noinline heap_consume(r::Base.RefValue{Float32}) = r[] + 1f0
+
+struct HeapAnyBox
+    x::Any
+end
+@noinline heap_consume(b::HeapAnyBox) = (b.x::Float32) * 2f0
+
+@testset "device heap" begin
+    # a Ref handed to a @noinline callee is heap-allocated
+    function ref_kernel(a)
+        i = get_global_id()
+        @inbounds a[i] = heap_consume(Ref(a[i]))
+        return
+    end
+    a = oneArray(Float32[41])
+    @oneapi ref_kernel(a)
+    @test Array(a) == [42]
+
+    # so is a struct whose `Any` field boxes its value
+    function anybox_kernel(a)
+        i = get_global_id()
+        @inbounds a[i] = heap_consume(HeapAnyBox(a[i]))
+        return
+    end
+    a = oneArray(Float32[21])
+    @oneapi anybox_kernel(a)
+    @test Array(a) == [42]
+
+    # every work-item has its own heap
+    n = 4096
+    a = oneArray(Float32.(1:n))
+    @oneapi items=256 groups=n÷256 ref_kernel(a)
+    @test Array(a) == Float32.(2:n+1)
+
+    # objects stay valid across later allocations by the same work-item
+    function select_kernel(a, idx)
+        i = get_global_id()
+        refs = ntuple(j -> Ref(a[i] * j), Val(4))
+        @inbounds a[i] = heap_consume(refs[idx])
+        return
+    end
+    a = oneArray(Float32[1, 2, 3, 4])
+    @oneapi items=4 select_kernel(a, 3)
+    @test Array(a) == Float32[4, 7, 10, 13]
+
+    # nothing is freed: a work-item that allocates more than the heap holds runs out of
+    # memory, which is reported, and exits without writing its result
+    function loop_kernel(a, n)
+        i = get_global_id()
+        @inbounds x = a[i]
+        for _ in 1:n
+            x = heap_consume(Ref(x))
+        end
+        @inbounds a[i] = x
+        return
+    end
+    fits = oneAPI.HEAP_SIZE ÷ oneAPI.HEAP_ALIGNMENT
+    a = oneArray(Float32.(1:256))
+    @oneapi items=256 loop_kernel(a, fits)
+    @test Array(a) == Float32.(1:256) .+ fits
+    a = oneArray(Float32[1])
+    _, out = @grab_output begin
+        @oneapi loop_kernel(a, fits + 1)
+        synchronize()
+    end
+    @test occursin("Out of dynamic GPU memory", out)
+    @test Array(a) == [1]
+
+    # only kernels that allocate carry a heap
+    function plain_kernel(a)
+        i = get_global_id()
+        @inbounds a[i] += 1f0
+        return
+    end
+    T = Tuple{oneDeviceVector{Float32, oneAPI.AS.CrossWorkgroup}}
+    @test occursin(r"%heap\d* = alloca", sprint(io -> oneAPI.code_llvm(io, ref_kernel, T; kernel=true)))
+    @test !occursin("alloca", sprint(io -> oneAPI.code_llvm(io, plain_kernel, T; kernel=true)))
+end