cleanups

* STORE doesn't have ranges anymore

* fix

.github/workflows/autogen.yml

@@ -48,7 +48,7 @@ jobs:
         python3 -c "from tinygrad.runtime.autogen import opencl"
         python3 -c "from tinygrad.runtime.autogen import cuda, nvrtc, nvjitlink, nv_570, nv_580, nv"
         python3 -c "from tinygrad.runtime.autogen import comgr_3, hsa, hip, amd_gpu, sqtt, rocprof, amdgpu_kd, amdgpu_drm"
-        python3 -c "from tinygrad.runtime.autogen.am import am, pm4_soc15, pm4_nv, sdma_4_0_0, sdma_5_0_0, sdma_6_0_0, smu_v13_0_0, smu_v13_0_6, smu_v13_0_12, smu_v14_0_2"
+        python3 -c "from tinygrad.runtime.autogen.am import am, pm4_soc15, pm4_nv, sdma_4_0_0, sdma_5_0_0, sdma_6_0_0, smu_v13_0_0, smu_v13_0_6, smu_v13_0_12, smu_v14_0_2, fw, navi_offsets, vega_offsets"
         python3 -c "from tinygrad.runtime.autogen import libc, kfd, io_uring, ib, pci, vfio"
         python3 -c "from tinygrad.runtime.autogen import llvm"
         python3 -c "from tinygrad.runtime.autogen import webgpu"

docs/abstractions4.py

@@ -105,7 +105,7 @@ def example_3_custom_uop(a:Tensor, correct):
 def example_5_custom_assembly(a:Tensor, correct):
   # Kernel class copied from amd_asm_matmul
   class Kernel:
-    def __init__(self, arch='gfx1100'): self.instructions, self.labels, self.pos, self.arch = [], {}, 0, arch
+    def __init__(self): self.instructions, self.labels, self.pos = [], {}, 0
     def label(self, name): self.labels[name] = self.pos
     def emit(self, inst, target=None):
       self.instructions.append(inst)

docs/tensor/elementwise.md

@@ -68,6 +68,7 @@ Elementwise ops operate on a per element basis. They don't change the shape of t
 ::: tinygrad.Tensor.div
 ::: tinygrad.Tensor.idiv
 ::: tinygrad.Tensor.mod
+::: tinygrad.Tensor.fmod
 ::: tinygrad.Tensor.bitwise_xor
 ::: tinygrad.Tensor.bitwise_and
 ::: tinygrad.Tensor.bitwise_or

examples/mlperf/model_train.py

@@ -1357,6 +1357,7 @@ def train_llama3():
       MLLOGGER.event(key=mllog_constants.OPT_LR_WARMUP_STEPS, value=WARMUP_STEPS)
       MLLOGGER.event(key=mllog_constants.NUM_WARMUP_STEPS, value=WARMUP_STEPS)
       MLLOGGER.event(key=mllog_constants.OPT_LR_DECAY_STEPS, value=MAX_STEPS - WARMUP_STEPS)
+      MLLOGGER.event(key=mllog_constants.OPT_LR_DECAY_SCHEDULE, value="cosine with linear warmup")
       MLLOGGER.event(key=mllog_constants.OPT_GRADIENT_CLIP_NORM, value=1.0)
   else:
     MLLOGGER = None
@@ -1433,6 +1434,7 @@ def train_llama3():
     load_state_dict(scheduler, safe_load(fn), realize=False)
 
   fp8_amax = [t for ts in model._fp8_amax.values() for t in ts]
+  fp8_grad_amax = [t for ts in model._fp8_grad_amax.values() for t in ts] if hasattr(model, "_fp8_grad_amax") else []
   fp8_inv_scales = list(model._fp8_inv_scale.values())
 
   from tinygrad.nn.state import get_state_dict
@@ -1444,6 +1446,10 @@ def train_llama3():
       idx = next(j for j, p in enumerate(optim.params) if p is w)
       optim.master_params[idx].assign((optim.master_params[idx] * w._inv_scale.reshape(-1, *([1]*(w.ndim-1)))).contiguous())
 
+  # realize everything here
+  if optim.master_params: Tensor.realize(*optim.master_params)
+  Tensor.realize(*optim.params, *fp8_inv_scales, *fp8_amax, *fp8_grad_amax)
+
   @TinyJit
   def minibatch(tokens:Tensor):
     if is_dp: tokens = tokens.to(None).shard(device, 0)
@@ -1460,7 +1466,7 @@ def train_llama3():
       apply_grad(g, new_g.uop)
 
     loss_cpu = loss.flatten().float().to("CPU")
-    return loss_cpu.realize(*grads, *fp8_amax)
+    return loss_cpu.realize(*grads, *fp8_amax, *fp8_grad_amax)
 
   @TinyJit
   def optim_step():
@@ -1635,7 +1641,6 @@ def train_llama3():
         tqdm.write(f"target achieved after {sequences_seen} sequences")
         if MLLOGGER and RUNMLPERF:
           MLLOGGER.end(key=mllog_constants.EPOCH_STOP, metadata={mllog_constants.SAMPLES_COUNT: sequences_seen})
-          MLLOGGER.event(key=mllog_constants.TRAIN_SAMPLES, value=sequences_seen)
           MLLOGGER.end(key=mllog_constants.RUN_STOP, metadata={mllog_constants.STATUS: mllog_constants.SUCCESS})
         if getenv("CKPT"):
           if not os.path.exists(ckpt_dir := "./ckpts"): os.mkdir(ckpt_dir)

examples/mlperf/models/flat_llama.py

@@ -20,55 +20,73 @@ from extra.llama_kernels.rmsnorm import rmsnorm
 from extra.llama_kernels import FP8_MAX, local_abs_max
 
 ASM_GEMM = getenv("ASM_GEMM", 0)
+FUSED_INPUT_QUANTIZE = getenv("FUSED_INPUT_QUANTIZE", 0)
+FUSED_ADD_NORM_MUL_QUANTIZE = getenv("FUSED_ADD_NORM_MUL_QUANTIZE", 0)
+FUSED_SILU_W13 = getenv("FUSED_SILU_W13", 0)
 
 FP8_DTYPE = dtypes.fp8e4m3
 FP8_GRAD_DTYPE = dtypes.fp8e5m2
 
 def quantize_fp8(x:Tensor, amax_state:Tensor|None=None):
-  new_amax = (local_abs_max(x) if isinstance(x.device, tuple) else x.abs().max()).detach()
+  new_amax = (local_abs_max(x) if isinstance(x.device, tuple) else x.abs().max()).detach().cast(dtypes.float32)
   scale = FP8_MAX / ((amax_state if amax_state is not None else new_amax) + 1e-8)
   x_scaled = x * scale
   x_clamped = x_scaled + (x_scaled.detach().clamp(-FP8_MAX, FP8_MAX) - x_scaled.detach())  # STE
   return x_clamped.cast(FP8_DTYPE), scale.float().reciprocal(), new_amax
 
 def matmul(x:Tensor, w:Tensor, fp8:bool=True, amax_x:Tensor|None=None, w_inv_scale:Tensor|None=None,
-           x_fp8:Tensor|None=None, x_scale:Tensor|None=None, x_new_amax:Tensor|None=None) -> tuple[Tensor,...]:
+           x_fp8:Tensor|None=None, x_scale:Tensor|None=None, x_new_amax:Tensor|None=None,
+           grad_amax_state:Tensor|None=None) -> tuple[Tensor,...]:
   if not fp8:
     if ASM_GEMM:
       from extra.gemm.cdna_asm_gemm import can_use_asm_gemm, asm_gemm
       if can_use_asm_gemm(x, w.T): return (asm_gemm(x, w.T),)
     return (x @ w.T,)
   assert w_inv_scale is not None, "fp8 matmul requires w_inv_scale (weights must be stored in fp8 with per-tensor scale)"
-  if x_fp8 is None: x_fp8, x_scale, x_new_amax = quantize_fp8(x, amax_state=amax_x)
+  if x_fp8 is None:
+    if FUSED_INPUT_QUANTIZE and amax_x is not None:
+      from extra.llama_kernels.quantize_fp8_delayed import quantize_fp8_delayed
+      x_fp8, x_scale, x_new_amax, _ = quantize_fp8_delayed(x, amax_x, FP8_DTYPE)
+    else:
+      x_fp8, x_scale, x_new_amax = quantize_fp8(x, amax_state=amax_x)
   if ASM_GEMM:
     from extra.gemm.cdna_asm_gemm import can_use_asm_gemm, asm_gemm
-    if can_use_asm_gemm(x_fp8, w.T): return asm_gemm(x_fp8, w.T, x_scale=x_scale, w_scale=w_inv_scale), x_new_amax, x_fp8, w
-  return (x_fp8.dot(w.T, dtype=dtypes.float) * x_scale * w_inv_scale).cast(dtypes.bfloat16), x_new_amax, x_fp8, w
+    if can_use_asm_gemm(x_fp8, w.T):
+      return asm_gemm(x_fp8, w.T, x_scale=x_scale, w_scale=w_inv_scale, grad_amax_state=grad_amax_state), x_new_amax, x_fp8, w
+  return x_fp8.dot(w.T, dtype=dtypes.float) * x_scale * w_inv_scale, x_new_amax, x_fp8, w
 
-def norm_mul_quantize_matmul(x:Tensor, norm:Tensor, amax_x, w_inv_scale, w:Tensor, eps:float):
-  FUSED_NORM_MUL_QUANTIZE = getenv("FUSED_NORM_MUL_QUANTIZE", 0)
-  normed, rrms = rmsnorm(x, eps)
-  if FUSED_NORM_MUL_QUANTIZE:
-    from extra.llama_kernels.fused_mul_quantize_fp8 import fused_mul_quantize_fp8
-    amax_s = amax_x if amax_x is not None else Tensor.full((), 1.0, dtype=dtypes.bfloat16, device=normed.device)
-    x_fp8, x_inv_scale, new_amax = fused_mul_quantize_fp8(normed, norm, amax_s, FP8_DTYPE)
+def norm_quantize_matmul(x:Tensor, norm:Tensor, w:Tensor, w_inv_scale:Tensor, eps:float, amax_x:Tensor, grad_amax_state:Tensor):
+  if FUSED_ADD_NORM_MUL_QUANTIZE:
+    from extra.llama_kernels.fused_rmsnorm_mul_quantize_fp8 import fused_rmsnorm_mul_quantize_fp8
+    x_fp8, x_inv_scale, new_amax, x_normed, rrms = fused_rmsnorm_mul_quantize_fp8(x, norm, amax_x, eps, FP8_DTYPE)
+    out, *ret = matmul(None, w, w_inv_scale=w_inv_scale, x_fp8=x_fp8, x_scale=x_inv_scale, x_new_amax=new_amax, grad_amax_state=grad_amax_state)
+    return out, x_normed, rrms, ret
+  x_normed, rrms = rmsnorm(x, eps)
+  out, *ret = matmul(x_normed * norm, w, amax_x=amax_x, w_inv_scale=w_inv_scale, grad_amax_state=grad_amax_state)
+  return out, x_normed, rrms, ret
+
+def add_norm_quantize_matmul(x:Tensor, residual:Tensor, norm:Tensor, w:Tensor, w_inv_scale:Tensor, eps:float, amax_x:Tensor):
+  if FUSED_ADD_NORM_MUL_QUANTIZE:
+    from extra.llama_kernels.fused_rmsnorm_mul_quantize_fp8 import fused_add_rmsnorm_mul_quantize_fp8
+    x_fp8, x_inv_scale, new_amax, h, x_normed, rrms = fused_add_rmsnorm_mul_quantize_fp8(x, residual, norm, amax_x, eps, FP8_DTYPE)
     out, *ret = matmul(None, w, w_inv_scale=w_inv_scale, x_fp8=x_fp8, x_scale=x_inv_scale, x_new_amax=new_amax)
-  else:
-    x = normed * norm
-    out, *ret = matmul(x, w, amax_x=amax_x, w_inv_scale=w_inv_scale)
-  return out, normed, rrms, ret
+    return out, h, x_normed, rrms, ret
+  h = x + residual
+  x_normed, rrms = rmsnorm(h, eps)
+  out, *ret = matmul(x_normed * norm, w, amax_x=amax_x, w_inv_scale=w_inv_scale)
+  return out, h, x_normed, rrms, ret
 
-def silu_w13_matmul(x_w13:Tensor, w2:Tensor, amax_x2, s_2):
-  FUSED_SILU_W13 = getenv("FUSED_SILU_W13", 0)
+def silu_w13_quantize_matmul(x_w13:Tensor, w2:Tensor, s_2:Tensor,
+                             amax_x2:Tensor,
+                             grad_amax_xw13:Tensor, grad_amax_xout:Tensor):
   if FUSED_SILU_W13:
     from extra.llama_kernels.cast_amax import fused_quantize_fp8_w13
-    amax_s = amax_x2 if amax_x2 is not None else Tensor.full((), 1.0, dtype=dtypes.bfloat16, device=x_w13.device)
-    x2_fp8, x2_inv_scale, new_amax_x2 = fused_quantize_fp8_w13(x_w13, amax_s, FP8_DTYPE)
-    out, *ret = matmul(None, w2, w_inv_scale=s_2, x_fp8=x2_fp8, x_scale=x2_inv_scale, x_new_amax=new_amax_x2)
-  else:
-    hidden_dim = x_w13.shape[-1] // 2
-    x_w1, x_w3 = x_w13[..., :hidden_dim], x_w13[..., hidden_dim:]
-    out, *ret = matmul(x_w1.silu() * x_w3, w2, amax_x=amax_x2, w_inv_scale=s_2)
+    x2_fp8, x2_inv_scale, new_amax_x2 = fused_quantize_fp8_w13(x_w13, amax_x2, FP8_DTYPE, grad_amax_state=grad_amax_xw13)
+    out, *ret = matmul(None, w2, w_inv_scale=s_2, x_fp8=x2_fp8, x_scale=x2_inv_scale, x_new_amax=new_amax_x2, grad_amax_state=grad_amax_xout)
+    return out, ret
+  hidden = x_w13.shape[-1] // 2
+  x_w1, x_w3 = x_w13[..., :hidden], x_w13[..., hidden:]
+  out, *ret = matmul(x_w1.silu() * x_w3, w2, amax_x=amax_x2, w_inv_scale=s_2, grad_amax_state=grad_amax_xout)
   return out, ret
 
 class FlatTransformer:
@@ -104,32 +122,33 @@ class FlatTransformer:
     self.output = Tensor.normal(1, vocab_size, dim, mean=0.0, std=0.02, dtype=dtypes.bfloat16)
     self.freqs_cis = precompute_freqs_cis(dim // n_heads, max_context * 2, rope_theta).contiguous().requires_grad_(False)
 
-    def _amax(): return Tensor.full((), FP8_MAX).contiguous().requires_grad_(False)
+    def _amax(): return Tensor.full((), FP8_MAX, dtype=dtypes.float32).contiguous().requires_grad_(False)
     names = ["xqkv", "xo", "x13", "x2"]
     self._fp8_amax = {name: [_amax() for _ in range(n_layers)] for name in names}
-    # per-weight inv_scale: single (n_layers,) float32 tensor per weight (kernel reads float* pointers)
+    grad_names = ["xqkv", "xo", "xw13", "xout"]
+    self._fp8_grad_amax = {name: [_amax() for _ in range(n_layers)] for name in grad_names}
     w_names = ["wqkv", "wo", "w13", "w2"]
-    self._fp8_inv_scale = {}
-    for wname, inv_scales in zip(w_names, self._init_inv_scales):
-      self._fp8_inv_scale[wname] = inv_scales.float().contiguous().requires_grad_(False)
+    self._fp8_inv_scale = {wname: inv_scales.float().contiguous().requires_grad_(False)
+                           for wname, inv_scales in zip(w_names, self._init_inv_scales)}
     del self._init_inv_scales
 
   def lin_per_layer(self, in_features:int, out_features:int, std:float=0.02):
-    if getenv("ZEROS", 0): w = Tensor.zeros(self.n_layers, out_features, in_features)
+    if getenv("ZEROS"): w = Tensor.zeros(self.n_layers, out_features, in_features)
     else: w = Tensor.normal(self.n_layers, out_features, in_features, mean=0.0, std=std)
-    # per-layer scaled fp8 cast: fill the fp8 range for best precision
     amax = w.abs().flatten(1).max(1).detach()
     scale = FP8_MAX / (amax + 1e-8)
-    self._init_inv_scales.append((amax + 1e-8) / FP8_MAX)  # save for inv_scale init
+    self._init_inv_scales.append((amax + 1e-8) / FP8_MAX)
     return (w * scale.reshape(-1, 1, 1)).clamp(-FP8_MAX, FP8_MAX).cast(FP8_DTYPE)
 
   def attention(self, x:Tensor, freqs_cis:Tensor, attention_norm:Tensor, wqkv:Tensor, wo:Tensor,
-                amax_xqkv=None, amax_xo=None, s_qkv=None, s_o=None):
+                amax_xqkv:Tensor, amax_xo:Tensor, s_qkv:Tensor, s_o:Tensor,
+                grad_amax_xqkv:Tensor, grad_amax_xo:Tensor):
     bsz, seqlen, _ = x.shape
     new_amaxs, saves = [], []
 
-    xqkv, normed, rrms, ret = norm_mul_quantize_matmul(x, attention_norm, amax_xqkv, s_qkv, wqkv, self.norm_eps)
-    saves.extend([normed, rrms])
+    xqkv, x_normed, rrms, ret = norm_quantize_matmul(x, attention_norm, wqkv, s_qkv, self.norm_eps,
+                                                     amax_x=amax_xqkv, grad_amax_state=grad_amax_xqkv)
+    saves.extend([x_normed, rrms])
     new_amaxs.extend(ret[:1])
     saves.extend(ret[1:] + [xqkv])
     xqkv = xqkv.reshape(bsz, seqlen, self.n_kv_heads, self.n_rep + 2, self.head_dim)
@@ -139,49 +158,52 @@ class FlatTransformer:
 
     xq, xk = apply_rotary_emb(xq, xk, freqs_cis)
     xq, xk, xv = xq.cast(dtypes.bfloat16), xk.cast(dtypes.bfloat16), xv.cast(dtypes.bfloat16)
-    xq, xk, xv = xq.transpose(1, 2), xk.transpose(1, 2), xv.transpose(1, 2)
     if getenv("HK_FLASH_ATTENTION"):
       from extra.thunder.amd.fa import flash_attention
       attn, *save = flash_attention(xq, xk, xv, is_causal=True)
       saves.extend(save)
     else:
-      attn = xq.scaled_dot_product_attention(xk, xv, is_causal=True, enable_gqa=True)
-    attn = attn.transpose(1, 2).reshape(bsz, seqlen, -1)
+      xq, xk, xv = xq.transpose(1, 2), xk.transpose(1, 2), xv.transpose(1, 2)
+      attn = xq.scaled_dot_product_attention(xk, xv, is_causal=True, enable_gqa=True).transpose(1, 2)
+    attn = attn.reshape(bsz, seqlen, -1)
 
-    out, *ret = matmul(attn, wo, amax_x=amax_xo, w_inv_scale=s_o)
+    out, *ret = matmul(attn, wo, amax_x=amax_xo, w_inv_scale=s_o, grad_amax_state=grad_amax_xo)
     new_amaxs.extend(ret[:1])
     saves.extend(ret[1:] + [out])
     return (out, *new_amaxs, *saves)
 
-  def feed_forward(self, x:Tensor, ffn_norm:Tensor, w13:Tensor, w2:Tensor,
-                   amax_x13=None, amax_x2=None, s_13=None, s_2=None):
+  def feed_forward(self, x:Tensor, residual:Tensor, ffn_norm:Tensor, w13:Tensor, w2:Tensor,
+                   amax_x13:Tensor, amax_x2:Tensor, s_13:Tensor, s_2:Tensor,
+                   grad_amax_xw13:Tensor, grad_amax_xout:Tensor):
     new_amaxs, saves = [], []
 
-    x_w13, normed, rrms, ret = norm_mul_quantize_matmul(x, ffn_norm, amax_x13, s_13, w13, self.norm_eps)
-    saves.extend([normed, rrms])
+    x_w13, h, x_normed, rrms, ret = add_norm_quantize_matmul(x, residual, ffn_norm, w13, s_13, self.norm_eps,
+                                                             amax_x=amax_x13)
+    saves.extend([x_normed, rrms])
     new_amaxs.extend(ret[:1])
     saves.extend(ret[1:] + [x_w13])
 
-    out, ret = silu_w13_matmul(x_w13, w2, amax_x2, s_2)
+    out, ret = silu_w13_quantize_matmul(x_w13, w2, s_2, amax_x2=amax_x2, grad_amax_xw13=grad_amax_xw13, grad_amax_xout=grad_amax_xout)
     new_amaxs.extend(ret[:1])
     saves.extend(ret[1:] + [out])
-    return (out, *new_amaxs, *saves)
+    return (out, h, *new_amaxs, *saves)
 
   @function(precompile=True, precompile_backward=True)
   def run_layer(self, x:Tensor, freqs_cis:Tensor,
                 attention_norm:Tensor, wqkv:Tensor, wo:Tensor,
                 ffn_norm:Tensor, w13:Tensor, w2:Tensor,
-                amax_xqkv=None, amax_xo=None,
-                amax_x13=None, amax_x2=None,
-                s_qkv=None, s_o=None, s_13=None, s_2=None):
+                amax_xqkv:Tensor, amax_xo:Tensor,
+                amax_x13:Tensor, amax_x2:Tensor,
+                s_qkv:Tensor, s_o:Tensor, s_13:Tensor, s_2:Tensor,
+                grad_amax_xqkv:Tensor, grad_amax_xo:Tensor,
+                grad_amax_xw13:Tensor, grad_amax_xout:Tensor):
     attn, *attn_ret = self.attention(x, freqs_cis, attention_norm, wqkv, wo,
-                                     amax_xqkv=amax_xqkv, amax_xo=amax_xo,
-                                     s_qkv=s_qkv, s_o=s_o)
+                                     amax_xqkv=amax_xqkv, amax_xo=amax_xo, s_qkv=s_qkv, s_o=s_o,
+                                     grad_amax_xqkv=grad_amax_xqkv, grad_amax_xo=grad_amax_xo)
     attn_amaxs, attn_saves = attn_ret[:2], attn_ret[2:]
-    h = x + attn
-    ffn, *ffn_ret = self.feed_forward(h, ffn_norm, w13, w2,
-                                      amax_x13=amax_x13, amax_x2=amax_x2,
-                                      s_13=s_13, s_2=s_2)
+    ffn, h, *ffn_ret = self.feed_forward(x, attn, ffn_norm, w13, w2,
+                                                   amax_x13=amax_x13, amax_x2=amax_x2, s_13=s_13, s_2=s_2,
+                                                   grad_amax_xw13=grad_amax_xw13, grad_amax_xout=grad_amax_xout)
     ffn_amaxs, ffn_saves = ffn_ret[:2], ffn_ret[2:]
     h = h + ffn
     return (h, *attn_amaxs, *ffn_amaxs, *attn_saves, *ffn_saves)
@@ -202,28 +224,31 @@ class FlatTransformer:
       self.tok_embeddings.weight.shard_(device, axis=0).realize()
       self.output.shard_(device, axis=1).realize()
       self.freqs_cis.shard_(device, axis=None).realize()
-      for name in self._fp8_amax:
-        for i in range(len(self._fp8_amax[name])):
-          self._fp8_amax[name][i] = self._fp8_amax[name][i].to(device).contiguous().requires_grad_(False)
+      for amax_dict in (self._fp8_amax, self._fp8_grad_amax):
+        for name in amax_dict:
+          for i in range(len(amax_dict[name])):
+            amax_dict[name][i] = amax_dict[name][i].to(device).contiguous().requires_grad_(False)
       for name in self._fp8_inv_scale:
         self._fp8_inv_scale[name] = self._fp8_inv_scale[name].to(device).contiguous().requires_grad_(False)
 
   def __call__(self, tokens:Tensor):
     h = self.tok_embeddings(tokens)
     freqs_cis = self.freqs_cis.cast(h.dtype)[:, :tokens.shape[1], :, :, :]
-    amaxs, inv_scales = self._fp8_amax, self._fp8_inv_scale
+    a, ga, s = self._fp8_amax, self._fp8_grad_amax, self._fp8_inv_scale
     for i in range(self.n_layers):
       h, *ret = self.run_layer(h, freqs_cis,
                                self.attention_norm[i], self.wqkv[i], self.wo[i],
                                self.ffn_norm[i], self.w13[i], self.w2[i],
-                               amax_xqkv=amaxs["xqkv"][i], amax_xo=amaxs["xo"][i],
-                               amax_x13=amaxs["x13"][i], amax_x2=amaxs["x2"][i],
-                               s_qkv=inv_scales["wqkv"][i], s_o=inv_scales["wo"][i],
-                               s_13=inv_scales["w13"][i], s_2=inv_scales["w2"][i])
+                               amax_xqkv=a["xqkv"][i], amax_xo=a["xo"][i],
+                               amax_x13=a["x13"][i], amax_x2=a["x2"][i],
+                               s_qkv=s["wqkv"][i], s_o=s["wo"][i],
+                               s_13=s["w13"][i], s_2=s["w2"][i],
+                               grad_amax_xqkv=ga["xqkv"][i], grad_amax_xo=ga["xo"][i],
+                               grad_amax_xw13=ga["xw13"][i], grad_amax_xout=ga["xout"][i])
       for name, new_val in zip(["xqkv", "xo", "x13", "x2"], ret[:5]):
-        amaxs[name][i].assign(new_val)
+        a[name][i].assign(new_val)
 
-    logits = matmul(self.norm(h).contiguous().contiguous_backward(), self.output[0], fp8=False)[0].contiguous_backward()
+    logits = matmul(self.norm(h), self.output[0], fp8=False)[0]
     return logits
 
 def _get_pads(uop:UOp) -> list[UOp]:
@@ -239,6 +264,11 @@ def apply_grad(grad_buf:Tensor, new_grad:UOp):
     return
   sorted_pads = sorted(pads, key=lambda p: p.marg[0][0] if p.op == Ops.PAD else 0)
   inners = [Tensor(p.src[0] if p.op == Ops.PAD else p, device=grad_buf.device).cast(grad_buf.dtype) for p in sorted_pads]
+  if getenv("FUSED_PAD_GRAD_ACCUM", 0):
+    from extra.llama_kernels.fused_pad_grad_accum import fused_pad_grad_accum, can_fused_pad_grad_accum
+    if can_fused_pad_grad_accum(grad_buf, inners):
+      grad_buf.uop = fused_pad_grad_accum(grad_buf, inners).uop
+      return
   grad_buf.assign(grad_buf + inners[0].cat(*inners[1:], dim=0))
 
 if __name__ == "__main__":

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama8b/implementations/tinybox_8xMI350X/dev_beam.sh

@@ -16,6 +16,10 @@ export MASTER_WEIGHTS=${MASTER_WEIGHTS:-1}
 export FP8=${FP8:-1}
 export ALLREDUCE_CAST=${ALLREDUCE_CAST:-1}
 export FAST_CE=${FASE_CE:-1}
+export FUSED_INPUT_QUANTIZE=${FUSED_INPUT_QUANTIZE:-1}
+export FUSED_ADD_NORM_MUL_QUANTIZE=${FUSED_ADD_NORM_MUL_QUANTIZE:-1}
+export FUSED_SILU_W13=${FUSED_SILU_W13:-1}
+export FUSED_PAD_GRAD_ACCUM=${FUSED_PAD_GRAD_ACCUM:-1}
 
 export DEFAULT_FLOAT="bfloat16" OPTIM_DTYPE="bfloat16"
 export DP=${DP:-8} MP=${MP:-1} BS=${BS:-16} EVAL_BS=${EVAL_BS:-8} GRADIENT_ACC_STEPS=${GRADIENT_ACC_STEPS:-2}

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama8b/implementations/tinybox_8xMI350X/dev_run.sh

@@ -16,6 +16,10 @@ export MASTER_WEIGHTS=${MASTER_WEIGHTS:-1}
 export FP8=${FP8:-1}
 export ALLREDUCE_CAST=${ALLREDUCE_CAST:-1}
 export FAST_CE=${FASE_CE:-1}
+export FUSED_INPUT_QUANTIZE=${FUSED_INPUT_QUANTIZE:-1}
+export FUSED_ADD_NORM_MUL_QUANTIZE=${FUSED_ADD_NORM_MUL_QUANTIZE:-1}
+export FUSED_SILU_W13=${FUSED_SILU_W13:-1}
+export FUSED_PAD_GRAD_ACCUM=${FUSED_PAD_GRAD_ACCUM:-1}
 
 export DEFAULT_FLOAT="bfloat16" OPTIM_DTYPE="bfloat16"
 export DP=${DP:-8} MP=${MP:-1} BS=${BS:-16} EVAL_BS=${EVAL_BS:-8} GRADIENT_ACC_STEPS=${GRADIENT_ACC_STEPS:-2}

examples/mlperf/training_submission_v6.0/tinycorp/benchmarks/llama8b/implementations/tinybox_8xMI350X/run_and_time.sh

@@ -17,6 +17,10 @@ export MASTER_WEIGHTS=1
 export FP8=1
 export ALLREDUCE_CAST=1
 export FAST_CE=1
+export FUSED_INPUT_QUANTIZE=1
+export FUSED_ADD_NORM_MUL_QUANTIZE=1
+export FUSED_SILU_W13=1
+export FUSED_PAD_GRAD_ACCUM=1
 
 export DEFAULT_FLOAT="bfloat16" OPTIM_DTYPE="bfloat16"
 export DP=8 MP=1 BS=16 EVAL_BS=8 GRADIENT_ACC_STEPS=2

extra/amdpci/am_smi.py

@@ -64,7 +64,7 @@ def get_bar0_size(pcibus):
 
 class AMSMI(AMDev):
   def __init__(self, pcibus, vram_bar:MMIOInterface, doorbell_bar:MMIOInterface, mmio_bar:MMIOInterface):
-    self.pcibus = pcibus
+    self.pcibus, self.devfmt = pcibus, pcibus
     self.vram, self.doorbell64, self.mmio = vram_bar, doorbell_bar, mmio_bar
     self.pci_state = self.read_pci_state()
     if self.pci_state == "D0": self._init_from_d0()
@@ -91,6 +91,7 @@ class SMICtx:
     self.prev_lines_cnt = 0
     self.prev_terminal_width = 0
     self.prev_terminal_height = 0
+    self.prev_metrics = {}
 
     remove_parts = ["Advanced Micro Devices, Inc. [AMD/ATI]", "VGA compatible controller:", "Processing accelerators:"]
     lspci = subprocess.check_output(["lspci"]).decode("utf-8").splitlines()
@@ -235,6 +236,29 @@ class SMICtx:
       case (13,0,12): return self._smuq10_round(metrics.SocketPower), self._smuq10_round(metrics.SocketPowerLimit)
       case _: return metrics.SmuMetrics.AverageSocketPower, metrics.SmuMetrics.dGPU_W_MAX
 
+  def get_throttle_info(self, dev, metrics):
+    match dev.ip_ver[am.MP1_HWIP]:
+      case (13,0,6)|(13,0,12):
+        throttle_fields = [('ProchotResidencyAcc', 'Prochot'), ('PptResidencyAcc', 'PPT'),
+                           ('SocketThmResidencyAcc', 'Socket Thm'), ('VrThmResidencyAcc', 'VR Thm'), ('HbmThmResidencyAcc', 'HBM Thm')]
+        prev = self.prev_metrics.get(dev.pcibus)
+        active = []
+        if prev is not None:
+          acc_delta = metrics.AccumulationCounter - prev.AccumulationCounter
+          if acc_delta > 0:
+            for field, name in throttle_fields:
+              delta = getattr(metrics, field) - getattr(prev, field)
+              if delta > 0 and (pct := min(100, (delta * 100 + acc_delta // 2) // acc_delta)) > 0: active.append((name, pct))
+        return active
+      case _:
+        smu_mod = dev.smu.smu_mod
+        throttler_names = {getattr(smu_mod, a): a[len('THROTTLER_'):-len('_BIT')]
+                           for a in dir(smu_mod) if a.startswith('THROTTLER_') and a.endswith('_BIT')}
+        active = []
+        for i, pct in enumerate(metrics.SmuMetrics.ThrottlingPercentage):
+          if pct > 0: active.append((throttler_names.get(i, f"UNK_{i}"), int(pct)))
+        return active
+
   def get_mem_usage(self, dev):
     usage = 0
     pt_stack = [dev.mm.root_page_table]
@@ -281,6 +305,13 @@ class SMICtx:
                     + [f"MEM Activity {draw_bar(self.get_mem_activity(dev, metrics) / 100, activity_line_width)}"] \
                     + [f"MEM Usage    {draw_bar(mem_used / mem_total, activity_line_width, opt_text=mem_fmt)}"] \
 
+      throttle_info = self.get_throttle_info(dev, metrics)
+      if throttle_info:
+        throttle_text = colored(', '.join(f"{name} {pct}%" for name, pct in throttle_info), "red")
+      else:
+        throttle_text = colored("None", "green")
+      activity_line += [f"Throttle     {throttle_text}" + " " * (activity_line_width + 2)]
+
       temps_data, temps_data_compact = self.get_temps(dev, metrics), self.get_temps(dev, metrics, compact=True)
       temps_table = ["=== Temps (°C) ==="] + [f"{name:<16}: {color_temp(val)}" for name, val in temps_data.items()]
       temps_table_compact = ["Temps (°C):" + '/'.join([f"{color_temp(val)} {name}" for name, val in temps_data_compact.items()])]
@@ -324,6 +355,8 @@ class SMICtx:
 
       dev_content.append(device_line + activity_line + same_line([temps_table, power_table, frequency_table]))
 
+    self.prev_metrics = {dev.pcibus: m for dev, m in dev_metrics.items() if m is not None}
+
     raw_text = 'AM Monitor'.center(terminal_width) + "\n" + "=" * terminal_width + "\n\n"
     for i in range(0, len(dev_content), 2):
       if i + 1 < len(dev_content): raw_text += '\n'.join(same_line([dev_content[i], dev_content[i+1]], split=padding))

extra/amdpci/headers/amdgpu_irq.h

@@ -28,15 +28,7 @@
 // #include "soc15_ih_clientid.h"
 // #include "amdgpu_ih.h"
 
-#define int32_t int
-#define uint32_t unsigned int
-#define int8_t signed char
-#define uint8_t unsigned char
-#define uint16_t unsigned short
-#define int16_t short
-#define uint64_t unsigned long long
 #define bool _Bool
-#define u32 unsigned int
 
 #define AMDGPU_MAX_IRQ_SRC_ID		0x100
 #define AMDGPU_MAX_IRQ_CLIENT_ID	0x100

extra/amdpci/headers/amdgpu_smu.h

@@ -22,15 +22,7 @@
 #ifndef __AMDGPU_SMU_H__
 #define __AMDGPU_SMU_H__
 
-#define int32_t int
-#define uint32_t unsigned int
-#define int8_t signed char
-#define uint8_t unsigned char
-#define uint16_t unsigned short
-#define int16_t short
-#define uint64_t unsigned long long
 #define bool _Bool
-#define u32 unsigned int
 
 #define SMU_THERMAL_MINIMUM_ALERT_TEMP		0
 #define SMU_THERMAL_MAXIMUM_ALERT_TEMP		255

extra/amdpci/headers/amdgpu_ucode.h

@@ -24,15 +24,7 @@
 #define __AMDGPU_UCODE_H__
 
 // #include "amdgpu_socbb.h"
-#define int32_t int
-#define uint32_t unsigned int
-#define int8_t signed char
-#define uint8_t unsigned char
-#define uint16_t unsigned short
-#define int16_t short
-#define uint64_t unsigned long long
 #define bool _Bool
-#define u32 unsigned int
 
 struct common_firmware_header {
 	uint32_t size_bytes; /* size of the entire header+image(s) in bytes */

extra/gemm/amd_asm_matmul.py

@@ -167,7 +167,7 @@ PREFETCH_LOADS = [(V_LDS_A_DATA[4+2*i], V_LDS_A_DATA[4+2*i+1], V_GLOBAL_B_ADDR,
 # =============================================================================
 
 class Kernel:
-  def __init__(self, arch='gfx1100'): self.instructions, self.labels, self.pos, self.arch = [], {}, 0, arch
+  def __init__(self): self.instructions, self.labels, self.pos = [], {}, 0
   def label(self, name): self.labels[name] = self.pos
 
   def emit(self, inst, target=None):
@@ -196,10 +196,10 @@ class Kernel:
 # Kernel builder
 # =============================================================================
 
-def build_kernel(N, arch='gfx1100'):
+def build_kernel(N):
   assert N % 128 == 0, f"N must be a multiple of 128 (tile size), got {N}"
   assert N >= 256, f"N must be >= 256 (prefetch pipeline requires at least 2 K-blocks), got {N}"
-  k = Kernel(arch)
+  k = Kernel()
 
   # ===========================================================================
   # PROLOGUE: Load kernel arguments, compute tile coordinates and addresses
@@ -443,7 +443,7 @@ def test_matmul():
   dev = Device[Device.DEFAULT]
   print(f"Device arch: {dev.renderer.target.arch}")
 
-  insts = build_kernel(N, dev.renderer.target.arch)
+  insts = build_kernel(N)
 
   rng = np.random.default_rng(42)
   a = Tensor(rng.random((N, N), dtype=np.float32) - 0.5)

extra/gemm/cdna_asm_gemm.py

@@ -2628,20 +2628,24 @@ def custom_asm_gemm(C:UOp, A:UOp, B:UOp, dname:str) -> UOp:
 # ** FP8 GEMM custom kernel
 
 @functools.cache
-def custom_hk_fp8_gemm(C:UOp, A:UOp, B:UOp, X_s:UOp, W_s:UOp, dname:str) -> UOp:
-  # A is (batch, M, K), B is (N, K) transposed, X_s is x_scale, W_s is w_scale — kernel multiplies by both
+def custom_hk_fp8_gemm(C:UOp, A:UOp, B:UOp, *args:UOp, dname:str, scale_mode:int=3) -> UOp:
+  # scale_mode: 0=no scale, 1=x only, 2=w only, 3=both
+  n_scales = (1 if scale_mode & 1 else 0) + (1 if scale_mode & 2 else 0)
+  scales, extra = args[:n_scales], args[n_scales:]
   M, K = A.shape[0]*A.shape[1], A.shape[2]
   N, K2 = B.shape[(1 if B.ndim == 3 else 0):]
   assert K == K2, f"{A.shape} {B.shape}"
   block_size = 256
   threads = UOp.special(64 * 8, "lidx0")
   workgroups = UOp.special((M // block_size) * (N // block_size), "gidx0")
-  sink = UOp.sink(C.base, A.base, B.base, X_s.base, W_s.base, threads, workgroups,
+  sink_inputs = (C.base, A.base, B.base) + tuple(s.base for s in scales) + (threads, workgroups)
+  sink = UOp.sink(*sink_inputs,
                   arg=KernelInfo(f"hk_fp8_gemm_{M}_{N}_{K}", estimates=Estimates(ops=2*M*N*K, mem=(M*K+N*K)*A.dtype.itemsize+M*N*C.dtype.itemsize)))
   kittens_path = pathlib.Path(__file__).parent.parent/"thunder"/"amd"
   src = (kittens_path/"gemm_fp8.cpp").read_text()
   lib = HIPCCCompiler("gfx950", [f"-I{(kittens_path/'include').as_posix()}", "-std=c++20", "-DKITTENS_CDNA4", "-ffast-math",
-                                 "-DHIP_ENABLE_WARP_SYNC_BUILTINS", f"-DGEMM_M={M}", f"-DGEMM_N={N}", f"-DGEMM_K={K}"]).compile_cached(src)
+                                 "-DHIP_ENABLE_WARP_SYNC_BUILTINS", f"-DGEMM_M={M}", f"-DGEMM_N={N}", f"-DGEMM_K={K}",
+                                 f"-DSCALE_MODE={scale_mode}"]).compile_cached(src)
   return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)), UOp(Ops.SOURCE, arg=src),
                                UOp(Ops.BINARY, arg=lib)))
 
@@ -2698,19 +2702,31 @@ def custom_uop_gemm(C:UOp, A:UOp, B:UOp) -> UOp:
 
 def custom_gemm_bw(gradient:UOp, kernel:UOp):
   inputs = kernel.src[1:]
-  # fp8 scaled gemm has 5 inputs (out, a, b, x_scale, w_scale), others have 3 (out, a, b)
-  if len(inputs) == 5:
-    out, a, b, s_x, s_w = inputs
+  if inputs[1].dtype == FP8_DTYPE:
+    grad_amax_state = inputs[5] if len(inputs) == 6 else None
+    out, a, b, s_x, s_w = inputs[:5]
     a_t, b_t, g_t = Tensor(a, device=a.device), Tensor(b, device=a.device), Tensor(gradient, device=a.device)
     s_x_t, s_w_t = Tensor(s_x, device=a.device), Tensor(s_w, device=a.device)
     g_t = g_t[:a.shape[0]]
-    g_fp8, g_scale, _ = quantize_fp8(g_t)
+    from extra.llama_kernels.cast_amax import _grad_fp8_mailbox
+    from extra.llama_kernels.quantize_fp8_delayed import quantize_fp8_delayed
+    gbase = gradient.base if hasattr(gradient, "base") else gradient
+    mailbox_entry = _grad_fp8_mailbox.pop(gbase, None) or _grad_fp8_mailbox.pop(gradient, None)
+    if mailbox_entry is not None:
+      g_fp8_u, inv_scale_u, _new_amax_u, store_effect = mailbox_entry
+      g_fp8 = Tensor(g_fp8_u, device=a.device)[:a.shape[0]]
+      g_scale = Tensor(inv_scale_u, device=a.device)
+    else:
+      assert grad_amax_state is not None, "fp8 matmul bwd needs either a mailbox entry or a grad_amax_state"
+      g_fp8, g_scale, _, store_effect = quantize_fp8_delayed(g_t, Tensor(grad_amax_state, device=a.device))
     # dgrad: uses g_scale * x_scale * w_scale
     grad_a = asm_gemm(g_fp8, b_t, x_scale=g_scale * s_x_t, w_scale=s_w_t)
     # wgrad: no w_scale
-    _one = Tensor(1.0, dtype=dtypes.float, device=a.device)
-    grad_b = asm_gemm(g_fp8.permute(2, 0, 1).reshape(g_t.shape[-1], -1), a_t.reshape(-1, a_t.shape[-1]), x_scale=g_scale * s_x_t, w_scale=_one)
-    return (None, grad_a.uop, grad_b.uop, None, None)
+    grad_b = asm_gemm(g_fp8.permute(2, 0, 1).reshape(g_t.shape[-1], -1), a_t.reshape(-1, a_t.shape[-1]), x_scale=g_scale * s_x_t)
+    # Attach the delayed-amax store effect (if any) to grad_a so realizing grads commits the amax update.
+    ret = (None, grad_a.uop.after(store_effect), grad_b.uop, None, None)
+    if len(inputs) == 6: ret = ret + (None,)
+    return ret
   else:
     out, a, b = inputs
     assert all_same([gradient.device, a.device, b.device, out.device])
@@ -2725,7 +2741,7 @@ def custom_gemm_bw(gradient:UOp, kernel:UOp):
 
 # ** main gemm function
 
-def asm_gemm(a:Tensor, b:Tensor, x_scale:Tensor|None=None, w_scale:Tensor|None=None) -> Tensor:
+def asm_gemm(a:Tensor, b:Tensor, x_scale:Tensor|None=None, w_scale:Tensor|None=None, grad_amax_state:Tensor|None=None) -> Tensor:
   assert can_use_asm_gemm(a, b), f"{counters['todos'][-1]}"
   counters["used"] += 1
   unfold_batch = a.ndim == 3 and isinstance(a.device, tuple) and a.uop.axis == 2 and b.uop.axis == 0
@@ -2759,10 +2775,11 @@ def asm_gemm(a:Tensor, b:Tensor, x_scale:Tensor|None=None, w_scale:Tensor|None=N
   if arch.startswith("gfx950") and getenv("USE_ASM", 1):
     # fp8 gemm computes [email protected], kernel multiplies output by x_scale * w_scale before bf16 store
     if a.dtype == FP8_DTYPE:
-      _one = lambda: Tensor(1.0, dtype=dtypes.float, device=a.device)
-      xs = x_scale if x_scale is not None else _one()
-      ws = w_scale if w_scale is not None else _one()
-      out = Tensor.custom_kernel(out, a, b.T, xs, ws, fxn=functools.partial(custom_hk_fp8_gemm, dname=dname), grad_fxn=custom_gemm_bw)[0]
+      scales = tuple(s for s in (x_scale, w_scale) if s is not None)
+      scale_mode = (1 if x_scale is not None else 0) | (2 if w_scale is not None else 0)
+      extra = [grad_amax_state] if grad_amax_state is not None else []
+      fxn = functools.partial(custom_hk_fp8_gemm, dname=dname, scale_mode=scale_mode)
+      out = Tensor.custom_kernel(out, a, b.T, *scales, *extra, fxn=fxn, grad_fxn=custom_gemm_bw)[0]
     else:
       out = Tensor.custom_kernel(out, a, b, fxn=functools.partial(custom_asm_gemm, dname=dname), grad_fxn=custom_gemm_bw)[0]
   else:

extra/gemm/triton_nv_matmul.py

@@ -4,7 +4,8 @@ import triton.language as tl
 from triton.compiler import AttrsDescriptor, ASTSource, compile as triton_compile
 import numpy as np
 from tinygrad import Tensor, dtypes, Device
-from tinygrad.engine.realize import CompiledRunner
+from tinygrad.engine.realize import get_runtime
+from tinygrad.codegen import to_program
 from tinygrad.uop.ops import Ops, UOp, KernelInfo, ProgramInfo
 from tinygrad.helpers import getenv
 np.set_printoptions(suppress=True)
@@ -92,13 +93,15 @@ if __name__ == "__main__":
   info = ProgramInfo(name="matmul_kernel",
                      global_size=(M//BLOCK_SIZE_M, N//BLOCK_SIZE_N, 1), local_size=(32*compiled.metadata.num_warps, 1, 1))
   sink = UOp.sink(arg=KernelInfo(name="matmul_kernel"))
-  prg_uop = UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=Device.DEFAULT), UOp(Ops.LINEAR), UOp(Ops.SOURCE, arg=src)), arg=info)
-  runner = CompiledRunner(prg_uop, Device.DEFAULT)
+  prg_uop = to_program(UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=Device.DEFAULT), UOp(Ops.LINEAR), UOp(Ops.SOURCE, arg=src)), arg=info),
+                       Device.default.renderer)
+  rt = get_runtime(Device.DEFAULT, prg_uop)
   all_bufs = [x.ensure_allocated() for x in bufs]
-  prg_bufs = [all_bufs[i] for i in runner.p.globals]
+  prg_bufs = [all_bufs[i] for i in info.globals]
+  gsize, lsize = info.launch_dims({})
   tflops = []
   for i in range(5):
-    tm = runner(prg_bufs, {}, wait=True)
+    tm = rt(*[b._buf for b in prg_bufs], global_size=gsize, local_size=lsize, vals=info.vals({}), wait=True)
     tflops.append((2*M*K*N/tm)*1e-12)
   print(f"TFLOPS: {max(tflops):.2f}")
 

extra/llama_kernels/__init__.py

@@ -29,7 +29,24 @@ def shard_shape(shape:tuple, axis:int, ndev:int) -> list:
   s[axis] //= ndev
   return s
 
+def dname_of(device) -> str:
+  if isinstance(device, tuple): return device[0].split(":")[0]
+  return device.split(":")[0] if isinstance(device, str) else device
+
+def alloc_like(shape, dtype, device, axis=None) -> Tensor:
+  if isinstance(device, tuple):
+    if axis is None: return Tensor(Tensor.invalids(*shape, dtype=dtype, device=device).uop.multi(0), device=device)
+    return Tensor(Tensor.invalids(*shard_shape(shape, axis, len(device)), dtype=dtype, device=device).uop.multi(axis), device=device)
+  return Tensor.invalids(*shape, dtype=dtype, device=device)
+
+def alloc_local(shape, dtype, device) -> Tensor:
+  if isinstance(device, tuple):
+    return Tensor(Tensor.invalids(*shape, dtype=dtype, device=device).uop.multi(0), device=device)
+  return Tensor.invalids(*shape, dtype=dtype, device=device)
+
+def compile_hip(src:str, defines:list[str]):
+  return HIPCCCompiler("gfx950", ["-std=c++20", "-ffast-math", *defines]).compile_cached(src)
+
 def compile_cpp(cpp_dir:pathlib.Path, cpp_name:str, n_elems:int, hidden:int):
   src = (cpp_dir/cpp_name).read_text()
-  defines = [f"-DN_ELEMS={n_elems}", f"-DHIDDEN={hidden}", f"-DNUM_WG={NUM_WG}", f"-DTHREADS_PER_WG={THREADS_PER_WG}"]
-  return src, HIPCCCompiler("gfx950", ["-std=c++20", "-ffast-math", *defines]).compile_cached(src)
+  return src, compile_hip(src, [f"-DN_ELEMS={n_elems}", f"-DHIDDEN={hidden}", f"-DNUM_WG={NUM_WG}", f"-DTHREADS_PER_WG={THREADS_PER_WG}"])

extra/llama_kernels/cast_amax/__init__.py

@@ -3,26 +3,34 @@ import functools, pathlib
 from tinygrad import Tensor, dtypes
 from tinygrad.uop.ops import UOp, Ops, KernelInfo
 from tinygrad.renderer import Estimates
-from extra.llama_kernels import FP8_MAX, NUM_WG, THREADS_PER_WG, compile_cpp, shard_shape, scalar_amax
+from extra.llama_kernels import FP8_MAX, NUM_WG, THREADS_PER_WG, compile_cpp, alloc_like, alloc_local, scalar_amax, dname_of
+
+# module-level mailbox: grad_xw13 UOp -> (grad_xw13_fp8 UOp, inv_scale UOp, new_amax UOp, store_effect)
+# lets cdna_asm_gemm's bwd reuse the fp8 companion produced by the fused silu_mul bwd kernel
+# instead of doing a redundant bf16 -> fp8 quantize.
+_grad_fp8_mailbox:dict = {}
 
 @functools.cache
-def _custom_fused_bwd_w13(grad_xw13:UOp, xw13:UOp, grad_x2:UOp, amax_state:UOp, dname:str) -> UOp:
+def _custom_fused_bwd_w13(grad_xw13:UOp, grad_xw13_fp8:UOp, grad_amax_buf:UOp,
+                          xw13:UOp, grad_x2:UOp, amax_state:UOp, grad_amax_state:UOp, dname:str) -> UOp:
   hidden = xw13.shape[2] // 2
   n_elems = xw13.shape[0] * xw13.shape[1] * hidden
   threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(NUM_WG, "gidx0")
-  mem = n_elems * 2 * 5
-  sink = UOp.sink(grad_xw13.base, xw13.base, grad_x2.base, amax_state.base, threads, workgroups,
-                  arg=KernelInfo(f"fused_silu_mul_bwd_w13_{n_elems}", estimates=Estimates(ops=8*n_elems, mem=mem)))
+  mem = n_elems * 2 * 5 + n_elems * 2 + NUM_WG * 4 + 4
+  sink = UOp.sink(grad_xw13.base, grad_xw13_fp8.base, grad_amax_buf.base,
+                  xw13.base, grad_x2.base, amax_state.base, grad_amax_state.base, threads, workgroups,
+                  arg=KernelInfo(f"fused_silu_mul_bwd_w13_{n_elems}", estimates=Estimates(ops=10*n_elems, mem=mem)))
   src, lib = compile_cpp(pathlib.Path(__file__).parent, "cast_amax_bwd_w13.cpp", n_elems, hidden)
   return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
                                UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=lib)))
 
 @functools.cache
-def _custom_fused_cast_amax_w13(fp8_out:UOp, amax_buf:UOp, xw13:UOp, amax_state:UOp, dname:str) -> UOp:
+def _custom_fused_cast_amax_w13(fp8_out:UOp, amax_buf:UOp, xw13:UOp, amax_state:UOp, grad_amax_state:UOp, dname:str) -> UOp:
+  # NOTE: grad_amax_state is plumbed through as an unused fwd input so the bwd kernel can read it via kernel.src
   hidden = xw13.shape[2] // 2
   n_elems = xw13.shape[0] * xw13.shape[1] * hidden
   threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(NUM_WG, "gidx0")
-  mem = n_elems * 2 * 2 + n_elems + NUM_WG * 2
+  mem = n_elems * 2 * 2 + n_elems + NUM_WG * 4
   sink = UOp.sink(fp8_out.base, amax_buf.base, xw13.base, amax_state.base, threads, workgroups,
                   arg=KernelInfo(f"fused_silu_mul_cast_amax_w13_{n_elems}", estimates=Estimates(ops=5*n_elems, mem=mem)))
   src, lib = compile_cpp(pathlib.Path(__file__).parent, "cast_amax_fwd_w13.cpp", n_elems, hidden)
@@ -30,44 +38,39 @@ def _custom_fused_cast_amax_w13(fp8_out:UOp, amax_buf:UOp, xw13:UOp, amax_state:
                                UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=lib)))
 
 def _fused_quantize_bwd_w13(gradient:UOp, kernel:UOp):
-  # NOTE: inputs are (fp8_out, amax_buf, xw13, amax_state); grad for xw13 only
-  _, _, xw13, amax_state = kernel.src[1:]
+  _, _, xw13, amax_state, grad_amax_state = kernel.src[1:]
   device = xw13.device
-  if isinstance(device, tuple):
-    axis, ndev = xw13.axis, len(device)
-    assert axis in (0, 1), f"unsupported sharding axis={axis}"
-    grad_xw13 = Tensor(Tensor.invalids(*shard_shape(xw13.shape, axis, ndev), dtype=dtypes.bfloat16,
-                                        device=device).uop.multi(axis), device=device)
-    dname = device[0].split(":")[0]
-  else:
-    grad_xw13 = Tensor.invalids(*xw13.shape, dtype=dtypes.bfloat16, device=device)
-    dname = device.split(":")[0] if isinstance(device, str) else device
-  grad_x2_t = Tensor(gradient, device=device).cast(dtypes.bfloat16)
-  fxn = functools.partial(_custom_fused_bwd_w13, dname=dname)
-  grad_xw13, *_ = Tensor.custom_kernel(grad_xw13, Tensor(xw13, device=device), grad_x2_t,
-                                        Tensor(amax_state, device=device), fxn=fxn)
-  return (None, None, grad_xw13.uop, None)
+  axis = xw13.axis if isinstance(device, tuple) else None
+  if isinstance(device, tuple): assert axis in (0, 1), f"unsupported sharding axis={axis}"
+  grad_xw13     = alloc_like(xw13.shape, dtypes.bfloat16, device, axis)
+  grad_xw13_fp8 = alloc_like(xw13.shape, dtypes.fp8e4m3,  device, axis)
+  grad_amax_buf = alloc_local((NUM_WG,), dtypes.float32,  device)
+  grad_amax_state_t = Tensor(grad_amax_state, device=device)
+  fxn = functools.partial(_custom_fused_bwd_w13, dname=dname_of(device))
+  grad_xw13, grad_xw13_fp8, grad_amax_buf, *_ = Tensor.custom_kernel(
+    grad_xw13, grad_xw13_fp8, grad_amax_buf,
+    Tensor(xw13, device=device), Tensor(gradient, device=device).cast(dtypes.bfloat16),
+    Tensor(amax_state, device=device), grad_amax_state_t, fxn=fxn)
+  inv_scale = (grad_amax_state_t.float() + 1e-8) / FP8_MAX
+  new_grad_amax = scalar_amax(grad_amax_buf)
+  store_effect = grad_amax_state_t.uop.store(new_grad_amax.uop)
+  # Stash fp8 companion + amax store for cdna_asm_gemm's bwd to attach to grad_a.
+  _grad_fp8_mailbox[grad_xw13.uop] = (grad_xw13_fp8.uop, inv_scale.uop, new_grad_amax.uop, store_effect)
+  return (None, None, grad_xw13.uop, None, None)
 
-def fused_quantize_fp8_w13(xw13:Tensor, amax_state:Tensor, fp8_dtype) -> tuple[Tensor, Tensor, Tensor]:
+def fused_quantize_fp8_w13(xw13:Tensor, amax_state:Tensor, fp8_dtype, grad_amax_state:Tensor) -> tuple[Tensor, Tensor, Tensor]:
   # NOTE: silu(xw1)*xw3 -> fp8 + amax over fused xw13 layout. Returns (fp8, inv_scale, new_amax)
+  # grad_amax_state: delayed amax for grad_xw13 fp8 quantization in the backward.
   assert xw13.dtype == dtypes.bfloat16, f"expected bf16, got {xw13.dtype}"
   MBS, SEQ, H2 = xw13.shape
   assert H2 % 2 == 0, f"w13 last-axis must be even, got {H2}"
   HIDDEN = H2 // 2
-  if isinstance(xw13.device, tuple):
-    axis, ndev = xw13.uop.axis, len(xw13.device)
-    assert axis in (0, 1), f"unsupported sharding axis={axis}"
-    fp8_out = Tensor(Tensor.invalids(*shard_shape((MBS, SEQ, HIDDEN), axis, ndev), dtype=fp8_dtype,
-                                     device=xw13.device).uop.multi(axis), device=xw13.device)
-    amax_buf = Tensor(Tensor.invalids(NUM_WG, dtype=dtypes.bfloat16, device=xw13.device).uop.multi(0),
-                      device=xw13.device)
-    dname = xw13.device[0].split(":")[0]
-  else:
-    fp8_out = Tensor.invalids(MBS, SEQ, HIDDEN, dtype=fp8_dtype, device=xw13.device)
-    amax_buf = Tensor.invalids(NUM_WG, dtype=dtypes.bfloat16, device=xw13.device)
-    dname = xw13.device.split(":")[0] if isinstance(xw13.device, str) else xw13.device
-  fxn = functools.partial(_custom_fused_cast_amax_w13, dname=dname)
-  fp8_out, amax_buf, *_ = Tensor.custom_kernel(fp8_out, amax_buf, xw13, amax_state, fxn=fxn,
-                                                grad_fxn=_fused_quantize_bwd_w13)
+  axis = xw13.uop.axis if isinstance(xw13.device, tuple) else None
+  if isinstance(xw13.device, tuple): assert axis in (0, 1), f"unsupported sharding axis={axis}"
+  fp8_out  = alloc_like((MBS, SEQ, HIDDEN), fp8_dtype,      xw13.device, axis)
+  amax_buf = alloc_local((NUM_WG,),         dtypes.float32, xw13.device)
+  fxn = functools.partial(_custom_fused_cast_amax_w13, dname=dname_of(xw13.device))
+  fp8_out, amax_buf, *_ = Tensor.custom_kernel(fp8_out, amax_buf, xw13, amax_state, grad_amax_state,
+                                                fxn=fxn, grad_fxn=_fused_quantize_bwd_w13)
   inv_scale = (amax_state.float() + 1e-8) / FP8_MAX
   return fp8_out, inv_scale, scalar_amax(amax_buf)

extra/llama_kernels/cast_amax/cast_amax_bwd_w13.cpp

@@ -1,5 +1,6 @@
 #include <hip/hip_runtime.h>
 #include <hip/hip_bf16.h>
+#include <hip/hip_fp8.h>
 
 #ifndef N_ELEMS
 #define N_ELEMS 234881024
@@ -20,19 +21,32 @@ constexpr float FP8_MAX = 448.0f;
 static_assert(N_ELEMS % VEC == 0, "N_ELEMS must be divisible by VEC");
 static_assert(HIDDEN % VEC == 0, "HIDDEN must be divisible by VEC");
 
+// fused silu*mul backward, three outputs in a single HBM pass:
+//   1) bf16 grad_xw13      — consumed by downstream bf16 autograd chain
+//   2) fp8  grad_xw13_fp8  — delayed-scale quantize using grad_amax_state (mailbox to matmul bwd)
+//   3) fp32 grad_amax_buf  — per-WG partial |grad_xw13|, reduced into next step's grad_amax_state
+// grad_amax_state is read for the fp8 scale. The store of new_grad_amax into grad_amax_state's
+// buffer is built in Python as a separate effect and threaded into grad_a via .after(store).
 extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
 fused_silu_mul_bwd_w13(
-    __hip_bfloat16*       __restrict__ grad_xw13_out,    // bf16, 2*N_ELEMS (interleaved layout)
-    const __hip_bfloat16* __restrict__ xw13,             // bf16, 2*N_ELEMS (interleaved)
-    const __hip_bfloat16* __restrict__ grad_x2,          // bf16, N_ELEMS
-    const __hip_bfloat16* __restrict__ amax_state)       // bf16 scalar
+    __hip_bfloat16*       __restrict__ grad_xw13_out,        // bf16, 2*N_ELEMS
+    __hip_fp8_storage_t*  __restrict__ grad_xw13_fp8_out,    // fp8,  2*N_ELEMS
+    float*                __restrict__ grad_amax_buf,        // fp32, NUM_WG per-WG partials
+    const __hip_bfloat16* __restrict__ xw13,                 // bf16, 2*N_ELEMS
+    const __hip_bfloat16* __restrict__ grad_x2,              // bf16, N_ELEMS
+    const float*          __restrict__ amax_state,           // fp32 scalar (fwd x2 amax)
+    const float*          __restrict__ grad_amax_state)      // fp32 scalar (delayed grad amax)
 {
+  __shared__ float sdata[THREADS_PER_WG];
+
   const int tid = threadIdx.x;
   const int wg  = blockIdx.x;
   const int gid = wg * THREADS_PER_WG + tid;
   const int stride_elems = NUM_WG * THREADS_PER_WG * VEC;
 
   const float scale = FP8_MAX / (static_cast<float>(*amax_state) + 1e-8f);
+  const float g_scale = FP8_MAX / (static_cast<float>(*grad_amax_state) + 1e-8f);
+  float local_max = 0.0f;
 
   for (int base = gid * VEC; base < N_ELEMS; base += stride_elems) {
     const int outer = base / HIDDEN;
@@ -49,6 +63,7 @@ fused_silu_mul_bwd_w13(
     const __hip_bfloat16 *gv = reinterpret_cast<const __hip_bfloat16*>(&g_raw);
 
     __hip_bfloat16 out1[VEC], out3[VEC];
+    __hip_fp8_storage_t fp8_1[VEC], fp8_3[VEC];
     #pragma unroll
     for (int i = 0; i < VEC; i++) {
       const float f1 = static_cast<float>(x1[i]);
@@ -58,11 +73,26 @@ fused_silu_mul_bwd_w13(
       const float silu = f1 * sig;
       const float silu_prime = sig + silu * (1.0f - sig);
       const float gs = fg * scale;
-      out1[i] = static_cast<__hip_bfloat16>(gs * silu_prime * f3);
-      out3[i] = static_cast<__hip_bfloat16>(gs * silu);
+      const float g1 = gs * silu_prime * f3;
+      const float g3 = gs * silu;
+      out1[i] = static_cast<__hip_bfloat16>(g1);
+      out3[i] = static_cast<__hip_bfloat16>(g3);
+      local_max = fmaxf(local_max, fmaxf(fabsf(g1), fabsf(g3)));
+      fp8_1[i] = __hip_cvt_float_to_fp8(fmaxf(-FP8_MAX, fminf(FP8_MAX, g1 * g_scale)), __HIP_SATFINITE, __HIP_E4M3);
+      fp8_3[i] = __hip_cvt_float_to_fp8(fmaxf(-FP8_MAX, fminf(FP8_MAX, g3 * g_scale)), __HIP_SATFINITE, __HIP_E4M3);
     }
 
     *reinterpret_cast<float4*>(&grad_xw13_out[xw1_off]) = *reinterpret_cast<float4*>(out1);
     *reinterpret_cast<float4*>(&grad_xw13_out[xw3_off]) = *reinterpret_cast<float4*>(out3);
+    *reinterpret_cast<uint64_t*>(&grad_xw13_fp8_out[xw1_off]) = *reinterpret_cast<uint64_t*>(fp8_1);
+    *reinterpret_cast<uint64_t*>(&grad_xw13_fp8_out[xw3_off]) = *reinterpret_cast<uint64_t*>(fp8_3);
   }
+
+  sdata[tid] = local_max;
+  __syncthreads();
+  for (int s = THREADS_PER_WG / 2; s > 0; s >>= 1) {
+    if (tid < s) sdata[tid] = fmaxf(sdata[tid], sdata[tid + s]);
+    __syncthreads();
+  }
+  if (tid == 0) grad_amax_buf[wg] = sdata[0];
 }

extra/llama_kernels/cast_amax/cast_amax_fwd_w13.cpp

@@ -24,9 +24,9 @@ static_assert(HIDDEN % VEC == 0, "HIDDEN must be divisible by VEC (so VEC loads
 extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
 fused_silu_mul_cast_amax_w13(
     __hip_fp8_storage_t*  __restrict__ fp8_out,         // fp8, N_ELEMS
-    __hip_bfloat16*       __restrict__ amax_buf,        // bf16, NUM_WG (per-WG amaxes)
+    float*                __restrict__ amax_buf,        // fp32, NUM_WG (per-WG amaxes)
     const __hip_bfloat16* __restrict__ xw13,            // bf16, 2*N_ELEMS
-    const __hip_bfloat16* __restrict__ amax_state)      // bf16 scalar
+    const float*          __restrict__ amax_state)      // fp32 scalar
 {
   __shared__ float sdata[THREADS_PER_WG];
 
@@ -75,5 +75,5 @@ fused_silu_mul_cast_amax_w13(
     __syncthreads();
   }
 
-  if (tid == 0) amax_buf[wg] = static_cast<__hip_bfloat16>(sdata[0]);
+  if (tid == 0) amax_buf[wg] = sdata[0];
 }

extra/llama_kernels/fused_ce/__init__.py

@@ -42,8 +42,7 @@ def _fused_ce_loss_bwd(gradient:UOp, kernel:UOp, label_smoothing:float):
   # gradient is the upstream grad w.r.t. per-row loss (shape: (rows,) fp32)
   _, _, lse_u, logits_u, targets_u = kernel.src[1:]
   device = logits_u.device
-  rows_vocab = logits_u.shape  # (rows, VOCAB) after reshape
-  rows, VOCAB = rows_vocab
+  rows, VOCAB = logits_u.shape  # (rows, VOCAB) after reshape
   if isinstance(device, tuple):
     axis = logits_u.axis
     ndev = len(device)
@@ -54,9 +53,8 @@ def _fused_ce_loss_bwd(gradient:UOp, kernel:UOp, label_smoothing:float):
     d_logits = Tensor.invalids(rows, VOCAB, dtype=dtypes.bfloat16, device=device)
     dname = device.split(":")[0] if isinstance(device, str) else device
     rows_per_dev = rows
-  grad_t = Tensor(gradient, device=device).float().reshape(-1)  # (rows,) fp32
   # NOTE: .mean() backward gives same grad per row (1/N), so broadcast is safe; take scalar
-  scale = grad_t[0:1].contiguous()
+  scale = Tensor(gradient, device=device).float().reshape(-1)[0:1].contiguous()
   logits_t = Tensor(logits_u.after(kernel), device=device)
   lse_t = Tensor(lse_u.after(kernel), device=device)
   targets_t = Tensor(targets_u, device=device)

extra/llama_kernels/fused_mul_quantize_fp8/__init__.py

@@ -1,54 +0,0 @@
-from __future__ import annotations
-import functools, pathlib
-from tinygrad import Tensor, dtypes
-from tinygrad.uop.ops import UOp, Ops, KernelInfo
-from tinygrad.renderer import Estimates
-from extra.llama_kernels import FP8_MAX, NUM_WG, THREADS_PER_WG, compile_cpp, shard_shape, scalar_amax
-
-@functools.cache
-def _custom_mul_quantize_fp8(fp8_out:UOp, amax_buf:UOp, x:UOp, weight:UOp, amax_state:UOp, dname:str) -> UOp:
-  MBS, SEQ, HIDDEN = x.shape
-  n_elems = MBS * SEQ * HIDDEN
-  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(NUM_WG, "gidx0")
-  mem = n_elems * 2 + HIDDEN * 2 + n_elems + NUM_WG * 2
-  sink = UOp.sink(fp8_out.base, amax_buf.base, x.base, weight.base, amax_state.base, threads, workgroups,
-                  arg=KernelInfo(f"fused_mul_quantize_fp8_{n_elems}_h{HIDDEN}", estimates=Estimates(ops=3*n_elems, mem=mem)))
-  src, lib = compile_cpp(pathlib.Path(__file__).parent, "fused_mul_quantize_fp8.cpp", n_elems, HIDDEN)
-  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
-                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=lib)))
-
-def _fused_mul_quantize_fp8_bwd(gradient:UOp, kernel:UOp):
-  # NOTE: inputs are (fp8_out, amax_buf, x, weight, amax_state); grads for x and weight
-  _, _, x_u, weight_u, amax_state_u = kernel.src[1:]
-  device = x_u.device
-  grad_t = Tensor(gradient, device=device).cast(dtypes.bfloat16)
-  x_t, weight_t = Tensor(x_u, device=device), Tensor(weight_u, device=device)
-  scale = FP8_MAX / (Tensor(amax_state_u, device=device).float() + 1e-8)
-  grad_scaled = grad_t.float() * scale
-  # NOTE: grad_x stays bf16 to avoid CSE materializing a (MBS, SEQ, HIDDEN) fp32 intermediate
-  grad_x = (grad_scaled * weight_t.float()).cast(dtypes.bfloat16)
-  grad_weight = (grad_scaled * x_t.float()).sum(axis=(0, 1)).cast(dtypes.bfloat16)
-  return (None, None, grad_x.uop, grad_weight.uop, None)
-
-def fused_mul_quantize_fp8(x:Tensor, weight:Tensor, amax_state:Tensor, fp8_dtype) -> tuple[Tensor, Tensor, Tensor]:
-  # NOTE: (x * weight) -> fp8 + amax, delayed scaling. Returns (fp8, inv_scale, new_amax)
-  assert x.dtype == dtypes.bfloat16 and weight.dtype == dtypes.bfloat16
-  assert x.shape[-1] == weight.shape[-1], f"HIDDEN mismatch: x={x.shape}, weight={weight.shape}"
-  MBS, SEQ, HIDDEN = x.shape
-  if isinstance(x.device, tuple):
-    axis, ndev = x.uop.axis, len(x.device)
-    assert axis in (0, 1), f"unsupported sharding axis={axis}"
-    fp8_out = Tensor(Tensor.invalids(*shard_shape((MBS, SEQ, HIDDEN), axis, ndev), dtype=fp8_dtype,
-                                     device=x.device).uop.multi(axis), device=x.device)
-    amax_buf = Tensor(Tensor.invalids(NUM_WG, dtype=dtypes.bfloat16, device=x.device).uop.multi(0), device=x.device)
-    dname = x.device[0].split(":")[0]
-  else:
-    fp8_out = Tensor.invalids(MBS, SEQ, HIDDEN, dtype=fp8_dtype, device=x.device)
-    amax_buf = Tensor.invalids(NUM_WG, dtype=dtypes.bfloat16, device=x.device)
-    dname = x.device.split(":")[0] if isinstance(x.device, str) else x.device
-  fxn = functools.partial(_custom_mul_quantize_fp8, dname=dname)
-  fp8_out, amax_buf, *_ = Tensor.custom_kernel(fp8_out, amax_buf, x, weight, amax_state, fxn=fxn,
-                                                grad_fxn=_fused_mul_quantize_fp8_bwd)
-  new_amax = scalar_amax(amax_buf)
-  inv_scale = (amax_state.float() + 1e-8) / FP8_MAX
-  return fp8_out, inv_scale, new_amax

extra/llama_kernels/fused_pad_grad_accum/__init__.py

@@ -0,0 +1,55 @@
+from __future__ import annotations
+import functools, pathlib
+from tinygrad import Tensor, dtypes
+from tinygrad.uop.ops import UOp, Ops, KernelInfo
+from tinygrad.renderer import Estimates
+from extra.llama_kernels import THREADS_PER_WG, dname_of, compile_hip
+
+ELEMS_PER_THREAD = 8  # vectorized 16-byte load (uint4 = 8 bf16)
+
+def _build_src(n_chunks:int) -> str:
+  template = (pathlib.Path(__file__).parent/"fused_pad_grad_accum.cpp").read_text()
+  params = "".join(f",\n    const __hip_bfloat16* __restrict__ chunk{i}" for i in range(n_chunks))
+  dispatch = "\n    ".join(f"case {i}: chunk_ptr = chunk{i}; break;" for i in range(n_chunks))
+  return (template.replace("__FUSED_PAD_GRAD_ACCUM_PARAMS", params)
+                  .replace("__FUSED_PAD_GRAD_ACCUM_DISPATCH", dispatch))
+
+@functools.cache
+def _custom_fused_pad_grad_accum(grad_buf:UOp, *chunk_uops, dname:str, n_chunks:int, chunk_size:int) -> UOp:
+  total = n_chunks * chunk_size
+  elems_per_block = THREADS_PER_WG * ELEMS_PER_THREAD
+  assert chunk_size % elems_per_block == 0, f"chunk_size {chunk_size} must be multiple of {elems_per_block}"
+  num_wg = total // elems_per_block
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(num_wg, "gidx0")
+  mem = total * 2 * 3
+  sink = UOp.sink(grad_buf.base, *(c.base for c in chunk_uops), threads, workgroups,
+                  arg=KernelInfo(f"fused_pad_grad_accum_n{n_chunks}_c{chunk_size}",
+                                 estimates=Estimates(ops=2*total, mem=mem)))
+  src = _build_src(n_chunks)
+  defines = [f"-DCHUNK_SIZE={chunk_size}", f"-DTHREADS_PER_WG={THREADS_PER_WG}", f"-DELEMS_PER_THREAD={ELEMS_PER_THREAD}"]
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=compile_hip(src, defines))))
+
+def can_fused_pad_grad_accum(grad_buf:Tensor, chunks:list[Tensor]) -> bool:
+  if not chunks or grad_buf.dtype != dtypes.bfloat16: return False
+  if any(c.dtype != dtypes.bfloat16 for c in chunks): return False
+  chunk_shape = chunks[0].shape
+  if any(c.shape != chunk_shape for c in chunks): return False
+  chunk_size, total = 1, 1
+  for d in chunk_shape: chunk_size *= d
+  for d in grad_buf.shape: total *= d
+  return total == len(chunks) * chunk_size and chunk_size % (THREADS_PER_WG * ELEMS_PER_THREAD) == 0
+
+def fused_pad_grad_accum(grad_buf:Tensor, chunks:list[Tensor]) -> Tensor:
+  # NOTE: grad_buf += cat(*chunks, dim=0) in one HBM pass (in-place add). Returns new grad_buf Tensor.
+  # Requires uniform chunk shapes and chunk_size % (THREADS_PER_WG*ELEMS_PER_THREAD) == 0.
+  assert chunks and grad_buf.dtype == dtypes.bfloat16
+  for c in chunks: assert c.dtype == dtypes.bfloat16, f"chunk dtype must be bf16, got {c.dtype}"
+  chunk_size, total = 1, 1
+  for d in chunks[0].shape: chunk_size *= d
+  for d in grad_buf.shape: total *= d
+  assert total == len(chunks) * chunk_size, f"grad_buf size {total} != n_chunks {len(chunks)} * chunk_size {chunk_size}"
+  fxn = functools.partial(_custom_fused_pad_grad_accum, dname=dname_of(grad_buf.device),
+                          n_chunks=len(chunks), chunk_size=chunk_size)
+  out, *_ = Tensor.custom_kernel(grad_buf, *chunks, fxn=fxn)
+  return out

extra/llama_kernels/fused_pad_grad_accum/fused_pad_grad_accum.cpp

@@ -0,0 +1,63 @@
+// Fused custom kernel: grad_buf += cat(*chunks, dim=0) in one HBM pass.
+//
+// Template source — chunk parameter list and switch dispatch are filled by codegen
+// in cast_amax.py:_build_fused_pad_grad_accum_src to support arbitrary N.
+//
+// Defines required at compile time:
+//   CHUNK_SIZE       elements per chunk (must be multiple of THREADS_PER_WG * ELEMS_PER_THREAD)
+//   THREADS_PER_WG
+//   ELEMS_PER_THREAD (8 = one uint4 per thread = 16-byte vectorized load)
+//
+// Layout: one block-per-(slice-of-chunk) — blockIdx.x / BLOCKS_PER_CHUNK selects the chunk.
+// All threads in a block read the same chunk → switch is uniform → no warp divergence.
+
+#include <hip/hip_runtime.h>
+#include <hip/hip_bf16.h>
+
+#ifndef THREADS_PER_WG
+#define THREADS_PER_WG 256
+#endif
+#ifndef ELEMS_PER_THREAD
+#define ELEMS_PER_THREAD 8
+#endif
+
+#define ELEMS_PER_BLOCK (THREADS_PER_WG * ELEMS_PER_THREAD)
+#define BLOCKS_PER_CHUNK (CHUNK_SIZE / ELEMS_PER_BLOCK)
+
+extern "C" __attribute__((global))
+__attribute__((amdgpu_flat_work_group_size(1, THREADS_PER_WG)))
+void fused_pad_grad_accum(
+    __hip_bfloat16* __restrict__ grad_buf
+    __FUSED_PAD_GRAD_ACCUM_PARAMS
+) {
+  const int bid = blockIdx.x;
+  const int chunk_idx = bid / BLOCKS_PER_CHUNK;
+  const int block_in_chunk = bid - chunk_idx * BLOCKS_PER_CHUNK;
+  const int tid = threadIdx.x;
+
+  const __hip_bfloat16* chunk_ptr;
+  switch (chunk_idx) {
+    __FUSED_PAD_GRAD_ACCUM_DISPATCH
+    default: chunk_ptr = (const __hip_bfloat16*)0; break;  // unreachable
+  }
+
+  // int64 for global_offset: at 32 chunks × 117M elements = 3.6B, int32 overflows → MEMVIOL.
+  const int local_offset = block_in_chunk * ELEMS_PER_BLOCK + tid * ELEMS_PER_THREAD;
+  const long long global_offset = (long long)chunk_idx * (long long)CHUNK_SIZE + (long long)local_offset;
+
+  // Vectorized 16-byte load (uint4 = 8 bf16). Requires CHUNK_SIZE % 8 == 0 and 16-byte alignment.
+  const uint4 chunk_v = *reinterpret_cast<const uint4*>(&chunk_ptr[local_offset]);
+  const uint4 grad_v  = *reinterpret_cast<const uint4*>(&grad_buf[global_offset]);
+  uint4 out_v;
+
+  const __hip_bfloat16* chunk_bf = reinterpret_cast<const __hip_bfloat16*>(&chunk_v);
+  const __hip_bfloat16* grad_bf  = reinterpret_cast<const __hip_bfloat16*>(&grad_v);
+  __hip_bfloat16*       out_bf   = reinterpret_cast<__hip_bfloat16*>(&out_v);
+
+  #pragma unroll
+  for (int i = 0; i < ELEMS_PER_THREAD; i++) {
+    out_bf[i] = (__hip_bfloat16)((float)grad_bf[i] + (float)chunk_bf[i]);
+  }
+
+  *reinterpret_cast<uint4*>(&grad_buf[global_offset]) = out_v;
+}

extra/llama_kernels/fused_rmsnorm_mul_quantize_fp8/__init__.py

@@ -0,0 +1,153 @@
+from __future__ import annotations
+import functools, pathlib
+from tinygrad import Tensor, dtypes
+from tinygrad.uop.ops import UOp, Ops, KernelInfo
+from tinygrad.renderer import Estimates
+from extra.llama_kernels import FP8_MAX, NUM_WG, THREADS_PER_WG, alloc_like, alloc_local, scalar_amax, dname_of, compile_hip
+
+def _src() -> str: return (pathlib.Path(__file__).parent/"fused_rmsnorm_mul_quantize_fp8.cpp").read_text()
+def _src_bwd() -> str: return (pathlib.Path(__file__).parent/"fused_rmsnorm_mul_quantize_fp8_bwd.cpp").read_text()
+
+@functools.cache
+def _custom_fwd(fp8_out:UOp, x_normed_out:UOp, rrms_out:UOp, amax_buf:UOp,
+                x:UOp, weight:UOp, amax_state:UOp, dname:str, eps_val:float) -> UOp:
+  MBS, SEQ, HIDDEN = x.shape
+  n_elems = MBS * SEQ * HIDDEN
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(NUM_WG, "gidx0")
+  mem = n_elems * 2 + n_elems + MBS * SEQ * 4 + n_elems + HIDDEN * 2 + NUM_WG * 4 + 4
+  sink = UOp.sink(fp8_out.base, x_normed_out.base, rrms_out.base, amax_buf.base,
+                  x.base, weight.base, amax_state.base, threads, workgroups,
+                  arg=KernelInfo(f"fused_rmsnorm_mul_quantize_fp8_{n_elems}_h{HIDDEN}_eps{eps_val:.0e}",
+                                 estimates=Estimates(ops=6*n_elems, mem=mem)))
+  defines = [f"-DN_ELEMS={n_elems}", f"-DHIDDEN={HIDDEN}", f"-DNUM_WG={NUM_WG}", f"-DTHREADS_PER_WG={THREADS_PER_WG}",
+             f"-DEPS_LITERAL={eps_val}f"]
+  src = _src()
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=compile_hip(src, defines))))
+
+@functools.cache
+def _custom_fwd_add(fp8_out:UOp, h_out:UOp, x_normed_out:UOp, rrms_out:UOp, amax_buf:UOp,
+                    x:UOp, residual:UOp, weight:UOp, amax_state:UOp, dname:str, eps_val:float) -> UOp:
+  MBS, SEQ, HIDDEN = x.shape
+  n_elems = MBS * SEQ * HIDDEN
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(NUM_WG, "gidx0")
+  mem = n_elems * 2 * 4 + MBS * SEQ * 4 + HIDDEN * 2 + NUM_WG * 4 + 4
+  sink = UOp.sink(fp8_out.base, h_out.base, x_normed_out.base, rrms_out.base, amax_buf.base,
+                  x.base, residual.base, weight.base, amax_state.base, threads, workgroups,
+                  arg=KernelInfo(f"fused_add_rmsnorm_mul_quantize_fp8_{n_elems}_h{HIDDEN}_eps{eps_val:.0e}",
+                                 estimates=Estimates(ops=7*n_elems, mem=mem)))
+  defines = [f"-DN_ELEMS={n_elems}", f"-DHIDDEN={HIDDEN}", f"-DNUM_WG={NUM_WG}", f"-DTHREADS_PER_WG={THREADS_PER_WG}",
+             f"-DEPS_LITERAL={eps_val}f", f"-DHAS_RESIDUAL=1"]
+  src = _src()
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=compile_hip(src, defines))))
+
+@functools.cache
+def _custom_bwd(grad_x:UOp, grad_weight_partial:UOp,
+                grad_fp8:UOp, x_normed:UOp, rrms:UOp, weight:UOp, amax_state:UOp, dname:str) -> UOp:
+  MBS, SEQ, HIDDEN = x_normed.shape
+  n_elems = MBS * SEQ * HIDDEN
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(NUM_WG, "gidx0")
+  mem = n_elems * 2 * 3 + NUM_WG * HIDDEN * 4 + MBS * SEQ * 4 + HIDDEN * 2 + 4
+  sink = UOp.sink(grad_x.base, grad_weight_partial.base,
+                  grad_fp8.base, x_normed.base, rrms.base, weight.base, amax_state.base, threads, workgroups,
+                  arg=KernelInfo(f"fused_rmsnorm_mul_quantize_fp8_bwd_{n_elems}_h{HIDDEN}",
+                                 estimates=Estimates(ops=8*n_elems, mem=mem)))
+  defines = [f"-DN_ELEMS={n_elems}", f"-DHIDDEN={HIDDEN}", f"-DNUM_WG={NUM_WG}", f"-DTHREADS_PER_WG={THREADS_PER_WG}"]
+  src = _src_bwd()
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=compile_hip(src, defines))))
+
+def _bwd_common(fp8_grad_u, h_grad_u, x_u, x_normed_u, rrms_u, weight_u, amax_state_u, kernel:UOp):
+  device = x_u.device
+  MBS, SEQ, HIDDEN = x_normed_u.shape
+  axis = x_normed_u.axis if isinstance(device, tuple) else None
+  grad_x = alloc_like((MBS, SEQ, HIDDEN), dtypes.bfloat16, device, axis)
+  grad_weight_partial = alloc_local((NUM_WG, HIDDEN), dtypes.float32, device)
+  grad_h_from_fp8 = None
+  grad_weight_uop = None
+  if fp8_grad_u is not None:
+    fxn = functools.partial(_custom_bwd, dname=dname_of(device))
+    grad_x_t, grad_weight_partial_t, *_ = Tensor.custom_kernel(
+      grad_x, grad_weight_partial,
+      Tensor(fp8_grad_u, device=device).cast(dtypes.bfloat16),
+      Tensor(x_normed_u.after(kernel), device=device),
+      Tensor(rrms_u.after(kernel), device=device),
+      Tensor(weight_u, device=device),
+      Tensor(amax_state_u, device=device), fxn=fxn)
+    grad_h_from_fp8 = grad_x_t
+    grad_weight_uop = grad_weight_partial_t.sum(axis=0).cast(dtypes.bfloat16).uop
+  if h_grad_u is not None:
+    h_grad_t = Tensor(h_grad_u, device=device).cast(dtypes.bfloat16)
+    grad_total = (grad_h_from_fp8 + h_grad_t) if grad_h_from_fp8 is not None else h_grad_t
+  else:
+    grad_total = grad_h_from_fp8
+  return grad_total.uop, grad_weight_uop
+
+def _fused_bwd(gradient:UOp, kernel:UOp):
+  # NOTE: fwd inputs (fp8_out, x_normed_out, rrms_out, amax_buf, x, weight, amax_state)
+  _, x_normed_u, rrms_u, _, x_u, weight_u, amax_state_u = kernel.src[1:]
+  grad_x, grad_w = _bwd_common(gradient, None, x_u, x_normed_u, rrms_u, weight_u, amax_state_u, kernel)
+  return (None, None, None, None, grad_x, grad_w, None)
+
+def _fused_add_bwd(*args, **kwargs):
+  # Two invocation modes: 1 grad => positional; >1 grads => kwarg `call=`.
+  # Outputs: (fp8_out, h_out, x_normed_out, rrms_out, amax_buf). Both fp8 and h may be consumed
+  # downstream — TUPLE order in gradient.py preserves kernel-output slot order.
+  # Don't dispatch by dtype: matmul's bwd emits fp8 grad as bf16 (no explicit cast), so
+  # dtype-detection collapses both into h_grad and silently drops the rmsnorm-bwd path.
+  if 'call' in kwargs:
+    kernel, all_grads = kwargs['call'], list(args)
+  else:
+    gradient, kernel = args
+    all_grads = [gradient]
+  fp8_grad_u = h_grad_u = None
+  if len(all_grads) >= 2:
+    fp8_grad_u, h_grad_u = all_grads[0], all_grads[1]
+  elif len(all_grads) == 1:
+    g = all_grads[0]
+    if g.dtype == dtypes.bfloat16: h_grad_u = g
+    else: fp8_grad_u = g
+  _, _, x_normed_u, rrms_u, _, x_u, _, weight_u, amax_state_u = kernel.src[1:]
+  grad_h, grad_w = _bwd_common(fp8_grad_u, h_grad_u, x_u, x_normed_u, rrms_u, weight_u, amax_state_u, kernel)
+  return (None, None, None, None, None, grad_h, grad_h, grad_w, None)
+
+def fused_rmsnorm_mul_quantize_fp8(x:Tensor, weight:Tensor, amax_state:Tensor, eps:float, fp8_dtype) -> tuple[Tensor, Tensor, Tensor, Tensor, Tensor]:
+  # NOTE: rmsnorm(x) * weight -> fp8 + amax. Returns (fp8, inv_scale, new_amax, x_normed, rrms).
+  # x_normed + rrms are saved for the rmsnorm backward (also recomputed here from x regs).
+  assert x.dtype == dtypes.bfloat16 and weight.dtype == dtypes.bfloat16
+  assert x.shape[-1] == weight.shape[-1], f"HIDDEN mismatch: x={x.shape}, weight={weight.shape}"
+  MBS, SEQ, HIDDEN = x.shape
+  axis = x.uop.axis if isinstance(x.device, tuple) else None
+  if isinstance(x.device, tuple): assert axis in (0, 1), f"unsupported sharding axis={axis}"
+  fp8_out      = alloc_like((MBS, SEQ, HIDDEN), fp8_dtype,       x.device, axis)
+  x_normed_out = alloc_like((MBS, SEQ, HIDDEN), dtypes.bfloat16, x.device, axis)
+  rrms_out     = alloc_like((MBS, SEQ),         dtypes.float32,  x.device, axis)
+  amax_buf     = alloc_local((NUM_WG,),         dtypes.float32,  x.device)
+  fxn = functools.partial(_custom_fwd, dname=dname_of(x.device), eps_val=eps)
+  fp8_out, x_normed_out, rrms_out, amax_buf, *_ = Tensor.custom_kernel(
+    fp8_out, x_normed_out, rrms_out, amax_buf, x, weight, amax_state, fxn=fxn, grad_fxn=_fused_bwd)
+  inv_scale = (amax_state.float() + 1e-8) / FP8_MAX
+  return fp8_out, inv_scale, scalar_amax(amax_buf), x_normed_out, rrms_out
+
+def fused_add_rmsnorm_mul_quantize_fp8(x:Tensor, residual:Tensor, weight:Tensor, amax_state:Tensor,
+                                       eps:float, fp8_dtype) -> tuple[Tensor, Tensor, Tensor, Tensor, Tensor, Tensor]:
+  # NOTE: h = x + residual; y_normed = rmsnorm(h); fp8 = quantize(y_normed * weight).
+  # Returns (fp8, inv_scale, new_amax, h, x_normed, rrms). h is also written so downstream can
+  # reuse it without recomputing x+residual — eliminates the separate residual-add kernel.
+  assert x.dtype == dtypes.bfloat16 and residual.dtype == dtypes.bfloat16 and weight.dtype == dtypes.bfloat16
+  assert x.shape == residual.shape
+  MBS, SEQ, HIDDEN = x.shape
+  axis = x.uop.axis if isinstance(x.device, tuple) else None
+  if isinstance(x.device, tuple): assert axis in (0, 1), f"unsupported sharding axis={axis}"
+  fp8_out      = alloc_like((MBS, SEQ, HIDDEN), fp8_dtype,       x.device, axis)
+  h_out        = alloc_like((MBS, SEQ, HIDDEN), dtypes.bfloat16, x.device, axis)
+  x_normed_out = alloc_like((MBS, SEQ, HIDDEN), dtypes.bfloat16, x.device, axis)
+  rrms_out     = alloc_like((MBS, SEQ),         dtypes.float32,  x.device, axis)
+  amax_buf     = alloc_local((NUM_WG,),         dtypes.float32,  x.device)
+  fxn = functools.partial(_custom_fwd_add, dname=dname_of(x.device), eps_val=eps)
+  fp8_out, h_out, x_normed_out, rrms_out, amax_buf, *_ = Tensor.custom_kernel(
+    fp8_out, h_out, x_normed_out, rrms_out, amax_buf, x, residual, weight, amax_state,
+    fxn=fxn, grad_fxn=_fused_add_bwd)
+  inv_scale = (amax_state.float() + 1e-8) / FP8_MAX
+  return fp8_out, inv_scale, scalar_amax(amax_buf), h_out, x_normed_out, rrms_out

extra/llama_kernels/fused_rmsnorm_mul_quantize_fp8/fused_rmsnorm_mul_quantize_fp8.cpp

@@ -0,0 +1,155 @@
+#include <hip/hip_runtime.h>
+#include <hip/hip_bf16.h>
+#include <hip/hip_fp8.h>
+
+// Fuses the full pre-matmul preparation for a layer into a single HBM pass:
+//   y = rmsnorm(x) * weight            (reduce-mean-square + rsqrt + per-elem mul)
+//   fp8 = fp8_sat(y * (FP8_MAX / amax_state))
+// Also writes:
+//   rrms[row]        — saved for the rmsnorm backward
+//   amax_buf[wg]     — per-WG |y| partials, reduced later to update amax_state
+//
+// Layout: one WG per row, ROWS_PER_WG rows per WG via grid-stride (ROWS = N_ELEMS / HIDDEN).
+// Each thread handles HIDDEN / THREADS_PER_WG elements per row.
+
+#ifndef N_ELEMS
+#define N_ELEMS 67108864
+#endif
+#ifndef HIDDEN
+#define HIDDEN 4096
+#endif
+#ifndef NUM_WG
+#define NUM_WG 1024
+#endif
+#ifndef THREADS_PER_WG
+#define THREADS_PER_WG 256
+#endif
+#ifndef EPS_LITERAL
+#define EPS_LITERAL 1e-5f
+#endif
+#ifndef HAS_RESIDUAL
+#define HAS_RESIDUAL 0
+#endif
+
+constexpr int VEC = 8;
+constexpr float FP8_MAX = 448.0f;
+
+static_assert(N_ELEMS % HIDDEN == 0, "N_ELEMS must be a multiple of HIDDEN");
+static_assert(HIDDEN % (THREADS_PER_WG * VEC) == 0, "HIDDEN must be divisible by THREADS_PER_WG*VEC");
+
+constexpr int ROWS = N_ELEMS / HIDDEN;
+constexpr int ELEMS_PER_THREAD = HIDDEN / THREADS_PER_WG;  // each thread sees this many elems per row
+constexpr int VECS_PER_THREAD = ELEMS_PER_THREAD / VEC;    // number of 8-wide vec loads
+
+#if HAS_RESIDUAL
+extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
+fused_add_rmsnorm_mul_quantize_fp8(
+    __hip_fp8_storage_t*  __restrict__ fp8_out,         // fp8, ROWS*HIDDEN
+    __hip_bfloat16*       __restrict__ h_out,           // bf16, ROWS*HIDDEN — x + residual (saved for downstream)
+    __hip_bfloat16*       __restrict__ x_normed_out,    // bf16, ROWS*HIDDEN
+    float*                __restrict__ rrms_out,        // fp32, ROWS
+    float*                __restrict__ amax_buf,        // fp32, NUM_WG
+    const __hip_bfloat16* __restrict__ x,               // bf16, ROWS*HIDDEN
+    const __hip_bfloat16* __restrict__ residual,        // bf16, ROWS*HIDDEN — added into x before rmsnorm
+    const __hip_bfloat16* __restrict__ weight,          // bf16, HIDDEN
+    const float*          __restrict__ amax_state)      // fp32 scalar
+{
+#else
+extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
+fused_rmsnorm_mul_quantize_fp8(
+    __hip_fp8_storage_t*  __restrict__ fp8_out,         // fp8, ROWS*HIDDEN
+    __hip_bfloat16*       __restrict__ x_normed_out,    // bf16, ROWS*HIDDEN (saved for rmsnorm bwd)
+    float*                __restrict__ rrms_out,        // fp32, ROWS (fp32 to match rmsnorm_bwd.cpp expectation)
+    float*                __restrict__ amax_buf,        // fp32, NUM_WG per-WG partials
+    const __hip_bfloat16* __restrict__ x,               // bf16, ROWS*HIDDEN
+    const __hip_bfloat16* __restrict__ weight,          // bf16, HIDDEN (per-hidden scale)
+    const float*          __restrict__ amax_state)      // fp32 scalar
+{
+#endif
+  __shared__ float sdata[THREADS_PER_WG];
+
+  const int tid = threadIdx.x;
+  const int wg  = blockIdx.x;
+
+  const float scale = FP8_MAX / (static_cast<float>(*amax_state) + 1e-8f);
+  const float inv_hidden = 1.0f / static_cast<float>(HIDDEN);
+  float local_max = 0.0f;
+
+  // Grid-stride over rows. Each WG processes rows (wg, wg+NUM_WG, wg+2*NUM_WG, ...).
+  for (int row = wg; row < ROWS; row += NUM_WG) {
+    const int row_off = row * HIDDEN;
+
+    // Load row (+ residual if present) into registers.
+    float regs[ELEMS_PER_THREAD];
+    float sum_sq = 0.0f;
+    #pragma unroll
+    for (int v = 0; v < VECS_PER_THREAD; v++) {
+      const int h_base = tid * VEC + v * THREADS_PER_WG * VEC;
+      float4 raw = *reinterpret_cast<const float4*>(&x[row_off + h_base]);
+      const __hip_bfloat16 *xi = reinterpret_cast<const __hip_bfloat16*>(&raw);
+#if HAS_RESIDUAL
+      float4 res_raw = *reinterpret_cast<const float4*>(&residual[row_off + h_base]);
+      const __hip_bfloat16 *ri = reinterpret_cast<const __hip_bfloat16*>(&res_raw);
+      __hip_bfloat16 h_buf[VEC];
+#endif
+      #pragma unroll
+      for (int i = 0; i < VEC; i++) {
+#if HAS_RESIDUAL
+        const float f = static_cast<float>(xi[i]) + static_cast<float>(ri[i]);
+        h_buf[i] = static_cast<__hip_bfloat16>(f);
+#else
+        const float f = static_cast<float>(xi[i]);
+#endif
+        regs[v * VEC + i] = f;
+        sum_sq += f * f;
+      }
+#if HAS_RESIDUAL
+      *reinterpret_cast<float4*>(&h_out[row_off + h_base]) = *reinterpret_cast<float4*>(h_buf);
+#endif
+    }
+
+    // LDS tree-reduce sum_sq across the WG.
+    sdata[tid] = sum_sq;
+    __syncthreads();
+    for (int s = THREADS_PER_WG / 2; s > 0; s >>= 1) {
+      if (tid < s) sdata[tid] = sdata[tid] + sdata[tid + s];
+      __syncthreads();
+    }
+    const float mean_sq = sdata[0] * inv_hidden;
+    const float rrms = 1.0f / sqrtf(mean_sq + EPS_LITERAL);
+
+    if (tid == 0) rrms_out[row] = rrms;
+
+    // Normalize, multiply by weight, quantize. Also write x_normed (for rmsnorm bwd).
+    #pragma unroll
+    for (int v = 0; v < VECS_PER_THREAD; v++) {
+      const int h_base = tid * VEC + v * THREADS_PER_WG * VEC;
+      float4 w_raw = *reinterpret_cast<const float4*>(&weight[h_base]);
+      const __hip_bfloat16 *wi = reinterpret_cast<const __hip_bfloat16*>(&w_raw);
+
+      __hip_fp8_storage_t out[VEC];
+      __hip_bfloat16 xn[VEC];
+      #pragma unroll
+      for (int i = 0; i < VEC; i++) {
+        const float x_normed = regs[v * VEC + i] * rrms;
+        xn[i] = static_cast<__hip_bfloat16>(x_normed);
+        const float y = x_normed * static_cast<float>(wi[i]);
+        local_max = fmaxf(local_max, fabsf(y));
+        const float scaled = fmaxf(-FP8_MAX, fminf(FP8_MAX, y * scale));
+        out[i] = __hip_cvt_float_to_fp8(scaled, __HIP_SATFINITE, __HIP_E4M3);
+      }
+      *reinterpret_cast<uint64_t*>(&fp8_out[row_off + h_base]) = *reinterpret_cast<uint64_t*>(out);
+      *reinterpret_cast<float4*>(&x_normed_out[row_off + h_base]) = *reinterpret_cast<float4*>(xn);
+    }
+    __syncthreads();  // before next row's sum_sq reduce reuses sdata
+  }
+
+  // Final per-WG amax reduce.
+  sdata[tid] = local_max;
+  __syncthreads();
+  for (int s = THREADS_PER_WG / 2; s > 0; s >>= 1) {
+    if (tid < s) sdata[tid] = fmaxf(sdata[tid], sdata[tid + s]);
+    __syncthreads();
+  }
+  if (tid == 0) amax_buf[wg] = sdata[0];
+}

extra/llama_kernels/fused_rmsnorm_mul_quantize_fp8/fused_rmsnorm_mul_quantize_fp8_bwd.cpp

@@ -0,0 +1,147 @@
+#include <hip/hip_runtime.h>
+#include <hip/hip_bf16.h>
+
+// Full backward for fused_rmsnorm_mul_quantize_fp8.cpp. One HBM pass per row produces:
+//   grad_x (bf16)              — gradient w.r.t. pre-rmsnorm x
+//   grad_weight_partial (fp32) — per-WG partial of the weight gradient, reduced later
+//
+// Input (all read):
+//   grad_fp8 (bf16)     — upstream grad w.r.t. fp8_out (bf16-typed gradient value)
+//   x_normed (bf16)     — saved from the fwd kernel, shape (ROWS, HIDDEN)
+//   rrms (fp32)         — saved rrms per row
+//   weight (bf16)       — per-HIDDEN rmsnorm weight
+//   amax_state (bf16)   — delayed amax used to compute the fp8 scale in fwd
+//
+// Chain: y = x_normed * weight; fp8 = sat(y * scale). Through STE: grad_y = grad_fp8 * scale.
+//        grad_x_normed = grad_y * weight.
+//        grad_weight   = sum_rows(grad_y * x_normed).
+//        grad_x        = rrms * (grad_x_normed - x_normed * mean(grad_x_normed * x_normed, last_dim)).
+
+#ifndef N_ELEMS
+#define N_ELEMS 67108864
+#endif
+#ifndef HIDDEN
+#define HIDDEN 4096
+#endif
+#ifndef NUM_WG
+#define NUM_WG 1024
+#endif
+#ifndef THREADS_PER_WG
+#define THREADS_PER_WG 256
+#endif
+
+constexpr int VEC = 8;
+constexpr float FP8_MAX = 448.0f;
+
+static_assert(N_ELEMS % HIDDEN == 0, "N_ELEMS must be a multiple of HIDDEN");
+static_assert(HIDDEN % (THREADS_PER_WG * VEC) == 0, "HIDDEN must be divisible by THREADS_PER_WG*VEC");
+
+constexpr int ROWS = N_ELEMS / HIDDEN;
+constexpr int ELEMS_PER_THREAD = HIDDEN / THREADS_PER_WG;
+constexpr int VECS_PER_THREAD = ELEMS_PER_THREAD / VEC;
+
+extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
+fused_rmsnorm_mul_quantize_fp8_bwd(
+    __hip_bfloat16*       __restrict__ grad_x,              // out: bf16, ROWS*HIDDEN
+    float*                __restrict__ grad_weight_partial, // out: fp32, NUM_WG*HIDDEN
+    const __hip_bfloat16* __restrict__ grad_fp8,            // in:  bf16, ROWS*HIDDEN (grad of fp8_out)
+    const __hip_bfloat16* __restrict__ x_normed,            // in:  bf16, ROWS*HIDDEN
+    const float*          __restrict__ rrms,                // in:  fp32, ROWS
+    const __hip_bfloat16* __restrict__ weight,              // in:  bf16, HIDDEN
+    const float*          __restrict__ amax_state)          // in:  fp32 scalar
+{
+  __shared__ float sdata[THREADS_PER_WG];
+
+  const int tid = threadIdx.x;
+  const int wg  = blockIdx.x;
+
+  const float scale = FP8_MAX / (static_cast<float>(*amax_state) + 1e-8f);
+  const float inv_hidden = 1.0f / static_cast<float>(HIDDEN);
+
+  // Per-thread accumulator for grad_weight (across all rows this WG touches).
+  float gw_accum[ELEMS_PER_THREAD];
+  #pragma unroll
+  for (int i = 0; i < ELEMS_PER_THREAD; i++) gw_accum[i] = 0.0f;
+
+  // Preload weight into registers (same across rows). Use ELEMS_PER_THREAD entries.
+  float w_regs[ELEMS_PER_THREAD];
+  #pragma unroll
+  for (int v = 0; v < VECS_PER_THREAD; v++) {
+    const int h_base = tid * VEC + v * THREADS_PER_WG * VEC;
+    float4 w_raw = *reinterpret_cast<const float4*>(&weight[h_base]);
+    const __hip_bfloat16 *wi = reinterpret_cast<const __hip_bfloat16*>(&w_raw);
+    #pragma unroll
+    for (int i = 0; i < VEC; i++) w_regs[v * VEC + i] = static_cast<float>(wi[i]);
+  }
+
+  for (int row = wg; row < ROWS; row += NUM_WG) {
+    const int row_off = row * HIDDEN;
+    const float rrms_v = rrms[row];
+
+    // Load grad_fp8 and x_normed rows into registers, compute grad_y and grad_x_normed.
+    float g_y_regs[ELEMS_PER_THREAD];
+    float xn_regs[ELEMS_PER_THREAD];
+    float g_xn_regs[ELEMS_PER_THREAD];  // grad_x_normed
+    float local_dot = 0.0f;  // sum(grad_x_normed * x_normed) for mean
+
+    #pragma unroll
+    for (int v = 0; v < VECS_PER_THREAD; v++) {
+      const int h_base = tid * VEC + v * THREADS_PER_WG * VEC;
+      float4 g_raw  = *reinterpret_cast<const float4*>(&grad_fp8[row_off + h_base]);
+      float4 xn_raw = *reinterpret_cast<const float4*>(&x_normed[row_off + h_base]);
+      const __hip_bfloat16 *gi  = reinterpret_cast<const __hip_bfloat16*>(&g_raw);
+      const __hip_bfloat16 *xni = reinterpret_cast<const __hip_bfloat16*>(&xn_raw);
+
+      #pragma unroll
+      for (int i = 0; i < VEC; i++) {
+        const int idx = v * VEC + i;
+        const float g_y = static_cast<float>(gi[i]) * scale;
+        const float xn  = static_cast<float>(xni[i]);
+        g_y_regs[idx] = g_y;
+        xn_regs[idx]  = xn;
+        g_xn_regs[idx] = g_y * w_regs[idx];           // grad_x_normed = grad_y * weight
+        gw_accum[idx] += g_y * xn;                    // grad_weight contrib
+        local_dot += g_xn_regs[idx] * xn;             // for mean
+      }
+    }
+
+    // LDS reduce local_dot to sdata[0].
+    sdata[tid] = local_dot;
+    __syncthreads();
+    for (int s = THREADS_PER_WG / 2; s > 0; s >>= 1) {
+      if (tid < s) sdata[tid] = sdata[tid] + sdata[tid + s];
+      __syncthreads();
+    }
+    const float mean_term = sdata[0] * inv_hidden;
+
+    // Compute grad_x = rrms * (grad_x_normed - x_normed * mean_term) and write.
+    #pragma unroll
+    for (int v = 0; v < VECS_PER_THREAD; v++) {
+      const int h_base = tid * VEC + v * THREADS_PER_WG * VEC;
+      __hip_bfloat16 out[VEC];
+      #pragma unroll
+      for (int i = 0; i < VEC; i++) {
+        const int idx = v * VEC + i;
+        const float dx = rrms_v * (g_xn_regs[idx] - xn_regs[idx] * mean_term);
+        out[i] = static_cast<__hip_bfloat16>(dx);
+      }
+      *reinterpret_cast<float4*>(&grad_x[row_off + h_base]) = *reinterpret_cast<float4*>(out);
+    }
+    __syncthreads();
+  }
+
+  // Write this WG's grad_weight partial to HBM (fp32, NUM_WG x HIDDEN layout).
+  const int gw_row_off = wg * HIDDEN;
+  #pragma unroll
+  for (int v = 0; v < VECS_PER_THREAD; v++) {
+    const int h_base = tid * VEC + v * THREADS_PER_WG * VEC;
+    // Write 8 fp32 values with two float4 stores.
+    float4 out_lo, out_hi;
+    out_lo.x = gw_accum[v * VEC + 0]; out_lo.y = gw_accum[v * VEC + 1];
+    out_lo.z = gw_accum[v * VEC + 2]; out_lo.w = gw_accum[v * VEC + 3];
+    out_hi.x = gw_accum[v * VEC + 4]; out_hi.y = gw_accum[v * VEC + 5];
+    out_hi.z = gw_accum[v * VEC + 6]; out_hi.w = gw_accum[v * VEC + 7];
+    *reinterpret_cast<float4*>(&grad_weight_partial[gw_row_off + h_base + 0]) = out_lo;
+    *reinterpret_cast<float4*>(&grad_weight_partial[gw_row_off + h_base + 4]) = out_hi;
+  }
+}

extra/llama_kernels/quantize_fp8_delayed/__init__.py

@@ -0,0 +1,67 @@
+from __future__ import annotations
+import functools, pathlib
+from tinygrad import Tensor, dtypes
+from tinygrad.uop.ops import UOp, Ops, KernelInfo
+from tinygrad.renderer import Estimates
+from extra.llama_kernels import FP8_MAX, NUM_WG, THREADS_PER_WG, alloc_like, alloc_local, scalar_amax, dname_of, compile_hip
+
+@functools.cache
+def _custom_quantize_fp8_with_amax(fp8_out:UOp, amax_partial:UOp, x:UOp, amax_state:UOp, dname:str) -> UOp:
+  n_elems = 1
+  for d in x.shape: n_elems *= d
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(NUM_WG, "gidx0")
+  mem = n_elems * 2 + n_elems + 4 + NUM_WG * 4
+  sink = UOp.sink(fp8_out.base, amax_partial.base, x.base, amax_state.base, threads, workgroups,
+                  arg=KernelInfo(f"quantize_fp8_with_amax_{n_elems}", estimates=Estimates(ops=3*n_elems, mem=mem)))
+  src = (pathlib.Path(__file__).parent/"quantize_fp8_with_amax.cpp").read_text()
+  defines = [f"-DN_ELEMS={n_elems}", f"-DNUM_WG={NUM_WG}", f"-DTHREADS_PER_WG={THREADS_PER_WG}"]
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=compile_hip(src, defines))))
+
+@functools.cache
+def _custom_quantize_fp8_scalar(fp8_out:UOp, x:UOp, amax_state:UOp, dname:str) -> UOp:
+  n_elems = 1
+  for d in x.shape: n_elems *= d
+  threads, workgroups = UOp.special(THREADS_PER_WG, "lidx0"), UOp.special(NUM_WG, "gidx0")
+  mem = n_elems * 2 + n_elems
+  sink = UOp.sink(fp8_out.base, x.base, amax_state.base, threads, workgroups,
+                  arg=KernelInfo(f"quantize_fp8_scalar_{n_elems}", estimates=Estimates(ops=2*n_elems, mem=mem)))
+  src = (pathlib.Path(__file__).parent/"quantize_fp8_scalar.cpp").read_text()
+  defines = [f"-DN_ELEMS={n_elems}", f"-DNUM_WG={NUM_WG}", f"-DTHREADS_PER_WG={THREADS_PER_WG}"]
+  return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg=dname), UOp(Ops.LINEAR, src=(*sink.src, sink)),
+                               UOp(Ops.SOURCE, arg=src), UOp(Ops.BINARY, arg=compile_hip(src, defines))))
+
+def _quantize_fp8_delayed_bwd(gradient:UOp, kernel:UOp):
+  # NOTE: STE-equivalent backward — grad_x = grad_fp8 * scale, scale = FP8_MAX / amax_state.
+  # `gradient` is bf16 grad w.r.t. fp8 output (asm_gemm bwd already applied x_scale).
+  _, _, x, amax_state = kernel.src[1:]
+  device = x.device
+  scale = FP8_MAX / (Tensor(amax_state, device=device).float() + 1e-8)
+  grad_x = (Tensor(gradient, device=device).float() * scale).cast(dtypes.bfloat16)
+  return (None, None, grad_x.uop, None)
+
+def quantize_fp8_delayed(x:Tensor, amax_state:Tensor, fp8_dtype=dtypes.fp8e4m3) -> tuple[Tensor, Tensor, Tensor, UOp]:
+  # NOTE: one-pass bf16 -> fp8 quantize with delayed scaling. Returns (fp8, inv_scale, new_amax, store_effect).
+  # Fused kernel reads x once and writes fp8 + per-WG |x| partials (then a small reduce produces scalar new_amax).
+  # store_effect writes new_amax into amax_state's buffer — the caller must thread it into a realized
+  # output via `.after(store_effect)`. Calling `amax_state.assign(new_amax)` inside a grad_fxn does
+  # NOT work because .assign mutates only the temp Tensor's .uop, not the original layer-owned buffer.
+  assert x.dtype == dtypes.bfloat16, f"expected bf16, got {x.dtype}"
+  axis = x.uop.axis if isinstance(x.device, tuple) else None
+  fp8_out      = alloc_like(x.shape,  fp8_dtype,      x.device, axis)
+  amax_partial = alloc_local((NUM_WG,), dtypes.float32, x.device)
+  fxn = functools.partial(_custom_quantize_fp8_with_amax, dname=dname_of(x.device))
+  fp8_out, amax_partial, *_ = Tensor.custom_kernel(fp8_out, amax_partial, x, amax_state,
+                                                    fxn=fxn, grad_fxn=_quantize_fp8_delayed_bwd)
+  new_amax = scalar_amax(amax_partial)
+  inv_scale = (amax_state.float() + 1e-8) / FP8_MAX
+  store_effect = amax_state.uop.store(new_amax.uop)
+  return fp8_out, inv_scale, new_amax, store_effect
+
+def quantize_fp8_scalar(x:Tensor, amax_state:Tensor, fp8_dtype=dtypes.fp8e4m3) -> Tensor:
+  # NOTE: pure one-pass bf16 -> fp8 quantize with delayed scalar scale. No amax computation.
+  axis = x.uop.axis if isinstance(x.device, tuple) else None
+  fp8_out = alloc_like(x.shape, fp8_dtype, x.device, axis)
+  fxn = functools.partial(_custom_quantize_fp8_scalar, dname=dname_of(x.device))
+  fp8_out, *_ = Tensor.custom_kernel(fp8_out, x, amax_state, fxn=fxn)
+  return fp8_out

extra/llama_kernels/quantize_fp8_delayed/quantize_fp8_scalar.cpp

@@ -0,0 +1,48 @@
+#include <hip/hip_runtime.h>
+#include <hip/hip_bf16.h>
+#include <hip/hip_fp8.h>
+
+// Pure one-pass bf16 -> fp8 quantize with delayed scalar scale. No amax computation.
+
+#ifndef N_ELEMS
+#define N_ELEMS 67108864
+#endif
+#ifndef NUM_WG
+#define NUM_WG 1024
+#endif
+#ifndef THREADS_PER_WG
+#define THREADS_PER_WG 256
+#endif
+
+constexpr int VEC = 8;
+constexpr float FP8_MAX = 448.0f;
+
+static_assert(N_ELEMS % VEC == 0, "N_ELEMS must be divisible by VEC");
+
+extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
+quantize_fp8_scalar(
+    __hip_fp8_storage_t*  __restrict__ fp8_out,         // fp8, N_ELEMS
+    const __hip_bfloat16* __restrict__ x,               // bf16, N_ELEMS
+    const float*          __restrict__ amax_state)      // fp32 scalar (delayed)
+{
+  const int tid = threadIdx.x;
+  const int wg  = blockIdx.x;
+  const int gid = wg * THREADS_PER_WG + tid;
+  const int stride_elems = NUM_WG * THREADS_PER_WG * VEC;
+
+  const float scale = FP8_MAX / (static_cast<float>(*amax_state) + 1e-8f);
+
+  for (int base = gid * VEC; base < N_ELEMS; base += stride_elems) {
+    float4 x_raw = *reinterpret_cast<const float4*>(&x[base]);
+    const __hip_bfloat16 *xi = reinterpret_cast<const __hip_bfloat16*>(&x_raw);
+
+    __hip_fp8_storage_t out[VEC];
+    #pragma unroll
+    for (int i = 0; i < VEC; i++) {
+      const float v = static_cast<float>(xi[i]);
+      const float scaled = fmaxf(-FP8_MAX, fminf(FP8_MAX, v * scale));
+      out[i] = __hip_cvt_float_to_fp8(scaled, __HIP_SATFINITE, __HIP_E4M3);
+    }
+    *reinterpret_cast<uint64_t*>(&fp8_out[base]) = *reinterpret_cast<uint64_t*>(out);
+  }
+}

extra/llama_kernels/quantize_fp8_delayed/quantize_fp8_with_amax.cpp

@@ -2,12 +2,13 @@
 #include <hip/hip_bf16.h>
 #include <hip/hip_fp8.h>
 
+// One-pass bf16 -> fp8 quantize using a scalar delayed amax state,
+// AND simultaneously computes per-WG |x| max partials for the next step's amax state.
+// Saves one full HBM pass over the grad tensor vs. doing quantize + separate abs().max().
+
 #ifndef N_ELEMS
 #define N_ELEMS 67108864
 #endif
-#ifndef HIDDEN
-#define HIDDEN 4096
-#endif
 #ifndef NUM_WG
 #define NUM_WG 1024
 #endif
@@ -19,15 +20,13 @@ constexpr int VEC = 8;
 constexpr float FP8_MAX = 448.0f;
 
 static_assert(N_ELEMS % VEC == 0, "N_ELEMS must be divisible by VEC");
-static_assert(HIDDEN % VEC == 0, "HIDDEN must be divisible by VEC");
 
 extern "C" __global__ __launch_bounds__(THREADS_PER_WG) void
-fused_mul_quantize_fp8(
-    __hip_fp8_storage_t*  __restrict__ fp8_out,         // fp8, N_ELEMS
-    __hip_bfloat16*       __restrict__ amax_buf,        // bf16, NUM_WG
-    const __hip_bfloat16* __restrict__ x,               // bf16, N_ELEMS
-    const __hip_bfloat16* __restrict__ weight,          // bf16, HIDDEN (per-hidden scale)
-    const __hip_bfloat16* __restrict__ amax_state)      // bf16 scalar
+quantize_fp8_with_amax(
+    __hip_fp8_storage_t*  __restrict__ fp8_out,       // out: fp8, N_ELEMS
+    float*                __restrict__ amax_partial,  // out: fp32, NUM_WG per-WG partials
+    const __hip_bfloat16* __restrict__ x,             // in:  bf16, N_ELEMS
+    const float*          __restrict__ amax_state)    // in:  fp32 scalar (delayed)
 {
   __shared__ float sdata[THREADS_PER_WG];
 
@@ -40,32 +39,25 @@ fused_mul_quantize_fp8(
   float local_max = 0.0f;
 
   for (int base = gid * VEC; base < N_ELEMS; base += stride_elems) {
-    const int h = base % HIDDEN;   // 0..HIDDEN-VEC, 8-aligned (since base is 8-aligned and HIDDEN divides VEC)
     float4 x_raw = *reinterpret_cast<const float4*>(&x[base]);
-    float4 w_raw = *reinterpret_cast<const float4*>(&weight[h]);
-
     const __hip_bfloat16 *xi = reinterpret_cast<const __hip_bfloat16*>(&x_raw);
-    const __hip_bfloat16 *wi = reinterpret_cast<const __hip_bfloat16*>(&w_raw);
 
     __hip_fp8_storage_t out[VEC];
     #pragma unroll
     for (int i = 0; i < VEC; i++) {
-      const float val = static_cast<float>(xi[i]) * static_cast<float>(wi[i]);
-      local_max = fmaxf(local_max, fabsf(val));
-      const float scaled = fmaxf(-FP8_MAX, fminf(FP8_MAX, val * scale));
+      const float v = static_cast<float>(xi[i]);
+      local_max = fmaxf(local_max, fabsf(v));
+      const float scaled = fmaxf(-FP8_MAX, fminf(FP8_MAX, v * scale));
       out[i] = __hip_cvt_float_to_fp8(scaled, __HIP_SATFINITE, __HIP_E4M3);
     }
-
     *reinterpret_cast<uint64_t*>(&fp8_out[base]) = *reinterpret_cast<uint64_t*>(out);
   }
 
-  // LDS tree-reduce per-WG amax
   sdata[tid] = local_max;
   __syncthreads();
   for (int s = THREADS_PER_WG / 2; s > 0; s >>= 1) {
     if (tid < s) sdata[tid] = fmaxf(sdata[tid], sdata[tid + s]);
     __syncthreads();
   }
-
-  if (tid == 0) amax_buf[wg] = static_cast<__hip_bfloat16>(sdata[0]);
+  if (tid == 0) amax_partial[wg] = sdata[0];
 }

extra/sqtt/roc.py

@@ -1,10 +1,13 @@
 #!/usr/bin/env python3
-import ctypes, pathlib, argparse, pickle, dataclasses, threading
+import ctypes, pathlib, argparse, pickle, dataclasses, threading, itertools
+from decimal import Decimal
 from typing import Generator
 from tinygrad.helpers import temp, unwrap, DEBUG
 from tinygrad.runtime.ops_amd import ProfileSQTTEvent
 from tinygrad.runtime.autogen import rocprof
 from tinygrad.renderer.amd.dsl import Inst
+from tinygrad.helpers import ProfileEvent, ProfileRangeEvent, ProfilePointEvent
+from tinygrad.device import ProfileProgramEvent
 from test.amd.disasm import disasm
 
 @dataclasses.dataclass(frozen=True)
@@ -126,6 +129,71 @@ def decode(sqtt_evs:list[ProfileSQTTEvent], disasms:dict[str, dict[int, Inst]])
     raise exc
   return ROCParseCtx
 
+def unpack_occ(viz_data, i:int, j:int, key:tuple[str, int], data:list, p:ProfileProgramEvent, target:str) -> dict:
+  from tinygrad.viz.serve import amd_decode, create_step, row_tuple
+  steps = viz_data.ctxs[i]["steps"]
+  if len(steps[j+1:]) > 0: return {"steps":[{k:v for k,v in s.items() if k != "data"} for s in steps[j+1:]]}
+  base = unwrap(p.base)
+  disasm:dict[int, Inst] = {addr+base:inst for addr,inst in amd_decode(unwrap(p.lib), target).items()}
+  rctx = decode(data, {p.tag:disasm})
+  cu_events:dict[str, list[ProfileEvent]] = {}
+  # ** inst traces
+  wave_insts:dict[str, dict[str, dict]] = {}
+  inst_units:dict[str, itertools.count] = {}
+  for w in rctx.inst_execs.get(key, []):
+    if (u:=w.wave_loc) not in inst_units: inst_units[u] = itertools.count(0)
+    n = next(inst_units[u])
+    if (events:=cu_events.get(w.cu_loc)) is None: cu_events[w.cu_loc] = events = []
+    events.append(ProfileRangeEvent(f"SIMD:{w.simd}", loc:=f"INST WAVE:{w.wave_id} N:{n}", Decimal(w.begin_time), Decimal(w.end_time)))
+    wave_insts.setdefault(w.cu_loc, {})[f"{u} N:{n}"] = {"wave":w, "disasm":disasm, "prg":p, "run_number":n, "loc":loc}
+  # ** occ traces (only WAVESTART/WAVEEND)
+  units:dict[str, itertools.count] = {}
+  wave_start:dict[str, int] = {}
+  for occ in rctx.occ_events.get(key, []):
+    if (u:=occ.wave_loc) not in units: units[u] = itertools.count(0)
+    if u in inst_units: continue
+    if occ.start: wave_start[u] = occ.time
+    else:
+      if (events:=cu_events.get(occ.cu_loc)) is None: cu_events[occ.cu_loc] = events = []
+      events.append(ProfileRangeEvent(f"SIMD:{occ.simd}", f"OCC WAVE:{occ.wave_id} N:{next(units[u])}", Decimal(wave_start.pop(u)),Decimal(occ.time)))
+  # ** split graph by CU
+  for cu in sorted(cu_events, key=row_tuple):
+    steps.append(create_step(f"{cu} {len(cu_events[cu])}", ("/cu-sqtt", i, len(steps)), depth=1,
+                             data=[ProfilePointEvent(unit, "start", unit, ts=Decimal(0)) for unit in units]+cu_events[cu]))
+    for k in sorted(wave_insts.get(cu, []), key=row_tuple):
+      wd = wave_insts[cu][k]
+      steps.append(create_step(k.replace(cu, ""), ("/amd-sqtt-insts", i, len(steps)), loc=wd["loc"], depth=2,
+                               data={"fxn":unpack_insts, "args":(wd,)}))
+  return {"steps":[{k:v for k,v in s.items() if k != "data"} for s in steps[j+1:]]}
+
+def unpack_insts(viz_data, i:int, j:int, data:dict) -> dict:
+  columns = ["PC", "Instruction", "Hits", "Cycles", "Stall", "Type"]
+  inst_columns = ["N", "Clk", "Idle", "Dur", "Stall"]
+  # Idle:     The total time gap between the completion of previous instruction and the beginning of the current instruction.
+  #           The idle time can be caused by:
+  #             * Arbiter loss
+  #             * Source or destination register dependency
+  #             * Instruction cache miss
+  # Stall:    The total number of cycles the hardware pipe couldn't issue an instruction.
+  # Duration: Total latency in cycles, defined as "Stall time + Issue time" for gfx9 or "Stall time + Execute time" for gfx10+.
+  prev_instr = (w:=data["wave"]).begin_time
+  pc_to_inst = data["disasm"]
+  start_pc = None
+  rows:dict[int, dict] = {}
+  for pc, inst in pc_to_inst.items():
+    if start_pc is None: start_pc = pc
+    rows[pc] = {"pc":pc-start_pc, "inst":str(inst), "hit_count":0, "dur":0, "stall":0, "type":"", "hits":{"cols":inst_columns, "rows":[]}}
+  for e in w.unpack_insts():
+    if not (inst:=rows[e.pc]).get("type"): inst["type"] = str(e.typ).split("_")[-1]
+    inst["hit_count"] += 1
+    inst["dur"] += e.dur
+    inst["stall"] += e.stall
+    inst["hits"]["rows"].append((inst["hit_count"]-1, e.time, max(0, e.time-prev_instr), e.dur, e.stall))
+    prev_instr = max(prev_instr, e.time + e.dur)
+  summary = [{"label":"Total Cycles", "value":w.end_time-w.begin_time}, {"label":"SE", "value":w.se}, {"label":"CU", "value":w.cu},
+             {"label":"SIMD", "value":w.simd}, {"label":"Wave ID", "value":w.wave_id}, {"label":"Run number", "value":data["run_number"]}]
+  return {"rows":[tuple(v.values()) for v in rows.values()], "cols":columns, "metadata":[summary], "ref":viz_data.ref_map.get(data["prg"].name)}
+
 def print_data(data:dict) -> None:
   from tabulate import tabulate
   # plaintext

extra/thunder/amd/fa.py

@@ -55,8 +55,6 @@ def flash_attention(xq, xk, xv, attn_mask:Tensor|None=None, is_causal:bool=False
   assert attn_mask is None, "attn_mask not supported"
   assert is_causal, "only causal attention supported"
 
-  xq, xk, xv = xq.transpose(1, 2), xk.transpose(1, 2), xv.transpose(1, 2)
-
   B, N, H, D = xq.shape
   H_KV = xk.shape[2]
   assert D == 128, "only D=128 supported"
@@ -81,7 +79,7 @@ def flash_attention(xq, xk, xv, attn_mask:Tensor|None=None, is_causal:bool=False
 
   attn, l_vec = Tensor.custom_kernel(attn, l_vec, xq, xk, xv, fxn=functools.partial(custom_fa_forward, device=single_device, arch=arch, B=B_local, N=N, H=H_local, H_KV=H_KV_local, D=D), grad_fxn=grad)[:2]
 
-  return attn.transpose(1, 2), attn, l_vec
+  return attn, attn, l_vec
 
 @functools.cache
 def custom_fa_forward(o:UOp, l_vec:UOp, q:UOp, k:UOp, v:UOp, device:str, arch:str, B:int, N:int, H:int, H_KV:int, D:int):

extra/thunder/amd/gemm_fp8.cpp

@@ -93,7 +93,20 @@ constexpr int NUM_WARPS = 8;
 
 using G = kittens::group<NUM_WARPS>;
 
-__global__ __launch_bounds__(512, 2) void hk_fp8_gemm(bf16 *C_ptr, fp8e4m3 *A_ptr, fp8e4m3 *B_ptr, float *x_scale_ptr, float *w_scale_ptr) {
+// scale_mode: 0=no scale, 1=x only, 2=w only, 3=both
+#ifndef SCALE_MODE
+#define SCALE_MODE 3
+#endif
+
+__global__ __launch_bounds__(512, 2) void hk_fp8_gemm(bf16 *C_ptr, fp8e4m3 *A_ptr, fp8e4m3 *B_ptr
+#if SCALE_MODE == 1
+    , float *x_scale_ptr
+#elif SCALE_MODE == 2
+    , float *w_scale_ptr
+#elif SCALE_MODE == 3
+    , float *x_scale_ptr, float *w_scale_ptr
+#endif
+) {
     constexpr int M = GEMM_M, N = GEMM_N, K = GEMM_K;
 
     kittens::gl<fp8e4m3, 1, 1, M, K> A{A_ptr, nullptr, nullptr, nullptr, nullptr};
@@ -333,11 +346,25 @@ __global__ __launch_bounds__(512, 2) void hk_fp8_gemm(bf16 *C_ptr, fp8e4m3 *A_pt
     }
 
     // apply x_scale * w_scale before bf16 store to prevent overflow
+#if SCALE_MODE == 1
+    float scale = *x_scale_ptr;
+    mul(cA, cA, scale);
+    mul(cB, cB, scale);
+    mul(cC, cC, scale);
+    mul(cD, cD, scale);
+#elif SCALE_MODE == 2
+    float scale = *w_scale_ptr;
+    mul(cA, cA, scale);
+    mul(cB, cB, scale);
+    mul(cC, cC, scale);
+    mul(cD, cD, scale);
+#elif SCALE_MODE == 3
     float scale = *x_scale_ptr * *w_scale_ptr;
     mul(cA, cA, scale);
     mul(cB, cB, scale);
     mul(cC, cC, scale);
     mul(cD, cD, scale);
+#endif
 
     store(C, cA, {0, 0, block_row * WARPS_ROW * 2 + warp_m, block_col * WARPS_COL * 2 + warp_n});
     store(C, cB, {0, 0, block_row * WARPS_ROW * 2 + warp_m, block_col * WARPS_COL * 2 + WARPS_COL + warp_n});

test/amd/test_custom_kernel.py

@@ -3,7 +3,7 @@ import functools
 import numpy as np
 from tinygrad import Tensor, Device, dtypes
 from tinygrad.uop.ops import UOp, Ops, KernelInfo
-from tinygrad.engine.realize import run_linear, estimate_uop
+from tinygrad.engine.realize import run_linear, estimate_uop, compile_linear
 from tinygrad.renderer import Estimates
 from tinygrad.dtype import AddrSpace
 from tinygrad.helpers import getenv
@@ -99,13 +99,13 @@ def custom_lds_sync(A:UOp, arch:str) -> UOp:
   sink = UOp.sink(A.base, lds, threads, wg, arg=KernelInfo("custom_lds_sync"))
   return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg="AMD"), UOp(Ops.LINEAR, src=tuple([UOp(Ops.INS, arg=x) for x in insts]))))
 
-def custom_handwritten(A:UOp, arch:str) -> UOp:
+def custom_handwritten(A:UOp) -> UOp:
   A = A.flatten()
   threads = UOp.special(128, "lidx0")
   wg = UOp.special(1, "gidx0")
   lds = UOp(Ops.DEFINE_LOCAL, dtypes.uint8.ptr(size=512, addrspace=AddrSpace.LOCAL), (), 'lds')  # 128 * 4 bytes
   pipes = {getenv("PIPE", "")} if getenv("PIPE", "") else {"SALU", "VALU", "TRANSCENDENTAL", "WMMA"}
-  k = Kernel(arch)
+  k = Kernel()
   # wrap in loop to filter out icache misses
   LOOP_N, UNROLL_N = 8, 5
   k.emit(r4.s_mov_b32(s[1], LOOP_N))
@@ -145,10 +145,10 @@ def custom_handwritten(A:UOp, arch:str) -> UOp:
   sink = UOp.sink(A.base, threads, wg, lds, arg=KernelInfo("custom_handwritten"))
   return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg="AMD"), UOp(Ops.LINEAR, src=tuple([UOp(Ops.INS, arg=x) for x in insts]))))
 
-def custom_data_deps(A:UOp, arch:str) -> UOp:
+def custom_data_deps(A:UOp) -> UOp:
   A = A.flatten()
   threads = UOp.special(A.numel(), "lidx0")
-  k = Kernel(arch)
+  k = Kernel()
   k.emit(s_load_b64(s[0:1], s[0:1], soffset=NULL))
   k.emit(s_waitcnt_lgkmcnt(sdst=NULL, simm16=0))
   k.emit(v_lshlrev_b32_e32(v[0], 2, v[0]))
@@ -169,7 +169,7 @@ class TestCustomKernel(unittest.TestCase):
     if self.arch != "rdna3": self.skipTest("only rdna3")
     a = Tensor.full((16, 16), 1.).contiguous().realize()
     a = Tensor.custom_kernel(a, fxn=custom_add_one)[0]
-    linear = a.schedule_linear()
+    linear = compile_linear(a.schedule_linear())
     est = estimate_uop(linear.src[-1])
     self.assertEqual(est.ops, a.numel())
     self.assertEqual(est.mem, a.nbytes()*2)
@@ -198,13 +198,13 @@ class TestCustomKernel(unittest.TestCase):
   def test_handwritten(self):
     if self.arch != "rdna4": self.skipTest("only tested on rdna4")
     a = Tensor.empty(1024, dtype=dtypes.int32).contiguous().realize()
-    a = Tensor.custom_kernel(a, fxn=functools.partial(custom_handwritten, arch=self.arch))[0]
+    a = Tensor.custom_kernel(a, fxn=custom_handwritten)[0]
     a.realize()
 
   def test_data_deps(self):
     if self.arch != "rdna3": self.skipTest("only tested on rdna3")
     a = Tensor(np.full(32, 5.0, dtype=np.float32)).realize()
-    a = Tensor.custom_kernel(a, fxn=functools.partial(custom_data_deps, arch=self.arch))[0]
+    a = Tensor.custom_kernel(a, fxn=custom_data_deps)[0]
     a.realize()
     self.assertTrue((a.numpy() == 6.0).all())
 

test/amd/test_mockgpu_invalid.py

@@ -8,17 +8,14 @@ class TestMockGPUInvalidInstruction(unittest.TestCase):
     test_code = '''
 import struct
 from tinygrad import Device, Tensor
-from tinygrad.engine.realize import get_runner
+from tinygrad.engine.realize import compile_linear
 from tinygrad.runtime.ops_amd import AMDProgram
 
 dev = Device["AMD"]
 a = Tensor([1.0]).realize()
 b = a + 1
-si = b.schedule_linear().src[-1]
-runner = get_runner(dev.device, si.src[0])
-
-prg = runner._prg
-lib = bytearray(prg.lib)
+linear = compile_linear(b.schedule_linear())
+lib = bytearray(linear.src[-1].src[0].src[4].arg)
 
 # Find s_endpgm (0xBFB00000) and replace with V_MOVRELD_B32 (op=66) which has no pcode
 # VOP1 encoding: bits[31:25]=0x7E, op=bits[16:9], so op=66 -> 66<<9 = 0x8400

test/backend/test_arange.py

@@ -2,14 +2,14 @@ import unittest
 import numpy as np
 from tinygrad import Tensor, GlobalCounters, dtypes, nn, Device, Variable
 from tinygrad.helpers import Context, getenv, DEV
-from tinygrad.engine.realize import run_linear, estimate_uop
+from tinygrad.engine.realize import run_linear, estimate_uop, compile_linear
 from tinygrad.renderer.ptx import PTXRenderer
 from test.helpers import needs_second_gpu
 
 class TestArange(unittest.TestCase):
   def _get_flops(self, tensor, desired):
     GlobalCounters.reset()
-    linear = tensor.schedule_linear()
+    linear = compile_linear(tensor.schedule_linear())
     self.assertEqual(len(linear.src), 1)
     run_linear(linear)
     np.testing.assert_equal(tensor.numpy(), desired)
@@ -36,7 +36,7 @@ class TestArange(unittest.TestCase):
   def test_tri_complexity(self):
     with Context(NOOPT=1):
       t = Tensor.ones(256, 256).contiguous().realize()
-      linear = t.triu().schedule_linear()
+      linear = compile_linear(t.triu().schedule_linear())
       self.assertLessEqual(estimate_uop(linear.src[-1]).ops, 4 * 256 * 256)
 
 DSET, DDIM = 2048, 32
@@ -229,7 +229,7 @@ class TestIndexing(unittest.TestCase):
     xq = xq.reshape(bs, seqlen, n_heads, head_dim)
     xq_rope, _ = apply_rotary_emb(xq, xq, freqs_cis)
     xq_rope.sum().backward()
-    linear = wq.grad.schedule_linear()
+    linear = compile_linear(wq.grad.schedule_linear())
     assert len(linear.src) == 1, f"expected one kernel for backward, got: {len(linear.src)}"
     bwd_ops = estimate_uop(linear.src[0]).ops
     # bfloat16 on non CDNA4 has ~10x ops overhead because of the software emulation

test/backend/test_custom_kernel.py

@@ -1,5 +1,5 @@
 import unittest
-from tinygrad import Tensor, UOp
+from tinygrad import Tensor, UOp, GlobalCounters
 from tinygrad.dtype import AddrSpace, dtypes
 from tinygrad.uop.ops import KernelInfo, AxisType
 
@@ -308,6 +308,36 @@ class TestCustomKernel(unittest.TestCase):
     expected = (3+2)*2+2
     assert all(x == expected for x in result), f"expected all {expected}, got {result}"
 
+  def test_custom_kernel_sched(self, use_custom=False):
+    x = Tensor.arange(32).reshape(8, 4).realize()
+    y = Tensor.empty_like(x)
+    y = Tensor.custom_kernel(y, x, fxn=custom_add_one_kernel)[0]
+    if use_custom:
+      z = Tensor.empty_like(x)
+      z = Tensor.custom_kernel(y, y.T.T, fxn=custom_add_one_kernel)[0]
+    else: z = y.T.T+1
+    GlobalCounters.reset()
+    z.realize()
+    self.assertEqual(GlobalCounters.kernel_count, 2)
+    self.assertEqual(z.tolist(), x.add(2).tolist())
+
+  @unittest.expectedFailure
+  def test_custom_kernel_sched_copy(self): self.test_custom_kernel_sched(use_custom=True)
+
+  @unittest.expectedFailure
+  def test_sliced_buffer_function(self):
+    x = Tensor.arange(32).reshape(8, 4).realize()
+    from tinygrad import function
+    @function(precompile=True)
+    def run(x:Tensor) -> Tensor:
+      y = Tensor.invalids(*x.shape, dtype=x.dtype)
+      return Tensor.custom_kernel(y, x, fxn=custom_add_one_kernel)[0]
+    GlobalCounters.reset()
+    y = run(x[0]).realize()
+    # it's copying the input and the output
+    self.assertEqual(GlobalCounters.kernel_count, 1)
+    self.assertEqual(y.tolist(), [1, 2, 3, 4])
+
 class TestUOpReduce(unittest.TestCase):
   def test_uop_sum(self):
     a = Tensor([1.0, 2, 3, 4, 5])

test/backend/test_linearizer.py

@@ -2,10 +2,10 @@ import numpy as np
 import unittest
 
 from tinygrad.codegen.opt import Opt, OptOps
-from tinygrad.uop.ops import UOp, Ops, GroupOp, AxisType
+from tinygrad.uop.ops import UOp, Ops, GroupOp, AxisType, buffers
 from tinygrad.device import Device, Buffer, is_dtype_supported
 from tinygrad.tensor import Tensor, _to_np_dtype
-from tinygrad.engine.realize import run_linear, CompiledRunner
+from tinygrad.engine.realize import run_linear
 from tinygrad.codegen import to_program
 from tinygrad.helpers import Context, flatten, dedup, TC_SELECT, TC_OPT, DEV
 from tinygrad.dtype import DType, dtypes, PtrDType, AddrSpace
@@ -14,7 +14,7 @@ from tinygrad.renderer.cstyle import CUDARenderer
 from test.helpers import replace_opts
 MOCKGPU = DEV.interface.startswith("MOCK")
 
-from tinygrad.uop.ops import print_uops # noqa: F401 # pylint: disable=unused-import
+from tinygrad.uop.render import print_uops # noqa: F401 # pylint: disable=unused-import
 
 class TestLinearizer(unittest.TestCase):
   def test_arg_dedup(self):
@@ -274,7 +274,7 @@ class TestLinearizer(unittest.TestCase):
     sched = [si for si in t.schedule_linear().src if si.src[0].op is Ops.SINK]
     # sum_collapse is a full collapse now
     assert len(sched) == 1
-    assert not any(u.op is Ops.REDUCE_AXIS for u in sched[0].src[0].toposort()), "found reduce in sum collapse"
+    assert not any(u.op is Ops.REDUCE and len(u.arg[1]) > 0 for u in sched[0].src[0].toposort()), "found reduce in sum collapse"
     #lin = Kernel(sched[0].ast)
     #assert not any(u.op is Ops.RANGE for u in lin.linearize().uops), "found loop in sum collapse"
 
@@ -424,30 +424,28 @@ def reset_bufs(bufs:list[Buffer]):
 def _helper_linearizer_opt_ast(realized_ast:UOp, real_bufs:list[Buffer], opts=[],
                                apply_tc=False, atol=1e-4, rtol=1e-4, color_sizes=[], wanna_output=[]):
   outbufs = real_bufs[:len(realized_ast.src)]
-  device = real_bufs[0].device
   wanna_output = [np.array(x).flatten() for x in wanna_output]
+  buf_uops = [UOp.new_buffer(b.device, b.size, b.dtype) for b in real_bufs]
+  for u,b in zip(buf_uops, real_bufs): buffers[u] = b
 
-  def get_prg(opts):
+  def run_prg(opts):
     ast = realized_ast if opts is None else replace_opts(realized_ast, list(opts))
-    return CompiledRunner(to_program(ast, renderer=Device[Device.DEFAULT].renderer), device)
+    run_linear(UOp(Ops.LINEAR, src=(ast.call(*buf_uops),)))
 
   def check_opt(opts):
-    prg = get_prg(opts=opts)
     reset_bufs(outbufs)
-    prg.exec(real_bufs)
+    run_prg(opts)
     for x,want in zip(copyout_outputs(outbufs), wanna_output): np.testing.assert_allclose(x, want, atol=atol, rtol=rtol)
 
   # Get baseline if it is not provided, which is not optimized at all.
-  prg = get_prg(opts=())
-  prg.exec(real_bufs)
+  run_prg(opts=())
   if len(wanna_output) == 0: wanna_output = copyout_outputs(outbufs)
   else:
     for buf,want in zip(copyout_outputs(outbufs), wanna_output): np.testing.assert_allclose(buf, want, atol=atol, rtol=rtol)
 
   # Check correctness of handcoded optimiztions.
-  prg = get_prg(opts=None)
   reset_bufs(outbufs)
-  prg.exec(real_bufs)
+  run_prg(opts=None)
   for buf,want in zip(copyout_outputs(outbufs), wanna_output): np.testing.assert_allclose(buf, want, atol=atol, rtol=rtol)
   for x in opts: # Check custom transformations if any.
     check_opt(([Opt(OptOps.TC, 0, (TC_SELECT.value, TC_OPT.value, 1))] if apply_tc else [])+x)

test/backend/test_ops.py

@@ -636,6 +636,17 @@ class TestOps(unittest.TestCase):
         helper_test_op(None, lambda x: 100%x, forward_only=True, vals=[va])
         helper_test_op(None, lambda x: 100.5%x, forward_only=True, vals=[va])
 
+  def test_fmod(self):
+    a = [-4, 7, 5, 4, -7, 8, -9]
+    b = [2, -3, 8, -2, 3, 5, -5]
+    for float_a in [True, False]:
+      for float_b in [True, False]:
+        va = [float(ai) for ai in a] if float_a else a
+        vb = [float(bi) for bi in b] if float_b else b
+        helper_test_op(None, lambda x,y: x.fmod(y), forward_only=True, vals=[va, vb])
+        helper_test_op(None, lambda x: x.fmod(2), forward_only=True, vals=[va])
+        helper_test_op(None, lambda x: x.fmod(3.5), forward_only=True, vals=[va])
+
   def test_mul_naninf(self):
     helper_test_op([(45,65)], lambda x: x*math.inf)
     helper_test_op([(45,65)], lambda x: x*-math.inf)

test/backend/test_profiler.py

@@ -3,7 +3,8 @@ from tinygrad import Device, Tensor, dtypes, TinyJit
 from tinygrad.helpers import CI, DEV, Context, ProfileRangeEvent, cpu_profile, cpu_events, ProfilePointEvent, dedup
 from tinygrad.device import Buffer, BufferSpec, Compiled, ProfileDeviceEvent, ProfileGraphEvent
 from tinygrad.runtime.support.hcq import HCQCompiled
-from tinygrad.engine.realize import get_runner
+from tinygrad.engine.realize import get_runtime
+from tinygrad.codegen import to_program
 
 MOCKGPU = DEV.interface.startswith("MOCK")
 def _dev_base(d):
@@ -46,13 +47,15 @@ class TestProfiler(unittest.TestCase):
     TestProfiler.b = self.a + 1
     si = self.b.schedule_linear().src[-1]
 
-    TestProfiler.runner = get_runner(TestProfiler.d0.device, si.src[0])
+    TestProfiler.prg = to_program(si.src[0], TestProfiler.d0.renderer)
+    TestProfiler.runtime = get_runtime(TestProfiler.d0.device, TestProfiler.prg)
     TestProfiler.b.uop.buffer.allocate()
 
-  def test_profile_kernel_run(self):
-    runner_name = TestProfiler.runner._prg.name
+  def test_profile_kernel_run(self, wait=False):
+    runner_name = TestProfiler.runtime.name
     with helper_collect_profile(TestProfiler.d0) as profile:
-      TestProfiler.runner([TestProfiler.b.uop.buffer, TestProfiler.a.uop.buffer], var_vals={})
+      gs, ls = TestProfiler.prg.arg.launch_dims({})
+      TestProfiler.runtime(TestProfiler.b.uop.buffer._buf, TestProfiler.a.uop.buffer._buf, global_size=gs, local_size=ls, wait=wait)
 
     profile, _ = helper_profile_filter_device(profile, TestProfiler.d0.device)
     kernel_runs = [x for x in profile if isinstance(x, ProfileRangeEvent)]
@@ -60,6 +63,9 @@ class TestProfiler(unittest.TestCase):
     assert kernel_runs[0].name == runner_name, "kernel name is not correct"
     assert _dev_base(kernel_runs[0].device) == kernel_runs[0].device, "kernel should not be on a sub-device"
 
+  def test_profile_kernel_run_wait(self):
+    self.test_profile_kernel_run(wait=True)
+
   def test_profile_copyin(self):
     buf1 = Buffer(Device.DEFAULT, 2, dtypes.float, options=BufferSpec(nolru=True)).ensure_allocated()
 
@@ -70,12 +76,13 @@ class TestProfiler(unittest.TestCase):
     assert len(kernel_runs) == 1, "one kernel run is expected"
 
   def test_profile_multiops(self):
-    runner_name = TestProfiler.runner._prg.name
+    runner_name = TestProfiler.runtime.name
     buf1 = Buffer(Device.DEFAULT, 2, dtypes.float, options=BufferSpec(nolru=True)).ensure_allocated()
 
     with helper_collect_profile(TestProfiler.d0) as profile:
       buf1.copyin(memoryview(bytearray(struct.pack("ff", 0, 1))))
-      TestProfiler.runner([buf1, TestProfiler.a.uop.buffer], var_vals={})
+      gs, ls = TestProfiler.prg.arg.launch_dims({})
+      TestProfiler.runtime(buf1._buf, TestProfiler.a.uop.buffer._buf, global_size=gs, local_size=ls)
       buf1.copyout(memoryview(bytearray(buf1.nbytes)))
 
     evs = [x for x in profile if isinstance(x, ProfileRangeEvent) and x.device.startswith(TestProfiler.d0.device)]

test/backend/test_randomness.py

@@ -307,17 +307,26 @@ class TestRandomness(unittest.TestCase):
     with self.assertRaises(TypeError): Tensor.randint((3, 4), low=0, high=3.5)
     with self.assertRaises(TypeError): Tensor.randint((3, 4), low=1, high=3, dtype="float")
     with self.assertRaises(TypeError): Tensor.randint((3, 4), low=0, high=3, dtype=dtypes.float32)
+    # check low < high
+    with self.assertRaises(ValueError): Tensor.randint((3, 4), low=10, high=5)
+    with self.assertRaises(ValueError): Tensor.randint((3, 4), low=10, high=10)
+    np.testing.assert_array_equal(Tensor.randint(16, low=5, high=6).numpy(), 5)
 
   def test_normal(self):
     self.assertTrue(normal_test(Tensor.normal))
     self.assertTrue(equal_distribution(Tensor.normal, lambda x: torch.nn.init.normal_(torch.empty(x), mean=0, std=1),
                                                       lambda x: np.random.normal(loc=0, scale=1, size=x)))
+    # check std >= 0
+    with self.assertRaises(ValueError): Tensor.normal((3, 4), mean=0, std=-1)
 
   def test_uniform(self):
     self.assertFalse(normal_test(Tensor.uniform))
     self.assertTrue(equal_distribution(Tensor.uniform, lambda x: torch.nn.init.uniform_(torch.empty(x)), lambda x: np.random.uniform(size=x)))
     self.assertTrue(equal_distribution(partial(Tensor.uniform, low=-100, high=100, dtype=dtypes.int32),
                                        numpy_func=lambda x: np.random.randint(low=-100, high=100, size=x)))
+    # check low < high
+    with self.assertRaises(ValueError): Tensor.uniform((3, 4), low=5.0, high=3.0)
+    with self.assertRaises(ValueError): Tensor.uniform((3, 4), low=1.0, high=1.0)
 
   def test_scaled_uniform(self):
     self.assertFalse(normal_test(Tensor.scaled_uniform))
@@ -352,7 +361,7 @@ class TestRandomness(unittest.TestCase):
     _check_with_torch(w=[0.231, 0., 1., 0.5], num_samples=300, replacement=True)
     _check_with_torch(w=[[0.2, 0.8]], num_samples=300, replacement=True)  # 2D but only 1 row
     _check_with_torch(w=[[0.453, 0., 1., 0.81], [0.1, 0.8, 0., 0.1]], num_samples=300, replacement=True)
-    # no-replacement isn't supported, unless taking only one sample
+    # no-replacement
     w = [0.1, 0.9]
     self.assertRaises(AssertionError, lambda: Tensor(w).multinomial(100, replacement=False))
 
@@ -363,6 +372,23 @@ class TestRandomness(unittest.TestCase):
     torch_samples = [torch.tensor(w).multinomial(1, replacement=False).item() for _ in range(1000)]
     self.assertTrue(equal_distribution(lambda *_: Tensor(tiny_samples), lambda _: torch.tensor(torch_samples)))
 
+    w = list(range(32))
+    s1 = Tensor(w).multinomial(5, replacement=False).numpy()
+    self.assertEqual(len(set(s1.tolist())), 5)
+    s2 = Tensor(w).multinomial(5, replacement=False).numpy()
+    self.assertFalse(np.array_equal(s1, s2))
+    full = Tensor(w).multinomial(len(w), replacement=False).numpy()
+    self.assertEqual(sorted(full.tolist()), w)
+
+    w = [0.1, 0.2, 0.3, 0.4]
+    @TinyJit
+    def sample_three(): return Tensor(w).multinomial(3, replacement=False).realize()
+
+    tiny_draws = np.array([sample_three().numpy() for _ in range(1000)])
+    torch_draws = np.array([torch.tensor(w).multinomial(3, replacement=False).numpy() for _ in range(1000)])
+    for pos in range(3):
+      self.assertTrue(equal_distribution(lambda *_: Tensor(tiny_draws[:, pos]), lambda _: torch.tensor(torch_draws[:, pos])))
+
   @unittest.skip("this test is flaky")
   def test_multinomial_counterexample(self):
     tiny_res = Tensor([0.3, 0.6, 0.1]).multinomial(4000, replacement=True)

test/backend/test_renderer_failures.py

@@ -1,9 +1,8 @@
 import unittest
 import numpy as np
-from dataclasses import replace
-from tinygrad.device import Buffer, Device, is_dtype_supported
+from tinygrad.device import Device, is_dtype_supported
 from tinygrad.dtype import dtypes, ConstType
-from tinygrad.engine.realize import CompiledRunner
+from tinygrad.engine.realize import run_linear
 from tinygrad.codegen import to_program
 from tinygrad.helpers import prod
 from tinygrad.renderer.cstyle import CStyleLanguage
@@ -13,17 +12,13 @@ from tinygrad.runtime.ops_python import PythonRenderer
 from tinygrad.uop.ops import UOp, Ops, KernelInfo, python_alu
 from tinygrad.tensor import Tensor, _to_np_dtype
 
-def _test_uop_result(inputs:list[Tensor], prg:UOp, local_size=None):
+def _test_uop_result(inputs:list[Tensor], sink:UOp, local_size=None):
   for x in inputs: x.realize()
-  uops = prg.src[2].src
-  outbufs = [Buffer(Device.DEFAULT, sz:=(1 if local_size is None else prod(local_size)), (dtype:=u.src[1].dtype), \
-      initial_value=np.zeros(sz, dtype=_to_np_dtype(dtype)).data) for u in uops if u.op is Ops.STORE]
-  inbufs = [x.uop.base.buffer for x in inputs]
-  info = prg.arg
-  if local_size is not None: info = replace(info, local_size=tuple(local_size))
-  ei = CompiledRunner(prg.replace(arg=info), Device.DEFAULT)
-  ei.exec(outbufs+inbufs)
-  return [np.frombuffer(x.as_memoryview(), _to_np_dtype(x.dtype)) for x in outbufs]
+  sz = 1 if local_size is None else prod(local_size)
+  outs = [UOp.new_buffer(Device.DEFAULT, sz, u.src[1].dtype) for u in sink.src if u.op is Ops.STORE]
+  for u in outs: u.buffer.allocate().copyin(np.zeros(sz, dtype=_to_np_dtype(u.dtype)).data)
+  run_linear(UOp(Ops.LINEAR, src=(sink.call(*outs, *(x.uop.base for x in inputs)),)))
+  return [u.buffer.numpy() for u in outs]
 
 def _setup_and_test_alu(alu_op:Ops, input_val:ConstType, *alu_src_uops:UOp):
   dtype = alu_src_uops[0].dtype
@@ -33,9 +28,7 @@ def _setup_and_test_alu(alu_op:Ops, input_val:ConstType, *alu_src_uops:UOp):
   ld = b.index(idx)
   alu = ld.alu(alu_op, *alu_src_uops)
   store = UOp.store(a.index(idx), alu)
-  sink = UOp(Ops.SINK, dtypes.void, (store,), arg=KernelInfo())
-  prg = to_program(sink, Device[Device.DEFAULT].renderer)
-  return _test_uop_result([Tensor([input_val])], prg)[0]
+  return _test_uop_result([Tensor([input_val])], UOp(Ops.SINK, dtypes.void, (store,), arg=KernelInfo()))[0]
 
 class TestRendererFailures(unittest.TestCase):
   @unittest.skipIf(not isinstance(Device[Device.DEFAULT].renderer, (PTXRenderer, PythonRenderer)), "test is for ptx or python renderer")
@@ -44,8 +37,7 @@ class TestRendererFailures(unittest.TestCase):
     gate_alu = (lidx0:=UOp(Ops.SPECIAL, dtypes.int, (UOp.const(dtypes.int, 4),), 'lidx0')).ne(0)
     gated_alu_store = UOp(Ops.STORE, dtypes.void, (a.index(lidx0.valid(gate_alu)), UOp.const(dtypes.int, 1)))
     sink = UOp(Ops.SINK, dtypes.void, (gated_alu_store,), arg=KernelInfo())
-    prg = to_program(sink, Device[Device.DEFAULT].renderer)
-    ret = _test_uop_result([], prg, local_size=[4, 1, 1])[0]
+    ret = _test_uop_result([], sink, local_size=[4, 1, 1])[0]
     np.testing.assert_equal(ret, [0, 1, 1, 1])
 
   @unittest.skipIf(not isinstance(Device[Device.DEFAULT].renderer, (PTXRenderer, PythonRenderer)), "test is for ptx or python renderer")
@@ -55,8 +47,7 @@ class TestRendererFailures(unittest.TestCase):
     gate_alu_1 = (lidx1:=UOp(Ops.SPECIAL, dtypes.int, (UOp.const(dtypes.int, 2),), 'lidx1')).ne(0)
     gated_alu_store = UOp(Ops.STORE, dtypes.void, (a.index((lidx0+lidx1*4).valid(gate_alu_0&gate_alu_1)), UOp.const(dtypes.int, 1)))
     sink = UOp(Ops.SINK, dtypes.void, (gated_alu_store,), arg=KernelInfo())
-    prg = to_program(sink, Device[Device.DEFAULT].renderer)
-    ret = _test_uop_result([], prg, local_size=[4, 2, 1])[0]
+    ret = _test_uop_result([], sink, local_size=[4, 2, 1])[0]
     np.testing.assert_equal(ret, [0, 0, 0, 0, 0, 1, 1, 1])
 
 @unittest.skipIf(not isinstance(Device[Device.DEFAULT].renderer, CStyleLanguage), "uops are for cstyle")
@@ -102,8 +93,7 @@ class TestPTXFailures(unittest.TestCase):
     if_uop = UOp(Ops.IF, dtypes.void, (gate_alu,))
     gated_alu_store = UOp(Ops.STORE, dtypes.void, (a.index(lidx0, if_uop), val))
     sink = UOp(Ops.SINK, dtypes.void, (gated_alu_store,), arg=KernelInfo())
-    prg = to_program(sink, Device[Device.DEFAULT].renderer)
-    ret = _test_uop_result([], prg, local_size=[4, 1, 1])[0]
+    ret = _test_uop_result([], sink, local_size=[4, 1, 1])[0]
     np.testing.assert_equal(ret, [0, 1, 1, 1])
 
   @unittest.skipUnless(is_dtype_supported(dtypes.half), "need half")

test/backend/test_uops.py

@@ -5,18 +5,19 @@ from tinygrad.tensor import Tensor, _to_np_dtype
 from tinygrad.helpers import CI, Context
 from tinygrad.dtype import dtypes, DType, AddrSpace, ConstFloat  # noqa: F401
 from tinygrad.device import Buffer, Device
-from tinygrad.uop.ops import Ops, UOp, KernelInfo, AxisType
+from tinygrad.uop.ops import Ops, UOp, KernelInfo, AxisType, buffers
 from tinygrad.renderer.cstyle import CStyleLanguage
-from tinygrad.engine.realize import CompiledRunner, run_linear
+from tinygrad.engine.realize import run_linear
 from tinygrad.codegen import to_program
 from tinygrad.device import is_dtype_supported
 from tinygrad.codegen.opt import Opt, OptOps
 from tinygrad.renderer.ptx import PTXRenderer
 from test.helpers import to_uops_list
 
-def _uops_to_prg(uops_list):
-  prg = to_program(UOp.sink(*uops_list, arg=KernelInfo()), Device[Device.DEFAULT].renderer)
-  return CompiledRunner(prg, Device.DEFAULT)
+def run_uops(uops_list:list[UOp], bufs:list[Buffer]):
+  buf_uops = [UOp.new_buffer(b.device, b.size, b.dtype) for b in bufs]
+  for u,b in zip(buf_uops, bufs): buffers[u] = b
+  run_linear(UOp(Ops.LINEAR, src=(UOp.sink(*uops_list, arg=KernelInfo()).call(*buf_uops),)))
 
 def uop(uops:list[UOp], op:Ops, dtype:Optional[DType], src:tuple[UOp, ...], arg:Any=None) -> UOp:
   if op is Ops.CONST: uops.append(UOp.const(dtype, arg))
@@ -33,8 +34,7 @@ def _test_single_value(vals, op, dts):
   out = uop(uops, Ops.STORE, dtypes.void, (buf_store.index(uop(uops, Ops.CONST, dtypes.int32, (), 0), ptr=True), alu))
   buf = Buffer(Device.DEFAULT, 1, output_dtype).allocate()
   buf2 = [Buffer(Device.DEFAULT, 1, dtype).allocate().copyin(np.array([a], dtype=_to_np_dtype(dtype)).data) for a,dtype in zip(vals, dts)]
-  prg = _uops_to_prg([out])
-  prg.exec([buf]+buf2)
+  run_uops([out], [buf]+buf2)
   ret = np.empty(1, _to_np_dtype(output_dtype))
   buf.copyout(ret.data)
   return ret[0]
@@ -47,8 +47,7 @@ def _test_single_value_const(vals, op, dts):
   alu = uop(uops, op, output_dtype, loads)
   out = buf_store[UOp.const(dtypes.int32, 0)].store(alu)
   buf = Buffer(Device.DEFAULT, 1, output_dtype).allocate()
-  prg = _uops_to_prg([out])
-  prg.exec([buf])
+  run_uops([out], [buf])
   ret = np.empty(1, _to_np_dtype(output_dtype))
   buf.copyout(ret.data)
   return ret[0]
@@ -59,8 +58,7 @@ def _test_uops_result(output_dtype, uops, res):
   # res = output_fn(uops)
   out = uop(uops, Ops.STORE, dtypes.void, (buf_store.index(uop(uops, Ops.CONST, dtypes.int32, (), 0)), res))
   buf = Buffer(Device.DEFAULT, 1, output_dtype).allocate()
-  prg = _uops_to_prg([out])
-  prg.exec([buf])
+  run_uops([out], [buf])
   ret = np.empty(1, _to_np_dtype(output_dtype))
   buf.copyout(ret.data)
   return ret[0]

test/device/test_hcq.py

@@ -6,7 +6,7 @@ from tinygrad.device import Buffer, BufferSpec
 from tinygrad.runtime.support.hcq import HCQCompiled, HCQBuffer
 from tinygrad.runtime.autogen import libc
 from tinygrad.runtime.support.system import PCIIfaceBase
-from tinygrad.engine.realize import get_runner, CompiledRunner
+from tinygrad.engine.realize import get_runtime
 from tinygrad.codegen import to_program
 from tinygrad.codegen.opt import Opt, OptOps
 from tinygrad import Variable
@@ -22,11 +22,12 @@ class TestHCQ(unittest.TestCase):
     TestHCQ.b = self.a + 1
     si = self.b.schedule_linear().src[-1]
 
-    TestHCQ.runner = get_runner(TestHCQ.d0.device, si.src[0])
+    TestHCQ.prg = to_program(si.src[0], TestHCQ.d0.renderer)
+    TestHCQ.runtime = get_runtime(TestHCQ.d0.device, TestHCQ.prg)
     TestHCQ.b.uop.buffer.allocate()
 
-    TestHCQ.kernargs_ba_ptr = TestHCQ.runner._prg.fill_kernargs([TestHCQ.b.uop.buffer._buf, TestHCQ.a.uop.buffer._buf])
-    TestHCQ.kernargs_ab_ptr = TestHCQ.runner._prg.fill_kernargs([TestHCQ.a.uop.buffer._buf, TestHCQ.b.uop.buffer._buf])
+    TestHCQ.kernargs_ba_ptr = TestHCQ.runtime.fill_kernargs([TestHCQ.b.uop.buffer._buf, TestHCQ.a.uop.buffer._buf])
+    TestHCQ.kernargs_ab_ptr = TestHCQ.runtime.fill_kernargs([TestHCQ.a.uop.buffer._buf, TestHCQ.b.uop.buffer._buf])
 
   def setUp(self):
     TestHCQ.d0.synchronize()
@@ -114,7 +115,7 @@ class TestHCQ(unittest.TestCase):
 
   # Test exec
   def test_exec_one_kernel(self):
-    TestHCQ.d0.hw_compute_queue_t().exec(TestHCQ.runner._prg, TestHCQ.kernargs_ba_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size) \
+    TestHCQ.d0.hw_compute_queue_t().exec(TestHCQ.runtime, TestHCQ.kernargs_ba_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size) \
                                    .signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value).submit(TestHCQ.d0)
 
     TestHCQ.d0.timeline_signal.wait(TestHCQ.d0.timeline_value)
@@ -128,8 +129,8 @@ class TestHCQ(unittest.TestCase):
 
     q = TestHCQ.d0.hw_compute_queue_t()
     q.wait(TestHCQ.d0.timeline_signal, virt_val - 1) \
-     .exec(TestHCQ.runner._prg, TestHCQ.kernargs_ba_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size) \
-     .exec(TestHCQ.runner._prg, TestHCQ.kernargs_ab_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size) \
+     .exec(TestHCQ.runtime, TestHCQ.kernargs_ba_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size) \
+     .exec(TestHCQ.runtime, TestHCQ.kernargs_ab_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size) \
      .signal(TestHCQ.d0.timeline_signal, virt_val)
 
     for _ in range(100):
@@ -141,11 +142,11 @@ class TestHCQ(unittest.TestCase):
 
   @unittest.skipIf(Device.DEFAULT in {"CPU"}, "No globals/locals on LLVM/CPU")
   def test_exec_update(self):
-    sint_global = (Variable("sint_global", 0, 0xffffffff, dtypes.uint32),) + tuple(TestHCQ.runner.p.global_size[1:])
-    sint_local = (Variable("sint_local", 0, 0xffffffff, dtypes.uint32),) + tuple(TestHCQ.runner.p.local_size[1:])
+    sint_global = (Variable("sint_global", 0, 0xffffffff, dtypes.uint32),) + tuple(TestHCQ.prg.arg.global_size[1:])
+    sint_local = (Variable("sint_local", 0, 0xffffffff, dtypes.uint32),) + tuple(TestHCQ.prg.arg.local_size[1:])
 
     q = TestHCQ.d0.hw_compute_queue_t()
-    q.exec(TestHCQ.runner._prg, TestHCQ.kernargs_ba_ptr, sint_global, sint_local) \
+    q.exec(TestHCQ.runtime, TestHCQ.kernargs_ba_ptr, sint_global, sint_local) \
      .signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value)
 
     q.submit(TestHCQ.d0, {sint_global[0].expr: 1, sint_local[0].expr: 1})
@@ -166,17 +167,17 @@ class TestHCQ(unittest.TestCase):
     b = a + 1
     si = b.schedule_linear().src[-1]
 
-    runner = CompiledRunner(to_program(replace_opts(si.src[0], [Opt(op=OptOps.LOCAL, axis=0, arg=3) for _ in range(3)]), TestHCQ.d0.renderer),
-                            Device.DEFAULT)
+    prg = to_program(replace_opts(si.src[0], [Opt(op=OptOps.LOCAL, axis=0, arg=3) for _ in range(3)]), TestHCQ.d0.renderer)
+    runtime = get_runtime(Device.DEFAULT, prg)
 
     zb = Buffer(Device.DEFAULT, 3 * 3 * 3, dtypes.int, options=BufferSpec(cpu_access=True, nolru=True)).ensure_allocated()
     zt = Buffer(Device.DEFAULT, 3 * 3 * 3, dtypes.int, options=BufferSpec(cpu_access=True, nolru=True)).ensure_allocated()
     ctypes.memset(zb._buf.va_addr, 0, zb.nbytes)
-    kernargs = runner._prg.fill_kernargs([zt._buf, zb._buf])
+    kernargs = runtime.fill_kernargs([zt._buf, zb._buf])
 
     q = TestHCQ.d0.hw_compute_queue_t()
     q.memory_barrier() \
-     .exec(runner._prg, kernargs, (1,1,1), virt_local) \
+     .exec(runtime, kernargs, (1,1,1), virt_local) \
      .signal(TestHCQ.d0.timeline_signal, virt_val)
 
     for x in range(1, 4):
@@ -330,7 +331,7 @@ class TestHCQ(unittest.TestCase):
   def test_speed_exec_time(self):
     sig_st, sig_en = TestHCQ.d0.new_signal(), TestHCQ.d0.new_signal()
     TestHCQ.d0.hw_compute_queue_t().timestamp(sig_st) \
-                                   .exec(TestHCQ.runner._prg, TestHCQ.kernargs_ba_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size) \
+                                   .exec(TestHCQ.runtime, TestHCQ.kernargs_ba_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size) \
                                    .timestamp(sig_en) \
                                    .signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value).submit(TestHCQ.d0)
 
@@ -470,12 +471,13 @@ class TestHCQ(unittest.TestCase):
   def test_memory_barrier(self):
     a = Tensor([0, 1], device=Device.DEFAULT, dtype=dtypes.int8).realize()
     b = a + 1
-    runner = get_runner(TestHCQ.d0.device, b.schedule_linear().src[-1].src[0])
+    prg = to_program(b.schedule_linear().src[-1].src[0], TestHCQ.d0.renderer)
+    runtime = get_runtime(TestHCQ.d0.device, prg)
 
     buf1 = Buffer(Device.DEFAULT, 2, dtypes.int8, options=BufferSpec(nolru=True)).ensure_allocated()
     buf2 = Buffer(Device.DEFAULT, 2, dtypes.int8, options=BufferSpec(cpu_access=True, nolru=True)).ensure_allocated()
 
-    kernargs_ptr = runner._prg.fill_kernargs([buf1._buf, buf2._buf])
+    kernargs_ptr = runtime.fill_kernargs([buf1._buf, buf2._buf])
 
     for i in range(255):
       ctypes.memset(buf2._buf.va_addr, i, 2)
@@ -483,7 +485,7 @@ class TestHCQ(unittest.TestCase):
       # Need memory_barrier after direct write to vram
       TestHCQ.d0.hw_compute_queue_t().wait(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value - 1) \
                                      .memory_barrier() \
-                                     .exec(runner._prg, kernargs_ptr, runner.p.global_size, runner.p.local_size) \
+                                     .exec(runtime, kernargs_ptr, prg.arg.global_size, prg.arg.local_size) \
                                      .signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value).submit(TestHCQ.d0)
       TestHCQ.d0.timeline_signal.wait(TestHCQ.d0.timeline_value)
       TestHCQ.d0.timeline_value += 1

test/device/test_metal.py

@@ -50,19 +50,6 @@ kernel void r_5(device int* data0, const device int* data1, uint3 gid [[threadgr
       compiled = compiled[:40] # corrupt the compiled program
       MetalProgram(device, "r_5", compiled)
 
-  def test_wait_skips_in_flight(self):
-    device = MetalDevice("metal")
-    compiled = MetalCompiler().compile("""
-#include <metal_stdlib>
-kernel void noop(uint3 gid [[threadgroup_position_in_grid]], uint3 lid [[thread_position_in_threadgroup]]) {}
-""")
-    prg = MetalProgram(device, "noop", compiled)
-    self.assertIsInstance(prg(wait=True), float)
-    self.assertEqual(device.mtl_buffers_in_flight, [])
-    self.assertIsNone(prg(wait=False))
-    self.assertEqual(len(device.mtl_buffers_in_flight), 1)
-    device.synchronize()
-
   def test_free(self):
     size = 2**16
     device = Device['METAL']

test/external/external_benchmark_op_conv.py

@@ -3,7 +3,7 @@ from dataclasses import replace
 from tinygrad import dtypes, Device
 from tinygrad.uop.ops import UOp, AxisType, Ops, KernelInfo
 from tinygrad.codegen.opt import Opt, OptOps # pylint: disable=unused-import
-from tinygrad.engine.realize import CompiledRunner
+from tinygrad.engine.realize import get_runtime
 from tinygrad.codegen import to_program
 from tinygrad.helpers import dedup, getenv
 from tinygrad.device import Buffer
@@ -90,12 +90,13 @@ renderer = Device.default.renderer
 allocator = Device.default.allocator
 
 ps = to_program(ast, renderer)
-cr = CompiledRunner(ps, Device.DEFAULT)
+rt = get_runtime(Device.DEFAULT, ps)
 
 gs = sorted(dedup([u for u in ast.toposort() if u.op is Ops.PARAM]), key=lambda u: u.arg)
 # print(len(gs))
 # print([g.dtype for g in gs])
 bufs = [Buffer(ps.arg.device, g.size, g.dtype if isinstance(g.dtype, ImageDType) else g.dtype._base).ensure_allocated() for g in gs]
 
-t = cr(bufs, wait=True)
+gsize, lsize = ps.arg.launch_dims({})
+t = rt(*[b._buf for b in bufs], global_size=gsize, local_size=lsize, vals=ps.arg.vals({}), wait=True)
 print(f"{t*1e6:.2f} us")

test/external/external_fuzz_beam_timeout_recovery.py

@@ -3,7 +3,7 @@
 Stress test for beam timeout + device recovery on AM devices.
 
 Usage:
-  DEV=AMD python test/external/external_test_beam_timeout_recovery.py
+  DEV=AMD python test/external/external_fuzz_beam_timeout_recovery.py
 """
 from tinygrad import Tensor, Device
 from tinygrad.helpers import Context

test/external/external_test_hcq.py

@@ -2,7 +2,8 @@ import unittest, ctypes, struct, time, array
 from tinygrad import Device, Tensor, dtypes
 from tinygrad.helpers import to_mv, CI
 from tinygrad.device import Buffer, BufferSpec
-from tinygrad.engine.realize import get_runner
+from tinygrad.engine.realize import get_runtime
+from tinygrad.codegen import to_program
 
 def _time_queue(q, d):
   st = time.perf_counter()
@@ -21,13 +22,14 @@ class TestHCQ(unittest.TestCase):
     TestHCQ.a = Tensor([0.,1.], device=Device.DEFAULT).realize()
     TestHCQ.b = self.a + 1
     linear = self.b.schedule_linear()
-    TestHCQ.runner = get_runner(TestHCQ.d0.device, linear.src[-1].src[0])
+    TestHCQ.prg = to_program(linear.src[-1].src[0], TestHCQ.d0.renderer)
+    TestHCQ.runtime = get_runtime(TestHCQ.d0.device, TestHCQ.prg)
     TestHCQ.b.uop.buffer.allocate()
     # wow that's a lot of abstraction layers
     TestHCQ.addr = struct.pack("QQ", TestHCQ.b.uop.buffer._buf, TestHCQ.a.uop.buffer._buf)
     TestHCQ.addr2 = struct.pack("QQ", TestHCQ.a.uop.buffer._buf, TestHCQ.b.uop.buffer._buf)
-    TestHCQ.kernargs_off = TestHCQ.runner._prg.kernargs_offset
-    TestHCQ.kernargs_size = TestHCQ.runner._prg.kernargs_alloc_size
+    TestHCQ.kernargs_off = TestHCQ.runtime.kernargs_offset
+    TestHCQ.kernargs_size = TestHCQ.runtime.kernargs_alloc_size
     ctypes.memmove(TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_off, TestHCQ.addr, len(TestHCQ.addr))
     ctypes.memmove(TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size+TestHCQ.kernargs_off, TestHCQ.addr2, len(TestHCQ.addr2))
 
@@ -38,8 +40,8 @@ class TestHCQ(unittest.TestCase):
     elif Device.DEFAULT == "NV":
       from tinygrad.runtime.ops_nv import HWQueue, HWQueue
       # nv need to copy constbuffer there as well
-      to_mv(TestHCQ.d0.kernargs_ptr, 0x160).cast('I')[:] = array.array('I', TestHCQ.runner._prg.constbuffer_0)
-      to_mv(TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, 0x160).cast('I')[:] = array.array('I', TestHCQ.runner._prg.constbuffer_0)
+      to_mv(TestHCQ.d0.kernargs_ptr, 0x160).cast('I')[:] = array.array('I', TestHCQ.runtime.constbuffer_0)
+      to_mv(TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, 0x160).cast('I')[:] = array.array('I', TestHCQ.runtime.constbuffer_0)
       TestHCQ.compute_queue = HWQueue
       TestHCQ.copy_queue = HWQueue
 
@@ -53,11 +55,11 @@ class TestHCQ(unittest.TestCase):
     temp_signal, temp_value = TestHCQ.d0._alloc_signal(value=0), 0
     q = TestHCQ.compute_queue()
     for _ in range(1000):
-      q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
+      q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
       q.signal(temp_signal, temp_value + 1).wait(temp_signal, temp_value + 1)
       temp_value += 1
 
-      q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
+      q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
       q.signal(temp_signal, temp_value + 1).wait(temp_signal, temp_value + 1)
       temp_value += 1
 
@@ -71,10 +73,10 @@ class TestHCQ(unittest.TestCase):
   def test_run_1000_times(self):
     temp_signal = TestHCQ.d0._alloc_signal(value=0)
     q = TestHCQ.compute_queue()
-    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
+    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
     q.signal(temp_signal, 2).wait(temp_signal, 2)
-    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, TestHCQ.runner.p.global_size,
-           TestHCQ.runner.p.local_size)
+    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, TestHCQ.prg.arg.global_size,
+           TestHCQ.prg.arg.local_size)
     for _ in range(1000):
       TestHCQ.d0._set_signal(temp_signal, 1)
       q.submit(TestHCQ.d0)
@@ -87,11 +89,11 @@ class TestHCQ(unittest.TestCase):
   def test_run_to_3(self):
     temp_signal = TestHCQ.d0._alloc_signal(value=0)
     q = TestHCQ.compute_queue()
-    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
+    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
     q.signal(temp_signal, 1).wait(temp_signal, 1)
-    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
+    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
     q.signal(temp_signal, 2).wait(temp_signal, 2)
-    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
+    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
     q.signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value).submit(TestHCQ.d0)
     TestHCQ.d0._wait_signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value)
     TestHCQ.d0.timeline_value += 1
@@ -101,7 +103,7 @@ class TestHCQ(unittest.TestCase):
   def test_update_exec(self):
     q = TestHCQ.compute_queue()
     exec_cmd_idx = len(q)
-    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
+    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
     q.update_exec(exec_cmd_idx, (1,1,1), (1,1,1))
     q.signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value).submit(TestHCQ.d0)
     TestHCQ.d0._wait_signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value)
@@ -115,10 +117,10 @@ class TestHCQ(unittest.TestCase):
   def test_bind_run(self):
     temp_signal = TestHCQ.d0._alloc_signal(value=0)
     q = TestHCQ.compute_queue()
-    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
+    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
     q.signal(temp_signal, 2).wait(temp_signal, 2)
-    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, TestHCQ.runner.p.global_size,
-           TestHCQ.runner.p.local_size)
+    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, TestHCQ.prg.arg.global_size,
+           TestHCQ.prg.arg.local_size)
     q.bind(TestHCQ.d0)
     for _ in range(1000):
       TestHCQ.d0._set_signal(temp_signal, 1)
@@ -133,7 +135,7 @@ class TestHCQ(unittest.TestCase):
   def test_update_exec_binded(self):
     q = TestHCQ.compute_queue()
     exec_ptr = q.ptr()
-    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
+    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
     q.signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value)
     q.bind(TestHCQ.d0)
 
@@ -170,7 +172,7 @@ class TestHCQ(unittest.TestCase):
 
   def test_run_normal(self):
     q = TestHCQ.compute_queue()
-    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
+    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
     q.signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value).submit(TestHCQ.d0)
     TestHCQ.d0._wait_signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value)
     TestHCQ.d0.timeline_value += 1
@@ -201,7 +203,7 @@ class TestHCQ(unittest.TestCase):
 
   def test_run_signal(self):
     q = TestHCQ.compute_queue()
-    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
+    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
     q.signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value)
     q.submit(TestHCQ.d0)
     TestHCQ.d0._wait_signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value)
@@ -278,7 +280,7 @@ class TestHCQ(unittest.TestCase):
   def test_interleave_compute_and_copy(self):
     q = TestHCQ.compute_queue()
     qc = TestHCQ.copy_queue()
-    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)  # b = [1, 2]
+    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)  # b = [1, 2]
     q.signal(sig:=TestHCQ.d0._alloc_signal(value=0), value=1)
     qc.wait(sig, value=1)
     qc.copy(TestHCQ.a.uop.buffer._buf, TestHCQ.b.uop.buffer._buf, 8)
@@ -315,7 +317,7 @@ class TestHCQ(unittest.TestCase):
     for _ in range(40):
       q = TestHCQ.compute_queue()
       q.wait(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value - 1)
-      q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
+      q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
       q.signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value).submit(TestHCQ.d0)
       TestHCQ.d0._wait_signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value)
       TestHCQ.d0.timeline_value += 1

test/external/external_test_speed_llama.py

@@ -5,7 +5,7 @@ from tinygrad.tensor import Tensor
 from tinygrad import Device
 from tinygrad.nn.state import get_state_dict
 from tinygrad.device import Allocator, Compiled
-from tinygrad.engine.realize import method_cache
+from tinygrad.codegen import to_program_cache
 from tinygrad.helpers import Profiling
 
 class FakeProgram:
@@ -31,8 +31,8 @@ class TestLLaMASpeed(unittest.TestCase):
     for v in get_state_dict(model).values(): v.assign(Tensor.empty(*v.shape, dtype=v.dtype))
     print("assigned empty tensors, doing warmup")
 
-    def run_llama(st, empty_method_cache=True):
-      if empty_method_cache: method_cache.clear()
+    def run_llama(st, empty_cache=True):
+      if empty_cache: to_program_cache.clear()
       tms = [time.perf_counter()]
       for i in range(5):
         model(Tensor([[1,2,3,4]]), i).realize()

test/external/external_uop_gc.py

@@ -1,7 +1,6 @@
 import gc
 from tinygrad import Tensor, UOp, Device, nn
 from tinygrad.schedule import schedule_cache
-from tinygrad.engine.realize import method_cache
 from tinygrad.codegen import to_program, to_program_cache
 from tinygrad.schedule.indexing import apply_movement_op, _apply_reshape
 from tinygrad.uop.divandmod import fold_divmod_general
@@ -71,7 +70,6 @@ if __name__ == "__main__":
 
     # these caches will keep uops alive
     schedule_cache.clear()
-    method_cache.clear()
     to_program_cache.clear()
     apply_movement_op.cache_clear()
     _apply_reshape.cache_clear()

test/mockgpu/amd/emu.py

@@ -53,10 +53,11 @@ class _MXCSRContext:
 
 from tinygrad.uop.ops import UOp, Ops, KernelInfo, AxisType
 from tinygrad.dtype import dtypes, AddrSpace
-from tinygrad.device import Buffer, BufferSpec
+from tinygrad.device import Buffer, BufferSpec, Device
 from tinygrad.runtime.autogen import hsa
 from tinygrad.helpers import Context, DEBUG, PROFILE, colored
-from tinygrad.engine.realize import get_runner
+from tinygrad.engine.realize import get_runtime
+from tinygrad.codegen import to_program
 
 from tinygrad.renderer.amd import decode_inst
 from tinygrad.runtime.autogen.amd.rdna3.str_pcode import PCODE as PCODE_RDNA3
@@ -374,22 +375,24 @@ def _mem_store(mem: UOp, addr: UOp, val: UOp, active: UOp, addr_bits: int = 32,
   """Conditional memory store with sub-word support. Returns list of store UOps."""
   adt = dtypes.uint64 if addr_bits == 64 else dtypes.uint32
   word_addr = addr >> UOp.const(adt, 2)
-  idx = mem.index(word_addr.cast(dtypes.int), active)
-  if data_bits == 32: return [idx.store(active.where(_to_u32(val), idx))]
+  bidx = mem.index(word_addr.cast(dtypes.int), ptr=True)
+  if data_bits == 32: return [UOp(Ops.STORE, dtypes.void, (bidx, _to_u32(val), active))]
   # Sub-word store: read-modify-write with mask
+  cur = bidx.load(active, _c(0, dtypes.uint32))
   byte_pos = addr.cast(dtypes.uint32) & _c(3)
   byte_shift = byte_pos * _c(8)
   val_u32, size_mask = val.cast(dtypes.uint32), _c(0xFF if data_bits == 8 else 0xFFFF)
   mask = size_mask << byte_shift
-  new_word = (idx & (mask ^ _c(0xFFFFFFFF))) | ((val_u32 & size_mask) << byte_shift)
-  if data_bits == 8: return [idx.store(active.where(new_word, idx))]
+  new_word = (cur & (mask ^ _c(0xFFFFFFFF))) | ((val_u32 & size_mask) << byte_shift)
+  if data_bits == 8: return [UOp(Ops.STORE, dtypes.void, (bidx, new_word, active))]
   # 16-bit cross-word case: byte_pos == 3 means value spans two words
   is_cross = byte_pos.eq(_c(3))
-  cross_word0 = (idx & _c(0x00FFFFFF)) | ((val_u32 & _c(0xFF)) << _c(24))
-  store0 = idx.store(active.where(is_cross.where(cross_word0, new_word), idx))
-  next_idx = mem.index((word_addr + UOp.const(adt, 1)).cast(dtypes.int), active & is_cross)
-  cross_word1 = (next_idx & _c(0xFFFFFF00)) | ((val_u32 >> _c(8)) & _c(0xFF))
-  return [store0, next_idx.store((active & is_cross).where(cross_word1, next_idx))]
+  cross_word0 = (cur & _c(0x00FFFFFF)) | ((val_u32 & _c(0xFF)) << _c(24))
+  store0 = UOp(Ops.STORE, dtypes.void, (bidx, is_cross.where(cross_word0, new_word), active))
+  next_bidx = mem.index((word_addr + UOp.const(adt, 1)).cast(dtypes.int), ptr=True)
+  next_cur = next_bidx.load(active & is_cross, _c(0, dtypes.uint32))
+  cross_word1 = (next_cur & _c(0xFFFFFF00)) | ((val_u32 >> _c(8)) & _c(0xFF))
+  return [store0, UOp(Ops.STORE, dtypes.void, (next_bidx, cross_word1, active & is_cross))]
 
 def _mem_store_bytes(mem: UOp, addr: UOp, val: UOp, active: UOp, data_bits: int = 32) -> list[UOp]:
   """Store to byte-addressable memory (scratch). addr is byte offset, mem is uint8 buffer."""
@@ -397,7 +400,8 @@ def _mem_store_bytes(mem: UOp, addr: UOp, val: UOp, active: UOp, data_bits: int
   val_u32 = val.cast(dtypes.uint32) if val.dtype != dtypes.uint32 else val
   for i in range(data_bits // 8):
     byte_val = (val_u32 >> UOp.const(dtypes.uint32, i * 8)) & UOp.const(dtypes.uint32, 0xFF)
-    stores.append(mem.index((addr + UOp.const(dtypes.uint64, i)).cast(dtypes.int), active).store(byte_val.cast(dtypes.uint8)))
+    bidx = mem.index((addr + UOp.const(dtypes.uint64, i)).cast(dtypes.int), ptr=True)
+    stores.append(UOp(Ops.STORE, dtypes.void, (bidx, byte_val.cast(dtypes.uint8), active)))
   return stores
 
 def _collect_data_slices(assigns: list[tuple[str, UOp]], data_prefix: str, pcode_vars: dict | None = None, op_name: str = "") -> dict[int, UOp]:
@@ -515,14 +519,15 @@ class _Ctx:
   # Dynamic register access (takes UOp index instead of int)
   def rsgpr_dyn(self, reg: UOp, valid: UOp | None = None) -> UOp:
     """Read SGPR with dynamic register index."""
-    if valid is not None: return self.sgpr.index(reg.cast(dtypes.int), valid, ptr=True).load()
+    if valid is not None: return self.sgpr.index(reg.cast(dtypes.int), ptr=True).load(valid, _c(0, dtypes.uint32))
     return self.sgpr.index(reg.cast(dtypes.int), ptr=True).load()
 
   def wsgpr_dyn(self, reg: UOp, val: UOp) -> UOp:
     """Write SGPR with dynamic register index. On RDNA, index 124 = NULL (writes discarded). On CDNA, index 124 = M0 (read/write)."""
     # RDNA: NULL (124) discards writes. CDNA: M0 (124) is writable.
     valid = None if self.wave_size == 64 else reg.ne(_c(124))
-    return self.sgpr.index(reg.cast(dtypes.int), valid).store(val.cast(dtypes.uint32))
+    bidx = self.sgpr.index(reg.cast(dtypes.int), ptr=True)
+    return UOp(Ops.STORE, dtypes.void, (bidx, val.cast(dtypes.uint32))+((valid,) if valid is not None else ()))
 
   def wmask(self, reg: UOp, val: UOp) -> list[UOp]:
     """Write a lane mask (VCC/EXEC). Splits into lo/hi for wave64."""
@@ -539,24 +544,26 @@ class _Ctx:
   def rvgpr_dyn(self, reg: UOp, lane: UOp, valid: UOp | None = None) -> UOp:
     """Read VGPR with dynamic register index."""
     idx = reg.cast(dtypes.int) * _c(self.wave_size, dtypes.int) + lane.cast(dtypes.int)
-    return self.vgpr.index(idx, valid, ptr=True).load() if valid is not None else self.vgpr.index(idx, ptr=True).load()
+    if valid is not None: return self.vgpr.index(idx, ptr=True).load(valid, _c(0, dtypes.uint32))
+    return self.vgpr.index(idx, ptr=True).load()
 
   def wvgpr_dyn(self, reg: UOp, lane: UOp, val: UOp, exec_mask: UOp, after: UOp | None = None) -> UOp:
     """Write VGPR with dynamic register index."""
     buf = self.vgpr.after(after) if after is not None else self.vgpr
     offset = reg.cast(dtypes.int) * _c(self.wave_size, dtypes.int) + lane.cast(dtypes.int)
-    return buf.index(offset, _lane_active(exec_mask, lane)).store(val.cast(dtypes.uint32))
+    return UOp(Ops.STORE, dtypes.void, (buf.index(offset, ptr=True), val.cast(dtypes.uint32), _lane_active(exec_mask, lane)))
 
   def raccvgpr_dyn(self, reg: UOp, lane: UOp, valid: UOp | None = None) -> UOp:
     """Read ACCVGPR with dynamic register index (CDNA only)."""
     idx = reg.cast(dtypes.int) * _c(self.wave_size, dtypes.int) + lane.cast(dtypes.int)
-    return self.accvgpr.index(idx, valid, ptr=True).load() if valid is not None else self.accvgpr.index(idx, ptr=True).load()
+    if valid is not None: return self.accvgpr.index(idx, ptr=True).load(valid, _c(0, dtypes.uint32))
+    return self.accvgpr.index(idx, ptr=True).load()
 
   def waccvgpr_dyn(self, reg: UOp, lane: UOp, val: UOp, exec_mask: UOp, after: UOp | None = None) -> UOp:
     """Write ACCVGPR with dynamic register index (CDNA only)."""
     buf = self.accvgpr.after(after) if after is not None else self.accvgpr
     offset = reg.cast(dtypes.int) * _c(self.wave_size, dtypes.int) + lane.cast(dtypes.int)
-    return buf.index(offset, _lane_active(exec_mask, lane)).store(val.cast(dtypes.uint32))
+    return UOp(Ops.STORE, dtypes.void, (buf.index(offset, ptr=True), val.cast(dtypes.uint32), _lane_active(exec_mask, lane)))
 
   def rsrc_dyn(self, off: UOp, lane: UOp | None, bits: int = 32, literal: UOp | None = None, is_f64: bool = False, do_cast: bool = True) -> UOp:
     """Read source operand with dynamic offset. Handles SGPR/inline constants (<256), VGPR (>=256).
@@ -712,7 +719,7 @@ class _Ctx:
         old = self.vgpr.index(val[0].cast(dtypes.int), ptr=True).load()
         new_val = _set_bits(old, _val_to_bits(val[1]), width, lo_bit).cast(dtypes.uint32)
         active = _lane_active(exec_mask, lane)
-        raw_stores.append(('vgpr_direct', self.vgpr.index(val[0].cast(dtypes.int), active).store(new_val)))
+        raw_stores.append(('vgpr_direct', UOp(Ops.STORE, dtypes.void, (self.vgpr.index(val[0].cast(dtypes.int), ptr=True), new_val, active))))
         continue
       if 'D0' in dest and '[laneId]' in dest:
         old_vcc = self.rmask(_c(VCC_LO.offset))
@@ -1846,15 +1853,16 @@ def _compile_mem_op(inst: ir3.DS|ir3.FLAT|ir3.GLOBAL|ir3.SCRATCH|ir4.DS|ir4.VFLA
     if data_bits < 32:
       # Sub-dword LDS write: read-modify-write within the uint32 slot
       word_addr = (addr >> addr_shift).cast(dtypes.int)
-      idx = mem.index(word_addr, active)
+      bidx = mem.index(word_addr, ptr=True)
+      cur = bidx.load(active, _c(0, dtypes.uint32))
       byte_pos = addr.cast(dtypes.uint32) & _c(3)
       byte_shift = byte_pos * _c(8)
       size_mask = _c(0xFF if data_bits == 8 else 0xFFFF)
       mask = size_mask << byte_shift
-      new_word = (idx & (mask ^ _c(0xFFFFFFFF))) | ((val.cast(dtypes.uint32) & size_mask) << byte_shift)
-      return idx.store(active.where(new_word, idx))
-    idx = mem.index((addr >> addr_shift).cast(dtypes.int))
-    return idx.store(active.where(val, idx.load()))
+      new_word = (cur & (mask ^ _c(0xFFFFFFFF))) | ((val.cast(dtypes.uint32) & size_mask) << byte_shift)
+      return UOp(Ops.STORE, dtypes.void, (bidx, new_word, active))
+    bidx = mem.index((addr >> addr_shift).cast(dtypes.int), ptr=True)
+    return UOp(Ops.STORE, dtypes.void, (bidx, val, active))
 
   def make_srcs(lane: UOp) -> dict:
     addr = make_addr(lane)
@@ -2004,17 +2012,18 @@ def _compile_mubuf(inst: irc.MUBUF, ctx: _Ctx) -> UOp:
       word_addr = (addr + UOp.const(dtypes.uint64, i * 4)) >> UOp.const(dtypes.uint64, 2)
       val = in_bounds.where(mem.index(word_addr.cast(dtypes.int64), ptr=True).load(), _c(0))
       lds_idx = ((lds_addr + _c(i * 4)) >> _c(2)).cast(dtypes.int)
-      stores.append(ctx.lds.index(lds_idx, active).store(active.where(val, ctx.lds.index(lds_idx, active))))
+      bidx = ctx.lds.index(lds_idx, ptr=True)
+      stores.append(UOp(Ops.STORE, dtypes.void, (bidx, val, active)))
   elif is_store:
     for i in range(n_dwords):
       word_addr = (addr + UOp.const(dtypes.uint64, i * 4)) >> UOp.const(dtypes.uint64, 2)
-      idx = mem.index(word_addr.cast(dtypes.int64), in_bounds)
+      idx = mem.index(word_addr.cast(dtypes.int64), ptr=True)
       val = (ctx.raccvgpr_dyn if use_acc else ctx.rvgpr_dyn)(vdata + _c(i), lane)
-      stores.append(idx.store(in_bounds.where(_to_u32(val), idx)))
+      stores.append(UOp(Ops.STORE, dtypes.void, (idx, _to_u32(val), in_bounds)))
   else:
     for i in range(n_dwords):
       word_addr = (addr + UOp.const(dtypes.uint64, i * 4)) >> UOp.const(dtypes.uint64, 2)
-      val = in_bounds.where(mem.index(word_addr.cast(dtypes.int64), in_bounds, ptr=True).load(), _c(0))
+      val = mem.index(word_addr.cast(dtypes.int64), ptr=True).load(in_bounds, _c(0, dtypes.uint32))
       stores.append((ctx.waccvgpr_dyn if use_acc else ctx.wvgpr_dyn)(vdata + _c(i), lane, val, exec_mask))
   return UOp.sink(UOp.group(*stores).end(lane), *ctx.inc_pc())
 
@@ -2045,18 +2054,18 @@ _INST_HANDLERS: dict[type, Callable[..., UOp]] = {
 # PROGRAM DECODE AND COMPILATION
 # ═══════════════════════════════════════════════════════════════════════════════
 
-_canonical_runner_cache: list[tuple[type, int, int, int, object]] = []  # [(inst_type, base, mask, size, runner), ...]
+_canonical_runner_cache: list[tuple[type, int, int, int, tuple[UOp, object]]] = []  # [(inst_type, base, mask, size, (prg, runtime)), ...]
 
 @functools.cache
 def _get_runner(inst_bytes: bytes, arch: str = "rdna3"):
-  """Build and compile instruction to CompiledRunner. Cached by instruction bytes, with canonical dedup."""
+  """Build and compile instruction to (prg, runtime). Cached by instruction bytes, with canonical dedup."""
   inst = decode_inst(inst_bytes, arch)
   inst_size = inst.size()
   inst_int = int.from_bytes(inst_bytes[:inst_size], 'little')
 
   # Check if instruction matches any cached canonical pattern (must also match instruction type to avoid variant conflicts)
-  for inst_type, base, mask, size, runner in _canonical_runner_cache:
-    if type(inst) is inst_type and inst_size == size and (inst_int & mask) == base: return runner
+  for inst_type, base, mask, size, entry in _canonical_runner_cache:
+    if type(inst) is inst_type and inst_size == size and (inst_int & mask) == base: return entry
 
   # Look up handler by type, falling back to base classes for _LIT variants
   handler = _INST_HANDLERS.get(type(inst))
@@ -2075,9 +2084,10 @@ def _get_runner(inst_bytes: bytes, arch: str = "rdna3"):
 
   # NOTE: renderer output is not reproducible because of _MXCSRContext. PROFILE=0 prevents emulator instruction runners from polluting profiling.
   with Context(NOOPT=1, CHECK_OOB=0, TUPLE_ORDER=0, EMULATED_DTYPES="", CAPTURE_PROCESS_REPLAY=0, PROFILE=0):
-    runner = get_runner('CPU', sink)
-  _canonical_runner_cache.append((type(inst), base, mask, size, runner))
-  return runner
+    prg = to_program(sink, Device['CPU'].renderer)
+    runtime = get_runtime('CPU', prg)
+  _canonical_runner_cache.append((type(inst), base, mask, size, (prg, runtime)))
+  return prg, runtime
 
 _BARRIER_OPS = {ir3.SOPPOp.S_BARRIER, irc.SOPPOp.S_BARRIER}
 if hasattr(ir4.SOPPOp, 'S_BARRIER_WAIT'): _BARRIER_OPS.add(ir4.SOPPOp.S_BARRIER_WAIT)
@@ -2208,10 +2218,10 @@ def run_asm(lib: int, lib_sz: int, gx: int, gy: int, gz: int, lx: int, ly: int,
   def _ensure_compiled(pc: int) -> tuple[Callable, list[int], bool, Inst]:
     if pc not in program:
       prev_len = len(_canonical_runner_cache)
-      runner, inst = _decode_at(pc, arch)
+      (prg, runtime), inst = _decode_at(pc, arch)
       is_barrier = (isinstance(inst, (ir3.SOPP, ir4.SOPP, irc.SOPP)) and inst.op in _BARRIER_OPS) or \
                    (isinstance(inst, (ir4.SOP1,)) and inst.op in _BARRIER_SOP1_OPS)
-      program[pc] = (runner._prg.fxn, runner.p.globals, is_barrier, inst)
+      program[pc] = (runtime.fxn, prg.arg.globals, is_barrier, inst)
       if DEBUG >= 3:
         msg = f"[emu] PC={pc - lib}: {inst!r}"
         print(colored(msg, 'green') if len(_canonical_runner_cache) > prev_len else msg)

test/mockgpu/amd/pcode.py

@@ -828,28 +828,30 @@ class Parser:
     assert mem is not None, "memory load requires _vmem or _lds"
     adt = dtypes.uint64 if addr.dtype == dtypes.uint64 else dtypes.uint32
     active = self.vars.get('_active')
-    gate = (active,) if active is not None else ()
+    # gate now lives on LOAD; helper to construct gated load with 0 alt
+    def _gload(bidx, dtype):
+      return bidx.load(active, _const(dtype.base, 0)) if active is not None else bidx.load()
     byte_mem = mem.dtype.base == dtypes.uint8
     if byte_mem:
       idx = addr.cast(dtypes.int)
       if dt in (dtypes.uint64, dtypes.int64, dtypes.float64):
         val = _u32(0).cast(dtypes.uint64)
-        for i in range(8): val = val | (mem.index(idx + _const(dtypes.int, i), *gate, ptr=True).load().cast(dtypes.uint64) << _u64(i * 8))
+        for i in range(8): val = val | (_gload(mem.index(idx + _const(dtypes.int, i), ptr=True), mem.dtype).cast(dtypes.uint64) << _u64(i * 8))
       elif dt in (dtypes.uint8, dtypes.int8):
-        val = mem.index(idx, *gate, ptr=True).load().cast(dt)
+        val = _gload(mem.index(idx, ptr=True), mem.dtype).cast(dt)
       elif dt in (dtypes.uint16, dtypes.int16, dtypes.short):
-        lo = mem.index(idx, *gate, ptr=True).load().cast(dtypes.uint32)
-        hi = mem.index(idx + _const(dtypes.int, 1), *gate, ptr=True).load().cast(dtypes.uint32)
+        lo = _gload(mem.index(idx, ptr=True), mem.dtype).cast(dtypes.uint32)
+        hi = _gload(mem.index(idx + _const(dtypes.int, 1), ptr=True), mem.dtype).cast(dtypes.uint32)
         val = (lo | (hi << _u32(8))).cast(dt)
       else:
         val = _u32(0)
-        for i in range(4): val = val | (mem.index(idx + _const(dtypes.int, i), *gate, ptr=True).load().cast(dtypes.uint32) << _u32(i * 8))
+        for i in range(4): val = val | (_gload(mem.index(idx + _const(dtypes.int, i), ptr=True), mem.dtype).cast(dtypes.uint32) << _u32(i * 8))
     else:
       idx = (addr >> _const(addr.dtype, 2)).cast(dtypes.int)
-      val = mem.index(idx, *gate)
+      val = _gload(mem.index(idx, ptr=True), mem.dtype)
       if dt in (dtypes.uint64, dtypes.int64, dtypes.float64):
         idx2 = ((addr + _const(adt, 4)) >> _const(adt, 2)).cast(dtypes.int)
-        val = val.cast(dtypes.uint64) | (mem.index(idx2, *gate).cast(dtypes.uint64) << _u64(32))
+        val = val.cast(dtypes.uint64) | (_gload(mem.index(idx2, ptr=True), mem.dtype).cast(dtypes.uint64) << _u64(32))
       elif dt in (dtypes.uint8, dtypes.int8): val = (val >> ((addr & _const(adt, 3)).cast(dtypes.uint32) * _u32(8))) & _u32(0xFF)
       elif dt in (dtypes.uint16, dtypes.int16):
         val = (val >> (((addr >> _const(adt, 1)) & _const(adt, 1)).cast(dtypes.uint32) * _u32(16))) & _u32(0xFFFF)
@@ -862,7 +864,7 @@ class Parser:
         idx_native = (addr >> _const(adt, 2)).cast(dtypes.int64)
         idx_hi_native = ((addr + _const(adt, 4)) >> _const(adt, 2)).cast(dtypes.int64)
         safe_idx_hi = is_unaligned.where(idx_hi_native, idx_native)
-        hi = mem.index(safe_idx_hi, *gate)
+        hi = _gload(mem.index(safe_idx_hi, ptr=True), mem.dtype)
         combined = val.cast(dtypes.uint64) | (hi.cast(dtypes.uint64) << UOp.const(dtypes.uint64, 32))
         val = is_unaligned.where((combined >> (byte_off.cast(dtypes.uint64) * UOp.const(dtypes.uint64, 8))).cast(dtypes.uint32), val)
     return _cast_to(val, dt)

test/null/test_schedule.py

@@ -380,7 +380,7 @@ class TestSchedule(unittest.TestCase):
     r1 = (x - r0).sum(axis=0).div(2)
     out = r0 + r1
     linear, _ = check_schedule(out, 2)
-    reduceops = [x for si in linear.src for x in si.src[0].toposort() if x.op in {Ops.REDUCE_AXIS, Ops.REDUCE}]
+    reduceops = [x for si in linear.src for x in si.src[0].toposort() if x.op is Ops.REDUCE]
     assert len(reduceops) == 2
 
   def test_cache_reduce_multiple_children(self):
@@ -391,7 +391,7 @@ class TestSchedule(unittest.TestCase):
     out0 = r0 + y
     out1 = r1 + y
     linear, _ = check_schedule([out0, out1], 3)
-    reduceops = [x for si in linear.src for x in si.src[0].toposort() if x.op in {Ops.REDUCE_AXIS, Ops.REDUCE}]
+    reduceops = [x for si in linear.src for x in si.src[0].toposort() if x.op is Ops.REDUCE]
     self.assertEqual(len(reduceops), 2) # why is RANGEIFY different?
 
   def test_dedup_assign(self):

test/null/test_simplify_valid_idx.py

@@ -49,7 +49,7 @@ class TestValidIdxSimplification(unittest.TestCase):
   def check(self, load, sidx, svalid, extra=()):
     with Context(NOOPT=1, SPEC=0):
       load = full_rewrite_to_sink(UOp.sink(load, *extra)).src[0]
-    idx, valid = load.src[0].src[1], load.src[0].src[2]
+    idx, valid = load.src[0].src[1], load.src[1]
     check_uop_against_string(self, idx, sidx)
     check_uop_against_string(self, valid, svalid)
 
@@ -157,7 +157,7 @@ class TestValidIdxSimplification(unittest.TestCase):
     valid = (ridx2<1)&(ridx1<6)
     load = get_gated_load_uop(valid, idx)
     # prevent ridx1 and ridx2 from being shrunk
-    red = UOp(Ops.REDUCE, dtypes.float, (load, ridx1, ridx2), Ops.ADD)
+    red = load.reduce(ridx1, ridx2, arg=Ops.ADD)
     self.check(load,
       "(r0*1568)",
       "((r2<1)&(r1<6))",
@@ -225,9 +225,11 @@ class TestImageSimplification(unittest.TestCase):
     check_uop_against_string(self, idx0, sidx0)
     check_uop_against_string(self, idx1, sidx1)
     if svalid is not None:
-      check_uop_against_string(self, load.src[0].src[2], svalid)
+      check_uop_against_string(self, load.src[1], svalid)
     else:
-      self.assertEqual(len(load.src[0].src), 2, "svalid is None but load still has a valid")
+      # gate is at LOAD.src[1] when present; if simplified away, src[1] should not be bool
+      self.assertFalse(len(load.src) >= 2 and load.src[1].dtype.scalar() == dtypes.bool,
+                       "svalid is None but load still has a valid")
 
   def test_idx_gt_c(self):
     # (idx1 < c+1).ne(True) ? (..., idx1-1+c) : 0 can drop the valid
@@ -512,18 +514,34 @@ class TestUnfoldableImage(unittest.TestCase):
       self.assertEqual(res.src[0].src[0].dtype, dtypes.float.ptr(400))
 
 class TestDropTrueGate(unittest.TestCase):
-  def test_drop_true_gate_on_index(self):
-    # test that INDEX with a constant True gate gets simplified to drop the gate
+  def test_drop_true_gate_on_load(self):
+    # test that LOAD with a constant True gate gets simplified to drop the gate
     from tinygrad.codegen.late.devectorizer import load_store_indexing
     from tinygrad.uop.ops import graph_rewrite
     buf = UOp(Ops.PARAM, dtypes.int.ptr(), arg=0)
     idx = UOp.const(dtypes.weakint, 0)
     true_gate = UOp.const(dtypes.bool, True)
-    index_with_gate = UOp(Ops.INDEX, dtypes.int.ptr(), (buf, idx, true_gate))
+    bidx = UOp(Ops.INDEX, dtypes.int.ptr(), (buf, idx))
+    load = UOp(Ops.LOAD, dtypes.int, (bidx, true_gate))
     # apply the optimization
-    result = graph_rewrite(index_with_gate, load_store_indexing)
-    # the True gate should be dropped (INDEX should only have 2 sources)
-    self.assertEqual(len(result.src), 2, "True gate should be dropped from INDEX")
+    result = graph_rewrite(load, load_store_indexing)
+    # the True gate should be dropped (LOAD should only have 1 source)
+    self.assertEqual(len(result.src), 1, "True gate should be dropped from LOAD")
+
+  def test_drop_true_gate_on_store(self):
+    # test that STORE with a constant True gate gets simplified to drop the gate
+    from tinygrad.codegen.late.devectorizer import load_store_indexing
+    from tinygrad.uop.ops import graph_rewrite
+    buf = UOp(Ops.PARAM, dtypes.int.ptr(), arg=0)
+    idx = UOp.const(dtypes.weakint, 0)
+    val = UOp.const(dtypes.int, 42)
+    true_gate = UOp.const(dtypes.bool, True)
+    bidx = UOp(Ops.INDEX, dtypes.int.ptr(), (buf, idx))
+    store = UOp(Ops.STORE, dtypes.void, (bidx, val, true_gate))
+    # apply the optimization
+    result = graph_rewrite(store, load_store_indexing)
+    # the True gate should be dropped (STORE should only have 2 sources)
+    self.assertEqual(len(result.src), 2, "True gate should be dropped from STORE")
 
 class TestRangeShrink(unittest.TestCase):
   def get_ranges(self, sink):
@@ -569,7 +587,7 @@ class TestRangeShrink(unittest.TestCase):
     # range used in both a gated load AND directly in the reduce expression -> no shrink
     r = Range(0, 204)
     gated_load = get_gated_load_uop(r < UOp.const(dtypes.weakint, 4), r)
-    red = UOp(Ops.REDUCE, dtypes.float, (r.cast(dtypes.float) + gated_load, r), Ops.ADD)
+    red = (r.cast(dtypes.float) + gated_load).reduce(r, arg=Ops.ADD)
     ranges = self.get_ranges(red.sink())
     self.assertEqual(len(ranges), 1)
     self.assertEqual(ranges[0].src[0].arg, 204)

test/null/test_tensor_uop_mixin.py

@@ -137,6 +137,22 @@ class TestTensorUOpAllclose(unittest.TestCase):
     a, b = _t(4).float(), _t(4).float()
     self.assertIs(_strip_unique(a.allclose(b).uop), _strip_unique(a.uop.allclose(b.uop)))
 
+class TestTensorUOpBitcast(unittest.TestCase):
+  def test_bitcast_same_dtype(self): _check(self, _t(4).float(), lambda x: x.bitcast(dtypes.float32))
+
+class TestTensorUOpRand(unittest.TestCase):
+  def test_random_bits(self):
+    k = UOp.empty((2,), dtype=dtypes.uint32)
+    c = UOp.zeros(2, dtype=dtypes.uint32)
+    for num in (1, 4, 7, 1024):
+      self.assertIs(_strip_unique(Tensor.random_bits(Tensor(k), Tensor(c), num).uop),
+                    _strip_unique(UOp.random_bits(k, c, num)))
+  def test_bits_to_rand_float32(self):
+    bits_uop = UOp.empty((8,), dtype=dtypes.uint32)
+    for shape in ((8,), (2, 4), (5,)):
+      self.assertIs(_strip_unique(Tensor._bits_to_rand(Tensor(bits_uop), shape, dtypes.float32).uop),
+                    _strip_unique(UOp._bits_to_rand(bits_uop, shape, dtypes.float32)))
+
 class TestTensorUOpGather(unittest.TestCase):
   def _check(self, t, dim, idx):
     self.assertIs(_strip_unique(t.gather(dim, idx).uop), _strip_unique(t.uop.gather(dim, idx.uop)))

test/null/test_uop_graph.py

@@ -424,11 +424,12 @@ class TestUOpGraph(unittest.TestCase):
     ld = d0.index(ridx0.valid(ridx0<50))
     w = (ridx0<50).where(ld, 5)
     # prevent ridx0 from being shrunk
-    red = UOp(Ops.REDUCE, dtypes.long, (ridx0.cast(dtypes.long), ridx0), Ops.ADD)
+    red = ridx0.cast(dtypes.long).reduce(ridx0, arg=Ops.ADD)
     uops = to_uops_list([w, red])
     for u in uops:
       assert u.op is not Ops.WHERE
-      if u.op is Ops.LOAD and u.src[0].src[0].op is Ops.PARAM: assert u.src[1].arg==5
+      # alt is at src[2] in new gated LOAD shape (idx, gate, alt)
+      if u.op is Ops.LOAD and u.src[0].src[0].op is Ops.PARAM: assert u.src[2].arg==5
 
   def test_where_on_gated_load_folds_swapped_branches(self):
     ridx0 = UOp.range(100, 0)
@@ -438,7 +439,7 @@ class TestUOpGraph(unittest.TestCase):
     uops = to_uops_list([w])
     for u in uops:
       assert u.op is not Ops.WHERE
-      if u.op is Ops.LOAD: assert u.src[1].arg==5
+      if u.op is Ops.LOAD: assert u.src[2].arg==5
 
   def test_where_on_gated_load_with_cast(self):
     ridx0 = UOp.range(100, 0)
@@ -447,11 +448,11 @@ class TestUOpGraph(unittest.TestCase):
     ld = d0.index(gate_idx).cast(dtypes.float)
     w = (ridx0<50).where(ld, 5.0)
     # prevent ridx0 from being shrunk
-    red = UOp(Ops.REDUCE, dtypes.long, (ridx0.cast(dtypes.long), ridx0), Ops.ADD)
+    red = ridx0.cast(dtypes.long).reduce(ridx0, arg=Ops.ADD)
     uops = to_uops_list([w, red])
     for u in uops:
       assert u.op is not Ops.WHERE
-      if u.op is Ops.LOAD and u.src[0].src[0].op is Ops.PARAM: assert u.src[1].arg == 5
+      if u.op is Ops.LOAD and u.src[0].src[0].op is Ops.PARAM: assert u.src[2].arg == 5
 
   def test_where_on_casted_gated_load_extra_cond(self):
     ridx0 = UOp.range(100, 0)
@@ -459,7 +460,7 @@ class TestUOpGraph(unittest.TestCase):
     ld = d0.index(ridx0.valid(ridx0<50))
     w = ((ridx0<50) & (ridx0>30)).where(ld, UOp.const(dtypes.float, 0)).cast(dtypes.half)
     # prevent ridx0 from being shrunk
-    red = UOp(Ops.REDUCE, dtypes.long, (ridx0.cast(dtypes.long), ridx0), Ops.ADD)
+    red = ridx0.cast(dtypes.long).reduce(ridx0, arg=Ops.ADD)
     uops = to_uops_list([w, red])
     for u in uops:
       assert u.op is not Ops.WHERE
@@ -470,7 +471,7 @@ class TestUOpGraph(unittest.TestCase):
     ld = d0.index(ridx0.valid(ridx0<50))
     w = ((ridx0<50) & (ridx0>30)).where(UOp.const(dtypes.float, 0), ld).cast(dtypes.half)
     # prevent ridx0 from being shrunk
-    red = UOp(Ops.REDUCE, dtypes.long, (ridx0.cast(dtypes.long), ridx0), Ops.ADD)
+    red = ridx0.cast(dtypes.long).reduce(ridx0, arg=Ops.ADD)
     uops = to_uops_list([w, red])
     for u in uops:
       assert u.op is not Ops.WHERE
@@ -679,7 +680,7 @@ class TestExpander(unittest.TestCase):
   @unittest.skip("no longer supported")
   def test_reduce_known_axis(self):
     e1 = UOp(Ops.UNROLL, dtypes.int, tuple(UOp.const(dtypes.int, x) for x in range(4)), ((1,4),))
-    sink = UOp(Ops.REDUCE, dtypes.int, (3*e1,e1), Ops.ADD)
+    sink = (3*e1).reduce(e1, arg=Ops.ADD)
     sink = expander_rewrite(sink)
     assert sink.op is Ops.CONST
     self.assertEqual(sink.arg, 3*(0+1+2+3))
@@ -687,7 +688,7 @@ class TestExpander(unittest.TestCase):
   @unittest.skip("no longer supported")
   def test_reduce_const(self):
     e1 = UOp(Ops.UNROLL, dtypes.int, tuple(UOp.const(dtypes.int, x) for x in range(4)), ((1,4),))
-    sink = UOp(Ops.REDUCE, dtypes.int, (UOp.const(dtypes.int, 3), e1), Ops.ADD)
+    sink = UOp.const(dtypes.int, 3).reduce(e1, arg=Ops.ADD)
     sink = expander_rewrite(sink)
     assert sink.op is Ops.CONST
     self.assertEqual(sink.arg, 3*4)
@@ -728,7 +729,7 @@ class TestExpander(unittest.TestCase):
   def test_reduce_different_axis(self):
     e1 = UOp(Ops.UNROLL, dtypes.int, tuple(UOp.const(dtypes.int, x) for x in range(4)), ((1,4),))
     e2 = UOp(Ops.UNROLL, dtypes.int, tuple(UOp.const(dtypes.int, x) for x in range(4)), ((2,4),))
-    sink = UOp(Ops.REDUCE, dtypes.int, (e1,e2), Ops.ADD)
+    sink = e1.reduce(e2, arg=Ops.ADD)
     sink = expander_rewrite(sink)
     print(sink)
 

test/null/test_uop_vmin_vmax.py

@@ -297,7 +297,7 @@ class TestVminVmaxVConst(unittest.TestCase):
     # vmin and vmax for a vector constant of bool values
     d1 = UOp(Ops.PARAM, dtypes.int.ptr(), (), 1)
     idx = UOp.const(dtypes.int, 0)
-    val = UOp(Ops.LOAD, dtypes.int.vec(2), (d1.index(idx),))
+    val = UOp(Ops.LOAD, dtypes.int.vec(2), (d1.index(idx).cast(dtypes.int.vec(2).ptr()),))
     uop = (val // 32).gep(0)
     self.assertEqual(uop.vmin, -67108864)
     self.assertEqual(uop.vmax, 67108863)

test/null/test_viz.py

@@ -11,7 +11,7 @@ from tinygrad.helpers import cpu_profile, ProfilePointEvent, unwrap
 from tinygrad.device import Buffer
 
 from tinygrad.uop.ops import tracked_keys, tracked_ctxs, uop_fields, active_rewrites, active_group, _name_cnt, RewriteTrace
-from tinygrad.viz.serve import load_rewrites, get_full_rewrite, uop_to_json, VizData
+from tinygrad.viz.serve import load_rewrites, get_full_rewrite, uop_to_json, VizData, get_render
 from tinygrad.codegen import to_program_cache
 from tinygrad.codegen import to_program
 
@@ -321,7 +321,6 @@ class TestVizGC(unittest.TestCase):
 # VIZ integrates with other parts of tinygrad
 
 from tinygrad import Tensor, Device
-from tinygrad.engine.realize import get_runner
 
 class TestVizIntegration(unittest.TestCase):
   # codegen supports rendering of code blocks
@@ -707,7 +706,6 @@ class TestVizMemoryLayout(unittest.TestCase):
     self.assertEqual(len(programs), len(set(users)), n)
 
 from tinygrad.uop.ops import KernelInfo
-from tinygrad.viz.serve import amdgpu_cfg
 from tinygrad.renderer.amd.dsl import s
 from tinygrad.runtime.autogen.amd.rdna3.ins import (s_add_u32, s_branch, s_cbranch_execz, s_cbranch_scc0, s_cbranch_scc1, s_cmp_eq_i32,
                                                     s_cmp_eq_u64, s_code_end, s_endpgm, s_mov_b32, s_nop)
@@ -723,13 +721,16 @@ class TestCfg(unittest.TestCase):
       gidx = UOp.special(1, "gidx0")
       sink = UOp.sink(out.base, lidx, gidx, arg=KernelInfo(name=name))
       return UOp(Ops.PROGRAM, src=(sink, UOp(Ops.DEVICE, arg="NULL"), UOp(Ops.LINEAR, src=tuple([UOp(Ops.INS, arg=x) for x in insts]))))
-    with Context(DEV=f"NULL::{self.arch}"):
-      out = Tensor.custom_kernel(Tensor.empty(1), fxn=fxn)[0]
-      runner = get_runner(out.device, out.schedule_linear().src[-1].src[0])
-      return amdgpu_cfg(runner.prg.src[4].arg, self.arch)
+    with save_viz() as viz:
+      with Context(DEV=f"NULL::{self.arch}"):
+        out = Tensor.custom_kernel(Tensor.empty(1), fxn=fxn)[0]
+        _ = to_program(out.schedule_linear().src[-1].src[0], Device[out.device].renderer)
+    codegen_rewrites = next(s for s in viz.list_items() if s["name"] == name)
+    disasm = next(s for s in codegen_rewrites["steps"] if s["name"] == "View Disassembly")
+    return get_render(viz.data, disasm["query"])
 
   def test_simple(self):
-    k = Kernel(arch=self.arch)
+    k = Kernel()
     k.label("entry")
     k.emit(s_branch(), target="bb1")
     k.label("bb1")
@@ -739,7 +740,7 @@ class TestCfg(unittest.TestCase):
     self.assertEqual(len(cfg["blocks"]), 2)
 
   def test_diamond(self):
-    k = Kernel(arch=self.arch)
+    k = Kernel()
     k.label("entry")
     k.emit(s_mov_b32(s[0], 0))
     k.emit(s_mov_b32(s[1], 0))
@@ -773,7 +774,7 @@ class TestCfg(unittest.TestCase):
       assert st.startswith("s_code_end") and st.endswith("x)"), st
 
   def test_loop(self):
-    k = Kernel(arch=self.arch)
+    k = Kernel()
     k.label("entry")
     k.emit(s_mov_b32(s[1], 4))
     k.label("loop")
@@ -785,7 +786,7 @@ class TestCfg(unittest.TestCase):
     self.get_cfg("simple_loop", k)
 
   def test_loop_branch(self):
-    k = Kernel(arch=self.arch)
+    k = Kernel()
     k.label("entry")
     k.emit(s_mov_b32(s[1], 4))
     k.label("loop")
@@ -803,7 +804,7 @@ class TestCfg(unittest.TestCase):
     self.get_cfg("loop_if", k)
 
   def test_loop_break(self):
-    k = Kernel(arch=self.arch)
+    k = Kernel()
     k.label("entry")
     k.emit(s_mov_b32(s[1], 8))
     k.label("loop")
@@ -818,7 +819,7 @@ class TestCfg(unittest.TestCase):
     self.get_cfg("loop_break", k)
 
   def test_switch(self):
-    k = Kernel(arch=self.arch)
+    k = Kernel()
     k.label("entry")
     k.emit(s_cmp_eq_i32(s[0], 0))
     k.emit(s_cbranch_scc1(), target="case0")
@@ -840,7 +841,7 @@ class TestCfg(unittest.TestCase):
     self.get_cfg("switch_case", k)
 
   def test_ping_pong(self):
-    k = Kernel(arch=self.arch)
+    k = Kernel()
     k.label("entry")
     k.emit(s_cmp_eq_i32(s[0], 0))
     k.emit(s_cbranch_scc1(), target="ping")
@@ -859,7 +860,7 @@ class TestCfg(unittest.TestCase):
 
   def test_colored_blocks(self):
     N = 10
-    k = Kernel(arch=self.arch)
+    k = Kernel()
     k.label("entry")
     k.emit(s_branch(), target="init0")
     for i in range(N):
@@ -879,7 +880,7 @@ class TestCfg(unittest.TestCase):
     self.get_cfg("test_colored_blocks", k)
 
   def test_jump_back_to_end(self):
-    k = Kernel(arch=self.arch)
+    k = Kernel()
     k.label("entry")
     k.emit(s_mov_b32(s[1], 2))
     k.emit(s_cbranch_execz(), target="loop")

test/opt/test_tensor_cores.py

@@ -3,12 +3,12 @@ import unittest
 
 from tinygrad import Device, Tensor, dtypes
 from tinygrad.tensor import _to_np_dtype
-from tinygrad.uop.ops import Ops
+from tinygrad.uop.ops import Ops, UOp, buffers
 from tinygrad.dtype import DType
-from tinygrad.device import is_dtype_supported
+from tinygrad.device import Buffer, is_dtype_supported
 from tinygrad.helpers import DEV, Context
 from test.helpers import slow, replace_opts
-from tinygrad.engine.realize import CompiledRunner
+from tinygrad.engine.realize import run_linear
 from tinygrad.codegen import to_program
 from tinygrad.codegen.opt import Opt, OptOps, KernelOptError
 from tinygrad.codegen.opt.tc import amd_cdna_1616128
@@ -20,6 +20,11 @@ from test.backend.test_linearizer import helper_realized_ast, helper_linearizer_
 
 AMX = "AMX" in DEV.arch
 
+def run_program(prg:UOp, bufs:list[Buffer]):
+  buf_uops = [UOp.new_buffer(b.device, b.size, b.dtype) for b in bufs]
+  for u,b in zip(buf_uops, bufs): buffers[u] = b
+  run_linear(UOp(Ops.LINEAR, src=(prg.call(*buf_uops),)))
+
 def helper_tc_ensure_uops_and_opts_count(N: int, M:int, K:int, dtype_in:DType, dtype_out:DType, axis:int=0, tc_select:int=-1, tc_opt:int=0,
                                          ensure_triggered:bool=True):
   a, b = Tensor.rand(M, K, dtype=dtype_in), Tensor.rand(K, N, dtype=dtype_in)
@@ -47,11 +52,11 @@ def helper_tc_allclose(N:int, M:int, K:int, dtype_in:DType, dtype_out:DType, axi
   if dtype_in == dtypes.bfloat16: r = r.float()
   realized_ast, bufs = helper_realized_ast(r)
   opts = [Opt(op=OptOps.TC, axis=axis, arg=(tc_select, tc_opt, use_tensor_cores))]
-  pu = to_program(replace_opts(realized_ast, opts), Device[Device.DEFAULT].renderer)
+  ast = replace_opts(realized_ast, opts)
+  pu = to_program(ast, Device[Device.DEFAULT].renderer)
   if use_tensor_cores == 1: assert len([uop for uop in pu.src[2].src if uop.op is Ops.WMMA]) > 0, "wmma not triggered"
   assert len([x for x in pu.src[0].arg.applied_opts if x.op is OptOps.TC]) == 1, "tensor core opt not included"
-  prg = CompiledRunner(pu, Device.DEFAULT)
-  prg.exec(bufs)
+  run_program(ast, bufs)
   if dtype_in == dtypes.half: tc_atol, tc_rtol = 1e-2, 1e-3
   elif dtype_in == dtypes.bfloat16: tc_atol, tc_rtol = (1e-1, 2e-2) if dtype_out == dtypes.bfloat16 else (1e-2, 1e-2)
   else: tc_atol, tc_rtol = 5e-3, 1e-4
@@ -145,15 +150,15 @@ class TestTensorCores(unittest.TestCase):
         c = a.conv2d(b, padding=1, dtype=tc.dtype_out)
         realized_ast, real_bufs = helper_realized_ast(c)
 
-        program = to_program(replace_opts(realized_ast, [Opt(OptOps.TC, axis, (-1, 2, 1))]), Device[Device.DEFAULT].renderer)
+        ast = replace_opts(realized_ast, [Opt(OptOps.TC, axis, (-1, 2, 1))])
+        program = to_program(ast, Device[Device.DEFAULT].renderer)
         assert len([uop for uop in tuple(program.src[2].src) if uop.op is Ops.WMMA]) > 0, "tensor core not triggered"
         assert len([x for x in program.src[0].arg.applied_opts if x.op is OptOps.TC]) == 1, "tensor core opt not included"
 
-        prg = CompiledRunner(program, Device.DEFAULT)
         # TODO: support this even if numpy doesn't
         if _to_np_dtype(real_bufs[0].dtype) is None: continue
         real_bufs[0].copyin(np.zeros((real_bufs[0].size, ), dtype=_to_np_dtype(real_bufs[0].dtype)).data) # Zero to check that all values are filled
-        prg.exec(real_bufs)
+        run_program(ast, real_bufs)
         result = np.frombuffer(real_bufs[0].as_memoryview(), _to_np_dtype(real_bufs[0].dtype))
 
         # ensure the results for each choice of axis matches

test/unit/test_assign.py

@@ -903,8 +903,8 @@ class TestAssignToUnrealizedView(unittest.TestCase):
 
   def test_reduce(self):
     a = Tensor([[1,2],[3,4]]).contiguous().realize()
-    r = a.sum(axis=0)  # unrealized REDUCE_AXIS
-    self.assertIs(r.uop.base.op, Ops.REDUCE_AXIS)
+    r = a.sum(axis=0)  # unrealized REDUCE
+    self.assertIs(r.uop.base.op, Ops.REDUCE)
     r[:1].assign(Tensor([99]).realize())
     try:
       self.assertEqual(r.tolist(), [99,6])

tinygrad/codegen/__init__.py

@@ -3,7 +3,8 @@ from dataclasses import replace
 import itertools
 from tinygrad.helpers import DISABLE_FAST_IDIV, DEVECTORIZE, TRANSCENDENTAL, SPEC, DEBUG, VIZ, IMAGE, NOOPT, EMULATED_DTYPES, NOLOCALS, USE_TC
 from tinygrad.helpers import ALLOW_TF32, TracingKey, Context, Target, panic
-from tinygrad.uop.ops import PatternMatcher, graph_rewrite, UOp, pm_lower_index_dtype, Ops, UPat, track_rewrites, KernelInfo, ProgramInfo, pyrender
+from tinygrad.uop.ops import PatternMatcher, graph_rewrite, UOp, pm_lower_index_dtype, Ops, UPat, track_rewrites, KernelInfo, ProgramInfo
+from tinygrad.uop.render import pyrender
 from tinygrad.uop.spec import type_verify, program_spec, kernel_spec
 from tinygrad.renderer import Renderer, Estimates
 from tinygrad.dtype import dtypes
@@ -107,9 +108,9 @@ def full_rewrite_to_sink(sink:UOp, ren:Renderer|None=None, optimize:bool=True, b
 pm_linearize_cleanups = PatternMatcher([
   # if statements are not allowed in the graph
   (UPat((Ops.IF, Ops.ENDIF)), lambda: panic(RuntimeError, "if not allowed in graph")),
-  # gated INDEX becomes IF-STORE-ENDIF. this is the only use of IF-ENDIF
-  (UPat(Ops.STORE, name="u", src=(UPat(Ops.INDEX, src=(UPat(), UPat(), UPat(name="gate", dtype=dtypes.bool))).or_casted(), UPat())),
-   lambda u, gate: (u, [mif:=UOp(Ops.IF, src=(gate, u.src[0])), u, UOp(Ops.ENDIF, src=(mif,))]))
+  # gated STORE becomes IF-STORE-ENDIF. this is the only use of IF-ENDIF
+  (UPat(Ops.STORE, name="u", src=(UPat(Ops.INDEX).or_casted(), UPat(), UPat(name="gate", dtype=dtypes.bool))),
+   lambda u, gate: (u.replace(src=u.src[:2]), [mif:=UOp(Ops.IF, src=(gate, u.src[0])), u.replace(src=u.src[:2]), UOp(Ops.ENDIF, src=(mif,))]))
 ])
 
 # requires lst be toposorted. like graph rewrite, but for lines
@@ -180,8 +181,6 @@ def do_to_program(ast:UOp, renderer:Renderer) -> UOp:
 
 to_program_cache: dict[tuple, UOp] = {}
 def to_program(ast:UOp, renderer:Renderer) -> UOp:
-  if ast.op is Ops.PROGRAM and len(ast.src) >= 5 and ast.src[4].op is Ops.BINARY:
-    return ast if isinstance(ast.arg, ProgramInfo) else ast.replace(arg=ProgramInfo.from_sink(ast.src[0]))
   config = (NOOPT, DEVECTORIZE, EMULATED_DTYPES, NOLOCALS, USE_TC, IMAGE, DISABLE_FAST_IDIV, TRANSCENDENTAL, ALLOW_TF32)
   key = (ast.key, type(renderer), renderer.target, *[x.value for x in config])
   if (prg:=to_program_cache.get(key)) is None: to_program_cache[key] = prg = do_to_program(ast, renderer)

tinygrad/codegen/late/devectorizer.py

@@ -53,10 +53,15 @@ def simplify_valid_load(buf:UOp, start_idx:UOp, valid:UOp) -> UOp|None:
 load_store_indexing = PatternMatcher([
   # image load valid idx simplification
   (UPat(Ops.INDEX, src=(UPat.var("buf"), invalid_gate)), lambda buf,x,i,cond: simplify_valid_load(buf, x, cond)),
-  # simplify away long after index has been lowered
-  (UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var("x", dtypes.long), UPat.var("c", dtypes.bool))), lambda buf,x,c: simplify_valid_load(buf, x, c)),
-  # drop true gate
-  (UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var("x"), UPat.const(dtypes.bool, True)),), lambda buf,x: buf.index(x, ptr=True)),
+  # drop true gate from gated LOAD with alt: also drop the now-unused alt
+  (UPat(Ops.LOAD, src=(UPat.var("idx"), UPat.const(dtypes.bool, True), UPat()), allow_any_len=True, name="ld"),
+   lambda ld,idx: ld.replace(src=(idx,)+ld.src[3:])),
+  # drop true gate from gated LOAD without alt
+  (UPat(Ops.LOAD, src=(UPat.var("idx"), UPat.const(dtypes.bool, True)), allow_any_len=True, name="ld"),
+   lambda ld,idx: ld.replace(src=(idx,)+ld.src[2:])),
+  # drop true gate from STORE
+  (UPat(Ops.STORE, src=(UPat.var("idx"), UPat.var("val"), UPat.const(dtypes.bool, True))),
+   lambda idx,val: idx.store(val)),
 ])
 
 # ***** load/store grouping *****
@@ -76,7 +81,7 @@ def expand_index(buf:UOp, vec:UOp):
     buf = buf.replace(dtype=(dtypes.imageh if dt.itemsize == 2 else dtypes.imagef)((h, w, 4)))
   if getenv("UNSAFE_DISABLE_MASK", 0): vec = vec.get_idx()
   # generate the individual indexes
-  return UOp(Ops.STACK, buf.dtype, tuple(buf.index(vec.gep(i), ptr=True) for i in range(vec.shape[0])))
+  return UOp(Ops.STACK, buf.dtype, tuple(buf.index(vec.gep(i), ptr=True) for i in range(vec.dtype.count)))
 
 def fold_expanded_index(midx:UOp):
   buf = midx.src[0].src[0]
@@ -104,7 +109,7 @@ def fold_expanded_index(midx:UOp):
     for grp in grouped_offsets:
       # get the index offset for this element. using [0] is okay, because they are the same
       lidx = midx.src[offsets[grp[0]][0]]
-      #if len(grp) > 1: lidx = lidx.cast(buf.ptrdtype.base.vec(len(grp)).ptr(size=buf.ptrdtype.size, addrspace=buf.ptrdtype.addrspace))
+      if len(grp) > 1: lidx = lidx.cast(buf.ptrdtype.base.vec(len(grp)).ptr(size=buf.ptrdtype.size, addrspace=buf.ptrdtype.addrspace))
       # set the idxs of the output
       for i,g in enumerate(grp):
         for oo in offsets[g]: idxs[oo] = global_offset+i
@@ -116,22 +121,22 @@ def fold_expanded_index(midx:UOp):
   post_cat = UOp(Ops.PTRCAT, buf.ptrdtype.base.ptr(size=buf.ptrdtype.size, addrspace=buf.ptrdtype.addrspace).vec(global_offset), tuple(ret))
   return post_cat.gep(tuple(cast(list[int], idxs)))
 
-def cat_after_store(cat:UOp, data:UOp, sto:UOp):
+def cat_after_store(cat:UOp, data:UOp):
   # TODO: this is written in many places
   offset = 0
   ret: list[UOp] = []
   for s in cat.src:
-    ret.append(s.store(data.gep(tuple(range(offset, offset+s.dtype.count))), *sto.src[2:]))
+    ret.append(s.store(data.gep(tuple(range(offset, offset+s.dtype.count)))))
     offset += s.dtype.count
   return UOp.group(*ret)
 
-def gep_on_store(gep:UOp, st:UOp, sto:UOp):
+def gep_on_store(gep:UOp, st:UOp):
   # NOTE: we need to invert the gep here, but it may be an expanding gep
   # fake argsort. TODO: handle duplicates
   a = {}
   for i,x in enumerate(gep.arg): a[x] = i
   new_arg = tuple(x[1] for x in sorted(a.items()))
-  return gep.src[0].store(st.gep(new_arg), *sto.src[2:])
+  return gep.src[0].store(st.gep(new_arg))
 
 load_store_folding = PatternMatcher([
   (UPat(Ops.INDEX, src=(UPat(Ops.STACK, src=UPat(GroupOp.Defines).or_after(name="buf")), UPat.var("vec"))), expand_index),
@@ -140,12 +145,12 @@ load_store_folding = PatternMatcher([
   (UPat(Ops.LOAD, src=(UPat(Ops.GEP, name="gep"),), name="ld", allow_any_len=True),
    lambda gep, ld: ld.replace(dtype=ld.dtype.scalar().vec(gep.dtype.count), src=(gep.src[0],)+ld.src[1:]).gep(gep.arg)),
   # GEP on data of STORE
-  (UPat(Ops.STORE, src=(UPat(Ops.GEP, name="gep"), UPat.var("st")), name="sto"), gep_on_store),
+  (UPat(Ops.STORE, src=(UPat(Ops.GEP, name="gep"), UPat.var("st"))), gep_on_store),
   # put PTRCAT after LOAD
   (UPat(Ops.LOAD, src=(UPat(Ops.PTRCAT, name="cat"),), name="ld", allow_any_len=True),
-   lambda cat,ld: UOp(Ops.VCAT, cat.dtype.base, tuple(ld.replace(dtype=x.dtype.base, src=(x,)+ld.src[1:]) for x in cat.src))),
+   lambda cat,ld: UOp(Ops.VCAT, cat.dtype.base.vec(cat.dtype.vcount), tuple(ld.replace(dtype=x.dtype.base, src=(x,)+ld.src[1:]) for x in cat.src))),
   # put PTRCAT after STORE
-  (UPat(Ops.STORE, src=(UPat(Ops.PTRCAT, name="cat"), UPat(name="data")), name="sto"), cat_after_store),
+  (UPat(Ops.STORE, src=(UPat(Ops.PTRCAT, name="cat"), UPat(name="data"))), cat_after_store),
 ])
 
 # *** correct load/store ***
@@ -187,7 +192,7 @@ def split_load_store(ctx:Renderer|None, ls:UOp, idx:UOp):
       if global_offset+fold_length > sz: continue
       lidx = buf.index((offset + global_offset).valid(mask), ptr=True)
       if fold_length > 1: lidx = lidx.cast(buf.ptrdtype.base.vec(fold_length).ptr(size=buf.ptrdtype.size, addrspace=buf.ptrdtype.addrspace))
-      if ls.op is Ops.STORE: ret.append(ls.replace(src=(lidx,ls.src[1].gep(tuple(range(global_offset, global_offset+fold_length))))+ls.src[2:]))
+      if ls.op is Ops.STORE: ret.append(ls.replace(src=(lidx,ls.src[1].gep(tuple(range(global_offset, global_offset+fold_length))))))
       else: ret.append(ls.replace(src=(lidx,)+ls.src[1:], dtype=ls.dtype.scalar().vec(fold_length)))
       global_offset += fold_length
       break
@@ -197,8 +202,9 @@ def split_load_store(ctx:Renderer|None, ls:UOp, idx:UOp):
   return UOp(Ops.VCAT, ls.dtype, tuple(ret)) if ls.op is Ops.LOAD else UOp.group(*ret)
 
 def get_image_idx(idx:UOp, width:int):
-  oidx = UOp(Ops.STACK, dtypes.weakint.vec(2), (((x:=idx.src[1].get_idx()) // 4) % width, (x // (4*width))))
-  return idx.replace(src=(idx.src[0], oidx.valid(idx.src[1].get_valid())))
+  x, valid = idx.src[1].get_idx(), idx.src[1].get_valid()
+  idx_x, idx_y = (x // 4) % width, x // (4*width)
+  return idx.replace(src=(idx.src[0], UOp.vectorize(idx_x, idx_y).valid(valid)))
 
 def image_fixup(ls:UOp):
   # normal image load or store, with the CAST from expand_index
@@ -280,18 +286,18 @@ pm_render = PatternMatcher([
   (UPat(Ops.GEP, name='gep'), lambda gep: UOp(Ops.STACK, gep.dtype, tuple(gep.src[0].gep(x) for x in gep.arg)) if len(gep.arg) > 1 else None),
   (UPat(Ops.GEP, name='gep'), lambda gep: gep.src[0] if gep.src[0].dtype.vcount == 1 and gep.arg == (0,) else None),
   (UPat(Ops.STACK, src=(UPat(name='x'),)), lambda x: x),
-  # give any loads that are masked an alt value
-  (UPat(Ops.LOAD, src=(UPat(Ops.INDEX, src=(UPat(), UPat(), UPat())).or_casted(),), allow_any_len=True, name="x"),
-    lambda x: x.replace(src=(x.src[0], x.const_like(0))+x.src[1:])
-      if len(x.src) == 1 or x.src[1].op in (Ops.CUSTOM, Ops.STORE, Ops.BARRIER) else None),
+  # give any gated loads (gate at src[1]) an alt value at src[2]
+  (UPat(Ops.LOAD, src=(UPat(Ops.INDEX).or_casted(), UPat(dtype=dtypes.bool)), allow_any_len=True, name="x"),
+    lambda x: x.replace(src=(x.src[0], x.src[1], x.const_like(0))+x.src[2:])
+      if len(x.src) == 2 or x.src[2].op in (Ops.CUSTOM, Ops.STORE, Ops.BARRIER) else None),
   # Where after gated load becomes alt value
   # NOTE: if a is CAST and a.src[0].dtype == l.dtype, use a.src[0] to avoid roundtrip cast (e.g. uint->float->uint)
-  (UPat.var("c").where(UPat(Ops.LOAD, src=(UPat().index(UPat(), UPat.var("c")).or_casted(),), allow_any_len=True, name="l").or_casted(),
-    UPat.var("a")), lambda c,l,a: l.replace(src=(l.src[0], a.src[0] if a.op is Ops.CAST and a.src[0].dtype == l.dtype else a.cast(l.dtype))+
-                                            l.src[2:]).cast(a.dtype)),
-  (UPat.var("c").where(UPat.var("a"), UPat(Ops.LOAD, src=(UPat().index(UPat(), UPat.var("c", dtype=dtypes.bool).logical_not()).or_casted(),),
-    allow_any_len=True, name="l").or_casted()), lambda c,l,a: l.replace(src=(l.src[0], a.src[0] if a.op is Ops.CAST and a.src[0].dtype == l.dtype
-                                                                             else a.cast(l.dtype))+l.src[2:]).cast(a.dtype)),
+  (UPat.var("c").where(UPat(Ops.LOAD, src=(UPat(Ops.INDEX).or_casted(), UPat.var("c")), allow_any_len=True, name="l").or_casted(),
+    UPat.var("a")), lambda c,l,a: l.replace(src=(l.src[0], l.src[1], a.src[0] if a.op is Ops.CAST and a.src[0].dtype == l.dtype else a.cast(l.dtype))+
+                                            l.src[3:]).cast(a.dtype)),
+  (UPat.var("c").where(UPat.var("a"), UPat(Ops.LOAD, src=(UPat(Ops.INDEX).or_casted(), UPat.var("c", dtype=dtypes.bool).logical_not()),
+    allow_any_len=True, name="l").or_casted()), lambda c,l,a: l.replace(src=(l.src[0], l.src[1], a.src[0] if a.op is Ops.CAST and a.src[0].dtype == l.dtype
+                                                                             else a.cast(l.dtype))+l.src[3:]).cast(a.dtype)),
 ])
 
 # *** Ops.REDUCE -> Ops.DEFINE_ACC ***
@@ -317,12 +323,12 @@ def reduce_to_acc(ctx:ReduceContext, red:UOp):
     topo = inp.toposort()
     ended_ranges = flatten([x.ended_ranges for x in topo if x.op is Ops.END])
     input_ranges = tuple([x for x in topo if x.op is Ops.RANGE and x not in reduce_range and x not in ended_ranges])
-    identity = red.const(red.dtype, identity_element(red.arg, red.dtype.scalar()))
+    identity = red.const(red.dtype, identity_element(red.arg[0], red.dtype.scalar()))
     acc = UOp.placeholder((1,), red.dtype, ctx.acc_num, AddrSpace.REG)
     acc_init = acc.after(*input_ranges).index(UOp.const(dtypes.weakint, 0)).store(identity)
     lst = [acc.after(acc_init, *reduce_range).index(UOp.const(dtypes.weakint, 0))] + lst  # put acc as the first element
     ctx.acc_num += 1
-  ret = functools.reduce(lambda x,y: x.alu(red.arg, y), lst)
+  ret = functools.reduce(lambda x,y: x.alu(red.arg[0], y), lst)
   if len(reduce_range) == 0: return ret
   end = acc.index(UOp.const(dtypes.weakint, 0)).store(ret).end(*reduce_range).rtag("mergeable")
   return acc.after(end).index(UOp.const(dtypes.weakint, 0))
@@ -358,10 +364,14 @@ pm_reduce = PatternMatcher([
 
 # add loads
 
+def add_load(idx:UOp):
+  if isinstance(idx.dtype, PtrDType): return None
+  assert isinstance(idx.src[0].dtype, PtrDType), f"param is not PtrDType {idx.src[0].dtype}"
+  return idx.replace(dtype=idx.src[0].dtype).load(dtype=idx.dtype.base)
+
 pm_add_loads = PatternMatcher([
   # add loads to non ptr index
-  (UPat(Ops.INDEX, name="idx"), lambda idx: None if isinstance(idx.dtype, PtrDType) else
-    idx.replace(dtype=idx.src[0].dtype).load(dtype=idx.dtype.base)),
+  (UPat(Ops.INDEX, name="idx"), add_load),
   # remove loads from stores
   (UPat(Ops.STORE, src=(UPat(Ops.LOAD), UPat(name="val")), name="s"), lambda s,val: s.replace(src=(s.src[0].src[0], val))),
 ])

tinygrad/codegen/late/expander.py

@@ -45,7 +45,7 @@ def do_expand(root:UOp):
     else:
       # non-UNROLL input
       if root.op in range_start and i >= range_start[root.op]:
-        # for any range args of STORE/REDUCE, pass them through
+        # for any range args of REDUCE/WMMA/END/etc., pass them through
         new_srcs.append(src)
       elif root.op is Ops.INDEX and i >= 1 and not isinstance(root.dtype, PtrDType):
         new_srcs.append(src)
@@ -71,8 +71,7 @@ def do_expand(root:UOp):
     assert root.dtype.count == 1
     # is this right?
     new_arg = tuple(range(root.arg[0], new_srcs[0].dtype.count, new_srcs[0].dtype.count // expand_sz))
-  #nsrc = UOp(root.op, root.dtype.scalar().vec(root.dtype.count*expand_sz), tuple(new_srcs), new_arg)
-  nsrc = UOp(root.op, root.dtype, tuple(new_srcs), new_arg)
+  nsrc = UOp(root.op, root.dtype.scalar().vec(root.dtype.count*expand_sz), tuple(new_srcs), new_arg)
   return UOp(Ops.UNROLL, root.dtype, (nsrc,), expand_args)
 
 def do_contract(con:UOp):
@@ -148,7 +147,7 @@ def fix_group_for_reduce(x:UOp):
 pm_pre_expander = PatternMatcher([
   # rewrite UPCAST/UNROLL range to something to be expanded
   (UPat(Ops.RANGE, name="r"),
-   lambda r: UOp(Ops.UNROLL, r.dtype, (UOp.const(r.dtype, tuple(range(s:=r.vmax+1))),), ((r.arg[0],s),)) \
+   lambda r: UOp(Ops.UNROLL, r.dtype, (UOp.const(r.dtype.vec(s:=r.vmax+1), tuple(range(s))),), ((r.arg[0],s),)) \
     if r.arg[1] in {AxisType.UNROLL, AxisType.UPCAST} else None),
   # fix REDUCEs with UNROLLs
   (UPat(Ops.REDUCE, name="x"), fix_reduce_unroll),

tinygrad/codegen/opt/postrange.py

@@ -219,7 +219,7 @@ class Scheduler:
   def _apply_tc_opt(self, use_tensor_cores:int, axis:int, tc_select:int, opt_level:int) -> None|list[UOp]:
     if not (reduceops := self.reduceops): raise KernelOptError("no reduce ops for TensorCore")
     reduceop = reduceops[0]
-    if use_tensor_cores and reduceop.arg is Ops.ADD:
+    if use_tensor_cores and reduceop.arg[0] is Ops.ADD:
       mul = reduceop.src[0] if reduceop.src[0].op is not Ops.CAST else reduceop.src[0].src[0]
       if mul.op is not Ops.MUL: return None
       in0, in1 = mul.src
@@ -305,7 +305,7 @@ class Scheduler:
 
             # preserve extra reduces
             reduce_ranges = [x for x in UOp.sink(*reduceop.src[1:]).toposort() if x.op is Ops.RANGE and x.arg[0] not in tc_reduce_axes]
-            if len(reduce_ranges): tc_uop = UOp(Ops.REDUCE, tc_uop.dtype, (tc_uop,)+tuple(reduce_ranges), Ops.ADD)
+            if len(reduce_ranges): tc_uop = UOp(Ops.REDUCE, tc_uop.dtype, (tc_uop,)+tuple(reduce_ranges), (Ops.ADD, ()))
             self.ast = self.ast.substitute({reduceop: tc_uop})
           self.tensor_core = tc
           return axes
@@ -317,7 +317,7 @@ class Scheduler:
   @property
   def reduceop(self) -> UOp|None:
     if not (red := self.reduceops): return None
-    return UOp(Ops.REDUCE_AXIS, red[0].dtype, red[0].src, (red[0].arg, ()))
+    return UOp(Ops.REDUCE, red[0].dtype, red[0].src, red[0].arg)
   @property
   def bufs(self) -> list[UOp]: return [x for x in self.ast.toposort() if x.op is Ops.INDEX][::-1]
   @property

tinygrad/codegen/opt/search.py

@@ -1,12 +1,13 @@
-import functools, math, time, multiprocessing, traceback, signal, atexit
+import math, time, multiprocessing, traceback, signal, atexit
 from dataclasses import replace
-from tinygrad.uop.ops import sym_infer, AxisType, pyrender, UOp, Ops
-from tinygrad.device import Device, Buffer, Compiler
+from tinygrad.uop.ops import sym_infer, AxisType, UOp
+from tinygrad.uop.render import pyrender
+from tinygrad.device import Device, Buffer
 from tinygrad.helpers import prod, flatten, DEBUG, CACHELEVEL, diskcache_get, diskcache_put, getenv, Context, colored, time_to_str, unwrap
 from tinygrad.helpers import IGNORE_BEAM_CACHE
 from tinygrad.codegen.opt import Opt, OptOps, KernelOptError
 from tinygrad.tensor import Tensor
-from tinygrad.engine.realize import CompiledRunner
+from tinygrad.engine.realize import get_runtime
 from tinygrad.codegen import to_program
 from tinygrad.codegen.opt.postrange import Scheduler
 
@@ -34,25 +35,24 @@ def get_test_global_size(global_size, max_global_size, var_vals):
         break
   return test_global_size, input_size / prod(test_global_size)
 
-def _time_program(prg:UOp, lib:bytes, var_vals:dict[str, int], rawbufs:list[Buffer], early_stop:float|None=None,
+def _time_program(prg:UOp, var_vals:dict[str, int], rawbufs:list[Buffer], early_stop:float|None=None,
                   allow_test_size:int=True, max_global_size:int|None=65536, clear_l2=False, cnt=3, name="test", dev_timeout=False) -> list[float]:
   timeout = int(early_stop * 1e3) if dev_timeout and early_stop is not None and early_stop < math.inf else None
   factor = 1
-  info = prg.arg
   if allow_test_size and max_global_size is not None:
-    global_size, factor = get_test_global_size(info.global_size, max_global_size, var_vals)
-    prg = prg.replace(arg=replace(info, global_size=tuple(global_size)))
-  if len(prg.src) <= 4 or prg.src[4].op is not Ops.BINARY: prg = prg.replace(src=prg.src + (UOp(Ops.BINARY, arg=lib),))
-  try: car = CompiledRunner(prg, prg.src[1].arg)
+    global_size, factor = get_test_global_size(prg.arg.global_size, max_global_size, var_vals)
+    prg = prg.replace(arg=replace(prg.arg, global_size=tuple(global_size)))
+  try: rt = get_runtime(prg.src[1].arg, prg)
   except AssertionError: return [math.inf] * cnt
+  global_size, local_size = prg.arg.launch_dims(var_vals)
+  bufs = [rawbufs[i]._buf for i in prg.arg.globals]
   tms = []
-  input_bufs = [rawbufs[i] for i in car.p.globals]
   for _ in range(cnt):
     if clear_l2:
       if hasattr(dev:=Device[prg.src[1].arg], 'invalidate_caches'): dev.invalidate_caches()
       else:
         with Context(DEBUG=0, BEAM=0, CAPTURING=0, TRACK_MATCH_STATS=0): Tensor.ones(1024,1024).contiguous().realize(do_update_stats=False)
-    tms.append(unwrap(car(input_bufs, var_vals, wait=True, timeout=timeout))*factor)
+    tms.append(unwrap(rt(*bufs, global_size=global_size, local_size=local_size, vals=prg.arg.vals(var_vals), wait=True, timeout=timeout))*factor)
     if early_stop is not None and early_stop < min(tms): break
   return tms
 
@@ -61,22 +61,21 @@ def timeout_handler(signum, frame):
   if DEBUG >= 2: print("*** BEAM COMPILE TIMEOUT")
   raise TimeoutException()
 
-def _try_compile(x:tuple[int,Scheduler], compiler:Compiler) -> tuple[int, tuple[UOp, bytes, float]|None]:
+def _try_compile(x:tuple[int,Scheduler]) -> tuple[int, tuple[UOp, float]|None]:
   if hasattr(signal, "alarm"):
     signal.signal(getattr(signal, 'SIGALRM'), timeout_handler)
     # set timeout
     signal.alarm(getenv("BEAM_TIMEOUT_SEC", 10))
   ret = None
   try:
+    st = time.perf_counter()
     prg = to_program(x[1].copy().get_optimized_ast(name_override="test"), x[1].ren)
+    et = time.perf_counter() - st
     uops = prg.src[2].src
     if len(uops) >= (uops_max:=getenv("BEAM_UOPS_MAX", 3000)) > 0:
       if getenv("BEAM_LOG_SURPASS_MAX"): print(f"too many uops. {len(uops)=}, {uops_max=}")
       raise RuntimeError("too many uops")
-    st = time.perf_counter()
-    prog = prg.src[4].arg if len(prg.src) > 4 and prg.src[4].op is Ops.BINARY else compiler.compile(prg.src[3].arg)
-    et = time.perf_counter() - st
-    ret = (prg, prog, et)
+    ret = (prg, et)
   except RuntimeError:
     if DEBUG >= 4: traceback.print_exc()
   except Exception as e:
@@ -150,12 +149,11 @@ def beam_search(s:Scheduler, rawbufs:list[Buffer], amt:int, allow_test_size=True
     while not exiting:
       candidates: list[Scheduler] = flatten([get_kernel_actions(si, include_0=False).values() for si,_ in beam])
       timed: list[tuple[Scheduler, float]] = []
-      _compile_fn = functools.partial(_try_compile, compiler=dev.compiler)
       least_compute_ops = math.inf
-      for i,proc in (map(_compile_fn, enumerate(candidates)) if beam_pool is None else beam_pool.imap_unordered(_compile_fn, enumerate(candidates))):
+      for i, proc in ((map if beam_pool is None else beam_pool.imap_unordered)(_try_compile, enumerate(candidates))):
         if proc is None: continue
-        prg, lib, compile_et = proc
-        if lib in seen_libs: continue
+        prg, compile_et = proc
+        if (lib:=prg.src[4].arg) in seen_libs: continue
         # filter out kernels that use 1000x more compute than the smallest
         estimates = prg.src[0].arg.estimates
         least_compute_ops = min(this_compute_ops:=sym_infer(estimates.ops if estimates is not None else 0, var_vals), least_compute_ops)
@@ -163,7 +161,7 @@ def beam_search(s:Scheduler, rawbufs:list[Buffer], amt:int, allow_test_size=True
           if getenv("BEAM_LOG_SURPASS_MAX"): print(f"too much compute. {this_compute_ops} when least is {least_compute_ops}")
           continue
         seen_libs.add(lib)
-        try: tms = _time_program(prg, lib, var_vals, rawbufs, early_stop=beam[0][1]*3 if len(beam) else 1.0,
+        try: tms = _time_program(prg, var_vals, rawbufs, early_stop=beam[0][1]*3 if len(beam) else 1.0,
                                  allow_test_size=allow_test_size, clear_l2=hasattr(dev, 'invalidate_caches'),
                                  dev_timeout=getenv("BEAM_DEV_TIMEOUT", 1))
         except Exception as e:

tinygrad/codegen/simplify.py

@@ -14,7 +14,7 @@ def flatten_range(r:UOp) -> UOp|None:
 
 pm_flatten_range = PatternMatcher([
   # real ranges only
-  (UPat((Ops.REDUCE, Ops.STORE, Ops.END), name="r"), flatten_range),
+  (UPat((Ops.REDUCE, Ops.END), name="r"), flatten_range),
 ])
 
 def count_divmod(x:UOp) -> int: return sum(u.op in {Ops.IDIV, Ops.MOD} for u in x.backward_slice)
@@ -75,14 +75,14 @@ pm_split_ranges = PatternMatcher([
 def no_range(u:UOp) -> bool: return not any(x.op is Ops.RANGE for x in u.backward_slice_with_self)
 
 def reduce_unparented(red:UOp) -> UOp|None:
-  if red.arg not in {Ops.ADD, Ops.MAX, Ops.MUL}: return None
+  if red.arg[0] not in {Ops.ADD, Ops.MAX, Ops.MUL}: return None
   assert all(x.op is Ops.RANGE for x in red.src[1:]), "some reduce srcs aren't ranges"
   reduce_parented, reduce_unparented = partition(red.src[1:], lambda x: x in red.src[0].ranges)
   if len(reduce_unparented) == 0: return None
   ret = red.replace(src=(red.src[0],)+tuple(reduce_parented)) if len(reduce_parented) or red.dtype != red.src[0].dtype else red.src[0]
-  if red.arg is Ops.ADD:
+  if red.arg[0] is Ops.ADD:
     for r in reduce_unparented: ret = ret * r.src[0].cast(ret.dtype.scalar()).broadcast(ret.dtype.count)
-  if red.arg is Ops.MUL:
+  if red.arg[0] is Ops.MUL:
     for r in reduce_unparented: ret = ret ** r.src[0].cast(ret.dtype.scalar()).broadcast(ret.dtype.count)
   return ret
 
@@ -145,12 +145,12 @@ def reduce_load_collapse(red:UOp, u:UOp) -> UOp|None: return reduce_collapse(red
 
 # remove REDUCE without loads (generic arange opt / indexing).
 pm_reduce_simplify = pm_reduce_unparented + PatternMatcher([
-  (UPat(Ops.REDUCE, src=(UPat.var("u"),), allow_any_len=True, arg=Ops.ADD, name="red"), reduce_collapse),
+  (UPat(Ops.REDUCE, src=(UPat.var("u"),), allow_any_len=True, arg=(Ops.ADD, ()), name="red"), reduce_collapse),
 ])
 # remove REDUCE on load, comes from indexing a tensor with another tensor
 def no_load(u:UOp) -> bool: return not any(x.op is Ops.INDEX for x in u.backward_slice_with_self)
 pm_load_collapse = PatternMatcher([
-  (UPat(Ops.REDUCE, arg=Ops.ADD, src=(UPat.var("u"), UPat()), name="red"), reduce_load_collapse),
+  (UPat(Ops.REDUCE, arg=(Ops.ADD, ()), src=(UPat.var("u"), UPat()), name="red"), reduce_load_collapse),
   # we want to make sure we dont do math on a loaded index since that can cause overflow, this undoes the rule in pm_reduce_load_collapse
   ((UPat.var("x", dtypes.weakint)+UPat.var("y"))<UPat.var("c"), lambda x,y,c: x < c-y if no_load(y) and no_load(c) and not no_load(x) else None),
 ])

tinygrad/dtype.py

@@ -91,7 +91,7 @@ class DType(metaclass=DTypeMetaClass):
     return float("inf") if dtypes.is_float(self) else True
   def const(self, val: tuple[ConstType, ...]|ConstType):
     if isinstance(val, tuple):
-      #assert len(val) == self.count, f"mismatch {val} {self}"
+      assert len(val) == self.count, f"mismatch {val} {self}"
       return tuple(map(self.const, val))
     if isinstance(val, InvalidType): return val
     # NOTE: float('nan') != float('nan'), so we canonicalize here

tinygrad/engine/jit.py

@@ -1,12 +1,12 @@
 from typing import TypeVar, Generic, Callable, Any
 import functools, collections
 from tinygrad.tensor import Tensor
-from tinygrad.helpers import flatten, merge_dicts, DEBUG, Context, BEAM, getenv, colored, JIT, JIT_BATCH_SIZE, dedup, pluralize, VIZ
+from tinygrad.helpers import flatten, merge_dicts, DEBUG, Context, BEAM, getenv, JIT, JIT_BATCH_SIZE, dedup, pluralize, VIZ
 from tinygrad.device import Buffer, Compiled, Device, MultiBuffer
 from tinygrad.dtype import DType, dtypes
 from tinygrad.uop.ops import UOp, PatternMatcher, Variable, sym_infer, Ops, buffers, track_rewrites, graph_rewrite
-from tinygrad.engine.realize import capturing, CompiledRunner, Runner, Estimates, compile_linear, run_linear, get_runner, graph_cache, estimate_uop
-from tinygrad.engine.realize import unwrap_multi, resolve_params
+from tinygrad.engine.realize import capturing, Estimates, compile_linear, run_linear, graph_cache, estimate_uop, get_runtime
+from tinygrad.engine.realize import unwrap_multi, resolve_params, get_call_arg_uops, get_call_outs_ins
 from tinygrad.schedule.memory import memory_plan_rewrite, _collect_bufs
 from tinygrad.nn.state import get_parameters
 from tinygrad.schedule.rangeify import mop_cleanup
@@ -59,14 +59,6 @@ def graph_split_rewrite(linear:UOp, max_batch_size:int=0) -> UOp:
   if current_batch: flush_batch()
   return linear.replace(src=tuple(new_src))
 
-def _call_outs_ins(call:UOp) -> tuple[set[int], set[int]]:
-  non_bind = [s for s in call.src[1:] if s.op is not Ops.BIND]
-  ast = call.src[0]
-  if ast.op is Ops.PROGRAM: return set(ast.arg.outs), set(ast.arg.ins)
-  if ast.op in (Ops.COPY, Ops.BUFFER_VIEW): return {0}, {1}
-  if ast.op is Ops.CUSTOM_FUNCTION and ast.arg == "encdec": return {0}, set(range(1, len(non_bind)))
-  return set(), set()
-
 def _copy_input(u:UOp) -> UOp:
   run_linear(UOp(Ops.LINEAR, src=(u.copy_to_device(u.device).call(new:=UOp.new_buffer(u.device, u.arg, u.dtype), u, metadata=()),)))
   return new
@@ -95,17 +87,17 @@ def _check_no_non_tensor_return(ret):
 
 def graph_class(dev): return dev.graph.func if isinstance(dev.graph, functools.partial) else dev.graph
 
-class GraphRunner(Runner):
+class GraphRunner:
   def __init__(self, linear:UOp, input_uops:tuple[UOp, ...]=()):
     self.linear = linear.src[0]
     self.calls: list[tuple[int, UOp, list[Buffer], dict[str, int]]] = []
-    self.progs: list[CompiledRunner|None] = []
+    self.runtimes: list[Any|None] = []
     self.uop_replace: list[list[tuple[int, int]]] = []
     for call in self.linear.src:
-      replace = [(p, b.arg) for p, b in enumerate(b for b in call.src[1:] if b.op is not Ops.BIND) if b.op is Ops.PARAM]
+      replace = [(p, b.arg) for p, b in enumerate(get_call_arg_uops(call)) if b.op is Ops.PARAM]
       for dev_idx, (bufs, device_vars) in enumerate(unwrap_multi(call, resolve_params(call, input_uops))):
         self.calls.append((dev_idx, call.src[0], [b.ensure_allocated() for b in bufs], device_vars))
-        self.progs.append(get_runner(bufs[0].device, call.src[0]) if call.src[0].op is Ops.PROGRAM else None)
+        self.runtimes.append(get_runtime(bufs[0].device, call.src[0]) if call.src[0].op is Ops.PROGRAM else None)
         self.uop_replace.append(replace)
 
     self.var_vals_replace:dict[int, list[tuple[int, int]]] = {}
@@ -114,20 +106,20 @@ class GraphRunner(Runner):
 
     def is_sym_dim(dim) -> bool: return not all(isinstance(d, (int, float)) for d in dim)
 
-    crs = [(j, p, self.calls[j][3]) for j,p in enumerate(self.progs) if isinstance(p, CompiledRunner)]
-    self.vars = sorted({v.expr for _,p,dv in crs for v in p.p.vars if v.expr not in dv | p.p.runtimevars})
-    self.symbolic_dims = dedup(tuple(d) for _,p,_ in crs for d in (p.p.local_size, p.p.global_size) if d and is_sym_dim(d))
+    crs = [(j, self.calls[j][1].arg, self.calls[j][3]) for j in range(len(self.calls)) if self.calls[j][1].op is Ops.PROGRAM]
+    self.vars = sorted({v.expr for _,p,dv in crs for v in p.vars if v.expr not in dv | p.runtimevars})
+    self.symbolic_dims = dedup(tuple(d) for _,p,_ in crs for d in (p.local_size, p.global_size) if d and is_sym_dim(d))
 
     def find_symbolic_dim(dim): return self.symbolic_dims.index(tuple(dim)) if dim is not None and tuple(dim) in self.symbolic_dims else None
 
     for j,p,dv in crs:
-      if (replace:=[(i, self.vars.index(v.expr)) for i, v in enumerate(p.p.vars) if v.expr not in dv | p.p.runtimevars]):
+      if (replace:=[(i, self.vars.index(v.expr)) for i, v in enumerate(p.vars) if v.expr not in dv | p.runtimevars]):
         self.var_vals_replace[j] = replace
-      global_dim_idx, local_dim_idx = find_symbolic_dim(p.p.global_size), find_symbolic_dim(p.p.local_size)
+      global_dim_idx, local_dim_idx = find_symbolic_dim(p.global_size), find_symbolic_dim(p.local_size)
       if global_dim_idx is not None or local_dim_idx is not None:
         self.launch_dims_replace[j] = (global_dim_idx, local_dim_idx)
-        assert p.p.local_size is not None
-        self.launch_dims_base[j] = (tuple(p.p.global_size), tuple(p.p.local_size))
+        assert p.local_size is not None
+        self.launch_dims_base[j] = (tuple(p.global_size), tuple(p.local_size))
 
     estimates = sum((estimate_uop(call) for call in self.linear.src), Estimates())
 
@@ -135,7 +127,9 @@ class GraphRunner(Runner):
     self.w_dependency_map: dict[int, list[tuple[int, int, Any]]] = collections.defaultdict(list)
     self.r_dependency_map: dict[int, list[tuple[int, int, Any]]] = collections.defaultdict(list)
 
-    super().__init__(colored(f"<batched {len(self.calls)}>", "cyan"), self.calls[0][2][0].device.split(":")[0], estimates.simplify())
+    self.device, self.estimates = self.calls[0][2][0].device.split(":")[0], estimates.simplify()
+
+  def __call__(self, input_uops:tuple[UOp, ...], var_vals:dict[str, int], wait=False) -> float|None: raise NotImplementedError("override this")
 
   def updated_vars(self, var_vals: dict[str, int]):
     vals = [var_vals[v] for v in self.vars]
@@ -168,7 +162,7 @@ class GraphRunner(Runner):
 
   @staticmethod
   def _all_devs(batch_devs:list[Compiled], new_call:UOp) -> list[Compiled]:
-    return dedup(batch_devs + [Device[x] for b in new_call.src[1:] if b.op is not Ops.BIND
+    return dedup(batch_devs + [Device[x] for b in get_call_arg_uops(new_call)
                  for x in (b.device if isinstance(b.device, tuple) else (b.device,))])
 
   @staticmethod
@@ -197,9 +191,9 @@ class CapturedJit(Generic[ReturnType]):
     out: set[UOp] = set()
     for call in self.linear.toposort():
       if call.op is not Ops.CALL: continue
-      non_bind = [s for s in call.src[1:] if s.op is not Ops.BIND]
-      outs, ins = _call_outs_ins(call)
-      out |= {non_bind[k] for k in outs - ins if non_bind[k].op in (Ops.BUFFER, Ops.BUFFER_VIEW)}
+      arg_uops = get_call_arg_uops(call)
+      outs, ins = get_call_outs_ins(call)
+      out |= {arg_uops[k] for k in set(outs) - set(ins) if arg_uops[k].op in (Ops.BUFFER, Ops.BUFFER_VIEW)}
     return out
 
   def __call__(self, input_uops:list[UOp], var_vals:dict[str, int]) -> ReturnType:

tinygrad/engine/realize.py

@@ -1,9 +1,9 @@
-from typing import cast, Iterator
+from __future__ import annotations
+from typing import cast, Iterator, Any
 import time, random, itertools, math, contextlib, weakref
 from dataclasses import dataclass, replace, field
-from tinygrad.helpers import colored, DEBUG, GlobalCounters, ansilen, NOOPT, all_int, Metadata, TRACEMETA, TracingKey
-from tinygrad.helpers import BEAM, DEVECTORIZE, size_to_str, time_to_str, VALIDATE_WITH_CPU, cpu_profile, PROFILE, ProfilePointEvent, cpu_events
-from tinygrad.helpers import prod, EMULATED_DTYPES, flatten
+from tinygrad.helpers import colored, DEBUG, GlobalCounters, ansilen, all_int, TRACEMETA, prod, flatten
+from tinygrad.helpers import BEAM, size_to_str, time_to_str, VALIDATE_WITH_CPU, PROFILE, ProfilePointEvent, cpu_events
 from tinygrad.dtype import dtypes
 from tinygrad.uop.ops import Ops, PatternMatcher, UOp, UPat, sym_infer, buffers, graph_rewrite, ProgramInfo
 from tinygrad.device import Device, Buffer, MultiBuffer
@@ -11,50 +11,74 @@ from tinygrad.renderer import Estimates
 from tinygrad.codegen import to_program
 from tinygrad.codegen.opt.postrange import bufs_from_ast
 
+# **************** Helpers ****************
+
+def get_call_arg_uops(call:UOp) -> tuple[UOp, ...]: return tuple(s for s in call.src[1:] if s.op is not Ops.BIND)
+
+def get_call_outs_ins(call:UOp) -> tuple[tuple[int, ...], tuple[int, ...]]:
+  ast = call.src[0]
+  if ast.op is Ops.PROGRAM: return tuple(ast.arg.outs), tuple(ast.arg.ins)
+  if ast.op in (Ops.COPY, Ops.BUFFER_VIEW): return (0,), (1,)
+  if ast.op is Ops.CUSTOM_FUNCTION and ast.arg == "encdec": return (0,), tuple(range(1, len(get_call_arg_uops(call))))
+  return (), ()
+
+def get_call_name(call:UOp, bufs:list[Buffer], var_vals:dict[str, int]|None=None) -> str:
+  def _uop_sz_to_str(uop:UOp) -> str: return size_to_str(sym_infer(prod(uop.shape) * uop.dtype.itemsize, var_vals or {}))
+
+  ast, arg_uops = call.src[0], get_call_arg_uops(call)
+  if ast.op is Ops.PROGRAM: return ast.arg.name
+  if ast.op is Ops.BUFFER_VIEW: return colored(f"view {_uop_sz_to_str(arg_uops[0]):>10} @ {ast.arg[1] * arg_uops[1].dtype.itemsize:<10d}", "yellow")
+  if ast.op is Ops.COPY: return colored(f"copy {_uop_sz_to_str(arg_uops[0]):>10}, {bufs[0].device[:7]:>7s} <- {bufs[1].device[:7]:7s}", "yellow")
+  if ast.op is Ops.CUSTOM_FUNCTION and ast.arg == "encdec": return colored(f"enc/dec {_uop_sz_to_str(arg_uops[0])}", "yellow")
+  if ast.op is Ops.CUSTOM_FUNCTION and ast.arg == "graph": return colored(f"batched {len(ast.src[0].src)}", "cyan")
+  raise NotImplementedError("get_call_name is not implemented")
+
 # **************** Stat ****************
 
 def estimate_uop(call:UOp) -> Estimates:
-  if call.src[0].op is Ops.SINK: call = pm_compile.rewrite(call)
-
   ast = call.src[0]
   if ast.op is Ops.PROGRAM: return ast.src[0].arg.estimates or Estimates()
   if ast.op is Ops.COPY or (ast.op is Ops.CUSTOM_FUNCTION and ast.arg == "encdec"):
     nbytes = prod(call.src[1].shape) * call.src[1].dtype.itemsize
     return Estimates(lds=nbytes, mem=nbytes)
-  if ast.op is Ops.CUSTOM_FUNCTION and ast.arg == "graph":
-    return runner.estimates if (runner:=graph_cache.get(ast)) is not None else Estimates()
+  if ast.op is Ops.CUSTOM_FUNCTION and ast.arg == "graph": return get_graph_runtime(ast).estimates
   return Estimates()
 
-def update_stats(display_name:str, device:str, estimates:Estimates, var_vals:dict[str, int], et:float|None, buf_count:int,
-                 jit=False, metadata:tuple[Metadata, ...]=(), first_run=False):
+first_run_cache:set[bytes] = set()
+@contextlib.contextmanager
+def track_stats(ctx:ExecContext, call:UOp, device:str, bufs:list[Buffer], var_vals:dict[str, int]):
+  if PROFILE:
+    outputs, inputs = get_call_outs_ins(call)
+    cpu_events.append(ProfilePointEvent(device, "exec", len(cpu_events), {"metadata": call.arg.metadata, "var_vals": var_vals,
+      "bufs": [b.trace_num for b in bufs], "name": get_call_name(call, bufs, var_vals), "outputs": outputs, "inputs": inputs}))
+  et: list[float|None] = [None]
+  if DEBUG >= 2: st = time.perf_counter()
+  yield et
+  if not ctx.do_update_stats: return
+
+  if DEBUG >= 2 and et[0] is None:
+    Device[device].synchronize()
+    et[0] = time.perf_counter() - st
+
+  estimates = estimate_uop(call)
   GlobalCounters.kernel_count += 1
   GlobalCounters.global_ops += (op_est:=sym_infer(estimates.ops, var_vals))
   GlobalCounters.global_mem += (mem_est:=sym_infer(estimates.mem, var_vals))
-  if et is not None: GlobalCounters.time_sum_s += et
+  if et[0] is not None: GlobalCounters.time_sum_s += et[0]
   if DEBUG >= 2:
+    display_name = get_call_name(call, bufs, var_vals)
     lds_est = sym_infer(estimates.lds, var_vals)
-    header_color = 'magenta' if jit else ('green' if first_run else None)
-    ptm = colored(time_to_str(et, w=9), "yellow" if et > 0.01 else None) if et is not None else ""
-    flops, membw, ldsbw = op_est/(et or 1e-20), mem_est/(et or 1e-20), lds_est/(et or 1e-20)
+    header_color = 'magenta' if ctx.jit else ('green' if call.src[0].key not in first_run_cache else None)
+    ptm = colored(time_to_str(et[0], w=9), "yellow" if et[0] > 0.01 else None) if et[0] is not None else ""
+    flops, membw, ldsbw = op_est/(et[0] or 1e-20), mem_est/(et[0] or 1e-20), lds_est/(et[0] or 1e-20)
     flops_str = f"{flops*1e-9:7.0f} GFLOPS" if flops < 1e14 else colored(f"{flops*1e-12:7.0f} TFLOPS", 'green')
     mem_str = f"{membw*1e-9:4.0f}|{ldsbw*1e-9:<6.0f} GB/s" if membw < 1e13 and ldsbw < 1e15 else \
       colored(f"{membw*1e-12:4.0f}|{ldsbw*1e-12:<6.0f} TB/s", 'green')
     print(f"{colored(f'*** {device[:7]:7s} {GlobalCounters.kernel_count:4d}', header_color)}"+
-      f" {display_name+' '*(46-ansilen(display_name))} arg {buf_count:2d} mem {GlobalCounters.mem_used/1e9:6.2f} GB"+
-      ("" if et is None else f" tm {ptm}/{GlobalCounters.time_sum_s*1e3:9.2f}ms ({flops_str} {mem_str})")+
-      f" {[repr(m) if TRACEMETA >= 2 else str(m) for m in metadata] if metadata else ''}")
-
-# **************** Runners ****************
-
-class Runner:
-  def __init__(self, display_name:str, device:str, estimates=Estimates()):
-    self.first_run, self.display_name, self.device, self.estimates = True, display_name, device, estimates
-  @property
-  def dev(self): return Device[self.device]
-  def exec(self, rawbufs:list[Buffer], var_vals:dict[str, int]|None=None) -> float|None:
-    return self(rawbufs, {} if var_vals is None else var_vals)
-  def __call__(self, rawbufs:list[Buffer], var_vals:dict[str, int], wait=False) -> float|None:
-    raise NotImplementedError("override this")
+      f" {display_name+' '*(46-ansilen(display_name))} arg {len(bufs):2d} mem {GlobalCounters.mem_used/1e9:6.2f} GB"+
+      ("" if et[0] is None else f" tm {ptm}/{GlobalCounters.time_sum_s*1e3:9.2f}ms ({flops_str} {mem_str})")+
+      f" {[repr(m) if TRACEMETA >= 2 else str(m) for m in call.arg.metadata] if call.arg.metadata else ''}")
+    first_run_cache.add(call.src[0].key)
 
 local_size_cache: dict[bytes, tuple[int, ...]] = {}
 def optimize_local_size(call:UOp, prg:UOp) -> UOp|None:
@@ -78,43 +102,24 @@ def optimize_local_size(call:UOp, prg:UOp) -> UOp|None:
   new_global = tuple(g//l if g%l == 0 else g/l for g,l in zip(prg.arg.global_size, local_size))
   return call.replace(src=(prg.replace(arg=replace(prg.arg, global_size=new_global, local_size=local_size)), *call.src[1:]))
 
-class CompiledRunner(Runner):
-  def __init__(self, prg:UOp, device:str):
-    info: ProgramInfo = prg.arg
-    sink = prg.src[0]
-    if DEBUG >= 3 and sink.arg.applied_opts: print(sink.arg.applied_opts)
-    if DEBUG >= 4: print(prg.src[3].arg)
-    if len(prg.src) <= 4 or prg.src[4].op is not Ops.BINARY:
-      with cpu_profile(TracingKey(f"compile {info.name}", (info.function_name,)), "TINY"):
-        lib = Device[device].compiler.compile_cached(prg.src[3].arg)
-      prg = prg.replace(src=prg.src + (UOp(Ops.BINARY, arg=lib),))
-    self.prg:UOp = prg
-    self.p:ProgramInfo = info
-    if DEBUG >= 7: Device[device].compiler.disassemble(prg.src[4].arg)
-    self._prg = Device[device].runtime(info.function_name, prg.src[4].arg, *info.aux, runtimevars=info.runtimevars)
-    super().__init__(info.name, device, sink.arg.estimates or Estimates())
+# **************** runtime cache ****************
 
-  def __call__(self, rawbufs:list[Buffer], var_vals:dict[str, int]|None=None, wait=False, timeout:int|None=None) -> float|None:
-    if var_vals is None: var_vals = {}
-    global_size, local_size = self.p.launch_dims(var_vals)
-    return self._prg(*[x._buf for x in rawbufs], global_size=tuple(global_size), local_size=tuple(local_size) if local_size else None,
-                     vals=tuple(var_vals[k.expr] if k.expr not in self.p.runtimevars else None for k in self.p.vars), wait=wait, timeout=timeout)
+runtime_cache: dict[tuple[bytes, str], Any] = {}
+def get_runtime(device:str, ast:UOp):
+  assert ast.op is Ops.PROGRAM and isinstance(ast.arg, ProgramInfo), "get_runtime should only be called with a PROGRAM ast"
+  if (runtime:=runtime_cache.get(key:=(ast.key, device))) is None:
+    if DEBUG >= 3 and ast.src[0].arg.applied_opts: print(ast.src[0].arg.applied_opts)
+    if DEBUG >= 4: print(ast.src[3].arg)
+    if DEBUG >= 7: Device[device].compiler.disassemble(ast.src[4].arg)
+    runtime = runtime_cache[key] = Device[device].runtime(ast.arg.function_name, ast.src[4].arg, *ast.arg.aux, runtimevars=ast.arg.runtimevars)
+  return runtime
 
-# **************** method cache ****************
-
-method_cache: dict[tuple[str, type, bytes, tuple, bool], CompiledRunner] = {}
-def get_runner(device:str, ast:UOp) -> CompiledRunner:
-  # TODO: this should be all context relevant to rendering
-  context = (NOOPT.value, DEVECTORIZE.value, EMULATED_DTYPES.value)
-  ckey = (device, type(Device[device].compiler), ast.key, context, False)
-  if cret:=method_cache.get(ckey): return cret
-  bkey = (device.split(":")[0], type(Device[device].compiler), ast.key, context, True)
-  if bret:=method_cache.get(bkey):
-    method_cache[ckey] = ret = CompiledRunner(bret.prg, device)
-  else:
-    prg = to_program(ast, Device[device].renderer)
-    method_cache[ckey] = method_cache[bkey] = ret = CompiledRunner(prg, device)
-  return ret
+graph_cache:weakref.WeakKeyDictionary[UOp, Any] = weakref.WeakKeyDictionary()
+def get_graph_runtime(ast:UOp, input_uops:tuple[UOp, ...]|None=None):
+  assert ast.op is Ops.CUSTOM_FUNCTION and ast.arg == "graph", "get_graph_runtime should only be called with a graph ast"
+  if (runtime:=graph_cache.get(ast)) is None and input_uops is not None:
+    graph_cache[ast] = runtime = Device[ast.device if isinstance(ast.device, str) else ast.device[0]].graph(ast, input_uops=input_uops)
+  return runtime
 
 # **************** run linear ****************
 
@@ -130,21 +135,7 @@ class ExecContext:
 def _resolve(b:UOp, inputs:tuple[UOp, ...]) -> UOp:
   if b.op in (Ops.BUFFER_VIEW, Ops.MSELECT) and b.src[0].op is Ops.PARAM: return b.replace(src=(inputs[b.src[0].arg], *b.src[1:]))
   return inputs[b.arg] if b.op is Ops.PARAM else b
-def resolve_params(call:UOp, inputs:tuple[UOp, ...]) -> list[UOp]: return [_resolve(b, inputs) for b in call.src[1:] if b.op is not Ops.BIND]
-
-@contextlib.contextmanager
-def track_stats(ctx:ExecContext, call:UOp, device:str, display_name:str, bufs:list[Buffer], var_vals:dict[str, int],
-                outputs=(0,), inputs=(1,), first_run=False):
-  if PROFILE: cpu_events.append(ProfilePointEvent(device, "exec", len(cpu_events), {"metadata": call.arg.metadata, "var_vals": var_vals,
-                                                  "bufs": [b.trace_num for b in bufs], "name": display_name, "outputs": outputs, "inputs": inputs}))
-  timing: list[float|None] = [None]
-  if DEBUG >= 2: st = time.perf_counter()
-  yield timing
-  if not ctx.do_update_stats: return
-  if DEBUG >= 2 and timing[0] is None:
-    Device[device].synchronize()
-    timing[0] = time.perf_counter() - st
-  update_stats(display_name, device, estimate_uop(call), var_vals, timing[0], len(bufs), jit=ctx.jit, metadata=call.arg.metadata, first_run=first_run)
+def resolve_params(call:UOp, inputs:tuple[UOp, ...]) -> list[UOp]: return [_resolve(b, inputs) for b in get_call_arg_uops(call)]
 
 def unwrap_multi(call:UOp, resolved:list[UOp]) -> Iterator[tuple[list[Buffer], dict[str, int]]]:
   bufs = [b.buffer for b in resolved]
@@ -157,16 +148,13 @@ def exec_view(ctx:ExecContext, call, ast):
   resolved = resolve_params(call, ctx.input_uops)
   bufs = [cast(Buffer, b.buffer) for b in resolved]
   bv = bufs[1].view(resolved[0].arg, ast.dtype, ast.arg[1]*bufs[1].dtype.itemsize)
-  with track_stats(ctx, call, bv.device, colored(f"view {bv.nbytes:8d} @ {bv.offset:<10d}", "yellow"), [bv, bufs[1]], ctx.var_vals):
-    buffers[resolved[0]] = bv
+  with track_stats(ctx, call, bv.device, [bv, bufs[1]], ctx.var_vals): buffers[resolved[0]] = bv
 
 def exec_copy(ctx:ExecContext, call, ast):
   for bufs, device_vars in unwrap_multi(call, resolve_params(call, ctx.input_uops)):
     dest, src = bufs[0].ensure_allocated(), bufs[1].ensure_allocated()
-    xfer = hasattr(dest.allocator,'_transfer') and dest.allocator.supports_transfer and dest.device.split(":")[0] == src.device.split(":")[0]
-    name = colored(f"{'xfer' if xfer else 'copy'} {size_to_str(bufs[0].nbytes):>10}, {dest.device[:7]:>7s} <- {src.device[:7]:7s}", "yellow")
-    with track_stats(ctx, call, dest.device, name, [dest, src], ctx.var_vals):
-      if xfer:
+    with track_stats(ctx, call, dest.device, [dest, src], ctx.var_vals):
+      if hasattr(dest.allocator,'_transfer') and dest.allocator.supports_transfer and dest.device.split(":")[0] == src.device.split(":")[0]:
         dest.allocator._transfer(dest._buf, src._buf, dest.nbytes, src_dev=src.allocator.dev, dest_dev=dest.allocator.dev) # type:ignore[attr-defined]
       elif src.device.startswith("DISK") and getattr(src.allocator.dev, 'fd', None) is not None \
            and hasattr(dest.allocator, 'copy_from_disk') and src.nbytes >= 4096 and dest.allocator.supports_copy_from_disk:
@@ -178,35 +166,31 @@ def exec_copy(ctx:ExecContext, call, ast):
 def exec_kernel(ctx:ExecContext, call, ast):
   for bufs, device_vars in unwrap_multi(call, resolve_params(call, ctx.input_uops)):
     var_vals = {**ctx.var_vals, **device_vars}
-    prg = get_runner(bufs[0].device, ast)
-    prg_bufs = [bufs[i].ensure_allocated() for i in prg.p.globals]
-
-    with track_stats(ctx, call, prg.device, prg.display_name, prg_bufs, var_vals,
-                     outputs=tuple(prg.p.outs), inputs=tuple(prg.p.ins), first_run=prg.first_run) as timing:
-      timing[0] = prg(prg_bufs, var_vals, wait=DEBUG >= 2)
-      prg.first_run = False
+    prg_bufs = [bufs[i].ensure_allocated() for i in ast.arg.globals]
+    rt = get_runtime(device:=bufs[0].device, ast)
+    global_size, local_size = ast.arg.launch_dims(var_vals)
+    with track_stats(ctx, call, device, prg_bufs, var_vals) as tm:
+      tm[0] = rt(*[b._buf for b in prg_bufs], global_size=global_size, local_size=local_size, vals=ast.arg.vals(var_vals), wait=DEBUG>=2)
 
 def exec_validate(ctx:ExecContext, call, ast):
   import numpy as np
   for bufs, device_vars in unwrap_multi(call, resolve_params(call, ctx.input_uops)):
-    cpu_bufs, dev_bufs = bufs[:len(bufs)//2], bufs[len(bufs)//2:]
-    cpu_prg = get_runner("CPU", ast.src[0])
-    cpu_prg([cpu_bufs[i].ensure_allocated() for i in cpu_prg.p.globals], {**ctx.var_vals, **device_vars}, wait=False)
-    for i in cpu_prg.p.outs: np.testing.assert_allclose(dev_bufs[i].ensure_allocated().numpy(), cpu_bufs[i].numpy(), rtol=1e-3, atol=1e-3)
+    bufs, dev_bufs = bufs[:len(bufs)//2], bufs[len(bufs)//2:]
+    var_vals = {**ctx.var_vals, **device_vars}
+    cpu_rt = get_runtime("CPU", prg:=to_program(ast.src[0], Device["CPU"].renderer))
+    global_size, local_size = prg.arg.launch_dims(var_vals)
+    cpu_rt(*[bufs[i].ensure_allocated()._buf for i in prg.arg.globals], global_size=global_size, local_size=local_size, vals=prg.arg.vals(var_vals))
+    for i in prg.arg.outs: np.testing.assert_allclose(dev_bufs[i].ensure_allocated().numpy(), bufs[i].numpy(), rtol=1e-3, atol=1e-3)
 
 def exec_encdec(ctx:ExecContext, call, ast):
   bufs = [cast(Buffer, b.buffer).ensure_allocated() for b in resolve_params(call, ctx.input_uops)]
   shape, pos_var = tuple(s.arg for s in ast.src if s.op is Ops.CONST), ast.variables()[0].expr
-  with track_stats(ctx, call, bufs[0].device, colored(f"enc/dec {size_to_str(bufs[0].nbytes)}", "yellow"), bufs, ctx.var_vals):
+  with track_stats(ctx, call, bufs[0].device, bufs, ctx.var_vals):
     bufs[0].allocator._encode_decode(bufs[0]._buf, bufs[1]._buf, bufs[2]._buf, [x._buf for x in bufs[3:]], shape, ctx.var_vals[pos_var])
 
-graph_cache:weakref.WeakKeyDictionary[UOp, Runner] = weakref.WeakKeyDictionary()
-def exec_graph(ctx:ExecContext, call, cf):
-  bufs = flatten([b.bufs if isinstance(b, MultiBuffer) else [b] for b in (u.buffer for u in resolve_params(call, ctx.input_uops))])
-  if (runner:=graph_cache.get(cf)) is None:
-    graph_cache[cf] = runner = Device[cf.device if isinstance(cf.device, str) else cf.device[0]].graph(cf, input_uops=ctx.input_uops)
-  with track_stats(ctx, call, runner.device, runner.display_name, bufs, ctx.var_vals) as t:
-    t[0] = runner(bufs, ctx.var_vals, wait=DEBUG >= 2, input_uops=ctx.input_uops) # type: ignore[call-arg]
+def exec_graph(ctx:ExecContext, call, ast):
+  rt = get_graph_runtime(ast, ctx.input_uops)
+  with track_stats(ctx, call, rt.device, [], ctx.var_vals) as t: t[0] = rt(ctx.input_uops, ctx.var_vals, wait=DEBUG>=2) # type: ignore[call-arg]
 
 # flatten LINEAR-in-LINEAR: any nested LINEAR child gets inlined into its parent's src
 pm_flatten_linear = PatternMatcher([
@@ -215,7 +199,7 @@ pm_flatten_linear = PatternMatcher([
 ])
 
 def _validate(call:UOp, sink:UOp) -> UOp:
-  params = tuple(p for p in call.src[1:] if p.op is not Ops.BIND)
+  params = get_call_arg_uops(call)
   shadows = tuple(UOp.new_buffer(("CPU",)*len(p.device) if isinstance(p.device, tuple) else "CPU", prod(p.max_shape), p.dtype.base) for p in params)
   copies = tuple(p.copy_to_device(s.device).call(s, p) for s, p in zip(shadows, params))
   return UOp(Ops.LINEAR, src=copies + (call, UOp(Ops.CUSTOM_FUNCTION, dtypes.void, src=(sink,), arg="validate").call(*shadows, *params)))
@@ -241,7 +225,7 @@ pm_exec = PatternMatcher([
   (UPat(Ops.CALL, src=(UPat(Ops.COPY, name="ast"),), name="call", allow_any_len=True), exec_copy),
   (UPat(Ops.CALL, src=(UPat(Ops.PROGRAM, name="ast"),), name="call", allow_any_len=True), exec_kernel),
   (UPat(Ops.CALL, src=(UPat(Ops.CUSTOM_FUNCTION, arg="encdec", name="ast"),), name="call", allow_any_len=True), exec_encdec),
-  (UPat(Ops.CALL, src=(UPat(Ops.CUSTOM_FUNCTION, arg="graph", name="cf"),), name="call", allow_any_len=True), exec_graph),
+  (UPat(Ops.CALL, src=(UPat(Ops.CUSTOM_FUNCTION, arg="graph", name="ast"),), name="call", allow_any_len=True), exec_graph),
   (UPat(Ops.CALL, src=(UPat(Ops.CUSTOM_FUNCTION, arg="validate", name="ast"),), name="call", allow_any_len=True), exec_validate),
 ])
 

tinygrad/gradient.py

@@ -8,7 +8,7 @@ def reduce_gradient(ctx:UOp, ret:UOp, op:Ops):
   def broadcast_to_input(x): return x.reshape(x.shape+(1,)*(len(ret.src[0].shape)-len(x.shape))).expand(ret.src[0].shape)
   if op == Ops.ADD: return (broadcast_to_input(ctx),)
   if op == Ops.MAX:
-    assert ret.op is Ops.REDUCE_AXIS, "only works on REDUCE_AXIS"
+    assert ret.op is Ops.REDUCE, "only works on REDUCE"
     mask = ret.src[0].eq(broadcast_to_input(ret)).cast(ctx.dtype)
     count = mask._rop(Ops.ADD, ret.arg[1])
     return ((mask/broadcast_to_input(count)) * broadcast_to_input(ctx),)
@@ -61,7 +61,7 @@ pm_gradient = PatternMatcher([
     ((x>y).where(ctx, (x.eq(y)).where(ctx * 0.5, 0)), (x<y).where(ctx, (x.eq(y)).where(ctx * 0.5, 0)))),
   (UPat(Ops.MUL, name="ret"), lambda ctx, ret: (ret.src[1]*ctx, ret.src[0]*ctx)),
   (UPat(Ops.WHERE, name="ret"), lambda ctx, ret: (None, ret.src[0].where(ctx, ctx.const_like(0)), ret.src[0].where(ctx.const_like(0), ctx))),
-  (UPat(Ops.REDUCE_AXIS, name="ret"), lambda ctx, ret: reduce_gradient(ctx, ret, ret.arg[0])),
+  (UPat(Ops.REDUCE, name="ret"), lambda ctx, ret: reduce_gradient(ctx, ret, ret.arg[0])),
   (UPat(Ops.CONTIGUOUS), lambda ctx: (ctx,)),
   (UPat(Ops.CONTIGUOUS_BACKWARD), lambda ctx: (ctx.contiguous(),)),
   (UPat(Ops.RESHAPE, name="ret"), lambda ctx, ret: (ctx.reshape(ret.src[0].shape), None)),

tinygrad/mixin/__init__.py

@@ -6,7 +6,7 @@ from tinygrad.mixin.movement import MovementMixin
 from tinygrad.mixin.reduce import ReduceMixin
 from tinygrad.uop import Ops
 from tinygrad.uop.ops import _broadcast_shape, resolve, smax, smin, identity_element
-from tinygrad.dtype import ConstType, DTypeLike, Invalid, InvalidType, PtrDType, PyConst, dtypes, least_upper_dtype, sum_acc_dtype, to_dtype
+from tinygrad.dtype import ConstType, DType, DTypeLike, Invalid, InvalidType, PtrDType, PyConst, dtypes, least_upper_dtype, sum_acc_dtype, to_dtype
 from tinygrad.helpers import all_int, argfix, ceildiv, flatten, flat_to_grouped, make_tuple, prod, resolve_pool_pads, round_up
 
 if TYPE_CHECKING:
@@ -191,6 +191,36 @@ class OpMixin(ElementwiseMixin, ReduceMixin):
     """
     return self._tri(self.shape[-2], self.shape[-1], diagonal+1, self.device).where(self.zeros_like(), self)
 
+  # ***** random *****
+
+  @staticmethod
+  def _threefry_random_bits(key, counts0, counts1):
+    x = (counts1.cast(dtypes.uint64) << 32) | counts0.cast(dtypes.uint64)
+    x = x.threefry((key[1]._broadcast_to(x.shape).cast(dtypes.uint64) << 32) | key[0]._broadcast_to(x.shape).cast(dtypes.uint64))
+    return (x & 0xffffffff).cast(dtypes.uint32).cat(((x >> 32) & 0xffffffff).cast(dtypes.uint32))
+
+  @classmethod
+  def random_bits(cls, key:Self, counter:Self, num:int) -> Self:
+    low, high = counter[0:1], counter[1:2]
+    bits = []
+    for i in range(0, num, dtypes.uint32.max):
+      chunk_num = min(num - i, dtypes.uint32.max)
+      c_low = low + (i & 0xffffffff)
+      c_high = high + (i >> 32) + (c_low < low).cast(dtypes.uint32)
+      new_key = cls._threefry_random_bits(key, c_low, c_high)
+      counts0 = cls.arange(ceildiv(chunk_num, 2), device=key.device, dtype=dtypes.uint32)
+      counts1 = counts0 + ceildiv(chunk_num, 2)
+      bits.append(cls._threefry_random_bits(new_key, counts0, counts1)[:chunk_num])
+    return bits[0].cat(*bits[1:])
+
+  @staticmethod
+  def _bits_to_rand(bits, shape:tuple[int, ...], dtype:DType):
+    _, nmant = dtypes.finfo(dtype)
+    uint_dtype = {1: dtypes.uint8, 2: dtypes.uint16, 4: dtypes.uint32, 8: dtypes.uint64}[dtype.itemsize]
+    uint_bits = bits.bitcast(uint_dtype)
+    float_one_bits = uint_bits.ones_like(dtype=dtype).bitcast(uint_dtype)
+    return uint_bits.rshift(dtype.bitsize - nmant).bitwise_or(float_one_bits).bitcast(dtype)[:prod(shape)].sub(1).reshape(shape)
+
   def _pad_constant(self, pX, value:float) -> Self:
     # shrink first for negative pads, then pad with only non-negative values
     pX = tuple((0, 0) if p is None else p for p in pX)

tinygrad/mixin/dtype.py

@@ -7,6 +7,8 @@ class DTypeMixin:
 
   def cast(self, dtype:DType) -> Self: raise NotImplementedError
 
+  def bitcast(self, dtype:DType) -> Self: raise NotImplementedError
+
   def element_size(self) -> int:
     """
     Returns the size in bytes of an individual element in the tensor.

tinygrad/nn/onnx.py

@@ -644,7 +644,7 @@ def get_onnx_ops() -> dict[str, types.FunctionType|dict[OpSetId, types.FunctionT
   def BitwiseOr(x:Tensor,y:Tensor): return x | y
   def BitwiseXor(x:Tensor,y:Tensor): return x ^ y
   def BitwiseNot(x:Tensor): return ~x
-  def Mod(x:Tensor,y:Tensor,fmod=0): return x - x.div(y, rounding_mode="trunc") * y if fmod else x % y
+  def Mod(x:Tensor,y:Tensor,fmod=0): return x.fmod(y) if fmod else x % y
 
   # ***** Casting Ops *****
   # NOTE: saturate only applies to FP8 types

tinygrad/renderer/__init__.py

@@ -3,7 +3,7 @@ from typing import Callable, cast
 from dataclasses import dataclass
 from tinygrad.helpers import prod, Target
 from tinygrad.uop.ops import Ops, UOp, sint, ssimplify, smin, GroupOp, PatternMatcher
-from tinygrad.dtype import AddrSpace, PtrDType
+from tinygrad.dtype import AddrSpace, PtrDType, dtypes
 from tinygrad.codegen.opt.tc import TensorCore
 from tinygrad.device import Compiler
 
@@ -31,8 +31,11 @@ class Estimates:
         if u.op in {Ops.LOAD, Ops.STORE}:
           # if u.src[0] is INDEX, we have to include the buffer since it might be an AFTER
           dont_count = dont_count.union((UOp.sink(*u.src[0].src[1:]) if u.src[0].op is Ops.INDEX else u.src[0]).toposort(range_gate))
-          # TODO: is this correct? this all needs to be cleaned up
-          if len(u.src) > 2: dont_count = dont_count.union(u.src[2].toposort())
+          # gate (bool-typed src) is part of indexing/predication, exclude its computation
+          # LOAD: gate at src[1]; STORE: gate at src[2]
+          gate_pos = 1 if u.op is Ops.LOAD else 2
+          if len(u.src) > gate_pos and u.src[gate_pos].dtype.scalar() == dtypes.bool:
+            dont_count = dont_count.union(u.src[gate_pos].toposort(range_gate))
         elif u.op is Ops.IF:
           dont_count = dont_count.union(u.src[0].toposort())
     for u in uops:

tinygrad/renderer/cstyle.py

@@ -44,9 +44,9 @@ base_rewrite = PatternMatcher([
   # default const render
   (UPat(Ops.CONST, name="x"), lambda ctx,x: str(x.arg)),
   # new load/store
-  (UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var('idx')), allow_any_len=True),
+  (UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var('idx'))),
    lambda ctx,buf,idx: f"({ctx[buf]}+{strip_parens(ctx[idx]) if idx.arg == Ops.ADD else ctx[idx]})"),
-  (UPat(Ops.LOAD, src=(UPat(Ops.INDEX, src=(UPat(), UPat(), UPat.var("gate"))).or_casted("bidx"), UPat.var("var"))),
+  (UPat(Ops.LOAD, src=(UPat.var('bidx'), UPat.var("gate", dtype=dtypes.bool), UPat.var("var"))),
    lambda ctx,bidx,var,gate: f"({ctx[gate]}?*{ctx[bidx]}:{ctx[var]})"),
   (UPat(Ops.LOAD, src=(UPat.var('bidx'),)), lambda ctx,bidx: f"(*{ctx[bidx]})"),
   (UPat(Ops.STORE, src=(UPat.var('bidx'), UPat.var("var"))), lambda ctx,bidx,var: f"*{ctx[bidx]} = {ctx[var]};"),
@@ -302,12 +302,13 @@ class OpenCLRenderer(CStyleLanguage):
     (UPat(Ops.CONST, dtypes.bfloat16, name="x"),
       lambda ctx,x: f"{(struct.unpack('I', struct.pack('f', float_to_bf16(x.arg)))[0] >> 16)}u"),
     # load/store image (OpenCL)
-    (UPat(Ops.LOAD, dtype=dtypes.float.vec(4), src=(UPat.var('buf').index(UPat.var('idx', dtypes.int.vec(2)), UPat.var("gate")), UPat.var("var"))),
+    (UPat(Ops.LOAD, dtype=dtypes.float.vec(4), src=(UPat.var('buf').index(UPat.var('idx', dtypes.int.vec(2))), UPat.var("gate", dtype=dtypes.bool),
+                                                    UPat.var("var"))),
       lambda ctx,buf,idx,var,gate: f"({ctx[gate]}?read_imagef({ctx[buf]}, smp, {ctx[idx]}):{ctx[var]})"),
     (UPat(Ops.LOAD, dtype=dtypes.float.vec(4), src=(UPat.var('buf').index(UPat.var('idx', dtypes.int.vec(2))),)),
       lambda ctx,buf,idx: f"read_imagef({ctx[buf]}, smp, {ctx[idx]})"),
-    (UPat(Ops.STORE, src=(UPat.var('buf').index(UPat.var('idx', dtypes.int.vec(2)), allow_any_len=True),
-                          UPat.var("var", dtypes.float.vec(4))), allow_any_len=True),
+    (UPat(Ops.STORE, src=(UPat.var('buf').index(UPat.var('idx', dtypes.int.vec(2))),
+                          UPat.var("var", dtypes.float.vec(4)))),
       lambda ctx,buf,idx,var: f"write_imagef({ctx[buf]}, {ctx[idx]}, {ctx[var]});"),
   ]) + base_rewrite
 
@@ -536,6 +537,7 @@ class HIPRenderer(CStyleLanguage):
     if any(dt.scalar() == dtypes.half for dt in used_dtypes): prefix.append("#define half _Float16")
     if any(dt.scalar() in dtypes.fp8s for dt in used_dtypes):
       prefix += ["typedef unsigned char hip_bf8;", "typedef unsigned char hip_fp8;"]
+    if any(u.op is Ops.CAST and u.dtype in dtypes.fp8s and u.src[0].dtype == dtypes.float for u in uops):
       prefix.append("""static inline __attribute__((device)) unsigned char f32_to_fp8(float v, int is_bf8) {
   v = (((*(unsigned*)&v)&0x7F800000)!=0x7F800000)?__builtin_amdgcn_fmed3f(v,is_bf8?57344.0f:448.0f,is_bf8?-57344.0f:-448.0f) : v;
   return (unsigned char)(is_bf8?__builtin_amdgcn_cvt_pk_bf8_f32(v,v,0,false):__builtin_amdgcn_cvt_pk_fp8_f32(v,v,0,false));\n}""")

tinygrad/renderer/llvmir.py

@@ -76,7 +76,7 @@ base_rewrite = PatternMatcher([
   # memory load/store
   (UPat(Ops.INDEX, name="x"), lambda ctx,x:
    f"  {ctx[x]} = getelementptr inbounds {ldt(x.dtype.base)}, {ldt(x.src[0].dtype)} {ctx[x.src[0]]}, {ldt(x.src[1].dtype)} {ctx[x.src[1]]}"),
-  (UPat(Ops.LOAD, src=(UPat(Ops.INDEX, src=(UPat(), UPat(), UPat.var("mask"))).or_casted("idx"), UPat.var("alt")), allow_any_len=True, name="x"),
+  (UPat(Ops.LOAD, src=(UPat(Ops.INDEX).or_casted("idx"), UPat.var("mask", dtype=dtypes.bool), UPat.var("alt")), allow_any_len=True, name="x"),
    lambda ctx,x,idx,alt,mask:
    f"  br label {ctx[x]}_entry\n{ctx[x][1:]}_entry:\n"
    f"  br i1 {ctx[mask]}, label {ctx[x]}_load, label {ctx[x]}_exit\n{ctx[x][1:]}_load:\n"

tinygrad/renderer/nir.py

@@ -125,19 +125,20 @@ class NIRRenderer(Renderer):
     (UPat.cvar("x", dtypes.uints), lambda x: UOp.const(x.dtype, x.dtype.max+x.arg+1) if x.arg < 0 else None),
     # from ptx
     (UPat.var('x', dtype=dtypes.bool)<UPat.var('y'), lambda x,y: (x^True)&y),
-    # load/store bool -> uint8
+    # load/store bool -> uint8 (alt at src[2] in new gated shape; preserve gate at src[1])
     (UPat(Ops.LOAD, dtypes.bool, name="x"),
-     lambda x: x.replace(dtype=dtypes.uint8, src=x.src[0:1]+((x.src[1].cast(dtypes.uint8),) if len(x.src)>=2 else ())+x.src[2:]).cast(dtypes.bool)),
+     lambda x: x.replace(dtype=dtypes.uint8, src=tuple(s.cast(dtypes.uint8) if i == 2 and s.dtype.scalar() == dtypes.bool else s
+                                                       for i,s in enumerate(x.src))).cast(dtypes.bool)),
     (UPat(Ops.STORE, src=(UPat(), UPat(dtype=dtypes.bool)), name="x", allow_any_len=True),
-     lambda x: x.replace(src=x.src[0:1] + (x.src[1].cast(dtypes.uint8),) + x.src[2:])),
+     lambda x: x.replace(src=(x.src[0], x.src[1].cast(dtypes.uint8))+x.src[2:])),
     # NIR requires shift amount to be 32 bit: https://docs.mesa3d.org/nir/alu.html#nir-alu-op-ishl
     (UPat((Ops.SHL, Ops.SHR), name="x"), lambda x: x.replace(src=(x.src[0], x.src[1].cast(dtypes.uint))) if x.src[1].dtype.bitsize != 32 else None),
     # OpConvertFToU is undefined if Result Type is not wide enough, cast through int32
     # ref: https://registry.khronos.org/SPIR-V/specs/unified1/SPIRV.html#OpConvertFToU
     (UPat(Ops.CAST, (dtypes.uchar, dtypes.ushort), src=(UPat.var("x", dtypes.floats),), name="c"), lambda x,c: x.cast(dtypes.int32).cast(c.dtype)),
     # load/store use pointer arithmetic, and the cast does nothing
-    (UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var("off")), allow_any_len=True, name="x"), lambda x,buf,off: x.replace(
-      src=(buf,off.cast(dtypes.long))+x.src[2:]) if buf.dtype.addrspace != AddrSpace.REG and off.op not in (Ops.CAST, Ops.STACK) else None),
+    (UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var("off")), name="x"), lambda x,buf,off: x.replace(
+      src=(buf,off.cast(dtypes.long))) if buf.dtype.addrspace != AddrSpace.REG and off.op not in (Ops.CAST, Ops.STACK) else None),
     (UPat(Ops.CAST, name="x"), lambda x: x.src[0] if isinstance(x.dtype, PtrDType) or x.src[0].dtype == dtypes.void else None),
   ])
 
@@ -146,9 +147,10 @@ class NIRRenderer(Renderer):
     (UPat(Ops.PARAM, name="x"), lambda ctx,x: ctx.param(ctx.b, x, 8)),
     (UPat(Ops.DEFINE_VAR, name="x"), lambda ctx,x: ctx.param(ctx.b, x, 4)),
     (UPat(Ops.SPECIAL, name="x"), lambda ctx,x: nchannel(ctx.b, {'g':ngid, 'l':nlid, 'i': nid}[x.arg[0]](ctx.b), int(x.arg[-1]))),
-    (UPat(Ops.STORE, src=(UPat(Ops.INDEX, src=(UPat.var("buf"),UPat.var("off")), allow_any_len=True), UPat.var("val")), allow_any_len=True),
+    (UPat(Ops.STORE, src=(UPat(Ops.INDEX, src=(UPat.var("buf"),UPat.var("off"))), UPat.var("val"))),
      lambda ctx,buf,off,val: nstore(ctx.b, buf.ptrdtype.addrspace, nidx(ctx.b, ctx.r[buf], ctx.r[off], buf.dtype), ctx.r[val], val.dtype)),
-    (UPat(Ops.LOAD, src=(UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var("off"), UPat.var("gate"))), UPat.var("alt")), allow_any_len=True, name="x"),
+    (UPat(Ops.LOAD, src=(UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var("off"))), UPat.var("gate", dtype=dtypes.bool), UPat.var("alt")),
+          allow_any_len=True, name="x"),
      lambda ctx,x,buf,off,alt,gate: if_phi(ctx.b, ctx.r[gate],
       lambda: nload(ctx.b, buf.ptrdtype.addrspace, nidx(ctx.b, ctx.r[buf], ctx.r[off], buf.dtype, ctx.r[gate]), x.dtype), lambda: ctx.r[alt])),
     (UPat(Ops.LOAD, src=(UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var("off"))),), allow_any_len=True, name="x"),
@@ -268,9 +270,9 @@ class IR3Renderer(NIRRenderer, OpenCLRenderer):
     return _nload_img(ctx.b, ctx.r[img], ctx.r[coord], img.dtype)
 
   def_rewrite = PatternMatcher([
-    (UPat(Ops.STORE, src=(UPat.var('img').index(UPat.var('coord', dtypes.int.vec(2)), allow_any_len=True), UPat.var("val")),
-          allow_any_len=True), lambda ctx,img,coord,val: nstore_img(ctx.b, ctx.r[img], ctx.r[coord], ctx.r[val], val.dtype)),
-    (UPat(Ops.LOAD, src=(UPat.var('img').index(UPat.var('coord', dtypes.int.vec(2)), UPat.var("gate")), UPat.var("alt"))),
+    (UPat(Ops.STORE, src=(UPat.var('img').index(UPat.var('coord', dtypes.int.vec(2))), UPat.var("val"))),
+     lambda ctx,img,coord,val: nstore_img(ctx.b, ctx.r[img], ctx.r[coord], ctx.r[val], val.dtype)),
+    (UPat(Ops.LOAD, src=(UPat.var('img').index(UPat.var('coord', dtypes.int.vec(2))), UPat.var("gate", dtype=dtypes.bool), UPat.var("alt"))),
      lambda ctx,img,coord,alt,gate: if_phi(ctx.b, ctx.r[gate], lambda: ctx.nload_img(img, coord), lambda: ctx.r[alt])),
     (UPat(Ops.LOAD, src=(UPat.var('img').index(UPat.var('coord', dtypes.int.vec(2))),)), nload_img),
   ]) + NIRRenderer.def_rewrite

tinygrad/renderer/ptx.py

@@ -45,14 +45,15 @@ ptx_matcher = PatternMatcher([
   # upcast to float32 all the ops that don't support half
   (UPat(doesnt_support_half, dtype=dtypes.half, name="x"),
     lambda x: (UOp(x.op, dtypes.float32, tuple(vv.cast(dtypes.float32) for vv in x.src), x.arg).cast(dtypes.half))),
-  # load/store bool -> uint8
+  # load/store bool -> uint8 (alt at src[2] in new gated shape; preserve gate at src[1])
   (UPat(Ops.LOAD, dtypes.bool, src=(UPat(dtype=dtypes.int64),), name="x", allow_any_len=True),
-   lambda x: UOp(x.op, dtypes.uint8, x.src[0:1] + ((x.src[1].cast(dtypes.uint8),) if len(x.src) >= 2 else ()) + x.src[2:]).cast(dtypes.bool)),
+   lambda x: UOp(x.op, dtypes.uint8, tuple(s.cast(dtypes.uint8) if i == 2 and s.dtype.scalar() == dtypes.bool else s
+                                           for i,s in enumerate(x.src))).cast(dtypes.bool)),
   (UPat(Ops.STORE, src=(UPat(dtype=dtypes.int64), UPat(dtype=dtypes.bool)), name="x", allow_any_len=True),
-   lambda x: UOp(x.op, dtypes.void, x.src[0:1] + (x.src[1].cast(dtypes.uint8),) + x.src[2:])),
+   lambda x: UOp(x.op, dtypes.void, (x.src[0], x.src[1].cast(dtypes.uint8))+x.src[2:])),
   # indexing on PTX is in uint64, we do the math while it's still in the graph
-  (UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var("idx")), name="op", allow_any_len=True), lambda buf,idx,op:
-    UOp(Ops.INDEX, dtype=dtypes.int64, src=(buf, buf.cast(dtypes.int64)+idx.cast(dtypes.int64)*buf.dtype.itemsize)+op.src[2:]) \
+  (UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var("idx")), name="op"), lambda buf,idx,op:
+    UOp(Ops.INDEX, dtype=dtypes.int64, src=(buf, buf.cast(dtypes.int64)+idx.cast(dtypes.int64)*buf.dtype.itemsize)) \
       if op.dtype != dtypes.int64 and buf.dtype.addrspace != AddrSpace.REG else None),
   # load/store use pointer arithmetic, and the cast does nothing
   (UPat(Ops.CAST, name="x"), lambda x: x.src[0] if isinstance(x.dtype, PtrDType) or x.src[0].dtype == dtypes.void else None),
@@ -102,18 +103,18 @@ string_rewrite = PatternMatcher([
   (UPat(Ops.CAST, name="x", src=(UPat.var("a"),)),
    lambda ctx, x, a: f"cvt{modifier(x.dtype, a.dtype)}.{ctx.cast_types[x.dtype]}.{ctx.cast_types[a.dtype]} {ctx.r[x]}, {ctx.r[a]};"),
   # store / gated load / load
-  (UPat(Ops.STORE, src=(UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var("loc")), allow_any_len=True), UPat.var("var"))),
+  (UPat(Ops.STORE, src=(UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var("loc"))), UPat.var("var"))),
    lambda ctx, loc, var, buf: f"st.{mem_type(buf)}" + \
     f"{f'.v{cnt}' if ((cnt:=var.dtype.count)>1) else ''}.{ctx.mem_types[var.dtype.scalar()]} " + \
     f"[{ctx.r[loc]}+0], {('{' + ', '.join(ctx.r[var]) + '}') if var.dtype.count > 1 else ctx.r[var]};"),
-  (UPat(Ops.LOAD, name="x", src=(UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var("loc"), UPat.var("gate"))), UPat.var("alt")), allow_any_len=True),
+  (UPat(Ops.LOAD, name="x", src=(UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var("loc"))), UPat.var("gate", dtype=dtypes.bool), UPat.var("alt"))),
     lambda ctx, x, loc, alt, gate, buf: flatten([
     [f"mov.{ctx.mem_types[x.dtype.scalar()]} {v}, {render_val(0, x.dtype.scalar())};" for v in ctx.r[x]],
     [f"@{ctx.r[gate]} ld.{mem_type(buf)}.v{x.dtype.count}.{ctx.mem_types[x.dtype.scalar()]} {{{', '.join(ctx.r[x])}}}, [{ctx.r[loc]}+0];"]
   ]) if alt.dtype.count > 1 else [
     f"@{ctx.r[gate]} ld.{mem_type(buf)}.{ctx.mem_types[x.dtype.scalar()]} {ctx.r[x]}, [{ctx.r[loc]}+0];",
     f"@!{ctx.r[gate]} mov.b{ctx.types[x.dtype.scalar()][1:]} {ctx.r[x]}, {ctx.r[alt]};"]),
-  (UPat(Ops.LOAD, name="x", src=(UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var("loc"))),), allow_any_len=True),
+  (UPat(Ops.LOAD, name="x", src=(UPat(Ops.INDEX, src=(UPat.var("buf"), UPat.var("loc"))),)),
     lambda ctx, x, loc, buf: f"ld.{mem_type(buf)}.v{x.dtype.count}.{ctx.mem_types[x.dtype.scalar()]} {{{', '.join(ctx.r[x])}}}, [{ctx.r[loc]}+0];" \
      if x.dtype.count > 1 else f"ld.{mem_type(buf)}.{ctx.mem_types[x.dtype]} {ctx.r[x]}, [{ctx.r[loc]}+0];"),
   # simple

tinygrad/renderer/wgsl.py

@@ -10,21 +10,27 @@ def sign_extend(val:UOp, sext_am:int):
         | val.bitcast(dtypes.uint32)).bitcast(dtypes.int)
 
 # store for char: buf[idx/4] <- (var << (idx%4)*8))
-def packed_store(bidx:UOp, var:UOp):
+def packed_store(bidx:UOp, var:UOp, *extra:UOp):
   elems, mask = 4//var.dtype.itemsize, _mask(var.dtype)
   shift_am, div_idx = (bidx.src[1].cast(dtypes.uint32) % elems) * (8*var.dtype.itemsize), bidx.src[1] // elems
   new_v, wmask = (var & mask).cast(dtypes.uint32) << shift_am, ((mask << shift_am) ^ 0xFFFFFFFF).cast(dtypes.uint32)
-  # preserve valid condition (bidx.src[2]) if it exists for gated stores
-  idx_src = (bidx.src[0], div_idx) if len(bidx.src) == 2 else (bidx.src[0], div_idx, bidx.src[2])
-  buf = UOp.load(UOp(Ops.INDEX, bidx.dtype, idx_src), dtype=dtypes.uint32)
-  return UOp.store(UOp(Ops.INDEX, bidx.dtype, idx_src), (buf & wmask) | new_v)
+  new_idx = UOp(Ops.INDEX, bidx.dtype, (bidx.src[0], div_idx))
+  buf = UOp.load(new_idx, dtype=dtypes.uint32)
+  # preserve trailing srcs (e.g. gate at src[2] for gated stores)
+  return UOp(Ops.STORE, dtypes.void, (new_idx, (buf & wmask) | new_v) + extra)
 
 # load for char: sign_extend(buf[idx/4] >> ((idx%4)*8))
 def packed_load(root:UOp, bidx:UOp, dtype:DType, var:UOp|None=None):
   elems, mask = 4//dtype.itemsize, _mask(dtype)
   shift_am, div_idx = (bidx.src[1].cast(dtypes.uint32) % elems) * (8*dtype.itemsize), bidx.src[1] // elems
-  idx = UOp(Ops.INDEX, bidx.dtype, (bidx.src[0], div_idx, bidx.src[2]) if var is not None else (bidx.src[0], div_idx))
-  load = UOp.load(idx, *([var] if var is not None else root.src[1:]), dtype=dtypes.uint32, arg=root.arg)
+  new_idx = UOp(Ops.INDEX, bidx.dtype, (bidx.src[0], div_idx))
+  # rebuild LOAD srcs preserving gate at src[1] (if bool) and replacing alt with var if provided
+  other_srcs = list(root.src[1:])
+  if var is not None:
+    alt_pos = 1 if (len(other_srcs) >= 1 and other_srcs[0].dtype.scalar() == dtypes.bool) else 0
+    if alt_pos < len(other_srcs): other_srcs[alt_pos] = var
+    else: other_srcs.append(var)
+  load = UOp.load(new_idx, *other_srcs, dtype=dtypes.uint32, arg=root.arg)
   val = (load.cast(dtypes.uint32) >> shift_am) & mask
   return sign_extend(val, 8*dtype.itemsize).cast(dtype) if dtype in [dtypes.char, dtypes.short] else val.cast(dtype)
 
@@ -40,12 +46,12 @@ def is_nan(a):
 wgsl_matcher = PatternMatcher([
   (UPat((Ops.CMPLT, Ops.XOR), src=(UPat(name="a", dtype=dtypes.bool), UPat.var("b")), name="c"),
    lambda a,b,c: a.cast(dtypes.int).alu(c.op, b.cast(dtypes.int)).cast(dtypes.bool)),
-  # TODO: load alt value doesnt have to be a const
-  (UPat.load(UPat.var("b"), UPat.cvar("c"), allow_any_len=True, name="l"),
-   lambda l,b,c: packed_load(l,b,l.dtype,c.cast(dtypes.uint32)) if is_packed(l.dtype, b.dtype) else None),
+  # TODO: load alt value doesnt have to be a const (alt is at src[2] in gated LOAD)
+  (UPat.load(UPat.var("b"), UPat.var("g", dtype=dtypes.bool), UPat.cvar("c"), name="l"),
+   lambda l,b,g,c: packed_load(l,b,l.dtype,c.cast(dtypes.uint32)) if is_packed(l.dtype, b.dtype) else None),
   (UPat.load(UPat.var("b"), name='l', allow_any_len=True), lambda l,b: packed_load(l, b, l.dtype) if is_packed(l.dtype, b.dtype) else None),
-  (UPat.store(UPat.var("bidx"), UPat.var("var"), allow_any_len=True),
-   lambda bidx,var: packed_store(bidx,var) if is_packed(var.dtype, bidx.dtype) else None),
+  (UPat.store(UPat.var("bidx"), UPat.var("var"), allow_any_len=True, name="sto"),
+   lambda bidx,var,sto: packed_store(bidx,var,*sto.src[2:]) if is_packed(var.dtype, bidx.dtype) else None),
   (UPat.var("a") << UPat.var("b"),lambda a,b:(a.bitcast(dtypes.uint32)<<b.cast(dtypes.uint32)).bitcast(a.dtype) if b.dtype!=dtypes.uint32 else None),
   (UPat.var("x") >> UPat.var("y"), lambda x,y: UOp(Ops.SHR, x.dtype, (x,y.cast(dtypes.uint))) if y.dtype != dtypes.uint else None),
   # fix nan check: 'a != a -> is_nan()'
@@ -81,15 +87,15 @@ class WGSLRenderer(CStyleLanguage):
     (UPat(Ops.BITCAST, dtype=dtypes.short, name="x"), lambda ctx,x: f"bitcast<i32>(vec2<f16>({ctx[x.src[0]]},0))" \
      if x.src[0].dtype == dtypes.half else f"((i32({ctx[x.src[0]]}&0xFFFF)<<16)>>16)"),
     (UPat(Ops.BITCAST, name="x"), lambda ctx,x: f"bitcast<{ctx.type_map[x.dtype]}>({ctx[x.src[0]]})"),
-    # TODO: load alt value doesnt have to be a const
-    (UPat.load(UPat.var("b"), UPat.cvar("v"), allow_any_len=True),
-      lambda ctx,b,v: f"select({ctx[v]}, {ctx.render_load(ctx[b],b.src[0].dtype)}, {ctx[b.src[2]]})"),
+    # TODO: load alt value doesnt have to be a const (gated load: src[1]=gate, src[2]=alt)
+    (UPat.load(UPat.var("b"), UPat.var("g", dtype=dtypes.bool), UPat.cvar("v")),
+      lambda ctx,b,g,v: f"select({ctx[v]}, {ctx.render_load(ctx[b],b.src[0].dtype)}, {ctx[g]})"),
     (UPat.load(UPat.var("b"), allow_any_len=True), lambda ctx, b: ctx.render_load(ctx[b], b.dtype)),
     (UPat.store(UPat.var("b"), UPat.var("v"), allow_any_len=True),lambda ctx,b,v:\
      # (load & mask) | var -> mask = v.src[0].src[1], var = v.src[1]
      f"atomicAnd(&{ctx[b]},{ctx[v.src[0].src[1]]});\n  atomicAdd(&{ctx[b]},{ctx[v.src[1]]});" if is_packed(b.src[0].dtype) \
       else f"{ctx[b]} = {ctx[v]};"),
-    (UPat(Ops.INDEX, src=(UPat.var("b"), UPat.var("idx")), allow_any_len=True),
+    (UPat(Ops.INDEX, src=(UPat.var("b"), UPat.var("idx"))),
      lambda ctx,b,idx: f"{ctx[b]}[{strip_parens(ctx[idx]) if idx.arg is Ops.ADD else ctx[idx]}]"),
   ]) + base_rewrite
 

tinygrad/runtime/autogen/__init__.py

@@ -12,15 +12,17 @@ liburing_src = "https://raw.githubusercontent.com/axboe/liburing/refs/tags/libur
 ggml_common_src = "https://raw.githubusercontent.com/ggml-org/ggml/d4fcfe88a8bcf5c9840be14be6c2fbf1f5b3b2db/src/ggml-common.h"
 macossdk = "/var/db/xcode_select_link/Platforms/MacOSX.platform/Developer/SDKs/MacOSX.sdk"
 
-llvm_lib = (r"'C:\\Program Files\\LLVM\\bin\\LLVM-C.dll' if WIN else '/opt/homebrew/opt/llvm@20/lib/libLLVM.dylib' if OSX else " +
-            repr(['LLVM'] + [f'LLVM-{i}' for i in reversed(range(14, 21+1))]))
-clang_lib = "'/opt/homebrew/opt/llvm@20/lib/libclang.dylib' if OSX else ['clang-20', 'clang']"
+llvm_lib = (
+  (win_llvm:=r"'C:\\Program Files\\LLVM\\bin\\LLVM-C.dll' if WIN else ") +
+  (mac_llvm:=repr([f'/opt/homebrew/opt/llvm@{i}/lib/libLLVM.dylib' for i in reversed(range(14, 21+1))]) + " if OSX else ") +
+  (other_llvm:=repr(['LLVM'] + [f'LLVM-{i}' for i in reversed(range(14, 21+1))])))
+clang_lib = win_llvm.replace("LLVM-C", "libclang") + (mac_llvm + other_llvm).replace("LLVM", "clang")
 
 webgpu_lib = "os.path.join(sysconfig.get_paths()['purelib'], 'pydawn', 'lib', 'libwebgpu_dawn.dll') if WIN else 'webgpu_dawn'"
 nv_lib_path = ("[f'/{pre}/cuda/targets/{tgt}/lib' for pre in ['opt', 'usr/local'] for tgt in "
                "[sysconfig.get_config_vars().get(\"MULTIARCH\", \"\").rsplit(\"-\", 1)[0], 'sbsa-linux']]")
 
-def load(name, dll, files, **kwargs):
+def load(name, files, **kwargs):
   if not (f:=(root/(path:=kwargs.pop("path", __name__)).replace('.','/')/f"{name}.py")).exists() or getenv('REGEN'):
     files, kwargs['args'] = files() if callable(files) else files, args() if callable(args:=kwargs.get('args', [])) else args
     if (srcs:=kwargs.pop('srcs', None)):
@@ -37,23 +39,24 @@ def load(name, dll, files, **kwargs):
       if (preprocess:=kwargs.pop('preprocess', None)): preprocess(srcpath)
     files = flatten(sorted(glob.glob(p, recursive=True)) if isinstance(p, str) and '*' in p else [p] for p in files)
     kwargs['epilog'] = (epi(srcpath) if srcs else epi()) if callable(epi:=kwargs.get('epilog', [])) else epi
-    f.write_text(importlib.import_module("tinygrad.runtime.support.autogen").gen(name, dll, files, **kwargs))
+    try: f.write_text(kwargs.pop("gen", importlib.import_module("tinygrad.runtime.support.autogen").gen)(name, files, **kwargs))
+    except Exception as e: raise RuntimeError(f"error while generating {name}") from e
     if srcs: td.cleanup()
   return importlib.import_module(f"{path}.{name.replace('/', '.')}")
 
 def __getattr__(nm):
   match nm:
-    case "libc": return load("libc", "'c'", lambda: (
+    case "libc": return load("libc", lambda: (
       [i for i in system("dpkg -L libc6-dev").split() if 'sys/mman.h' in i or 'sys/syscall.h' in i] +
-      ["/usr/include/string.h", "/usr/include/elf.h", "/usr/include/unistd.h", "/usr/include/asm-generic/mman-common.h"]), errno=True)
-    case "avcodec": return load("avcodec", None, ["{}/libavcodec/hevc/hevc.h", "{}/libavcodec/cbs_h265.h"], srcs=ffmpeg_src)
-    case "opencl": return load("opencl", "'OpenCL'", ["/usr/include/CL/cl.h"])
-    case "cuda": return load("cuda", "'cuda'", ["/usr/include/cuda.h"], args=["-D__CUDA_API_VERSION_INTERNAL"], parse_macros=False)
-    case "nvrtc": return load("nvrtc", "'nvrtc'", ["/usr/include/nvrtc.h"], paths=nv_lib_path, prolog=["import sysconfig"])
-    case "nvjitlink": load("nvjitlink", "'nvJitLink'", [root/"extra/nvJitLink.h"], paths=nv_lib_path, prolog=["import sysconfig"])
-    case "kfd": return load("kfd", None, [root/"extra/hip_gpu_driver/kfd_ioctl.h"])
+      ["/usr/include/string.h", "/usr/include/elf.h", "/usr/include/unistd.h", "/usr/include/asm-generic/mman-common.h"]), dll="'c'", errno=True)
+    case "avcodec": return load("avcodec", ["{}/libavcodec/hevc/hevc.h", "{}/libavcodec/cbs_h265.h"], srcs=ffmpeg_src)
+    case "opencl": return load("opencl", ["/usr/include/CL/cl.h"], dll="'OpenCL'")
+    case "cuda": return load("cuda", ["/usr/include/cuda.h"], dll="'cuda'", args=["-D__CUDA_API_VERSION_INTERNAL"], macros=False)
+    case "nvrtc": return load("nvrtc", ["/usr/include/nvrtc.h"], dll="'nvrtc'", paths=nv_lib_path, prolog=["import sysconfig"])
+    case "nvjitlink": load("nvjitlink", [root/"extra/nvJitLink.h"], dll="'nvJitLink'", paths=nv_lib_path, prolog=["import sysconfig"])
+    case "kfd": return load("kfd", [root/"extra/hip_gpu_driver/kfd_ioctl.h"])
     case "nv_570" | "nv_580":
-      return load(nm, None, [
+      return load(nm, [
         *[root/"extra/nv_gpu_driver"/s for s in ["clc9b0.h", "clc6c0qmd.h","clcec0qmd.h", "nvdec_drv.h"]], "{}/kernel-open/common/inc/nvmisc.h",
         *[f"{{}}/src/common/sdk/nvidia/inc/class/cl{s}.h" for s in ["0000", "0070", "0080", "2080", "2080_notification", "c56f", "c86f", "c96f", "c761",
                                                                     "83de", "b2cc", "c6c0", "cdc0"]],
@@ -68,7 +71,7 @@ def __getattr__(nm):
         "-include", "{}/src/common/sdk/nvidia/inc/nvtypes.h", "-I{}/src/common/inc", "-I{}/kernel-open/nvidia-uvm", "-I{}/kernel-open/common/inc",
         "-I{}/src/common/sdk/nvidia/inc", "-I{}/src/nvidia/arch/nvalloc/unix/include", "-I{}/src/common/sdk/nvidia/inc/ctrl"
       ], rules=[(r'MW\(([^:]+):(.+)\)',r'(\1, \2)'), (r'(\d+):(\d+)', r'(\1, \2)')], srcs=nv_src[nm], anon_names={"{}/kernel-open/common/inc/nvstatus.h:37":"nv_status_codes"})
-    case "nv": return load("nv", None, [
+    case "nv": return load("nv", [
       *[f"{{}}/src/nvidia/inc/kernel/gpu/{s}.h" for s in ["fsp/kern_fsp_cot_payload", "gsp/gsp_init_args"]],
       *[f"{{}}/src/nvidia/arch/nvalloc/common/inc/{s}.h" for s in ["gsp/gspifpub", "gsp/gsp_fw_wpr_meta", "gsp/gsp_fw_sr_meta", "rmRiscvUcode",
                                                                    "fsp/fsp_nvdm_format"]],
@@ -87,49 +90,50 @@ def __getattr__(nm):
     })
     # this defines all syscall numbers. should probably unify linux autogen?
     case "io_uring":
-      return load("io_uring", None, ["{}/liburing.h", "{}/usr/include/linux/io_uring.h", "{}/usr/include/asm-generic/unistd.h"],
+      return load("io_uring", ["{}/liburing.h", "{}/usr/include/linux/io_uring.h", "{}/usr/include/asm-generic/unistd.h"],
                   args=["-I{}/usr/include"], srcs=[linux_headers_deb, liburing_src], rules=[('__NR', 'NR')],
                   preprocess=lambda path: subprocess.run(f"ar x {linux_headers_deb.split('/')[-1]} && tar xf data.tar.xz", cwd=path, shell=True, check=True))
-    case "ib": return load("ib", "'ibverbs'", ["/usr/include/infiniband/verbs.h", "/usr/include/infiniband/verbs_api.h",
-                                               "/usr/include/infiniband/ib_user_ioctl_verbs.h","/usr/include/rdma/ib_user_verbs.h"], errno=True)
-    case "llvm": return load("llvm", llvm_lib, lambda: [system("llvm-config-20 --includedir")+"/llvm-c/**/*.h"],
+    case "ib": return load("ib", ["/usr/include/infiniband/verbs.h", "/usr/include/infiniband/verbs_api.h",
+                                  "/usr/include/infiniband/ib_user_ioctl_verbs.h", "/usr/include/rdma/ib_user_verbs.h"], dll="'ibverbs'", errno=True)
+    case "llvm": return load("llvm", lambda: [system("llvm-config-20 --includedir")+"/llvm-c/**/*.h"], dll=llvm_lib,
                              args=lambda: system("llvm-config-20 --cflags").split(), recsym=True, prolog=["from tinygrad.helpers import WIN, OSX"])
-    case "pci": return load("pci", None, ["{}/usr/include/linux/pci_regs.h"], srcs=linux_headers_deb,
+    case "pci": return load("pci", ["{}/usr/include/linux/pci_regs.h"], srcs=linux_headers_deb,
                              preprocess=lambda path: subprocess.run(f"ar x {linux_headers_deb.split('/')[-1]} && tar xf data.tar.xz", cwd=path, shell=True, check=True))
-    case "vfio": return load("vfio", None, ["{}/usr/include/linux/vfio.h"], args=["-I{}/usr/include"], srcs=linux_headers_deb,
+    case "vfio": return load("vfio", ["{}/usr/include/linux/vfio.h"], args=["-I{}/usr/include"], srcs=linux_headers_deb,
                              preprocess=lambda path: subprocess.run(f"ar x {linux_headers_deb.split('/')[-1]} && tar xf data.tar.xz", cwd=path, shell=True, check=True))
     # could add rule: WGPU_COMMA -> ','
-    case "webgpu": return load("webgpu", webgpu_lib, [root/"extra/webgpu/webgpu.h"],
+    case "webgpu": return load("webgpu", [root/"extra/webgpu/webgpu.h"], dll=webgpu_lib,
                                prolog=["from tinygrad.helpers import WIN, OSX", "import sysconfig, os"])
-    case "libusb": return load("libusb", "'usb-1.0'", ["/usr/include/libusb-1.0/libusb.h"])
-    case "hip": return load("hip", "os.getenv('ROCM_PATH', '/opt/rocm')+'/lib/libamdhip64.so'", ["/opt/rocm/include/hip/hip_ext.h",
-                            "/opt/rocm/include/hip/hiprtc.h", "/opt/rocm/include/hip/hip_runtime_api.h", "/opt/rocm/include/hip/driver_types.h"],
+    case "libusb": return load("libusb", ["/usr/include/libusb-1.0/libusb.h"], dll="'usb-1.0'")
+    case "hip": return load("hip", ["/opt/rocm/include/hip/hip_ext.h", "/opt/rocm/include/hip/hiprtc.h",
+                                    "/opt/rocm/include/hip/hip_runtime_api.h", "/opt/rocm/include/hip/driver_types.h"],
+                            dll="os.getenv('ROCM_PATH', '/opt/rocm')+'/lib/libamdhip64.so'",
                             args=["-D__HIP_PLATFORM_AMD__", "-I/opt/rocm/include", "-x", "c++"], prolog=["import os"])
     case "comgr" | "comgr_3":
-      return load("comgr_3" if nm == "comgr_3" else "comgr", "[os.getenv('ROCM_PATH', '/opt/rocm')+'/lib/libamd_comgr.so', 'amd_comgr']",
-                  ["/opt/rocm/include/amd_comgr/amd_comgr.h"], args=["-D__HIP_PLATFORM_AMD__", "-I/opt/rocm/include", "-x", "c++"],
-                  prolog=["import os"])
-    case "hsa": return load("hsa", "[os.getenv('ROCM_PATH', '/opt/rocm')+'/lib/libhsa-runtime64.so', 'hsa-runtime64']", [
-      *[f"{{}}/projects/rocr-runtime/runtime/hsa-runtime/core/inc/{s}.h" for s in ["registers"]],
-      *[f"{{}}/projects/rocr-runtime/runtime/hsa-runtime/inc/{s}.h" for s in ["hsa", "hsa_ext_amd", "amd_hsa_signal", "amd_hsa_queue",
-                                                                              "amd_hsa_kernel_code", "hsa_ext_finalize",
-                                                                              "hsa_ext_image", "hsa_ven_amd_aqlprofile"]]],
+      return load("comgr_3" if nm == "comgr_3" else "comgr", ["/opt/rocm/include/amd_comgr/amd_comgr.h"],
+                  dll= "[os.getenv('ROCM_PATH', '/opt/rocm')+'/lib/libamd_comgr.so', 'amd_comgr']",
+                  args=["-D__HIP_PLATFORM_AMD__", "-I/opt/rocm/include", "-x", "c++"], prolog=["import os"])
+    case "hsa": return load("hsa", [*[f"{{}}/projects/rocr-runtime/runtime/hsa-runtime/core/inc/{s}.h" for s in ["registers"]],
+                                    *[f"{{}}/projects/rocr-runtime/runtime/hsa-runtime/inc/{s}.h" for s in [
+                                        "hsa", "hsa_ext_amd", "amd_hsa_signal", "amd_hsa_queue", "amd_hsa_kernel_code",
+                                        "hsa_ext_finalize", "hsa_ext_image", "hsa_ven_amd_aqlprofile"]]],
+      dll="[os.getenv('ROCM_PATH', '/opt/rocm')+'/lib/libhsa-runtime64.so', 'hsa-runtime64']",
       srcs=rocr_src, args=["-DLITTLEENDIAN_CPU"], prolog=["import os"])
-    case "amdgpu_kd": return load("amdgpu_kd", None, lambda: [f"{system('llvm-config-20 --includedir')}/llvm/Support/AMDHSAKernelDescriptor.h"],
-                                  args=lambda: system("llvm-config-20 --cflags").split() + ["-x", "c++"], recsym=True, parse_macros=False)
-    case "amd_gpu": return load("amd_gpu", None, [root/f"extra/hip_gpu_driver/{s}.h" for s in ["sdma_registers", "nvd", "gc_11_0_0_offset",
+    case "amdgpu_kd": return load("amdgpu_kd", lambda: [f"{system('llvm-config-20 --includedir')}/llvm/Support/AMDHSAKernelDescriptor.h"],
+                                  args=lambda: system("llvm-config-20 --cflags").split() + ["-x", "c++"], recsym=True, macros=False)
+    case "amd_gpu": return load("amd_gpu", [root/f"extra/hip_gpu_driver/{s}.h" for s in ["sdma_registers", "nvd", "gc_11_0_0_offset",
                                                                                                "sienna_cichlid_ip_offset"]],
                                 args=["-I/opt/rocm/include", "-x", "c++"])
-    case "amdgpu_drm": return load("amdgpu_drm", None, [ "/usr/include/drm/drm.h", *[root/f"extra/hip_gpu_driver/{s}.h" for s in ["amdgpu_drm"]]])
-    case "kgsl": return load("kgsl", None, [root/"extra/qcom_gpu_driver/msm_kgsl.h"], args=["-D__user="])
+    case "amdgpu_drm": return load("amdgpu_drm", [ "/usr/include/drm/drm.h", *[root/f"extra/hip_gpu_driver/{s}.h" for s in ["amdgpu_drm"]]])
+    case "kgsl": return load("kgsl", [root/"extra/qcom_gpu_driver/msm_kgsl.h"], args=["-D__user="])
     case "qcom_dsp":
-      return load("qcom_dsp", None, [root/f"extra/dsp/include/{s}.h" for s in ["ion", "msm_ion", "adsprpc_shared", "remote_default", "apps_std"]])
-    case "sqtt": return load("sqtt", None, [root/"extra/sqtt/sqtt.h"])
+      return load("qcom_dsp", [root/f"extra/dsp/include/{s}.h" for s in ["ion", "msm_ion", "adsprpc_shared", "remote_default", "apps_std"]])
+    case "sqtt": return load("sqtt", [root/"extra/sqtt/sqtt.h"])
     case "rocprof":
-      return load("rocprof", "['rocprof-trace-decoder', p:='/usr/local/lib/rocprof-trace-decoder.so', p.replace('so','dylib')]",
-                  [f"{{}}/include/{s}.h" for s in ["rocprof_trace_decoder", "trace_decoder_instrument", "trace_decoder_types"]],
+      return load("rocprof", [f"{{}}/include/{s}.h" for s in ["rocprof_trace_decoder", "trace_decoder_instrument", "trace_decoder_types"]],
+                  dll= "['rocprof-trace-decoder', p:='/usr/local/lib/rocprof-trace-decoder.so', p.replace('so','dylib')]",
                   srcs="https://github.com/ROCm/rocprof-trace-decoder/archive/dd0485100971522cc4cd8ae136bdda431061a04d.tar.gz")
-    case "mesa": return load("mesa", "([] if DEV.renderer == 'LVP' else ['tinymesa']) + ['tinymesa_cpu']", [
+    case "mesa": return load("mesa", [
         *[f"{{}}/src/compiler/nir/{s}.h" for s in ["nir", "nir_builder", "nir_shader_compiler_options", "nir_serialize"]], "{}/gen/nir_intrinsics.h",
         *[f"{{}}/src/nouveau/{s}.h" for s in ["headers/nv_device_info", "compiler/nak"]],
         *[f"{{}}/src/gallium/auxiliary/gallivm/lp_bld{s}.h" for s in ["", "_passmgr", "_misc", "_type", "_init", "_nir", "_struct", "_jit_types",
@@ -148,28 +152,28 @@ def __getattr__(nm):
           *[f"python3 src/compiler/{s}_h.py > gen/{s.split('/')[-1]}.h" for s in ["nir/nir_opcodes", "nir/nir_builder_opcodes"]],
           *[f"python3 src/compiler/nir/nir_{s}_h.py --outdir gen" for s in ["intrinsics", "intrinsics_indices"]]]), cwd=path, shell=True, check=True),
   srcs="https://gitlab.freedesktop.org/mesa/mesa/-/archive/mesa-25.2.7/mesa-25.2.7.tar.gz",
+  dll="([] if DEV.renderer == 'LVP' else ['tinymesa']) + ['tinymesa_cpu']",
   prolog=["from tinygrad.helpers import DEV", "import gzip, base64"],
   epilog=lambda path: [system(f"{root}/extra/mesa/lvp_nir_options.sh {path}")])
     case "libclang":
-      return load("libclang", clang_lib,
+      return load("libclang",
                   lambda: [f"{system('llvm-config-20 --includedir')}/clang-c/{s}.h" for s in ["Index", "CXString", "CXSourceLocation", "CXFile"]],
-                  prolog=["from tinygrad.helpers import OSX"], args=lambda: system("llvm-config-20 --cflags").split())
+                  dll=clang_lib, prolog=["from tinygrad.helpers import WIN, OSX"], args=lambda: system("llvm-config-20 --cflags").split())
     case "metal":
-      return load("metal", "'Metal'", [f"{macossdk}/System/Library/Frameworks/Metal.framework/Headers/MTL{s}.h" for s in
+      return load("metal", [f"{macossdk}/System/Library/Frameworks/Metal.framework/Headers/MTL{s}.h" for s in
                   ["ComputeCommandEncoder", "ComputePipeline", "CommandQueue", "Device", "IndirectCommandBuffer", "Resource", "CommandEncoder"]],
-                  args=["-xobjective-c","-isysroot",macossdk], types={"dispatch_data_t":"objc.id_"})
-    case "iokit": return load("iokit", "'IOKit'", [f"{macossdk}/System/Library/Frameworks/IOKit.framework/Headers/IOKitLib.h"],
+                  dll="'Metal'", args=["-xobjective-c","-isysroot",macossdk], types={"dispatch_data_t":"objc.id_"})
+    case "iokit": return load("iokit", [f"{macossdk}/System/Library/Frameworks/IOKit.framework/Headers/IOKitLib.h"], dll="'IOKit'",
                               args=["-isysroot", macossdk])
-    case "corefoundation": return load("corefoundation", "'CoreFoundation'",
+    case "corefoundation": return load("corefoundation",
                                        [f"{macossdk}/System/Library/Frameworks/CoreFoundation.framework/Headers/CF{s}.h" for s in ["String", "Data"]],
-                                       args=["-isysroot", macossdk])
-    case "llvm_qcom": return load("llvm_qcom", "'llvm-qcom'", [root/"extra/tinydreno.h"])
-    case "ggml_common":
-      return load("ggml_common", None, ["{}/ggml-common.h"], srcs=ggml_common_src,
-                  args=["-DGGML_COMMON_DECL_C", "-DGGML_COMMON_IMPL_C"], parse_macros=False)
+                                       dll="'CoreFoundation'",args=["-isysroot", macossdk])
+    case "llvm_qcom": return load("llvm_qcom", [root/"extra/tinydreno.h"], dll="'llvm-qcom'")
+    case "ggml_common": return load("ggml_common", ["{}/ggml-common.h"], srcs=ggml_common_src,
+                                    args=["-DGGML_COMMON_DECL_C", "-DGGML_COMMON_IMPL_C"], macros=False)
     case "mlx5":
       kh = "{}/usr/src/linux-headers-6.18.9+deb14-common/include/linux/mlx5"
-      return load("mlx5", None, [root/"extra/mlx_driver/mlx5.h", f"{kh}/mlx5_ifc.h"], srcs=linux_headers_kern_deb,
+      return load("mlx5", [root/"extra/mlx_driver/mlx5.h", f"{kh}/mlx5_ifc.h"], srcs=linux_headers_kern_deb,
                   args=["-Du8=unsigned char", "-Du16=unsigned short", "-Du32=unsigned int", "-Du64=unsigned long long",
                         "-D__be16=unsigned short", "-D__be32=unsigned int", "-D__be64=unsigned long long", f"-I{kh}"],
                   preprocess=lambda path: subprocess.run(f"ar x {linux_headers_kern_deb.split('/')[-1]} && tar xf data.tar.xz",

tinygrad/runtime/autogen/am/__init__.py

@@ -1,29 +1,39 @@
+import pathlib, hashlib
 from tinygrad.runtime.autogen import load, root
 
 am_src="https://github.com/ROCm/ROCK-Kernel-Driver/archive/33970e1351f5e511029602454979f3de7e22260f.tar.gz"
 AMD, AMDINC = "{}/drivers/gpu/drm/amd", "{}/drivers/gpu/drm/amd/include"
 inc, kern_rules = ["-include", "stdint.h"], [(r'le32_to_cpu', ''),]
+fw_src="https://gitlab.com/kernel-firmware/linux-firmware/-/archive/1e2c15348485939baf1b6d1f5a7a3b799d80703d/1e2c15348485939baf1b6d1f5a7a3b799d80703d.tar.gz"
 
 def __getattr__(nm):
   match nm:
-    case "am": return load("am/am", [], [root/f"extra/amdpci/headers/{s}.h" for s in ["v11_structs", "v12_structs", "amdgpu_vm",
+    case "am": return load("am/am", [root/f"extra/amdpci/headers/{s}.h" for s in ["v11_structs", "v12_structs", "amdgpu_vm",
       "discovery", "amdgpu_ucode", "psp_gfx_if", "amdgpu_psp", "amdgpu_irq", "amdgpu_doorbell"]] + [f"{AMD}/amdkfd/soc15_int.h"] + \
       [f"{AMDINC}/ivsrcid/{s}.h" for s in [f"gfx/irqsrcs_gfx_{x}_0" for x in ('9','11_0','12_0')] + [f"sdma0/irqsrcs_sdma0_{x}_0" for x in (4,5)]] + \
       [f"{AMDINC}/{s}.h" for s in ["v9_structs", "soc15_ih_clientid"]], args=inc, srcs=am_src, rules=kern_rules)
-    case "pm4_soc15": return load("am/pm4_soc15", [], [f"{AMD}/amdkfd/kfd_pm4_headers_ai.h", f"{AMD}/amdgpu/soc15d.h"], srcs=am_src)
-    case "pm4_nv": return load("am/pm4_nv", [], [f"{AMD}/amdkfd/kfd_pm4_headers_ai.h", f"{AMD}/amdgpu/nvd.h"], srcs=am_src)
-    case "sdma_4_0_0": return load("am/sdma_4_0_0", [], [root/"extra/hip_gpu_driver/sdma_registers.h", f"{AMD}/amdgpu/vega10_sdma_pkt_open.h"],
+    case "pm4_soc15": return load("am/pm4_soc15", [f"{AMD}/amdkfd/kfd_pm4_headers_ai.h", f"{AMD}/amdgpu/soc15d.h"], srcs=am_src)
+    case "pm4_nv": return load("am/pm4_nv", [f"{AMD}/amdkfd/kfd_pm4_headers_ai.h", f"{AMD}/amdgpu/nvd.h"], srcs=am_src)
+    case "sdma_4_0_0": return load("am/sdma_4_0_0", [root/"extra/hip_gpu_driver/sdma_registers.h", f"{AMD}/amdgpu/vega10_sdma_pkt_open.h"],
                                    args=["-I/opt/rocm/include", "-x", "c++"], srcs=am_src)
-    case "sdma_5_0_0": return load("am/sdma_5_0_0", [], [root/"extra/hip_gpu_driver/sdma_registers.h", f"{AMD}/amdgpu/navi10_sdma_pkt_open.h"],
+    case "sdma_5_0_0": return load("am/sdma_5_0_0", [root/"extra/hip_gpu_driver/sdma_registers.h", f"{AMD}/amdgpu/navi10_sdma_pkt_open.h"],
                                    args=["-I/opt/rocm/include", "-x", "c++"], srcs=am_src)
-    case "sdma_6_0_0": return load("am/sdma_6_0_0", [], [root/"extra/hip_gpu_driver/sdma_registers.h", f"{AMD}/amdgpu/sdma_v6_0_0_pkt_open.h"],
+    case "sdma_6_0_0": return load("am/sdma_6_0_0", [root/"extra/hip_gpu_driver/sdma_registers.h", f"{AMD}/amdgpu/sdma_v6_0_0_pkt_open.h"],
                                    args=["-I/opt/rocm/include", "-x", "c++"], srcs=am_src)
-    case "smu_v13_0_0": return load("am/smu_v13_0_0",[],[f"{AMD}/pm/swsmu/inc/pmfw_if/{s}.h" for s in ["smu_v13_0_0_ppsmc","smu13_driver_if_v13_0_0"]]
+    case "smu_v13_0_0": return load("am/smu_v13_0_0", [f"{AMD}/pm/swsmu/inc/pmfw_if/{s}.h" for s in ["smu_v13_0_0_ppsmc","smu13_driver_if_v13_0_0"]]
                                     +[root/"extra/amdpci/headers/amdgpu_smu.h"], args=inc, srcs=am_src)
-    case "smu_v13_0_6": return load("am/smu_v13_0_6",[],[f"{AMD}/pm/swsmu/inc/pmfw_if/{s}.h" for s in ["smu_v13_0_6_ppsmc","smu_v13_0_6_pmfw", \
+    case "smu_v13_0_6": return load("am/smu_v13_0_6", [f"{AMD}/pm/swsmu/inc/pmfw_if/{s}.h" for s in ["smu_v13_0_6_ppsmc","smu_v13_0_6_pmfw", \
       "smu13_driver_if_v13_0_6"]] +[root/"extra/amdpci/headers/amdgpu_smu.h"], args=inc, srcs=am_src)
-    case "smu_v13_0_12": return load("am/smu_v13_0_12",[],[f"{AMD}/pm/swsmu/inc/pmfw_if/{s}.h" for s in ["smu_v13_0_12_ppsmc","smu_v13_0_12_pmfw",
+    case "smu_v13_0_12": return load("am/smu_v13_0_12", [f"{AMD}/pm/swsmu/inc/pmfw_if/{s}.h" for s in ["smu_v13_0_12_ppsmc","smu_v13_0_12_pmfw",
       "smu13_driver_if_v13_0_6"]] +[root/"extra/amdpci/headers/amdgpu_smu.h"], args=inc, srcs=am_src)
-    case "smu_v14_0_2": return load("am/smu_v14_0_2", [], [f"{AMD}/pm/swsmu/inc/pmfw_if/{s}.h" for s in ["smu_v14_0_0_pmfw", "smu_v14_0_2_ppsmc",
+    case "smu_v14_0_2": return load("am/smu_v14_0_2", [f"{AMD}/pm/swsmu/inc/pmfw_if/{s}.h" for s in ["smu_v14_0_0_pmfw", "smu_v14_0_2_ppsmc",
                                     "smu14_driver_if_v14_0"]]+[root/"extra/amdpci/headers/amdgpu_smu.h"], args=inc, srcs=am_src)
+    # firmware hashes
+    case "fw":
+      def genfw(name, files, **kwargs): return "\n".join(["hashes = {"] + [f"  {p.name!r}: {hashlib.sha256(p.read_bytes()).hexdigest()!r},"
+                                                                           for f in files if (p:=pathlib.Path(f)).is_file()] + ["}"])
+      return load("am/fw", ["{}/amdgpu/psp_*_sos.bin", "{}/amdgpu/smu_*.bin", "{}/amdgpu/sdma_*.bin"] +
+                           [f"{{}}/amdgpu/gc_*_{x}.bin" for x in ["pfp", "me", "mec", "imu", "rlc"]], srcs=fw_src, gen=genfw)
+    case "navi_offsets": return load("am/navi_offsets", [f"{AMD}/include/sienna_cichlid_ip_offset.h"], srcs=am_src)
+    case "vega_offsets": return load("am/vega_offsets", [f"{AMD}/include/vega20_ip_offset.h"], srcs=am_src)
     case _: raise AttributeError(f"no such autogen: {nm}")

tinygrad/runtime/autogen/am/am.py

@@ -2750,48 +2750,48 @@ class struct_common_firmware_header(c.Struct):
   ucode_size_bytes: int
   ucode_array_offset_bytes: int
   crc32: int
-struct_common_firmware_header.register_fields([('size_bytes', ctypes.c_uint32, 0), ('header_size_bytes', ctypes.c_uint32, 4), ('header_version_major', ctypes.c_uint16, 8), ('header_version_minor', ctypes.c_uint16, 10), ('ip_version_major', ctypes.c_uint16, 12), ('ip_version_minor', ctypes.c_uint16, 14), ('ucode_version', ctypes.c_uint32, 16), ('ucode_size_bytes', ctypes.c_uint32, 20), ('ucode_array_offset_bytes', ctypes.c_uint32, 24), ('crc32', ctypes.c_uint32, 28)])
+struct_common_firmware_header.register_fields([('size_bytes', uint32_t, 0), ('header_size_bytes', uint32_t, 4), ('header_version_major', uint16_t, 8), ('header_version_minor', uint16_t, 10), ('ip_version_major', uint16_t, 12), ('ip_version_minor', uint16_t, 14), ('ucode_version', uint32_t, 16), ('ucode_size_bytes', uint32_t, 20), ('ucode_array_offset_bytes', uint32_t, 24), ('crc32', uint32_t, 28)])
 @c.record
 class struct_mc_firmware_header_v1_0(c.Struct):
   SIZE = 40
   header: struct_common_firmware_header
   io_debug_size_bytes: int
   io_debug_array_offset_bytes: int
-struct_mc_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('io_debug_size_bytes', ctypes.c_uint32, 32), ('io_debug_array_offset_bytes', ctypes.c_uint32, 36)])
+struct_mc_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('io_debug_size_bytes', uint32_t, 32), ('io_debug_array_offset_bytes', uint32_t, 36)])
 @c.record
 class struct_smc_firmware_header_v1_0(c.Struct):
   SIZE = 36
   header: struct_common_firmware_header
   ucode_start_addr: int
-struct_smc_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_start_addr', ctypes.c_uint32, 32)])
+struct_smc_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_start_addr', uint32_t, 32)])
 @c.record
 class struct_smc_firmware_header_v2_0(c.Struct):
   SIZE = 44
   v1_0: struct_smc_firmware_header_v1_0
   ppt_offset_bytes: int
   ppt_size_bytes: int
-struct_smc_firmware_header_v2_0.register_fields([('v1_0', struct_smc_firmware_header_v1_0, 0), ('ppt_offset_bytes', ctypes.c_uint32, 36), ('ppt_size_bytes', ctypes.c_uint32, 40)])
+struct_smc_firmware_header_v2_0.register_fields([('v1_0', struct_smc_firmware_header_v1_0, 0), ('ppt_offset_bytes', uint32_t, 36), ('ppt_size_bytes', uint32_t, 40)])
 @c.record
 class struct_smc_soft_pptable_entry(c.Struct):
   SIZE = 12
   id: int
   ppt_offset_bytes: int
   ppt_size_bytes: int
-struct_smc_soft_pptable_entry.register_fields([('id', ctypes.c_uint32, 0), ('ppt_offset_bytes', ctypes.c_uint32, 4), ('ppt_size_bytes', ctypes.c_uint32, 8)])
+struct_smc_soft_pptable_entry.register_fields([('id', uint32_t, 0), ('ppt_offset_bytes', uint32_t, 4), ('ppt_size_bytes', uint32_t, 8)])
 @c.record
 class struct_smc_firmware_header_v2_1(c.Struct):
   SIZE = 44
   v1_0: struct_smc_firmware_header_v1_0
   pptable_count: int
   pptable_entry_offset: int
-struct_smc_firmware_header_v2_1.register_fields([('v1_0', struct_smc_firmware_header_v1_0, 0), ('pptable_count', ctypes.c_uint32, 36), ('pptable_entry_offset', ctypes.c_uint32, 40)])
+struct_smc_firmware_header_v2_1.register_fields([('v1_0', struct_smc_firmware_header_v1_0, 0), ('pptable_count', uint32_t, 36), ('pptable_entry_offset', uint32_t, 40)])
 @c.record
 class struct_psp_fw_legacy_bin_desc(c.Struct):
   SIZE = 12
   fw_version: int
   offset_bytes: int
   size_bytes: int
-struct_psp_fw_legacy_bin_desc.register_fields([('fw_version', ctypes.c_uint32, 0), ('offset_bytes', ctypes.c_uint32, 4), ('size_bytes', ctypes.c_uint32, 8)])
+struct_psp_fw_legacy_bin_desc.register_fields([('fw_version', uint32_t, 0), ('offset_bytes', uint32_t, 4), ('size_bytes', uint32_t, 8)])
 @c.record
 class struct_psp_firmware_header_v1_0(c.Struct):
   SIZE = 44
@@ -2828,7 +2828,7 @@ class struct_psp_fw_bin_desc(c.Struct):
   fw_version: int
   offset_bytes: int
   size_bytes: int
-struct_psp_fw_bin_desc.register_fields([('fw_type', ctypes.c_uint32, 0), ('fw_version', ctypes.c_uint32, 4), ('offset_bytes', ctypes.c_uint32, 8), ('size_bytes', ctypes.c_uint32, 12)])
+struct_psp_fw_bin_desc.register_fields([('fw_type', uint32_t, 0), ('fw_version', uint32_t, 4), ('offset_bytes', uint32_t, 8), ('size_bytes', uint32_t, 12)])
 enum_psp_fw_type: dict[int, str] = {(PSP_FW_TYPE_UNKOWN:=0): 'PSP_FW_TYPE_UNKOWN', (PSP_FW_TYPE_PSP_SOS:=1): 'PSP_FW_TYPE_PSP_SOS', (PSP_FW_TYPE_PSP_SYS_DRV:=2): 'PSP_FW_TYPE_PSP_SYS_DRV', (PSP_FW_TYPE_PSP_KDB:=3): 'PSP_FW_TYPE_PSP_KDB', (PSP_FW_TYPE_PSP_TOC:=4): 'PSP_FW_TYPE_PSP_TOC', (PSP_FW_TYPE_PSP_SPL:=5): 'PSP_FW_TYPE_PSP_SPL', (PSP_FW_TYPE_PSP_RL:=6): 'PSP_FW_TYPE_PSP_RL', (PSP_FW_TYPE_PSP_SOC_DRV:=7): 'PSP_FW_TYPE_PSP_SOC_DRV', (PSP_FW_TYPE_PSP_INTF_DRV:=8): 'PSP_FW_TYPE_PSP_INTF_DRV', (PSP_FW_TYPE_PSP_DBG_DRV:=9): 'PSP_FW_TYPE_PSP_DBG_DRV', (PSP_FW_TYPE_PSP_RAS_DRV:=10): 'PSP_FW_TYPE_PSP_RAS_DRV', (PSP_FW_TYPE_PSP_IPKEYMGR_DRV:=11): 'PSP_FW_TYPE_PSP_IPKEYMGR_DRV', (PSP_FW_TYPE_MAX_INDEX:=12): 'PSP_FW_TYPE_MAX_INDEX'}
 @c.record
 class struct_psp_firmware_header_v2_0(c.Struct):
@@ -2836,7 +2836,7 @@ class struct_psp_firmware_header_v2_0(c.Struct):
   header: struct_common_firmware_header
   psp_fw_bin_count: int
   psp_fw_bin: c.Array[struct_psp_fw_bin_desc, Literal[1]]
-struct_psp_firmware_header_v2_0.register_fields([('header', struct_common_firmware_header, 0), ('psp_fw_bin_count', ctypes.c_uint32, 32), ('psp_fw_bin', c.Array[struct_psp_fw_bin_desc, Literal[1]], 36)])
+struct_psp_firmware_header_v2_0.register_fields([('header', struct_common_firmware_header, 0), ('psp_fw_bin_count', uint32_t, 32), ('psp_fw_bin', c.Array[struct_psp_fw_bin_desc, Literal[1]], 36)])
 @c.record
 class struct_psp_firmware_header_v2_1(c.Struct):
   SIZE = 56
@@ -2844,7 +2844,7 @@ class struct_psp_firmware_header_v2_1(c.Struct):
   psp_fw_bin_count: int
   psp_aux_fw_bin_index: int
   psp_fw_bin: c.Array[struct_psp_fw_bin_desc, Literal[1]]
-struct_psp_firmware_header_v2_1.register_fields([('header', struct_common_firmware_header, 0), ('psp_fw_bin_count', ctypes.c_uint32, 32), ('psp_aux_fw_bin_index', ctypes.c_uint32, 36), ('psp_fw_bin', c.Array[struct_psp_fw_bin_desc, Literal[1]], 40)])
+struct_psp_firmware_header_v2_1.register_fields([('header', struct_common_firmware_header, 0), ('psp_fw_bin_count', uint32_t, 32), ('psp_aux_fw_bin_index', uint32_t, 36), ('psp_fw_bin', c.Array[struct_psp_fw_bin_desc, Literal[1]], 40)])
 @c.record
 class struct_ta_firmware_header_v1_0(c.Struct):
   SIZE = 92
@@ -2862,7 +2862,7 @@ class struct_ta_firmware_header_v2_0(c.Struct):
   header: struct_common_firmware_header
   ta_fw_bin_count: int
   ta_fw_bin: c.Array[struct_psp_fw_bin_desc, Literal[1]]
-struct_ta_firmware_header_v2_0.register_fields([('header', struct_common_firmware_header, 0), ('ta_fw_bin_count', ctypes.c_uint32, 32), ('ta_fw_bin', c.Array[struct_psp_fw_bin_desc, Literal[1]], 36)])
+struct_ta_firmware_header_v2_0.register_fields([('header', struct_common_firmware_header, 0), ('ta_fw_bin_count', uint32_t, 32), ('ta_fw_bin', c.Array[struct_psp_fw_bin_desc, Literal[1]], 36)])
 @c.record
 class struct_gfx_firmware_header_v1_0(c.Struct):
   SIZE = 44
@@ -2870,7 +2870,7 @@ class struct_gfx_firmware_header_v1_0(c.Struct):
   ucode_feature_version: int
   jt_offset: int
   jt_size: int
-struct_gfx_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_feature_version', ctypes.c_uint32, 32), ('jt_offset', ctypes.c_uint32, 36), ('jt_size', ctypes.c_uint32, 40)])
+struct_gfx_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_feature_version', uint32_t, 32), ('jt_offset', uint32_t, 36), ('jt_size', uint32_t, 40)])
 @c.record
 class struct_gfx_firmware_header_v2_0(c.Struct):
   SIZE = 60
@@ -2882,7 +2882,7 @@ class struct_gfx_firmware_header_v2_0(c.Struct):
   data_offset_bytes: int
   ucode_start_addr_lo: int
   ucode_start_addr_hi: int
-struct_gfx_firmware_header_v2_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_feature_version', ctypes.c_uint32, 32), ('ucode_size_bytes', ctypes.c_uint32, 36), ('ucode_offset_bytes', ctypes.c_uint32, 40), ('data_size_bytes', ctypes.c_uint32, 44), ('data_offset_bytes', ctypes.c_uint32, 48), ('ucode_start_addr_lo', ctypes.c_uint32, 52), ('ucode_start_addr_hi', ctypes.c_uint32, 56)])
+struct_gfx_firmware_header_v2_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_feature_version', uint32_t, 32), ('ucode_size_bytes', uint32_t, 36), ('ucode_offset_bytes', uint32_t, 40), ('data_size_bytes', uint32_t, 44), ('data_offset_bytes', uint32_t, 48), ('ucode_start_addr_lo', uint32_t, 52), ('ucode_start_addr_hi', uint32_t, 56)])
 @c.record
 class struct_mes_firmware_header_v1_0(c.Struct):
   SIZE = 72
@@ -2897,7 +2897,7 @@ class struct_mes_firmware_header_v1_0(c.Struct):
   mes_uc_start_addr_hi: int
   mes_data_start_addr_lo: int
   mes_data_start_addr_hi: int
-struct_mes_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('mes_ucode_version', ctypes.c_uint32, 32), ('mes_ucode_size_bytes', ctypes.c_uint32, 36), ('mes_ucode_offset_bytes', ctypes.c_uint32, 40), ('mes_ucode_data_version', ctypes.c_uint32, 44), ('mes_ucode_data_size_bytes', ctypes.c_uint32, 48), ('mes_ucode_data_offset_bytes', ctypes.c_uint32, 52), ('mes_uc_start_addr_lo', ctypes.c_uint32, 56), ('mes_uc_start_addr_hi', ctypes.c_uint32, 60), ('mes_data_start_addr_lo', ctypes.c_uint32, 64), ('mes_data_start_addr_hi', ctypes.c_uint32, 68)])
+struct_mes_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('mes_ucode_version', uint32_t, 32), ('mes_ucode_size_bytes', uint32_t, 36), ('mes_ucode_offset_bytes', uint32_t, 40), ('mes_ucode_data_version', uint32_t, 44), ('mes_ucode_data_size_bytes', uint32_t, 48), ('mes_ucode_data_offset_bytes', uint32_t, 52), ('mes_uc_start_addr_lo', uint32_t, 56), ('mes_uc_start_addr_hi', uint32_t, 60), ('mes_data_start_addr_lo', uint32_t, 64), ('mes_data_start_addr_hi', uint32_t, 68)])
 @c.record
 class struct_rlc_firmware_header_v1_0(c.Struct):
   SIZE = 52
@@ -2907,7 +2907,7 @@ class struct_rlc_firmware_header_v1_0(c.Struct):
   clear_state_descriptor_offset: int
   avail_scratch_ram_locations: int
   master_pkt_description_offset: int
-struct_rlc_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_feature_version', ctypes.c_uint32, 32), ('save_and_restore_offset', ctypes.c_uint32, 36), ('clear_state_descriptor_offset', ctypes.c_uint32, 40), ('avail_scratch_ram_locations', ctypes.c_uint32, 44), ('master_pkt_description_offset', ctypes.c_uint32, 48)])
+struct_rlc_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_feature_version', uint32_t, 32), ('save_and_restore_offset', uint32_t, 36), ('clear_state_descriptor_offset', uint32_t, 40), ('avail_scratch_ram_locations', uint32_t, 44), ('master_pkt_description_offset', uint32_t, 48)])
 @c.record
 class struct_rlc_firmware_header_v2_0(c.Struct):
   SIZE = 104
@@ -2930,7 +2930,7 @@ class struct_rlc_firmware_header_v2_0(c.Struct):
   reg_list_format_separate_array_offset_bytes: int
   reg_list_separate_size_bytes: int
   reg_list_separate_array_offset_bytes: int
-struct_rlc_firmware_header_v2_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_feature_version', ctypes.c_uint32, 32), ('jt_offset', ctypes.c_uint32, 36), ('jt_size', ctypes.c_uint32, 40), ('save_and_restore_offset', ctypes.c_uint32, 44), ('clear_state_descriptor_offset', ctypes.c_uint32, 48), ('avail_scratch_ram_locations', ctypes.c_uint32, 52), ('reg_restore_list_size', ctypes.c_uint32, 56), ('reg_list_format_start', ctypes.c_uint32, 60), ('reg_list_format_separate_start', ctypes.c_uint32, 64), ('starting_offsets_start', ctypes.c_uint32, 68), ('reg_list_format_size_bytes', ctypes.c_uint32, 72), ('reg_list_format_array_offset_bytes', ctypes.c_uint32, 76), ('reg_list_size_bytes', ctypes.c_uint32, 80), ('reg_list_array_offset_bytes', ctypes.c_uint32, 84), ('reg_list_format_separate_size_bytes', ctypes.c_uint32, 88), ('reg_list_format_separate_array_offset_bytes', ctypes.c_uint32, 92), ('reg_list_separate_size_bytes', ctypes.c_uint32, 96), ('reg_list_separate_array_offset_bytes', ctypes.c_uint32, 100)])
+struct_rlc_firmware_header_v2_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_feature_version', uint32_t, 32), ('jt_offset', uint32_t, 36), ('jt_size', uint32_t, 40), ('save_and_restore_offset', uint32_t, 44), ('clear_state_descriptor_offset', uint32_t, 48), ('avail_scratch_ram_locations', uint32_t, 52), ('reg_restore_list_size', uint32_t, 56), ('reg_list_format_start', uint32_t, 60), ('reg_list_format_separate_start', uint32_t, 64), ('starting_offsets_start', uint32_t, 68), ('reg_list_format_size_bytes', uint32_t, 72), ('reg_list_format_array_offset_bytes', uint32_t, 76), ('reg_list_size_bytes', uint32_t, 80), ('reg_list_array_offset_bytes', uint32_t, 84), ('reg_list_format_separate_size_bytes', uint32_t, 88), ('reg_list_format_separate_array_offset_bytes', uint32_t, 92), ('reg_list_separate_size_bytes', uint32_t, 96), ('reg_list_separate_array_offset_bytes', uint32_t, 100)])
 @c.record
 class struct_rlc_firmware_header_v2_1(c.Struct):
   SIZE = 156
@@ -2948,7 +2948,7 @@ class struct_rlc_firmware_header_v2_1(c.Struct):
   save_restore_list_srm_feature_ver: int
   save_restore_list_srm_size_bytes: int
   save_restore_list_srm_offset_bytes: int
-struct_rlc_firmware_header_v2_1.register_fields([('v2_0', struct_rlc_firmware_header_v2_0, 0), ('reg_list_format_direct_reg_list_length', ctypes.c_uint32, 104), ('save_restore_list_cntl_ucode_ver', ctypes.c_uint32, 108), ('save_restore_list_cntl_feature_ver', ctypes.c_uint32, 112), ('save_restore_list_cntl_size_bytes', ctypes.c_uint32, 116), ('save_restore_list_cntl_offset_bytes', ctypes.c_uint32, 120), ('save_restore_list_gpm_ucode_ver', ctypes.c_uint32, 124), ('save_restore_list_gpm_feature_ver', ctypes.c_uint32, 128), ('save_restore_list_gpm_size_bytes', ctypes.c_uint32, 132), ('save_restore_list_gpm_offset_bytes', ctypes.c_uint32, 136), ('save_restore_list_srm_ucode_ver', ctypes.c_uint32, 140), ('save_restore_list_srm_feature_ver', ctypes.c_uint32, 144), ('save_restore_list_srm_size_bytes', ctypes.c_uint32, 148), ('save_restore_list_srm_offset_bytes', ctypes.c_uint32, 152)])
+struct_rlc_firmware_header_v2_1.register_fields([('v2_0', struct_rlc_firmware_header_v2_0, 0), ('reg_list_format_direct_reg_list_length', uint32_t, 104), ('save_restore_list_cntl_ucode_ver', uint32_t, 108), ('save_restore_list_cntl_feature_ver', uint32_t, 112), ('save_restore_list_cntl_size_bytes', uint32_t, 116), ('save_restore_list_cntl_offset_bytes', uint32_t, 120), ('save_restore_list_gpm_ucode_ver', uint32_t, 124), ('save_restore_list_gpm_feature_ver', uint32_t, 128), ('save_restore_list_gpm_size_bytes', uint32_t, 132), ('save_restore_list_gpm_offset_bytes', uint32_t, 136), ('save_restore_list_srm_ucode_ver', uint32_t, 140), ('save_restore_list_srm_feature_ver', uint32_t, 144), ('save_restore_list_srm_size_bytes', uint32_t, 148), ('save_restore_list_srm_offset_bytes', uint32_t, 152)])
 @c.record
 class struct_rlc_firmware_header_v2_2(c.Struct):
   SIZE = 172
@@ -2957,7 +2957,7 @@ class struct_rlc_firmware_header_v2_2(c.Struct):
   rlc_iram_ucode_offset_bytes: int
   rlc_dram_ucode_size_bytes: int
   rlc_dram_ucode_offset_bytes: int
-struct_rlc_firmware_header_v2_2.register_fields([('v2_1', struct_rlc_firmware_header_v2_1, 0), ('rlc_iram_ucode_size_bytes', ctypes.c_uint32, 156), ('rlc_iram_ucode_offset_bytes', ctypes.c_uint32, 160), ('rlc_dram_ucode_size_bytes', ctypes.c_uint32, 164), ('rlc_dram_ucode_offset_bytes', ctypes.c_uint32, 168)])
+struct_rlc_firmware_header_v2_2.register_fields([('v2_1', struct_rlc_firmware_header_v2_1, 0), ('rlc_iram_ucode_size_bytes', uint32_t, 156), ('rlc_iram_ucode_offset_bytes', uint32_t, 160), ('rlc_dram_ucode_size_bytes', uint32_t, 164), ('rlc_dram_ucode_offset_bytes', uint32_t, 168)])
 @c.record
 class struct_rlc_firmware_header_v2_3(c.Struct):
   SIZE = 204
@@ -2970,7 +2970,7 @@ class struct_rlc_firmware_header_v2_3(c.Struct):
   rlcv_ucode_feature_version: int
   rlcv_ucode_size_bytes: int
   rlcv_ucode_offset_bytes: int
-struct_rlc_firmware_header_v2_3.register_fields([('v2_2', struct_rlc_firmware_header_v2_2, 0), ('rlcp_ucode_version', ctypes.c_uint32, 172), ('rlcp_ucode_feature_version', ctypes.c_uint32, 176), ('rlcp_ucode_size_bytes', ctypes.c_uint32, 180), ('rlcp_ucode_offset_bytes', ctypes.c_uint32, 184), ('rlcv_ucode_version', ctypes.c_uint32, 188), ('rlcv_ucode_feature_version', ctypes.c_uint32, 192), ('rlcv_ucode_size_bytes', ctypes.c_uint32, 196), ('rlcv_ucode_offset_bytes', ctypes.c_uint32, 200)])
+struct_rlc_firmware_header_v2_3.register_fields([('v2_2', struct_rlc_firmware_header_v2_2, 0), ('rlcp_ucode_version', uint32_t, 172), ('rlcp_ucode_feature_version', uint32_t, 176), ('rlcp_ucode_size_bytes', uint32_t, 180), ('rlcp_ucode_offset_bytes', uint32_t, 184), ('rlcv_ucode_version', uint32_t, 188), ('rlcv_ucode_feature_version', uint32_t, 192), ('rlcv_ucode_size_bytes', uint32_t, 196), ('rlcv_ucode_offset_bytes', uint32_t, 200)])
 @c.record
 class struct_rlc_firmware_header_v2_4(c.Struct):
   SIZE = 244
@@ -2985,7 +2985,7 @@ class struct_rlc_firmware_header_v2_4(c.Struct):
   se2_tap_delays_ucode_offset_bytes: int
   se3_tap_delays_ucode_size_bytes: int
   se3_tap_delays_ucode_offset_bytes: int
-struct_rlc_firmware_header_v2_4.register_fields([('v2_3', struct_rlc_firmware_header_v2_3, 0), ('global_tap_delays_ucode_size_bytes', ctypes.c_uint32, 204), ('global_tap_delays_ucode_offset_bytes', ctypes.c_uint32, 208), ('se0_tap_delays_ucode_size_bytes', ctypes.c_uint32, 212), ('se0_tap_delays_ucode_offset_bytes', ctypes.c_uint32, 216), ('se1_tap_delays_ucode_size_bytes', ctypes.c_uint32, 220), ('se1_tap_delays_ucode_offset_bytes', ctypes.c_uint32, 224), ('se2_tap_delays_ucode_size_bytes', ctypes.c_uint32, 228), ('se2_tap_delays_ucode_offset_bytes', ctypes.c_uint32, 232), ('se3_tap_delays_ucode_size_bytes', ctypes.c_uint32, 236), ('se3_tap_delays_ucode_offset_bytes', ctypes.c_uint32, 240)])
+struct_rlc_firmware_header_v2_4.register_fields([('v2_3', struct_rlc_firmware_header_v2_3, 0), ('global_tap_delays_ucode_size_bytes', uint32_t, 204), ('global_tap_delays_ucode_offset_bytes', uint32_t, 208), ('se0_tap_delays_ucode_size_bytes', uint32_t, 212), ('se0_tap_delays_ucode_offset_bytes', uint32_t, 216), ('se1_tap_delays_ucode_size_bytes', uint32_t, 220), ('se1_tap_delays_ucode_offset_bytes', uint32_t, 224), ('se2_tap_delays_ucode_size_bytes', uint32_t, 228), ('se2_tap_delays_ucode_offset_bytes', uint32_t, 232), ('se3_tap_delays_ucode_size_bytes', uint32_t, 236), ('se3_tap_delays_ucode_offset_bytes', uint32_t, 240)])
 @c.record
 class struct_sdma_firmware_header_v1_0(c.Struct):
   SIZE = 48
@@ -2994,13 +2994,13 @@ class struct_sdma_firmware_header_v1_0(c.Struct):
   ucode_change_version: int
   jt_offset: int
   jt_size: int
-struct_sdma_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_feature_version', ctypes.c_uint32, 32), ('ucode_change_version', ctypes.c_uint32, 36), ('jt_offset', ctypes.c_uint32, 40), ('jt_size', ctypes.c_uint32, 44)])
+struct_sdma_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_feature_version', uint32_t, 32), ('ucode_change_version', uint32_t, 36), ('jt_offset', uint32_t, 40), ('jt_size', uint32_t, 44)])
 @c.record
 class struct_sdma_firmware_header_v1_1(c.Struct):
   SIZE = 52
   v1_0: struct_sdma_firmware_header_v1_0
   digest_size: int
-struct_sdma_firmware_header_v1_1.register_fields([('v1_0', struct_sdma_firmware_header_v1_0, 0), ('digest_size', ctypes.c_uint32, 48)])
+struct_sdma_firmware_header_v1_1.register_fields([('v1_0', struct_sdma_firmware_header_v1_0, 0), ('digest_size', uint32_t, 48)])
 @c.record
 class struct_sdma_firmware_header_v2_0(c.Struct):
   SIZE = 64
@@ -3013,7 +3013,7 @@ class struct_sdma_firmware_header_v2_0(c.Struct):
   ctl_ucode_size_bytes: int
   ctl_jt_offset: int
   ctl_jt_size: int
-struct_sdma_firmware_header_v2_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_feature_version', ctypes.c_uint32, 32), ('ctx_ucode_size_bytes', ctypes.c_uint32, 36), ('ctx_jt_offset', ctypes.c_uint32, 40), ('ctx_jt_size', ctypes.c_uint32, 44), ('ctl_ucode_offset', ctypes.c_uint32, 48), ('ctl_ucode_size_bytes', ctypes.c_uint32, 52), ('ctl_jt_offset', ctypes.c_uint32, 56), ('ctl_jt_size', ctypes.c_uint32, 60)])
+struct_sdma_firmware_header_v2_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_feature_version', uint32_t, 32), ('ctx_ucode_size_bytes', uint32_t, 36), ('ctx_jt_offset', uint32_t, 40), ('ctx_jt_size', uint32_t, 44), ('ctl_ucode_offset', uint32_t, 48), ('ctl_ucode_size_bytes', uint32_t, 52), ('ctl_jt_offset', uint32_t, 56), ('ctl_jt_size', uint32_t, 60)])
 @c.record
 class struct_vpe_firmware_header_v1_0(c.Struct):
   SIZE = 64
@@ -3026,7 +3026,7 @@ class struct_vpe_firmware_header_v1_0(c.Struct):
   ctl_ucode_size_bytes: int
   ctl_jt_offset: int
   ctl_jt_size: int
-struct_vpe_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_feature_version', ctypes.c_uint32, 32), ('ctx_ucode_size_bytes', ctypes.c_uint32, 36), ('ctx_jt_offset', ctypes.c_uint32, 40), ('ctx_jt_size', ctypes.c_uint32, 44), ('ctl_ucode_offset', ctypes.c_uint32, 48), ('ctl_ucode_size_bytes', ctypes.c_uint32, 52), ('ctl_jt_offset', ctypes.c_uint32, 56), ('ctl_jt_size', ctypes.c_uint32, 60)])
+struct_vpe_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_feature_version', uint32_t, 32), ('ctx_ucode_size_bytes', uint32_t, 36), ('ctx_jt_offset', uint32_t, 40), ('ctx_jt_size', uint32_t, 44), ('ctl_ucode_offset', uint32_t, 48), ('ctl_ucode_size_bytes', uint32_t, 52), ('ctl_jt_offset', uint32_t, 56), ('ctl_jt_size', uint32_t, 60)])
 @c.record
 class struct_umsch_mm_firmware_header_v1_0(c.Struct):
   SIZE = 80
@@ -3043,7 +3043,7 @@ class struct_umsch_mm_firmware_header_v1_0(c.Struct):
   umsch_mm_uc_start_addr_hi: int
   umsch_mm_data_start_addr_lo: int
   umsch_mm_data_start_addr_hi: int
-struct_umsch_mm_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('umsch_mm_ucode_version', ctypes.c_uint32, 32), ('umsch_mm_ucode_size_bytes', ctypes.c_uint32, 36), ('umsch_mm_ucode_offset_bytes', ctypes.c_uint32, 40), ('umsch_mm_ucode_data_version', ctypes.c_uint32, 44), ('umsch_mm_ucode_data_size_bytes', ctypes.c_uint32, 48), ('umsch_mm_ucode_data_offset_bytes', ctypes.c_uint32, 52), ('umsch_mm_irq_start_addr_lo', ctypes.c_uint32, 56), ('umsch_mm_irq_start_addr_hi', ctypes.c_uint32, 60), ('umsch_mm_uc_start_addr_lo', ctypes.c_uint32, 64), ('umsch_mm_uc_start_addr_hi', ctypes.c_uint32, 68), ('umsch_mm_data_start_addr_lo', ctypes.c_uint32, 72), ('umsch_mm_data_start_addr_hi', ctypes.c_uint32, 76)])
+struct_umsch_mm_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('umsch_mm_ucode_version', uint32_t, 32), ('umsch_mm_ucode_size_bytes', uint32_t, 36), ('umsch_mm_ucode_offset_bytes', uint32_t, 40), ('umsch_mm_ucode_data_version', uint32_t, 44), ('umsch_mm_ucode_data_size_bytes', uint32_t, 48), ('umsch_mm_ucode_data_offset_bytes', uint32_t, 52), ('umsch_mm_irq_start_addr_lo', uint32_t, 56), ('umsch_mm_irq_start_addr_hi', uint32_t, 60), ('umsch_mm_uc_start_addr_lo', uint32_t, 64), ('umsch_mm_uc_start_addr_hi', uint32_t, 68), ('umsch_mm_data_start_addr_lo', uint32_t, 72), ('umsch_mm_data_start_addr_hi', uint32_t, 76)])
 @c.record
 class struct_sdma_firmware_header_v3_0(c.Struct):
   SIZE = 44
@@ -3051,7 +3051,7 @@ class struct_sdma_firmware_header_v3_0(c.Struct):
   ucode_feature_version: int
   ucode_offset_bytes: int
   ucode_size_bytes: int
-struct_sdma_firmware_header_v3_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_feature_version', ctypes.c_uint32, 32), ('ucode_offset_bytes', ctypes.c_uint32, 36), ('ucode_size_bytes', ctypes.c_uint32, 40)])
+struct_sdma_firmware_header_v3_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_feature_version', uint32_t, 32), ('ucode_offset_bytes', uint32_t, 36), ('ucode_size_bytes', uint32_t, 40)])
 @c.record
 class struct_gpu_info_firmware_v1_0(c.Struct):
   SIZE = 60
@@ -3070,35 +3070,35 @@ class struct_gpu_info_firmware_v1_0(c.Struct):
   gc_max_waves_per_simd: int
   gc_max_scratch_slots_per_cu: int
   gc_lds_size: int
-struct_gpu_info_firmware_v1_0.register_fields([('gc_num_se', ctypes.c_uint32, 0), ('gc_num_cu_per_sh', ctypes.c_uint32, 4), ('gc_num_sh_per_se', ctypes.c_uint32, 8), ('gc_num_rb_per_se', ctypes.c_uint32, 12), ('gc_num_tccs', ctypes.c_uint32, 16), ('gc_num_gprs', ctypes.c_uint32, 20), ('gc_num_max_gs_thds', ctypes.c_uint32, 24), ('gc_gs_table_depth', ctypes.c_uint32, 28), ('gc_gsprim_buff_depth', ctypes.c_uint32, 32), ('gc_parameter_cache_depth', ctypes.c_uint32, 36), ('gc_double_offchip_lds_buffer', ctypes.c_uint32, 40), ('gc_wave_size', ctypes.c_uint32, 44), ('gc_max_waves_per_simd', ctypes.c_uint32, 48), ('gc_max_scratch_slots_per_cu', ctypes.c_uint32, 52), ('gc_lds_size', ctypes.c_uint32, 56)])
+struct_gpu_info_firmware_v1_0.register_fields([('gc_num_se', uint32_t, 0), ('gc_num_cu_per_sh', uint32_t, 4), ('gc_num_sh_per_se', uint32_t, 8), ('gc_num_rb_per_se', uint32_t, 12), ('gc_num_tccs', uint32_t, 16), ('gc_num_gprs', uint32_t, 20), ('gc_num_max_gs_thds', uint32_t, 24), ('gc_gs_table_depth', uint32_t, 28), ('gc_gsprim_buff_depth', uint32_t, 32), ('gc_parameter_cache_depth', uint32_t, 36), ('gc_double_offchip_lds_buffer', uint32_t, 40), ('gc_wave_size', uint32_t, 44), ('gc_max_waves_per_simd', uint32_t, 48), ('gc_max_scratch_slots_per_cu', uint32_t, 52), ('gc_lds_size', uint32_t, 56)])
 @c.record
 class struct_gpu_info_firmware_v1_1(c.Struct):
   SIZE = 68
   v1_0: struct_gpu_info_firmware_v1_0
   num_sc_per_sh: int
   num_packer_per_sc: int
-struct_gpu_info_firmware_v1_1.register_fields([('v1_0', struct_gpu_info_firmware_v1_0, 0), ('num_sc_per_sh', ctypes.c_uint32, 60), ('num_packer_per_sc', ctypes.c_uint32, 64)])
+struct_gpu_info_firmware_v1_1.register_fields([('v1_0', struct_gpu_info_firmware_v1_0, 0), ('num_sc_per_sh', uint32_t, 60), ('num_packer_per_sc', uint32_t, 64)])
 @c.record
 class struct_gpu_info_firmware_header_v1_0(c.Struct):
   SIZE = 36
   header: struct_common_firmware_header
   version_major: int
   version_minor: int
-struct_gpu_info_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('version_major', ctypes.c_uint16, 32), ('version_minor', ctypes.c_uint16, 34)])
+struct_gpu_info_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('version_major', uint16_t, 32), ('version_minor', uint16_t, 34)])
 @c.record
 class struct_dmcu_firmware_header_v1_0(c.Struct):
   SIZE = 40
   header: struct_common_firmware_header
   intv_offset_bytes: int
   intv_size_bytes: int
-struct_dmcu_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('intv_offset_bytes', ctypes.c_uint32, 32), ('intv_size_bytes', ctypes.c_uint32, 36)])
+struct_dmcu_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('intv_offset_bytes', uint32_t, 32), ('intv_size_bytes', uint32_t, 36)])
 @c.record
 class struct_dmcub_firmware_header_v1_0(c.Struct):
   SIZE = 40
   header: struct_common_firmware_header
   inst_const_bytes: int
   bss_data_bytes: int
-struct_dmcub_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('inst_const_bytes', ctypes.c_uint32, 32), ('bss_data_bytes', ctypes.c_uint32, 36)])
+struct_dmcub_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('inst_const_bytes', uint32_t, 32), ('bss_data_bytes', uint32_t, 36)])
 @c.record
 class struct_imu_firmware_header_v1_0(c.Struct):
   SIZE = 48
@@ -3107,7 +3107,7 @@ class struct_imu_firmware_header_v1_0(c.Struct):
   imu_iram_ucode_offset_bytes: int
   imu_dram_ucode_size_bytes: int
   imu_dram_ucode_offset_bytes: int
-struct_imu_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('imu_iram_ucode_size_bytes', ctypes.c_uint32, 32), ('imu_iram_ucode_offset_bytes', ctypes.c_uint32, 36), ('imu_dram_ucode_size_bytes', ctypes.c_uint32, 40), ('imu_dram_ucode_offset_bytes', ctypes.c_uint32, 44)])
+struct_imu_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('imu_iram_ucode_size_bytes', uint32_t, 32), ('imu_iram_ucode_offset_bytes', uint32_t, 36), ('imu_dram_ucode_size_bytes', uint32_t, 40), ('imu_dram_ucode_offset_bytes', uint32_t, 44)])
 @c.record
 class union_amdgpu_firmware_header(c.Struct):
   SIZE = 256
@@ -3139,7 +3139,7 @@ class union_amdgpu_firmware_header(c.Struct):
   dmcub: struct_dmcub_firmware_header_v1_0
   imu: struct_imu_firmware_header_v1_0
   raw: c.Array[ctypes.c_ubyte, Literal[256]]
-union_amdgpu_firmware_header.register_fields([('common', struct_common_firmware_header, 0), ('mc', struct_mc_firmware_header_v1_0, 0), ('smc', struct_smc_firmware_header_v1_0, 0), ('smc_v2_0', struct_smc_firmware_header_v2_0, 0), ('psp', struct_psp_firmware_header_v1_0, 0), ('psp_v1_1', struct_psp_firmware_header_v1_1, 0), ('psp_v1_3', struct_psp_firmware_header_v1_3, 0), ('psp_v2_0', struct_psp_firmware_header_v2_0, 0), ('psp_v2_1', struct_psp_firmware_header_v2_0, 0), ('ta', struct_ta_firmware_header_v1_0, 0), ('ta_v2_0', struct_ta_firmware_header_v2_0, 0), ('gfx', struct_gfx_firmware_header_v1_0, 0), ('gfx_v2_0', struct_gfx_firmware_header_v2_0, 0), ('rlc', struct_rlc_firmware_header_v1_0, 0), ('rlc_v2_0', struct_rlc_firmware_header_v2_0, 0), ('rlc_v2_1', struct_rlc_firmware_header_v2_1, 0), ('rlc_v2_2', struct_rlc_firmware_header_v2_2, 0), ('rlc_v2_3', struct_rlc_firmware_header_v2_3, 0), ('rlc_v2_4', struct_rlc_firmware_header_v2_4, 0), ('sdma', struct_sdma_firmware_header_v1_0, 0), ('sdma_v1_1', struct_sdma_firmware_header_v1_1, 0), ('sdma_v2_0', struct_sdma_firmware_header_v2_0, 0), ('sdma_v3_0', struct_sdma_firmware_header_v3_0, 0), ('gpu_info', struct_gpu_info_firmware_header_v1_0, 0), ('dmcu', struct_dmcu_firmware_header_v1_0, 0), ('dmcub', struct_dmcub_firmware_header_v1_0, 0), ('imu', struct_imu_firmware_header_v1_0, 0), ('raw', c.Array[ctypes.c_ubyte, Literal[256]], 0)])
+union_amdgpu_firmware_header.register_fields([('common', struct_common_firmware_header, 0), ('mc', struct_mc_firmware_header_v1_0, 0), ('smc', struct_smc_firmware_header_v1_0, 0), ('smc_v2_0', struct_smc_firmware_header_v2_0, 0), ('psp', struct_psp_firmware_header_v1_0, 0), ('psp_v1_1', struct_psp_firmware_header_v1_1, 0), ('psp_v1_3', struct_psp_firmware_header_v1_3, 0), ('psp_v2_0', struct_psp_firmware_header_v2_0, 0), ('psp_v2_1', struct_psp_firmware_header_v2_0, 0), ('ta', struct_ta_firmware_header_v1_0, 0), ('ta_v2_0', struct_ta_firmware_header_v2_0, 0), ('gfx', struct_gfx_firmware_header_v1_0, 0), ('gfx_v2_0', struct_gfx_firmware_header_v2_0, 0), ('rlc', struct_rlc_firmware_header_v1_0, 0), ('rlc_v2_0', struct_rlc_firmware_header_v2_0, 0), ('rlc_v2_1', struct_rlc_firmware_header_v2_1, 0), ('rlc_v2_2', struct_rlc_firmware_header_v2_2, 0), ('rlc_v2_3', struct_rlc_firmware_header_v2_3, 0), ('rlc_v2_4', struct_rlc_firmware_header_v2_4, 0), ('sdma', struct_sdma_firmware_header_v1_0, 0), ('sdma_v1_1', struct_sdma_firmware_header_v1_1, 0), ('sdma_v2_0', struct_sdma_firmware_header_v2_0, 0), ('sdma_v3_0', struct_sdma_firmware_header_v3_0, 0), ('gpu_info', struct_gpu_info_firmware_header_v1_0, 0), ('dmcu', struct_dmcu_firmware_header_v1_0, 0), ('dmcub', struct_dmcub_firmware_header_v1_0, 0), ('imu', struct_imu_firmware_header_v1_0, 0), ('raw', c.Array[uint8_t, Literal[256]], 0)])
 enum_AMDGPU_UCODE_ID: dict[int, str] = {(AMDGPU_UCODE_ID_CAP:=0): 'AMDGPU_UCODE_ID_CAP', (AMDGPU_UCODE_ID_SDMA0:=1): 'AMDGPU_UCODE_ID_SDMA0', (AMDGPU_UCODE_ID_SDMA1:=2): 'AMDGPU_UCODE_ID_SDMA1', (AMDGPU_UCODE_ID_SDMA2:=3): 'AMDGPU_UCODE_ID_SDMA2', (AMDGPU_UCODE_ID_SDMA3:=4): 'AMDGPU_UCODE_ID_SDMA3', (AMDGPU_UCODE_ID_SDMA4:=5): 'AMDGPU_UCODE_ID_SDMA4', (AMDGPU_UCODE_ID_SDMA5:=6): 'AMDGPU_UCODE_ID_SDMA5', (AMDGPU_UCODE_ID_SDMA6:=7): 'AMDGPU_UCODE_ID_SDMA6', (AMDGPU_UCODE_ID_SDMA7:=8): 'AMDGPU_UCODE_ID_SDMA7', (AMDGPU_UCODE_ID_SDMA_UCODE_TH0:=9): 'AMDGPU_UCODE_ID_SDMA_UCODE_TH0', (AMDGPU_UCODE_ID_SDMA_UCODE_TH1:=10): 'AMDGPU_UCODE_ID_SDMA_UCODE_TH1', (AMDGPU_UCODE_ID_SDMA_RS64:=11): 'AMDGPU_UCODE_ID_SDMA_RS64', (AMDGPU_UCODE_ID_CP_CE:=12): 'AMDGPU_UCODE_ID_CP_CE', (AMDGPU_UCODE_ID_CP_PFP:=13): 'AMDGPU_UCODE_ID_CP_PFP', (AMDGPU_UCODE_ID_CP_ME:=14): 'AMDGPU_UCODE_ID_CP_ME', (AMDGPU_UCODE_ID_CP_RS64_PFP:=15): 'AMDGPU_UCODE_ID_CP_RS64_PFP', (AMDGPU_UCODE_ID_CP_RS64_ME:=16): 'AMDGPU_UCODE_ID_CP_RS64_ME', (AMDGPU_UCODE_ID_CP_RS64_MEC:=17): 'AMDGPU_UCODE_ID_CP_RS64_MEC', (AMDGPU_UCODE_ID_CP_RS64_PFP_P0_STACK:=18): 'AMDGPU_UCODE_ID_CP_RS64_PFP_P0_STACK', (AMDGPU_UCODE_ID_CP_RS64_PFP_P1_STACK:=19): 'AMDGPU_UCODE_ID_CP_RS64_PFP_P1_STACK', (AMDGPU_UCODE_ID_CP_RS64_ME_P0_STACK:=20): 'AMDGPU_UCODE_ID_CP_RS64_ME_P0_STACK', (AMDGPU_UCODE_ID_CP_RS64_ME_P1_STACK:=21): 'AMDGPU_UCODE_ID_CP_RS64_ME_P1_STACK', (AMDGPU_UCODE_ID_CP_RS64_MEC_P0_STACK:=22): 'AMDGPU_UCODE_ID_CP_RS64_MEC_P0_STACK', (AMDGPU_UCODE_ID_CP_RS64_MEC_P1_STACK:=23): 'AMDGPU_UCODE_ID_CP_RS64_MEC_P1_STACK', (AMDGPU_UCODE_ID_CP_RS64_MEC_P2_STACK:=24): 'AMDGPU_UCODE_ID_CP_RS64_MEC_P2_STACK', (AMDGPU_UCODE_ID_CP_RS64_MEC_P3_STACK:=25): 'AMDGPU_UCODE_ID_CP_RS64_MEC_P3_STACK', (AMDGPU_UCODE_ID_CP_MEC1:=26): 'AMDGPU_UCODE_ID_CP_MEC1', (AMDGPU_UCODE_ID_CP_MEC1_JT:=27): 'AMDGPU_UCODE_ID_CP_MEC1_JT', (AMDGPU_UCODE_ID_CP_MEC2:=28): 'AMDGPU_UCODE_ID_CP_MEC2', (AMDGPU_UCODE_ID_CP_MEC2_JT:=29): 'AMDGPU_UCODE_ID_CP_MEC2_JT', (AMDGPU_UCODE_ID_CP_MES:=30): 'AMDGPU_UCODE_ID_CP_MES', (AMDGPU_UCODE_ID_CP_MES_DATA:=31): 'AMDGPU_UCODE_ID_CP_MES_DATA', (AMDGPU_UCODE_ID_CP_MES1:=32): 'AMDGPU_UCODE_ID_CP_MES1', (AMDGPU_UCODE_ID_CP_MES1_DATA:=33): 'AMDGPU_UCODE_ID_CP_MES1_DATA', (AMDGPU_UCODE_ID_IMU_I:=34): 'AMDGPU_UCODE_ID_IMU_I', (AMDGPU_UCODE_ID_IMU_D:=35): 'AMDGPU_UCODE_ID_IMU_D', (AMDGPU_UCODE_ID_GLOBAL_TAP_DELAYS:=36): 'AMDGPU_UCODE_ID_GLOBAL_TAP_DELAYS', (AMDGPU_UCODE_ID_SE0_TAP_DELAYS:=37): 'AMDGPU_UCODE_ID_SE0_TAP_DELAYS', (AMDGPU_UCODE_ID_SE1_TAP_DELAYS:=38): 'AMDGPU_UCODE_ID_SE1_TAP_DELAYS', (AMDGPU_UCODE_ID_SE2_TAP_DELAYS:=39): 'AMDGPU_UCODE_ID_SE2_TAP_DELAYS', (AMDGPU_UCODE_ID_SE3_TAP_DELAYS:=40): 'AMDGPU_UCODE_ID_SE3_TAP_DELAYS', (AMDGPU_UCODE_ID_RLC_RESTORE_LIST_CNTL:=41): 'AMDGPU_UCODE_ID_RLC_RESTORE_LIST_CNTL', (AMDGPU_UCODE_ID_RLC_RESTORE_LIST_GPM_MEM:=42): 'AMDGPU_UCODE_ID_RLC_RESTORE_LIST_GPM_MEM', (AMDGPU_UCODE_ID_RLC_RESTORE_LIST_SRM_MEM:=43): 'AMDGPU_UCODE_ID_RLC_RESTORE_LIST_SRM_MEM', (AMDGPU_UCODE_ID_RLC_IRAM:=44): 'AMDGPU_UCODE_ID_RLC_IRAM', (AMDGPU_UCODE_ID_RLC_DRAM:=45): 'AMDGPU_UCODE_ID_RLC_DRAM', (AMDGPU_UCODE_ID_RLC_P:=46): 'AMDGPU_UCODE_ID_RLC_P', (AMDGPU_UCODE_ID_RLC_V:=47): 'AMDGPU_UCODE_ID_RLC_V', (AMDGPU_UCODE_ID_RLC_G:=48): 'AMDGPU_UCODE_ID_RLC_G', (AMDGPU_UCODE_ID_STORAGE:=49): 'AMDGPU_UCODE_ID_STORAGE', (AMDGPU_UCODE_ID_SMC:=50): 'AMDGPU_UCODE_ID_SMC', (AMDGPU_UCODE_ID_PPTABLE:=51): 'AMDGPU_UCODE_ID_PPTABLE', (AMDGPU_UCODE_ID_UVD:=52): 'AMDGPU_UCODE_ID_UVD', (AMDGPU_UCODE_ID_UVD1:=53): 'AMDGPU_UCODE_ID_UVD1', (AMDGPU_UCODE_ID_VCE:=54): 'AMDGPU_UCODE_ID_VCE', (AMDGPU_UCODE_ID_VCN:=55): 'AMDGPU_UCODE_ID_VCN', (AMDGPU_UCODE_ID_VCN1:=56): 'AMDGPU_UCODE_ID_VCN1', (AMDGPU_UCODE_ID_DMCU_ERAM:=57): 'AMDGPU_UCODE_ID_DMCU_ERAM', (AMDGPU_UCODE_ID_DMCU_INTV:=58): 'AMDGPU_UCODE_ID_DMCU_INTV', (AMDGPU_UCODE_ID_VCN0_RAM:=59): 'AMDGPU_UCODE_ID_VCN0_RAM', (AMDGPU_UCODE_ID_VCN1_RAM:=60): 'AMDGPU_UCODE_ID_VCN1_RAM', (AMDGPU_UCODE_ID_DMCUB:=61): 'AMDGPU_UCODE_ID_DMCUB', (AMDGPU_UCODE_ID_VPE_CTX:=62): 'AMDGPU_UCODE_ID_VPE_CTX', (AMDGPU_UCODE_ID_VPE_CTL:=63): 'AMDGPU_UCODE_ID_VPE_CTL', (AMDGPU_UCODE_ID_VPE:=64): 'AMDGPU_UCODE_ID_VPE', (AMDGPU_UCODE_ID_UMSCH_MM_UCODE:=65): 'AMDGPU_UCODE_ID_UMSCH_MM_UCODE', (AMDGPU_UCODE_ID_UMSCH_MM_DATA:=66): 'AMDGPU_UCODE_ID_UMSCH_MM_DATA', (AMDGPU_UCODE_ID_UMSCH_MM_CMD_BUFFER:=67): 'AMDGPU_UCODE_ID_UMSCH_MM_CMD_BUFFER', (AMDGPU_UCODE_ID_P2S_TABLE:=68): 'AMDGPU_UCODE_ID_P2S_TABLE', (AMDGPU_UCODE_ID_JPEG_RAM:=69): 'AMDGPU_UCODE_ID_JPEG_RAM', (AMDGPU_UCODE_ID_ISP:=70): 'AMDGPU_UCODE_ID_ISP', (AMDGPU_UCODE_ID_MAXIMUM:=71): 'AMDGPU_UCODE_ID_MAXIMUM'}
 enum_AMDGPU_UCODE_STATUS: dict[int, str] = {(AMDGPU_UCODE_STATUS_INVALID:=0): 'AMDGPU_UCODE_STATUS_INVALID', (AMDGPU_UCODE_STATUS_NOT_LOADED:=1): 'AMDGPU_UCODE_STATUS_NOT_LOADED', (AMDGPU_UCODE_STATUS_LOADED:=2): 'AMDGPU_UCODE_STATUS_LOADED'}
 enum_amdgpu_firmware_load_type: dict[int, str] = {(AMDGPU_FW_LOAD_DIRECT:=0): 'AMDGPU_FW_LOAD_DIRECT', (AMDGPU_FW_LOAD_PSP:=1): 'AMDGPU_FW_LOAD_PSP', (AMDGPU_FW_LOAD_SMU:=2): 'AMDGPU_FW_LOAD_SMU', (AMDGPU_FW_LOAD_RLC_BACKDOOR_AUTO:=3): 'AMDGPU_FW_LOAD_RLC_BACKDOOR_AUTO'}
@@ -3154,7 +3154,7 @@ class struct_amdgpu_firmware_info(c.Struct):
   tmr_mc_addr_lo: int
   tmr_mc_addr_hi: int
 class struct_firmware(c.Struct): pass
-struct_amdgpu_firmware_info.register_fields([('ucode_id', ctypes.c_uint32, 0), ('fw', c.POINTER[struct_firmware], 8), ('mc_addr', ctypes.c_uint64, 16), ('kaddr', ctypes.c_void_p, 24), ('ucode_size', ctypes.c_uint32, 32), ('tmr_mc_addr_lo', ctypes.c_uint32, 36), ('tmr_mc_addr_hi', ctypes.c_uint32, 40)])
+struct_amdgpu_firmware_info.register_fields([('ucode_id', ctypes.c_uint32, 0), ('fw', c.POINTER[struct_firmware], 8), ('mc_addr', uint64_t, 16), ('kaddr', ctypes.c_void_p, 24), ('ucode_size', uint32_t, 32), ('tmr_mc_addr_lo', uint32_t, 36), ('tmr_mc_addr_hi', uint32_t, 40)])
 enum_psp_gfx_crtl_cmd_id: dict[int, str] = {(GFX_CTRL_CMD_ID_INIT_RBI_RING:=65536): 'GFX_CTRL_CMD_ID_INIT_RBI_RING', (GFX_CTRL_CMD_ID_INIT_GPCOM_RING:=131072): 'GFX_CTRL_CMD_ID_INIT_GPCOM_RING', (GFX_CTRL_CMD_ID_DESTROY_RINGS:=196608): 'GFX_CTRL_CMD_ID_DESTROY_RINGS', (GFX_CTRL_CMD_ID_CAN_INIT_RINGS:=262144): 'GFX_CTRL_CMD_ID_CAN_INIT_RINGS', (GFX_CTRL_CMD_ID_ENABLE_INT:=327680): 'GFX_CTRL_CMD_ID_ENABLE_INT', (GFX_CTRL_CMD_ID_DISABLE_INT:=393216): 'GFX_CTRL_CMD_ID_DISABLE_INT', (GFX_CTRL_CMD_ID_MODE1_RST:=458752): 'GFX_CTRL_CMD_ID_MODE1_RST', (GFX_CTRL_CMD_ID_GBR_IH_SET:=524288): 'GFX_CTRL_CMD_ID_GBR_IH_SET', (GFX_CTRL_CMD_ID_CONSUME_CMD:=589824): 'GFX_CTRL_CMD_ID_CONSUME_CMD', (GFX_CTRL_CMD_ID_DESTROY_GPCOM_RING:=786432): 'GFX_CTRL_CMD_ID_DESTROY_GPCOM_RING', (GFX_CTRL_CMD_ID_MAX:=983040): 'GFX_CTRL_CMD_ID_MAX'}
 @c.record
 class struct_psp_gfx_ctrl(c.Struct):
@@ -3383,7 +3383,7 @@ class struct_amdgpu_iv_entry(c.Struct):
   node_id: int
   src_data: c.Array[ctypes.c_uint32, Literal[4]]
   iv_entry: c.POINTER[ctypes.c_uint32]
-struct_amdgpu_iv_entry.register_fields([('client_id', ctypes.c_uint32, 0), ('src_id', ctypes.c_uint32, 4), ('ring_id', ctypes.c_uint32, 8), ('vmid', ctypes.c_uint32, 12), ('vmid_src', ctypes.c_uint32, 16), ('timestamp', ctypes.c_uint64, 24), ('timestamp_src', ctypes.c_uint32, 32), ('pasid', ctypes.c_uint32, 36), ('node_id', ctypes.c_uint32, 40), ('src_data', c.Array[ctypes.c_uint32, Literal[4]], 44), ('iv_entry', c.POINTER[ctypes.c_uint32], 64)])
+struct_amdgpu_iv_entry.register_fields([('client_id', ctypes.c_uint32, 0), ('src_id', ctypes.c_uint32, 4), ('ring_id', ctypes.c_uint32, 8), ('vmid', ctypes.c_uint32, 12), ('vmid_src', ctypes.c_uint32, 16), ('timestamp', uint64_t, 24), ('timestamp_src', ctypes.c_uint32, 32), ('pasid', ctypes.c_uint32, 36), ('node_id', ctypes.c_uint32, 40), ('src_data', c.Array[ctypes.c_uint32, Literal[4]], 44), ('iv_entry', c.POINTER[uint32_t], 64)])
 enum_interrupt_node_id_per_aid: dict[int, str] = {(AID0_NODEID:=0): 'AID0_NODEID', (XCD0_NODEID:=1): 'XCD0_NODEID', (XCD1_NODEID:=2): 'XCD1_NODEID', (AID1_NODEID:=4): 'AID1_NODEID', (XCD2_NODEID:=5): 'XCD2_NODEID', (XCD3_NODEID:=6): 'XCD3_NODEID', (AID2_NODEID:=8): 'AID2_NODEID', (XCD4_NODEID:=9): 'XCD4_NODEID', (XCD5_NODEID:=10): 'XCD5_NODEID', (AID3_NODEID:=12): 'AID3_NODEID', (XCD6_NODEID:=13): 'XCD6_NODEID', (XCD7_NODEID:=14): 'XCD7_NODEID', (NODEID_MAX:=15): 'NODEID_MAX'}
 enum_AMDGPU_DOORBELL_ASSIGNMENT: dict[int, str] = {(AMDGPU_DOORBELL_KIQ:=0): 'AMDGPU_DOORBELL_KIQ', (AMDGPU_DOORBELL_HIQ:=1): 'AMDGPU_DOORBELL_HIQ', (AMDGPU_DOORBELL_DIQ:=2): 'AMDGPU_DOORBELL_DIQ', (AMDGPU_DOORBELL_MEC_RING0:=16): 'AMDGPU_DOORBELL_MEC_RING0', (AMDGPU_DOORBELL_MEC_RING1:=17): 'AMDGPU_DOORBELL_MEC_RING1', (AMDGPU_DOORBELL_MEC_RING2:=18): 'AMDGPU_DOORBELL_MEC_RING2', (AMDGPU_DOORBELL_MEC_RING3:=19): 'AMDGPU_DOORBELL_MEC_RING3', (AMDGPU_DOORBELL_MEC_RING4:=20): 'AMDGPU_DOORBELL_MEC_RING4', (AMDGPU_DOORBELL_MEC_RING5:=21): 'AMDGPU_DOORBELL_MEC_RING5', (AMDGPU_DOORBELL_MEC_RING6:=22): 'AMDGPU_DOORBELL_MEC_RING6', (AMDGPU_DOORBELL_MEC_RING7:=23): 'AMDGPU_DOORBELL_MEC_RING7', (AMDGPU_DOORBELL_GFX_RING0:=32): 'AMDGPU_DOORBELL_GFX_RING0', (AMDGPU_DOORBELL_sDMA_ENGINE0:=480): 'AMDGPU_DOORBELL_sDMA_ENGINE0', (AMDGPU_DOORBELL_sDMA_ENGINE1:=481): 'AMDGPU_DOORBELL_sDMA_ENGINE1', (AMDGPU_DOORBELL_IH:=488): 'AMDGPU_DOORBELL_IH', (AMDGPU_DOORBELL_MAX_ASSIGNMENT:=1023): 'AMDGPU_DOORBELL_MAX_ASSIGNMENT', (AMDGPU_DOORBELL_INVALID:=65535): 'AMDGPU_DOORBELL_INVALID'}
 enum_AMDGPU_VEGA20_DOORBELL_ASSIGNMENT: dict[int, str] = {(AMDGPU_VEGA20_DOORBELL_KIQ:=0): 'AMDGPU_VEGA20_DOORBELL_KIQ', (AMDGPU_VEGA20_DOORBELL_HIQ:=1): 'AMDGPU_VEGA20_DOORBELL_HIQ', (AMDGPU_VEGA20_DOORBELL_DIQ:=2): 'AMDGPU_VEGA20_DOORBELL_DIQ', (AMDGPU_VEGA20_DOORBELL_MEC_RING0:=3): 'AMDGPU_VEGA20_DOORBELL_MEC_RING0', (AMDGPU_VEGA20_DOORBELL_MEC_RING1:=4): 'AMDGPU_VEGA20_DOORBELL_MEC_RING1', (AMDGPU_VEGA20_DOORBELL_MEC_RING2:=5): 'AMDGPU_VEGA20_DOORBELL_MEC_RING2', (AMDGPU_VEGA20_DOORBELL_MEC_RING3:=6): 'AMDGPU_VEGA20_DOORBELL_MEC_RING3', (AMDGPU_VEGA20_DOORBELL_MEC_RING4:=7): 'AMDGPU_VEGA20_DOORBELL_MEC_RING4', (AMDGPU_VEGA20_DOORBELL_MEC_RING5:=8): 'AMDGPU_VEGA20_DOORBELL_MEC_RING5', (AMDGPU_VEGA20_DOORBELL_MEC_RING6:=9): 'AMDGPU_VEGA20_DOORBELL_MEC_RING6', (AMDGPU_VEGA20_DOORBELL_MEC_RING7:=10): 'AMDGPU_VEGA20_DOORBELL_MEC_RING7', (AMDGPU_VEGA20_DOORBELL_USERQUEUE_START:=11): 'AMDGPU_VEGA20_DOORBELL_USERQUEUE_START', (AMDGPU_VEGA20_DOORBELL_USERQUEUE_END:=138): 'AMDGPU_VEGA20_DOORBELL_USERQUEUE_END', (AMDGPU_VEGA20_DOORBELL_GFX_RING0:=139): 'AMDGPU_VEGA20_DOORBELL_GFX_RING0', (AMDGPU_VEGA20_DOORBELL_sDMA_ENGINE0:=256): 'AMDGPU_VEGA20_DOORBELL_sDMA_ENGINE0', (AMDGPU_VEGA20_DOORBELL_sDMA_ENGINE1:=266): 'AMDGPU_VEGA20_DOORBELL_sDMA_ENGINE1', (AMDGPU_VEGA20_DOORBELL_sDMA_ENGINE2:=276): 'AMDGPU_VEGA20_DOORBELL_sDMA_ENGINE2', (AMDGPU_VEGA20_DOORBELL_sDMA_ENGINE3:=286): 'AMDGPU_VEGA20_DOORBELL_sDMA_ENGINE3', (AMDGPU_VEGA20_DOORBELL_sDMA_ENGINE4:=296): 'AMDGPU_VEGA20_DOORBELL_sDMA_ENGINE4', (AMDGPU_VEGA20_DOORBELL_sDMA_ENGINE5:=306): 'AMDGPU_VEGA20_DOORBELL_sDMA_ENGINE5', (AMDGPU_VEGA20_DOORBELL_sDMA_ENGINE6:=316): 'AMDGPU_VEGA20_DOORBELL_sDMA_ENGINE6', (AMDGPU_VEGA20_DOORBELL_sDMA_ENGINE7:=326): 'AMDGPU_VEGA20_DOORBELL_sDMA_ENGINE7', (AMDGPU_VEGA20_DOORBELL_IH:=376): 'AMDGPU_VEGA20_DOORBELL_IH', (AMDGPU_VEGA20_DOORBELL64_VCN0_1:=392): 'AMDGPU_VEGA20_DOORBELL64_VCN0_1', (AMDGPU_VEGA20_DOORBELL64_VCN2_3:=393): 'AMDGPU_VEGA20_DOORBELL64_VCN2_3', (AMDGPU_VEGA20_DOORBELL64_VCN4_5:=394): 'AMDGPU_VEGA20_DOORBELL64_VCN4_5', (AMDGPU_VEGA20_DOORBELL64_VCN6_7:=395): 'AMDGPU_VEGA20_DOORBELL64_VCN6_7', (AMDGPU_VEGA20_DOORBELL64_VCN8_9:=396): 'AMDGPU_VEGA20_DOORBELL64_VCN8_9', (AMDGPU_VEGA20_DOORBELL64_VCNa_b:=397): 'AMDGPU_VEGA20_DOORBELL64_VCNa_b', (AMDGPU_VEGA20_DOORBELL64_VCNc_d:=398): 'AMDGPU_VEGA20_DOORBELL64_VCNc_d', (AMDGPU_VEGA20_DOORBELL64_VCNe_f:=399): 'AMDGPU_VEGA20_DOORBELL64_VCNe_f', (AMDGPU_VEGA20_DOORBELL64_UVD_RING0_1:=392): 'AMDGPU_VEGA20_DOORBELL64_UVD_RING0_1', (AMDGPU_VEGA20_DOORBELL64_UVD_RING2_3:=393): 'AMDGPU_VEGA20_DOORBELL64_UVD_RING2_3', (AMDGPU_VEGA20_DOORBELL64_UVD_RING4_5:=394): 'AMDGPU_VEGA20_DOORBELL64_UVD_RING4_5', (AMDGPU_VEGA20_DOORBELL64_UVD_RING6_7:=395): 'AMDGPU_VEGA20_DOORBELL64_UVD_RING6_7', (AMDGPU_VEGA20_DOORBELL64_VCE_RING0_1:=396): 'AMDGPU_VEGA20_DOORBELL64_VCE_RING0_1', (AMDGPU_VEGA20_DOORBELL64_VCE_RING2_3:=397): 'AMDGPU_VEGA20_DOORBELL64_VCE_RING2_3', (AMDGPU_VEGA20_DOORBELL64_VCE_RING4_5:=398): 'AMDGPU_VEGA20_DOORBELL64_VCE_RING4_5', (AMDGPU_VEGA20_DOORBELL64_VCE_RING6_7:=399): 'AMDGPU_VEGA20_DOORBELL64_VCE_RING6_7', (AMDGPU_VEGA20_DOORBELL64_FIRST_NON_CP:=256): 'AMDGPU_VEGA20_DOORBELL64_FIRST_NON_CP', (AMDGPU_VEGA20_DOORBELL64_LAST_NON_CP:=399): 'AMDGPU_VEGA20_DOORBELL64_LAST_NON_CP', (AMDGPU_VEGA20_DOORBELL_XCC1_KIQ_START:=400): 'AMDGPU_VEGA20_DOORBELL_XCC1_KIQ_START', (AMDGPU_VEGA20_DOORBELL_XCC1_MEC_RING0_START:=407): 'AMDGPU_VEGA20_DOORBELL_XCC1_MEC_RING0_START', (AMDGPU_VEGA20_DOORBELL_AID1_sDMA_START:=464): 'AMDGPU_VEGA20_DOORBELL_AID1_sDMA_START', (AMDGPU_VEGA20_DOORBELL_MAX_ASSIGNMENT:=503): 'AMDGPU_VEGA20_DOORBELL_MAX_ASSIGNMENT', (AMDGPU_VEGA20_DOORBELL_INVALID:=65535): 'AMDGPU_VEGA20_DOORBELL_INVALID'}
@@ -4110,201 +4110,199 @@ class struct_v9_gfx_meta_data(c.Struct):
 struct_v9_gfx_meta_data.register_fields([('ce_payload', struct_v9_ce_ib_state, 0), ('reserved1', c.Array[uint32_t, Literal[54]], 40), ('de_payload', struct_v9_de_ib_state, 256), ('DeIbBaseAddrLo', uint32_t, 364), ('DeIbBaseAddrHi', uint32_t, 368), ('reserved2', c.Array[uint32_t, Literal[931]], 372)])
 enum_soc15_ih_clientid: dict[int, str] = {(SOC15_IH_CLIENTID_IH:=0): 'SOC15_IH_CLIENTID_IH', (SOC15_IH_CLIENTID_ACP:=1): 'SOC15_IH_CLIENTID_ACP', (SOC15_IH_CLIENTID_ATHUB:=2): 'SOC15_IH_CLIENTID_ATHUB', (SOC15_IH_CLIENTID_BIF:=3): 'SOC15_IH_CLIENTID_BIF', (SOC15_IH_CLIENTID_DCE:=4): 'SOC15_IH_CLIENTID_DCE', (SOC15_IH_CLIENTID_ISP:=5): 'SOC15_IH_CLIENTID_ISP', (SOC15_IH_CLIENTID_PCIE0:=6): 'SOC15_IH_CLIENTID_PCIE0', (SOC15_IH_CLIENTID_RLC:=7): 'SOC15_IH_CLIENTID_RLC', (SOC15_IH_CLIENTID_SDMA0:=8): 'SOC15_IH_CLIENTID_SDMA0', (SOC15_IH_CLIENTID_SDMA1:=9): 'SOC15_IH_CLIENTID_SDMA1', (SOC15_IH_CLIENTID_SE0SH:=10): 'SOC15_IH_CLIENTID_SE0SH', (SOC15_IH_CLIENTID_SE1SH:=11): 'SOC15_IH_CLIENTID_SE1SH', (SOC15_IH_CLIENTID_SE2SH:=12): 'SOC15_IH_CLIENTID_SE2SH', (SOC15_IH_CLIENTID_SE3SH:=13): 'SOC15_IH_CLIENTID_SE3SH', (SOC15_IH_CLIENTID_UVD1:=14): 'SOC15_IH_CLIENTID_UVD1', (SOC15_IH_CLIENTID_THM:=15): 'SOC15_IH_CLIENTID_THM', (SOC15_IH_CLIENTID_UVD:=16): 'SOC15_IH_CLIENTID_UVD', (SOC15_IH_CLIENTID_VCE0:=17): 'SOC15_IH_CLIENTID_VCE0', (SOC15_IH_CLIENTID_VMC:=18): 'SOC15_IH_CLIENTID_VMC', (SOC15_IH_CLIENTID_XDMA:=19): 'SOC15_IH_CLIENTID_XDMA', (SOC15_IH_CLIENTID_GRBM_CP:=20): 'SOC15_IH_CLIENTID_GRBM_CP', (SOC15_IH_CLIENTID_ATS:=21): 'SOC15_IH_CLIENTID_ATS', (SOC15_IH_CLIENTID_ROM_SMUIO:=22): 'SOC15_IH_CLIENTID_ROM_SMUIO', (SOC15_IH_CLIENTID_DF:=23): 'SOC15_IH_CLIENTID_DF', (SOC15_IH_CLIENTID_VCE1:=24): 'SOC15_IH_CLIENTID_VCE1', (SOC15_IH_CLIENTID_PWR:=25): 'SOC15_IH_CLIENTID_PWR', (SOC15_IH_CLIENTID_RESERVED:=26): 'SOC15_IH_CLIENTID_RESERVED', (SOC15_IH_CLIENTID_UTCL2:=27): 'SOC15_IH_CLIENTID_UTCL2', (SOC15_IH_CLIENTID_EA:=28): 'SOC15_IH_CLIENTID_EA', (SOC15_IH_CLIENTID_UTCL2LOG:=29): 'SOC15_IH_CLIENTID_UTCL2LOG', (SOC15_IH_CLIENTID_MP0:=30): 'SOC15_IH_CLIENTID_MP0', (SOC15_IH_CLIENTID_MP1:=31): 'SOC15_IH_CLIENTID_MP1', (SOC15_IH_CLIENTID_MAX:=32): 'SOC15_IH_CLIENTID_MAX', (SOC15_IH_CLIENTID_VCN:=16): 'SOC15_IH_CLIENTID_VCN', (SOC15_IH_CLIENTID_VCN1:=14): 'SOC15_IH_CLIENTID_VCN1', (SOC15_IH_CLIENTID_SDMA2:=1): 'SOC15_IH_CLIENTID_SDMA2', (SOC15_IH_CLIENTID_SDMA3:=4): 'SOC15_IH_CLIENTID_SDMA3', (SOC15_IH_CLIENTID_SDMA3_Sienna_Cichlid:=5): 'SOC15_IH_CLIENTID_SDMA3_Sienna_Cichlid', (SOC15_IH_CLIENTID_SDMA4:=5): 'SOC15_IH_CLIENTID_SDMA4', (SOC15_IH_CLIENTID_SDMA5:=17): 'SOC15_IH_CLIENTID_SDMA5', (SOC15_IH_CLIENTID_SDMA6:=19): 'SOC15_IH_CLIENTID_SDMA6', (SOC15_IH_CLIENTID_SDMA7:=24): 'SOC15_IH_CLIENTID_SDMA7', (SOC15_IH_CLIENTID_VMC1:=6): 'SOC15_IH_CLIENTID_VMC1'}
 enum_soc21_ih_clientid: dict[int, str] = {(SOC21_IH_CLIENTID_IH:=0): 'SOC21_IH_CLIENTID_IH', (SOC21_IH_CLIENTID_ATHUB:=2): 'SOC21_IH_CLIENTID_ATHUB', (SOC21_IH_CLIENTID_BIF:=3): 'SOC21_IH_CLIENTID_BIF', (SOC21_IH_CLIENTID_DCN:=4): 'SOC21_IH_CLIENTID_DCN', (SOC21_IH_CLIENTID_ISP:=5): 'SOC21_IH_CLIENTID_ISP', (SOC21_IH_CLIENTID_MP3:=6): 'SOC21_IH_CLIENTID_MP3', (SOC21_IH_CLIENTID_RLC:=7): 'SOC21_IH_CLIENTID_RLC', (SOC21_IH_CLIENTID_GFX:=10): 'SOC21_IH_CLIENTID_GFX', (SOC21_IH_CLIENTID_IMU:=11): 'SOC21_IH_CLIENTID_IMU', (SOC21_IH_CLIENTID_VCN1:=14): 'SOC21_IH_CLIENTID_VCN1', (SOC21_IH_CLIENTID_THM:=15): 'SOC21_IH_CLIENTID_THM', (SOC21_IH_CLIENTID_VCN:=16): 'SOC21_IH_CLIENTID_VCN', (SOC21_IH_CLIENTID_VPE1:=17): 'SOC21_IH_CLIENTID_VPE1', (SOC21_IH_CLIENTID_VMC:=18): 'SOC21_IH_CLIENTID_VMC', (SOC21_IH_CLIENTID_GRBM_CP:=20): 'SOC21_IH_CLIENTID_GRBM_CP', (SOC21_IH_CLIENTID_ROM_SMUIO:=22): 'SOC21_IH_CLIENTID_ROM_SMUIO', (SOC21_IH_CLIENTID_DF:=23): 'SOC21_IH_CLIENTID_DF', (SOC21_IH_CLIENTID_VPE:=24): 'SOC21_IH_CLIENTID_VPE', (SOC21_IH_CLIENTID_PWR:=25): 'SOC21_IH_CLIENTID_PWR', (SOC21_IH_CLIENTID_LSDMA:=26): 'SOC21_IH_CLIENTID_LSDMA', (SOC21_IH_CLIENTID_MP0:=30): 'SOC21_IH_CLIENTID_MP0', (SOC21_IH_CLIENTID_MP1:=31): 'SOC21_IH_CLIENTID_MP1', (SOC21_IH_CLIENTID_MAX:=32): 'SOC21_IH_CLIENTID_MAX'}
-AMDGPU_VM_MAX_UPDATE_SIZE = 0x3FFFF # type: ignore
-AMDGPU_PTE_VALID = (1 << 0) # type: ignore
-AMDGPU_PTE_SYSTEM = (1 << 1) # type: ignore
-AMDGPU_PTE_SNOOPED = (1 << 2) # type: ignore
-AMDGPU_PTE_TMZ = (1 << 3) # type: ignore
-AMDGPU_PTE_EXECUTABLE = (1 << 4) # type: ignore
-AMDGPU_PTE_READABLE = (1 << 5) # type: ignore
-AMDGPU_PTE_WRITEABLE = (1 << 6) # type: ignore
+AMDGPU_VM_MAX_UPDATE_SIZE = 0x3FFFF
+AMDGPU_PTE_VALID = (1 << 0)
+AMDGPU_PTE_SYSTEM = (1 << 1)
+AMDGPU_PTE_SNOOPED = (1 << 2)
+AMDGPU_PTE_TMZ = (1 << 3)
+AMDGPU_PTE_EXECUTABLE = (1 << 4)
+AMDGPU_PTE_READABLE = (1 << 5)
+AMDGPU_PTE_WRITEABLE = (1 << 6)
 AMDGPU_PTE_FRAG = lambda x: ((x & 0x1f) << 7) # type: ignore
-AMDGPU_PTE_PRT = (1 << 51) # type: ignore
-AMDGPU_PDE_PTE = (1 << 54) # type: ignore
-AMDGPU_PTE_LOG = (1 << 55) # type: ignore
-AMDGPU_PTE_TF = (1 << 56) # type: ignore
-AMDGPU_PTE_NOALLOC = (1 << 58) # type: ignore
+AMDGPU_PTE_PRT = (1 << 51)
+AMDGPU_PDE_PTE = (1 << 54)
+AMDGPU_PTE_LOG = (1 << 55)
+AMDGPU_PTE_TF = (1 << 56)
+AMDGPU_PTE_NOALLOC = (1 << 58)
 AMDGPU_PDE_BFS = lambda a: (a << 59) # type: ignore
-AMDGPU_VM_NORETRY_FLAGS = (AMDGPU_PTE_EXECUTABLE | AMDGPU_PDE_PTE | AMDGPU_PTE_TF) # type: ignore
-AMDGPU_VM_NORETRY_FLAGS_TF = (AMDGPU_PTE_VALID | AMDGPU_PTE_SYSTEM | AMDGPU_PTE_PRT) # type: ignore
+AMDGPU_VM_NORETRY_FLAGS = (AMDGPU_PTE_EXECUTABLE | AMDGPU_PDE_PTE | AMDGPU_PTE_TF)
+AMDGPU_VM_NORETRY_FLAGS_TF = (AMDGPU_PTE_VALID | AMDGPU_PTE_SYSTEM | AMDGPU_PTE_PRT)
 AMDGPU_PTE_MTYPE_VG10_SHIFT = lambda mtype: ((mtype) << 57) # type: ignore
-AMDGPU_PTE_MTYPE_VG10_MASK = AMDGPU_PTE_MTYPE_VG10_SHIFT(3) # type: ignore
+AMDGPU_PTE_MTYPE_VG10_MASK = AMDGPU_PTE_MTYPE_VG10_SHIFT(3)
 AMDGPU_PTE_MTYPE_VG10 = lambda flags,mtype: (((flags) & (~AMDGPU_PTE_MTYPE_VG10_MASK)) | AMDGPU_PTE_MTYPE_VG10_SHIFT(mtype)) # type: ignore
-AMDGPU_MTYPE_NC = 0 # type: ignore
-AMDGPU_MTYPE_CC = 2 # type: ignore
+AMDGPU_MTYPE_NC = 0
+AMDGPU_MTYPE_CC = 2
 AMDGPU_PTE_MTYPE_NV10_SHIFT = lambda mtype: ((mtype) << 48) # type: ignore
-AMDGPU_PTE_MTYPE_NV10_MASK = AMDGPU_PTE_MTYPE_NV10_SHIFT(7) # type: ignore
+AMDGPU_PTE_MTYPE_NV10_MASK = AMDGPU_PTE_MTYPE_NV10_SHIFT(7)
 AMDGPU_PTE_MTYPE_NV10 = lambda flags,mtype: (((flags) & (~AMDGPU_PTE_MTYPE_NV10_MASK)) | AMDGPU_PTE_MTYPE_NV10_SHIFT(mtype)) # type: ignore
-AMDGPU_PTE_PRT_GFX12 = (1 << 56) # type: ignore
+AMDGPU_PTE_PRT_GFX12 = (1 << 56)
 AMDGPU_PTE_MTYPE_GFX12_SHIFT = lambda mtype: ((mtype) << 54) # type: ignore
-AMDGPU_PTE_MTYPE_GFX12_MASK = AMDGPU_PTE_MTYPE_GFX12_SHIFT(3) # type: ignore
+AMDGPU_PTE_MTYPE_GFX12_MASK = AMDGPU_PTE_MTYPE_GFX12_SHIFT(3)
 AMDGPU_PTE_MTYPE_GFX12 = lambda flags,mtype: (((flags) & (~AMDGPU_PTE_MTYPE_GFX12_MASK)) | AMDGPU_PTE_MTYPE_GFX12_SHIFT(mtype)) # type: ignore
-AMDGPU_PTE_IS_PTE = (1 << 63) # type: ignore
+AMDGPU_PTE_IS_PTE = (1 << 63)
 AMDGPU_PDE_BFS_GFX12 = lambda a: (((a) & 0x1f) << 58) # type: ignore
-AMDGPU_PDE_PTE_GFX12 = (1 << 63) # type: ignore
-AMDGPU_VM_FAULT_STOP_NEVER = 0 # type: ignore
-AMDGPU_VM_FAULT_STOP_FIRST = 1 # type: ignore
-AMDGPU_VM_FAULT_STOP_ALWAYS = 2 # type: ignore
-AMDGPU_VM_RESERVED_VRAM = (8 << 20) # type: ignore
-AMDGPU_MAX_VMHUBS = 13 # type: ignore
-AMDGPU_GFXHUB_START = 0 # type: ignore
-AMDGPU_MMHUB0_START = 8 # type: ignore
-AMDGPU_MMHUB1_START = 12 # type: ignore
+AMDGPU_PDE_PTE_GFX12 = (1 << 63)
+AMDGPU_VM_FAULT_STOP_NEVER = 0
+AMDGPU_VM_FAULT_STOP_FIRST = 1
+AMDGPU_VM_FAULT_STOP_ALWAYS = 2
+AMDGPU_VM_RESERVED_VRAM = (8 << 20)
+AMDGPU_MAX_VMHUBS = 13
+AMDGPU_GFXHUB_START = 0
+AMDGPU_MMHUB0_START = 8
+AMDGPU_MMHUB1_START = 12
 AMDGPU_GFXHUB = lambda x: (AMDGPU_GFXHUB_START + (x)) # type: ignore
 AMDGPU_MMHUB0 = lambda x: (AMDGPU_MMHUB0_START + (x)) # type: ignore
 AMDGPU_MMHUB1 = lambda x: (AMDGPU_MMHUB1_START + (x)) # type: ignore
 AMDGPU_IS_GFXHUB = lambda x: ((x) >= AMDGPU_GFXHUB_START and (x) < AMDGPU_MMHUB0_START) # type: ignore
 AMDGPU_IS_MMHUB0 = lambda x: ((x) >= AMDGPU_MMHUB0_START and (x) < AMDGPU_MMHUB1_START) # type: ignore
 AMDGPU_IS_MMHUB1 = lambda x: ((x) >= AMDGPU_MMHUB1_START and (x) < AMDGPU_MAX_VMHUBS) # type: ignore
-AMDGPU_VA_RESERVED_CSA_SIZE = (2 << 20) # type: ignore
-AMDGPU_VA_RESERVED_SEQ64_SIZE = (2 << 20) # type: ignore
+AMDGPU_VA_RESERVED_CSA_SIZE = (2 << 20)
+AMDGPU_VA_RESERVED_SEQ64_SIZE = (2 << 20)
 AMDGPU_VA_RESERVED_SEQ64_START = lambda adev: (AMDGPU_VA_RESERVED_CSA_START(adev) - AMDGPU_VA_RESERVED_SEQ64_SIZE) # type: ignore
-AMDGPU_VA_RESERVED_TRAP_SIZE = (2 << 12) # type: ignore
+AMDGPU_VA_RESERVED_TRAP_SIZE = (2 << 12)
 AMDGPU_VA_RESERVED_TRAP_START = lambda adev: (AMDGPU_VA_RESERVED_SEQ64_START(adev) - AMDGPU_VA_RESERVED_TRAP_SIZE) # type: ignore
-AMDGPU_VA_RESERVED_BOTTOM = (1 << 16) # type: ignore
-AMDGPU_VA_RESERVED_TOP = (AMDGPU_VA_RESERVED_TRAP_SIZE + AMDGPU_VA_RESERVED_SEQ64_SIZE + AMDGPU_VA_RESERVED_CSA_SIZE) # type: ignore
-AMDGPU_VM_USE_CPU_FOR_GFX = (1 << 0) # type: ignore
-AMDGPU_VM_USE_CPU_FOR_COMPUTE = (1 << 1) # type: ignore
-PSP_HEADER_SIZE = 256 # type: ignore
-BINARY_SIGNATURE = 0x28211407 # type: ignore
-DISCOVERY_TABLE_SIGNATURE = 0x53445049 # type: ignore
-GC_TABLE_ID = 0x4347 # type: ignore
-HARVEST_TABLE_SIGNATURE = 0x56524148 # type: ignore
-VCN_INFO_TABLE_ID = 0x004E4356 # type: ignore
-MALL_INFO_TABLE_ID = 0x4C4C414D # type: ignore
-NPS_INFO_TABLE_ID = 0x0053504E # type: ignore
-VCN_INFO_TABLE_MAX_NUM_INSTANCES = 4 # type: ignore
-NPS_INFO_TABLE_MAX_NUM_INSTANCES = 12 # type: ignore
-HWIP_MAX_INSTANCE = 44 # type: ignore
-HW_ID_MAX = 300 # type: ignore
-MP1_HWID = 1 # type: ignore
-MP2_HWID = 2 # type: ignore
-THM_HWID = 3 # type: ignore
-SMUIO_HWID = 4 # type: ignore
-FUSE_HWID = 5 # type: ignore
-CLKA_HWID = 6 # type: ignore
-PWR_HWID = 10 # type: ignore
-GC_HWID = 11 # type: ignore
-UVD_HWID = 12 # type: ignore
-VCN_HWID = UVD_HWID # type: ignore
-AUDIO_AZ_HWID = 13 # type: ignore
-ACP_HWID = 14 # type: ignore
-DCI_HWID = 15 # type: ignore
-DMU_HWID = 271 # type: ignore
-DCO_HWID = 16 # type: ignore
-DIO_HWID = 272 # type: ignore
-XDMA_HWID = 17 # type: ignore
-DCEAZ_HWID = 18 # type: ignore
-DAZ_HWID = 274 # type: ignore
-SDPMUX_HWID = 19 # type: ignore
-NTB_HWID = 20 # type: ignore
-VPE_HWID = 21 # type: ignore
-IOHC_HWID = 24 # type: ignore
-L2IMU_HWID = 28 # type: ignore
-VCE_HWID = 32 # type: ignore
-MMHUB_HWID = 34 # type: ignore
-ATHUB_HWID = 35 # type: ignore
-DBGU_NBIO_HWID = 36 # type: ignore
-DFX_HWID = 37 # type: ignore
-DBGU0_HWID = 38 # type: ignore
-DBGU1_HWID = 39 # type: ignore
-OSSSYS_HWID = 40 # type: ignore
-HDP_HWID = 41 # type: ignore
-SDMA0_HWID = 42 # type: ignore
-SDMA1_HWID = 43 # type: ignore
-ISP_HWID = 44 # type: ignore
-DBGU_IO_HWID = 45 # type: ignore
-DF_HWID = 46 # type: ignore
-CLKB_HWID = 47 # type: ignore
-FCH_HWID = 48 # type: ignore
-DFX_DAP_HWID = 49 # type: ignore
-L1IMU_PCIE_HWID = 50 # type: ignore
-L1IMU_NBIF_HWID = 51 # type: ignore
-L1IMU_IOAGR_HWID = 52 # type: ignore
-L1IMU3_HWID = 53 # type: ignore
-L1IMU4_HWID = 54 # type: ignore
-L1IMU5_HWID = 55 # type: ignore
-L1IMU6_HWID = 56 # type: ignore
-L1IMU7_HWID = 57 # type: ignore
-L1IMU8_HWID = 58 # type: ignore
-L1IMU9_HWID = 59 # type: ignore
-L1IMU10_HWID = 60 # type: ignore
-L1IMU11_HWID = 61 # type: ignore
-L1IMU12_HWID = 62 # type: ignore
-L1IMU13_HWID = 63 # type: ignore
-L1IMU14_HWID = 64 # type: ignore
-L1IMU15_HWID = 65 # type: ignore
-WAFLC_HWID = 66 # type: ignore
-FCH_USB_PD_HWID = 67 # type: ignore
-SDMA2_HWID = 68 # type: ignore
-SDMA3_HWID = 69 # type: ignore
-PCIE_HWID = 70 # type: ignore
-PCS_HWID = 80 # type: ignore
-DDCL_HWID = 89 # type: ignore
-SST_HWID = 90 # type: ignore
-LSDMA_HWID = 91 # type: ignore
-IOAGR_HWID = 100 # type: ignore
-NBIF_HWID = 108 # type: ignore
-IOAPIC_HWID = 124 # type: ignore
-SYSTEMHUB_HWID = 128 # type: ignore
-NTBCCP_HWID = 144 # type: ignore
-UMC_HWID = 150 # type: ignore
-SATA_HWID = 168 # type: ignore
-USB_HWID = 170 # type: ignore
-CCXSEC_HWID = 176 # type: ignore
-XGMI_HWID = 200 # type: ignore
-XGBE_HWID = 216 # type: ignore
-MP0_HWID = 255 # type: ignore
-hw_id_map = {GC_HWIP:GC_HWID,HDP_HWIP:HDP_HWID,SDMA0_HWIP:SDMA0_HWID,SDMA1_HWIP:SDMA1_HWID,SDMA2_HWIP:SDMA2_HWID,SDMA3_HWIP:SDMA3_HWID,LSDMA_HWIP:LSDMA_HWID,MMHUB_HWIP:MMHUB_HWID,ATHUB_HWIP:ATHUB_HWID,NBIO_HWIP:NBIF_HWID,MP0_HWIP:MP0_HWID,MP1_HWIP:MP1_HWID,UVD_HWIP:UVD_HWID,VCE_HWIP:VCE_HWID,DF_HWIP:DF_HWID,DCE_HWIP:DMU_HWID,OSSSYS_HWIP:OSSSYS_HWID,SMUIO_HWIP:SMUIO_HWID,PWR_HWIP:PWR_HWID,NBIF_HWIP:NBIF_HWID,THM_HWIP:THM_HWID,CLK_HWIP:CLKA_HWID,UMC_HWIP:UMC_HWID,XGMI_HWIP:XGMI_HWID,DCI_HWIP:DCI_HWID,PCIE_HWIP:PCIE_HWID,VPE_HWIP:VPE_HWID,ISP_HWIP:ISP_HWID} # type: ignore
-int32_t = int # type: ignore
-AMDGPU_SDMA0_UCODE_LOADED = 0x00000001 # type: ignore
-AMDGPU_SDMA1_UCODE_LOADED = 0x00000002 # type: ignore
-AMDGPU_CPCE_UCODE_LOADED = 0x00000004 # type: ignore
-AMDGPU_CPPFP_UCODE_LOADED = 0x00000008 # type: ignore
-AMDGPU_CPME_UCODE_LOADED = 0x00000010 # type: ignore
-AMDGPU_CPMEC1_UCODE_LOADED = 0x00000020 # type: ignore
-AMDGPU_CPMEC2_UCODE_LOADED = 0x00000040 # type: ignore
-AMDGPU_CPRLC_UCODE_LOADED = 0x00000100 # type: ignore
-PSP_GFX_CMD_BUF_VERSION = 0x00000001 # type: ignore
-GFX_CMD_STATUS_MASK = 0x0000FFFF # type: ignore
-GFX_CMD_ID_MASK = 0x000F0000 # type: ignore
-GFX_CMD_RESERVED_MASK = 0x7FF00000 # type: ignore
-GFX_CMD_RESPONSE_MASK = 0x80000000 # type: ignore
-C2PMSG_CMD_GFX_USB_PD_FW_VER = 0x2000000 # type: ignore
-GFX_FLAG_RESPONSE = 0x80000000 # type: ignore
-GFX_BUF_MAX_DESC = 64 # type: ignore
-FRAME_TYPE_DESTROY = 1 # type: ignore
-PSP_ERR_UNKNOWN_COMMAND = 0x00000100 # type: ignore
-PSP_FENCE_BUFFER_SIZE = 0x1000 # type: ignore
-PSP_CMD_BUFFER_SIZE = 0x1000 # type: ignore
-PSP_1_MEG = 0x100000 # type: ignore
-PSP_TMR_ALIGNMENT = 0x100000 # type: ignore
-PSP_FW_NAME_LEN = 0x24 # type: ignore
-AMDGPU_XGMI_MAX_CONNECTED_NODES = 64 # type: ignore
-MEM_TRAIN_SYSTEM_SIGNATURE = 0x54534942 # type: ignore
-GDDR6_MEM_TRAINING_DATA_SIZE_IN_BYTES = 0x1000 # type: ignore
-GDDR6_MEM_TRAINING_OFFSET = 0x8000 # type: ignore
-BIST_MEM_TRAINING_ENCROACHED_SIZE = 0x2000000 # type: ignore
-PSP_RUNTIME_DB_SIZE_IN_BYTES = 0x10000 # type: ignore
-PSP_RUNTIME_DB_OFFSET = 0x100000 # type: ignore
-PSP_RUNTIME_DB_COOKIE_ID = 0x0ed5 # type: ignore
-PSP_RUNTIME_DB_VER_1 = 0x0100 # type: ignore
-PSP_RUNTIME_DB_DIAG_ENTRY_MAX_COUNT = 0x40 # type: ignore
-int32_t = int # type: ignore
-AMDGPU_MAX_IRQ_SRC_ID = 0x100 # type: ignore
-AMDGPU_MAX_IRQ_CLIENT_ID = 0x100 # type: ignore
-AMDGPU_IRQ_CLIENTID_LEGACY = 0 # type: ignore
-AMDGPU_IRQ_CLIENTID_MAX = SOC15_IH_CLIENTID_MAX # type: ignore
-AMDGPU_IRQ_SRC_DATA_MAX_SIZE_DW = 4 # type: ignore
-SOC15_INTSRC_CP_END_OF_PIPE = 181 # type: ignore
-SOC15_INTSRC_CP_BAD_OPCODE = 183 # type: ignore
-SOC15_INTSRC_SQ_INTERRUPT_MSG = 239 # type: ignore
-SOC15_INTSRC_VMC_FAULT = 0 # type: ignore
-SOC15_INTSRC_VMC_UTCL2_POISON = 1 # type: ignore
-SOC15_INTSRC_SDMA_TRAP = 224 # type: ignore
-SOC15_INTSRC_SDMA_ECC = 220 # type: ignore
-SOC21_INTSRC_SDMA_TRAP = 49 # type: ignore
-SOC21_INTSRC_SDMA_ECC = 62 # type: ignore
+AMDGPU_VA_RESERVED_BOTTOM = (1 << 16)
+AMDGPU_VA_RESERVED_TOP = (AMDGPU_VA_RESERVED_TRAP_SIZE + AMDGPU_VA_RESERVED_SEQ64_SIZE + AMDGPU_VA_RESERVED_CSA_SIZE)
+AMDGPU_VM_USE_CPU_FOR_GFX = (1 << 0)
+AMDGPU_VM_USE_CPU_FOR_COMPUTE = (1 << 1)
+PSP_HEADER_SIZE = 256
+BINARY_SIGNATURE = 0x28211407
+DISCOVERY_TABLE_SIGNATURE = 0x53445049
+GC_TABLE_ID = 0x4347
+HARVEST_TABLE_SIGNATURE = 0x56524148
+VCN_INFO_TABLE_ID = 0x004E4356
+MALL_INFO_TABLE_ID = 0x4C4C414D
+NPS_INFO_TABLE_ID = 0x0053504E
+VCN_INFO_TABLE_MAX_NUM_INSTANCES = 4
+NPS_INFO_TABLE_MAX_NUM_INSTANCES = 12
+HWIP_MAX_INSTANCE = 44
+HW_ID_MAX = 300
+MP1_HWID = 1
+MP2_HWID = 2
+THM_HWID = 3
+SMUIO_HWID = 4
+FUSE_HWID = 5
+CLKA_HWID = 6
+PWR_HWID = 10
+GC_HWID = 11
+UVD_HWID = 12
+VCN_HWID = UVD_HWID
+AUDIO_AZ_HWID = 13
+ACP_HWID = 14
+DCI_HWID = 15
+DMU_HWID = 271
+DCO_HWID = 16
+DIO_HWID = 272
+XDMA_HWID = 17
+DCEAZ_HWID = 18
+DAZ_HWID = 274
+SDPMUX_HWID = 19
+NTB_HWID = 20
+VPE_HWID = 21
+IOHC_HWID = 24
+L2IMU_HWID = 28
+VCE_HWID = 32
+MMHUB_HWID = 34
+ATHUB_HWID = 35
+DBGU_NBIO_HWID = 36
+DFX_HWID = 37
+DBGU0_HWID = 38
+DBGU1_HWID = 39
+OSSSYS_HWID = 40
+HDP_HWID = 41
+SDMA0_HWID = 42
+SDMA1_HWID = 43
+ISP_HWID = 44
+DBGU_IO_HWID = 45
+DF_HWID = 46
+CLKB_HWID = 47
+FCH_HWID = 48
+DFX_DAP_HWID = 49
+L1IMU_PCIE_HWID = 50
+L1IMU_NBIF_HWID = 51
+L1IMU_IOAGR_HWID = 52
+L1IMU3_HWID = 53
+L1IMU4_HWID = 54
+L1IMU5_HWID = 55
+L1IMU6_HWID = 56
+L1IMU7_HWID = 57
+L1IMU8_HWID = 58
+L1IMU9_HWID = 59
+L1IMU10_HWID = 60
+L1IMU11_HWID = 61
+L1IMU12_HWID = 62
+L1IMU13_HWID = 63
+L1IMU14_HWID = 64
+L1IMU15_HWID = 65
+WAFLC_HWID = 66
+FCH_USB_PD_HWID = 67
+SDMA2_HWID = 68
+SDMA3_HWID = 69
+PCIE_HWID = 70
+PCS_HWID = 80
+DDCL_HWID = 89
+SST_HWID = 90
+LSDMA_HWID = 91
+IOAGR_HWID = 100
+NBIF_HWID = 108
+IOAPIC_HWID = 124
+SYSTEMHUB_HWID = 128
+NTBCCP_HWID = 144
+UMC_HWID = 150
+SATA_HWID = 168
+USB_HWID = 170
+CCXSEC_HWID = 176
+XGMI_HWID = 200
+XGBE_HWID = 216
+MP0_HWID = 255
+hw_id_map = {GC_HWIP:GC_HWID,HDP_HWIP:HDP_HWID,SDMA0_HWIP:SDMA0_HWID,SDMA1_HWIP:SDMA1_HWID,SDMA2_HWIP:SDMA2_HWID,SDMA3_HWIP:SDMA3_HWID,LSDMA_HWIP:LSDMA_HWID,MMHUB_HWIP:MMHUB_HWID,ATHUB_HWIP:ATHUB_HWID,NBIO_HWIP:NBIF_HWID,MP0_HWIP:MP0_HWID,MP1_HWIP:MP1_HWID,UVD_HWIP:UVD_HWID,VCE_HWIP:VCE_HWID,DF_HWIP:DF_HWID,DCE_HWIP:DMU_HWID,OSSSYS_HWIP:OSSSYS_HWID,SMUIO_HWIP:SMUIO_HWID,PWR_HWIP:PWR_HWID,NBIF_HWIP:NBIF_HWID,THM_HWIP:THM_HWID,CLK_HWIP:CLKA_HWID,UMC_HWIP:UMC_HWID,XGMI_HWIP:XGMI_HWID,DCI_HWIP:DCI_HWID,PCIE_HWIP:PCIE_HWID,VPE_HWIP:VPE_HWID,ISP_HWIP:ISP_HWID}
+AMDGPU_SDMA0_UCODE_LOADED = 0x00000001
+AMDGPU_SDMA1_UCODE_LOADED = 0x00000002
+AMDGPU_CPCE_UCODE_LOADED = 0x00000004
+AMDGPU_CPPFP_UCODE_LOADED = 0x00000008
+AMDGPU_CPME_UCODE_LOADED = 0x00000010
+AMDGPU_CPMEC1_UCODE_LOADED = 0x00000020
+AMDGPU_CPMEC2_UCODE_LOADED = 0x00000040
+AMDGPU_CPRLC_UCODE_LOADED = 0x00000100
+PSP_GFX_CMD_BUF_VERSION = 0x00000001
+GFX_CMD_STATUS_MASK = 0x0000FFFF
+GFX_CMD_ID_MASK = 0x000F0000
+GFX_CMD_RESERVED_MASK = 0x7FF00000
+GFX_CMD_RESPONSE_MASK = 0x80000000
+C2PMSG_CMD_GFX_USB_PD_FW_VER = 0x2000000
+GFX_FLAG_RESPONSE = 0x80000000
+GFX_BUF_MAX_DESC = 64
+FRAME_TYPE_DESTROY = 1
+PSP_ERR_UNKNOWN_COMMAND = 0x00000100
+PSP_FENCE_BUFFER_SIZE = 0x1000
+PSP_CMD_BUFFER_SIZE = 0x1000
+PSP_1_MEG = 0x100000
+PSP_TMR_ALIGNMENT = 0x100000
+PSP_FW_NAME_LEN = 0x24
+AMDGPU_XGMI_MAX_CONNECTED_NODES = 64
+MEM_TRAIN_SYSTEM_SIGNATURE = 0x54534942
+GDDR6_MEM_TRAINING_DATA_SIZE_IN_BYTES = 0x1000
+GDDR6_MEM_TRAINING_OFFSET = 0x8000
+BIST_MEM_TRAINING_ENCROACHED_SIZE = 0x2000000
+PSP_RUNTIME_DB_SIZE_IN_BYTES = 0x10000
+PSP_RUNTIME_DB_OFFSET = 0x100000
+PSP_RUNTIME_DB_COOKIE_ID = 0x0ed5
+PSP_RUNTIME_DB_VER_1 = 0x0100
+PSP_RUNTIME_DB_DIAG_ENTRY_MAX_COUNT = 0x40
+AMDGPU_MAX_IRQ_SRC_ID = 0x100
+AMDGPU_MAX_IRQ_CLIENT_ID = 0x100
+AMDGPU_IRQ_CLIENTID_LEGACY = 0
+AMDGPU_IRQ_CLIENTID_MAX = SOC15_IH_CLIENTID_MAX
+AMDGPU_IRQ_SRC_DATA_MAX_SIZE_DW = 4
+SOC15_INTSRC_CP_END_OF_PIPE = 181
+SOC15_INTSRC_CP_BAD_OPCODE = 183
+SOC15_INTSRC_SQ_INTERRUPT_MSG = 239
+SOC15_INTSRC_VMC_FAULT = 0
+SOC15_INTSRC_VMC_UTCL2_POISON = 1
+SOC15_INTSRC_SDMA_TRAP = 224
+SOC15_INTSRC_SDMA_ECC = 220
+SOC21_INTSRC_SDMA_TRAP = 49
+SOC21_INTSRC_SDMA_ECC = 62
 SOC15_CLIENT_ID_FROM_IH_ENTRY = lambda entry: ((entry[0]) & 0xff) # type: ignore
 SOC15_SOURCE_ID_FROM_IH_ENTRY = lambda entry: ((entry[0]) >> 8 & 0xff) # type: ignore
 SOC15_RING_ID_FROM_IH_ENTRY = lambda entry: ((entry[0]) >> 16 & 0xff) # type: ignore
@@ -4316,155 +4314,155 @@ SOC15_CONTEXT_ID0_FROM_IH_ENTRY = lambda entry: ((entry[4])) # type: ignore
 SOC15_CONTEXT_ID1_FROM_IH_ENTRY = lambda entry: ((entry[5])) # type: ignore
 SOC15_CONTEXT_ID2_FROM_IH_ENTRY = lambda entry: ((entry[6])) # type: ignore
 SOC15_CONTEXT_ID3_FROM_IH_ENTRY = lambda entry: ((entry[7])) # type: ignore
-GFX_9_0__SRCID__CP_RB_INTERRUPT_PKT = 176 # type: ignore
-GFX_9_0__SRCID__CP_IB1_INTERRUPT_PKT = 177 # type: ignore
-GFX_9_0__SRCID__CP_IB2_INTERRUPT_PKT = 178 # type: ignore
-GFX_9_0__SRCID__CP_PM4_PKT_RSVD_BIT_ERROR = 180 # type: ignore
-GFX_9_0__SRCID__CP_EOP_INTERRUPT = 181 # type: ignore
-GFX_9_0__SRCID__CP_BAD_OPCODE_ERROR = 183 # type: ignore
-GFX_9_0__SRCID__CP_PRIV_REG_FAULT = 184 # type: ignore
-GFX_9_0__SRCID__CP_PRIV_INSTR_FAULT = 185 # type: ignore
-GFX_9_0__SRCID__CP_WAIT_MEM_SEM_FAULT = 186 # type: ignore
-GFX_9_0__SRCID__CP_CTX_EMPTY_INTERRUPT = 187 # type: ignore
-GFX_9_0__SRCID__CP_CTX_BUSY_INTERRUPT = 188 # type: ignore
-GFX_9_0__SRCID__CP_ME_WAIT_REG_MEM_POLL_TIMEOUT = 192 # type: ignore
-GFX_9_0__SRCID__CP_SIG_INCOMPLETE = 193 # type: ignore
-GFX_9_0__SRCID__CP_PREEMPT_ACK = 194 # type: ignore
-GFX_9_0__SRCID__CP_GPF = 195 # type: ignore
-GFX_9_0__SRCID__CP_GDS_ALLOC_ERROR = 196 # type: ignore
-GFX_9_0__SRCID__CP_ECC_ERROR = 197 # type: ignore
-GFX_9_0__SRCID__CP_COMPUTE_QUERY_STATUS = 199 # type: ignore
-GFX_9_0__SRCID__CP_VM_DOORBELL = 200 # type: ignore
-GFX_9_0__SRCID__CP_FUE_ERROR = 201 # type: ignore
-GFX_9_0__SRCID__RLC_STRM_PERF_MONITOR_INTERRUPT = 202 # type: ignore
-GFX_9_0__SRCID__GRBM_RD_TIMEOUT_ERROR = 232 # type: ignore
-GFX_9_0__SRCID__GRBM_REG_GUI_IDLE = 233 # type: ignore
-GFX_9_0__SRCID__SQ_INTERRUPT_ID = 239 # type: ignore
-GFX_11_0_0__SRCID__UTCL2_FAULT = 0 # type: ignore
-GFX_11_0_0__SRCID__UTCL2_DATA_POISONING = 1 # type: ignore
-GFX_11_0_0__SRCID__MEM_ACCES_MON = 10 # type: ignore
-GFX_11_0_0__SRCID__SDMA_ATOMIC_RTN_DONE = 48 # type: ignore
-GFX_11_0_0__SRCID__SDMA_TRAP = 49 # type: ignore
-GFX_11_0_0__SRCID__SDMA_SRBMWRITE = 50 # type: ignore
-GFX_11_0_0__SRCID__SDMA_CTXEMPTY = 51 # type: ignore
-GFX_11_0_0__SRCID__SDMA_PREEMPT = 52 # type: ignore
-GFX_11_0_0__SRCID__SDMA_IB_PREEMPT = 53 # type: ignore
-GFX_11_0_0__SRCID__SDMA_DOORBELL_INVALID = 54 # type: ignore
-GFX_11_0_0__SRCID__SDMA_QUEUE_HANG = 55 # type: ignore
-GFX_11_0_0__SRCID__SDMA_ATOMIC_TIMEOUT = 56 # type: ignore
-GFX_11_0_0__SRCID__SDMA_POLL_TIMEOUT = 57 # type: ignore
-GFX_11_0_0__SRCID__SDMA_PAGE_TIMEOUT = 58 # type: ignore
-GFX_11_0_0__SRCID__SDMA_PAGE_NULL = 59 # type: ignore
-GFX_11_0_0__SRCID__SDMA_PAGE_FAULT = 60 # type: ignore
-GFX_11_0_0__SRCID__SDMA_VM_HOLE = 61 # type: ignore
-GFX_11_0_0__SRCID__SDMA_ECC = 62 # type: ignore
-GFX_11_0_0__SRCID__SDMA_FROZEN = 63 # type: ignore
-GFX_11_0_0__SRCID__SDMA_SRAM_ECC = 64 # type: ignore
-GFX_11_0_0__SRCID__SDMA_SEM_INCOMPLETE_TIMEOUT = 65 # type: ignore
-GFX_11_0_0__SRCID__SDMA_SEM_WAIT_FAIL_TIMEOUT = 66 # type: ignore
-GFX_11_0_0__SRCID__SDMA_FENCE = 67 # type: ignore
-GFX_11_0_0__SRCID__RLC_GC_FED_INTERRUPT = 128 # type: ignore
-GFX_11_0_0__SRCID__CP_GENERIC_INT = 177 # type: ignore
-GFX_11_0_0__SRCID__CP_PM4_PKT_RSVD_BIT_ERROR = 180 # type: ignore
-GFX_11_0_0__SRCID__CP_EOP_INTERRUPT = 181 # type: ignore
-GFX_11_0_0__SRCID__CP_BAD_OPCODE_ERROR = 183 # type: ignore
-GFX_11_0_0__SRCID__CP_PRIV_REG_FAULT = 184 # type: ignore
-GFX_11_0_0__SRCID__CP_PRIV_INSTR_FAULT = 185 # type: ignore
-GFX_11_0_0__SRCID__CP_WAIT_MEM_SEM_FAULT = 186 # type: ignore
-GFX_11_0_0__SRCID__CP_CTX_EMPTY_INTERRUPT = 187 # type: ignore
-GFX_11_0_0__SRCID__CP_CTX_BUSY_INTERRUPT = 188 # type: ignore
-GFX_11_0_0__SRCID__CP_ME_WAIT_REG_MEM_POLL_TIMEOUT = 192 # type: ignore
-GFX_11_0_0__SRCID__CP_SIG_INCOMPLETE = 193 # type: ignore
-GFX_11_0_0__SRCID__CP_PREEMPT_ACK = 194 # type: ignore
-GFX_11_0_0__SRCID__CP_GPF = 195 # type: ignore
-GFX_11_0_0__SRCID__CP_GDS_ALLOC_ERROR = 196 # type: ignore
-GFX_11_0_0__SRCID__CP_ECC_ERROR = 197 # type: ignore
-GFX_11_0_0__SRCID__CP_COMPUTE_QUERY_STATUS = 199 # type: ignore
-GFX_11_0_0__SRCID__CP_VM_DOORBELL = 200 # type: ignore
-GFX_11_0_0__SRCID__CP_FUE_ERROR = 201 # type: ignore
-GFX_11_0_0__SRCID__RLC_STRM_PERF_MONITOR_INTERRUPT = 202 # type: ignore
-GFX_11_0_0__SRCID__GRBM_RD_TIMEOUT_ERROR = 232 # type: ignore
-GFX_11_0_0__SRCID__GRBM_REG_GUI_IDLE = 233 # type: ignore
-GFX_11_0_0__SRCID__SQ_INTERRUPT_ID = 239 # type: ignore
-GFX_12_0_0__SRCID__UTCL2_FAULT = 0 # type: ignore
-GFX_12_0_0__SRCID__UTCL2_DATA_POISONING = 1 # type: ignore
-GFX_12_0_0__SRCID__MEM_ACCES_MON = 10 # type: ignore
-GFX_12_0_0__SRCID__SDMA_ATOMIC_RTN_DONE = 48 # type: ignore
-GFX_12_0_0__SRCID__SDMA_TRAP = 49 # type: ignore
-GFX_12_0_0__SRCID__SDMA_SRBMWRITE = 50 # type: ignore
-GFX_12_0_0__SRCID__SDMA_CTXEMPTY = 51 # type: ignore
-GFX_12_0_0__SRCID__SDMA_PREEMPT = 52 # type: ignore
-GFX_12_0_0__SRCID__SDMA_IB_PREEMPT = 53 # type: ignore
-GFX_12_0_0__SRCID__SDMA_DOORBELL_INVALID = 54 # type: ignore
-GFX_12_0_0__SRCID__SDMA_QUEUE_HANG = 55 # type: ignore
-GFX_12_0_0__SRCID__SDMA_ATOMIC_TIMEOUT = 56 # type: ignore
-GFX_12_0_0__SRCID__SDMA_POLL_TIMEOUT = 57 # type: ignore
-GFX_12_0_0__SRCID__SDMA_PAGE_TIMEOUT = 58 # type: ignore
-GFX_12_0_0__SRCID__SDMA_PAGE_NULL = 59 # type: ignore
-GFX_12_0_0__SRCID__SDMA_PAGE_FAULT = 60 # type: ignore
-GFX_12_0_0__SRCID__SDMA_VM_HOLE = 61 # type: ignore
-GFX_12_0_0__SRCID__SDMA_ECC = 62 # type: ignore
-GFX_12_0_0__SRCID__SDMA_FROZEN = 63 # type: ignore
-GFX_12_0_0__SRCID__SDMA_SRAM_ECC = 64 # type: ignore
-GFX_12_0_0__SRCID__SDMA_SEM_INCOMPLETE_TIMEOUT = 65 # type: ignore
-GFX_12_0_0__SRCID__SDMA_SEM_WAIT_FAIL_TIMEOUT = 66 # type: ignore
-GFX_12_0_0__SRCID__SDMA_FENCE = 70 # type: ignore
-GFX_12_0_0__SRCID__RLC_GC_FED_INTERRUPT = 128 # type: ignore
-GFX_12_0_0__SRCID__CP_GENERIC_INT = 177 # type: ignore
-GFX_12_0_0__SRCID__CP_PM4_PKT_RSVD_BIT_ERROR = 180 # type: ignore
-GFX_12_0_0__SRCID__CP_EOP_INTERRUPT = 181 # type: ignore
-GFX_12_0_0__SRCID__CP_BAD_OPCODE_ERROR = 183 # type: ignore
-GFX_12_0_0__SRCID__CP_PRIV_REG_FAULT = 184 # type: ignore
-GFX_12_0_0__SRCID__CP_PRIV_INSTR_FAULT = 185 # type: ignore
-GFX_12_0_0__SRCID__CP_WAIT_MEM_SEM_FAULT = 186 # type: ignore
-GFX_12_0_0__SRCID__CP_CTX_EMPTY_INTERRUPT = 187 # type: ignore
-GFX_12_0_0__SRCID__CP_CTX_BUSY_INTERRUPT = 188 # type: ignore
-GFX_12_0_0__SRCID__CP_ME_WAIT_REG_MEM_POLL_TIMEOUT = 192 # type: ignore
-GFX_12_0_0__SRCID__CP_SIG_INCOMPLETE = 193 # type: ignore
-GFX_12_0_0__SRCID__CP_PREEMPT_ACK = 194 # type: ignore
-GFX_12_0_0__SRCID__CP_GPF = 195 # type: ignore
-GFX_12_0_0__SRCID__CP_GDS_ALLOC_ERROR = 196 # type: ignore
-GFX_12_0_0__SRCID__CP_ECC_ERROR = 197 # type: ignore
-GFX_12_0_0__SRCID__CP_COMPUTE_QUERY_STATUS = 199 # type: ignore
-GFX_12_0_0__SRCID__CP_VM_DOORBELL = 200 # type: ignore
-GFX_12_0_0__SRCID__CP_FUE_ERROR = 201 # type: ignore
-GFX_12_0_0__SRCID__RLC_STRM_PERF_MONITOR_INTERRUPT = 202 # type: ignore
-GFX_12_0_0__SRCID__GRBM_RD_TIMEOUT_ERROR = 232 # type: ignore
-GFX_12_0_0__SRCID__GRBM_REG_GUI_IDLE = 233 # type: ignore
-GFX_12_0_0__SRCID__SQ_INTERRUPT_ID = 239 # type: ignore
-SDMA0_4_0__SRCID__SDMA_ATOMIC_RTN_DONE = 217 # type: ignore
-SDMA0_4_0__SRCID__SDMA_ATOMIC_TIMEOUT = 218 # type: ignore
-SDMA0_4_0__SRCID__SDMA_IB_PREEMPT = 219 # type: ignore
-SDMA0_4_0__SRCID__SDMA_ECC = 220 # type: ignore
-SDMA0_4_0__SRCID__SDMA_PAGE_FAULT = 221 # type: ignore
-SDMA0_4_0__SRCID__SDMA_PAGE_NULL = 222 # type: ignore
-SDMA0_4_0__SRCID__SDMA_XNACK = 223 # type: ignore
-SDMA0_4_0__SRCID__SDMA_TRAP = 224 # type: ignore
-SDMA0_4_0__SRCID__SDMA_SEM_INCOMPLETE_TIMEOUT = 225 # type: ignore
-SDMA0_4_0__SRCID__SDMA_SEM_WAIT_FAIL_TIMEOUT = 226 # type: ignore
-SDMA0_4_0__SRCID__SDMA_SRAM_ECC = 228 # type: ignore
-SDMA0_4_0__SRCID__SDMA_PREEMPT = 240 # type: ignore
-SDMA0_4_0__SRCID__SDMA_VM_HOLE = 242 # type: ignore
-SDMA0_4_0__SRCID__SDMA_CTXEMPTY = 243 # type: ignore
-SDMA0_4_0__SRCID__SDMA_DOORBELL_INVALID = 244 # type: ignore
-SDMA0_4_0__SRCID__SDMA_FROZEN = 245 # type: ignore
-SDMA0_4_0__SRCID__SDMA_POLL_TIMEOUT = 246 # type: ignore
-SDMA0_4_0__SRCID__SDMA_SRBMWRITE = 247 # type: ignore
-SDMA0_5_0__SRCID__SDMA_ATOMIC_RTN_DONE = 217 # type: ignore
-SDMA0_5_0__SRCID__SDMA_ATOMIC_TIMEOUT = 218 # type: ignore
-SDMA0_5_0__SRCID__SDMA_IB_PREEMPT = 219 # type: ignore
-SDMA0_5_0__SRCID__SDMA_ECC = 220 # type: ignore
-SDMA0_5_0__SRCID__SDMA_PAGE_FAULT = 221 # type: ignore
-SDMA0_5_0__SRCID__SDMA_PAGE_NULL = 222 # type: ignore
-SDMA0_5_0__SRCID__SDMA_XNACK = 223 # type: ignore
-SDMA0_5_0__SRCID__SDMA_TRAP = 224 # type: ignore
-SDMA0_5_0__SRCID__SDMA_SEM_INCOMPLETE_TIMEOUT = 225 # type: ignore
-SDMA0_5_0__SRCID__SDMA_SEM_WAIT_FAIL_TIMEOUT = 226 # type: ignore
-SDMA0_5_0__SRCID__SDMA_SRAM_ECC = 228 # type: ignore
-SDMA0_5_0__SRCID__SDMA_PREEMPT = 240 # type: ignore
-SDMA0_5_0__SRCID__SDMA_VM_HOLE = 242 # type: ignore
-SDMA0_5_0__SRCID__SDMA_CTXEMPTY = 243 # type: ignore
-SDMA0_5_0__SRCID__SDMA_DOORBELL_INVALID = 244 # type: ignore
-SDMA0_5_0__SRCID__SDMA_FROZEN = 245 # type: ignore
-SDMA0_5_0__SRCID__SDMA_POLL_TIMEOUT = 246 # type: ignore
-SDMA0_5_0__SRCID__SDMA_SRBMWRITE = 247 # type: ignore
+GFX_9_0__SRCID__CP_RB_INTERRUPT_PKT = 176
+GFX_9_0__SRCID__CP_IB1_INTERRUPT_PKT = 177
+GFX_9_0__SRCID__CP_IB2_INTERRUPT_PKT = 178
+GFX_9_0__SRCID__CP_PM4_PKT_RSVD_BIT_ERROR = 180
+GFX_9_0__SRCID__CP_EOP_INTERRUPT = 181
+GFX_9_0__SRCID__CP_BAD_OPCODE_ERROR = 183
+GFX_9_0__SRCID__CP_PRIV_REG_FAULT = 184
+GFX_9_0__SRCID__CP_PRIV_INSTR_FAULT = 185
+GFX_9_0__SRCID__CP_WAIT_MEM_SEM_FAULT = 186
+GFX_9_0__SRCID__CP_CTX_EMPTY_INTERRUPT = 187
+GFX_9_0__SRCID__CP_CTX_BUSY_INTERRUPT = 188
+GFX_9_0__SRCID__CP_ME_WAIT_REG_MEM_POLL_TIMEOUT = 192
+GFX_9_0__SRCID__CP_SIG_INCOMPLETE = 193
+GFX_9_0__SRCID__CP_PREEMPT_ACK = 194
+GFX_9_0__SRCID__CP_GPF = 195
+GFX_9_0__SRCID__CP_GDS_ALLOC_ERROR = 196
+GFX_9_0__SRCID__CP_ECC_ERROR = 197
+GFX_9_0__SRCID__CP_COMPUTE_QUERY_STATUS = 199
+GFX_9_0__SRCID__CP_VM_DOORBELL = 200
+GFX_9_0__SRCID__CP_FUE_ERROR = 201
+GFX_9_0__SRCID__RLC_STRM_PERF_MONITOR_INTERRUPT = 202
+GFX_9_0__SRCID__GRBM_RD_TIMEOUT_ERROR = 232
+GFX_9_0__SRCID__GRBM_REG_GUI_IDLE = 233
+GFX_9_0__SRCID__SQ_INTERRUPT_ID = 239
+GFX_11_0_0__SRCID__UTCL2_FAULT = 0
+GFX_11_0_0__SRCID__UTCL2_DATA_POISONING = 1
+GFX_11_0_0__SRCID__MEM_ACCES_MON = 10
+GFX_11_0_0__SRCID__SDMA_ATOMIC_RTN_DONE = 48
+GFX_11_0_0__SRCID__SDMA_TRAP = 49
+GFX_11_0_0__SRCID__SDMA_SRBMWRITE = 50
+GFX_11_0_0__SRCID__SDMA_CTXEMPTY = 51
+GFX_11_0_0__SRCID__SDMA_PREEMPT = 52
+GFX_11_0_0__SRCID__SDMA_IB_PREEMPT = 53
+GFX_11_0_0__SRCID__SDMA_DOORBELL_INVALID = 54
+GFX_11_0_0__SRCID__SDMA_QUEUE_HANG = 55
+GFX_11_0_0__SRCID__SDMA_ATOMIC_TIMEOUT = 56
+GFX_11_0_0__SRCID__SDMA_POLL_TIMEOUT = 57
+GFX_11_0_0__SRCID__SDMA_PAGE_TIMEOUT = 58
+GFX_11_0_0__SRCID__SDMA_PAGE_NULL = 59
+GFX_11_0_0__SRCID__SDMA_PAGE_FAULT = 60
+GFX_11_0_0__SRCID__SDMA_VM_HOLE = 61
+GFX_11_0_0__SRCID__SDMA_ECC = 62
+GFX_11_0_0__SRCID__SDMA_FROZEN = 63
+GFX_11_0_0__SRCID__SDMA_SRAM_ECC = 64
+GFX_11_0_0__SRCID__SDMA_SEM_INCOMPLETE_TIMEOUT = 65
+GFX_11_0_0__SRCID__SDMA_SEM_WAIT_FAIL_TIMEOUT = 66
+GFX_11_0_0__SRCID__SDMA_FENCE = 67
+GFX_11_0_0__SRCID__RLC_GC_FED_INTERRUPT = 128
+GFX_11_0_0__SRCID__CP_GENERIC_INT = 177
+GFX_11_0_0__SRCID__CP_PM4_PKT_RSVD_BIT_ERROR = 180
+GFX_11_0_0__SRCID__CP_EOP_INTERRUPT = 181
+GFX_11_0_0__SRCID__CP_BAD_OPCODE_ERROR = 183
+GFX_11_0_0__SRCID__CP_PRIV_REG_FAULT = 184
+GFX_11_0_0__SRCID__CP_PRIV_INSTR_FAULT = 185
+GFX_11_0_0__SRCID__CP_WAIT_MEM_SEM_FAULT = 186
+GFX_11_0_0__SRCID__CP_CTX_EMPTY_INTERRUPT = 187
+GFX_11_0_0__SRCID__CP_CTX_BUSY_INTERRUPT = 188
+GFX_11_0_0__SRCID__CP_ME_WAIT_REG_MEM_POLL_TIMEOUT = 192
+GFX_11_0_0__SRCID__CP_SIG_INCOMPLETE = 193
+GFX_11_0_0__SRCID__CP_PREEMPT_ACK = 194
+GFX_11_0_0__SRCID__CP_GPF = 195
+GFX_11_0_0__SRCID__CP_GDS_ALLOC_ERROR = 196
+GFX_11_0_0__SRCID__CP_ECC_ERROR = 197
+GFX_11_0_0__SRCID__CP_COMPUTE_QUERY_STATUS = 199
+GFX_11_0_0__SRCID__CP_VM_DOORBELL = 200
+GFX_11_0_0__SRCID__CP_FUE_ERROR = 201
+GFX_11_0_0__SRCID__RLC_STRM_PERF_MONITOR_INTERRUPT = 202
+GFX_11_0_0__SRCID__GRBM_RD_TIMEOUT_ERROR = 232
+GFX_11_0_0__SRCID__GRBM_REG_GUI_IDLE = 233
+GFX_11_0_0__SRCID__SQ_INTERRUPT_ID = 239
+GFX_12_0_0__SRCID__UTCL2_FAULT = 0
+GFX_12_0_0__SRCID__UTCL2_DATA_POISONING = 1
+GFX_12_0_0__SRCID__MEM_ACCES_MON = 10
+GFX_12_0_0__SRCID__SDMA_ATOMIC_RTN_DONE = 48
+GFX_12_0_0__SRCID__SDMA_TRAP = 49
+GFX_12_0_0__SRCID__SDMA_SRBMWRITE = 50
+GFX_12_0_0__SRCID__SDMA_CTXEMPTY = 51
+GFX_12_0_0__SRCID__SDMA_PREEMPT = 52
+GFX_12_0_0__SRCID__SDMA_IB_PREEMPT = 53
+GFX_12_0_0__SRCID__SDMA_DOORBELL_INVALID = 54
+GFX_12_0_0__SRCID__SDMA_QUEUE_HANG = 55
+GFX_12_0_0__SRCID__SDMA_ATOMIC_TIMEOUT = 56
+GFX_12_0_0__SRCID__SDMA_POLL_TIMEOUT = 57
+GFX_12_0_0__SRCID__SDMA_PAGE_TIMEOUT = 58
+GFX_12_0_0__SRCID__SDMA_PAGE_NULL = 59
+GFX_12_0_0__SRCID__SDMA_PAGE_FAULT = 60
+GFX_12_0_0__SRCID__SDMA_VM_HOLE = 61
+GFX_12_0_0__SRCID__SDMA_ECC = 62
+GFX_12_0_0__SRCID__SDMA_FROZEN = 63
+GFX_12_0_0__SRCID__SDMA_SRAM_ECC = 64
+GFX_12_0_0__SRCID__SDMA_SEM_INCOMPLETE_TIMEOUT = 65
+GFX_12_0_0__SRCID__SDMA_SEM_WAIT_FAIL_TIMEOUT = 66
+GFX_12_0_0__SRCID__SDMA_FENCE = 70
+GFX_12_0_0__SRCID__RLC_GC_FED_INTERRUPT = 128
+GFX_12_0_0__SRCID__CP_GENERIC_INT = 177
+GFX_12_0_0__SRCID__CP_PM4_PKT_RSVD_BIT_ERROR = 180
+GFX_12_0_0__SRCID__CP_EOP_INTERRUPT = 181
+GFX_12_0_0__SRCID__CP_BAD_OPCODE_ERROR = 183
+GFX_12_0_0__SRCID__CP_PRIV_REG_FAULT = 184
+GFX_12_0_0__SRCID__CP_PRIV_INSTR_FAULT = 185
+GFX_12_0_0__SRCID__CP_WAIT_MEM_SEM_FAULT = 186
+GFX_12_0_0__SRCID__CP_CTX_EMPTY_INTERRUPT = 187
+GFX_12_0_0__SRCID__CP_CTX_BUSY_INTERRUPT = 188
+GFX_12_0_0__SRCID__CP_ME_WAIT_REG_MEM_POLL_TIMEOUT = 192
+GFX_12_0_0__SRCID__CP_SIG_INCOMPLETE = 193
+GFX_12_0_0__SRCID__CP_PREEMPT_ACK = 194
+GFX_12_0_0__SRCID__CP_GPF = 195
+GFX_12_0_0__SRCID__CP_GDS_ALLOC_ERROR = 196
+GFX_12_0_0__SRCID__CP_ECC_ERROR = 197
+GFX_12_0_0__SRCID__CP_COMPUTE_QUERY_STATUS = 199
+GFX_12_0_0__SRCID__CP_VM_DOORBELL = 200
+GFX_12_0_0__SRCID__CP_FUE_ERROR = 201
+GFX_12_0_0__SRCID__RLC_STRM_PERF_MONITOR_INTERRUPT = 202
+GFX_12_0_0__SRCID__GRBM_RD_TIMEOUT_ERROR = 232
+GFX_12_0_0__SRCID__GRBM_REG_GUI_IDLE = 233
+GFX_12_0_0__SRCID__SQ_INTERRUPT_ID = 239
+SDMA0_4_0__SRCID__SDMA_ATOMIC_RTN_DONE = 217
+SDMA0_4_0__SRCID__SDMA_ATOMIC_TIMEOUT = 218
+SDMA0_4_0__SRCID__SDMA_IB_PREEMPT = 219
+SDMA0_4_0__SRCID__SDMA_ECC = 220
+SDMA0_4_0__SRCID__SDMA_PAGE_FAULT = 221
+SDMA0_4_0__SRCID__SDMA_PAGE_NULL = 222
+SDMA0_4_0__SRCID__SDMA_XNACK = 223
+SDMA0_4_0__SRCID__SDMA_TRAP = 224
+SDMA0_4_0__SRCID__SDMA_SEM_INCOMPLETE_TIMEOUT = 225
+SDMA0_4_0__SRCID__SDMA_SEM_WAIT_FAIL_TIMEOUT = 226
+SDMA0_4_0__SRCID__SDMA_SRAM_ECC = 228
+SDMA0_4_0__SRCID__SDMA_PREEMPT = 240
+SDMA0_4_0__SRCID__SDMA_VM_HOLE = 242
+SDMA0_4_0__SRCID__SDMA_CTXEMPTY = 243
+SDMA0_4_0__SRCID__SDMA_DOORBELL_INVALID = 244
+SDMA0_4_0__SRCID__SDMA_FROZEN = 245
+SDMA0_4_0__SRCID__SDMA_POLL_TIMEOUT = 246
+SDMA0_4_0__SRCID__SDMA_SRBMWRITE = 247
+SDMA0_5_0__SRCID__SDMA_ATOMIC_RTN_DONE = 217
+SDMA0_5_0__SRCID__SDMA_ATOMIC_TIMEOUT = 218
+SDMA0_5_0__SRCID__SDMA_IB_PREEMPT = 219
+SDMA0_5_0__SRCID__SDMA_ECC = 220
+SDMA0_5_0__SRCID__SDMA_PAGE_FAULT = 221
+SDMA0_5_0__SRCID__SDMA_PAGE_NULL = 222
+SDMA0_5_0__SRCID__SDMA_XNACK = 223
+SDMA0_5_0__SRCID__SDMA_TRAP = 224
+SDMA0_5_0__SRCID__SDMA_SEM_INCOMPLETE_TIMEOUT = 225
+SDMA0_5_0__SRCID__SDMA_SEM_WAIT_FAIL_TIMEOUT = 226
+SDMA0_5_0__SRCID__SDMA_SRAM_ECC = 228
+SDMA0_5_0__SRCID__SDMA_PREEMPT = 240
+SDMA0_5_0__SRCID__SDMA_VM_HOLE = 242
+SDMA0_5_0__SRCID__SDMA_CTXEMPTY = 243
+SDMA0_5_0__SRCID__SDMA_DOORBELL_INVALID = 244
+SDMA0_5_0__SRCID__SDMA_FROZEN = 245
+SDMA0_5_0__SRCID__SDMA_POLL_TIMEOUT = 246
+SDMA0_5_0__SRCID__SDMA_SRBMWRITE = 247

tinygrad/runtime/autogen/am/fw.py

@@ -0,0 +1,106 @@
+hashes = {
+  'psp_13_0_0_sos.bin': 'b5592f46885585b935e013f46c949db8ff2f15c0b346caf70e7fcd2776623d13',
+  'psp_13_0_10_sos.bin': '0bcaaad9cd8578d3841ae69155a6bd4fc3ceae8f4fb5a6ba4f576e7ace94d1d9',
+  'psp_13_0_12_sos.bin': '89da90bf4286b38678b1fd175c78462a426afa3d258d15872cd14072d7098b9b',
+  'psp_13_0_14_sos.bin': 'a4f0d5f76d27b77409ec0b71d7cc6a848ddfd29f8c84f3003edf74ad3999fb7d',
+  'psp_13_0_6_sos.bin': '27657daa0f91ad8095d3610224a7de748b8b348a4cb211ecb5fccabe47369716',
+  'psp_13_0_7_sos.bin': 'ef1af0ecea38abbac6f85cce71789f19848c498d0cb8ef13748dab2d65b23c31',
+  'psp_14_0_2_sos.bin': '7b538448b57d4f9dd06b2eea90d4f86a16e65e3027cdecee8db71c2c5f1fa243',
+  'psp_14_0_3_sos.bin': '23bea01a0c6f36d00759d0765d46cb4cb4aa87398b2fbccacbf547a890c0bf51',
+  'smu_13_0_0.bin': '2ffac37fd8534965eeba19755db0e5ec80278213487dc4af0fbc8453befb64b1',
+  'smu_13_0_0_kicker.bin': '7f83656a2a89b7fce1c8a85e96d91cd8265a91fe883a7027f1a0ed18ced501de',
+  'smu_13_0_10.bin': 'daedb9cbdf48942be7ffe00d31b7c16bb36e11ff5a9d7495f218e95c07717b71',
+  'smu_13_0_14.bin': 'a4f36de75fdcecd8000246762e027b4be489b6787afea57675225b0b39d35625',
+  'smu_13_0_6.bin': 'ad7232264e8c57c2094244fbdd5a55d7a4575ffe9b44d229884bc0b6a44fb0b1',
+  'smu_13_0_7.bin': 'ccecc0fd0196b9613c920a51c2fd9436e739ff19dda5bdf74d97562387231732',
+  'smu_14_0_2.bin': '6951995d1d606f4dc60c895f19d34ed18aa40e62129f83d8510c45e8aa9ae2fc',
+  'smu_14_0_3.bin': 'df230947ddb7bbfd6e77d1280001db886e69adf2b2a448b47fa668a48bc0009f',
+  'smu_14_0_3_kicker.bin': '8ddc1da5b4e1619796c2cc81f19f388a35bf7d78bfe476cee559625589cb4dc7',
+  'sdma_4_4_2.bin': '456061b814268425843537da6f2191c8861d4e1a18d4c5d90c44ea6be18c78ff',
+  'sdma_4_4_4.bin': 'af47a2940e72b932d3e3a7e8f34f7a182624e5e433f7c56dff939ca5549cd33a',
+  'sdma_4_4_5.bin': '6127baabea3de7b18db3868c983b02c0fbf2cd75997f7f11241a5b1be27e5134',
+  'sdma_5_2_6.bin': '3a163db00eb7e4752be8adbd61cf7dd8f08d924e59a6f798ced7dfcd89f340ed',
+  'sdma_5_2_7.bin': '16fe80dc866b323e15a06f51646ef0f036878ad34da66921fcdb8167207d6b2b',
+  'sdma_6_0_0.bin': '0f3da6b211f376356335b41be07149f650c10cfa4e23f7e25d53836006ed11f5',
+  'sdma_6_0_1.bin': 'ff565d3c215a30737560d4e3df6fc2c637738407e91d212fb200fdfb185b6744',
+  'sdma_6_0_2.bin': '398380184bb69113ef4c8964a3b55f6184deb0c1ffd96c9683490a3eec3ba8f3',
+  'sdma_6_0_3.bin': '0e8a83513087db865ba926f8b65cfb003fd41098f707e178d7a7ae2941fed0b1',
+  'sdma_6_1_0.bin': '22e55d0ad5f0247a7f0fffc67cfd3161b39f24ad6062ff3c91ec7ff38bd7e1e1',
+  'sdma_6_1_1.bin': '74533a581b8e3e2743b3c9c803d0666405e80898c4a630acefed82cb6b516ba2',
+  'sdma_6_1_2.bin': '4fe04b0286ec739b0414e8aee17e62e85e691f0246d1d9b56bc18a1219072314',
+  'sdma_6_1_3.bin': '35c9ed7e3a237c0d4a83b4975c63b62488f72aeafbb648342f384618e103f66b',
+  'sdma_7_0_0.bin': 'beaafb53993a106edd392392d5896245ae2a957c6d0f495d0002eec72ad8ad38',
+  'sdma_7_0_1.bin': '73c29e1c1714ebc95d2221ba56e187910902891593010653bf9518937e414a59',
+  'gc_10_3_6_pfp.bin': '793d678427887a0e724c79e356440aec33e6d1301f2a4e63543500249ebec064',
+  'gc_10_3_7_pfp.bin': '3ae29aac3f424f7de97f82ce7158beba69509afb2dcbf1a428dc315df474a524',
+  'gc_11_0_0_pfp.bin': 'e175cb0f580a38c961a6f7366142c08e413995f57f78f39795368b15442df8a3',
+  'gc_11_0_1_pfp.bin': 'f5bf21dfbd9e72a30b4caf4704282c27854710c1b7c4affbb2a19530466b12a8',
+  'gc_11_0_2_pfp.bin': '001c4dec1119e29314d725cc1280fc4f0cd9cabdf61ea5ee2260cfd4e62ec141',
+  'gc_11_0_3_pfp.bin': '0488034c85be97125e39e860308d33c3f76a01df8250092a32d4d55acb2526fd',
+  'gc_11_0_4_pfp.bin': '5ae8b7bb6316f87ae8b978354c088e3bd8c890959382d72886377cda25b1ffd1',
+  'gc_11_5_0_pfp.bin': '0124f540871a7759fa8aaae046d458dfb34aeea12a1183ff962c3f1a33067d5a',
+  'gc_11_5_1_pfp.bin': '7794ea46d0d3cf9cb3f7938affbdf09dd7a9970340da5cd02b774cb393436d24',
+  'gc_11_5_2_pfp.bin': '55e64741de28c506524959f7f696713a72aafe46f49ccd827781d67a9475b386',
+  'gc_11_5_3_pfp.bin': 'ce805040fb347fddbc89b2715e66b446865dda9e2056a9b233269b72bc09c387',
+  'gc_12_0_0_pfp.bin': '16bfd64c10fe73b5e760055069a60e5841dba16c0ed4edb56c20d675e23901f6',
+  'gc_12_0_1_pfp.bin': '49efb319305c5fffd90ac1eef7d7a0bdec72998ecb5cf4526996311788a53dc3',
+  'gc_10_3_6_me.bin': '141b59faad3f2f1be16a2178833b7ca8e97519e1e844c8fda6689572c3767902',
+  'gc_10_3_7_me.bin': '9eb0b56e9bcc9dad5d53437b162226fcb37e5df102832260f1232832f3658edf',
+  'gc_11_0_0_me.bin': 'f8fba8a63dd4293b8fc1e4aab78b6fac630e575d1d62838c7996d9210f82aea1',
+  'gc_11_0_1_me.bin': '5030040b00955de94876341ec64ea43b96640413d7a03dc460a83c8386bf76e0',
+  'gc_11_0_2_me.bin': '0f21fd43f1dfbc6ccced9a2b3774de25c993c61a689aabab8b45333937b7945e',
+  'gc_11_0_3_me.bin': '3acb5061dba342ade81d329d1932f19ec01f0c5bf44e6e3568008a951a351bac',
+  'gc_11_0_4_me.bin': 'e4f1f6abcd213d54ad9e885d9f550083b0e2f67d983566015e8a53981e1cb155',
+  'gc_11_5_0_me.bin': '8f906b64d0a29503daa662c93ec44d076fcac11b78f70cd50ce0af2b500a05a6',
+  'gc_11_5_1_me.bin': '7e42602bcbaf1e511f8b4f6ed2246844ad1f6e351ce2b663d89062a7be263663',
+  'gc_11_5_2_me.bin': 'aae26255d8efff81e0e3bbcb727efb8b837d8e25fe85c708545f5328f1077b50',
+  'gc_11_5_3_me.bin': '93cd588348b16fe432609fe8da6e6b5da0a52da5c5884882aecf7b1001f72700',
+  'gc_12_0_0_me.bin': 'd7eba5197f2580f32b8256b1d9cb68e723e9e644293a34446a7913e3c093cba5',
+  'gc_12_0_1_me.bin': '365e7f193b39cbb10d3af44905fefaca0e9844721801755276baebac7b19c1ea',
+  'gc_10_3_6_mec.bin': '247943415658159704a21f670dd7b3e7cb2d2fc0c17b000a5098715979c8d95e',
+  'gc_10_3_7_mec.bin': 'ee58a523375bcf5b89400b32b801f95e182b632a26bce4f2bed5c07928d486dc',
+  'gc_11_0_0_mec.bin': '801a09c9bf06188260db9b51ad8f978f15d84c72ca91b90643a2ef8af4074776',
+  'gc_11_0_1_mec.bin': '6afadcb7504bb11bcc9d4a205cdf73f7934a615e28f178fcf7285971df2ccd05',
+  'gc_11_0_2_mec.bin': '0da0edee28c73a6fa1191f77853d380ec2503cbf43e0aaae4617f32f1f8a48fa',
+  'gc_11_0_3_mec.bin': '323cfa6658b6b5169830f852e2ff0552acae8dfb9e44b42c63de7b2900d3fd9e',
+  'gc_11_0_4_mec.bin': '5d89cf6b60354f3746c2cbd1ff0cb1a741556ca20d72745242cb69b553d0985c',
+  'gc_11_5_0_mec.bin': 'a01c324ab14ec89792449a621a541829b9af26865019027a411a14b910145dfa',
+  'gc_11_5_1_mec.bin': 'eab05719371caa68df09d4f7574e3958a3c4f5044ab3c7b0d2b214add0c6d1c4',
+  'gc_11_5_2_mec.bin': 'a374b2335802e24f8b9a3ce40000a1d37a52a14eb87099bebcc6680c27cc93e5',
+  'gc_11_5_3_mec.bin': '165025437cba80dd32c19ebbc83b756fa7adac7053ff7780ba4aa2f8089c6a3f',
+  'gc_12_0_0_mec.bin': '1931593440b8f9423580d9e2cdc5b34e7c682cdffe1ca4b74b0c2f6a0420236d',
+  'gc_12_0_1_mec.bin': 'f57541688a5108730bf210663f1137ffc2121f3acfe614a6de09ec1982c69a2f',
+  'gc_9_4_3_mec.bin': '3159176e72301fb88dc416721fb3d0ab82ece484cf93a43c3f37430c7e6673a1',
+  'gc_9_4_3_sjt_mec.bin': 'd19468dbb47849640bd0e6cdc8d7e25a3c8442c7ca2ca81357702e0d6baab50f',
+  'gc_9_4_4_mec.bin': '5004f73e43db2dd45e77d65942e33d4a69e7157618cfd23944c30f801c77a0f3',
+  'gc_9_4_4_sjt_mec.bin': '627a9e98102e70fe3bf0947eb764187f29f5e775d1130c7310e0ba5fc0502dbe',
+  'gc_9_5_0_mec.bin': 'c5eca4311a6f6e8f81cf41c2c46941d5dcf90789ee8326901da2dfc86ac14c31',
+  'gc_9_5_0_sjt_mec.bin': 'f162e509379288e3f3b1eead541b315c2262d625d433287ecd34ca185614d312',
+  'gc_11_0_0_imu.bin': 'b4f8fc056b45709a6abf48e7885fb1b4ab8d3cc092cbfa2c554a78564a6403bc',
+  'gc_11_0_1_imu.bin': 'ac71f4eec713fc35b4a1fe27531e3eb04edd81eeac2cef64df01ac50d8510805',
+  'gc_11_0_2_imu.bin': '9befca62b0b0cfd252c3df4a9edca295526f4d43821cd99a6326454995a6ca2d',
+  'gc_11_0_3_imu.bin': 'beaf704d5acdf4623456b0d0cbcea8b8e428058340cd922a259a9045f5c457a3',
+  'gc_11_0_4_imu.bin': 'ac71f4eec713fc35b4a1fe27531e3eb04edd81eeac2cef64df01ac50d8510805',
+  'gc_11_5_0_imu.bin': '469add57cafead90ab1953d6039cd8e39bc50dfd78aa5fd78f019ccf66a0af41',
+  'gc_11_5_1_imu.bin': '0aaca8a01b2237fca1b3c0cd082b5e12a271df334ff368bbf8e2be17f192b785',
+  'gc_11_5_2_imu.bin': 'fb684842839c61a0706a19df8e15eb8afc17e20c14267eb71e7d7d824c180acf',
+  'gc_11_5_3_imu.bin': 'fb684842839c61a0706a19df8e15eb8afc17e20c14267eb71e7d7d824c180acf',
+  'gc_12_0_0_imu.bin': 'aa15e5b3156bffc45e0c50bccbcd364fbd3f958531b695b7487a803d780b8328',
+  'gc_12_0_1_imu.bin': 'b3b301fb636efc77b63ce4d2ced0f90c851d03c19681852faa45598e6f5773fd',
+  'gc_10_3_6_rlc.bin': 'acfbac75c0dcfbfe40e222640ef17eb3dc8d206d30bc3863f275f2dd1cb132a5',
+  'gc_10_3_7_rlc.bin': 'a02585ebe3b36d942e883057119572d9497600c52fc65b8a523487eb65d874f2',
+  'gc_11_0_0_rlc.bin': 'dabd49039772d02f5fd5e48dc21d35ad52a6b1283b470dabca86ca159c4c7c8e',
+  'gc_11_0_1_rlc.bin': '86145719a58e9428562930c6b5ee3b6ced4701d34a80d0b4d84d6026c93134f2',
+  'gc_11_0_2_rlc.bin': 'b43eb2fd0600f50a1a5796bc9983d6b39b5c20960234920f5e89cb362193e0b8',
+  'gc_11_0_3_rlc.bin': '29b0b456f5b53076ddffa6f09de3bb697219e8e7b33504bf6c197e8b858426dc',
+  'gc_11_0_4_rlc.bin': '823573078b608108fbe4dd8176c396ec582632913db9c59a512d82b068f8eba0',
+  'gc_11_5_0_rlc.bin': '68cd85567f4f2f8d6b80db294988806d956bf826979c3597daccb71c7ee6aadd',
+  'gc_11_5_1_rlc.bin': '92731ecabbeb77865fb71787b4268dc738a58779f1190bdc2056482cb88a08f6',
+  'gc_11_5_2_rlc.bin': 'ef3a9209d3eccfbe18fce9e972c146ac283719798bb788096c176b796dc9aee5',
+  'gc_11_5_3_rlc.bin': '10a68940c6258d5818d9c05fd98eb0ccc8d5aee99b2769fbad30e5abd0d9327e',
+  'gc_12_0_0_rlc.bin': '6436b582734a413456fff3d3c7195e71cc9e78a7ed31ee21c83ffd6fae1ad186',
+  'gc_12_0_1_rlc.bin': '6ba4459532246a5c415d3cb33c9b1248294e48f67b827e2accb292a8d1a5c0ec',
+  'gc_9_4_3_rlc.bin': '5345d388712d547b0ae16f199ad5ccadb65643584b3efa7817049ddeb3fdcd12',
+  'gc_9_4_4_rlc.bin': 'e0c3585c72f8136670ca63e607fba32c1ae4948f493f13e33fc4d466bd6318a8',
+  'gc_9_5_0_rlc.bin': '9b1268f5751153fe57f527c9acb417bfa53ed42c9bc083c9d3da2ba61fe5fdc4',
+}

tinygrad/runtime/autogen/am/navi_offsets.py

@@ -0,0 +1,823 @@
+# mypy: disable-error-code="empty-body"
+from __future__ import annotations
+import ctypes
+from typing import Literal, TypeAlias
+from tinygrad.runtime.support.c import _IO, _IOW, _IOR, _IOWR
+from tinygrad.runtime.support import c
+@c.record
+class struct_IP_BASE_INSTANCE(c.Struct):
+  SIZE = 20
+  segment: c.Array[ctypes.c_uint32, Literal[5]]
+struct_IP_BASE_INSTANCE.register_fields([('segment', c.Array[ctypes.c_uint32, Literal[5]], 0)])
+@c.record
+class struct_IP_BASE(c.Struct):
+  SIZE = 140
+  instance: c.Array[struct_IP_BASE_INSTANCE, Literal[7]]
+struct_IP_BASE.register_fields([('instance', c.Array[struct_IP_BASE_INSTANCE, Literal[7]], 0)])
+MAX_INSTANCE = 7
+MAX_SEGMENT = 5
+ATHUB_BASE__INST0_SEG0 = 0x00000C00
+ATHUB_BASE__INST0_SEG1 = 0x02408C00
+ATHUB_BASE__INST0_SEG2 = 0
+ATHUB_BASE__INST0_SEG3 = 0
+ATHUB_BASE__INST0_SEG4 = 0
+ATHUB_BASE__INST1_SEG0 = 0
+ATHUB_BASE__INST1_SEG1 = 0
+ATHUB_BASE__INST1_SEG2 = 0
+ATHUB_BASE__INST1_SEG3 = 0
+ATHUB_BASE__INST1_SEG4 = 0
+ATHUB_BASE__INST2_SEG0 = 0
+ATHUB_BASE__INST2_SEG1 = 0
+ATHUB_BASE__INST2_SEG2 = 0
+ATHUB_BASE__INST2_SEG3 = 0
+ATHUB_BASE__INST2_SEG4 = 0
+ATHUB_BASE__INST3_SEG0 = 0
+ATHUB_BASE__INST3_SEG1 = 0
+ATHUB_BASE__INST3_SEG2 = 0
+ATHUB_BASE__INST3_SEG3 = 0
+ATHUB_BASE__INST3_SEG4 = 0
+ATHUB_BASE__INST4_SEG0 = 0
+ATHUB_BASE__INST4_SEG1 = 0
+ATHUB_BASE__INST4_SEG2 = 0
+ATHUB_BASE__INST4_SEG3 = 0
+ATHUB_BASE__INST4_SEG4 = 0
+ATHUB_BASE__INST5_SEG0 = 0
+ATHUB_BASE__INST5_SEG1 = 0
+ATHUB_BASE__INST5_SEG2 = 0
+ATHUB_BASE__INST5_SEG3 = 0
+ATHUB_BASE__INST5_SEG4 = 0
+ATHUB_BASE__INST6_SEG0 = 0
+ATHUB_BASE__INST6_SEG1 = 0
+ATHUB_BASE__INST6_SEG2 = 0
+ATHUB_BASE__INST6_SEG3 = 0
+ATHUB_BASE__INST6_SEG4 = 0
+CLK_BASE__INST0_SEG0 = 0x00016C00
+CLK_BASE__INST0_SEG1 = 0x02401800
+CLK_BASE__INST0_SEG2 = 0
+CLK_BASE__INST0_SEG3 = 0
+CLK_BASE__INST0_SEG4 = 0
+CLK_BASE__INST1_SEG0 = 0x00016E00
+CLK_BASE__INST1_SEG1 = 0x02401C00
+CLK_BASE__INST1_SEG2 = 0
+CLK_BASE__INST1_SEG3 = 0
+CLK_BASE__INST1_SEG4 = 0
+CLK_BASE__INST2_SEG0 = 0x00017000
+CLK_BASE__INST2_SEG1 = 0x02402000
+CLK_BASE__INST2_SEG2 = 0
+CLK_BASE__INST2_SEG3 = 0
+CLK_BASE__INST2_SEG4 = 0
+CLK_BASE__INST3_SEG0 = 0x00017200
+CLK_BASE__INST3_SEG1 = 0x02402400
+CLK_BASE__INST3_SEG2 = 0
+CLK_BASE__INST3_SEG3 = 0
+CLK_BASE__INST3_SEG4 = 0
+CLK_BASE__INST4_SEG0 = 0x0001B000
+CLK_BASE__INST4_SEG1 = 0x0242D800
+CLK_BASE__INST4_SEG2 = 0
+CLK_BASE__INST4_SEG3 = 0
+CLK_BASE__INST4_SEG4 = 0
+CLK_BASE__INST5_SEG0 = 0x0001B200
+CLK_BASE__INST5_SEG1 = 0x0242DC00
+CLK_BASE__INST5_SEG2 = 0
+CLK_BASE__INST5_SEG3 = 0
+CLK_BASE__INST5_SEG4 = 0
+CLK_BASE__INST6_SEG0 = 0x0001B400
+CLK_BASE__INST6_SEG1 = 0x0242E000
+CLK_BASE__INST6_SEG2 = 0
+CLK_BASE__INST6_SEG3 = 0
+CLK_BASE__INST6_SEG4 = 0
+DF_BASE__INST0_SEG0 = 0x00007000
+DF_BASE__INST0_SEG1 = 0x0240B800
+DF_BASE__INST0_SEG2 = 0
+DF_BASE__INST0_SEG3 = 0
+DF_BASE__INST0_SEG4 = 0
+DF_BASE__INST1_SEG0 = 0
+DF_BASE__INST1_SEG1 = 0
+DF_BASE__INST1_SEG2 = 0
+DF_BASE__INST1_SEG3 = 0
+DF_BASE__INST1_SEG4 = 0
+DF_BASE__INST2_SEG0 = 0
+DF_BASE__INST2_SEG1 = 0
+DF_BASE__INST2_SEG2 = 0
+DF_BASE__INST2_SEG3 = 0
+DF_BASE__INST2_SEG4 = 0
+DF_BASE__INST3_SEG0 = 0
+DF_BASE__INST3_SEG1 = 0
+DF_BASE__INST3_SEG2 = 0
+DF_BASE__INST3_SEG3 = 0
+DF_BASE__INST3_SEG4 = 0
+DF_BASE__INST4_SEG0 = 0
+DF_BASE__INST4_SEG1 = 0
+DF_BASE__INST4_SEG2 = 0
+DF_BASE__INST4_SEG3 = 0
+DF_BASE__INST4_SEG4 = 0
+DF_BASE__INST5_SEG0 = 0
+DF_BASE__INST5_SEG1 = 0
+DF_BASE__INST5_SEG2 = 0
+DF_BASE__INST5_SEG3 = 0
+DF_BASE__INST5_SEG4 = 0
+DF_BASE__INST6_SEG0 = 0
+DF_BASE__INST6_SEG1 = 0
+DF_BASE__INST6_SEG2 = 0
+DF_BASE__INST6_SEG3 = 0
+DF_BASE__INST6_SEG4 = 0
+DIO_BASE__INST0_SEG0 = 0x02404000
+DIO_BASE__INST0_SEG1 = 0
+DIO_BASE__INST0_SEG2 = 0
+DIO_BASE__INST0_SEG3 = 0
+DIO_BASE__INST0_SEG4 = 0
+DIO_BASE__INST1_SEG0 = 0
+DIO_BASE__INST1_SEG1 = 0
+DIO_BASE__INST1_SEG2 = 0
+DIO_BASE__INST1_SEG3 = 0
+DIO_BASE__INST1_SEG4 = 0
+DIO_BASE__INST2_SEG0 = 0
+DIO_BASE__INST2_SEG1 = 0
+DIO_BASE__INST2_SEG2 = 0
+DIO_BASE__INST2_SEG3 = 0
+DIO_BASE__INST2_SEG4 = 0
+DIO_BASE__INST3_SEG0 = 0
+DIO_BASE__INST3_SEG1 = 0
+DIO_BASE__INST3_SEG2 = 0
+DIO_BASE__INST3_SEG3 = 0
+DIO_BASE__INST3_SEG4 = 0
+DIO_BASE__INST4_SEG0 = 0
+DIO_BASE__INST4_SEG1 = 0
+DIO_BASE__INST4_SEG2 = 0
+DIO_BASE__INST4_SEG3 = 0
+DIO_BASE__INST4_SEG4 = 0
+DIO_BASE__INST5_SEG0 = 0
+DIO_BASE__INST5_SEG1 = 0
+DIO_BASE__INST5_SEG2 = 0
+DIO_BASE__INST5_SEG3 = 0
+DIO_BASE__INST5_SEG4 = 0
+DIO_BASE__INST6_SEG0 = 0
+DIO_BASE__INST6_SEG1 = 0
+DIO_BASE__INST6_SEG2 = 0
+DIO_BASE__INST6_SEG3 = 0
+DIO_BASE__INST6_SEG4 = 0
+DCN_BASE__INST0_SEG0 = 0x00000012
+DCN_BASE__INST0_SEG1 = 0x000000C0
+DCN_BASE__INST0_SEG2 = 0x000034C0
+DCN_BASE__INST0_SEG3 = 0x00009000
+DCN_BASE__INST0_SEG4 = 0x02403C00
+DCN_BASE__INST1_SEG0 = 0
+DCN_BASE__INST1_SEG1 = 0
+DCN_BASE__INST1_SEG2 = 0
+DCN_BASE__INST1_SEG3 = 0
+DCN_BASE__INST1_SEG4 = 0
+DCN_BASE__INST2_SEG0 = 0
+DCN_BASE__INST2_SEG1 = 0
+DCN_BASE__INST2_SEG2 = 0
+DCN_BASE__INST2_SEG3 = 0
+DCN_BASE__INST2_SEG4 = 0
+DCN_BASE__INST3_SEG0 = 0
+DCN_BASE__INST3_SEG1 = 0
+DCN_BASE__INST3_SEG2 = 0
+DCN_BASE__INST3_SEG3 = 0
+DCN_BASE__INST3_SEG4 = 0
+DCN_BASE__INST4_SEG0 = 0
+DCN_BASE__INST4_SEG1 = 0
+DCN_BASE__INST4_SEG2 = 0
+DCN_BASE__INST4_SEG3 = 0
+DCN_BASE__INST4_SEG4 = 0
+DCN_BASE__INST5_SEG0 = 0
+DCN_BASE__INST5_SEG1 = 0
+DCN_BASE__INST5_SEG2 = 0
+DCN_BASE__INST5_SEG3 = 0
+DCN_BASE__INST5_SEG4 = 0
+DCN_BASE__INST6_SEG0 = 0
+DCN_BASE__INST6_SEG1 = 0
+DCN_BASE__INST6_SEG2 = 0
+DCN_BASE__INST6_SEG3 = 0
+DCN_BASE__INST6_SEG4 = 0
+DPCS_BASE__INST0_SEG0 = 0x00000012
+DPCS_BASE__INST0_SEG1 = 0x000000C0
+DPCS_BASE__INST0_SEG2 = 0x000034C0
+DPCS_BASE__INST0_SEG3 = 0x00009000
+DPCS_BASE__INST0_SEG4 = 0x02403C00
+DPCS_BASE__INST1_SEG0 = 0
+DPCS_BASE__INST1_SEG1 = 0
+DPCS_BASE__INST1_SEG2 = 0
+DPCS_BASE__INST1_SEG3 = 0
+DPCS_BASE__INST1_SEG4 = 0
+DPCS_BASE__INST2_SEG0 = 0
+DPCS_BASE__INST2_SEG1 = 0
+DPCS_BASE__INST2_SEG2 = 0
+DPCS_BASE__INST2_SEG3 = 0
+DPCS_BASE__INST2_SEG4 = 0
+DPCS_BASE__INST3_SEG0 = 0
+DPCS_BASE__INST3_SEG1 = 0
+DPCS_BASE__INST3_SEG2 = 0
+DPCS_BASE__INST3_SEG3 = 0
+DPCS_BASE__INST3_SEG4 = 0
+DPCS_BASE__INST4_SEG0 = 0
+DPCS_BASE__INST4_SEG1 = 0
+DPCS_BASE__INST4_SEG2 = 0
+DPCS_BASE__INST4_SEG3 = 0
+DPCS_BASE__INST4_SEG4 = 0
+DPCS_BASE__INST5_SEG0 = 0
+DPCS_BASE__INST5_SEG1 = 0
+DPCS_BASE__INST5_SEG2 = 0
+DPCS_BASE__INST5_SEG3 = 0
+DPCS_BASE__INST5_SEG4 = 0
+DPCS_BASE__INST6_SEG0 = 0
+DPCS_BASE__INST6_SEG1 = 0
+DPCS_BASE__INST6_SEG2 = 0
+DPCS_BASE__INST6_SEG3 = 0
+DPCS_BASE__INST6_SEG4 = 0
+FUSE_BASE__INST0_SEG0 = 0x00017400
+FUSE_BASE__INST0_SEG1 = 0x02401400
+FUSE_BASE__INST0_SEG2 = 0
+FUSE_BASE__INST0_SEG3 = 0
+FUSE_BASE__INST0_SEG4 = 0
+FUSE_BASE__INST1_SEG0 = 0
+FUSE_BASE__INST1_SEG1 = 0
+FUSE_BASE__INST1_SEG2 = 0
+FUSE_BASE__INST1_SEG3 = 0
+FUSE_BASE__INST1_SEG4 = 0
+FUSE_BASE__INST2_SEG0 = 0
+FUSE_BASE__INST2_SEG1 = 0
+FUSE_BASE__INST2_SEG2 = 0
+FUSE_BASE__INST2_SEG3 = 0
+FUSE_BASE__INST2_SEG4 = 0
+FUSE_BASE__INST3_SEG0 = 0
+FUSE_BASE__INST3_SEG1 = 0
+FUSE_BASE__INST3_SEG2 = 0
+FUSE_BASE__INST3_SEG3 = 0
+FUSE_BASE__INST3_SEG4 = 0
+FUSE_BASE__INST4_SEG0 = 0
+FUSE_BASE__INST4_SEG1 = 0
+FUSE_BASE__INST4_SEG2 = 0
+FUSE_BASE__INST4_SEG3 = 0
+FUSE_BASE__INST4_SEG4 = 0
+FUSE_BASE__INST5_SEG0 = 0
+FUSE_BASE__INST5_SEG1 = 0
+FUSE_BASE__INST5_SEG2 = 0
+FUSE_BASE__INST5_SEG3 = 0
+FUSE_BASE__INST5_SEG4 = 0
+FUSE_BASE__INST6_SEG0 = 0
+FUSE_BASE__INST6_SEG1 = 0
+FUSE_BASE__INST6_SEG2 = 0
+FUSE_BASE__INST6_SEG3 = 0
+FUSE_BASE__INST6_SEG4 = 0
+GC_BASE__INST0_SEG0 = 0x00001260
+GC_BASE__INST0_SEG1 = 0x0000A000
+GC_BASE__INST0_SEG2 = 0x0001C000
+GC_BASE__INST0_SEG3 = 0x02402C00
+GC_BASE__INST0_SEG4 = 0
+GC_BASE__INST1_SEG0 = 0
+GC_BASE__INST1_SEG1 = 0
+GC_BASE__INST1_SEG2 = 0
+GC_BASE__INST1_SEG3 = 0
+GC_BASE__INST1_SEG4 = 0
+GC_BASE__INST2_SEG0 = 0
+GC_BASE__INST2_SEG1 = 0
+GC_BASE__INST2_SEG2 = 0
+GC_BASE__INST2_SEG3 = 0
+GC_BASE__INST2_SEG4 = 0
+GC_BASE__INST3_SEG0 = 0
+GC_BASE__INST3_SEG1 = 0
+GC_BASE__INST3_SEG2 = 0
+GC_BASE__INST3_SEG3 = 0
+GC_BASE__INST3_SEG4 = 0
+GC_BASE__INST4_SEG0 = 0
+GC_BASE__INST4_SEG1 = 0
+GC_BASE__INST4_SEG2 = 0
+GC_BASE__INST4_SEG3 = 0
+GC_BASE__INST4_SEG4 = 0
+GC_BASE__INST5_SEG0 = 0
+GC_BASE__INST5_SEG1 = 0
+GC_BASE__INST5_SEG2 = 0
+GC_BASE__INST5_SEG3 = 0
+GC_BASE__INST5_SEG4 = 0
+GC_BASE__INST6_SEG0 = 0
+GC_BASE__INST6_SEG1 = 0
+GC_BASE__INST6_SEG2 = 0
+GC_BASE__INST6_SEG3 = 0
+GC_BASE__INST6_SEG4 = 0
+HDA_BASE__INST0_SEG0 = 0x004C0000
+HDA_BASE__INST0_SEG1 = 0x02404800
+HDA_BASE__INST0_SEG2 = 0
+HDA_BASE__INST0_SEG3 = 0
+HDA_BASE__INST0_SEG4 = 0
+HDA_BASE__INST1_SEG0 = 0
+HDA_BASE__INST1_SEG1 = 0
+HDA_BASE__INST1_SEG2 = 0
+HDA_BASE__INST1_SEG3 = 0
+HDA_BASE__INST1_SEG4 = 0
+HDA_BASE__INST2_SEG0 = 0
+HDA_BASE__INST2_SEG1 = 0
+HDA_BASE__INST2_SEG2 = 0
+HDA_BASE__INST2_SEG3 = 0
+HDA_BASE__INST2_SEG4 = 0
+HDA_BASE__INST3_SEG0 = 0
+HDA_BASE__INST3_SEG1 = 0
+HDA_BASE__INST3_SEG2 = 0
+HDA_BASE__INST3_SEG3 = 0
+HDA_BASE__INST3_SEG4 = 0
+HDA_BASE__INST4_SEG0 = 0
+HDA_BASE__INST4_SEG1 = 0
+HDA_BASE__INST4_SEG2 = 0
+HDA_BASE__INST4_SEG3 = 0
+HDA_BASE__INST4_SEG4 = 0
+HDA_BASE__INST5_SEG0 = 0
+HDA_BASE__INST5_SEG1 = 0
+HDA_BASE__INST5_SEG2 = 0
+HDA_BASE__INST5_SEG3 = 0
+HDA_BASE__INST5_SEG4 = 0
+HDA_BASE__INST6_SEG0 = 0
+HDA_BASE__INST6_SEG1 = 0
+HDA_BASE__INST6_SEG2 = 0
+HDA_BASE__INST6_SEG3 = 0
+HDA_BASE__INST6_SEG4 = 0
+HDP_BASE__INST0_SEG0 = 0x00000F20
+HDP_BASE__INST0_SEG1 = 0x0240A400
+HDP_BASE__INST0_SEG2 = 0
+HDP_BASE__INST0_SEG3 = 0
+HDP_BASE__INST0_SEG4 = 0
+HDP_BASE__INST1_SEG0 = 0
+HDP_BASE__INST1_SEG1 = 0
+HDP_BASE__INST1_SEG2 = 0
+HDP_BASE__INST1_SEG3 = 0
+HDP_BASE__INST1_SEG4 = 0
+HDP_BASE__INST2_SEG0 = 0
+HDP_BASE__INST2_SEG1 = 0
+HDP_BASE__INST2_SEG2 = 0
+HDP_BASE__INST2_SEG3 = 0
+HDP_BASE__INST2_SEG4 = 0
+HDP_BASE__INST3_SEG0 = 0
+HDP_BASE__INST3_SEG1 = 0
+HDP_BASE__INST3_SEG2 = 0
+HDP_BASE__INST3_SEG3 = 0
+HDP_BASE__INST3_SEG4 = 0
+HDP_BASE__INST4_SEG0 = 0
+HDP_BASE__INST4_SEG1 = 0
+HDP_BASE__INST4_SEG2 = 0
+HDP_BASE__INST4_SEG3 = 0
+HDP_BASE__INST4_SEG4 = 0
+HDP_BASE__INST5_SEG0 = 0
+HDP_BASE__INST5_SEG1 = 0
+HDP_BASE__INST5_SEG2 = 0
+HDP_BASE__INST5_SEG3 = 0
+HDP_BASE__INST5_SEG4 = 0
+HDP_BASE__INST6_SEG0 = 0
+HDP_BASE__INST6_SEG1 = 0
+HDP_BASE__INST6_SEG2 = 0
+HDP_BASE__INST6_SEG3 = 0
+HDP_BASE__INST6_SEG4 = 0
+MMHUB_BASE__INST0_SEG0 = 0x0001A000
+MMHUB_BASE__INST0_SEG1 = 0x02408800
+MMHUB_BASE__INST0_SEG2 = 0
+MMHUB_BASE__INST0_SEG3 = 0
+MMHUB_BASE__INST0_SEG4 = 0
+MMHUB_BASE__INST1_SEG0 = 0
+MMHUB_BASE__INST1_SEG1 = 0
+MMHUB_BASE__INST1_SEG2 = 0
+MMHUB_BASE__INST1_SEG3 = 0
+MMHUB_BASE__INST1_SEG4 = 0
+MMHUB_BASE__INST2_SEG0 = 0
+MMHUB_BASE__INST2_SEG1 = 0
+MMHUB_BASE__INST2_SEG2 = 0
+MMHUB_BASE__INST2_SEG3 = 0
+MMHUB_BASE__INST2_SEG4 = 0
+MMHUB_BASE__INST3_SEG0 = 0
+MMHUB_BASE__INST3_SEG1 = 0
+MMHUB_BASE__INST3_SEG2 = 0
+MMHUB_BASE__INST3_SEG3 = 0
+MMHUB_BASE__INST3_SEG4 = 0
+MMHUB_BASE__INST4_SEG0 = 0
+MMHUB_BASE__INST4_SEG1 = 0
+MMHUB_BASE__INST4_SEG2 = 0
+MMHUB_BASE__INST4_SEG3 = 0
+MMHUB_BASE__INST4_SEG4 = 0
+MMHUB_BASE__INST5_SEG0 = 0
+MMHUB_BASE__INST5_SEG1 = 0
+MMHUB_BASE__INST5_SEG2 = 0
+MMHUB_BASE__INST5_SEG3 = 0
+MMHUB_BASE__INST5_SEG4 = 0
+MMHUB_BASE__INST6_SEG0 = 0
+MMHUB_BASE__INST6_SEG1 = 0
+MMHUB_BASE__INST6_SEG2 = 0
+MMHUB_BASE__INST6_SEG3 = 0
+MMHUB_BASE__INST6_SEG4 = 0
+MP0_BASE__INST0_SEG0 = 0x00016000
+MP0_BASE__INST0_SEG1 = 0x00DC0000
+MP0_BASE__INST0_SEG2 = 0x00E00000
+MP0_BASE__INST0_SEG3 = 0x00E40000
+MP0_BASE__INST0_SEG4 = 0x0243FC00
+MP0_BASE__INST1_SEG0 = 0
+MP0_BASE__INST1_SEG1 = 0
+MP0_BASE__INST1_SEG2 = 0
+MP0_BASE__INST1_SEG3 = 0
+MP0_BASE__INST1_SEG4 = 0
+MP0_BASE__INST2_SEG0 = 0
+MP0_BASE__INST2_SEG1 = 0
+MP0_BASE__INST2_SEG2 = 0
+MP0_BASE__INST2_SEG3 = 0
+MP0_BASE__INST2_SEG4 = 0
+MP0_BASE__INST3_SEG0 = 0
+MP0_BASE__INST3_SEG1 = 0
+MP0_BASE__INST3_SEG2 = 0
+MP0_BASE__INST3_SEG3 = 0
+MP0_BASE__INST3_SEG4 = 0
+MP0_BASE__INST4_SEG0 = 0
+MP0_BASE__INST4_SEG1 = 0
+MP0_BASE__INST4_SEG2 = 0
+MP0_BASE__INST4_SEG3 = 0
+MP0_BASE__INST4_SEG4 = 0
+MP0_BASE__INST5_SEG0 = 0
+MP0_BASE__INST5_SEG1 = 0
+MP0_BASE__INST5_SEG2 = 0
+MP0_BASE__INST5_SEG3 = 0
+MP0_BASE__INST5_SEG4 = 0
+MP0_BASE__INST6_SEG0 = 0
+MP0_BASE__INST6_SEG1 = 0
+MP0_BASE__INST6_SEG2 = 0
+MP0_BASE__INST6_SEG3 = 0
+MP0_BASE__INST6_SEG4 = 0
+MP1_BASE__INST0_SEG0 = 0x00016000
+MP1_BASE__INST0_SEG1 = 0x00DC0000
+MP1_BASE__INST0_SEG2 = 0x00E00000
+MP1_BASE__INST0_SEG3 = 0x00E40000
+MP1_BASE__INST0_SEG4 = 0x0243FC00
+MP1_BASE__INST1_SEG0 = 0
+MP1_BASE__INST1_SEG1 = 0
+MP1_BASE__INST1_SEG2 = 0
+MP1_BASE__INST1_SEG3 = 0
+MP1_BASE__INST1_SEG4 = 0
+MP1_BASE__INST2_SEG0 = 0
+MP1_BASE__INST2_SEG1 = 0
+MP1_BASE__INST2_SEG2 = 0
+MP1_BASE__INST2_SEG3 = 0
+MP1_BASE__INST2_SEG4 = 0
+MP1_BASE__INST3_SEG0 = 0
+MP1_BASE__INST3_SEG1 = 0
+MP1_BASE__INST3_SEG2 = 0
+MP1_BASE__INST3_SEG3 = 0
+MP1_BASE__INST3_SEG4 = 0
+MP1_BASE__INST4_SEG0 = 0
+MP1_BASE__INST4_SEG1 = 0
+MP1_BASE__INST4_SEG2 = 0
+MP1_BASE__INST4_SEG3 = 0
+MP1_BASE__INST4_SEG4 = 0
+MP1_BASE__INST5_SEG0 = 0
+MP1_BASE__INST5_SEG1 = 0
+MP1_BASE__INST5_SEG2 = 0
+MP1_BASE__INST5_SEG3 = 0
+MP1_BASE__INST5_SEG4 = 0
+MP1_BASE__INST6_SEG0 = 0
+MP1_BASE__INST6_SEG1 = 0
+MP1_BASE__INST6_SEG2 = 0
+MP1_BASE__INST6_SEG3 = 0
+MP1_BASE__INST6_SEG4 = 0
+NBIO_BASE__INST0_SEG0 = 0x00000000
+NBIO_BASE__INST0_SEG1 = 0x00000014
+NBIO_BASE__INST0_SEG2 = 0x00000D20
+NBIO_BASE__INST0_SEG3 = 0x00010400
+NBIO_BASE__INST0_SEG4 = 0x0241B000
+NBIO_BASE__INST1_SEG0 = 0
+NBIO_BASE__INST1_SEG1 = 0
+NBIO_BASE__INST1_SEG2 = 0
+NBIO_BASE__INST1_SEG3 = 0
+NBIO_BASE__INST1_SEG4 = 0
+NBIO_BASE__INST2_SEG0 = 0
+NBIO_BASE__INST2_SEG1 = 0
+NBIO_BASE__INST2_SEG2 = 0
+NBIO_BASE__INST2_SEG3 = 0
+NBIO_BASE__INST2_SEG4 = 0
+NBIO_BASE__INST3_SEG0 = 0
+NBIO_BASE__INST3_SEG1 = 0
+NBIO_BASE__INST3_SEG2 = 0
+NBIO_BASE__INST3_SEG3 = 0
+NBIO_BASE__INST3_SEG4 = 0
+NBIO_BASE__INST4_SEG0 = 0
+NBIO_BASE__INST4_SEG1 = 0
+NBIO_BASE__INST4_SEG2 = 0
+NBIO_BASE__INST4_SEG3 = 0
+NBIO_BASE__INST4_SEG4 = 0
+NBIO_BASE__INST5_SEG0 = 0
+NBIO_BASE__INST5_SEG1 = 0
+NBIO_BASE__INST5_SEG2 = 0
+NBIO_BASE__INST5_SEG3 = 0
+NBIO_BASE__INST5_SEG4 = 0
+NBIO_BASE__INST6_SEG0 = 0
+NBIO_BASE__INST6_SEG1 = 0
+NBIO_BASE__INST6_SEG2 = 0
+NBIO_BASE__INST6_SEG3 = 0
+NBIO_BASE__INST6_SEG4 = 0
+OSSSYS_BASE__INST0_SEG0 = 0x000010A0
+OSSSYS_BASE__INST0_SEG1 = 0x0240A000
+OSSSYS_BASE__INST0_SEG2 = 0
+OSSSYS_BASE__INST0_SEG3 = 0
+OSSSYS_BASE__INST0_SEG4 = 0
+OSSSYS_BASE__INST1_SEG0 = 0
+OSSSYS_BASE__INST1_SEG1 = 0
+OSSSYS_BASE__INST1_SEG2 = 0
+OSSSYS_BASE__INST1_SEG3 = 0
+OSSSYS_BASE__INST1_SEG4 = 0
+OSSSYS_BASE__INST2_SEG0 = 0
+OSSSYS_BASE__INST2_SEG1 = 0
+OSSSYS_BASE__INST2_SEG2 = 0
+OSSSYS_BASE__INST2_SEG3 = 0
+OSSSYS_BASE__INST2_SEG4 = 0
+OSSSYS_BASE__INST3_SEG0 = 0
+OSSSYS_BASE__INST3_SEG1 = 0
+OSSSYS_BASE__INST3_SEG2 = 0
+OSSSYS_BASE__INST3_SEG3 = 0
+OSSSYS_BASE__INST3_SEG4 = 0
+OSSSYS_BASE__INST4_SEG0 = 0
+OSSSYS_BASE__INST4_SEG1 = 0
+OSSSYS_BASE__INST4_SEG2 = 0
+OSSSYS_BASE__INST4_SEG3 = 0
+OSSSYS_BASE__INST4_SEG4 = 0
+OSSSYS_BASE__INST5_SEG0 = 0
+OSSSYS_BASE__INST5_SEG1 = 0
+OSSSYS_BASE__INST5_SEG2 = 0
+OSSSYS_BASE__INST5_SEG3 = 0
+OSSSYS_BASE__INST5_SEG4 = 0
+OSSSYS_BASE__INST6_SEG0 = 0
+OSSSYS_BASE__INST6_SEG1 = 0
+OSSSYS_BASE__INST6_SEG2 = 0
+OSSSYS_BASE__INST6_SEG3 = 0
+OSSSYS_BASE__INST6_SEG4 = 0
+PCIE0_BASE__INST0_SEG0 = 0x00000000
+PCIE0_BASE__INST0_SEG1 = 0x00000014
+PCIE0_BASE__INST0_SEG2 = 0x00000D20
+PCIE0_BASE__INST0_SEG3 = 0x00010400
+PCIE0_BASE__INST0_SEG4 = 0x0241B000
+PCIE0_BASE__INST1_SEG0 = 0
+PCIE0_BASE__INST1_SEG1 = 0
+PCIE0_BASE__INST1_SEG2 = 0
+PCIE0_BASE__INST1_SEG3 = 0
+PCIE0_BASE__INST1_SEG4 = 0
+PCIE0_BASE__INST2_SEG0 = 0
+PCIE0_BASE__INST2_SEG1 = 0
+PCIE0_BASE__INST2_SEG2 = 0
+PCIE0_BASE__INST2_SEG3 = 0
+PCIE0_BASE__INST2_SEG4 = 0
+PCIE0_BASE__INST3_SEG0 = 0
+PCIE0_BASE__INST3_SEG1 = 0
+PCIE0_BASE__INST3_SEG2 = 0
+PCIE0_BASE__INST3_SEG3 = 0
+PCIE0_BASE__INST3_SEG4 = 0
+PCIE0_BASE__INST4_SEG0 = 0
+PCIE0_BASE__INST4_SEG1 = 0
+PCIE0_BASE__INST4_SEG2 = 0
+PCIE0_BASE__INST4_SEG3 = 0
+PCIE0_BASE__INST4_SEG4 = 0
+PCIE0_BASE__INST5_SEG0 = 0
+PCIE0_BASE__INST5_SEG1 = 0
+PCIE0_BASE__INST5_SEG2 = 0
+PCIE0_BASE__INST5_SEG3 = 0
+PCIE0_BASE__INST5_SEG4 = 0
+PCIE0_BASE__INST6_SEG0 = 0
+PCIE0_BASE__INST6_SEG1 = 0
+PCIE0_BASE__INST6_SEG2 = 0
+PCIE0_BASE__INST6_SEG3 = 0
+PCIE0_BASE__INST6_SEG4 = 0
+SDMA0_BASE__INST0_SEG0 = 0x00001260
+SDMA0_BASE__INST0_SEG1 = 0x0000A000
+SDMA0_BASE__INST0_SEG2 = 0x0001C000
+SDMA0_BASE__INST0_SEG3 = 0x02402C00
+SDMA0_BASE__INST0_SEG4 = 0
+SDMA0_BASE__INST1_SEG0 = 0
+SDMA0_BASE__INST1_SEG1 = 0
+SDMA0_BASE__INST1_SEG2 = 0
+SDMA0_BASE__INST1_SEG3 = 0
+SDMA0_BASE__INST1_SEG4 = 0
+SDMA0_BASE__INST2_SEG0 = 0
+SDMA0_BASE__INST2_SEG1 = 0
+SDMA0_BASE__INST2_SEG2 = 0
+SDMA0_BASE__INST2_SEG3 = 0
+SDMA0_BASE__INST2_SEG4 = 0
+SDMA0_BASE__INST3_SEG0 = 0
+SDMA0_BASE__INST3_SEG1 = 0
+SDMA0_BASE__INST3_SEG2 = 0
+SDMA0_BASE__INST3_SEG3 = 0
+SDMA0_BASE__INST3_SEG4 = 0
+SDMA0_BASE__INST4_SEG0 = 0
+SDMA0_BASE__INST4_SEG1 = 0
+SDMA0_BASE__INST4_SEG2 = 0
+SDMA0_BASE__INST4_SEG3 = 0
+SDMA0_BASE__INST4_SEG4 = 0
+SDMA0_BASE__INST5_SEG0 = 0
+SDMA0_BASE__INST5_SEG1 = 0
+SDMA0_BASE__INST5_SEG2 = 0
+SDMA0_BASE__INST5_SEG3 = 0
+SDMA0_BASE__INST5_SEG4 = 0
+SDMA0_BASE__INST6_SEG0 = 0
+SDMA0_BASE__INST6_SEG1 = 0
+SDMA0_BASE__INST6_SEG2 = 0
+SDMA0_BASE__INST6_SEG3 = 0
+SDMA0_BASE__INST6_SEG4 = 0
+SDMA1_BASE__INST0_SEG0 = 0x00001260
+SDMA1_BASE__INST0_SEG1 = 0x0000A000
+SDMA1_BASE__INST0_SEG2 = 0x0001C000
+SDMA1_BASE__INST0_SEG3 = 0x02402C00
+SDMA1_BASE__INST0_SEG4 = 0
+SDMA1_BASE__INST1_SEG0 = 0
+SDMA1_BASE__INST1_SEG1 = 0
+SDMA1_BASE__INST1_SEG2 = 0
+SDMA1_BASE__INST1_SEG3 = 0
+SDMA1_BASE__INST1_SEG4 = 0
+SDMA1_BASE__INST2_SEG0 = 0
+SDMA1_BASE__INST2_SEG1 = 0
+SDMA1_BASE__INST2_SEG2 = 0
+SDMA1_BASE__INST2_SEG3 = 0
+SDMA1_BASE__INST2_SEG4 = 0
+SDMA1_BASE__INST3_SEG0 = 0
+SDMA1_BASE__INST3_SEG1 = 0
+SDMA1_BASE__INST3_SEG2 = 0
+SDMA1_BASE__INST3_SEG3 = 0
+SDMA1_BASE__INST3_SEG4 = 0
+SDMA1_BASE__INST4_SEG0 = 0
+SDMA1_BASE__INST4_SEG1 = 0
+SDMA1_BASE__INST4_SEG2 = 0
+SDMA1_BASE__INST4_SEG3 = 0
+SDMA1_BASE__INST4_SEG4 = 0
+SDMA1_BASE__INST5_SEG0 = 0
+SDMA1_BASE__INST5_SEG1 = 0
+SDMA1_BASE__INST5_SEG2 = 0
+SDMA1_BASE__INST5_SEG3 = 0
+SDMA1_BASE__INST5_SEG4 = 0
+SDMA1_BASE__INST6_SEG0 = 0
+SDMA1_BASE__INST6_SEG1 = 0
+SDMA1_BASE__INST6_SEG2 = 0
+SDMA1_BASE__INST6_SEG3 = 0
+SDMA1_BASE__INST6_SEG4 = 0
+SMUIO_BASE__INST0_SEG0 = 0x00016800
+SMUIO_BASE__INST0_SEG1 = 0x00016A00
+SMUIO_BASE__INST0_SEG2 = 0x00440000
+SMUIO_BASE__INST0_SEG3 = 0x02401000
+SMUIO_BASE__INST0_SEG4 = 0
+SMUIO_BASE__INST1_SEG0 = 0
+SMUIO_BASE__INST1_SEG1 = 0
+SMUIO_BASE__INST1_SEG2 = 0
+SMUIO_BASE__INST1_SEG3 = 0
+SMUIO_BASE__INST1_SEG4 = 0
+SMUIO_BASE__INST2_SEG0 = 0
+SMUIO_BASE__INST2_SEG1 = 0
+SMUIO_BASE__INST2_SEG2 = 0
+SMUIO_BASE__INST2_SEG3 = 0
+SMUIO_BASE__INST2_SEG4 = 0
+SMUIO_BASE__INST3_SEG0 = 0
+SMUIO_BASE__INST3_SEG1 = 0
+SMUIO_BASE__INST3_SEG2 = 0
+SMUIO_BASE__INST3_SEG3 = 0
+SMUIO_BASE__INST3_SEG4 = 0
+SMUIO_BASE__INST4_SEG0 = 0
+SMUIO_BASE__INST4_SEG1 = 0
+SMUIO_BASE__INST4_SEG2 = 0
+SMUIO_BASE__INST4_SEG3 = 0
+SMUIO_BASE__INST4_SEG4 = 0
+SMUIO_BASE__INST5_SEG0 = 0
+SMUIO_BASE__INST5_SEG1 = 0
+SMUIO_BASE__INST5_SEG2 = 0
+SMUIO_BASE__INST5_SEG3 = 0
+SMUIO_BASE__INST5_SEG4 = 0
+SMUIO_BASE__INST6_SEG0 = 0
+SMUIO_BASE__INST6_SEG1 = 0
+SMUIO_BASE__INST6_SEG2 = 0
+SMUIO_BASE__INST6_SEG3 = 0
+SMUIO_BASE__INST6_SEG4 = 0
+THM_BASE__INST0_SEG0 = 0x00016600
+THM_BASE__INST0_SEG1 = 0x02400C00
+THM_BASE__INST0_SEG2 = 0
+THM_BASE__INST0_SEG3 = 0
+THM_BASE__INST0_SEG4 = 0
+THM_BASE__INST1_SEG0 = 0
+THM_BASE__INST1_SEG1 = 0
+THM_BASE__INST1_SEG2 = 0
+THM_BASE__INST1_SEG3 = 0
+THM_BASE__INST1_SEG4 = 0
+THM_BASE__INST2_SEG0 = 0
+THM_BASE__INST2_SEG1 = 0
+THM_BASE__INST2_SEG2 = 0
+THM_BASE__INST2_SEG3 = 0
+THM_BASE__INST2_SEG4 = 0
+THM_BASE__INST3_SEG0 = 0
+THM_BASE__INST3_SEG1 = 0
+THM_BASE__INST3_SEG2 = 0
+THM_BASE__INST3_SEG3 = 0
+THM_BASE__INST3_SEG4 = 0
+THM_BASE__INST4_SEG0 = 0
+THM_BASE__INST4_SEG1 = 0
+THM_BASE__INST4_SEG2 = 0
+THM_BASE__INST4_SEG3 = 0
+THM_BASE__INST4_SEG4 = 0
+THM_BASE__INST5_SEG0 = 0
+THM_BASE__INST5_SEG1 = 0
+THM_BASE__INST5_SEG2 = 0
+THM_BASE__INST5_SEG3 = 0
+THM_BASE__INST5_SEG4 = 0
+THM_BASE__INST6_SEG0 = 0
+THM_BASE__INST6_SEG1 = 0
+THM_BASE__INST6_SEG2 = 0
+THM_BASE__INST6_SEG3 = 0
+THM_BASE__INST6_SEG4 = 0
+UMC_BASE__INST0_SEG0 = 0x00014000
+UMC_BASE__INST0_SEG1 = 0x02425800
+UMC_BASE__INST0_SEG2 = 0
+UMC_BASE__INST0_SEG3 = 0
+UMC_BASE__INST0_SEG4 = 0
+UMC_BASE__INST1_SEG0 = 0x00054000
+UMC_BASE__INST1_SEG1 = 0x02425C00
+UMC_BASE__INST1_SEG2 = 0
+UMC_BASE__INST1_SEG3 = 0
+UMC_BASE__INST1_SEG4 = 0
+UMC_BASE__INST2_SEG0 = 0x00094000
+UMC_BASE__INST2_SEG1 = 0x02426000
+UMC_BASE__INST2_SEG2 = 0
+UMC_BASE__INST2_SEG3 = 0
+UMC_BASE__INST2_SEG4 = 0
+UMC_BASE__INST3_SEG0 = 0x000D4000
+UMC_BASE__INST3_SEG1 = 0x02426400
+UMC_BASE__INST3_SEG2 = 0
+UMC_BASE__INST3_SEG3 = 0
+UMC_BASE__INST3_SEG4 = 0
+UMC_BASE__INST4_SEG0 = 0x00114000
+UMC_BASE__INST4_SEG1 = 0x02426800
+UMC_BASE__INST4_SEG2 = 0
+UMC_BASE__INST4_SEG3 = 0
+UMC_BASE__INST4_SEG4 = 0
+UMC_BASE__INST5_SEG0 = 0x00154000
+UMC_BASE__INST5_SEG1 = 0x02426C00
+UMC_BASE__INST5_SEG2 = 0
+UMC_BASE__INST5_SEG3 = 0
+UMC_BASE__INST5_SEG4 = 0
+UMC_BASE__INST6_SEG0 = 0x00194000
+UMC_BASE__INST6_SEG1 = 0x02427000
+UMC_BASE__INST6_SEG2 = 0
+UMC_BASE__INST6_SEG3 = 0
+UMC_BASE__INST6_SEG4 = 0
+USB0_BASE__INST0_SEG0 = 0x0242A800
+USB0_BASE__INST0_SEG1 = 0x05B00000
+USB0_BASE__INST0_SEG2 = 0
+USB0_BASE__INST0_SEG3 = 0
+USB0_BASE__INST0_SEG4 = 0
+USB0_BASE__INST1_SEG0 = 0
+USB0_BASE__INST1_SEG1 = 0
+USB0_BASE__INST1_SEG2 = 0
+USB0_BASE__INST1_SEG3 = 0
+USB0_BASE__INST1_SEG4 = 0
+USB0_BASE__INST2_SEG0 = 0
+USB0_BASE__INST2_SEG1 = 0
+USB0_BASE__INST2_SEG2 = 0
+USB0_BASE__INST2_SEG3 = 0
+USB0_BASE__INST2_SEG4 = 0
+USB0_BASE__INST3_SEG0 = 0
+USB0_BASE__INST3_SEG1 = 0
+USB0_BASE__INST3_SEG2 = 0
+USB0_BASE__INST3_SEG3 = 0
+USB0_BASE__INST3_SEG4 = 0
+USB0_BASE__INST4_SEG0 = 0
+USB0_BASE__INST4_SEG1 = 0
+USB0_BASE__INST4_SEG2 = 0
+USB0_BASE__INST4_SEG3 = 0
+USB0_BASE__INST4_SEG4 = 0
+USB0_BASE__INST5_SEG0 = 0
+USB0_BASE__INST5_SEG1 = 0
+USB0_BASE__INST5_SEG2 = 0
+USB0_BASE__INST5_SEG3 = 0
+USB0_BASE__INST5_SEG4 = 0
+USB0_BASE__INST6_SEG0 = 0
+USB0_BASE__INST6_SEG1 = 0
+USB0_BASE__INST6_SEG2 = 0
+USB0_BASE__INST6_SEG3 = 0
+USB0_BASE__INST6_SEG4 = 0
+VCN_BASE__INST0_SEG0 = 0x00007800
+VCN_BASE__INST0_SEG1 = 0x00007E00
+VCN_BASE__INST0_SEG2 = 0x02403000
+VCN_BASE__INST0_SEG3 = 0
+VCN_BASE__INST0_SEG4 = 0
+VCN_BASE__INST1_SEG0 = 0x00007B00
+VCN_BASE__INST1_SEG1 = 0x00012000
+VCN_BASE__INST1_SEG2 = 0x02445000
+VCN_BASE__INST1_SEG3 = 0
+VCN_BASE__INST1_SEG4 = 0
+VCN_BASE__INST2_SEG0 = 0
+VCN_BASE__INST2_SEG1 = 0
+VCN_BASE__INST2_SEG2 = 0
+VCN_BASE__INST2_SEG3 = 0
+VCN_BASE__INST2_SEG4 = 0
+VCN_BASE__INST3_SEG0 = 0
+VCN_BASE__INST3_SEG1 = 0
+VCN_BASE__INST3_SEG2 = 0
+VCN_BASE__INST3_SEG3 = 0
+VCN_BASE__INST3_SEG4 = 0
+VCN_BASE__INST4_SEG0 = 0
+VCN_BASE__INST4_SEG1 = 0
+VCN_BASE__INST4_SEG2 = 0
+VCN_BASE__INST4_SEG3 = 0
+VCN_BASE__INST4_SEG4 = 0
+VCN_BASE__INST5_SEG0 = 0
+VCN_BASE__INST5_SEG1 = 0
+VCN_BASE__INST5_SEG2 = 0
+VCN_BASE__INST5_SEG3 = 0
+VCN_BASE__INST5_SEG4 = 0
+VCN_BASE__INST6_SEG0 = 0
+VCN_BASE__INST6_SEG1 = 0
+VCN_BASE__INST6_SEG2 = 0
+VCN_BASE__INST6_SEG3 = 0
+VCN_BASE__INST6_SEG4 = 0

tinygrad/runtime/autogen/am/pm4_nv.py

@@ -37,33 +37,33 @@ enum_WRITE_DATA_wr_confirm_enum: dict[int, str] = {(wr_confirm___write_data__do_
 enum_WRITE_DATA_cache_policy_enum: dict[int, str] = {(cache_policy___write_data__lru:=0): 'cache_policy___write_data__lru', (cache_policy___write_data__stream:=1): 'cache_policy___write_data__stream'}
 class struct_pm4_mec_write_data_mmio(c.Struct): pass
 _anonenum0: dict[int, str] = {(CACHE_FLUSH_AND_INV_TS_EVENT:=20): 'CACHE_FLUSH_AND_INV_TS_EVENT'}
-PACKET_TYPE0 = 0 # type: ignore
-PACKET_TYPE1 = 1 # type: ignore
-PACKET_TYPE2 = 2 # type: ignore
-PACKET_TYPE3 = 3 # type: ignore
+PACKET_TYPE0 = 0
+PACKET_TYPE1 = 1
+PACKET_TYPE2 = 2
+PACKET_TYPE3 = 3
 CP_PACKET_GET_TYPE = lambda h: (((h) >> 30) & 3) # type: ignore
 CP_PACKET_GET_COUNT = lambda h: (((h) >> 16) & 0x3FFF) # type: ignore
 CP_PACKET0_GET_REG = lambda h: ((h) & 0xFFFF) # type: ignore
 CP_PACKET3_GET_OPCODE = lambda h: (((h) >> 8) & 0xFF) # type: ignore
 PACKET0 = lambda reg,n: ((PACKET_TYPE0 << 30) | ((reg) & 0xFFFF) | ((n) & 0x3FFF) << 16) # type: ignore
-CP_PACKET2 = 0x80000000 # type: ignore
-PACKET2_PAD_SHIFT = 0 # type: ignore
-PACKET2_PAD_MASK = (0x3fffffff << 0) # type: ignore
+CP_PACKET2 = 0x80000000
+PACKET2_PAD_SHIFT = 0
+PACKET2_PAD_MASK = (0x3fffffff << 0)
 PACKET2 = lambda v: (CP_PACKET2 | REG_SET(PACKET2_PAD, (v))) # type: ignore
 PACKET3 = lambda op,n: ((PACKET_TYPE3 << 30) | (((op) & 0xFF) << 8) | ((n) & 0x3FFF) << 16) # type: ignore
 PACKET3_COMPUTE = lambda op,n: (PACKET3(op, n) | 1 << 1) # type: ignore
-PACKET3_NOP = 0x10 # type: ignore
-PACKET3_SET_BASE = 0x11 # type: ignore
+PACKET3_NOP = 0x10
+PACKET3_SET_BASE = 0x11
 PACKET3_BASE_INDEX = lambda x: ((x) << 0) # type: ignore
-CE_PARTITION_BASE = 3 # type: ignore
-PACKET3_CLEAR_STATE = 0x12 # type: ignore
-PACKET3_INDEX_BUFFER_SIZE = 0x13 # type: ignore
-PACKET3_DISPATCH_DIRECT = 0x15 # type: ignore
-PACKET3_DISPATCH_INDIRECT = 0x16 # type: ignore
-PACKET3_INDIRECT_BUFFER_END = 0x17 # type: ignore
-PACKET3_INDIRECT_BUFFER_CNST_END = 0x19 # type: ignore
-PACKET3_ATOMIC_GDS = 0x1D # type: ignore
-PACKET3_ATOMIC_MEM = 0x1E # type: ignore
+CE_PARTITION_BASE = 3
+PACKET3_CLEAR_STATE = 0x12
+PACKET3_INDEX_BUFFER_SIZE = 0x13
+PACKET3_DISPATCH_DIRECT = 0x15
+PACKET3_DISPATCH_INDIRECT = 0x16
+PACKET3_INDIRECT_BUFFER_END = 0x17
+PACKET3_INDIRECT_BUFFER_CNST_END = 0x19
+PACKET3_ATOMIC_GDS = 0x1D
+PACKET3_ATOMIC_MEM = 0x1E
 PACKET3_ATOMIC_MEM__ATOMIC = lambda x: ((((unsigned)(x)) & 0x7F) << 0) # type: ignore
 PACKET3_ATOMIC_MEM__COMMAND = lambda x: ((((unsigned)(x)) & 0xF) << 8) # type: ignore
 PACKET3_ATOMIC_MEM__CACHE_POLICY = lambda x: ((((unsigned)(x)) & 0x3) << 25) # type: ignore
@@ -74,39 +74,39 @@ PACKET3_ATOMIC_MEM__SRC_DATA_HI = lambda x: (((unsigned)(x))) # type: ignore
 PACKET3_ATOMIC_MEM__CMP_DATA_LO = lambda x: (((unsigned)(x))) # type: ignore
 PACKET3_ATOMIC_MEM__CMP_DATA_HI = lambda x: (((unsigned)(x))) # type: ignore
 PACKET3_ATOMIC_MEM__LOOP_INTERVAL = lambda x: ((((unsigned)(x)) & 0x1FFF) << 0) # type: ignore
-PACKET3_ATOMIC_MEM__COMMAND__SINGLE_PASS_ATOMIC = 0 # type: ignore
-PACKET3_ATOMIC_MEM__COMMAND__LOOP_UNTIL_COMPARE_SATISFIED = 1 # type: ignore
-PACKET3_ATOMIC_MEM__COMMAND__WAIT_FOR_WRITE_CONFIRMATION = 2 # type: ignore
-PACKET3_ATOMIC_MEM__COMMAND__SEND_AND_CONTINUE = 3 # type: ignore
-PACKET3_ATOMIC_MEM__CACHE_POLICY__LRU = 0 # type: ignore
-PACKET3_ATOMIC_MEM__CACHE_POLICY__STREAM = 1 # type: ignore
-PACKET3_ATOMIC_MEM__CACHE_POLICY__NOA = 2 # type: ignore
-PACKET3_ATOMIC_MEM__CACHE_POLICY__BYPASS = 3 # type: ignore
-PACKET3_OCCLUSION_QUERY = 0x1F # type: ignore
-PACKET3_SET_PREDICATION = 0x20 # type: ignore
-PACKET3_REG_RMW = 0x21 # type: ignore
-PACKET3_COND_EXEC = 0x22 # type: ignore
-PACKET3_PRED_EXEC = 0x23 # type: ignore
-PACKET3_DRAW_INDIRECT = 0x24 # type: ignore
-PACKET3_DRAW_INDEX_INDIRECT = 0x25 # type: ignore
-PACKET3_INDEX_BASE = 0x26 # type: ignore
-PACKET3_DRAW_INDEX_2 = 0x27 # type: ignore
-PACKET3_CONTEXT_CONTROL = 0x28 # type: ignore
-PACKET3_INDEX_TYPE = 0x2A # type: ignore
-PACKET3_DRAW_INDIRECT_MULTI = 0x2C # type: ignore
-PACKET3_DRAW_INDEX_AUTO = 0x2D # type: ignore
-PACKET3_NUM_INSTANCES = 0x2F # type: ignore
-PACKET3_DRAW_INDEX_MULTI_AUTO = 0x30 # type: ignore
-PACKET3_INDIRECT_BUFFER_PRIV = 0x32 # type: ignore
-PACKET3_INDIRECT_BUFFER_CNST = 0x33 # type: ignore
-PACKET3_COND_INDIRECT_BUFFER_CNST = 0x33 # type: ignore
-PACKET3_STRMOUT_BUFFER_UPDATE = 0x34 # type: ignore
-PACKET3_DRAW_INDEX_OFFSET_2 = 0x35 # type: ignore
-PACKET3_DRAW_PREAMBLE = 0x36 # type: ignore
-PACKET3_WRITE_DATA = 0x37 # type: ignore
+PACKET3_ATOMIC_MEM__COMMAND__SINGLE_PASS_ATOMIC = 0
+PACKET3_ATOMIC_MEM__COMMAND__LOOP_UNTIL_COMPARE_SATISFIED = 1
+PACKET3_ATOMIC_MEM__COMMAND__WAIT_FOR_WRITE_CONFIRMATION = 2
+PACKET3_ATOMIC_MEM__COMMAND__SEND_AND_CONTINUE = 3
+PACKET3_ATOMIC_MEM__CACHE_POLICY__LRU = 0
+PACKET3_ATOMIC_MEM__CACHE_POLICY__STREAM = 1
+PACKET3_ATOMIC_MEM__CACHE_POLICY__NOA = 2
+PACKET3_ATOMIC_MEM__CACHE_POLICY__BYPASS = 3
+PACKET3_OCCLUSION_QUERY = 0x1F
+PACKET3_SET_PREDICATION = 0x20
+PACKET3_REG_RMW = 0x21
+PACKET3_COND_EXEC = 0x22
+PACKET3_PRED_EXEC = 0x23
+PACKET3_DRAW_INDIRECT = 0x24
+PACKET3_DRAW_INDEX_INDIRECT = 0x25
+PACKET3_INDEX_BASE = 0x26
+PACKET3_DRAW_INDEX_2 = 0x27
+PACKET3_CONTEXT_CONTROL = 0x28
+PACKET3_INDEX_TYPE = 0x2A
+PACKET3_DRAW_INDIRECT_MULTI = 0x2C
+PACKET3_DRAW_INDEX_AUTO = 0x2D
+PACKET3_NUM_INSTANCES = 0x2F
+PACKET3_DRAW_INDEX_MULTI_AUTO = 0x30
+PACKET3_INDIRECT_BUFFER_PRIV = 0x32
+PACKET3_INDIRECT_BUFFER_CNST = 0x33
+PACKET3_COND_INDIRECT_BUFFER_CNST = 0x33
+PACKET3_STRMOUT_BUFFER_UPDATE = 0x34
+PACKET3_DRAW_INDEX_OFFSET_2 = 0x35
+PACKET3_DRAW_PREAMBLE = 0x36
+PACKET3_WRITE_DATA = 0x37
 WRITE_DATA_DST_SEL = lambda x: ((x) << 8) # type: ignore
-WR_ONE_ADDR = (1 << 16) # type: ignore
-WR_CONFIRM = (1 << 20) # type: ignore
+WR_ONE_ADDR = (1 << 16)
+WR_CONFIRM = (1 << 20)
 WRITE_DATA_CACHE_POLICY = lambda x: ((x) << 25) # type: ignore
 WRITE_DATA_ENGINE_SEL = lambda x: ((x) << 30) # type: ignore
 PACKET3_WRITE_DATA__DST_SEL = lambda x: ((((unsigned)(x)) & 0xF) << 8) # type: ignore
@@ -122,34 +122,34 @@ PACKET3_WRITE_DATA__AID_ID = lambda x: ((((unsigned)(x)) & 0x3) << 22) # type: i
 PACKET3_WRITE_DATA__TEMPORAL = lambda x: ((((unsigned)(x)) & 0x3) << 24) # type: ignore
 PACKET3_WRITE_DATA__DST_MMREG_ADDR_LO = lambda x: ((unsigned)(x)) # type: ignore
 PACKET3_WRITE_DATA__DST_MMREG_ADDR_HI = lambda x: ((((unsigned)(x)) & 0xFF) << 0) # type: ignore
-PACKET3_WRITE_DATA__DST_SEL__MEM_MAPPED_REGISTER = 0 # type: ignore
-PACKET3_WRITE_DATA__DST_SEL__TC_L2 = 2 # type: ignore
-PACKET3_WRITE_DATA__DST_SEL__GDS = 3 # type: ignore
-PACKET3_WRITE_DATA__DST_SEL__MEMORY = 5 # type: ignore
-PACKET3_WRITE_DATA__DST_SEL__MEMORY_MAPPED_ADC_PERSISTENT_STATE = 6 # type: ignore
-PACKET3_WRITE_DATA__ADDR_INCR__INCREMENT_ADDRESS = 0 # type: ignore
-PACKET3_WRITE_DATA__ADDR_INCR__DO_NOT_INCREMENT_ADDRESS = 1 # type: ignore
-PACKET3_WRITE_DATA__WR_CONFIRM__DO_NOT_WAIT_FOR_WRITE_CONFIRMATION = 0 # type: ignore
-PACKET3_WRITE_DATA__WR_CONFIRM__WAIT_FOR_WRITE_CONFIRMATION = 1 # type: ignore
-PACKET3_WRITE_DATA__MODE__PF_VF_DISABLED = 0 # type: ignore
-PACKET3_WRITE_DATA__MODE__PF_VF_ENABLED = 1 # type: ignore
-PACKET3_WRITE_DATA__TEMPORAL__RT = 0 # type: ignore
-PACKET3_WRITE_DATA__TEMPORAL__NT = 1 # type: ignore
-PACKET3_WRITE_DATA__TEMPORAL__HT = 2 # type: ignore
-PACKET3_WRITE_DATA__TEMPORAL__LU = 3 # type: ignore
-PACKET3_WRITE_DATA__CACHE_POLICY__LRU = 0 # type: ignore
-PACKET3_WRITE_DATA__CACHE_POLICY__STREAM = 1 # type: ignore
-PACKET3_WRITE_DATA__CACHE_POLICY__NOA = 2 # type: ignore
-PACKET3_WRITE_DATA__CACHE_POLICY__BYPASS = 3 # type: ignore
-PACKET3_DRAW_INDEX_INDIRECT_MULTI = 0x38 # type: ignore
-PACKET3_MEM_SEMAPHORE = 0x39 # type: ignore
-PACKET3_SEM_USE_MAILBOX = (0x1 << 16) # type: ignore
-PACKET3_SEM_SEL_SIGNAL_TYPE = (0x1 << 20) # type: ignore
-PACKET3_SEM_SEL_SIGNAL = (0x6 << 29) # type: ignore
-PACKET3_SEM_SEL_WAIT = (0x7 << 29) # type: ignore
-PACKET3_DRAW_INDEX_MULTI_INST = 0x3A # type: ignore
-PACKET3_COPY_DW = 0x3B # type: ignore
-PACKET3_WAIT_REG_MEM = 0x3C # type: ignore
+PACKET3_WRITE_DATA__DST_SEL__MEM_MAPPED_REGISTER = 0
+PACKET3_WRITE_DATA__DST_SEL__TC_L2 = 2
+PACKET3_WRITE_DATA__DST_SEL__GDS = 3
+PACKET3_WRITE_DATA__DST_SEL__MEMORY = 5
+PACKET3_WRITE_DATA__DST_SEL__MEMORY_MAPPED_ADC_PERSISTENT_STATE = 6
+PACKET3_WRITE_DATA__ADDR_INCR__INCREMENT_ADDRESS = 0
+PACKET3_WRITE_DATA__ADDR_INCR__DO_NOT_INCREMENT_ADDRESS = 1
+PACKET3_WRITE_DATA__WR_CONFIRM__DO_NOT_WAIT_FOR_WRITE_CONFIRMATION = 0
+PACKET3_WRITE_DATA__WR_CONFIRM__WAIT_FOR_WRITE_CONFIRMATION = 1
+PACKET3_WRITE_DATA__MODE__PF_VF_DISABLED = 0
+PACKET3_WRITE_DATA__MODE__PF_VF_ENABLED = 1
+PACKET3_WRITE_DATA__TEMPORAL__RT = 0
+PACKET3_WRITE_DATA__TEMPORAL__NT = 1
+PACKET3_WRITE_DATA__TEMPORAL__HT = 2
+PACKET3_WRITE_DATA__TEMPORAL__LU = 3
+PACKET3_WRITE_DATA__CACHE_POLICY__LRU = 0
+PACKET3_WRITE_DATA__CACHE_POLICY__STREAM = 1
+PACKET3_WRITE_DATA__CACHE_POLICY__NOA = 2
+PACKET3_WRITE_DATA__CACHE_POLICY__BYPASS = 3
+PACKET3_DRAW_INDEX_INDIRECT_MULTI = 0x38
+PACKET3_MEM_SEMAPHORE = 0x39
+PACKET3_SEM_USE_MAILBOX = (0x1 << 16)
+PACKET3_SEM_SEL_SIGNAL_TYPE = (0x1 << 20)
+PACKET3_SEM_SEL_SIGNAL = (0x6 << 29)
+PACKET3_SEM_SEL_WAIT = (0x7 << 29)
+PACKET3_DRAW_INDEX_MULTI_INST = 0x3A
+PACKET3_COPY_DW = 0x3B
+PACKET3_WAIT_REG_MEM = 0x3C
 WAIT_REG_MEM_FUNCTION = lambda x: ((x) << 0) # type: ignore
 WAIT_REG_MEM_MEM_SPACE = lambda x: ((x) << 4) # type: ignore
 WAIT_REG_MEM_OPERATION = lambda x: ((x) << 6) # type: ignore
@@ -170,28 +170,28 @@ PACKET3_WAIT_REG_MEM__REFERENCE = lambda x: ((unsigned)(x)) # type: ignore
 PACKET3_WAIT_REG_MEM__MASK = lambda x: ((unsigned)(x)) # type: ignore
 PACKET3_WAIT_REG_MEM__POLL_INTERVAL = lambda x: ((((unsigned)(x)) & 0xFFFF) << 0) # type: ignore
 PACKET3_WAIT_REG_MEM__OPTIMIZE_ACE_OFFLOAD_MODE = lambda x: ((((unsigned)(x)) & 0x1) << 31) # type: ignore
-PACKET3_WAIT_REG_MEM__FUNCTION__ALWAYS_PASS = 0 # type: ignore
-PACKET3_WAIT_REG_MEM__FUNCTION__LESS_THAN_REF_VALUE = 1 # type: ignore
-PACKET3_WAIT_REG_MEM__FUNCTION__LESS_THAN_EQUAL_TO_THE_REF_VALUE = 2 # type: ignore
-PACKET3_WAIT_REG_MEM__FUNCTION__EQUAL_TO_THE_REFERENCE_VALUE = 3 # type: ignore
-PACKET3_WAIT_REG_MEM__FUNCTION__NOT_EQUAL_REFERENCE_VALUE = 4 # type: ignore
-PACKET3_WAIT_REG_MEM__FUNCTION__GREATER_THAN_OR_EQUAL_REFERENCE_VALUE = 5 # type: ignore
-PACKET3_WAIT_REG_MEM__FUNCTION__GREATER_THAN_REFERENCE_VALUE = 6 # type: ignore
-PACKET3_WAIT_REG_MEM__MEM_SPACE__REGISTER_SPACE = 0 # type: ignore
-PACKET3_WAIT_REG_MEM__MEM_SPACE__MEMORY_SPACE = 1 # type: ignore
-PACKET3_WAIT_REG_MEM__OPERATION__WAIT_REG_MEM = 0 # type: ignore
-PACKET3_WAIT_REG_MEM__OPERATION__WR_WAIT_WR_REG = 1 # type: ignore
-PACKET3_WAIT_REG_MEM__OPERATION__WAIT_MEM_PREEMPTABLE = 3 # type: ignore
-PACKET3_WAIT_REG_MEM__CACHE_POLICY__LRU = 0 # type: ignore
-PACKET3_WAIT_REG_MEM__CACHE_POLICY__STREAM = 1 # type: ignore
-PACKET3_WAIT_REG_MEM__CACHE_POLICY__NOA = 2 # type: ignore
-PACKET3_WAIT_REG_MEM__CACHE_POLICY__BYPASS = 3 # type: ignore
-PACKET3_WAIT_REG_MEM__TEMPORAL__RT = 0 # type: ignore
-PACKET3_WAIT_REG_MEM__TEMPORAL__NT = 1 # type: ignore
-PACKET3_WAIT_REG_MEM__TEMPORAL__HT = 2 # type: ignore
-PACKET3_WAIT_REG_MEM__TEMPORAL__LU = 3 # type: ignore
-PACKET3_INDIRECT_BUFFER = 0x3F # type: ignore
-INDIRECT_BUFFER_VALID = (1 << 23) # type: ignore
+PACKET3_WAIT_REG_MEM__FUNCTION__ALWAYS_PASS = 0
+PACKET3_WAIT_REG_MEM__FUNCTION__LESS_THAN_REF_VALUE = 1
+PACKET3_WAIT_REG_MEM__FUNCTION__LESS_THAN_EQUAL_TO_THE_REF_VALUE = 2
+PACKET3_WAIT_REG_MEM__FUNCTION__EQUAL_TO_THE_REFERENCE_VALUE = 3
+PACKET3_WAIT_REG_MEM__FUNCTION__NOT_EQUAL_REFERENCE_VALUE = 4
+PACKET3_WAIT_REG_MEM__FUNCTION__GREATER_THAN_OR_EQUAL_REFERENCE_VALUE = 5
+PACKET3_WAIT_REG_MEM__FUNCTION__GREATER_THAN_REFERENCE_VALUE = 6
+PACKET3_WAIT_REG_MEM__MEM_SPACE__REGISTER_SPACE = 0
+PACKET3_WAIT_REG_MEM__MEM_SPACE__MEMORY_SPACE = 1
+PACKET3_WAIT_REG_MEM__OPERATION__WAIT_REG_MEM = 0
+PACKET3_WAIT_REG_MEM__OPERATION__WR_WAIT_WR_REG = 1
+PACKET3_WAIT_REG_MEM__OPERATION__WAIT_MEM_PREEMPTABLE = 3
+PACKET3_WAIT_REG_MEM__CACHE_POLICY__LRU = 0
+PACKET3_WAIT_REG_MEM__CACHE_POLICY__STREAM = 1
+PACKET3_WAIT_REG_MEM__CACHE_POLICY__NOA = 2
+PACKET3_WAIT_REG_MEM__CACHE_POLICY__BYPASS = 3
+PACKET3_WAIT_REG_MEM__TEMPORAL__RT = 0
+PACKET3_WAIT_REG_MEM__TEMPORAL__NT = 1
+PACKET3_WAIT_REG_MEM__TEMPORAL__HT = 2
+PACKET3_WAIT_REG_MEM__TEMPORAL__LU = 3
+PACKET3_INDIRECT_BUFFER = 0x3F
+INDIRECT_BUFFER_VALID = (1 << 23)
 INDIRECT_BUFFER_CACHE_POLICY = lambda x: ((x) << 28) # type: ignore
 INDIRECT_BUFFER_PRE_ENB = lambda x: ((x) << 21) # type: ignore
 INDIRECT_BUFFER_PRE_RESUME = lambda x: ((x) << 30) # type: ignore
@@ -205,16 +205,16 @@ PACKET3_INDIRECT_BUFFER__VMID = lambda x: ((((unsigned)(x)) & 0xF) << 24) # type
 PACKET3_INDIRECT_BUFFER__CACHE_POLICY = lambda x: ((((unsigned)(x)) & 0x3) << 28) # type: ignore
 PACKET3_INDIRECT_BUFFER__TEMPORAL = lambda x: ((((unsigned)(x)) & 0x3) << 28) # type: ignore
 PACKET3_INDIRECT_BUFFER__PRIV = lambda x: ((((unsigned)(x)) & 0x1) << 31) # type: ignore
-PACKET3_INDIRECT_BUFFER__TEMPORAL__RT = 0 # type: ignore
-PACKET3_INDIRECT_BUFFER__TEMPORAL__NT = 1 # type: ignore
-PACKET3_INDIRECT_BUFFER__TEMPORAL__HT = 2 # type: ignore
-PACKET3_INDIRECT_BUFFER__TEMPORAL__LU = 3 # type: ignore
-PACKET3_INDIRECT_BUFFER__CACHE_POLICY__LRU = 0 # type: ignore
-PACKET3_INDIRECT_BUFFER__CACHE_POLICY__STREAM = 1 # type: ignore
-PACKET3_INDIRECT_BUFFER__CACHE_POLICY__NOA = 2 # type: ignore
-PACKET3_INDIRECT_BUFFER__CACHE_POLICY__BYPASS = 3 # type: ignore
-PACKET3_COND_INDIRECT_BUFFER = 0x3F # type: ignore
-PACKET3_COPY_DATA = 0x40 # type: ignore
+PACKET3_INDIRECT_BUFFER__TEMPORAL__RT = 0
+PACKET3_INDIRECT_BUFFER__TEMPORAL__NT = 1
+PACKET3_INDIRECT_BUFFER__TEMPORAL__HT = 2
+PACKET3_INDIRECT_BUFFER__TEMPORAL__LU = 3
+PACKET3_INDIRECT_BUFFER__CACHE_POLICY__LRU = 0
+PACKET3_INDIRECT_BUFFER__CACHE_POLICY__STREAM = 1
+PACKET3_INDIRECT_BUFFER__CACHE_POLICY__NOA = 2
+PACKET3_INDIRECT_BUFFER__CACHE_POLICY__BYPASS = 3
+PACKET3_COND_INDIRECT_BUFFER = 0x3F
+PACKET3_COPY_DATA = 0x40
 PACKET3_COPY_DATA__SRC_SEL = lambda x: ((((unsigned)(x)) & 0xF) << 0) # type: ignore
 PACKET3_COPY_DATA__DST_SEL = lambda x: ((((unsigned)(x)) & 0xF) << 8) # type: ignore
 PACKET3_COPY_DATA__SRC_CACHE_POLICY = lambda x: ((((unsigned)(x)) & 0x3) << 13) # type: ignore
@@ -242,53 +242,53 @@ PACKET3_COPY_DATA__SRC_REG_OFFSET_LO = lambda x: ((unsigned)(x)) # type: ignore
 PACKET3_COPY_DATA__SRC_REG_OFFSET_HI = lambda x: ((((unsigned)(x)) & 0xFF) << 0) # type: ignore
 PACKET3_COPY_DATA__DST_REG_OFFSET_LO = lambda x: ((unsigned)(x)) # type: ignore
 PACKET3_COPY_DATA__DST_REG_OFFSET_HI = lambda x: ((((unsigned)(x)) & 0xFF) << 0) # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__MEM_MAPPED_REGISTER = 0 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__TC_L2_OBSOLETE = 1 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__TC_L2 = 2 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__GDS = 3 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__PERFCOUNTERS = 4 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__IMMEDIATE_DATA = 5 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__ATOMIC_RETURN_DATA = 6 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__GDS_ATOMIC_RETURN_DATA0 = 7 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__GDS_ATOMIC_RETURN_DATA1 = 8 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__GPU_CLOCK_COUNT = 9 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__SYSTEM_CLOCK_COUNT = 10 # type: ignore
-PACKET3_COPY_DATA__DST_SEL__MEM_MAPPED_REGISTER = 0 # type: ignore
-PACKET3_COPY_DATA__DST_SEL__TC_L2 = 2 # type: ignore
-PACKET3_COPY_DATA__DST_SEL__GDS = 3 # type: ignore
-PACKET3_COPY_DATA__DST_SEL__PERFCOUNTERS = 4 # type: ignore
-PACKET3_COPY_DATA__DST_SEL__TC_L2_OBSOLETE = 5 # type: ignore
-PACKET3_COPY_DATA__DST_SEL__MEM_MAPPED_REG_DC = 6 # type: ignore
-PACKET3_COPY_DATA__SRC_TEMPORAL__RT = 0 # type: ignore
-PACKET3_COPY_DATA__SRC_TEMPORAL__NT = 1 # type: ignore
-PACKET3_COPY_DATA__SRC_TEMPORAL__HT = 2 # type: ignore
-PACKET3_COPY_DATA__SRC_TEMPORAL__LU = 3 # type: ignore
-PACKET3_COPY_DATA__SRC_CACHE_POLICY__LRU = 0 # type: ignore
-PACKET3_COPY_DATA__SRC_CACHE_POLICY__STREAM = 1 # type: ignore
-PACKET3_COPY_DATA__SRC_CACHE_POLICY__NOA = 2 # type: ignore
-PACKET3_COPY_DATA__SRC_CACHE_POLICY__BYPASS = 3 # type: ignore
-PACKET3_COPY_DATA__COUNT_SEL__32_BITS_OF_DATA = 0 # type: ignore
-PACKET3_COPY_DATA__COUNT_SEL__64_BITS_OF_DATA = 1 # type: ignore
-PACKET3_COPY_DATA__WR_CONFIRM__DO_NOT_WAIT_FOR_CONFIRMATION = 0 # type: ignore
-PACKET3_COPY_DATA__WR_CONFIRM__WAIT_FOR_CONFIRMATION = 1 # type: ignore
-PACKET3_COPY_DATA__MODE__PF_VF_DISABLED = 0 # type: ignore
-PACKET3_COPY_DATA__MODE__PF_VF_ENABLED = 1 # type: ignore
-PACKET3_COPY_DATA__DST_TEMPORAL__RT = 0 # type: ignore
-PACKET3_COPY_DATA__DST_TEMPORAL__NT = 1 # type: ignore
-PACKET3_COPY_DATA__DST_TEMPORAL__HT = 2 # type: ignore
-PACKET3_COPY_DATA__DST_TEMPORAL__LU = 3 # type: ignore
-PACKET3_COPY_DATA__DST_CACHE_POLICY__LRU = 0 # type: ignore
-PACKET3_COPY_DATA__DST_CACHE_POLICY__STREAM = 1 # type: ignore
-PACKET3_COPY_DATA__DST_CACHE_POLICY__NOA = 2 # type: ignore
-PACKET3_COPY_DATA__DST_CACHE_POLICY__BYPASS = 3 # type: ignore
-PACKET3_COPY_DATA__PQ_EXE_STATUS__DEFAULT = 0 # type: ignore
-PACKET3_COPY_DATA__PQ_EXE_STATUS__PHASE_UPDATE = 1 # type: ignore
-PACKET3_CP_DMA = 0x41 # type: ignore
-PACKET3_PFP_SYNC_ME = 0x42 # type: ignore
-PACKET3_SURFACE_SYNC = 0x43 # type: ignore
-PACKET3_ME_INITIALIZE = 0x44 # type: ignore
-PACKET3_COND_WRITE = 0x45 # type: ignore
-PACKET3_EVENT_WRITE = 0x46 # type: ignore
+PACKET3_COPY_DATA__SRC_SEL__MEM_MAPPED_REGISTER = 0
+PACKET3_COPY_DATA__SRC_SEL__TC_L2_OBSOLETE = 1
+PACKET3_COPY_DATA__SRC_SEL__TC_L2 = 2
+PACKET3_COPY_DATA__SRC_SEL__GDS = 3
+PACKET3_COPY_DATA__SRC_SEL__PERFCOUNTERS = 4
+PACKET3_COPY_DATA__SRC_SEL__IMMEDIATE_DATA = 5
+PACKET3_COPY_DATA__SRC_SEL__ATOMIC_RETURN_DATA = 6
+PACKET3_COPY_DATA__SRC_SEL__GDS_ATOMIC_RETURN_DATA0 = 7
+PACKET3_COPY_DATA__SRC_SEL__GDS_ATOMIC_RETURN_DATA1 = 8
+PACKET3_COPY_DATA__SRC_SEL__GPU_CLOCK_COUNT = 9
+PACKET3_COPY_DATA__SRC_SEL__SYSTEM_CLOCK_COUNT = 10
+PACKET3_COPY_DATA__DST_SEL__MEM_MAPPED_REGISTER = 0
+PACKET3_COPY_DATA__DST_SEL__TC_L2 = 2
+PACKET3_COPY_DATA__DST_SEL__GDS = 3
+PACKET3_COPY_DATA__DST_SEL__PERFCOUNTERS = 4
+PACKET3_COPY_DATA__DST_SEL__TC_L2_OBSOLETE = 5
+PACKET3_COPY_DATA__DST_SEL__MEM_MAPPED_REG_DC = 6
+PACKET3_COPY_DATA__SRC_TEMPORAL__RT = 0
+PACKET3_COPY_DATA__SRC_TEMPORAL__NT = 1
+PACKET3_COPY_DATA__SRC_TEMPORAL__HT = 2
+PACKET3_COPY_DATA__SRC_TEMPORAL__LU = 3
+PACKET3_COPY_DATA__SRC_CACHE_POLICY__LRU = 0
+PACKET3_COPY_DATA__SRC_CACHE_POLICY__STREAM = 1
+PACKET3_COPY_DATA__SRC_CACHE_POLICY__NOA = 2
+PACKET3_COPY_DATA__SRC_CACHE_POLICY__BYPASS = 3
+PACKET3_COPY_DATA__COUNT_SEL__32_BITS_OF_DATA = 0
+PACKET3_COPY_DATA__COUNT_SEL__64_BITS_OF_DATA = 1
+PACKET3_COPY_DATA__WR_CONFIRM__DO_NOT_WAIT_FOR_CONFIRMATION = 0
+PACKET3_COPY_DATA__WR_CONFIRM__WAIT_FOR_CONFIRMATION = 1
+PACKET3_COPY_DATA__MODE__PF_VF_DISABLED = 0
+PACKET3_COPY_DATA__MODE__PF_VF_ENABLED = 1
+PACKET3_COPY_DATA__DST_TEMPORAL__RT = 0
+PACKET3_COPY_DATA__DST_TEMPORAL__NT = 1
+PACKET3_COPY_DATA__DST_TEMPORAL__HT = 2
+PACKET3_COPY_DATA__DST_TEMPORAL__LU = 3
+PACKET3_COPY_DATA__DST_CACHE_POLICY__LRU = 0
+PACKET3_COPY_DATA__DST_CACHE_POLICY__STREAM = 1
+PACKET3_COPY_DATA__DST_CACHE_POLICY__NOA = 2
+PACKET3_COPY_DATA__DST_CACHE_POLICY__BYPASS = 3
+PACKET3_COPY_DATA__PQ_EXE_STATUS__DEFAULT = 0
+PACKET3_COPY_DATA__PQ_EXE_STATUS__PHASE_UPDATE = 1
+PACKET3_CP_DMA = 0x41
+PACKET3_PFP_SYNC_ME = 0x42
+PACKET3_SURFACE_SYNC = 0x43
+PACKET3_ME_INITIALIZE = 0x44
+PACKET3_COND_WRITE = 0x45
+PACKET3_EVENT_WRITE = 0x46
 EVENT_TYPE = lambda x: ((x) << 0) # type: ignore
 EVENT_INDEX = lambda x: ((x) << 8) # type: ignore
 PACKET3_EVENT_WRITE__EVENT_TYPE = lambda x: ((((unsigned)(x)) & 0x3F) << 0) # type: ignore
@@ -297,57 +297,57 @@ PACKET3_EVENT_WRITE__SAMP_PLST_CNTR_MODE = lambda x: ((((unsigned)(x)) & 0x3) <<
 PACKET3_EVENT_WRITE__OFFLOAD_ENABLE = lambda x: ((((unsigned)(x)) & 0x1) << 0) # type: ignore
 PACKET3_EVENT_WRITE__ADDRESS_LO = lambda x: ((((unsigned)(x)) & 0x1FFFFFFF) << 3) # type: ignore
 PACKET3_EVENT_WRITE__ADDRESS_HI = lambda x: ((unsigned)(x)) # type: ignore
-PACKET3_EVENT_WRITE__EVENT_INDEX__OTHER = 0 # type: ignore
-PACKET3_EVENT_WRITE__EVENT_INDEX__SAMPLE_PIPELINESTAT = 2 # type: ignore
-PACKET3_EVENT_WRITE__EVENT_INDEX__CS_PARTIAL_FLUSH = 4 # type: ignore
-PACKET3_EVENT_WRITE__EVENT_INDEX__SAMPLE_STREAMOUTSTATS = 8 # type: ignore
-PACKET3_EVENT_WRITE__EVENT_INDEX__SAMPLE_STREAMOUTSTATS1 = 9 # type: ignore
-PACKET3_EVENT_WRITE__EVENT_INDEX__SAMPLE_STREAMOUTSTATS2 = 10 # type: ignore
-PACKET3_EVENT_WRITE__EVENT_INDEX__SAMPLE_STREAMOUTSTATS3 = 11 # type: ignore
-PACKET3_EVENT_WRITE__SAMP_PLST_CNTR_MODE__LEGACY_MODE = 0 # type: ignore
-PACKET3_EVENT_WRITE__SAMP_PLST_CNTR_MODE__MIXED_MODE1 = 1 # type: ignore
-PACKET3_EVENT_WRITE__SAMP_PLST_CNTR_MODE__NEW_MODE = 2 # type: ignore
-PACKET3_EVENT_WRITE__SAMP_PLST_CNTR_MODE__MIXED_MODE3 = 3 # type: ignore
-PACKET3_EVENT_WRITE_EOP = 0x47 # type: ignore
-PACKET3_EVENT_WRITE_EOS = 0x48 # type: ignore
-PACKET3_RELEASE_MEM = 0x49 # type: ignore
+PACKET3_EVENT_WRITE__EVENT_INDEX__OTHER = 0
+PACKET3_EVENT_WRITE__EVENT_INDEX__SAMPLE_PIPELINESTAT = 2
+PACKET3_EVENT_WRITE__EVENT_INDEX__CS_PARTIAL_FLUSH = 4
+PACKET3_EVENT_WRITE__EVENT_INDEX__SAMPLE_STREAMOUTSTATS = 8
+PACKET3_EVENT_WRITE__EVENT_INDEX__SAMPLE_STREAMOUTSTATS1 = 9
+PACKET3_EVENT_WRITE__EVENT_INDEX__SAMPLE_STREAMOUTSTATS2 = 10
+PACKET3_EVENT_WRITE__EVENT_INDEX__SAMPLE_STREAMOUTSTATS3 = 11
+PACKET3_EVENT_WRITE__SAMP_PLST_CNTR_MODE__LEGACY_MODE = 0
+PACKET3_EVENT_WRITE__SAMP_PLST_CNTR_MODE__MIXED_MODE1 = 1
+PACKET3_EVENT_WRITE__SAMP_PLST_CNTR_MODE__NEW_MODE = 2
+PACKET3_EVENT_WRITE__SAMP_PLST_CNTR_MODE__MIXED_MODE3 = 3
+PACKET3_EVENT_WRITE_EOP = 0x47
+PACKET3_EVENT_WRITE_EOS = 0x48
+PACKET3_RELEASE_MEM = 0x49
 PACKET3_RELEASE_MEM_EVENT_TYPE = lambda x: ((x) << 0) # type: ignore
 PACKET3_RELEASE_MEM_EVENT_INDEX = lambda x: ((x) << 8) # type: ignore
-PACKET3_RELEASE_MEM_GCR_GLM_WB = (1 << 12) # type: ignore
-PACKET3_RELEASE_MEM_GCR_GLM_INV = (1 << 13) # type: ignore
-PACKET3_RELEASE_MEM_GCR_GLV_INV = (1 << 14) # type: ignore
-PACKET3_RELEASE_MEM_GCR_GL1_INV = (1 << 15) # type: ignore
-PACKET3_RELEASE_MEM_GCR_GL2_US = (1 << 16) # type: ignore
-PACKET3_RELEASE_MEM_GCR_GL2_RANGE = (1 << 17) # type: ignore
-PACKET3_RELEASE_MEM_GCR_GL2_DISCARD = (1 << 19) # type: ignore
-PACKET3_RELEASE_MEM_GCR_GL2_INV = (1 << 20) # type: ignore
-PACKET3_RELEASE_MEM_GCR_GL2_WB = (1 << 21) # type: ignore
-PACKET3_RELEASE_MEM_GCR_SEQ = (1 << 22) # type: ignore
+PACKET3_RELEASE_MEM_GCR_GLM_WB = (1 << 12)
+PACKET3_RELEASE_MEM_GCR_GLM_INV = (1 << 13)
+PACKET3_RELEASE_MEM_GCR_GLV_INV = (1 << 14)
+PACKET3_RELEASE_MEM_GCR_GL1_INV = (1 << 15)
+PACKET3_RELEASE_MEM_GCR_GL2_US = (1 << 16)
+PACKET3_RELEASE_MEM_GCR_GL2_RANGE = (1 << 17)
+PACKET3_RELEASE_MEM_GCR_GL2_DISCARD = (1 << 19)
+PACKET3_RELEASE_MEM_GCR_GL2_INV = (1 << 20)
+PACKET3_RELEASE_MEM_GCR_GL2_WB = (1 << 21)
+PACKET3_RELEASE_MEM_GCR_SEQ = (1 << 22)
 PACKET3_RELEASE_MEM_CACHE_POLICY = lambda x: ((x) << 25) # type: ignore
-PACKET3_RELEASE_MEM_EXECUTE = (1 << 28) # type: ignore
+PACKET3_RELEASE_MEM_EXECUTE = (1 << 28)
 PACKET3_RELEASE_MEM_DATA_SEL = lambda x: ((x) << 29) # type: ignore
 PACKET3_RELEASE_MEM_INT_SEL = lambda x: ((x) << 24) # type: ignore
 PACKET3_RELEASE_MEM_DST_SEL = lambda x: ((x) << 16) # type: ignore
-PACKET3_PREAMBLE_CNTL = 0x4A # type: ignore
-PACKET3_PREAMBLE_BEGIN_CLEAR_STATE = (2 << 28) # type: ignore
-PACKET3_PREAMBLE_END_CLEAR_STATE = (3 << 28) # type: ignore
-PACKET3_DMA_DATA = 0x50 # type: ignore
+PACKET3_PREAMBLE_CNTL = 0x4A
+PACKET3_PREAMBLE_BEGIN_CLEAR_STATE = (2 << 28)
+PACKET3_PREAMBLE_END_CLEAR_STATE = (3 << 28)
+PACKET3_DMA_DATA = 0x50
 PACKET3_DMA_DATA_ENGINE = lambda x: ((x) << 0) # type: ignore
 PACKET3_DMA_DATA_SRC_CACHE_POLICY = lambda x: ((x) << 13) # type: ignore
 PACKET3_DMA_DATA_DST_SEL = lambda x: ((x) << 20) # type: ignore
 PACKET3_DMA_DATA_DST_CACHE_POLICY = lambda x: ((x) << 25) # type: ignore
 PACKET3_DMA_DATA_SRC_SEL = lambda x: ((x) << 29) # type: ignore
-PACKET3_DMA_DATA_CP_SYNC = (1 << 31) # type: ignore
-PACKET3_DMA_DATA_CMD_SAS = (1 << 26) # type: ignore
-PACKET3_DMA_DATA_CMD_DAS = (1 << 27) # type: ignore
-PACKET3_DMA_DATA_CMD_SAIC = (1 << 28) # type: ignore
-PACKET3_DMA_DATA_CMD_DAIC = (1 << 29) # type: ignore
-PACKET3_DMA_DATA_CMD_RAW_WAIT = (1 << 30) # type: ignore
-PACKET3_CONTEXT_REG_RMW = 0x51 # type: ignore
-PACKET3_GFX_CNTX_UPDATE = 0x52 # type: ignore
-PACKET3_BLK_CNTX_UPDATE = 0x53 # type: ignore
-PACKET3_INCR_UPDT_STATE = 0x55 # type: ignore
-PACKET3_ACQUIRE_MEM = 0x58 # type: ignore
+PACKET3_DMA_DATA_CP_SYNC = (1 << 31)
+PACKET3_DMA_DATA_CMD_SAS = (1 << 26)
+PACKET3_DMA_DATA_CMD_DAS = (1 << 27)
+PACKET3_DMA_DATA_CMD_SAIC = (1 << 28)
+PACKET3_DMA_DATA_CMD_DAIC = (1 << 29)
+PACKET3_DMA_DATA_CMD_RAW_WAIT = (1 << 30)
+PACKET3_CONTEXT_REG_RMW = 0x51
+PACKET3_GFX_CNTX_UPDATE = 0x52
+PACKET3_BLK_CNTX_UPDATE = 0x53
+PACKET3_INCR_UPDT_STATE = 0x55
+PACKET3_ACQUIRE_MEM = 0x58
 PACKET3_ACQUIRE_MEM_GCR_CNTL_GLI_INV = lambda x: ((x) << 0) # type: ignore
 PACKET3_ACQUIRE_MEM_GCR_CNTL_GL1_RANGE = lambda x: ((x) << 2) # type: ignore
 PACKET3_ACQUIRE_MEM_GCR_CNTL_GLM_WB = lambda x: ((x) << 4) # type: ignore
@@ -362,94 +362,94 @@ PACKET3_ACQUIRE_MEM_GCR_CNTL_GL2_DISCARD = lambda x: ((x) << 13) # type: ignore
 PACKET3_ACQUIRE_MEM_GCR_CNTL_GL2_INV = lambda x: ((x) << 14) # type: ignore
 PACKET3_ACQUIRE_MEM_GCR_CNTL_GL2_WB = lambda x: ((x) << 15) # type: ignore
 PACKET3_ACQUIRE_MEM_GCR_CNTL_SEQ = lambda x: ((x) << 16) # type: ignore
-PACKET3_ACQUIRE_MEM_GCR_RANGE_IS_PA = (1 << 18) # type: ignore
+PACKET3_ACQUIRE_MEM_GCR_RANGE_IS_PA = (1 << 18)
 PACKET3_ACQUIRE_MEM__COHER_SIZE = lambda x: ((unsigned)(x)) # type: ignore
 PACKET3_ACQUIRE_MEM__COHER_SIZE_HI = lambda x: ((((unsigned)(x)) & 0xFF) << 0) # type: ignore
 PACKET3_ACQUIRE_MEM__COHER_BASE_LO = lambda x: ((unsigned)(x)) # type: ignore
 PACKET3_ACQUIRE_MEM__COHER_BASE_HI = lambda x: ((((unsigned)(x)) & 0xFFFFFF) << 0) # type: ignore
 PACKET3_ACQUIRE_MEM__POLL_INTERVAL = lambda x: ((((unsigned)(x)) & 0xFFFF) << 0) # type: ignore
 PACKET3_ACQUIRE_MEM__GCR_CNTL = lambda x: ((((unsigned)(x)) & 0x7FFFF) << 0) # type: ignore
-PACKET3_REWIND = 0x59 # type: ignore
-PACKET3_INTERRUPT = 0x5A # type: ignore
-PACKET3_GEN_PDEPTE = 0x5B # type: ignore
-PACKET3_INDIRECT_BUFFER_PASID = 0x5C # type: ignore
-PACKET3_PRIME_UTCL2 = 0x5D # type: ignore
-PACKET3_LOAD_UCONFIG_REG = 0x5E # type: ignore
-PACKET3_LOAD_SH_REG = 0x5F # type: ignore
-PACKET3_LOAD_CONFIG_REG = 0x60 # type: ignore
-PACKET3_LOAD_CONTEXT_REG = 0x61 # type: ignore
-PACKET3_LOAD_COMPUTE_STATE = 0x62 # type: ignore
-PACKET3_LOAD_SH_REG_INDEX = 0x63 # type: ignore
-PACKET3_SET_CONFIG_REG = 0x68 # type: ignore
-PACKET3_SET_CONFIG_REG_START = 0x00002000 # type: ignore
-PACKET3_SET_CONFIG_REG_END = 0x00002c00 # type: ignore
-PACKET3_SET_CONTEXT_REG = 0x69 # type: ignore
-PACKET3_SET_CONTEXT_REG_START = 0x0000a000 # type: ignore
-PACKET3_SET_CONTEXT_REG_END = 0x0000a400 # type: ignore
-PACKET3_SET_CONTEXT_REG_INDEX = 0x6A # type: ignore
-PACKET3_SET_VGPR_REG_DI_MULTI = 0x71 # type: ignore
-PACKET3_SET_SH_REG_DI = 0x72 # type: ignore
-PACKET3_SET_CONTEXT_REG_INDIRECT = 0x73 # type: ignore
-PACKET3_SET_SH_REG_DI_MULTI = 0x74 # type: ignore
-PACKET3_GFX_PIPE_LOCK = 0x75 # type: ignore
-PACKET3_SET_SH_REG = 0x76 # type: ignore
-PACKET3_SET_SH_REG_START = 0x00002c00 # type: ignore
-PACKET3_SET_SH_REG_END = 0x00003000 # type: ignore
+PACKET3_REWIND = 0x59
+PACKET3_INTERRUPT = 0x5A
+PACKET3_GEN_PDEPTE = 0x5B
+PACKET3_INDIRECT_BUFFER_PASID = 0x5C
+PACKET3_PRIME_UTCL2 = 0x5D
+PACKET3_LOAD_UCONFIG_REG = 0x5E
+PACKET3_LOAD_SH_REG = 0x5F
+PACKET3_LOAD_CONFIG_REG = 0x60
+PACKET3_LOAD_CONTEXT_REG = 0x61
+PACKET3_LOAD_COMPUTE_STATE = 0x62
+PACKET3_LOAD_SH_REG_INDEX = 0x63
+PACKET3_SET_CONFIG_REG = 0x68
+PACKET3_SET_CONFIG_REG_START = 0x00002000
+PACKET3_SET_CONFIG_REG_END = 0x00002c00
+PACKET3_SET_CONTEXT_REG = 0x69
+PACKET3_SET_CONTEXT_REG_START = 0x0000a000
+PACKET3_SET_CONTEXT_REG_END = 0x0000a400
+PACKET3_SET_CONTEXT_REG_INDEX = 0x6A
+PACKET3_SET_VGPR_REG_DI_MULTI = 0x71
+PACKET3_SET_SH_REG_DI = 0x72
+PACKET3_SET_CONTEXT_REG_INDIRECT = 0x73
+PACKET3_SET_SH_REG_DI_MULTI = 0x74
+PACKET3_GFX_PIPE_LOCK = 0x75
+PACKET3_SET_SH_REG = 0x76
+PACKET3_SET_SH_REG_START = 0x00002c00
+PACKET3_SET_SH_REG_END = 0x00003000
 PACKET3_SET_SH_REG__REG_OFFSET = lambda x: ((((unsigned)(x)) & 0xFFFF) << 0) # type: ignore
 PACKET3_SET_SH_REG__VMID_SHIFT = lambda x: ((((unsigned)(x)) & 0x1F) << 23) # type: ignore
 PACKET3_SET_SH_REG__INDEX = lambda x: ((((unsigned)(x)) & 0xF) << 28) # type: ignore
-PACKET3_SET_SH_REG__INDEX__DEFAULT = 0 # type: ignore
-PACKET3_SET_SH_REG__INDEX__INSERT_VMID = 1 # type: ignore
-PACKET3_SET_SH_REG_OFFSET = 0x77 # type: ignore
-PACKET3_SET_QUEUE_REG = 0x78 # type: ignore
-PACKET3_SET_UCONFIG_REG = 0x79 # type: ignore
-PACKET3_SET_UCONFIG_REG_START = 0x0000c000 # type: ignore
-PACKET3_SET_UCONFIG_REG_END = 0x0000c400 # type: ignore
+PACKET3_SET_SH_REG__INDEX__DEFAULT = 0
+PACKET3_SET_SH_REG__INDEX__INSERT_VMID = 1
+PACKET3_SET_SH_REG_OFFSET = 0x77
+PACKET3_SET_QUEUE_REG = 0x78
+PACKET3_SET_UCONFIG_REG = 0x79
+PACKET3_SET_UCONFIG_REG_START = 0x0000c000
+PACKET3_SET_UCONFIG_REG_END = 0x0000c400
 PACKET3_SET_UCONFIG_REG__REG_OFFSET = lambda x: ((((unsigned)(x)) & 0xFFFF) << 0) # type: ignore
-PACKET3_SET_UCONFIG_REG_INDEX = 0x7A # type: ignore
-PACKET3_FORWARD_HEADER = 0x7C # type: ignore
-PACKET3_SCRATCH_RAM_WRITE = 0x7D # type: ignore
-PACKET3_SCRATCH_RAM_READ = 0x7E # type: ignore
-PACKET3_LOAD_CONST_RAM = 0x80 # type: ignore
-PACKET3_WRITE_CONST_RAM = 0x81 # type: ignore
-PACKET3_DUMP_CONST_RAM = 0x83 # type: ignore
-PACKET3_INCREMENT_CE_COUNTER = 0x84 # type: ignore
-PACKET3_INCREMENT_DE_COUNTER = 0x85 # type: ignore
-PACKET3_WAIT_ON_CE_COUNTER = 0x86 # type: ignore
-PACKET3_WAIT_ON_DE_COUNTER_DIFF = 0x88 # type: ignore
-PACKET3_SWITCH_BUFFER = 0x8B # type: ignore
-PACKET3_DISPATCH_DRAW_PREAMBLE = 0x8C # type: ignore
-PACKET3_DISPATCH_DRAW_PREAMBLE_ACE = 0x8C # type: ignore
-PACKET3_DISPATCH_DRAW = 0x8D # type: ignore
-PACKET3_DISPATCH_DRAW_ACE = 0x8D # type: ignore
-PACKET3_GET_LOD_STATS = 0x8E # type: ignore
-PACKET3_DRAW_MULTI_PREAMBLE = 0x8F # type: ignore
-PACKET3_FRAME_CONTROL = 0x90 # type: ignore
-FRAME_TMZ = (1 << 0) # type: ignore
+PACKET3_SET_UCONFIG_REG_INDEX = 0x7A
+PACKET3_FORWARD_HEADER = 0x7C
+PACKET3_SCRATCH_RAM_WRITE = 0x7D
+PACKET3_SCRATCH_RAM_READ = 0x7E
+PACKET3_LOAD_CONST_RAM = 0x80
+PACKET3_WRITE_CONST_RAM = 0x81
+PACKET3_DUMP_CONST_RAM = 0x83
+PACKET3_INCREMENT_CE_COUNTER = 0x84
+PACKET3_INCREMENT_DE_COUNTER = 0x85
+PACKET3_WAIT_ON_CE_COUNTER = 0x86
+PACKET3_WAIT_ON_DE_COUNTER_DIFF = 0x88
+PACKET3_SWITCH_BUFFER = 0x8B
+PACKET3_DISPATCH_DRAW_PREAMBLE = 0x8C
+PACKET3_DISPATCH_DRAW_PREAMBLE_ACE = 0x8C
+PACKET3_DISPATCH_DRAW = 0x8D
+PACKET3_DISPATCH_DRAW_ACE = 0x8D
+PACKET3_GET_LOD_STATS = 0x8E
+PACKET3_DRAW_MULTI_PREAMBLE = 0x8F
+PACKET3_FRAME_CONTROL = 0x90
+FRAME_TMZ = (1 << 0)
 FRAME_CMD = lambda x: ((x) << 28) # type: ignore
-PACKET3_INDEX_ATTRIBUTES_INDIRECT = 0x91 # type: ignore
-PACKET3_WAIT_REG_MEM64 = 0x93 # type: ignore
-PACKET3_COND_PREEMPT = 0x94 # type: ignore
-PACKET3_HDP_FLUSH = 0x95 # type: ignore
-PACKET3_COPY_DATA_RB = 0x96 # type: ignore
-PACKET3_INVALIDATE_TLBS = 0x98 # type: ignore
+PACKET3_INDEX_ATTRIBUTES_INDIRECT = 0x91
+PACKET3_WAIT_REG_MEM64 = 0x93
+PACKET3_COND_PREEMPT = 0x94
+PACKET3_HDP_FLUSH = 0x95
+PACKET3_COPY_DATA_RB = 0x96
+PACKET3_INVALIDATE_TLBS = 0x98
 PACKET3_INVALIDATE_TLBS_DST_SEL = lambda x: ((x) << 0) # type: ignore
 PACKET3_INVALIDATE_TLBS_ALL_HUB = lambda x: ((x) << 4) # type: ignore
 PACKET3_INVALIDATE_TLBS_PASID = lambda x: ((x) << 5) # type: ignore
 PACKET3_INVALIDATE_TLBS_FLUSH_TYPE = lambda x: ((x) << 29) # type: ignore
-PACKET3_AQL_PACKET = 0x99 # type: ignore
-PACKET3_DMA_DATA_FILL_MULTI = 0x9A # type: ignore
-PACKET3_SET_SH_REG_INDEX = 0x9B # type: ignore
-PACKET3_DRAW_INDIRECT_COUNT_MULTI = 0x9C # type: ignore
-PACKET3_DRAW_INDEX_INDIRECT_COUNT_MULTI = 0x9D # type: ignore
-PACKET3_DUMP_CONST_RAM_OFFSET = 0x9E # type: ignore
-PACKET3_LOAD_CONTEXT_REG_INDEX = 0x9F # type: ignore
-PACKET3_SET_RESOURCES = 0xA0 # type: ignore
+PACKET3_AQL_PACKET = 0x99
+PACKET3_DMA_DATA_FILL_MULTI = 0x9A
+PACKET3_SET_SH_REG_INDEX = 0x9B
+PACKET3_DRAW_INDIRECT_COUNT_MULTI = 0x9C
+PACKET3_DRAW_INDEX_INDIRECT_COUNT_MULTI = 0x9D
+PACKET3_DUMP_CONST_RAM_OFFSET = 0x9E
+PACKET3_LOAD_CONTEXT_REG_INDEX = 0x9F
+PACKET3_SET_RESOURCES = 0xA0
 PACKET3_SET_RESOURCES_VMID_MASK = lambda x: ((x) << 0) # type: ignore
 PACKET3_SET_RESOURCES_UNMAP_LATENTY = lambda x: ((x) << 16) # type: ignore
 PACKET3_SET_RESOURCES_QUEUE_TYPE = lambda x: ((x) << 29) # type: ignore
-PACKET3_MAP_PROCESS = 0xA1 # type: ignore
-PACKET3_MAP_QUEUES = 0xA2 # type: ignore
+PACKET3_MAP_PROCESS = 0xA1
+PACKET3_MAP_QUEUES = 0xA2
 PACKET3_MAP_QUEUES_QUEUE_SEL = lambda x: ((x) << 4) # type: ignore
 PACKET3_MAP_QUEUES_VMID = lambda x: ((x) << 8) # type: ignore
 PACKET3_MAP_QUEUES_QUEUE = lambda x: ((x) << 13) # type: ignore
@@ -461,7 +461,7 @@ PACKET3_MAP_QUEUES_ENGINE_SEL = lambda x: ((x) << 26) # type: ignore
 PACKET3_MAP_QUEUES_NUM_QUEUES = lambda x: ((x) << 29) # type: ignore
 PACKET3_MAP_QUEUES_CHECK_DISABLE = lambda x: ((x) << 1) # type: ignore
 PACKET3_MAP_QUEUES_DOORBELL_OFFSET = lambda x: ((x) << 2) # type: ignore
-PACKET3_UNMAP_QUEUES = 0xA3 # type: ignore
+PACKET3_UNMAP_QUEUES = 0xA3
 PACKET3_UNMAP_QUEUES_ACTION = lambda x: ((x) << 0) # type: ignore
 PACKET3_UNMAP_QUEUES_QUEUE_SEL = lambda x: ((x) << 4) # type: ignore
 PACKET3_UNMAP_QUEUES_ENGINE_SEL = lambda x: ((x) << 26) # type: ignore
@@ -472,16 +472,16 @@ PACKET3_UNMAP_QUEUES_DOORBELL_OFFSET1 = lambda x: ((x) << 2) # type: ignore
 PACKET3_UNMAP_QUEUES_RB_WPTR = lambda x: ((x) << 0) # type: ignore
 PACKET3_UNMAP_QUEUES_DOORBELL_OFFSET2 = lambda x: ((x) << 2) # type: ignore
 PACKET3_UNMAP_QUEUES_DOORBELL_OFFSET3 = lambda x: ((x) << 2) # type: ignore
-PACKET3_QUERY_STATUS = 0xA4 # type: ignore
+PACKET3_QUERY_STATUS = 0xA4
 PACKET3_QUERY_STATUS_CONTEXT_ID = lambda x: ((x) << 0) # type: ignore
 PACKET3_QUERY_STATUS_INTERRUPT_SEL = lambda x: ((x) << 28) # type: ignore
 PACKET3_QUERY_STATUS_COMMAND = lambda x: ((x) << 30) # type: ignore
 PACKET3_QUERY_STATUS_PASID = lambda x: ((x) << 0) # type: ignore
 PACKET3_QUERY_STATUS_DOORBELL_OFFSET = lambda x: ((x) << 2) # type: ignore
 PACKET3_QUERY_STATUS_ENG_SEL = lambda x: ((x) << 25) # type: ignore
-PACKET3_RUN_LIST = 0xA5 # type: ignore
-PACKET3_MAP_PROCESS_VM = 0xA6 # type: ignore
-PACKET3_RUN_CLEANER_SHADER = 0xD2 # type: ignore
-PACKET3_SET_Q_PREEMPTION_MODE = 0xF0 # type: ignore
+PACKET3_RUN_LIST = 0xA5
+PACKET3_MAP_PROCESS_VM = 0xA6
+PACKET3_RUN_CLEANER_SHADER = 0xD2
+PACKET3_SET_Q_PREEMPTION_MODE = 0xF0
 PACKET3_SET_Q_PREEMPTION_MODE_IB_VMID = lambda x: ((x) << 0) # type: ignore
-PACKET3_SET_Q_PREEMPTION_MODE_INIT_SHADOW_MEM = (1 << 0) # type: ignore
+PACKET3_SET_Q_PREEMPTION_MODE_INIT_SHADOW_MEM = (1 << 0)

tinygrad/runtime/autogen/am/pm4_soc15.py

@@ -37,55 +37,55 @@ enum_WRITE_DATA_wr_confirm_enum: dict[int, str] = {(wr_confirm___write_data__do_
 enum_WRITE_DATA_cache_policy_enum: dict[int, str] = {(cache_policy___write_data__lru:=0): 'cache_policy___write_data__lru', (cache_policy___write_data__stream:=1): 'cache_policy___write_data__stream'}
 class struct_pm4_mec_write_data_mmio(c.Struct): pass
 _anonenum0: dict[int, str] = {(CACHE_FLUSH_AND_INV_TS_EVENT:=20): 'CACHE_FLUSH_AND_INV_TS_EVENT'}
-GFX9_NUM_GFX_RINGS = 1 # type: ignore
-GFX9_NUM_COMPUTE_RINGS = 8 # type: ignore
-PACKET_TYPE0 = 0 # type: ignore
-PACKET_TYPE1 = 1 # type: ignore
-PACKET_TYPE2 = 2 # type: ignore
-PACKET_TYPE3 = 3 # type: ignore
+GFX9_NUM_GFX_RINGS = 1
+GFX9_NUM_COMPUTE_RINGS = 8
+PACKET_TYPE0 = 0
+PACKET_TYPE1 = 1
+PACKET_TYPE2 = 2
+PACKET_TYPE3 = 3
 CP_PACKET_GET_TYPE = lambda h: (((h) >> 30) & 3) # type: ignore
 CP_PACKET_GET_COUNT = lambda h: (((h) >> 16) & 0x3FFF) # type: ignore
 CP_PACKET0_GET_REG = lambda h: ((h) & 0xFFFF) # type: ignore
 CP_PACKET3_GET_OPCODE = lambda h: (((h) >> 8) & 0xFF) # type: ignore
 PACKET0 = lambda reg,n: ((PACKET_TYPE0 << 30) | ((reg) & 0xFFFF) | ((n) & 0x3FFF) << 16) # type: ignore
-CP_PACKET2 = 0x80000000 # type: ignore
-PACKET2_PAD_SHIFT = 0 # type: ignore
-PACKET2_PAD_MASK = (0x3fffffff << 0) # type: ignore
+CP_PACKET2 = 0x80000000
+PACKET2_PAD_SHIFT = 0
+PACKET2_PAD_MASK = (0x3fffffff << 0)
 PACKET2 = lambda v: (CP_PACKET2 | REG_SET(PACKET2_PAD, (v))) # type: ignore
 PACKET3 = lambda op,n: ((PACKET_TYPE3 << 30) | (((op) & 0xFF) << 8) | ((n) & 0x3FFF) << 16) # type: ignore
 PACKET3_COMPUTE = lambda op,n: (PACKET3(op, n) | 1 << 1) # type: ignore
-PACKETJ_CONDITION_CHECK0 = 0 # type: ignore
-PACKETJ_CONDITION_CHECK1 = 1 # type: ignore
-PACKETJ_CONDITION_CHECK2 = 2 # type: ignore
-PACKETJ_CONDITION_CHECK3 = 3 # type: ignore
-PACKETJ_CONDITION_CHECK4 = 4 # type: ignore
-PACKETJ_CONDITION_CHECK5 = 5 # type: ignore
-PACKETJ_CONDITION_CHECK6 = 6 # type: ignore
-PACKETJ_CONDITION_CHECK7 = 7 # type: ignore
-PACKETJ_TYPE0 = 0 # type: ignore
-PACKETJ_TYPE1 = 1 # type: ignore
-PACKETJ_TYPE2 = 2 # type: ignore
-PACKETJ_TYPE3 = 3 # type: ignore
-PACKETJ_TYPE4 = 4 # type: ignore
-PACKETJ_TYPE5 = 5 # type: ignore
-PACKETJ_TYPE6 = 6 # type: ignore
-PACKETJ_TYPE7 = 7 # type: ignore
+PACKETJ_CONDITION_CHECK0 = 0
+PACKETJ_CONDITION_CHECK1 = 1
+PACKETJ_CONDITION_CHECK2 = 2
+PACKETJ_CONDITION_CHECK3 = 3
+PACKETJ_CONDITION_CHECK4 = 4
+PACKETJ_CONDITION_CHECK5 = 5
+PACKETJ_CONDITION_CHECK6 = 6
+PACKETJ_CONDITION_CHECK7 = 7
+PACKETJ_TYPE0 = 0
+PACKETJ_TYPE1 = 1
+PACKETJ_TYPE2 = 2
+PACKETJ_TYPE3 = 3
+PACKETJ_TYPE4 = 4
+PACKETJ_TYPE5 = 5
+PACKETJ_TYPE6 = 6
+PACKETJ_TYPE7 = 7
 PACKETJ = lambda reg,r,cond,type: ((reg & 0x3FFFF) | ((r & 0x3F) << 18) | ((cond & 0xF) << 24) | ((type & 0xF) << 28)) # type: ignore
-CP_PACKETJ_NOP = 0x60000000 # type: ignore
+CP_PACKETJ_NOP = 0x60000000
 CP_PACKETJ_GET_REG = lambda x: ((x) & 0x3FFFF) # type: ignore
 CP_PACKETJ_GET_RES = lambda x: (((x) >> 18) & 0x3F) # type: ignore
 CP_PACKETJ_GET_COND = lambda x: (((x) >> 24) & 0xF) # type: ignore
 CP_PACKETJ_GET_TYPE = lambda x: (((x) >> 28) & 0xF) # type: ignore
-PACKET3_NOP = 0x10 # type: ignore
-PACKET3_SET_BASE = 0x11 # type: ignore
+PACKET3_NOP = 0x10
+PACKET3_SET_BASE = 0x11
 PACKET3_BASE_INDEX = lambda x: ((x) << 0) # type: ignore
-CE_PARTITION_BASE = 3 # type: ignore
-PACKET3_CLEAR_STATE = 0x12 # type: ignore
-PACKET3_INDEX_BUFFER_SIZE = 0x13 # type: ignore
-PACKET3_DISPATCH_DIRECT = 0x15 # type: ignore
-PACKET3_DISPATCH_INDIRECT = 0x16 # type: ignore
-PACKET3_ATOMIC_GDS = 0x1D # type: ignore
-PACKET3_ATOMIC_MEM = 0x1E # type: ignore
+CE_PARTITION_BASE = 3
+PACKET3_CLEAR_STATE = 0x12
+PACKET3_INDEX_BUFFER_SIZE = 0x13
+PACKET3_DISPATCH_DIRECT = 0x15
+PACKET3_DISPATCH_INDIRECT = 0x16
+PACKET3_ATOMIC_GDS = 0x1D
+PACKET3_ATOMIC_MEM = 0x1E
 PACKET3_ATOMIC_MEM__ATOMIC = lambda x: ((((unsigned)(x)) & 0x3F) << 0) # type: ignore
 PACKET3_ATOMIC_MEM__COMMAND = lambda x: ((((unsigned)(x)) & 0xF) << 8) # type: ignore
 PACKET3_ATOMIC_MEM__CACHE_POLICY = lambda x: ((((unsigned)(x)) & 0x3) << 25) # type: ignore
@@ -96,33 +96,33 @@ PACKET3_ATOMIC_MEM__SRC_DATA_HI = lambda x: (((unsigned)(x)) << 0) # type: ignor
 PACKET3_ATOMIC_MEM__CMP_DATA_LO = lambda x: (((unsigned)(x)) << 0) # type: ignore
 PACKET3_ATOMIC_MEM__CMP_DATA_HI = lambda x: (((unsigned)(x)) << 0) # type: ignore
 PACKET3_ATOMIC_MEM__LOOP_INTERVAL = lambda x: ((((unsigned)(x)) & 0x1FFF) << 0) # type: ignore
-PACKET3_ATOMIC_MEM__COMMAND__SINGLE_PASS_ATOMIC = 0 # type: ignore
-PACKET3_ATOMIC_MEM__COMMAND__LOOP_UNTIL_COMPARE_SATISFIED = 1 # type: ignore
-PACKET3_OCCLUSION_QUERY = 0x1F # type: ignore
-PACKET3_SET_PREDICATION = 0x20 # type: ignore
-PACKET3_REG_RMW = 0x21 # type: ignore
-PACKET3_COND_EXEC = 0x22 # type: ignore
-PACKET3_PRED_EXEC = 0x23 # type: ignore
+PACKET3_ATOMIC_MEM__COMMAND__SINGLE_PASS_ATOMIC = 0
+PACKET3_ATOMIC_MEM__COMMAND__LOOP_UNTIL_COMPARE_SATISFIED = 1
+PACKET3_OCCLUSION_QUERY = 0x1F
+PACKET3_SET_PREDICATION = 0x20
+PACKET3_REG_RMW = 0x21
+PACKET3_COND_EXEC = 0x22
+PACKET3_PRED_EXEC = 0x23
 PACKET3_PRED_EXEC__EXEC_COUNT = lambda x: ((((unsigned)(x)) & 0x3FFF) << 0) # type: ignore
 PACKET3_PRED_EXEC__VIRTUAL_XCC_ID_SELECT = lambda x: ((((unsigned)(x)) & 0xFF) << 24) # type: ignore
-PACKET3_DRAW_INDIRECT = 0x24 # type: ignore
-PACKET3_DRAW_INDEX_INDIRECT = 0x25 # type: ignore
-PACKET3_INDEX_BASE = 0x26 # type: ignore
-PACKET3_DRAW_INDEX_2 = 0x27 # type: ignore
-PACKET3_CONTEXT_CONTROL = 0x28 # type: ignore
-PACKET3_INDEX_TYPE = 0x2A # type: ignore
-PACKET3_DRAW_INDIRECT_MULTI = 0x2C # type: ignore
-PACKET3_DRAW_INDEX_AUTO = 0x2D # type: ignore
-PACKET3_NUM_INSTANCES = 0x2F # type: ignore
-PACKET3_DRAW_INDEX_MULTI_AUTO = 0x30 # type: ignore
-PACKET3_INDIRECT_BUFFER_CONST = 0x33 # type: ignore
-PACKET3_STRMOUT_BUFFER_UPDATE = 0x34 # type: ignore
-PACKET3_DRAW_INDEX_OFFSET_2 = 0x35 # type: ignore
-PACKET3_DRAW_PREAMBLE = 0x36 # type: ignore
-PACKET3_WRITE_DATA = 0x37 # type: ignore
+PACKET3_DRAW_INDIRECT = 0x24
+PACKET3_DRAW_INDEX_INDIRECT = 0x25
+PACKET3_INDEX_BASE = 0x26
+PACKET3_DRAW_INDEX_2 = 0x27
+PACKET3_CONTEXT_CONTROL = 0x28
+PACKET3_INDEX_TYPE = 0x2A
+PACKET3_DRAW_INDIRECT_MULTI = 0x2C
+PACKET3_DRAW_INDEX_AUTO = 0x2D
+PACKET3_NUM_INSTANCES = 0x2F
+PACKET3_DRAW_INDEX_MULTI_AUTO = 0x30
+PACKET3_INDIRECT_BUFFER_CONST = 0x33
+PACKET3_STRMOUT_BUFFER_UPDATE = 0x34
+PACKET3_DRAW_INDEX_OFFSET_2 = 0x35
+PACKET3_DRAW_PREAMBLE = 0x36
+PACKET3_WRITE_DATA = 0x37
 WRITE_DATA_DST_SEL = lambda x: ((x) << 8) # type: ignore
-WR_ONE_ADDR = (1 << 16) # type: ignore
-WR_CONFIRM = (1 << 20) # type: ignore
+WR_ONE_ADDR = (1 << 16)
+WR_CONFIRM = (1 << 20)
 WRITE_DATA_CACHE_POLICY = lambda x: ((x) << 25) # type: ignore
 WRITE_DATA_ENGINE_SEL = lambda x: ((x) << 30) # type: ignore
 PACKET3_WRITE_DATA__DST_SEL = lambda x: ((((unsigned)(x)) & 0xF) << 8) # type: ignore
@@ -134,26 +134,26 @@ PACKET3_WRITE_DATA__DST_MMREG_ADDR = lambda x: ((((unsigned)(x)) & 0x3FFFF) << 0
 PACKET3_WRITE_DATA__DST_GDS_ADDR = lambda x: ((((unsigned)(x)) & 0xFFFF) << 0) # type: ignore
 PACKET3_WRITE_DATA__DST_MEM_ADDR_LO = lambda x: ((((unsigned)(x)) & 0x3FFFFFFF) << 2) # type: ignore
 PACKET3_WRITE_DATA__DST_MEM_ADDR_HI = lambda x: ((unsigned)(x)) # type: ignore
-PACKET3_WRITE_DATA__DST_SEL__MEM_MAPPED_REGISTER = 0 # type: ignore
-PACKET3_WRITE_DATA__DST_SEL__TC_L2 = 2 # type: ignore
-PACKET3_WRITE_DATA__DST_SEL__GDS = 3 # type: ignore
-PACKET3_WRITE_DATA__DST_SEL__MEMORY = 5 # type: ignore
-PACKET3_WRITE_DATA__DST_SEL__MEMORY_MAPPED_ADC_PERSISTENT_STATE = 6 # type: ignore
-PACKET3_WRITE_DATA__ADDR_INCR__INCREMENT_ADDRESS = 0 # type: ignore
-PACKET3_WRITE_DATA__ADDR_INCR__DO_NOT_INCREMENT_ADDRESS = 1 # type: ignore
-PACKET3_WRITE_DATA__WR_CONFIRM__DO_NOT_WAIT_FOR_WRITE_CONFIRMATION = 0 # type: ignore
-PACKET3_WRITE_DATA__WR_CONFIRM__WAIT_FOR_WRITE_CONFIRMATION = 1 # type: ignore
-PACKET3_WRITE_DATA__CACHE_POLICY__LRU = 0 # type: ignore
-PACKET3_WRITE_DATA__CACHE_POLICY__STREAM = 1 # type: ignore
-PACKET3_WRITE_DATA__CACHE_POLICY__NOA = 2 # type: ignore
-PACKET3_WRITE_DATA__CACHE_POLICY__BYPASS = 3 # type: ignore
-PACKET3_DRAW_INDEX_INDIRECT_MULTI = 0x38 # type: ignore
-PACKET3_MEM_SEMAPHORE = 0x39 # type: ignore
-PACKET3_SEM_USE_MAILBOX = (0x1 << 16) # type: ignore
-PACKET3_SEM_SEL_SIGNAL_TYPE = (0x1 << 20) # type: ignore
-PACKET3_SEM_SEL_SIGNAL = (0x6 << 29) # type: ignore
-PACKET3_SEM_SEL_WAIT = (0x7 << 29) # type: ignore
-PACKET3_WAIT_REG_MEM = 0x3C # type: ignore
+PACKET3_WRITE_DATA__DST_SEL__MEM_MAPPED_REGISTER = 0
+PACKET3_WRITE_DATA__DST_SEL__TC_L2 = 2
+PACKET3_WRITE_DATA__DST_SEL__GDS = 3
+PACKET3_WRITE_DATA__DST_SEL__MEMORY = 5
+PACKET3_WRITE_DATA__DST_SEL__MEMORY_MAPPED_ADC_PERSISTENT_STATE = 6
+PACKET3_WRITE_DATA__ADDR_INCR__INCREMENT_ADDRESS = 0
+PACKET3_WRITE_DATA__ADDR_INCR__DO_NOT_INCREMENT_ADDRESS = 1
+PACKET3_WRITE_DATA__WR_CONFIRM__DO_NOT_WAIT_FOR_WRITE_CONFIRMATION = 0
+PACKET3_WRITE_DATA__WR_CONFIRM__WAIT_FOR_WRITE_CONFIRMATION = 1
+PACKET3_WRITE_DATA__CACHE_POLICY__LRU = 0
+PACKET3_WRITE_DATA__CACHE_POLICY__STREAM = 1
+PACKET3_WRITE_DATA__CACHE_POLICY__NOA = 2
+PACKET3_WRITE_DATA__CACHE_POLICY__BYPASS = 3
+PACKET3_DRAW_INDEX_INDIRECT_MULTI = 0x38
+PACKET3_MEM_SEMAPHORE = 0x39
+PACKET3_SEM_USE_MAILBOX = (0x1 << 16)
+PACKET3_SEM_SEL_SIGNAL_TYPE = (0x1 << 20)
+PACKET3_SEM_SEL_SIGNAL = (0x6 << 29)
+PACKET3_SEM_SEL_WAIT = (0x7 << 29)
+PACKET3_WAIT_REG_MEM = 0x3C
 WAIT_REG_MEM_FUNCTION = lambda x: ((x) << 0) # type: ignore
 WAIT_REG_MEM_MEM_SPACE = lambda x: ((x) << 4) # type: ignore
 WAIT_REG_MEM_OPERATION = lambda x: ((x) << 6) # type: ignore
@@ -173,20 +173,20 @@ PACKET3_WAIT_REG_MEM__REFERENCE = lambda x: ((unsigned)(x)) # type: ignore
 PACKET3_WAIT_REG_MEM__MASK = lambda x: ((unsigned)(x)) # type: ignore
 PACKET3_WAIT_REG_MEM__POLL_INTERVAL = lambda x: ((((unsigned)(x)) & 0xFFFF) << 0) # type: ignore
 PACKET3_WAIT_REG_MEM__OPTIMIZE_ACE_OFFLOAD_MODE = lambda x: ((((unsigned)(x)) & 0x1) << 31) # type: ignore
-PACKET3_WAIT_REG_MEM__FUNCTION__ALWAYS_PASS = 0 # type: ignore
-PACKET3_WAIT_REG_MEM__FUNCTION__LESS_THAN_REF_VALUE = 1 # type: ignore
-PACKET3_WAIT_REG_MEM__FUNCTION__LESS_THAN_EQUAL_TO_THE_REF_VALUE = 2 # type: ignore
-PACKET3_WAIT_REG_MEM__FUNCTION__EQUAL_TO_THE_REFERENCE_VALUE = 3 # type: ignore
-PACKET3_WAIT_REG_MEM__FUNCTION__NOT_EQUAL_REFERENCE_VALUE = 4 # type: ignore
-PACKET3_WAIT_REG_MEM__FUNCTION__GREATER_THAN_OR_EQUAL_REFERENCE_VALUE = 5 # type: ignore
-PACKET3_WAIT_REG_MEM__FUNCTION__GREATER_THAN_REFERENCE_VALUE = 6 # type: ignore
-PACKET3_WAIT_REG_MEM__MEM_SPACE__REGISTER_SPACE = 0 # type: ignore
-PACKET3_WAIT_REG_MEM__MEM_SPACE__MEMORY_SPACE = 1 # type: ignore
-PACKET3_WAIT_REG_MEM__OPERATION__WAIT_REG_MEM = 0 # type: ignore
-PACKET3_WAIT_REG_MEM__OPERATION__WR_WAIT_WR_REG = 1 # type: ignore
-PACKET3_WAIT_REG_MEM__OPERATION__WAIT_MEM_PREEMPTABLE = 3 # type: ignore
-PACKET3_INDIRECT_BUFFER = 0x3F # type: ignore
-INDIRECT_BUFFER_VALID = (1 << 23) # type: ignore
+PACKET3_WAIT_REG_MEM__FUNCTION__ALWAYS_PASS = 0
+PACKET3_WAIT_REG_MEM__FUNCTION__LESS_THAN_REF_VALUE = 1
+PACKET3_WAIT_REG_MEM__FUNCTION__LESS_THAN_EQUAL_TO_THE_REF_VALUE = 2
+PACKET3_WAIT_REG_MEM__FUNCTION__EQUAL_TO_THE_REFERENCE_VALUE = 3
+PACKET3_WAIT_REG_MEM__FUNCTION__NOT_EQUAL_REFERENCE_VALUE = 4
+PACKET3_WAIT_REG_MEM__FUNCTION__GREATER_THAN_OR_EQUAL_REFERENCE_VALUE = 5
+PACKET3_WAIT_REG_MEM__FUNCTION__GREATER_THAN_REFERENCE_VALUE = 6
+PACKET3_WAIT_REG_MEM__MEM_SPACE__REGISTER_SPACE = 0
+PACKET3_WAIT_REG_MEM__MEM_SPACE__MEMORY_SPACE = 1
+PACKET3_WAIT_REG_MEM__OPERATION__WAIT_REG_MEM = 0
+PACKET3_WAIT_REG_MEM__OPERATION__WR_WAIT_WR_REG = 1
+PACKET3_WAIT_REG_MEM__OPERATION__WAIT_MEM_PREEMPTABLE = 3
+PACKET3_INDIRECT_BUFFER = 0x3F
+INDIRECT_BUFFER_VALID = (1 << 23)
 INDIRECT_BUFFER_CACHE_POLICY = lambda x: ((x) << 28) # type: ignore
 INDIRECT_BUFFER_PRE_ENB = lambda x: ((x) << 21) # type: ignore
 INDIRECT_BUFFER_PRE_RESUME = lambda x: ((x) << 30) # type: ignore
@@ -199,9 +199,9 @@ PACKET3_INDIRECT_BUFFER__VALID = lambda x: ((((unsigned)(x)) & 0x1) << 23) # typ
 PACKET3_INDIRECT_BUFFER__VMID = lambda x: ((((unsigned)(x)) & 0xF) << 24) # type: ignore
 PACKET3_INDIRECT_BUFFER__CACHE_POLICY = lambda x: ((((unsigned)(x)) & 0x3) << 28) # type: ignore
 PACKET3_INDIRECT_BUFFER__PRIV = lambda x: ((((unsigned)(x)) & 0x1) << 31) # type: ignore
-PACKET3_INDIRECT_BUFFER__CACHE_POLICY__LRU = 0 # type: ignore
-PACKET3_INDIRECT_BUFFER__CACHE_POLICY__STREAM = 1 # type: ignore
-PACKET3_COPY_DATA = 0x40 # type: ignore
+PACKET3_INDIRECT_BUFFER__CACHE_POLICY__LRU = 0
+PACKET3_INDIRECT_BUFFER__CACHE_POLICY__STREAM = 1
+PACKET3_COPY_DATA = 0x40
 PACKET3_COPY_DATA__SRC_SEL = lambda x: ((((unsigned)(x)) & 0xF) << 0) # type: ignore
 PACKET3_COPY_DATA__DST_SEL = lambda x: ((((unsigned)(x)) & 0xF) << 8) # type: ignore
 PACKET3_COPY_DATA__SRC_CACHE_POLICY = lambda x: ((((unsigned)(x)) & 0x3) << 13) # type: ignore
@@ -221,35 +221,35 @@ PACKET3_COPY_DATA__DST_32B_ADDR_LO = lambda x: ((((unsigned)(x)) & 0x3FFFFFFF) <
 PACKET3_COPY_DATA__DST_64B_ADDR_LO = lambda x: ((((unsigned)(x)) & 0x1FFFFFFF) << 3) # type: ignore
 PACKET3_COPY_DATA__DST_GDS_ADDR_LO = lambda x: ((((unsigned)(x)) & 0xFFFF) << 0) # type: ignore
 PACKET3_COPY_DATA__DST_ADDR_HI = lambda x: ((unsigned)(x)) # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__MEM_MAPPED_REGISTER = 0 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__MEMORY = 1 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__TC_L2 = 2 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__GDS = 3 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__PERFCOUNTERS = 4 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__IMMEDIATE_DATA = 5 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__ATOMIC_RETURN_DATA = 6 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__GDS_ATOMIC_RETURN_DATA0 = 7 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__GDS_ATOMIC_RETURN_DATA1 = 8 # type: ignore
-PACKET3_COPY_DATA__SRC_SEL__GPU_CLOCK_COUNT = 9 # type: ignore
-PACKET3_COPY_DATA__DST_SEL__MEM_MAPPED_REGISTER = 0 # type: ignore
-PACKET3_COPY_DATA__DST_SEL__TC_L2 = 2 # type: ignore
-PACKET3_COPY_DATA__DST_SEL__GDS = 3 # type: ignore
-PACKET3_COPY_DATA__DST_SEL__PERFCOUNTERS = 4 # type: ignore
-PACKET3_COPY_DATA__DST_SEL__MEMORY = 5 # type: ignore
-PACKET3_COPY_DATA__DST_SEL__MEM_MAPPED_REG_DC = 6 # type: ignore
-PACKET3_COPY_DATA__SRC_CACHE_POLICY__LRU = 0 # type: ignore
-PACKET3_COPY_DATA__SRC_CACHE_POLICY__STREAM = 1 # type: ignore
-PACKET3_COPY_DATA__COUNT_SEL__32_BITS_OF_DATA = 0 # type: ignore
-PACKET3_COPY_DATA__COUNT_SEL__64_BITS_OF_DATA = 1 # type: ignore
-PACKET3_COPY_DATA__WR_CONFIRM__DO_NOT_WAIT_FOR_CONFIRMATION = 0 # type: ignore
-PACKET3_COPY_DATA__WR_CONFIRM__WAIT_FOR_CONFIRMATION = 1 # type: ignore
-PACKET3_COPY_DATA__DST_CACHE_POLICY__LRU = 0 # type: ignore
-PACKET3_COPY_DATA__DST_CACHE_POLICY__STREAM = 1 # type: ignore
-PACKET3_COPY_DATA__PQ_EXE_STATUS__DEFAULT = 0 # type: ignore
-PACKET3_COPY_DATA__PQ_EXE_STATUS__PHASE_UPDATE = 1 # type: ignore
-PACKET3_PFP_SYNC_ME = 0x42 # type: ignore
-PACKET3_COND_WRITE = 0x45 # type: ignore
-PACKET3_EVENT_WRITE = 0x46 # type: ignore
+PACKET3_COPY_DATA__SRC_SEL__MEM_MAPPED_REGISTER = 0
+PACKET3_COPY_DATA__SRC_SEL__MEMORY = 1
+PACKET3_COPY_DATA__SRC_SEL__TC_L2 = 2
+PACKET3_COPY_DATA__SRC_SEL__GDS = 3
+PACKET3_COPY_DATA__SRC_SEL__PERFCOUNTERS = 4
+PACKET3_COPY_DATA__SRC_SEL__IMMEDIATE_DATA = 5
+PACKET3_COPY_DATA__SRC_SEL__ATOMIC_RETURN_DATA = 6
+PACKET3_COPY_DATA__SRC_SEL__GDS_ATOMIC_RETURN_DATA0 = 7
+PACKET3_COPY_DATA__SRC_SEL__GDS_ATOMIC_RETURN_DATA1 = 8
+PACKET3_COPY_DATA__SRC_SEL__GPU_CLOCK_COUNT = 9
+PACKET3_COPY_DATA__DST_SEL__MEM_MAPPED_REGISTER = 0
+PACKET3_COPY_DATA__DST_SEL__TC_L2 = 2
+PACKET3_COPY_DATA__DST_SEL__GDS = 3
+PACKET3_COPY_DATA__DST_SEL__PERFCOUNTERS = 4
+PACKET3_COPY_DATA__DST_SEL__MEMORY = 5
+PACKET3_COPY_DATA__DST_SEL__MEM_MAPPED_REG_DC = 6
+PACKET3_COPY_DATA__SRC_CACHE_POLICY__LRU = 0
+PACKET3_COPY_DATA__SRC_CACHE_POLICY__STREAM = 1
+PACKET3_COPY_DATA__COUNT_SEL__32_BITS_OF_DATA = 0
+PACKET3_COPY_DATA__COUNT_SEL__64_BITS_OF_DATA = 1
+PACKET3_COPY_DATA__WR_CONFIRM__DO_NOT_WAIT_FOR_CONFIRMATION = 0
+PACKET3_COPY_DATA__WR_CONFIRM__WAIT_FOR_CONFIRMATION = 1
+PACKET3_COPY_DATA__DST_CACHE_POLICY__LRU = 0
+PACKET3_COPY_DATA__DST_CACHE_POLICY__STREAM = 1
+PACKET3_COPY_DATA__PQ_EXE_STATUS__DEFAULT = 0
+PACKET3_COPY_DATA__PQ_EXE_STATUS__PHASE_UPDATE = 1
+PACKET3_PFP_SYNC_ME = 0x42
+PACKET3_COND_WRITE = 0x45
+PACKET3_EVENT_WRITE = 0x46
 EVENT_TYPE = lambda x: ((x) << 0) # type: ignore
 EVENT_INDEX = lambda x: ((x) << 8) # type: ignore
 PACKET3_EVENT_WRITE__EVENT_TYPE = lambda x: ((((unsigned)(x)) & 0x3F) << 0) # type: ignore
@@ -258,39 +258,39 @@ PACKET3_EVENT_WRITE__OFFLOAD_ENABLE = lambda x: ((((unsigned)(x)) & 0x1) << 31)
 PACKET3_EVENT_WRITE__SAMP_PLST_CNTR_MODE = lambda x: ((((unsigned)(x)) & 0x3) << 29) # type: ignore
 PACKET3_EVENT_WRITE__ADDRESS_LO = lambda x: ((((unsigned)(x)) & 0x1FFFFFFF) << 3) # type: ignore
 PACKET3_EVENT_WRITE__ADDRESS_HI = lambda x: (((unsigned)(x)) << 0) # type: ignore
-PACKET3_EVENT_WRITE__EVENT_INDEX__OTHER = 0 # type: ignore
-PACKET3_EVENT_WRITE__EVENT_INDEX__SAMPLE_PIPELINESTATS = 2 # type: ignore
-PACKET3_EVENT_WRITE__EVENT_INDEX__CS_PARTIAL_FLUSH = 4 # type: ignore
-PACKET3_RELEASE_MEM = 0x49 # type: ignore
+PACKET3_EVENT_WRITE__EVENT_INDEX__OTHER = 0
+PACKET3_EVENT_WRITE__EVENT_INDEX__SAMPLE_PIPELINESTATS = 2
+PACKET3_EVENT_WRITE__EVENT_INDEX__CS_PARTIAL_FLUSH = 4
+PACKET3_RELEASE_MEM = 0x49
 EVENT_TYPE = lambda x: ((x) << 0) # type: ignore
 EVENT_INDEX = lambda x: ((x) << 8) # type: ignore
-EOP_TCL1_VOL_ACTION_EN = (1 << 12) # type: ignore
-EOP_TC_VOL_ACTION_EN = (1 << 13) # type: ignore
-EOP_TC_WB_ACTION_EN = (1 << 15) # type: ignore
-EOP_TCL1_ACTION_EN = (1 << 16) # type: ignore
-EOP_TC_ACTION_EN = (1 << 17) # type: ignore
-EOP_TC_NC_ACTION_EN = (1 << 19) # type: ignore
-EOP_TC_MD_ACTION_EN = (1 << 21) # type: ignore
-EOP_EXEC = (1 << 28) # type: ignore
+EOP_TCL1_VOL_ACTION_EN = (1 << 12)
+EOP_TC_VOL_ACTION_EN = (1 << 13)
+EOP_TC_WB_ACTION_EN = (1 << 15)
+EOP_TCL1_ACTION_EN = (1 << 16)
+EOP_TC_ACTION_EN = (1 << 17)
+EOP_TC_NC_ACTION_EN = (1 << 19)
+EOP_TC_MD_ACTION_EN = (1 << 21)
+EOP_EXEC = (1 << 28)
 DATA_SEL = lambda x: ((x) << 29) # type: ignore
 INT_SEL = lambda x: ((x) << 24) # type: ignore
 DST_SEL = lambda x: ((x) << 16) # type: ignore
-PACKET3_PREAMBLE_CNTL = 0x4A # type: ignore
-PACKET3_PREAMBLE_BEGIN_CLEAR_STATE = (2 << 28) # type: ignore
-PACKET3_PREAMBLE_END_CLEAR_STATE = (3 << 28) # type: ignore
-PACKET3_DMA_DATA = 0x50 # type: ignore
+PACKET3_PREAMBLE_CNTL = 0x4A
+PACKET3_PREAMBLE_BEGIN_CLEAR_STATE = (2 << 28)
+PACKET3_PREAMBLE_END_CLEAR_STATE = (3 << 28)
+PACKET3_DMA_DATA = 0x50
 PACKET3_DMA_DATA_ENGINE = lambda x: ((x) << 0) # type: ignore
 PACKET3_DMA_DATA_SRC_CACHE_POLICY = lambda x: ((x) << 13) # type: ignore
 PACKET3_DMA_DATA_DST_SEL = lambda x: ((x) << 20) # type: ignore
 PACKET3_DMA_DATA_DST_CACHE_POLICY = lambda x: ((x) << 25) # type: ignore
 PACKET3_DMA_DATA_SRC_SEL = lambda x: ((x) << 29) # type: ignore
-PACKET3_DMA_DATA_CP_SYNC = (1 << 31) # type: ignore
-PACKET3_DMA_DATA_CMD_SAS = (1 << 26) # type: ignore
-PACKET3_DMA_DATA_CMD_DAS = (1 << 27) # type: ignore
-PACKET3_DMA_DATA_CMD_SAIC = (1 << 28) # type: ignore
-PACKET3_DMA_DATA_CMD_DAIC = (1 << 29) # type: ignore
-PACKET3_DMA_DATA_CMD_RAW_WAIT = (1 << 30) # type: ignore
-PACKET3_ACQUIRE_MEM = 0x58 # type: ignore
+PACKET3_DMA_DATA_CP_SYNC = (1 << 31)
+PACKET3_DMA_DATA_CMD_SAS = (1 << 26)
+PACKET3_DMA_DATA_CMD_DAS = (1 << 27)
+PACKET3_DMA_DATA_CMD_SAIC = (1 << 28)
+PACKET3_DMA_DATA_CMD_DAIC = (1 << 29)
+PACKET3_DMA_DATA_CMD_RAW_WAIT = (1 << 30)
+PACKET3_ACQUIRE_MEM = 0x58
 PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_TC_NC_ACTION_ENA = lambda x: ((x) << 3) # type: ignore
 PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_TC_WC_ACTION_ENA = lambda x: ((x) << 4) # type: ignore
 PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_TC_INV_METADATA_ACTION_ENA = lambda x: ((x) << 5) # type: ignore
@@ -304,7 +304,7 @@ PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_SH_KCACHE_ACTION_ENA = lambda x: ((x) << 27) #
 PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_SH_KCACHE_VOL_ACTION_ENA = lambda x: ((x) << 28) # type: ignore
 PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_SH_ICACHE_ACTION_ENA = lambda x: ((x) << 29) # type: ignore
 PACKET3_ACQUIRE_MEM_CP_COHER_CNTL_SH_KCACHE_WB_ACTION_ENA = lambda x: ((x) << 30) # type: ignore
-PACKET3_REWIND = 0x59 # type: ignore
+PACKET3_REWIND = 0x59
 PACKET3_ACQUIRE_MEM__COHER_SIZE = lambda x: ((unsigned)(x)) # type: ignore
 PACKET3_ACQUIRE_MEM__COHER_SIZE_HI = lambda x: ((((unsigned)(x)) & 0xFF) << 0) # type: ignore
 PACKET3_ACQUIRE_MEM__COHER_SIZE_HI_VG10 = lambda x: ((((unsigned)(x)) & 0xFFFFFF) << 0) # type: ignore
@@ -312,53 +312,53 @@ PACKET3_ACQUIRE_MEM__COHER_BASE_LO = lambda x: ((unsigned)(x)) # type: ignore
 PACKET3_ACQUIRE_MEM__COHER_BASE_HI = lambda x: ((((unsigned)(x)) & 0xFFFFFF) << 0) # type: ignore
 PACKET3_ACQUIRE_MEM__POLL_INTERVAL = lambda x: ((((unsigned)(x)) & 0xFFFF) << 0) # type: ignore
 PACKET3_ACQUIRE_MEM__GCR_CNTL = lambda x: ((((unsigned)(x)) & 0x7FF) << 0) # type: ignore
-PACKET3_LOAD_UCONFIG_REG = 0x5E # type: ignore
-PACKET3_LOAD_SH_REG = 0x5F # type: ignore
-PACKET3_LOAD_CONFIG_REG = 0x60 # type: ignore
-PACKET3_LOAD_CONTEXT_REG = 0x61 # type: ignore
-PACKET3_SET_CONFIG_REG = 0x68 # type: ignore
-PACKET3_SET_CONFIG_REG_START = 0x00002000 # type: ignore
-PACKET3_SET_CONFIG_REG_END = 0x00002c00 # type: ignore
-PACKET3_SET_CONTEXT_REG = 0x69 # type: ignore
-PACKET3_SET_CONTEXT_REG_START = 0x0000a000 # type: ignore
-PACKET3_SET_CONTEXT_REG_END = 0x0000a400 # type: ignore
-PACKET3_SET_CONTEXT_REG_INDIRECT = 0x73 # type: ignore
-PACKET3_SET_SH_REG = 0x76 # type: ignore
-PACKET3_SET_SH_REG_START = 0x00002c00 # type: ignore
-PACKET3_SET_SH_REG_END = 0x00003000 # type: ignore
+PACKET3_LOAD_UCONFIG_REG = 0x5E
+PACKET3_LOAD_SH_REG = 0x5F
+PACKET3_LOAD_CONFIG_REG = 0x60
+PACKET3_LOAD_CONTEXT_REG = 0x61
+PACKET3_SET_CONFIG_REG = 0x68
+PACKET3_SET_CONFIG_REG_START = 0x00002000
+PACKET3_SET_CONFIG_REG_END = 0x00002c00
+PACKET3_SET_CONTEXT_REG = 0x69
+PACKET3_SET_CONTEXT_REG_START = 0x0000a000
+PACKET3_SET_CONTEXT_REG_END = 0x0000a400
+PACKET3_SET_CONTEXT_REG_INDIRECT = 0x73
+PACKET3_SET_SH_REG = 0x76
+PACKET3_SET_SH_REG_START = 0x00002c00
+PACKET3_SET_SH_REG_END = 0x00003000
 PACKET3_SET_SH_REG__REG_OFFSET = lambda x: ((((unsigned)(x)) & 0xFFFF) << 0) # type: ignore
 PACKET3_SET_SH_REG__VMID_SHIFT = lambda x: ((((unsigned)(x)) & 0x1F) << 23) # type: ignore
 PACKET3_SET_SH_REG__INDEX = lambda x: ((((unsigned)(x)) & 0xF) << 28) # type: ignore
-PACKET3_SET_SH_REG_OFFSET = 0x77 # type: ignore
-PACKET3_SET_QUEUE_REG = 0x78 # type: ignore
-PACKET3_SET_UCONFIG_REG = 0x79 # type: ignore
-PACKET3_SET_UCONFIG_REG_START = 0x0000c000 # type: ignore
-PACKET3_SET_UCONFIG_REG_END = 0x0000c400 # type: ignore
-PACKET3_SET_UCONFIG_REG_INDEX_TYPE = (2 << 28) # type: ignore
+PACKET3_SET_SH_REG_OFFSET = 0x77
+PACKET3_SET_QUEUE_REG = 0x78
+PACKET3_SET_UCONFIG_REG = 0x79
+PACKET3_SET_UCONFIG_REG_START = 0x0000c000
+PACKET3_SET_UCONFIG_REG_END = 0x0000c400
+PACKET3_SET_UCONFIG_REG_INDEX_TYPE = (2 << 28)
 PACKET3_SET_UCONFIG_REG__REG_OFFSET = lambda x: ((((unsigned)(x)) & 0xFFFF) << 0) # type: ignore
-PACKET3_SCRATCH_RAM_WRITE = 0x7D # type: ignore
-PACKET3_SCRATCH_RAM_READ = 0x7E # type: ignore
-PACKET3_LOAD_CONST_RAM = 0x80 # type: ignore
-PACKET3_WRITE_CONST_RAM = 0x81 # type: ignore
-PACKET3_DUMP_CONST_RAM = 0x83 # type: ignore
-PACKET3_INCREMENT_CE_COUNTER = 0x84 # type: ignore
-PACKET3_INCREMENT_DE_COUNTER = 0x85 # type: ignore
-PACKET3_WAIT_ON_CE_COUNTER = 0x86 # type: ignore
-PACKET3_WAIT_ON_DE_COUNTER_DIFF = 0x88 # type: ignore
-PACKET3_SWITCH_BUFFER = 0x8B # type: ignore
-PACKET3_FRAME_CONTROL = 0x90 # type: ignore
-FRAME_TMZ = (1 << 0) # type: ignore
+PACKET3_SCRATCH_RAM_WRITE = 0x7D
+PACKET3_SCRATCH_RAM_READ = 0x7E
+PACKET3_LOAD_CONST_RAM = 0x80
+PACKET3_WRITE_CONST_RAM = 0x81
+PACKET3_DUMP_CONST_RAM = 0x83
+PACKET3_INCREMENT_CE_COUNTER = 0x84
+PACKET3_INCREMENT_DE_COUNTER = 0x85
+PACKET3_WAIT_ON_CE_COUNTER = 0x86
+PACKET3_WAIT_ON_DE_COUNTER_DIFF = 0x88
+PACKET3_SWITCH_BUFFER = 0x8B
+PACKET3_FRAME_CONTROL = 0x90
+FRAME_TMZ = (1 << 0)
 FRAME_CMD = lambda x: ((x) << 28) # type: ignore
-PACKET3_INVALIDATE_TLBS = 0x98 # type: ignore
+PACKET3_INVALIDATE_TLBS = 0x98
 PACKET3_INVALIDATE_TLBS_DST_SEL = lambda x: ((x) << 0) # type: ignore
 PACKET3_INVALIDATE_TLBS_ALL_HUB = lambda x: ((x) << 4) # type: ignore
 PACKET3_INVALIDATE_TLBS_PASID = lambda x: ((x) << 5) # type: ignore
 PACKET3_INVALIDATE_TLBS_FLUSH_TYPE = lambda x: ((x) << 29) # type: ignore
-PACKET3_SET_RESOURCES = 0xA0 # type: ignore
+PACKET3_SET_RESOURCES = 0xA0
 PACKET3_SET_RESOURCES_VMID_MASK = lambda x: ((x) << 0) # type: ignore
 PACKET3_SET_RESOURCES_UNMAP_LATENTY = lambda x: ((x) << 16) # type: ignore
 PACKET3_SET_RESOURCES_QUEUE_TYPE = lambda x: ((x) << 29) # type: ignore
-PACKET3_MAP_QUEUES = 0xA2 # type: ignore
+PACKET3_MAP_QUEUES = 0xA2
 PACKET3_MAP_QUEUES_QUEUE_SEL = lambda x: ((x) << 4) # type: ignore
 PACKET3_MAP_QUEUES_VMID = lambda x: ((x) << 8) # type: ignore
 PACKET3_MAP_QUEUES_QUEUE = lambda x: ((x) << 13) # type: ignore
@@ -370,7 +370,7 @@ PACKET3_MAP_QUEUES_ENGINE_SEL = lambda x: ((x) << 26) # type: ignore
 PACKET3_MAP_QUEUES_NUM_QUEUES = lambda x: ((x) << 29) # type: ignore
 PACKET3_MAP_QUEUES_CHECK_DISABLE = lambda x: ((x) << 1) # type: ignore
 PACKET3_MAP_QUEUES_DOORBELL_OFFSET = lambda x: ((x) << 2) # type: ignore
-PACKET3_UNMAP_QUEUES = 0xA3 # type: ignore
+PACKET3_UNMAP_QUEUES = 0xA3
 PACKET3_UNMAP_QUEUES_ACTION = lambda x: ((x) << 0) # type: ignore
 PACKET3_UNMAP_QUEUES_QUEUE_SEL = lambda x: ((x) << 4) # type: ignore
 PACKET3_UNMAP_QUEUES_ENGINE_SEL = lambda x: ((x) << 26) # type: ignore
@@ -381,32 +381,32 @@ PACKET3_UNMAP_QUEUES_DOORBELL_OFFSET1 = lambda x: ((x) << 2) # type: ignore
 PACKET3_UNMAP_QUEUES_RB_WPTR = lambda x: ((x) << 0) # type: ignore
 PACKET3_UNMAP_QUEUES_DOORBELL_OFFSET2 = lambda x: ((x) << 2) # type: ignore
 PACKET3_UNMAP_QUEUES_DOORBELL_OFFSET3 = lambda x: ((x) << 2) # type: ignore
-PACKET3_QUERY_STATUS = 0xA4 # type: ignore
+PACKET3_QUERY_STATUS = 0xA4
 PACKET3_QUERY_STATUS_CONTEXT_ID = lambda x: ((x) << 0) # type: ignore
 PACKET3_QUERY_STATUS_INTERRUPT_SEL = lambda x: ((x) << 28) # type: ignore
 PACKET3_QUERY_STATUS_COMMAND = lambda x: ((x) << 30) # type: ignore
 PACKET3_QUERY_STATUS_PASID = lambda x: ((x) << 0) # type: ignore
 PACKET3_QUERY_STATUS_DOORBELL_OFFSET = lambda x: ((x) << 2) # type: ignore
 PACKET3_QUERY_STATUS_ENG_SEL = lambda x: ((x) << 25) # type: ignore
-PACKET3_RUN_CLEANER_SHADER_9_0 = 0xD7 # type: ignore
-PACKET3_RUN_CLEANER_SHADER = 0xD2 # type: ignore
-VCE_CMD_NO_OP = 0x00000000 # type: ignore
-VCE_CMD_END = 0x00000001 # type: ignore
-VCE_CMD_IB = 0x00000002 # type: ignore
-VCE_CMD_FENCE = 0x00000003 # type: ignore
-VCE_CMD_TRAP = 0x00000004 # type: ignore
-VCE_CMD_IB_AUTO = 0x00000005 # type: ignore
-VCE_CMD_SEMAPHORE = 0x00000006 # type: ignore
-VCE_CMD_IB_VM = 0x00000102 # type: ignore
-VCE_CMD_WAIT_GE = 0x00000106 # type: ignore
-VCE_CMD_UPDATE_PTB = 0x00000107 # type: ignore
-VCE_CMD_FLUSH_TLB = 0x00000108 # type: ignore
-VCE_CMD_REG_WRITE = 0x00000109 # type: ignore
-VCE_CMD_REG_WAIT = 0x0000010a # type: ignore
-HEVC_ENC_CMD_NO_OP = 0x00000000 # type: ignore
-HEVC_ENC_CMD_END = 0x00000001 # type: ignore
-HEVC_ENC_CMD_FENCE = 0x00000003 # type: ignore
-HEVC_ENC_CMD_TRAP = 0x00000004 # type: ignore
-HEVC_ENC_CMD_IB_VM = 0x00000102 # type: ignore
-HEVC_ENC_CMD_REG_WRITE = 0x00000109 # type: ignore
-HEVC_ENC_CMD_REG_WAIT = 0x0000010a # type: ignore
+PACKET3_RUN_CLEANER_SHADER_9_0 = 0xD7
+PACKET3_RUN_CLEANER_SHADER = 0xD2
+VCE_CMD_NO_OP = 0x00000000
+VCE_CMD_END = 0x00000001
+VCE_CMD_IB = 0x00000002
+VCE_CMD_FENCE = 0x00000003
+VCE_CMD_TRAP = 0x00000004
+VCE_CMD_IB_AUTO = 0x00000005
+VCE_CMD_SEMAPHORE = 0x00000006
+VCE_CMD_IB_VM = 0x00000102
+VCE_CMD_WAIT_GE = 0x00000106
+VCE_CMD_UPDATE_PTB = 0x00000107
+VCE_CMD_FLUSH_TLB = 0x00000108
+VCE_CMD_REG_WRITE = 0x00000109
+VCE_CMD_REG_WAIT = 0x0000010a
+HEVC_ENC_CMD_NO_OP = 0x00000000
+HEVC_ENC_CMD_END = 0x00000001
+HEVC_ENC_CMD_FENCE = 0x00000003
+HEVC_ENC_CMD_TRAP = 0x00000004
+HEVC_ENC_CMD_IB_VM = 0x00000102
+HEVC_ENC_CMD_REG_WRITE = 0x00000109
+HEVC_ENC_CMD_REG_WAIT = 0x0000010a

tinygrad/runtime/autogen/am/smu_v13_0_0.py

@@ -760,7 +760,7 @@ class struct_smu_state_memory_block(c.Struct):
   dll_off: bool
   m3arb: int
   unused: c.Array[ctypes.c_ubyte, Literal[3]]
-struct_smu_state_memory_block.register_fields([('dll_off', ctypes.c_bool, 0), ('m3arb', ctypes.c_ubyte, 1), ('unused', c.Array[ctypes.c_ubyte, Literal[3]], 2)])
+struct_smu_state_memory_block.register_fields([('dll_off', ctypes.c_bool, 0), ('m3arb', uint8_t, 1), ('unused', c.Array[uint8_t, Literal[3]], 2)])
 @c.record
 class struct_smu_state_software_algorithm_block(c.Struct):
   SIZE = 2
@@ -788,13 +788,13 @@ class struct_smu_state_validation_block(c.Struct):
   single_display_only: bool
   disallow_on_dc: bool
   supported_power_levels: int
-struct_smu_state_validation_block.register_fields([('single_display_only', ctypes.c_bool, 0), ('disallow_on_dc', ctypes.c_bool, 1), ('supported_power_levels', ctypes.c_ubyte, 2)])
+struct_smu_state_validation_block.register_fields([('single_display_only', ctypes.c_bool, 0), ('disallow_on_dc', ctypes.c_bool, 1), ('supported_power_levels', uint8_t, 2)])
 @c.record
 class struct_smu_uvd_clocks(c.Struct):
   SIZE = 8
   vclk: int
   dclk: int
-struct_smu_uvd_clocks.register_fields([('vclk', ctypes.c_uint32, 0), ('dclk', ctypes.c_uint32, 4)])
+struct_smu_uvd_clocks.register_fields([('vclk', uint32_t, 0), ('dclk', uint32_t, 4)])
 enum_smu_power_src_type: dict[int, str] = {(SMU_POWER_SOURCE_AC:=0): 'SMU_POWER_SOURCE_AC', (SMU_POWER_SOURCE_DC:=1): 'SMU_POWER_SOURCE_DC', (SMU_POWER_SOURCE_COUNT:=2): 'SMU_POWER_SOURCE_COUNT'}
 enum_smu_ppt_limit_type: dict[int, str] = {(SMU_DEFAULT_PPT_LIMIT:=0): 'SMU_DEFAULT_PPT_LIMIT', (SMU_FAST_PPT_LIMIT:=1): 'SMU_FAST_PPT_LIMIT'}
 enum_smu_ppt_limit_level: dict[int, str] = {(SMU_PPT_LIMIT_MIN:=-1): 'SMU_PPT_LIMIT_MIN', (SMU_PPT_LIMIT_CURRENT:=0): 'SMU_PPT_LIMIT_CURRENT', (SMU_PPT_LIMIT_DEFAULT:=1): 'SMU_PPT_LIMIT_DEFAULT', (SMU_PPT_LIMIT_MAX:=2): 'SMU_PPT_LIMIT_MAX'}
@@ -811,7 +811,7 @@ class struct_smu_user_dpm_profile(c.Struct):
   user_od: int
   clk_mask: c.Array[ctypes.c_uint32, Literal[28]]
   clk_dependency: int
-struct_smu_user_dpm_profile.register_fields([('fan_mode', ctypes.c_uint32, 0), ('power_limit', ctypes.c_uint32, 4), ('fan_speed_pwm', ctypes.c_uint32, 8), ('fan_speed_rpm', ctypes.c_uint32, 12), ('flags', ctypes.c_uint32, 16), ('user_od', ctypes.c_uint32, 20), ('clk_mask', c.Array[ctypes.c_uint32, Literal[28]], 24), ('clk_dependency', ctypes.c_uint32, 136)])
+struct_smu_user_dpm_profile.register_fields([('fan_mode', uint32_t, 0), ('power_limit', uint32_t, 4), ('fan_speed_pwm', uint32_t, 8), ('fan_speed_rpm', uint32_t, 12), ('flags', uint32_t, 16), ('user_od', uint32_t, 20), ('clk_mask', c.Array[uint32_t, Literal[28]], 24), ('clk_dependency', uint32_t, 136)])
 @c.record
 class struct_smu_table(c.Struct):
   SIZE = 48
@@ -823,7 +823,7 @@ class struct_smu_table(c.Struct):
   bo: c.POINTER[struct_amdgpu_bo]
   version: int
 class struct_amdgpu_bo(c.Struct): pass
-struct_smu_table.register_fields([('size', ctypes.c_uint64, 0), ('align', ctypes.c_uint32, 8), ('domain', ctypes.c_ubyte, 12), ('mc_address', ctypes.c_uint64, 16), ('cpu_addr', ctypes.c_void_p, 24), ('bo', c.POINTER[struct_amdgpu_bo], 32), ('version', ctypes.c_uint32, 40)])
+struct_smu_table.register_fields([('size', uint64_t, 0), ('align', uint32_t, 8), ('domain', uint8_t, 12), ('mc_address', uint64_t, 16), ('cpu_addr', ctypes.c_void_p, 24), ('bo', c.POINTER[struct_amdgpu_bo], 32), ('version', uint32_t, 40)])
 enum_smu_perf_level_designation: dict[int, str] = {(PERF_LEVEL_ACTIVITY:=0): 'PERF_LEVEL_ACTIVITY', (PERF_LEVEL_POWER_CONTAINMENT:=1): 'PERF_LEVEL_POWER_CONTAINMENT'}
 @c.record
 class struct_smu_performance_level(c.Struct):
@@ -834,7 +834,7 @@ class struct_smu_performance_level(c.Struct):
   vddci: int
   non_local_mem_freq: int
   non_local_mem_width: int
-struct_smu_performance_level.register_fields([('core_clock', ctypes.c_uint32, 0), ('memory_clock', ctypes.c_uint32, 4), ('vddc', ctypes.c_uint32, 8), ('vddci', ctypes.c_uint32, 12), ('non_local_mem_freq', ctypes.c_uint32, 16), ('non_local_mem_width', ctypes.c_uint32, 20)])
+struct_smu_performance_level.register_fields([('core_clock', uint32_t, 0), ('memory_clock', uint32_t, 4), ('vddc', uint32_t, 8), ('vddci', uint32_t, 12), ('non_local_mem_freq', uint32_t, 16), ('non_local_mem_width', uint32_t, 20)])
 @c.record
 class struct_smu_clock_info(c.Struct):
   SIZE = 24
@@ -844,7 +844,7 @@ class struct_smu_clock_info(c.Struct):
   max_eng_clk: int
   min_bus_bandwidth: int
   max_bus_bandwidth: int
-struct_smu_clock_info.register_fields([('min_mem_clk', ctypes.c_uint32, 0), ('max_mem_clk', ctypes.c_uint32, 4), ('min_eng_clk', ctypes.c_uint32, 8), ('max_eng_clk', ctypes.c_uint32, 12), ('min_bus_bandwidth', ctypes.c_uint32, 16), ('max_bus_bandwidth', ctypes.c_uint32, 20)])
+struct_smu_clock_info.register_fields([('min_mem_clk', uint32_t, 0), ('max_mem_clk', uint32_t, 4), ('min_eng_clk', uint32_t, 8), ('max_eng_clk', uint32_t, 12), ('min_bus_bandwidth', uint32_t, 16), ('max_bus_bandwidth', uint32_t, 20)])
 @c.record
 class struct_smu_bios_boot_up_values(c.Struct):
   SIZE = 68
@@ -867,360 +867,359 @@ class struct_smu_bios_boot_up_values(c.Struct):
   fclk: int
   lclk: int
   firmware_caps: int
-struct_smu_bios_boot_up_values.register_fields([('revision', ctypes.c_uint32, 0), ('gfxclk', ctypes.c_uint32, 4), ('uclk', ctypes.c_uint32, 8), ('socclk', ctypes.c_uint32, 12), ('dcefclk', ctypes.c_uint32, 16), ('eclk', ctypes.c_uint32, 20), ('vclk', ctypes.c_uint32, 24), ('dclk', ctypes.c_uint32, 28), ('vddc', ctypes.c_uint16, 32), ('vddci', ctypes.c_uint16, 34), ('mvddc', ctypes.c_uint16, 36), ('vdd_gfx', ctypes.c_uint16, 38), ('cooling_id', ctypes.c_ubyte, 40), ('pp_table_id', ctypes.c_uint32, 44), ('format_revision', ctypes.c_uint32, 48), ('content_revision', ctypes.c_uint32, 52), ('fclk', ctypes.c_uint32, 56), ('lclk', ctypes.c_uint32, 60), ('firmware_caps', ctypes.c_uint32, 64)])
+struct_smu_bios_boot_up_values.register_fields([('revision', uint32_t, 0), ('gfxclk', uint32_t, 4), ('uclk', uint32_t, 8), ('socclk', uint32_t, 12), ('dcefclk', uint32_t, 16), ('eclk', uint32_t, 20), ('vclk', uint32_t, 24), ('dclk', uint32_t, 28), ('vddc', uint16_t, 32), ('vddci', uint16_t, 34), ('mvddc', uint16_t, 36), ('vdd_gfx', uint16_t, 38), ('cooling_id', uint8_t, 40), ('pp_table_id', uint32_t, 44), ('format_revision', uint32_t, 48), ('content_revision', uint32_t, 52), ('fclk', uint32_t, 56), ('lclk', uint32_t, 60), ('firmware_caps', uint32_t, 64)])
 enum_smu_table_id: dict[int, str] = {(SMU_TABLE_PPTABLE:=0): 'SMU_TABLE_PPTABLE', (SMU_TABLE_WATERMARKS:=1): 'SMU_TABLE_WATERMARKS', (SMU_TABLE_CUSTOM_DPM:=2): 'SMU_TABLE_CUSTOM_DPM', (SMU_TABLE_DPMCLOCKS:=3): 'SMU_TABLE_DPMCLOCKS', (SMU_TABLE_AVFS:=4): 'SMU_TABLE_AVFS', (SMU_TABLE_AVFS_PSM_DEBUG:=5): 'SMU_TABLE_AVFS_PSM_DEBUG', (SMU_TABLE_AVFS_FUSE_OVERRIDE:=6): 'SMU_TABLE_AVFS_FUSE_OVERRIDE', (SMU_TABLE_PMSTATUSLOG:=7): 'SMU_TABLE_PMSTATUSLOG', (SMU_TABLE_SMU_METRICS:=8): 'SMU_TABLE_SMU_METRICS', (SMU_TABLE_DRIVER_SMU_CONFIG:=9): 'SMU_TABLE_DRIVER_SMU_CONFIG', (SMU_TABLE_ACTIVITY_MONITOR_COEFF:=10): 'SMU_TABLE_ACTIVITY_MONITOR_COEFF', (SMU_TABLE_OVERDRIVE:=11): 'SMU_TABLE_OVERDRIVE', (SMU_TABLE_I2C_COMMANDS:=12): 'SMU_TABLE_I2C_COMMANDS', (SMU_TABLE_PACE:=13): 'SMU_TABLE_PACE', (SMU_TABLE_ECCINFO:=14): 'SMU_TABLE_ECCINFO', (SMU_TABLE_COMBO_PPTABLE:=15): 'SMU_TABLE_COMBO_PPTABLE', (SMU_TABLE_WIFIBAND:=16): 'SMU_TABLE_WIFIBAND', (SMU_TABLE_COUNT:=17): 'SMU_TABLE_COUNT'}
-PPSMC_VERSION = 0x1 # type: ignore
-DEBUGSMC_VERSION = 0x1 # type: ignore
-PPSMC_Result_OK = 0x1 # type: ignore
-PPSMC_Result_Failed = 0xFF # type: ignore
-PPSMC_Result_UnknownCmd = 0xFE # type: ignore
-PPSMC_Result_CmdRejectedPrereq = 0xFD # type: ignore
-PPSMC_Result_CmdRejectedBusy = 0xFC # type: ignore
-PPSMC_MSG_TestMessage = 0x1 # type: ignore
-PPSMC_MSG_GetSmuVersion = 0x2 # type: ignore
-PPSMC_MSG_GetDriverIfVersion = 0x3 # type: ignore
-PPSMC_MSG_SetAllowedFeaturesMaskLow = 0x4 # type: ignore
-PPSMC_MSG_SetAllowedFeaturesMaskHigh = 0x5 # type: ignore
-PPSMC_MSG_EnableAllSmuFeatures = 0x6 # type: ignore
-PPSMC_MSG_DisableAllSmuFeatures = 0x7 # type: ignore
-PPSMC_MSG_EnableSmuFeaturesLow = 0x8 # type: ignore
-PPSMC_MSG_EnableSmuFeaturesHigh = 0x9 # type: ignore
-PPSMC_MSG_DisableSmuFeaturesLow = 0xA # type: ignore
-PPSMC_MSG_DisableSmuFeaturesHigh = 0xB # type: ignore
-PPSMC_MSG_GetRunningSmuFeaturesLow = 0xC # type: ignore
-PPSMC_MSG_GetRunningSmuFeaturesHigh = 0xD # type: ignore
-PPSMC_MSG_SetDriverDramAddrHigh = 0xE # type: ignore
-PPSMC_MSG_SetDriverDramAddrLow = 0xF # type: ignore
-PPSMC_MSG_SetToolsDramAddrHigh = 0x10 # type: ignore
-PPSMC_MSG_SetToolsDramAddrLow = 0x11 # type: ignore
-PPSMC_MSG_TransferTableSmu2Dram = 0x12 # type: ignore
-PPSMC_MSG_TransferTableDram2Smu = 0x13 # type: ignore
-PPSMC_MSG_UseDefaultPPTable = 0x14 # type: ignore
-PPSMC_MSG_EnterBaco = 0x15 # type: ignore
-PPSMC_MSG_ExitBaco = 0x16 # type: ignore
-PPSMC_MSG_ArmD3 = 0x17 # type: ignore
-PPSMC_MSG_BacoAudioD3PME = 0x18 # type: ignore
-PPSMC_MSG_SetSoftMinByFreq = 0x19 # type: ignore
-PPSMC_MSG_SetSoftMaxByFreq = 0x1A # type: ignore
-PPSMC_MSG_SetHardMinByFreq = 0x1B # type: ignore
-PPSMC_MSG_SetHardMaxByFreq = 0x1C # type: ignore
-PPSMC_MSG_GetMinDpmFreq = 0x1D # type: ignore
-PPSMC_MSG_GetMaxDpmFreq = 0x1E # type: ignore
-PPSMC_MSG_GetDpmFreqByIndex = 0x1F # type: ignore
-PPSMC_MSG_OverridePcieParameters = 0x20 # type: ignore
-PPSMC_MSG_DramLogSetDramAddrHigh = 0x21 # type: ignore
-PPSMC_MSG_DramLogSetDramAddrLow = 0x22 # type: ignore
-PPSMC_MSG_DramLogSetDramSize = 0x23 # type: ignore
-PPSMC_MSG_SetWorkloadMask = 0x24 # type: ignore
-PPSMC_MSG_GetVoltageByDpm = 0x25 # type: ignore
-PPSMC_MSG_SetVideoFps = 0x26 # type: ignore
-PPSMC_MSG_GetDcModeMaxDpmFreq = 0x27 # type: ignore
-PPSMC_MSG_AllowGfxOff = 0x28 # type: ignore
-PPSMC_MSG_DisallowGfxOff = 0x29 # type: ignore
-PPSMC_MSG_PowerUpVcn = 0x2A # type: ignore
-PPSMC_MSG_PowerDownVcn = 0x2B # type: ignore
-PPSMC_MSG_PowerUpJpeg = 0x2C # type: ignore
-PPSMC_MSG_PowerDownJpeg = 0x2D # type: ignore
-PPSMC_MSG_PrepareMp1ForUnload = 0x2E # type: ignore
-PPSMC_MSG_Mode1Reset = 0x2F # type: ignore
-PPSMC_MSG_Mode2Reset = 0x4F # type: ignore
-PPSMC_MSG_SetSystemVirtualDramAddrHigh = 0x30 # type: ignore
-PPSMC_MSG_SetSystemVirtualDramAddrLow = 0x31 # type: ignore
-PPSMC_MSG_SetPptLimit = 0x32 # type: ignore
-PPSMC_MSG_GetPptLimit = 0x33 # type: ignore
-PPSMC_MSG_ReenableAcDcInterrupt = 0x34 # type: ignore
-PPSMC_MSG_NotifyPowerSource = 0x35 # type: ignore
-PPSMC_MSG_RunDcBtc = 0x36 # type: ignore
-PPSMC_MSG_GetDebugData = 0x37 # type: ignore
-PPSMC_MSG_SetTemperatureInputSelect = 0x38 # type: ignore
-PPSMC_MSG_SetFwDstatesMask = 0x39 # type: ignore
-PPSMC_MSG_SetThrottlerMask = 0x3A # type: ignore
-PPSMC_MSG_SetExternalClientDfCstateAllow = 0x3B # type: ignore
-PPSMC_MSG_SetMGpuFanBoostLimitRpm = 0x3C # type: ignore
-PPSMC_MSG_DumpSTBtoDram = 0x3D # type: ignore
-PPSMC_MSG_STBtoDramLogSetDramAddrHigh = 0x3E # type: ignore
-PPSMC_MSG_STBtoDramLogSetDramAddrLow = 0x3F # type: ignore
-PPSMC_MSG_STBtoDramLogSetDramSize = 0x40 # type: ignore
-PPSMC_MSG_SetGpoAllow = 0x41 # type: ignore
-PPSMC_MSG_AllowGfxDcs = 0x42 # type: ignore
-PPSMC_MSG_DisallowGfxDcs = 0x43 # type: ignore
-PPSMC_MSG_EnableAudioStutterWA = 0x44 # type: ignore
-PPSMC_MSG_PowerUpUmsch = 0x45 # type: ignore
-PPSMC_MSG_PowerDownUmsch = 0x46 # type: ignore
-PPSMC_MSG_SetDcsArch = 0x47 # type: ignore
-PPSMC_MSG_TriggerVFFLR = 0x48 # type: ignore
-PPSMC_MSG_SetNumBadMemoryPagesRetired = 0x49 # type: ignore
-PPSMC_MSG_SetBadMemoryPagesRetiredFlagsPerChannel = 0x4A # type: ignore
-PPSMC_MSG_SetPriorityDeltaGain = 0x4B # type: ignore
-PPSMC_MSG_AllowIHHostInterrupt = 0x4C # type: ignore
-PPSMC_MSG_DALNotPresent = 0x4E # type: ignore
-PPSMC_MSG_EnableUCLKShadow = 0x51 # type: ignore
-PPSMC_Message_Count = 0x52 # type: ignore
-DEBUGSMC_MSG_TestMessage = 0x1 # type: ignore
-DEBUGSMC_MSG_GetDebugData = 0x2 # type: ignore
-DEBUGSMC_MSG_DebugDumpExit = 0x3 # type: ignore
-DEBUGSMC_Message_Count = 0x4 # type: ignore
-SMU13_0_0_DRIVER_IF_VERSION = 0x3D # type: ignore
-PPTABLE_VERSION = 0x2B # type: ignore
-NUM_GFXCLK_DPM_LEVELS = 16 # type: ignore
-NUM_SOCCLK_DPM_LEVELS = 8 # type: ignore
-NUM_MP0CLK_DPM_LEVELS = 2 # type: ignore
-NUM_DCLK_DPM_LEVELS = 8 # type: ignore
-NUM_VCLK_DPM_LEVELS = 8 # type: ignore
-NUM_DISPCLK_DPM_LEVELS = 8 # type: ignore
-NUM_DPPCLK_DPM_LEVELS = 8 # type: ignore
-NUM_DPREFCLK_DPM_LEVELS = 8 # type: ignore
-NUM_DCFCLK_DPM_LEVELS = 8 # type: ignore
-NUM_DTBCLK_DPM_LEVELS = 8 # type: ignore
-NUM_UCLK_DPM_LEVELS = 4 # type: ignore
-NUM_LINK_LEVELS = 3 # type: ignore
-NUM_FCLK_DPM_LEVELS = 8 # type: ignore
-NUM_OD_FAN_MAX_POINTS = 6 # type: ignore
-FEATURE_FW_DATA_READ_BIT = 0 # type: ignore
-FEATURE_DPM_GFXCLK_BIT = 1 # type: ignore
-FEATURE_DPM_GFX_POWER_OPTIMIZER_BIT = 2 # type: ignore
-FEATURE_DPM_UCLK_BIT = 3 # type: ignore
-FEATURE_DPM_FCLK_BIT = 4 # type: ignore
-FEATURE_DPM_SOCCLK_BIT = 5 # type: ignore
-FEATURE_DPM_MP0CLK_BIT = 6 # type: ignore
-FEATURE_DPM_LINK_BIT = 7 # type: ignore
-FEATURE_DPM_DCN_BIT = 8 # type: ignore
-FEATURE_VMEMP_SCALING_BIT = 9 # type: ignore
-FEATURE_VDDIO_MEM_SCALING_BIT = 10 # type: ignore
-FEATURE_DS_GFXCLK_BIT = 11 # type: ignore
-FEATURE_DS_SOCCLK_BIT = 12 # type: ignore
-FEATURE_DS_FCLK_BIT = 13 # type: ignore
-FEATURE_DS_LCLK_BIT = 14 # type: ignore
-FEATURE_DS_DCFCLK_BIT = 15 # type: ignore
-FEATURE_DS_UCLK_BIT = 16 # type: ignore
-FEATURE_GFX_ULV_BIT = 17 # type: ignore
-FEATURE_FW_DSTATE_BIT = 18 # type: ignore
-FEATURE_GFXOFF_BIT = 19 # type: ignore
-FEATURE_BACO_BIT = 20 # type: ignore
-FEATURE_MM_DPM_BIT = 21 # type: ignore
-FEATURE_SOC_MPCLK_DS_BIT = 22 # type: ignore
-FEATURE_BACO_MPCLK_DS_BIT = 23 # type: ignore
-FEATURE_THROTTLERS_BIT = 24 # type: ignore
-FEATURE_SMARTSHIFT_BIT = 25 # type: ignore
-FEATURE_GTHR_BIT = 26 # type: ignore
-FEATURE_ACDC_BIT = 27 # type: ignore
-FEATURE_VR0HOT_BIT = 28 # type: ignore
-FEATURE_FW_CTF_BIT = 29 # type: ignore
-FEATURE_FAN_CONTROL_BIT = 30 # type: ignore
-FEATURE_GFX_DCS_BIT = 31 # type: ignore
-FEATURE_GFX_READ_MARGIN_BIT = 32 # type: ignore
-FEATURE_LED_DISPLAY_BIT = 33 # type: ignore
-FEATURE_GFXCLK_SPREAD_SPECTRUM_BIT = 34 # type: ignore
-FEATURE_OUT_OF_BAND_MONITOR_BIT = 35 # type: ignore
-FEATURE_OPTIMIZED_VMIN_BIT = 36 # type: ignore
-FEATURE_GFX_IMU_BIT = 37 # type: ignore
-FEATURE_BOOT_TIME_CAL_BIT = 38 # type: ignore
-FEATURE_GFX_PCC_DFLL_BIT = 39 # type: ignore
-FEATURE_SOC_CG_BIT = 40 # type: ignore
-FEATURE_DF_CSTATE_BIT = 41 # type: ignore
-FEATURE_GFX_EDC_BIT = 42 # type: ignore
-FEATURE_BOOT_POWER_OPT_BIT = 43 # type: ignore
-FEATURE_CLOCK_POWER_DOWN_BYPASS_BIT = 44 # type: ignore
-FEATURE_DS_VCN_BIT = 45 # type: ignore
-FEATURE_BACO_CG_BIT = 46 # type: ignore
-FEATURE_MEM_TEMP_READ_BIT = 47 # type: ignore
-FEATURE_ATHUB_MMHUB_PG_BIT = 48 # type: ignore
-FEATURE_SOC_PCC_BIT = 49 # type: ignore
-FEATURE_EDC_PWRBRK_BIT = 50 # type: ignore
-FEATURE_BOMXCO_SVI3_PROG_BIT = 51 # type: ignore
-FEATURE_SPARE_52_BIT = 52 # type: ignore
-FEATURE_SPARE_53_BIT = 53 # type: ignore
-FEATURE_SPARE_54_BIT = 54 # type: ignore
-FEATURE_SPARE_55_BIT = 55 # type: ignore
-FEATURE_SPARE_56_BIT = 56 # type: ignore
-FEATURE_SPARE_57_BIT = 57 # type: ignore
-FEATURE_SPARE_58_BIT = 58 # type: ignore
-FEATURE_SPARE_59_BIT = 59 # type: ignore
-FEATURE_SPARE_60_BIT = 60 # type: ignore
-FEATURE_SPARE_61_BIT = 61 # type: ignore
-FEATURE_SPARE_62_BIT = 62 # type: ignore
-FEATURE_SPARE_63_BIT = 63 # type: ignore
-NUM_FEATURES = 64 # type: ignore
-ALLOWED_FEATURE_CTRL_DEFAULT = 0xFFFFFFFFFFFFFFFF # type: ignore
-ALLOWED_FEATURE_CTRL_SCPM = ((1 << FEATURE_DPM_GFXCLK_BIT) | (1 << FEATURE_DPM_GFX_POWER_OPTIMIZER_BIT) | (1 << FEATURE_DPM_UCLK_BIT) | (1 << FEATURE_DPM_FCLK_BIT) | (1 << FEATURE_DPM_SOCCLK_BIT) | (1 << FEATURE_DPM_MP0CLK_BIT) | (1 << FEATURE_DPM_LINK_BIT) | (1 << FEATURE_DPM_DCN_BIT) | (1 << FEATURE_DS_GFXCLK_BIT) | (1 << FEATURE_DS_SOCCLK_BIT) | (1 << FEATURE_DS_FCLK_BIT) | (1 << FEATURE_DS_LCLK_BIT) | (1 << FEATURE_DS_DCFCLK_BIT) | (1 << FEATURE_DS_UCLK_BIT) | (1 << FEATURE_DS_VCN_BIT)) # type: ignore
-DEBUG_OVERRIDE_DISABLE_VOLT_LINK_VCN_FCLK = 0x00000001 # type: ignore
-DEBUG_OVERRIDE_DISABLE_VOLT_LINK_DCN_FCLK = 0x00000002 # type: ignore
-DEBUG_OVERRIDE_DISABLE_VOLT_LINK_MP0_FCLK = 0x00000004 # type: ignore
-DEBUG_OVERRIDE_DISABLE_VOLT_LINK_VCN_DCFCLK = 0x00000008 # type: ignore
-DEBUG_OVERRIDE_DISABLE_FAST_FCLK_TIMER = 0x00000010 # type: ignore
-DEBUG_OVERRIDE_DISABLE_VCN_PG = 0x00000020 # type: ignore
-DEBUG_OVERRIDE_DISABLE_FMAX_VMAX = 0x00000040 # type: ignore
-DEBUG_OVERRIDE_DISABLE_IMU_FW_CHECKS = 0x00000080 # type: ignore
-DEBUG_OVERRIDE_DISABLE_D0i2_REENTRY_HSR_TIMER_CHECK = 0x00000100 # type: ignore
-DEBUG_OVERRIDE_DISABLE_DFLL = 0x00000200 # type: ignore
-DEBUG_OVERRIDE_ENABLE_RLC_VF_BRINGUP_MODE = 0x00000400 # type: ignore
-DEBUG_OVERRIDE_DFLL_MASTER_MODE = 0x00000800 # type: ignore
-DEBUG_OVERRIDE_ENABLE_PROFILING_MODE = 0x00001000 # type: ignore
-VR_MAPPING_VR_SELECT_MASK = 0x01 # type: ignore
-VR_MAPPING_VR_SELECT_SHIFT = 0x00 # type: ignore
-VR_MAPPING_PLANE_SELECT_MASK = 0x02 # type: ignore
-VR_MAPPING_PLANE_SELECT_SHIFT = 0x01 # type: ignore
-PSI_SEL_VR0_PLANE0_PSI0 = 0x01 # type: ignore
-PSI_SEL_VR0_PLANE0_PSI1 = 0x02 # type: ignore
-PSI_SEL_VR0_PLANE1_PSI0 = 0x04 # type: ignore
-PSI_SEL_VR0_PLANE1_PSI1 = 0x08 # type: ignore
-PSI_SEL_VR1_PLANE0_PSI0 = 0x10 # type: ignore
-PSI_SEL_VR1_PLANE0_PSI1 = 0x20 # type: ignore
-PSI_SEL_VR1_PLANE1_PSI0 = 0x40 # type: ignore
-PSI_SEL_VR1_PLANE1_PSI1 = 0x80 # type: ignore
-THROTTLER_TEMP_EDGE_BIT = 0 # type: ignore
-THROTTLER_TEMP_HOTSPOT_BIT = 1 # type: ignore
-THROTTLER_TEMP_HOTSPOT_G_BIT = 2 # type: ignore
-THROTTLER_TEMP_HOTSPOT_M_BIT = 3 # type: ignore
-THROTTLER_TEMP_MEM_BIT = 4 # type: ignore
-THROTTLER_TEMP_VR_GFX_BIT = 5 # type: ignore
-THROTTLER_TEMP_VR_MEM0_BIT = 6 # type: ignore
-THROTTLER_TEMP_VR_MEM1_BIT = 7 # type: ignore
-THROTTLER_TEMP_VR_SOC_BIT = 8 # type: ignore
-THROTTLER_TEMP_VR_U_BIT = 9 # type: ignore
-THROTTLER_TEMP_LIQUID0_BIT = 10 # type: ignore
-THROTTLER_TEMP_LIQUID1_BIT = 11 # type: ignore
-THROTTLER_TEMP_PLX_BIT = 12 # type: ignore
-THROTTLER_TDC_GFX_BIT = 13 # type: ignore
-THROTTLER_TDC_SOC_BIT = 14 # type: ignore
-THROTTLER_TDC_U_BIT = 15 # type: ignore
-THROTTLER_PPT0_BIT = 16 # type: ignore
-THROTTLER_PPT1_BIT = 17 # type: ignore
-THROTTLER_PPT2_BIT = 18 # type: ignore
-THROTTLER_PPT3_BIT = 19 # type: ignore
-THROTTLER_FIT_BIT = 20 # type: ignore
-THROTTLER_GFX_APCC_PLUS_BIT = 21 # type: ignore
-THROTTLER_COUNT = 22 # type: ignore
-FW_DSTATE_SOC_ULV_BIT = 0 # type: ignore
-FW_DSTATE_G6_HSR_BIT = 1 # type: ignore
-FW_DSTATE_G6_PHY_VMEMP_OFF_BIT = 2 # type: ignore
-FW_DSTATE_SMN_DS_BIT = 3 # type: ignore
-FW_DSTATE_MP1_WHISPER_MODE_BIT = 4 # type: ignore
-FW_DSTATE_SOC_LIV_MIN_BIT = 5 # type: ignore
-FW_DSTATE_SOC_PLL_PWRDN_BIT = 6 # type: ignore
-FW_DSTATE_MEM_PLL_PWRDN_BIT = 7 # type: ignore
-FW_DSTATE_MALL_ALLOC_BIT = 8 # type: ignore
-FW_DSTATE_MEM_PSI_BIT = 9 # type: ignore
-FW_DSTATE_HSR_NON_STROBE_BIT = 10 # type: ignore
-FW_DSTATE_MP0_ENTER_WFI_BIT = 11 # type: ignore
-FW_DSTATE_U_ULV_BIT = 12 # type: ignore
-FW_DSTATE_MALL_FLUSH_BIT = 13 # type: ignore
-FW_DSTATE_SOC_PSI_BIT = 14 # type: ignore
-FW_DSTATE_U_PSI_BIT = 15 # type: ignore
-FW_DSTATE_UCP_DS_BIT = 16 # type: ignore
-FW_DSTATE_CSRCLK_DS_BIT = 17 # type: ignore
-FW_DSTATE_MMHUB_INTERLOCK_BIT = 18 # type: ignore
-FW_DSTATE_D0i3_2_QUIET_FW_BIT = 19 # type: ignore
-FW_DSTATE_CLDO_PRG_BIT = 20 # type: ignore
-FW_DSTATE_DF_PLL_PWRDN_BIT = 21 # type: ignore
-FW_DSTATE_U_LOW_PWR_MODE_EN_BIT = 22 # type: ignore
-FW_DSTATE_GFX_PSI6_BIT = 23 # type: ignore
-FW_DSTATE_GFX_VR_PWR_STAGE_BIT = 24 # type: ignore
-LED_DISPLAY_GFX_DPM_BIT = 0 # type: ignore
-LED_DISPLAY_PCIE_BIT = 1 # type: ignore
-LED_DISPLAY_ERROR_BIT = 2 # type: ignore
-MEM_TEMP_READ_OUT_OF_BAND_BIT = 0 # type: ignore
-MEM_TEMP_READ_IN_BAND_REFRESH_BIT = 1 # type: ignore
-MEM_TEMP_READ_IN_BAND_DUMMY_PSTATE_BIT = 2 # type: ignore
-NUM_I2C_CONTROLLERS = 8 # type: ignore
-I2C_CONTROLLER_ENABLED = 1 # type: ignore
-I2C_CONTROLLER_DISABLED = 0 # type: ignore
-MAX_SW_I2C_COMMANDS = 24 # type: ignore
-CMDCONFIG_STOP_BIT = 0 # type: ignore
-CMDCONFIG_RESTART_BIT = 1 # type: ignore
-CMDCONFIG_READWRITE_BIT = 2 # type: ignore
-CMDCONFIG_STOP_MASK = (1 << CMDCONFIG_STOP_BIT) # type: ignore
-CMDCONFIG_RESTART_MASK = (1 << CMDCONFIG_RESTART_BIT) # type: ignore
-CMDCONFIG_READWRITE_MASK = (1 << CMDCONFIG_READWRITE_BIT) # type: ignore
-PP_NUM_RTAVFS_PWL_ZONES = 5 # type: ignore
-PP_OD_FEATURE_GFX_VF_CURVE_BIT = 0 # type: ignore
-PP_OD_FEATURE_PPT_BIT = 2 # type: ignore
-PP_OD_FEATURE_FAN_CURVE_BIT = 3 # type: ignore
-PP_OD_FEATURE_GFXCLK_BIT = 7 # type: ignore
-PP_OD_FEATURE_UCLK_BIT = 8 # type: ignore
-PP_OD_FEATURE_ZERO_FAN_BIT = 9 # type: ignore
-PP_OD_FEATURE_TEMPERATURE_BIT = 10 # type: ignore
-PP_OD_FEATURE_COUNT = 13 # type: ignore
-PP_NUM_OD_VF_CURVE_POINTS = PP_NUM_RTAVFS_PWL_ZONES + 1 # type: ignore
-INVALID_BOARD_GPIO = 0xFF # type: ignore
-MARKETING_BASE_CLOCKS = 0 # type: ignore
-MARKETING_GAME_CLOCKS = 1 # type: ignore
-MARKETING_BOOST_CLOCKS = 2 # type: ignore
-NUM_WM_RANGES = 4 # type: ignore
-WORKLOAD_PPLIB_DEFAULT_BIT = 0 # type: ignore
-WORKLOAD_PPLIB_FULL_SCREEN_3D_BIT = 1 # type: ignore
-WORKLOAD_PPLIB_POWER_SAVING_BIT = 2 # type: ignore
-WORKLOAD_PPLIB_VIDEO_BIT = 3 # type: ignore
-WORKLOAD_PPLIB_VR_BIT = 4 # type: ignore
-WORKLOAD_PPLIB_COMPUTE_BIT = 5 # type: ignore
-WORKLOAD_PPLIB_CUSTOM_BIT = 6 # type: ignore
-WORKLOAD_PPLIB_WINDOW_3D_BIT = 7 # type: ignore
-WORKLOAD_PPLIB_COUNT = 8 # type: ignore
-TABLE_TRANSFER_OK = 0x0 # type: ignore
-TABLE_TRANSFER_FAILED = 0xFF # type: ignore
-TABLE_TRANSFER_PENDING = 0xAB # type: ignore
-TABLE_PPTABLE = 0 # type: ignore
-TABLE_COMBO_PPTABLE = 1 # type: ignore
-TABLE_WATERMARKS = 2 # type: ignore
-TABLE_AVFS_PSM_DEBUG = 3 # type: ignore
-TABLE_PMSTATUSLOG = 4 # type: ignore
-TABLE_SMU_METRICS = 5 # type: ignore
-TABLE_DRIVER_SMU_CONFIG = 6 # type: ignore
-TABLE_ACTIVITY_MONITOR_COEFF = 7 # type: ignore
-TABLE_OVERDRIVE = 8 # type: ignore
-TABLE_I2C_COMMANDS = 9 # type: ignore
-TABLE_DRIVER_INFO = 10 # type: ignore
-TABLE_ECCINFO = 11 # type: ignore
-TABLE_WIFIBAND = 12 # type: ignore
-TABLE_COUNT = 13 # type: ignore
-IH_INTERRUPT_ID_TO_DRIVER = 0xFE # type: ignore
-IH_INTERRUPT_CONTEXT_ID_BACO = 0x2 # type: ignore
-IH_INTERRUPT_CONTEXT_ID_AC = 0x3 # type: ignore
-IH_INTERRUPT_CONTEXT_ID_DC = 0x4 # type: ignore
-IH_INTERRUPT_CONTEXT_ID_AUDIO_D0 = 0x5 # type: ignore
-IH_INTERRUPT_CONTEXT_ID_AUDIO_D3 = 0x6 # type: ignore
-IH_INTERRUPT_CONTEXT_ID_THERMAL_THROTTLING = 0x7 # type: ignore
-IH_INTERRUPT_CONTEXT_ID_FAN_ABNORMAL = 0x8 # type: ignore
-IH_INTERRUPT_CONTEXT_ID_FAN_RECOVERY = 0x9 # type: ignore
-int32_t = int # type: ignore
-SMU_THERMAL_MINIMUM_ALERT_TEMP = 0 # type: ignore
-SMU_THERMAL_MAXIMUM_ALERT_TEMP = 255 # type: ignore
-SMU_TEMPERATURE_UNITS_PER_CENTIGRADES = 1000 # type: ignore
-SMU_FW_NAME_LEN = 0x24 # type: ignore
-SMU_DPM_USER_PROFILE_RESTORE = (1 << 0) # type: ignore
-SMU_CUSTOM_FAN_SPEED_RPM = (1 << 1) # type: ignore
-SMU_CUSTOM_FAN_SPEED_PWM = (1 << 2) # type: ignore
-SMU_THROTTLER_PPT0_BIT = 0 # type: ignore
-SMU_THROTTLER_PPT1_BIT = 1 # type: ignore
-SMU_THROTTLER_PPT2_BIT = 2 # type: ignore
-SMU_THROTTLER_PPT3_BIT = 3 # type: ignore
-SMU_THROTTLER_SPL_BIT = 4 # type: ignore
-SMU_THROTTLER_FPPT_BIT = 5 # type: ignore
-SMU_THROTTLER_SPPT_BIT = 6 # type: ignore
-SMU_THROTTLER_SPPT_APU_BIT = 7 # type: ignore
-SMU_THROTTLER_TDC_GFX_BIT = 16 # type: ignore
-SMU_THROTTLER_TDC_SOC_BIT = 17 # type: ignore
-SMU_THROTTLER_TDC_MEM_BIT = 18 # type: ignore
-SMU_THROTTLER_TDC_VDD_BIT = 19 # type: ignore
-SMU_THROTTLER_TDC_CVIP_BIT = 20 # type: ignore
-SMU_THROTTLER_EDC_CPU_BIT = 21 # type: ignore
-SMU_THROTTLER_EDC_GFX_BIT = 22 # type: ignore
-SMU_THROTTLER_APCC_BIT = 23 # type: ignore
-SMU_THROTTLER_TEMP_GPU_BIT = 32 # type: ignore
-SMU_THROTTLER_TEMP_CORE_BIT = 33 # type: ignore
-SMU_THROTTLER_TEMP_MEM_BIT = 34 # type: ignore
-SMU_THROTTLER_TEMP_EDGE_BIT = 35 # type: ignore
-SMU_THROTTLER_TEMP_HOTSPOT_BIT = 36 # type: ignore
-SMU_THROTTLER_TEMP_SOC_BIT = 37 # type: ignore
-SMU_THROTTLER_TEMP_VR_GFX_BIT = 38 # type: ignore
-SMU_THROTTLER_TEMP_VR_SOC_BIT = 39 # type: ignore
-SMU_THROTTLER_TEMP_VR_MEM0_BIT = 40 # type: ignore
-SMU_THROTTLER_TEMP_VR_MEM1_BIT = 41 # type: ignore
-SMU_THROTTLER_TEMP_LIQUID0_BIT = 42 # type: ignore
-SMU_THROTTLER_TEMP_LIQUID1_BIT = 43 # type: ignore
-SMU_THROTTLER_VRHOT0_BIT = 44 # type: ignore
-SMU_THROTTLER_VRHOT1_BIT = 45 # type: ignore
-SMU_THROTTLER_PROCHOT_CPU_BIT = 46 # type: ignore
-SMU_THROTTLER_PROCHOT_GFX_BIT = 47 # type: ignore
-SMU_THROTTLER_PPM_BIT = 56 # type: ignore
-SMU_THROTTLER_FIT_BIT = 57 # type: ignore
+PPSMC_VERSION = 0x1
+DEBUGSMC_VERSION = 0x1
+PPSMC_Result_OK = 0x1
+PPSMC_Result_Failed = 0xFF
+PPSMC_Result_UnknownCmd = 0xFE
+PPSMC_Result_CmdRejectedPrereq = 0xFD
+PPSMC_Result_CmdRejectedBusy = 0xFC
+PPSMC_MSG_TestMessage = 0x1
+PPSMC_MSG_GetSmuVersion = 0x2
+PPSMC_MSG_GetDriverIfVersion = 0x3
+PPSMC_MSG_SetAllowedFeaturesMaskLow = 0x4
+PPSMC_MSG_SetAllowedFeaturesMaskHigh = 0x5
+PPSMC_MSG_EnableAllSmuFeatures = 0x6
+PPSMC_MSG_DisableAllSmuFeatures = 0x7
+PPSMC_MSG_EnableSmuFeaturesLow = 0x8
+PPSMC_MSG_EnableSmuFeaturesHigh = 0x9
+PPSMC_MSG_DisableSmuFeaturesLow = 0xA
+PPSMC_MSG_DisableSmuFeaturesHigh = 0xB
+PPSMC_MSG_GetRunningSmuFeaturesLow = 0xC
+PPSMC_MSG_GetRunningSmuFeaturesHigh = 0xD
+PPSMC_MSG_SetDriverDramAddrHigh = 0xE
+PPSMC_MSG_SetDriverDramAddrLow = 0xF
+PPSMC_MSG_SetToolsDramAddrHigh = 0x10
+PPSMC_MSG_SetToolsDramAddrLow = 0x11
+PPSMC_MSG_TransferTableSmu2Dram = 0x12
+PPSMC_MSG_TransferTableDram2Smu = 0x13
+PPSMC_MSG_UseDefaultPPTable = 0x14
+PPSMC_MSG_EnterBaco = 0x15
+PPSMC_MSG_ExitBaco = 0x16
+PPSMC_MSG_ArmD3 = 0x17
+PPSMC_MSG_BacoAudioD3PME = 0x18
+PPSMC_MSG_SetSoftMinByFreq = 0x19
+PPSMC_MSG_SetSoftMaxByFreq = 0x1A
+PPSMC_MSG_SetHardMinByFreq = 0x1B
+PPSMC_MSG_SetHardMaxByFreq = 0x1C
+PPSMC_MSG_GetMinDpmFreq = 0x1D
+PPSMC_MSG_GetMaxDpmFreq = 0x1E
+PPSMC_MSG_GetDpmFreqByIndex = 0x1F
+PPSMC_MSG_OverridePcieParameters = 0x20
+PPSMC_MSG_DramLogSetDramAddrHigh = 0x21
+PPSMC_MSG_DramLogSetDramAddrLow = 0x22
+PPSMC_MSG_DramLogSetDramSize = 0x23
+PPSMC_MSG_SetWorkloadMask = 0x24
+PPSMC_MSG_GetVoltageByDpm = 0x25
+PPSMC_MSG_SetVideoFps = 0x26
+PPSMC_MSG_GetDcModeMaxDpmFreq = 0x27
+PPSMC_MSG_AllowGfxOff = 0x28
+PPSMC_MSG_DisallowGfxOff = 0x29
+PPSMC_MSG_PowerUpVcn = 0x2A
+PPSMC_MSG_PowerDownVcn = 0x2B
+PPSMC_MSG_PowerUpJpeg = 0x2C
+PPSMC_MSG_PowerDownJpeg = 0x2D
+PPSMC_MSG_PrepareMp1ForUnload = 0x2E
+PPSMC_MSG_Mode1Reset = 0x2F
+PPSMC_MSG_Mode2Reset = 0x4F
+PPSMC_MSG_SetSystemVirtualDramAddrHigh = 0x30
+PPSMC_MSG_SetSystemVirtualDramAddrLow = 0x31
+PPSMC_MSG_SetPptLimit = 0x32
+PPSMC_MSG_GetPptLimit = 0x33
+PPSMC_MSG_ReenableAcDcInterrupt = 0x34
+PPSMC_MSG_NotifyPowerSource = 0x35
+PPSMC_MSG_RunDcBtc = 0x36
+PPSMC_MSG_GetDebugData = 0x37
+PPSMC_MSG_SetTemperatureInputSelect = 0x38
+PPSMC_MSG_SetFwDstatesMask = 0x39
+PPSMC_MSG_SetThrottlerMask = 0x3A
+PPSMC_MSG_SetExternalClientDfCstateAllow = 0x3B
+PPSMC_MSG_SetMGpuFanBoostLimitRpm = 0x3C
+PPSMC_MSG_DumpSTBtoDram = 0x3D
+PPSMC_MSG_STBtoDramLogSetDramAddrHigh = 0x3E
+PPSMC_MSG_STBtoDramLogSetDramAddrLow = 0x3F
+PPSMC_MSG_STBtoDramLogSetDramSize = 0x40
+PPSMC_MSG_SetGpoAllow = 0x41
+PPSMC_MSG_AllowGfxDcs = 0x42
+PPSMC_MSG_DisallowGfxDcs = 0x43
+PPSMC_MSG_EnableAudioStutterWA = 0x44
+PPSMC_MSG_PowerUpUmsch = 0x45
+PPSMC_MSG_PowerDownUmsch = 0x46
+PPSMC_MSG_SetDcsArch = 0x47
+PPSMC_MSG_TriggerVFFLR = 0x48
+PPSMC_MSG_SetNumBadMemoryPagesRetired = 0x49
+PPSMC_MSG_SetBadMemoryPagesRetiredFlagsPerChannel = 0x4A
+PPSMC_MSG_SetPriorityDeltaGain = 0x4B
+PPSMC_MSG_AllowIHHostInterrupt = 0x4C
+PPSMC_MSG_DALNotPresent = 0x4E
+PPSMC_MSG_EnableUCLKShadow = 0x51
+PPSMC_Message_Count = 0x52
+DEBUGSMC_MSG_TestMessage = 0x1
+DEBUGSMC_MSG_GetDebugData = 0x2
+DEBUGSMC_MSG_DebugDumpExit = 0x3
+DEBUGSMC_Message_Count = 0x4
+SMU13_0_0_DRIVER_IF_VERSION = 0x3D
+PPTABLE_VERSION = 0x2B
+NUM_GFXCLK_DPM_LEVELS = 16
+NUM_SOCCLK_DPM_LEVELS = 8
+NUM_MP0CLK_DPM_LEVELS = 2
+NUM_DCLK_DPM_LEVELS = 8
+NUM_VCLK_DPM_LEVELS = 8
+NUM_DISPCLK_DPM_LEVELS = 8
+NUM_DPPCLK_DPM_LEVELS = 8
+NUM_DPREFCLK_DPM_LEVELS = 8
+NUM_DCFCLK_DPM_LEVELS = 8
+NUM_DTBCLK_DPM_LEVELS = 8
+NUM_UCLK_DPM_LEVELS = 4
+NUM_LINK_LEVELS = 3
+NUM_FCLK_DPM_LEVELS = 8
+NUM_OD_FAN_MAX_POINTS = 6
+FEATURE_FW_DATA_READ_BIT = 0
+FEATURE_DPM_GFXCLK_BIT = 1
+FEATURE_DPM_GFX_POWER_OPTIMIZER_BIT = 2
+FEATURE_DPM_UCLK_BIT = 3
+FEATURE_DPM_FCLK_BIT = 4
+FEATURE_DPM_SOCCLK_BIT = 5
+FEATURE_DPM_MP0CLK_BIT = 6
+FEATURE_DPM_LINK_BIT = 7
+FEATURE_DPM_DCN_BIT = 8
+FEATURE_VMEMP_SCALING_BIT = 9
+FEATURE_VDDIO_MEM_SCALING_BIT = 10
+FEATURE_DS_GFXCLK_BIT = 11
+FEATURE_DS_SOCCLK_BIT = 12
+FEATURE_DS_FCLK_BIT = 13
+FEATURE_DS_LCLK_BIT = 14
+FEATURE_DS_DCFCLK_BIT = 15
+FEATURE_DS_UCLK_BIT = 16
+FEATURE_GFX_ULV_BIT = 17
+FEATURE_FW_DSTATE_BIT = 18
+FEATURE_GFXOFF_BIT = 19
+FEATURE_BACO_BIT = 20
+FEATURE_MM_DPM_BIT = 21
+FEATURE_SOC_MPCLK_DS_BIT = 22
+FEATURE_BACO_MPCLK_DS_BIT = 23
+FEATURE_THROTTLERS_BIT = 24
+FEATURE_SMARTSHIFT_BIT = 25
+FEATURE_GTHR_BIT = 26
+FEATURE_ACDC_BIT = 27
+FEATURE_VR0HOT_BIT = 28
+FEATURE_FW_CTF_BIT = 29
+FEATURE_FAN_CONTROL_BIT = 30
+FEATURE_GFX_DCS_BIT = 31
+FEATURE_GFX_READ_MARGIN_BIT = 32
+FEATURE_LED_DISPLAY_BIT = 33
+FEATURE_GFXCLK_SPREAD_SPECTRUM_BIT = 34
+FEATURE_OUT_OF_BAND_MONITOR_BIT = 35
+FEATURE_OPTIMIZED_VMIN_BIT = 36
+FEATURE_GFX_IMU_BIT = 37
+FEATURE_BOOT_TIME_CAL_BIT = 38
+FEATURE_GFX_PCC_DFLL_BIT = 39
+FEATURE_SOC_CG_BIT = 40
+FEATURE_DF_CSTATE_BIT = 41
+FEATURE_GFX_EDC_BIT = 42
+FEATURE_BOOT_POWER_OPT_BIT = 43
+FEATURE_CLOCK_POWER_DOWN_BYPASS_BIT = 44
+FEATURE_DS_VCN_BIT = 45
+FEATURE_BACO_CG_BIT = 46
+FEATURE_MEM_TEMP_READ_BIT = 47
+FEATURE_ATHUB_MMHUB_PG_BIT = 48
+FEATURE_SOC_PCC_BIT = 49
+FEATURE_EDC_PWRBRK_BIT = 50
+FEATURE_BOMXCO_SVI3_PROG_BIT = 51
+FEATURE_SPARE_52_BIT = 52
+FEATURE_SPARE_53_BIT = 53
+FEATURE_SPARE_54_BIT = 54
+FEATURE_SPARE_55_BIT = 55
+FEATURE_SPARE_56_BIT = 56
+FEATURE_SPARE_57_BIT = 57
+FEATURE_SPARE_58_BIT = 58
+FEATURE_SPARE_59_BIT = 59
+FEATURE_SPARE_60_BIT = 60
+FEATURE_SPARE_61_BIT = 61
+FEATURE_SPARE_62_BIT = 62
+FEATURE_SPARE_63_BIT = 63
+NUM_FEATURES = 64
+ALLOWED_FEATURE_CTRL_DEFAULT = 0xFFFFFFFFFFFFFFFF
+ALLOWED_FEATURE_CTRL_SCPM = ((1 << FEATURE_DPM_GFXCLK_BIT) | (1 << FEATURE_DPM_GFX_POWER_OPTIMIZER_BIT) | (1 << FEATURE_DPM_UCLK_BIT) | (1 << FEATURE_DPM_FCLK_BIT) | (1 << FEATURE_DPM_SOCCLK_BIT) | (1 << FEATURE_DPM_MP0CLK_BIT) | (1 << FEATURE_DPM_LINK_BIT) | (1 << FEATURE_DPM_DCN_BIT) | (1 << FEATURE_DS_GFXCLK_BIT) | (1 << FEATURE_DS_SOCCLK_BIT) | (1 << FEATURE_DS_FCLK_BIT) | (1 << FEATURE_DS_LCLK_BIT) | (1 << FEATURE_DS_DCFCLK_BIT) | (1 << FEATURE_DS_UCLK_BIT) | (1 << FEATURE_DS_VCN_BIT))
+DEBUG_OVERRIDE_DISABLE_VOLT_LINK_VCN_FCLK = 0x00000001
+DEBUG_OVERRIDE_DISABLE_VOLT_LINK_DCN_FCLK = 0x00000002
+DEBUG_OVERRIDE_DISABLE_VOLT_LINK_MP0_FCLK = 0x00000004
+DEBUG_OVERRIDE_DISABLE_VOLT_LINK_VCN_DCFCLK = 0x00000008
+DEBUG_OVERRIDE_DISABLE_FAST_FCLK_TIMER = 0x00000010
+DEBUG_OVERRIDE_DISABLE_VCN_PG = 0x00000020
+DEBUG_OVERRIDE_DISABLE_FMAX_VMAX = 0x00000040
+DEBUG_OVERRIDE_DISABLE_IMU_FW_CHECKS = 0x00000080
+DEBUG_OVERRIDE_DISABLE_D0i2_REENTRY_HSR_TIMER_CHECK = 0x00000100
+DEBUG_OVERRIDE_DISABLE_DFLL = 0x00000200
+DEBUG_OVERRIDE_ENABLE_RLC_VF_BRINGUP_MODE = 0x00000400
+DEBUG_OVERRIDE_DFLL_MASTER_MODE = 0x00000800
+DEBUG_OVERRIDE_ENABLE_PROFILING_MODE = 0x00001000
+VR_MAPPING_VR_SELECT_MASK = 0x01
+VR_MAPPING_VR_SELECT_SHIFT = 0x00
+VR_MAPPING_PLANE_SELECT_MASK = 0x02
+VR_MAPPING_PLANE_SELECT_SHIFT = 0x01
+PSI_SEL_VR0_PLANE0_PSI0 = 0x01
+PSI_SEL_VR0_PLANE0_PSI1 = 0x02
+PSI_SEL_VR0_PLANE1_PSI0 = 0x04
+PSI_SEL_VR0_PLANE1_PSI1 = 0x08
+PSI_SEL_VR1_PLANE0_PSI0 = 0x10
+PSI_SEL_VR1_PLANE0_PSI1 = 0x20
+PSI_SEL_VR1_PLANE1_PSI0 = 0x40
+PSI_SEL_VR1_PLANE1_PSI1 = 0x80
+THROTTLER_TEMP_EDGE_BIT = 0
+THROTTLER_TEMP_HOTSPOT_BIT = 1
+THROTTLER_TEMP_HOTSPOT_G_BIT = 2
+THROTTLER_TEMP_HOTSPOT_M_BIT = 3
+THROTTLER_TEMP_MEM_BIT = 4
+THROTTLER_TEMP_VR_GFX_BIT = 5
+THROTTLER_TEMP_VR_MEM0_BIT = 6
+THROTTLER_TEMP_VR_MEM1_BIT = 7
+THROTTLER_TEMP_VR_SOC_BIT = 8
+THROTTLER_TEMP_VR_U_BIT = 9
+THROTTLER_TEMP_LIQUID0_BIT = 10
+THROTTLER_TEMP_LIQUID1_BIT = 11
+THROTTLER_TEMP_PLX_BIT = 12
+THROTTLER_TDC_GFX_BIT = 13
+THROTTLER_TDC_SOC_BIT = 14
+THROTTLER_TDC_U_BIT = 15
+THROTTLER_PPT0_BIT = 16
+THROTTLER_PPT1_BIT = 17
+THROTTLER_PPT2_BIT = 18
+THROTTLER_PPT3_BIT = 19
+THROTTLER_FIT_BIT = 20
+THROTTLER_GFX_APCC_PLUS_BIT = 21
+THROTTLER_COUNT = 22
+FW_DSTATE_SOC_ULV_BIT = 0
+FW_DSTATE_G6_HSR_BIT = 1
+FW_DSTATE_G6_PHY_VMEMP_OFF_BIT = 2
+FW_DSTATE_SMN_DS_BIT = 3
+FW_DSTATE_MP1_WHISPER_MODE_BIT = 4
+FW_DSTATE_SOC_LIV_MIN_BIT = 5
+FW_DSTATE_SOC_PLL_PWRDN_BIT = 6
+FW_DSTATE_MEM_PLL_PWRDN_BIT = 7
+FW_DSTATE_MALL_ALLOC_BIT = 8
+FW_DSTATE_MEM_PSI_BIT = 9
+FW_DSTATE_HSR_NON_STROBE_BIT = 10
+FW_DSTATE_MP0_ENTER_WFI_BIT = 11
+FW_DSTATE_U_ULV_BIT = 12
+FW_DSTATE_MALL_FLUSH_BIT = 13
+FW_DSTATE_SOC_PSI_BIT = 14
+FW_DSTATE_U_PSI_BIT = 15
+FW_DSTATE_UCP_DS_BIT = 16
+FW_DSTATE_CSRCLK_DS_BIT = 17
+FW_DSTATE_MMHUB_INTERLOCK_BIT = 18
+FW_DSTATE_D0i3_2_QUIET_FW_BIT = 19
+FW_DSTATE_CLDO_PRG_BIT = 20
+FW_DSTATE_DF_PLL_PWRDN_BIT = 21
+FW_DSTATE_U_LOW_PWR_MODE_EN_BIT = 22
+FW_DSTATE_GFX_PSI6_BIT = 23
+FW_DSTATE_GFX_VR_PWR_STAGE_BIT = 24
+LED_DISPLAY_GFX_DPM_BIT = 0
+LED_DISPLAY_PCIE_BIT = 1
+LED_DISPLAY_ERROR_BIT = 2
+MEM_TEMP_READ_OUT_OF_BAND_BIT = 0
+MEM_TEMP_READ_IN_BAND_REFRESH_BIT = 1
+MEM_TEMP_READ_IN_BAND_DUMMY_PSTATE_BIT = 2
+NUM_I2C_CONTROLLERS = 8
+I2C_CONTROLLER_ENABLED = 1
+I2C_CONTROLLER_DISABLED = 0
+MAX_SW_I2C_COMMANDS = 24
+CMDCONFIG_STOP_BIT = 0
+CMDCONFIG_RESTART_BIT = 1
+CMDCONFIG_READWRITE_BIT = 2
+CMDCONFIG_STOP_MASK = (1 << CMDCONFIG_STOP_BIT)
+CMDCONFIG_RESTART_MASK = (1 << CMDCONFIG_RESTART_BIT)
+CMDCONFIG_READWRITE_MASK = (1 << CMDCONFIG_READWRITE_BIT)
+PP_NUM_RTAVFS_PWL_ZONES = 5
+PP_OD_FEATURE_GFX_VF_CURVE_BIT = 0
+PP_OD_FEATURE_PPT_BIT = 2
+PP_OD_FEATURE_FAN_CURVE_BIT = 3
+PP_OD_FEATURE_GFXCLK_BIT = 7
+PP_OD_FEATURE_UCLK_BIT = 8
+PP_OD_FEATURE_ZERO_FAN_BIT = 9
+PP_OD_FEATURE_TEMPERATURE_BIT = 10
+PP_OD_FEATURE_COUNT = 13
+PP_NUM_OD_VF_CURVE_POINTS = PP_NUM_RTAVFS_PWL_ZONES + 1
+INVALID_BOARD_GPIO = 0xFF
+MARKETING_BASE_CLOCKS = 0
+MARKETING_GAME_CLOCKS = 1
+MARKETING_BOOST_CLOCKS = 2
+NUM_WM_RANGES = 4
+WORKLOAD_PPLIB_DEFAULT_BIT = 0
+WORKLOAD_PPLIB_FULL_SCREEN_3D_BIT = 1
+WORKLOAD_PPLIB_POWER_SAVING_BIT = 2
+WORKLOAD_PPLIB_VIDEO_BIT = 3
+WORKLOAD_PPLIB_VR_BIT = 4
+WORKLOAD_PPLIB_COMPUTE_BIT = 5
+WORKLOAD_PPLIB_CUSTOM_BIT = 6
+WORKLOAD_PPLIB_WINDOW_3D_BIT = 7
+WORKLOAD_PPLIB_COUNT = 8
+TABLE_TRANSFER_OK = 0x0
+TABLE_TRANSFER_FAILED = 0xFF
+TABLE_TRANSFER_PENDING = 0xAB
+TABLE_PPTABLE = 0
+TABLE_COMBO_PPTABLE = 1
+TABLE_WATERMARKS = 2
+TABLE_AVFS_PSM_DEBUG = 3
+TABLE_PMSTATUSLOG = 4
+TABLE_SMU_METRICS = 5
+TABLE_DRIVER_SMU_CONFIG = 6
+TABLE_ACTIVITY_MONITOR_COEFF = 7
+TABLE_OVERDRIVE = 8
+TABLE_I2C_COMMANDS = 9
+TABLE_DRIVER_INFO = 10
+TABLE_ECCINFO = 11
+TABLE_WIFIBAND = 12
+TABLE_COUNT = 13
+IH_INTERRUPT_ID_TO_DRIVER = 0xFE
+IH_INTERRUPT_CONTEXT_ID_BACO = 0x2
+IH_INTERRUPT_CONTEXT_ID_AC = 0x3
+IH_INTERRUPT_CONTEXT_ID_DC = 0x4
+IH_INTERRUPT_CONTEXT_ID_AUDIO_D0 = 0x5
+IH_INTERRUPT_CONTEXT_ID_AUDIO_D3 = 0x6
+IH_INTERRUPT_CONTEXT_ID_THERMAL_THROTTLING = 0x7
+IH_INTERRUPT_CONTEXT_ID_FAN_ABNORMAL = 0x8
+IH_INTERRUPT_CONTEXT_ID_FAN_RECOVERY = 0x9
+SMU_THERMAL_MINIMUM_ALERT_TEMP = 0
+SMU_THERMAL_MAXIMUM_ALERT_TEMP = 255
+SMU_TEMPERATURE_UNITS_PER_CENTIGRADES = 1000
+SMU_FW_NAME_LEN = 0x24
+SMU_DPM_USER_PROFILE_RESTORE = (1 << 0)
+SMU_CUSTOM_FAN_SPEED_RPM = (1 << 1)
+SMU_CUSTOM_FAN_SPEED_PWM = (1 << 2)
+SMU_THROTTLER_PPT0_BIT = 0
+SMU_THROTTLER_PPT1_BIT = 1
+SMU_THROTTLER_PPT2_BIT = 2
+SMU_THROTTLER_PPT3_BIT = 3
+SMU_THROTTLER_SPL_BIT = 4
+SMU_THROTTLER_FPPT_BIT = 5
+SMU_THROTTLER_SPPT_BIT = 6
+SMU_THROTTLER_SPPT_APU_BIT = 7
+SMU_THROTTLER_TDC_GFX_BIT = 16
+SMU_THROTTLER_TDC_SOC_BIT = 17
+SMU_THROTTLER_TDC_MEM_BIT = 18
+SMU_THROTTLER_TDC_VDD_BIT = 19
+SMU_THROTTLER_TDC_CVIP_BIT = 20
+SMU_THROTTLER_EDC_CPU_BIT = 21
+SMU_THROTTLER_EDC_GFX_BIT = 22
+SMU_THROTTLER_APCC_BIT = 23
+SMU_THROTTLER_TEMP_GPU_BIT = 32
+SMU_THROTTLER_TEMP_CORE_BIT = 33
+SMU_THROTTLER_TEMP_MEM_BIT = 34
+SMU_THROTTLER_TEMP_EDGE_BIT = 35
+SMU_THROTTLER_TEMP_HOTSPOT_BIT = 36
+SMU_THROTTLER_TEMP_SOC_BIT = 37
+SMU_THROTTLER_TEMP_VR_GFX_BIT = 38
+SMU_THROTTLER_TEMP_VR_SOC_BIT = 39
+SMU_THROTTLER_TEMP_VR_MEM0_BIT = 40
+SMU_THROTTLER_TEMP_VR_MEM1_BIT = 41
+SMU_THROTTLER_TEMP_LIQUID0_BIT = 42
+SMU_THROTTLER_TEMP_LIQUID1_BIT = 43
+SMU_THROTTLER_VRHOT0_BIT = 44
+SMU_THROTTLER_VRHOT1_BIT = 45
+SMU_THROTTLER_PROCHOT_CPU_BIT = 46
+SMU_THROTTLER_PROCHOT_GFX_BIT = 47
+SMU_THROTTLER_PPM_BIT = 56
+SMU_THROTTLER_FIT_BIT = 57

tinygrad/runtime/autogen/am/smu_v13_0_12.py

@@ -230,7 +230,7 @@ class struct_smu_state_memory_block(c.Struct):
   dll_off: bool
   m3arb: int
   unused: c.Array[ctypes.c_ubyte, Literal[3]]
-struct_smu_state_memory_block.register_fields([('dll_off', ctypes.c_bool, 0), ('m3arb', ctypes.c_ubyte, 1), ('unused', c.Array[ctypes.c_ubyte, Literal[3]], 2)])
+struct_smu_state_memory_block.register_fields([('dll_off', ctypes.c_bool, 0), ('m3arb', uint8_t, 1), ('unused', c.Array[uint8_t, Literal[3]], 2)])
 @c.record
 class struct_smu_state_software_algorithm_block(c.Struct):
   SIZE = 2
@@ -258,13 +258,13 @@ class struct_smu_state_validation_block(c.Struct):
   single_display_only: bool
   disallow_on_dc: bool
   supported_power_levels: int
-struct_smu_state_validation_block.register_fields([('single_display_only', ctypes.c_bool, 0), ('disallow_on_dc', ctypes.c_bool, 1), ('supported_power_levels', ctypes.c_ubyte, 2)])
+struct_smu_state_validation_block.register_fields([('single_display_only', ctypes.c_bool, 0), ('disallow_on_dc', ctypes.c_bool, 1), ('supported_power_levels', uint8_t, 2)])
 @c.record
 class struct_smu_uvd_clocks(c.Struct):
   SIZE = 8
   vclk: int
   dclk: int
-struct_smu_uvd_clocks.register_fields([('vclk', ctypes.c_uint32, 0), ('dclk', ctypes.c_uint32, 4)])
+struct_smu_uvd_clocks.register_fields([('vclk', uint32_t, 0), ('dclk', uint32_t, 4)])
 enum_smu_power_src_type: dict[int, str] = {(SMU_POWER_SOURCE_AC:=0): 'SMU_POWER_SOURCE_AC', (SMU_POWER_SOURCE_DC:=1): 'SMU_POWER_SOURCE_DC', (SMU_POWER_SOURCE_COUNT:=2): 'SMU_POWER_SOURCE_COUNT'}
 enum_smu_ppt_limit_type: dict[int, str] = {(SMU_DEFAULT_PPT_LIMIT:=0): 'SMU_DEFAULT_PPT_LIMIT', (SMU_FAST_PPT_LIMIT:=1): 'SMU_FAST_PPT_LIMIT'}
 enum_smu_ppt_limit_level: dict[int, str] = {(SMU_PPT_LIMIT_MIN:=-1): 'SMU_PPT_LIMIT_MIN', (SMU_PPT_LIMIT_CURRENT:=0): 'SMU_PPT_LIMIT_CURRENT', (SMU_PPT_LIMIT_DEFAULT:=1): 'SMU_PPT_LIMIT_DEFAULT', (SMU_PPT_LIMIT_MAX:=2): 'SMU_PPT_LIMIT_MAX'}
@@ -281,7 +281,7 @@ class struct_smu_user_dpm_profile(c.Struct):
   user_od: int
   clk_mask: c.Array[ctypes.c_uint32, Literal[28]]
   clk_dependency: int
-struct_smu_user_dpm_profile.register_fields([('fan_mode', ctypes.c_uint32, 0), ('power_limit', ctypes.c_uint32, 4), ('fan_speed_pwm', ctypes.c_uint32, 8), ('fan_speed_rpm', ctypes.c_uint32, 12), ('flags', ctypes.c_uint32, 16), ('user_od', ctypes.c_uint32, 20), ('clk_mask', c.Array[ctypes.c_uint32, Literal[28]], 24), ('clk_dependency', ctypes.c_uint32, 136)])
+struct_smu_user_dpm_profile.register_fields([('fan_mode', uint32_t, 0), ('power_limit', uint32_t, 4), ('fan_speed_pwm', uint32_t, 8), ('fan_speed_rpm', uint32_t, 12), ('flags', uint32_t, 16), ('user_od', uint32_t, 20), ('clk_mask', c.Array[uint32_t, Literal[28]], 24), ('clk_dependency', uint32_t, 136)])
 @c.record
 class struct_smu_table(c.Struct):
   SIZE = 48
@@ -293,7 +293,7 @@ class struct_smu_table(c.Struct):
   bo: c.POINTER[struct_amdgpu_bo]
   version: int
 class struct_amdgpu_bo(c.Struct): pass
-struct_smu_table.register_fields([('size', ctypes.c_uint64, 0), ('align', ctypes.c_uint32, 8), ('domain', ctypes.c_ubyte, 12), ('mc_address', ctypes.c_uint64, 16), ('cpu_addr', ctypes.c_void_p, 24), ('bo', c.POINTER[struct_amdgpu_bo], 32), ('version', ctypes.c_uint32, 40)])
+struct_smu_table.register_fields([('size', uint64_t, 0), ('align', uint32_t, 8), ('domain', uint8_t, 12), ('mc_address', uint64_t, 16), ('cpu_addr', ctypes.c_void_p, 24), ('bo', c.POINTER[struct_amdgpu_bo], 32), ('version', uint32_t, 40)])
 enum_smu_perf_level_designation: dict[int, str] = {(PERF_LEVEL_ACTIVITY:=0): 'PERF_LEVEL_ACTIVITY', (PERF_LEVEL_POWER_CONTAINMENT:=1): 'PERF_LEVEL_POWER_CONTAINMENT'}
 @c.record
 class struct_smu_performance_level(c.Struct):
@@ -304,7 +304,7 @@ class struct_smu_performance_level(c.Struct):
   vddci: int
   non_local_mem_freq: int
   non_local_mem_width: int
-struct_smu_performance_level.register_fields([('core_clock', ctypes.c_uint32, 0), ('memory_clock', ctypes.c_uint32, 4), ('vddc', ctypes.c_uint32, 8), ('vddci', ctypes.c_uint32, 12), ('non_local_mem_freq', ctypes.c_uint32, 16), ('non_local_mem_width', ctypes.c_uint32, 20)])
+struct_smu_performance_level.register_fields([('core_clock', uint32_t, 0), ('memory_clock', uint32_t, 4), ('vddc', uint32_t, 8), ('vddci', uint32_t, 12), ('non_local_mem_freq', uint32_t, 16), ('non_local_mem_width', uint32_t, 20)])
 @c.record
 class struct_smu_clock_info(c.Struct):
   SIZE = 24
@@ -314,7 +314,7 @@ class struct_smu_clock_info(c.Struct):
   max_eng_clk: int
   min_bus_bandwidth: int
   max_bus_bandwidth: int
-struct_smu_clock_info.register_fields([('min_mem_clk', ctypes.c_uint32, 0), ('max_mem_clk', ctypes.c_uint32, 4), ('min_eng_clk', ctypes.c_uint32, 8), ('max_eng_clk', ctypes.c_uint32, 12), ('min_bus_bandwidth', ctypes.c_uint32, 16), ('max_bus_bandwidth', ctypes.c_uint32, 20)])
+struct_smu_clock_info.register_fields([('min_mem_clk', uint32_t, 0), ('max_mem_clk', uint32_t, 4), ('min_eng_clk', uint32_t, 8), ('max_eng_clk', uint32_t, 12), ('min_bus_bandwidth', uint32_t, 16), ('max_bus_bandwidth', uint32_t, 20)])
 @c.record
 class struct_smu_bios_boot_up_values(c.Struct):
   SIZE = 68
@@ -337,197 +337,196 @@ class struct_smu_bios_boot_up_values(c.Struct):
   fclk: int
   lclk: int
   firmware_caps: int
-struct_smu_bios_boot_up_values.register_fields([('revision', ctypes.c_uint32, 0), ('gfxclk', ctypes.c_uint32, 4), ('uclk', ctypes.c_uint32, 8), ('socclk', ctypes.c_uint32, 12), ('dcefclk', ctypes.c_uint32, 16), ('eclk', ctypes.c_uint32, 20), ('vclk', ctypes.c_uint32, 24), ('dclk', ctypes.c_uint32, 28), ('vddc', ctypes.c_uint16, 32), ('vddci', ctypes.c_uint16, 34), ('mvddc', ctypes.c_uint16, 36), ('vdd_gfx', ctypes.c_uint16, 38), ('cooling_id', ctypes.c_ubyte, 40), ('pp_table_id', ctypes.c_uint32, 44), ('format_revision', ctypes.c_uint32, 48), ('content_revision', ctypes.c_uint32, 52), ('fclk', ctypes.c_uint32, 56), ('lclk', ctypes.c_uint32, 60), ('firmware_caps', ctypes.c_uint32, 64)])
+struct_smu_bios_boot_up_values.register_fields([('revision', uint32_t, 0), ('gfxclk', uint32_t, 4), ('uclk', uint32_t, 8), ('socclk', uint32_t, 12), ('dcefclk', uint32_t, 16), ('eclk', uint32_t, 20), ('vclk', uint32_t, 24), ('dclk', uint32_t, 28), ('vddc', uint16_t, 32), ('vddci', uint16_t, 34), ('mvddc', uint16_t, 36), ('vdd_gfx', uint16_t, 38), ('cooling_id', uint8_t, 40), ('pp_table_id', uint32_t, 44), ('format_revision', uint32_t, 48), ('content_revision', uint32_t, 52), ('fclk', uint32_t, 56), ('lclk', uint32_t, 60), ('firmware_caps', uint32_t, 64)])
 enum_smu_table_id: dict[int, str] = {(SMU_TABLE_PPTABLE:=0): 'SMU_TABLE_PPTABLE', (SMU_TABLE_WATERMARKS:=1): 'SMU_TABLE_WATERMARKS', (SMU_TABLE_CUSTOM_DPM:=2): 'SMU_TABLE_CUSTOM_DPM', (SMU_TABLE_DPMCLOCKS:=3): 'SMU_TABLE_DPMCLOCKS', (SMU_TABLE_AVFS:=4): 'SMU_TABLE_AVFS', (SMU_TABLE_AVFS_PSM_DEBUG:=5): 'SMU_TABLE_AVFS_PSM_DEBUG', (SMU_TABLE_AVFS_FUSE_OVERRIDE:=6): 'SMU_TABLE_AVFS_FUSE_OVERRIDE', (SMU_TABLE_PMSTATUSLOG:=7): 'SMU_TABLE_PMSTATUSLOG', (SMU_TABLE_SMU_METRICS:=8): 'SMU_TABLE_SMU_METRICS', (SMU_TABLE_DRIVER_SMU_CONFIG:=9): 'SMU_TABLE_DRIVER_SMU_CONFIG', (SMU_TABLE_ACTIVITY_MONITOR_COEFF:=10): 'SMU_TABLE_ACTIVITY_MONITOR_COEFF', (SMU_TABLE_OVERDRIVE:=11): 'SMU_TABLE_OVERDRIVE', (SMU_TABLE_I2C_COMMANDS:=12): 'SMU_TABLE_I2C_COMMANDS', (SMU_TABLE_PACE:=13): 'SMU_TABLE_PACE', (SMU_TABLE_ECCINFO:=14): 'SMU_TABLE_ECCINFO', (SMU_TABLE_COMBO_PPTABLE:=15): 'SMU_TABLE_COMBO_PPTABLE', (SMU_TABLE_WIFIBAND:=16): 'SMU_TABLE_WIFIBAND', (SMU_TABLE_COUNT:=17): 'SMU_TABLE_COUNT'}
-PPSMC_Result_OK = 0x1 # type: ignore
-PPSMC_Result_Failed = 0xFF # type: ignore
-PPSMC_Result_UnknownCmd = 0xFE # type: ignore
-PPSMC_Result_CmdRejectedPrereq = 0xFD # type: ignore
-PPSMC_Result_CmdRejectedBusy = 0xFC # type: ignore
-PPSMC_MSG_TestMessage = 0x1 # type: ignore
-PPSMC_MSG_GetSmuVersion = 0x2 # type: ignore
-PPSMC_MSG_GfxDriverReset = 0x3 # type: ignore
-PPSMC_MSG_GetDriverIfVersion = 0x4 # type: ignore
-PPSMC_MSG_EnableAllSmuFeatures = 0x5 # type: ignore
-PPSMC_MSG_DisableAllSmuFeatures = 0x6 # type: ignore
-PPSMC_MSG_RequestI2cTransaction = 0x7 # type: ignore
-PPSMC_MSG_GetMetricsVersion = 0x8 # type: ignore
-PPSMC_MSG_GetMetricsTable = 0x9 # type: ignore
-PPSMC_MSG_GetEccInfoTable = 0xA # type: ignore
-PPSMC_MSG_GetEnabledSmuFeaturesLow = 0xB # type: ignore
-PPSMC_MSG_GetEnabledSmuFeaturesHigh = 0xC # type: ignore
-PPSMC_MSG_SetDriverDramAddrHigh = 0xD # type: ignore
-PPSMC_MSG_SetDriverDramAddrLow = 0xE # type: ignore
-PPSMC_MSG_SetToolsDramAddrHigh = 0xF # type: ignore
-PPSMC_MSG_SetToolsDramAddrLow = 0x10 # type: ignore
-PPSMC_MSG_SetSystemVirtualDramAddrHigh = 0x11 # type: ignore
-PPSMC_MSG_SetSystemVirtualDramAddrLow = 0x12 # type: ignore
-PPSMC_MSG_SetSoftMinByFreq = 0x13 # type: ignore
-PPSMC_MSG_SetSoftMaxByFreq = 0x14 # type: ignore
-PPSMC_MSG_GetMinDpmFreq = 0x15 # type: ignore
-PPSMC_MSG_GetMaxDpmFreq = 0x16 # type: ignore
-PPSMC_MSG_GetDpmFreqByIndex = 0x17 # type: ignore
-PPSMC_MSG_SetPptLimit = 0x18 # type: ignore
-PPSMC_MSG_GetPptLimit = 0x19 # type: ignore
-PPSMC_MSG_DramLogSetDramAddrHigh = 0x1A # type: ignore
-PPSMC_MSG_DramLogSetDramAddrLow = 0x1B # type: ignore
-PPSMC_MSG_DramLogSetDramSize = 0x1C # type: ignore
-PPSMC_MSG_GetDebugData = 0x1D # type: ignore
-PPSMC_MSG_HeavySBR = 0x1E # type: ignore
-PPSMC_MSG_SetNumBadHbmPagesRetired = 0x1F # type: ignore
-PPSMC_MSG_DFCstateControl = 0x20 # type: ignore
-PPSMC_MSG_GetGmiPwrDnHyst = 0x21 # type: ignore
-PPSMC_MSG_SetGmiPwrDnHyst = 0x22 # type: ignore
-PPSMC_MSG_GmiPwrDnControl = 0x23 # type: ignore
-PPSMC_MSG_EnterGfxoff = 0x24 # type: ignore
-PPSMC_MSG_ExitGfxoff = 0x25 # type: ignore
-PPSMC_MSG_EnableDeterminism = 0x26 # type: ignore
-PPSMC_MSG_DisableDeterminism = 0x27 # type: ignore
-PPSMC_MSG_DumpSTBtoDram = 0x28 # type: ignore
-PPSMC_MSG_STBtoDramLogSetDramAddrHigh = 0x29 # type: ignore
-PPSMC_MSG_STBtoDramLogSetDramAddrLow = 0x2A # type: ignore
-PPSMC_MSG_STBtoDramLogSetDramSize = 0x2B # type: ignore
-PPSMC_MSG_SetSystemVirtualSTBtoDramAddrHigh = 0x2C # type: ignore
-PPSMC_MSG_SetSystemVirtualSTBtoDramAddrLow = 0x2D # type: ignore
-PPSMC_MSG_GfxDriverResetRecovery = 0x2E # type: ignore
-PPSMC_MSG_TriggerVFFLR = 0x2F # type: ignore
-PPSMC_MSG_SetSoftMinGfxClk = 0x30 # type: ignore
-PPSMC_MSG_SetSoftMaxGfxClk = 0x31 # type: ignore
-PPSMC_MSG_GetMinGfxDpmFreq = 0x32 # type: ignore
-PPSMC_MSG_GetMaxGfxDpmFreq = 0x33 # type: ignore
-PPSMC_MSG_PrepareForDriverUnload = 0x34 # type: ignore
-PPSMC_MSG_ReadThrottlerLimit = 0x35 # type: ignore
-PPSMC_MSG_QueryValidMcaCount = 0x36 # type: ignore
-PPSMC_MSG_McaBankDumpDW = 0x37 # type: ignore
-PPSMC_MSG_GetCTFLimit = 0x38 # type: ignore
-PPSMC_MSG_ClearMcaOnRead = 0x39 # type: ignore
-PPSMC_MSG_QueryValidMcaCeCount = 0x3A # type: ignore
-PPSMC_MSG_McaBankCeDumpDW = 0x3B # type: ignore
-PPSMC_MSG_SelectPLPDMode = 0x40 # type: ignore
-PPSMC_MSG_PmLogReadSample = 0x41 # type: ignore
-PPSMC_MSG_PmLogGetTableVersion = 0x42 # type: ignore
-PPSMC_MSG_RmaDueToBadPageThreshold = 0x43 # type: ignore
-PPSMC_MSG_SetThrottlingPolicy = 0x44 # type: ignore
-PPSMC_MSG_SetPhaseDetectCSBWThreshold = 0x45 # type: ignore
-PPSMC_MSG_SetPhaseDetectFreqHigh = 0x46 # type: ignore
-PPSMC_MSG_SetPhaseDetectFreqLow = 0x47 # type: ignore
-PPSMC_MSG_SetPhaseDetectDownHysterisis = 0x48 # type: ignore
-PPSMC_MSG_SetPhaseDetectAlphaX1e6 = 0x49 # type: ignore
-PPSMC_MSG_SetPhaseDetectOnOff = 0x4A # type: ignore
-PPSMC_MSG_GetPhaseDetectResidency = 0x4B # type: ignore
-PPSMC_MSG_UpdatePccWaitDecMaxStr = 0x4C # type: ignore
-PPSMC_MSG_ResetSDMA = 0x4D # type: ignore
-PPSMC_MSG_GetRasTableVersion = 0x4E # type: ignore
-PPSMC_MSG_GetBadPageCount = 0x50 # type: ignore
-PPSMC_MSG_GetBadPageMcaAddress = 0x51 # type: ignore
-PPSMC_MSG_SetTimestamp = 0x53 # type: ignore
-PPSMC_MSG_SetTimestampHi = 0x54 # type: ignore
-PPSMC_MSG_GetTimestamp = 0x55 # type: ignore
-PPSMC_MSG_GetBadPageIpIdLoHi = 0x57 # type: ignore
-PPSMC_MSG_EraseRasTable = 0x58 # type: ignore
-PPSMC_MSG_GetStaticMetricsTable = 0x59 # type: ignore
-PPSMC_MSG_ResetVfArbitersByIndex = 0x5A # type: ignore
-PPSMC_MSG_GetSystemMetricsTable = 0x5C # type: ignore
-PPSMC_MSG_GetSystemMetricsVersion = 0x5D # type: ignore
-PPSMC_MSG_ResetVCN = 0x5E # type: ignore
-PPSMC_MSG_SetFastPptLimit = 0x5F # type: ignore
-PPSMC_MSG_GetFastPptLimit = 0x60 # type: ignore
-PPSMC_Message_Count = 0x61 # type: ignore
-PPSMC_RESET_TYPE_DRIVER_MODE_1_RESET = 0x1 # type: ignore
-PPSMC_RESET_TYPE_DRIVER_MODE_2_RESET = 0x2 # type: ignore
-PPSMC_RESET_TYPE_DRIVER_MODE_3_RESET = 0x3 # type: ignore
-PPSMC_THROTTLING_LIMIT_TYPE_SOCKET = 0x1 # type: ignore
-PPSMC_THROTTLING_LIMIT_TYPE_HBM = 0x2 # type: ignore
-PPSMC_AID_THM_TYPE = 0x1 # type: ignore
-PPSMC_CCD_THM_TYPE = 0x2 # type: ignore
-PPSMC_XCD_THM_TYPE = 0x3 # type: ignore
-PPSMC_HBM_THM_TYPE = 0x4 # type: ignore
-PPSMC_PLPD_MODE_DEFAULT = 0x1 # type: ignore
-PPSMC_PLPD_MODE_OPTIMIZED = 0x2 # type: ignore
-NUM_VCLK_DPM_LEVELS = 4 # type: ignore
-NUM_DCLK_DPM_LEVELS = 4 # type: ignore
-NUM_SOCCLK_DPM_LEVELS = 4 # type: ignore
-NUM_LCLK_DPM_LEVELS = 4 # type: ignore
-NUM_UCLK_DPM_LEVELS = 4 # type: ignore
-NUM_FCLK_DPM_LEVELS = 4 # type: ignore
-NUM_XGMI_DPM_LEVELS = 2 # type: ignore
-NUM_CXL_BITRATES = 4 # type: ignore
-NUM_PCIE_BITRATES = 4 # type: ignore
-NUM_XGMI_BITRATES = 4 # type: ignore
-NUM_XGMI_WIDTHS = 3 # type: ignore
-NUM_TDP_GROUPS = 4 # type: ignore
-NUM_SOC_P2S_TABLES = 6 # type: ignore
-NUM_GFX_P2S_TABLES = 8 # type: ignore
-NUM_PSM_DIDT_THRESHOLDS = 3 # type: ignore
-NUM_XVMIN_VMIN_THRESHOLDS = 3 # type: ignore
-PRODUCT_MODEL_NUMBER_LEN = 20 # type: ignore
-PRODUCT_NAME_LEN = 64 # type: ignore
-PRODUCT_SERIAL_LEN = 20 # type: ignore
-PRODUCT_MANUFACTURER_NAME_LEN = 32 # type: ignore
-PRODUCT_FRU_ID_LEN = 32 # type: ignore
-SMU_METRICS_TABLE_VERSION = 0x15 # type: ignore
-SMU_SYSTEM_METRICS_TABLE_VERSION = 0x1 # type: ignore
-SMU_VF_METRICS_TABLE_MASK = (1 << 31) # type: ignore
-SMU_VF_METRICS_TABLE_VERSION = (0x6 | SMU_VF_METRICS_TABLE_MASK) # type: ignore
-SMU13_0_6_DRIVER_IF_VERSION = 0x08042024 # type: ignore
-NUM_I2C_CONTROLLERS = 8 # type: ignore
-I2C_CONTROLLER_ENABLED = 1 # type: ignore
-I2C_CONTROLLER_DISABLED = 0 # type: ignore
-MAX_SW_I2C_COMMANDS = 24 # type: ignore
-CMDCONFIG_STOP_BIT = 0 # type: ignore
-CMDCONFIG_RESTART_BIT = 1 # type: ignore
-CMDCONFIG_READWRITE_BIT = 2 # type: ignore
-CMDCONFIG_STOP_MASK = (1 << CMDCONFIG_STOP_BIT) # type: ignore
-CMDCONFIG_RESTART_MASK = (1 << CMDCONFIG_RESTART_BIT) # type: ignore
-CMDCONFIG_READWRITE_MASK = (1 << CMDCONFIG_READWRITE_BIT) # type: ignore
-IH_INTERRUPT_ID_TO_DRIVER = 0xFE # type: ignore
-IH_INTERRUPT_CONTEXT_ID_THERMAL_THROTTLING = 0x7 # type: ignore
-THROTTLER_PROCHOT_BIT = 0 # type: ignore
-THROTTLER_PPT_BIT = 1 # type: ignore
-THROTTLER_THERMAL_SOCKET_BIT = 2 # type: ignore
-THROTTLER_THERMAL_VR_BIT = 3 # type: ignore
-THROTTLER_THERMAL_HBM_BIT = 4 # type: ignore
-ClearMcaOnRead_UE_FLAG_MASK = 0x1 # type: ignore
-ClearMcaOnRead_CE_POLL_MASK = 0x2 # type: ignore
-int32_t = int # type: ignore
-SMU_THERMAL_MINIMUM_ALERT_TEMP = 0 # type: ignore
-SMU_THERMAL_MAXIMUM_ALERT_TEMP = 255 # type: ignore
-SMU_TEMPERATURE_UNITS_PER_CENTIGRADES = 1000 # type: ignore
-SMU_FW_NAME_LEN = 0x24 # type: ignore
-SMU_DPM_USER_PROFILE_RESTORE = (1 << 0) # type: ignore
-SMU_CUSTOM_FAN_SPEED_RPM = (1 << 1) # type: ignore
-SMU_CUSTOM_FAN_SPEED_PWM = (1 << 2) # type: ignore
-SMU_THROTTLER_PPT0_BIT = 0 # type: ignore
-SMU_THROTTLER_PPT1_BIT = 1 # type: ignore
-SMU_THROTTLER_PPT2_BIT = 2 # type: ignore
-SMU_THROTTLER_PPT3_BIT = 3 # type: ignore
-SMU_THROTTLER_SPL_BIT = 4 # type: ignore
-SMU_THROTTLER_FPPT_BIT = 5 # type: ignore
-SMU_THROTTLER_SPPT_BIT = 6 # type: ignore
-SMU_THROTTLER_SPPT_APU_BIT = 7 # type: ignore
-SMU_THROTTLER_TDC_GFX_BIT = 16 # type: ignore
-SMU_THROTTLER_TDC_SOC_BIT = 17 # type: ignore
-SMU_THROTTLER_TDC_MEM_BIT = 18 # type: ignore
-SMU_THROTTLER_TDC_VDD_BIT = 19 # type: ignore
-SMU_THROTTLER_TDC_CVIP_BIT = 20 # type: ignore
-SMU_THROTTLER_EDC_CPU_BIT = 21 # type: ignore
-SMU_THROTTLER_EDC_GFX_BIT = 22 # type: ignore
-SMU_THROTTLER_APCC_BIT = 23 # type: ignore
-SMU_THROTTLER_TEMP_GPU_BIT = 32 # type: ignore
-SMU_THROTTLER_TEMP_CORE_BIT = 33 # type: ignore
-SMU_THROTTLER_TEMP_MEM_BIT = 34 # type: ignore
-SMU_THROTTLER_TEMP_EDGE_BIT = 35 # type: ignore
-SMU_THROTTLER_TEMP_HOTSPOT_BIT = 36 # type: ignore
-SMU_THROTTLER_TEMP_SOC_BIT = 37 # type: ignore
-SMU_THROTTLER_TEMP_VR_GFX_BIT = 38 # type: ignore
-SMU_THROTTLER_TEMP_VR_SOC_BIT = 39 # type: ignore
-SMU_THROTTLER_TEMP_VR_MEM0_BIT = 40 # type: ignore
-SMU_THROTTLER_TEMP_VR_MEM1_BIT = 41 # type: ignore
-SMU_THROTTLER_TEMP_LIQUID0_BIT = 42 # type: ignore
-SMU_THROTTLER_TEMP_LIQUID1_BIT = 43 # type: ignore
-SMU_THROTTLER_VRHOT0_BIT = 44 # type: ignore
-SMU_THROTTLER_VRHOT1_BIT = 45 # type: ignore
-SMU_THROTTLER_PROCHOT_CPU_BIT = 46 # type: ignore
-SMU_THROTTLER_PROCHOT_GFX_BIT = 47 # type: ignore
-SMU_THROTTLER_PPM_BIT = 56 # type: ignore
-SMU_THROTTLER_FIT_BIT = 57 # type: ignore
+PPSMC_Result_OK = 0x1
+PPSMC_Result_Failed = 0xFF
+PPSMC_Result_UnknownCmd = 0xFE
+PPSMC_Result_CmdRejectedPrereq = 0xFD
+PPSMC_Result_CmdRejectedBusy = 0xFC
+PPSMC_MSG_TestMessage = 0x1
+PPSMC_MSG_GetSmuVersion = 0x2
+PPSMC_MSG_GfxDriverReset = 0x3
+PPSMC_MSG_GetDriverIfVersion = 0x4
+PPSMC_MSG_EnableAllSmuFeatures = 0x5
+PPSMC_MSG_DisableAllSmuFeatures = 0x6
+PPSMC_MSG_RequestI2cTransaction = 0x7
+PPSMC_MSG_GetMetricsVersion = 0x8
+PPSMC_MSG_GetMetricsTable = 0x9
+PPSMC_MSG_GetEccInfoTable = 0xA
+PPSMC_MSG_GetEnabledSmuFeaturesLow = 0xB
+PPSMC_MSG_GetEnabledSmuFeaturesHigh = 0xC
+PPSMC_MSG_SetDriverDramAddrHigh = 0xD
+PPSMC_MSG_SetDriverDramAddrLow = 0xE
+PPSMC_MSG_SetToolsDramAddrHigh = 0xF
+PPSMC_MSG_SetToolsDramAddrLow = 0x10
+PPSMC_MSG_SetSystemVirtualDramAddrHigh = 0x11
+PPSMC_MSG_SetSystemVirtualDramAddrLow = 0x12
+PPSMC_MSG_SetSoftMinByFreq = 0x13
+PPSMC_MSG_SetSoftMaxByFreq = 0x14
+PPSMC_MSG_GetMinDpmFreq = 0x15
+PPSMC_MSG_GetMaxDpmFreq = 0x16
+PPSMC_MSG_GetDpmFreqByIndex = 0x17
+PPSMC_MSG_SetPptLimit = 0x18
+PPSMC_MSG_GetPptLimit = 0x19
+PPSMC_MSG_DramLogSetDramAddrHigh = 0x1A
+PPSMC_MSG_DramLogSetDramAddrLow = 0x1B
+PPSMC_MSG_DramLogSetDramSize = 0x1C
+PPSMC_MSG_GetDebugData = 0x1D
+PPSMC_MSG_HeavySBR = 0x1E
+PPSMC_MSG_SetNumBadHbmPagesRetired = 0x1F
+PPSMC_MSG_DFCstateControl = 0x20
+PPSMC_MSG_GetGmiPwrDnHyst = 0x21
+PPSMC_MSG_SetGmiPwrDnHyst = 0x22
+PPSMC_MSG_GmiPwrDnControl = 0x23
+PPSMC_MSG_EnterGfxoff = 0x24
+PPSMC_MSG_ExitGfxoff = 0x25
+PPSMC_MSG_EnableDeterminism = 0x26
+PPSMC_MSG_DisableDeterminism = 0x27
+PPSMC_MSG_DumpSTBtoDram = 0x28
+PPSMC_MSG_STBtoDramLogSetDramAddrHigh = 0x29
+PPSMC_MSG_STBtoDramLogSetDramAddrLow = 0x2A
+PPSMC_MSG_STBtoDramLogSetDramSize = 0x2B
+PPSMC_MSG_SetSystemVirtualSTBtoDramAddrHigh = 0x2C
+PPSMC_MSG_SetSystemVirtualSTBtoDramAddrLow = 0x2D
+PPSMC_MSG_GfxDriverResetRecovery = 0x2E
+PPSMC_MSG_TriggerVFFLR = 0x2F
+PPSMC_MSG_SetSoftMinGfxClk = 0x30
+PPSMC_MSG_SetSoftMaxGfxClk = 0x31
+PPSMC_MSG_GetMinGfxDpmFreq = 0x32
+PPSMC_MSG_GetMaxGfxDpmFreq = 0x33
+PPSMC_MSG_PrepareForDriverUnload = 0x34
+PPSMC_MSG_ReadThrottlerLimit = 0x35
+PPSMC_MSG_QueryValidMcaCount = 0x36
+PPSMC_MSG_McaBankDumpDW = 0x37
+PPSMC_MSG_GetCTFLimit = 0x38
+PPSMC_MSG_ClearMcaOnRead = 0x39
+PPSMC_MSG_QueryValidMcaCeCount = 0x3A
+PPSMC_MSG_McaBankCeDumpDW = 0x3B
+PPSMC_MSG_SelectPLPDMode = 0x40
+PPSMC_MSG_PmLogReadSample = 0x41
+PPSMC_MSG_PmLogGetTableVersion = 0x42
+PPSMC_MSG_RmaDueToBadPageThreshold = 0x43
+PPSMC_MSG_SetThrottlingPolicy = 0x44
+PPSMC_MSG_SetPhaseDetectCSBWThreshold = 0x45
+PPSMC_MSG_SetPhaseDetectFreqHigh = 0x46
+PPSMC_MSG_SetPhaseDetectFreqLow = 0x47
+PPSMC_MSG_SetPhaseDetectDownHysterisis = 0x48
+PPSMC_MSG_SetPhaseDetectAlphaX1e6 = 0x49
+PPSMC_MSG_SetPhaseDetectOnOff = 0x4A
+PPSMC_MSG_GetPhaseDetectResidency = 0x4B
+PPSMC_MSG_UpdatePccWaitDecMaxStr = 0x4C
+PPSMC_MSG_ResetSDMA = 0x4D
+PPSMC_MSG_GetRasTableVersion = 0x4E
+PPSMC_MSG_GetBadPageCount = 0x50
+PPSMC_MSG_GetBadPageMcaAddress = 0x51
+PPSMC_MSG_SetTimestamp = 0x53
+PPSMC_MSG_SetTimestampHi = 0x54
+PPSMC_MSG_GetTimestamp = 0x55
+PPSMC_MSG_GetBadPageIpIdLoHi = 0x57
+PPSMC_MSG_EraseRasTable = 0x58
+PPSMC_MSG_GetStaticMetricsTable = 0x59
+PPSMC_MSG_ResetVfArbitersByIndex = 0x5A
+PPSMC_MSG_GetSystemMetricsTable = 0x5C
+PPSMC_MSG_GetSystemMetricsVersion = 0x5D
+PPSMC_MSG_ResetVCN = 0x5E
+PPSMC_MSG_SetFastPptLimit = 0x5F
+PPSMC_MSG_GetFastPptLimit = 0x60
+PPSMC_Message_Count = 0x61
+PPSMC_RESET_TYPE_DRIVER_MODE_1_RESET = 0x1
+PPSMC_RESET_TYPE_DRIVER_MODE_2_RESET = 0x2
+PPSMC_RESET_TYPE_DRIVER_MODE_3_RESET = 0x3
+PPSMC_THROTTLING_LIMIT_TYPE_SOCKET = 0x1
+PPSMC_THROTTLING_LIMIT_TYPE_HBM = 0x2
+PPSMC_AID_THM_TYPE = 0x1
+PPSMC_CCD_THM_TYPE = 0x2
+PPSMC_XCD_THM_TYPE = 0x3
+PPSMC_HBM_THM_TYPE = 0x4
+PPSMC_PLPD_MODE_DEFAULT = 0x1
+PPSMC_PLPD_MODE_OPTIMIZED = 0x2
+NUM_VCLK_DPM_LEVELS = 4
+NUM_DCLK_DPM_LEVELS = 4
+NUM_SOCCLK_DPM_LEVELS = 4
+NUM_LCLK_DPM_LEVELS = 4
+NUM_UCLK_DPM_LEVELS = 4
+NUM_FCLK_DPM_LEVELS = 4
+NUM_XGMI_DPM_LEVELS = 2
+NUM_CXL_BITRATES = 4
+NUM_PCIE_BITRATES = 4
+NUM_XGMI_BITRATES = 4
+NUM_XGMI_WIDTHS = 3
+NUM_TDP_GROUPS = 4
+NUM_SOC_P2S_TABLES = 6
+NUM_GFX_P2S_TABLES = 8
+NUM_PSM_DIDT_THRESHOLDS = 3
+NUM_XVMIN_VMIN_THRESHOLDS = 3
+PRODUCT_MODEL_NUMBER_LEN = 20
+PRODUCT_NAME_LEN = 64
+PRODUCT_SERIAL_LEN = 20
+PRODUCT_MANUFACTURER_NAME_LEN = 32
+PRODUCT_FRU_ID_LEN = 32
+SMU_METRICS_TABLE_VERSION = 0x15
+SMU_SYSTEM_METRICS_TABLE_VERSION = 0x1
+SMU_VF_METRICS_TABLE_MASK = (1 << 31)
+SMU_VF_METRICS_TABLE_VERSION = (0x6 | SMU_VF_METRICS_TABLE_MASK)
+SMU13_0_6_DRIVER_IF_VERSION = 0x08042024
+NUM_I2C_CONTROLLERS = 8
+I2C_CONTROLLER_ENABLED = 1
+I2C_CONTROLLER_DISABLED = 0
+MAX_SW_I2C_COMMANDS = 24
+CMDCONFIG_STOP_BIT = 0
+CMDCONFIG_RESTART_BIT = 1
+CMDCONFIG_READWRITE_BIT = 2
+CMDCONFIG_STOP_MASK = (1 << CMDCONFIG_STOP_BIT)
+CMDCONFIG_RESTART_MASK = (1 << CMDCONFIG_RESTART_BIT)
+CMDCONFIG_READWRITE_MASK = (1 << CMDCONFIG_READWRITE_BIT)
+IH_INTERRUPT_ID_TO_DRIVER = 0xFE
+IH_INTERRUPT_CONTEXT_ID_THERMAL_THROTTLING = 0x7
+THROTTLER_PROCHOT_BIT = 0
+THROTTLER_PPT_BIT = 1
+THROTTLER_THERMAL_SOCKET_BIT = 2
+THROTTLER_THERMAL_VR_BIT = 3
+THROTTLER_THERMAL_HBM_BIT = 4
+ClearMcaOnRead_UE_FLAG_MASK = 0x1
+ClearMcaOnRead_CE_POLL_MASK = 0x2
+SMU_THERMAL_MINIMUM_ALERT_TEMP = 0
+SMU_THERMAL_MAXIMUM_ALERT_TEMP = 255
+SMU_TEMPERATURE_UNITS_PER_CENTIGRADES = 1000
+SMU_FW_NAME_LEN = 0x24
+SMU_DPM_USER_PROFILE_RESTORE = (1 << 0)
+SMU_CUSTOM_FAN_SPEED_RPM = (1 << 1)
+SMU_CUSTOM_FAN_SPEED_PWM = (1 << 2)
+SMU_THROTTLER_PPT0_BIT = 0
+SMU_THROTTLER_PPT1_BIT = 1
+SMU_THROTTLER_PPT2_BIT = 2
+SMU_THROTTLER_PPT3_BIT = 3
+SMU_THROTTLER_SPL_BIT = 4
+SMU_THROTTLER_FPPT_BIT = 5
+SMU_THROTTLER_SPPT_BIT = 6
+SMU_THROTTLER_SPPT_APU_BIT = 7
+SMU_THROTTLER_TDC_GFX_BIT = 16
+SMU_THROTTLER_TDC_SOC_BIT = 17
+SMU_THROTTLER_TDC_MEM_BIT = 18
+SMU_THROTTLER_TDC_VDD_BIT = 19
+SMU_THROTTLER_TDC_CVIP_BIT = 20
+SMU_THROTTLER_EDC_CPU_BIT = 21
+SMU_THROTTLER_EDC_GFX_BIT = 22
+SMU_THROTTLER_APCC_BIT = 23
+SMU_THROTTLER_TEMP_GPU_BIT = 32
+SMU_THROTTLER_TEMP_CORE_BIT = 33
+SMU_THROTTLER_TEMP_MEM_BIT = 34
+SMU_THROTTLER_TEMP_EDGE_BIT = 35
+SMU_THROTTLER_TEMP_HOTSPOT_BIT = 36
+SMU_THROTTLER_TEMP_SOC_BIT = 37
+SMU_THROTTLER_TEMP_VR_GFX_BIT = 38
+SMU_THROTTLER_TEMP_VR_SOC_BIT = 39
+SMU_THROTTLER_TEMP_VR_MEM0_BIT = 40
+SMU_THROTTLER_TEMP_VR_MEM1_BIT = 41
+SMU_THROTTLER_TEMP_LIQUID0_BIT = 42
+SMU_THROTTLER_TEMP_LIQUID1_BIT = 43
+SMU_THROTTLER_VRHOT0_BIT = 44
+SMU_THROTTLER_VRHOT1_BIT = 45
+SMU_THROTTLER_PROCHOT_CPU_BIT = 46
+SMU_THROTTLER_PROCHOT_GFX_BIT = 47
+SMU_THROTTLER_PPM_BIT = 56
+SMU_THROTTLER_FIT_BIT = 57

tinygrad/runtime/autogen/am/smu_v13_0_6.py

@@ -341,7 +341,7 @@ class struct_smu_state_memory_block(c.Struct):
   dll_off: bool
   m3arb: int
   unused: c.Array[ctypes.c_ubyte, Literal[3]]
-struct_smu_state_memory_block.register_fields([('dll_off', ctypes.c_bool, 0), ('m3arb', ctypes.c_ubyte, 1), ('unused', c.Array[ctypes.c_ubyte, Literal[3]], 2)])
+struct_smu_state_memory_block.register_fields([('dll_off', ctypes.c_bool, 0), ('m3arb', uint8_t, 1), ('unused', c.Array[uint8_t, Literal[3]], 2)])
 @c.record
 class struct_smu_state_software_algorithm_block(c.Struct):
   SIZE = 2
@@ -369,13 +369,13 @@ class struct_smu_state_validation_block(c.Struct):
   single_display_only: bool
   disallow_on_dc: bool
   supported_power_levels: int
-struct_smu_state_validation_block.register_fields([('single_display_only', ctypes.c_bool, 0), ('disallow_on_dc', ctypes.c_bool, 1), ('supported_power_levels', ctypes.c_ubyte, 2)])
+struct_smu_state_validation_block.register_fields([('single_display_only', ctypes.c_bool, 0), ('disallow_on_dc', ctypes.c_bool, 1), ('supported_power_levels', uint8_t, 2)])
 @c.record
 class struct_smu_uvd_clocks(c.Struct):
   SIZE = 8
   vclk: int
   dclk: int
-struct_smu_uvd_clocks.register_fields([('vclk', ctypes.c_uint32, 0), ('dclk', ctypes.c_uint32, 4)])
+struct_smu_uvd_clocks.register_fields([('vclk', uint32_t, 0), ('dclk', uint32_t, 4)])
 enum_smu_power_src_type: dict[int, str] = {(SMU_POWER_SOURCE_AC:=0): 'SMU_POWER_SOURCE_AC', (SMU_POWER_SOURCE_DC:=1): 'SMU_POWER_SOURCE_DC', (SMU_POWER_SOURCE_COUNT:=2): 'SMU_POWER_SOURCE_COUNT'}
 enum_smu_ppt_limit_type: dict[int, str] = {(SMU_DEFAULT_PPT_LIMIT:=0): 'SMU_DEFAULT_PPT_LIMIT', (SMU_FAST_PPT_LIMIT:=1): 'SMU_FAST_PPT_LIMIT'}
 enum_smu_ppt_limit_level: dict[int, str] = {(SMU_PPT_LIMIT_MIN:=-1): 'SMU_PPT_LIMIT_MIN', (SMU_PPT_LIMIT_CURRENT:=0): 'SMU_PPT_LIMIT_CURRENT', (SMU_PPT_LIMIT_DEFAULT:=1): 'SMU_PPT_LIMIT_DEFAULT', (SMU_PPT_LIMIT_MAX:=2): 'SMU_PPT_LIMIT_MAX'}
@@ -392,7 +392,7 @@ class struct_smu_user_dpm_profile(c.Struct):
   user_od: int
   clk_mask: c.Array[ctypes.c_uint32, Literal[28]]
   clk_dependency: int
-struct_smu_user_dpm_profile.register_fields([('fan_mode', ctypes.c_uint32, 0), ('power_limit', ctypes.c_uint32, 4), ('fan_speed_pwm', ctypes.c_uint32, 8), ('fan_speed_rpm', ctypes.c_uint32, 12), ('flags', ctypes.c_uint32, 16), ('user_od', ctypes.c_uint32, 20), ('clk_mask', c.Array[ctypes.c_uint32, Literal[28]], 24), ('clk_dependency', ctypes.c_uint32, 136)])
+struct_smu_user_dpm_profile.register_fields([('fan_mode', uint32_t, 0), ('power_limit', uint32_t, 4), ('fan_speed_pwm', uint32_t, 8), ('fan_speed_rpm', uint32_t, 12), ('flags', uint32_t, 16), ('user_od', uint32_t, 20), ('clk_mask', c.Array[uint32_t, Literal[28]], 24), ('clk_dependency', uint32_t, 136)])
 @c.record
 class struct_smu_table(c.Struct):
   SIZE = 48
@@ -404,7 +404,7 @@ class struct_smu_table(c.Struct):
   bo: c.POINTER[struct_amdgpu_bo]
   version: int
 class struct_amdgpu_bo(c.Struct): pass
-struct_smu_table.register_fields([('size', ctypes.c_uint64, 0), ('align', ctypes.c_uint32, 8), ('domain', ctypes.c_ubyte, 12), ('mc_address', ctypes.c_uint64, 16), ('cpu_addr', ctypes.c_void_p, 24), ('bo', c.POINTER[struct_amdgpu_bo], 32), ('version', ctypes.c_uint32, 40)])
+struct_smu_table.register_fields([('size', uint64_t, 0), ('align', uint32_t, 8), ('domain', uint8_t, 12), ('mc_address', uint64_t, 16), ('cpu_addr', ctypes.c_void_p, 24), ('bo', c.POINTER[struct_amdgpu_bo], 32), ('version', uint32_t, 40)])
 enum_smu_perf_level_designation: dict[int, str] = {(PERF_LEVEL_ACTIVITY:=0): 'PERF_LEVEL_ACTIVITY', (PERF_LEVEL_POWER_CONTAINMENT:=1): 'PERF_LEVEL_POWER_CONTAINMENT'}
 @c.record
 class struct_smu_performance_level(c.Struct):
@@ -415,7 +415,7 @@ class struct_smu_performance_level(c.Struct):
   vddci: int
   non_local_mem_freq: int
   non_local_mem_width: int
-struct_smu_performance_level.register_fields([('core_clock', ctypes.c_uint32, 0), ('memory_clock', ctypes.c_uint32, 4), ('vddc', ctypes.c_uint32, 8), ('vddci', ctypes.c_uint32, 12), ('non_local_mem_freq', ctypes.c_uint32, 16), ('non_local_mem_width', ctypes.c_uint32, 20)])
+struct_smu_performance_level.register_fields([('core_clock', uint32_t, 0), ('memory_clock', uint32_t, 4), ('vddc', uint32_t, 8), ('vddci', uint32_t, 12), ('non_local_mem_freq', uint32_t, 16), ('non_local_mem_width', uint32_t, 20)])
 @c.record
 class struct_smu_clock_info(c.Struct):
   SIZE = 24
@@ -425,7 +425,7 @@ class struct_smu_clock_info(c.Struct):
   max_eng_clk: int
   min_bus_bandwidth: int
   max_bus_bandwidth: int
-struct_smu_clock_info.register_fields([('min_mem_clk', ctypes.c_uint32, 0), ('max_mem_clk', ctypes.c_uint32, 4), ('min_eng_clk', ctypes.c_uint32, 8), ('max_eng_clk', ctypes.c_uint32, 12), ('min_bus_bandwidth', ctypes.c_uint32, 16), ('max_bus_bandwidth', ctypes.c_uint32, 20)])
+struct_smu_clock_info.register_fields([('min_mem_clk', uint32_t, 0), ('max_mem_clk', uint32_t, 4), ('min_eng_clk', uint32_t, 8), ('max_eng_clk', uint32_t, 12), ('min_bus_bandwidth', uint32_t, 16), ('max_bus_bandwidth', uint32_t, 20)])
 @c.record
 class struct_smu_bios_boot_up_values(c.Struct):
   SIZE = 68
@@ -448,171 +448,170 @@ class struct_smu_bios_boot_up_values(c.Struct):
   fclk: int
   lclk: int
   firmware_caps: int
-struct_smu_bios_boot_up_values.register_fields([('revision', ctypes.c_uint32, 0), ('gfxclk', ctypes.c_uint32, 4), ('uclk', ctypes.c_uint32, 8), ('socclk', ctypes.c_uint32, 12), ('dcefclk', ctypes.c_uint32, 16), ('eclk', ctypes.c_uint32, 20), ('vclk', ctypes.c_uint32, 24), ('dclk', ctypes.c_uint32, 28), ('vddc', ctypes.c_uint16, 32), ('vddci', ctypes.c_uint16, 34), ('mvddc', ctypes.c_uint16, 36), ('vdd_gfx', ctypes.c_uint16, 38), ('cooling_id', ctypes.c_ubyte, 40), ('pp_table_id', ctypes.c_uint32, 44), ('format_revision', ctypes.c_uint32, 48), ('content_revision', ctypes.c_uint32, 52), ('fclk', ctypes.c_uint32, 56), ('lclk', ctypes.c_uint32, 60), ('firmware_caps', ctypes.c_uint32, 64)])
+struct_smu_bios_boot_up_values.register_fields([('revision', uint32_t, 0), ('gfxclk', uint32_t, 4), ('uclk', uint32_t, 8), ('socclk', uint32_t, 12), ('dcefclk', uint32_t, 16), ('eclk', uint32_t, 20), ('vclk', uint32_t, 24), ('dclk', uint32_t, 28), ('vddc', uint16_t, 32), ('vddci', uint16_t, 34), ('mvddc', uint16_t, 36), ('vdd_gfx', uint16_t, 38), ('cooling_id', uint8_t, 40), ('pp_table_id', uint32_t, 44), ('format_revision', uint32_t, 48), ('content_revision', uint32_t, 52), ('fclk', uint32_t, 56), ('lclk', uint32_t, 60), ('firmware_caps', uint32_t, 64)])
 enum_smu_table_id: dict[int, str] = {(SMU_TABLE_PPTABLE:=0): 'SMU_TABLE_PPTABLE', (SMU_TABLE_WATERMARKS:=1): 'SMU_TABLE_WATERMARKS', (SMU_TABLE_CUSTOM_DPM:=2): 'SMU_TABLE_CUSTOM_DPM', (SMU_TABLE_DPMCLOCKS:=3): 'SMU_TABLE_DPMCLOCKS', (SMU_TABLE_AVFS:=4): 'SMU_TABLE_AVFS', (SMU_TABLE_AVFS_PSM_DEBUG:=5): 'SMU_TABLE_AVFS_PSM_DEBUG', (SMU_TABLE_AVFS_FUSE_OVERRIDE:=6): 'SMU_TABLE_AVFS_FUSE_OVERRIDE', (SMU_TABLE_PMSTATUSLOG:=7): 'SMU_TABLE_PMSTATUSLOG', (SMU_TABLE_SMU_METRICS:=8): 'SMU_TABLE_SMU_METRICS', (SMU_TABLE_DRIVER_SMU_CONFIG:=9): 'SMU_TABLE_DRIVER_SMU_CONFIG', (SMU_TABLE_ACTIVITY_MONITOR_COEFF:=10): 'SMU_TABLE_ACTIVITY_MONITOR_COEFF', (SMU_TABLE_OVERDRIVE:=11): 'SMU_TABLE_OVERDRIVE', (SMU_TABLE_I2C_COMMANDS:=12): 'SMU_TABLE_I2C_COMMANDS', (SMU_TABLE_PACE:=13): 'SMU_TABLE_PACE', (SMU_TABLE_ECCINFO:=14): 'SMU_TABLE_ECCINFO', (SMU_TABLE_COMBO_PPTABLE:=15): 'SMU_TABLE_COMBO_PPTABLE', (SMU_TABLE_WIFIBAND:=16): 'SMU_TABLE_WIFIBAND', (SMU_TABLE_COUNT:=17): 'SMU_TABLE_COUNT'}
-PPSMC_Result_OK = 0x1 # type: ignore
-PPSMC_Result_Failed = 0xFF # type: ignore
-PPSMC_Result_UnknownCmd = 0xFE # type: ignore
-PPSMC_Result_CmdRejectedPrereq = 0xFD # type: ignore
-PPSMC_Result_CmdRejectedBusy = 0xFC # type: ignore
-PPSMC_MSG_TestMessage = 0x1 # type: ignore
-PPSMC_MSG_GetSmuVersion = 0x2 # type: ignore
-PPSMC_MSG_GfxDriverReset = 0x3 # type: ignore
-PPSMC_MSG_GetDriverIfVersion = 0x4 # type: ignore
-PPSMC_MSG_EnableAllSmuFeatures = 0x5 # type: ignore
-PPSMC_MSG_DisableAllSmuFeatures = 0x6 # type: ignore
-PPSMC_MSG_RequestI2cTransaction = 0x7 # type: ignore
-PPSMC_MSG_GetMetricsVersion = 0x8 # type: ignore
-PPSMC_MSG_GetMetricsTable = 0x9 # type: ignore
-PPSMC_MSG_GetEccInfoTable = 0xA # type: ignore
-PPSMC_MSG_GetEnabledSmuFeaturesLow = 0xB # type: ignore
-PPSMC_MSG_GetEnabledSmuFeaturesHigh = 0xC # type: ignore
-PPSMC_MSG_SetDriverDramAddrHigh = 0xD # type: ignore
-PPSMC_MSG_SetDriverDramAddrLow = 0xE # type: ignore
-PPSMC_MSG_SetToolsDramAddrHigh = 0xF # type: ignore
-PPSMC_MSG_SetToolsDramAddrLow = 0x10 # type: ignore
-PPSMC_MSG_SetSystemVirtualDramAddrHigh = 0x11 # type: ignore
-PPSMC_MSG_SetSystemVirtualDramAddrLow = 0x12 # type: ignore
-PPSMC_MSG_SetSoftMinByFreq = 0x13 # type: ignore
-PPSMC_MSG_SetSoftMaxByFreq = 0x14 # type: ignore
-PPSMC_MSG_GetMinDpmFreq = 0x15 # type: ignore
-PPSMC_MSG_GetMaxDpmFreq = 0x16 # type: ignore
-PPSMC_MSG_GetDpmFreqByIndex = 0x17 # type: ignore
-PPSMC_MSG_SetPptLimit = 0x18 # type: ignore
-PPSMC_MSG_GetPptLimit = 0x19 # type: ignore
-PPSMC_MSG_DramLogSetDramAddrHigh = 0x1A # type: ignore
-PPSMC_MSG_DramLogSetDramAddrLow = 0x1B # type: ignore
-PPSMC_MSG_DramLogSetDramSize = 0x1C # type: ignore
-PPSMC_MSG_GetDebugData = 0x1D # type: ignore
-PPSMC_MSG_HeavySBR = 0x1E # type: ignore
-PPSMC_MSG_SetNumBadHbmPagesRetired = 0x1F # type: ignore
-PPSMC_MSG_DFCstateControl = 0x20 # type: ignore
-PPSMC_MSG_GetGmiPwrDnHyst = 0x21 # type: ignore
-PPSMC_MSG_SetGmiPwrDnHyst = 0x22 # type: ignore
-PPSMC_MSG_GmiPwrDnControl = 0x23 # type: ignore
-PPSMC_MSG_EnterGfxoff = 0x24 # type: ignore
-PPSMC_MSG_ExitGfxoff = 0x25 # type: ignore
-PPSMC_MSG_EnableDeterminism = 0x26 # type: ignore
-PPSMC_MSG_DisableDeterminism = 0x27 # type: ignore
-PPSMC_MSG_DumpSTBtoDram = 0x28 # type: ignore
-PPSMC_MSG_STBtoDramLogSetDramAddrHigh = 0x29 # type: ignore
-PPSMC_MSG_STBtoDramLogSetDramAddrLow = 0x2A # type: ignore
-PPSMC_MSG_STBtoDramLogSetDramSize = 0x2B # type: ignore
-PPSMC_MSG_SetSystemVirtualSTBtoDramAddrHigh = 0x2C # type: ignore
-PPSMC_MSG_SetSystemVirtualSTBtoDramAddrLow = 0x2D # type: ignore
-PPSMC_MSG_GfxDriverResetRecovery = 0x2E # type: ignore
-PPSMC_MSG_TriggerVFFLR = 0x2F # type: ignore
-PPSMC_MSG_SetSoftMinGfxClk = 0x30 # type: ignore
-PPSMC_MSG_SetSoftMaxGfxClk = 0x31 # type: ignore
-PPSMC_MSG_GetMinGfxDpmFreq = 0x32 # type: ignore
-PPSMC_MSG_GetMaxGfxDpmFreq = 0x33 # type: ignore
-PPSMC_MSG_PrepareForDriverUnload = 0x34 # type: ignore
-PPSMC_MSG_ReadThrottlerLimit = 0x35 # type: ignore
-PPSMC_MSG_QueryValidMcaCount = 0x36 # type: ignore
-PPSMC_MSG_McaBankDumpDW = 0x37 # type: ignore
-PPSMC_MSG_GetCTFLimit = 0x38 # type: ignore
-PPSMC_MSG_ClearMcaOnRead = 0x39 # type: ignore
-PPSMC_MSG_QueryValidMcaCeCount = 0x3A # type: ignore
-PPSMC_MSG_McaBankCeDumpDW = 0x3B # type: ignore
-PPSMC_MSG_SelectPLPDMode = 0x40 # type: ignore
-PPSMC_MSG_RmaDueToBadPageThreshold = 0x43 # type: ignore
-PPSMC_MSG_SetThrottlingPolicy = 0x44 # type: ignore
-PPSMC_MSG_SetPhsDetWRbwThreshold = 0x45 # type: ignore
-PPSMC_MSG_SetPhsDetWRbwFreqHigh = 0x46 # type: ignore
-PPSMC_MSG_SetPhsDetWRbwFreqLow = 0x47 # type: ignore
-PPSMC_MSG_SetPhsDetWRbwHystDown = 0x48 # type: ignore
-PPSMC_MSG_SetPhsDetWRbwAlpha = 0x49 # type: ignore
-PPSMC_MSG_SetPhsDetOnOff = 0x4A # type: ignore
-PPSMC_MSG_GetPhsDetResidency = 0x4B # type: ignore
-PPSMC_MSG_ResetSDMA = 0x4D # type: ignore
-PPSMC_MSG_GetStaticMetricsTable = 0x59 # type: ignore
-PPSMC_MSG_ResetVCN = 0x5B # type: ignore
-PPSMC_Message_Count = 0x5C # type: ignore
-PPSMC_RESET_TYPE_DRIVER_MODE_1_RESET = 0x1 # type: ignore
-PPSMC_RESET_TYPE_DRIVER_MODE_2_RESET = 0x2 # type: ignore
-PPSMC_RESET_TYPE_DRIVER_MODE_3_RESET = 0x3 # type: ignore
-PPSMC_THROTTLING_LIMIT_TYPE_SOCKET = 0x1 # type: ignore
-PPSMC_THROTTLING_LIMIT_TYPE_HBM = 0x2 # type: ignore
-PPSMC_AID_THM_TYPE = 0x1 # type: ignore
-PPSMC_CCD_THM_TYPE = 0x2 # type: ignore
-PPSMC_XCD_THM_TYPE = 0x3 # type: ignore
-PPSMC_HBM_THM_TYPE = 0x4 # type: ignore
-PPSMC_PLPD_MODE_DEFAULT = 0x1 # type: ignore
-PPSMC_PLPD_MODE_OPTIMIZED = 0x2 # type: ignore
-NUM_VCLK_DPM_LEVELS = 4 # type: ignore
-NUM_DCLK_DPM_LEVELS = 4 # type: ignore
-NUM_SOCCLK_DPM_LEVELS = 4 # type: ignore
-NUM_LCLK_DPM_LEVELS = 4 # type: ignore
-NUM_UCLK_DPM_LEVELS = 4 # type: ignore
-NUM_FCLK_DPM_LEVELS = 4 # type: ignore
-NUM_XGMI_DPM_LEVELS = 2 # type: ignore
-NUM_CXL_BITRATES = 4 # type: ignore
-NUM_PCIE_BITRATES = 4 # type: ignore
-NUM_XGMI_BITRATES = 4 # type: ignore
-NUM_XGMI_WIDTHS = 3 # type: ignore
-NUM_SOC_P2S_TABLES = 3 # type: ignore
-NUM_TDP_GROUPS = 4 # type: ignore
-SMU_METRICS_TABLE_VERSION = 0x11 # type: ignore
-SMU_VF_METRICS_TABLE_VERSION = 0x5 # type: ignore
-SMU13_0_6_DRIVER_IF_VERSION = 0x08042024 # type: ignore
-NUM_I2C_CONTROLLERS = 8 # type: ignore
-I2C_CONTROLLER_ENABLED = 1 # type: ignore
-I2C_CONTROLLER_DISABLED = 0 # type: ignore
-MAX_SW_I2C_COMMANDS = 24 # type: ignore
-CMDCONFIG_STOP_BIT = 0 # type: ignore
-CMDCONFIG_RESTART_BIT = 1 # type: ignore
-CMDCONFIG_READWRITE_BIT = 2 # type: ignore
-CMDCONFIG_STOP_MASK = (1 << CMDCONFIG_STOP_BIT) # type: ignore
-CMDCONFIG_RESTART_MASK = (1 << CMDCONFIG_RESTART_BIT) # type: ignore
-CMDCONFIG_READWRITE_MASK = (1 << CMDCONFIG_READWRITE_BIT) # type: ignore
-IH_INTERRUPT_ID_TO_DRIVER = 0xFE # type: ignore
-IH_INTERRUPT_CONTEXT_ID_THERMAL_THROTTLING = 0x7 # type: ignore
-THROTTLER_PROCHOT_BIT = 0 # type: ignore
-THROTTLER_PPT_BIT = 1 # type: ignore
-THROTTLER_THERMAL_SOCKET_BIT = 2 # type: ignore
-THROTTLER_THERMAL_VR_BIT = 3 # type: ignore
-THROTTLER_THERMAL_HBM_BIT = 4 # type: ignore
-ClearMcaOnRead_UE_FLAG_MASK = 0x1 # type: ignore
-ClearMcaOnRead_CE_POLL_MASK = 0x2 # type: ignore
-int32_t = int # type: ignore
-SMU_THERMAL_MINIMUM_ALERT_TEMP = 0 # type: ignore
-SMU_THERMAL_MAXIMUM_ALERT_TEMP = 255 # type: ignore
-SMU_TEMPERATURE_UNITS_PER_CENTIGRADES = 1000 # type: ignore
-SMU_FW_NAME_LEN = 0x24 # type: ignore
-SMU_DPM_USER_PROFILE_RESTORE = (1 << 0) # type: ignore
-SMU_CUSTOM_FAN_SPEED_RPM = (1 << 1) # type: ignore
-SMU_CUSTOM_FAN_SPEED_PWM = (1 << 2) # type: ignore
-SMU_THROTTLER_PPT0_BIT = 0 # type: ignore
-SMU_THROTTLER_PPT1_BIT = 1 # type: ignore
-SMU_THROTTLER_PPT2_BIT = 2 # type: ignore
-SMU_THROTTLER_PPT3_BIT = 3 # type: ignore
-SMU_THROTTLER_SPL_BIT = 4 # type: ignore
-SMU_THROTTLER_FPPT_BIT = 5 # type: ignore
-SMU_THROTTLER_SPPT_BIT = 6 # type: ignore
-SMU_THROTTLER_SPPT_APU_BIT = 7 # type: ignore
-SMU_THROTTLER_TDC_GFX_BIT = 16 # type: ignore
-SMU_THROTTLER_TDC_SOC_BIT = 17 # type: ignore
-SMU_THROTTLER_TDC_MEM_BIT = 18 # type: ignore
-SMU_THROTTLER_TDC_VDD_BIT = 19 # type: ignore
-SMU_THROTTLER_TDC_CVIP_BIT = 20 # type: ignore
-SMU_THROTTLER_EDC_CPU_BIT = 21 # type: ignore
-SMU_THROTTLER_EDC_GFX_BIT = 22 # type: ignore
-SMU_THROTTLER_APCC_BIT = 23 # type: ignore
-SMU_THROTTLER_TEMP_GPU_BIT = 32 # type: ignore
-SMU_THROTTLER_TEMP_CORE_BIT = 33 # type: ignore
-SMU_THROTTLER_TEMP_MEM_BIT = 34 # type: ignore
-SMU_THROTTLER_TEMP_EDGE_BIT = 35 # type: ignore
-SMU_THROTTLER_TEMP_HOTSPOT_BIT = 36 # type: ignore
-SMU_THROTTLER_TEMP_SOC_BIT = 37 # type: ignore
-SMU_THROTTLER_TEMP_VR_GFX_BIT = 38 # type: ignore
-SMU_THROTTLER_TEMP_VR_SOC_BIT = 39 # type: ignore
-SMU_THROTTLER_TEMP_VR_MEM0_BIT = 40 # type: ignore
-SMU_THROTTLER_TEMP_VR_MEM1_BIT = 41 # type: ignore
-SMU_THROTTLER_TEMP_LIQUID0_BIT = 42 # type: ignore
-SMU_THROTTLER_TEMP_LIQUID1_BIT = 43 # type: ignore
-SMU_THROTTLER_VRHOT0_BIT = 44 # type: ignore
-SMU_THROTTLER_VRHOT1_BIT = 45 # type: ignore
-SMU_THROTTLER_PROCHOT_CPU_BIT = 46 # type: ignore
-SMU_THROTTLER_PROCHOT_GFX_BIT = 47 # type: ignore
-SMU_THROTTLER_PPM_BIT = 56 # type: ignore
-SMU_THROTTLER_FIT_BIT = 57 # type: ignore
+PPSMC_Result_OK = 0x1
+PPSMC_Result_Failed = 0xFF
+PPSMC_Result_UnknownCmd = 0xFE
+PPSMC_Result_CmdRejectedPrereq = 0xFD
+PPSMC_Result_CmdRejectedBusy = 0xFC
+PPSMC_MSG_TestMessage = 0x1
+PPSMC_MSG_GetSmuVersion = 0x2
+PPSMC_MSG_GfxDriverReset = 0x3
+PPSMC_MSG_GetDriverIfVersion = 0x4
+PPSMC_MSG_EnableAllSmuFeatures = 0x5
+PPSMC_MSG_DisableAllSmuFeatures = 0x6
+PPSMC_MSG_RequestI2cTransaction = 0x7
+PPSMC_MSG_GetMetricsVersion = 0x8
+PPSMC_MSG_GetMetricsTable = 0x9
+PPSMC_MSG_GetEccInfoTable = 0xA
+PPSMC_MSG_GetEnabledSmuFeaturesLow = 0xB
+PPSMC_MSG_GetEnabledSmuFeaturesHigh = 0xC
+PPSMC_MSG_SetDriverDramAddrHigh = 0xD
+PPSMC_MSG_SetDriverDramAddrLow = 0xE
+PPSMC_MSG_SetToolsDramAddrHigh = 0xF
+PPSMC_MSG_SetToolsDramAddrLow = 0x10
+PPSMC_MSG_SetSystemVirtualDramAddrHigh = 0x11
+PPSMC_MSG_SetSystemVirtualDramAddrLow = 0x12
+PPSMC_MSG_SetSoftMinByFreq = 0x13
+PPSMC_MSG_SetSoftMaxByFreq = 0x14
+PPSMC_MSG_GetMinDpmFreq = 0x15
+PPSMC_MSG_GetMaxDpmFreq = 0x16
+PPSMC_MSG_GetDpmFreqByIndex = 0x17
+PPSMC_MSG_SetPptLimit = 0x18
+PPSMC_MSG_GetPptLimit = 0x19
+PPSMC_MSG_DramLogSetDramAddrHigh = 0x1A
+PPSMC_MSG_DramLogSetDramAddrLow = 0x1B
+PPSMC_MSG_DramLogSetDramSize = 0x1C
+PPSMC_MSG_GetDebugData = 0x1D
+PPSMC_MSG_HeavySBR = 0x1E
+PPSMC_MSG_SetNumBadHbmPagesRetired = 0x1F
+PPSMC_MSG_DFCstateControl = 0x20
+PPSMC_MSG_GetGmiPwrDnHyst = 0x21
+PPSMC_MSG_SetGmiPwrDnHyst = 0x22
+PPSMC_MSG_GmiPwrDnControl = 0x23
+PPSMC_MSG_EnterGfxoff = 0x24
+PPSMC_MSG_ExitGfxoff = 0x25
+PPSMC_MSG_EnableDeterminism = 0x26
+PPSMC_MSG_DisableDeterminism = 0x27
+PPSMC_MSG_DumpSTBtoDram = 0x28
+PPSMC_MSG_STBtoDramLogSetDramAddrHigh = 0x29
+PPSMC_MSG_STBtoDramLogSetDramAddrLow = 0x2A
+PPSMC_MSG_STBtoDramLogSetDramSize = 0x2B
+PPSMC_MSG_SetSystemVirtualSTBtoDramAddrHigh = 0x2C
+PPSMC_MSG_SetSystemVirtualSTBtoDramAddrLow = 0x2D
+PPSMC_MSG_GfxDriverResetRecovery = 0x2E
+PPSMC_MSG_TriggerVFFLR = 0x2F
+PPSMC_MSG_SetSoftMinGfxClk = 0x30
+PPSMC_MSG_SetSoftMaxGfxClk = 0x31
+PPSMC_MSG_GetMinGfxDpmFreq = 0x32
+PPSMC_MSG_GetMaxGfxDpmFreq = 0x33
+PPSMC_MSG_PrepareForDriverUnload = 0x34
+PPSMC_MSG_ReadThrottlerLimit = 0x35
+PPSMC_MSG_QueryValidMcaCount = 0x36
+PPSMC_MSG_McaBankDumpDW = 0x37
+PPSMC_MSG_GetCTFLimit = 0x38
+PPSMC_MSG_ClearMcaOnRead = 0x39
+PPSMC_MSG_QueryValidMcaCeCount = 0x3A
+PPSMC_MSG_McaBankCeDumpDW = 0x3B
+PPSMC_MSG_SelectPLPDMode = 0x40
+PPSMC_MSG_RmaDueToBadPageThreshold = 0x43
+PPSMC_MSG_SetThrottlingPolicy = 0x44
+PPSMC_MSG_SetPhsDetWRbwThreshold = 0x45
+PPSMC_MSG_SetPhsDetWRbwFreqHigh = 0x46
+PPSMC_MSG_SetPhsDetWRbwFreqLow = 0x47
+PPSMC_MSG_SetPhsDetWRbwHystDown = 0x48
+PPSMC_MSG_SetPhsDetWRbwAlpha = 0x49
+PPSMC_MSG_SetPhsDetOnOff = 0x4A
+PPSMC_MSG_GetPhsDetResidency = 0x4B
+PPSMC_MSG_ResetSDMA = 0x4D
+PPSMC_MSG_GetStaticMetricsTable = 0x59
+PPSMC_MSG_ResetVCN = 0x5B
+PPSMC_Message_Count = 0x5C
+PPSMC_RESET_TYPE_DRIVER_MODE_1_RESET = 0x1
+PPSMC_RESET_TYPE_DRIVER_MODE_2_RESET = 0x2
+PPSMC_RESET_TYPE_DRIVER_MODE_3_RESET = 0x3
+PPSMC_THROTTLING_LIMIT_TYPE_SOCKET = 0x1
+PPSMC_THROTTLING_LIMIT_TYPE_HBM = 0x2
+PPSMC_AID_THM_TYPE = 0x1
+PPSMC_CCD_THM_TYPE = 0x2
+PPSMC_XCD_THM_TYPE = 0x3
+PPSMC_HBM_THM_TYPE = 0x4
+PPSMC_PLPD_MODE_DEFAULT = 0x1
+PPSMC_PLPD_MODE_OPTIMIZED = 0x2
+NUM_VCLK_DPM_LEVELS = 4
+NUM_DCLK_DPM_LEVELS = 4
+NUM_SOCCLK_DPM_LEVELS = 4
+NUM_LCLK_DPM_LEVELS = 4
+NUM_UCLK_DPM_LEVELS = 4
+NUM_FCLK_DPM_LEVELS = 4
+NUM_XGMI_DPM_LEVELS = 2
+NUM_CXL_BITRATES = 4
+NUM_PCIE_BITRATES = 4
+NUM_XGMI_BITRATES = 4
+NUM_XGMI_WIDTHS = 3
+NUM_SOC_P2S_TABLES = 3
+NUM_TDP_GROUPS = 4
+SMU_METRICS_TABLE_VERSION = 0x11
+SMU_VF_METRICS_TABLE_VERSION = 0x5
+SMU13_0_6_DRIVER_IF_VERSION = 0x08042024
+NUM_I2C_CONTROLLERS = 8
+I2C_CONTROLLER_ENABLED = 1
+I2C_CONTROLLER_DISABLED = 0
+MAX_SW_I2C_COMMANDS = 24
+CMDCONFIG_STOP_BIT = 0
+CMDCONFIG_RESTART_BIT = 1
+CMDCONFIG_READWRITE_BIT = 2
+CMDCONFIG_STOP_MASK = (1 << CMDCONFIG_STOP_BIT)
+CMDCONFIG_RESTART_MASK = (1 << CMDCONFIG_RESTART_BIT)
+CMDCONFIG_READWRITE_MASK = (1 << CMDCONFIG_READWRITE_BIT)
+IH_INTERRUPT_ID_TO_DRIVER = 0xFE
+IH_INTERRUPT_CONTEXT_ID_THERMAL_THROTTLING = 0x7
+THROTTLER_PROCHOT_BIT = 0
+THROTTLER_PPT_BIT = 1
+THROTTLER_THERMAL_SOCKET_BIT = 2
+THROTTLER_THERMAL_VR_BIT = 3
+THROTTLER_THERMAL_HBM_BIT = 4
+ClearMcaOnRead_UE_FLAG_MASK = 0x1
+ClearMcaOnRead_CE_POLL_MASK = 0x2
+SMU_THERMAL_MINIMUM_ALERT_TEMP = 0
+SMU_THERMAL_MAXIMUM_ALERT_TEMP = 255
+SMU_TEMPERATURE_UNITS_PER_CENTIGRADES = 1000
+SMU_FW_NAME_LEN = 0x24
+SMU_DPM_USER_PROFILE_RESTORE = (1 << 0)
+SMU_CUSTOM_FAN_SPEED_RPM = (1 << 1)
+SMU_CUSTOM_FAN_SPEED_PWM = (1 << 2)
+SMU_THROTTLER_PPT0_BIT = 0
+SMU_THROTTLER_PPT1_BIT = 1
+SMU_THROTTLER_PPT2_BIT = 2
+SMU_THROTTLER_PPT3_BIT = 3
+SMU_THROTTLER_SPL_BIT = 4
+SMU_THROTTLER_FPPT_BIT = 5
+SMU_THROTTLER_SPPT_BIT = 6
+SMU_THROTTLER_SPPT_APU_BIT = 7
+SMU_THROTTLER_TDC_GFX_BIT = 16
+SMU_THROTTLER_TDC_SOC_BIT = 17
+SMU_THROTTLER_TDC_MEM_BIT = 18
+SMU_THROTTLER_TDC_VDD_BIT = 19
+SMU_THROTTLER_TDC_CVIP_BIT = 20
+SMU_THROTTLER_EDC_CPU_BIT = 21
+SMU_THROTTLER_EDC_GFX_BIT = 22
+SMU_THROTTLER_APCC_BIT = 23
+SMU_THROTTLER_TEMP_GPU_BIT = 32
+SMU_THROTTLER_TEMP_CORE_BIT = 33
+SMU_THROTTLER_TEMP_MEM_BIT = 34
+SMU_THROTTLER_TEMP_EDGE_BIT = 35
+SMU_THROTTLER_TEMP_HOTSPOT_BIT = 36
+SMU_THROTTLER_TEMP_SOC_BIT = 37
+SMU_THROTTLER_TEMP_VR_GFX_BIT = 38
+SMU_THROTTLER_TEMP_VR_SOC_BIT = 39
+SMU_THROTTLER_TEMP_VR_MEM0_BIT = 40
+SMU_THROTTLER_TEMP_VR_MEM1_BIT = 41
+SMU_THROTTLER_TEMP_LIQUID0_BIT = 42
+SMU_THROTTLER_TEMP_LIQUID1_BIT = 43
+SMU_THROTTLER_VRHOT0_BIT = 44
+SMU_THROTTLER_VRHOT1_BIT = 45
+SMU_THROTTLER_PROCHOT_CPU_BIT = 46
+SMU_THROTTLER_PROCHOT_GFX_BIT = 47
+SMU_THROTTLER_PPM_BIT = 56
+SMU_THROTTLER_FIT_BIT = 57

tinygrad/runtime/autogen/am/smu_v14_0_2.py

@@ -1014,7 +1014,7 @@ class struct_smu_state_memory_block(c.Struct):
   dll_off: bool
   m3arb: int
   unused: c.Array[ctypes.c_ubyte, Literal[3]]
-struct_smu_state_memory_block.register_fields([('dll_off', ctypes.c_bool, 0), ('m3arb', ctypes.c_ubyte, 1), ('unused', c.Array[ctypes.c_ubyte, Literal[3]], 2)])
+struct_smu_state_memory_block.register_fields([('dll_off', ctypes.c_bool, 0), ('m3arb', uint8_t, 1), ('unused', c.Array[uint8_t, Literal[3]], 2)])
 @c.record
 class struct_smu_state_software_algorithm_block(c.Struct):
   SIZE = 2
@@ -1042,13 +1042,13 @@ class struct_smu_state_validation_block(c.Struct):
   single_display_only: bool
   disallow_on_dc: bool
   supported_power_levels: int
-struct_smu_state_validation_block.register_fields([('single_display_only', ctypes.c_bool, 0), ('disallow_on_dc', ctypes.c_bool, 1), ('supported_power_levels', ctypes.c_ubyte, 2)])
+struct_smu_state_validation_block.register_fields([('single_display_only', ctypes.c_bool, 0), ('disallow_on_dc', ctypes.c_bool, 1), ('supported_power_levels', uint8_t, 2)])
 @c.record
 class struct_smu_uvd_clocks(c.Struct):
   SIZE = 8
   vclk: int
   dclk: int
-struct_smu_uvd_clocks.register_fields([('vclk', ctypes.c_uint32, 0), ('dclk', ctypes.c_uint32, 4)])
+struct_smu_uvd_clocks.register_fields([('vclk', uint32_t, 0), ('dclk', uint32_t, 4)])
 enum_smu_power_src_type: dict[int, str] = {(SMU_POWER_SOURCE_AC:=0): 'SMU_POWER_SOURCE_AC', (SMU_POWER_SOURCE_DC:=1): 'SMU_POWER_SOURCE_DC', (SMU_POWER_SOURCE_COUNT:=2): 'SMU_POWER_SOURCE_COUNT'}
 enum_smu_ppt_limit_type: dict[int, str] = {(SMU_DEFAULT_PPT_LIMIT:=0): 'SMU_DEFAULT_PPT_LIMIT', (SMU_FAST_PPT_LIMIT:=1): 'SMU_FAST_PPT_LIMIT'}
 enum_smu_ppt_limit_level: dict[int, str] = {(SMU_PPT_LIMIT_MIN:=-1): 'SMU_PPT_LIMIT_MIN', (SMU_PPT_LIMIT_CURRENT:=0): 'SMU_PPT_LIMIT_CURRENT', (SMU_PPT_LIMIT_DEFAULT:=1): 'SMU_PPT_LIMIT_DEFAULT', (SMU_PPT_LIMIT_MAX:=2): 'SMU_PPT_LIMIT_MAX'}
@@ -1065,7 +1065,7 @@ class struct_smu_user_dpm_profile(c.Struct):
   user_od: int
   clk_mask: c.Array[ctypes.c_uint32, Literal[28]]
   clk_dependency: int
-struct_smu_user_dpm_profile.register_fields([('fan_mode', ctypes.c_uint32, 0), ('power_limit', ctypes.c_uint32, 4), ('fan_speed_pwm', ctypes.c_uint32, 8), ('fan_speed_rpm', ctypes.c_uint32, 12), ('flags', ctypes.c_uint32, 16), ('user_od', ctypes.c_uint32, 20), ('clk_mask', c.Array[ctypes.c_uint32, Literal[28]], 24), ('clk_dependency', ctypes.c_uint32, 136)])
+struct_smu_user_dpm_profile.register_fields([('fan_mode', uint32_t, 0), ('power_limit', uint32_t, 4), ('fan_speed_pwm', uint32_t, 8), ('fan_speed_rpm', uint32_t, 12), ('flags', uint32_t, 16), ('user_od', uint32_t, 20), ('clk_mask', c.Array[uint32_t, Literal[28]], 24), ('clk_dependency', uint32_t, 136)])
 @c.record
 class struct_smu_table(c.Struct):
   SIZE = 48
@@ -1077,7 +1077,7 @@ class struct_smu_table(c.Struct):
   bo: c.POINTER[struct_amdgpu_bo]
   version: int
 class struct_amdgpu_bo(c.Struct): pass
-struct_smu_table.register_fields([('size', ctypes.c_uint64, 0), ('align', ctypes.c_uint32, 8), ('domain', ctypes.c_ubyte, 12), ('mc_address', ctypes.c_uint64, 16), ('cpu_addr', ctypes.c_void_p, 24), ('bo', c.POINTER[struct_amdgpu_bo], 32), ('version', ctypes.c_uint32, 40)])
+struct_smu_table.register_fields([('size', uint64_t, 0), ('align', uint32_t, 8), ('domain', uint8_t, 12), ('mc_address', uint64_t, 16), ('cpu_addr', ctypes.c_void_p, 24), ('bo', c.POINTER[struct_amdgpu_bo], 32), ('version', uint32_t, 40)])
 enum_smu_perf_level_designation: dict[int, str] = {(PERF_LEVEL_ACTIVITY:=0): 'PERF_LEVEL_ACTIVITY', (PERF_LEVEL_POWER_CONTAINMENT:=1): 'PERF_LEVEL_POWER_CONTAINMENT'}
 @c.record
 class struct_smu_performance_level(c.Struct):
@@ -1088,7 +1088,7 @@ class struct_smu_performance_level(c.Struct):
   vddci: int
   non_local_mem_freq: int
   non_local_mem_width: int
-struct_smu_performance_level.register_fields([('core_clock', ctypes.c_uint32, 0), ('memory_clock', ctypes.c_uint32, 4), ('vddc', ctypes.c_uint32, 8), ('vddci', ctypes.c_uint32, 12), ('non_local_mem_freq', ctypes.c_uint32, 16), ('non_local_mem_width', ctypes.c_uint32, 20)])
+struct_smu_performance_level.register_fields([('core_clock', uint32_t, 0), ('memory_clock', uint32_t, 4), ('vddc', uint32_t, 8), ('vddci', uint32_t, 12), ('non_local_mem_freq', uint32_t, 16), ('non_local_mem_width', uint32_t, 20)])
 @c.record
 class struct_smu_clock_info(c.Struct):
   SIZE = 24
@@ -1098,7 +1098,7 @@ class struct_smu_clock_info(c.Struct):
   max_eng_clk: int
   min_bus_bandwidth: int
   max_bus_bandwidth: int
-struct_smu_clock_info.register_fields([('min_mem_clk', ctypes.c_uint32, 0), ('max_mem_clk', ctypes.c_uint32, 4), ('min_eng_clk', ctypes.c_uint32, 8), ('max_eng_clk', ctypes.c_uint32, 12), ('min_bus_bandwidth', ctypes.c_uint32, 16), ('max_bus_bandwidth', ctypes.c_uint32, 20)])
+struct_smu_clock_info.register_fields([('min_mem_clk', uint32_t, 0), ('max_mem_clk', uint32_t, 4), ('min_eng_clk', uint32_t, 8), ('max_eng_clk', uint32_t, 12), ('min_bus_bandwidth', uint32_t, 16), ('max_bus_bandwidth', uint32_t, 20)])
 @c.record
 class struct_smu_bios_boot_up_values(c.Struct):
   SIZE = 68
@@ -1121,444 +1121,443 @@ class struct_smu_bios_boot_up_values(c.Struct):
   fclk: int
   lclk: int
   firmware_caps: int
-struct_smu_bios_boot_up_values.register_fields([('revision', ctypes.c_uint32, 0), ('gfxclk', ctypes.c_uint32, 4), ('uclk', ctypes.c_uint32, 8), ('socclk', ctypes.c_uint32, 12), ('dcefclk', ctypes.c_uint32, 16), ('eclk', ctypes.c_uint32, 20), ('vclk', ctypes.c_uint32, 24), ('dclk', ctypes.c_uint32, 28), ('vddc', ctypes.c_uint16, 32), ('vddci', ctypes.c_uint16, 34), ('mvddc', ctypes.c_uint16, 36), ('vdd_gfx', ctypes.c_uint16, 38), ('cooling_id', ctypes.c_ubyte, 40), ('pp_table_id', ctypes.c_uint32, 44), ('format_revision', ctypes.c_uint32, 48), ('content_revision', ctypes.c_uint32, 52), ('fclk', ctypes.c_uint32, 56), ('lclk', ctypes.c_uint32, 60), ('firmware_caps', ctypes.c_uint32, 64)])
+struct_smu_bios_boot_up_values.register_fields([('revision', uint32_t, 0), ('gfxclk', uint32_t, 4), ('uclk', uint32_t, 8), ('socclk', uint32_t, 12), ('dcefclk', uint32_t, 16), ('eclk', uint32_t, 20), ('vclk', uint32_t, 24), ('dclk', uint32_t, 28), ('vddc', uint16_t, 32), ('vddci', uint16_t, 34), ('mvddc', uint16_t, 36), ('vdd_gfx', uint16_t, 38), ('cooling_id', uint8_t, 40), ('pp_table_id', uint32_t, 44), ('format_revision', uint32_t, 48), ('content_revision', uint32_t, 52), ('fclk', uint32_t, 56), ('lclk', uint32_t, 60), ('firmware_caps', uint32_t, 64)])
 enum_smu_table_id: dict[int, str] = {(SMU_TABLE_PPTABLE:=0): 'SMU_TABLE_PPTABLE', (SMU_TABLE_WATERMARKS:=1): 'SMU_TABLE_WATERMARKS', (SMU_TABLE_CUSTOM_DPM:=2): 'SMU_TABLE_CUSTOM_DPM', (SMU_TABLE_DPMCLOCKS:=3): 'SMU_TABLE_DPMCLOCKS', (SMU_TABLE_AVFS:=4): 'SMU_TABLE_AVFS', (SMU_TABLE_AVFS_PSM_DEBUG:=5): 'SMU_TABLE_AVFS_PSM_DEBUG', (SMU_TABLE_AVFS_FUSE_OVERRIDE:=6): 'SMU_TABLE_AVFS_FUSE_OVERRIDE', (SMU_TABLE_PMSTATUSLOG:=7): 'SMU_TABLE_PMSTATUSLOG', (SMU_TABLE_SMU_METRICS:=8): 'SMU_TABLE_SMU_METRICS', (SMU_TABLE_DRIVER_SMU_CONFIG:=9): 'SMU_TABLE_DRIVER_SMU_CONFIG', (SMU_TABLE_ACTIVITY_MONITOR_COEFF:=10): 'SMU_TABLE_ACTIVITY_MONITOR_COEFF', (SMU_TABLE_OVERDRIVE:=11): 'SMU_TABLE_OVERDRIVE', (SMU_TABLE_I2C_COMMANDS:=12): 'SMU_TABLE_I2C_COMMANDS', (SMU_TABLE_PACE:=13): 'SMU_TABLE_PACE', (SMU_TABLE_ECCINFO:=14): 'SMU_TABLE_ECCINFO', (SMU_TABLE_COMBO_PPTABLE:=15): 'SMU_TABLE_COMBO_PPTABLE', (SMU_TABLE_WIFIBAND:=16): 'SMU_TABLE_WIFIBAND', (SMU_TABLE_COUNT:=17): 'SMU_TABLE_COUNT'}
-FEATURE_CCLK_DPM_BIT = 0 # type: ignore
-FEATURE_FAN_CONTROLLER_BIT = 1 # type: ignore
-FEATURE_DATA_CALCULATION_BIT = 2 # type: ignore
-FEATURE_PPT_BIT = 3 # type: ignore
-FEATURE_TDC_BIT = 4 # type: ignore
-FEATURE_THERMAL_BIT = 5 # type: ignore
-FEATURE_FIT_BIT = 6 # type: ignore
-FEATURE_EDC_BIT = 7 # type: ignore
-FEATURE_PLL_POWER_DOWN_BIT = 8 # type: ignore
-FEATURE_VDDOFF_BIT = 9 # type: ignore
-FEATURE_VCN_DPM_BIT = 10 # type: ignore
-FEATURE_DS_MPM_BIT = 11 # type: ignore
-FEATURE_FCLK_DPM_BIT = 12 # type: ignore
-FEATURE_SOCCLK_DPM_BIT = 13 # type: ignore
-FEATURE_DS_MPIO_BIT = 14 # type: ignore
-FEATURE_LCLK_DPM_BIT = 15 # type: ignore
-FEATURE_SHUBCLK_DPM_BIT = 16 # type: ignore
-FEATURE_DCFCLK_DPM_BIT = 17 # type: ignore
-FEATURE_ISP_DPM_BIT = 18 # type: ignore
-FEATURE_IPU_DPM_BIT = 19 # type: ignore
-FEATURE_GFX_DPM_BIT = 20 # type: ignore
-FEATURE_DS_GFXCLK_BIT = 21 # type: ignore
-FEATURE_DS_SOCCLK_BIT = 22 # type: ignore
-FEATURE_DS_LCLK_BIT = 23 # type: ignore
-FEATURE_LOW_POWER_DCNCLKS_BIT = 24 # type: ignore
-FEATURE_DS_SHUBCLK_BIT = 25 # type: ignore
-FEATURE_RESERVED0_BIT = 26 # type: ignore
-FEATURE_ZSTATES_BIT = 27 # type: ignore
-FEATURE_IOMMUL2_PG_BIT = 28 # type: ignore
-FEATURE_DS_FCLK_BIT = 29 # type: ignore
-FEATURE_DS_SMNCLK_BIT = 30 # type: ignore
-FEATURE_DS_MP1CLK_BIT = 31 # type: ignore
-FEATURE_WHISPER_MODE_BIT = 32 # type: ignore
-FEATURE_SMU_LOW_POWER_BIT = 33 # type: ignore
-FEATURE_RESERVED1_BIT = 34 # type: ignore
-FEATURE_GFX_DEM_BIT = 35 # type: ignore
-FEATURE_PSI_BIT = 36 # type: ignore
-FEATURE_PROCHOT_BIT = 37 # type: ignore
-FEATURE_CPUOFF_BIT = 38 # type: ignore
-FEATURE_STAPM_BIT = 39 # type: ignore
-FEATURE_S0I3_BIT = 40 # type: ignore
-FEATURE_DF_LIGHT_CSTATE = 41 # type: ignore
-FEATURE_PERF_LIMIT_BIT = 42 # type: ignore
-FEATURE_CORE_DLDO_BIT = 43 # type: ignore
-FEATURE_DVO_BIT = 44 # type: ignore
-FEATURE_DS_VCN_BIT = 45 # type: ignore
-FEATURE_CPPC_BIT = 46 # type: ignore
-FEATURE_CPPC_PREFERRED_CORES = 47 # type: ignore
-FEATURE_DF_CSTATES_BIT = 48 # type: ignore
-FEATURE_FAST_PSTATE_CLDO_BIT = 49 # type: ignore
-FEATURE_ATHUB_PG_BIT = 50 # type: ignore
-FEATURE_VDDOFF_ECO_BIT = 51 # type: ignore
-FEATURE_ZSTATES_ECO_BIT = 52 # type: ignore
-FEATURE_CC6_BIT = 53 # type: ignore
-FEATURE_DS_UMCCLK_BIT = 54 # type: ignore
-FEATURE_DS_ISPCLK_BIT = 55 # type: ignore
-FEATURE_DS_HSPCLK_BIT = 56 # type: ignore
-FEATURE_P3T_BIT = 57 # type: ignore
-FEATURE_DS_IPUCLK_BIT = 58 # type: ignore
-FEATURE_DS_VPECLK_BIT = 59 # type: ignore
-FEATURE_VPE_DPM_BIT = 60 # type: ignore
-FEATURE_SMART_L3_RINSER_BIT = 61 # type: ignore
-FEATURE_PCC_BIT = 62 # type: ignore
-NUM_FEATURES = 63 # type: ignore
-PPSMC_VERSION = 0x1 # type: ignore
-PPSMC_Result_OK = 0x1 # type: ignore
-PPSMC_Result_Failed = 0xFF # type: ignore
-PPSMC_Result_UnknownCmd = 0xFE # type: ignore
-PPSMC_Result_CmdRejectedPrereq = 0xFD # type: ignore
-PPSMC_Result_CmdRejectedBusy = 0xFC # type: ignore
-PPSMC_MSG_TestMessage = 0x1 # type: ignore
-PPSMC_MSG_GetSmuVersion = 0x2 # type: ignore
-PPSMC_MSG_GetDriverIfVersion = 0x3 # type: ignore
-PPSMC_MSG_SetAllowedFeaturesMaskLow = 0x4 # type: ignore
-PPSMC_MSG_SetAllowedFeaturesMaskHigh = 0x5 # type: ignore
-PPSMC_MSG_EnableAllSmuFeatures = 0x6 # type: ignore
-PPSMC_MSG_DisableAllSmuFeatures = 0x7 # type: ignore
-PPSMC_MSG_EnableSmuFeaturesLow = 0x8 # type: ignore
-PPSMC_MSG_EnableSmuFeaturesHigh = 0x9 # type: ignore
-PPSMC_MSG_DisableSmuFeaturesLow = 0xA # type: ignore
-PPSMC_MSG_DisableSmuFeaturesHigh = 0xB # type: ignore
-PPSMC_MSG_GetRunningSmuFeaturesLow = 0xC # type: ignore
-PPSMC_MSG_GetRunningSmuFeaturesHigh = 0xD # type: ignore
-PPSMC_MSG_SetDriverDramAddrHigh = 0xE # type: ignore
-PPSMC_MSG_SetDriverDramAddrLow = 0xF # type: ignore
-PPSMC_MSG_SetToolsDramAddrHigh = 0x10 # type: ignore
-PPSMC_MSG_SetToolsDramAddrLow = 0x11 # type: ignore
-PPSMC_MSG_TransferTableSmu2Dram = 0x12 # type: ignore
-PPSMC_MSG_TransferTableDram2Smu = 0x13 # type: ignore
-PPSMC_MSG_UseDefaultPPTable = 0x14 # type: ignore
-PPSMC_MSG_EnterBaco = 0x15 # type: ignore
-PPSMC_MSG_ExitBaco = 0x16 # type: ignore
-PPSMC_MSG_ArmD3 = 0x17 # type: ignore
-PPSMC_MSG_BacoAudioD3PME = 0x18 # type: ignore
-PPSMC_MSG_SetSoftMinByFreq = 0x19 # type: ignore
-PPSMC_MSG_SetSoftMaxByFreq = 0x1A # type: ignore
-PPSMC_MSG_SetHardMinByFreq = 0x1B # type: ignore
-PPSMC_MSG_SetHardMaxByFreq = 0x1C # type: ignore
-PPSMC_MSG_GetMinDpmFreq = 0x1D # type: ignore
-PPSMC_MSG_GetMaxDpmFreq = 0x1E # type: ignore
-PPSMC_MSG_GetDpmFreqByIndex = 0x1F # type: ignore
-PPSMC_MSG_OverridePcieParameters = 0x20 # type: ignore
-PPSMC_MSG_DramLogSetDramAddrHigh = 0x21 # type: ignore
-PPSMC_MSG_DramLogSetDramAddrLow = 0x22 # type: ignore
-PPSMC_MSG_DramLogSetDramSize = 0x23 # type: ignore
-PPSMC_MSG_SetWorkloadMask = 0x24 # type: ignore
-PPSMC_MSG_GetVoltageByDpm = 0x25 # type: ignore
-PPSMC_MSG_SetVideoFps = 0x26 # type: ignore
-PPSMC_MSG_GetDcModeMaxDpmFreq = 0x27 # type: ignore
-PPSMC_MSG_AllowGfxOff = 0x28 # type: ignore
-PPSMC_MSG_DisallowGfxOff = 0x29 # type: ignore
-PPSMC_MSG_PowerUpVcn = 0x2A # type: ignore
-PPSMC_MSG_PowerDownVcn = 0x2B # type: ignore
-PPSMC_MSG_PowerUpJpeg = 0x2C # type: ignore
-PPSMC_MSG_PowerDownJpeg = 0x2D # type: ignore
-PPSMC_MSG_PrepareMp1ForUnload = 0x2E # type: ignore
-PPSMC_MSG_SetSystemVirtualDramAddrHigh = 0x30 # type: ignore
-PPSMC_MSG_SetSystemVirtualDramAddrLow = 0x31 # type: ignore
-PPSMC_MSG_SetPptLimit = 0x32 # type: ignore
-PPSMC_MSG_GetPptLimit = 0x33 # type: ignore
-PPSMC_MSG_ReenableAcDcInterrupt = 0x34 # type: ignore
-PPSMC_MSG_NotifyPowerSource = 0x35 # type: ignore
-PPSMC_MSG_RunDcBtc = 0x36 # type: ignore
-PPSMC_MSG_SetTemperatureInputSelect = 0x38 # type: ignore
-PPSMC_MSG_SetFwDstatesMask = 0x39 # type: ignore
-PPSMC_MSG_SetThrottlerMask = 0x3A # type: ignore
-PPSMC_MSG_SetExternalClientDfCstateAllow = 0x3B # type: ignore
-PPSMC_MSG_SetMGpuFanBoostLimitRpm = 0x3C # type: ignore
-PPSMC_MSG_DumpSTBtoDram = 0x3D # type: ignore
-PPSMC_MSG_STBtoDramLogSetDramAddress = 0x3E # type: ignore
-PPSMC_MSG_DummyUndefined = 0x3F # type: ignore
-PPSMC_MSG_STBtoDramLogSetDramSize = 0x40 # type: ignore
-PPSMC_MSG_SetOBMTraceBufferLogging = 0x41 # type: ignore
-PPSMC_MSG_UseProfilingMode = 0x42 # type: ignore
-PPSMC_MSG_AllowGfxDcs = 0x43 # type: ignore
-PPSMC_MSG_DisallowGfxDcs = 0x44 # type: ignore
-PPSMC_MSG_EnableAudioStutterWA = 0x45 # type: ignore
-PPSMC_MSG_PowerUpUmsch = 0x46 # type: ignore
-PPSMC_MSG_PowerDownUmsch = 0x47 # type: ignore
-PPSMC_MSG_SetDcsArch = 0x48 # type: ignore
-PPSMC_MSG_TriggerVFFLR = 0x49 # type: ignore
-PPSMC_MSG_SetNumBadMemoryPagesRetired = 0x4A # type: ignore
-PPSMC_MSG_SetBadMemoryPagesRetiredFlagsPerChannel = 0x4B # type: ignore
-PPSMC_MSG_SetPriorityDeltaGain = 0x4C # type: ignore
-PPSMC_MSG_AllowIHHostInterrupt = 0x4D # type: ignore
-PPSMC_MSG_EnableShadowDpm = 0x4E # type: ignore
-PPSMC_MSG_Mode3Reset = 0x4F # type: ignore
-PPSMC_MSG_SetDriverDramAddr = 0x50 # type: ignore
-PPSMC_MSG_SetToolsDramAddr = 0x51 # type: ignore
-PPSMC_MSG_TransferTableSmu2DramWithAddr = 0x52 # type: ignore
-PPSMC_MSG_TransferTableDram2SmuWithAddr = 0x53 # type: ignore
-PPSMC_MSG_GetAllRunningSmuFeatures = 0x54 # type: ignore
-PPSMC_MSG_GetSvi3Voltage = 0x55 # type: ignore
-PPSMC_MSG_UpdatePolicy = 0x56 # type: ignore
-PPSMC_MSG_ExtPwrConnSupport = 0x57 # type: ignore
-PPSMC_MSG_PreloadSwPstateForUclkOverDrive = 0x58 # type: ignore
-PPSMC_Message_Count = 0x59 # type: ignore
-PPTABLE_VERSION = 0x1B # type: ignore
-NUM_GFXCLK_DPM_LEVELS = 16 # type: ignore
-NUM_SOCCLK_DPM_LEVELS = 8 # type: ignore
-NUM_MP0CLK_DPM_LEVELS = 2 # type: ignore
-NUM_DCLK_DPM_LEVELS = 8 # type: ignore
-NUM_VCLK_DPM_LEVELS = 8 # type: ignore
-NUM_DISPCLK_DPM_LEVELS = 8 # type: ignore
-NUM_DPPCLK_DPM_LEVELS = 8 # type: ignore
-NUM_DPREFCLK_DPM_LEVELS = 8 # type: ignore
-NUM_DCFCLK_DPM_LEVELS = 8 # type: ignore
-NUM_DTBCLK_DPM_LEVELS = 8 # type: ignore
-NUM_UCLK_DPM_LEVELS = 6 # type: ignore
-NUM_LINK_LEVELS = 3 # type: ignore
-NUM_FCLK_DPM_LEVELS = 8 # type: ignore
-NUM_OD_FAN_MAX_POINTS = 6 # type: ignore
-FEATURE_FW_DATA_READ_BIT = 0 # type: ignore
-FEATURE_DPM_GFXCLK_BIT = 1 # type: ignore
-FEATURE_DPM_GFX_POWER_OPTIMIZER_BIT = 2 # type: ignore
-FEATURE_DPM_UCLK_BIT = 3 # type: ignore
-FEATURE_DPM_FCLK_BIT = 4 # type: ignore
-FEATURE_DPM_SOCCLK_BIT = 5 # type: ignore
-FEATURE_DPM_LINK_BIT = 6 # type: ignore
-FEATURE_DPM_DCN_BIT = 7 # type: ignore
-FEATURE_VMEMP_SCALING_BIT = 8 # type: ignore
-FEATURE_VDDIO_MEM_SCALING_BIT = 9 # type: ignore
-FEATURE_DS_GFXCLK_BIT = 10 # type: ignore
-FEATURE_DS_SOCCLK_BIT = 11 # type: ignore
-FEATURE_DS_FCLK_BIT = 12 # type: ignore
-FEATURE_DS_LCLK_BIT = 13 # type: ignore
-FEATURE_DS_DCFCLK_BIT = 14 # type: ignore
-FEATURE_DS_UCLK_BIT = 15 # type: ignore
-FEATURE_GFX_ULV_BIT = 16 # type: ignore
-FEATURE_FW_DSTATE_BIT = 17 # type: ignore
-FEATURE_GFXOFF_BIT = 18 # type: ignore
-FEATURE_BACO_BIT = 19 # type: ignore
-FEATURE_MM_DPM_BIT = 20 # type: ignore
-FEATURE_SOC_MPCLK_DS_BIT = 21 # type: ignore
-FEATURE_BACO_MPCLK_DS_BIT = 22 # type: ignore
-FEATURE_THROTTLERS_BIT = 23 # type: ignore
-FEATURE_SMARTSHIFT_BIT = 24 # type: ignore
-FEATURE_GTHR_BIT = 25 # type: ignore
-FEATURE_ACDC_BIT = 26 # type: ignore
-FEATURE_VR0HOT_BIT = 27 # type: ignore
-FEATURE_FW_CTF_BIT = 28 # type: ignore
-FEATURE_FAN_CONTROL_BIT = 29 # type: ignore
-FEATURE_GFX_DCS_BIT = 30 # type: ignore
-FEATURE_GFX_READ_MARGIN_BIT = 31 # type: ignore
-FEATURE_LED_DISPLAY_BIT = 32 # type: ignore
-FEATURE_GFXCLK_SPREAD_SPECTRUM_BIT = 33 # type: ignore
-FEATURE_OUT_OF_BAND_MONITOR_BIT = 34 # type: ignore
-FEATURE_OPTIMIZED_VMIN_BIT = 35 # type: ignore
-FEATURE_GFX_IMU_BIT = 36 # type: ignore
-FEATURE_BOOT_TIME_CAL_BIT = 37 # type: ignore
-FEATURE_GFX_PCC_DFLL_BIT = 38 # type: ignore
-FEATURE_SOC_CG_BIT = 39 # type: ignore
-FEATURE_DF_CSTATE_BIT = 40 # type: ignore
-FEATURE_GFX_EDC_BIT = 41 # type: ignore
-FEATURE_BOOT_POWER_OPT_BIT = 42 # type: ignore
-FEATURE_CLOCK_POWER_DOWN_BYPASS_BIT = 43 # type: ignore
-FEATURE_DS_VCN_BIT = 44 # type: ignore
-FEATURE_BACO_CG_BIT = 45 # type: ignore
-FEATURE_MEM_TEMP_READ_BIT = 46 # type: ignore
-FEATURE_ATHUB_MMHUB_PG_BIT = 47 # type: ignore
-FEATURE_SOC_PCC_BIT = 48 # type: ignore
-FEATURE_EDC_PWRBRK_BIT = 49 # type: ignore
-FEATURE_SOC_EDC_XVMIN_BIT = 50 # type: ignore
-FEATURE_GFX_PSM_DIDT_BIT = 51 # type: ignore
-FEATURE_APT_ALL_ENABLE_BIT = 52 # type: ignore
-FEATURE_APT_SQ_THROTTLE_BIT = 53 # type: ignore
-FEATURE_APT_PF_DCS_BIT = 54 # type: ignore
-FEATURE_GFX_EDC_XVMIN_BIT = 55 # type: ignore
-FEATURE_GFX_DIDT_XVMIN_BIT = 56 # type: ignore
-FEATURE_FAN_ABNORMAL_BIT = 57 # type: ignore
-FEATURE_CLOCK_STRETCH_COMPENSATOR = 58 # type: ignore
-FEATURE_SPARE_59_BIT = 59 # type: ignore
-FEATURE_SPARE_60_BIT = 60 # type: ignore
-FEATURE_SPARE_61_BIT = 61 # type: ignore
-FEATURE_SPARE_62_BIT = 62 # type: ignore
-FEATURE_SPARE_63_BIT = 63 # type: ignore
-NUM_FEATURES = 64 # type: ignore
-ALLOWED_FEATURE_CTRL_DEFAULT = 0xFFFFFFFFFFFFFFFF # type: ignore
-ALLOWED_FEATURE_CTRL_SCPM = (1 << FEATURE_DPM_GFXCLK_BIT) | (1 << FEATURE_DPM_GFX_POWER_OPTIMIZER_BIT) | (1 << FEATURE_DPM_UCLK_BIT) | (1 << FEATURE_DPM_FCLK_BIT) | (1 << FEATURE_DPM_SOCCLK_BIT) | (1 << FEATURE_DPM_LINK_BIT) | (1 << FEATURE_DPM_DCN_BIT) | (1 << FEATURE_DS_GFXCLK_BIT) | (1 << FEATURE_DS_SOCCLK_BIT) | (1 << FEATURE_DS_FCLK_BIT) | (1 << FEATURE_DS_LCLK_BIT) | (1 << FEATURE_DS_DCFCLK_BIT) | (1 << FEATURE_DS_UCLK_BIT) | (1 << FEATURE_DS_VCN_BIT) # type: ignore
-DEBUG_OVERRIDE_NOT_USE = 0x00000001 # type: ignore
-DEBUG_OVERRIDE_DISABLE_VOLT_LINK_DCN_FCLK = 0x00000002 # type: ignore
-DEBUG_OVERRIDE_DISABLE_VOLT_LINK_MP0_FCLK = 0x00000004 # type: ignore
-DEBUG_OVERRIDE_DISABLE_VOLT_LINK_VCN_DCFCLK = 0x00000008 # type: ignore
-DEBUG_OVERRIDE_DISABLE_FAST_FCLK_TIMER = 0x00000010 # type: ignore
-DEBUG_OVERRIDE_DISABLE_VCN_PG = 0x00000020 # type: ignore
-DEBUG_OVERRIDE_DISABLE_FMAX_VMAX = 0x00000040 # type: ignore
-DEBUG_OVERRIDE_DISABLE_IMU_FW_CHECKS = 0x00000080 # type: ignore
-DEBUG_OVERRIDE_DISABLE_D0i2_REENTRY_HSR_TIMER_CHECK = 0x00000100 # type: ignore
-DEBUG_OVERRIDE_DISABLE_DFLL = 0x00000200 # type: ignore
-DEBUG_OVERRIDE_ENABLE_RLC_VF_BRINGUP_MODE = 0x00000400 # type: ignore
-DEBUG_OVERRIDE_DFLL_MASTER_MODE = 0x00000800 # type: ignore
-DEBUG_OVERRIDE_ENABLE_PROFILING_MODE = 0x00001000 # type: ignore
-DEBUG_OVERRIDE_ENABLE_SOC_VF_BRINGUP_MODE = 0x00002000 # type: ignore
-DEBUG_OVERRIDE_ENABLE_PER_WGP_RESIENCY = 0x00004000 # type: ignore
-DEBUG_OVERRIDE_DISABLE_MEMORY_VOLTAGE_SCALING = 0x00008000 # type: ignore
-DEBUG_OVERRIDE_DFLL_BTC_FCW_LOG = 0x00010000 # type: ignore
-VR_MAPPING_VR_SELECT_MASK = 0x01 # type: ignore
-VR_MAPPING_VR_SELECT_SHIFT = 0x00 # type: ignore
-VR_MAPPING_PLANE_SELECT_MASK = 0x02 # type: ignore
-VR_MAPPING_PLANE_SELECT_SHIFT = 0x01 # type: ignore
-PSI_SEL_VR0_PLANE0_PSI0 = 0x01 # type: ignore
-PSI_SEL_VR0_PLANE0_PSI1 = 0x02 # type: ignore
-PSI_SEL_VR0_PLANE1_PSI0 = 0x04 # type: ignore
-PSI_SEL_VR0_PLANE1_PSI1 = 0x08 # type: ignore
-PSI_SEL_VR1_PLANE0_PSI0 = 0x10 # type: ignore
-PSI_SEL_VR1_PLANE0_PSI1 = 0x20 # type: ignore
-PSI_SEL_VR1_PLANE1_PSI0 = 0x40 # type: ignore
-PSI_SEL_VR1_PLANE1_PSI1 = 0x80 # type: ignore
-THROTTLER_TEMP_EDGE_BIT = 0 # type: ignore
-THROTTLER_TEMP_HOTSPOT_BIT = 1 # type: ignore
-THROTTLER_TEMP_HOTSPOT_GFX_BIT = 2 # type: ignore
-THROTTLER_TEMP_HOTSPOT_SOC_BIT = 3 # type: ignore
-THROTTLER_TEMP_MEM_BIT = 4 # type: ignore
-THROTTLER_TEMP_VR_GFX_BIT = 5 # type: ignore
-THROTTLER_TEMP_VR_SOC_BIT = 6 # type: ignore
-THROTTLER_TEMP_VR_MEM0_BIT = 7 # type: ignore
-THROTTLER_TEMP_VR_MEM1_BIT = 8 # type: ignore
-THROTTLER_TEMP_LIQUID0_BIT = 9 # type: ignore
-THROTTLER_TEMP_LIQUID1_BIT = 10 # type: ignore
-THROTTLER_TEMP_PLX_BIT = 11 # type: ignore
-THROTTLER_TDC_GFX_BIT = 12 # type: ignore
-THROTTLER_TDC_SOC_BIT = 13 # type: ignore
-THROTTLER_PPT0_BIT = 14 # type: ignore
-THROTTLER_PPT1_BIT = 15 # type: ignore
-THROTTLER_PPT2_BIT = 16 # type: ignore
-THROTTLER_PPT3_BIT = 17 # type: ignore
-THROTTLER_FIT_BIT = 18 # type: ignore
-THROTTLER_GFX_APCC_PLUS_BIT = 19 # type: ignore
-THROTTLER_GFX_DVO_BIT = 20 # type: ignore
-THROTTLER_COUNT = 21 # type: ignore
-FW_DSTATE_SOC_ULV_BIT = 0 # type: ignore
-FW_DSTATE_G6_HSR_BIT = 1 # type: ignore
-FW_DSTATE_G6_PHY_VMEMP_OFF_BIT = 2 # type: ignore
-FW_DSTATE_SMN_DS_BIT = 3 # type: ignore
-FW_DSTATE_MP1_WHISPER_MODE_BIT = 4 # type: ignore
-FW_DSTATE_SOC_LIV_MIN_BIT = 5 # type: ignore
-FW_DSTATE_SOC_PLL_PWRDN_BIT = 6 # type: ignore
-FW_DSTATE_MEM_PLL_PWRDN_BIT = 7 # type: ignore
-FW_DSTATE_MALL_ALLOC_BIT = 8 # type: ignore
-FW_DSTATE_MEM_PSI_BIT = 9 # type: ignore
-FW_DSTATE_HSR_NON_STROBE_BIT = 10 # type: ignore
-FW_DSTATE_MP0_ENTER_WFI_BIT = 11 # type: ignore
-FW_DSTATE_MALL_FLUSH_BIT = 12 # type: ignore
-FW_DSTATE_SOC_PSI_BIT = 13 # type: ignore
-FW_DSTATE_MMHUB_INTERLOCK_BIT = 14 # type: ignore
-FW_DSTATE_D0i3_2_QUIET_FW_BIT = 15 # type: ignore
-FW_DSTATE_CLDO_PRG_BIT = 16 # type: ignore
-FW_DSTATE_DF_PLL_PWRDN_BIT = 17 # type: ignore
-LED_DISPLAY_GFX_DPM_BIT = 0 # type: ignore
-LED_DISPLAY_PCIE_BIT = 1 # type: ignore
-LED_DISPLAY_ERROR_BIT = 2 # type: ignore
-MEM_TEMP_READ_OUT_OF_BAND_BIT = 0 # type: ignore
-MEM_TEMP_READ_IN_BAND_REFRESH_BIT = 1 # type: ignore
-MEM_TEMP_READ_IN_BAND_DUMMY_PSTATE_BIT = 2 # type: ignore
-NUM_I2C_CONTROLLERS = 8 # type: ignore
-I2C_CONTROLLER_ENABLED = 1 # type: ignore
-I2C_CONTROLLER_DISABLED = 0 # type: ignore
-MAX_SW_I2C_COMMANDS = 24 # type: ignore
-CMDCONFIG_STOP_BIT = 0 # type: ignore
-CMDCONFIG_RESTART_BIT = 1 # type: ignore
-CMDCONFIG_READWRITE_BIT = 2 # type: ignore
-CMDCONFIG_STOP_MASK = (1 << CMDCONFIG_STOP_BIT) # type: ignore
-CMDCONFIG_RESTART_MASK = (1 << CMDCONFIG_RESTART_BIT) # type: ignore
-CMDCONFIG_READWRITE_MASK = (1 << CMDCONFIG_READWRITE_BIT) # type: ignore
-EPCS_HIGH_POWER = 600 # type: ignore
-EPCS_NORMAL_POWER = 450 # type: ignore
-EPCS_LOW_POWER = 300 # type: ignore
-EPCS_SHORTED_POWER = 150 # type: ignore
-EPCS_NO_BOOTUP = 0 # type: ignore
-PP_NUM_RTAVFS_PWL_ZONES = 5 # type: ignore
-PP_NUM_PSM_DIDT_PWL_ZONES = 3 # type: ignore
-PP_NUM_OD_VF_CURVE_POINTS = PP_NUM_RTAVFS_PWL_ZONES + 1 # type: ignore
-PP_OD_FEATURE_GFX_VF_CURVE_BIT = 0 # type: ignore
-PP_OD_FEATURE_GFX_VMAX_BIT = 1 # type: ignore
-PP_OD_FEATURE_SOC_VMAX_BIT = 2 # type: ignore
-PP_OD_FEATURE_PPT_BIT = 3 # type: ignore
-PP_OD_FEATURE_FAN_CURVE_BIT = 4 # type: ignore
-PP_OD_FEATURE_FAN_LEGACY_BIT = 5 # type: ignore
-PP_OD_FEATURE_FULL_CTRL_BIT = 6 # type: ignore
-PP_OD_FEATURE_TDC_BIT = 7 # type: ignore
-PP_OD_FEATURE_GFXCLK_BIT = 8 # type: ignore
-PP_OD_FEATURE_UCLK_BIT = 9 # type: ignore
-PP_OD_FEATURE_FCLK_BIT = 10 # type: ignore
-PP_OD_FEATURE_ZERO_FAN_BIT = 11 # type: ignore
-PP_OD_FEATURE_TEMPERATURE_BIT = 12 # type: ignore
-PP_OD_FEATURE_EDC_BIT = 13 # type: ignore
-PP_OD_FEATURE_COUNT = 14 # type: ignore
-INVALID_BOARD_GPIO = 0xFF # type: ignore
-NUM_WM_RANGES = 4 # type: ignore
-WORKLOAD_PPLIB_DEFAULT_BIT = 0 # type: ignore
-WORKLOAD_PPLIB_FULL_SCREEN_3D_BIT = 1 # type: ignore
-WORKLOAD_PPLIB_POWER_SAVING_BIT = 2 # type: ignore
-WORKLOAD_PPLIB_VIDEO_BIT = 3 # type: ignore
-WORKLOAD_PPLIB_VR_BIT = 4 # type: ignore
-WORKLOAD_PPLIB_COMPUTE_BIT = 5 # type: ignore
-WORKLOAD_PPLIB_CUSTOM_BIT = 6 # type: ignore
-WORKLOAD_PPLIB_WINDOW_3D_BIT = 7 # type: ignore
-WORKLOAD_PPLIB_DIRECT_ML_BIT = 8 # type: ignore
-WORKLOAD_PPLIB_CGVDI_BIT = 9 # type: ignore
-WORKLOAD_PPLIB_COUNT = 10 # type: ignore
-TABLE_TRANSFER_OK = 0x0 # type: ignore
-TABLE_TRANSFER_FAILED = 0xFF # type: ignore
-TABLE_TRANSFER_PENDING = 0xAB # type: ignore
-TABLE_PPT_FAILED = 0x100 # type: ignore
-TABLE_TDC_FAILED = 0x200 # type: ignore
-TABLE_TEMP_FAILED = 0x400 # type: ignore
-TABLE_FAN_TARGET_TEMP_FAILED = 0x800 # type: ignore
-TABLE_FAN_STOP_TEMP_FAILED = 0x1000 # type: ignore
-TABLE_FAN_START_TEMP_FAILED = 0x2000 # type: ignore
-TABLE_FAN_PWM_MIN_FAILED = 0x4000 # type: ignore
-TABLE_ACOUSTIC_TARGET_RPM_FAILED = 0x8000 # type: ignore
-TABLE_ACOUSTIC_LIMIT_RPM_FAILED = 0x10000 # type: ignore
-TABLE_MGPU_ACOUSTIC_TARGET_RPM_FAILED = 0x20000 # type: ignore
-TABLE_PPTABLE = 0 # type: ignore
-TABLE_COMBO_PPTABLE = 1 # type: ignore
-TABLE_WATERMARKS = 2 # type: ignore
-TABLE_AVFS_PSM_DEBUG = 3 # type: ignore
-TABLE_PMSTATUSLOG = 4 # type: ignore
-TABLE_SMU_METRICS = 5 # type: ignore
-TABLE_DRIVER_SMU_CONFIG = 6 # type: ignore
-TABLE_ACTIVITY_MONITOR_COEFF = 7 # type: ignore
-TABLE_OVERDRIVE = 8 # type: ignore
-TABLE_I2C_COMMANDS = 9 # type: ignore
-TABLE_DRIVER_INFO = 10 # type: ignore
-TABLE_ECCINFO = 11 # type: ignore
-TABLE_CUSTOM_SKUTABLE = 12 # type: ignore
-TABLE_COUNT = 13 # type: ignore
-IH_INTERRUPT_ID_TO_DRIVER = 0xFE # type: ignore
-IH_INTERRUPT_CONTEXT_ID_BACO = 0x2 # type: ignore
-IH_INTERRUPT_CONTEXT_ID_AC = 0x3 # type: ignore
-IH_INTERRUPT_CONTEXT_ID_DC = 0x4 # type: ignore
-IH_INTERRUPT_CONTEXT_ID_AUDIO_D0 = 0x5 # type: ignore
-IH_INTERRUPT_CONTEXT_ID_AUDIO_D3 = 0x6 # type: ignore
-IH_INTERRUPT_CONTEXT_ID_THERMAL_THROTTLING = 0x7 # type: ignore
-IH_INTERRUPT_CONTEXT_ID_FAN_ABNORMAL = 0x8 # type: ignore
-IH_INTERRUPT_CONTEXT_ID_FAN_RECOVERY = 0x9 # type: ignore
-IH_INTERRUPT_CONTEXT_ID_DYNAMIC_TABLE = 0xA # type: ignore
-int32_t = int # type: ignore
-SMU_THERMAL_MINIMUM_ALERT_TEMP = 0 # type: ignore
-SMU_THERMAL_MAXIMUM_ALERT_TEMP = 255 # type: ignore
-SMU_TEMPERATURE_UNITS_PER_CENTIGRADES = 1000 # type: ignore
-SMU_FW_NAME_LEN = 0x24 # type: ignore
-SMU_DPM_USER_PROFILE_RESTORE = (1 << 0) # type: ignore
-SMU_CUSTOM_FAN_SPEED_RPM = (1 << 1) # type: ignore
-SMU_CUSTOM_FAN_SPEED_PWM = (1 << 2) # type: ignore
-SMU_THROTTLER_PPT0_BIT = 0 # type: ignore
-SMU_THROTTLER_PPT1_BIT = 1 # type: ignore
-SMU_THROTTLER_PPT2_BIT = 2 # type: ignore
-SMU_THROTTLER_PPT3_BIT = 3 # type: ignore
-SMU_THROTTLER_SPL_BIT = 4 # type: ignore
-SMU_THROTTLER_FPPT_BIT = 5 # type: ignore
-SMU_THROTTLER_SPPT_BIT = 6 # type: ignore
-SMU_THROTTLER_SPPT_APU_BIT = 7 # type: ignore
-SMU_THROTTLER_TDC_GFX_BIT = 16 # type: ignore
-SMU_THROTTLER_TDC_SOC_BIT = 17 # type: ignore
-SMU_THROTTLER_TDC_MEM_BIT = 18 # type: ignore
-SMU_THROTTLER_TDC_VDD_BIT = 19 # type: ignore
-SMU_THROTTLER_TDC_CVIP_BIT = 20 # type: ignore
-SMU_THROTTLER_EDC_CPU_BIT = 21 # type: ignore
-SMU_THROTTLER_EDC_GFX_BIT = 22 # type: ignore
-SMU_THROTTLER_APCC_BIT = 23 # type: ignore
-SMU_THROTTLER_TEMP_GPU_BIT = 32 # type: ignore
-SMU_THROTTLER_TEMP_CORE_BIT = 33 # type: ignore
-SMU_THROTTLER_TEMP_MEM_BIT = 34 # type: ignore
-SMU_THROTTLER_TEMP_EDGE_BIT = 35 # type: ignore
-SMU_THROTTLER_TEMP_HOTSPOT_BIT = 36 # type: ignore
-SMU_THROTTLER_TEMP_SOC_BIT = 37 # type: ignore
-SMU_THROTTLER_TEMP_VR_GFX_BIT = 38 # type: ignore
-SMU_THROTTLER_TEMP_VR_SOC_BIT = 39 # type: ignore
-SMU_THROTTLER_TEMP_VR_MEM0_BIT = 40 # type: ignore
-SMU_THROTTLER_TEMP_VR_MEM1_BIT = 41 # type: ignore
-SMU_THROTTLER_TEMP_LIQUID0_BIT = 42 # type: ignore
-SMU_THROTTLER_TEMP_LIQUID1_BIT = 43 # type: ignore
-SMU_THROTTLER_VRHOT0_BIT = 44 # type: ignore
-SMU_THROTTLER_VRHOT1_BIT = 45 # type: ignore
-SMU_THROTTLER_PROCHOT_CPU_BIT = 46 # type: ignore
-SMU_THROTTLER_PROCHOT_GFX_BIT = 47 # type: ignore
-SMU_THROTTLER_PPM_BIT = 56 # type: ignore
-SMU_THROTTLER_FIT_BIT = 57 # type: ignore
+FEATURE_CCLK_DPM_BIT = 0
+FEATURE_FAN_CONTROLLER_BIT = 1
+FEATURE_DATA_CALCULATION_BIT = 2
+FEATURE_PPT_BIT = 3
+FEATURE_TDC_BIT = 4
+FEATURE_THERMAL_BIT = 5
+FEATURE_FIT_BIT = 6
+FEATURE_EDC_BIT = 7
+FEATURE_PLL_POWER_DOWN_BIT = 8
+FEATURE_VDDOFF_BIT = 9
+FEATURE_VCN_DPM_BIT = 10
+FEATURE_DS_MPM_BIT = 11
+FEATURE_FCLK_DPM_BIT = 12
+FEATURE_SOCCLK_DPM_BIT = 13
+FEATURE_DS_MPIO_BIT = 14
+FEATURE_LCLK_DPM_BIT = 15
+FEATURE_SHUBCLK_DPM_BIT = 16
+FEATURE_DCFCLK_DPM_BIT = 17
+FEATURE_ISP_DPM_BIT = 18
+FEATURE_IPU_DPM_BIT = 19
+FEATURE_GFX_DPM_BIT = 20
+FEATURE_DS_GFXCLK_BIT = 21
+FEATURE_DS_SOCCLK_BIT = 22
+FEATURE_DS_LCLK_BIT = 23
+FEATURE_LOW_POWER_DCNCLKS_BIT = 24
+FEATURE_DS_SHUBCLK_BIT = 25
+FEATURE_RESERVED0_BIT = 26
+FEATURE_ZSTATES_BIT = 27
+FEATURE_IOMMUL2_PG_BIT = 28
+FEATURE_DS_FCLK_BIT = 29
+FEATURE_DS_SMNCLK_BIT = 30
+FEATURE_DS_MP1CLK_BIT = 31
+FEATURE_WHISPER_MODE_BIT = 32
+FEATURE_SMU_LOW_POWER_BIT = 33
+FEATURE_RESERVED1_BIT = 34
+FEATURE_GFX_DEM_BIT = 35
+FEATURE_PSI_BIT = 36
+FEATURE_PROCHOT_BIT = 37
+FEATURE_CPUOFF_BIT = 38
+FEATURE_STAPM_BIT = 39
+FEATURE_S0I3_BIT = 40
+FEATURE_DF_LIGHT_CSTATE = 41
+FEATURE_PERF_LIMIT_BIT = 42
+FEATURE_CORE_DLDO_BIT = 43
+FEATURE_DVO_BIT = 44
+FEATURE_DS_VCN_BIT = 45
+FEATURE_CPPC_BIT = 46
+FEATURE_CPPC_PREFERRED_CORES = 47
+FEATURE_DF_CSTATES_BIT = 48
+FEATURE_FAST_PSTATE_CLDO_BIT = 49
+FEATURE_ATHUB_PG_BIT = 50
+FEATURE_VDDOFF_ECO_BIT = 51
+FEATURE_ZSTATES_ECO_BIT = 52
+FEATURE_CC6_BIT = 53
+FEATURE_DS_UMCCLK_BIT = 54
+FEATURE_DS_ISPCLK_BIT = 55
+FEATURE_DS_HSPCLK_BIT = 56
+FEATURE_P3T_BIT = 57
+FEATURE_DS_IPUCLK_BIT = 58
+FEATURE_DS_VPECLK_BIT = 59
+FEATURE_VPE_DPM_BIT = 60
+FEATURE_SMART_L3_RINSER_BIT = 61
+FEATURE_PCC_BIT = 62
+NUM_FEATURES = 63
+PPSMC_VERSION = 0x1
+PPSMC_Result_OK = 0x1
+PPSMC_Result_Failed = 0xFF
+PPSMC_Result_UnknownCmd = 0xFE
+PPSMC_Result_CmdRejectedPrereq = 0xFD
+PPSMC_Result_CmdRejectedBusy = 0xFC
+PPSMC_MSG_TestMessage = 0x1
+PPSMC_MSG_GetSmuVersion = 0x2
+PPSMC_MSG_GetDriverIfVersion = 0x3
+PPSMC_MSG_SetAllowedFeaturesMaskLow = 0x4
+PPSMC_MSG_SetAllowedFeaturesMaskHigh = 0x5
+PPSMC_MSG_EnableAllSmuFeatures = 0x6
+PPSMC_MSG_DisableAllSmuFeatures = 0x7
+PPSMC_MSG_EnableSmuFeaturesLow = 0x8
+PPSMC_MSG_EnableSmuFeaturesHigh = 0x9
+PPSMC_MSG_DisableSmuFeaturesLow = 0xA
+PPSMC_MSG_DisableSmuFeaturesHigh = 0xB
+PPSMC_MSG_GetRunningSmuFeaturesLow = 0xC
+PPSMC_MSG_GetRunningSmuFeaturesHigh = 0xD
+PPSMC_MSG_SetDriverDramAddrHigh = 0xE
+PPSMC_MSG_SetDriverDramAddrLow = 0xF
+PPSMC_MSG_SetToolsDramAddrHigh = 0x10
+PPSMC_MSG_SetToolsDramAddrLow = 0x11
+PPSMC_MSG_TransferTableSmu2Dram = 0x12
+PPSMC_MSG_TransferTableDram2Smu = 0x13
+PPSMC_MSG_UseDefaultPPTable = 0x14
+PPSMC_MSG_EnterBaco = 0x15
+PPSMC_MSG_ExitBaco = 0x16
+PPSMC_MSG_ArmD3 = 0x17
+PPSMC_MSG_BacoAudioD3PME = 0x18
+PPSMC_MSG_SetSoftMinByFreq = 0x19
+PPSMC_MSG_SetSoftMaxByFreq = 0x1A
+PPSMC_MSG_SetHardMinByFreq = 0x1B
+PPSMC_MSG_SetHardMaxByFreq = 0x1C
+PPSMC_MSG_GetMinDpmFreq = 0x1D
+PPSMC_MSG_GetMaxDpmFreq = 0x1E
+PPSMC_MSG_GetDpmFreqByIndex = 0x1F
+PPSMC_MSG_OverridePcieParameters = 0x20
+PPSMC_MSG_DramLogSetDramAddrHigh = 0x21
+PPSMC_MSG_DramLogSetDramAddrLow = 0x22
+PPSMC_MSG_DramLogSetDramSize = 0x23
+PPSMC_MSG_SetWorkloadMask = 0x24
+PPSMC_MSG_GetVoltageByDpm = 0x25
+PPSMC_MSG_SetVideoFps = 0x26
+PPSMC_MSG_GetDcModeMaxDpmFreq = 0x27
+PPSMC_MSG_AllowGfxOff = 0x28
+PPSMC_MSG_DisallowGfxOff = 0x29
+PPSMC_MSG_PowerUpVcn = 0x2A
+PPSMC_MSG_PowerDownVcn = 0x2B
+PPSMC_MSG_PowerUpJpeg = 0x2C
+PPSMC_MSG_PowerDownJpeg = 0x2D
+PPSMC_MSG_PrepareMp1ForUnload = 0x2E
+PPSMC_MSG_SetSystemVirtualDramAddrHigh = 0x30
+PPSMC_MSG_SetSystemVirtualDramAddrLow = 0x31
+PPSMC_MSG_SetPptLimit = 0x32
+PPSMC_MSG_GetPptLimit = 0x33
+PPSMC_MSG_ReenableAcDcInterrupt = 0x34
+PPSMC_MSG_NotifyPowerSource = 0x35
+PPSMC_MSG_RunDcBtc = 0x36
+PPSMC_MSG_SetTemperatureInputSelect = 0x38
+PPSMC_MSG_SetFwDstatesMask = 0x39
+PPSMC_MSG_SetThrottlerMask = 0x3A
+PPSMC_MSG_SetExternalClientDfCstateAllow = 0x3B
+PPSMC_MSG_SetMGpuFanBoostLimitRpm = 0x3C
+PPSMC_MSG_DumpSTBtoDram = 0x3D
+PPSMC_MSG_STBtoDramLogSetDramAddress = 0x3E
+PPSMC_MSG_DummyUndefined = 0x3F
+PPSMC_MSG_STBtoDramLogSetDramSize = 0x40
+PPSMC_MSG_SetOBMTraceBufferLogging = 0x41
+PPSMC_MSG_UseProfilingMode = 0x42
+PPSMC_MSG_AllowGfxDcs = 0x43
+PPSMC_MSG_DisallowGfxDcs = 0x44
+PPSMC_MSG_EnableAudioStutterWA = 0x45
+PPSMC_MSG_PowerUpUmsch = 0x46
+PPSMC_MSG_PowerDownUmsch = 0x47
+PPSMC_MSG_SetDcsArch = 0x48
+PPSMC_MSG_TriggerVFFLR = 0x49
+PPSMC_MSG_SetNumBadMemoryPagesRetired = 0x4A
+PPSMC_MSG_SetBadMemoryPagesRetiredFlagsPerChannel = 0x4B
+PPSMC_MSG_SetPriorityDeltaGain = 0x4C
+PPSMC_MSG_AllowIHHostInterrupt = 0x4D
+PPSMC_MSG_EnableShadowDpm = 0x4E
+PPSMC_MSG_Mode3Reset = 0x4F
+PPSMC_MSG_SetDriverDramAddr = 0x50
+PPSMC_MSG_SetToolsDramAddr = 0x51
+PPSMC_MSG_TransferTableSmu2DramWithAddr = 0x52
+PPSMC_MSG_TransferTableDram2SmuWithAddr = 0x53
+PPSMC_MSG_GetAllRunningSmuFeatures = 0x54
+PPSMC_MSG_GetSvi3Voltage = 0x55
+PPSMC_MSG_UpdatePolicy = 0x56
+PPSMC_MSG_ExtPwrConnSupport = 0x57
+PPSMC_MSG_PreloadSwPstateForUclkOverDrive = 0x58
+PPSMC_Message_Count = 0x59
+PPTABLE_VERSION = 0x1B
+NUM_GFXCLK_DPM_LEVELS = 16
+NUM_SOCCLK_DPM_LEVELS = 8
+NUM_MP0CLK_DPM_LEVELS = 2
+NUM_DCLK_DPM_LEVELS = 8
+NUM_VCLK_DPM_LEVELS = 8
+NUM_DISPCLK_DPM_LEVELS = 8
+NUM_DPPCLK_DPM_LEVELS = 8
+NUM_DPREFCLK_DPM_LEVELS = 8
+NUM_DCFCLK_DPM_LEVELS = 8
+NUM_DTBCLK_DPM_LEVELS = 8
+NUM_UCLK_DPM_LEVELS = 6
+NUM_LINK_LEVELS = 3
+NUM_FCLK_DPM_LEVELS = 8
+NUM_OD_FAN_MAX_POINTS = 6
+FEATURE_FW_DATA_READ_BIT = 0
+FEATURE_DPM_GFXCLK_BIT = 1
+FEATURE_DPM_GFX_POWER_OPTIMIZER_BIT = 2
+FEATURE_DPM_UCLK_BIT = 3
+FEATURE_DPM_FCLK_BIT = 4
+FEATURE_DPM_SOCCLK_BIT = 5
+FEATURE_DPM_LINK_BIT = 6
+FEATURE_DPM_DCN_BIT = 7
+FEATURE_VMEMP_SCALING_BIT = 8
+FEATURE_VDDIO_MEM_SCALING_BIT = 9
+FEATURE_DS_GFXCLK_BIT = 10
+FEATURE_DS_SOCCLK_BIT = 11
+FEATURE_DS_FCLK_BIT = 12
+FEATURE_DS_LCLK_BIT = 13
+FEATURE_DS_DCFCLK_BIT = 14
+FEATURE_DS_UCLK_BIT = 15
+FEATURE_GFX_ULV_BIT = 16
+FEATURE_FW_DSTATE_BIT = 17
+FEATURE_GFXOFF_BIT = 18
+FEATURE_BACO_BIT = 19
+FEATURE_MM_DPM_BIT = 20
+FEATURE_SOC_MPCLK_DS_BIT = 21
+FEATURE_BACO_MPCLK_DS_BIT = 22
+FEATURE_THROTTLERS_BIT = 23
+FEATURE_SMARTSHIFT_BIT = 24
+FEATURE_GTHR_BIT = 25
+FEATURE_ACDC_BIT = 26
+FEATURE_VR0HOT_BIT = 27
+FEATURE_FW_CTF_BIT = 28
+FEATURE_FAN_CONTROL_BIT = 29
+FEATURE_GFX_DCS_BIT = 30
+FEATURE_GFX_READ_MARGIN_BIT = 31
+FEATURE_LED_DISPLAY_BIT = 32
+FEATURE_GFXCLK_SPREAD_SPECTRUM_BIT = 33
+FEATURE_OUT_OF_BAND_MONITOR_BIT = 34
+FEATURE_OPTIMIZED_VMIN_BIT = 35
+FEATURE_GFX_IMU_BIT = 36
+FEATURE_BOOT_TIME_CAL_BIT = 37
+FEATURE_GFX_PCC_DFLL_BIT = 38
+FEATURE_SOC_CG_BIT = 39
+FEATURE_DF_CSTATE_BIT = 40
+FEATURE_GFX_EDC_BIT = 41
+FEATURE_BOOT_POWER_OPT_BIT = 42
+FEATURE_CLOCK_POWER_DOWN_BYPASS_BIT = 43
+FEATURE_DS_VCN_BIT = 44
+FEATURE_BACO_CG_BIT = 45
+FEATURE_MEM_TEMP_READ_BIT = 46
+FEATURE_ATHUB_MMHUB_PG_BIT = 47
+FEATURE_SOC_PCC_BIT = 48
+FEATURE_EDC_PWRBRK_BIT = 49
+FEATURE_SOC_EDC_XVMIN_BIT = 50
+FEATURE_GFX_PSM_DIDT_BIT = 51
+FEATURE_APT_ALL_ENABLE_BIT = 52
+FEATURE_APT_SQ_THROTTLE_BIT = 53
+FEATURE_APT_PF_DCS_BIT = 54
+FEATURE_GFX_EDC_XVMIN_BIT = 55
+FEATURE_GFX_DIDT_XVMIN_BIT = 56
+FEATURE_FAN_ABNORMAL_BIT = 57
+FEATURE_CLOCK_STRETCH_COMPENSATOR = 58
+FEATURE_SPARE_59_BIT = 59
+FEATURE_SPARE_60_BIT = 60
+FEATURE_SPARE_61_BIT = 61
+FEATURE_SPARE_62_BIT = 62
+FEATURE_SPARE_63_BIT = 63
+NUM_FEATURES = 64
+ALLOWED_FEATURE_CTRL_DEFAULT = 0xFFFFFFFFFFFFFFFF
+ALLOWED_FEATURE_CTRL_SCPM = (1 << FEATURE_DPM_GFXCLK_BIT) | (1 << FEATURE_DPM_GFX_POWER_OPTIMIZER_BIT) | (1 << FEATURE_DPM_UCLK_BIT) | (1 << FEATURE_DPM_FCLK_BIT) | (1 << FEATURE_DPM_SOCCLK_BIT) | (1 << FEATURE_DPM_LINK_BIT) | (1 << FEATURE_DPM_DCN_BIT) | (1 << FEATURE_DS_GFXCLK_BIT) | (1 << FEATURE_DS_SOCCLK_BIT) | (1 << FEATURE_DS_FCLK_BIT) | (1 << FEATURE_DS_LCLK_BIT) | (1 << FEATURE_DS_DCFCLK_BIT) | (1 << FEATURE_DS_UCLK_BIT) | (1 << FEATURE_DS_VCN_BIT)
+DEBUG_OVERRIDE_NOT_USE = 0x00000001
+DEBUG_OVERRIDE_DISABLE_VOLT_LINK_DCN_FCLK = 0x00000002
+DEBUG_OVERRIDE_DISABLE_VOLT_LINK_MP0_FCLK = 0x00000004
+DEBUG_OVERRIDE_DISABLE_VOLT_LINK_VCN_DCFCLK = 0x00000008
+DEBUG_OVERRIDE_DISABLE_FAST_FCLK_TIMER = 0x00000010
+DEBUG_OVERRIDE_DISABLE_VCN_PG = 0x00000020
+DEBUG_OVERRIDE_DISABLE_FMAX_VMAX = 0x00000040
+DEBUG_OVERRIDE_DISABLE_IMU_FW_CHECKS = 0x00000080
+DEBUG_OVERRIDE_DISABLE_D0i2_REENTRY_HSR_TIMER_CHECK = 0x00000100
+DEBUG_OVERRIDE_DISABLE_DFLL = 0x00000200
+DEBUG_OVERRIDE_ENABLE_RLC_VF_BRINGUP_MODE = 0x00000400
+DEBUG_OVERRIDE_DFLL_MASTER_MODE = 0x00000800
+DEBUG_OVERRIDE_ENABLE_PROFILING_MODE = 0x00001000
+DEBUG_OVERRIDE_ENABLE_SOC_VF_BRINGUP_MODE = 0x00002000
+DEBUG_OVERRIDE_ENABLE_PER_WGP_RESIENCY = 0x00004000
+DEBUG_OVERRIDE_DISABLE_MEMORY_VOLTAGE_SCALING = 0x00008000
+DEBUG_OVERRIDE_DFLL_BTC_FCW_LOG = 0x00010000
+VR_MAPPING_VR_SELECT_MASK = 0x01
+VR_MAPPING_VR_SELECT_SHIFT = 0x00
+VR_MAPPING_PLANE_SELECT_MASK = 0x02
+VR_MAPPING_PLANE_SELECT_SHIFT = 0x01
+PSI_SEL_VR0_PLANE0_PSI0 = 0x01
+PSI_SEL_VR0_PLANE0_PSI1 = 0x02
+PSI_SEL_VR0_PLANE1_PSI0 = 0x04
+PSI_SEL_VR0_PLANE1_PSI1 = 0x08
+PSI_SEL_VR1_PLANE0_PSI0 = 0x10
+PSI_SEL_VR1_PLANE0_PSI1 = 0x20
+PSI_SEL_VR1_PLANE1_PSI0 = 0x40
+PSI_SEL_VR1_PLANE1_PSI1 = 0x80
+THROTTLER_TEMP_EDGE_BIT = 0
+THROTTLER_TEMP_HOTSPOT_BIT = 1
+THROTTLER_TEMP_HOTSPOT_GFX_BIT = 2
+THROTTLER_TEMP_HOTSPOT_SOC_BIT = 3
+THROTTLER_TEMP_MEM_BIT = 4
+THROTTLER_TEMP_VR_GFX_BIT = 5
+THROTTLER_TEMP_VR_SOC_BIT = 6
+THROTTLER_TEMP_VR_MEM0_BIT = 7
+THROTTLER_TEMP_VR_MEM1_BIT = 8
+THROTTLER_TEMP_LIQUID0_BIT = 9
+THROTTLER_TEMP_LIQUID1_BIT = 10
+THROTTLER_TEMP_PLX_BIT = 11
+THROTTLER_TDC_GFX_BIT = 12
+THROTTLER_TDC_SOC_BIT = 13
+THROTTLER_PPT0_BIT = 14
+THROTTLER_PPT1_BIT = 15
+THROTTLER_PPT2_BIT = 16
+THROTTLER_PPT3_BIT = 17
+THROTTLER_FIT_BIT = 18
+THROTTLER_GFX_APCC_PLUS_BIT = 19
+THROTTLER_GFX_DVO_BIT = 20
+THROTTLER_COUNT = 21
+FW_DSTATE_SOC_ULV_BIT = 0
+FW_DSTATE_G6_HSR_BIT = 1
+FW_DSTATE_G6_PHY_VMEMP_OFF_BIT = 2
+FW_DSTATE_SMN_DS_BIT = 3
+FW_DSTATE_MP1_WHISPER_MODE_BIT = 4
+FW_DSTATE_SOC_LIV_MIN_BIT = 5
+FW_DSTATE_SOC_PLL_PWRDN_BIT = 6
+FW_DSTATE_MEM_PLL_PWRDN_BIT = 7
+FW_DSTATE_MALL_ALLOC_BIT = 8
+FW_DSTATE_MEM_PSI_BIT = 9
+FW_DSTATE_HSR_NON_STROBE_BIT = 10
+FW_DSTATE_MP0_ENTER_WFI_BIT = 11
+FW_DSTATE_MALL_FLUSH_BIT = 12
+FW_DSTATE_SOC_PSI_BIT = 13
+FW_DSTATE_MMHUB_INTERLOCK_BIT = 14
+FW_DSTATE_D0i3_2_QUIET_FW_BIT = 15
+FW_DSTATE_CLDO_PRG_BIT = 16
+FW_DSTATE_DF_PLL_PWRDN_BIT = 17
+LED_DISPLAY_GFX_DPM_BIT = 0
+LED_DISPLAY_PCIE_BIT = 1
+LED_DISPLAY_ERROR_BIT = 2
+MEM_TEMP_READ_OUT_OF_BAND_BIT = 0
+MEM_TEMP_READ_IN_BAND_REFRESH_BIT = 1
+MEM_TEMP_READ_IN_BAND_DUMMY_PSTATE_BIT = 2
+NUM_I2C_CONTROLLERS = 8
+I2C_CONTROLLER_ENABLED = 1
+I2C_CONTROLLER_DISABLED = 0
+MAX_SW_I2C_COMMANDS = 24
+CMDCONFIG_STOP_BIT = 0
+CMDCONFIG_RESTART_BIT = 1
+CMDCONFIG_READWRITE_BIT = 2
+CMDCONFIG_STOP_MASK = (1 << CMDCONFIG_STOP_BIT)
+CMDCONFIG_RESTART_MASK = (1 << CMDCONFIG_RESTART_BIT)
+CMDCONFIG_READWRITE_MASK = (1 << CMDCONFIG_READWRITE_BIT)
+EPCS_HIGH_POWER = 600
+EPCS_NORMAL_POWER = 450
+EPCS_LOW_POWER = 300
+EPCS_SHORTED_POWER = 150
+EPCS_NO_BOOTUP = 0
+PP_NUM_RTAVFS_PWL_ZONES = 5
+PP_NUM_PSM_DIDT_PWL_ZONES = 3
+PP_NUM_OD_VF_CURVE_POINTS = PP_NUM_RTAVFS_PWL_ZONES + 1
+PP_OD_FEATURE_GFX_VF_CURVE_BIT = 0
+PP_OD_FEATURE_GFX_VMAX_BIT = 1
+PP_OD_FEATURE_SOC_VMAX_BIT = 2
+PP_OD_FEATURE_PPT_BIT = 3
+PP_OD_FEATURE_FAN_CURVE_BIT = 4
+PP_OD_FEATURE_FAN_LEGACY_BIT = 5
+PP_OD_FEATURE_FULL_CTRL_BIT = 6
+PP_OD_FEATURE_TDC_BIT = 7
+PP_OD_FEATURE_GFXCLK_BIT = 8
+PP_OD_FEATURE_UCLK_BIT = 9
+PP_OD_FEATURE_FCLK_BIT = 10
+PP_OD_FEATURE_ZERO_FAN_BIT = 11
+PP_OD_FEATURE_TEMPERATURE_BIT = 12
+PP_OD_FEATURE_EDC_BIT = 13
+PP_OD_FEATURE_COUNT = 14
+INVALID_BOARD_GPIO = 0xFF
+NUM_WM_RANGES = 4
+WORKLOAD_PPLIB_DEFAULT_BIT = 0
+WORKLOAD_PPLIB_FULL_SCREEN_3D_BIT = 1
+WORKLOAD_PPLIB_POWER_SAVING_BIT = 2
+WORKLOAD_PPLIB_VIDEO_BIT = 3
+WORKLOAD_PPLIB_VR_BIT = 4
+WORKLOAD_PPLIB_COMPUTE_BIT = 5
+WORKLOAD_PPLIB_CUSTOM_BIT = 6
+WORKLOAD_PPLIB_WINDOW_3D_BIT = 7
+WORKLOAD_PPLIB_DIRECT_ML_BIT = 8
+WORKLOAD_PPLIB_CGVDI_BIT = 9
+WORKLOAD_PPLIB_COUNT = 10
+TABLE_TRANSFER_OK = 0x0
+TABLE_TRANSFER_FAILED = 0xFF
+TABLE_TRANSFER_PENDING = 0xAB
+TABLE_PPT_FAILED = 0x100
+TABLE_TDC_FAILED = 0x200
+TABLE_TEMP_FAILED = 0x400
+TABLE_FAN_TARGET_TEMP_FAILED = 0x800
+TABLE_FAN_STOP_TEMP_FAILED = 0x1000
+TABLE_FAN_START_TEMP_FAILED = 0x2000
+TABLE_FAN_PWM_MIN_FAILED = 0x4000
+TABLE_ACOUSTIC_TARGET_RPM_FAILED = 0x8000
+TABLE_ACOUSTIC_LIMIT_RPM_FAILED = 0x10000
+TABLE_MGPU_ACOUSTIC_TARGET_RPM_FAILED = 0x20000
+TABLE_PPTABLE = 0
+TABLE_COMBO_PPTABLE = 1
+TABLE_WATERMARKS = 2
+TABLE_AVFS_PSM_DEBUG = 3
+TABLE_PMSTATUSLOG = 4
+TABLE_SMU_METRICS = 5
+TABLE_DRIVER_SMU_CONFIG = 6
+TABLE_ACTIVITY_MONITOR_COEFF = 7
+TABLE_OVERDRIVE = 8
+TABLE_I2C_COMMANDS = 9
+TABLE_DRIVER_INFO = 10
+TABLE_ECCINFO = 11
+TABLE_CUSTOM_SKUTABLE = 12
+TABLE_COUNT = 13
+IH_INTERRUPT_ID_TO_DRIVER = 0xFE
+IH_INTERRUPT_CONTEXT_ID_BACO = 0x2
+IH_INTERRUPT_CONTEXT_ID_AC = 0x3
+IH_INTERRUPT_CONTEXT_ID_DC = 0x4
+IH_INTERRUPT_CONTEXT_ID_AUDIO_D0 = 0x5
+IH_INTERRUPT_CONTEXT_ID_AUDIO_D3 = 0x6
+IH_INTERRUPT_CONTEXT_ID_THERMAL_THROTTLING = 0x7
+IH_INTERRUPT_CONTEXT_ID_FAN_ABNORMAL = 0x8
+IH_INTERRUPT_CONTEXT_ID_FAN_RECOVERY = 0x9
+IH_INTERRUPT_CONTEXT_ID_DYNAMIC_TABLE = 0xA
+SMU_THERMAL_MINIMUM_ALERT_TEMP = 0
+SMU_THERMAL_MAXIMUM_ALERT_TEMP = 255
+SMU_TEMPERATURE_UNITS_PER_CENTIGRADES = 1000
+SMU_FW_NAME_LEN = 0x24
+SMU_DPM_USER_PROFILE_RESTORE = (1 << 0)
+SMU_CUSTOM_FAN_SPEED_RPM = (1 << 1)
+SMU_CUSTOM_FAN_SPEED_PWM = (1 << 2)
+SMU_THROTTLER_PPT0_BIT = 0
+SMU_THROTTLER_PPT1_BIT = 1
+SMU_THROTTLER_PPT2_BIT = 2
+SMU_THROTTLER_PPT3_BIT = 3
+SMU_THROTTLER_SPL_BIT = 4
+SMU_THROTTLER_FPPT_BIT = 5
+SMU_THROTTLER_SPPT_BIT = 6
+SMU_THROTTLER_SPPT_APU_BIT = 7
+SMU_THROTTLER_TDC_GFX_BIT = 16
+SMU_THROTTLER_TDC_SOC_BIT = 17
+SMU_THROTTLER_TDC_MEM_BIT = 18
+SMU_THROTTLER_TDC_VDD_BIT = 19
+SMU_THROTTLER_TDC_CVIP_BIT = 20
+SMU_THROTTLER_EDC_CPU_BIT = 21
+SMU_THROTTLER_EDC_GFX_BIT = 22
+SMU_THROTTLER_APCC_BIT = 23
+SMU_THROTTLER_TEMP_GPU_BIT = 32
+SMU_THROTTLER_TEMP_CORE_BIT = 33
+SMU_THROTTLER_TEMP_MEM_BIT = 34
+SMU_THROTTLER_TEMP_EDGE_BIT = 35
+SMU_THROTTLER_TEMP_HOTSPOT_BIT = 36
+SMU_THROTTLER_TEMP_SOC_BIT = 37
+SMU_THROTTLER_TEMP_VR_GFX_BIT = 38
+SMU_THROTTLER_TEMP_VR_SOC_BIT = 39
+SMU_THROTTLER_TEMP_VR_MEM0_BIT = 40
+SMU_THROTTLER_TEMP_VR_MEM1_BIT = 41
+SMU_THROTTLER_TEMP_LIQUID0_BIT = 42
+SMU_THROTTLER_TEMP_LIQUID1_BIT = 43
+SMU_THROTTLER_VRHOT0_BIT = 44
+SMU_THROTTLER_VRHOT1_BIT = 45
+SMU_THROTTLER_PROCHOT_CPU_BIT = 46
+SMU_THROTTLER_PROCHOT_GFX_BIT = 47
+SMU_THROTTLER_PPM_BIT = 56
+SMU_THROTTLER_FIT_BIT = 57

tinygrad/runtime/autogen/am/vega_offsets.py

@@ -0,0 +1,774 @@
+# mypy: disable-error-code="empty-body"
+from __future__ import annotations
+import ctypes
+from typing import Literal, TypeAlias
+from tinygrad.runtime.support.c import _IO, _IOW, _IOR, _IOWR
+from tinygrad.runtime.support import c
+@c.record
+class struct_IP_BASE_INSTANCE(c.Struct):
+  SIZE = 24
+  segment: c.Array[ctypes.c_uint32, Literal[6]]
+struct_IP_BASE_INSTANCE.register_fields([('segment', c.Array[ctypes.c_uint32, Literal[6]], 0)])
+@c.record
+class struct_IP_BASE(c.Struct):
+  SIZE = 144
+  instance: c.Array[struct_IP_BASE_INSTANCE, Literal[6]]
+struct_IP_BASE.register_fields([('instance', c.Array[struct_IP_BASE_INSTANCE, Literal[6]], 0)])
+MAX_INSTANCE = 6
+MAX_SEGMENT = 6
+ATHUB_BASE__INST0_SEG0 = 0x00000C20
+ATHUB_BASE__INST0_SEG1 = 0
+ATHUB_BASE__INST0_SEG2 = 0
+ATHUB_BASE__INST0_SEG3 = 0
+ATHUB_BASE__INST0_SEG4 = 0
+ATHUB_BASE__INST0_SEG5 = 0
+ATHUB_BASE__INST1_SEG0 = 0
+ATHUB_BASE__INST1_SEG1 = 0
+ATHUB_BASE__INST1_SEG2 = 0
+ATHUB_BASE__INST1_SEG3 = 0
+ATHUB_BASE__INST1_SEG4 = 0
+ATHUB_BASE__INST1_SEG5 = 0
+ATHUB_BASE__INST2_SEG0 = 0
+ATHUB_BASE__INST2_SEG1 = 0
+ATHUB_BASE__INST2_SEG2 = 0
+ATHUB_BASE__INST2_SEG3 = 0
+ATHUB_BASE__INST2_SEG4 = 0
+ATHUB_BASE__INST2_SEG5 = 0
+ATHUB_BASE__INST3_SEG0 = 0
+ATHUB_BASE__INST3_SEG1 = 0
+ATHUB_BASE__INST3_SEG2 = 0
+ATHUB_BASE__INST3_SEG3 = 0
+ATHUB_BASE__INST3_SEG4 = 0
+ATHUB_BASE__INST3_SEG5 = 0
+ATHUB_BASE__INST4_SEG0 = 0
+ATHUB_BASE__INST4_SEG1 = 0
+ATHUB_BASE__INST4_SEG2 = 0
+ATHUB_BASE__INST4_SEG3 = 0
+ATHUB_BASE__INST4_SEG4 = 0
+ATHUB_BASE__INST4_SEG5 = 0
+ATHUB_BASE__INST5_SEG0 = 0
+ATHUB_BASE__INST5_SEG1 = 0
+ATHUB_BASE__INST5_SEG2 = 0
+ATHUB_BASE__INST5_SEG3 = 0
+ATHUB_BASE__INST5_SEG4 = 0
+ATHUB_BASE__INST5_SEG5 = 0
+CLK_BASE__INST0_SEG0 = 0x00016C00
+CLK_BASE__INST0_SEG1 = 0x00016E00
+CLK_BASE__INST0_SEG2 = 0x00017000
+CLK_BASE__INST0_SEG3 = 0x00017200
+CLK_BASE__INST0_SEG4 = 0x0001B000
+CLK_BASE__INST0_SEG5 = 0x0001B200
+CLK_BASE__INST1_SEG0 = 0
+CLK_BASE__INST1_SEG1 = 0
+CLK_BASE__INST1_SEG2 = 0
+CLK_BASE__INST1_SEG3 = 0
+CLK_BASE__INST1_SEG4 = 0
+CLK_BASE__INST1_SEG5 = 0
+CLK_BASE__INST2_SEG0 = 0
+CLK_BASE__INST2_SEG1 = 0
+CLK_BASE__INST2_SEG2 = 0
+CLK_BASE__INST2_SEG3 = 0
+CLK_BASE__INST2_SEG4 = 0
+CLK_BASE__INST2_SEG5 = 0
+CLK_BASE__INST3_SEG0 = 0
+CLK_BASE__INST3_SEG1 = 0
+CLK_BASE__INST3_SEG2 = 0
+CLK_BASE__INST3_SEG3 = 0
+CLK_BASE__INST3_SEG4 = 0
+CLK_BASE__INST3_SEG5 = 0
+CLK_BASE__INST4_SEG0 = 0
+CLK_BASE__INST4_SEG1 = 0
+CLK_BASE__INST4_SEG2 = 0
+CLK_BASE__INST4_SEG3 = 0
+CLK_BASE__INST4_SEG4 = 0
+CLK_BASE__INST4_SEG5 = 0
+CLK_BASE__INST5_SEG0 = 0
+CLK_BASE__INST5_SEG1 = 0
+CLK_BASE__INST5_SEG2 = 0
+CLK_BASE__INST5_SEG3 = 0
+CLK_BASE__INST5_SEG4 = 0
+CLK_BASE__INST5_SEG5 = 0
+DCE_BASE__INST0_SEG0 = 0x00000012
+DCE_BASE__INST0_SEG1 = 0x000000C0
+DCE_BASE__INST0_SEG2 = 0x000034C0
+DCE_BASE__INST0_SEG3 = 0
+DCE_BASE__INST0_SEG4 = 0
+DCE_BASE__INST0_SEG5 = 0
+DCE_BASE__INST1_SEG0 = 0
+DCE_BASE__INST1_SEG1 = 0
+DCE_BASE__INST1_SEG2 = 0
+DCE_BASE__INST1_SEG3 = 0
+DCE_BASE__INST1_SEG4 = 0
+DCE_BASE__INST1_SEG5 = 0
+DCE_BASE__INST2_SEG0 = 0
+DCE_BASE__INST2_SEG1 = 0
+DCE_BASE__INST2_SEG2 = 0
+DCE_BASE__INST2_SEG3 = 0
+DCE_BASE__INST2_SEG4 = 0
+DCE_BASE__INST2_SEG5 = 0
+DCE_BASE__INST3_SEG0 = 0
+DCE_BASE__INST3_SEG1 = 0
+DCE_BASE__INST3_SEG2 = 0
+DCE_BASE__INST3_SEG3 = 0
+DCE_BASE__INST3_SEG4 = 0
+DCE_BASE__INST3_SEG5 = 0
+DCE_BASE__INST4_SEG0 = 0
+DCE_BASE__INST4_SEG1 = 0
+DCE_BASE__INST4_SEG2 = 0
+DCE_BASE__INST4_SEG3 = 0
+DCE_BASE__INST4_SEG4 = 0
+DCE_BASE__INST4_SEG5 = 0
+DCE_BASE__INST5_SEG0 = 0
+DCE_BASE__INST5_SEG1 = 0
+DCE_BASE__INST5_SEG2 = 0
+DCE_BASE__INST5_SEG3 = 0
+DCE_BASE__INST5_SEG4 = 0
+DCE_BASE__INST5_SEG5 = 0
+DF_BASE__INST0_SEG0 = 0x00007000
+DF_BASE__INST0_SEG1 = 0
+DF_BASE__INST0_SEG2 = 0
+DF_BASE__INST0_SEG3 = 0
+DF_BASE__INST0_SEG4 = 0
+DF_BASE__INST0_SEG5 = 0
+DF_BASE__INST1_SEG0 = 0
+DF_BASE__INST1_SEG1 = 0
+DF_BASE__INST1_SEG2 = 0
+DF_BASE__INST1_SEG3 = 0
+DF_BASE__INST1_SEG4 = 0
+DF_BASE__INST1_SEG5 = 0
+DF_BASE__INST2_SEG0 = 0
+DF_BASE__INST2_SEG1 = 0
+DF_BASE__INST2_SEG2 = 0
+DF_BASE__INST2_SEG3 = 0
+DF_BASE__INST2_SEG4 = 0
+DF_BASE__INST2_SEG5 = 0
+DF_BASE__INST3_SEG0 = 0
+DF_BASE__INST3_SEG1 = 0
+DF_BASE__INST3_SEG2 = 0
+DF_BASE__INST3_SEG3 = 0
+DF_BASE__INST3_SEG4 = 0
+DF_BASE__INST3_SEG5 = 0
+DF_BASE__INST4_SEG0 = 0
+DF_BASE__INST4_SEG1 = 0
+DF_BASE__INST4_SEG2 = 0
+DF_BASE__INST4_SEG3 = 0
+DF_BASE__INST4_SEG4 = 0
+DF_BASE__INST4_SEG5 = 0
+DF_BASE__INST5_SEG0 = 0
+DF_BASE__INST5_SEG1 = 0
+DF_BASE__INST5_SEG2 = 0
+DF_BASE__INST5_SEG3 = 0
+DF_BASE__INST5_SEG4 = 0
+DF_BASE__INST5_SEG5 = 0
+FUSE_BASE__INST0_SEG0 = 0x00017400
+FUSE_BASE__INST0_SEG1 = 0
+FUSE_BASE__INST0_SEG2 = 0
+FUSE_BASE__INST0_SEG3 = 0
+FUSE_BASE__INST0_SEG4 = 0
+FUSE_BASE__INST0_SEG5 = 0
+FUSE_BASE__INST1_SEG0 = 0
+FUSE_BASE__INST1_SEG1 = 0
+FUSE_BASE__INST1_SEG2 = 0
+FUSE_BASE__INST1_SEG3 = 0
+FUSE_BASE__INST1_SEG4 = 0
+FUSE_BASE__INST1_SEG5 = 0
+FUSE_BASE__INST2_SEG0 = 0
+FUSE_BASE__INST2_SEG1 = 0
+FUSE_BASE__INST2_SEG2 = 0
+FUSE_BASE__INST2_SEG3 = 0
+FUSE_BASE__INST2_SEG4 = 0
+FUSE_BASE__INST2_SEG5 = 0
+FUSE_BASE__INST3_SEG0 = 0
+FUSE_BASE__INST3_SEG1 = 0
+FUSE_BASE__INST3_SEG2 = 0
+FUSE_BASE__INST3_SEG3 = 0
+FUSE_BASE__INST3_SEG4 = 0
+FUSE_BASE__INST3_SEG5 = 0
+FUSE_BASE__INST4_SEG0 = 0
+FUSE_BASE__INST4_SEG1 = 0
+FUSE_BASE__INST4_SEG2 = 0
+FUSE_BASE__INST4_SEG3 = 0
+FUSE_BASE__INST4_SEG4 = 0
+FUSE_BASE__INST4_SEG5 = 0
+FUSE_BASE__INST5_SEG0 = 0
+FUSE_BASE__INST5_SEG1 = 0
+FUSE_BASE__INST5_SEG2 = 0
+FUSE_BASE__INST5_SEG3 = 0
+FUSE_BASE__INST5_SEG4 = 0
+FUSE_BASE__INST5_SEG5 = 0
+GC_BASE__INST0_SEG0 = 0x00002000
+GC_BASE__INST0_SEG1 = 0x0000A000
+GC_BASE__INST0_SEG2 = 0
+GC_BASE__INST0_SEG3 = 0
+GC_BASE__INST0_SEG4 = 0
+GC_BASE__INST0_SEG5 = 0
+GC_BASE__INST1_SEG0 = 0
+GC_BASE__INST1_SEG1 = 0
+GC_BASE__INST1_SEG2 = 0
+GC_BASE__INST1_SEG3 = 0
+GC_BASE__INST1_SEG4 = 0
+GC_BASE__INST1_SEG5 = 0
+GC_BASE__INST2_SEG0 = 0
+GC_BASE__INST2_SEG1 = 0
+GC_BASE__INST2_SEG2 = 0
+GC_BASE__INST2_SEG3 = 0
+GC_BASE__INST2_SEG4 = 0
+GC_BASE__INST2_SEG5 = 0
+GC_BASE__INST3_SEG0 = 0
+GC_BASE__INST3_SEG1 = 0
+GC_BASE__INST3_SEG2 = 0
+GC_BASE__INST3_SEG3 = 0
+GC_BASE__INST3_SEG4 = 0
+GC_BASE__INST3_SEG5 = 0
+GC_BASE__INST4_SEG0 = 0
+GC_BASE__INST4_SEG1 = 0
+GC_BASE__INST4_SEG2 = 0
+GC_BASE__INST4_SEG3 = 0
+GC_BASE__INST4_SEG4 = 0
+GC_BASE__INST4_SEG5 = 0
+GC_BASE__INST5_SEG0 = 0
+GC_BASE__INST5_SEG1 = 0
+GC_BASE__INST5_SEG2 = 0
+GC_BASE__INST5_SEG3 = 0
+GC_BASE__INST5_SEG4 = 0
+GC_BASE__INST5_SEG5 = 0
+HDP_BASE__INST0_SEG0 = 0x00000F20
+HDP_BASE__INST0_SEG1 = 0
+HDP_BASE__INST0_SEG2 = 0
+HDP_BASE__INST0_SEG3 = 0
+HDP_BASE__INST0_SEG4 = 0
+HDP_BASE__INST0_SEG5 = 0
+HDP_BASE__INST1_SEG0 = 0
+HDP_BASE__INST1_SEG1 = 0
+HDP_BASE__INST1_SEG2 = 0
+HDP_BASE__INST1_SEG3 = 0
+HDP_BASE__INST1_SEG4 = 0
+HDP_BASE__INST1_SEG5 = 0
+HDP_BASE__INST2_SEG0 = 0
+HDP_BASE__INST2_SEG1 = 0
+HDP_BASE__INST2_SEG2 = 0
+HDP_BASE__INST2_SEG3 = 0
+HDP_BASE__INST2_SEG4 = 0
+HDP_BASE__INST2_SEG5 = 0
+HDP_BASE__INST3_SEG0 = 0
+HDP_BASE__INST3_SEG1 = 0
+HDP_BASE__INST3_SEG2 = 0
+HDP_BASE__INST3_SEG3 = 0
+HDP_BASE__INST3_SEG4 = 0
+HDP_BASE__INST3_SEG5 = 0
+HDP_BASE__INST4_SEG0 = 0
+HDP_BASE__INST4_SEG1 = 0
+HDP_BASE__INST4_SEG2 = 0
+HDP_BASE__INST4_SEG3 = 0
+HDP_BASE__INST4_SEG4 = 0
+HDP_BASE__INST4_SEG5 = 0
+HDP_BASE__INST5_SEG0 = 0
+HDP_BASE__INST5_SEG1 = 0
+HDP_BASE__INST5_SEG2 = 0
+HDP_BASE__INST5_SEG3 = 0
+HDP_BASE__INST5_SEG4 = 0
+HDP_BASE__INST5_SEG5 = 0
+MMHUB_BASE__INST0_SEG0 = 0x0001A000
+MMHUB_BASE__INST0_SEG1 = 0
+MMHUB_BASE__INST0_SEG2 = 0
+MMHUB_BASE__INST0_SEG3 = 0
+MMHUB_BASE__INST0_SEG4 = 0
+MMHUB_BASE__INST0_SEG5 = 0
+MMHUB_BASE__INST1_SEG0 = 0
+MMHUB_BASE__INST1_SEG1 = 0
+MMHUB_BASE__INST1_SEG2 = 0
+MMHUB_BASE__INST1_SEG3 = 0
+MMHUB_BASE__INST1_SEG4 = 0
+MMHUB_BASE__INST1_SEG5 = 0
+MMHUB_BASE__INST2_SEG0 = 0
+MMHUB_BASE__INST2_SEG1 = 0
+MMHUB_BASE__INST2_SEG2 = 0
+MMHUB_BASE__INST2_SEG3 = 0
+MMHUB_BASE__INST2_SEG4 = 0
+MMHUB_BASE__INST2_SEG5 = 0
+MMHUB_BASE__INST3_SEG0 = 0
+MMHUB_BASE__INST3_SEG1 = 0
+MMHUB_BASE__INST3_SEG2 = 0
+MMHUB_BASE__INST3_SEG3 = 0
+MMHUB_BASE__INST3_SEG4 = 0
+MMHUB_BASE__INST3_SEG5 = 0
+MMHUB_BASE__INST4_SEG0 = 0
+MMHUB_BASE__INST4_SEG1 = 0
+MMHUB_BASE__INST4_SEG2 = 0
+MMHUB_BASE__INST4_SEG3 = 0
+MMHUB_BASE__INST4_SEG4 = 0
+MMHUB_BASE__INST4_SEG5 = 0
+MMHUB_BASE__INST5_SEG0 = 0
+MMHUB_BASE__INST5_SEG1 = 0
+MMHUB_BASE__INST5_SEG2 = 0
+MMHUB_BASE__INST5_SEG3 = 0
+MMHUB_BASE__INST5_SEG4 = 0
+MMHUB_BASE__INST5_SEG5 = 0
+MP0_BASE__INST0_SEG0 = 0x00016000
+MP0_BASE__INST0_SEG1 = 0
+MP0_BASE__INST0_SEG2 = 0
+MP0_BASE__INST0_SEG3 = 0
+MP0_BASE__INST0_SEG4 = 0
+MP0_BASE__INST0_SEG5 = 0
+MP0_BASE__INST1_SEG0 = 0
+MP0_BASE__INST1_SEG1 = 0
+MP0_BASE__INST1_SEG2 = 0
+MP0_BASE__INST1_SEG3 = 0
+MP0_BASE__INST1_SEG4 = 0
+MP0_BASE__INST1_SEG5 = 0
+MP0_BASE__INST2_SEG0 = 0
+MP0_BASE__INST2_SEG1 = 0
+MP0_BASE__INST2_SEG2 = 0
+MP0_BASE__INST2_SEG3 = 0
+MP0_BASE__INST2_SEG4 = 0
+MP0_BASE__INST2_SEG5 = 0
+MP0_BASE__INST3_SEG0 = 0
+MP0_BASE__INST3_SEG1 = 0
+MP0_BASE__INST3_SEG2 = 0
+MP0_BASE__INST3_SEG3 = 0
+MP0_BASE__INST3_SEG4 = 0
+MP0_BASE__INST3_SEG5 = 0
+MP0_BASE__INST4_SEG0 = 0
+MP0_BASE__INST4_SEG1 = 0
+MP0_BASE__INST4_SEG2 = 0
+MP0_BASE__INST4_SEG3 = 0
+MP0_BASE__INST4_SEG4 = 0
+MP0_BASE__INST4_SEG5 = 0
+MP0_BASE__INST5_SEG0 = 0
+MP0_BASE__INST5_SEG1 = 0
+MP0_BASE__INST5_SEG2 = 0
+MP0_BASE__INST5_SEG3 = 0
+MP0_BASE__INST5_SEG4 = 0
+MP0_BASE__INST5_SEG5 = 0
+MP1_BASE__INST0_SEG0 = 0x00016000
+MP1_BASE__INST0_SEG1 = 0
+MP1_BASE__INST0_SEG2 = 0
+MP1_BASE__INST0_SEG3 = 0
+MP1_BASE__INST0_SEG4 = 0
+MP1_BASE__INST0_SEG5 = 0
+MP1_BASE__INST1_SEG0 = 0
+MP1_BASE__INST1_SEG1 = 0
+MP1_BASE__INST1_SEG2 = 0
+MP1_BASE__INST1_SEG3 = 0
+MP1_BASE__INST1_SEG4 = 0
+MP1_BASE__INST1_SEG5 = 0
+MP1_BASE__INST2_SEG0 = 0
+MP1_BASE__INST2_SEG1 = 0
+MP1_BASE__INST2_SEG2 = 0
+MP1_BASE__INST2_SEG3 = 0
+MP1_BASE__INST2_SEG4 = 0
+MP1_BASE__INST2_SEG5 = 0
+MP1_BASE__INST3_SEG0 = 0
+MP1_BASE__INST3_SEG1 = 0
+MP1_BASE__INST3_SEG2 = 0
+MP1_BASE__INST3_SEG3 = 0
+MP1_BASE__INST3_SEG4 = 0
+MP1_BASE__INST3_SEG5 = 0
+MP1_BASE__INST4_SEG0 = 0
+MP1_BASE__INST4_SEG1 = 0
+MP1_BASE__INST4_SEG2 = 0
+MP1_BASE__INST4_SEG3 = 0
+MP1_BASE__INST4_SEG4 = 0
+MP1_BASE__INST4_SEG5 = 0
+MP1_BASE__INST5_SEG0 = 0
+MP1_BASE__INST5_SEG1 = 0
+MP1_BASE__INST5_SEG2 = 0
+MP1_BASE__INST5_SEG3 = 0
+MP1_BASE__INST5_SEG4 = 0
+MP1_BASE__INST5_SEG5 = 0
+NBIO_BASE__INST0_SEG0 = 0x00000000
+NBIO_BASE__INST0_SEG1 = 0x00000014
+NBIO_BASE__INST0_SEG2 = 0x00000D20
+NBIO_BASE__INST0_SEG3 = 0x00010400
+NBIO_BASE__INST0_SEG4 = 0
+NBIO_BASE__INST0_SEG5 = 0
+NBIO_BASE__INST1_SEG0 = 0
+NBIO_BASE__INST1_SEG1 = 0
+NBIO_BASE__INST1_SEG2 = 0
+NBIO_BASE__INST1_SEG3 = 0
+NBIO_BASE__INST1_SEG4 = 0
+NBIO_BASE__INST1_SEG5 = 0
+NBIO_BASE__INST2_SEG0 = 0
+NBIO_BASE__INST2_SEG1 = 0
+NBIO_BASE__INST2_SEG2 = 0
+NBIO_BASE__INST2_SEG3 = 0
+NBIO_BASE__INST2_SEG4 = 0
+NBIO_BASE__INST2_SEG5 = 0
+NBIO_BASE__INST3_SEG0 = 0
+NBIO_BASE__INST3_SEG1 = 0
+NBIO_BASE__INST3_SEG2 = 0
+NBIO_BASE__INST3_SEG3 = 0
+NBIO_BASE__INST3_SEG4 = 0
+NBIO_BASE__INST3_SEG5 = 0
+NBIO_BASE__INST4_SEG0 = 0
+NBIO_BASE__INST4_SEG1 = 0
+NBIO_BASE__INST4_SEG2 = 0
+NBIO_BASE__INST4_SEG3 = 0
+NBIO_BASE__INST4_SEG4 = 0
+NBIO_BASE__INST4_SEG5 = 0
+NBIO_BASE__INST5_SEG0 = 0
+NBIO_BASE__INST5_SEG1 = 0
+NBIO_BASE__INST5_SEG2 = 0
+NBIO_BASE__INST5_SEG3 = 0
+NBIO_BASE__INST5_SEG4 = 0
+NBIO_BASE__INST5_SEG5 = 0
+OSSSYS_BASE__INST0_SEG0 = 0x000010A0
+OSSSYS_BASE__INST0_SEG1 = 0
+OSSSYS_BASE__INST0_SEG2 = 0
+OSSSYS_BASE__INST0_SEG3 = 0
+OSSSYS_BASE__INST0_SEG4 = 0
+OSSSYS_BASE__INST0_SEG5 = 0
+OSSSYS_BASE__INST1_SEG0 = 0
+OSSSYS_BASE__INST1_SEG1 = 0
+OSSSYS_BASE__INST1_SEG2 = 0
+OSSSYS_BASE__INST1_SEG3 = 0
+OSSSYS_BASE__INST1_SEG4 = 0
+OSSSYS_BASE__INST1_SEG5 = 0
+OSSSYS_BASE__INST2_SEG0 = 0
+OSSSYS_BASE__INST2_SEG1 = 0
+OSSSYS_BASE__INST2_SEG2 = 0
+OSSSYS_BASE__INST2_SEG3 = 0
+OSSSYS_BASE__INST2_SEG4 = 0
+OSSSYS_BASE__INST2_SEG5 = 0
+OSSSYS_BASE__INST3_SEG0 = 0
+OSSSYS_BASE__INST3_SEG1 = 0
+OSSSYS_BASE__INST3_SEG2 = 0
+OSSSYS_BASE__INST3_SEG3 = 0
+OSSSYS_BASE__INST3_SEG4 = 0
+OSSSYS_BASE__INST3_SEG5 = 0
+OSSSYS_BASE__INST4_SEG0 = 0
+OSSSYS_BASE__INST4_SEG1 = 0
+OSSSYS_BASE__INST4_SEG2 = 0
+OSSSYS_BASE__INST4_SEG3 = 0
+OSSSYS_BASE__INST4_SEG4 = 0
+OSSSYS_BASE__INST4_SEG5 = 0
+OSSSYS_BASE__INST5_SEG0 = 0
+OSSSYS_BASE__INST5_SEG1 = 0
+OSSSYS_BASE__INST5_SEG2 = 0
+OSSSYS_BASE__INST5_SEG3 = 0
+OSSSYS_BASE__INST5_SEG4 = 0
+OSSSYS_BASE__INST5_SEG5 = 0
+SDMA0_BASE__INST0_SEG0 = 0x00001260
+SDMA0_BASE__INST0_SEG1 = 0
+SDMA0_BASE__INST0_SEG2 = 0
+SDMA0_BASE__INST0_SEG3 = 0
+SDMA0_BASE__INST0_SEG4 = 0
+SDMA0_BASE__INST0_SEG5 = 0
+SDMA0_BASE__INST1_SEG0 = 0
+SDMA0_BASE__INST1_SEG1 = 0
+SDMA0_BASE__INST1_SEG2 = 0
+SDMA0_BASE__INST1_SEG3 = 0
+SDMA0_BASE__INST1_SEG4 = 0
+SDMA0_BASE__INST1_SEG5 = 0
+SDMA0_BASE__INST2_SEG0 = 0
+SDMA0_BASE__INST2_SEG1 = 0
+SDMA0_BASE__INST2_SEG2 = 0
+SDMA0_BASE__INST2_SEG3 = 0
+SDMA0_BASE__INST2_SEG4 = 0
+SDMA0_BASE__INST2_SEG5 = 0
+SDMA0_BASE__INST3_SEG0 = 0
+SDMA0_BASE__INST3_SEG1 = 0
+SDMA0_BASE__INST3_SEG2 = 0
+SDMA0_BASE__INST3_SEG3 = 0
+SDMA0_BASE__INST3_SEG4 = 0
+SDMA0_BASE__INST3_SEG5 = 0
+SDMA0_BASE__INST4_SEG0 = 0
+SDMA0_BASE__INST4_SEG1 = 0
+SDMA0_BASE__INST4_SEG2 = 0
+SDMA0_BASE__INST4_SEG3 = 0
+SDMA0_BASE__INST4_SEG4 = 0
+SDMA0_BASE__INST4_SEG5 = 0
+SDMA0_BASE__INST5_SEG0 = 0
+SDMA0_BASE__INST5_SEG1 = 0
+SDMA0_BASE__INST5_SEG2 = 0
+SDMA0_BASE__INST5_SEG3 = 0
+SDMA0_BASE__INST5_SEG4 = 0
+SDMA0_BASE__INST5_SEG5 = 0
+SDMA1_BASE__INST0_SEG0 = 0x00001860
+SDMA1_BASE__INST0_SEG1 = 0
+SDMA1_BASE__INST0_SEG2 = 0
+SDMA1_BASE__INST0_SEG3 = 0
+SDMA1_BASE__INST0_SEG4 = 0
+SDMA1_BASE__INST0_SEG5 = 0
+SDMA1_BASE__INST1_SEG0 = 0
+SDMA1_BASE__INST1_SEG1 = 0
+SDMA1_BASE__INST1_SEG2 = 0
+SDMA1_BASE__INST1_SEG3 = 0
+SDMA1_BASE__INST1_SEG4 = 0
+SDMA1_BASE__INST1_SEG5 = 0
+SDMA1_BASE__INST2_SEG0 = 0
+SDMA1_BASE__INST2_SEG1 = 0
+SDMA1_BASE__INST2_SEG2 = 0
+SDMA1_BASE__INST2_SEG3 = 0
+SDMA1_BASE__INST2_SEG4 = 0
+SDMA1_BASE__INST2_SEG5 = 0
+SDMA1_BASE__INST3_SEG0 = 0
+SDMA1_BASE__INST3_SEG1 = 0
+SDMA1_BASE__INST3_SEG2 = 0
+SDMA1_BASE__INST3_SEG3 = 0
+SDMA1_BASE__INST3_SEG4 = 0
+SDMA1_BASE__INST3_SEG5 = 0
+SDMA1_BASE__INST4_SEG0 = 0
+SDMA1_BASE__INST4_SEG1 = 0
+SDMA1_BASE__INST4_SEG2 = 0
+SDMA1_BASE__INST4_SEG3 = 0
+SDMA1_BASE__INST4_SEG4 = 0
+SDMA1_BASE__INST4_SEG5 = 0
+SDMA1_BASE__INST5_SEG0 = 0
+SDMA1_BASE__INST5_SEG1 = 0
+SDMA1_BASE__INST5_SEG2 = 0
+SDMA1_BASE__INST5_SEG3 = 0
+SDMA1_BASE__INST5_SEG4 = 0
+SDMA1_BASE__INST5_SEG5 = 0
+SMUIO_BASE__INST0_SEG0 = 0x00016800
+SMUIO_BASE__INST0_SEG1 = 0x00016A00
+SMUIO_BASE__INST0_SEG2 = 0
+SMUIO_BASE__INST0_SEG3 = 0
+SMUIO_BASE__INST0_SEG4 = 0
+SMUIO_BASE__INST0_SEG5 = 0
+SMUIO_BASE__INST1_SEG0 = 0
+SMUIO_BASE__INST1_SEG1 = 0
+SMUIO_BASE__INST1_SEG2 = 0
+SMUIO_BASE__INST1_SEG3 = 0
+SMUIO_BASE__INST1_SEG4 = 0
+SMUIO_BASE__INST1_SEG5 = 0
+SMUIO_BASE__INST2_SEG0 = 0
+SMUIO_BASE__INST2_SEG1 = 0
+SMUIO_BASE__INST2_SEG2 = 0
+SMUIO_BASE__INST2_SEG3 = 0
+SMUIO_BASE__INST2_SEG4 = 0
+SMUIO_BASE__INST2_SEG5 = 0
+SMUIO_BASE__INST3_SEG0 = 0
+SMUIO_BASE__INST3_SEG1 = 0
+SMUIO_BASE__INST3_SEG2 = 0
+SMUIO_BASE__INST3_SEG3 = 0
+SMUIO_BASE__INST3_SEG4 = 0
+SMUIO_BASE__INST3_SEG5 = 0
+SMUIO_BASE__INST4_SEG0 = 0
+SMUIO_BASE__INST4_SEG1 = 0
+SMUIO_BASE__INST4_SEG2 = 0
+SMUIO_BASE__INST4_SEG3 = 0
+SMUIO_BASE__INST4_SEG4 = 0
+SMUIO_BASE__INST4_SEG5 = 0
+SMUIO_BASE__INST5_SEG0 = 0
+SMUIO_BASE__INST5_SEG1 = 0
+SMUIO_BASE__INST5_SEG2 = 0
+SMUIO_BASE__INST5_SEG3 = 0
+SMUIO_BASE__INST5_SEG4 = 0
+SMUIO_BASE__INST5_SEG5 = 0
+THM_BASE__INST0_SEG0 = 0x00016600
+THM_BASE__INST0_SEG1 = 0
+THM_BASE__INST0_SEG2 = 0
+THM_BASE__INST0_SEG3 = 0
+THM_BASE__INST0_SEG4 = 0
+THM_BASE__INST0_SEG5 = 0
+THM_BASE__INST1_SEG0 = 0
+THM_BASE__INST1_SEG1 = 0
+THM_BASE__INST1_SEG2 = 0
+THM_BASE__INST1_SEG3 = 0
+THM_BASE__INST1_SEG4 = 0
+THM_BASE__INST1_SEG5 = 0
+THM_BASE__INST2_SEG0 = 0
+THM_BASE__INST2_SEG1 = 0
+THM_BASE__INST2_SEG2 = 0
+THM_BASE__INST2_SEG3 = 0
+THM_BASE__INST2_SEG4 = 0
+THM_BASE__INST2_SEG5 = 0
+THM_BASE__INST3_SEG0 = 0
+THM_BASE__INST3_SEG1 = 0
+THM_BASE__INST3_SEG2 = 0
+THM_BASE__INST3_SEG3 = 0
+THM_BASE__INST3_SEG4 = 0
+THM_BASE__INST3_SEG5 = 0
+THM_BASE__INST4_SEG0 = 0
+THM_BASE__INST4_SEG1 = 0
+THM_BASE__INST4_SEG2 = 0
+THM_BASE__INST4_SEG3 = 0
+THM_BASE__INST4_SEG4 = 0
+THM_BASE__INST4_SEG5 = 0
+THM_BASE__INST5_SEG0 = 0
+THM_BASE__INST5_SEG1 = 0
+THM_BASE__INST5_SEG2 = 0
+THM_BASE__INST5_SEG3 = 0
+THM_BASE__INST5_SEG4 = 0
+THM_BASE__INST5_SEG5 = 0
+UMC_BASE__INST0_SEG0 = 0x00014000
+UMC_BASE__INST0_SEG1 = 0
+UMC_BASE__INST0_SEG2 = 0
+UMC_BASE__INST0_SEG3 = 0
+UMC_BASE__INST0_SEG4 = 0
+UMC_BASE__INST0_SEG5 = 0
+UMC_BASE__INST1_SEG0 = 0
+UMC_BASE__INST1_SEG1 = 0
+UMC_BASE__INST1_SEG2 = 0
+UMC_BASE__INST1_SEG3 = 0
+UMC_BASE__INST1_SEG4 = 0
+UMC_BASE__INST1_SEG5 = 0
+UMC_BASE__INST2_SEG0 = 0
+UMC_BASE__INST2_SEG1 = 0
+UMC_BASE__INST2_SEG2 = 0
+UMC_BASE__INST2_SEG3 = 0
+UMC_BASE__INST2_SEG4 = 0
+UMC_BASE__INST2_SEG5 = 0
+UMC_BASE__INST3_SEG0 = 0
+UMC_BASE__INST3_SEG1 = 0
+UMC_BASE__INST3_SEG2 = 0
+UMC_BASE__INST3_SEG3 = 0
+UMC_BASE__INST3_SEG4 = 0
+UMC_BASE__INST3_SEG5 = 0
+UMC_BASE__INST4_SEG0 = 0
+UMC_BASE__INST4_SEG1 = 0
+UMC_BASE__INST4_SEG2 = 0
+UMC_BASE__INST4_SEG3 = 0
+UMC_BASE__INST4_SEG4 = 0
+UMC_BASE__INST4_SEG5 = 0
+UMC_BASE__INST5_SEG0 = 0
+UMC_BASE__INST5_SEG1 = 0
+UMC_BASE__INST5_SEG2 = 0
+UMC_BASE__INST5_SEG3 = 0
+UMC_BASE__INST5_SEG4 = 0
+UMC_BASE__INST5_SEG5 = 0
+UVD_BASE__INST0_SEG0 = 0x00007800
+UVD_BASE__INST0_SEG1 = 0x00007E00
+UVD_BASE__INST0_SEG2 = 0
+UVD_BASE__INST0_SEG3 = 0
+UVD_BASE__INST0_SEG4 = 0
+UVD_BASE__INST0_SEG5 = 0
+UVD_BASE__INST1_SEG0 = 0
+UVD_BASE__INST1_SEG1 = 0x00009000
+UVD_BASE__INST1_SEG2 = 0
+UVD_BASE__INST1_SEG3 = 0
+UVD_BASE__INST1_SEG4 = 0
+UVD_BASE__INST1_SEG5 = 0
+UVD_BASE__INST2_SEG0 = 0
+UVD_BASE__INST2_SEG1 = 0
+UVD_BASE__INST2_SEG2 = 0
+UVD_BASE__INST2_SEG3 = 0
+UVD_BASE__INST2_SEG4 = 0
+UVD_BASE__INST2_SEG5 = 0
+UVD_BASE__INST3_SEG0 = 0
+UVD_BASE__INST3_SEG1 = 0
+UVD_BASE__INST3_SEG2 = 0
+UVD_BASE__INST3_SEG3 = 0
+UVD_BASE__INST3_SEG4 = 0
+UVD_BASE__INST3_SEG5 = 0
+UVD_BASE__INST4_SEG0 = 0
+UVD_BASE__INST4_SEG1 = 0
+UVD_BASE__INST4_SEG2 = 0
+UVD_BASE__INST4_SEG3 = 0
+UVD_BASE__INST4_SEG4 = 0
+UVD_BASE__INST4_SEG5 = 0
+UVD_BASE__INST5_SEG0 = 0
+UVD_BASE__INST5_SEG1 = 0
+UVD_BASE__INST5_SEG2 = 0
+UVD_BASE__INST5_SEG3 = 0
+UVD_BASE__INST5_SEG4 = 0
+UVD_BASE__INST5_SEG5 = 0
+VCE_BASE__INST0_SEG0 = 0x00008800
+VCE_BASE__INST0_SEG1 = 0
+VCE_BASE__INST0_SEG2 = 0
+VCE_BASE__INST0_SEG3 = 0
+VCE_BASE__INST0_SEG4 = 0
+VCE_BASE__INST0_SEG5 = 0
+VCE_BASE__INST1_SEG0 = 0
+VCE_BASE__INST1_SEG1 = 0
+VCE_BASE__INST1_SEG2 = 0
+VCE_BASE__INST1_SEG3 = 0
+VCE_BASE__INST1_SEG4 = 0
+VCE_BASE__INST1_SEG5 = 0
+VCE_BASE__INST2_SEG0 = 0
+VCE_BASE__INST2_SEG1 = 0
+VCE_BASE__INST2_SEG2 = 0
+VCE_BASE__INST2_SEG3 = 0
+VCE_BASE__INST2_SEG4 = 0
+VCE_BASE__INST2_SEG5 = 0
+VCE_BASE__INST3_SEG0 = 0
+VCE_BASE__INST3_SEG1 = 0
+VCE_BASE__INST3_SEG2 = 0
+VCE_BASE__INST3_SEG3 = 0
+VCE_BASE__INST3_SEG4 = 0
+VCE_BASE__INST3_SEG5 = 0
+VCE_BASE__INST4_SEG0 = 0
+VCE_BASE__INST4_SEG1 = 0
+VCE_BASE__INST4_SEG2 = 0
+VCE_BASE__INST4_SEG3 = 0
+VCE_BASE__INST4_SEG4 = 0
+VCE_BASE__INST4_SEG5 = 0
+VCE_BASE__INST5_SEG0 = 0
+VCE_BASE__INST5_SEG1 = 0
+VCE_BASE__INST5_SEG2 = 0
+VCE_BASE__INST5_SEG3 = 0
+VCE_BASE__INST5_SEG4 = 0
+VCE_BASE__INST5_SEG5 = 0
+XDMA_BASE__INST0_SEG0 = 0x00003400
+XDMA_BASE__INST0_SEG1 = 0
+XDMA_BASE__INST0_SEG2 = 0
+XDMA_BASE__INST0_SEG3 = 0
+XDMA_BASE__INST0_SEG4 = 0
+XDMA_BASE__INST0_SEG5 = 0
+XDMA_BASE__INST1_SEG0 = 0
+XDMA_BASE__INST1_SEG1 = 0
+XDMA_BASE__INST1_SEG2 = 0
+XDMA_BASE__INST1_SEG3 = 0
+XDMA_BASE__INST1_SEG4 = 0
+XDMA_BASE__INST1_SEG5 = 0
+XDMA_BASE__INST2_SEG0 = 0
+XDMA_BASE__INST2_SEG1 = 0
+XDMA_BASE__INST2_SEG2 = 0
+XDMA_BASE__INST2_SEG3 = 0
+XDMA_BASE__INST2_SEG4 = 0
+XDMA_BASE__INST2_SEG5 = 0
+XDMA_BASE__INST3_SEG0 = 0
+XDMA_BASE__INST3_SEG1 = 0
+XDMA_BASE__INST3_SEG2 = 0
+XDMA_BASE__INST3_SEG3 = 0
+XDMA_BASE__INST3_SEG4 = 0
+XDMA_BASE__INST3_SEG5 = 0
+XDMA_BASE__INST4_SEG0 = 0
+XDMA_BASE__INST4_SEG1 = 0
+XDMA_BASE__INST4_SEG2 = 0
+XDMA_BASE__INST4_SEG3 = 0
+XDMA_BASE__INST4_SEG4 = 0
+XDMA_BASE__INST4_SEG5 = 0
+XDMA_BASE__INST5_SEG0 = 0
+XDMA_BASE__INST5_SEG1 = 0
+XDMA_BASE__INST5_SEG2 = 0
+XDMA_BASE__INST5_SEG3 = 0
+XDMA_BASE__INST5_SEG4 = 0
+XDMA_BASE__INST5_SEG5 = 0
+RSMU_BASE__INST0_SEG0 = 0x00012000
+RSMU_BASE__INST0_SEG1 = 0
+RSMU_BASE__INST0_SEG2 = 0
+RSMU_BASE__INST0_SEG3 = 0
+RSMU_BASE__INST0_SEG4 = 0
+RSMU_BASE__INST0_SEG5 = 0
+RSMU_BASE__INST1_SEG0 = 0
+RSMU_BASE__INST1_SEG1 = 0
+RSMU_BASE__INST1_SEG2 = 0
+RSMU_BASE__INST1_SEG3 = 0
+RSMU_BASE__INST1_SEG4 = 0
+RSMU_BASE__INST1_SEG5 = 0
+RSMU_BASE__INST2_SEG0 = 0
+RSMU_BASE__INST2_SEG1 = 0
+RSMU_BASE__INST2_SEG2 = 0
+RSMU_BASE__INST2_SEG3 = 0
+RSMU_BASE__INST2_SEG4 = 0
+RSMU_BASE__INST2_SEG5 = 0
+RSMU_BASE__INST3_SEG0 = 0
+RSMU_BASE__INST3_SEG1 = 0
+RSMU_BASE__INST3_SEG2 = 0
+RSMU_BASE__INST3_SEG3 = 0
+RSMU_BASE__INST3_SEG4 = 0
+RSMU_BASE__INST3_SEG5 = 0
+RSMU_BASE__INST4_SEG0 = 0
+RSMU_BASE__INST4_SEG1 = 0
+RSMU_BASE__INST4_SEG2 = 0
+RSMU_BASE__INST4_SEG3 = 0
+RSMU_BASE__INST4_SEG4 = 0
+RSMU_BASE__INST4_SEG5 = 0
+RSMU_BASE__INST5_SEG0 = 0
+RSMU_BASE__INST5_SEG1 = 0
+RSMU_BASE__INST5_SEG2 = 0
+RSMU_BASE__INST5_SEG3 = 0
+RSMU_BASE__INST5_SEG4 = 0
+RSMU_BASE__INST5_SEG5 = 0

tinygrad/runtime/autogen/comgr.py

@@ -201,5 +201,5 @@ def amd_comgr_lookup_code_object(data:amd_comgr_data_t, info_list:c.POINTER[amd_
 @dll.bind(amd_comgr_status_t, amd_comgr_data_t, uint64_t, c.POINTER[uint64_t], c.POINTER[uint64_t], c.POINTER[ctypes.c_bool])
 def amd_comgr_map_elf_virtual_address_to_code_object_offset(data:amd_comgr_data_t, elf_virtual_address:uint64_t, code_object_offset:c.POINTER[uint64_t], slice_size:c.POINTER[uint64_t], nobits:c.POINTER[ctypes.c_bool]) -> amd_comgr_status_t: ...
 AMD_COMGR_DEPRECATED = lambda msg: __attribute__((deprecated(msg))) # type: ignore
-AMD_COMGR_INTERFACE_VERSION_MAJOR = 2 # type: ignore
-AMD_COMGR_INTERFACE_VERSION_MINOR = 8 # type: ignore
+AMD_COMGR_INTERFACE_VERSION_MAJOR = 2
+AMD_COMGR_INTERFACE_VERSION_MINOR = 8

tinygrad/runtime/autogen/comgr_3.py

@@ -201,5 +201,5 @@ def amd_comgr_lookup_code_object(data:amd_comgr_data_t, info_list:c.POINTER[amd_
 @dll.bind(amd_comgr_status_t, amd_comgr_data_t, uint64_t, c.POINTER[uint64_t], c.POINTER[uint64_t], c.POINTER[ctypes.c_bool])
 def amd_comgr_map_elf_virtual_address_to_code_object_offset(data:amd_comgr_data_t, elf_virtual_address:uint64_t, code_object_offset:c.POINTER[uint64_t], slice_size:c.POINTER[uint64_t], nobits:c.POINTER[ctypes.c_bool]) -> amd_comgr_status_t: ...
 AMD_COMGR_DEPRECATED = lambda msg: __attribute__((deprecated(msg))) # type: ignore
-AMD_COMGR_INTERFACE_VERSION_MAJOR = 3 # type: ignore
-AMD_COMGR_INTERFACE_VERSION_MINOR = 0 # type: ignore
+AMD_COMGR_INTERFACE_VERSION_MAJOR = 3
+AMD_COMGR_INTERFACE_VERSION_MINOR = 0

tinygrad/runtime/autogen/corefoundation.py

@@ -305,9 +305,9 @@ CFDataSearchFlags: TypeAlias = ctypes.c_uint64
 _anonenum3: dict[int, str] = {(kCFDataSearchBackwards:=1): 'kCFDataSearchBackwards', (kCFDataSearchAnchored:=2): 'kCFDataSearchAnchored'}
 @dll.bind(CFRange, CFDataRef, CFDataRef, CFRange, CFDataSearchFlags)
 def CFDataFind(theData:CFDataRef, dataToFind:CFDataRef, searchRange:CFRange, compareOptions:CFDataSearchFlags) -> CFRange: ...
-__COREFOUNDATION_CFSTRING__ = 1 # type: ignore
-kCFStringEncodingInvalidId = (0xffffffff) # type: ignore
+__COREFOUNDATION_CFSTRING__ = 1
+kCFStringEncodingInvalidId = (0xffffffff)
 CF_FORMAT_FUNCTION = lambda F,A: __attribute__((format(CFString, F, A))) # type: ignore
 CF_FORMAT_ARGUMENT = lambda A: __attribute__((format_arg(A))) # type: ignore
-__kCFStringInlineBufferLength = 64 # type: ignore
-__COREFOUNDATION_CFDATA__ = 1 # type: ignore
+__kCFStringInlineBufferLength = 64
+__COREFOUNDATION_CFDATA__ = 1

tinygrad/runtime/autogen/hip.py

@@ -2097,145 +2097,145 @@ def hipDestroySurfaceObject(surfaceObject:hipSurfaceObject_t) -> ctypes.c_uint32
 hipmipmappedArray: TypeAlias = c.POINTER[hipMipmappedArray]
 hipResourcetype: TypeAlias = ctypes.c_uint32
 hipMemcpyFlags: dict[int, str] = {(hipMemcpyFlagDefault:=0): 'hipMemcpyFlagDefault', (hipMemcpyFlagPreferOverlapWithCompute:=1): 'hipMemcpyFlagPreferOverlapWithCompute'}
-hiprtcJIT_option = hipJitOption # type: ignore
-HIPRTC_JIT_MAX_REGISTERS = hipJitOptionMaxRegisters # type: ignore
-HIPRTC_JIT_THREADS_PER_BLOCK = hipJitOptionThreadsPerBlock # type: ignore
-HIPRTC_JIT_WALL_TIME = hipJitOptionWallTime # type: ignore
-HIPRTC_JIT_INFO_LOG_BUFFER = hipJitOptionInfoLogBuffer # type: ignore
-HIPRTC_JIT_INFO_LOG_BUFFER_SIZE_BYTES = hipJitOptionInfoLogBufferSizeBytes # type: ignore
-HIPRTC_JIT_ERROR_LOG_BUFFER = hipJitOptionErrorLogBuffer # type: ignore
-HIPRTC_JIT_ERROR_LOG_BUFFER_SIZE_BYTES = hipJitOptionErrorLogBufferSizeBytes # type: ignore
-HIPRTC_JIT_OPTIMIZATION_LEVEL = hipJitOptionOptimizationLevel # type: ignore
-HIPRTC_JIT_TARGET_FROM_HIPCONTEXT = hipJitOptionTargetFromContext # type: ignore
-HIPRTC_JIT_TARGET = hipJitOptionTarget # type: ignore
-HIPRTC_JIT_FALLBACK_STRATEGY = hipJitOptionFallbackStrategy # type: ignore
-HIPRTC_JIT_GENERATE_DEBUG_INFO = hipJitOptionGenerateDebugInfo # type: ignore
-HIPRTC_JIT_LOG_VERBOSE = hipJitOptionLogVerbose # type: ignore
-HIPRTC_JIT_GENERATE_LINE_INFO = hipJitOptionGenerateLineInfo # type: ignore
-HIPRTC_JIT_CACHE_MODE = hipJitOptionCacheMode # type: ignore
-HIPRTC_JIT_NEW_SM3X_OPT = hipJitOptionSm3xOpt # type: ignore
-HIPRTC_JIT_FAST_COMPILE = hipJitOptionFastCompile # type: ignore
-HIPRTC_JIT_GLOBAL_SYMBOL_NAMES = hipJitOptionGlobalSymbolNames # type: ignore
-HIPRTC_JIT_GLOBAL_SYMBOL_ADDRESS = hipJitOptionGlobalSymbolAddresses # type: ignore
-HIPRTC_JIT_GLOBAL_SYMBOL_COUNT = hipJitOptionGlobalSymbolCount # type: ignore
-HIPRTC_JIT_LTO = hipJitOptionLto # type: ignore
-HIPRTC_JIT_FTZ = hipJitOptionFtz # type: ignore
-HIPRTC_JIT_PREC_DIV = hipJitOptionPrecDiv # type: ignore
-HIPRTC_JIT_PREC_SQRT = hipJitOptionPrecSqrt # type: ignore
-HIPRTC_JIT_FMA = hipJitOptionFma # type: ignore
-HIPRTC_JIT_POSITION_INDEPENDENT_CODE = hipJitOptionPositionIndependentCode # type: ignore
-HIPRTC_JIT_MIN_CTA_PER_SM = hipJitOptionMinCTAPerSM # type: ignore
-HIPRTC_JIT_MAX_THREADS_PER_BLOCK = hipJitOptionMaxThreadsPerBlock # type: ignore
-HIPRTC_JIT_OVERRIDE_DIRECT_VALUES = hipJitOptionOverrideDirectiveValues # type: ignore
-HIPRTC_JIT_NUM_OPTIONS = hipJitOptionNumOptions # type: ignore
-HIPRTC_JIT_IR_TO_ISA_OPT_EXT = hipJitOptionIRtoISAOptExt # type: ignore
-HIPRTC_JIT_IR_TO_ISA_OPT_COUNT_EXT = hipJitOptionIRtoISAOptCountExt # type: ignore
-hiprtcJITInputType = hipJitInputType # type: ignore
-HIPRTC_JIT_INPUT_CUBIN = hipJitInputCubin # type: ignore
-HIPRTC_JIT_INPUT_PTX = hipJitInputPtx # type: ignore
-HIPRTC_JIT_INPUT_FATBINARY = hipJitInputFatBinary # type: ignore
-HIPRTC_JIT_INPUT_OBJECT = hipJitInputObject # type: ignore
-HIPRTC_JIT_INPUT_LIBRARY = hipJitInputLibrary # type: ignore
-HIPRTC_JIT_INPUT_NVVM = hipJitInputNvvm # type: ignore
-HIPRTC_JIT_NUM_LEGACY_INPUT_TYPES = hipJitNumLegacyInputTypes # type: ignore
-HIPRTC_JIT_INPUT_LLVM_BITCODE = hipJitInputLLVMBitcode # type: ignore
-HIPRTC_JIT_INPUT_LLVM_BUNDLED_BITCODE = hipJitInputLLVMBundledBitcode # type: ignore
-HIPRTC_JIT_INPUT_LLVM_ARCHIVES_OF_BUNDLED_BITCODE = hipJitInputLLVMArchivesOfBundledBitcode # type: ignore
-HIPRTC_JIT_INPUT_SPIRV = hipJitInputSpirv # type: ignore
-HIPRTC_JIT_NUM_INPUT_TYPES = hipJitNumInputTypes # type: ignore
-hipGetDeviceProperties = hipGetDevicePropertiesR0600 # type: ignore
-hipDeviceProp_t = hipDeviceProp_tR0600 # type: ignore
-hipChooseDevice = hipChooseDeviceR0600 # type: ignore
-GENERIC_GRID_LAUNCH = 1 # type: ignore
+hiprtcJIT_option = hipJitOption
+HIPRTC_JIT_MAX_REGISTERS = hipJitOptionMaxRegisters
+HIPRTC_JIT_THREADS_PER_BLOCK = hipJitOptionThreadsPerBlock
+HIPRTC_JIT_WALL_TIME = hipJitOptionWallTime
+HIPRTC_JIT_INFO_LOG_BUFFER = hipJitOptionInfoLogBuffer
+HIPRTC_JIT_INFO_LOG_BUFFER_SIZE_BYTES = hipJitOptionInfoLogBufferSizeBytes
+HIPRTC_JIT_ERROR_LOG_BUFFER = hipJitOptionErrorLogBuffer
+HIPRTC_JIT_ERROR_LOG_BUFFER_SIZE_BYTES = hipJitOptionErrorLogBufferSizeBytes
+HIPRTC_JIT_OPTIMIZATION_LEVEL = hipJitOptionOptimizationLevel
+HIPRTC_JIT_TARGET_FROM_HIPCONTEXT = hipJitOptionTargetFromContext
+HIPRTC_JIT_TARGET = hipJitOptionTarget
+HIPRTC_JIT_FALLBACK_STRATEGY = hipJitOptionFallbackStrategy
+HIPRTC_JIT_GENERATE_DEBUG_INFO = hipJitOptionGenerateDebugInfo
+HIPRTC_JIT_LOG_VERBOSE = hipJitOptionLogVerbose
+HIPRTC_JIT_GENERATE_LINE_INFO = hipJitOptionGenerateLineInfo
+HIPRTC_JIT_CACHE_MODE = hipJitOptionCacheMode
+HIPRTC_JIT_NEW_SM3X_OPT = hipJitOptionSm3xOpt
+HIPRTC_JIT_FAST_COMPILE = hipJitOptionFastCompile
+HIPRTC_JIT_GLOBAL_SYMBOL_NAMES = hipJitOptionGlobalSymbolNames
+HIPRTC_JIT_GLOBAL_SYMBOL_ADDRESS = hipJitOptionGlobalSymbolAddresses
+HIPRTC_JIT_GLOBAL_SYMBOL_COUNT = hipJitOptionGlobalSymbolCount
+HIPRTC_JIT_LTO = hipJitOptionLto
+HIPRTC_JIT_FTZ = hipJitOptionFtz
+HIPRTC_JIT_PREC_DIV = hipJitOptionPrecDiv
+HIPRTC_JIT_PREC_SQRT = hipJitOptionPrecSqrt
+HIPRTC_JIT_FMA = hipJitOptionFma
+HIPRTC_JIT_POSITION_INDEPENDENT_CODE = hipJitOptionPositionIndependentCode
+HIPRTC_JIT_MIN_CTA_PER_SM = hipJitOptionMinCTAPerSM
+HIPRTC_JIT_MAX_THREADS_PER_BLOCK = hipJitOptionMaxThreadsPerBlock
+HIPRTC_JIT_OVERRIDE_DIRECT_VALUES = hipJitOptionOverrideDirectiveValues
+HIPRTC_JIT_NUM_OPTIONS = hipJitOptionNumOptions
+HIPRTC_JIT_IR_TO_ISA_OPT_EXT = hipJitOptionIRtoISAOptExt
+HIPRTC_JIT_IR_TO_ISA_OPT_COUNT_EXT = hipJitOptionIRtoISAOptCountExt
+hiprtcJITInputType = hipJitInputType
+HIPRTC_JIT_INPUT_CUBIN = hipJitInputCubin
+HIPRTC_JIT_INPUT_PTX = hipJitInputPtx
+HIPRTC_JIT_INPUT_FATBINARY = hipJitInputFatBinary
+HIPRTC_JIT_INPUT_OBJECT = hipJitInputObject
+HIPRTC_JIT_INPUT_LIBRARY = hipJitInputLibrary
+HIPRTC_JIT_INPUT_NVVM = hipJitInputNvvm
+HIPRTC_JIT_NUM_LEGACY_INPUT_TYPES = hipJitNumLegacyInputTypes
+HIPRTC_JIT_INPUT_LLVM_BITCODE = hipJitInputLLVMBitcode
+HIPRTC_JIT_INPUT_LLVM_BUNDLED_BITCODE = hipJitInputLLVMBundledBitcode
+HIPRTC_JIT_INPUT_LLVM_ARCHIVES_OF_BUNDLED_BITCODE = hipJitInputLLVMArchivesOfBundledBitcode
+HIPRTC_JIT_INPUT_SPIRV = hipJitInputSpirv
+HIPRTC_JIT_NUM_INPUT_TYPES = hipJitNumInputTypes
+hipGetDeviceProperties = hipGetDevicePropertiesR0600
+hipDeviceProp_t = hipDeviceProp_tR0600
+hipChooseDevice = hipChooseDeviceR0600
+GENERIC_GRID_LAUNCH = 1
 HIP_DEPRECATED = lambda msg: __attribute__((deprecated(msg))) # type: ignore
-hipIpcMemLazyEnablePeerAccess = 0x01 # type: ignore
-HIP_IPC_HANDLE_SIZE = 64 # type: ignore
-hipStreamDefault = 0x00 # type: ignore
-hipStreamNonBlocking = 0x01 # type: ignore
-hipEventDefault = 0x0 # type: ignore
-hipEventBlockingSync = 0x1 # type: ignore
-hipEventDisableTiming = 0x2 # type: ignore
-hipEventInterprocess = 0x4 # type: ignore
-hipEventRecordDefault = 0x00 # type: ignore
-hipEventRecordExternal = 0x01 # type: ignore
-hipEventWaitDefault = 0x00 # type: ignore
-hipEventWaitExternal = 0x01 # type: ignore
-hipEventDisableSystemFence = 0x20000000 # type: ignore
-hipEventReleaseToDevice = 0x40000000 # type: ignore
-hipEventReleaseToSystem = 0x80000000 # type: ignore
-hipEnableDefault = 0x0 # type: ignore
-hipEnableLegacyStream = 0x1 # type: ignore
-hipEnablePerThreadDefaultStream = 0x2 # type: ignore
-hipHostAllocDefault = 0x0 # type: ignore
-hipHostMallocDefault = 0x0 # type: ignore
-hipHostAllocPortable = 0x1 # type: ignore
-hipHostMallocPortable = 0x1 # type: ignore
-hipHostAllocMapped = 0x2 # type: ignore
-hipHostMallocMapped = 0x2 # type: ignore
-hipHostAllocWriteCombined = 0x4 # type: ignore
-hipHostMallocWriteCombined = 0x4 # type: ignore
-hipHostMallocUncached = 0x10000000 # type: ignore
-hipHostAllocUncached = hipHostMallocUncached # type: ignore
-hipHostMallocNumaUser = 0x20000000 # type: ignore
-hipHostMallocCoherent = 0x40000000 # type: ignore
-hipHostMallocNonCoherent = 0x80000000 # type: ignore
-hipMemAttachGlobal = 0x01 # type: ignore
-hipMemAttachHost = 0x02 # type: ignore
-hipMemAttachSingle = 0x04 # type: ignore
-hipDeviceMallocDefault = 0x0 # type: ignore
-hipDeviceMallocFinegrained = 0x1 # type: ignore
-hipMallocSignalMemory = 0x2 # type: ignore
-hipDeviceMallocUncached = 0x3 # type: ignore
-hipDeviceMallocContiguous = 0x4 # type: ignore
-hipHostRegisterDefault = 0x0 # type: ignore
-hipHostRegisterPortable = 0x1 # type: ignore
-hipHostRegisterMapped = 0x2 # type: ignore
-hipHostRegisterIoMemory = 0x4 # type: ignore
-hipHostRegisterReadOnly = 0x08 # type: ignore
-hipExtHostRegisterCoarseGrained = 0x8 # type: ignore
-hipExtHostRegisterUncached = 0x80000000 # type: ignore
-hipDeviceScheduleAuto = 0x0 # type: ignore
-hipDeviceScheduleSpin = 0x1 # type: ignore
-hipDeviceScheduleYield = 0x2 # type: ignore
-hipDeviceScheduleBlockingSync = 0x4 # type: ignore
-hipDeviceScheduleMask = 0x7 # type: ignore
-hipDeviceMapHost = 0x8 # type: ignore
-hipDeviceLmemResizeToMax = 0x10 # type: ignore
-hipArrayDefault = 0x00 # type: ignore
-hipArrayLayered = 0x01 # type: ignore
-hipArraySurfaceLoadStore = 0x02 # type: ignore
-hipArrayCubemap = 0x04 # type: ignore
-hipArrayTextureGather = 0x08 # type: ignore
-hipOccupancyDefault = 0x00 # type: ignore
-hipOccupancyDisableCachingOverride = 0x01 # type: ignore
-hipCooperativeLaunchMultiDeviceNoPreSync = 0x01 # type: ignore
-hipCooperativeLaunchMultiDeviceNoPostSync = 0x02 # type: ignore
-hipExtAnyOrderLaunch = 0x01 # type: ignore
-hipStreamWaitValueGte = 0x0 # type: ignore
-hipStreamWaitValueEq = 0x1 # type: ignore
-hipStreamWaitValueAnd = 0x2 # type: ignore
-hipStreamWaitValueNor = 0x3 # type: ignore
-hipExternalMemoryDedicated = 0x1 # type: ignore
-hipStreamAttrID = hipLaunchAttributeID # type: ignore
-hipStreamAttributeAccessPolicyWindow = hipLaunchAttributeAccessPolicyWindow # type: ignore
-hipStreamAttributeSynchronizationPolicy = hipLaunchAttributeSynchronizationPolicy # type: ignore
-hipStreamAttributeMemSyncDomainMap = hipLaunchAttributeMemSyncDomainMap # type: ignore
-hipStreamAttributeMemSyncDomain = hipLaunchAttributeMemSyncDomain # type: ignore
-hipStreamAttributePriority = hipLaunchAttributePriority # type: ignore
-hipStreamAttrValue = hipLaunchAttributeValue # type: ignore
-hipKernelNodeAttrID = hipLaunchAttributeID # type: ignore
-hipKernelNodeAttributeAccessPolicyWindow = hipLaunchAttributeAccessPolicyWindow # type: ignore
-hipKernelNodeAttributeCooperative = hipLaunchAttributeCooperative # type: ignore
-hipKernelNodeAttributePriority = hipLaunchAttributePriority # type: ignore
-hipKernelNodeAttrValue = hipLaunchAttributeValue # type: ignore
-hipDrvLaunchAttributeCooperative = hipLaunchAttributeCooperative # type: ignore
-hipDrvLaunchAttributeID = hipLaunchAttributeID # type: ignore
-hipDrvLaunchAttributeValue = hipLaunchAttributeValue # type: ignore
-hipDrvLaunchAttribute = hipLaunchAttribute # type: ignore
-hipGraphKernelNodePortDefault = 0 # type: ignore
-hipGraphKernelNodePortLaunchCompletion = 2 # type: ignore
-hipGraphKernelNodePortProgrammatic = 1 # type: ignore
-HIP_TRSA_OVERRIDE_FORMAT = 0x01 # type: ignore
-HIP_TRSF_READ_AS_INTEGER = 0x01 # type: ignore
-HIP_TRSF_NORMALIZED_COORDINATES = 0x02 # type: ignore
-HIP_TRSF_SRGB = 0x10 # type: ignore
+hipIpcMemLazyEnablePeerAccess = 0x01
+HIP_IPC_HANDLE_SIZE = 64
+hipStreamDefault = 0x00
+hipStreamNonBlocking = 0x01
+hipEventDefault = 0x0
+hipEventBlockingSync = 0x1
+hipEventDisableTiming = 0x2
+hipEventInterprocess = 0x4
+hipEventRecordDefault = 0x00
+hipEventRecordExternal = 0x01
+hipEventWaitDefault = 0x00
+hipEventWaitExternal = 0x01
+hipEventDisableSystemFence = 0x20000000
+hipEventReleaseToDevice = 0x40000000
+hipEventReleaseToSystem = 0x80000000
+hipEnableDefault = 0x0
+hipEnableLegacyStream = 0x1
+hipEnablePerThreadDefaultStream = 0x2
+hipHostAllocDefault = 0x0
+hipHostMallocDefault = 0x0
+hipHostAllocPortable = 0x1
+hipHostMallocPortable = 0x1
+hipHostAllocMapped = 0x2
+hipHostMallocMapped = 0x2
+hipHostAllocWriteCombined = 0x4
+hipHostMallocWriteCombined = 0x4
+hipHostMallocUncached = 0x10000000
+hipHostAllocUncached = hipHostMallocUncached
+hipHostMallocNumaUser = 0x20000000
+hipHostMallocCoherent = 0x40000000
+hipHostMallocNonCoherent = 0x80000000
+hipMemAttachGlobal = 0x01
+hipMemAttachHost = 0x02
+hipMemAttachSingle = 0x04
+hipDeviceMallocDefault = 0x0
+hipDeviceMallocFinegrained = 0x1
+hipMallocSignalMemory = 0x2
+hipDeviceMallocUncached = 0x3
+hipDeviceMallocContiguous = 0x4
+hipHostRegisterDefault = 0x0
+hipHostRegisterPortable = 0x1
+hipHostRegisterMapped = 0x2
+hipHostRegisterIoMemory = 0x4
+hipHostRegisterReadOnly = 0x08
+hipExtHostRegisterCoarseGrained = 0x8
+hipExtHostRegisterUncached = 0x80000000
+hipDeviceScheduleAuto = 0x0
+hipDeviceScheduleSpin = 0x1
+hipDeviceScheduleYield = 0x2
+hipDeviceScheduleBlockingSync = 0x4
+hipDeviceScheduleMask = 0x7
+hipDeviceMapHost = 0x8
+hipDeviceLmemResizeToMax = 0x10
+hipArrayDefault = 0x00
+hipArrayLayered = 0x01
+hipArraySurfaceLoadStore = 0x02
+hipArrayCubemap = 0x04
+hipArrayTextureGather = 0x08
+hipOccupancyDefault = 0x00
+hipOccupancyDisableCachingOverride = 0x01
+hipCooperativeLaunchMultiDeviceNoPreSync = 0x01
+hipCooperativeLaunchMultiDeviceNoPostSync = 0x02
+hipExtAnyOrderLaunch = 0x01
+hipStreamWaitValueGte = 0x0
+hipStreamWaitValueEq = 0x1
+hipStreamWaitValueAnd = 0x2
+hipStreamWaitValueNor = 0x3
+hipExternalMemoryDedicated = 0x1
+hipStreamAttrID = hipLaunchAttributeID
+hipStreamAttributeAccessPolicyWindow = hipLaunchAttributeAccessPolicyWindow
+hipStreamAttributeSynchronizationPolicy = hipLaunchAttributeSynchronizationPolicy
+hipStreamAttributeMemSyncDomainMap = hipLaunchAttributeMemSyncDomainMap
+hipStreamAttributeMemSyncDomain = hipLaunchAttributeMemSyncDomain
+hipStreamAttributePriority = hipLaunchAttributePriority
+hipStreamAttrValue = hipLaunchAttributeValue
+hipKernelNodeAttrID = hipLaunchAttributeID
+hipKernelNodeAttributeAccessPolicyWindow = hipLaunchAttributeAccessPolicyWindow
+hipKernelNodeAttributeCooperative = hipLaunchAttributeCooperative
+hipKernelNodeAttributePriority = hipLaunchAttributePriority
+hipKernelNodeAttrValue = hipLaunchAttributeValue
+hipDrvLaunchAttributeCooperative = hipLaunchAttributeCooperative
+hipDrvLaunchAttributeID = hipLaunchAttributeID
+hipDrvLaunchAttributeValue = hipLaunchAttributeValue
+hipDrvLaunchAttribute = hipLaunchAttribute
+hipGraphKernelNodePortDefault = 0
+hipGraphKernelNodePortLaunchCompletion = 2
+hipGraphKernelNodePortProgrammatic = 1
+HIP_TRSA_OVERRIDE_FORMAT = 0x01
+HIP_TRSF_READ_AS_INTEGER = 0x01
+HIP_TRSF_NORMALIZED_COORDINATES = 0x02
+HIP_TRSF_SRGB = 0x10