PTX carriers up the ML stack on RTX (FFN/attention/block/llama) + kernel-graph runtime + Android SPIR-V

NameDescription

Extend the proven form-ptx lane (matvec f32/f16/bf16 + affine-train, four-way verdict 63, bit-exact on RTX via driver JIT -O0) up the layer stack to reach Metal parity and enable full models on all three backends. Three tiers: (1) RTX/PTX carriers for the layers Metal already has but PTX lacks — FFN/MLP (fwd+backprop, single CUDA block + bar.sync between phases, gelu composed from the recipe's 14-term Taylor fexp/ftanh NOT ex2.approx so it stays bit-exact to tn-gelu), attention (scores+softmax+weighted-sum, serial softmax for bit-exact), transformer block-fwd, llama-block (fwd/causal/decode+KV); plus f16/bf16 across the block-level kernels. (2) Cross-cutting runtime: a kernel-graph/scheduler with persistent device buffers (resident weights + KV-cache) so hundreds of per-layer dispatches run end-to-end — today each audit allocates per-call; plus weight-load (Q4_K/Q6_K dequant already exist) onto device. Decide serial(bit-exact) vs block/warp-reduced(epsilon) per kernel. (3) Android has NO GPU compute backend — hati-os-targets declares gpu:compute as Vulkan+OpenGL-ES and dsp as NNAPI but no SPIR-V/GLSL emitter exists; add a SPIR-V compute emit target (documented IR, same Form-emits-the-ISA pattern as PTX/MSL). Algorithm-layer gaps for diffusion: conv2d/GroupNorm/UNet/VAE (only Conv1D today), flash-attention, sampling (top-k/top-p/temperature/beam), cross-entropy, generic fp8/fp4. Starting point this session: FFN/MLP forward PTX, bit-exact, with the recipe's Taylor gelu as a reusable PTX device function.

Potential value (USD)Estimated cost (USD)Confidence (0–1)