ember_ir Namespace Reference

Namespaces
namespace	cpu
namespace	expr
namespace	rec

Classes
struct	DigestResult
class	Interpreter
struct	Node

Functions
void	linear_attention_row_cpu (const float *q, const float *k, const float *v, float *num, float *den, float *out, int width, float den_floor)
void	ssm_scan_row_cpu (const float *u, const float *A, const float *B, const float *C, float D, float *s, float *scratch, float *y, int n)
bool	has_recurrent_state (const std::string &ir_json)
bool	has_windowed_attention (const std::string &ir_json)
bool	has_state_tokens (const std::string &ir_json)
bool	tuned_llama_match (const std::string &ir_json, const GGUFData &gguf, std::string *why)
bool	device_supported (const std::string &ir_json)
bool	has_ir_builder (const std::string &arch)
std::string	find_ir (const GGUFData &gguf)
std::string	build_ir_from_gguf (const GGUFData &g, bool log)
static void	cpu_only (bool gpu, const char *what)
void	ir_rows_gather_cuda (const float *src, const int *idx, float *dst, int n, int w, cudaStream_t stream)
void	ir_rows_scatter_cuda (const float *src, const int *idx, float *dst, int n, int w, cudaStream_t stream)
void	ir_mor_gate_cuda (const float *x_in, float *x, const float *g, int n, int w, cudaStream_t stream)
void	ir_rope_at_cuda (float *q, float *k, int T, int hq, int hk, int hd, const int *pos, const float *freqs, cudaStream_t stream)
void	ir_rope_rows_at_cuda (float *x, int T, int heads, int hd, int n_rot, bool neox, const int *pos, const float *freqs, cudaStream_t stream)
void	ir_turn_state_attention_cuda (const float *q, const float *K, const float *V, float *out, int T, int row0, const int *codes, int hq, int hkv, int hd, float scale, cudaStream_t stream, const int *drow)
void	ir_kv_compact_cuda (float *K, float *V, const int *idx, int w0, int n, int row, cudaStream_t stream)
void	ir_embedding_state_cuda (const float *table, const float *state, int vocab, const int *tokens, float *out, int T, int d, cudaStream_t stream)
void	ir_window_attention_cuda (const float *q, const float *K, const float *V, float *out, int T, int start_pos, int hq, int hkv, int hd, float scale, int window, int sinks, cudaStream_t stream, const int *dpos, float softcap)
void	ir_kv_write_cuda (float *K, float *V, const float *k, const float *v, int T, int row, int start_pos, cudaStream_t stream, const int *dpos)
void	ir_relu_cuda (const float *x, float *out, size_t n, cudaStream_t stream)
void	ir_sigmoid_cuda (const float *x, float *out, size_t n, cudaStream_t stream)
void	ir_affine_cuda (const float *x, float *out, size_t n, float scale, float shift, cudaStream_t stream)
void	ir_mul_col_cuda (const float *x, const float *c, float *out, int T, int w, cudaStream_t stream)
void	ir_scale_cuda (const float *x, const float *g, const float *b, float *out, int T, int w, cudaStream_t stream)
void	ir_rms_row_cuda (const float *x, float *out, int T, int w, float eps, cudaStream_t stream)
void	ir_depth_attend_cuda (const float *const *sources, int S, const float *w, float *out, int T, int d, float eps, cudaStream_t stream)
void	ir_linear_attention_cuda (const float *q, const float *k, const float *v, float *num, float *den, float *out, int T, int width, float den_floor, cudaStream_t stream)
void	ir_ssm_scan_cuda (const float *u, const float *A, const float *B, const float *C, const float *D, float *s, float *y, int T, int n, cudaStream_t stream)
void	ir_embedding_cuda (const float *table, const int *tokens, float *out, int T, int d, cudaStream_t stream)
void	ir_pos_embedding_cuda (const float *x, const float *P, float *out, int T, int d, int start_pos, cudaStream_t stream, const int *dpos)
void	ir_layernorm_cuda (const float *x, const float *w, const float *b, float *out, int T, int d, float eps, cudaStream_t stream)
void	ir_bias_add_cuda (float *y, const float *b, int T, int O, cudaStream_t stream)
void	ir_binary_cuda (const float *a, const float *b, float *out, size_t n, int op, cudaStream_t stream)
void	ir_silu_cuda (const float *x, float *out, size_t n, cudaStream_t stream)
void	ir_gelu_cuda (const float *x, float *out, size_t n, bool tanh_form, cudaStream_t stream)
void	ir_rope_cuda (float *q, float *k, int T, int hq, int hk, int hd, int start_pos, const float *freqs, cudaStream_t stream, const int *dpos)
void	ir_rope_rows_cuda (float *x, int T, int heads, int hd, int n_rot, bool neox, int start_pos, const float *freqs, cudaStream_t stream, const int *dpos)
void	ir_conv1d_causal_cuda (const float *x, const float *w, const float *bias, float *state, float *out, int T, int C, int K, cudaStream_t stream)
void	ir_selective_scan_cuda (const float *x, const float *dt, const float *B, const float *C, const float *A, const float *d, float *s, float *y, int T, int D, int N, bool softplus, cudaStream_t stream)
size_t	ir_expr_scratch_floats (const std::vector< expr::Stage > &stages, int T, int w)
void	ir_expr_cuda (const std::vector< expr::Stage > &stages, const float *const *ins, int n_ins, float *scratch, float *out, int T, int w, cudaStream_t stream)
void	ir_mla_kv_write_cuda (float *KV, const float *c, const float *kr, int T, int r, int dr, int start_pos, cudaStream_t stream, const int *dpos)
void	ir_mla_attention_cuda (const float *qn, const float *qr, const float *KV, const float *Kb, const float *Vb, float *out, int T, int start_pos, int H, int r, int dn, int dr, int dv, float scale, int window, int sinks, cudaStream_t stream, const int *dpos)
json	parse_ir_document (const std::string &text)
std::vector< std::string >	ir_schema_errors (const json &doc)
std::vector< std::string >	ir_document_errors (const json &doc)
std::vector< std::string >	ir_container_errors (const json &doc, const GGUFData &gguf)
void	validate_ir (const json &doc, const GGUFData &gguf)
std::string	sha256_hex_bytes (const std::string &bytes)
std::string	es_number_string (double v)
DigestResult	ir_digest (const std::string &text, bool prose)

Variables
const char *const	kIrSchemaPieces []
const char *const	patternProperties

Function Documentation

build_ir_from_gguf()

std::string ember_ir::build_ir_from_gguf	(	const GGUFData &	g,
		bool	log
	)

