Raising: read constant globals as tensors

[wsmoses]

mfem's VectorDiffusionIntegrator::AssemblePA keeps a kernel-local const int map[9] = {0,3,6,1,4,7,2,5,8} that the optimizer promotes to a host .rodata global; the raised kernel captured its address as a buffer argument, which compiled fine and failed only at runtime ("pointer does not belong to any reactant allocation" — found via gdb info symbol on the failing pointer: __const.<AssemblePA lambda>.map).

A constant global with a dense initializer is a tensor the module already holds, so the raiser now reads it as one instead of rewriting the loads:

• llvm.mlir.addressof of such a global raises to a stablehlo.constant of the (flattened) initializer. The pointer2memref view and the loads through it then go through the existing buffer raising unchanged (slices for constant indices, gathers for runtime ones).
• A wrapper operand rooted at such a global (the view taken outside the gpu_wrapper, or the raw pointer with the view inside, which is mfem's shape) is cloned into the region, like the index casts already are, instead of becoming a kernel operand; the global being constant, nothing in the region can have written through it.
• A pointer2memref of a static tensor to a memref<?x...> view now sizes the view statically instead of reading the dimension back at runtime; without it the constant-index load raises through the dynamic-extent path (get_dimension_size + dynamic_pad + slice of a tensor<?xi32>) instead of a plain slice.

Part of #2968.

🤖 Generated with Claude Code

https://claude.ai/code/session_016zErYp7upmqr4NHfhod9UD