work

.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, fw"
+        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 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): self.instructions, self.labels, self.pos = [], {}, 0
+    def __init__(self, arch='gfx1100'): self.instructions, self.labels, self.pos, self.arch = [], {}, 0, arch
     def label(self, name): self.labels[name] = self.pos
     def emit(self, inst, target=None):
       self.instructions.append(inst)

examples/mlperf/model_train.py

@@ -1357,7 +1357,6 @@ 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
@@ -1434,7 +1433,6 @@ 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
@@ -1462,7 +1460,7 @@ def train_llama3():
       apply_grad(g, new_g.uop)
 
     loss_cpu = loss.flatten().float().to("CPU")
-    return loss_cpu.realize(*grads, *fp8_amax, *fp8_grad_amax)
+    return loss_cpu.realize(*grads, *fp8_amax)
 
   @TinyJit
   def optim_step():
@@ -1637,6 +1635,7 @@ 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,73 +20,55 @@ 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().cast(dtypes.float32)
+  new_amax = (local_abs_max(x) if isinstance(x.device, tuple) else x.abs().max()).detach()
   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,
-           grad_amax_state:Tensor|None=None) -> tuple[Tensor,...]:
+           x_fp8:Tensor|None=None, x_scale:Tensor|None=None, x_new_amax: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:
-    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 x_fp8 is None: 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, 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
+    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
 
-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)
+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)
     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)
-    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
+  else:
+    x = normed * norm
+    out, *ret = matmul(x, w, amax_x=amax_x, w_inv_scale=w_inv_scale)
+  return out, normed, rrms, ret
 
-def silu_w13_quantize_matmul(x_w13:Tensor, w2:Tensor, s_2:Tensor,
-                             amax_x2:Tensor,
-                             grad_amax_xw13:Tensor, grad_amax_xout:Tensor):
+def silu_w13_matmul(x_w13:Tensor, w2:Tensor, amax_x2, s_2):
+  FUSED_SILU_W13 = getenv("FUSED_SILU_W13", 0)
   if FUSED_SILU_W13:
     from extra.llama_kernels.cast_amax import fused_quantize_fp8_w13
-    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)
+    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)
   return out, ret
 
 class FlatTransformer:
@@ -122,33 +104,32 @@ 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, dtype=dtypes.float32).contiguous().requires_grad_(False)
+    def _amax(): return Tensor.full((), FP8_MAX).contiguous().requires_grad_(False)
     names = ["xqkv", "xo", "x13", "x2"]
     self._fp8_amax = {name: [_amax() for _ in range(n_layers)] for name in names}
-    grad_names = ["xqkv", "xo", "xw13", "xout"]
-    self._fp8_grad_amax = {name: [_amax() for _ in range(n_layers)] for name in grad_names}
+    # per-weight inv_scale: single (n_layers,) float32 tensor per weight (kernel reads float* pointers)
     w_names = ["wqkv", "wo", "w13", "w2"]
-    self._fp8_inv_scale = {wname: inv_scales.float().contiguous().requires_grad_(False)
-                           for wname, inv_scales in zip(w_names, self._init_inv_scales)}
+    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)
     del self._init_inv_scales
 
   def lin_per_layer(self, in_features:int, out_features:int, std:float=0.02):
-    if getenv("ZEROS"): w = Tensor.zeros(self.n_layers, out_features, in_features)
+    if getenv("ZEROS", 0): 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)
+    self._init_inv_scales.append((amax + 1e-8) / FP8_MAX)  # save for inv_scale init
     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:Tensor, amax_xo:Tensor, s_qkv:Tensor, s_o:Tensor,
-                grad_amax_xqkv:Tensor, grad_amax_xo:Tensor):
+                amax_xqkv=None, amax_xo=None, s_qkv=None, s_o=None):
     bsz, seqlen, _ = x.shape
     new_amaxs, saves = [], []
 
-    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])
+    xqkv, normed, rrms, ret = norm_mul_quantize_matmul(x, attention_norm, amax_xqkv, s_qkv, wqkv, self.norm_eps)
+    saves.extend([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)
@@ -167,43 +148,40 @@ class FlatTransformer:
       attn = xq.scaled_dot_product_attention(xk, xv, is_causal=True, enable_gqa=True)
     attn = attn.transpose(1, 2).reshape(bsz, seqlen, -1)
 
-    out, *ret = matmul(attn, wo, amax_x=amax_xo, w_inv_scale=s_o, grad_amax_state=grad_amax_xo)
+    out, *ret = matmul(attn, wo, amax_x=amax_xo, w_inv_scale=s_o)
     new_amaxs.extend(ret[:1])
     saves.extend(ret[1:] + [out])
     return (out, *new_amaxs, *saves)
 
-  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):
+  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):
     new_amaxs, saves = [], []
 
-    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])
+    x_w13, normed, rrms, ret = norm_mul_quantize_matmul(x, ffn_norm, amax_x13, s_13, w13, self.norm_eps)
+    saves.extend([normed, rrms])
     new_amaxs.extend(ret[:1])
     saves.extend(ret[1:] + [x_w13])
 
-    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)
+    out, ret = silu_w13_matmul(x_w13, w2, amax_x2, s_2)
     new_amaxs.extend(ret[:1])
     saves.extend(ret[1:] + [out])
-    return (out, h, *new_amaxs, *saves)
+    return (out, *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: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):
+                amax_xqkv=None, amax_xo=None,
+                amax_x13=None, amax_x2=None,
+                s_qkv=None, s_o=None, s_13=None, s_2=None):
     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,
-                                     grad_amax_xqkv=grad_amax_xqkv, grad_amax_xo=grad_amax_xo)
+                                     amax_xqkv=amax_xqkv, amax_xo=amax_xo,
+                                     s_qkv=s_qkv, s_o=s_o)
     attn_amaxs, attn_saves = attn_ret[:2], attn_ret[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)
+    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_amaxs, ffn_saves = ffn_ret[:2], ffn_ret[2:]
     h = h + ffn
     return (h, *attn_amaxs, *ffn_amaxs, *attn_saves, *ffn_saves)
@@ -224,31 +202,28 @@ 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 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_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 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], :, :, :]
-    a, ga, s = self._fp8_amax, self._fp8_grad_amax, self._fp8_inv_scale
+    amaxs, inv_scales = self._fp8_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=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])
+                               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])
       for name, new_val in zip(["xqkv", "xo", "x13", "x2"], ret[:5]):
-        a[name][i].assign(new_val)
+        amaxs[name][i].assign(new_val)
 
-    logits = matmul(self.norm(h), self.output[0], fp8=False)[0]
+    logits = matmul(self.norm(h).contiguous().contiguous_backward(), self.output[0], fp8=False)[0].contiguous_backward()
     return logits
 
 def _get_pads(uop:UOp) -> list[UOp]:
@@ -264,11 +239,6 @@ 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,10 +16,6 @@ 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,10 +16,6 @@ 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,10 +17,6 @@ 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/headers/amdgpu_irq.h

@@ -28,7 +28,15 @@
 // #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,7 +22,15 @@
 #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,7 +24,15 @@
 #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): self.instructions, self.labels, self.pos = [], {}, 0
+  def __init__(self, arch='gfx1100'): self.instructions, self.labels, self.pos, self.arch = [], {}, 0, arch
   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):
+def build_kernel(N, arch='gfx1100'):
   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()
+  k = Kernel(arch)
 
   # ===========================================================================
   # 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)
+  insts = build_kernel(N, dev.renderer.target.arch)
 
   rng = np.random.default_rng(42)
   a = Tensor(rng.random((N, N), dtype=np.float32) - 0.5)

extra/gemm/cdna_asm_gemm.py

@@ -2628,9 +2628,8 @@ 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, *extra: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.
-  # extra is unused fwd inputs (e.g. grad_amax_state) plumbed through so the bwd can read them via kernel.src.
+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
   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}"
@@ -2699,33 +2698,19 @@ 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) optionally plus grad_amax_state (6 total); plain gemm has 3
-  if len(inputs) >= 5:
-    grad_amax_state = inputs[5] if len(inputs) == 6 else None
-    out, a, b, s_x, s_w = inputs[:5]
+  # 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
     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]]
-    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))
+    g_fp8, g_scale, _ = quantize_fp8(g_t)
     # 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)
-    # 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
+    return (None, grad_a.uop, grad_b.uop, None, None)
   else:
     out, a, b = inputs
     assert all_same([gradient.device, a.device, b.device, out.device])
@@ -2740,7 +2725,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, grad_amax_state:Tensor|None=None) -> Tensor:
+def asm_gemm(a:Tensor, b:Tensor, x_scale:Tensor|None=None, w_scale: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
@@ -2777,8 +2762,7 @@ def asm_gemm(a:Tensor, b:Tensor, x_scale:Tensor|None=None, w_scale:Tensor|None=N
       _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()
-      extra = [grad_amax_state] if grad_amax_state is not None else []
-      out = Tensor.custom_kernel(out, a, b.T, xs, ws, *extra, fxn=functools.partial(custom_hk_fp8_gemm, dname=dname), grad_fxn=custom_gemm_bw)[0]
+      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]
     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,8 +4,7 @@ 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 get_runtime
-from tinygrad.codegen import to_program
+from tinygrad.engine.realize import CompiledRunner
 from tinygrad.uop.ops import Ops, UOp, KernelInfo, ProgramInfo
 from tinygrad.helpers import getenv
 np.set_printoptions(suppress=True)
@@ -93,15 +92,13 @@ 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 = 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)
+  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)
   all_bufs = [x.ensure_allocated() for x in bufs]
-  prg_bufs = [all_bufs[i] for i in info.globals]
-  gsize, lsize = info.launch_dims({})
+  prg_bufs = [all_bufs[i] for i in runner.p.globals]
   tflops = []
   for i in range(5):
-    tm = rt(*[b._buf for b in prg_bufs], global_size=gsize, local_size=lsize, vals=info.vals({}), wait=True)
+    tm = runner(prg_bufs, {}, wait=True)
     tflops.append((2*M*K*N/tm)*1e-12)
   print(f"TFLOPS: {max(tflops):.2f}")
 

extra/llama_kernels/__init__.py

@@ -29,24 +29,7 @@ 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()
-  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}"])
+  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)

extra/llama_kernels/cast_amax/__init__.py

@@ -3,34 +3,26 @@ 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, 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 = {}
+from extra.llama_kernels import FP8_MAX, NUM_WG, THREADS_PER_WG, compile_cpp, shard_shape, scalar_amax
 
 @functools.cache
-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:
+def _custom_fused_bwd_w13(grad_xw13:UOp, xw13:UOp, grad_x2:UOp, 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 + 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)))
+  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)))
   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, 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
+def _custom_fused_cast_amax_w13(fp8_out:UOp, amax_buf:UOp, xw13:UOp, 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 * 2 + n_elems + NUM_WG * 4
+  mem = n_elems * 2 * 2 + n_elems + NUM_WG * 2
   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)
@@ -38,39 +30,44 @@ 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):
-  _, _, xw13, amax_state, grad_amax_state = kernel.src[1:]
+  # NOTE: inputs are (fp8_out, amax_buf, xw13, amax_state); grad for xw13 only
+  _, _, xw13, amax_state = kernel.src[1:]
   device = xw13.device
-  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)
+  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)
 
-def fused_quantize_fp8_w13(xw13:Tensor, amax_state:Tensor, fp8_dtype, grad_amax_state:Tensor) -> tuple[Tensor, Tensor, Tensor]:
+def fused_quantize_fp8_w13(xw13:Tensor, amax_state:Tensor, fp8_dtype) -> 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
-  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)
+  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)
   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,6 +1,5 @@
 #include <hip/hip_runtime.h>
 #include <hip/hip_bf16.h>
-#include <hip/hip_fp8.h>
 
 #ifndef N_ELEMS
 #define N_ELEMS 234881024
@@ -21,32 +20,19 @@ 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
-    __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)
+    __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
 {
-  __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;
@@ -63,7 +49,6 @@ 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]);
@@ -73,26 +58,11 @@ fused_silu_mul_bwd_w13(
       const float silu = f1 * sig;
       const float silu_prime = sig + silu * (1.0f - sig);
       const float gs = fg * scale;
-      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);
+      out1[i] = static_cast<__hip_bfloat16>(gs * silu_prime * f3);
+      out3[i] = static_cast<__hip_bfloat16>(gs * silu);
     }
 
     *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
-    float*                __restrict__ amax_buf,        // fp32, NUM_WG (per-WG amaxes)
+    __hip_bfloat16*       __restrict__ amax_buf,        // bf16, NUM_WG (per-WG amaxes)
     const __hip_bfloat16* __restrict__ xw13,            // bf16, 2*N_ELEMS
-    const float*          __restrict__ amax_state)      // fp32 scalar
+    const __hip_bfloat16* __restrict__ amax_state)      // bf16 scalar
 {
   __shared__ float sdata[THREADS_PER_WG];
 
@@ -75,5 +75,5 @@ fused_silu_mul_cast_amax_w13(
     __syncthreads();
   }
 
-  if (tid == 0) amax_buf[wg] = sdata[0];
+  if (tid == 0) amax_buf[wg] = static_cast<__hip_bfloat16>(sdata[0]);
 }

extra/llama_kernels/fused_ce/__init__.py

@@ -42,7 +42,8 @@ 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 = logits_u.shape  # (rows, VOCAB) after reshape
+  rows, VOCAB = rows_vocab
   if isinstance(device, tuple):
     axis = logits_u.axis
     ndev = len(device)
@@ -53,8 +54,9 @@ 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 = Tensor(gradient, device=device).float().reshape(-1)[0:1].contiguous()
+  scale = grad_t[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

@@ -0,0 +1,54 @@
+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_mul_quantize_fp8/fused_mul_quantize_fp8.cpp

@@ -2,13 +2,12 @@
 #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
@@ -20,13 +19,15 @@ 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
-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)
+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
 {
   __shared__ float sdata[THREADS_PER_WG];
 
@@ -39,25 +40,32 @@ quantize_fp8_with_amax(
   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 v = static_cast<float>(xi[i]);
-      local_max = fmaxf(local_max, fabsf(v));
-      const float scaled = fmaxf(-FP8_MAX, fminf(FP8_MAX, v * scale));
+      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));
       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_partial[wg] = sdata[0];
+
+  if (tid == 0) amax_buf[wg] = static_cast<__hip_bfloat16>(sdata[0]);
 }

extra/llama_kernels/fused_pad_grad_accum/__init__.py

@@ -1,55 +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 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

@@ -1,63 +0,0 @@
-// 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

@@ -1,153 +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, 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

@@ -1,155 +0,0 @@
-#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

@@ -1,147 +0,0 @@
-#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

@@ -1,67 +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, 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

@@ -1,48 +0,0 @@
-#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/sqtt/roc.py

@@ -1,13 +1,10 @@
 #!/usr/bin/env python3
-import ctypes, pathlib, argparse, pickle, dataclasses, threading, itertools
-from decimal import Decimal
+import ctypes, pathlib, argparse, pickle, dataclasses, threading
 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)
@@ -129,71 +126,6 @@ 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

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, compile_linear
+from tinygrad.engine.realize import run_linear, estimate_uop
 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) -> UOp:
+def custom_handwritten(A:UOp, arch:str) -> 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()
+  k = Kernel(arch)
   # 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) -> 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) -> UOp:
+def custom_data_deps(A:UOp, arch:str) -> UOp:
   A = A.flatten()
   threads = UOp.special(A.numel(), "lidx0")
-  k = Kernel()
+  k = Kernel(arch)
   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 = compile_linear(a.schedule_linear())
+    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=custom_handwritten)[0]
+    a = Tensor.custom_kernel(a, fxn=functools.partial(custom_handwritten, arch=self.arch))[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=custom_data_deps)[0]
+    a = Tensor.custom_kernel(a, fxn=functools.partial(custom_data_deps, arch=self.arch))[0]
     a.realize()
     self.assertTrue((a.numpy() == 6.0).all())
 

test/amd/test_mockgpu_invalid.py

@@ -8,14 +8,17 @@ class TestMockGPUInvalidInstruction(unittest.TestCase):
     test_code = '''
 import struct
 from tinygrad import Device, Tensor
-from tinygrad.engine.realize import compile_linear
+from tinygrad.engine.realize import get_runner
 from tinygrad.runtime.ops_amd import AMDProgram
 
 dev = Device["AMD"]
 a = Tensor([1.0]).realize()
 b = a + 1
-linear = compile_linear(b.schedule_linear())
-lib = bytearray(linear.src[-1].src[0].src[4].arg)
+si = b.schedule_linear().src[-1]
+runner = get_runner(dev.device, si.src[0])
+
+prg = runner._prg
+lib = bytearray(prg.lib)
 
 # 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, compile_linear
+from tinygrad.engine.realize import run_linear, estimate_uop
 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 = compile_linear(tensor.schedule_linear())
+    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 = compile_linear(t.triu().schedule_linear())
+      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 = compile_linear(wq.grad.schedule_linear())
+    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, GlobalCounters
+from tinygrad import Tensor, UOp
 from tinygrad.dtype import AddrSpace, dtypes
 from tinygrad.uop.ops import KernelInfo, AxisType
 
@@ -308,22 +308,6 @@ 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)
-
 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, buffers
+from tinygrad.uop.ops import UOp, Ops, GroupOp, AxisType
 from tinygrad.device import Device, Buffer, is_dtype_supported
 from tinygrad.tensor import Tensor, _to_np_dtype
-from tinygrad.engine.realize import run_linear
+from tinygrad.engine.realize import run_linear, CompiledRunner
 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
@@ -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 and len(u.arg[1]) > 0 for u in sched[0].src[0].toposort()), "found reduce in sum collapse"
+    assert not any(u.op is Ops.REDUCE_AXIS 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,28 +424,30 @@ 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 run_prg(opts):
+  def get_prg(opts):
     ast = realized_ast if opts is None else replace_opts(realized_ast, list(opts))
-    run_linear(UOp(Ops.LINEAR, src=(ast.call(*buf_uops),)))
+    return CompiledRunner(to_program(ast, renderer=Device[Device.DEFAULT].renderer), device)
 
   def check_opt(opts):
+    prg = get_prg(opts=opts)
     reset_bufs(outbufs)
-    run_prg(opts)
+    prg.exec(real_bufs)
     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.
-  run_prg(opts=())
+  prg = get_prg(opts=())
+  prg.exec(real_bufs)
   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)
-  run_prg(opts=None)
+  prg.exec(real_bufs)
   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_profiler.py

@@ -3,8 +3,7 @@ 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_runtime
-from tinygrad.codegen import to_program
+from tinygrad.engine.realize import get_runner
 
 MOCKGPU = DEV.interface.startswith("MOCK")
 def _dev_base(d):
@@ -47,15 +46,13 @@ class TestProfiler(unittest.TestCase):
     TestProfiler.b = self.a + 1
     si = self.b.schedule_linear().src[-1]
 
-    TestProfiler.prg = to_program(si.src[0], TestProfiler.d0.renderer)
-    TestProfiler.runtime = get_runtime(TestProfiler.d0.device, TestProfiler.prg)
+    TestProfiler.runner = get_runner(TestProfiler.d0.device, si.src[0])
     TestProfiler.b.uop.buffer.allocate()
 
   def test_profile_kernel_run(self):
-    runner_name = TestProfiler.runtime.name
+    runner_name = TestProfiler.runner._prg.name
     with helper_collect_profile(TestProfiler.d0) as profile:
-      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)
+      TestProfiler.runner([TestProfiler.b.uop.buffer, TestProfiler.a.uop.buffer], var_vals={})
 
     profile, _ = helper_profile_filter_device(profile, TestProfiler.d0.device)
     kernel_runs = [x for x in profile if isinstance(x, ProfileRangeEvent)]
@@ -73,13 +70,12 @@ class TestProfiler(unittest.TestCase):
     assert len(kernel_runs) == 1, "one kernel run is expected"
 
   def test_profile_multiops(self):
-    runner_name = TestProfiler.runtime.name
+    runner_name = TestProfiler.runner._prg.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))))
-      gs, ls = TestProfiler.prg.arg.launch_dims({})
-      TestProfiler.runtime(buf1._buf, TestProfiler.a.uop.buffer._buf, global_size=gs, local_size=ls)
+      TestProfiler.runner([buf1, TestProfiler.a.uop.buffer], var_vals={})
       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_renderer_failures.py

@@ -1,8 +1,9 @@
 import unittest
 import numpy as np
-from tinygrad.device import Device, is_dtype_supported
+from dataclasses import replace
+from tinygrad.device import Buffer, Device, is_dtype_supported
 from tinygrad.dtype import dtypes, ConstType
-from tinygrad.engine.realize import run_linear
+from tinygrad.engine.realize import CompiledRunner
 from tinygrad.codegen import to_program
 from tinygrad.helpers import prod
 from tinygrad.renderer.cstyle import CStyleLanguage
@@ -12,13 +13,17 @@ 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], sink:UOp, local_size=None):
+def _test_uop_result(inputs:list[Tensor], prg:UOp, local_size=None):
   for x in inputs: x.realize()
-  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]
+  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]
 
 def _setup_and_test_alu(alu_op:Ops, input_val:ConstType, *alu_src_uops:UOp):
   dtype = alu_src_uops[0].dtype
@@ -28,7 +33,9 @@ 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)
-  return _test_uop_result([Tensor([input_val])], UOp(Ops.SINK, dtypes.void, (store,), arg=KernelInfo()))[0]
+  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]
 
 class TestRendererFailures(unittest.TestCase):
   @unittest.skipIf(not isinstance(Device[Device.DEFAULT].renderer, (PTXRenderer, PythonRenderer)), "test is for ptx or python renderer")
@@ -37,7 +44,8 @@ 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())
-    ret = _test_uop_result([], sink, local_size=[4, 1, 1])[0]
+    prg = to_program(sink, Device[Device.DEFAULT].renderer)
+    ret = _test_uop_result([], prg, 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")
@@ -47,7 +55,8 @@ 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())
-    ret = _test_uop_result([], sink, local_size=[4, 2, 1])[0]
+    prg = to_program(sink, Device[Device.DEFAULT].renderer)
+    ret = _test_uop_result([], prg, 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")
@@ -93,7 +102,8 @@ 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())
-    ret = _test_uop_result([], sink, local_size=[4, 1, 1])[0]
+    prg = to_program(sink, Device[Device.DEFAULT].renderer)
+    ret = _test_uop_result([], prg, 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,19 +5,18 @@ 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, buffers
+from tinygrad.uop.ops import Ops, UOp, KernelInfo, AxisType
 from tinygrad.renderer.cstyle import CStyleLanguage
-from tinygrad.engine.realize import run_linear
+from tinygrad.engine.realize import CompiledRunner, 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 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 _uops_to_prg(uops_list):
+  prg = to_program(UOp.sink(*uops_list, arg=KernelInfo()), Device[Device.DEFAULT].renderer)
+  return CompiledRunner(prg, Device.DEFAULT)
 
 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))
@@ -34,7 +33,8 @@ 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)]
-  run_uops([out], [buf]+buf2)
+  prg = _uops_to_prg([out])
+  prg.exec([buf]+buf2)
   ret = np.empty(1, _to_np_dtype(output_dtype))
   buf.copyout(ret.data)
   return ret[0]
@@ -47,7 +47,8 @@ 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()
-  run_uops([out], [buf])
+  prg = _uops_to_prg([out])
+  prg.exec([buf])
   ret = np.empty(1, _to_np_dtype(output_dtype))
   buf.copyout(ret.data)
   return ret[0]
@@ -58,7 +59,8 @@ 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()
-  run_uops([out], [buf])
+  prg = _uops_to_prg([out])
+  prg.exec([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_runtime
+from tinygrad.engine.realize import get_runner, CompiledRunner
 from tinygrad.codegen import to_program
 from tinygrad.codegen.opt import Opt, OptOps
 from tinygrad import Variable
@@ -22,12 +22,11 @@ class TestHCQ(unittest.TestCase):
     TestHCQ.b = self.a + 1
     si = self.b.schedule_linear().src[-1]
 
-    TestHCQ.prg = to_program(si.src[0], TestHCQ.d0.renderer)
-    TestHCQ.runtime = get_runtime(TestHCQ.d0.device, TestHCQ.prg)
+    TestHCQ.runner = get_runner(TestHCQ.d0.device, si.src[0])
     TestHCQ.b.uop.buffer.allocate()
 
-    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])
+    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])
 
   def setUp(self):
     TestHCQ.d0.synchronize()
@@ -115,7 +114,7 @@ class TestHCQ(unittest.TestCase):
 
   # Test exec
   def test_exec_one_kernel(self):
-    TestHCQ.d0.hw_compute_queue_t().exec(TestHCQ.runtime, TestHCQ.kernargs_ba_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size) \
+    TestHCQ.d0.hw_compute_queue_t().exec(TestHCQ.runner._prg, TestHCQ.kernargs_ba_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size) \
                                    .signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value).submit(TestHCQ.d0)
 
     TestHCQ.d0.timeline_signal.wait(TestHCQ.d0.timeline_value)
@@ -129,8 +128,8 @@ class TestHCQ(unittest.TestCase):
 
     q = TestHCQ.d0.hw_compute_queue_t()
     q.wait(TestHCQ.d0.timeline_signal, virt_val - 1) \
-     .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) \
+     .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) \
      .signal(TestHCQ.d0.timeline_signal, virt_val)
 
     for _ in range(100):
@@ -142,11 +141,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.prg.arg.global_size[1:])
-    sint_local = (Variable("sint_local", 0, 0xffffffff, dtypes.uint32),) + tuple(TestHCQ.prg.arg.local_size[1:])
+    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:])
 
     q = TestHCQ.d0.hw_compute_queue_t()
-    q.exec(TestHCQ.runtime, TestHCQ.kernargs_ba_ptr, sint_global, sint_local) \
+    q.exec(TestHCQ.runner._prg, 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})
@@ -167,17 +166,17 @@ class TestHCQ(unittest.TestCase):
     b = a + 1
     si = b.schedule_linear().src[-1]
 
-    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)
+    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)
 
     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 = runtime.fill_kernargs([zt._buf, zb._buf])
+    kernargs = runner._prg.fill_kernargs([zt._buf, zb._buf])
 
     q = TestHCQ.d0.hw_compute_queue_t()
     q.memory_barrier() \
-     .exec(runtime, kernargs, (1,1,1), virt_local) \
+     .exec(runner._prg, kernargs, (1,1,1), virt_local) \
      .signal(TestHCQ.d0.timeline_signal, virt_val)
 
     for x in range(1, 4):
@@ -331,7 +330,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.runtime, TestHCQ.kernargs_ba_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size) \
+                                   .exec(TestHCQ.runner._prg, TestHCQ.kernargs_ba_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size) \
                                    .timestamp(sig_en) \
                                    .signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value).submit(TestHCQ.d0)
 
@@ -471,13 +470,12 @@ class TestHCQ(unittest.TestCase):
   def test_memory_barrier(self):
     a = Tensor([0, 1], device=Device.DEFAULT, dtype=dtypes.int8).realize()
     b = a + 1
-    prg = to_program(b.schedule_linear().src[-1].src[0], TestHCQ.d0.renderer)
-    runtime = get_runtime(TestHCQ.d0.device, prg)
+    runner = get_runner(TestHCQ.d0.device, b.schedule_linear().src[-1].src[0])
 
     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 = runtime.fill_kernargs([buf1._buf, buf2._buf])
+    kernargs_ptr = runner._prg.fill_kernargs([buf1._buf, buf2._buf])
 
     for i in range(255):
       ctypes.memset(buf2._buf.va_addr, i, 2)
@@ -485,7 +483,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(runtime, kernargs_ptr, prg.arg.global_size, prg.arg.local_size) \
+                                     .exec(runner._prg, kernargs_ptr, runner.p.global_size, runner.p.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/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 get_runtime
+from tinygrad.engine.realize import CompiledRunner
 from tinygrad.codegen import to_program
 from tinygrad.helpers import dedup, getenv
 from tinygrad.device import Buffer
@@ -90,13 +90,12 @@ renderer = Device.default.renderer
 allocator = Device.default.allocator
 
 ps = to_program(ast, renderer)
-rt = get_runtime(Device.DEFAULT, ps)
+cr = CompiledRunner(ps, Device.DEFAULT)
 
 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]
 
-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)
+t = cr(bufs, wait=True)
 print(f"{t*1e6:.2f} us")

test/external/external_test_hcq.py

@@ -2,8 +2,7 @@ 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_runtime
-from tinygrad.codegen import to_program
+from tinygrad.engine.realize import get_runner
 
 def _time_queue(q, d):
   st = time.perf_counter()
@@ -22,14 +21,13 @@ class TestHCQ(unittest.TestCase):
     TestHCQ.a = Tensor([0.,1.], device=Device.DEFAULT).realize()
     TestHCQ.b = self.a + 1
     linear = self.b.schedule_linear()
-    TestHCQ.prg = to_program(linear.src[-1].src[0], TestHCQ.d0.renderer)
-    TestHCQ.runtime = get_runtime(TestHCQ.d0.device, TestHCQ.prg)
+    TestHCQ.runner = get_runner(TestHCQ.d0.device, linear.src[-1].src[0])
     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.runtime.kernargs_offset
-    TestHCQ.kernargs_size = TestHCQ.runtime.kernargs_alloc_size
+    TestHCQ.kernargs_off = TestHCQ.runner._prg.kernargs_offset
+    TestHCQ.kernargs_size = TestHCQ.runner._prg.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))
 
@@ -40,8 +38,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.runtime.constbuffer_0)
-      to_mv(TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, 0x160).cast('I')[:] = array.array('I', TestHCQ.runtime.constbuffer_0)
+      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)
       TestHCQ.compute_queue = HWQueue
       TestHCQ.copy_queue = HWQueue
 
@@ -55,11 +53,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.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
+      q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
       q.signal(temp_signal, temp_value + 1).wait(temp_signal, temp_value + 1)
       temp_value += 1
 
-      q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
+      q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
       q.signal(temp_signal, temp_value + 1).wait(temp_signal, temp_value + 1)
       temp_value += 1
 
@@ -73,10 +71,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.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
+    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
     q.signal(temp_signal, 2).wait(temp_signal, 2)
-    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, TestHCQ.prg.arg.global_size,
-           TestHCQ.prg.arg.local_size)
+    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, TestHCQ.runner.p.global_size,
+           TestHCQ.runner.p.local_size)
     for _ in range(1000):
       TestHCQ.d0._set_signal(temp_signal, 1)
       q.submit(TestHCQ.d0)
@@ -89,11 +87,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.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
+    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
     q.signal(temp_signal, 1).wait(temp_signal, 1)
-    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
+    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
     q.signal(temp_signal, 2).wait(temp_signal, 2)
-    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
+    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.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
@@ -103,7 +101,7 @@ class TestHCQ(unittest.TestCase):
   def test_update_exec(self):
     q = TestHCQ.compute_queue()
     exec_cmd_idx = len(q)
-    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
+    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.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)
@@ -117,10 +115,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.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
+    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
     q.signal(temp_signal, 2).wait(temp_signal, 2)
-    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, TestHCQ.prg.arg.global_size,
-           TestHCQ.prg.arg.local_size)
+    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr+TestHCQ.kernargs_size, TestHCQ.runner.p.global_size,
+           TestHCQ.runner.p.local_size)
     q.bind(TestHCQ.d0)
     for _ in range(1000):
       TestHCQ.d0._set_signal(temp_signal, 1)
@@ -135,7 +133,7 @@ class TestHCQ(unittest.TestCase):
   def test_update_exec_binded(self):
     q = TestHCQ.compute_queue()
     exec_ptr = q.ptr()
-    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
+    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.local_size)
     q.signal(TestHCQ.d0.timeline_signal, TestHCQ.d0.timeline_value)
     q.bind(TestHCQ.d0)
 
@@ -172,7 +170,7 @@ class TestHCQ(unittest.TestCase):
 
   def test_run_normal(self):
     q = TestHCQ.compute_queue()
-    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
+    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.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
@@ -203,7 +201,7 @@ class TestHCQ(unittest.TestCase):
 
   def test_run_signal(self):
     q = TestHCQ.compute_queue()
-    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
+    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.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)
@@ -280,7 +278,7 @@ class TestHCQ(unittest.TestCase):
   def test_interleave_compute_and_copy(self):
     q = TestHCQ.compute_queue()
     qc = TestHCQ.copy_queue()
-    q.exec(TestHCQ.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)  # b = [1, 2]
+    q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.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)
@@ -317,7 +315,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.runtime, TestHCQ.d0.kernargs_ptr, TestHCQ.prg.arg.global_size, TestHCQ.prg.arg.local_size)
+      q.exec(TestHCQ.runner._prg, TestHCQ.d0.kernargs_ptr, TestHCQ.runner.p.global_size, TestHCQ.runner.p.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.codegen import to_program_cache
+from tinygrad.engine.realize import method_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_cache=True):
-      if empty_cache: to_program_cache.clear()
+    def run_llama(st, empty_method_cache=True):
+      if empty_method_cache: method_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,6 +1,7 @@
 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
@@ -70,6 +71,7 @@ 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,11 +53,10 @@ class _MXCSRContext:
 
 from tinygrad.uop.ops import UOp, Ops, KernelInfo, AxisType
 from tinygrad.dtype import dtypes, AddrSpace
-from tinygrad.device import Buffer, BufferSpec, Device
+from tinygrad.device import Buffer, BufferSpec
 from tinygrad.runtime.autogen import hsa
 from tinygrad.helpers import Context, DEBUG, PROFILE, colored
-from tinygrad.engine.realize import get_runtime
-from tinygrad.codegen import to_program
+from tinygrad.engine.realize import get_runner
 
 from tinygrad.renderer.amd import decode_inst
 from tinygrad.runtime.autogen.amd.rdna3.str_pcode import PCODE as PCODE_RDNA3
@@ -2046,18 +2045,18 @@ _INST_HANDLERS: dict[type, Callable[..., UOp]] = {
 # PROGRAM DECODE AND COMPILATION
 # ═══════════════════════════════════════════════════════════════════════════════
 
-_canonical_runner_cache: list[tuple[type, int, int, int, tuple[UOp, object]]] = []  # [(inst_type, base, mask, size, (prg, runtime)), ...]
+_canonical_runner_cache: list[tuple[type, int, int, int, object]] = []  # [(inst_type, base, mask, size, runner), ...]
 
 @functools.cache
 def _get_runner(inst_bytes: bytes, arch: str = "rdna3"):
-  """Build and compile instruction to (prg, runtime). Cached by instruction bytes, with canonical dedup."""
+  """Build and compile instruction to CompiledRunner. 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, entry in _canonical_runner_cache:
-    if type(inst) is inst_type and inst_size == size and (inst_int & mask) == base: return entry
+  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
 
   # Look up handler by type, falling back to base classes for _LIT variants
   handler = _INST_HANDLERS.get(type(inst))
@@ -2076,10 +2075,9 @@ 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):
-    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
+    runner = get_runner('CPU', sink)
+  _canonical_runner_cache.append((type(inst), base, mask, size, runner))
+  return runner
 
 _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)
@@ -2210,10 +2208,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)
-      (prg, runtime), inst = _decode_at(pc, arch)
+      runner, 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] = (runtime.fxn, prg.arg.globals, is_barrier, inst)
+      program[pc] = (runner._prg.fxn, runner.p.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/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 is Ops.REDUCE]
+    reduceops = [x for si in linear.src for x in si.src[0].toposort() if x.op in {Ops.REDUCE_AXIS, 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 is Ops.REDUCE]
+    reduceops = [x for si in linear.src for x in si.src[0].toposort() if x.op in {Ops.REDUCE_AXIS, Ops.REDUCE}]
     self.assertEqual(len(reduceops), 2) # why is RANGEIFY different?
 
   def test_dedup_assign(self):

test/null/test_simplify_valid_idx.py

@@ -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 = load.reduce(ridx1, ridx2, arg=Ops.ADD)
+    red = UOp(Ops.REDUCE, dtypes.float, (load, ridx1, ridx2), Ops.ADD)
     self.check(load,
       "(r0*1568)",
       "((r2<1)&(r1<6))",
@@ -569,7 +569,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 = (r.cast(dtypes.float) + gated_load).reduce(r, arg=Ops.ADD)
+    red = UOp(Ops.REDUCE, dtypes.float, (r.cast(dtypes.float) + gated_load, r), 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,22 +137,6 @@ 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,7 +424,7 @@ class TestUOpGraph(unittest.TestCase):
     ld = d0.index(ridx0.valid(ridx0<50))
     w = (ridx0<50).where(ld, 5)
     # prevent ridx0 from being shrunk
-    red = ridx0.cast(dtypes.long).reduce(ridx0, arg=Ops.ADD)
+    red = UOp(Ops.REDUCE, dtypes.long, (ridx0.cast(dtypes.long), ridx0), Ops.ADD)
     uops = to_uops_list([w, red])
     for u in uops:
       assert u.op is not Ops.WHERE
@@ -447,7 +447,7 @@ 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 = ridx0.cast(dtypes.long).reduce(ridx0, arg=Ops.ADD)
+    red = UOp(Ops.REDUCE, dtypes.long, (ridx0.cast(dtypes.long), ridx0), Ops.ADD)
     uops = to_uops_list([w, red])
     for u in uops:
       assert u.op is not Ops.WHERE
@@ -459,7 +459,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 = ridx0.cast(dtypes.long).reduce(ridx0, arg=Ops.ADD)
+    red = UOp(Ops.REDUCE, dtypes.long, (ridx0.cast(dtypes.long), ridx0), Ops.ADD)
     uops = to_uops_list([w, red])
     for u in uops:
       assert u.op is not Ops.WHERE
@@ -470,7 +470,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 = ridx0.cast(dtypes.long).reduce(ridx0, arg=Ops.ADD)
+    red = UOp(Ops.REDUCE, dtypes.long, (ridx0.cast(dtypes.long), ridx0), Ops.ADD)
     uops = to_uops_list([w, red])
     for u in uops:
       assert u.op is not Ops.WHERE
@@ -679,7 +679,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 = (3*e1).reduce(e1, arg=Ops.ADD)
+    sink = UOp(Ops.REDUCE, dtypes.int, (3*e1,e1), Ops.ADD)
     sink = expander_rewrite(sink)
     assert sink.op is Ops.CONST
     self.assertEqual(sink.arg, 3*(0+1+2+3))
@@ -687,7 +687,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.const(dtypes.int, 3).reduce(e1, arg=Ops.ADD)
+    sink = UOp(Ops.REDUCE, dtypes.int, (UOp.const(dtypes.int, 3), e1), Ops.ADD)
     sink = expander_rewrite(sink)
     assert sink.op is Ops.CONST
     self.assertEqual(sink.arg, 3*4)
@@ -728,7 +728,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 = e1.reduce(e2, arg=Ops.ADD)
+    sink = UOp(Ops.REDUCE, dtypes.int, (e1,e2), 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).cast(dtypes.int.vec(2).ptr()),))
+    val = UOp(Ops.LOAD, dtypes.int.vec(2), (d1.index(idx),))
     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, get_render
+from tinygrad.viz.serve import load_rewrites, get_full_rewrite, uop_to_json, VizData
 from tinygrad.codegen import to_program_cache
 from tinygrad.codegen import to_program
 
@@ -321,6 +321,7 @@ 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
@@ -706,6 +707,7 @@ 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)
@@ -721,16 +723,13 @@ 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 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"])
+    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)
 
   def test_simple(self):
-    k = Kernel()
+    k = Kernel(arch=self.arch)
     k.label("entry")
     k.emit(s_branch(), target="bb1")
     k.label("bb1")
@@ -740,7 +739,7 @@ class TestCfg(unittest.TestCase):
     self.assertEqual(len(cfg["blocks"]), 2)
 
   def test_diamond(self):
-    k = Kernel()
+    k = Kernel(arch=self.arch)
     k.label("entry")
     k.emit(s_mov_b32(s[0], 0))
     k.emit(s_mov_b32(s[1], 0))
@@ -774,7 +773,7 @@ class TestCfg(unittest.TestCase):
       assert st.startswith("s_code_end") and st.endswith("x)"), st
 
   def test_loop(self):
-    k = Kernel()
+    k = Kernel(arch=self.arch)
     k.label("entry")
     k.emit(s_mov_b32(s[1], 4))
     k.label("loop")
@@ -786,7 +785,7 @@ class TestCfg(unittest.TestCase):
     self.get_cfg("simple_loop", k)
 
   def test_loop_branch(self):
-    k = Kernel()
+    k = Kernel(arch=self.arch)
     k.label("entry")
     k.emit(s_mov_b32(s[1], 4))
     k.label("loop")
@@ -804,7 +803,7 @@ class TestCfg(unittest.TestCase):
     self.get_cfg("loop_if", k)
 
   def test_loop_break(self):
-    k = Kernel()
+    k = Kernel(arch=self.arch)
     k.label("entry")
     k.emit(s_mov_b32(s[1], 8))
     k.label("loop")
@@ -819,7 +818,7 @@ class TestCfg(unittest.TestCase):
     self.get_cfg("loop_break", k)
 
   def test_switch(self):
-    k = Kernel()
+    k = Kernel(arch=self.arch)
     k.label("entry")
     k.emit(s_cmp_eq_i32(s[0], 0))
     k.emit(s_cbranch_scc1(), target="case0")
@@ -841,7 +840,7 @@ class TestCfg(unittest.TestCase):
     self.get_cfg("switch_case", k)
 
   def test_ping_pong(self):
-    k = Kernel()
+    k = Kernel(arch=self.arch)
     k.label("entry")
     k.emit(s_cmp_eq_i32(s[0], 0))
     k.emit(s_cbranch_scc1(), target="ping")
@@ -860,7 +859,7 @@ class TestCfg(unittest.TestCase):
 
   def test_colored_blocks(self):
     N = 10
-    k = Kernel()
+    k = Kernel(arch=self.arch)
     k.label("entry")
     k.emit(s_branch(), target="init0")
     for i in range(N):
@@ -880,7 +879,7 @@ class TestCfg(unittest.TestCase):
     self.get_cfg("test_colored_blocks", k)
 
   def test_jump_back_to_end(self):
-    k = Kernel()
+    k = Kernel(arch=self.arch)
     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, UOp, buffers
+from tinygrad.uop.ops import Ops
 from tinygrad.dtype import DType
-from tinygrad.device import Buffer, is_dtype_supported
+from tinygrad.device import is_dtype_supported
 from tinygrad.helpers import DEV, Context
 from test.helpers import slow, replace_opts
-from tinygrad.engine.realize import run_linear
+from tinygrad.engine.realize import CompiledRunner
 from tinygrad.codegen import to_program
 from tinygrad.codegen.opt import Opt, OptOps, KernelOptError
 from tinygrad.codegen.opt.tc import amd_cdna_1616128
@@ -20,11 +20,6 @@ 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)
@@ -52,11 +47,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))]
-  ast = replace_opts(realized_ast, opts)
-  pu = to_program(ast, Device[Device.DEFAULT].renderer)
+  pu = to_program(replace_opts(realized_ast, opts), 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"
-  run_program(ast, bufs)
+  prg = CompiledRunner(pu, Device.DEFAULT)
+  prg.exec(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
@@ -150,15 +145,15 @@ class TestTensorCores(unittest.TestCase):
         c = a.conv2d(b, padding=1, dtype=tc.dtype_out)
         realized_ast, real_bufs = helper_realized_ast(c)
 
-        ast = replace_opts(realized_ast, [Opt(OptOps.TC, axis, (-1, 2, 1))])
-        program = to_program(ast, Device[Device.DEFAULT].renderer)
+        program = to_program(replace_opts(realized_ast, [Opt(OptOps.TC, axis, (-1, 2, 1))]), 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
-        run_program(ast, real_bufs)
+        prg.exec(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
-    self.assertIs(r.uop.base.op, Ops.REDUCE)
+    r = a.sum(axis=0)  # unrealized REDUCE_AXIS
+    self.assertIs(r.uop.base.op, Ops.REDUCE_AXIS)
     r[:1].assign(Tensor([99]).realize())
     try:
       self.assertEqual(r.tolist(), [99,6])

tinygrad/codegen/__init__.py

@@ -180,6 +180,8 @@ 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

@@ -76,7 +76,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.dtype.count)))
+  return UOp(Ops.STACK, buf.dtype, tuple(buf.index(vec.gep(i), ptr=True) for i in range(vec.shape[0])))
 
 def fold_expanded_index(midx:UOp):
   buf = midx.src[0].src[0]
@@ -104,7 +104,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
@@ -143,7 +143,7 @@ load_store_folding = PatternMatcher([
   (UPat(Ops.STORE, src=(UPat(Ops.GEP, name="gep"), UPat.var("st")), name="sto"), 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.vec(cat.dtype.vcount), 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, 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),
 ])
@@ -317,12 +317,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[0], red.dtype.scalar()))
+    identity = red.const(red.dtype, identity_element(red.arg, 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[0], y), lst)
+  ret = functools.reduce(lambda x,y: x.alu(red.arg, 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,14 +358,10 @@ 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"), add_load),
+  (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)),
   # 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

@@ -71,7 +71,8 @@ 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.scalar().vec(root.dtype.count*expand_sz), tuple(new_srcs), new_arg)
+  nsrc = UOp(root.op, root.dtype, tuple(new_srcs), new_arg)
   return UOp(Ops.UNROLL, root.dtype, (nsrc,), expand_args)
 
 def do_contract(con:UOp):
@@ -147,7 +148,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.vec(s:=r.vmax+1), tuple(range(s))),), ((r.arg[0],s),)) \
+   lambda r: UOp(Ops.UNROLL, r.dtype, (UOp.const(r.dtype, tuple(range(s:=r.vmax+1))),), ((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[0] is Ops.ADD:
+    if use_tensor_cores and reduceop.arg 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, red[0].dtype, red[0].src, red[0].arg)
+    return UOp(Ops.REDUCE_AXIS, 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,12 @@
-import math, time, multiprocessing, traceback, signal, atexit
+import functools, math, time, multiprocessing, traceback, signal, atexit
 from dataclasses import replace
-from tinygrad.uop.ops import sym_infer, AxisType, pyrender, UOp
-from tinygrad.device import Device, Buffer
+from tinygrad.uop.ops import sym_infer, AxisType, pyrender, UOp, Ops
+from tinygrad.device import Device, Buffer, Compiler
 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 get_runtime
+from tinygrad.engine.realize import CompiledRunner
 from tinygrad.codegen import to_program
 from tinygrad.codegen.opt.postrange import Scheduler
 
@@ -34,24 +34,25 @@ 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, var_vals:dict[str, int], rawbufs:list[Buffer], early_stop:float|None=None,
+def _time_program(prg:UOp, lib:bytes, 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(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)
+    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)
   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(rt(*bufs, global_size=global_size, local_size=local_size, vals=prg.arg.vals(var_vals), wait=True, timeout=timeout))*factor)
+    tms.append(unwrap(car(input_bufs, var_vals, wait=True, timeout=timeout))*factor)
     if early_stop is not None and early_stop < min(tms): break
   return tms
 
@@ -60,21 +61,22 @@ def timeout_handler(signum, frame):
   if DEBUG >= 2: print("*** BEAM COMPILE TIMEOUT")
   raise TimeoutException()
 
-def _try_compile(x:tuple[int,Scheduler]) -> tuple[int, tuple[UOp, float]|None]:
+def _try_compile(x:tuple[int,Scheduler], compiler:Compiler) -> tuple[int, tuple[UOp, bytes, 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")
-    ret = (prg, et)
+    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)
   except RuntimeError:
     if DEBUG >= 4: traceback.print_exc()
   except Exception as e:
@@ -148,11 +150,12 @@ 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 if beam_pool is None else beam_pool.imap_unordered)(_try_compile, enumerate(candidates))):
+      for i,proc in (map(_compile_fn, enumerate(candidates)) if beam_pool is None else beam_pool.imap_unordered(_compile_fn, enumerate(candidates))):
         if proc is None: continue
-        prg, compile_et = proc
-        if (lib:=prg.src[4].arg) in seen_libs: continue
+        prg, lib, compile_et = proc
+        if lib 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)
@@ -160,7 +163,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, var_vals, rawbufs, early_stop=beam[0][1]*3 if len(beam) else 1.0,
+        try: tms = _time_program(prg, lib, 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

@@ -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[0] not in {Ops.ADD, Ops.MAX, Ops.MUL}: return None
+  if red.arg 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[0] is Ops.ADD:
+  if red.arg is Ops.ADD:
     for r in reduce_unparented: ret = ret * r.src[0].cast(ret.dtype.scalar()).broadcast(ret.dtype.count)
-  if red.arg[0] is Ops.MUL:
+  if red.arg 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, JIT, JIT_BATCH_SIZE, dedup, pluralize, VIZ
+from tinygrad.helpers import flatten, merge_dicts, DEBUG, Context, BEAM, getenv, colored, 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, 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.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.schedule.memory import memory_plan_rewrite, _collect_bufs
 from tinygrad.nn.state import get_parameters
 from tinygrad.schedule.rangeify import mop_cleanup
@@ -59,6 +59,14 @@ 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
@@ -87,17 +95,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:
+class GraphRunner(Runner):
   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.runtimes: list[Any|None] = []
+    self.progs: list[CompiledRunner|None] = []
     self.uop_replace: list[list[tuple[int, int]]] = []
     for call in self.linear.src:
-      replace = [(p, b.arg) for p, b in enumerate(get_call_arg_uops(call)) if b.op is Ops.PARAM]
+      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]
       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.runtimes.append(get_runtime(bufs[0].device, call.src[0]) if call.src[0].op is Ops.PROGRAM else None)
+        self.progs.append(get_runner(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]]] = {}
@@ -106,20 +114,20 @@ class GraphRunner:
 
     def is_sym_dim(dim) -> bool: return not all(isinstance(d, (int, float)) for d in dim)
 
-    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))
+    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))
 
     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.vars) if v.expr not in dv | p.runtimevars]):
+      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]):
         self.var_vals_replace[j] = replace
-      global_dim_idx, local_dim_idx = find_symbolic_dim(p.global_size), find_symbolic_dim(p.local_size)
+      global_dim_idx, local_dim_idx = find_symbolic_dim(p.p.global_size), find_symbolic_dim(p.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.local_size is not None
-        self.launch_dims_base[j] = (tuple(p.global_size), tuple(p.local_size))
+        assert p.p.local_size is not None
+        self.launch_dims_base[j] = (tuple(p.p.global_size), tuple(p.p.local_size))
 
     estimates = sum((estimate_uop(call) for call in self.linear.src), Estimates())
 
@@ -127,9 +135,7 @@ class GraphRunner:
     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)
 
-    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")
+    super().__init__(colored(f"<batched {len(self.calls)}>", "cyan"), self.calls[0][2][0].device.split(":")[0], estimates.simplify())
 
   def updated_vars(self, var_vals: dict[str, int]):
     vals = [var_vals[v] for v in self.vars]
@@ -162,7 +168,7 @@ class GraphRunner:
 
   @staticmethod
   def _all_devs(batch_devs:list[Compiled], new_call:UOp) -> list[Compiled]:
-    return dedup(batch_devs + [Device[x] for b in get_call_arg_uops(new_call)
+    return dedup(batch_devs + [Device[x] for b in new_call.src[1:] if b.op is not Ops.BIND
                  for x in (b.device if isinstance(b.device, tuple) else (b.device,))])
 
   @staticmethod
@@ -191,9 +197,9 @@ class CapturedJit(Generic[ReturnType]):
     out: set[UOp] = set()
     for call in self.linear.toposort():
       if call.op is not Ops.CALL: continue
-      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)}
+      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)}
     return out
 
   def __call__(self, input_uops:list[UOp], var_vals:dict[str, int]) -> ReturnType:

tinygrad/engine/realize.py

@@ -1,9 +1,9 @@
-from __future__ import annotations
-from typing import cast, Iterator, Any
+from typing import cast, Iterator
 import time, random, itertools, math, contextlib, weakref
 from dataclasses import dataclass, replace, field
-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.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.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,74 +11,50 @@ 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 get_graph_runtime(ast).estimates
+  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()
   return Estimates()
 
-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)
+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):
   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[0] is not None: GlobalCounters.time_sum_s += et[0]
+  if et is not None: GlobalCounters.time_sum_s += et
   if DEBUG >= 2:
-    display_name = get_call_name(call, bufs, var_vals)
     lds_est = sym_infer(estimates.lds, var_vals)
-    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)
+    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)
     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 {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)
+      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")
 
 local_size_cache: dict[bytes, tuple[int, ...]] = {}
 def optimize_local_size(call:UOp, prg:UOp) -> UOp|None:
@@ -102,24 +78,43 @@ 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:]))
 
-# **************** runtime cache ****************
+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: 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
+  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)
 
-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
+# **************** 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
 
 # **************** run linear ****************
 
@@ -135,7 +130,21 @@ 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 get_call_arg_uops(call)]
+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 unwrap_multi(call:UOp, resolved:list[UOp]) -> Iterator[tuple[list[Buffer], dict[str, int]]]:
   bufs = [b.buffer for b in resolved]
@@ -148,13 +157,16 @@ 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, [bv, bufs[1]], ctx.var_vals): buffers[resolved[0]] = bv
+  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
 
 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()
-    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]:
+    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:
         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:
@@ -166,31 +178,35 @@ 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_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)
+    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
 
 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)):
-    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)
+    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)
 
 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, bufs, ctx.var_vals):
+  with track_stats(ctx, call, bufs[0].device, colored(f"enc/dec {size_to_str(bufs[0].nbytes)}", "yellow"), 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])
 
-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]
+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]
 
 # flatten LINEAR-in-LINEAR: any nested LINEAR child gets inlined into its parent's src
 pm_flatten_linear = PatternMatcher([
@@ -199,7 +215,7 @@ pm_flatten_linear = PatternMatcher([
 ])
 
 def _validate(call:UOp, sink:UOp) -> UOp:
-  params = get_call_arg_uops(call)
+  params = tuple(p for p in call.src[1:] if p.op is not Ops.BIND)
   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)))
@@ -225,7 +241,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="ast"),), name="call", allow_any_len=True), exec_graph),
+  (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="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, "only works on REDUCE"
+    assert ret.op is Ops.REDUCE_AXIS, "only works on REDUCE_AXIS"
     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, name="ret"), lambda ctx, ret: reduce_gradient(ctx, ret, ret.arg[0])),
+  (UPat(Ops.REDUCE_AXIS, 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, DType, DTypeLike, Invalid, InvalidType, PtrDType, PyConst, dtypes, least_upper_dtype, sum_acc_dtype, to_dtype
+from tinygrad.dtype import ConstType, 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,36 +191,6 @@ 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,8 +7,6 @@ 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/renderer/cstyle.py

@@ -536,7 +536,6 @@ 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/runtime/autogen/__init__.py

@@ -12,17 +12,15 @@ 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 = (
-  (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")
+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']"
 
 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, files, **kwargs):
+def load(name, dll, 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)):
@@ -39,24 +37,23 @@ def load(name, 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
-    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
+    f.write_text(importlib.import_module("tinygrad.runtime.support.autogen").gen(name, dll, files, **kwargs))
     if srcs: td.cleanup()
   return importlib.import_module(f"{path}.{name.replace('/', '.')}")
 
 def __getattr__(nm):
   match nm:
-    case "libc": return load("libc", lambda: (
+    case "libc": return load("libc", "'c'", 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"]), 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"])
+      ["/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"])
     case "nv_570" | "nv_580":
-      return load(nm, [
+      return load(nm, None, [
         *[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"]],
@@ -71,7 +68,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", [
+    case "nv": return load("nv", None, [
       *[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"]],
@@ -90,50 +87,49 @@ def __getattr__(nm):
     })
     # this defines all syscall numbers. should probably unify linux autogen?
     case "io_uring":
-      return load("io_uring", ["{}/liburing.h", "{}/usr/include/linux/io_uring.h", "{}/usr/include/asm-generic/unistd.h"],
+      return load("io_uring", None, ["{}/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", ["/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,
+    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"],
                              args=lambda: system("llvm-config-20 --cflags").split(), recsym=True, prolog=["from tinygrad.helpers import WIN, OSX"])
-    case "pci": return load("pci", ["{}/usr/include/linux/pci_regs.h"], srcs=linux_headers_deb,
+    case "pci": return load("pci", None, ["{}/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", ["{}/usr/include/linux/vfio.h"], args=["-I{}/usr/include"], srcs=linux_headers_deb,
+    case "vfio": return load("vfio", None, ["{}/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", [root/"extra/webgpu/webgpu.h"], dll=webgpu_lib,
+    case "webgpu": return load("webgpu", webgpu_lib, [root/"extra/webgpu/webgpu.h"],
                                prolog=["from tinygrad.helpers import WIN, OSX", "import sysconfig, os"])
-    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'",
+    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"],
                             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", ["/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']",
+      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"]]],
       srcs=rocr_src, args=["-DLITTLEENDIAN_CPU"], prolog=["import os"])
-    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",
+    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",
                                                                                                "sienna_cichlid_ip_offset"]],
                                 args=["-I/opt/rocm/include", "-x", "c++"])
-    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 "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 "qcom_dsp":
-      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"])
+      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"])
     case "rocprof":
-      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')]",
+      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"]],
                   srcs="https://github.com/ROCm/rocprof-trace-decoder/archive/dd0485100971522cc4cd8ae136bdda431061a04d.tar.gz")
-    case "mesa": return load("mesa", [
+    case "mesa": return load("mesa", "([] if DEV.renderer == 'LVP' else ['tinymesa']) + ['tinymesa_cpu']", [
         *[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",
@@ -152,28 +148,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",
+      return load("libclang", clang_lib,
                   lambda: [f"{system('llvm-config-20 --includedir')}/clang-c/{s}.h" for s in ["Index", "CXString", "CXSourceLocation", "CXFile"]],
-                  dll=clang_lib, prolog=["from tinygrad.helpers import WIN, OSX"], args=lambda: system("llvm-config-20 --cflags").split())
+                  prolog=["from tinygrad.helpers import OSX"], args=lambda: system("llvm-config-20 --cflags").split())
     case "metal":
-      return load("metal", [f"{macossdk}/System/Library/Frameworks/Metal.framework/Headers/MTL{s}.h" for s in
+      return load("metal", "'Metal'", [f"{macossdk}/System/Library/Frameworks/Metal.framework/Headers/MTL{s}.h" for s in
                   ["ComputeCommandEncoder", "ComputePipeline", "CommandQueue", "Device", "IndirectCommandBuffer", "Resource", "CommandEncoder"]],
-                  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=["-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"],
                               args=["-isysroot", macossdk])
-    case "corefoundation": return load("corefoundation",
+    case "corefoundation": return load("corefoundation", "'CoreFoundation'",
                                        [f"{macossdk}/System/Library/Frameworks/CoreFoundation.framework/Headers/CF{s}.h" for s in ["String", "Data"]],
-                                       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)
+                                       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)
     case "mlx5":
       kh = "{}/usr/src/linux-headers-6.18.9+deb14-common/include/linux/mlx5"
-      return load("mlx5", [root/"extra/mlx_driver/mlx5.h", f"{kh}/mlx5_ifc.h"], srcs=linux_headers_kern_deb,
+      return load("mlx5", None, [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,37 +1,29 @@
-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 _: 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', 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)])
+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)])
 @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', uint32_t, 32), ('io_debug_array_offset_bytes', uint32_t, 36)])
+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)])
 @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', uint32_t, 32)])
+struct_smc_firmware_header_v1_0.register_fields([('header', struct_common_firmware_header, 0), ('ucode_start_addr', ctypes.c_uint32, 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', uint32_t, 36), ('ppt_size_bytes', uint32_t, 40)])
+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)])
 @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', uint32_t, 0), ('ppt_offset_bytes', uint32_t, 4), ('ppt_size_bytes', uint32_t, 8)])
+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)])
 @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', uint32_t, 36), ('pptable_entry_offset', uint32_t, 40)])
+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)])
 @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', uint32_t, 0), ('offset_bytes', uint32_t, 4), ('size_bytes', uint32_t, 8)])
+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)])
 @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', uint32_t, 0), ('fw_version', uint32_t, 4), ('offset_bytes', uint32_t, 8), ('size_bytes', uint32_t, 12)])
+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)])
 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', uint32_t, 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', ctypes.c_uint32, 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', uint32_t, 32), ('psp_aux_fw_bin_index', uint32_t, 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', 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)])
 @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', uint32_t, 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', ctypes.c_uint32, 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', uint32_t, 32), ('jt_offset', uint32_t, 36), ('jt_size', uint32_t, 40)])
+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)])
 @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', 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)])
+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)])
 @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', 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)])
+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)])
 @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', 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)])
+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)])
 @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', 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)])
+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)])
 @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', 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)])
+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)])
 @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', 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)])
+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)])
 @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', 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)])
+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)])
 @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', 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)])
+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)])
 @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', uint32_t, 32), ('ucode_change_version', uint32_t, 36), ('jt_offset', uint32_t, 40), ('jt_size', uint32_t, 44)])
+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)])
 @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', uint32_t, 48)])
+struct_sdma_firmware_header_v1_1.register_fields([('v1_0', struct_sdma_firmware_header_v1_0, 0), ('digest_size', ctypes.c_uint32, 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', 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)])
+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)])
 @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', 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)])
+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)])
 @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', 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)])
+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)])
 @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', uint32_t, 32), ('ucode_offset_bytes', uint32_t, 36), ('ucode_size_bytes', uint32_t, 40)])
+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)])
 @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', 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)])
+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)])
 @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', uint32_t, 60), ('num_packer_per_sc', uint32_t, 64)])
+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)])
 @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', uint16_t, 32), ('version_minor', uint16_t, 34)])
+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)])
 @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', uint32_t, 32), ('intv_size_bytes', uint32_t, 36)])
+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)])
 @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', uint32_t, 32), ('bss_data_bytes', uint32_t, 36)])
+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)])
 @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', 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)])
+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)])
 @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[uint8_t, 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[ctypes.c_ubyte, 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', 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)])
+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)])
 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', 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)])
+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)])
 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,199 +4110,201 @@ 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
-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_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_PTE_FRAG = lambda x: ((x & 0x1f) << 7) # 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_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_PDE_BFS = lambda a: (a << 59) # 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_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_PTE_MTYPE_VG10_SHIFT = lambda mtype: ((mtype) << 57) # type: ignore
-AMDGPU_PTE_MTYPE_VG10_MASK = AMDGPU_PTE_MTYPE_VG10_SHIFT(3)
+AMDGPU_PTE_MTYPE_VG10_MASK = AMDGPU_PTE_MTYPE_VG10_SHIFT(3) # type: ignore
 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
-AMDGPU_MTYPE_CC = 2
+AMDGPU_MTYPE_NC = 0 # type: ignore
+AMDGPU_MTYPE_CC = 2 # type: ignore
 AMDGPU_PTE_MTYPE_NV10_SHIFT = lambda mtype: ((mtype) << 48) # type: ignore
-AMDGPU_PTE_MTYPE_NV10_MASK = AMDGPU_PTE_MTYPE_NV10_SHIFT(7)
+AMDGPU_PTE_MTYPE_NV10_MASK = AMDGPU_PTE_MTYPE_NV10_SHIFT(7) # type: ignore
 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)
+AMDGPU_PTE_PRT_GFX12 = (1 << 56) # type: ignore
 AMDGPU_PTE_MTYPE_GFX12_SHIFT = lambda mtype: ((mtype) << 54) # type: ignore
-AMDGPU_PTE_MTYPE_GFX12_MASK = AMDGPU_PTE_MTYPE_GFX12_SHIFT(3)
+AMDGPU_PTE_MTYPE_GFX12_MASK = AMDGPU_PTE_MTYPE_GFX12_SHIFT(3) # type: ignore
 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)
+AMDGPU_PTE_IS_PTE = (1 << 63) # type: ignore
 AMDGPU_PDE_BFS_GFX12 = lambda a: (((a) & 0x1f) << 58) # 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_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_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)
-AMDGPU_VA_RESERVED_SEQ64_SIZE = (2 << 20)
+AMDGPU_VA_RESERVED_CSA_SIZE = (2 << 20) # type: ignore
+AMDGPU_VA_RESERVED_SEQ64_SIZE = (2 << 20) # type: ignore
 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)
+AMDGPU_VA_RESERVED_TRAP_SIZE = (2 << 12) # type: ignore
 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)
-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
+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
 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
@@ -4314,155 +4316,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
-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
+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

tinygrad/runtime/autogen/am/fw.py

@@ -1,106 +0,0 @@
-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/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
-PACKET_TYPE1 = 1
-PACKET_TYPE2 = 2
-PACKET_TYPE3 = 3
+PACKET_TYPE0 = 0 # type: ignore
+PACKET_TYPE1 = 1 # type: ignore
+PACKET_TYPE2 = 2 # type: ignore
+PACKET_TYPE3 = 3 # type: ignore
 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
-PACKET2_PAD_SHIFT = 0
-PACKET2_PAD_MASK = (0x3fffffff << 0)
+CP_PACKET2 = 0x80000000 # type: ignore
+PACKET2_PAD_SHIFT = 0 # type: ignore
+PACKET2_PAD_MASK = (0x3fffffff << 0) # type: ignore
 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
-PACKET3_SET_BASE = 0x11
+PACKET3_NOP = 0x10 # type: ignore
+PACKET3_SET_BASE = 0x11 # type: ignore
 PACKET3_BASE_INDEX = lambda x: ((x) << 0) # 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
+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
 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
-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
+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
 WRITE_DATA_DST_SEL = lambda x: ((x) << 8) # type: ignore
-WR_ONE_ADDR = (1 << 16)
-WR_CONFIRM = (1 << 20)
+WR_ONE_ADDR = (1 << 16) # type: ignore
+WR_CONFIRM = (1 << 20) # type: ignore
 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
-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
+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
 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
-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)
+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
 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
-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_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_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
-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
+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
 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
-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_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_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)
-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_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_CACHE_POLICY = lambda x: ((x) << 25) # type: ignore
-PACKET3_RELEASE_MEM_EXECUTE = (1 << 28)
+PACKET3_RELEASE_MEM_EXECUTE = (1 << 28) # type: ignore
 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
-PACKET3_PREAMBLE_BEGIN_CLEAR_STATE = (2 << 28)
-PACKET3_PREAMBLE_END_CLEAR_STATE = (3 << 28)
-PACKET3_DMA_DATA = 0x50
+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_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)
-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_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_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)
+PACKET3_ACQUIRE_MEM_GCR_RANGE_IS_PA = (1 << 18) # type: ignore
 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
-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_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_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
-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_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_UCONFIG_REG__REG_OFFSET = lambda x: ((((unsigned)(x)) & 0xFFFF) << 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)
+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
 FRAME_CMD = lambda x: ((x) << 28) # 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_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_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
-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_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_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
-PACKET3_MAP_QUEUES = 0xA2
+PACKET3_MAP_PROCESS = 0xA1 # type: ignore
+PACKET3_MAP_QUEUES = 0xA2 # type: ignore
 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
+PACKET3_UNMAP_QUEUES = 0xA3 # type: ignore
 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
+PACKET3_QUERY_STATUS = 0xA4 # type: ignore
 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
-PACKET3_MAP_PROCESS_VM = 0xA6
-PACKET3_RUN_CLEANER_SHADER = 0xD2
-PACKET3_SET_Q_PREEMPTION_MODE = 0xF0
+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_SET_Q_PREEMPTION_MODE_IB_VMID = lambda x: ((x) << 0) # type: ignore
-PACKET3_SET_Q_PREEMPTION_MODE_INIT_SHADOW_MEM = (1 << 0)
+PACKET3_SET_Q_PREEMPTION_MODE_INIT_SHADOW_MEM = (1 << 0) # type: ignore

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
-GFX9_NUM_COMPUTE_RINGS = 8
-PACKET_TYPE0 = 0
-PACKET_TYPE1 = 1
-PACKET_TYPE2 = 2
-PACKET_TYPE3 = 3
+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
 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
-PACKET2_PAD_SHIFT = 0
-PACKET2_PAD_MASK = (0x3fffffff << 0)
+CP_PACKET2 = 0x80000000 # type: ignore
+PACKET2_PAD_SHIFT = 0 # type: ignore
+PACKET2_PAD_MASK = (0x3fffffff << 0) # type: ignore
 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
-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_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 = lambda reg,r,cond,type: ((reg & 0x3FFFF) | ((r & 0x3F) << 18) | ((cond & 0xF) << 24) | ((type & 0xF) << 28)) # type: ignore
-CP_PACKETJ_NOP = 0x60000000
+CP_PACKETJ_NOP = 0x60000000 # type: ignore
 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
-PACKET3_SET_BASE = 0x11
+PACKET3_NOP = 0x10 # type: ignore
+PACKET3_SET_BASE = 0x11 # type: ignore
 PACKET3_BASE_INDEX = lambda x: ((x) << 0) # 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
+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
 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
-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_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_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
-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
+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
 WRITE_DATA_DST_SEL = lambda x: ((x) << 8) # type: ignore
-WR_ONE_ADDR = (1 << 16)
-WR_CONFIRM = (1 << 20)
+WR_ONE_ADDR = (1 << 16) # type: ignore
+WR_CONFIRM = (1 << 20) # type: ignore
 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
-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
+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
 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
-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)
+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
 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
-PACKET3_INDIRECT_BUFFER__CACHE_POLICY__STREAM = 1
-PACKET3_COPY_DATA = 0x40
+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_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
-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
+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
 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
-PACKET3_EVENT_WRITE__EVENT_INDEX__SAMPLE_PIPELINESTATS = 2
-PACKET3_EVENT_WRITE__EVENT_INDEX__CS_PARTIAL_FLUSH = 4
-PACKET3_RELEASE_MEM = 0x49
+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
 EVENT_TYPE = lambda x: ((x) << 0) # type: ignore
 EVENT_INDEX = lambda x: ((x) << 8) # 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)
+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
 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
-PACKET3_PREAMBLE_BEGIN_CLEAR_STATE = (2 << 28)
-PACKET3_PREAMBLE_END_CLEAR_STATE = (3 << 28)
-PACKET3_DMA_DATA = 0x50
+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_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)
-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_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_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
+PACKET3_REWIND = 0x59 # type: ignore
 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
-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_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_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
-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_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_UCONFIG_REG__REG_OFFSET = lambda x: ((((unsigned)(x)) & 0xFFFF) << 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)
+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
 FRAME_CMD = lambda x: ((x) << 28) # type: ignore
-PACKET3_INVALIDATE_TLBS = 0x98
+PACKET3_INVALIDATE_TLBS = 0x98 # type: ignore
 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
+PACKET3_SET_RESOURCES = 0xA0 # type: ignore
 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
+PACKET3_MAP_QUEUES = 0xA2 # type: ignore
 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
+PACKET3_UNMAP_QUEUES = 0xA3 # type: ignore
 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
+PACKET3_QUERY_STATUS = 0xA4 # type: ignore
 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
-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
+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

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', uint8_t, 1), ('unused', c.Array[uint8_t, Literal[3]], 2)])
+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)])
 @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', uint8_t, 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', ctypes.c_ubyte, 2)])
 @c.record
 class struct_smu_uvd_clocks(c.Struct):
   SIZE = 8
   vclk: int
   dclk: int
-struct_smu_uvd_clocks.register_fields([('vclk', uint32_t, 0), ('dclk', uint32_t, 4)])
+struct_smu_uvd_clocks.register_fields([('vclk', ctypes.c_uint32, 0), ('dclk', ctypes.c_uint32, 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', 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)])
+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)])
 @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', 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)])
+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)])
 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', 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)])
+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)])
 @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', 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)])
+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)])
 @c.record
 class struct_smu_bios_boot_up_values(c.Struct):
   SIZE = 68
@@ -867,359 +867,360 @@ 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', 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)])
+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)])
 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
-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
+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

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', uint8_t, 1), ('unused', c.Array[uint8_t, Literal[3]], 2)])
+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)])
 @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', uint8_t, 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', ctypes.c_ubyte, 2)])
 @c.record
 class struct_smu_uvd_clocks(c.Struct):
   SIZE = 8
   vclk: int
   dclk: int
-struct_smu_uvd_clocks.register_fields([('vclk', uint32_t, 0), ('dclk', uint32_t, 4)])
+struct_smu_uvd_clocks.register_fields([('vclk', ctypes.c_uint32, 0), ('dclk', ctypes.c_uint32, 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', 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)])
+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)])
 @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', 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)])
+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)])
 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', 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)])
+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)])
 @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', 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)])
+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)])
 @c.record
 class struct_smu_bios_boot_up_values(c.Struct):
   SIZE = 68
@@ -337,196 +337,197 @@ 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', 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)])
+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)])
 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
-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
+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

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', uint8_t, 1), ('unused', c.Array[uint8_t, Literal[3]], 2)])
+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)])
 @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', uint8_t, 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', ctypes.c_ubyte, 2)])
 @c.record
 class struct_smu_uvd_clocks(c.Struct):
   SIZE = 8
   vclk: int
   dclk: int
-struct_smu_uvd_clocks.register_fields([('vclk', uint32_t, 0), ('dclk', uint32_t, 4)])
+struct_smu_uvd_clocks.register_fields([('vclk', ctypes.c_uint32, 0), ('dclk', ctypes.c_uint32, 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', 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)])
+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)])
 @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', 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)])
+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)])
 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', 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)])
+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)])
 @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', 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)])
+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)])
 @c.record
 class struct_smu_bios_boot_up_values(c.Struct):
   SIZE = 68
@@ -448,170 +448,171 @@ 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', 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)])
+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)])
 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
-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
+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

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', uint8_t, 1), ('unused', c.Array[uint8_t, Literal[3]], 2)])
+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)])
 @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', uint8_t, 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', ctypes.c_ubyte, 2)])
 @c.record
 class struct_smu_uvd_clocks(c.Struct):
   SIZE = 8
   vclk: int
   dclk: int
-struct_smu_uvd_clocks.register_fields([('vclk', uint32_t, 0), ('dclk', uint32_t, 4)])
+struct_smu_uvd_clocks.register_fields([('vclk', ctypes.c_uint32, 0), ('dclk', ctypes.c_uint32, 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', 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)])
+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)])
 @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', 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)])
+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)])
 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', 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)])
+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)])
 @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', 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)])
+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)])
 @c.record
 class struct_smu_bios_boot_up_values(c.Struct):
   SIZE = 68
@@ -1121,443 +1121,444 @@ 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', 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)])
+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)])
 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
-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
+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

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
-AMD_COMGR_INTERFACE_VERSION_MINOR = 8
+AMD_COMGR_INTERFACE_VERSION_MAJOR = 2 # type: ignore
+AMD_COMGR_INTERFACE_VERSION_MINOR = 8 # type: ignore

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
-AMD_COMGR_INTERFACE_VERSION_MINOR = 0
+AMD_COMGR_INTERFACE_VERSION_MAJOR = 3 # type: ignore
+AMD_COMGR_INTERFACE_VERSION_MINOR = 0 # type: ignore

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
-kCFStringEncodingInvalidId = (0xffffffff)
+__COREFOUNDATION_CFSTRING__ = 1 # type: ignore
+kCFStringEncodingInvalidId = (0xffffffff) # type: ignore
 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
-__COREFOUNDATION_CFDATA__ = 1
+__kCFStringInlineBufferLength = 64 # type: ignore
+__COREFOUNDATION_CFDATA__ = 1 # type: ignore

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
-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
+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
 HIP_DEPRECATED = lambda msg: __attribute__((deprecated(msg))) # 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
+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

tinygrad/runtime/autogen/hsa.py

@@ -1629,15 +1629,15 @@ class struct_hsa_ven_amd_aqlprofile_1_00_pfn_s(c.Struct):
 struct_hsa_ven_amd_aqlprofile_1_00_pfn_s.register_fields([('hsa_ven_amd_aqlprofile_version_major', c.CFUNCTYPE[uint32_t, []], 0), ('hsa_ven_amd_aqlprofile_version_minor', c.CFUNCTYPE[uint32_t, []], 8), ('hsa_ven_amd_aqlprofile_error_string', c.CFUNCTYPE[ctypes.c_uint32, [c.POINTER[c.POINTER[ctypes.c_char]]]], 16), ('hsa_ven_amd_aqlprofile_validate_event', c.CFUNCTYPE[ctypes.c_uint32, [hsa_agent_t, c.POINTER[hsa_ven_amd_aqlprofile_event_t], c.POINTER[ctypes.c_bool]]], 24), ('hsa_ven_amd_aqlprofile_start', c.CFUNCTYPE[ctypes.c_uint32, [c.POINTER[hsa_ven_amd_aqlprofile_profile_t], c.POINTER[hsa_ext_amd_aql_pm4_packet_t]]], 32), ('hsa_ven_amd_aqlprofile_stop', c.CFUNCTYPE[ctypes.c_uint32, [c.POINTER[hsa_ven_amd_aqlprofile_profile_t], c.POINTER[hsa_ext_amd_aql_pm4_packet_t]]], 40), ('hsa_ven_amd_aqlprofile_read', c.CFUNCTYPE[ctypes.c_uint32, [c.POINTER[hsa_ven_amd_aqlprofile_profile_t], c.POINTER[hsa_ext_amd_aql_pm4_packet_t]]], 48), ('hsa_ven_amd_aqlprofile_legacy_get_pm4', c.CFUNCTYPE[ctypes.c_uint32, [c.POINTER[hsa_ext_amd_aql_pm4_packet_t], ctypes.c_void_p]], 56), ('hsa_ven_amd_aqlprofile_get_info', c.CFUNCTYPE[ctypes.c_uint32, [c.POINTER[hsa_ven_amd_aqlprofile_profile_t], ctypes.c_uint32, ctypes.c_void_p]], 64), ('hsa_ven_amd_aqlprofile_iterate_data', c.CFUNCTYPE[ctypes.c_uint32, [c.POINTER[hsa_ven_amd_aqlprofile_profile_t], hsa_ven_amd_aqlprofile_data_callback_t, ctypes.c_void_p]], 72), ('hsa_ven_amd_aqlprofile_iterate_event_ids', c.CFUNCTYPE[ctypes.c_uint32, [hsa_ven_amd_aqlprofile_eventname_callback_t]], 80), ('hsa_ven_amd_aqlprofile_iterate_event_coord', c.CFUNCTYPE[ctypes.c_uint32, [hsa_agent_t, hsa_ven_amd_aqlprofile_event_t, uint32_t, hsa_ven_amd_aqlprofile_coordinate_callback_t, ctypes.c_void_p]], 88), ('hsa_ven_amd_aqlprofile_att_marker', c.CFUNCTYPE[ctypes.c_uint32, [c.POINTER[hsa_ven_amd_aqlprofile_profile_t], c.POINTER[hsa_ext_amd_aql_pm4_packet_t], uint32_t, ctypes.c_uint32]], 96)])
 hsa_ven_amd_aqlprofile_1_00_pfn_t: TypeAlias = struct_hsa_ven_amd_aqlprofile_1_00_pfn_s
 hsa_ven_amd_aqlprofile_pfn_t: TypeAlias = struct_hsa_ven_amd_aqlprofile_1_00_pfn_s
-HSA_VERSION_1_0 = 1
-HSA_AMD_INTERFACE_VERSION_MAJOR = 1
-HSA_AMD_INTERFACE_VERSION_MINOR = 14
-AMD_SIGNAL_ALIGN_BYTES = 64
-AMD_QUEUE_ALIGN_BYTES = 64
-MAX_NUM_XCC = 128
-AMD_CONTROL_DIRECTIVES_ALIGN_BYTES = 64
-AMD_ISA_ALIGN_BYTES = 256
-AMD_KERNEL_CODE_ALIGN_BYTES = 64
-HSA_AQLPROFILE_VERSION_MAJOR = 2
-HSA_AQLPROFILE_VERSION_MINOR = 0
-hsa_ven_amd_aqlprofile_VERSION_MAJOR = 1
+HSA_VERSION_1_0 = 1 # type: ignore
+HSA_AMD_INTERFACE_VERSION_MAJOR = 1 # type: ignore
+HSA_AMD_INTERFACE_VERSION_MINOR = 14 # type: ignore
+AMD_SIGNAL_ALIGN_BYTES = 64 # type: ignore
+AMD_QUEUE_ALIGN_BYTES = 64 # type: ignore
+MAX_NUM_XCC = 128 # type: ignore
+AMD_CONTROL_DIRECTIVES_ALIGN_BYTES = 64 # type: ignore
+AMD_ISA_ALIGN_BYTES = 256 # type: ignore
+AMD_KERNEL_CODE_ALIGN_BYTES = 64 # type: ignore
+HSA_AQLPROFILE_VERSION_MAJOR = 2 # type: ignore
+HSA_AQLPROFILE_VERSION_MINOR = 0 # type: ignore
+hsa_ven_amd_aqlprofile_VERSION_MAJOR = 1 # type: ignore

tinygrad/runtime/autogen/ib.py

@@ -2685,37 +2685,37 @@ struct_ib_uverbs_ex_modify_cq.register_fields([('cq_handle', ctypes.c_uint32, 0)
 enum_ib_uverbs_device_cap_flags: dict[int, str] = {(IB_UVERBS_DEVICE_RESIZE_MAX_WR:=1): 'IB_UVERBS_DEVICE_RESIZE_MAX_WR', (IB_UVERBS_DEVICE_BAD_PKEY_CNTR:=2): 'IB_UVERBS_DEVICE_BAD_PKEY_CNTR', (IB_UVERBS_DEVICE_BAD_QKEY_CNTR:=4): 'IB_UVERBS_DEVICE_BAD_QKEY_CNTR', (IB_UVERBS_DEVICE_RAW_MULTI:=8): 'IB_UVERBS_DEVICE_RAW_MULTI', (IB_UVERBS_DEVICE_AUTO_PATH_MIG:=16): 'IB_UVERBS_DEVICE_AUTO_PATH_MIG', (IB_UVERBS_DEVICE_CHANGE_PHY_PORT:=32): 'IB_UVERBS_DEVICE_CHANGE_PHY_PORT', (IB_UVERBS_DEVICE_UD_AV_PORT_ENFORCE:=64): 'IB_UVERBS_DEVICE_UD_AV_PORT_ENFORCE', (IB_UVERBS_DEVICE_CURR_QP_STATE_MOD:=128): 'IB_UVERBS_DEVICE_CURR_QP_STATE_MOD', (IB_UVERBS_DEVICE_SHUTDOWN_PORT:=256): 'IB_UVERBS_DEVICE_SHUTDOWN_PORT', (IB_UVERBS_DEVICE_PORT_ACTIVE_EVENT:=1024): 'IB_UVERBS_DEVICE_PORT_ACTIVE_EVENT', (IB_UVERBS_DEVICE_SYS_IMAGE_GUID:=2048): 'IB_UVERBS_DEVICE_SYS_IMAGE_GUID', (IB_UVERBS_DEVICE_RC_RNR_NAK_GEN:=4096): 'IB_UVERBS_DEVICE_RC_RNR_NAK_GEN', (IB_UVERBS_DEVICE_SRQ_RESIZE:=8192): 'IB_UVERBS_DEVICE_SRQ_RESIZE', (IB_UVERBS_DEVICE_N_NOTIFY_CQ:=16384): 'IB_UVERBS_DEVICE_N_NOTIFY_CQ', (IB_UVERBS_DEVICE_MEM_WINDOW:=131072): 'IB_UVERBS_DEVICE_MEM_WINDOW', (IB_UVERBS_DEVICE_UD_IP_CSUM:=262144): 'IB_UVERBS_DEVICE_UD_IP_CSUM', (IB_UVERBS_DEVICE_XRC:=1048576): 'IB_UVERBS_DEVICE_XRC', (IB_UVERBS_DEVICE_MEM_MGT_EXTENSIONS:=2097152): 'IB_UVERBS_DEVICE_MEM_MGT_EXTENSIONS', (IB_UVERBS_DEVICE_MEM_WINDOW_TYPE_2A:=8388608): 'IB_UVERBS_DEVICE_MEM_WINDOW_TYPE_2A', (IB_UVERBS_DEVICE_MEM_WINDOW_TYPE_2B:=16777216): 'IB_UVERBS_DEVICE_MEM_WINDOW_TYPE_2B', (IB_UVERBS_DEVICE_RC_IP_CSUM:=33554432): 'IB_UVERBS_DEVICE_RC_IP_CSUM', (IB_UVERBS_DEVICE_RAW_IP_CSUM:=67108864): 'IB_UVERBS_DEVICE_RAW_IP_CSUM', (IB_UVERBS_DEVICE_MANAGED_FLOW_STEERING:=536870912): 'IB_UVERBS_DEVICE_MANAGED_FLOW_STEERING', (IB_UVERBS_DEVICE_RAW_SCATTER_FCS:=17179869184): 'IB_UVERBS_DEVICE_RAW_SCATTER_FCS', (IB_UVERBS_DEVICE_PCI_WRITE_END_PADDING:=68719476736): 'IB_UVERBS_DEVICE_PCI_WRITE_END_PADDING', (IB_UVERBS_DEVICE_FLUSH_GLOBAL:=274877906944): 'IB_UVERBS_DEVICE_FLUSH_GLOBAL', (IB_UVERBS_DEVICE_FLUSH_PERSISTENT:=549755813888): 'IB_UVERBS_DEVICE_FLUSH_PERSISTENT', (IB_UVERBS_DEVICE_ATOMIC_WRITE:=1099511627776): 'IB_UVERBS_DEVICE_ATOMIC_WRITE'}
 enum_ib_uverbs_raw_packet_caps: dict[int, str] = {(IB_UVERBS_RAW_PACKET_CAP_CVLAN_STRIPPING:=1): 'IB_UVERBS_RAW_PACKET_CAP_CVLAN_STRIPPING', (IB_UVERBS_RAW_PACKET_CAP_SCATTER_FCS:=2): 'IB_UVERBS_RAW_PACKET_CAP_SCATTER_FCS', (IB_UVERBS_RAW_PACKET_CAP_IP_CSUM:=4): 'IB_UVERBS_RAW_PACKET_CAP_IP_CSUM', (IB_UVERBS_RAW_PACKET_CAP_DELAY_DROP:=8): 'IB_UVERBS_RAW_PACKET_CAP_DELAY_DROP'}
 vext_field_avail = lambda type,fld,sz: (offsetof(type, fld) < (sz)) # type: ignore
-IBV_DEVICE_RAW_SCATTER_FCS = (1 << 34)
-IBV_DEVICE_PCI_WRITE_END_PADDING = (1 << 36)
+IBV_DEVICE_RAW_SCATTER_FCS = (1 << 34) # type: ignore
+IBV_DEVICE_PCI_WRITE_END_PADDING = (1 << 36) # type: ignore
 ibv_query_port = lambda context,port_num,port_attr: ___ibv_query_port(context, port_num, port_attr) # type: ignore
 ibv_reg_mr = lambda pd,addr,length,access: __ibv_reg_mr(pd, addr, length, access, __builtin_constant_p( ((int)(access) & IBV_ACCESS_OPTIONAL_RANGE) == 0)) # type: ignore
 ibv_reg_mr_iova = lambda pd,addr,length,iova,access: __ibv_reg_mr_iova(pd, addr, length, iova, access, __builtin_constant_p( ((access) & IBV_ACCESS_OPTIONAL_RANGE) == 0)) # type: ignore
-ETHERNET_LL_SIZE = 6
-IB_ROCE_UDP_ENCAP_VALID_PORT_MIN = (0xC000)
-IB_ROCE_UDP_ENCAP_VALID_PORT_MAX = (0xFFFF)
-IB_GRH_FLOWLABEL_MASK = (0x000FFFFF)
-IBV_FLOW_ACTION_ESP_KEYMAT_AES_GCM = IB_UVERBS_FLOW_ACTION_ESP_KEYMAT_AES_GCM
-IBV_FLOW_ACTION_IV_ALGO_SEQ = IB_UVERBS_FLOW_ACTION_IV_ALGO_SEQ
-IBV_FLOW_ACTION_ESP_REPLAY_NONE = IB_UVERBS_FLOW_ACTION_ESP_REPLAY_NONE
-IBV_FLOW_ACTION_ESP_REPLAY_BMP = IB_UVERBS_FLOW_ACTION_ESP_REPLAY_BMP
-IBV_FLOW_ACTION_ESP_FLAGS_INLINE_CRYPTO = IB_UVERBS_FLOW_ACTION_ESP_FLAGS_INLINE_CRYPTO
-IBV_FLOW_ACTION_ESP_FLAGS_FULL_OFFLOAD = IB_UVERBS_FLOW_ACTION_ESP_FLAGS_FULL_OFFLOAD
-IBV_FLOW_ACTION_ESP_FLAGS_TUNNEL = IB_UVERBS_FLOW_ACTION_ESP_FLAGS_TUNNEL
-IBV_FLOW_ACTION_ESP_FLAGS_TRANSPORT = IB_UVERBS_FLOW_ACTION_ESP_FLAGS_TRANSPORT
-IBV_FLOW_ACTION_ESP_FLAGS_DECRYPT = IB_UVERBS_FLOW_ACTION_ESP_FLAGS_DECRYPT
-IBV_FLOW_ACTION_ESP_FLAGS_ENCRYPT = IB_UVERBS_FLOW_ACTION_ESP_FLAGS_ENCRYPT
-IBV_FLOW_ACTION_ESP_FLAGS_ESN_NEW_WINDOW = IB_UVERBS_FLOW_ACTION_ESP_FLAGS_ESN_NEW_WINDOW
-IBV_ADVISE_MR_ADVICE_PREFETCH = IB_UVERBS_ADVISE_MR_ADVICE_PREFETCH
-IBV_ADVISE_MR_ADVICE_PREFETCH_WRITE = IB_UVERBS_ADVISE_MR_ADVICE_PREFETCH_WRITE
-IBV_ADVISE_MR_ADVICE_PREFETCH_NO_FAULT = IB_UVERBS_ADVISE_MR_ADVICE_PREFETCH_NO_FAULT
-IBV_ADVISE_MR_FLAG_FLUSH = IB_UVERBS_ADVISE_MR_FLAG_FLUSH
-IBV_QPF_GRH_REQUIRED = IB_UVERBS_QPF_GRH_REQUIRED
-IBV_ACCESS_OPTIONAL_RANGE = IB_UVERBS_ACCESS_OPTIONAL_RANGE
-IB_UVERBS_ACCESS_OPTIONAL_FIRST = (1 << 20)
-IB_UVERBS_ACCESS_OPTIONAL_LAST = (1 << 29)
-IB_USER_VERBS_ABI_VERSION = 6
-IB_USER_VERBS_CMD_THRESHOLD = 50
-IB_USER_VERBS_CMD_COMMAND_MASK = 0xff
-IB_USER_VERBS_CMD_FLAG_EXTENDED = 0x80000000
-IB_USER_VERBS_MAX_LOG_IND_TBL_SIZE = 0x0d
-IB_DEVICE_NAME_MAX = 64
+ETHERNET_LL_SIZE = 6 # type: ignore
+IB_ROCE_UDP_ENCAP_VALID_PORT_MIN = (0xC000) # type: ignore
+IB_ROCE_UDP_ENCAP_VALID_PORT_MAX = (0xFFFF) # type: ignore
+IB_GRH_FLOWLABEL_MASK = (0x000FFFFF) # type: ignore
+IBV_FLOW_ACTION_ESP_KEYMAT_AES_GCM = IB_UVERBS_FLOW_ACTION_ESP_KEYMAT_AES_GCM # type: ignore
+IBV_FLOW_ACTION_IV_ALGO_SEQ = IB_UVERBS_FLOW_ACTION_IV_ALGO_SEQ # type: ignore
+IBV_FLOW_ACTION_ESP_REPLAY_NONE = IB_UVERBS_FLOW_ACTION_ESP_REPLAY_NONE # type: ignore
+IBV_FLOW_ACTION_ESP_REPLAY_BMP = IB_UVERBS_FLOW_ACTION_ESP_REPLAY_BMP # type: ignore
+IBV_FLOW_ACTION_ESP_FLAGS_INLINE_CRYPTO = IB_UVERBS_FLOW_ACTION_ESP_FLAGS_INLINE_CRYPTO # type: ignore
+IBV_FLOW_ACTION_ESP_FLAGS_FULL_OFFLOAD = IB_UVERBS_FLOW_ACTION_ESP_FLAGS_FULL_OFFLOAD # type: ignore
+IBV_FLOW_ACTION_ESP_FLAGS_TUNNEL = IB_UVERBS_FLOW_ACTION_ESP_FLAGS_TUNNEL # type: ignore
+IBV_FLOW_ACTION_ESP_FLAGS_TRANSPORT = IB_UVERBS_FLOW_ACTION_ESP_FLAGS_TRANSPORT # type: ignore
+IBV_FLOW_ACTION_ESP_FLAGS_DECRYPT = IB_UVERBS_FLOW_ACTION_ESP_FLAGS_DECRYPT # type: ignore
+IBV_FLOW_ACTION_ESP_FLAGS_ENCRYPT = IB_UVERBS_FLOW_ACTION_ESP_FLAGS_ENCRYPT # type: ignore
+IBV_FLOW_ACTION_ESP_FLAGS_ESN_NEW_WINDOW = IB_UVERBS_FLOW_ACTION_ESP_FLAGS_ESN_NEW_WINDOW # type: ignore
+IBV_ADVISE_MR_ADVICE_PREFETCH = IB_UVERBS_ADVISE_MR_ADVICE_PREFETCH # type: ignore
+IBV_ADVISE_MR_ADVICE_PREFETCH_WRITE = IB_UVERBS_ADVISE_MR_ADVICE_PREFETCH_WRITE # type: ignore
+IBV_ADVISE_MR_ADVICE_PREFETCH_NO_FAULT = IB_UVERBS_ADVISE_MR_ADVICE_PREFETCH_NO_FAULT # type: ignore
+IBV_ADVISE_MR_FLAG_FLUSH = IB_UVERBS_ADVISE_MR_FLAG_FLUSH # type: ignore
+IBV_QPF_GRH_REQUIRED = IB_UVERBS_QPF_GRH_REQUIRED # type: ignore
+IBV_ACCESS_OPTIONAL_RANGE = IB_UVERBS_ACCESS_OPTIONAL_RANGE # type: ignore
+IB_UVERBS_ACCESS_OPTIONAL_FIRST = (1 << 20) # type: ignore
+IB_UVERBS_ACCESS_OPTIONAL_LAST = (1 << 29) # type: ignore
+IB_USER_VERBS_ABI_VERSION = 6 # type: ignore
+IB_USER_VERBS_CMD_THRESHOLD = 50 # type: ignore
+IB_USER_VERBS_CMD_COMMAND_MASK = 0xff # type: ignore
+IB_USER_VERBS_CMD_FLAG_EXTENDED = 0x80000000 # type: ignore
+IB_USER_VERBS_MAX_LOG_IND_TBL_SIZE = 0x0d # type: ignore
+IB_DEVICE_NAME_MAX = 64 # type: ignore

tinygrad/runtime/autogen/io_uring.py

@@ -444,461 +444,461 @@ class struct_io_uring_zcrx_ifq_reg(c.Struct):
 struct_io_uring_zcrx_ifq_reg.register_fields([('if_idx', ctypes.c_uint32, 0), ('if_rxq', ctypes.c_uint32, 4), ('rq_entries', ctypes.c_uint32, 8), ('flags', ctypes.c_uint32, 12), ('area_ptr', ctypes.c_uint64, 16), ('region_ptr', ctypes.c_uint64, 24), ('offsets', struct_io_uring_zcrx_offsets, 32), ('zcrx_id', ctypes.c_uint32, 64), ('__resv2', ctypes.c_uint32, 68), ('__resv', c.Array[ctypes.c_uint64, Literal[3]], 72)])
 uring_unlikely = lambda cond: __builtin_expect( not  not (cond), 0) # type: ignore
 uring_likely = lambda cond: __builtin_expect( not  not (cond), 1) # type: ignore
-NR_io_uring_setup = 425
-NR_io_uring_enter = 426
-NR_io_uring_register = 427
+NR_io_uring_setup = 425 # type: ignore
+NR_io_uring_enter = 426 # type: ignore
+NR_io_uring_register = 427 # type: ignore
 IO_URING_CHECK_VERSION = lambda major,minor: (major > IO_URING_VERSION_MAJOR or (major == IO_URING_VERSION_MAJOR and minor > IO_URING_VERSION_MINOR)) # type: ignore
-IORING_RW_ATTR_FLAG_PI = (1 << 0)
-IORING_FILE_INDEX_ALLOC = (~0)
-IOSQE_FIXED_FILE = (1 << IOSQE_FIXED_FILE_BIT)
-IOSQE_IO_DRAIN = (1 << IOSQE_IO_DRAIN_BIT)
-IOSQE_IO_LINK = (1 << IOSQE_IO_LINK_BIT)
-IOSQE_IO_HARDLINK = (1 << IOSQE_IO_HARDLINK_BIT)
-IOSQE_ASYNC = (1 << IOSQE_ASYNC_BIT)
-IOSQE_BUFFER_SELECT = (1 << IOSQE_BUFFER_SELECT_BIT)
-IOSQE_CQE_SKIP_SUCCESS = (1 << IOSQE_CQE_SKIP_SUCCESS_BIT)
-IORING_SETUP_IOPOLL = (1 << 0)
-IORING_SETUP_SQPOLL = (1 << 1)
-IORING_SETUP_SQ_AFF = (1 << 2)
-IORING_SETUP_CQSIZE = (1 << 3)
-IORING_SETUP_CLAMP = (1 << 4)
-IORING_SETUP_ATTACH_WQ = (1 << 5)
-IORING_SETUP_R_DISABLED = (1 << 6)
-IORING_SETUP_SUBMIT_ALL = (1 << 7)
-IORING_SETUP_COOP_TASKRUN = (1 << 8)
-IORING_SETUP_TASKRUN_FLAG = (1 << 9)
-IORING_SETUP_SQE128 = (1 << 10)
-IORING_SETUP_CQE32 = (1 << 11)
-IORING_SETUP_SINGLE_ISSUER = (1 << 12)
-IORING_SETUP_DEFER_TASKRUN = (1 << 13)
-IORING_SETUP_NO_MMAP = (1 << 14)
-IORING_SETUP_REGISTERED_FD_ONLY = (1 << 15)
-IORING_SETUP_NO_SQARRAY = (1 << 16)
-IORING_SETUP_HYBRID_IOPOLL = (1 << 17)
-IORING_SETUP_CQE_MIXED = (1 << 18)
-IORING_URING_CMD_FIXED = (1 << 0)
-IORING_URING_CMD_MULTISHOT = (1 << 1)
-IORING_URING_CMD_MASK = (IORING_URING_CMD_FIXED | IORING_URING_CMD_MULTISHOT)
-IORING_FSYNC_DATASYNC = (1 << 0)
-IORING_TIMEOUT_ABS = (1 << 0)
-IORING_TIMEOUT_UPDATE = (1 << 1)
-IORING_TIMEOUT_BOOTTIME = (1 << 2)
-IORING_TIMEOUT_REALTIME = (1 << 3)
-IORING_LINK_TIMEOUT_UPDATE = (1 << 4)
-IORING_TIMEOUT_ETIME_SUCCESS = (1 << 5)
-IORING_TIMEOUT_MULTISHOT = (1 << 6)
-IORING_TIMEOUT_CLOCK_MASK = (IORING_TIMEOUT_BOOTTIME | IORING_TIMEOUT_REALTIME)
-IORING_TIMEOUT_UPDATE_MASK = (IORING_TIMEOUT_UPDATE | IORING_LINK_TIMEOUT_UPDATE)
-SPLICE_F_FD_IN_FIXED = (1 << 31)
-IORING_POLL_ADD_MULTI = (1 << 0)
-IORING_POLL_UPDATE_EVENTS = (1 << 1)
-IORING_POLL_UPDATE_USER_DATA = (1 << 2)
-IORING_POLL_ADD_LEVEL = (1 << 3)
-IORING_ASYNC_CANCEL_ALL = (1 << 0)
-IORING_ASYNC_CANCEL_FD = (1 << 1)
-IORING_ASYNC_CANCEL_ANY = (1 << 2)
-IORING_ASYNC_CANCEL_FD_FIXED = (1 << 3)
-IORING_ASYNC_CANCEL_USERDATA = (1 << 4)
-IORING_ASYNC_CANCEL_OP = (1 << 5)
-IORING_RECVSEND_POLL_FIRST = (1 << 0)
-IORING_RECV_MULTISHOT = (1 << 1)
-IORING_RECVSEND_FIXED_BUF = (1 << 2)
-IORING_SEND_ZC_REPORT_USAGE = (1 << 3)
-IORING_RECVSEND_BUNDLE = (1 << 4)
-IORING_SEND_VECTORIZED = (1 << 5)
-IORING_NOTIF_USAGE_ZC_COPIED = (1 << 31)
-IORING_ACCEPT_MULTISHOT = (1 << 0)
-IORING_ACCEPT_DONTWAIT = (1 << 1)
-IORING_ACCEPT_POLL_FIRST = (1 << 2)
-IORING_MSG_RING_CQE_SKIP = (1 << 0)
-IORING_MSG_RING_FLAGS_PASS = (1 << 1)
-IORING_FIXED_FD_NO_CLOEXEC = (1 << 0)
-IORING_NOP_INJECT_RESULT = (1 << 0)
-IORING_NOP_FILE = (1 << 1)
-IORING_NOP_FIXED_FILE = (1 << 2)
-IORING_NOP_FIXED_BUFFER = (1 << 3)
-IORING_NOP_TW = (1 << 4)
-IORING_NOP_CQE32 = (1 << 5)
-IORING_CQE_F_BUFFER = (1 << 0)
-IORING_CQE_F_MORE = (1 << 1)
-IORING_CQE_F_SOCK_NONEMPTY = (1 << 2)
-IORING_CQE_F_NOTIF = (1 << 3)
-IORING_CQE_F_BUF_MORE = (1 << 4)
-IORING_CQE_F_SKIP = (1 << 5)
-IORING_CQE_F_32 = (1 << 15)
-IORING_CQE_BUFFER_SHIFT = 16
-IORING_OFF_SQ_RING = 0
-IORING_OFF_CQ_RING = 0x8000000
-IORING_OFF_SQES = 0x10000000
-IORING_OFF_PBUF_RING = 0x80000000
-IORING_OFF_PBUF_SHIFT = 16
-IORING_OFF_MMAP_MASK = 0xf8000000
-IORING_SQ_NEED_WAKEUP = (1 << 0)
-IORING_SQ_CQ_OVERFLOW = (1 << 1)
-IORING_SQ_TASKRUN = (1 << 2)
-IORING_CQ_EVENTFD_DISABLED = (1 << 0)
-IORING_ENTER_GETEVENTS = (1 << 0)
-IORING_ENTER_SQ_WAKEUP = (1 << 1)
-IORING_ENTER_SQ_WAIT = (1 << 2)
-IORING_ENTER_EXT_ARG = (1 << 3)
-IORING_ENTER_REGISTERED_RING = (1 << 4)
-IORING_ENTER_ABS_TIMER = (1 << 5)
-IORING_ENTER_EXT_ARG_REG = (1 << 6)
-IORING_ENTER_NO_IOWAIT = (1 << 7)
-IORING_FEAT_SINGLE_MMAP = (1 << 0)
-IORING_FEAT_NODROP = (1 << 1)
-IORING_FEAT_SUBMIT_STABLE = (1 << 2)
-IORING_FEAT_RW_CUR_POS = (1 << 3)
-IORING_FEAT_CUR_PERSONALITY = (1 << 4)
-IORING_FEAT_FAST_POLL = (1 << 5)
-IORING_FEAT_POLL_32BITS = (1 << 6)
-IORING_FEAT_SQPOLL_NONFIXED = (1 << 7)
-IORING_FEAT_EXT_ARG = (1 << 8)
-IORING_FEAT_NATIVE_WORKERS = (1 << 9)
-IORING_FEAT_RSRC_TAGS = (1 << 10)
-IORING_FEAT_CQE_SKIP = (1 << 11)
-IORING_FEAT_LINKED_FILE = (1 << 12)
-IORING_FEAT_REG_REG_RING = (1 << 13)
-IORING_FEAT_RECVSEND_BUNDLE = (1 << 14)
-IORING_FEAT_MIN_TIMEOUT = (1 << 15)
-IORING_FEAT_RW_ATTR = (1 << 16)
-IORING_FEAT_NO_IOWAIT = (1 << 17)
-IORING_RSRC_REGISTER_SPARSE = (1 << 0)
-IORING_REGISTER_FILES_SKIP = (-2)
-IO_URING_OP_SUPPORTED = (1 << 0)
-IORING_TIMESTAMP_HW_SHIFT = 16
-IORING_TIMESTAMP_TYPE_SHIFT = (IORING_TIMESTAMP_HW_SHIFT + 1)
-IORING_ZCRX_AREA_SHIFT = 48
+IORING_RW_ATTR_FLAG_PI = (1 << 0) # type: ignore
+IORING_FILE_INDEX_ALLOC = (~0) # type: ignore
+IOSQE_FIXED_FILE = (1 << IOSQE_FIXED_FILE_BIT) # type: ignore
+IOSQE_IO_DRAIN = (1 << IOSQE_IO_DRAIN_BIT) # type: ignore
+IOSQE_IO_LINK = (1 << IOSQE_IO_LINK_BIT) # type: ignore
+IOSQE_IO_HARDLINK = (1 << IOSQE_IO_HARDLINK_BIT) # type: ignore
+IOSQE_ASYNC = (1 << IOSQE_ASYNC_BIT) # type: ignore
+IOSQE_BUFFER_SELECT = (1 << IOSQE_BUFFER_SELECT_BIT) # type: ignore
+IOSQE_CQE_SKIP_SUCCESS = (1 << IOSQE_CQE_SKIP_SUCCESS_BIT) # type: ignore
+IORING_SETUP_IOPOLL = (1 << 0) # type: ignore
+IORING_SETUP_SQPOLL = (1 << 1) # type: ignore
+IORING_SETUP_SQ_AFF = (1 << 2) # type: ignore
+IORING_SETUP_CQSIZE = (1 << 3) # type: ignore
+IORING_SETUP_CLAMP = (1 << 4) # type: ignore
+IORING_SETUP_ATTACH_WQ = (1 << 5) # type: ignore
+IORING_SETUP_R_DISABLED = (1 << 6) # type: ignore
+IORING_SETUP_SUBMIT_ALL = (1 << 7) # type: ignore
+IORING_SETUP_COOP_TASKRUN = (1 << 8) # type: ignore
+IORING_SETUP_TASKRUN_FLAG = (1 << 9) # type: ignore
+IORING_SETUP_SQE128 = (1 << 10) # type: ignore
+IORING_SETUP_CQE32 = (1 << 11) # type: ignore
+IORING_SETUP_SINGLE_ISSUER = (1 << 12) # type: ignore
+IORING_SETUP_DEFER_TASKRUN = (1 << 13) # type: ignore
+IORING_SETUP_NO_MMAP = (1 << 14) # type: ignore
+IORING_SETUP_REGISTERED_FD_ONLY = (1 << 15) # type: ignore
+IORING_SETUP_NO_SQARRAY = (1 << 16) # type: ignore
+IORING_SETUP_HYBRID_IOPOLL = (1 << 17) # type: ignore
+IORING_SETUP_CQE_MIXED = (1 << 18) # type: ignore
+IORING_URING_CMD_FIXED = (1 << 0) # type: ignore
+IORING_URING_CMD_MULTISHOT = (1 << 1) # type: ignore
+IORING_URING_CMD_MASK = (IORING_URING_CMD_FIXED | IORING_URING_CMD_MULTISHOT) # type: ignore
+IORING_FSYNC_DATASYNC = (1 << 0) # type: ignore
+IORING_TIMEOUT_ABS = (1 << 0) # type: ignore
+IORING_TIMEOUT_UPDATE = (1 << 1) # type: ignore
+IORING_TIMEOUT_BOOTTIME = (1 << 2) # type: ignore
+IORING_TIMEOUT_REALTIME = (1 << 3) # type: ignore
+IORING_LINK_TIMEOUT_UPDATE = (1 << 4) # type: ignore
+IORING_TIMEOUT_ETIME_SUCCESS = (1 << 5) # type: ignore
+IORING_TIMEOUT_MULTISHOT = (1 << 6) # type: ignore
+IORING_TIMEOUT_CLOCK_MASK = (IORING_TIMEOUT_BOOTTIME | IORING_TIMEOUT_REALTIME) # type: ignore
+IORING_TIMEOUT_UPDATE_MASK = (IORING_TIMEOUT_UPDATE | IORING_LINK_TIMEOUT_UPDATE) # type: ignore
+SPLICE_F_FD_IN_FIXED = (1 << 31) # type: ignore
+IORING_POLL_ADD_MULTI = (1 << 0) # type: ignore
+IORING_POLL_UPDATE_EVENTS = (1 << 1) # type: ignore
+IORING_POLL_UPDATE_USER_DATA = (1 << 2) # type: ignore
+IORING_POLL_ADD_LEVEL = (1 << 3) # type: ignore
+IORING_ASYNC_CANCEL_ALL = (1 << 0) # type: ignore
+IORING_ASYNC_CANCEL_FD = (1 << 1) # type: ignore
+IORING_ASYNC_CANCEL_ANY = (1 << 2) # type: ignore
+IORING_ASYNC_CANCEL_FD_FIXED = (1 << 3) # type: ignore
+IORING_ASYNC_CANCEL_USERDATA = (1 << 4) # type: ignore
+IORING_ASYNC_CANCEL_OP = (1 << 5) # type: ignore
+IORING_RECVSEND_POLL_FIRST = (1 << 0) # type: ignore
+IORING_RECV_MULTISHOT = (1 << 1) # type: ignore
+IORING_RECVSEND_FIXED_BUF = (1 << 2) # type: ignore
+IORING_SEND_ZC_REPORT_USAGE = (1 << 3) # type: ignore
+IORING_RECVSEND_BUNDLE = (1 << 4) # type: ignore
+IORING_SEND_VECTORIZED = (1 << 5) # type: ignore
+IORING_NOTIF_USAGE_ZC_COPIED = (1 << 31) # type: ignore
+IORING_ACCEPT_MULTISHOT = (1 << 0) # type: ignore
+IORING_ACCEPT_DONTWAIT = (1 << 1) # type: ignore
+IORING_ACCEPT_POLL_FIRST = (1 << 2) # type: ignore
+IORING_MSG_RING_CQE_SKIP = (1 << 0) # type: ignore
+IORING_MSG_RING_FLAGS_PASS = (1 << 1) # type: ignore
+IORING_FIXED_FD_NO_CLOEXEC = (1 << 0) # type: ignore
+IORING_NOP_INJECT_RESULT = (1 << 0) # type: ignore
+IORING_NOP_FILE = (1 << 1) # type: ignore
+IORING_NOP_FIXED_FILE = (1 << 2) # type: ignore
+IORING_NOP_FIXED_BUFFER = (1 << 3) # type: ignore
+IORING_NOP_TW = (1 << 4) # type: ignore
+IORING_NOP_CQE32 = (1 << 5) # type: ignore
+IORING_CQE_F_BUFFER = (1 << 0) # type: ignore
+IORING_CQE_F_MORE = (1 << 1) # type: ignore
+IORING_CQE_F_SOCK_NONEMPTY = (1 << 2) # type: ignore
+IORING_CQE_F_NOTIF = (1 << 3) # type: ignore
+IORING_CQE_F_BUF_MORE = (1 << 4) # type: ignore
+IORING_CQE_F_SKIP = (1 << 5) # type: ignore
+IORING_CQE_F_32 = (1 << 15) # type: ignore
+IORING_CQE_BUFFER_SHIFT = 16 # type: ignore
+IORING_OFF_SQ_RING = 0 # type: ignore
+IORING_OFF_CQ_RING = 0x8000000 # type: ignore
+IORING_OFF_SQES = 0x10000000 # type: ignore
+IORING_OFF_PBUF_RING = 0x80000000 # type: ignore
+IORING_OFF_PBUF_SHIFT = 16 # type: ignore
+IORING_OFF_MMAP_MASK = 0xf8000000 # type: ignore
+IORING_SQ_NEED_WAKEUP = (1 << 0) # type: ignore
+IORING_SQ_CQ_OVERFLOW = (1 << 1) # type: ignore
+IORING_SQ_TASKRUN = (1 << 2) # type: ignore
+IORING_CQ_EVENTFD_DISABLED = (1 << 0) # type: ignore
+IORING_ENTER_GETEVENTS = (1 << 0) # type: ignore
+IORING_ENTER_SQ_WAKEUP = (1 << 1) # type: ignore
+IORING_ENTER_SQ_WAIT = (1 << 2) # type: ignore
+IORING_ENTER_EXT_ARG = (1 << 3) # type: ignore
+IORING_ENTER_REGISTERED_RING = (1 << 4) # type: ignore
+IORING_ENTER_ABS_TIMER = (1 << 5) # type: ignore
+IORING_ENTER_EXT_ARG_REG = (1 << 6) # type: ignore
+IORING_ENTER_NO_IOWAIT = (1 << 7) # type: ignore
+IORING_FEAT_SINGLE_MMAP = (1 << 0) # type: ignore
+IORING_FEAT_NODROP = (1 << 1) # type: ignore
+IORING_FEAT_SUBMIT_STABLE = (1 << 2) # type: ignore
+IORING_FEAT_RW_CUR_POS = (1 << 3) # type: ignore
+IORING_FEAT_CUR_PERSONALITY = (1 << 4) # type: ignore
+IORING_FEAT_FAST_POLL = (1 << 5) # type: ignore
+IORING_FEAT_POLL_32BITS = (1 << 6) # type: ignore
+IORING_FEAT_SQPOLL_NONFIXED = (1 << 7) # type: ignore
+IORING_FEAT_EXT_ARG = (1 << 8) # type: ignore
+IORING_FEAT_NATIVE_WORKERS = (1 << 9) # type: ignore
+IORING_FEAT_RSRC_TAGS = (1 << 10) # type: ignore
+IORING_FEAT_CQE_SKIP = (1 << 11) # type: ignore
+IORING_FEAT_LINKED_FILE = (1 << 12) # type: ignore
+IORING_FEAT_REG_REG_RING = (1 << 13) # type: ignore
+IORING_FEAT_RECVSEND_BUNDLE = (1 << 14) # type: ignore
+IORING_FEAT_MIN_TIMEOUT = (1 << 15) # type: ignore
+IORING_FEAT_RW_ATTR = (1 << 16) # type: ignore
+IORING_FEAT_NO_IOWAIT = (1 << 17) # type: ignore
+IORING_RSRC_REGISTER_SPARSE = (1 << 0) # type: ignore
+IORING_REGISTER_FILES_SKIP = (-2) # type: ignore
+IO_URING_OP_SUPPORTED = (1 << 0) # type: ignore
+IORING_TIMESTAMP_HW_SHIFT = 16 # type: ignore
+IORING_TIMESTAMP_TYPE_SHIFT = (IORING_TIMESTAMP_HW_SHIFT + 1) # type: ignore
+IORING_ZCRX_AREA_SHIFT = 48 # type: ignore
 __SC_3264 = lambda _nr,_32,_64: __SYSCALL(_nr, _64) # type: ignore
 __SC_COMP = lambda _nr,_sys,_comp: __SYSCALL(_nr, _sys) # type: ignore
 __SC_COMP_3264 = lambda _nr,_32,_64,_comp: __SC_3264(_nr, _32, _64) # type: ignore
-NR_io_setup = 0
-NR_io_destroy = 1
-NR_io_submit = 2
-NR_io_cancel = 3
-NR_io_getevents = 4
-NR_setxattr = 5
-NR_lsetxattr = 6
-NR_fsetxattr = 7
-NR_getxattr = 8
-NR_lgetxattr = 9
-NR_fgetxattr = 10
-NR_listxattr = 11
-NR_llistxattr = 12
-NR_flistxattr = 13
-NR_removexattr = 14
-NR_lremovexattr = 15
-NR_fremovexattr = 16
-NR_getcwd = 17
-NR_lookup_dcookie = 18
-NR_eventfd2 = 19
-NR_epoll_create1 = 20
-NR_epoll_ctl = 21
-NR_epoll_pwait = 22
-NR_dup = 23
-NR_dup3 = 24
-NR3264_fcntl = 25
-NR_inotify_init1 = 26
-NR_inotify_add_watch = 27
-NR_inotify_rm_watch = 28
-NR_ioctl = 29
-NR_ioprio_set = 30
-NR_ioprio_get = 31
-NR_flock = 32
-NR_mknodat = 33
-NR_mkdirat = 34
-NR_unlinkat = 35
-NR_symlinkat = 36
-NR_linkat = 37
-NR_umount2 = 39
-NR_mount = 40
-NR_pivot_root = 41
-NR_nfsservctl = 42
-NR3264_statfs = 43
-NR3264_fstatfs = 44
-NR3264_truncate = 45
-NR3264_ftruncate = 46
-NR_fallocate = 47
-NR_faccessat = 48
-NR_chdir = 49
-NR_fchdir = 50
-NR_chroot = 51
-NR_fchmod = 52
-NR_fchmodat = 53
-NR_fchownat = 54
-NR_fchown = 55
-NR_openat = 56
-NR_close = 57
-NR_vhangup = 58
-NR_pipe2 = 59
-NR_quotactl = 60
-NR_getdents64 = 61
-NR3264_lseek = 62
-NR_read = 63
-NR_write = 64
-NR_readv = 65
-NR_writev = 66
-NR_pread64 = 67
-NR_pwrite64 = 68
-NR_preadv = 69
-NR_pwritev = 70
-NR3264_sendfile = 71
-NR_pselect6 = 72
-NR_ppoll = 73
-NR_signalfd4 = 74
-NR_vmsplice = 75
-NR_splice = 76
-NR_tee = 77
-NR_readlinkat = 78
-NR_sync = 81
-NR_fsync = 82
-NR_fdatasync = 83
-NR_sync_file_range = 84
-NR_timerfd_create = 85
-NR_timerfd_settime = 86
-NR_timerfd_gettime = 87
-NR_utimensat = 88
-NR_acct = 89
-NR_capget = 90
-NR_capset = 91
-NR_personality = 92
-NR_exit = 93
-NR_exit_group = 94
-NR_waitid = 95
-NR_set_tid_address = 96
-NR_unshare = 97
-NR_futex = 98
-NR_set_robust_list = 99
-NR_get_robust_list = 100
-NR_nanosleep = 101
-NR_getitimer = 102
-NR_setitimer = 103
-NR_kexec_load = 104
-NR_init_module = 105
-NR_delete_module = 106
-NR_timer_create = 107
-NR_timer_gettime = 108
-NR_timer_getoverrun = 109
-NR_timer_settime = 110
-NR_timer_delete = 111
-NR_clock_settime = 112
-NR_clock_gettime = 113
-NR_clock_getres = 114
-NR_clock_nanosleep = 115
-NR_syslog = 116
-NR_ptrace = 117
-NR_sched_setparam = 118
-NR_sched_setscheduler = 119
-NR_sched_getscheduler = 120
-NR_sched_getparam = 121
-NR_sched_setaffinity = 122
-NR_sched_getaffinity = 123
-NR_sched_yield = 124
-NR_sched_get_priority_max = 125
-NR_sched_get_priority_min = 126
-NR_sched_rr_get_interval = 127
-NR_restart_syscall = 128
-NR_kill = 129
-NR_tkill = 130
-NR_tgkill = 131
-NR_sigaltstack = 132
-NR_rt_sigsuspend = 133
-NR_rt_sigaction = 134
-NR_rt_sigprocmask = 135
-NR_rt_sigpending = 136
-NR_rt_sigtimedwait = 137
-NR_rt_sigqueueinfo = 138
-NR_rt_sigreturn = 139
-NR_setpriority = 140
-NR_getpriority = 141
-NR_reboot = 142
-NR_setregid = 143
-NR_setgid = 144
-NR_setreuid = 145
-NR_setuid = 146
-NR_setresuid = 147
-NR_getresuid = 148
-NR_setresgid = 149
-NR_getresgid = 150
-NR_setfsuid = 151
-NR_setfsgid = 152
-NR_times = 153
-NR_setpgid = 154
-NR_getpgid = 155
-NR_getsid = 156
-NR_setsid = 157
-NR_getgroups = 158
-NR_setgroups = 159
-NR_uname = 160
-NR_sethostname = 161
-NR_setdomainname = 162
-NR_getrusage = 165
-NR_umask = 166
-NR_prctl = 167
-NR_getcpu = 168
-NR_gettimeofday = 169
-NR_settimeofday = 170
-NR_adjtimex = 171
-NR_getpid = 172
-NR_getppid = 173
-NR_getuid = 174
-NR_geteuid = 175
-NR_getgid = 176
-NR_getegid = 177
-NR_gettid = 178
-NR_sysinfo = 179
-NR_mq_open = 180
-NR_mq_unlink = 181
-NR_mq_timedsend = 182
-NR_mq_timedreceive = 183
-NR_mq_notify = 184
-NR_mq_getsetattr = 185
-NR_msgget = 186
-NR_msgctl = 187
-NR_msgrcv = 188
-NR_msgsnd = 189
-NR_semget = 190
-NR_semctl = 191
-NR_semtimedop = 192
-NR_semop = 193
-NR_shmget = 194
-NR_shmctl = 195
-NR_shmat = 196
-NR_shmdt = 197
-NR_socket = 198
-NR_socketpair = 199
-NR_bind = 200
-NR_listen = 201
-NR_accept = 202
-NR_connect = 203
-NR_getsockname = 204
-NR_getpeername = 205
-NR_sendto = 206
-NR_recvfrom = 207
-NR_setsockopt = 208
-NR_getsockopt = 209
-NR_shutdown = 210
-NR_sendmsg = 211
-NR_recvmsg = 212
-NR_readahead = 213
-NR_brk = 214
-NR_munmap = 215
-NR_mremap = 216
-NR_add_key = 217
-NR_request_key = 218
-NR_keyctl = 219
-NR_clone = 220
-NR_execve = 221
-NR3264_mmap = 222
-NR3264_fadvise64 = 223
-NR_swapon = 224
-NR_swapoff = 225
-NR_mprotect = 226
-NR_msync = 227
-NR_mlock = 228
-NR_munlock = 229
-NR_mlockall = 230
-NR_munlockall = 231
-NR_mincore = 232
-NR_madvise = 233
-NR_remap_file_pages = 234
-NR_mbind = 235
-NR_get_mempolicy = 236
-NR_set_mempolicy = 237
-NR_migrate_pages = 238
-NR_move_pages = 239
-NR_rt_tgsigqueueinfo = 240
-NR_perf_event_open = 241
-NR_accept4 = 242
-NR_recvmmsg = 243
-NR_arch_specific_syscall = 244
-NR_wait4 = 260
-NR_prlimit64 = 261
-NR_fanotify_init = 262
-NR_fanotify_mark = 263
-NR_name_to_handle_at = 264
-NR_open_by_handle_at = 265
-NR_clock_adjtime = 266
-NR_syncfs = 267
-NR_setns = 268
-NR_sendmmsg = 269
-NR_process_vm_readv = 270
-NR_process_vm_writev = 271
-NR_kcmp = 272
-NR_finit_module = 273
-NR_sched_setattr = 274
-NR_sched_getattr = 275
-NR_renameat2 = 276
-NR_seccomp = 277
-NR_getrandom = 278
-NR_memfd_create = 279
-NR_bpf = 280
-NR_execveat = 281
-NR_userfaultfd = 282
-NR_membarrier = 283
-NR_mlock2 = 284
-NR_copy_file_range = 285
-NR_preadv2 = 286
-NR_pwritev2 = 287
-NR_pkey_mprotect = 288
-NR_pkey_alloc = 289
-NR_pkey_free = 290
-NR_statx = 291
-NR_io_pgetevents = 292
-NR_rseq = 293
-NR_kexec_file_load = 294
-NR_pidfd_send_signal = 424
-NR_io_uring_setup = 425
-NR_io_uring_enter = 426
-NR_io_uring_register = 427
-NR_open_tree = 428
-NR_move_mount = 429
-NR_fsopen = 430
-NR_fsconfig = 431
-NR_fsmount = 432
-NR_fspick = 433
-NR_pidfd_open = 434
-NR_clone3 = 435
-NR_close_range = 436
-NR_openat2 = 437
-NR_pidfd_getfd = 438
-NR_faccessat2 = 439
-NR_process_madvise = 440
-NR_epoll_pwait2 = 441
-NR_mount_setattr = 442
-NR_quotactl_fd = 443
-NR_landlock_create_ruleset = 444
-NR_landlock_add_rule = 445
-NR_landlock_restrict_self = 446
-NR_process_mrelease = 448
-NR_futex_waitv = 449
-NR_set_mempolicy_home_node = 450
-NR_cachestat = 451
-NR_fchmodat2 = 452
-NR_map_shadow_stack = 453
-NR_futex_wake = 454
-NR_futex_wait = 455
-NR_futex_requeue = 456
-NR_statmount = 457
-NR_listmount = 458
-NR_lsm_get_self_attr = 459
-NR_lsm_set_self_attr = 460
-NR_lsm_list_modules = 461
-NR_mseal = 462
-NR_setxattrat = 463
-NR_getxattrat = 464
-NR_listxattrat = 465
-NR_removexattrat = 466
-NR_open_tree_attr = 467
-NR_file_getattr = 468
-NR_file_setattr = 469
-NR_syscalls = 470
-NR_fcntl = NR3264_fcntl
-NR_statfs = NR3264_statfs
-NR_fstatfs = NR3264_fstatfs
-NR_truncate = NR3264_truncate
-NR_ftruncate = NR3264_ftruncate
-NR_lseek = NR3264_lseek
-NR_sendfile = NR3264_sendfile
-NR_mmap = NR3264_mmap
-NR_fadvise64 = NR3264_fadvise64
+NR_io_setup = 0 # type: ignore
+NR_io_destroy = 1 # type: ignore
+NR_io_submit = 2 # type: ignore
+NR_io_cancel = 3 # type: ignore
+NR_io_getevents = 4 # type: ignore
+NR_setxattr = 5 # type: ignore
+NR_lsetxattr = 6 # type: ignore
+NR_fsetxattr = 7 # type: ignore
+NR_getxattr = 8 # type: ignore
+NR_lgetxattr = 9 # type: ignore
+NR_fgetxattr = 10 # type: ignore
+NR_listxattr = 11 # type: ignore
+NR_llistxattr = 12 # type: ignore
+NR_flistxattr = 13 # type: ignore
+NR_removexattr = 14 # type: ignore
+NR_lremovexattr = 15 # type: ignore
+NR_fremovexattr = 16 # type: ignore
+NR_getcwd = 17 # type: ignore
+NR_lookup_dcookie = 18 # type: ignore
+NR_eventfd2 = 19 # type: ignore
+NR_epoll_create1 = 20 # type: ignore
+NR_epoll_ctl = 21 # type: ignore
+NR_epoll_pwait = 22 # type: ignore
+NR_dup = 23 # type: ignore
+NR_dup3 = 24 # type: ignore
+NR3264_fcntl = 25 # type: ignore
+NR_inotify_init1 = 26 # type: ignore
+NR_inotify_add_watch = 27 # type: ignore
+NR_inotify_rm_watch = 28 # type: ignore
+NR_ioctl = 29 # type: ignore
+NR_ioprio_set = 30 # type: ignore
+NR_ioprio_get = 31 # type: ignore
+NR_flock = 32 # type: ignore
+NR_mknodat = 33 # type: ignore
+NR_mkdirat = 34 # type: ignore
+NR_unlinkat = 35 # type: ignore
+NR_symlinkat = 36 # type: ignore
+NR_linkat = 37 # type: ignore
+NR_umount2 = 39 # type: ignore
+NR_mount = 40 # type: ignore
+NR_pivot_root = 41 # type: ignore
+NR_nfsservctl = 42 # type: ignore
+NR3264_statfs = 43 # type: ignore
+NR3264_fstatfs = 44 # type: ignore
+NR3264_truncate = 45 # type: ignore
+NR3264_ftruncate = 46 # type: ignore
+NR_fallocate = 47 # type: ignore
+NR_faccessat = 48 # type: ignore
+NR_chdir = 49 # type: ignore
+NR_fchdir = 50 # type: ignore
+NR_chroot = 51 # type: ignore
+NR_fchmod = 52 # type: ignore
+NR_fchmodat = 53 # type: ignore
+NR_fchownat = 54 # type: ignore
+NR_fchown = 55 # type: ignore
+NR_openat = 56 # type: ignore
+NR_close = 57 # type: ignore
+NR_vhangup = 58 # type: ignore
+NR_pipe2 = 59 # type: ignore
+NR_quotactl = 60 # type: ignore
+NR_getdents64 = 61 # type: ignore
+NR3264_lseek = 62 # type: ignore
+NR_read = 63 # type: ignore
+NR_write = 64 # type: ignore
+NR_readv = 65 # type: ignore
+NR_writev = 66 # type: ignore
+NR_pread64 = 67 # type: ignore
+NR_pwrite64 = 68 # type: ignore
+NR_preadv = 69 # type: ignore
+NR_pwritev = 70 # type: ignore
+NR3264_sendfile = 71 # type: ignore
+NR_pselect6 = 72 # type: ignore
+NR_ppoll = 73 # type: ignore
+NR_signalfd4 = 74 # type: ignore
+NR_vmsplice = 75 # type: ignore
+NR_splice = 76 # type: ignore
+NR_tee = 77 # type: ignore
+NR_readlinkat = 78 # type: ignore
+NR_sync = 81 # type: ignore
+NR_fsync = 82 # type: ignore
+NR_fdatasync = 83 # type: ignore
+NR_sync_file_range = 84 # type: ignore
+NR_timerfd_create = 85 # type: ignore
+NR_timerfd_settime = 86 # type: ignore
+NR_timerfd_gettime = 87 # type: ignore
+NR_utimensat = 88 # type: ignore
+NR_acct = 89 # type: ignore
+NR_capget = 90 # type: ignore
+NR_capset = 91 # type: ignore
+NR_personality = 92 # type: ignore
+NR_exit = 93 # type: ignore
+NR_exit_group = 94 # type: ignore
+NR_waitid = 95 # type: ignore
+NR_set_tid_address = 96 # type: ignore
+NR_unshare = 97 # type: ignore
+NR_futex = 98 # type: ignore
+NR_set_robust_list = 99 # type: ignore
+NR_get_robust_list = 100 # type: ignore
+NR_nanosleep = 101 # type: ignore
+NR_getitimer = 102 # type: ignore
+NR_setitimer = 103 # type: ignore
+NR_kexec_load = 104 # type: ignore
+NR_init_module = 105 # type: ignore
+NR_delete_module = 106 # type: ignore
+NR_timer_create = 107 # type: ignore
+NR_timer_gettime = 108 # type: ignore
+NR_timer_getoverrun = 109 # type: ignore
+NR_timer_settime = 110 # type: ignore
+NR_timer_delete = 111 # type: ignore
+NR_clock_settime = 112 # type: ignore
+NR_clock_gettime = 113 # type: ignore
+NR_clock_getres = 114 # type: ignore
+NR_clock_nanosleep = 115 # type: ignore
+NR_syslog = 116 # type: ignore
+NR_ptrace = 117 # type: ignore
+NR_sched_setparam = 118 # type: ignore
+NR_sched_setscheduler = 119 # type: ignore
+NR_sched_getscheduler = 120 # type: ignore
+NR_sched_getparam = 121 # type: ignore
+NR_sched_setaffinity = 122 # type: ignore
+NR_sched_getaffinity = 123 # type: ignore
+NR_sched_yield = 124 # type: ignore
+NR_sched_get_priority_max = 125 # type: ignore
+NR_sched_get_priority_min = 126 # type: ignore
+NR_sched_rr_get_interval = 127 # type: ignore
+NR_restart_syscall = 128 # type: ignore
+NR_kill = 129 # type: ignore
+NR_tkill = 130 # type: ignore
+NR_tgkill = 131 # type: ignore
+NR_sigaltstack = 132 # type: ignore
+NR_rt_sigsuspend = 133 # type: ignore
+NR_rt_sigaction = 134 # type: ignore
+NR_rt_sigprocmask = 135 # type: ignore
+NR_rt_sigpending = 136 # type: ignore
+NR_rt_sigtimedwait = 137 # type: ignore
+NR_rt_sigqueueinfo = 138 # type: ignore
+NR_rt_sigreturn = 139 # type: ignore
+NR_setpriority = 140 # type: ignore
+NR_getpriority = 141 # type: ignore
+NR_reboot = 142 # type: ignore
+NR_setregid = 143 # type: ignore
+NR_setgid = 144 # type: ignore
+NR_setreuid = 145 # type: ignore
+NR_setuid = 146 # type: ignore
+NR_setresuid = 147 # type: ignore
+NR_getresuid = 148 # type: ignore
+NR_setresgid = 149 # type: ignore
+NR_getresgid = 150 # type: ignore
+NR_setfsuid = 151 # type: ignore
+NR_setfsgid = 152 # type: ignore
+NR_times = 153 # type: ignore
+NR_setpgid = 154 # type: ignore
+NR_getpgid = 155 # type: ignore
+NR_getsid = 156 # type: ignore
+NR_setsid = 157 # type: ignore
+NR_getgroups = 158 # type: ignore
+NR_setgroups = 159 # type: ignore
+NR_uname = 160 # type: ignore
+NR_sethostname = 161 # type: ignore
+NR_setdomainname = 162 # type: ignore
+NR_getrusage = 165 # type: ignore
+NR_umask = 166 # type: ignore
+NR_prctl = 167 # type: ignore
+NR_getcpu = 168 # type: ignore
+NR_gettimeofday = 169 # type: ignore
+NR_settimeofday = 170 # type: ignore
+NR_adjtimex = 171 # type: ignore
+NR_getpid = 172 # type: ignore
+NR_getppid = 173 # type: ignore
+NR_getuid = 174 # type: ignore
+NR_geteuid = 175 # type: ignore
+NR_getgid = 176 # type: ignore
+NR_getegid = 177 # type: ignore
+NR_gettid = 178 # type: ignore
+NR_sysinfo = 179 # type: ignore
+NR_mq_open = 180 # type: ignore
+NR_mq_unlink = 181 # type: ignore
+NR_mq_timedsend = 182 # type: ignore
+NR_mq_timedreceive = 183 # type: ignore
+NR_mq_notify = 184 # type: ignore
+NR_mq_getsetattr = 185 # type: ignore
+NR_msgget = 186 # type: ignore
+NR_msgctl = 187 # type: ignore
+NR_msgrcv = 188 # type: ignore
+NR_msgsnd = 189 # type: ignore
+NR_semget = 190 # type: ignore
+NR_semctl = 191 # type: ignore
+NR_semtimedop = 192 # type: ignore
+NR_semop = 193 # type: ignore
+NR_shmget = 194 # type: ignore
+NR_shmctl = 195 # type: ignore
+NR_shmat = 196 # type: ignore
+NR_shmdt = 197 # type: ignore
+NR_socket = 198 # type: ignore
+NR_socketpair = 199 # type: ignore
+NR_bind = 200 # type: ignore
+NR_listen = 201 # type: ignore
+NR_accept = 202 # type: ignore
+NR_connect = 203 # type: ignore
+NR_getsockname = 204 # type: ignore
+NR_getpeername = 205 # type: ignore
+NR_sendto = 206 # type: ignore
+NR_recvfrom = 207 # type: ignore
+NR_setsockopt = 208 # type: ignore
+NR_getsockopt = 209 # type: ignore
+NR_shutdown = 210 # type: ignore
+NR_sendmsg = 211 # type: ignore
+NR_recvmsg = 212 # type: ignore
+NR_readahead = 213 # type: ignore
+NR_brk = 214 # type: ignore
+NR_munmap = 215 # type: ignore
+NR_mremap = 216 # type: ignore
+NR_add_key = 217 # type: ignore
+NR_request_key = 218 # type: ignore
+NR_keyctl = 219 # type: ignore
+NR_clone = 220 # type: ignore
+NR_execve = 221 # type: ignore
+NR3264_mmap = 222 # type: ignore
+NR3264_fadvise64 = 223 # type: ignore
+NR_swapon = 224 # type: ignore
+NR_swapoff = 225 # type: ignore
+NR_mprotect = 226 # type: ignore
+NR_msync = 227 # type: ignore
+NR_mlock = 228 # type: ignore
+NR_munlock = 229 # type: ignore
+NR_mlockall = 230 # type: ignore
+NR_munlockall = 231 # type: ignore
+NR_mincore = 232 # type: ignore
+NR_madvise = 233 # type: ignore
+NR_remap_file_pages = 234 # type: ignore
+NR_mbind = 235 # type: ignore
+NR_get_mempolicy = 236 # type: ignore
+NR_set_mempolicy = 237 # type: ignore
+NR_migrate_pages = 238 # type: ignore
+NR_move_pages = 239 # type: ignore
+NR_rt_tgsigqueueinfo = 240 # type: ignore
+NR_perf_event_open = 241 # type: ignore
+NR_accept4 = 242 # type: ignore
+NR_recvmmsg = 243 # type: ignore
+NR_arch_specific_syscall = 244 # type: ignore
+NR_wait4 = 260 # type: ignore
+NR_prlimit64 = 261 # type: ignore
+NR_fanotify_init = 262 # type: ignore
+NR_fanotify_mark = 263 # type: ignore
+NR_name_to_handle_at = 264 # type: ignore
+NR_open_by_handle_at = 265 # type: ignore
+NR_clock_adjtime = 266 # type: ignore
+NR_syncfs = 267 # type: ignore
+NR_setns = 268 # type: ignore
+NR_sendmmsg = 269 # type: ignore
+NR_process_vm_readv = 270 # type: ignore
+NR_process_vm_writev = 271 # type: ignore
+NR_kcmp = 272 # type: ignore
+NR_finit_module = 273 # type: ignore
+NR_sched_setattr = 274 # type: ignore
+NR_sched_getattr = 275 # type: ignore
+NR_renameat2 = 276 # type: ignore
+NR_seccomp = 277 # type: ignore
+NR_getrandom = 278 # type: ignore
+NR_memfd_create = 279 # type: ignore
+NR_bpf = 280 # type: ignore
+NR_execveat = 281 # type: ignore
+NR_userfaultfd = 282 # type: ignore
+NR_membarrier = 283 # type: ignore
+NR_mlock2 = 284 # type: ignore
+NR_copy_file_range = 285 # type: ignore
+NR_preadv2 = 286 # type: ignore
+NR_pwritev2 = 287 # type: ignore
+NR_pkey_mprotect = 288 # type: ignore
+NR_pkey_alloc = 289 # type: ignore
+NR_pkey_free = 290 # type: ignore
+NR_statx = 291 # type: ignore
+NR_io_pgetevents = 292 # type: ignore
+NR_rseq = 293 # type: ignore
+NR_kexec_file_load = 294 # type: ignore
+NR_pidfd_send_signal = 424 # type: ignore
+NR_io_uring_setup = 425 # type: ignore
+NR_io_uring_enter = 426 # type: ignore
+NR_io_uring_register = 427 # type: ignore
+NR_open_tree = 428 # type: ignore
+NR_move_mount = 429 # type: ignore
+NR_fsopen = 430 # type: ignore
+NR_fsconfig = 431 # type: ignore
+NR_fsmount = 432 # type: ignore
+NR_fspick = 433 # type: ignore
+NR_pidfd_open = 434 # type: ignore
+NR_clone3 = 435 # type: ignore
+NR_close_range = 436 # type: ignore
+NR_openat2 = 437 # type: ignore
+NR_pidfd_getfd = 438 # type: ignore
+NR_faccessat2 = 439 # type: ignore
+NR_process_madvise = 440 # type: ignore
+NR_epoll_pwait2 = 441 # type: ignore
+NR_mount_setattr = 442 # type: ignore
+NR_quotactl_fd = 443 # type: ignore
+NR_landlock_create_ruleset = 444 # type: ignore
+NR_landlock_add_rule = 445 # type: ignore
+NR_landlock_restrict_self = 446 # type: ignore
+NR_process_mrelease = 448 # type: ignore
+NR_futex_waitv = 449 # type: ignore
+NR_set_mempolicy_home_node = 450 # type: ignore
+NR_cachestat = 451 # type: ignore
+NR_fchmodat2 = 452 # type: ignore
+NR_map_shadow_stack = 453 # type: ignore
+NR_futex_wake = 454 # type: ignore
+NR_futex_wait = 455 # type: ignore
+NR_futex_requeue = 456 # type: ignore
+NR_statmount = 457 # type: ignore
+NR_listmount = 458 # type: ignore
+NR_lsm_get_self_attr = 459 # type: ignore
+NR_lsm_set_self_attr = 460 # type: ignore
+NR_lsm_list_modules = 461 # type: ignore
+NR_mseal = 462 # type: ignore
+NR_setxattrat = 463 # type: ignore
+NR_getxattrat = 464 # type: ignore
+NR_listxattrat = 465 # type: ignore
+NR_removexattrat = 466 # type: ignore
+NR_open_tree_attr = 467 # type: ignore
+NR_file_getattr = 468 # type: ignore
+NR_file_setattr = 469 # type: ignore
+NR_syscalls = 470 # type: ignore
+NR_fcntl = NR3264_fcntl # type: ignore
+NR_statfs = NR3264_statfs # type: ignore
+NR_fstatfs = NR3264_fstatfs # type: ignore
+NR_truncate = NR3264_truncate # type: ignore
+NR_ftruncate = NR3264_ftruncate # type: ignore
+NR_lseek = NR3264_lseek # type: ignore
+NR_sendfile = NR3264_sendfile # type: ignore
+NR_mmap = NR3264_mmap # type: ignore
+NR_fadvise64 = NR3264_fadvise64 # type: ignore

tinygrad/runtime/autogen/kfd.py

@@ -683,137 +683,137 @@ class struct_kfd_ioctl_profiler_args(c.Struct):
   pmc: struct_kfd_ioctl_pmc_settings
   version: int
 struct_kfd_ioctl_profiler_args.register_fields([('op', ctypes.c_uint32, 0), ('pc_sample', struct_kfd_ioctl_pc_sample_args, 8), ('pmc', struct_kfd_ioctl_pmc_settings, 8), ('version', ctypes.c_uint32, 8)])
-KFD_IOCTL_MAJOR_VERSION = 1
-KFD_IOCTL_MINOR_VERSION = 17
-KFD_IOC_QUEUE_TYPE_COMPUTE = 0x0
-KFD_IOC_QUEUE_TYPE_SDMA = 0x1
-KFD_IOC_QUEUE_TYPE_COMPUTE_AQL = 0x2
-KFD_IOC_QUEUE_TYPE_SDMA_XGMI = 0x3
-KFD_IOC_QUEUE_TYPE_SDMA_BY_ENG_ID = 0x4
-KFD_MAX_QUEUE_PERCENTAGE = 100
-KFD_MAX_QUEUE_PRIORITY = 15
-KFD_IOC_CACHE_POLICY_COHERENT = 0
-KFD_IOC_CACHE_POLICY_NONCOHERENT = 1
-NUM_OF_SUPPORTED_GPUS = 7
-MAX_ALLOWED_NUM_POINTS = 100
-MAX_ALLOWED_AW_BUFF_SIZE = 4096
-MAX_ALLOWED_WAC_BUFF_SIZE = 128
-KFD_INVALID_FD = 0xffffffff
-KFD_IOC_EVENT_SIGNAL = 0
-KFD_IOC_EVENT_NODECHANGE = 1
-KFD_IOC_EVENT_DEVICESTATECHANGE = 2
-KFD_IOC_EVENT_HW_EXCEPTION = 3
-KFD_IOC_EVENT_SYSTEM_EVENT = 4
-KFD_IOC_EVENT_DEBUG_EVENT = 5
-KFD_IOC_EVENT_PROFILE_EVENT = 6
-KFD_IOC_EVENT_QUEUE_EVENT = 7
-KFD_IOC_EVENT_MEMORY = 8
-KFD_IOC_WAIT_RESULT_COMPLETE = 0
-KFD_IOC_WAIT_RESULT_TIMEOUT = 1
-KFD_IOC_WAIT_RESULT_FAIL = 2
-KFD_SIGNAL_EVENT_LIMIT = 4096
-KFD_HW_EXCEPTION_WHOLE_GPU_RESET = 0
-KFD_HW_EXCEPTION_PER_ENGINE_RESET = 1
-KFD_HW_EXCEPTION_GPU_HANG = 0
-KFD_HW_EXCEPTION_ECC = 1
-KFD_MEM_ERR_NO_RAS = 0
-KFD_MEM_ERR_SRAM_ECC = 1
-KFD_MEM_ERR_POISON_CONSUMED = 2
-KFD_MEM_ERR_GPU_HANG = 3
-KFD_IOC_ALLOC_MEM_FLAGS_VRAM = (1 << 0)
-KFD_IOC_ALLOC_MEM_FLAGS_GTT = (1 << 1)
-KFD_IOC_ALLOC_MEM_FLAGS_USERPTR = (1 << 2)
-KFD_IOC_ALLOC_MEM_FLAGS_DOORBELL = (1 << 3)
-KFD_IOC_ALLOC_MEM_FLAGS_MMIO_REMAP = (1 << 4)
-KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE = (1 << 31)
-KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE = (1 << 30)
-KFD_IOC_ALLOC_MEM_FLAGS_PUBLIC = (1 << 29)
-KFD_IOC_ALLOC_MEM_FLAGS_NO_SUBSTITUTE = (1 << 28)
-KFD_IOC_ALLOC_MEM_FLAGS_AQL_QUEUE_MEM = (1 << 27)
-KFD_IOC_ALLOC_MEM_FLAGS_COHERENT = (1 << 26)
-KFD_IOC_ALLOC_MEM_FLAGS_UNCACHED = (1 << 25)
-KFD_IOC_ALLOC_MEM_FLAGS_EXT_COHERENT = (1 << 24)
-KFD_IOC_ALLOC_MEM_FLAGS_CONTIGUOUS = (1 << 23)
+KFD_IOCTL_MAJOR_VERSION = 1 # type: ignore
+KFD_IOCTL_MINOR_VERSION = 17 # type: ignore
+KFD_IOC_QUEUE_TYPE_COMPUTE = 0x0 # type: ignore
+KFD_IOC_QUEUE_TYPE_SDMA = 0x1 # type: ignore
+KFD_IOC_QUEUE_TYPE_COMPUTE_AQL = 0x2 # type: ignore
+KFD_IOC_QUEUE_TYPE_SDMA_XGMI = 0x3 # type: ignore
+KFD_IOC_QUEUE_TYPE_SDMA_BY_ENG_ID = 0x4 # type: ignore
+KFD_MAX_QUEUE_PERCENTAGE = 100 # type: ignore
+KFD_MAX_QUEUE_PRIORITY = 15 # type: ignore
+KFD_IOC_CACHE_POLICY_COHERENT = 0 # type: ignore
+KFD_IOC_CACHE_POLICY_NONCOHERENT = 1 # type: ignore
+NUM_OF_SUPPORTED_GPUS = 7 # type: ignore
+MAX_ALLOWED_NUM_POINTS = 100 # type: ignore
+MAX_ALLOWED_AW_BUFF_SIZE = 4096 # type: ignore
+MAX_ALLOWED_WAC_BUFF_SIZE = 128 # type: ignore
+KFD_INVALID_FD = 0xffffffff # type: ignore
+KFD_IOC_EVENT_SIGNAL = 0 # type: ignore
+KFD_IOC_EVENT_NODECHANGE = 1 # type: ignore
+KFD_IOC_EVENT_DEVICESTATECHANGE = 2 # type: ignore
+KFD_IOC_EVENT_HW_EXCEPTION = 3 # type: ignore
+KFD_IOC_EVENT_SYSTEM_EVENT = 4 # type: ignore
+KFD_IOC_EVENT_DEBUG_EVENT = 5 # type: ignore
+KFD_IOC_EVENT_PROFILE_EVENT = 6 # type: ignore
+KFD_IOC_EVENT_QUEUE_EVENT = 7 # type: ignore
+KFD_IOC_EVENT_MEMORY = 8 # type: ignore
+KFD_IOC_WAIT_RESULT_COMPLETE = 0 # type: ignore
+KFD_IOC_WAIT_RESULT_TIMEOUT = 1 # type: ignore
+KFD_IOC_WAIT_RESULT_FAIL = 2 # type: ignore
+KFD_SIGNAL_EVENT_LIMIT = 4096 # type: ignore
+KFD_HW_EXCEPTION_WHOLE_GPU_RESET = 0 # type: ignore
+KFD_HW_EXCEPTION_PER_ENGINE_RESET = 1 # type: ignore
+KFD_HW_EXCEPTION_GPU_HANG = 0 # type: ignore
+KFD_HW_EXCEPTION_ECC = 1 # type: ignore
+KFD_MEM_ERR_NO_RAS = 0 # type: ignore
+KFD_MEM_ERR_SRAM_ECC = 1 # type: ignore
+KFD_MEM_ERR_POISON_CONSUMED = 2 # type: ignore
+KFD_MEM_ERR_GPU_HANG = 3 # type: ignore
+KFD_IOC_ALLOC_MEM_FLAGS_VRAM = (1 << 0) # type: ignore
+KFD_IOC_ALLOC_MEM_FLAGS_GTT = (1 << 1) # type: ignore
+KFD_IOC_ALLOC_MEM_FLAGS_USERPTR = (1 << 2) # type: ignore
+KFD_IOC_ALLOC_MEM_FLAGS_DOORBELL = (1 << 3) # type: ignore
+KFD_IOC_ALLOC_MEM_FLAGS_MMIO_REMAP = (1 << 4) # type: ignore
+KFD_IOC_ALLOC_MEM_FLAGS_WRITABLE = (1 << 31) # type: ignore
+KFD_IOC_ALLOC_MEM_FLAGS_EXECUTABLE = (1 << 30) # type: ignore
+KFD_IOC_ALLOC_MEM_FLAGS_PUBLIC = (1 << 29) # type: ignore
+KFD_IOC_ALLOC_MEM_FLAGS_NO_SUBSTITUTE = (1 << 28) # type: ignore
+KFD_IOC_ALLOC_MEM_FLAGS_AQL_QUEUE_MEM = (1 << 27) # type: ignore
+KFD_IOC_ALLOC_MEM_FLAGS_COHERENT = (1 << 26) # type: ignore
+KFD_IOC_ALLOC_MEM_FLAGS_UNCACHED = (1 << 25) # type: ignore
+KFD_IOC_ALLOC_MEM_FLAGS_EXT_COHERENT = (1 << 24) # type: ignore
+KFD_IOC_ALLOC_MEM_FLAGS_CONTIGUOUS = (1 << 23) # type: ignore
 KFD_SMI_EVENT_MASK_FROM_INDEX = lambda i: (1 << ((i) - 1)) # type: ignore
-KFD_SMI_EVENT_MSG_SIZE = 96
-KFD_IOCTL_SVM_FLAG_HOST_ACCESS = 0x00000001
-KFD_IOCTL_SVM_FLAG_COHERENT = 0x00000002
-KFD_IOCTL_SVM_FLAG_HIVE_LOCAL = 0x00000004
-KFD_IOCTL_SVM_FLAG_GPU_RO = 0x00000008
-KFD_IOCTL_SVM_FLAG_GPU_EXEC = 0x00000010
-KFD_IOCTL_SVM_FLAG_GPU_READ_MOSTLY = 0x00000020
-KFD_IOCTL_SVM_FLAG_GPU_ALWAYS_MAPPED = 0x00000040
-KFD_IOCTL_SVM_FLAG_EXT_COHERENT = 0x00000080
+KFD_SMI_EVENT_MSG_SIZE = 96 # type: ignore
+KFD_IOCTL_SVM_FLAG_HOST_ACCESS = 0x00000001 # type: ignore
+KFD_IOCTL_SVM_FLAG_COHERENT = 0x00000002 # type: ignore
+KFD_IOCTL_SVM_FLAG_HIVE_LOCAL = 0x00000004 # type: ignore
+KFD_IOCTL_SVM_FLAG_GPU_RO = 0x00000008 # type: ignore
+KFD_IOCTL_SVM_FLAG_GPU_EXEC = 0x00000010 # type: ignore
+KFD_IOCTL_SVM_FLAG_GPU_READ_MOSTLY = 0x00000020 # type: ignore
+KFD_IOCTL_SVM_FLAG_GPU_ALWAYS_MAPPED = 0x00000040 # type: ignore
+KFD_IOCTL_SVM_FLAG_EXT_COHERENT = 0x00000080 # type: ignore
 KFD_EC_MASK = lambda ecode: (1 << (ecode - 1)) # type: ignore
-KFD_EC_MASK_QUEUE = (KFD_EC_MASK(EC_QUEUE_WAVE_ABORT) | KFD_EC_MASK(EC_QUEUE_WAVE_TRAP) | KFD_EC_MASK(EC_QUEUE_WAVE_MATH_ERROR) | KFD_EC_MASK(EC_QUEUE_WAVE_ILLEGAL_INSTRUCTION) | KFD_EC_MASK(EC_QUEUE_WAVE_MEMORY_VIOLATION) | KFD_EC_MASK(EC_QUEUE_WAVE_APERTURE_VIOLATION) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_DIM_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_GROUP_SEGMENT_SIZE_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_CODE_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_RESERVED) | KFD_EC_MASK(EC_QUEUE_PACKET_UNSUPPORTED) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_WORK_GROUP_SIZE_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_REGISTER_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_VENDOR_UNSUPPORTED)	| KFD_EC_MASK(EC_QUEUE_PREEMPTION_ERROR)	| KFD_EC_MASK(EC_QUEUE_NEW))
-KFD_EC_MASK_DEVICE = (KFD_EC_MASK(EC_DEVICE_QUEUE_DELETE) | KFD_EC_MASK(EC_DEVICE_RAS_ERROR) | KFD_EC_MASK(EC_DEVICE_FATAL_HALT) | KFD_EC_MASK(EC_DEVICE_MEMORY_VIOLATION) | KFD_EC_MASK(EC_DEVICE_NEW))
-KFD_EC_MASK_PROCESS = (KFD_EC_MASK(EC_PROCESS_RUNTIME) | KFD_EC_MASK(EC_PROCESS_DEVICE_REMOVE))
-KFD_EC_MASK_PACKET = (KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_DIM_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_GROUP_SEGMENT_SIZE_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_CODE_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_RESERVED) | KFD_EC_MASK(EC_QUEUE_PACKET_UNSUPPORTED) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_WORK_GROUP_SIZE_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_REGISTER_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_VENDOR_UNSUPPORTED))
+KFD_EC_MASK_QUEUE = (KFD_EC_MASK(EC_QUEUE_WAVE_ABORT) | KFD_EC_MASK(EC_QUEUE_WAVE_TRAP) | KFD_EC_MASK(EC_QUEUE_WAVE_MATH_ERROR) | KFD_EC_MASK(EC_QUEUE_WAVE_ILLEGAL_INSTRUCTION) | KFD_EC_MASK(EC_QUEUE_WAVE_MEMORY_VIOLATION) | KFD_EC_MASK(EC_QUEUE_WAVE_APERTURE_VIOLATION) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_DIM_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_GROUP_SEGMENT_SIZE_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_CODE_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_RESERVED) | KFD_EC_MASK(EC_QUEUE_PACKET_UNSUPPORTED) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_WORK_GROUP_SIZE_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_REGISTER_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_VENDOR_UNSUPPORTED)	| KFD_EC_MASK(EC_QUEUE_PREEMPTION_ERROR)	| KFD_EC_MASK(EC_QUEUE_NEW)) # type: ignore
+KFD_EC_MASK_DEVICE = (KFD_EC_MASK(EC_DEVICE_QUEUE_DELETE) | KFD_EC_MASK(EC_DEVICE_RAS_ERROR) | KFD_EC_MASK(EC_DEVICE_FATAL_HALT) | KFD_EC_MASK(EC_DEVICE_MEMORY_VIOLATION) | KFD_EC_MASK(EC_DEVICE_NEW)) # type: ignore
+KFD_EC_MASK_PROCESS = (KFD_EC_MASK(EC_PROCESS_RUNTIME) | KFD_EC_MASK(EC_PROCESS_DEVICE_REMOVE)) # type: ignore
+KFD_EC_MASK_PACKET = (KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_DIM_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_GROUP_SEGMENT_SIZE_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_CODE_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_RESERVED) | KFD_EC_MASK(EC_QUEUE_PACKET_UNSUPPORTED) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_WORK_GROUP_SIZE_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_DISPATCH_REGISTER_INVALID) | KFD_EC_MASK(EC_QUEUE_PACKET_VENDOR_UNSUPPORTED)) # type: ignore
 KFD_DBG_EC_IS_VALID = lambda ecode: (ecode > EC_NONE and ecode < EC_MAX) # type: ignore
 KFD_DBG_EC_TYPE_IS_QUEUE = lambda ecode: (KFD_DBG_EC_IS_VALID(ecode) and not  not (KFD_EC_MASK(ecode) & KFD_EC_MASK_QUEUE)) # type: ignore
 KFD_DBG_EC_TYPE_IS_DEVICE = lambda ecode: (KFD_DBG_EC_IS_VALID(ecode) and not  not (KFD_EC_MASK(ecode) & KFD_EC_MASK_DEVICE)) # type: ignore
 KFD_DBG_EC_TYPE_IS_PROCESS = lambda ecode: (KFD_DBG_EC_IS_VALID(ecode) and not  not (KFD_EC_MASK(ecode) & KFD_EC_MASK_PROCESS)) # type: ignore
 KFD_DBG_EC_TYPE_IS_PACKET = lambda ecode: (KFD_DBG_EC_IS_VALID(ecode) and not  not (KFD_EC_MASK(ecode) & KFD_EC_MASK_PACKET)) # type: ignore
-KFD_RUNTIME_ENABLE_MODE_ENABLE_MASK = 1
-KFD_RUNTIME_ENABLE_MODE_TTMP_SAVE_MASK = 2
-KFD_DBG_QUEUE_ERROR_BIT = 30
-KFD_DBG_QUEUE_INVALID_BIT = 31
-KFD_DBG_QUEUE_ERROR_MASK = (1 << KFD_DBG_QUEUE_ERROR_BIT)
-KFD_DBG_QUEUE_INVALID_MASK = (1 << KFD_DBG_QUEUE_INVALID_BIT)
-KFD_IOCTL_PCS_FLAG_POWER_OF_2 = 0x00000001
-KFD_IOCTL_PCS_QUERY_TYPE_FULL = (1 << 0)
-KFD_IOC_PROFILER_VERSION_NUM = 1
-AMDKFD_IOCTL_BASE = 'K'
+KFD_RUNTIME_ENABLE_MODE_ENABLE_MASK = 1 # type: ignore
+KFD_RUNTIME_ENABLE_MODE_TTMP_SAVE_MASK = 2 # type: ignore
+KFD_DBG_QUEUE_ERROR_BIT = 30 # type: ignore
+KFD_DBG_QUEUE_INVALID_BIT = 31 # type: ignore
+KFD_DBG_QUEUE_ERROR_MASK = (1 << KFD_DBG_QUEUE_ERROR_BIT) # type: ignore
+KFD_DBG_QUEUE_INVALID_MASK = (1 << KFD_DBG_QUEUE_INVALID_BIT) # type: ignore
+KFD_IOCTL_PCS_FLAG_POWER_OF_2 = 0x00000001 # type: ignore
+KFD_IOCTL_PCS_QUERY_TYPE_FULL = (1 << 0) # type: ignore
+KFD_IOC_PROFILER_VERSION_NUM = 1 # type: ignore
+AMDKFD_IOCTL_BASE = 'K' # type: ignore
 AMDKFD_IO = lambda nr: _IO(AMDKFD_IOCTL_BASE, nr) # type: ignore
 AMDKFD_IOR = lambda nr,type: _IOR(AMDKFD_IOCTL_BASE, nr, type) # type: ignore
 AMDKFD_IOW = lambda nr,type: _IOW(AMDKFD_IOCTL_BASE, nr, type) # type: ignore
 AMDKFD_IOWR = lambda nr,type: _IOWR(AMDKFD_IOCTL_BASE, nr, type) # type: ignore
-AMDKFD_IOC_GET_VERSION = AMDKFD_IOR(0x01, struct_kfd_ioctl_get_version_args)
-AMDKFD_IOC_CREATE_QUEUE = AMDKFD_IOWR(0x02, struct_kfd_ioctl_create_queue_args)
-AMDKFD_IOC_DESTROY_QUEUE = AMDKFD_IOWR(0x03, struct_kfd_ioctl_destroy_queue_args)
-AMDKFD_IOC_SET_MEMORY_POLICY = AMDKFD_IOW(0x04, struct_kfd_ioctl_set_memory_policy_args)
-AMDKFD_IOC_GET_CLOCK_COUNTERS = AMDKFD_IOWR(0x05, struct_kfd_ioctl_get_clock_counters_args)
-AMDKFD_IOC_GET_PROCESS_APERTURES = AMDKFD_IOR(0x06, struct_kfd_ioctl_get_process_apertures_args)
-AMDKFD_IOC_UPDATE_QUEUE = AMDKFD_IOW(0x07, struct_kfd_ioctl_update_queue_args)
-AMDKFD_IOC_CREATE_EVENT = AMDKFD_IOWR(0x08, struct_kfd_ioctl_create_event_args)
-AMDKFD_IOC_DESTROY_EVENT = AMDKFD_IOW(0x09, struct_kfd_ioctl_destroy_event_args)
-AMDKFD_IOC_SET_EVENT = AMDKFD_IOW(0x0A, struct_kfd_ioctl_set_event_args)
-AMDKFD_IOC_RESET_EVENT = AMDKFD_IOW(0x0B, struct_kfd_ioctl_reset_event_args)
-AMDKFD_IOC_WAIT_EVENTS = AMDKFD_IOWR(0x0C, struct_kfd_ioctl_wait_events_args)
-AMDKFD_IOC_DBG_REGISTER_DEPRECATED = AMDKFD_IOW(0x0D, struct_kfd_ioctl_dbg_register_args)
-AMDKFD_IOC_DBG_UNREGISTER_DEPRECATED = AMDKFD_IOW(0x0E, struct_kfd_ioctl_dbg_unregister_args)
-AMDKFD_IOC_DBG_ADDRESS_WATCH_DEPRECATED = AMDKFD_IOW(0x0F, struct_kfd_ioctl_dbg_address_watch_args)
-AMDKFD_IOC_DBG_WAVE_CONTROL_DEPRECATED = AMDKFD_IOW(0x10, struct_kfd_ioctl_dbg_wave_control_args)
-AMDKFD_IOC_SET_SCRATCH_BACKING_VA = AMDKFD_IOWR(0x11, struct_kfd_ioctl_set_scratch_backing_va_args)
-AMDKFD_IOC_GET_TILE_CONFIG = AMDKFD_IOWR(0x12, struct_kfd_ioctl_get_tile_config_args)
-AMDKFD_IOC_SET_TRAP_HANDLER = AMDKFD_IOW(0x13, struct_kfd_ioctl_set_trap_handler_args)
-AMDKFD_IOC_GET_PROCESS_APERTURES_NEW = AMDKFD_IOWR(0x14, struct_kfd_ioctl_get_process_apertures_new_args)
-AMDKFD_IOC_ACQUIRE_VM = AMDKFD_IOW(0x15, struct_kfd_ioctl_acquire_vm_args)
-AMDKFD_IOC_ALLOC_MEMORY_OF_GPU = AMDKFD_IOWR(0x16, struct_kfd_ioctl_alloc_memory_of_gpu_args)
-AMDKFD_IOC_FREE_MEMORY_OF_GPU = AMDKFD_IOW(0x17, struct_kfd_ioctl_free_memory_of_gpu_args)
-AMDKFD_IOC_MAP_MEMORY_TO_GPU = AMDKFD_IOWR(0x18, struct_kfd_ioctl_map_memory_to_gpu_args)
-AMDKFD_IOC_UNMAP_MEMORY_FROM_GPU = AMDKFD_IOWR(0x19, struct_kfd_ioctl_unmap_memory_from_gpu_args)
-AMDKFD_IOC_SET_CU_MASK = AMDKFD_IOW(0x1A, struct_kfd_ioctl_set_cu_mask_args)
-AMDKFD_IOC_GET_QUEUE_WAVE_STATE = AMDKFD_IOWR(0x1B, struct_kfd_ioctl_get_queue_wave_state_args)
-AMDKFD_IOC_GET_DMABUF_INFO = AMDKFD_IOWR(0x1C, struct_kfd_ioctl_get_dmabuf_info_args)
-AMDKFD_IOC_IMPORT_DMABUF = AMDKFD_IOWR(0x1D, struct_kfd_ioctl_import_dmabuf_args)
-AMDKFD_IOC_ALLOC_QUEUE_GWS = AMDKFD_IOWR(0x1E, struct_kfd_ioctl_alloc_queue_gws_args)
-AMDKFD_IOC_SMI_EVENTS = AMDKFD_IOWR(0x1F, struct_kfd_ioctl_smi_events_args)
-AMDKFD_IOC_SVM = AMDKFD_IOWR(0x20, struct_kfd_ioctl_svm_args)
-AMDKFD_IOC_SET_XNACK_MODE = AMDKFD_IOWR(0x21, struct_kfd_ioctl_set_xnack_mode_args)
-AMDKFD_IOC_CRIU_OP = AMDKFD_IOWR(0x22, struct_kfd_ioctl_criu_args)
-AMDKFD_IOC_AVAILABLE_MEMORY = AMDKFD_IOWR(0x23, struct_kfd_ioctl_get_available_memory_args)
-AMDKFD_IOC_EXPORT_DMABUF = AMDKFD_IOWR(0x24, struct_kfd_ioctl_export_dmabuf_args)
-AMDKFD_IOC_RUNTIME_ENABLE = AMDKFD_IOWR(0x25, struct_kfd_ioctl_runtime_enable_args)
-AMDKFD_IOC_DBG_TRAP = AMDKFD_IOWR(0x26, struct_kfd_ioctl_dbg_trap_args)
-AMDKFD_COMMAND_START = 0x01
-AMDKFD_COMMAND_END = 0x27
-AMDKFD_IOC_IPC_IMPORT_HANDLE = AMDKFD_IOWR(0x80, struct_kfd_ioctl_ipc_import_handle_args)
-AMDKFD_IOC_IPC_EXPORT_HANDLE = AMDKFD_IOWR(0x81, struct_kfd_ioctl_ipc_export_handle_args)
-AMDKFD_IOC_DBG_TRAP_DEPRECATED = AMDKFD_IOWR(0x82, struct_kfd_ioctl_dbg_trap_args_deprecated)
-AMDKFD_IOC_CROSS_MEMORY_COPY_DEPRECATED = AMDKFD_IOWR(0x83, struct_kfd_ioctl_cross_memory_copy_deprecated_args)
-AMDKFD_IOC_RLC_SPM = AMDKFD_IOWR(0x84, struct_kfd_ioctl_spm_args)
-AMDKFD_IOC_PC_SAMPLE = AMDKFD_IOWR(0x85, struct_kfd_ioctl_pc_sample_args)
-AMDKFD_IOC_PROFILER = AMDKFD_IOWR(0x86, struct_kfd_ioctl_profiler_args)
-AMDKFD_COMMAND_START_2 = 0x80
-AMDKFD_COMMAND_END_2 = 0x87
+AMDKFD_IOC_GET_VERSION = AMDKFD_IOR(0x01, struct_kfd_ioctl_get_version_args) # type: ignore
+AMDKFD_IOC_CREATE_QUEUE = AMDKFD_IOWR(0x02, struct_kfd_ioctl_create_queue_args) # type: ignore
+AMDKFD_IOC_DESTROY_QUEUE = AMDKFD_IOWR(0x03, struct_kfd_ioctl_destroy_queue_args) # type: ignore
+AMDKFD_IOC_SET_MEMORY_POLICY = AMDKFD_IOW(0x04, struct_kfd_ioctl_set_memory_policy_args) # type: ignore
+AMDKFD_IOC_GET_CLOCK_COUNTERS = AMDKFD_IOWR(0x05, struct_kfd_ioctl_get_clock_counters_args) # type: ignore
+AMDKFD_IOC_GET_PROCESS_APERTURES = AMDKFD_IOR(0x06, struct_kfd_ioctl_get_process_apertures_args) # type: ignore
+AMDKFD_IOC_UPDATE_QUEUE = AMDKFD_IOW(0x07, struct_kfd_ioctl_update_queue_args) # type: ignore
+AMDKFD_IOC_CREATE_EVENT = AMDKFD_IOWR(0x08, struct_kfd_ioctl_create_event_args) # type: ignore
+AMDKFD_IOC_DESTROY_EVENT = AMDKFD_IOW(0x09, struct_kfd_ioctl_destroy_event_args) # type: ignore
+AMDKFD_IOC_SET_EVENT = AMDKFD_IOW(0x0A, struct_kfd_ioctl_set_event_args) # type: ignore
+AMDKFD_IOC_RESET_EVENT = AMDKFD_IOW(0x0B, struct_kfd_ioctl_reset_event_args) # type: ignore
+AMDKFD_IOC_WAIT_EVENTS = AMDKFD_IOWR(0x0C, struct_kfd_ioctl_wait_events_args) # type: ignore
+AMDKFD_IOC_DBG_REGISTER_DEPRECATED = AMDKFD_IOW(0x0D, struct_kfd_ioctl_dbg_register_args) # type: ignore
+AMDKFD_IOC_DBG_UNREGISTER_DEPRECATED = AMDKFD_IOW(0x0E, struct_kfd_ioctl_dbg_unregister_args) # type: ignore
+AMDKFD_IOC_DBG_ADDRESS_WATCH_DEPRECATED = AMDKFD_IOW(0x0F, struct_kfd_ioctl_dbg_address_watch_args) # type: ignore
+AMDKFD_IOC_DBG_WAVE_CONTROL_DEPRECATED = AMDKFD_IOW(0x10, struct_kfd_ioctl_dbg_wave_control_args) # type: ignore
+AMDKFD_IOC_SET_SCRATCH_BACKING_VA = AMDKFD_IOWR(0x11, struct_kfd_ioctl_set_scratch_backing_va_args) # type: ignore
+AMDKFD_IOC_GET_TILE_CONFIG = AMDKFD_IOWR(0x12, struct_kfd_ioctl_get_tile_config_args) # type: ignore
+AMDKFD_IOC_SET_TRAP_HANDLER = AMDKFD_IOW(0x13, struct_kfd_ioctl_set_trap_handler_args) # type: ignore
+AMDKFD_IOC_GET_PROCESS_APERTURES_NEW = AMDKFD_IOWR(0x14, struct_kfd_ioctl_get_process_apertures_new_args) # type: ignore
+AMDKFD_IOC_ACQUIRE_VM = AMDKFD_IOW(0x15, struct_kfd_ioctl_acquire_vm_args) # type: ignore
+AMDKFD_IOC_ALLOC_MEMORY_OF_GPU = AMDKFD_IOWR(0x16, struct_kfd_ioctl_alloc_memory_of_gpu_args) # type: ignore
+AMDKFD_IOC_FREE_MEMORY_OF_GPU = AMDKFD_IOW(0x17, struct_kfd_ioctl_free_memory_of_gpu_args) # type: ignore
+AMDKFD_IOC_MAP_MEMORY_TO_GPU = AMDKFD_IOWR(0x18, struct_kfd_ioctl_map_memory_to_gpu_args) # type: ignore
+AMDKFD_IOC_UNMAP_MEMORY_FROM_GPU = AMDKFD_IOWR(0x19, struct_kfd_ioctl_unmap_memory_from_gpu_args) # type: ignore
+AMDKFD_IOC_SET_CU_MASK = AMDKFD_IOW(0x1A, struct_kfd_ioctl_set_cu_mask_args) # type: ignore
+AMDKFD_IOC_GET_QUEUE_WAVE_STATE = AMDKFD_IOWR(0x1B, struct_kfd_ioctl_get_queue_wave_state_args) # type: ignore
+AMDKFD_IOC_GET_DMABUF_INFO = AMDKFD_IOWR(0x1C, struct_kfd_ioctl_get_dmabuf_info_args) # type: ignore
+AMDKFD_IOC_IMPORT_DMABUF = AMDKFD_IOWR(0x1D, struct_kfd_ioctl_import_dmabuf_args) # type: ignore
+AMDKFD_IOC_ALLOC_QUEUE_GWS = AMDKFD_IOWR(0x1E, struct_kfd_ioctl_alloc_queue_gws_args) # type: ignore
+AMDKFD_IOC_SMI_EVENTS = AMDKFD_IOWR(0x1F, struct_kfd_ioctl_smi_events_args) # type: ignore
+AMDKFD_IOC_SVM = AMDKFD_IOWR(0x20, struct_kfd_ioctl_svm_args) # type: ignore
+AMDKFD_IOC_SET_XNACK_MODE = AMDKFD_IOWR(0x21, struct_kfd_ioctl_set_xnack_mode_args) # type: ignore
+AMDKFD_IOC_CRIU_OP = AMDKFD_IOWR(0x22, struct_kfd_ioctl_criu_args) # type: ignore
+AMDKFD_IOC_AVAILABLE_MEMORY = AMDKFD_IOWR(0x23, struct_kfd_ioctl_get_available_memory_args) # type: ignore
+AMDKFD_IOC_EXPORT_DMABUF = AMDKFD_IOWR(0x24, struct_kfd_ioctl_export_dmabuf_args) # type: ignore
+AMDKFD_IOC_RUNTIME_ENABLE = AMDKFD_IOWR(0x25, struct_kfd_ioctl_runtime_enable_args) # type: ignore
+AMDKFD_IOC_DBG_TRAP = AMDKFD_IOWR(0x26, struct_kfd_ioctl_dbg_trap_args) # type: ignore
+AMDKFD_COMMAND_START = 0x01 # type: ignore
+AMDKFD_COMMAND_END = 0x27 # type: ignore
+AMDKFD_IOC_IPC_IMPORT_HANDLE = AMDKFD_IOWR(0x80, struct_kfd_ioctl_ipc_import_handle_args) # type: ignore
+AMDKFD_IOC_IPC_EXPORT_HANDLE = AMDKFD_IOWR(0x81, struct_kfd_ioctl_ipc_export_handle_args) # type: ignore
+AMDKFD_IOC_DBG_TRAP_DEPRECATED = AMDKFD_IOWR(0x82, struct_kfd_ioctl_dbg_trap_args_deprecated) # type: ignore
+AMDKFD_IOC_CROSS_MEMORY_COPY_DEPRECATED = AMDKFD_IOWR(0x83, struct_kfd_ioctl_cross_memory_copy_deprecated_args) # type: ignore
+AMDKFD_IOC_RLC_SPM = AMDKFD_IOWR(0x84, struct_kfd_ioctl_spm_args) # type: ignore
+AMDKFD_IOC_PC_SAMPLE = AMDKFD_IOWR(0x85, struct_kfd_ioctl_pc_sample_args) # type: ignore
+AMDKFD_IOC_PROFILER = AMDKFD_IOWR(0x86, struct_kfd_ioctl_profiler_args) # type: ignore
+AMDKFD_COMMAND_START_2 = 0x80 # type: ignore
+AMDKFD_COMMAND_END_2 = 0x87 # type: ignore

tinygrad/runtime/autogen/kgsl.py

@@ -552,236 +552,236 @@ class struct_kgsl_gpuobj_set_info(c.Struct):
   metadata_len: int
   type: int
 struct_kgsl_gpuobj_set_info.register_fields([('flags', ctypes.c_uint64, 0), ('metadata', ctypes.c_uint64, 8), ('id', ctypes.c_uint32, 16), ('metadata_len', ctypes.c_uint32, 20), ('type', ctypes.c_uint32, 24)])
-KGSL_VERSION_MAJOR = 3
-KGSL_VERSION_MINOR = 14
-KGSL_CONTEXT_SAVE_GMEM = 0x00000001
-KGSL_CONTEXT_NO_GMEM_ALLOC = 0x00000002
-KGSL_CONTEXT_SUBMIT_IB_LIST = 0x00000004
-KGSL_CONTEXT_CTX_SWITCH = 0x00000008
-KGSL_CONTEXT_PREAMBLE = 0x00000010
-KGSL_CONTEXT_TRASH_STATE = 0x00000020
-KGSL_CONTEXT_PER_CONTEXT_TS = 0x00000040
-KGSL_CONTEXT_USER_GENERATED_TS = 0x00000080
-KGSL_CONTEXT_END_OF_FRAME = 0x00000100
-KGSL_CONTEXT_NO_FAULT_TOLERANCE = 0x00000200
-KGSL_CONTEXT_SYNC = 0x00000400
-KGSL_CONTEXT_PWR_CONSTRAINT = 0x00000800
-KGSL_CONTEXT_PRIORITY_MASK = 0x0000F000
-KGSL_CONTEXT_PRIORITY_SHIFT = 12
-KGSL_CONTEXT_PRIORITY_UNDEF = 0
-KGSL_CONTEXT_IFH_NOP = 0x00010000
-KGSL_CONTEXT_SECURE = 0x00020000
-KGSL_CONTEXT_PREEMPT_STYLE_MASK = 0x0E000000
-KGSL_CONTEXT_PREEMPT_STYLE_SHIFT = 25
-KGSL_CONTEXT_PREEMPT_STYLE_DEFAULT = 0x0
-KGSL_CONTEXT_PREEMPT_STYLE_RINGBUFFER = 0x1
-KGSL_CONTEXT_PREEMPT_STYLE_FINEGRAIN = 0x2
-KGSL_CONTEXT_TYPE_MASK = 0x01F00000
-KGSL_CONTEXT_TYPE_SHIFT = 20
-KGSL_CONTEXT_TYPE_ANY = 0
-KGSL_CONTEXT_TYPE_GL = 1
-KGSL_CONTEXT_TYPE_CL = 2
-KGSL_CONTEXT_TYPE_C2D = 3
-KGSL_CONTEXT_TYPE_RS = 4
-KGSL_CONTEXT_TYPE_UNKNOWN = 0x1E
-KGSL_CONTEXT_INVALID = 0xffffffff
-KGSL_CMDBATCH_MEMLIST = 0x00000001
-KGSL_CMDBATCH_MARKER = 0x00000002
-KGSL_CMDBATCH_SUBMIT_IB_LIST = KGSL_CONTEXT_SUBMIT_IB_LIST
-KGSL_CMDBATCH_CTX_SWITCH = KGSL_CONTEXT_CTX_SWITCH
-KGSL_CMDBATCH_PROFILING = 0x00000010
-KGSL_CMDBATCH_PROFILING_KTIME = 0x00000020
-KGSL_CMDBATCH_END_OF_FRAME = KGSL_CONTEXT_END_OF_FRAME
-KGSL_CMDBATCH_SYNC = KGSL_CONTEXT_SYNC
-KGSL_CMDBATCH_PWR_CONSTRAINT = KGSL_CONTEXT_PWR_CONSTRAINT
-KGSL_CMDLIST_IB = 0x00000001
-KGSL_CMDLIST_CTXTSWITCH_PREAMBLE = 0x00000002
-KGSL_CMDLIST_IB_PREAMBLE = 0x00000004
-KGSL_OBJLIST_MEMOBJ = 0x00000008
-KGSL_OBJLIST_PROFILE = 0x00000010
-KGSL_CMD_SYNCPOINT_TYPE_TIMESTAMP = 0
-KGSL_CMD_SYNCPOINT_TYPE_FENCE = 1
-KGSL_MEMFLAGS_SECURE = 0x00000008
-KGSL_MEMFLAGS_GPUREADONLY = 0x01000000
-KGSL_MEMFLAGS_GPUWRITEONLY = 0x02000000
-KGSL_MEMFLAGS_FORCE_32BIT = 0x100000000
-KGSL_CACHEMODE_MASK = 0x0C000000
-KGSL_CACHEMODE_SHIFT = 26
-KGSL_CACHEMODE_WRITECOMBINE = 0
-KGSL_CACHEMODE_UNCACHED = 1
-KGSL_CACHEMODE_WRITETHROUGH = 2
-KGSL_CACHEMODE_WRITEBACK = 3
-KGSL_MEMFLAGS_USE_CPU_MAP = 0x10000000
-KGSL_MEMTYPE_MASK = 0x0000FF00
-KGSL_MEMTYPE_SHIFT = 8
-KGSL_MEMTYPE_OBJECTANY = 0
-KGSL_MEMTYPE_FRAMEBUFFER = 1
-KGSL_MEMTYPE_RENDERBUFFER = 2
-KGSL_MEMTYPE_ARRAYBUFFER = 3
-KGSL_MEMTYPE_ELEMENTARRAYBUFFER = 4
-KGSL_MEMTYPE_VERTEXARRAYBUFFER = 5
-KGSL_MEMTYPE_TEXTURE = 6
-KGSL_MEMTYPE_SURFACE = 7
-KGSL_MEMTYPE_EGL_SURFACE = 8
-KGSL_MEMTYPE_GL = 9
-KGSL_MEMTYPE_CL = 10
-KGSL_MEMTYPE_CL_BUFFER_MAP = 11
-KGSL_MEMTYPE_CL_BUFFER_NOMAP = 12
-KGSL_MEMTYPE_CL_IMAGE_MAP = 13
-KGSL_MEMTYPE_CL_IMAGE_NOMAP = 14
-KGSL_MEMTYPE_CL_KERNEL_STACK = 15
-KGSL_MEMTYPE_COMMAND = 16
-KGSL_MEMTYPE_2D = 17
-KGSL_MEMTYPE_EGL_IMAGE = 18
-KGSL_MEMTYPE_EGL_SHADOW = 19
-KGSL_MEMTYPE_MULTISAMPLE = 20
-KGSL_MEMTYPE_KERNEL = 255
-KGSL_MEMALIGN_MASK = 0x00FF0000
-KGSL_MEMALIGN_SHIFT = 16
-KGSL_MEMFLAGS_USERMEM_MASK = 0x000000e0
-KGSL_MEMFLAGS_USERMEM_SHIFT = 5
+KGSL_VERSION_MAJOR = 3 # type: ignore
+KGSL_VERSION_MINOR = 14 # type: ignore
+KGSL_CONTEXT_SAVE_GMEM = 0x00000001 # type: ignore
+KGSL_CONTEXT_NO_GMEM_ALLOC = 0x00000002 # type: ignore
+KGSL_CONTEXT_SUBMIT_IB_LIST = 0x00000004 # type: ignore
+KGSL_CONTEXT_CTX_SWITCH = 0x00000008 # type: ignore
+KGSL_CONTEXT_PREAMBLE = 0x00000010 # type: ignore
+KGSL_CONTEXT_TRASH_STATE = 0x00000020 # type: ignore
+KGSL_CONTEXT_PER_CONTEXT_TS = 0x00000040 # type: ignore
+KGSL_CONTEXT_USER_GENERATED_TS = 0x00000080 # type: ignore
+KGSL_CONTEXT_END_OF_FRAME = 0x00000100 # type: ignore
+KGSL_CONTEXT_NO_FAULT_TOLERANCE = 0x00000200 # type: ignore
+KGSL_CONTEXT_SYNC = 0x00000400 # type: ignore
+KGSL_CONTEXT_PWR_CONSTRAINT = 0x00000800 # type: ignore
+KGSL_CONTEXT_PRIORITY_MASK = 0x0000F000 # type: ignore
+KGSL_CONTEXT_PRIORITY_SHIFT = 12 # type: ignore
+KGSL_CONTEXT_PRIORITY_UNDEF = 0 # type: ignore
+KGSL_CONTEXT_IFH_NOP = 0x00010000 # type: ignore
+KGSL_CONTEXT_SECURE = 0x00020000 # type: ignore
+KGSL_CONTEXT_PREEMPT_STYLE_MASK = 0x0E000000 # type: ignore
+KGSL_CONTEXT_PREEMPT_STYLE_SHIFT = 25 # type: ignore
+KGSL_CONTEXT_PREEMPT_STYLE_DEFAULT = 0x0 # type: ignore
+KGSL_CONTEXT_PREEMPT_STYLE_RINGBUFFER = 0x1 # type: ignore
+KGSL_CONTEXT_PREEMPT_STYLE_FINEGRAIN = 0x2 # type: ignore
+KGSL_CONTEXT_TYPE_MASK = 0x01F00000 # type: ignore
+KGSL_CONTEXT_TYPE_SHIFT = 20 # type: ignore
+KGSL_CONTEXT_TYPE_ANY = 0 # type: ignore
+KGSL_CONTEXT_TYPE_GL = 1 # type: ignore
+KGSL_CONTEXT_TYPE_CL = 2 # type: ignore
+KGSL_CONTEXT_TYPE_C2D = 3 # type: ignore
+KGSL_CONTEXT_TYPE_RS = 4 # type: ignore
+KGSL_CONTEXT_TYPE_UNKNOWN = 0x1E # type: ignore
+KGSL_CONTEXT_INVALID = 0xffffffff # type: ignore
+KGSL_CMDBATCH_MEMLIST = 0x00000001 # type: ignore
+KGSL_CMDBATCH_MARKER = 0x00000002 # type: ignore
+KGSL_CMDBATCH_SUBMIT_IB_LIST = KGSL_CONTEXT_SUBMIT_IB_LIST # type: ignore
+KGSL_CMDBATCH_CTX_SWITCH = KGSL_CONTEXT_CTX_SWITCH # type: ignore
+KGSL_CMDBATCH_PROFILING = 0x00000010 # type: ignore
+KGSL_CMDBATCH_PROFILING_KTIME = 0x00000020 # type: ignore
+KGSL_CMDBATCH_END_OF_FRAME = KGSL_CONTEXT_END_OF_FRAME # type: ignore
+KGSL_CMDBATCH_SYNC = KGSL_CONTEXT_SYNC # type: ignore
+KGSL_CMDBATCH_PWR_CONSTRAINT = KGSL_CONTEXT_PWR_CONSTRAINT # type: ignore
+KGSL_CMDLIST_IB = 0x00000001 # type: ignore
+KGSL_CMDLIST_CTXTSWITCH_PREAMBLE = 0x00000002 # type: ignore
+KGSL_CMDLIST_IB_PREAMBLE = 0x00000004 # type: ignore
+KGSL_OBJLIST_MEMOBJ = 0x00000008 # type: ignore
+KGSL_OBJLIST_PROFILE = 0x00000010 # type: ignore
+KGSL_CMD_SYNCPOINT_TYPE_TIMESTAMP = 0 # type: ignore
+KGSL_CMD_SYNCPOINT_TYPE_FENCE = 1 # type: ignore
+KGSL_MEMFLAGS_SECURE = 0x00000008 # type: ignore
+KGSL_MEMFLAGS_GPUREADONLY = 0x01000000 # type: ignore
+KGSL_MEMFLAGS_GPUWRITEONLY = 0x02000000 # type: ignore
+KGSL_MEMFLAGS_FORCE_32BIT = 0x100000000 # type: ignore
+KGSL_CACHEMODE_MASK = 0x0C000000 # type: ignore
+KGSL_CACHEMODE_SHIFT = 26 # type: ignore
+KGSL_CACHEMODE_WRITECOMBINE = 0 # type: ignore
+KGSL_CACHEMODE_UNCACHED = 1 # type: ignore
+KGSL_CACHEMODE_WRITETHROUGH = 2 # type: ignore
+KGSL_CACHEMODE_WRITEBACK = 3 # type: ignore
+KGSL_MEMFLAGS_USE_CPU_MAP = 0x10000000 # type: ignore
+KGSL_MEMTYPE_MASK = 0x0000FF00 # type: ignore
+KGSL_MEMTYPE_SHIFT = 8 # type: ignore
+KGSL_MEMTYPE_OBJECTANY = 0 # type: ignore
+KGSL_MEMTYPE_FRAMEBUFFER = 1 # type: ignore
+KGSL_MEMTYPE_RENDERBUFFER = 2 # type: ignore
+KGSL_MEMTYPE_ARRAYBUFFER = 3 # type: ignore
+KGSL_MEMTYPE_ELEMENTARRAYBUFFER = 4 # type: ignore
+KGSL_MEMTYPE_VERTEXARRAYBUFFER = 5 # type: ignore
+KGSL_MEMTYPE_TEXTURE = 6 # type: ignore
+KGSL_MEMTYPE_SURFACE = 7 # type: ignore
+KGSL_MEMTYPE_EGL_SURFACE = 8 # type: ignore
+KGSL_MEMTYPE_GL = 9 # type: ignore
+KGSL_MEMTYPE_CL = 10 # type: ignore
+KGSL_MEMTYPE_CL_BUFFER_MAP = 11 # type: ignore
+KGSL_MEMTYPE_CL_BUFFER_NOMAP = 12 # type: ignore
+KGSL_MEMTYPE_CL_IMAGE_MAP = 13 # type: ignore
+KGSL_MEMTYPE_CL_IMAGE_NOMAP = 14 # type: ignore
+KGSL_MEMTYPE_CL_KERNEL_STACK = 15 # type: ignore
+KGSL_MEMTYPE_COMMAND = 16 # type: ignore
+KGSL_MEMTYPE_2D = 17 # type: ignore
+KGSL_MEMTYPE_EGL_IMAGE = 18 # type: ignore
+KGSL_MEMTYPE_EGL_SHADOW = 19 # type: ignore
+KGSL_MEMTYPE_MULTISAMPLE = 20 # type: ignore
+KGSL_MEMTYPE_KERNEL = 255 # type: ignore
+KGSL_MEMALIGN_MASK = 0x00FF0000 # type: ignore
+KGSL_MEMALIGN_SHIFT = 16 # type: ignore
+KGSL_MEMFLAGS_USERMEM_MASK = 0x000000e0 # type: ignore
+KGSL_MEMFLAGS_USERMEM_SHIFT = 5 # type: ignore
 KGSL_USERMEM_FLAG = lambda x: (((x) + 1) << KGSL_MEMFLAGS_USERMEM_SHIFT) # type: ignore
-KGSL_MEMFLAGS_NOT_USERMEM = 0
-KGSL_MEMFLAGS_USERMEM_PMEM = KGSL_USERMEM_FLAG(KGSL_USER_MEM_TYPE_PMEM)
-KGSL_MEMFLAGS_USERMEM_ASHMEM = KGSL_USERMEM_FLAG(KGSL_USER_MEM_TYPE_ASHMEM)
-KGSL_MEMFLAGS_USERMEM_ADDR = KGSL_USERMEM_FLAG(KGSL_USER_MEM_TYPE_ADDR)
-KGSL_MEMFLAGS_USERMEM_ION = KGSL_USERMEM_FLAG(KGSL_USER_MEM_TYPE_ION)
-KGSL_FLAGS_NORMALMODE = 0x00000000
-KGSL_FLAGS_SAFEMODE = 0x00000001
-KGSL_FLAGS_INITIALIZED0 = 0x00000002
-KGSL_FLAGS_INITIALIZED = 0x00000004
-KGSL_FLAGS_STARTED = 0x00000008
-KGSL_FLAGS_ACTIVE = 0x00000010
-KGSL_FLAGS_RESERVED0 = 0x00000020
-KGSL_FLAGS_RESERVED1 = 0x00000040
-KGSL_FLAGS_RESERVED2 = 0x00000080
-KGSL_FLAGS_SOFT_RESET = 0x00000100
-KGSL_FLAGS_PER_CONTEXT_TIMESTAMPS = 0x00000200
-KGSL_SYNCOBJ_SERVER_TIMEOUT = 2000
+KGSL_MEMFLAGS_NOT_USERMEM = 0 # type: ignore
+KGSL_MEMFLAGS_USERMEM_PMEM = KGSL_USERMEM_FLAG(KGSL_USER_MEM_TYPE_PMEM) # type: ignore
+KGSL_MEMFLAGS_USERMEM_ASHMEM = KGSL_USERMEM_FLAG(KGSL_USER_MEM_TYPE_ASHMEM) # type: ignore
+KGSL_MEMFLAGS_USERMEM_ADDR = KGSL_USERMEM_FLAG(KGSL_USER_MEM_TYPE_ADDR) # type: ignore
+KGSL_MEMFLAGS_USERMEM_ION = KGSL_USERMEM_FLAG(KGSL_USER_MEM_TYPE_ION) # type: ignore
+KGSL_FLAGS_NORMALMODE = 0x00000000 # type: ignore
+KGSL_FLAGS_SAFEMODE = 0x00000001 # type: ignore
+KGSL_FLAGS_INITIALIZED0 = 0x00000002 # type: ignore
+KGSL_FLAGS_INITIALIZED = 0x00000004 # type: ignore
+KGSL_FLAGS_STARTED = 0x00000008 # type: ignore
+KGSL_FLAGS_ACTIVE = 0x00000010 # type: ignore
+KGSL_FLAGS_RESERVED0 = 0x00000020 # type: ignore
+KGSL_FLAGS_RESERVED1 = 0x00000040 # type: ignore
+KGSL_FLAGS_RESERVED2 = 0x00000080 # type: ignore
+KGSL_FLAGS_SOFT_RESET = 0x00000100 # type: ignore
+KGSL_FLAGS_PER_CONTEXT_TIMESTAMPS = 0x00000200 # type: ignore
+KGSL_SYNCOBJ_SERVER_TIMEOUT = 2000 # type: ignore
 KGSL_CONVERT_TO_MBPS = lambda val: (val*1000*1000) # type: ignore
 KGSL_MEMSTORE_OFFSET = lambda ctxt_id,field: ((ctxt_id)*sizeof(struct_kgsl_devmemstore) + offsetof(struct_kgsl_devmemstore, field)) # type: ignore
-KGSL_PROP_DEVICE_INFO = 0x1
-KGSL_PROP_DEVICE_SHADOW = 0x2
-KGSL_PROP_DEVICE_POWER = 0x3
-KGSL_PROP_SHMEM = 0x4
-KGSL_PROP_SHMEM_APERTURES = 0x5
-KGSL_PROP_MMU_ENABLE = 0x6
-KGSL_PROP_INTERRUPT_WAITS = 0x7
-KGSL_PROP_VERSION = 0x8
-KGSL_PROP_GPU_RESET_STAT = 0x9
-KGSL_PROP_PWRCTRL = 0xE
-KGSL_PROP_PWR_CONSTRAINT = 0x12
-KGSL_PROP_UCHE_GMEM_VADDR = 0x13
-KGSL_PROP_SP_GENERIC_MEM = 0x14
-KGSL_PROP_UCODE_VERSION = 0x15
-KGSL_PROP_GPMU_VERSION = 0x16
-KGSL_PROP_DEVICE_BITNESS = 0x18
-KGSL_PERFCOUNTER_GROUP_CP = 0x0
-KGSL_PERFCOUNTER_GROUP_RBBM = 0x1
-KGSL_PERFCOUNTER_GROUP_PC = 0x2
-KGSL_PERFCOUNTER_GROUP_VFD = 0x3
-KGSL_PERFCOUNTER_GROUP_HLSQ = 0x4
-KGSL_PERFCOUNTER_GROUP_VPC = 0x5
-KGSL_PERFCOUNTER_GROUP_TSE = 0x6
-KGSL_PERFCOUNTER_GROUP_RAS = 0x7
-KGSL_PERFCOUNTER_GROUP_UCHE = 0x8
-KGSL_PERFCOUNTER_GROUP_TP = 0x9
-KGSL_PERFCOUNTER_GROUP_SP = 0xA
-KGSL_PERFCOUNTER_GROUP_RB = 0xB
-KGSL_PERFCOUNTER_GROUP_PWR = 0xC
-KGSL_PERFCOUNTER_GROUP_VBIF = 0xD
-KGSL_PERFCOUNTER_GROUP_VBIF_PWR = 0xE
-KGSL_PERFCOUNTER_GROUP_MH = 0xF
-KGSL_PERFCOUNTER_GROUP_PA_SU = 0x10
-KGSL_PERFCOUNTER_GROUP_SQ = 0x11
-KGSL_PERFCOUNTER_GROUP_SX = 0x12
-KGSL_PERFCOUNTER_GROUP_TCF = 0x13
-KGSL_PERFCOUNTER_GROUP_TCM = 0x14
-KGSL_PERFCOUNTER_GROUP_TCR = 0x15
-KGSL_PERFCOUNTER_GROUP_L2 = 0x16
-KGSL_PERFCOUNTER_GROUP_VSC = 0x17
-KGSL_PERFCOUNTER_GROUP_CCU = 0x18
-KGSL_PERFCOUNTER_GROUP_LRZ = 0x19
-KGSL_PERFCOUNTER_GROUP_CMP = 0x1A
-KGSL_PERFCOUNTER_GROUP_ALWAYSON = 0x1B
-KGSL_PERFCOUNTER_GROUP_SP_PWR = 0x1C
-KGSL_PERFCOUNTER_GROUP_TP_PWR = 0x1D
-KGSL_PERFCOUNTER_GROUP_RB_PWR = 0x1E
-KGSL_PERFCOUNTER_GROUP_CCU_PWR = 0x1F
-KGSL_PERFCOUNTER_GROUP_UCHE_PWR = 0x20
-KGSL_PERFCOUNTER_GROUP_CP_PWR = 0x21
-KGSL_PERFCOUNTER_GROUP_GPMU_PWR = 0x22
-KGSL_PERFCOUNTER_GROUP_ALWAYSON_PWR = 0x23
-KGSL_PERFCOUNTER_GROUP_MAX = 0x24
-KGSL_PERFCOUNTER_NOT_USED = 0xFFFFFFFF
-KGSL_PERFCOUNTER_BROKEN = 0xFFFFFFFE
-KGSL_IOC_TYPE = 0x09
-IOCTL_KGSL_DEVICE_GETPROPERTY = _IOWR(KGSL_IOC_TYPE, 0x2, struct_kgsl_device_getproperty)
-IOCTL_KGSL_DEVICE_WAITTIMESTAMP = _IOW(KGSL_IOC_TYPE, 0x6, struct_kgsl_device_waittimestamp)
-IOCTL_KGSL_DEVICE_WAITTIMESTAMP_CTXTID = _IOW(KGSL_IOC_TYPE, 0x7, struct_kgsl_device_waittimestamp_ctxtid)
-IOCTL_KGSL_RINGBUFFER_ISSUEIBCMDS = _IOWR(KGSL_IOC_TYPE, 0x10, struct_kgsl_ringbuffer_issueibcmds)
-IOCTL_KGSL_CMDSTREAM_READTIMESTAMP_OLD = _IOR(KGSL_IOC_TYPE, 0x11, struct_kgsl_cmdstream_readtimestamp)
-IOCTL_KGSL_CMDSTREAM_READTIMESTAMP = _IOWR(KGSL_IOC_TYPE, 0x11, struct_kgsl_cmdstream_readtimestamp)
-IOCTL_KGSL_CMDSTREAM_FREEMEMONTIMESTAMP = _IOW(KGSL_IOC_TYPE, 0x12, struct_kgsl_cmdstream_freememontimestamp)
-IOCTL_KGSL_CMDSTREAM_FREEMEMONTIMESTAMP_OLD = _IOR(KGSL_IOC_TYPE, 0x12, struct_kgsl_cmdstream_freememontimestamp)
-IOCTL_KGSL_DRAWCTXT_CREATE = _IOWR(KGSL_IOC_TYPE, 0x13, struct_kgsl_drawctxt_create)
-IOCTL_KGSL_DRAWCTXT_DESTROY = _IOW(KGSL_IOC_TYPE, 0x14, struct_kgsl_drawctxt_destroy)
-IOCTL_KGSL_MAP_USER_MEM = _IOWR(KGSL_IOC_TYPE, 0x15, struct_kgsl_map_user_mem)
-IOCTL_KGSL_CMDSTREAM_READTIMESTAMP_CTXTID = _IOWR(KGSL_IOC_TYPE, 0x16, struct_kgsl_cmdstream_readtimestamp_ctxtid)
-IOCTL_KGSL_CMDSTREAM_FREEMEMONTIMESTAMP_CTXTID = _IOW(KGSL_IOC_TYPE, 0x17, struct_kgsl_cmdstream_freememontimestamp_ctxtid)
-IOCTL_KGSL_SHAREDMEM_FROM_PMEM = _IOWR(KGSL_IOC_TYPE, 0x20, struct_kgsl_sharedmem_from_pmem)
-IOCTL_KGSL_SHAREDMEM_FREE = _IOW(KGSL_IOC_TYPE, 0x21, struct_kgsl_sharedmem_free)
-IOCTL_KGSL_CFF_USER_EVENT = _IOW(KGSL_IOC_TYPE, 0x31, struct_kgsl_cff_user_event)
-IOCTL_KGSL_DRAWCTXT_BIND_GMEM_SHADOW = _IOW(KGSL_IOC_TYPE, 0x22, struct_kgsl_bind_gmem_shadow)
-IOCTL_KGSL_SHAREDMEM_FROM_VMALLOC = _IOWR(KGSL_IOC_TYPE, 0x23, struct_kgsl_sharedmem_from_vmalloc)
-IOCTL_KGSL_SHAREDMEM_FLUSH_CACHE = _IOW(KGSL_IOC_TYPE, 0x24, struct_kgsl_sharedmem_free)
-IOCTL_KGSL_DRAWCTXT_SET_BIN_BASE_OFFSET = _IOW(KGSL_IOC_TYPE, 0x25, struct_kgsl_drawctxt_set_bin_base_offset)
-IOCTL_KGSL_CMDWINDOW_WRITE = _IOW(KGSL_IOC_TYPE, 0x2e, struct_kgsl_cmdwindow_write)
-IOCTL_KGSL_GPUMEM_ALLOC = _IOWR(KGSL_IOC_TYPE, 0x2f, struct_kgsl_gpumem_alloc)
-IOCTL_KGSL_CFF_SYNCMEM = _IOW(KGSL_IOC_TYPE, 0x30, struct_kgsl_cff_syncmem)
-IOCTL_KGSL_TIMESTAMP_EVENT_OLD = _IOW(KGSL_IOC_TYPE, 0x31, struct_kgsl_timestamp_event)
-KGSL_TIMESTAMP_EVENT_GENLOCK = 1
-KGSL_TIMESTAMP_EVENT_FENCE = 2
-IOCTL_KGSL_SETPROPERTY = _IOW(KGSL_IOC_TYPE, 0x32, struct_kgsl_device_getproperty)
-IOCTL_KGSL_TIMESTAMP_EVENT = _IOWR(KGSL_IOC_TYPE, 0x33, struct_kgsl_timestamp_event)
-IOCTL_KGSL_GPUMEM_ALLOC_ID = _IOWR(KGSL_IOC_TYPE, 0x34, struct_kgsl_gpumem_alloc_id)
-IOCTL_KGSL_GPUMEM_FREE_ID = _IOWR(KGSL_IOC_TYPE, 0x35, struct_kgsl_gpumem_free_id)
-IOCTL_KGSL_GPUMEM_GET_INFO = _IOWR(KGSL_IOC_TYPE, 0x36, struct_kgsl_gpumem_get_info)
-KGSL_GPUMEM_CACHE_CLEAN = (1 << 0)
-KGSL_GPUMEM_CACHE_TO_GPU = KGSL_GPUMEM_CACHE_CLEAN
-KGSL_GPUMEM_CACHE_INV = (1 << 1)
-KGSL_GPUMEM_CACHE_FROM_GPU = KGSL_GPUMEM_CACHE_INV
-KGSL_GPUMEM_CACHE_FLUSH = (KGSL_GPUMEM_CACHE_CLEAN | KGSL_GPUMEM_CACHE_INV)
-KGSL_GPUMEM_CACHE_RANGE = (1 << 31)
-IOCTL_KGSL_GPUMEM_SYNC_CACHE = _IOW(KGSL_IOC_TYPE, 0x37, struct_kgsl_gpumem_sync_cache)
-IOCTL_KGSL_PERFCOUNTER_GET = _IOWR(KGSL_IOC_TYPE, 0x38, struct_kgsl_perfcounter_get)
-IOCTL_KGSL_PERFCOUNTER_PUT = _IOW(KGSL_IOC_TYPE, 0x39, struct_kgsl_perfcounter_put)
-IOCTL_KGSL_PERFCOUNTER_QUERY = _IOWR(KGSL_IOC_TYPE, 0x3A, struct_kgsl_perfcounter_query)
-IOCTL_KGSL_PERFCOUNTER_READ = _IOWR(KGSL_IOC_TYPE, 0x3B, struct_kgsl_perfcounter_read)
-IOCTL_KGSL_GPUMEM_SYNC_CACHE_BULK = _IOWR(KGSL_IOC_TYPE, 0x3C, struct_kgsl_gpumem_sync_cache_bulk)
-KGSL_IBDESC_MEMLIST = 0x1
-KGSL_IBDESC_PROFILING_BUFFER = 0x2
-IOCTL_KGSL_SUBMIT_COMMANDS = _IOWR(KGSL_IOC_TYPE, 0x3D, struct_kgsl_submit_commands)
-KGSL_CONSTRAINT_NONE = 0
-KGSL_CONSTRAINT_PWRLEVEL = 1
-KGSL_CONSTRAINT_PWR_MIN = 0
-KGSL_CONSTRAINT_PWR_MAX = 1
-IOCTL_KGSL_SYNCSOURCE_CREATE = _IOWR(KGSL_IOC_TYPE, 0x40, struct_kgsl_syncsource_create)
-IOCTL_KGSL_SYNCSOURCE_DESTROY = _IOWR(KGSL_IOC_TYPE, 0x41, struct_kgsl_syncsource_destroy)
-IOCTL_KGSL_SYNCSOURCE_CREATE_FENCE = _IOWR(KGSL_IOC_TYPE, 0x42, struct_kgsl_syncsource_create_fence)
-IOCTL_KGSL_SYNCSOURCE_SIGNAL_FENCE = _IOWR(KGSL_IOC_TYPE, 0x43, struct_kgsl_syncsource_signal_fence)
-IOCTL_KGSL_CFF_SYNC_GPUOBJ = _IOW(KGSL_IOC_TYPE, 0x44, struct_kgsl_cff_sync_gpuobj)
-KGSL_GPUOBJ_ALLOC_METADATA_MAX = 64
-IOCTL_KGSL_GPUOBJ_ALLOC = _IOWR(KGSL_IOC_TYPE, 0x45, struct_kgsl_gpuobj_alloc)
-KGSL_GPUOBJ_FREE_ON_EVENT = 1
-KGSL_GPU_EVENT_TIMESTAMP = 1
-KGSL_GPU_EVENT_FENCE = 2
-IOCTL_KGSL_GPUOBJ_FREE = _IOW(KGSL_IOC_TYPE, 0x46, struct_kgsl_gpuobj_free)
-IOCTL_KGSL_GPUOBJ_INFO = _IOWR(KGSL_IOC_TYPE, 0x47, struct_kgsl_gpuobj_info)
-IOCTL_KGSL_GPUOBJ_IMPORT = _IOWR(KGSL_IOC_TYPE, 0x48, struct_kgsl_gpuobj_import)
-IOCTL_KGSL_GPUOBJ_SYNC = _IOW(KGSL_IOC_TYPE, 0x49, struct_kgsl_gpuobj_sync)
-IOCTL_KGSL_GPU_COMMAND = _IOWR(KGSL_IOC_TYPE, 0x4A, struct_kgsl_gpu_command)
-IOCTL_KGSL_PREEMPTIONCOUNTER_QUERY = _IOWR(KGSL_IOC_TYPE, 0x4B, struct_kgsl_preemption_counters_query)
-KGSL_GPUOBJ_SET_INFO_METADATA = (1 << 0)
-KGSL_GPUOBJ_SET_INFO_TYPE = (1 << 1)
-IOCTL_KGSL_GPUOBJ_SET_INFO = _IOW(KGSL_IOC_TYPE, 0x4C, struct_kgsl_gpuobj_set_info)
+KGSL_PROP_DEVICE_INFO = 0x1 # type: ignore
+KGSL_PROP_DEVICE_SHADOW = 0x2 # type: ignore
+KGSL_PROP_DEVICE_POWER = 0x3 # type: ignore
+KGSL_PROP_SHMEM = 0x4 # type: ignore
+KGSL_PROP_SHMEM_APERTURES = 0x5 # type: ignore
+KGSL_PROP_MMU_ENABLE = 0x6 # type: ignore
+KGSL_PROP_INTERRUPT_WAITS = 0x7 # type: ignore
+KGSL_PROP_VERSION = 0x8 # type: ignore
+KGSL_PROP_GPU_RESET_STAT = 0x9 # type: ignore
+KGSL_PROP_PWRCTRL = 0xE # type: ignore
+KGSL_PROP_PWR_CONSTRAINT = 0x12 # type: ignore
+KGSL_PROP_UCHE_GMEM_VADDR = 0x13 # type: ignore
+KGSL_PROP_SP_GENERIC_MEM = 0x14 # type: ignore
+KGSL_PROP_UCODE_VERSION = 0x15 # type: ignore
+KGSL_PROP_GPMU_VERSION = 0x16 # type: ignore
+KGSL_PROP_DEVICE_BITNESS = 0x18 # type: ignore
+KGSL_PERFCOUNTER_GROUP_CP = 0x0 # type: ignore
+KGSL_PERFCOUNTER_GROUP_RBBM = 0x1 # type: ignore
+KGSL_PERFCOUNTER_GROUP_PC = 0x2 # type: ignore
+KGSL_PERFCOUNTER_GROUP_VFD = 0x3 # type: ignore
+KGSL_PERFCOUNTER_GROUP_HLSQ = 0x4 # type: ignore
+KGSL_PERFCOUNTER_GROUP_VPC = 0x5 # type: ignore
+KGSL_PERFCOUNTER_GROUP_TSE = 0x6 # type: ignore
+KGSL_PERFCOUNTER_GROUP_RAS = 0x7 # type: ignore
+KGSL_PERFCOUNTER_GROUP_UCHE = 0x8 # type: ignore
+KGSL_PERFCOUNTER_GROUP_TP = 0x9 # type: ignore
+KGSL_PERFCOUNTER_GROUP_SP = 0xA # type: ignore
+KGSL_PERFCOUNTER_GROUP_RB = 0xB # type: ignore
+KGSL_PERFCOUNTER_GROUP_PWR = 0xC # type: ignore
+KGSL_PERFCOUNTER_GROUP_VBIF = 0xD # type: ignore
+KGSL_PERFCOUNTER_GROUP_VBIF_PWR = 0xE # type: ignore
+KGSL_PERFCOUNTER_GROUP_MH = 0xF # type: ignore
+KGSL_PERFCOUNTER_GROUP_PA_SU = 0x10 # type: ignore
+KGSL_PERFCOUNTER_GROUP_SQ = 0x11 # type: ignore
+KGSL_PERFCOUNTER_GROUP_SX = 0x12 # type: ignore
+KGSL_PERFCOUNTER_GROUP_TCF = 0x13 # type: ignore
+KGSL_PERFCOUNTER_GROUP_TCM = 0x14 # type: ignore
+KGSL_PERFCOUNTER_GROUP_TCR = 0x15 # type: ignore
+KGSL_PERFCOUNTER_GROUP_L2 = 0x16 # type: ignore
+KGSL_PERFCOUNTER_GROUP_VSC = 0x17 # type: ignore
+KGSL_PERFCOUNTER_GROUP_CCU = 0x18 # type: ignore
+KGSL_PERFCOUNTER_GROUP_LRZ = 0x19 # type: ignore
+KGSL_PERFCOUNTER_GROUP_CMP = 0x1A # type: ignore
+KGSL_PERFCOUNTER_GROUP_ALWAYSON = 0x1B # type: ignore
+KGSL_PERFCOUNTER_GROUP_SP_PWR = 0x1C # type: ignore
+KGSL_PERFCOUNTER_GROUP_TP_PWR = 0x1D # type: ignore
+KGSL_PERFCOUNTER_GROUP_RB_PWR = 0x1E # type: ignore
+KGSL_PERFCOUNTER_GROUP_CCU_PWR = 0x1F # type: ignore
+KGSL_PERFCOUNTER_GROUP_UCHE_PWR = 0x20 # type: ignore
+KGSL_PERFCOUNTER_GROUP_CP_PWR = 0x21 # type: ignore
+KGSL_PERFCOUNTER_GROUP_GPMU_PWR = 0x22 # type: ignore
+KGSL_PERFCOUNTER_GROUP_ALWAYSON_PWR = 0x23 # type: ignore
+KGSL_PERFCOUNTER_GROUP_MAX = 0x24 # type: ignore
+KGSL_PERFCOUNTER_NOT_USED = 0xFFFFFFFF # type: ignore
+KGSL_PERFCOUNTER_BROKEN = 0xFFFFFFFE # type: ignore
+KGSL_IOC_TYPE = 0x09 # type: ignore
+IOCTL_KGSL_DEVICE_GETPROPERTY = _IOWR(KGSL_IOC_TYPE, 0x2, struct_kgsl_device_getproperty) # type: ignore
+IOCTL_KGSL_DEVICE_WAITTIMESTAMP = _IOW(KGSL_IOC_TYPE, 0x6, struct_kgsl_device_waittimestamp) # type: ignore
+IOCTL_KGSL_DEVICE_WAITTIMESTAMP_CTXTID = _IOW(KGSL_IOC_TYPE, 0x7, struct_kgsl_device_waittimestamp_ctxtid) # type: ignore
+IOCTL_KGSL_RINGBUFFER_ISSUEIBCMDS = _IOWR(KGSL_IOC_TYPE, 0x10, struct_kgsl_ringbuffer_issueibcmds) # type: ignore
+IOCTL_KGSL_CMDSTREAM_READTIMESTAMP_OLD = _IOR(KGSL_IOC_TYPE, 0x11, struct_kgsl_cmdstream_readtimestamp) # type: ignore
+IOCTL_KGSL_CMDSTREAM_READTIMESTAMP = _IOWR(KGSL_IOC_TYPE, 0x11, struct_kgsl_cmdstream_readtimestamp) # type: ignore
+IOCTL_KGSL_CMDSTREAM_FREEMEMONTIMESTAMP = _IOW(KGSL_IOC_TYPE, 0x12, struct_kgsl_cmdstream_freememontimestamp) # type: ignore
+IOCTL_KGSL_CMDSTREAM_FREEMEMONTIMESTAMP_OLD = _IOR(KGSL_IOC_TYPE, 0x12, struct_kgsl_cmdstream_freememontimestamp) # type: ignore
+IOCTL_KGSL_DRAWCTXT_CREATE = _IOWR(KGSL_IOC_TYPE, 0x13, struct_kgsl_drawctxt_create) # type: ignore
+IOCTL_KGSL_DRAWCTXT_DESTROY = _IOW(KGSL_IOC_TYPE, 0x14, struct_kgsl_drawctxt_destroy) # type: ignore
+IOCTL_KGSL_MAP_USER_MEM = _IOWR(KGSL_IOC_TYPE, 0x15, struct_kgsl_map_user_mem) # type: ignore
+IOCTL_KGSL_CMDSTREAM_READTIMESTAMP_CTXTID = _IOWR(KGSL_IOC_TYPE, 0x16, struct_kgsl_cmdstream_readtimestamp_ctxtid) # type: ignore
+IOCTL_KGSL_CMDSTREAM_FREEMEMONTIMESTAMP_CTXTID = _IOW(KGSL_IOC_TYPE, 0x17, struct_kgsl_cmdstream_freememontimestamp_ctxtid) # type: ignore
+IOCTL_KGSL_SHAREDMEM_FROM_PMEM = _IOWR(KGSL_IOC_TYPE, 0x20, struct_kgsl_sharedmem_from_pmem) # type: ignore
+IOCTL_KGSL_SHAREDMEM_FREE = _IOW(KGSL_IOC_TYPE, 0x21, struct_kgsl_sharedmem_free) # type: ignore
+IOCTL_KGSL_CFF_USER_EVENT = _IOW(KGSL_IOC_TYPE, 0x31, struct_kgsl_cff_user_event) # type: ignore
+IOCTL_KGSL_DRAWCTXT_BIND_GMEM_SHADOW = _IOW(KGSL_IOC_TYPE, 0x22, struct_kgsl_bind_gmem_shadow) # type: ignore
+IOCTL_KGSL_SHAREDMEM_FROM_VMALLOC = _IOWR(KGSL_IOC_TYPE, 0x23, struct_kgsl_sharedmem_from_vmalloc) # type: ignore
+IOCTL_KGSL_SHAREDMEM_FLUSH_CACHE = _IOW(KGSL_IOC_TYPE, 0x24, struct_kgsl_sharedmem_free) # type: ignore
+IOCTL_KGSL_DRAWCTXT_SET_BIN_BASE_OFFSET = _IOW(KGSL_IOC_TYPE, 0x25, struct_kgsl_drawctxt_set_bin_base_offset) # type: ignore
+IOCTL_KGSL_CMDWINDOW_WRITE = _IOW(KGSL_IOC_TYPE, 0x2e, struct_kgsl_cmdwindow_write) # type: ignore
+IOCTL_KGSL_GPUMEM_ALLOC = _IOWR(KGSL_IOC_TYPE, 0x2f, struct_kgsl_gpumem_alloc) # type: ignore
+IOCTL_KGSL_CFF_SYNCMEM = _IOW(KGSL_IOC_TYPE, 0x30, struct_kgsl_cff_syncmem) # type: ignore
+IOCTL_KGSL_TIMESTAMP_EVENT_OLD = _IOW(KGSL_IOC_TYPE, 0x31, struct_kgsl_timestamp_event) # type: ignore
+KGSL_TIMESTAMP_EVENT_GENLOCK = 1 # type: ignore
+KGSL_TIMESTAMP_EVENT_FENCE = 2 # type: ignore
+IOCTL_KGSL_SETPROPERTY = _IOW(KGSL_IOC_TYPE, 0x32, struct_kgsl_device_getproperty) # type: ignore
+IOCTL_KGSL_TIMESTAMP_EVENT = _IOWR(KGSL_IOC_TYPE, 0x33, struct_kgsl_timestamp_event) # type: ignore
+IOCTL_KGSL_GPUMEM_ALLOC_ID = _IOWR(KGSL_IOC_TYPE, 0x34, struct_kgsl_gpumem_alloc_id) # type: ignore
+IOCTL_KGSL_GPUMEM_FREE_ID = _IOWR(KGSL_IOC_TYPE, 0x35, struct_kgsl_gpumem_free_id) # type: ignore
+IOCTL_KGSL_GPUMEM_GET_INFO = _IOWR(KGSL_IOC_TYPE, 0x36, struct_kgsl_gpumem_get_info) # type: ignore
+KGSL_GPUMEM_CACHE_CLEAN = (1 << 0) # type: ignore
+KGSL_GPUMEM_CACHE_TO_GPU = KGSL_GPUMEM_CACHE_CLEAN # type: ignore
+KGSL_GPUMEM_CACHE_INV = (1 << 1) # type: ignore
+KGSL_GPUMEM_CACHE_FROM_GPU = KGSL_GPUMEM_CACHE_INV # type: ignore
+KGSL_GPUMEM_CACHE_FLUSH = (KGSL_GPUMEM_CACHE_CLEAN | KGSL_GPUMEM_CACHE_INV) # type: ignore
+KGSL_GPUMEM_CACHE_RANGE = (1 << 31) # type: ignore
+IOCTL_KGSL_GPUMEM_SYNC_CACHE = _IOW(KGSL_IOC_TYPE, 0x37, struct_kgsl_gpumem_sync_cache) # type: ignore
+IOCTL_KGSL_PERFCOUNTER_GET = _IOWR(KGSL_IOC_TYPE, 0x38, struct_kgsl_perfcounter_get) # type: ignore
+IOCTL_KGSL_PERFCOUNTER_PUT = _IOW(KGSL_IOC_TYPE, 0x39, struct_kgsl_perfcounter_put) # type: ignore
+IOCTL_KGSL_PERFCOUNTER_QUERY = _IOWR(KGSL_IOC_TYPE, 0x3A, struct_kgsl_perfcounter_query) # type: ignore
+IOCTL_KGSL_PERFCOUNTER_READ = _IOWR(KGSL_IOC_TYPE, 0x3B, struct_kgsl_perfcounter_read) # type: ignore
+IOCTL_KGSL_GPUMEM_SYNC_CACHE_BULK = _IOWR(KGSL_IOC_TYPE, 0x3C, struct_kgsl_gpumem_sync_cache_bulk) # type: ignore
+KGSL_IBDESC_MEMLIST = 0x1 # type: ignore
+KGSL_IBDESC_PROFILING_BUFFER = 0x2 # type: ignore
+IOCTL_KGSL_SUBMIT_COMMANDS = _IOWR(KGSL_IOC_TYPE, 0x3D, struct_kgsl_submit_commands) # type: ignore
+KGSL_CONSTRAINT_NONE = 0 # type: ignore
+KGSL_CONSTRAINT_PWRLEVEL = 1 # type: ignore
+KGSL_CONSTRAINT_PWR_MIN = 0 # type: ignore
+KGSL_CONSTRAINT_PWR_MAX = 1 # type: ignore
+IOCTL_KGSL_SYNCSOURCE_CREATE = _IOWR(KGSL_IOC_TYPE, 0x40, struct_kgsl_syncsource_create) # type: ignore
+IOCTL_KGSL_SYNCSOURCE_DESTROY = _IOWR(KGSL_IOC_TYPE, 0x41, struct_kgsl_syncsource_destroy) # type: ignore
+IOCTL_KGSL_SYNCSOURCE_CREATE_FENCE = _IOWR(KGSL_IOC_TYPE, 0x42, struct_kgsl_syncsource_create_fence) # type: ignore
+IOCTL_KGSL_SYNCSOURCE_SIGNAL_FENCE = _IOWR(KGSL_IOC_TYPE, 0x43, struct_kgsl_syncsource_signal_fence) # type: ignore
+IOCTL_KGSL_CFF_SYNC_GPUOBJ = _IOW(KGSL_IOC_TYPE, 0x44, struct_kgsl_cff_sync_gpuobj) # type: ignore
+KGSL_GPUOBJ_ALLOC_METADATA_MAX = 64 # type: ignore
+IOCTL_KGSL_GPUOBJ_ALLOC = _IOWR(KGSL_IOC_TYPE, 0x45, struct_kgsl_gpuobj_alloc) # type: ignore
+KGSL_GPUOBJ_FREE_ON_EVENT = 1 # type: ignore
+KGSL_GPU_EVENT_TIMESTAMP = 1 # type: ignore
+KGSL_GPU_EVENT_FENCE = 2 # type: ignore
+IOCTL_KGSL_GPUOBJ_FREE = _IOW(KGSL_IOC_TYPE, 0x46, struct_kgsl_gpuobj_free) # type: ignore
+IOCTL_KGSL_GPUOBJ_INFO = _IOWR(KGSL_IOC_TYPE, 0x47, struct_kgsl_gpuobj_info) # type: ignore
+IOCTL_KGSL_GPUOBJ_IMPORT = _IOWR(KGSL_IOC_TYPE, 0x48, struct_kgsl_gpuobj_import) # type: ignore
+IOCTL_KGSL_GPUOBJ_SYNC = _IOW(KGSL_IOC_TYPE, 0x49, struct_kgsl_gpuobj_sync) # type: ignore
+IOCTL_KGSL_GPU_COMMAND = _IOWR(KGSL_IOC_TYPE, 0x4A, struct_kgsl_gpu_command) # type: ignore
+IOCTL_KGSL_PREEMPTIONCOUNTER_QUERY = _IOWR(KGSL_IOC_TYPE, 0x4B, struct_kgsl_preemption_counters_query) # type: ignore
+KGSL_GPUOBJ_SET_INFO_METADATA = (1 << 0) # type: ignore
+KGSL_GPUOBJ_SET_INFO_TYPE = (1 << 1) # type: ignore
+IOCTL_KGSL_GPUOBJ_SET_INFO = _IOW(KGSL_IOC_TYPE, 0x4C, struct_kgsl_gpuobj_set_info) # type: ignore

tinygrad/runtime/autogen/libclang.py

@@ -4,8 +4,8 @@ import ctypes
 from typing import Literal, TypeAlias
 from tinygrad.runtime.support.c import _IO, _IOW, _IOR, _IOWR
 from tinygrad.runtime.support import c
-from tinygrad.helpers import WIN, OSX
-dll = c.DLL('libclang', 'C:\\Program Files\\LLVM\\bin\\libclang.dll' if WIN else ['/opt/homebrew/opt/llvm@21/lib/libclang.dylib', '/opt/homebrew/opt/llvm@20/lib/libclang.dylib', '/opt/homebrew/opt/llvm@19/lib/libclang.dylib', '/opt/homebrew/opt/llvm@18/lib/libclang.dylib', '/opt/homebrew/opt/llvm@17/lib/libclang.dylib', '/opt/homebrew/opt/llvm@16/lib/libclang.dylib', '/opt/homebrew/opt/llvm@15/lib/libclang.dylib', '/opt/homebrew/opt/llvm@14/lib/libclang.dylib'] if OSX else ['clang', 'clang-21', 'clang-20', 'clang-19', 'clang-18', 'clang-17', 'clang-16', 'clang-15', 'clang-14'])
+from tinygrad.helpers import OSX
+dll = c.DLL('libclang', '/opt/homebrew/opt/llvm@20/lib/libclang.dylib' if OSX else ['clang-20', 'clang'])
 CXIndex: TypeAlias = ctypes.c_void_p
 class struct_CXTargetInfoImpl(c.Struct): pass
 CXTargetInfo: TypeAlias = c.POINTER[struct_CXTargetInfoImpl]
@@ -1012,8 +1012,8 @@ def clang_getFileUniqueID(file:CXFile, outID:c.POINTER[CXFileUniqueID]) -> int:
 def clang_File_isEqual(file1:CXFile, file2:CXFile) -> int: ...
 @dll.bind(CXString, CXFile)
 def clang_File_tryGetRealPathName(file:CXFile) -> CXString: ...
-CINDEX_VERSION_MAJOR = 0
-CINDEX_VERSION_MINOR = 64
+CINDEX_VERSION_MAJOR = 0 # type: ignore
+CINDEX_VERSION_MINOR = 64 # type: ignore
 CINDEX_VERSION_ENCODE = lambda major,minor: (((major)*10000) + ((minor)*1)) # type: ignore
-CINDEX_VERSION = CINDEX_VERSION_ENCODE(CINDEX_VERSION_MAJOR, CINDEX_VERSION_MINOR)
+CINDEX_VERSION = CINDEX_VERSION_ENCODE(CINDEX_VERSION_MAJOR, CINDEX_VERSION_MINOR) # type: ignore
 CINDEX_VERSION_STRINGIZE = lambda major,minor: CINDEX_VERSION_STRINGIZE_(major, minor) # type: ignore

tinygrad/runtime/autogen/libusb.py

@@ -463,28 +463,28 @@ def libusb_hotplug_get_user_data(ctx:c.POINTER[libusb_context], callback_handle:
 @dll.bind(ctypes.c_int32, c.POINTER[libusb_context], ctypes.c_uint32)
 def libusb_set_option(ctx:c.POINTER[libusb_context], option:ctypes.c_uint32) -> int: ...
 LIBUSB_DEPRECATED_FOR = lambda f: __attribute__ ((deprecated)) # type: ignore
-LIBUSB_API_VERSION = 0x0100010A
-LIBUSBX_API_VERSION = LIBUSB_API_VERSION
-LIBUSB_DT_DEVICE_SIZE = 18
-LIBUSB_DT_CONFIG_SIZE = 9
-LIBUSB_DT_INTERFACE_SIZE = 9
-LIBUSB_DT_ENDPOINT_SIZE = 7
-LIBUSB_DT_ENDPOINT_AUDIO_SIZE = 9
-LIBUSB_DT_HUB_NONVAR_SIZE = 7
-LIBUSB_DT_SS_ENDPOINT_COMPANION_SIZE = 6
-LIBUSB_DT_BOS_SIZE = 5
-LIBUSB_DT_DEVICE_CAPABILITY_SIZE = 3
-LIBUSB_BT_USB_2_0_EXTENSION_SIZE = 7
-LIBUSB_BT_SS_USB_DEVICE_CAPABILITY_SIZE = 10
-LIBUSB_BT_CONTAINER_ID_SIZE = 20
-LIBUSB_BT_PLATFORM_DESCRIPTOR_MIN_SIZE = 20
-LIBUSB_DT_BOS_MAX_SIZE = (LIBUSB_DT_BOS_SIZE + LIBUSB_BT_USB_2_0_EXTENSION_SIZE + LIBUSB_BT_SS_USB_DEVICE_CAPABILITY_SIZE + LIBUSB_BT_CONTAINER_ID_SIZE)
-LIBUSB_ENDPOINT_ADDRESS_MASK = 0x0f
-LIBUSB_ENDPOINT_DIR_MASK = 0x80
-LIBUSB_TRANSFER_TYPE_MASK = 0x03
-LIBUSB_ISO_SYNC_TYPE_MASK = 0x0c
-LIBUSB_ISO_USAGE_TYPE_MASK = 0x30
-LIBUSB_ERROR_COUNT = 14
-LIBUSB_OPTION_WEAK_AUTHORITY = LIBUSB_OPTION_NO_DEVICE_DISCOVERY
-LIBUSB_HOTPLUG_NO_FLAGS = 0
-LIBUSB_HOTPLUG_MATCH_ANY = -1
+LIBUSB_API_VERSION = 0x0100010A # type: ignore
+LIBUSBX_API_VERSION = LIBUSB_API_VERSION # type: ignore
+LIBUSB_DT_DEVICE_SIZE = 18 # type: ignore
+LIBUSB_DT_CONFIG_SIZE = 9 # type: ignore
+LIBUSB_DT_INTERFACE_SIZE = 9 # type: ignore
+LIBUSB_DT_ENDPOINT_SIZE = 7 # type: ignore
+LIBUSB_DT_ENDPOINT_AUDIO_SIZE = 9 # type: ignore
+LIBUSB_DT_HUB_NONVAR_SIZE = 7 # type: ignore
+LIBUSB_DT_SS_ENDPOINT_COMPANION_SIZE = 6 # type: ignore
+LIBUSB_DT_BOS_SIZE = 5 # type: ignore
+LIBUSB_DT_DEVICE_CAPABILITY_SIZE = 3 # type: ignore
+LIBUSB_BT_USB_2_0_EXTENSION_SIZE = 7 # type: ignore
+LIBUSB_BT_SS_USB_DEVICE_CAPABILITY_SIZE = 10 # type: ignore
+LIBUSB_BT_CONTAINER_ID_SIZE = 20 # type: ignore
+LIBUSB_BT_PLATFORM_DESCRIPTOR_MIN_SIZE = 20 # type: ignore
+LIBUSB_DT_BOS_MAX_SIZE = (LIBUSB_DT_BOS_SIZE + LIBUSB_BT_USB_2_0_EXTENSION_SIZE + LIBUSB_BT_SS_USB_DEVICE_CAPABILITY_SIZE + LIBUSB_BT_CONTAINER_ID_SIZE) # type: ignore
+LIBUSB_ENDPOINT_ADDRESS_MASK = 0x0f # type: ignore
+LIBUSB_ENDPOINT_DIR_MASK = 0x80 # type: ignore
+LIBUSB_TRANSFER_TYPE_MASK = 0x03 # type: ignore
+LIBUSB_ISO_SYNC_TYPE_MASK = 0x0c # type: ignore
+LIBUSB_ISO_USAGE_TYPE_MASK = 0x30 # type: ignore
+LIBUSB_ERROR_COUNT = 14 # type: ignore
+LIBUSB_OPTION_WEAK_AUTHORITY = LIBUSB_OPTION_NO_DEVICE_DISCOVERY # type: ignore
+LIBUSB_HOTPLUG_NO_FLAGS = 0 # type: ignore
+LIBUSB_HOTPLUG_MATCH_ANY = -1 # type: ignore

tinygrad/runtime/autogen/llvm.py

@@ -5,7 +5,7 @@ from typing import Literal, TypeAlias
 from tinygrad.runtime.support.c import _IO, _IOW, _IOR, _IOWR
 from tinygrad.runtime.support import c
 from tinygrad.helpers import WIN, OSX
-dll = c.DLL('llvm', 'C:\\Program Files\\LLVM\\bin\\LLVM-C.dll' if WIN else ['/opt/homebrew/opt/llvm@21/lib/libLLVM.dylib', '/opt/homebrew/opt/llvm@20/lib/libLLVM.dylib', '/opt/homebrew/opt/llvm@19/lib/libLLVM.dylib', '/opt/homebrew/opt/llvm@18/lib/libLLVM.dylib', '/opt/homebrew/opt/llvm@17/lib/libLLVM.dylib', '/opt/homebrew/opt/llvm@16/lib/libLLVM.dylib', '/opt/homebrew/opt/llvm@15/lib/libLLVM.dylib', '/opt/homebrew/opt/llvm@14/lib/libLLVM.dylib'] if OSX else ['LLVM', 'LLVM-21', 'LLVM-20', 'LLVM-19', 'LLVM-18', 'LLVM-17', 'LLVM-16', 'LLVM-15', 'LLVM-14'])
+dll = c.DLL('llvm', 'C:\\Program Files\\LLVM\\bin\\LLVM-C.dll' if WIN else '/opt/homebrew/opt/llvm@20/lib/libLLVM.dylib' if OSX else ['LLVM', 'LLVM-21', 'LLVM-20', 'LLVM-19', 'LLVM-18', 'LLVM-17', 'LLVM-16', 'LLVM-15', 'LLVM-14'])
 intmax_t: TypeAlias = ctypes.c_int64
 @dll.bind(intmax_t, intmax_t)
 def imaxabs(__n:intmax_t) -> intmax_t: ...
@@ -3148,44 +3148,44 @@ def thinlto_codegen_set_cache_size_bytes(cg:thinlto_code_gen_t, max_size_bytes:i
 def thinlto_codegen_set_cache_size_megabytes(cg:thinlto_code_gen_t, max_size_megabytes:int) -> None: ...
 @dll.bind(None, thinlto_code_gen_t, ctypes.c_uint32)
 def thinlto_codegen_set_cache_size_files(cg:thinlto_code_gen_t, max_size_files:int) -> None: ...
-LLVMDisassembler_Option_UseMarkup = 1
-LLVMDisassembler_Option_PrintImmHex = 2
-LLVMDisassembler_Option_AsmPrinterVariant = 4
-LLVMDisassembler_Option_SetInstrComments = 8
-LLVMDisassembler_Option_PrintLatency = 16
-LLVMDisassembler_Option_Color = 32
-LLVMDisassembler_VariantKind_None = 0
-LLVMDisassembler_VariantKind_ARM_HI16 = 1
-LLVMDisassembler_VariantKind_ARM_LO16 = 2
-LLVMDisassembler_VariantKind_ARM64_PAGE = 1
-LLVMDisassembler_VariantKind_ARM64_PAGEOFF = 2
-LLVMDisassembler_VariantKind_ARM64_GOTPAGE = 3
-LLVMDisassembler_VariantKind_ARM64_GOTPAGEOFF = 4
-LLVMDisassembler_VariantKind_ARM64_TLVP = 5
-LLVMDisassembler_VariantKind_ARM64_TLVOFF = 6
-LLVMDisassembler_ReferenceType_InOut_None = 0
-LLVMDisassembler_ReferenceType_In_Branch = 1
-LLVMDisassembler_ReferenceType_In_PCrel_Load = 2
-LLVMDisassembler_ReferenceType_In_ARM64_ADRP = 0x100000001
-LLVMDisassembler_ReferenceType_In_ARM64_ADDXri = 0x100000002
-LLVMDisassembler_ReferenceType_In_ARM64_LDRXui = 0x100000003
-LLVMDisassembler_ReferenceType_In_ARM64_LDRXl = 0x100000004
-LLVMDisassembler_ReferenceType_In_ARM64_ADR = 0x100000005
-LLVMDisassembler_ReferenceType_Out_SymbolStub = 1
-LLVMDisassembler_ReferenceType_Out_LitPool_SymAddr = 2
-LLVMDisassembler_ReferenceType_Out_LitPool_CstrAddr = 3
-LLVMDisassembler_ReferenceType_Out_Objc_CFString_Ref = 4
-LLVMDisassembler_ReferenceType_Out_Objc_Message = 5
-LLVMDisassembler_ReferenceType_Out_Objc_Message_Ref = 6
-LLVMDisassembler_ReferenceType_Out_Objc_Selector_Ref = 7
-LLVMDisassembler_ReferenceType_Out_Objc_Class_Ref = 8
-LLVMDisassembler_ReferenceType_DeMangled_Name = 9
-LLVMErrorSuccess = 0
-REMARKS_API_VERSION = 1
-LLVM_BLAKE3_VERSION_STRING = "1.3.1"
-LLVM_BLAKE3_KEY_LEN = 32
-LLVM_BLAKE3_OUT_LEN = 32
-LLVM_BLAKE3_BLOCK_LEN = 64
-LLVM_BLAKE3_CHUNK_LEN = 1024
-LLVM_BLAKE3_MAX_DEPTH = 54
-LTO_API_VERSION = 29
+LLVMDisassembler_Option_UseMarkup = 1 # type: ignore
+LLVMDisassembler_Option_PrintImmHex = 2 # type: ignore
+LLVMDisassembler_Option_AsmPrinterVariant = 4 # type: ignore
+LLVMDisassembler_Option_SetInstrComments = 8 # type: ignore
+LLVMDisassembler_Option_PrintLatency = 16 # type: ignore
+LLVMDisassembler_Option_Color = 32 # type: ignore
+LLVMDisassembler_VariantKind_None = 0 # type: ignore
+LLVMDisassembler_VariantKind_ARM_HI16 = 1 # type: ignore
+LLVMDisassembler_VariantKind_ARM_LO16 = 2 # type: ignore
+LLVMDisassembler_VariantKind_ARM64_PAGE = 1 # type: ignore
+LLVMDisassembler_VariantKind_ARM64_PAGEOFF = 2 # type: ignore
+LLVMDisassembler_VariantKind_ARM64_GOTPAGE = 3 # type: ignore
+LLVMDisassembler_VariantKind_ARM64_GOTPAGEOFF = 4 # type: ignore
+LLVMDisassembler_VariantKind_ARM64_TLVP = 5 # type: ignore
+LLVMDisassembler_VariantKind_ARM64_TLVOFF = 6 # type: ignore
+LLVMDisassembler_ReferenceType_InOut_None = 0 # type: ignore
+LLVMDisassembler_ReferenceType_In_Branch = 1 # type: ignore
+LLVMDisassembler_ReferenceType_In_PCrel_Load = 2 # type: ignore
+LLVMDisassembler_ReferenceType_In_ARM64_ADRP = 0x100000001 # type: ignore
+LLVMDisassembler_ReferenceType_In_ARM64_ADDXri = 0x100000002 # type: ignore
+LLVMDisassembler_ReferenceType_In_ARM64_LDRXui = 0x100000003 # type: ignore
+LLVMDisassembler_ReferenceType_In_ARM64_LDRXl = 0x100000004 # type: ignore
+LLVMDisassembler_ReferenceType_In_ARM64_ADR = 0x100000005 # type: ignore
+LLVMDisassembler_ReferenceType_Out_SymbolStub = 1 # type: ignore
+LLVMDisassembler_ReferenceType_Out_LitPool_SymAddr = 2 # type: ignore
+LLVMDisassembler_ReferenceType_Out_LitPool_CstrAddr = 3 # type: ignore
+LLVMDisassembler_ReferenceType_Out_Objc_CFString_Ref = 4 # type: ignore
+LLVMDisassembler_ReferenceType_Out_Objc_Message = 5 # type: ignore
+LLVMDisassembler_ReferenceType_Out_Objc_Message_Ref = 6 # type: ignore
+LLVMDisassembler_ReferenceType_Out_Objc_Selector_Ref = 7 # type: ignore
+LLVMDisassembler_ReferenceType_Out_Objc_Class_Ref = 8 # type: ignore
+LLVMDisassembler_ReferenceType_DeMangled_Name = 9 # type: ignore
+LLVMErrorSuccess = 0 # type: ignore
+REMARKS_API_VERSION = 1 # type: ignore
+LLVM_BLAKE3_VERSION_STRING = "1.3.1" # type: ignore
+LLVM_BLAKE3_KEY_LEN = 32 # type: ignore
+LLVM_BLAKE3_OUT_LEN = 32 # type: ignore
+LLVM_BLAKE3_BLOCK_LEN = 64 # type: ignore
+LLVM_BLAKE3_CHUNK_LEN = 1024 # type: ignore
+LLVM_BLAKE3_MAX_DEPTH = 54 # type: ignore
+LTO_API_VERSION = 29 # type: ignore

tinygrad/runtime/autogen/llvm_qcom.py

@@ -93,12 +93,12 @@ cl_lib_img_desc.register_fields([('_unk0', c.Array[ctypes.c_char, Literal[196]],
 def cl_compiler_free_handle(handle:c.POINTER[cl_handle]) -> None: ...
 @dll.bind(None, ctypes.c_void_p)
 def cl_compiler_free_assembly(ptr:ctypes.c_void_p) -> None: ...
-CL_MODE_32BIT = 0
-CL_MODE_64BIT = 1
-CL_SRC_STR = 0
-CL_SRC_BLOB = 1
-CL_LIB_PROGRAM = 0
-CL_LIB_CONSTS = 6
-CL_LIB_IMAGE = 7
-CL_LIB_CODE = 10
-CL_LIB_IMAGE_DESC = 11
+CL_MODE_32BIT = 0 # type: ignore
+CL_MODE_64BIT = 1 # type: ignore
+CL_SRC_STR = 0 # type: ignore
+CL_SRC_BLOB = 1 # type: ignore
+CL_LIB_PROGRAM = 0 # type: ignore
+CL_LIB_CONSTS = 6 # type: ignore
+CL_LIB_IMAGE = 7 # type: ignore
+CL_LIB_CODE = 10 # type: ignore
+CL_LIB_IMAGE_DESC = 11 # type: ignore

tinygrad/runtime/autogen/metal.py

@@ -1457,9 +1457,9 @@ MTLCommandEncoder._methods_ = [
   ('label', NSString, []),
   ('setLabel:', None, [NSString]),
 ]
-MTLResourceCPUCacheModeShift = 0
-MTLResourceCPUCacheModeMask = (0xf << MTLResourceCPUCacheModeShift)
-MTLResourceStorageModeShift = 4
-MTLResourceStorageModeMask = (0xf << MTLResourceStorageModeShift)
-MTLResourceHazardTrackingModeShift = 8
-MTLResourceHazardTrackingModeMask = (0x3 << MTLResourceHazardTrackingModeShift)
+MTLResourceCPUCacheModeShift = 0 # type: ignore
+MTLResourceCPUCacheModeMask = (0xf << MTLResourceCPUCacheModeShift) # type: ignore
+MTLResourceStorageModeShift = 4 # type: ignore
+MTLResourceStorageModeMask = (0xf << MTLResourceStorageModeShift) # type: ignore
+MTLResourceHazardTrackingModeShift = 8 # type: ignore
+MTLResourceHazardTrackingModeMask = (0x3 << MTLResourceHazardTrackingModeShift) # type: ignore

tinygrad/runtime/autogen/mlx5.py

@@ -10424,101 +10424,101 @@ class struct_mlx5_ifc_psp_gen_spi_out_bits(c.Struct):
   reserved_at_60: c.Array[ctypes.c_ubyte, Literal[32]]
   key_spi: c.Array[struct_mlx5_ifc_key_spi_bits, Literal[0]]
 struct_mlx5_ifc_psp_gen_spi_out_bits.register_fields([('status', c.Array[ctypes.c_ubyte, Literal[8]], 0), ('reserved_at_8', c.Array[ctypes.c_ubyte, Literal[24]], 8), ('syndrome', c.Array[ctypes.c_ubyte, Literal[32]], 32), ('reserved_at_40', c.Array[ctypes.c_ubyte, Literal[16]], 64), ('num_of_spi', c.Array[ctypes.c_ubyte, Literal[16]], 80), ('reserved_at_60', c.Array[ctypes.c_ubyte, Literal[32]], 96), ('key_spi', c.Array[struct_mlx5_ifc_key_spi_bits, Literal[0]], 128)])
-MLX5_CMD_OP_QUERY_HCA_CAP = 0x100
-MLX5_CMD_OP_QUERY_ADAPTER = 0x101
-MLX5_CMD_OP_INIT_HCA = 0x102
-MLX5_CMD_OP_TEARDOWN_HCA = 0x103
-MLX5_CMD_OP_ENABLE_HCA = 0x104
-MLX5_CMD_OP_DISABLE_HCA = 0x105
-MLX5_CMD_OP_QUERY_PAGES = 0x107
-MLX5_CMD_OP_MANAGE_PAGES = 0x108
-MLX5_CMD_OP_SET_HCA_CAP = 0x109
-MLX5_CMD_OP_QUERY_ISSI = 0x10a
-MLX5_CMD_OP_SET_ISSI = 0x10b
-MLX5_CMD_OP_SET_DRIVER_VERSION = 0x10d
-MLX5_CMD_OP_CREATE_MKEY = 0x200
-MLX5_CMD_OP_QUERY_SPECIAL_CONTEXTS = 0x203
-MLX5_CMD_OP_CREATE_EQ = 0x301
-MLX5_CMD_OP_DESTROY_EQ = 0x302
-MLX5_CMD_OP_CREATE_CQ = 0x400
-MLX5_CMD_OP_DESTROY_CQ = 0x401
-MLX5_CMD_OP_CREATE_QP = 0x500
-MLX5_CMD_OP_DESTROY_QP = 0x501
-MLX5_CMD_OP_RST2INIT_QP = 0x502
-MLX5_CMD_OP_INIT2RTR_QP = 0x503
-MLX5_CMD_OP_RTR2RTS_QP = 0x504
-MLX5_CMD_OP_QUERY_NIC_VPORT_CONTEXT = 0x754
-MLX5_CMD_OP_MODIFY_NIC_VPORT_CONTEXT = 0x755
-MLX5_CMD_OP_SET_ROCE_ADDRESS = 0x761
-MLX5_CMD_OP_ALLOC_PD = 0x800
-MLX5_CMD_OP_ALLOC_UAR = 0x802
-MLX5_CMD_OP_ACCESS_REG = 0x805
-MLX5_CMD_OP_ALLOC_TRANSPORT_DOMAIN = 0x816
-MLX5_CMD_STAT_OK = 0x0
-MLX5_CMD_STAT_INT_ERR = 0x1
-MLX5_CMD_STAT_BAD_OP_ERR = 0x2
-MLX5_CMD_STAT_BAD_PARAM_ERR = 0x3
-MLX5_CMD_STAT_BAD_SYS_STATE_ERR = 0x4
-MLX5_CMD_STAT_BAD_RES_ERR = 0x5
-MLX5_CMD_STAT_RES_BUSY = 0x6
-MLX5_CMD_STAT_LIM_ERR = 0x8
-MLX5_CMD_STAT_BAD_RES_STATE_ERR = 0x9
-MLX5_CMD_STAT_NO_RES_ERR = 0xf
-MLX5_CMD_STAT_BAD_INP_LEN_ERR = 0x50
-MLX5_CMD_STAT_BAD_OUTP_LEN_ERR = 0x51
-MLX5_CAP_GENERAL = 0x0
-MLX5_CAP_ODP = 0x2
-MLX5_CAP_ATOMIC = 0x3
-MLX5_CAP_ROCE = 0x4
-HCA_CAP_OPMOD_GET_MAX = 0
-HCA_CAP_OPMOD_GET_CUR = 1
-MLX5_PAGES_GIVE = 1
-MLX5_PAGES_TAKE = 2
-MLX5_BOOT_PAGES = 1
-MLX5_INIT_PAGES = 2
-MLX5_REG_HOST_ENDIANNESS = 0x7004
-MLX5_REG_DTOR = 0xC00E
-MLX5_PCI_CMD_XPORT = 0x07
-MLX5_CMD_DATA_BLOCK_SIZE = 512
-CMD_OWNER_HW = 0x01
-CAP_GEN_ABS_NATIVE_PORT_NUM = 0x007
-CAP_GEN_HCA_CAP_2 = 0x020
-CAP_GEN_EVENT_ON_VHCA_STATE_ALLOCATED = 0x023
-CAP_GEN_EVENT_ON_VHCA_STATE_ACTIVE = 0x024
-CAP_GEN_EVENT_ON_VHCA_STATE_IN_USE = 0x025
-CAP_GEN_EVENT_ON_VHCA_STATE_TEARDOWN_REQUEST = 0x026
-CAP_GEN_LOG_MAX_QP = 0x09B
-CAP_GEN_LOG_MAX_CQ = 0x0DB
-CAP_GEN_RELEASE_ALL_PAGES = 0x145
-CAP_GEN_CACHE_LINE_128BYTE = 0x164
-CAP_GEN_NUM_PORTS = 0x1B8
-CAP_GEN_PKEY_TABLE_SIZE = 0x190
-CAP_GEN_PCI_SYNC_FOR_FW_UPDATE_EVENT = 0x1F1
-CAP_GEN_CMDIF_CHECKSUM = 0x210
-CAP_GEN_DCT = 0x21A
-CAP_GEN_ROCE = 0x21D
-CAP_GEN_ATOMIC = 0x21E
-CAP_GEN_ODP = 0x227
-CAP_GEN_MKEY_BY_NAME = 0x266
-CAP_GEN_LOG_MAX_PD = 0x32B
-CAP_GEN_PCIE_RESET_USING_HOTRESET = 0x335
-CAP_GEN_PCI_SYNC_FOR_FW_UPDATE_WITH_DRIVER_UNLOAD = 0x336
-CAP_GEN_VHCA_STATE = 0x3EA
-CAP_GEN_ROCE_RW_SUPPORTED = 0x3A1
-CAP_GEN_LOG_MAX_CURRENT_UC_LIST = 0x3FB
-CAP_GEN_LOG_UAR_PAGE_SZ = 0x490
-CAP_GEN_NUM_VHCA_PORTS = 0x610
-CAP_GEN_SW_OWNER_ID = 0x61E
-CAP_GEN_NUM_TOTAL_DYNAMIC_VF_MSIX = 0x708
-MLX5_FC_BULK_SIZE_FACTOR = 128
+MLX5_CMD_OP_QUERY_HCA_CAP = 0x100 # type: ignore
+MLX5_CMD_OP_QUERY_ADAPTER = 0x101 # type: ignore
+MLX5_CMD_OP_INIT_HCA = 0x102 # type: ignore
+MLX5_CMD_OP_TEARDOWN_HCA = 0x103 # type: ignore
+MLX5_CMD_OP_ENABLE_HCA = 0x104 # type: ignore
+MLX5_CMD_OP_DISABLE_HCA = 0x105 # type: ignore
+MLX5_CMD_OP_QUERY_PAGES = 0x107 # type: ignore
+MLX5_CMD_OP_MANAGE_PAGES = 0x108 # type: ignore
+MLX5_CMD_OP_SET_HCA_CAP = 0x109 # type: ignore
+MLX5_CMD_OP_QUERY_ISSI = 0x10a # type: ignore
+MLX5_CMD_OP_SET_ISSI = 0x10b # type: ignore
+MLX5_CMD_OP_SET_DRIVER_VERSION = 0x10d # type: ignore
+MLX5_CMD_OP_CREATE_MKEY = 0x200 # type: ignore
+MLX5_CMD_OP_QUERY_SPECIAL_CONTEXTS = 0x203 # type: ignore
+MLX5_CMD_OP_CREATE_EQ = 0x301 # type: ignore
+MLX5_CMD_OP_DESTROY_EQ = 0x302 # type: ignore
+MLX5_CMD_OP_CREATE_CQ = 0x400 # type: ignore
+MLX5_CMD_OP_DESTROY_CQ = 0x401 # type: ignore
+MLX5_CMD_OP_CREATE_QP = 0x500 # type: ignore
+MLX5_CMD_OP_DESTROY_QP = 0x501 # type: ignore
+MLX5_CMD_OP_RST2INIT_QP = 0x502 # type: ignore
+MLX5_CMD_OP_INIT2RTR_QP = 0x503 # type: ignore
+MLX5_CMD_OP_RTR2RTS_QP = 0x504 # type: ignore
+MLX5_CMD_OP_QUERY_NIC_VPORT_CONTEXT = 0x754 # type: ignore
+MLX5_CMD_OP_MODIFY_NIC_VPORT_CONTEXT = 0x755 # type: ignore
+MLX5_CMD_OP_SET_ROCE_ADDRESS = 0x761 # type: ignore
+MLX5_CMD_OP_ALLOC_PD = 0x800 # type: ignore
+MLX5_CMD_OP_ALLOC_UAR = 0x802 # type: ignore
+MLX5_CMD_OP_ACCESS_REG = 0x805 # type: ignore
+MLX5_CMD_OP_ALLOC_TRANSPORT_DOMAIN = 0x816 # type: ignore
+MLX5_CMD_STAT_OK = 0x0 # type: ignore
+MLX5_CMD_STAT_INT_ERR = 0x1 # type: ignore
+MLX5_CMD_STAT_BAD_OP_ERR = 0x2 # type: ignore
+MLX5_CMD_STAT_BAD_PARAM_ERR = 0x3 # type: ignore
+MLX5_CMD_STAT_BAD_SYS_STATE_ERR = 0x4 # type: ignore
+MLX5_CMD_STAT_BAD_RES_ERR = 0x5 # type: ignore
+MLX5_CMD_STAT_RES_BUSY = 0x6 # type: ignore
+MLX5_CMD_STAT_LIM_ERR = 0x8 # type: ignore
+MLX5_CMD_STAT_BAD_RES_STATE_ERR = 0x9 # type: ignore
+MLX5_CMD_STAT_NO_RES_ERR = 0xf # type: ignore
+MLX5_CMD_STAT_BAD_INP_LEN_ERR = 0x50 # type: ignore
+MLX5_CMD_STAT_BAD_OUTP_LEN_ERR = 0x51 # type: ignore
+MLX5_CAP_GENERAL = 0x0 # type: ignore
+MLX5_CAP_ODP = 0x2 # type: ignore
+MLX5_CAP_ATOMIC = 0x3 # type: ignore
+MLX5_CAP_ROCE = 0x4 # type: ignore
+HCA_CAP_OPMOD_GET_MAX = 0 # type: ignore
+HCA_CAP_OPMOD_GET_CUR = 1 # type: ignore
+MLX5_PAGES_GIVE = 1 # type: ignore
+MLX5_PAGES_TAKE = 2 # type: ignore
+MLX5_BOOT_PAGES = 1 # type: ignore
+MLX5_INIT_PAGES = 2 # type: ignore
+MLX5_REG_HOST_ENDIANNESS = 0x7004 # type: ignore
+MLX5_REG_DTOR = 0xC00E # type: ignore
+MLX5_PCI_CMD_XPORT = 0x07 # type: ignore
+MLX5_CMD_DATA_BLOCK_SIZE = 512 # type: ignore
+CMD_OWNER_HW = 0x01 # type: ignore
+CAP_GEN_ABS_NATIVE_PORT_NUM = 0x007 # type: ignore
+CAP_GEN_HCA_CAP_2 = 0x020 # type: ignore
+CAP_GEN_EVENT_ON_VHCA_STATE_ALLOCATED = 0x023 # type: ignore
+CAP_GEN_EVENT_ON_VHCA_STATE_ACTIVE = 0x024 # type: ignore
+CAP_GEN_EVENT_ON_VHCA_STATE_IN_USE = 0x025 # type: ignore
+CAP_GEN_EVENT_ON_VHCA_STATE_TEARDOWN_REQUEST = 0x026 # type: ignore
+CAP_GEN_LOG_MAX_QP = 0x09B # type: ignore
+CAP_GEN_LOG_MAX_CQ = 0x0DB # type: ignore
+CAP_GEN_RELEASE_ALL_PAGES = 0x145 # type: ignore
+CAP_GEN_CACHE_LINE_128BYTE = 0x164 # type: ignore
+CAP_GEN_NUM_PORTS = 0x1B8 # type: ignore
+CAP_GEN_PKEY_TABLE_SIZE = 0x190 # type: ignore
+CAP_GEN_PCI_SYNC_FOR_FW_UPDATE_EVENT = 0x1F1 # type: ignore
+CAP_GEN_CMDIF_CHECKSUM = 0x210 # type: ignore
+CAP_GEN_DCT = 0x21A # type: ignore
+CAP_GEN_ROCE = 0x21D # type: ignore
+CAP_GEN_ATOMIC = 0x21E # type: ignore
+CAP_GEN_ODP = 0x227 # type: ignore
+CAP_GEN_MKEY_BY_NAME = 0x266 # type: ignore
+CAP_GEN_LOG_MAX_PD = 0x32B # type: ignore
+CAP_GEN_PCIE_RESET_USING_HOTRESET = 0x335 # type: ignore
+CAP_GEN_PCI_SYNC_FOR_FW_UPDATE_WITH_DRIVER_UNLOAD = 0x336 # type: ignore
+CAP_GEN_VHCA_STATE = 0x3EA # type: ignore
+CAP_GEN_ROCE_RW_SUPPORTED = 0x3A1 # type: ignore
+CAP_GEN_LOG_MAX_CURRENT_UC_LIST = 0x3FB # type: ignore
+CAP_GEN_LOG_UAR_PAGE_SZ = 0x490 # type: ignore
+CAP_GEN_NUM_VHCA_PORTS = 0x610 # type: ignore
+CAP_GEN_SW_OWNER_ID = 0x61E # type: ignore
+CAP_GEN_NUM_TOTAL_DYNAMIC_VF_MSIX = 0x708 # type: ignore
+MLX5_FC_BULK_SIZE_FACTOR = 128 # type: ignore
 MLX5_FC_BULK_NUM_FCS = lambda fc_enum: (MLX5_FC_BULK_SIZE_FACTOR * (fc_enum)) # type: ignore
-MLX5_FT_MAX_MULTIPATH_LEVEL = 63
-MLX5_CMD_SET_MONITOR_NUM_PPCNT_COUNTER_SET1 = (6)
-MLX5_CMD_SET_MONITOR_NUM_Q_COUNTERS_SET1 = (1)
-MLX5_CMD_SET_MONITOR_NUM_COUNTER = (MLX5_CMD_SET_MONITOR_NUM_PPCNT_COUNTER_SET1 + MLX5_CMD_SET_MONITOR_NUM_Q_COUNTERS_SET1)
-MLX5_IFC_DEFINER_FORMAT_OFFSET_UNUSED = 0x0
-MLX5_IFC_DEFINER_FORMAT_OFFSET_OUTER_ETH_PKT_LEN = 0x48
-MLX5_IFC_DEFINER_DW_SELECTORS_NUM = 9
-MLX5_IFC_DEFINER_BYTE_SELECTORS_NUM = 8
-MLX5_MACSEC_ASO_INC_SN = 0x2
-MLX5_MACSEC_ASO_REG_C_4_5 = 0x2
+MLX5_FT_MAX_MULTIPATH_LEVEL = 63 # type: ignore
+MLX5_CMD_SET_MONITOR_NUM_PPCNT_COUNTER_SET1 = (6) # type: ignore
+MLX5_CMD_SET_MONITOR_NUM_Q_COUNTERS_SET1 = (1) # type: ignore
+MLX5_CMD_SET_MONITOR_NUM_COUNTER = (MLX5_CMD_SET_MONITOR_NUM_PPCNT_COUNTER_SET1 + MLX5_CMD_SET_MONITOR_NUM_Q_COUNTERS_SET1) # type: ignore
+MLX5_IFC_DEFINER_FORMAT_OFFSET_UNUSED = 0x0 # type: ignore
+MLX5_IFC_DEFINER_FORMAT_OFFSET_OUTER_ETH_PKT_LEN = 0x48 # type: ignore
+MLX5_IFC_DEFINER_DW_SELECTORS_NUM = 9 # type: ignore
+MLX5_IFC_DEFINER_BYTE_SELECTORS_NUM = 8 # type: ignore
+MLX5_MACSEC_ASO_INC_SN = 0x2 # type: ignore
+MLX5_MACSEC_ASO_REG_C_4_5 = 0x2 # type: ignore

tinygrad/runtime/autogen/nv.py

@@ -4683,173 +4683,173 @@ class struct__NV_PCI_DATA_EXT_STRUCT(c.Struct):
 struct__NV_PCI_DATA_EXT_STRUCT.register_fields([('signature', NvU32, 0), ('nvPciDataExtRev', NvU16, 4), ('nvPciDataExtLen', NvU16, 6), ('subimageLen', NvU16, 8), ('privLastImage', NvU8, 10), ('flags', NvU8, 11)])
 NV_PCI_DATA_EXT_STRUCT: TypeAlias = struct__NV_PCI_DATA_EXT_STRUCT
 PNV_PCI_DATA_EXT_STRUCT: TypeAlias = c.POINTER[struct__NV_PCI_DATA_EXT_STRUCT]
-GSP_FW_WPR_META_VERIFIED = 0xa0a0a0a0a0a0a0a0
-GSP_FW_WPR_META_REVISION = 1
-GSP_FW_WPR_META_MAGIC = 0xdc3aae21371a60b3
-GSP_FW_WPR_HEAP_FREE_REGION_COUNT = 128
-GSP_FW_HEAP_FREE_LIST_MAGIC = 0x4845415046524545
-GSP_FW_SR_META_MAGIC = 0x8a3bb9e6c6c39d93
-GSP_FW_SR_META_REVISION = 2
-GSP_FW_SR_META_INTERNAL_SIZE = 128
-NVDM_TYPE_HULK = 0x11
-NVDM_TYPE_FIRMWARE_UPDATE = 0x12
-NVDM_TYPE_PRC = 0x13
-NVDM_TYPE_COT = 0x14
-NVDM_TYPE_FSP_RESPONSE = 0x15
-NVDM_TYPE_CAPS_QUERY = 0x16
-NVDM_TYPE_INFOROM = 0x17
-NVDM_TYPE_SMBPBI = 0x18
-NVDM_TYPE_ROMREAD = 0x1A
-NVDM_TYPE_UEFI_RM = 0x1C
-NVDM_TYPE_UEFI_XTL_DEBUG_INTR = 0x1D
-NVDM_TYPE_TNVL = 0x1F
-NVDM_TYPE_CLOCK_BOOST = 0x20
-NVDM_TYPE_FSP_GSP_COMM = 0x21
-MAX_GPC_COUNT = 32
-VGPU_MAX_REGOPS_PER_RPC = 100
-VGPU_RESERVED_HANDLE_BASE = 0xCAF3F000
-VGPU_RESERVED_HANDLE_RANGE = 0x1000
+GSP_FW_WPR_META_VERIFIED = 0xa0a0a0a0a0a0a0a0 # type: ignore
+GSP_FW_WPR_META_REVISION = 1 # type: ignore
+GSP_FW_WPR_META_MAGIC = 0xdc3aae21371a60b3 # type: ignore
+GSP_FW_WPR_HEAP_FREE_REGION_COUNT = 128 # type: ignore
+GSP_FW_HEAP_FREE_LIST_MAGIC = 0x4845415046524545 # type: ignore
+GSP_FW_SR_META_MAGIC = 0x8a3bb9e6c6c39d93 # type: ignore
+GSP_FW_SR_META_REVISION = 2 # type: ignore
+GSP_FW_SR_META_INTERNAL_SIZE = 128 # type: ignore
+NVDM_TYPE_HULK = 0x11 # type: ignore
+NVDM_TYPE_FIRMWARE_UPDATE = 0x12 # type: ignore
+NVDM_TYPE_PRC = 0x13 # type: ignore
+NVDM_TYPE_COT = 0x14 # type: ignore
+NVDM_TYPE_FSP_RESPONSE = 0x15 # type: ignore
+NVDM_TYPE_CAPS_QUERY = 0x16 # type: ignore
+NVDM_TYPE_INFOROM = 0x17 # type: ignore
+NVDM_TYPE_SMBPBI = 0x18 # type: ignore
+NVDM_TYPE_ROMREAD = 0x1A # type: ignore
+NVDM_TYPE_UEFI_RM = 0x1C # type: ignore
+NVDM_TYPE_UEFI_XTL_DEBUG_INTR = 0x1D # type: ignore
+NVDM_TYPE_TNVL = 0x1F # type: ignore
+NVDM_TYPE_CLOCK_BOOST = 0x20 # type: ignore
+NVDM_TYPE_FSP_GSP_COMM = 0x21 # type: ignore
+MAX_GPC_COUNT = 32 # type: ignore
+VGPU_MAX_REGOPS_PER_RPC = 100 # type: ignore
+VGPU_RESERVED_HANDLE_BASE = 0xCAF3F000 # type: ignore
+VGPU_RESERVED_HANDLE_RANGE = 0x1000 # type: ignore
 VGPU_CALC_PARAM_OFFSET = lambda prev_offset,prev_params: (prev_offset + NV_ALIGN_UP(sizeof(prev_params), sizeof(NvU32))) # type: ignore
-NV_VGPU_MSG_HEADER_VERSION_MAJOR_TOT = 0x00000003
-NV_VGPU_MSG_HEADER_VERSION_MINOR_TOT = 0x00000000
-NV_VGPU_MSG_SIGNATURE_VALID = 0x43505256
-NV_VGPU_MSG_RESULT_VMIOP_INVAL = 0xFF000001
-NV_VGPU_MSG_RESULT_VMIOP_RESOURCE = 0xFF000002
-NV_VGPU_MSG_RESULT_VMIOP_RANGE = 0xFF000003
-NV_VGPU_MSG_RESULT_VMIOP_READ_ONLY = 0xFF000004
-NV_VGPU_MSG_RESULT_VMIOP_NOT_FOUND = 0xFF000005
-NV_VGPU_MSG_RESULT_VMIOP_NO_ADDRESS_SPACE = 0xFF000006
-NV_VGPU_MSG_RESULT_VMIOP_TIMEOUT = 0xFF000007
-NV_VGPU_MSG_RESULT_VMIOP_NOT_ALLOWED_IN_CALLBACK = 0xFF000008
-NV_VGPU_MSG_RESULT_VMIOP_ECC_MISMATCH = 0xFF000009
-NV_VGPU_MSG_RESULT_VMIOP_NOT_SUPPORTED = 0xFF00000a
-NV_VGPU_MSG_RESULT_RPC_UNKNOWN_FUNCTION = 0xFF100001
-NV_VGPU_MSG_RESULT_RPC_INVALID_MESSAGE_FORMAT = 0xFF100002
-NV_VGPU_MSG_RESULT_RPC_HANDLE_NOT_FOUND = 0xFF100003
-NV_VGPU_MSG_RESULT_RPC_HANDLE_EXISTS = 0xFF100004
-NV_VGPU_MSG_RESULT_RPC_UNKNOWN_RM_ERROR = 0xFF100005
-NV_VGPU_MSG_RESULT_RPC_UNKNOWN_VMIOP_ERROR = 0xFF100006
-NV_VGPU_MSG_RESULT_RPC_RESERVED_HANDLE = 0xFF100007
-NV_VGPU_MSG_RESULT_RPC_CUDA_PROFILING_DISABLED = 0xFF100008
-NV_VGPU_MSG_RESULT_RPC_API_CONTROL_NOT_SUPPORTED = 0xFF100009
-NV_VGPU_MSG_RESULT_RPC_PENDING = 0xFFFFFFFF
-NV_VGPU_MSG_UNION_INIT = 0x00000000
-NV_VGPU_PTEDESC_INIT = 0x00000000
-NV_VGPU_PTEDESC__PROD = 0x00000000
-NV_VGPU_PTEDESC_IDR_NONE = 0x00000000
-NV_VGPU_PTEDESC_IDR_SINGLE = 0x00000001
-NV_VGPU_PTEDESC_IDR_DOUBLE = 0x00000002
-NV_VGPU_PTEDESC_IDR_TRIPLE = 0x00000003
-NV_VGPU_PTE_PAGE_SIZE = 0x1000
-NV_VGPU_PTE_SIZE = 4
-NV_VGPU_PTE_INDEX_SHIFT = 10
-NV_VGPU_PTE_INDEX_MASK = 0x3FF
-NV_VGPU_PTE_64_PAGE_SIZE = 0x1000
-NV_VGPU_PTE_64_SIZE = 8
-NV_VGPU_PTE_64_INDEX_SHIFT = 9
-NV_VGPU_PTE_64_INDEX_MASK = 0x1FF
-NV_VGPU_LOG_LEVEL_FATAL = 0x00000000
-NV_VGPU_LOG_LEVEL_ERROR = 0x00000001
-NV_VGPU_LOG_LEVEL_NOTICE = 0x00000002
-NV_VGPU_LOG_LEVEL_STATUS = 0x00000003
-NV_VGPU_LOG_LEVEL_DEBUG = 0x00000004
-VGPU_RPC_GET_P2P_CAPS_V2_MAX_GPUS_SQUARED_PER_RPC = 512
-GR_MAX_RPC_CTX_BUFFER_COUNT = 32
-VGPU_RPC_CTRL_DEBUG_READ_ALL_SM_ERROR_STATES_PER_RPC_v21_06 = 80
-LIBOS_MEMORY_REGION_INIT_ARGUMENTS_MAX = 4096
-LIBOS_MEMORY_REGION_RADIX_PAGE_SIZE = 4096
-LIBOS_MEMORY_REGION_RADIX_PAGE_LOG2 = 12
-MSGQ_VERSION = 0
-MAX_DSM_SUPPORTED_FUNCS_RTN_LEN = 8
-NV_ACPI_GENERIC_FUNC_COUNT = 8
-REGISTRY_TABLE_ENTRY_TYPE_UNKNOWN = 0
-REGISTRY_TABLE_ENTRY_TYPE_DWORD = 1
-REGISTRY_TABLE_ENTRY_TYPE_BINARY = 2
-REGISTRY_TABLE_ENTRY_TYPE_STRING = 3
-MAX_GROUP_COUNT = 2
-RM_ENGINE_TYPE_GRAPHICS = RM_ENGINE_TYPE_GR0
-RM_ENGINE_TYPE_BSP = RM_ENGINE_TYPE_NVDEC0
-RM_ENGINE_TYPE_MSENC = RM_ENGINE_TYPE_NVENC0
-RM_ENGINE_TYPE_CIPHER = RM_ENGINE_TYPE_TSEC
-RM_ENGINE_TYPE_NVJPG = RM_ENGINE_TYPE_NVJPEG0
-RM_ENGINE_TYPE_COPY_SIZE = 20
-RM_ENGINE_TYPE_NVENC_SIZE = 4
-RM_ENGINE_TYPE_NVJPEG_SIZE = 8
-RM_ENGINE_TYPE_NVDEC_SIZE = 8
-RM_ENGINE_TYPE_OFA_SIZE = 2
-RM_ENGINE_TYPE_GR_SIZE = 8
-NVGPU_ENGINE_CAPS_MASK_BITS = 32
-NVGPU_ENGINE_CAPS_MASK_ARRAY_MAX = ((RM_ENGINE_TYPE_LAST-1)/NVGPU_ENGINE_CAPS_MASK_BITS + 1)
+NV_VGPU_MSG_HEADER_VERSION_MAJOR_TOT = 0x00000003 # type: ignore
+NV_VGPU_MSG_HEADER_VERSION_MINOR_TOT = 0x00000000 # type: ignore
+NV_VGPU_MSG_SIGNATURE_VALID = 0x43505256 # type: ignore
+NV_VGPU_MSG_RESULT_VMIOP_INVAL = 0xFF000001 # type: ignore
+NV_VGPU_MSG_RESULT_VMIOP_RESOURCE = 0xFF000002 # type: ignore
+NV_VGPU_MSG_RESULT_VMIOP_RANGE = 0xFF000003 # type: ignore
+NV_VGPU_MSG_RESULT_VMIOP_READ_ONLY = 0xFF000004 # type: ignore
+NV_VGPU_MSG_RESULT_VMIOP_NOT_FOUND = 0xFF000005 # type: ignore
+NV_VGPU_MSG_RESULT_VMIOP_NO_ADDRESS_SPACE = 0xFF000006 # type: ignore
+NV_VGPU_MSG_RESULT_VMIOP_TIMEOUT = 0xFF000007 # type: ignore
+NV_VGPU_MSG_RESULT_VMIOP_NOT_ALLOWED_IN_CALLBACK = 0xFF000008 # type: ignore
+NV_VGPU_MSG_RESULT_VMIOP_ECC_MISMATCH = 0xFF000009 # type: ignore
+NV_VGPU_MSG_RESULT_VMIOP_NOT_SUPPORTED = 0xFF00000a # type: ignore
+NV_VGPU_MSG_RESULT_RPC_UNKNOWN_FUNCTION = 0xFF100001 # type: ignore
+NV_VGPU_MSG_RESULT_RPC_INVALID_MESSAGE_FORMAT = 0xFF100002 # type: ignore
+NV_VGPU_MSG_RESULT_RPC_HANDLE_NOT_FOUND = 0xFF100003 # type: ignore
+NV_VGPU_MSG_RESULT_RPC_HANDLE_EXISTS = 0xFF100004 # type: ignore
+NV_VGPU_MSG_RESULT_RPC_UNKNOWN_RM_ERROR = 0xFF100005 # type: ignore
+NV_VGPU_MSG_RESULT_RPC_UNKNOWN_VMIOP_ERROR = 0xFF100006 # type: ignore
+NV_VGPU_MSG_RESULT_RPC_RESERVED_HANDLE = 0xFF100007 # type: ignore
+NV_VGPU_MSG_RESULT_RPC_CUDA_PROFILING_DISABLED = 0xFF100008 # type: ignore
+NV_VGPU_MSG_RESULT_RPC_API_CONTROL_NOT_SUPPORTED = 0xFF100009 # type: ignore
+NV_VGPU_MSG_RESULT_RPC_PENDING = 0xFFFFFFFF # type: ignore
+NV_VGPU_MSG_UNION_INIT = 0x00000000 # type: ignore
+NV_VGPU_PTEDESC_INIT = 0x00000000 # type: ignore
+NV_VGPU_PTEDESC__PROD = 0x00000000 # type: ignore
+NV_VGPU_PTEDESC_IDR_NONE = 0x00000000 # type: ignore
+NV_VGPU_PTEDESC_IDR_SINGLE = 0x00000001 # type: ignore
+NV_VGPU_PTEDESC_IDR_DOUBLE = 0x00000002 # type: ignore
+NV_VGPU_PTEDESC_IDR_TRIPLE = 0x00000003 # type: ignore
+NV_VGPU_PTE_PAGE_SIZE = 0x1000 # type: ignore
+NV_VGPU_PTE_SIZE = 4 # type: ignore
+NV_VGPU_PTE_INDEX_SHIFT = 10 # type: ignore
+NV_VGPU_PTE_INDEX_MASK = 0x3FF # type: ignore
+NV_VGPU_PTE_64_PAGE_SIZE = 0x1000 # type: ignore
+NV_VGPU_PTE_64_SIZE = 8 # type: ignore
+NV_VGPU_PTE_64_INDEX_SHIFT = 9 # type: ignore
+NV_VGPU_PTE_64_INDEX_MASK = 0x1FF # type: ignore
+NV_VGPU_LOG_LEVEL_FATAL = 0x00000000 # type: ignore
+NV_VGPU_LOG_LEVEL_ERROR = 0x00000001 # type: ignore
+NV_VGPU_LOG_LEVEL_NOTICE = 0x00000002 # type: ignore
+NV_VGPU_LOG_LEVEL_STATUS = 0x00000003 # type: ignore
+NV_VGPU_LOG_LEVEL_DEBUG = 0x00000004 # type: ignore
+VGPU_RPC_GET_P2P_CAPS_V2_MAX_GPUS_SQUARED_PER_RPC = 512 # type: ignore
+GR_MAX_RPC_CTX_BUFFER_COUNT = 32 # type: ignore
+VGPU_RPC_CTRL_DEBUG_READ_ALL_SM_ERROR_STATES_PER_RPC_v21_06 = 80 # type: ignore
+LIBOS_MEMORY_REGION_INIT_ARGUMENTS_MAX = 4096 # type: ignore
+LIBOS_MEMORY_REGION_RADIX_PAGE_SIZE = 4096 # type: ignore
+LIBOS_MEMORY_REGION_RADIX_PAGE_LOG2 = 12 # type: ignore
+MSGQ_VERSION = 0 # type: ignore
+MAX_DSM_SUPPORTED_FUNCS_RTN_LEN = 8 # type: ignore
+NV_ACPI_GENERIC_FUNC_COUNT = 8 # type: ignore
+REGISTRY_TABLE_ENTRY_TYPE_UNKNOWN = 0 # type: ignore
+REGISTRY_TABLE_ENTRY_TYPE_DWORD = 1 # type: ignore
+REGISTRY_TABLE_ENTRY_TYPE_BINARY = 2 # type: ignore
+REGISTRY_TABLE_ENTRY_TYPE_STRING = 3 # type: ignore
+MAX_GROUP_COUNT = 2 # type: ignore
+RM_ENGINE_TYPE_GRAPHICS = RM_ENGINE_TYPE_GR0 # type: ignore
+RM_ENGINE_TYPE_BSP = RM_ENGINE_TYPE_NVDEC0 # type: ignore
+RM_ENGINE_TYPE_MSENC = RM_ENGINE_TYPE_NVENC0 # type: ignore
+RM_ENGINE_TYPE_CIPHER = RM_ENGINE_TYPE_TSEC # type: ignore
+RM_ENGINE_TYPE_NVJPG = RM_ENGINE_TYPE_NVJPEG0 # type: ignore
+RM_ENGINE_TYPE_COPY_SIZE = 20 # type: ignore
+RM_ENGINE_TYPE_NVENC_SIZE = 4 # type: ignore
+RM_ENGINE_TYPE_NVJPEG_SIZE = 8 # type: ignore
+RM_ENGINE_TYPE_NVDEC_SIZE = 8 # type: ignore
+RM_ENGINE_TYPE_OFA_SIZE = 2 # type: ignore
+RM_ENGINE_TYPE_GR_SIZE = 8 # type: ignore
+NVGPU_ENGINE_CAPS_MASK_BITS = 32 # type: ignore
+NVGPU_ENGINE_CAPS_MASK_ARRAY_MAX = ((RM_ENGINE_TYPE_LAST-1)/NVGPU_ENGINE_CAPS_MASK_BITS + 1) # type: ignore
 NVGPU_GET_ENGINE_CAPS_MASK = lambda caps,id: (caps[(id)/NVGPU_ENGINE_CAPS_MASK_BITS] & NVBIT((id) % NVGPU_ENGINE_CAPS_MASK_BITS)) # type: ignore
-FALCON_APPLICATION_INTERFACE_ENTRY_ID_DMEMMAPPER = (0x4)
-FALCON_APPLICATION_INTERFACE_DMEM_MAPPER_V3_CMD_FRTS = (0x15)
-FALCON_APPLICATION_INTERFACE_DMEM_MAPPER_V3_CMD_SB = (0x19)
-BIT_HEADER_ID = 0xB8FF
-BIT_HEADER_SIGNATURE = 0x00544942
-BIT_HEADER_SIZE_OFFSET = 8
-BIT_HEADER_V1_00_FMT = "1w1d1w4b"
-BIT_TOKEN_V1_00_SIZE_6 = 6
-BIT_TOKEN_V1_00_SIZE_8 = 8
-BIT_TOKEN_V1_00_FMT_SIZE_6 = "2b2w"
-BIT_TOKEN_V1_00_FMT_SIZE_8 = "2b1w1d"
-BIT_TOKEN_BIOSDATA = 0x42
-BIT_DATA_BIOSDATA_VERSION_1 = 0x1
-BIT_DATA_BIOSDATA_VERSION_2 = 0x2
-BIT_DATA_BIOSDATA_BINVER_FMT = "1d1b"
-BIT_DATA_BIOSDATA_BINVER_SIZE_5 = 5
-BIT_TOKEN_FALCON_DATA = 0x70
-BIT_DATA_FALCON_DATA_V2_4_FMT = "1d"
-BIT_DATA_FALCON_DATA_V2_SIZE_4 = 4
-FALCON_UCODE_TABLE_HDR_V1_VERSION = 1
-FALCON_UCODE_TABLE_HDR_V1_SIZE_6 = 6
-FALCON_UCODE_TABLE_HDR_V1_6_FMT = "6b"
-FALCON_UCODE_TABLE_ENTRY_V1_VERSION = 1
-FALCON_UCODE_TABLE_ENTRY_V1_SIZE_6 = 6
-FALCON_UCODE_TABLE_ENTRY_V1_6_FMT = "2b1d"
-FALCON_UCODE_ENTRY_APPID_FIRMWARE_SEC_LIC = 0x05
-FALCON_UCODE_ENTRY_APPID_FWSEC_DBG = 0x45
-FALCON_UCODE_ENTRY_APPID_FWSEC_PROD = 0x85
-NV_BIT_FALCON_UCODE_DESC_HEADER_VDESC_FLAGS_VERSION_UNAVAILABLE = 0x00
-NV_BIT_FALCON_UCODE_DESC_HEADER_VDESC_FLAGS_VERSION_AVAILABLE = 0x01
-NV_BIT_FALCON_UCODE_DESC_HEADER_VDESC_VERSION_V1 = 0x01
-NV_BIT_FALCON_UCODE_DESC_HEADER_VDESC_VERSION_V2 = 0x02
-NV_BIT_FALCON_UCODE_DESC_HEADER_VDESC_VERSION_V3 = 0x03
-NV_BIT_FALCON_UCODE_DESC_HEADER_VDESC_VERSION_V4 = 0x04
-FALCON_UCODE_DESC_HEADER_FORMAT = "1d"
-FALCON_UCODE_DESC_V3_SIZE_44 = 44
-FALCON_UCODE_DESC_V3_44_FMT = "9d1w2b2w"
-BCRT30_RSA3K_SIG_SIZE = 384
-FWSECLIC_READ_VBIOS_STRUCT_FLAGS = (2)
-FWSECLIC_FRTS_REGION_MEDIA_FB = (2)
-FWSECLIC_FRTS_REGION_SIZE_1MB_IN_4K = (0x100)
-NV_BCRT_HASH_INFO_BASE_CODE_TYPE_VBIOS_BASE = 0x00
-NV_BCRT_HASH_INFO_BASE_CODE_TYPE_VBIOS_EXT = 0xE0
-PCI_EXP_ROM_SIGNATURE = 0xaa55
-PCI_EXP_ROM_SIGNATURE_NV = 0x4e56
-PCI_EXP_ROM_SIGNATURE_NV2 = 0xbb77
+FALCON_APPLICATION_INTERFACE_ENTRY_ID_DMEMMAPPER = (0x4) # type: ignore
+FALCON_APPLICATION_INTERFACE_DMEM_MAPPER_V3_CMD_FRTS = (0x15) # type: ignore
+FALCON_APPLICATION_INTERFACE_DMEM_MAPPER_V3_CMD_SB = (0x19) # type: ignore
+BIT_HEADER_ID = 0xB8FF # type: ignore
+BIT_HEADER_SIGNATURE = 0x00544942 # type: ignore
+BIT_HEADER_SIZE_OFFSET = 8 # type: ignore
+BIT_HEADER_V1_00_FMT = "1w1d1w4b" # type: ignore
+BIT_TOKEN_V1_00_SIZE_6 = 6 # type: ignore
+BIT_TOKEN_V1_00_SIZE_8 = 8 # type: ignore
+BIT_TOKEN_V1_00_FMT_SIZE_6 = "2b2w" # type: ignore
+BIT_TOKEN_V1_00_FMT_SIZE_8 = "2b1w1d" # type: ignore
+BIT_TOKEN_BIOSDATA = 0x42 # type: ignore
+BIT_DATA_BIOSDATA_VERSION_1 = 0x1 # type: ignore
+BIT_DATA_BIOSDATA_VERSION_2 = 0x2 # type: ignore
+BIT_DATA_BIOSDATA_BINVER_FMT = "1d1b" # type: ignore
+BIT_DATA_BIOSDATA_BINVER_SIZE_5 = 5 # type: ignore
+BIT_TOKEN_FALCON_DATA = 0x70 # type: ignore
+BIT_DATA_FALCON_DATA_V2_4_FMT = "1d" # type: ignore
+BIT_DATA_FALCON_DATA_V2_SIZE_4 = 4 # type: ignore
+FALCON_UCODE_TABLE_HDR_V1_VERSION = 1 # type: ignore
+FALCON_UCODE_TABLE_HDR_V1_SIZE_6 = 6 # type: ignore
+FALCON_UCODE_TABLE_HDR_V1_6_FMT = "6b" # type: ignore
+FALCON_UCODE_TABLE_ENTRY_V1_VERSION = 1 # type: ignore
+FALCON_UCODE_TABLE_ENTRY_V1_SIZE_6 = 6 # type: ignore
+FALCON_UCODE_TABLE_ENTRY_V1_6_FMT = "2b1d" # type: ignore
+FALCON_UCODE_ENTRY_APPID_FIRMWARE_SEC_LIC = 0x05 # type: ignore
+FALCON_UCODE_ENTRY_APPID_FWSEC_DBG = 0x45 # type: ignore
+FALCON_UCODE_ENTRY_APPID_FWSEC_PROD = 0x85 # type: ignore
+NV_BIT_FALCON_UCODE_DESC_HEADER_VDESC_FLAGS_VERSION_UNAVAILABLE = 0x00 # type: ignore
+NV_BIT_FALCON_UCODE_DESC_HEADER_VDESC_FLAGS_VERSION_AVAILABLE = 0x01 # type: ignore
+NV_BIT_FALCON_UCODE_DESC_HEADER_VDESC_VERSION_V1 = 0x01 # type: ignore
+NV_BIT_FALCON_UCODE_DESC_HEADER_VDESC_VERSION_V2 = 0x02 # type: ignore
+NV_BIT_FALCON_UCODE_DESC_HEADER_VDESC_VERSION_V3 = 0x03 # type: ignore
+NV_BIT_FALCON_UCODE_DESC_HEADER_VDESC_VERSION_V4 = 0x04 # type: ignore
+FALCON_UCODE_DESC_HEADER_FORMAT = "1d" # type: ignore
+FALCON_UCODE_DESC_V3_SIZE_44 = 44 # type: ignore
+FALCON_UCODE_DESC_V3_44_FMT = "9d1w2b2w" # type: ignore
+BCRT30_RSA3K_SIG_SIZE = 384 # type: ignore
+FWSECLIC_READ_VBIOS_STRUCT_FLAGS = (2) # type: ignore
+FWSECLIC_FRTS_REGION_MEDIA_FB = (2) # type: ignore
+FWSECLIC_FRTS_REGION_SIZE_1MB_IN_4K = (0x100) # type: ignore
+NV_BCRT_HASH_INFO_BASE_CODE_TYPE_VBIOS_BASE = 0x00 # type: ignore
+NV_BCRT_HASH_INFO_BASE_CODE_TYPE_VBIOS_EXT = 0xE0 # type: ignore
+PCI_EXP_ROM_SIGNATURE = 0xaa55 # type: ignore
+PCI_EXP_ROM_SIGNATURE_NV = 0x4e56 # type: ignore
+PCI_EXP_ROM_SIGNATURE_NV2 = 0xbb77 # type: ignore
 IS_VALID_PCI_ROM_SIG = lambda sig: ((sig == PCI_EXP_ROM_SIGNATURE) or (sig == PCI_EXP_ROM_SIGNATURE_NV) or (sig == PCI_EXP_ROM_SIGNATURE_NV2)) # type: ignore
-OFFSETOF_PCI_EXP_ROM_SIG = 0x0
-OFFSETOF_PCI_EXP_ROM_NBSI_DATA_OFFSET = 0x16
-OFFSETOF_PCI_EXP_ROM_PCI_DATA_STRUCT_PTR = 0x18
-PCI_DATA_STRUCT_SIGNATURE = 0x52494350
-PCI_DATA_STRUCT_SIGNATURE_NV = 0x5344504E
-PCI_DATA_STRUCT_SIGNATURE_NV2 = 0x53494752
+OFFSETOF_PCI_EXP_ROM_SIG = 0x0 # type: ignore
+OFFSETOF_PCI_EXP_ROM_NBSI_DATA_OFFSET = 0x16 # type: ignore
+OFFSETOF_PCI_EXP_ROM_PCI_DATA_STRUCT_PTR = 0x18 # type: ignore
+PCI_DATA_STRUCT_SIGNATURE = 0x52494350 # type: ignore
+PCI_DATA_STRUCT_SIGNATURE_NV = 0x5344504E # type: ignore
+PCI_DATA_STRUCT_SIGNATURE_NV2 = 0x53494752 # type: ignore
 IS_VALID_PCI_DATA_SIG = lambda sig: ((sig == PCI_DATA_STRUCT_SIGNATURE) or (sig == PCI_DATA_STRUCT_SIGNATURE_NV) or (sig == PCI_DATA_STRUCT_SIGNATURE_NV2)) # type: ignore
-PCI_ROM_IMAGE_BLOCK_SIZE = 512
-OFFSETOF_PCI_DATA_STRUCT_SIG = 0x0
-OFFSETOF_PCI_DATA_STRUCT_VENDOR_ID = 0x4
-OFFSETOF_PCI_DATA_STRUCT_LEN = 0xa
-OFFSETOF_PCI_DATA_STRUCT_CLASS_CODE = 0xd
-OFFSETOF_PCI_DATA_STRUCT_CODE_TYPE = 0x14
-OFFSETOF_PCI_DATA_STRUCT_IMAGE_LEN = 0x10
-OFFSETOF_PCI_DATA_STRUCT_LAST_IMAGE = 0x15
-NV_PCI_DATA_EXT_SIG = 0x4544504E
-NV_PCI_DATA_EXT_REV_10 = 0x100
-NV_PCI_DATA_EXT_REV_11 = 0x101
-OFFSETOF_PCI_DATA_EXT_STRUCT_SIG = 0x0
-OFFSETOF_PCI_DATA_EXT_STRUCT_LEN = 0x6
-OFFSETOF_PCI_DATA_EXT_STRUCT_REV = 0x4
-OFFSETOF_PCI_DATA_EXT_STRUCT_SUBIMAGE_LEN = 0x8
-OFFSETOF_PCI_DATA_EXT_STRUCT_LAST_IMAGE = 0xa
-OFFSETOF_PCI_DATA_EXT_STRUCT_FLAGS = 0xb
-PCI_DATA_EXT_STRUCT_FLAGS_CHECKSUM_DISABLED = 0x04
+PCI_ROM_IMAGE_BLOCK_SIZE = 512 # type: ignore
+OFFSETOF_PCI_DATA_STRUCT_SIG = 0x0 # type: ignore
+OFFSETOF_PCI_DATA_STRUCT_VENDOR_ID = 0x4 # type: ignore
+OFFSETOF_PCI_DATA_STRUCT_LEN = 0xa # type: ignore
+OFFSETOF_PCI_DATA_STRUCT_CLASS_CODE = 0xd # type: ignore
+OFFSETOF_PCI_DATA_STRUCT_CODE_TYPE = 0x14 # type: ignore
+OFFSETOF_PCI_DATA_STRUCT_IMAGE_LEN = 0x10 # type: ignore
+OFFSETOF_PCI_DATA_STRUCT_LAST_IMAGE = 0x15 # type: ignore
+NV_PCI_DATA_EXT_SIG = 0x4544504E # type: ignore
+NV_PCI_DATA_EXT_REV_10 = 0x100 # type: ignore
+NV_PCI_DATA_EXT_REV_11 = 0x101 # type: ignore
+OFFSETOF_PCI_DATA_EXT_STRUCT_SIG = 0x0 # type: ignore
+OFFSETOF_PCI_DATA_EXT_STRUCT_LEN = 0x6 # type: ignore
+OFFSETOF_PCI_DATA_EXT_STRUCT_REV = 0x4 # type: ignore
+OFFSETOF_PCI_DATA_EXT_STRUCT_SUBIMAGE_LEN = 0x8 # type: ignore
+OFFSETOF_PCI_DATA_EXT_STRUCT_LAST_IMAGE = 0xa # type: ignore
+OFFSETOF_PCI_DATA_EXT_STRUCT_FLAGS = 0xb # type: ignore
+PCI_DATA_EXT_STRUCT_FLAGS_CHECKSUM_DISABLED = 0x04 # type: ignore

tinygrad/runtime/autogen/opencl.py

@@ -344,468 +344,468 @@ def clCreateCommandQueue(context:cl_context, device:cl_device_id, properties:cl_
 def clCreateSampler(context:cl_context, normalized_coords:cl_bool, addressing_mode:cl_addressing_mode, filter_mode:cl_filter_mode, errcode_ret:c.POINTER[cl_int]) -> cl_sampler: ...
 @dll.bind(cl_int, cl_command_queue, cl_kernel, cl_uint, c.POINTER[cl_event], c.POINTER[cl_event])
 def clEnqueueTask(command_queue:cl_command_queue, kernel:cl_kernel, num_events_in_wait_list:cl_uint, event_wait_list:c.POINTER[cl_event], event:c.POINTER[cl_event]) -> cl_int: ...
-CL_NAME_VERSION_MAX_NAME_SIZE = 64
-CL_SUCCESS = 0
-CL_DEVICE_NOT_FOUND = -1
-CL_DEVICE_NOT_AVAILABLE = -2
-CL_COMPILER_NOT_AVAILABLE = -3
-CL_MEM_OBJECT_ALLOCATION_FAILURE = -4
-CL_OUT_OF_RESOURCES = -5
-CL_OUT_OF_HOST_MEMORY = -6
-CL_PROFILING_INFO_NOT_AVAILABLE = -7
-CL_MEM_COPY_OVERLAP = -8
-CL_IMAGE_FORMAT_MISMATCH = -9
-CL_IMAGE_FORMAT_NOT_SUPPORTED = -10
-CL_BUILD_PROGRAM_FAILURE = -11
-CL_MAP_FAILURE = -12
-CL_MISALIGNED_SUB_BUFFER_OFFSET = -13
-CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST = -14
-CL_COMPILE_PROGRAM_FAILURE = -15
-CL_LINKER_NOT_AVAILABLE = -16
-CL_LINK_PROGRAM_FAILURE = -17
-CL_DEVICE_PARTITION_FAILED = -18
-CL_KERNEL_ARG_INFO_NOT_AVAILABLE = -19
-CL_INVALID_VALUE = -30
-CL_INVALID_DEVICE_TYPE = -31
-CL_INVALID_PLATFORM = -32
-CL_INVALID_DEVICE = -33
-CL_INVALID_CONTEXT = -34
-CL_INVALID_QUEUE_PROPERTIES = -35
-CL_INVALID_COMMAND_QUEUE = -36
-CL_INVALID_HOST_PTR = -37
-CL_INVALID_MEM_OBJECT = -38
-CL_INVALID_IMAGE_FORMAT_DESCRIPTOR = -39
-CL_INVALID_IMAGE_SIZE = -40
-CL_INVALID_SAMPLER = -41
-CL_INVALID_BINARY = -42
-CL_INVALID_BUILD_OPTIONS = -43
-CL_INVALID_PROGRAM = -44
-CL_INVALID_PROGRAM_EXECUTABLE = -45
-CL_INVALID_KERNEL_NAME = -46
-CL_INVALID_KERNEL_DEFINITION = -47
-CL_INVALID_KERNEL = -48
-CL_INVALID_ARG_INDEX = -49
-CL_INVALID_ARG_VALUE = -50
-CL_INVALID_ARG_SIZE = -51
-CL_INVALID_KERNEL_ARGS = -52
-CL_INVALID_WORK_DIMENSION = -53
-CL_INVALID_WORK_GROUP_SIZE = -54
-CL_INVALID_WORK_ITEM_SIZE = -55
-CL_INVALID_GLOBAL_OFFSET = -56
-CL_INVALID_EVENT_WAIT_LIST = -57
-CL_INVALID_EVENT = -58
-CL_INVALID_OPERATION = -59
-CL_INVALID_GL_OBJECT = -60
-CL_INVALID_BUFFER_SIZE = -61
-CL_INVALID_MIP_LEVEL = -62
-CL_INVALID_GLOBAL_WORK_SIZE = -63
-CL_INVALID_PROPERTY = -64
-CL_INVALID_IMAGE_DESCRIPTOR = -65
-CL_INVALID_COMPILER_OPTIONS = -66
-CL_INVALID_LINKER_OPTIONS = -67
-CL_INVALID_DEVICE_PARTITION_COUNT = -68
-CL_INVALID_PIPE_SIZE = -69
-CL_INVALID_DEVICE_QUEUE = -70
-CL_INVALID_SPEC_ID = -71
-CL_MAX_SIZE_RESTRICTION_EXCEEDED = -72
-CL_FALSE = 0
-CL_TRUE = 1
-CL_BLOCKING = CL_TRUE
-CL_NON_BLOCKING = CL_FALSE
-CL_PLATFORM_PROFILE = 0x0900
-CL_PLATFORM_VERSION = 0x0901
-CL_PLATFORM_NAME = 0x0902
-CL_PLATFORM_VENDOR = 0x0903
-CL_PLATFORM_EXTENSIONS = 0x0904
-CL_PLATFORM_HOST_TIMER_RESOLUTION = 0x0905
-CL_PLATFORM_NUMERIC_VERSION = 0x0906
-CL_PLATFORM_EXTENSIONS_WITH_VERSION = 0x0907
-CL_DEVICE_TYPE_DEFAULT = (1 << 0)
-CL_DEVICE_TYPE_CPU = (1 << 1)
-CL_DEVICE_TYPE_GPU = (1 << 2)
-CL_DEVICE_TYPE_ACCELERATOR = (1 << 3)
-CL_DEVICE_TYPE_CUSTOM = (1 << 4)
-CL_DEVICE_TYPE_ALL = 0xFFFFFFFF
-CL_DEVICE_TYPE = 0x1000
-CL_DEVICE_VENDOR_ID = 0x1001
-CL_DEVICE_MAX_COMPUTE_UNITS = 0x1002
-CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS = 0x1003
-CL_DEVICE_MAX_WORK_GROUP_SIZE = 0x1004
-CL_DEVICE_MAX_WORK_ITEM_SIZES = 0x1005
-CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR = 0x1006
-CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT = 0x1007
-CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT = 0x1008
-CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG = 0x1009
-CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT = 0x100A
-CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE = 0x100B
-CL_DEVICE_MAX_CLOCK_FREQUENCY = 0x100C
-CL_DEVICE_ADDRESS_BITS = 0x100D
-CL_DEVICE_MAX_READ_IMAGE_ARGS = 0x100E
-CL_DEVICE_MAX_WRITE_IMAGE_ARGS = 0x100F
-CL_DEVICE_MAX_MEM_ALLOC_SIZE = 0x1010
-CL_DEVICE_IMAGE2D_MAX_WIDTH = 0x1011
-CL_DEVICE_IMAGE2D_MAX_HEIGHT = 0x1012
-CL_DEVICE_IMAGE3D_MAX_WIDTH = 0x1013
-CL_DEVICE_IMAGE3D_MAX_HEIGHT = 0x1014
-CL_DEVICE_IMAGE3D_MAX_DEPTH = 0x1015
-CL_DEVICE_IMAGE_SUPPORT = 0x1016
-CL_DEVICE_MAX_PARAMETER_SIZE = 0x1017
-CL_DEVICE_MAX_SAMPLERS = 0x1018
-CL_DEVICE_MEM_BASE_ADDR_ALIGN = 0x1019
-CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE = 0x101A
-CL_DEVICE_SINGLE_FP_CONFIG = 0x101B
-CL_DEVICE_GLOBAL_MEM_CACHE_TYPE = 0x101C
-CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE = 0x101D
-CL_DEVICE_GLOBAL_MEM_CACHE_SIZE = 0x101E
-CL_DEVICE_GLOBAL_MEM_SIZE = 0x101F
-CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE = 0x1020
-CL_DEVICE_MAX_CONSTANT_ARGS = 0x1021
-CL_DEVICE_LOCAL_MEM_TYPE = 0x1022
-CL_DEVICE_LOCAL_MEM_SIZE = 0x1023
-CL_DEVICE_ERROR_CORRECTION_SUPPORT = 0x1024
-CL_DEVICE_PROFILING_TIMER_RESOLUTION = 0x1025
-CL_DEVICE_ENDIAN_LITTLE = 0x1026
-CL_DEVICE_AVAILABLE = 0x1027
-CL_DEVICE_COMPILER_AVAILABLE = 0x1028
-CL_DEVICE_EXECUTION_CAPABILITIES = 0x1029
-CL_DEVICE_QUEUE_PROPERTIES = 0x102A
-CL_DEVICE_QUEUE_ON_HOST_PROPERTIES = 0x102A
-CL_DEVICE_NAME = 0x102B
-CL_DEVICE_VENDOR = 0x102C
-CL_DRIVER_VERSION = 0x102D
-CL_DEVICE_PROFILE = 0x102E
-CL_DEVICE_VERSION = 0x102F
-CL_DEVICE_EXTENSIONS = 0x1030
-CL_DEVICE_PLATFORM = 0x1031
-CL_DEVICE_DOUBLE_FP_CONFIG = 0x1032
-CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF = 0x1034
-CL_DEVICE_HOST_UNIFIED_MEMORY = 0x1035
-CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR = 0x1036
-CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT = 0x1037
-CL_DEVICE_NATIVE_VECTOR_WIDTH_INT = 0x1038
-CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG = 0x1039
-CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT = 0x103A
-CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE = 0x103B
-CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF = 0x103C
-CL_DEVICE_OPENCL_C_VERSION = 0x103D
-CL_DEVICE_LINKER_AVAILABLE = 0x103E
-CL_DEVICE_BUILT_IN_KERNELS = 0x103F
-CL_DEVICE_IMAGE_MAX_BUFFER_SIZE = 0x1040
-CL_DEVICE_IMAGE_MAX_ARRAY_SIZE = 0x1041
-CL_DEVICE_PARENT_DEVICE = 0x1042
-CL_DEVICE_PARTITION_MAX_SUB_DEVICES = 0x1043
-CL_DEVICE_PARTITION_PROPERTIES = 0x1044
-CL_DEVICE_PARTITION_AFFINITY_DOMAIN = 0x1045
-CL_DEVICE_PARTITION_TYPE = 0x1046
-CL_DEVICE_REFERENCE_COUNT = 0x1047
-CL_DEVICE_PREFERRED_INTEROP_USER_SYNC = 0x1048
-CL_DEVICE_PRINTF_BUFFER_SIZE = 0x1049
-CL_DEVICE_IMAGE_PITCH_ALIGNMENT = 0x104A
-CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT = 0x104B
-CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS = 0x104C
-CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE = 0x104D
-CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES = 0x104E
-CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE = 0x104F
-CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE = 0x1050
-CL_DEVICE_MAX_ON_DEVICE_QUEUES = 0x1051
-CL_DEVICE_MAX_ON_DEVICE_EVENTS = 0x1052
-CL_DEVICE_SVM_CAPABILITIES = 0x1053
-CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE = 0x1054
-CL_DEVICE_MAX_PIPE_ARGS = 0x1055
-CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS = 0x1056
-CL_DEVICE_PIPE_MAX_PACKET_SIZE = 0x1057
-CL_DEVICE_PREFERRED_PLATFORM_ATOMIC_ALIGNMENT = 0x1058
-CL_DEVICE_PREFERRED_GLOBAL_ATOMIC_ALIGNMENT = 0x1059
-CL_DEVICE_PREFERRED_LOCAL_ATOMIC_ALIGNMENT = 0x105A
-CL_DEVICE_IL_VERSION = 0x105B
-CL_DEVICE_MAX_NUM_SUB_GROUPS = 0x105C
-CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS = 0x105D
-CL_DEVICE_NUMERIC_VERSION = 0x105E
-CL_DEVICE_EXTENSIONS_WITH_VERSION = 0x1060
-CL_DEVICE_ILS_WITH_VERSION = 0x1061
-CL_DEVICE_BUILT_IN_KERNELS_WITH_VERSION = 0x1062
-CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES = 0x1063
-CL_DEVICE_ATOMIC_FENCE_CAPABILITIES = 0x1064
-CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT = 0x1065
-CL_DEVICE_OPENCL_C_ALL_VERSIONS = 0x1066
-CL_DEVICE_PREFERRED_WORK_GROUP_SIZE_MULTIPLE = 0x1067
-CL_DEVICE_WORK_GROUP_COLLECTIVE_FUNCTIONS_SUPPORT = 0x1068
-CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT = 0x1069
-CL_DEVICE_OPENCL_C_FEATURES = 0x106F
-CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES = 0x1070
-CL_DEVICE_PIPE_SUPPORT = 0x1071
-CL_DEVICE_LATEST_CONFORMANCE_VERSION_PASSED = 0x1072
-CL_FP_DENORM = (1 << 0)
-CL_FP_INF_NAN = (1 << 1)
-CL_FP_ROUND_TO_NEAREST = (1 << 2)
-CL_FP_ROUND_TO_ZERO = (1 << 3)
-CL_FP_ROUND_TO_INF = (1 << 4)
-CL_FP_FMA = (1 << 5)
-CL_FP_SOFT_FLOAT = (1 << 6)
-CL_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT = (1 << 7)
-CL_NONE = 0x0
-CL_READ_ONLY_CACHE = 0x1
-CL_READ_WRITE_CACHE = 0x2
-CL_LOCAL = 0x1
-CL_GLOBAL = 0x2
-CL_EXEC_KERNEL = (1 << 0)
-CL_EXEC_NATIVE_KERNEL = (1 << 1)
-CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE = (1 << 0)
-CL_QUEUE_PROFILING_ENABLE = (1 << 1)
-CL_QUEUE_ON_DEVICE = (1 << 2)
-CL_QUEUE_ON_DEVICE_DEFAULT = (1 << 3)
-CL_CONTEXT_REFERENCE_COUNT = 0x1080
-CL_CONTEXT_DEVICES = 0x1081
-CL_CONTEXT_PROPERTIES = 0x1082
-CL_CONTEXT_NUM_DEVICES = 0x1083
-CL_CONTEXT_PLATFORM = 0x1084
-CL_CONTEXT_INTEROP_USER_SYNC = 0x1085
-CL_DEVICE_PARTITION_EQUALLY = 0x1086
-CL_DEVICE_PARTITION_BY_COUNTS = 0x1087
-CL_DEVICE_PARTITION_BY_COUNTS_LIST_END = 0x0
-CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN = 0x1088
-CL_DEVICE_AFFINITY_DOMAIN_NUMA = (1 << 0)
-CL_DEVICE_AFFINITY_DOMAIN_L4_CACHE = (1 << 1)
-CL_DEVICE_AFFINITY_DOMAIN_L3_CACHE = (1 << 2)
-CL_DEVICE_AFFINITY_DOMAIN_L2_CACHE = (1 << 3)
-CL_DEVICE_AFFINITY_DOMAIN_L1_CACHE = (1 << 4)
-CL_DEVICE_AFFINITY_DOMAIN_NEXT_PARTITIONABLE = (1 << 5)
-CL_DEVICE_SVM_COARSE_GRAIN_BUFFER = (1 << 0)
-CL_DEVICE_SVM_FINE_GRAIN_BUFFER = (1 << 1)
-CL_DEVICE_SVM_FINE_GRAIN_SYSTEM = (1 << 2)
-CL_DEVICE_SVM_ATOMICS = (1 << 3)
-CL_QUEUE_CONTEXT = 0x1090
-CL_QUEUE_DEVICE = 0x1091
-CL_QUEUE_REFERENCE_COUNT = 0x1092
-CL_QUEUE_PROPERTIES = 0x1093
-CL_QUEUE_SIZE = 0x1094
-CL_QUEUE_DEVICE_DEFAULT = 0x1095
-CL_QUEUE_PROPERTIES_ARRAY = 0x1098
-CL_MEM_READ_WRITE = (1 << 0)
-CL_MEM_WRITE_ONLY = (1 << 1)
-CL_MEM_READ_ONLY = (1 << 2)
-CL_MEM_USE_HOST_PTR = (1 << 3)
-CL_MEM_ALLOC_HOST_PTR = (1 << 4)
-CL_MEM_COPY_HOST_PTR = (1 << 5)
-CL_MEM_HOST_WRITE_ONLY = (1 << 7)
-CL_MEM_HOST_READ_ONLY = (1 << 8)
-CL_MEM_HOST_NO_ACCESS = (1 << 9)
-CL_MEM_SVM_FINE_GRAIN_BUFFER = (1 << 10)
-CL_MEM_SVM_ATOMICS = (1 << 11)
-CL_MEM_KERNEL_READ_AND_WRITE = (1 << 12)
-CL_MIGRATE_MEM_OBJECT_HOST = (1 << 0)
-CL_MIGRATE_MEM_OBJECT_CONTENT_UNDEFINED = (1 << 1)
-CL_R = 0x10B0
-CL_A = 0x10B1
-CL_RG = 0x10B2
-CL_RA = 0x10B3
-CL_RGB = 0x10B4
-CL_RGBA = 0x10B5
-CL_BGRA = 0x10B6
-CL_ARGB = 0x10B7
-CL_INTENSITY = 0x10B8
-CL_LUMINANCE = 0x10B9
-CL_Rx = 0x10BA
-CL_RGx = 0x10BB
-CL_RGBx = 0x10BC
-CL_DEPTH = 0x10BD
-CL_sRGB = 0x10BF
-CL_sRGBx = 0x10C0
-CL_sRGBA = 0x10C1
-CL_sBGRA = 0x10C2
-CL_ABGR = 0x10C3
-CL_SNORM_INT8 = 0x10D0
-CL_SNORM_INT16 = 0x10D1
-CL_UNORM_INT8 = 0x10D2
-CL_UNORM_INT16 = 0x10D3
-CL_UNORM_SHORT_565 = 0x10D4
-CL_UNORM_SHORT_555 = 0x10D5
-CL_UNORM_INT_101010 = 0x10D6
-CL_SIGNED_INT8 = 0x10D7
-CL_SIGNED_INT16 = 0x10D8
-CL_SIGNED_INT32 = 0x10D9
-CL_UNSIGNED_INT8 = 0x10DA
-CL_UNSIGNED_INT16 = 0x10DB
-CL_UNSIGNED_INT32 = 0x10DC
-CL_HALF_FLOAT = 0x10DD
-CL_FLOAT = 0x10DE
-CL_UNORM_INT_101010_2 = 0x10E0
-CL_MEM_OBJECT_BUFFER = 0x10F0
-CL_MEM_OBJECT_IMAGE2D = 0x10F1
-CL_MEM_OBJECT_IMAGE3D = 0x10F2
-CL_MEM_OBJECT_IMAGE2D_ARRAY = 0x10F3
-CL_MEM_OBJECT_IMAGE1D = 0x10F4
-CL_MEM_OBJECT_IMAGE1D_ARRAY = 0x10F5
-CL_MEM_OBJECT_IMAGE1D_BUFFER = 0x10F6
-CL_MEM_OBJECT_PIPE = 0x10F7
-CL_MEM_TYPE = 0x1100
-CL_MEM_FLAGS = 0x1101
-CL_MEM_SIZE = 0x1102
-CL_MEM_HOST_PTR = 0x1103
-CL_MEM_MAP_COUNT = 0x1104
-CL_MEM_REFERENCE_COUNT = 0x1105
-CL_MEM_CONTEXT = 0x1106
-CL_MEM_ASSOCIATED_MEMOBJECT = 0x1107
-CL_MEM_OFFSET = 0x1108
-CL_MEM_USES_SVM_POINTER = 0x1109
-CL_MEM_PROPERTIES = 0x110A
-CL_IMAGE_FORMAT = 0x1110
-CL_IMAGE_ELEMENT_SIZE = 0x1111
-CL_IMAGE_ROW_PITCH = 0x1112
-CL_IMAGE_SLICE_PITCH = 0x1113
-CL_IMAGE_WIDTH = 0x1114
-CL_IMAGE_HEIGHT = 0x1115
-CL_IMAGE_DEPTH = 0x1116
-CL_IMAGE_ARRAY_SIZE = 0x1117
-CL_IMAGE_BUFFER = 0x1118
-CL_IMAGE_NUM_MIP_LEVELS = 0x1119
-CL_IMAGE_NUM_SAMPLES = 0x111A
-CL_PIPE_PACKET_SIZE = 0x1120
-CL_PIPE_MAX_PACKETS = 0x1121
-CL_PIPE_PROPERTIES = 0x1122
-CL_ADDRESS_NONE = 0x1130
-CL_ADDRESS_CLAMP_TO_EDGE = 0x1131
-CL_ADDRESS_CLAMP = 0x1132
-CL_ADDRESS_REPEAT = 0x1133
-CL_ADDRESS_MIRRORED_REPEAT = 0x1134
-CL_FILTER_NEAREST = 0x1140
-CL_FILTER_LINEAR = 0x1141
-CL_SAMPLER_REFERENCE_COUNT = 0x1150
-CL_SAMPLER_CONTEXT = 0x1151
-CL_SAMPLER_NORMALIZED_COORDS = 0x1152
-CL_SAMPLER_ADDRESSING_MODE = 0x1153
-CL_SAMPLER_FILTER_MODE = 0x1154
-CL_SAMPLER_MIP_FILTER_MODE = 0x1155
-CL_SAMPLER_LOD_MIN = 0x1156
-CL_SAMPLER_LOD_MAX = 0x1157
-CL_SAMPLER_PROPERTIES = 0x1158
-CL_MAP_READ = (1 << 0)
-CL_MAP_WRITE = (1 << 1)
-CL_MAP_WRITE_INVALIDATE_REGION = (1 << 2)
-CL_PROGRAM_REFERENCE_COUNT = 0x1160
-CL_PROGRAM_CONTEXT = 0x1161
-CL_PROGRAM_NUM_DEVICES = 0x1162
-CL_PROGRAM_DEVICES = 0x1163
-CL_PROGRAM_SOURCE = 0x1164
-CL_PROGRAM_BINARY_SIZES = 0x1165
-CL_PROGRAM_BINARIES = 0x1166
-CL_PROGRAM_NUM_KERNELS = 0x1167
-CL_PROGRAM_KERNEL_NAMES = 0x1168
-CL_PROGRAM_IL = 0x1169
-CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT = 0x116A
-CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT = 0x116B
-CL_PROGRAM_BUILD_STATUS = 0x1181
-CL_PROGRAM_BUILD_OPTIONS = 0x1182
-CL_PROGRAM_BUILD_LOG = 0x1183
-CL_PROGRAM_BINARY_TYPE = 0x1184
-CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE = 0x1185
-CL_PROGRAM_BINARY_TYPE_NONE = 0x0
-CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT = 0x1
-CL_PROGRAM_BINARY_TYPE_LIBRARY = 0x2
-CL_PROGRAM_BINARY_TYPE_EXECUTABLE = 0x4
-CL_BUILD_SUCCESS = 0
-CL_BUILD_NONE = -1
-CL_BUILD_ERROR = -2
-CL_BUILD_IN_PROGRESS = -3
-CL_KERNEL_FUNCTION_NAME = 0x1190
-CL_KERNEL_NUM_ARGS = 0x1191
-CL_KERNEL_REFERENCE_COUNT = 0x1192
-CL_KERNEL_CONTEXT = 0x1193
-CL_KERNEL_PROGRAM = 0x1194
-CL_KERNEL_ATTRIBUTES = 0x1195
-CL_KERNEL_ARG_ADDRESS_QUALIFIER = 0x1196
-CL_KERNEL_ARG_ACCESS_QUALIFIER = 0x1197
-CL_KERNEL_ARG_TYPE_NAME = 0x1198
-CL_KERNEL_ARG_TYPE_QUALIFIER = 0x1199
-CL_KERNEL_ARG_NAME = 0x119A
-CL_KERNEL_ARG_ADDRESS_GLOBAL = 0x119B
-CL_KERNEL_ARG_ADDRESS_LOCAL = 0x119C
-CL_KERNEL_ARG_ADDRESS_CONSTANT = 0x119D
-CL_KERNEL_ARG_ADDRESS_PRIVATE = 0x119E
-CL_KERNEL_ARG_ACCESS_READ_ONLY = 0x11A0
-CL_KERNEL_ARG_ACCESS_WRITE_ONLY = 0x11A1
-CL_KERNEL_ARG_ACCESS_READ_WRITE = 0x11A2
-CL_KERNEL_ARG_ACCESS_NONE = 0x11A3
-CL_KERNEL_ARG_TYPE_NONE = 0
-CL_KERNEL_ARG_TYPE_CONST = (1 << 0)
-CL_KERNEL_ARG_TYPE_RESTRICT = (1 << 1)
-CL_KERNEL_ARG_TYPE_VOLATILE = (1 << 2)
-CL_KERNEL_ARG_TYPE_PIPE = (1 << 3)
-CL_KERNEL_WORK_GROUP_SIZE = 0x11B0
-CL_KERNEL_COMPILE_WORK_GROUP_SIZE = 0x11B1
-CL_KERNEL_LOCAL_MEM_SIZE = 0x11B2
-CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE = 0x11B3
-CL_KERNEL_PRIVATE_MEM_SIZE = 0x11B4
-CL_KERNEL_GLOBAL_WORK_SIZE = 0x11B5
-CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE = 0x2033
-CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE = 0x2034
-CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT = 0x11B8
-CL_KERNEL_MAX_NUM_SUB_GROUPS = 0x11B9
-CL_KERNEL_COMPILE_NUM_SUB_GROUPS = 0x11BA
-CL_KERNEL_EXEC_INFO_SVM_PTRS = 0x11B6
-CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM = 0x11B7
-CL_EVENT_COMMAND_QUEUE = 0x11D0
-CL_EVENT_COMMAND_TYPE = 0x11D1
-CL_EVENT_REFERENCE_COUNT = 0x11D2
-CL_EVENT_COMMAND_EXECUTION_STATUS = 0x11D3
-CL_EVENT_CONTEXT = 0x11D4
-CL_COMMAND_NDRANGE_KERNEL = 0x11F0
-CL_COMMAND_TASK = 0x11F1
-CL_COMMAND_NATIVE_KERNEL = 0x11F2
-CL_COMMAND_READ_BUFFER = 0x11F3
-CL_COMMAND_WRITE_BUFFER = 0x11F4
-CL_COMMAND_COPY_BUFFER = 0x11F5
-CL_COMMAND_READ_IMAGE = 0x11F6
-CL_COMMAND_WRITE_IMAGE = 0x11F7
-CL_COMMAND_COPY_IMAGE = 0x11F8
-CL_COMMAND_COPY_IMAGE_TO_BUFFER = 0x11F9
-CL_COMMAND_COPY_BUFFER_TO_IMAGE = 0x11FA
-CL_COMMAND_MAP_BUFFER = 0x11FB
-CL_COMMAND_MAP_IMAGE = 0x11FC
-CL_COMMAND_UNMAP_MEM_OBJECT = 0x11FD
-CL_COMMAND_MARKER = 0x11FE
-CL_COMMAND_ACQUIRE_GL_OBJECTS = 0x11FF
-CL_COMMAND_RELEASE_GL_OBJECTS = 0x1200
-CL_COMMAND_READ_BUFFER_RECT = 0x1201
-CL_COMMAND_WRITE_BUFFER_RECT = 0x1202
-CL_COMMAND_COPY_BUFFER_RECT = 0x1203
-CL_COMMAND_USER = 0x1204
-CL_COMMAND_BARRIER = 0x1205
-CL_COMMAND_MIGRATE_MEM_OBJECTS = 0x1206
-CL_COMMAND_FILL_BUFFER = 0x1207
-CL_COMMAND_FILL_IMAGE = 0x1208
-CL_COMMAND_SVM_FREE = 0x1209
-CL_COMMAND_SVM_MEMCPY = 0x120A
-CL_COMMAND_SVM_MEMFILL = 0x120B
-CL_COMMAND_SVM_MAP = 0x120C
-CL_COMMAND_SVM_UNMAP = 0x120D
-CL_COMMAND_SVM_MIGRATE_MEM = 0x120E
-CL_COMPLETE = 0x0
-CL_RUNNING = 0x1
-CL_SUBMITTED = 0x2
-CL_QUEUED = 0x3
-CL_BUFFER_CREATE_TYPE_REGION = 0x1220
-CL_PROFILING_COMMAND_QUEUED = 0x1280
-CL_PROFILING_COMMAND_SUBMIT = 0x1281
-CL_PROFILING_COMMAND_START = 0x1282
-CL_PROFILING_COMMAND_END = 0x1283
-CL_PROFILING_COMMAND_COMPLETE = 0x1284
-CL_DEVICE_ATOMIC_ORDER_RELAXED = (1 << 0)
-CL_DEVICE_ATOMIC_ORDER_ACQ_REL = (1 << 1)
-CL_DEVICE_ATOMIC_ORDER_SEQ_CST = (1 << 2)
-CL_DEVICE_ATOMIC_SCOPE_WORK_ITEM = (1 << 3)
-CL_DEVICE_ATOMIC_SCOPE_WORK_GROUP = (1 << 4)
-CL_DEVICE_ATOMIC_SCOPE_DEVICE = (1 << 5)
-CL_DEVICE_ATOMIC_SCOPE_ALL_DEVICES = (1 << 6)
-CL_DEVICE_QUEUE_SUPPORTED = (1 << 0)
-CL_DEVICE_QUEUE_REPLACEABLE_DEFAULT = (1 << 1)
-CL_KHRONOS_VENDOR_ID_CODEPLAY = 0x10004
-CL_VERSION_MAJOR_BITS = (10)
-CL_VERSION_MINOR_BITS = (10)
-CL_VERSION_PATCH_BITS = (12)
-CL_VERSION_MAJOR_MASK = ((1 << CL_VERSION_MAJOR_BITS) - 1)
-CL_VERSION_MINOR_MASK = ((1 << CL_VERSION_MINOR_BITS) - 1)
-CL_VERSION_PATCH_MASK = ((1 << CL_VERSION_PATCH_BITS) - 1)
+CL_NAME_VERSION_MAX_NAME_SIZE = 64 # type: ignore
+CL_SUCCESS = 0 # type: ignore
+CL_DEVICE_NOT_FOUND = -1 # type: ignore
+CL_DEVICE_NOT_AVAILABLE = -2 # type: ignore
+CL_COMPILER_NOT_AVAILABLE = -3 # type: ignore
+CL_MEM_OBJECT_ALLOCATION_FAILURE = -4 # type: ignore
+CL_OUT_OF_RESOURCES = -5 # type: ignore
+CL_OUT_OF_HOST_MEMORY = -6 # type: ignore
+CL_PROFILING_INFO_NOT_AVAILABLE = -7 # type: ignore
+CL_MEM_COPY_OVERLAP = -8 # type: ignore
+CL_IMAGE_FORMAT_MISMATCH = -9 # type: ignore
+CL_IMAGE_FORMAT_NOT_SUPPORTED = -10 # type: ignore
+CL_BUILD_PROGRAM_FAILURE = -11 # type: ignore
+CL_MAP_FAILURE = -12 # type: ignore
+CL_MISALIGNED_SUB_BUFFER_OFFSET = -13 # type: ignore
+CL_EXEC_STATUS_ERROR_FOR_EVENTS_IN_WAIT_LIST = -14 # type: ignore
+CL_COMPILE_PROGRAM_FAILURE = -15 # type: ignore
+CL_LINKER_NOT_AVAILABLE = -16 # type: ignore
+CL_LINK_PROGRAM_FAILURE = -17 # type: ignore
+CL_DEVICE_PARTITION_FAILED = -18 # type: ignore
+CL_KERNEL_ARG_INFO_NOT_AVAILABLE = -19 # type: ignore
+CL_INVALID_VALUE = -30 # type: ignore
+CL_INVALID_DEVICE_TYPE = -31 # type: ignore
+CL_INVALID_PLATFORM = -32 # type: ignore
+CL_INVALID_DEVICE = -33 # type: ignore
+CL_INVALID_CONTEXT = -34 # type: ignore
+CL_INVALID_QUEUE_PROPERTIES = -35 # type: ignore
+CL_INVALID_COMMAND_QUEUE = -36 # type: ignore
+CL_INVALID_HOST_PTR = -37 # type: ignore
+CL_INVALID_MEM_OBJECT = -38 # type: ignore
+CL_INVALID_IMAGE_FORMAT_DESCRIPTOR = -39 # type: ignore
+CL_INVALID_IMAGE_SIZE = -40 # type: ignore
+CL_INVALID_SAMPLER = -41 # type: ignore
+CL_INVALID_BINARY = -42 # type: ignore
+CL_INVALID_BUILD_OPTIONS = -43 # type: ignore
+CL_INVALID_PROGRAM = -44 # type: ignore
+CL_INVALID_PROGRAM_EXECUTABLE = -45 # type: ignore
+CL_INVALID_KERNEL_NAME = -46 # type: ignore
+CL_INVALID_KERNEL_DEFINITION = -47 # type: ignore
+CL_INVALID_KERNEL = -48 # type: ignore
+CL_INVALID_ARG_INDEX = -49 # type: ignore
+CL_INVALID_ARG_VALUE = -50 # type: ignore
+CL_INVALID_ARG_SIZE = -51 # type: ignore
+CL_INVALID_KERNEL_ARGS = -52 # type: ignore
+CL_INVALID_WORK_DIMENSION = -53 # type: ignore
+CL_INVALID_WORK_GROUP_SIZE = -54 # type: ignore
+CL_INVALID_WORK_ITEM_SIZE = -55 # type: ignore
+CL_INVALID_GLOBAL_OFFSET = -56 # type: ignore
+CL_INVALID_EVENT_WAIT_LIST = -57 # type: ignore
+CL_INVALID_EVENT = -58 # type: ignore
+CL_INVALID_OPERATION = -59 # type: ignore
+CL_INVALID_GL_OBJECT = -60 # type: ignore
+CL_INVALID_BUFFER_SIZE = -61 # type: ignore
+CL_INVALID_MIP_LEVEL = -62 # type: ignore
+CL_INVALID_GLOBAL_WORK_SIZE = -63 # type: ignore
+CL_INVALID_PROPERTY = -64 # type: ignore
+CL_INVALID_IMAGE_DESCRIPTOR = -65 # type: ignore
+CL_INVALID_COMPILER_OPTIONS = -66 # type: ignore
+CL_INVALID_LINKER_OPTIONS = -67 # type: ignore
+CL_INVALID_DEVICE_PARTITION_COUNT = -68 # type: ignore
+CL_INVALID_PIPE_SIZE = -69 # type: ignore
+CL_INVALID_DEVICE_QUEUE = -70 # type: ignore
+CL_INVALID_SPEC_ID = -71 # type: ignore
+CL_MAX_SIZE_RESTRICTION_EXCEEDED = -72 # type: ignore
+CL_FALSE = 0 # type: ignore
+CL_TRUE = 1 # type: ignore
+CL_BLOCKING = CL_TRUE # type: ignore
+CL_NON_BLOCKING = CL_FALSE # type: ignore
+CL_PLATFORM_PROFILE = 0x0900 # type: ignore
+CL_PLATFORM_VERSION = 0x0901 # type: ignore
+CL_PLATFORM_NAME = 0x0902 # type: ignore
+CL_PLATFORM_VENDOR = 0x0903 # type: ignore
+CL_PLATFORM_EXTENSIONS = 0x0904 # type: ignore
+CL_PLATFORM_HOST_TIMER_RESOLUTION = 0x0905 # type: ignore
+CL_PLATFORM_NUMERIC_VERSION = 0x0906 # type: ignore
+CL_PLATFORM_EXTENSIONS_WITH_VERSION = 0x0907 # type: ignore
+CL_DEVICE_TYPE_DEFAULT = (1 << 0) # type: ignore
+CL_DEVICE_TYPE_CPU = (1 << 1) # type: ignore
+CL_DEVICE_TYPE_GPU = (1 << 2) # type: ignore
+CL_DEVICE_TYPE_ACCELERATOR = (1 << 3) # type: ignore
+CL_DEVICE_TYPE_CUSTOM = (1 << 4) # type: ignore
+CL_DEVICE_TYPE_ALL = 0xFFFFFFFF # type: ignore
+CL_DEVICE_TYPE = 0x1000 # type: ignore
+CL_DEVICE_VENDOR_ID = 0x1001 # type: ignore
+CL_DEVICE_MAX_COMPUTE_UNITS = 0x1002 # type: ignore
+CL_DEVICE_MAX_WORK_ITEM_DIMENSIONS = 0x1003 # type: ignore
+CL_DEVICE_MAX_WORK_GROUP_SIZE = 0x1004 # type: ignore
+CL_DEVICE_MAX_WORK_ITEM_SIZES = 0x1005 # type: ignore
+CL_DEVICE_PREFERRED_VECTOR_WIDTH_CHAR = 0x1006 # type: ignore
+CL_DEVICE_PREFERRED_VECTOR_WIDTH_SHORT = 0x1007 # type: ignore
+CL_DEVICE_PREFERRED_VECTOR_WIDTH_INT = 0x1008 # type: ignore
+CL_DEVICE_PREFERRED_VECTOR_WIDTH_LONG = 0x1009 # type: ignore
+CL_DEVICE_PREFERRED_VECTOR_WIDTH_FLOAT = 0x100A # type: ignore
+CL_DEVICE_PREFERRED_VECTOR_WIDTH_DOUBLE = 0x100B # type: ignore
+CL_DEVICE_MAX_CLOCK_FREQUENCY = 0x100C # type: ignore
+CL_DEVICE_ADDRESS_BITS = 0x100D # type: ignore
+CL_DEVICE_MAX_READ_IMAGE_ARGS = 0x100E # type: ignore
+CL_DEVICE_MAX_WRITE_IMAGE_ARGS = 0x100F # type: ignore
+CL_DEVICE_MAX_MEM_ALLOC_SIZE = 0x1010 # type: ignore
+CL_DEVICE_IMAGE2D_MAX_WIDTH = 0x1011 # type: ignore
+CL_DEVICE_IMAGE2D_MAX_HEIGHT = 0x1012 # type: ignore
+CL_DEVICE_IMAGE3D_MAX_WIDTH = 0x1013 # type: ignore
+CL_DEVICE_IMAGE3D_MAX_HEIGHT = 0x1014 # type: ignore
+CL_DEVICE_IMAGE3D_MAX_DEPTH = 0x1015 # type: ignore
+CL_DEVICE_IMAGE_SUPPORT = 0x1016 # type: ignore
+CL_DEVICE_MAX_PARAMETER_SIZE = 0x1017 # type: ignore
+CL_DEVICE_MAX_SAMPLERS = 0x1018 # type: ignore
+CL_DEVICE_MEM_BASE_ADDR_ALIGN = 0x1019 # type: ignore
+CL_DEVICE_MIN_DATA_TYPE_ALIGN_SIZE = 0x101A # type: ignore
+CL_DEVICE_SINGLE_FP_CONFIG = 0x101B # type: ignore
+CL_DEVICE_GLOBAL_MEM_CACHE_TYPE = 0x101C # type: ignore
+CL_DEVICE_GLOBAL_MEM_CACHELINE_SIZE = 0x101D # type: ignore
+CL_DEVICE_GLOBAL_MEM_CACHE_SIZE = 0x101E # type: ignore
+CL_DEVICE_GLOBAL_MEM_SIZE = 0x101F # type: ignore
+CL_DEVICE_MAX_CONSTANT_BUFFER_SIZE = 0x1020 # type: ignore
+CL_DEVICE_MAX_CONSTANT_ARGS = 0x1021 # type: ignore
+CL_DEVICE_LOCAL_MEM_TYPE = 0x1022 # type: ignore
+CL_DEVICE_LOCAL_MEM_SIZE = 0x1023 # type: ignore
+CL_DEVICE_ERROR_CORRECTION_SUPPORT = 0x1024 # type: ignore
+CL_DEVICE_PROFILING_TIMER_RESOLUTION = 0x1025 # type: ignore
+CL_DEVICE_ENDIAN_LITTLE = 0x1026 # type: ignore
+CL_DEVICE_AVAILABLE = 0x1027 # type: ignore
+CL_DEVICE_COMPILER_AVAILABLE = 0x1028 # type: ignore
+CL_DEVICE_EXECUTION_CAPABILITIES = 0x1029 # type: ignore
+CL_DEVICE_QUEUE_PROPERTIES = 0x102A # type: ignore
+CL_DEVICE_QUEUE_ON_HOST_PROPERTIES = 0x102A # type: ignore
+CL_DEVICE_NAME = 0x102B # type: ignore
+CL_DEVICE_VENDOR = 0x102C # type: ignore
+CL_DRIVER_VERSION = 0x102D # type: ignore
+CL_DEVICE_PROFILE = 0x102E # type: ignore
+CL_DEVICE_VERSION = 0x102F # type: ignore
+CL_DEVICE_EXTENSIONS = 0x1030 # type: ignore
+CL_DEVICE_PLATFORM = 0x1031 # type: ignore
+CL_DEVICE_DOUBLE_FP_CONFIG = 0x1032 # type: ignore
+CL_DEVICE_PREFERRED_VECTOR_WIDTH_HALF = 0x1034 # type: ignore
+CL_DEVICE_HOST_UNIFIED_MEMORY = 0x1035 # type: ignore
+CL_DEVICE_NATIVE_VECTOR_WIDTH_CHAR = 0x1036 # type: ignore
+CL_DEVICE_NATIVE_VECTOR_WIDTH_SHORT = 0x1037 # type: ignore
+CL_DEVICE_NATIVE_VECTOR_WIDTH_INT = 0x1038 # type: ignore
+CL_DEVICE_NATIVE_VECTOR_WIDTH_LONG = 0x1039 # type: ignore
+CL_DEVICE_NATIVE_VECTOR_WIDTH_FLOAT = 0x103A # type: ignore
+CL_DEVICE_NATIVE_VECTOR_WIDTH_DOUBLE = 0x103B # type: ignore
+CL_DEVICE_NATIVE_VECTOR_WIDTH_HALF = 0x103C # type: ignore
+CL_DEVICE_OPENCL_C_VERSION = 0x103D # type: ignore
+CL_DEVICE_LINKER_AVAILABLE = 0x103E # type: ignore
+CL_DEVICE_BUILT_IN_KERNELS = 0x103F # type: ignore
+CL_DEVICE_IMAGE_MAX_BUFFER_SIZE = 0x1040 # type: ignore
+CL_DEVICE_IMAGE_MAX_ARRAY_SIZE = 0x1041 # type: ignore
+CL_DEVICE_PARENT_DEVICE = 0x1042 # type: ignore
+CL_DEVICE_PARTITION_MAX_SUB_DEVICES = 0x1043 # type: ignore
+CL_DEVICE_PARTITION_PROPERTIES = 0x1044 # type: ignore
+CL_DEVICE_PARTITION_AFFINITY_DOMAIN = 0x1045 # type: ignore
+CL_DEVICE_PARTITION_TYPE = 0x1046 # type: ignore
+CL_DEVICE_REFERENCE_COUNT = 0x1047 # type: ignore
+CL_DEVICE_PREFERRED_INTEROP_USER_SYNC = 0x1048 # type: ignore
+CL_DEVICE_PRINTF_BUFFER_SIZE = 0x1049 # type: ignore
+CL_DEVICE_IMAGE_PITCH_ALIGNMENT = 0x104A # type: ignore
+CL_DEVICE_IMAGE_BASE_ADDRESS_ALIGNMENT = 0x104B # type: ignore
+CL_DEVICE_MAX_READ_WRITE_IMAGE_ARGS = 0x104C # type: ignore
+CL_DEVICE_MAX_GLOBAL_VARIABLE_SIZE = 0x104D # type: ignore
+CL_DEVICE_QUEUE_ON_DEVICE_PROPERTIES = 0x104E # type: ignore
+CL_DEVICE_QUEUE_ON_DEVICE_PREFERRED_SIZE = 0x104F # type: ignore
+CL_DEVICE_QUEUE_ON_DEVICE_MAX_SIZE = 0x1050 # type: ignore
+CL_DEVICE_MAX_ON_DEVICE_QUEUES = 0x1051 # type: ignore
+CL_DEVICE_MAX_ON_DEVICE_EVENTS = 0x1052 # type: ignore
+CL_DEVICE_SVM_CAPABILITIES = 0x1053 # type: ignore
+CL_DEVICE_GLOBAL_VARIABLE_PREFERRED_TOTAL_SIZE = 0x1054 # type: ignore
+CL_DEVICE_MAX_PIPE_ARGS = 0x1055 # type: ignore
+CL_DEVICE_PIPE_MAX_ACTIVE_RESERVATIONS = 0x1056 # type: ignore
+CL_DEVICE_PIPE_MAX_PACKET_SIZE = 0x1057 # type: ignore
+CL_DEVICE_PREFERRED_PLATFORM_ATOMIC_ALIGNMENT = 0x1058 # type: ignore
+CL_DEVICE_PREFERRED_GLOBAL_ATOMIC_ALIGNMENT = 0x1059 # type: ignore
+CL_DEVICE_PREFERRED_LOCAL_ATOMIC_ALIGNMENT = 0x105A # type: ignore
+CL_DEVICE_IL_VERSION = 0x105B # type: ignore
+CL_DEVICE_MAX_NUM_SUB_GROUPS = 0x105C # type: ignore
+CL_DEVICE_SUB_GROUP_INDEPENDENT_FORWARD_PROGRESS = 0x105D # type: ignore
+CL_DEVICE_NUMERIC_VERSION = 0x105E # type: ignore
+CL_DEVICE_EXTENSIONS_WITH_VERSION = 0x1060 # type: ignore
+CL_DEVICE_ILS_WITH_VERSION = 0x1061 # type: ignore
+CL_DEVICE_BUILT_IN_KERNELS_WITH_VERSION = 0x1062 # type: ignore
+CL_DEVICE_ATOMIC_MEMORY_CAPABILITIES = 0x1063 # type: ignore
+CL_DEVICE_ATOMIC_FENCE_CAPABILITIES = 0x1064 # type: ignore
+CL_DEVICE_NON_UNIFORM_WORK_GROUP_SUPPORT = 0x1065 # type: ignore
+CL_DEVICE_OPENCL_C_ALL_VERSIONS = 0x1066 # type: ignore
+CL_DEVICE_PREFERRED_WORK_GROUP_SIZE_MULTIPLE = 0x1067 # type: ignore
+CL_DEVICE_WORK_GROUP_COLLECTIVE_FUNCTIONS_SUPPORT = 0x1068 # type: ignore
+CL_DEVICE_GENERIC_ADDRESS_SPACE_SUPPORT = 0x1069 # type: ignore
+CL_DEVICE_OPENCL_C_FEATURES = 0x106F # type: ignore
+CL_DEVICE_DEVICE_ENQUEUE_CAPABILITIES = 0x1070 # type: ignore
+CL_DEVICE_PIPE_SUPPORT = 0x1071 # type: ignore
+CL_DEVICE_LATEST_CONFORMANCE_VERSION_PASSED = 0x1072 # type: ignore
+CL_FP_DENORM = (1 << 0) # type: ignore
+CL_FP_INF_NAN = (1 << 1) # type: ignore
+CL_FP_ROUND_TO_NEAREST = (1 << 2) # type: ignore
+CL_FP_ROUND_TO_ZERO = (1 << 3) # type: ignore
+CL_FP_ROUND_TO_INF = (1 << 4) # type: ignore
+CL_FP_FMA = (1 << 5) # type: ignore
+CL_FP_SOFT_FLOAT = (1 << 6) # type: ignore
+CL_FP_CORRECTLY_ROUNDED_DIVIDE_SQRT = (1 << 7) # type: ignore
+CL_NONE = 0x0 # type: ignore
+CL_READ_ONLY_CACHE = 0x1 # type: ignore
+CL_READ_WRITE_CACHE = 0x2 # type: ignore
+CL_LOCAL = 0x1 # type: ignore
+CL_GLOBAL = 0x2 # type: ignore
+CL_EXEC_KERNEL = (1 << 0) # type: ignore
+CL_EXEC_NATIVE_KERNEL = (1 << 1) # type: ignore
+CL_QUEUE_OUT_OF_ORDER_EXEC_MODE_ENABLE = (1 << 0) # type: ignore
+CL_QUEUE_PROFILING_ENABLE = (1 << 1) # type: ignore
+CL_QUEUE_ON_DEVICE = (1 << 2) # type: ignore
+CL_QUEUE_ON_DEVICE_DEFAULT = (1 << 3) # type: ignore
+CL_CONTEXT_REFERENCE_COUNT = 0x1080 # type: ignore
+CL_CONTEXT_DEVICES = 0x1081 # type: ignore
+CL_CONTEXT_PROPERTIES = 0x1082 # type: ignore
+CL_CONTEXT_NUM_DEVICES = 0x1083 # type: ignore
+CL_CONTEXT_PLATFORM = 0x1084 # type: ignore
+CL_CONTEXT_INTEROP_USER_SYNC = 0x1085 # type: ignore
+CL_DEVICE_PARTITION_EQUALLY = 0x1086 # type: ignore
+CL_DEVICE_PARTITION_BY_COUNTS = 0x1087 # type: ignore
+CL_DEVICE_PARTITION_BY_COUNTS_LIST_END = 0x0 # type: ignore
+CL_DEVICE_PARTITION_BY_AFFINITY_DOMAIN = 0x1088 # type: ignore
+CL_DEVICE_AFFINITY_DOMAIN_NUMA = (1 << 0) # type: ignore
+CL_DEVICE_AFFINITY_DOMAIN_L4_CACHE = (1 << 1) # type: ignore
+CL_DEVICE_AFFINITY_DOMAIN_L3_CACHE = (1 << 2) # type: ignore
+CL_DEVICE_AFFINITY_DOMAIN_L2_CACHE = (1 << 3) # type: ignore
+CL_DEVICE_AFFINITY_DOMAIN_L1_CACHE = (1 << 4) # type: ignore
+CL_DEVICE_AFFINITY_DOMAIN_NEXT_PARTITIONABLE = (1 << 5) # type: ignore
+CL_DEVICE_SVM_COARSE_GRAIN_BUFFER = (1 << 0) # type: ignore
+CL_DEVICE_SVM_FINE_GRAIN_BUFFER = (1 << 1) # type: ignore
+CL_DEVICE_SVM_FINE_GRAIN_SYSTEM = (1 << 2) # type: ignore
+CL_DEVICE_SVM_ATOMICS = (1 << 3) # type: ignore
+CL_QUEUE_CONTEXT = 0x1090 # type: ignore
+CL_QUEUE_DEVICE = 0x1091 # type: ignore
+CL_QUEUE_REFERENCE_COUNT = 0x1092 # type: ignore
+CL_QUEUE_PROPERTIES = 0x1093 # type: ignore
+CL_QUEUE_SIZE = 0x1094 # type: ignore
+CL_QUEUE_DEVICE_DEFAULT = 0x1095 # type: ignore
+CL_QUEUE_PROPERTIES_ARRAY = 0x1098 # type: ignore
+CL_MEM_READ_WRITE = (1 << 0) # type: ignore
+CL_MEM_WRITE_ONLY = (1 << 1) # type: ignore
+CL_MEM_READ_ONLY = (1 << 2) # type: ignore
+CL_MEM_USE_HOST_PTR = (1 << 3) # type: ignore
+CL_MEM_ALLOC_HOST_PTR = (1 << 4) # type: ignore
+CL_MEM_COPY_HOST_PTR = (1 << 5) # type: ignore
+CL_MEM_HOST_WRITE_ONLY = (1 << 7) # type: ignore
+CL_MEM_HOST_READ_ONLY = (1 << 8) # type: ignore
+CL_MEM_HOST_NO_ACCESS = (1 << 9) # type: ignore
+CL_MEM_SVM_FINE_GRAIN_BUFFER = (1 << 10) # type: ignore
+CL_MEM_SVM_ATOMICS = (1 << 11) # type: ignore
+CL_MEM_KERNEL_READ_AND_WRITE = (1 << 12) # type: ignore
+CL_MIGRATE_MEM_OBJECT_HOST = (1 << 0) # type: ignore
+CL_MIGRATE_MEM_OBJECT_CONTENT_UNDEFINED = (1 << 1) # type: ignore
+CL_R = 0x10B0 # type: ignore
+CL_A = 0x10B1 # type: ignore
+CL_RG = 0x10B2 # type: ignore
+CL_RA = 0x10B3 # type: ignore
+CL_RGB = 0x10B4 # type: ignore
+CL_RGBA = 0x10B5 # type: ignore
+CL_BGRA = 0x10B6 # type: ignore
+CL_ARGB = 0x10B7 # type: ignore
+CL_INTENSITY = 0x10B8 # type: ignore
+CL_LUMINANCE = 0x10B9 # type: ignore
+CL_Rx = 0x10BA # type: ignore
+CL_RGx = 0x10BB # type: ignore
+CL_RGBx = 0x10BC # type: ignore
+CL_DEPTH = 0x10BD # type: ignore
+CL_sRGB = 0x10BF # type: ignore
+CL_sRGBx = 0x10C0 # type: ignore
+CL_sRGBA = 0x10C1 # type: ignore
+CL_sBGRA = 0x10C2 # type: ignore
+CL_ABGR = 0x10C3 # type: ignore
+CL_SNORM_INT8 = 0x10D0 # type: ignore
+CL_SNORM_INT16 = 0x10D1 # type: ignore
+CL_UNORM_INT8 = 0x10D2 # type: ignore
+CL_UNORM_INT16 = 0x10D3 # type: ignore
+CL_UNORM_SHORT_565 = 0x10D4 # type: ignore
+CL_UNORM_SHORT_555 = 0x10D5 # type: ignore
+CL_UNORM_INT_101010 = 0x10D6 # type: ignore
+CL_SIGNED_INT8 = 0x10D7 # type: ignore
+CL_SIGNED_INT16 = 0x10D8 # type: ignore
+CL_SIGNED_INT32 = 0x10D9 # type: ignore
+CL_UNSIGNED_INT8 = 0x10DA # type: ignore
+CL_UNSIGNED_INT16 = 0x10DB # type: ignore
+CL_UNSIGNED_INT32 = 0x10DC # type: ignore
+CL_HALF_FLOAT = 0x10DD # type: ignore
+CL_FLOAT = 0x10DE # type: ignore
+CL_UNORM_INT_101010_2 = 0x10E0 # type: ignore
+CL_MEM_OBJECT_BUFFER = 0x10F0 # type: ignore
+CL_MEM_OBJECT_IMAGE2D = 0x10F1 # type: ignore
+CL_MEM_OBJECT_IMAGE3D = 0x10F2 # type: ignore
+CL_MEM_OBJECT_IMAGE2D_ARRAY = 0x10F3 # type: ignore
+CL_MEM_OBJECT_IMAGE1D = 0x10F4 # type: ignore
+CL_MEM_OBJECT_IMAGE1D_ARRAY = 0x10F5 # type: ignore
+CL_MEM_OBJECT_IMAGE1D_BUFFER = 0x10F6 # type: ignore
+CL_MEM_OBJECT_PIPE = 0x10F7 # type: ignore
+CL_MEM_TYPE = 0x1100 # type: ignore
+CL_MEM_FLAGS = 0x1101 # type: ignore
+CL_MEM_SIZE = 0x1102 # type: ignore
+CL_MEM_HOST_PTR = 0x1103 # type: ignore
+CL_MEM_MAP_COUNT = 0x1104 # type: ignore
+CL_MEM_REFERENCE_COUNT = 0x1105 # type: ignore
+CL_MEM_CONTEXT = 0x1106 # type: ignore
+CL_MEM_ASSOCIATED_MEMOBJECT = 0x1107 # type: ignore
+CL_MEM_OFFSET = 0x1108 # type: ignore
+CL_MEM_USES_SVM_POINTER = 0x1109 # type: ignore
+CL_MEM_PROPERTIES = 0x110A # type: ignore
+CL_IMAGE_FORMAT = 0x1110 # type: ignore
+CL_IMAGE_ELEMENT_SIZE = 0x1111 # type: ignore
+CL_IMAGE_ROW_PITCH = 0x1112 # type: ignore
+CL_IMAGE_SLICE_PITCH = 0x1113 # type: ignore
+CL_IMAGE_WIDTH = 0x1114 # type: ignore
+CL_IMAGE_HEIGHT = 0x1115 # type: ignore
+CL_IMAGE_DEPTH = 0x1116 # type: ignore
+CL_IMAGE_ARRAY_SIZE = 0x1117 # type: ignore
+CL_IMAGE_BUFFER = 0x1118 # type: ignore
+CL_IMAGE_NUM_MIP_LEVELS = 0x1119 # type: ignore
+CL_IMAGE_NUM_SAMPLES = 0x111A # type: ignore
+CL_PIPE_PACKET_SIZE = 0x1120 # type: ignore
+CL_PIPE_MAX_PACKETS = 0x1121 # type: ignore
+CL_PIPE_PROPERTIES = 0x1122 # type: ignore
+CL_ADDRESS_NONE = 0x1130 # type: ignore
+CL_ADDRESS_CLAMP_TO_EDGE = 0x1131 # type: ignore
+CL_ADDRESS_CLAMP = 0x1132 # type: ignore
+CL_ADDRESS_REPEAT = 0x1133 # type: ignore
+CL_ADDRESS_MIRRORED_REPEAT = 0x1134 # type: ignore
+CL_FILTER_NEAREST = 0x1140 # type: ignore
+CL_FILTER_LINEAR = 0x1141 # type: ignore
+CL_SAMPLER_REFERENCE_COUNT = 0x1150 # type: ignore
+CL_SAMPLER_CONTEXT = 0x1151 # type: ignore
+CL_SAMPLER_NORMALIZED_COORDS = 0x1152 # type: ignore
+CL_SAMPLER_ADDRESSING_MODE = 0x1153 # type: ignore
+CL_SAMPLER_FILTER_MODE = 0x1154 # type: ignore
+CL_SAMPLER_MIP_FILTER_MODE = 0x1155 # type: ignore
+CL_SAMPLER_LOD_MIN = 0x1156 # type: ignore
+CL_SAMPLER_LOD_MAX = 0x1157 # type: ignore
+CL_SAMPLER_PROPERTIES = 0x1158 # type: ignore
+CL_MAP_READ = (1 << 0) # type: ignore
+CL_MAP_WRITE = (1 << 1) # type: ignore
+CL_MAP_WRITE_INVALIDATE_REGION = (1 << 2) # type: ignore
+CL_PROGRAM_REFERENCE_COUNT = 0x1160 # type: ignore
+CL_PROGRAM_CONTEXT = 0x1161 # type: ignore
+CL_PROGRAM_NUM_DEVICES = 0x1162 # type: ignore
+CL_PROGRAM_DEVICES = 0x1163 # type: ignore
+CL_PROGRAM_SOURCE = 0x1164 # type: ignore
+CL_PROGRAM_BINARY_SIZES = 0x1165 # type: ignore
+CL_PROGRAM_BINARIES = 0x1166 # type: ignore
+CL_PROGRAM_NUM_KERNELS = 0x1167 # type: ignore
+CL_PROGRAM_KERNEL_NAMES = 0x1168 # type: ignore
+CL_PROGRAM_IL = 0x1169 # type: ignore
+CL_PROGRAM_SCOPE_GLOBAL_CTORS_PRESENT = 0x116A # type: ignore
+CL_PROGRAM_SCOPE_GLOBAL_DTORS_PRESENT = 0x116B # type: ignore
+CL_PROGRAM_BUILD_STATUS = 0x1181 # type: ignore
+CL_PROGRAM_BUILD_OPTIONS = 0x1182 # type: ignore
+CL_PROGRAM_BUILD_LOG = 0x1183 # type: ignore
+CL_PROGRAM_BINARY_TYPE = 0x1184 # type: ignore
+CL_PROGRAM_BUILD_GLOBAL_VARIABLE_TOTAL_SIZE = 0x1185 # type: ignore
+CL_PROGRAM_BINARY_TYPE_NONE = 0x0 # type: ignore
+CL_PROGRAM_BINARY_TYPE_COMPILED_OBJECT = 0x1 # type: ignore
+CL_PROGRAM_BINARY_TYPE_LIBRARY = 0x2 # type: ignore
+CL_PROGRAM_BINARY_TYPE_EXECUTABLE = 0x4 # type: ignore
+CL_BUILD_SUCCESS = 0 # type: ignore
+CL_BUILD_NONE = -1 # type: ignore
+CL_BUILD_ERROR = -2 # type: ignore
+CL_BUILD_IN_PROGRESS = -3 # type: ignore
+CL_KERNEL_FUNCTION_NAME = 0x1190 # type: ignore
+CL_KERNEL_NUM_ARGS = 0x1191 # type: ignore
+CL_KERNEL_REFERENCE_COUNT = 0x1192 # type: ignore
+CL_KERNEL_CONTEXT = 0x1193 # type: ignore
+CL_KERNEL_PROGRAM = 0x1194 # type: ignore
+CL_KERNEL_ATTRIBUTES = 0x1195 # type: ignore
+CL_KERNEL_ARG_ADDRESS_QUALIFIER = 0x1196 # type: ignore
+CL_KERNEL_ARG_ACCESS_QUALIFIER = 0x1197 # type: ignore
+CL_KERNEL_ARG_TYPE_NAME = 0x1198 # type: ignore
+CL_KERNEL_ARG_TYPE_QUALIFIER = 0x1199 # type: ignore
+CL_KERNEL_ARG_NAME = 0x119A # type: ignore
+CL_KERNEL_ARG_ADDRESS_GLOBAL = 0x119B # type: ignore
+CL_KERNEL_ARG_ADDRESS_LOCAL = 0x119C # type: ignore
+CL_KERNEL_ARG_ADDRESS_CONSTANT = 0x119D # type: ignore
+CL_KERNEL_ARG_ADDRESS_PRIVATE = 0x119E # type: ignore
+CL_KERNEL_ARG_ACCESS_READ_ONLY = 0x11A0 # type: ignore
+CL_KERNEL_ARG_ACCESS_WRITE_ONLY = 0x11A1 # type: ignore
+CL_KERNEL_ARG_ACCESS_READ_WRITE = 0x11A2 # type: ignore
+CL_KERNEL_ARG_ACCESS_NONE = 0x11A3 # type: ignore
+CL_KERNEL_ARG_TYPE_NONE = 0 # type: ignore
+CL_KERNEL_ARG_TYPE_CONST = (1 << 0) # type: ignore
+CL_KERNEL_ARG_TYPE_RESTRICT = (1 << 1) # type: ignore
+CL_KERNEL_ARG_TYPE_VOLATILE = (1 << 2) # type: ignore
+CL_KERNEL_ARG_TYPE_PIPE = (1 << 3) # type: ignore
+CL_KERNEL_WORK_GROUP_SIZE = 0x11B0 # type: ignore
+CL_KERNEL_COMPILE_WORK_GROUP_SIZE = 0x11B1 # type: ignore
+CL_KERNEL_LOCAL_MEM_SIZE = 0x11B2 # type: ignore
+CL_KERNEL_PREFERRED_WORK_GROUP_SIZE_MULTIPLE = 0x11B3 # type: ignore
+CL_KERNEL_PRIVATE_MEM_SIZE = 0x11B4 # type: ignore
+CL_KERNEL_GLOBAL_WORK_SIZE = 0x11B5 # type: ignore
+CL_KERNEL_MAX_SUB_GROUP_SIZE_FOR_NDRANGE = 0x2033 # type: ignore
+CL_KERNEL_SUB_GROUP_COUNT_FOR_NDRANGE = 0x2034 # type: ignore
+CL_KERNEL_LOCAL_SIZE_FOR_SUB_GROUP_COUNT = 0x11B8 # type: ignore
+CL_KERNEL_MAX_NUM_SUB_GROUPS = 0x11B9 # type: ignore
+CL_KERNEL_COMPILE_NUM_SUB_GROUPS = 0x11BA # type: ignore
+CL_KERNEL_EXEC_INFO_SVM_PTRS = 0x11B6 # type: ignore
+CL_KERNEL_EXEC_INFO_SVM_FINE_GRAIN_SYSTEM = 0x11B7 # type: ignore
+CL_EVENT_COMMAND_QUEUE = 0x11D0 # type: ignore
+CL_EVENT_COMMAND_TYPE = 0x11D1 # type: ignore
+CL_EVENT_REFERENCE_COUNT = 0x11D2 # type: ignore
+CL_EVENT_COMMAND_EXECUTION_STATUS = 0x11D3 # type: ignore
+CL_EVENT_CONTEXT = 0x11D4 # type: ignore
+CL_COMMAND_NDRANGE_KERNEL = 0x11F0 # type: ignore
+CL_COMMAND_TASK = 0x11F1 # type: ignore
+CL_COMMAND_NATIVE_KERNEL = 0x11F2 # type: ignore
+CL_COMMAND_READ_BUFFER = 0x11F3 # type: ignore
+CL_COMMAND_WRITE_BUFFER = 0x11F4 # type: ignore
+CL_COMMAND_COPY_BUFFER = 0x11F5 # type: ignore
+CL_COMMAND_READ_IMAGE = 0x11F6 # type: ignore
+CL_COMMAND_WRITE_IMAGE = 0x11F7 # type: ignore
+CL_COMMAND_COPY_IMAGE = 0x11F8 # type: ignore
+CL_COMMAND_COPY_IMAGE_TO_BUFFER = 0x11F9 # type: ignore
+CL_COMMAND_COPY_BUFFER_TO_IMAGE = 0x11FA # type: ignore
+CL_COMMAND_MAP_BUFFER = 0x11FB # type: ignore
+CL_COMMAND_MAP_IMAGE = 0x11FC # type: ignore
+CL_COMMAND_UNMAP_MEM_OBJECT = 0x11FD # type: ignore
+CL_COMMAND_MARKER = 0x11FE # type: ignore
+CL_COMMAND_ACQUIRE_GL_OBJECTS = 0x11FF # type: ignore
+CL_COMMAND_RELEASE_GL_OBJECTS = 0x1200 # type: ignore
+CL_COMMAND_READ_BUFFER_RECT = 0x1201 # type: ignore
+CL_COMMAND_WRITE_BUFFER_RECT = 0x1202 # type: ignore
+CL_COMMAND_COPY_BUFFER_RECT = 0x1203 # type: ignore
+CL_COMMAND_USER = 0x1204 # type: ignore
+CL_COMMAND_BARRIER = 0x1205 # type: ignore
+CL_COMMAND_MIGRATE_MEM_OBJECTS = 0x1206 # type: ignore
+CL_COMMAND_FILL_BUFFER = 0x1207 # type: ignore
+CL_COMMAND_FILL_IMAGE = 0x1208 # type: ignore
+CL_COMMAND_SVM_FREE = 0x1209 # type: ignore
+CL_COMMAND_SVM_MEMCPY = 0x120A # type: ignore
+CL_COMMAND_SVM_MEMFILL = 0x120B # type: ignore
+CL_COMMAND_SVM_MAP = 0x120C # type: ignore
+CL_COMMAND_SVM_UNMAP = 0x120D # type: ignore
+CL_COMMAND_SVM_MIGRATE_MEM = 0x120E # type: ignore
+CL_COMPLETE = 0x0 # type: ignore
+CL_RUNNING = 0x1 # type: ignore
+CL_SUBMITTED = 0x2 # type: ignore
+CL_QUEUED = 0x3 # type: ignore
+CL_BUFFER_CREATE_TYPE_REGION = 0x1220 # type: ignore
+CL_PROFILING_COMMAND_QUEUED = 0x1280 # type: ignore
+CL_PROFILING_COMMAND_SUBMIT = 0x1281 # type: ignore
+CL_PROFILING_COMMAND_START = 0x1282 # type: ignore
+CL_PROFILING_COMMAND_END = 0x1283 # type: ignore
+CL_PROFILING_COMMAND_COMPLETE = 0x1284 # type: ignore
+CL_DEVICE_ATOMIC_ORDER_RELAXED = (1 << 0) # type: ignore
+CL_DEVICE_ATOMIC_ORDER_ACQ_REL = (1 << 1) # type: ignore
+CL_DEVICE_ATOMIC_ORDER_SEQ_CST = (1 << 2) # type: ignore
+CL_DEVICE_ATOMIC_SCOPE_WORK_ITEM = (1 << 3) # type: ignore
+CL_DEVICE_ATOMIC_SCOPE_WORK_GROUP = (1 << 4) # type: ignore
+CL_DEVICE_ATOMIC_SCOPE_DEVICE = (1 << 5) # type: ignore
+CL_DEVICE_ATOMIC_SCOPE_ALL_DEVICES = (1 << 6) # type: ignore
+CL_DEVICE_QUEUE_SUPPORTED = (1 << 0) # type: ignore
+CL_DEVICE_QUEUE_REPLACEABLE_DEFAULT = (1 << 1) # type: ignore
+CL_KHRONOS_VENDOR_ID_CODEPLAY = 0x10004 # type: ignore
+CL_VERSION_MAJOR_BITS = (10) # type: ignore
+CL_VERSION_MINOR_BITS = (10) # type: ignore
+CL_VERSION_PATCH_BITS = (12) # type: ignore
+CL_VERSION_MAJOR_MASK = ((1 << CL_VERSION_MAJOR_BITS) - 1) # type: ignore
+CL_VERSION_MINOR_MASK = ((1 << CL_VERSION_MINOR_BITS) - 1) # type: ignore
+CL_VERSION_PATCH_MASK = ((1 << CL_VERSION_PATCH_BITS) - 1) # type: ignore
 CL_VERSION_MAJOR = lambda version: ((version) >> (CL_VERSION_MINOR_BITS + CL_VERSION_PATCH_BITS)) # type: ignore
 CL_VERSION_MINOR = lambda version: (((version) >> CL_VERSION_PATCH_BITS) & CL_VERSION_MINOR_MASK) # type: ignore
 CL_VERSION_PATCH = lambda version: ((version) & CL_VERSION_PATCH_MASK) # type: ignore