Definition at line 553 of file graph_ir.cpp.

                                                          {
  const std::string arch = meta_string(g, "general.architecture").value_or("");
  if (!has_ir_builder(arch)) throw std::runtime_error("no IR builder for general.architecture = '" + arch + "'");
  const std::string p = arch + ".";
  auto get = [&](const std::string& key, std::optional<double> dflt = std::nullopt) -> double {
    auto v = meta_number(g, p + key);
    if (!v) v = dflt;
    if (!v) throw std::runtime_error("IR builder: " + arch + " file has no " + p + key);
    return *v;
  };
  auto shape_has = [&](const std::string& t) { return g.tensor_infos_map.count(t) > 0; };
  auto no_rope_scaling = [&]() {
    const auto scaling = meta_string(g, p + "rope.scaling.type");
    if ((scaling && *scaling != "none") || shape_has("rope_freqs.weight"))
      throw std::runtime_error("IR builder: " + arch + " with RoPE scaling is not built yet");
  };
  const int d = static_cast<int>(get("embedding_length"));
  const int n_layers = static_cast<int>(get("block_count"));
  int vocab = static_cast<int>(g.tokenizer_tokens.size());
  if (vocab == 0) {
    auto it = g.tensor_infos_map.find("token_embd.weight");
    if (it == g.tensor_infos_map.end() || it->second.shape.size() < 2)
      throw std::runtime_error("IR builder: no tokens and no token_embd.weight");
    vocab = static_cast<int>(it->second.shape[1]);
  }
  const std::string B = "blk.{i}.";
 
  auto linear = [](const std::string& x, const std::string& w, const std::string& out, const std::string& b = "",
                   const std::string& tied = "") {
    json n = {{"op", "linear"}, {"in", {x}}, {"weight", w}, {"bias", nullptr}, {"out", out}};
    if (!b.empty()) n["bias"] = b;
    if (!tied.empty()) n["tied_to"] = tied;
    return n;
  };
  auto rmsnorm = [](const std::string& x, const std::string& w, double eps, const std::string& out, int group = 0) {
    json n = {{"op", "rmsnorm"}, {"in", {x}}, {"weight", w}, {"eps", f32_value(eps)}, {"out", out}};
    if (group) n["group"] = group;
    return n;
  };
  auto layernorm = [](const std::string& x, const std::string& name, double eps, const std::string& out) {
    return json{{"op", "layernorm"}, {"in", {x}},           {"weight", name + ".weight"},
                {"bias", name + ".bias"}, {"eps", f32_value(eps)}, {"out", out}};
  };
  auto op = [](const std::string& name, std::vector<std::string> in, const json& out, json attrs = json::object()) {
    attrs["op"] = name;
    attrs["in"] = in;
    attrs["out"] = out;
    return attrs;
  };
  auto rope = [](int H, int KV, int hd, int n_rot, const json& theta, const std::string& pairing,
                 const json& freq_scale = nullptr) {
    json n = {{"op", "rope"},     {"in", {"q", "k", "pos"}}, {"heads", H},          {"kv_heads", KV},
              {"head_dim", hd},   {"n_rot", n_rot},          {"theta", theta},      {"rope_heads", "all"},
              {"pairing", pairing}, {"out", {"q", "k"}}};
    if (!freq_scale.is_null()) n["freq_scale"] = freq_scale;
    return n;
  };
  auto attention = [](int H, int KV, int hd, double scale, const json& window = 0) {
    return json{{"op", "attention"}, {"in", {"q", "k", "v"}}, {"state", "kv"},     {"heads", H},
                {"kv_heads", KV},    {"head_dim", hd},        {"causal", true},    {"scale", f32_value(scale)},
                {"window", window},  {"sinks", 0},            {"softmax", "max-shifted, masked entries exactly zero"},
                {"out", "a"}};
  };
  auto kv_state = [](int H, int KV, int hd, int ctx, const json& window = 0) {
    return json{{"name", "kv"},     {"kind", "kv"}, {"heads", H}, {"kv_heads", KV}, {"head_dim", hd},
                {"max_len", ctx},   {"window", window}, {"sinks", 0}, {"dtype", "f32"}};
  };
  const bool tied = !shape_has("output.weight");
  auto tied_head = [&](json norm_nodes) {
    norm_nodes.push_back(linear("h", "output.weight", "logits", "", tied ? "token_embd.weight" : ""));
    return norm_nodes;
  };
  json embed = json::array({{{"op", "embedding"}, {"in", {"tokens"}}, {"weight", "token_embd.weight"}, {"out", "x"}}});
  json state, nodes, head;
  int ctx = 0;
  std::string summary;
 
  // A truth value (IR_SPEC D.1): a bool, or a number that is true when non-zero.
  auto truth = [&](const std::string& key, bool dflt) {
    if (const auto b = meta_bool(g, p + key)) return *b;
    if (const auto v = meta_number(g, p + key)) return *v != 0.0;
    return dflt;
  };
  // moe_ffn in llama.cpp's expert layout (IR_SPEC D.3, D.10), its gates normalised as gate_norm says.
  auto moe_node = [&](const std::string& gate_norm) {
    return op("moe_ffn", {"h2"}, "f",
              {{"router", B + "ffn_gate_inp.weight"},
               {"experts", {{"gate", B + "ffn_gate_exps.weight"}, {"up", B + "ffn_up_exps.weight"},
                            {"down", B + "ffn_down_exps.weight"}}},
               {"n_experts", static_cast<int>(get("expert_count"))},
               {"top_k", static_cast<int>(get("expert_used_count"))},
               {"gate_norm", gate_norm},
               {"activation", "silu"}});
  };
  // The shared expert's SwiGLU over ffn_{gate,up,down}_shexp (IR_SPEC D.10, D.11) -> sf.
  auto shared_expert = [&](json& nodes) {
    for (json n : {linear("h2", B + "ffn_gate_shexp.weight", "sg"), linear("h2", B + "ffn_up_shexp.weight", "su"),
                   op("silu", {"sg"}, "sg"), op("mul", {"sg", "su"}, "sgu"),
                   linear("sgu", B + "ffn_down_shexp.weight", "sf")})
      nodes.push_back(n);
  };
  const bool qwen_moe = arch == "qwen2moe" || arch == "qwen3moe";
 
  if (arch == "llama" || arch == "mistral" || arch == "qwen2" || arch == "qwen3" || qwen_moe) {
    const int H = static_cast<int>(get("attention.head_count"));
    const int KV = static_cast<int>(get("attention.head_count_kv", H));
    const int hd = static_cast<int>(get("attention.key_length", d / H));
    const int n_rot = static_cast<int>(get("rope.dimension_count", hd));
    ctx = static_cast<int>(get("context_length"));
    const double eps = get("attention.layer_norm_rms_epsilon");
    const double theta = get("rope.freq_base", 10000.0);
    no_rope_scaling();
    // HF Mistral's window (key j visible from query p iff p - w < j <= p); qwen2's
    // GGUF does not say which layers a window would cover, so none.
    int window = 0;
    if (arch == "llama" || arch == "mistral")
      window = static_cast<int>(meta_number(g, p + "attention.sliding_window").value_or(0));
    const bool bias = shape_has("blk.0.attn_q.bias"), qk_norm = shape_has("blk.0.attn_q_norm.weight");
    const std::string pairing = arch == "qwen2" || arch == "qwen3" || qwen_moe ? "neox" : "interleaved";
    auto proj = [&](const std::string& name, const std::string& out) {
      return linear("h", B + name + ".weight", out, bias ? B + name + ".bias" : "");
    };
    // Mixtral's layout (llama.cpp's llama graph; IR_SPEC 9.3, revision 1.3): llama.expert_count and ffn_gate_inp
    // put a moe_ffn in place of the SwiGLU.
    const bool moe = arch == "llama" && meta_number(g, p + "expert_count").value_or(0) > 0 &&
                     shape_has("blk.0.ffn_gate_inp.weight");
    nodes = json::array({rmsnorm("x", B + "attn_norm.weight", eps, "h"), proj("attn_q", "q"), proj("attn_k", "k"),
                         proj("attn_v", "v")});
    if (qk_norm) {
      nodes.push_back(rmsnorm("q", B + "attn_q_norm.weight", eps, "q", hd));
      nodes.push_back(rmsnorm("k", B + "attn_k_norm.weight", eps, "k", hd));
    }
    for (json n : {rope(H, KV, hd, n_rot, f32_value(theta), pairing),
                   attention(H, KV, hd, 1.0 / std::sqrt(static_cast<double>(hd)), window),
                   linear("a", B + "attn_output.weight", "o"), op("add", {"x", "o"}, "x"),
                   rmsnorm("x", B + "ffn_norm.weight", eps, "h2")})
      nodes.push_back(n);
    // Qwen-MoE (IR_SPEC D.10, revision 1.5): qwen2's / qwen3's attention, moe_ffn with gate_norm raw for qwen2moe
    // (llama.cpp b6000's graph ignores norm_topk_prob) and renorm for qwen3moe; qwen2moe's shared expert times
    // sigmoid(ffn_gate_inp_shexp h2) per token, added to f.
    const bool shexp = arch == "qwen2moe" && shape_has("blk.0.ffn_gate_inp_shexp.weight");
    if (moe) {
      nodes.push_back(moe_node("renorm"));
    } else if (qwen_moe) {
      nodes.push_back(moe_node(arch == "qwen2moe" ? "raw" : "renorm"));
      if (shexp) {
        shared_expert(nodes);
        for (json n : {linear("h2", B + "ffn_gate_inp_shexp.weight", "sgate"), op("sigmoid", {"sgate"}, "sgate"),
                       op("mul_col", {"sf", "sgate"}, "sf"), op("add", {"f", "sf"}, "f")})
          nodes.push_back(n);
      }
    } else {
      for (json n : {linear("h2", B + "ffn_gate.weight", "g"), linear("h2", B + "ffn_up.weight", "u"),
                     op("silu", {"g"}, "g"), op("mul", {"g", "u"}, "gu"), linear("gu", B + "ffn_down.weight", "f")})
        nodes.push_back(n);
    }
    nodes.push_back(op("add", {"x", "f"}, "x"));
    head = tied_head(json::array({rmsnorm("x", "output_norm.weight", eps, "h")}));
    state = json::array({kv_state(H, KV, hd, ctx, window)});
    summary = "heads " + std::to_string(H) + "/" + std::to_string(KV) + " x " + std::to_string(hd) + ", n_rot " +
              std::to_string(n_rot) + ", " + pairing + " RoPE" + (window ? ", window " + std::to_string(window) : "") +
              (bias ? ", q/k/v biases" : "") + (qk_norm ? ", per-head q/k RMSNorm" : "") +
              (moe || qwen_moe ? ", moe_ffn " + std::to_string(static_cast<int>(get("expert_count"))) + " experts" : "") +
              (shexp ? " + a gated shared expert" : "");
  } else if (arch == "gpt2" || arch == "gptneox") {
    // LayerNorm blocks over a fused, biased qkv (q | k | v as llama.cpp's convert
    // writes it; for gptneox it has already de-interleaved HF's per-head [q k v] rows).
    const bool neox = arch == "gptneox";
    const int H = static_cast<int>(get("attention.head_count"));
    const int KV = static_cast<int>(get("attention.head_count_kv", H));
    const int hd = static_cast<int>(get("attention.key_length", d / H));
    ctx = static_cast<int>(get("context_length"));
    const double eps = get("attention.layer_norm_epsilon");
    if (!neox)
      embed.push_back({{"op", "pos_embedding"}, {"in", {"x", "pos"}}, {"weight", "position_embd.weight"}, {"out", "x"}});
    nodes = json::array({layernorm("x", B + "attn_norm", eps, "h"),
                         linear("h", B + "attn_qkv.weight", "qkv", B + "attn_qkv.bias"),
                         op("split", {"qkv"}, {"q", "k", "v"}, {{"sizes", {H * hd, KV * hd, KV * hd}}})});
    int n_rot = 0;
    if (neox) {
      no_rope_scaling();
      n_rot = static_cast<int>(get("rope.dimension_count", hd));
      nodes.push_back(rope(H, KV, hd, n_rot, f32_value(get("rope.freq_base", 10000.0)), "neox"));
    }
    nodes.push_back(attention(H, KV, hd, 1.0 / std::sqrt(static_cast<double>(hd))));
    nodes.push_back(linear("a", B + "attn_output.weight", "o", B + "attn_output.bias"));
    const json ffn = json::array({layernorm("x", B + "ffn_norm", eps, "h2"),
                                  linear("h2", B + "ffn_up.weight", "u", B + "ffn_up.bias"),
                                  op("gelu", {"u"}, "u", {{"approx", neox ? "none" : "tanh"}}),
                                  linear("u", B + "ffn_down.weight", "f", B + "ffn_down.bias")});
    const bool parallel = neox && meta_bool(g, p + "use_parallel_residual").value_or(true);
    if (parallel) {
      // llama.cpp: cur = ffn(ln2 x) + x, then + attn_out; HF sums mlp + attn + x (ROUNDING apart).
      for (const auto& n : ffn) nodes.push_back(n);
      nodes.push_back(op("add", {"f", "x"}, "x"));
      nodes.push_back(op("add", {"x", "o"}, "x"));
    } else {
      nodes.push_back(op("add", {"x", "o"}, "x"));
      for (const auto& n : ffn) nodes.push_back(n);
      nodes.push_back(op("add", {"x", "f"}, "x"));
    }
    head = json::array({layernorm("x", "output_norm", eps, "h")});
    if (neox)
      head.push_back(linear("h", "output.weight", "logits"));
    else
      head = tied_head(head);
    state = json::array({kv_state(H, KV, hd, ctx)});
    summary = "heads " + std::to_string(H) + "/" + std::to_string(KV) + " x " + std::to_string(hd) + ", LayerNorm" +
              (neox ? ", NEOX RoPE n_rot " + std::to_string(n_rot) + (parallel ? ", parallel residual" : "")
                    : std::string(", learned positions"));
  } else if (arch == "gemma3") {
    // llama.cpp b6000 hparams.set_swa_pattern(6): layer i is local (windowed,
    // theta 10000) iff i % 6 < 5; the rest are global (the file's theta, linear scaling).
    constexpr int kSwaPattern = 6;
    constexpr double kLocalTheta = 10000.0;
    const int H = static_cast<int>(get("attention.head_count"));
    const int KV = static_cast<int>(get("attention.head_count_kv", H));
    const int hd = static_cast<int>(get("attention.key_length", d / H));
    const int n_rot = static_cast<int>(get("rope.dimension_count", hd));
    ctx = static_cast<int>(get("context_length"));
    const double eps = get("attention.layer_norm_rms_epsilon");
    if (shape_has("rope_freqs.weight")) throw std::runtime_error("IR builder: gemma3 with rope_freqs is not built yet");
    const auto stype = meta_string(g, p + "rope.scaling.type");
    if (stype && *stype != "none" && *stype != "linear")
      throw std::runtime_error("IR builder: gemma3 RoPE scaling '" + *stype + "' is not built yet");
    double factor = 1.0;
    if (stype && *stype == "linear") {
      factor = meta_number(g, p + "rope.scaling.factor").value_or(1.0);
      if (factor == 0.0) factor = 1.0;
    }
    const int win = static_cast<int>(get("attention.sliding_window"));
    const double theta_global = get("rope.freq_base");
    json windows = json::array(), thetas = json::array(), scales = json::array();
    for (int i = 0; i < n_layers; ++i) {
      const bool local = i % kSwaPattern < kSwaPattern - 1;
      windows.push_back(local ? win : 0);
      thetas.push_back(f32_value(local ? kLocalTheta : theta_global));
      scales.push_back(local ? 1.0 : f32_value(1.0 / factor));
    }
    windows = {{"per_layer", windows}};
    thetas = {{"per_layer", thetas}};
    scales = {{"per_layer", scales}};
    // llama.cpp: f_attention_scale = 1/sqrt(n_embd_head_k), except 1/sqrt(n_embd / n_head) for the 27B.
    const double scale = n_layers == 62 ? 1.0 / std::sqrt(static_cast<double>(d) / H)
                                        : 1.0 / std::sqrt(static_cast<double>(hd));
    embed.push_back(op("affine", {"x"}, "x", {{"scale", f32_value(std::sqrt(static_cast<double>(d)))}, {"shift", 0.0}}));
    nodes = json::array({rmsnorm("x", B + "attn_norm.weight", eps, "h"), linear("h", B + "attn_q.weight", "q"),
                         linear("h", B + "attn_k.weight", "k"), linear("h", B + "attn_v.weight", "v"),
                         rmsnorm("q", B + "attn_q_norm.weight", eps, "q", hd),
                         rmsnorm("k", B + "attn_k_norm.weight", eps, "k", hd),
                         rope(H, KV, hd, n_rot, thetas, "neox", factor != 1.0 ? scales : json(nullptr)),
                         attention(H, KV, hd, scale, windows), linear("a", B + "attn_output.weight", "o"),
                         rmsnorm("o", B + "post_attention_norm.weight", eps, "o"), op("add", {"x", "o"}, "x"),
                         rmsnorm("x", B + "ffn_norm.weight", eps, "h2"), linear("h2", B + "ffn_gate.weight", "g"),
                         linear("h2", B + "ffn_up.weight", "u"), op("gelu", {"g"}, "g", {{"approx", "tanh"}}),
                         op("mul", {"g", "u"}, "gu"), linear("gu", B + "ffn_down.weight", "f"),
                         rmsnorm("f", B + "post_ffw_norm.weight", eps, "f"), op("add", {"x", "f"}, "x")});
    head = tied_head(json::array({rmsnorm("x", "output_norm.weight", eps, "h")}));
    state = json::array({kv_state(H, KV, hd, ctx, windows)});
    summary = "heads " + std::to_string(H) + "/" + std::to_string(KV) + " x " + std::to_string(hd) +
              ", sandwich RMSNorms, GeGLU, window " + std::to_string(win) + " on 5 of every 6 layers" +
              (factor != 1.0 ? ", global RoPE scaled 1/" + std::to_string(factor) : "");
  } else if (arch == "deepseek2") {
    // DeepSeek-V2 and V3 in llama.cpp's MLA layout (IR_SPEC Appendix D.11; revision 1.6's leading dense blocks, V3's
    // router and YaRN): the query from attn_q (V2-Lite) or the low-rank attn_q_a / attn_q_b; attn_kv_a_mqa h split
    // into the latent c and the shared rotary key kr; mla_attention over a kv slot holding [c | kr]; the first
    // leading_dense_block_count applications a dense SwiGLU, the rest moe_ffn (times expert_weights_scale) plus
    // ungated shared experts, the two branches selected by apply. Files converted before llama.cpp's MLA keys,
    // other gating functions and other RoPE scaling build no graph.
    const int H = static_cast<int>(get("attention.head_count"));
    ctx = static_cast<int>(get("context_length"));
    const double eps = get("attention.layer_norm_rms_epsilon");
    if (!shape_has("blk.0.attn_k_b.weight") || !meta_number(g, p + "attention.key_length_mla"))
      throw std::runtime_error("IR builder: deepseek2 file without attn_k_b / attention.key_length_mla (converted "
                               "before llama.cpp's MLA layout) builds no graph");
    double gating = get("expert_gating_func", 1.0);
    if (gating == 0.0) gating = 1.0;
    if (gating != 1.0 && gating != 2.0)
      throw std::runtime_error("IR builder: deepseek2 with expert_gating_func " + std::to_string(gating) +
                               " (neither softmax 1 nor sigmoid 2) builds no graph");
    const auto stype = meta_string(g, p + "rope.scaling.type");
    if (shape_has("rope_freqs.weight") || (stype && *stype != "none" && *stype != "yarn"))
      throw std::runtime_error("IR builder: deepseek2 with RoPE scaling other than YaRN builds no graph");
    const bool yarn = stype && *stype == "yarn";
    const int k = static_cast<int>(get("leading_dense_block_count", 0.0));
    if (k < 0 || k > n_layers)
      throw std::runtime_error("IR builder: deepseek2 leading_dense_block_count " + std::to_string(k) + " exceeds block_count");
    const std::string M = "blk." + std::to_string(k) + ".";  // the first mixture application's tensors
    const int dr = static_cast<int>(get("rope.dimension_count"));
    const int dn = static_cast<int>(get("attention.key_length_mla")) - dr;
    const int dv = static_cast<int>(get("attention.value_length_mla"));
    const int r = static_cast<int>(get("attention.kv_lora_rank"));
    const double theta = get("rope.freq_base", 10000.0);
    const std::string gate_norm = truth("expert_weights_norm", false) ? "renorm" : "raw";
    double s = get("expert_weights_scale", 1.0);
    if (s == 0.0) s = 1.0;
    // YaRN (IR_SPEC 5.14, revision 1.6): pair j's divisor f_j = 1 / ((1 - e_j) / c + e_j), e_j = 1 - ramp(j), the
    // ramp linear from the correction dimension of beta_fast = 32 rotations over the original context (floor) to
    // that of beta_slow = 1 (ceiling), within 0 .. dr - 1; in binary64, each value rounded to f32. The softmax
    // temperature m = 1 + yarn_log_multiplier ln(c) enters the attention scale as m^2.
    json yarn_f = nullptr;
    double mscale = 1.0;
    if (yarn) {
      const double c = get("rope.scaling.factor"), o = get("rope.scaling.original_context_length");
      const double pi = 3.14159265358979323846;
      auto corr = [&](double beta) { return dr * std::log(o / (beta * 2.0 * pi)) / (2.0 * std::log(theta)); };
      double lo = std::max(std::floor(corr(32.0)), 0.0), hi = std::min(std::ceil(corr(1.0)), static_cast<double>(dr - 1));
      if (lo == hi) hi += 0.001;
      yarn_f = json::array();
      for (int j = 0; j < dr / 2; ++j) {
        const double ramp = std::min(std::max((j - lo) / (hi - lo), 0.0), 1.0);
        const double e = 1.0 - ramp;
        yarn_f.push_back(f32_value(1.0 / ((1.0 - e) / c + e)));
      }
      mscale = 1.0 + get("rope.scaling.yarn_log_multiplier", 0.0) * std::log(c);
    }
    nodes = json::array({rmsnorm("x", B + "attn_norm.weight", eps, "h")});
    const bool q_lora = shape_has("blk.0.attn_q_a.weight");
    if (q_lora) {
      for (json n : {linear("h", B + "attn_q_a.weight", "qa"), rmsnorm("qa", B + "attn_q_a_norm.weight", eps, "qa"),
                     linear("qa", B + "attn_q_b.weight", "q")})
        nodes.push_back(n);
    } else {
      nodes.push_back(linear("h", B + "attn_q.weight", "q"));
    }
    json rope_node = json{{"op", "rope"}, {"in", {"qr", "kr", "pos"}}, {"heads", H}, {"kv_heads", 1}, {"head_dim", dr},
                          {"n_rot", dr}, {"theta", f32_value(theta)}, {"rope_heads", "all"}, {"pairing", "interleaved"},
                          {"out", {"qr", "kr"}}};
    if (yarn) rope_node["freq_factors"] = yarn_f;
    for (json n : {op("split", {"q"}, {"qn", "qr"}, {{"sizes", {dn, dr}}, {"groups", H}}),
                   linear("h", B + "attn_kv_a_mqa.weight", "ckr"), op("split", {"ckr"}, {"c", "kr"}, {{"sizes", {r, dr}}}),
                   rmsnorm("c", B + "attn_kv_a_norm.weight", eps, "c"), rope_node,
                   json{{"op", "mla_attention"}, {"in", {"qn", "qr", "c", "kr"}}, {"state", "kv"}, {"heads", H},
                        {"kv_rank", r}, {"qk_nope", dn}, {"qk_rope", dr}, {"v_head_dim", dv},
                        {"k_b", B + "attn_k_b.weight"}, {"v_b", B + "attn_v_b.weight"},
                        {"scale", f32_value(mscale * mscale / std::sqrt(static_cast<double>(dn + dr)))},
                        {"causal", true}, {"window", 0}, {"sinks", 0}, {"out", "a"}},
                   linear("a", B + "attn_output.weight", "o"), op("add", {"x", "o"}, "x"),
                   rmsnorm("x", B + "ffn_norm.weight", eps, "h2")})
      nodes.push_back(n);
    // apply (IR_SPEC 3.8): with 0 < k < L the dense branch runs in applications i < k, the mixture in i >= k.
    const bool split_layers = k > 0 && k < n_layers;
    auto applied = [&](json n, bool dense) {
      if (split_layers) {
        json pl = json::array();
        for (int i = 0; i < n_layers; ++i) pl.push_back(dense ? i < k : i >= k);
        n["apply"] = {{"per_layer", pl}};
      }
      return n;
    };
    if (k > 0)
      for (json n : {linear("h2", B + "ffn_gate.weight", "g"), linear("h2", B + "ffn_up.weight", "u"),
                     op("silu", {"g"}, "g"), op("mul", {"g", "u"}, "gu"), linear("gu", B + "ffn_down.weight", "f")})
        nodes.push_back(applied(n, true));
    const bool shexp = k < n_layers && shape_has(M + "ffn_gate_shexp.weight");
    std::string router_note;
    if (k < n_layers) {
      json moe = moe_node(gate_norm);
      if (gating == 2.0) moe["scoring"] = "sigmoid";
      if (shape_has(M + "exp_probs_b.bias")) moe["selection_bias"] = B + "exp_probs_b.bias";
      const int gcount = static_cast<int>(get("expert_group_count", 1.0));
      if (gcount > 1) moe["groups"] = {{"n", gcount}, {"top", static_cast<int>(get("expert_group_used_count", 1.0))}};
      router_note = (moe.contains("scoring") ? ", sigmoid scores" : "") +
                    std::string(moe.contains("selection_bias") ? ", selection bias" : "") +
                    (moe.contains("groups") ? ", " + std::to_string(gcount) + " groups" : std::string());
      nodes.push_back(applied(moe, false));
      if (s != 1.0) nodes.push_back(applied(op("affine", {"f"}, "f", {{"scale", f32_value(s)}, {"shift", 0.0}}), false));
      if (shexp) {
        json sh = json::array();
        shared_expert(sh);
        for (const auto& n : sh) nodes.push_back(applied(n, false));
        nodes.push_back(applied(op("add", {"f", "sf"}, "f"), false));
      }
    }
    nodes.push_back(op("add", {"x", "f"}, "x"));
    head = tied_head(json::array({rmsnorm("x", "output_norm.weight", eps, "h")}));
    state = json::array({{{"name", "kv"}, {"kind", "kv"}, {"heads", H}, {"kv_heads", 1}, {"head_dim", r + dr},
                          {"max_len", ctx}, {"window", 0}, {"sinks", 0}, {"dtype", "f32"}}});
    summary = "MLA: heads " + std::to_string(H) + ", kv_rank " + std::to_string(r) + ", qk_nope " + std::to_string(dn) +
              ", qk_rope " + std::to_string(dr) + ", v_head_dim " + std::to_string(dv) +
              (q_lora ? ", low-rank query" : "") + (yarn ? ", YaRN" : "") +
              (k > 0 ? ", " + std::to_string(k) + " leading dense blocks" : std::string()) +
              (k < n_layers ? ", moe_ffn " + std::to_string(static_cast<int>(get("expert_count"))) + " experts (" +
                                  gate_norm + router_note + (s != 1.0 ? ", x " + json(f32_value(s)).dump() : std::string()) +
                                  ")" + (shexp ? " + shared experts" : "")
                            : std::string());
  } else if (arch == "phi3") {
    // IR_SPEC Appendix D.8 (revision 1.5): pre-RMSNorm, fused (or separate) q/k/v, NEOX RoPE over
    // rope.dimension_count dims with LongRoPE's factor sets when the file has them, HF's sliding window, SwiGLU
    // from the fused ffn_up (gate | up), output or the tied embedding.
    const int H = static_cast<int>(get("attention.head_count"));
    const int KV = static_cast<int>(get("attention.head_count_kv", H));
    const int hd = static_cast<int>(get("attention.key_length", d / H));
    const int n_rot = static_cast<int>(get("rope.dimension_count", hd));
    const int ff = static_cast<int>(get("feed_forward_length"));
    ctx = static_cast<int>(get("context_length"));
    const double eps = get("attention.layer_norm_rms_epsilon");
    const double theta = get("rope.freq_base", 10000.0);
    const int window = static_cast<int>(get("attention.sliding_window", 0.0));
    no_rope_scaling();
    if (shape_has("blk.0.ffn_gate_inp.weight"))
      throw std::runtime_error("IR builder: phi3 with a mixture of experts (Phi-3.5-MoE's layout) is not built");
    const bool fs = shape_has("rope_factors_short.weight"), fl = shape_has("rope_factors_long.weight");
    if (fs != fl)
      throw std::runtime_error("IR builder: phi3 with one LongRoPE factor set (rope_factors_short / _long) builds no graph");
    json r = rope(H, KV, hd, n_rot, f32_value(theta), "neox");
    double m = 1.0;
    if (fs) {
      r["freq_factors"] = {{"short", "rope_factors_short.weight"},
                           {"long", "rope_factors_long.weight"},
                           {"original_ctx", static_cast<int>(get("rope.scaling.original_context_length"))}};
      m = meta_number(g, p + "rope.scaling.attn_factor").value_or(1.0);
      if (m == 0.0) m = 1.0;
      if (m != 1.0) r["attn_factor"] = f32_value(m);
    }
    auto bias = [&](const std::string& name) { return shape_has("blk.0." + name + ".bias") ? B + name + ".bias" : ""; };
    nodes = json::array({rmsnorm("x", B + "attn_norm.weight", eps, "h")});
    if (shape_has("blk.0.attn_qkv.weight")) {
      nodes.push_back(linear("h", B + "attn_qkv.weight", "qkv", bias("attn_qkv")));
      nodes.push_back(op("split", {"qkv"}, {"q", "k", "v"}, {{"sizes", {H * hd, KV * hd, KV * hd}}}));
    } else {
      for (const char* t : {"q", "k", "v"})
        nodes.push_back(linear("h", B + "attn_" + t + ".weight", t, bias(std::string("attn_") + t)));
    }
    for (json n : {r, attention(H, KV, hd, 1.0 / std::sqrt(static_cast<double>(hd)), window),
                   linear("a", B + "attn_output.weight", "o", bias("attn_output")), op("add", {"x", "o"}, "x"),
                   rmsnorm("x", B + "ffn_norm.weight", eps, "h2"), linear("h2", B + "ffn_up.weight", "gu2"),
                   op("split", {"gu2"}, {"g", "u"}, {{"sizes", {ff, ff}}}), op("silu", {"g"}, "g"),
                   op("mul", {"g", "u"}, "gu"), linear("gu", B + "ffn_down.weight", "f"), op("add", {"x", "f"}, "x")})
      nodes.push_back(n);
    head = tied_head(json::array({rmsnorm("x", "output_norm.weight", eps, "h")}));
    state = json::array({kv_state(H, KV, hd, ctx, window)});
    summary = "heads " + std::to_string(H) + "/" + std::to_string(KV) + " x " + std::to_string(hd) + ", NEOX RoPE n_rot " +
              std::to_string(n_rot) + (fs ? ", LongRoPE (original_ctx " + r["freq_factors"]["original_ctx"].dump() + ")" : "") +
              (m != 1.0 ? ", attn_factor " + std::to_string(m) : "") + (window ? ", window " + std::to_string(window) : "");
  } else if (arch == "gemma2") {
    // IR_SPEC Appendix D.9 (revision 1.5): gemma3's block without q/k norms; one RoPE base; local applications
    // i % 2 == 0 windowed; the attention scores soft-capped at attn_logit_softcapping, the logits through
    // tanh(x / C) * C with C = final_logit_softcapping.
    const int H = static_cast<int>(get("attention.head_count"));
    const int KV = static_cast<int>(get("attention.head_count_kv", H));
    const int hd = static_cast<int>(get("attention.key_length", d / H));
    const int n_rot = static_cast<int>(get("rope.dimension_count", hd));
    ctx = static_cast<int>(get("context_length"));
    const double eps = get("attention.layer_norm_rms_epsilon");
    const double theta = get("rope.freq_base", 10000.0);
    no_rope_scaling();
    auto nonzero = [&](const std::string& key, double dflt) {
      const double v = get(key, dflt);
      return v == 0.0 ? dflt : v;
    };
    const int win = static_cast<int>(nonzero("attention.sliding_window", 4096.0));
    const double cap = nonzero("attn_logit_softcapping", 50.0), fcap = nonzero("final_logit_softcapping", 30.0);
    json windows = json::array();
    for (int i = 0; i < n_layers; ++i) windows.push_back(i % 2 == 0 ? win : 0);
    windows = {{"per_layer", windows}};
    const double scale = n_layers == 46 ? 1.0 / std::sqrt(static_cast<double>(d) / H)
                                        : 1.0 / std::sqrt(static_cast<double>(hd));
    json att = attention(H, KV, hd, scale, windows);
    att["softcap"] = f32_value(cap);
    // C: the shortest decimal that reads back as the f32 nearest fcap, written as Python's repr of that float
    // (an integral value with ".0").
    std::string C;
    {
      char buf[64];
      const auto res = std::to_chars(buf, buf + sizeof(buf), f32_value(fcap));
      C.assign(buf, res.ptr);
      if (C.find_first_of(".eEn") == std::string::npos) C += ".0";
    }
    embed.push_back(op("affine", {"x"}, "x", {{"scale", f32_value(std::sqrt(static_cast<double>(d)))}, {"shift", 0.0}}));
    nodes = json::array({rmsnorm("x", B + "attn_norm.weight", eps, "h"), linear("h", B + "attn_q.weight", "q"),
                         linear("h", B + "attn_k.weight", "k"), linear("h", B + "attn_v.weight", "v"),
                         rope(H, KV, hd, n_rot, f32_value(theta), "neox"), att,
                         linear("a", B + "attn_output.weight", "o"),
                         rmsnorm("o", B + "post_attention_norm.weight", eps, "o"), op("add", {"x", "o"}, "x"),
                         rmsnorm("x", B + "ffn_norm.weight", eps, "h2"), linear("h2", B + "ffn_gate.weight", "g"),
                         linear("h2", B + "ffn_up.weight", "u"), op("gelu", {"g"}, "g", {{"approx", "tanh"}}),
                         op("mul", {"g", "u"}, "gu"), linear("gu", B + "ffn_down.weight", "f"),
                         rmsnorm("f", B + "post_ffw_norm.weight", eps, "f"), op("add", {"x", "f"}, "x")});
    head = json::array({rmsnorm("x", "output_norm.weight", eps, "h"),
                        linear("h", "output.weight", "z", "", tied ? "token_embd.weight" : ""),
                        op("expr", {"z"}, "logits", {{"expr", "tanh(x / " + C + ") * " + C}})});
    state = json::array({kv_state(H, KV, hd, ctx, windows)});
    summary = "heads " + std::to_string(H) + "/" + std::to_string(KV) + " x " + std::to_string(hd) +
              ", sandwich RMSNorms, GeGLU, window " + std::to_string(win) + " on even layers, attention softcap " +
              std::to_string(cap) + ", final softcap " + C;
  } else {  // mamba
    const int K = static_cast<int>(get("ssm.conv_kernel")), Di = static_cast<int>(get("ssm.inner_size"));
    const int N = static_cast<int>(get("ssm.state_size")), R = static_cast<int>(get("ssm.time_step_rank"));
    const double eps = get("attention.layer_norm_rms_epsilon");
    ctx = static_cast<int>(get("context_length", 0.0));
    if (ctx == 0) ctx = 1048576;  // llama.cpp writes 1M: a recurrence has no length limit
    if (!shape_has("blk.0.ssm_dt.weight") || !shape_has("blk.0.ssm_d"))
      throw std::runtime_error("IR builder: mamba file lacks ssm_dt.weight or ssm_d");
    if (meta_bool(g, p + "ssm.dt_b_c_rms").value_or(false) || shape_has("blk.0.ssm_dt_norm.weight"))
      throw std::runtime_error("IR builder: mamba with dt/B/C RMSNorms (Falcon-Mamba) is not built yet");
    nodes = json::array(
        {rmsnorm("x", B + "attn_norm.weight", eps, "h"), linear("h", B + "ssm_in.weight", "xz"),
         op("split", {"xz"}, {"xs", "z"}, {{"sizes", {Di, Di}}}),
         op("conv1d_causal", {"xs"}, "xs",
            {{"state", "conv"}, {"weight", B + "ssm_conv1d.weight"}, {"bias", B + "ssm_conv1d.bias"},
             {"channels", Di}, {"kernel", K}}),
         op("silu", {"xs"}, "xs"), linear("xs", B + "ssm_x.weight", "xdb"),
         op("split", {"xdb"}, {"dt", "bm", "cm"}, {{"sizes", {R, N, N}}}),
         linear("dt", B + "ssm_dt.weight", "dt", B + "ssm_dt.bias"),
         op("selective_scan", {"xs", "dt", "bm", "cm"}, "y",
            {{"state", "ssm"}, {"a", B + "ssm_a"}, {"d", B + "ssm_d"}, {"channels", Di}, {"d_state", N},
             {"dt_softplus", true}}),
         op("silu", {"z"}, "z"), op("mul", {"y", "z"}, "y"), linear("y", B + "ssm_out.weight", "o"),
         op("add", {"x", "o"}, "x")});
    state = json::array({{{"name", "conv"}, {"kind", "recurrent"}, {"op", "conv1d_causal"}, {"dtype", "f32"},
                          {"tensors", json::array({{{"name", "x"}, {"shape", {K - 1, Di}}, {"init", 0}}})}},
                         {{"name", "ssm"}, {"kind", "recurrent"}, {"op", "selective_scan"}, {"dtype", "f32"},
                          {"tensors", json::array({{{"name", "s"}, {"shape", {Di, N}}, {"init", 0}}})}}});
    head = tied_head(json::array({rmsnorm("x", "output_norm.weight", eps, "h")}));
    summary = "selective scan d_inner " + std::to_string(Di) + ", d_state " + std::to_string(N) + ", dt_rank " +
              std::to_string(R) + ", conv kernel " + std::to_string(K);
  }
 
  const auto eos = meta_number(g, "tokenizer.ggml.eos_token_id");
  const auto bos = meta_number(g, "tokenizer.ggml.bos_token_id");
  json doc;
  doc["_comment"] = "built by ember's build_ir_from_gguf (graph_ir.cpp; paperkiln tools/gguf_ir.py's rules) from "
                    "general.architecture = " + arch + " and its metadata";
  doc["ir_version"] = 1;  // IR_SPEC 2.3: as tools/gguf_ir.py's IR_VERSION from paperkiln 5349ee19 on
  doc["family_hint"] = arch;
  doc["model"] = {{"vocab", vocab}, {"d", d}, {"n_layers", n_layers}, {"ctx", ctx}, {"dtype", "f32"}};
  doc["tokenizer"] = {{"type", "gguf"},
                      {"source", "gguf"},
                      {"specials", eos ? json{{"eot", static_cast<int64_t>(*eos)}} : json::object()},
                      {"bos", bos ? json(static_cast<int64_t>(*bos)) : json(nullptr)}};
  doc["inputs"] = json::array({{{"name", "tokens"}, {"kind", "token_ids"}}, {{"name", "pos"}, {"kind", "positions"}}});
  if (arch == "mamba") doc["inputs"] = json::array({{{"name", "tokens"}, {"kind", "token_ids"}}});
  doc["embed"] = {{"nodes", embed}};
  doc["block"] = {{"repeat", n_layers}, {"state", state}, {"nodes", nodes}};
  doc["head"] = {{"nodes", head}};
  doc["outputs"] = json::array({{{"name", "logits"}, {"kind", "logits"}}});
  doc["modes"] = {{"prefill", "all nodes over T tokens"},
                  {"decode", "the same graph with T = 1, attention reads and appends to the kv state, a recurrent op "
                             "reads and replaces its recurrent state"}};
  if (log)
    Logger::info("[IR] built the IR for general.architecture = " + arch + " from its metadata: d " + std::to_string(d) +
                 ", " + std::to_string(n_layers) + " layers, " + summary + (tied ? ", tied head" : ""));
  return doc.dump();
}

