Vectorized global/shared load and store words (v2/v4)

[tetsuo-cpp]

Summary

Add words for vectorized memory access (e.g. 2- and 4-element loads/stores) that lower to wide PTX loads/stores, improving bandwidth without requiring full tensor-core MMA.

Motivation

Many bandwidth-bound kernels (copies, projections, epilogues) benefit from ld.global.v4 / st.global.v4-style accesses. WarpForth currently exposes only scalar @/! and narrow scalar variants. Vector ops are a practical step between scalar memory and MMA (#11).

Proposed words (sketch)

Word	Stack effect	Description
V2@	( addr -- x y )	Load 2×i64 (or typed pair)
V4@	( addr -- a b c d )	Load 4×i64
V2!	( x y addr -- )	Store 2×i64
V4!	( a b c d addr -- )	Store 4×i64
FV2@ / FV4@	float variants	f32×2/×4 or f64×2 as appropriate
SV2@ …	shared-memory variants	address space 3

Alignment requirements must be documented (natural alignment for vector width).

Lowering

• Prefer LLVM/NVVM vector load/store or memref vector types
• Verify PTX contains vector forms on a simple copy kernel

Acceptance criteria

• At least V2/V4 for global i64 or f32 path implemented end-to-end
• LIT checks on lowered IR/PTX shape
• GPU test: vector copy or saxpy-like kernel matches scalar reference
• Language reference: alignment + stack effects
• Shared variants or explicit non-goal documented

Related

• Reduced-width memory (Add load/store words for reduced-width memory types #52 closed)
• MMA (Tensor core / MMA intrinsics #11) — complementary, not a substitute
• Element-typed params (cleaner addresses for vectors)