[vchuravy]

This is a prototype for discussion; it is not ready to merge.

What

StaticDimsArray{T,N,D,A} <: AbstractGPUArray{T,N} views a GPU vector parent::A with the dimensions dims::D. D is a tuple of Ints and Static.jl StaticInts:

using GPUArrays.Static: static
u = StaticDimsArray(CuArray{Float32}(undef, 5 * 4^3 * nelem), (static(5), static(4), static(4), static(4), nelem))

Trixi.jl stores its solution this way: the number of variables and nodes is fixed, and only the number of elements changes at runtime. The mixed-tuple design follows StrideArraysCore. The static entries have zero size, so the device array stays isbits (40 bytes over a CuDeviceVector). In kernels, linear indices, strides and bounds checks for the static dimensions become constants. Unlike FixedSizeArrays.jl#210, which makes all dimensions but the last static, any dimension can be static or dynamic here.

• Host type: StaticDimsArray forwards the backend, broadcast style, storage/memory_location (alias detection), derive, and CPU copies to the parent. All the generic GPUArrays functionality then works unchanged: broadcasting, AK reductions, fill!, views, scalar indexing. similar keeps the static dims when the size does not change.
• Device type: adapt turns the array into a StaticDimsDeviceArray. Its adapt_structure adapts the parent, so KernelAbstractions' @Const reaches the backend's read-only loads (ld.global.nc on CUDA).
• StaticArrayInterface extension: static_size and known_size.
• JLArrays: a StaticDimsArray over a JLArray is reduced on the host, like a JLArray.

Performance (CUDA, Quadro RTX 4000)

Tested against CUDA.jl main (02ddd063e). The kernel is a Trixi-style DGSEM volume term on u[v, i, j, k, e], with 32768 elements at N=4 and 16384 at N=5. Times are ncu measurements at locked base clocks.

	CuArray	StaticDimsArray, static	StaticDimsArray, all Int
N=4, @cuda, node index decomposed from the array size	583 µs	460 µs	584 µs
N=4, KernelAbstractions with @Const	498 µs	347 µs	499 µs
N=5, @cuda	706 µs	610 µs	707 µs
N=5, KernelAbstractions with @Const	624 µs	453 µs	624 µs
warp instructions (N=4, @cuda)	90.2M	17.3M	90.2M

With all-Int dims, the kernels have the same instruction counts as with a CuArray.

DRAM traffic is the same in every case (83 MB). The gain comes from removing 64-bit integer divisions and their exception paths from the CartesianIndices decomposition, and from fully unrolling the inner loop. The static kernels are now limited by the load/store queue (lg_throttle). Sustained runs show the same ordering, with 10–30% less energy per launch.

Costs

• New dependencies: Static.jl and, through it, IfElse, SciMLPublic and CommonWorldInvalidations.
• using GPUArrays takes about 30 ms longer, and loading it invalidates about 190 more method instances (2116 → 2306). Static.jl alone invalidates 836.
• Precompiling GPUArrays and time to first broadcast or reduction are unchanged.

Open questions / known issues

• Static.jl dependency: keep it, or define a minimal StaticInt in GPUArraysCore and convert in the StaticArrayInterface extension.
• All-Int overhead: fixed in 165710d. The Horner-form linear index i1 + d1*((i2-1) + d2*(i3-1)) nested the loop index in every multiplication, so LLVM could not hoist the loop-invariant terms. A generated sum of strided terms, as in Base, makes all-Int dims match CuArray. Static strides still fold to constants.
• Static dims are dropped by out-of-place broadcasts, reductions, views and reshapes, which return backend arrays.
• Unrelated finding: GPUArrays always uses the Cartesian broadcast kernel for N-D broadcasts, because Base only reports IndexLinear for 1-D Broadcasteds. For same-shape 5-D arrays that reaches about 160 GB/s, versus 379 GB/s for the same broadcast over vecs. That belongs in a separate PR.
• CI coverage: the testsuite covers JLArrays; CUDA was only tested locally.

🤖 Generated with Claude Code

[vchuravy]

Have to think about this.

[vchuravy]

Two ways forward.

1. I am going to try and use the FixedSizedArray with static dims only at the kernel callsite. The static dimensions only matter there.
2. Using optionally static dims, like Add static dimensions to FixedSizeArray via FixedSizeArrayDims JuliaArrays/FixedSizeArrays.jl#213 might be also a way forward, but would potentially require a GPUArrays v13