References c, has_ir_builder(), in, Logger::info(), j, k, layernorm, linear, n, name, out, p, rmsnorm, scale, state, GGUFData::tensor_infos_map, GGUFData::tokenizer_tokens, and v.

Referenced by find_ir().

cpu_only()

static void ember_ir::cpu_only ( bool  gpu, const char *  what  )

static

Definition at line 3053 of file graph_ir.cpp.

                                                 {
  if (gpu) throw std::runtime_error(std::string("IR ") + what + " is implemented for the CPU path only");
}

References cpu_only(), and j.

Referenced by cpu_only(), ember_ir::Interpreter::kv_positions(), ember_ir::Interpreter::load_state(), ember_ir::Interpreter::recurrent_values(), ember_ir::Interpreter::save_state(), and ember_ir::Interpreter::truncate().

device_supported()

bool ember_ir::device_supported

(

const std::string &

ir_json

)

Definition at line 463 of file graph_ir.cpp.

                                                {
  // Per-turn state tokens (state_weight, mask turn_state, evict turn_state) run on either path.
  const json j = json::parse(ir_json);
  // depth_attend reads at most every tensor ever put in a list (list inputs
  // and appends, the block's once per repeat) plus its own extra inputs.
  // Per-layer attributes (the ones the interpreter accepts: attention and
  // kv-state window and sinks, RoPE theta and freq_scale) run on either path.
  // Mixture-of-Recursions (depth_route) runs a step on its active rows at their own positions: rope and
  // attention take them (mor_step_device); a pos_embedding in its block stays on the CPU.
  size_t list_tensors = 0, extra = 0;
  bool mor = false, block_pos_embedding = false, softcap = false, turn_state = false;
  for (const char* sub : {"embed", "block", "head"}) {
    if (!j.contains(sub)) continue;
    const size_t times = std::string(sub) == "block" ? j.at(sub).value("repeat", 1) : 1;
    for (const auto& n : j.at(sub).value("nodes", json::array())) {
      const std::string op = n.value("op", "");
      if (!kCudaOps.count(op)) return false;
      if (std::string(sub) == "block") {
        mor = mor || op == "depth_route";
        block_pos_embedding = block_pos_embedding || op == "pos_embedding";
      }
      if (op == "attention" && n.contains("position_bias")) return false;  // CPU only (IR_SPEC 5.15, revision 1.3)
      if (op == "attention" && n.contains("relative_kv")) return false;    // CPU only (IR_SPEC 5.15, revision 1.4)
      // revision 1.6's apply and DeepSeek-V3's router have no device path: a CUDA build serves the document on the CPU
      if (n.contains("apply")) return false;
      if (op == "moe_ffn" && (n.contains("scoring") || n.contains("selection_bias") || n.contains("groups"))) return false;
      // softcap (revision 1.5) runs in the window kernel; with turn_state or depth_route it stays on the CPU
      if (op == "attention" && n.contains("softcap")) softcap = true;
      if (op == "attention" && n.contains("mask") && n.at("mask") != "causal") turn_state = true;
      if (op == "expr") {  // the CUDA form's limits (graph_ir_expr.h lower)
        try {
          expr::lower(expr::cached(n.at("expr").get<std::string>(), n.at("in").size()));
        } catch (const std::exception&) {
          return false;
        }
      }
      const size_t ins = n.contains("in") && n.at("in").is_array() ? n.at("in").size() : 1;
      if (op == "list") list_tensors += ins * times;
      if (op == "append") list_tensors += times;
      if (op == "depth_attend") extra = std::max(extra, ins - 1);
#ifdef HAS_CUDA
      if (op == "ssm_scan" && n.value("n", 0) > kIrSsmMax) return false;
#endif
    }
  }
  if (mor && block_pos_embedding) return false;
  // A window under depth_route (IR_SPEC 5.15, v1.4.2) counts positions over an application's compact cache: the
  // CPU interpreter's (attention under attn_row_); a CUDA build serves such a document on the CPU.
  if (mor && has_windowed_attention(ir_json)) return false;
  if (softcap && (mor || turn_state)) return false;
#ifdef HAS_CUDA
  if (list_tensors + extra > static_cast<size_t>(kIrDepthSourcesMax)) return false;
#endif
  return true;
}

