PRAXIS

03 // PRAXIS
AUTONOMOUS GPU
KERNEL OPTIMIZER
NO AGENT NEEDED
MINIMAL GPU TIME
AGENT-AUTHORABLE
PER-CHIP BOUNDS
PRECISION RELAXATION
CONTRACT VERIFIED
COPYRIGHT ©2026
ASKUS GMBH
VADUZ, LIECHTENSTEIN

The same CUDA kernel hits a different physical wall on every GPU. Memory hierarchy, shared SRAM, and warp geometry shift across chip generations, while Praxis re-specializes execution for the exact silicon you run.

TARGET SILICON LINEUP
ARCH=sm_90a // HOPPER HBM3e

Decouple bulk memory transfers from asynchronous warpgroup math.

On Hopper H200, thread-by-thread global loads starve tensor pipelines. Praxis isolates memory movement onto a single DMA producer warp driving hardware Tensor Memory Accelerator descriptors while consumer warpgroups execute asynchronous matrix instructions continuously.

Memory Movement
HBM3e via TMA bulk async
Warp Schedule
1 DMA + 2 WGMMA warpgroups
On-Chip Workspace
228 KB SRAM + cluster DSMEM
Blocker Bypassed
Register & barrier stalls
DECLARED SOURCE // PORTABLE POLICY KERNEL AUTHORED ONCE 
template <class Policy>
__global__ void step_field(praxis::View<const float> x,
                           praxis::View<const float> w,
                           praxis::View<float> y) {
  auto tile = Policy::Memory::stage(x, w);
  auto acc  = Policy::Compute::contract(tile);
  Policy::Publish::commit_verified(y, acc);
}
LOWERED SILICON PLAN // sm_90a PTX & WARPGROUP SCHEDULE MINIMAL GPU FOOTPRINT 
// Producer warp 0..3: Hardware TMA global-to-cluster transfer
cp.async.bulk.tensor.2d.shared::cluster.global [smem_x], [tma_desc_x];
mbarrier.expect_tx.relaxed.cluster.shared      [bar_ready], 32768;

// Consumer warpgroups 4..11: Asynchronous BF16 -> FP32 matrix pipe
wgmma.mma_async.sync.aligned.m64n128k16.f32.bf16.bf16 acc, desc_a, desc_b;
wgmma.commit_group.sync.aligned;
st.global.wb.v4.f32 [y_ptr], {acc0, acc1, acc2, acc3};
ON-CHIP EXECUTION RESIDENCY ALLOCATION HOPPER SM_90A PARTITION
REGISTERS 30% SHARED / DSMEM 44% WGMMA TENSOR PIPE 26%

Specialist kernel rewrites, found autonomously. Bypassing structural stalls and tuning stage-level precision without manual kernel engineering.

Standard Softmax executes three full-width FP32 passes separated by block barriers and off-chip memory traffic. Praxis keeps the row-maximum stability anchor in FP32 while relaxing internal exponent and normalization transport to FP16 in one pass.

UNOPTIMIZED CUDA · MULTI-PASS FP32 BASELINE 3 Kernel Passes · Global Memory & Barrier Stalls
PRAXIS PASS · 0 EXTERNAL AGENT CALLS Single On-Chip Stream · Safe Intermediate Bit-Width Relaxation
NUMERICAL FIDELITY Verified Against Reference Baseline · PASS
FP32 ANCHOR + FP16 PATH
SIX DETERMINISTIC GATES PTX Adaptation · Memory Sanitizer · Run Determinism · Input Sensitivity · Non-Degeneracy · Reference Equivalence

Standalone by default. Native to CUDA-writing agents. Optimize raw kernels with minimal GPU time, or let coding agents author portable Praxis policies from day one.

01 // STANDALONE ENGINE

Hand existing CUDA directly to Praxis.

Submit raw .cu, .cuh, or .ptx source without wrapping it in an agent loop. With minimal GPU profiling, Praxis maps register pressure, memory round-trips, and precision slack on its own, returning a hardware-verified rewrite tuned for your target chip. 

$ praxis specialize causal_softmax.cu --arch sm_90a
[gates] sanitizer · determinism · sensitivity · reference: PASS
[emit]  sm_90a PTX specialized (0 external agent calls)
02 // AGENT-AUTHORABLE POLICIES

Give CUDA-fluent agents portable primitives.

Any coding agent that knows CUDA can author kernels against Praxis headers, policies, and traits. Your agent declares layouts, ownership, and numerical obligations once; Praxis explores the legal policy lattice to specialize execution for every GPU lineup you deploy. 

#include <praxis/policies.cuh>
#include <praxis/contracts.cuh>
using Spec = praxis::Program<FieldKernel, praxis::LegalPolicies>;
VALIDATED GPU WORKLOAD FAMILIES
01

Particle Analysis & Dynamics

Eliminates global neighbor-list round-trips and warp divergence across uneven particle densities by sizing shared-SRAM interaction halos and register force accumulators to each GPU's SM bounds.

SPATIAL HALO TILING · WARP-UNIFORM TRAVERSAL · REGISTER ACCUMULATION
02

Solar Fluid & Magnetohydrodynamics

Fuses multi-pass solar convection and plasma-advection stencils into persistent register wavefronts, bypassing global synchronization barriers while preserving strict conservation invariants across time steps.

PERSISTENT STENCIL WAVEFRONTS · BARRIER ELISION · CONSERVATION CONTRACT
03

Water Rendering & Fresnel Transport

Folds spectral water surface evolution, dielectric Fresnel reflectance-refraction terms, and separable post-processing into single-launch tile passes, relaxing intermediate shading bitrates within pixel tolerances.

UNIFIED TILE COMPOSITING · FRESNEL CONTRACTION · MIXED-BITRATE SHADING
04

LLM Inference & Muse Challenge

Re-partitions paged attention decode, expert routing, and quantized matrix projections across Hopper H200 warpgroups and workstation GPUs while preserving exact logit and token-distribution obligations.

PAGED KV DECODE · ROUTED MOE · HOPPER sm_90a MUSE KERNELS
Next: 04 // Data
Coming Soon 03 // PRAXIS