#pragma once
#include <cuda_bf16.h>
#include <cuda_fp16.h>
#include <cuda_fp8.h>
#include <cstdint>

using state_t = {{ state_dtype }};
using input_t = {{ input_dtype }};
// dt accepted in its native dtype (e.g. bf16) — converted to f32 internally.
// Eliminates the need for a separate dtype conversion kernel launch on the host side.
using dt_t = {{ dt_dtype }};
using weight_t = {{ weight_dtype }};
using matrixA_t = {{ matrixA_dtype }};
using stateIndex_t = {{ stateIndex_dtype }};
// Type for block-scale decode factors (e.g. float, __half).
// void = no scaling (state_t is used as-is).
using state_scale_t = {{ state_scale_type }};

constexpr int DIM = {{ dim }};
constexpr int DSTATE = {{ dstate }};
constexpr int NPREDICTED = {{ npredicted }};
constexpr int MAX_WINDOW = {{ max_window }};
// nheads / ngroups — JIT-stamped so the kernel compiles only one
// (HEADS_PER_GROUP, NGROUPS) specialization per .so (was 7 via `dispatchRatio`).
// The wrapper computes these from the runtime tensors and selects the matching
// JIT URI; the launcher reads them directly without a runtime dispatch step.
// NGROUPS is compile-time so the persistent main's unflatten (NHEADS = NGROUPS *
// HEADS_PER_GROUP) divides by compile-time constants only — no runtime div/mod.
constexpr int HEADS_PER_GROUP = {{ heads_per_group }};
constexpr int NGROUPS = {{ num_groups }};
// Philox PRNG rounds for stochastic rounding of fp16 state stores.
// 0 = no stochastic rounding; typical value = 10.
constexpr int PHILOX_ROUNDS = {{ philox_rounds }};
// Programmatic Dependent Launch.  When true, the kernel emits the
// griddepcontrol.{wait,launch_dependents} PTX and the load is split around
// `gdc_wait` for cache-load-during-wait overlap.  When false, a single-pass
// load_data path is used (no PDL barriers) — matches v21.0 register profile
// and load order.  Stamped as a JIT URI key so each (enable_pdl) value
// compiles its own .so; no runtime branch in the kernel binary.
constexpr bool ENABLE_PDL = {{ enable_pdl }};