References ember_ir::expr::cached(), has_windowed_attention(), j, ember_ir::expr::lower(), n, and sub.

Referenced by ember_ir::Interpreter::init(), and tinyllama::TinyLlamaSession::TinyLlamaSession().

es_number_string()

std::string ember_ir::es_number_string

(

double

v

)

Definition at line 472 of file ir_digest.cpp.

                                     {
  if (v == 0.0) return "0";
  if (v < 0) return "-" + es_number_string(-v);
  char buf[64];
  const auto r = std::to_chars(buf, buf + sizeof(buf), v, std::chars_format::scientific);
  if (r.ec != std::errc()) throw std::runtime_error("es_number_string: to_chars failed");
  const std::string sci(buf, r.ptr);
  const size_t e = sci.find('e');
  std::string digits;
  for (size_t i = 0; i < e; ++i)
    if (sci[i] != '.') digits += sci[i];
  const int exp10 = std::atoi(sci.c_str() + e + 1);
  const int k = static_cast<int>(digits.size()), n = exp10 + 1;
  if (k <= n && n <= 21) return digits + std::string(static_cast<size_t>(n - k), '0');
  if (0 < n && n <= 21) return digits.substr(0, n) + "." + digits.substr(n);
  if (-6 < n && n <= 0) return "0." + std::string(static_cast<size_t>(-n), '0') + digits;
  const int ex = n - 1;
  std::string out = digits.substr(0, 1);
  if (k > 1) out += "." + digits.substr(1);
  out += ex >= 0 ? "e+" : "e-";
  out += std::to_string(ex >= 0 ? ex : -ex);
  return out;
}

References es_number_string(), j, k, n, out, and v.

Referenced by es_number_string().

find_ir()

std::string ember_ir::find_ir

(

const GGUFData &

gguf

)

Definition at line 525 of file graph_ir.cpp.

                                        {
  if (const char* path = std::getenv("EMBER_IR")) {
    std::ifstream f(path);
    if (!f) throw std::runtime_error(std::string("EMBER_IR: cannot read ") + path);
    std::stringstream ss;
    ss << f.rdbuf();
    return ss.str();
  }
  auto it = gguf.metadata.find("paperkiln.ir");
  if (it != gguf.metadata.end()) {
    if (const auto* s = std::get_if<std::string>(&it->second)) return *s;
  }
  // No IR in the file: every built architecture gets one from its metadata,
  // except that llama and mistral without a sliding window keep the llama path
  // unless EMBER_IR_BUILD=1 (a window needs the IR: the llama path has none).
  const std::string arch = meta_string(gguf, "general.architecture").value_or("llama");
  if (arch == "llama" || arch == "mistral") {
    const char* b = std::getenv("EMBER_IR_BUILD");
    const bool windowed = meta_number(gguf, arch + ".attention.sliding_window").value_or(0) > 0;
    return windowed || (b && std::string(b) == "1") ? build_ir_from_gguf(gguf) : std::string();
  }
  if (has_ir_builder(arch)) return build_ir_from_gguf(gguf);
  throw std::runtime_error("ember cannot run general.architecture = '" + arch +
                           "': the GGUF has no paperkiln.ir and ember has no IR builder for that architecture "
                           "(built: llama, mistral, qwen2, qwen3, gpt2, gptneox, gemma3, mamba, qwen2moe, qwen3moe, "
                           "deepseek2, phi3, gemma2)");
}

References build_ir_from_gguf(), has_ir_builder(), and j.

Referenced by TinyLlamaModel::TinyLlamaModel(), and tinyllama::TinyLlamaSession::TinyLlamaSession().

has_ir_builder()

bool ember_ir::has_ir_builder

(

const std::string &

arch

)

Definition at line 519 of file graph_ir.cpp.

                                           {
  static const std::set<std::string> kBuilt = {"llama", "mistral", "qwen2", "qwen3", "gpt2", "gptneox", "gemma3",
                                               "mamba", "qwen2moe", "qwen3moe", "deepseek2", "phi3", "gemma2"};
  return kBuilt.count(arch) > 0;
}

References j.

Referenced by build_ir_from_gguf(), and find_ir().

has_recurrent_state()

bool ember_ir::has_recurrent_state

(

const std::string &

ir_json

)

Definition at line 237 of file graph_ir.cpp.

                                                   {
  const json j = json::parse(ir_json);
  if (!j.contains("block")) return false;
  for (const auto& s : j.at("block").value("state", json::array()))
    if (s.value("kind", "") == "recurrent") return true;
  return false;
}

References j.

has_state_tokens()

bool ember_ir::has_state_tokens

(

const std::string &

ir_json

)

Definition at line 257 of file graph_ir.cpp.

                                                {
  const json j = json::parse(ir_json);
  for (const char* sub : {"embed", "block", "head"}) {
    if (!j.contains(sub)) continue;
    for (const auto& n : j.at(sub).value("nodes", json::array())) {
      const std::string op = n.value("op", "");
      if (op == "embedding" && has(n, "state_weight")) return true;
      if (op == "attention" && n.contains("mask") && n.at("mask") != "causal") return true;
    }
  }
  if (j.contains("block"))
    for (const auto& s : j.at("block").value("state", json::array()))
      if (has(s, "evict")) return true;
  return false;
}

References j, n, and sub.

has_windowed_attention()

bool ember_ir::has_windowed_attention

(

const std::string &

ir_json

)

Definition at line 245 of file graph_ir.cpp.

                                                      {
  const json j = json::parse(ir_json);
  if (!j.contains("block")) return false;
  const json& blk = j.at("block");
  for (const auto& s : blk.value("state", json::array()))
    if (any_nonzero(s, "window")) return true;
  for (const auto& n : blk.value("nodes", json::array()))
    if ((n.value("op", "") == "attention" || n.value("op", "") == "mla_attention") && any_nonzero(n, "window"))
      return true;
  return false;
}

References j, and n.

Referenced by device_supported().

ir_affine_cuda()

void ember_ir::ir_affine_cuda	(	const float *	x,
		float *	out,
		size_t	n,
		float	scale,
		float	shift,
		cudaStream_t	stream
	)

Definition at line 737 of file graph_ir_cuda.cu.

                                                                                                         {
  unary_kernel<<<blocks_for(n), kThreads, 0, stream>>>(x, out, n, 2, scale, shift);
  check_launch("affine");
}

References j, n, out, and scale.

ir_bias_add_cuda()

void ember_ir::ir_bias_add_cuda	(	float *	y,
		const float *	b,
		int	T,
		int	O,
		cudaStream_t	stream
	)

Definition at line 803 of file graph_ir_cuda.cu.

                                                                                   {
  const size_t n = (size_t)T * O;
  bias_add_kernel<<<blocks_for(n), kThreads, 0, stream>>>(y, b, n, O);
  check_launch("bias");
}

References j, and n.

ir_binary_cuda()

void ember_ir::ir_binary_cuda	(	const float *	a,
		const float *	b,
		float *	out,
		size_t	n,
		int	op,
		cudaStream_t	stream
	)

Definition at line 809 of file graph_ir_cuda.cu.

                                                                                                       {
  binary_kernel<<<blocks_for(n), kThreads, 0, stream>>>(a, b, out, n, op);
  check_launch(op == 0 ? "add" : op == 1 ? "mul" : "sub");
}

References a, j, n, and out.

ir_container_errors()

std::vector< std::string > ember_ir::ir_container_errors	(	const json &	doc,
		const GGUFData &	gguf
	)

Definition at line 1609 of file graph_ir_validate.cpp.

                                                                                {
  Errs errs;
  container_errors(doc, gguf, errs);
  mla_container_errors(doc, gguf, errs);
  architecture_errors(gguf, errs);
  std::vector<std::string> out;
  std::set<std::string> seen;
  for (const auto& m : errs) add_unique(out, seen, m);
  return out;
}

References ir_container_errors(), j, and out.

Referenced by ir_container_errors(), and validate_ir().

ir_conv1d_causal_cuda()

void ember_ir::ir_conv1d_causal_cuda	(	const float *	x,
		const float *	w,
		const float *	bias,
		float *	state,
		float *	out,
		int	T,
		int	C,
		int	K,
		cudaStream_t	stream
	)

Definition at line 838 of file graph_ir_cuda.cu.

                                                       {
  if (K < 1 || C <= 0) throw std::runtime_error("IR CUDA conv1d_causal: kernel " + std::to_string(K));
  conv1d_causal_kernel<<<(C + 127) / 128, 128, 0, stream>>>(x, w, bias, state, out, T, C, K);
  check_launch("conv1d_causal");
}

References j, out, and state.

ir_depth_attend_cuda()

void ember_ir::ir_depth_attend_cuda	(	const float *const *	sources,
		int	S,
		const float *	w,
		float *	out,
		int	T,
		int	d,
		float	eps,
		cudaStream_t	stream
	)

Definition at line 759 of file graph_ir_cuda.cu.

                                               {
  if (S < 2 || S > kIrDepthSourcesMax)
    throw std::runtime_error("IR CUDA depth_attend: " + std::to_string(S) + " sources (2.." +
                             std::to_string(kIrDepthSourcesMax) + " supported)");
  DepthSources src{};
  for (int i = 0; i < S; ++i) src.p[i] = sources[i];
  depth_attend_kernel<<<T, 128, 0, stream>>>(src, S, w, out, d, eps);
  check_launch("depth_attend");
}

References j, and out.

ir_digest()

DigestResult ember_ir::ir_digest	(	const std::string &	text,
		bool	prose
	)

Definition at line 496 of file ir_digest.cpp.

                                                          {
  DigestResult r;
  try {
    const Value v = Parser(text).document();
    if (holds_negative_zero(v)) throw NoDigest{"a negative zero (RFC 8785 writes it as 0, so it would share +0's digest)"};
    write(r.canonical, v, prose);
  } catch (const NoDigest& e) {
    r.refused = e.why;
    r.canonical.clear();
    return r;
  }
  r.sha256 = sha256_hex_bytes(r.canonical);
  r.ok = true;
  return r;
}

References j, sha256_hex_bytes(), text, and v.

ir_document_errors()

std::vector< std::string > ember_ir::ir_document_errors

(

const json &

doc

)

Definition at line 1592 of file graph_ir_validate.cpp.

                                                         {
  if (!doc.is_object()) return {E("10.1 rule 1", "the document is not a JSON object")};
  Errs errs = ir_schema_errors(doc);
  const json& v = get(doc, "ir_version");
  if (is_int(v) && v.get<int64_t>() != 0 && v.get<int64_t>() != 1)
    errs.push_back(E("2.3", "ir_version " + v.dump() + " is not supported (ember reads 0, 1)"));
  node_errors(doc, errs);
  state_errors(doc, errs);
  depth_route_errors(doc, errs);
  state_token_doc_errors(doc, errs);
  read_errors(doc, errs);
  std::vector<std::string> out;
  std::set<std::string> seen;
  for (const auto& m : errs) add_unique(out, seen, m);
  return out;
}

References ir_document_errors(), ir_schema_errors(), j, out, and v.

Referenced by ir_document_errors(), and validate_ir().

ir_embedding_cuda()

void ember_ir::ir_embedding_cuda	(	const float *	table,
		const int *	tokens,
		float *	out,
		int	T,
		int	d,
		cudaStream_t	stream
	)

Definition at line 786 of file graph_ir_cuda.cu.

                                                                                                             {
  embedding_kernel<<<T, 128, 0, stream>>>(table, tokens, out, d);
  check_launch("embedding");
}

References j, and out.

ir_embedding_state_cuda()

void ember_ir::ir_embedding_state_cuda	(	const float *	table,
		const float *	state,
		int	vocab,
		const int *	tokens,
		float *	out,
		int	T,
		int	d,
		cudaStream_t	stream
	)

Definition at line 702 of file graph_ir_cuda.cu.

                                                         {
  embedding_state_kernel<<<T, 128, 0, stream>>>(table, state, vocab, tokens, out, d);
  check_launch("embedding (state rows)");
}

References j, out, and state.

ir_expr_cuda()

void ember_ir::ir_expr_cuda	(	const std::vector< expr::Stage > &	stages,
		const float *const *	ins,
		int	n_ins,
		float *	scratch,
		float *	out,
		int	T,
		int	w,
		cudaStream_t	stream
	)

Definition at line 857 of file graph_ir_cuda.cu.

                                                                 {
  if (n_ins <= 0 || n_ins > (int)expr::Expr::kMaxInputs || stages.empty() || stages.size() - 1 > expr::kMaxColumns)
    throw std::runtime_error("IR CUDA expr: " + std::to_string(n_ins) + " inputs, " + std::to_string(stages.size()) +
                             " stages");
  ExprPtrs p{};
  for (int k = 0; k < n_ins; ++k) p.in[k] = ins[k];
  float* full = scratch;  // a reduction's argument, [T, w], reused stage by stage
  for (size_t k = 0; k + 1 < stages.size(); ++k) p.col[k] = scratch + (size_t)T * w + k * (size_t)T;
  const size_t n = (size_t)T * w;
  for (size_t k = 0; k < stages.size(); ++k) {
    const expr::Stage& s = stages[k];
    if (s.code.empty() || s.code.size() > expr::kMaxCode) throw std::runtime_error("IR CUDA expr: a stage's code");
    ExprCode code{};
    code.n = (int)s.code.size();
    for (size_t i = 0; i < s.code.size(); ++i) code.code[i] = s.code[i];
    const bool last = k + 1 == stages.size();
    expr_kernel<<<blocks_for(n), kThreads, 0, stream>>>(code, p, last ? out : full, n, w);
    check_launch("expr");
    if (last) break;
    row_reduce_kernel<<<(T + 63) / 64, 64, 0, stream>>>(full, const_cast<float*>(p.col[k]), T, w, s.reduce == 2);
    check_launch("expr row reduction");
  }
}

References code, j, k, ember_ir::expr::kMaxCode, ember_ir::expr::kMaxColumns, ember_ir::expr::Expr::kMaxInputs, n, out, and p.

ir_expr_scratch_floats()

size_t ember_ir::ir_expr_scratch_floats	(	const std::vector< expr::Stage > &	stages,
		int	T,
		int	w
	)

Definition at line 852 of file graph_ir_cuda.cu.

                                                                                {
  if (stages.size() <= 1) return 0;
  return (size_t)T * w + (stages.size() - 1) * (size_t)T;
}

References j.

ir_gelu_cuda()

void ember_ir::ir_gelu_cuda	(	const float *	x,
		float *	out,
		size_t	n,
		bool	tanh_form,
		cudaStream_t	stream
	)

Definition at line 819 of file graph_ir_cuda.cu.

                                                                                             {
  gelu_kernel<<<blocks_for(n), kThreads, 0, stream>>>(x, out, n, tanh_form ? 1 : 0);
  check_launch("gelu");
}

References j, n, and out.

ir_kv_compact_cuda()

void ember_ir::ir_kv_compact_cuda	(	float *	K,
		float *	V,
		const int *	idx,
		int	w0,
		int	n,
		int	row,
		cudaStream_t	stream
	)

Definition at line 696 of file graph_ir_cuda.cu.

                                                                                                         {
  if (w0 >= n) return;
  kv_compact_kernel<<<1, kThreads, 0, stream>>>(K, V, idx, w0, n, row);
  check_launch("kv compact");
}

References j, and n.

Referenced by ember_ir::Interpreter::ts_evict_below().

ir_kv_write_cuda()

void ember_ir::ir_kv_write_cuda	(	float *	K,
		float *	V,
		const float *	k,
		const float *	v,
		int	T,
		int	row,
		int	start_pos,
		cudaStream_t	stream,
		const int *	dpos
	)

Definition at line 720 of file graph_ir_cuda.cu.

                                                            {
  const size_t n = (size_t)T * row;
  kv_write_kernel<<<blocks_for(n), kThreads, 0, stream>>>(K, V, k, v, n, row, start_pos, dpos);
  check_launch("kv write");
}

References j, k, n, and v.

ir_layernorm_cuda()

void ember_ir::ir_layernorm_cuda	(	const float *	x,
		const float *	w,
		const float *	b,
		float *	out,
		int	T,
		int	d,
		float	eps,
		cudaStream_t	stream
	)

Definition at line 797 of file graph_ir_cuda.cu.

                                            {
  layernorm_kernel<<<T, kThreads, 0, stream>>>(x, w, b, out, d, eps);
  check_launch("layernorm");
}

References j, and out.

ir_linear_attention_cuda()

void ember_ir::ir_linear_attention_cuda	(	const float *	q,
		const float *	k,
		const float *	v,
		float *	num,
		float *	den,
		float *	out,
		int	T,
		int	width,
		float	den_floor,
		cudaStream_t	stream
	)

Definition at line 770 of file graph_ir_cuda.cu.

                                                                                      {
  linear_attention_kernel<<<(width + 127) / 128, 128, 0, stream>>>(q, k, v, num, den, out, T, width, den_floor);
  check_launch("linear_attention");
}

References j, k, out, and v.

ir_mla_attention_cuda()

void ember_ir::ir_mla_attention_cuda	(	const float *	qn,
		const float *	qr,
		const float *	KV,
		const float *	Kb,
		const float *	Vb,
		float *	out,
		int	T,
		int	start_pos,
		int	H,
		int	r,
		int	dn,
		int	dr,
		int	dv,
		float	scale,
		int	window,
		int	sinks,
		cudaStream_t	stream,
		const int *	dpos
	)

Definition at line 995 of file graph_ir_cuda.cu.

                                                                                        {
  const size_t smem = (static_cast<size_t>(r) * 2 + dr + kMlaChunk) * sizeof(float);
  if (smem > 48 * 1024)
    throw std::runtime_error("IR mla_attention on CUDA: kv_rank " + std::to_string(r) + " needs " + std::to_string(smem) +
                             " bytes of shared memory (at most 48 KiB)");
  mla_attention_kernel<<<static_cast<unsigned>(T) * H, 128, smem, stream>>>(qn, qr, KV, Kb, Vb, out, H, r, dn, dr, dv,
                                                                           scale, window, sinks, start_pos, dpos);
  check_launch("mla attention");
}

References j, out, and scale.

ir_mla_kv_write_cuda()

void ember_ir::ir_mla_kv_write_cuda	(	float *	KV,
		const float *	c,
		const float *	kr,
		int	T,
		int	r,
		int	dr,
		int	start_pos,
		cudaStream_t	stream,
		const int *	dpos
	)

Definition at line 988 of file graph_ir_cuda.cu.

                                                                {
  const size_t n = (size_t)T * (r + dr);
  mla_kv_write_kernel<<<blocks_for(n), kThreads, 0, stream>>>(KV, c, kr, T, r, dr, start_pos, dpos);
  check_launch("mla kv write");
}

References c, j, and n.

ir_mor_gate_cuda()

void ember_ir::ir_mor_gate_cuda	(	const float *	x_in,
		float *	x,
		const float *	g,
		int	n,
		int	w,
		cudaStream_t	stream
	)

Definition at line 665 of file graph_ir_cuda.cu.

                                                                                                      {
  if (n <= 0) return;
  mor_gate_kernel<<<n, 128, 0, stream>>>(x_in, x, g, w);
  check_launch("mor gate");
}

References j, and n.

ir_mul_col_cuda()

void ember_ir::ir_mul_col_cuda	(	const float *	x,
		const float *	c,
		float *	out,
		int	T,
		int	w,
		cudaStream_t	stream
	)

Definition at line 742 of file graph_ir_cuda.cu.

                                                                                                    {
  const size_t n = (size_t)T * w;
  mul_col_kernel<<<blocks_for(n), kThreads, 0, stream>>>(x, c, out, n, w);
  check_launch("mul_col");
}

References c, j, n, and out.

ir_pos_embedding_cuda()

void ember_ir::ir_pos_embedding_cuda	(	const float *	x,
		const float *	P,
		float *	out,
		int	T,
		int	d,
		int	start_pos,
		cudaStream_t	stream,
		const int *	dpos
	)

Definition at line 791 of file graph_ir_cuda.cu.

                                                                 {
  pos_embedding_kernel<<<T, 128, 0, stream>>>(x, P, out, d, start_pos, dpos);
  check_launch("pos_embedding");
}

References j, and out.

ir_relu_cuda()

void ember_ir::ir_relu_cuda	(	const float *	x,
		float *	out,
		size_t	n,
		cudaStream_t	stream
	)

Definition at line 727 of file graph_ir_cuda.cu.

                                                                             {
  unary_kernel<<<blocks_for(n), kThreads, 0, stream>>>(x, out, n, 0, 0.0f, 0.0f);
  check_launch("relu");
}

References j, n, and out.

ir_rms_row_cuda()

void ember_ir::ir_rms_row_cuda	(	const float *	x,
		float *	out,
		int	T,
		int	w,
		float	eps,
		cudaStream_t	stream
	)

Definition at line 754 of file graph_ir_cuda.cu.

                                                                                               {
  rms_row_kernel<<<(T + 63) / 64, 64, 0, stream>>>(x, out, T, w, eps);
  check_launch("rms_row");
}

References j, and out.

ir_rope_at_cuda()

void ember_ir::ir_rope_at_cuda	(	float *	q,
		float *	k,
		int	T,
		int	hq,
		int	hk,
		int	hd,
		const int *	pos,
		const float *	freqs,
		cudaStream_t	stream
	)

Definition at line 671 of file graph_ir_cuda.cu.

                                          {
  rope_at_kernel<<<T, 128, 0, stream>>>(q, k, hq, hk, hd, pos, freqs);
  check_launch("rope (positions)");
}

References j, k, and pos.

ir_rope_cuda()

void ember_ir::ir_rope_cuda	(	float *	q,
		float *	k,
		int	T,
		int	hq,
		int	hk,
		int	hd,
		int	start_pos,
		const float *	freqs,
		cudaStream_t	stream,
		const int *	dpos
	)

Definition at line 824 of file graph_ir_cuda.cu.

                                                        {
  rope_kernel<<<T, 128, 0, stream>>>(q, k, hq, hk, hd, start_pos, freqs, dpos);
  check_launch("rope");
}

References j, and k.

ir_rope_rows_at_cuda()

void ember_ir::ir_rope_rows_at_cuda	(	float *	x,
		int	T,
		int	heads,
		int	hd,
		int	n_rot,
		bool	neox,
		const int *	pos,
		const float *	freqs,
		cudaStream_t	stream
	)

Definition at line 677 of file graph_ir_cuda.cu.

                                               {
  if (n_rot <= 0 || n_rot % 2 != 0 || n_rot > hd)
    throw std::runtime_error("IR CUDA rope: n_rot " + std::to_string(n_rot) + " / head_dim " + std::to_string(hd));
  rope_rows_at_kernel<<<T, 128, 0, stream>>>(x, heads, hd, n_rot, neox ? 1 : 0, pos, freqs);
  check_launch("rope (rows, positions)");
}

References j, and pos.

ir_rope_rows_cuda()

void ember_ir::ir_rope_rows_cuda	(	float *	x,
		int	T,
		int	heads,
		int	hd,
		int	n_rot,
		bool	neox,
		int	start_pos,
		const float *	freqs,
		cudaStream_t	stream,
		const int *	dpos
	)

Definition at line 830 of file graph_ir_cuda.cu.

                                                             {
  if (n_rot <= 0 || n_rot % 2 != 0 || n_rot > hd)
    throw std::runtime_error("IR CUDA rope: n_rot " + std::to_string(n_rot) + " / head_dim " + std::to_string(hd));
  rope_rows_kernel<<<T, 128, 0, stream>>>(x, heads, hd, n_rot, neox ? 1 : 0, start_pos, freqs, dpos);
  check_launch("rope (rows)");
}

References j.

ir_rows_gather_cuda()

void ember_ir::ir_rows_gather_cuda	(	const float *	src,
		const int *	idx,
		float *	dst,
		int	n,
		int	w,
		cudaStream_t	stream
	)

Definition at line 653 of file graph_ir_cuda.cu.

                                                                                                          {
  if (n <= 0) return;
  rows_gather_kernel<<<n, 128, 0, stream>>>(src, idx, dst, w);
  check_launch("rows gather");
}

References j, and n.

ir_rows_scatter_cuda()

void ember_ir::ir_rows_scatter_cuda	(	const float *	src,
		const int *	idx,
		float *	dst,
		int	n,
		int	w,
		cudaStream_t	stream
	)

Definition at line 659 of file graph_ir_cuda.cu.

                                                                                                           {
  if (n <= 0) return;
  rows_scatter_kernel<<<n, 128, 0, stream>>>(src, idx, dst, w);
  check_launch("rows scatter");
}

References j, and n.

ir_scale_cuda()

void ember_ir::ir_scale_cuda	(	const float *	x,
		const float *	g,
		const float *	b,
		float *	out,
		int	T,
		int	w,
		cudaStream_t	stream
	)

Definition at line 748 of file graph_ir_cuda.cu.

                                                                                                                  {
  const size_t n = (size_t)T * w;
  scale_kernel<<<blocks_for(n), kThreads, 0, stream>>>(x, g, b, out, n, w);
  check_launch("scale");
}

References j, n, and out.

ir_schema_errors()

std::vector< std::string > ember_ir::ir_schema_errors

(

const json &

doc

)

Definition at line 1583 of file graph_ir_validate.cpp.

                                                       {
  Failures f;
  eval(schema(), doc, "doc", f);
  std::vector<std::string> out;
  std::set<std::string> seen;
  for (const auto& [p, m] : f) add_unique(out, seen, E(schema_rule(doc, p, m), p + " " + m));
  return out;
}

References ir_schema_errors(), j, out, and p.

Referenced by ir_document_errors(), and ir_schema_errors().

ir_selective_scan_cuda()

void ember_ir::ir_selective_scan_cuda	(	const float *	x,
		const float *	dt,
		const float *	B,
		const float *	C,
		const float *	A,
		const float *	d,
		float *	s,
		float *	y,
		int	T,
		int	D,
		int	N,
		bool	softplus,
		cudaStream_t	stream
	)

Definition at line 845 of file graph_ir_cuda.cu.

                                                 {
  selective_scan_kernel<<<(D + 127) / 128, 128, 0, stream>>>(x, dt, B, C, A, d, s, y, T, D, N, softplus ? 1 : 0);
  check_launch("selective_scan");
}

References j.

ir_sigmoid_cuda()

void ember_ir::ir_sigmoid_cuda	(	const float *	x,
		float *	out,
		size_t	n,
		cudaStream_t	stream
	)

Definition at line 732 of file graph_ir_cuda.cu.

                                                                                {
  unary_kernel<<<blocks_for(n), kThreads, 0, stream>>>(x, out, n, 1, 0.0f, 0.0f);
  check_launch("sigmoid");
}

References j, n, and out.

ir_silu_cuda()

void ember_ir::ir_silu_cuda	(	const float *	x,
		float *	out,
		size_t	n,
		cudaStream_t	stream
	)

Definition at line 814 of file graph_ir_cuda.cu.

                                                                             {
  silu_kernel<<<blocks_for(n), kThreads, 0, stream>>>(x, out, n);
  check_launch("silu");
}

References j, n, out, and silu_kernel().

ir_ssm_scan_cuda()

void ember_ir::ir_ssm_scan_cuda	(	const float *	u,
		const float *	A,
		const float *	B,
		const float *	C,
		const float *	D,
		float *	s,
		float *	y,
		int	T,
		int	n,
		cudaStream_t	stream
	)

Definition at line 776 of file graph_ir_cuda.cu.

                                                                   {
  if (n <= 0 || n > kIrSsmMax)
    throw std::runtime_error("IR CUDA ssm_scan: n " + std::to_string(n) + " (at most " + std::to_string(kIrSsmMax) +
                             ")");
  const int threads = n < kThreads ? (n + 31) / 32 * 32 : kThreads;
  ssm_scan_kernel<<<1, threads, (size_t)2 * n * sizeof(float), stream>>>(u, A, B, C, D, s, y, T, n);
  check_launch("ssm_scan");
}

References j, and n.

ir_turn_state_attention_cuda()

void ember_ir::ir_turn_state_attention_cuda	(	const float *	q,
		const float *	K,
		const float *	V,
		float *	out,
		int	T,
		int	row0,
		const int *	codes,
		int	hq,
		int	hkv,
		int	hd,
		float	scale,
		cudaStream_t	stream,
		const int *	drow
	)

Definition at line 685 of file graph_ir_cuda.cu.

                                                   {
  if (hkv <= 0 || hq % hkv != 0 || hd > kThreads * kWinDimsPerThread)
    throw std::runtime_error("IR CUDA turn_state attention: unsupported shape (head_dim " + std::to_string(hd) + ")");
  const size_t shm = (size_t)(hd + kWinChunk) * sizeof(float);
  turn_state_attention_kernel<<<dim3(T, hq), kThreads, shm, stream>>>(q, K, V, out, hq, hkv, hd, row0, codes, scale,
                                                                      drow);
  check_launch("turn_state attention");
}

References j, out, and scale.

ir_window_attention_cuda()

void ember_ir::ir_window_attention_cuda	(	const float *	q,
		const float *	K,
		const float *	V,
		float *	out,
		int	T,
		int	start_pos,
		int	hq,
		int	hkv,
		int	hd,
		float	scale,
		int	window,
		int	sinks,
		cudaStream_t	stream,
		const int *	dpos,
		float	softcap
	)

Definition at line 708 of file graph_ir_cuda.cu.

                                                              {
  if (window <= 0 || sinks < 0 || hkv <= 0 || hq % hkv != 0 || hd > kThreads * kWinDimsPerThread)
    throw std::runtime_error("IR CUDA window attention: unsupported shape (window " + std::to_string(window) +
                             ", head_dim " + std::to_string(hd) + ")");
  const size_t shm = (size_t)(hd + kWinChunk) * sizeof(float);
  window_attention_kernel<<<dim3(T, hq), kThreads, shm, stream>>>(q, K, V, out, hq, hkv, hd, start_pos, scale, window,
                                                                  sinks, dpos, softcap);
  check_launch("window attention");
}

References j, out, and scale.

linear_attention_row_cpu()

void ember_ir::linear_attention_row_cpu	(	const float *	q,
		const float *	k,
		const float *	v,
		float *	num,
		float *	den,
		float *	out,
		int	width,
		float	den_floor
	)

Definition at line 213 of file graph_ir.cpp.

                                                          {
  // num and den take token t before out[t] is formed (token t sees keys 0..t).
  for (int c = 0; c < width; ++c) {
    const float pk = elu_plus_one(k[c]);
    num[c] = num[c] + pk * v[c];
    den[c] = den[c] + pk;
    out[c] = elu_plus_one(q[c]) * num[c] / std::max(den[c], den_floor);
  }
}

References c, j, k, out, and v.

Referenced by ember_ir::Interpreter::run_nodes().

parse_ir_document()

nlohmann::json ember_ir::parse_ir_document

(

const std::string &

text

)

Definition at line 1474 of file graph_ir_validate.cpp.

                                              {
  // nlohmann's parse keeps the last of two members with one name; its callback sees every key as it is read,
  // so each open object keeps the names it has seen.
  std::vector<std::set<std::string>> open;
  std::string dup;
  int dup_depth = -1;
  const json::parser_callback_t cb = [&](int depth, json::parse_event_t ev, json& parsed) {
    if (ev == json::parse_event_t::object_start) {
      open.emplace_back();
    } else if (ev == json::parse_event_t::object_end) {
      if (!open.empty()) open.pop_back();
    } else if (ev == json::parse_event_t::key && dup_depth < 0 && !open.empty() && parsed.is_string() &&
               !open.back().insert(parsed.get<std::string>()).second) {
      dup = parsed.get<std::string>();
      dup_depth = depth;
    }
    return true;
  };
  json j = json::parse(text, cb);
  if (dup_depth >= 0)
    throw std::runtime_error(E("3.2 (rule 35)", "an object holds the member name " +
                                                    json(dup).dump(-1, ' ', false, json::error_handler_t::replace) +
                                                    " twice (nesting depth " + std::to_string(dup_depth) +
                                                    "): a document names each member of an object once"));
  return j;
}

References j, parse_ir_document(), and text.

Referenced by ember_ir::Interpreter::Interpreter(), parse_ir_document(), and tuned_llama_match().

sha256_hex_bytes()

std::string ember_ir::sha256_hex_bytes

(

const std::string &

bytes

)

Definition at line 454 of file ir_digest.cpp.

                                                   {
  uint32_t h[8] = {0x6a09e667, 0xbb67ae85, 0x3c6ef372, 0xa54ff53a, 0x510e527f, 0x9b05688c, 0x1f83d9ab, 0x5be0cd19};
  const uint64_t bits = static_cast<uint64_t>(bytes.size()) * 8;
  std::string msg = bytes;
  msg += static_cast<char>(0x80);
  while (msg.size() % 64 != 56) msg += '\0';
  for (int k = 7; k >= 0; --k) msg += static_cast<char>((bits >> (8 * k)) & 0xFF);
  for (size_t off = 0; off < msg.size(); off += 64) sha256_block(h, reinterpret_cast<const unsigned char*>(msg.data()) + off);
  static const char* hex = "0123456789abcdef";
  std::string out;
  for (uint32_t x : h)
    for (int k = 7; k >= 0; --k) out += hex[(x >> (4 * k)) & 15];
  return out;
}

References j, k, off, and out.

Referenced by ir_digest().

ssm_scan_row_cpu()

void ember_ir::ssm_scan_row_cpu	(	const float *	u,
		const float *	A,
		const float *	B,
		const float *	C,
		float	D,
		float *	s,
		float *	scratch,
		float *	y,
		int	n
	)

Definition at line 224 of file graph_ir.cpp.

                                                       {
  // ops::ssm_scan: all of s_t from s_{t-1}, then y from s_t.
  for (int i = 0; i < n; ++i) {
    float acc = 0;
    for (int j = 0; j < n; ++j) acc += A[static_cast<size_t>(i) * n + j] * s[j];
    acc += B[i] * u[i];
    scratch[i] = acc;
  }
  std::memcpy(s, scratch, static_cast<size_t>(n) * sizeof(float));
  for (int i = 0; i < n; ++i) y[i] = C[i] * s[i] + D * u[i];
}

References j, and n.

Referenced by ember_ir::Interpreter::run_nodes().

tuned_llama_match()

bool ember_ir::tuned_llama_match	(	const std::string &	ir_json,
		const GGUFData &	gguf,
		std::string *	why
	)

Definition at line 440 of file graph_ir.cpp.

                                                                                       {
  auto no = [&](const std::string& m) {
    if (why) *why = m;
    return false;
  };
  json j;
  try {
    j = parse_ir_document(ir_json);  // rule 35 too: the interpreter then refuses the document
  } catch (const std::exception& e) {
    return no(std::string("the document does not parse: ") + e.what());
  }
  if (!j.is_object() || j.value("family_hint", json()) != "llama") return no("family_hint is not llama");
  try {
    validate_ir(j, gguf);  // an invalid document goes to the interpreter, which refuses it
    LlamaMatch m(j, parse_model_config_from_gguf(gguf));
    if (!m.run()) return no(m.why);
  } catch (const std::exception& e) {
    return no(e.what());
  }
  if (why) why->clear();
  return true;
}

References j, parse_ir_document(), parse_model_config_from_gguf(), validate_ir(), and why.

Referenced by TinyLlamaModel::TinyLlamaModel(), and tinyllama::TinyLlamaSession::TinyLlamaSession().

validate_ir()

void ember_ir::validate_ir	(	const json &	doc,
		const GGUFData &	gguf
	)

Definition at line 1620 of file graph_ir_validate.cpp.

                                                        {
  std::vector<std::string> errs = ir_document_errors(doc);
  if (doc.is_object())
    for (const auto& m : ir_container_errors(doc, gguf)) errs.push_back(m);
  if (errs.empty()) return;
  std::string msg;
  for (const auto& e : errs) msg += (msg.empty() ? "" : "; ") + e;
  throw std::runtime_error(msg);
}

References ir_container_errors(), ir_document_errors(), j, and validate_ir().

Referenced by ember_ir::Interpreter::Interpreter(), tuned_llama_match(), and validate_ir().

Variable Documentation

kIrSchemaPieces

const char* const ember_ir::kIrSchemaPieces[]

inline

Initial value:

= {
    R"IRS({"$schema":"https:
    R"IRS(e":true,"kind":true,"op":{"enum":["linear_attention","ssm_scan","conv1d_causal","selective_scan","recurrence"]},"dtype":{"enum":["f32"]},"tensors":{"type":"array","minItems":1,"items":{"type":"object","required":["name","shape","init"],"properties":{"name":{"$ref":"#/$defs/name"},"shape":{"type":"array","minItems":1,"items":{"$ref":"#/$defs/posint"}},"init":{"enum":[0,"body"]}},"patternProperties":{"^_":true},"additionalProperties":false}}}

Definition at line 7 of file graph_ir_schema.h.

                                             {
    R"IRS({"$schema":"https://json-schema.org/draft/2020-12/schema","$id":"urn:paperkiln:ir:1","title":"paperkiln IR document, version 1","description":"The structure of an IR document (spec/IR_SPEC.md sections 3, 5, 6 and 7). Keys starting with '_' are prose and allowed anywhere. The semantic rules (section 10.1) are checked by tools/ir_validate.py.","type":"object","required":["ir_version","family_hint","model","embed","block","head"],"properties":{"ir_version":{"type":"integer","minimum":0},"family_hint":{"type":"string"},"model":{"$ref":"#/$defs/model"},"tokenizer":{"$ref":"#/$defs/tokenizer"},"inputs":{"type":"array","items":{"anyOf":[{"type":"string"},{"$ref":"#/$defs/io"}]}},"outputs":{"type":"array","items":{"anyOf":[{"type":"string"},{"$ref":"#/$defs/io"}]}},"modes":{"type":"object","additionalProperties":{"type":"string"}},"embed":{"$ref":"#/$defs/section"},"block":{"$ref":"#/$defs/block"},"head":{"$ref":"#/$defs/section"}},"patternProperties":{"^_":true},"additionalProperties":false,"$defs":{"name":{"type":"string","minLength":1},"names":{"type":"array","items":{"$ref":"#/$defs/name"}},"posint":{"type":"integer","minimum":1},"nonneg":{"type":"integer","minimum":0},"per_layer":{"type":"object","required":["per_layer"],"properties":{"per_layer":{"type":"array"}},"patternProperties":{"^_":true},"additionalProperties":false},"apply":{"anyOf":[{"type":"boolean"},{"$ref":"#/$defs/per_layer"}]},"int_or_layers":{"anyOf":[{"$ref":"#/$defs/nonneg"},{"$ref":"#/$defs/per_layer"}]},"num_or_layers":{"anyOf":[{"type":"number"},{"$ref":"#/$defs/per_layer"}]},"io":{"type":"object","required":["name"],"properties":{"name":{"type":"string"},"kind":{"type":"string"},"note":{"type":"string"}},"patternProperties":{"^_":true},"additionalProperties":false},"model":{"type":"object","required":["vocab","d","n_layers","ctx"],"properties":{"vocab":{"$ref":"#/$defs/posint"},"d":{"$ref":"#/$defs/posint"},"n_layers":{"$ref":"#/$defs/posint"},"ctx":{"$ref":"#/$defs/posint"},"dtype":{"enum":["f32"]}},"patternProperties":{"^_":true},"additionalProperties":false},"tokenizer":{"type":"object","properties":{"type":{"type":"string"},"source":{"type":"string"},"bos":{"type":["integer","null"]},"specials":{"type":"object","additionalProperties":{"type":"integer"}}},"patternProperties":{"^_":true},"additionalProperties":false},"section":{"type":"object","required":["nodes"],"properties":{"nodes":{"type":"array","items":{"$ref":"#/$defs/node"}}},"patternProperties":{"^_":true},"additionalProperties":false},"block":{"type":"object","required":["repeat","nodes"],"properties":{"repeat":{"$ref":"#/$defs/posint"},"state":{"type":"array","items":{"$ref":"#/$defs/slot"}},"nodes":{"type":"array","items":{"$ref":"#/$defs/node"}},"weight_index":{"$ref":"#/$defs/per_layer"}},"patternProperties":{"^_":true},"additionalProperties":false},"slot":{"type":"object","required":["name","kind"],"properties":{"name":{"$ref":"#/$defs/name"},"kind":{"enum":["kv","recurrent"]}},"allOf":[{"if":{"properties":{"kind":{"const":"kv"}},"required":["kind"]},"then":{"$ref":"#/$defs/slot_kv"}},{"if":{"properties":{"kind":{"const":"recurrent"}},"required":["kind"]},"then":{"$ref":"#/$defs/slot_recurrent"}}]},"slot_kv":{"type":"object","required":["name","kind","heads","kv_heads","head_dim","max_len","window","sinks"],"properties":{"name":true,"kind":true,"heads":{"anyOf":[{"$ref":"#/$defs/posint"},{"$ref":"#/$defs/per_layer"}]},"kv_heads":{"anyOf":[{"$ref":"#/$defs/posint"},{"$ref":"#/$defs/per_layer"}]},"head_dim":{"anyOf":[{"$ref":"#/$defs/posint"},{"$ref":"#/$defs/per_layer"}]},"max_len":{"anyOf":[{"$ref":"#/$defs/posint"},{"$ref":"#/$defs/per_layer"}]},"window":{"$ref":"#/$defs/int_or_layers"},"sinks":{"$ref":"#/$defs/int_or_layers"},"dtype":{"enum":["f32"]},"evict":{"enum":["turn_state"]},"kv_rows":{"enum":["active"]},"kv_from":true},"patternProperties":{"^_":true},"additionalProperties":false},"slot_recurrent":{"type":"object","required":["name","kind","op","dtype","tensors"],"properties":{"nam)IRS",
    R"IRS(e":true,"kind":true,"op":{"enum":["linear_attention","ssm_scan","conv1d_causal","selective_scan","recurrence"]},"dtype":{"enum":["f32"]},"tensors":{"type":"array","minItems":1,"items":{"type":"object","required":["name","shape","init"],"properties":{"name":{"$ref":"#/$defs/name"},"shape":{"type":"array","minItems":1,"items":{"$ref":"#/$defs/posint"}},"init":{"enum":[0,"body"]}},"patternProperties":{"^_":true},"additionalProperties":false}}},"patternProperties":{"^_":true},"additionalProperties":false},"node":{"type":"object","required":["op","in","out"],"properties":{"op":{"enum":["embedding","pos_embedding","linear","scale","split","rmsnorm","layernorm","gelu","silu","sigmoid","relu","add","mul","sub","affine","mul_col","rms_row","rope","attention","moe_ffn","list","append","depth_attend","expr","linear_attention","ssm_scan","conv1d_causal","selective_scan","depth_route","recurrence","mla_attention"]},"in":{"$ref":"#/$defs/names"}},"allOf":[{"if":{"properties":{"op":{"const":"embedding"}}},"then":{"$ref":"#/$defs/op_embedding"}},{"if":{"properties":{"op":{"const":"pos_embedding"}}},"then":{"$ref":"#/$defs/op_weight_only"}},{"if":{"properties":{"op":{"const":"linear"}}},"then":{"$ref":"#/$defs/op_linear"}},{"if":{"properties":{"op":{"const":"scale"}}},"then":{"$ref":"#/$defs/op_scale"}},{"if":{"properties":{"op":{"const":"split"}}},"then":{"$ref":"#/$defs/op_split"}},{"if":{"properties":{"op":{"const":"rmsnorm"}}},"then":{"$ref":"#/$defs/op_rmsnorm"}},{"if":{"properties":{"op":{"const":"layernorm"}}},"then":{"$ref":"#/$defs/op_layernorm"}},{"if":{"properties":{"op":{"const":"gelu"}}},"then":{"$ref":"#/$defs/op_gelu"}},{"if":{"properties":{"op":{"enum":["silu","sigmoid","relu","add","mul","sub","mul_col","list","append"]}}},"then":{"$ref":"#/$defs/op_plain"}},{"if":{"properties":{"op":{"const":"affine"}}},"then":{"$ref":"#/$defs/op_affine"}},{"if":{"properties":{"op":{"const":"rms_row"}}},"then":{"$ref":"#/$defs/op_rms_row"}},{"if":{"properties":{"op":{"const":"rope"}}},"then":{"$ref":"#/$defs/op_rope"}},{"if":{"properties":{"op":{"const":"attention"}}},"then":{"$ref":"#/$defs/op_attention"}},{"if":{"properties":{"op":{"const":"moe_ffn"}}},"then":{"$ref":"#/$defs/op_moe_ffn"}},{"if":{"properties":{"op":{"const":"depth_attend"}}},"then":{"$ref":"#/$defs/op_depth_attend"}},{"if":{"properties":{"op":{"const":"expr"}}},"then":{"$ref":"#/$defs/op_expr"}},{"if":{"properties":{"op":{"const":"linear_attention"}}},"then":{"$ref":"#/$defs/op_linear_attention"}},{"if":{"properties":{"op":{"const":"ssm_scan"}}},"then":{"$ref":"#/$defs/op_ssm_scan"}},{"if":{"properties":{"op":{"const":"conv1d_causal"}}},"then":{"$ref":"#/$defs/op_conv1d_causal"}},{"if":{"properties":{"op":{"const":"selective_scan"}}},"then":{"$ref":"#/$defs/op_selective_scan"}},{"if":{"properties":{"op":{"const":"depth_route"}}},"then":{"$ref":"#/$defs/op_depth_route"}},{"if":{"properties":{"op":{"const":"recurrence"}}},"then":{"$ref":"#/$defs/op_recurrence"}},{"if":{"properties":{"op":{"const":"mla_attention"}}},"then":{"$ref":"#/$defs/op_mla_attention"}}]},"op_embedding":{"required":["weight"],"properties":{"op":true,"in":true,"apply":{"$ref":"#/$defs/apply"},"out":{"$ref":"#/$defs/name"},"tensor_index":{"enum":["application"]},"weight":{"$ref":"#/$defs/name"},"state_weight":{"$ref":"#/$defs/name"}},"patternProperties":{"^_":true},"additionalProperties":false},"op_weight_only":{"required":["weight"],"properties":{"op":true,"in":true,"apply":{"$ref":"#/$defs/apply"},"out":{"$ref":"#/$defs/name"},"tensor_index":{"enum":["application"]},"weight":{"$ref":"#/$defs/name"}},"patternProperties":{"^_":true},"additionalProperties":false},"op_linear":{"required":["weight"],"properties":{"op":true,"in":true,"apply":{"$ref":"#/$defs/apply"},"out":{"$ref":"#/$defs/name"},"tensor_index":{"enum":["application"]},"weight":{"$ref":"#/$defs/name"},"bias":{"type":["string","null"]},"tied_to":{"$ref":"#/$defs/name"}},"patternProperties":{"^_":true},"additionalProperties":false},"op_scale":{"r)IRS",

patternProperties

const char* const ember_ir::patternProperties

Definition at line 9 of file graph_ir_schema.h.