#include <graph_ir.h>

Collaboration diagram for ember_ir::Interpreter:

Classes
struct	AppIndex

struct	Buf

struct	RecSlot

struct	RecTensor

struct	RecurrentTensorInfo

struct	State

struct	Weight

Public Member Functions
	Interpreter (const std::string &ir_json, const GGUFData &gguf, bool use_gpu=false)

	~Interpreter ()

	Interpreter (const Interpreter &)=delete

Interpreter &	operator= (const Interpreter &)=delete

std::vector< float >	forward (const std::vector< int > &tokens, int start_pos, bool all_rows=false, int tail_rows=0)

bool	has_recurrent () const

int	resume_point (int want) const

int	vocab () const

int	max_len () const

int	n_layers () const

const std::string &	family_hint () const

const std::vector< int > &	last_depths () const

std::vector< int >	held_positions (int application) const

bool	on_gpu () const

const char *	kv_format () const

size_t	kv_cache_bytes () const

size_t	kv_scratch_bytes () const

int	state_k () const

bool	turn_state () const

bool	evicts () const

int	kv_rows_held () const

const std::string &	kv_slot () const

bool	evict_declared () const

int	seq_len () const

bool	has_depth_route () const

std::vector< int >	kv_positions (int application) const

std::vector< RecurrentTensorInfo >	recurrent_tensors () const

std::vector< float >	recurrent_values (const std::string &slot, const std::string &tensor, int application) const

void	truncate (int n)

State	save_state () const

void	load_state (const State &s)

Private Types
using	Env = std::unordered_map< std::string, Buf >

Private Member Functions
void	init (const nlohmann::json &j, const std::string &ir_json, const GGUFData &gguf, bool use_gpu)

int	kv_rows_in_use (int application) const

void	run_nodes (const std::vector< Node > &nodes, int layer, Env &env, int T, int start_pos, bool head_last_only)

void	attention_window (const Node &n, int layer, int &window, int &sinks) const

const std::vector< std::pair< float, float > > &	rope_table (int layer) const

const std::vector< float > &	weight (const std::string &name) const

const Weight &	bound (const std::string &name) const

std::string	layer_name (const std::string &pattern, int layer) const

void	bind (const std::vector< Node > &nodes, int layer, const GGUFData &gguf)

void	bind_tensor (const std::string &name, const GGUFData &gguf, bool keep_quantised=false)

void	mor_step (int layer, Env &env, int T, int start_pos)

void	ts_begin (const std::vector< int > &tokens, int start_pos)

void	ts_evict_below (int seg)

int	ts_seg_after (int p) const

int	ts_segment_start (int m) const

void	build_rope_tables (const GGUFData &gguf)

void	set_rope_long (bool v)

std::vector< float >	forward_impl (const std::vector< int > &tokens, int start_pos, bool all_rows, int tail_rows)

const RecSlot &	rec_slot (const Node &n, const char *op) const

float *	rec_tensor (const RecSlot &slot, const char *tensor, int layer, size_t floats)

void	rec_begin (int start_pos, int T, bool all_rows)

void	rec_row_done (const RecSlot &slot, int layer, int t)

void	rec_end (int start_pos, int T)

Private Attributes
std::string	kv_slot_

bool	evict_declared_ = false

int	seq_len_ = 0

bool	app_index_ = false

std::vector< Node >	embed_

std::vector< Node >	block_

std::vector< Node >	head_

std::map< std::pair< const Node *, int >, int >	da_sources_

int	n_layers_ = 0

int	vocab_ = 0

int	max_len_ = 0

std::vector< int >	weight_map_

std::vector< int >	mor_step_

std::vector< std::vector< int > >	mor_rows_

int	mor_R_ = 0

bool	mor_gate_score_ = false

std::string	mor_x_

std::string	mor_out_

std::string	mor_router_

std::vector< int >	mor_depth_

std::vector< float >	mor_g_

std::vector< float >	mor_xin_

int	attn_row_ = -1

int	state_k_ = 0

std::string	state_weight_

bool	turn_state_ = false

bool	evict_ = false

std::vector< int >	ts_seg_

std::vector< int >	ts_sidx_

std::vector< int >	kv_pos_

std::vector< std::vector< int > >	ts_vis_

int	kv_row0_ = 0

int	evicted_below_ = 0

std::string	family_hint_

bool	dequant_all_ = false

std::unordered_map< std::string, Weight >	weights_

std::unordered_map< std::string, std::vector< int64_t > >	shapes_

std::vector< KVCacheLayer >	kv_

int	kv_heads_ = 0

int	head_dim_ = 0

int	kv_fmt_ = 0

int	window_ = 0

int	sinks_ = 0

float	rope_theta_ = 0.0f

int	rope_dim_ = 0

bool	rope_neox_ = false

float	rope_freq_scale_ = 1.0f

std::vector< int >	window_layer_

std::vector< float >	rope_theta_layer_

std::vector< float >	rope_scale_layer_

std::vector< float >	rope_attn_layer_

std::vector< std::vector< std::pair< float, float > > >	rope_tabs_

std::vector< int >	rope_tab_short_

std::vector< int >	rope_tab_long_

nlohmann::json	rope_ff_

int	rope_ff_ctx_ = 0

bool	rope_long_ = false

std::vector< int >	ids_

std::set< std::string >	offset_applied_

bool	gpu_ = false

std::vector< RecSlot >	rec_slots_

size_t	rec_floats_ = 0

std::vector< float >	rec_init_

std::map< const Node *, std::shared_ptr< rec::Recurrence > >	recs_

std::vector< float >	rec_

int	rec_pos_ = 0

std::map< int, std::vector< float > >	rec_snaps_

size_t	rec_max_snaps_ = 0

std::vector< int >	rec_snap_rows_

std::vector< float >	rec_rows_

Detailed Description

Definition at line 102 of file graph_ir.h.

Member Typedef Documentation

◆ Env

using ember_ir::Interpreter::Env = std::unordered_map<std::string, Buf>

private

Definition at line 230 of file graph_ir.h.

Constructor & Destructor Documentation

◆ Interpreter() [1/2]

ember_ir::Interpreter::Interpreter	(	const std::string &	ir_json,
		const GGUFData &	gguf,
		bool	use_gpu = `false`
	)

Definition at line 1099 of file graph_ir.cpp.

                                                                                     {
  const json j = parse_ir_document(ir_json);  // a member name twice is refused (IR_SPEC 3.2, rule 35)
  try {
    init(j, ir_json, gguf, use_gpu);
  } catch (const json::exception&) {
    // ember's own checks read the document as its schema shapes it: a document that is not is
    // refused naming the rules it breaks (IR_SPEC 10.1), not a JSON accessor's message.
    validate_ir(j, gguf);
    throw;
  }
}

References init(), j, ember_ir::parse_ir_document(), and ember_ir::validate_ir().

◆ ~Interpreter()

ember_ir::Interpreter::~Interpreter ( )

Definition at line 1759 of file graph_ir.cpp.

                          {
#ifdef HAS_CUDA
  free_device();
#endif
}

References j.

◆ Interpreter() [2/2]

ember_ir::Interpreter::Interpreter ( const Interpreter & )

delete

Member Function Documentation

◆ attention_window()

void ember_ir::Interpreter::attention_window	(	const Node &	n,
		int	layer,
		int &	window,
		int &	sinks
	)		const

private

Definition at line 1783 of file graph_ir.cpp.

                                                                                          {
  window = has(n.attrs, "window") ? at_layer(n.attrs.at("window"), layer).get<int>()
                                  : (window_layer_.empty() ? window_ : window_layer_.at(layer));
  sinks = has(n.attrs, "sinks") ? at_layer(n.attrs.at("sinks"), layer).get<int>() : sinks_;
}

References j, and n.

Referenced by init().

◆ bind()

void ember_ir::Interpreter::bind	(	const std::vector< Node > &	nodes,
		int	layer,
		const GGUFData &	gguf
	)

private

Definition at line 1948 of file graph_ir.cpp.

                                                                                    {
  for (const auto& n : nodes) {
    if (!applies(n, layer)) continue;  // its tensors are needed only where it runs (IR_SPEC 3.8, revision 1.6)
    AppIndex ti(app_index_, n);
    // depth_attend over one source never reads its pseudo-query (IR_SPEC 5.18), so a file may leave it
    // out; with S >= 2 sources a missing one is refused here, at load (rule 25).
    if (n.op == "depth_attend" && has(n.attrs, "weight")) {
      const std::string w = layer_name(n.attrs.at("weight").get<std::string>(), layer);
      if (!gguf.tensor_infos_map.count(w)) {
        const auto it = da_sources_.find({&n, layer});
        const int S = it == da_sources_.end() ? 0 : it->second;
        if (S == 1) continue;
        throw std::runtime_error("IR_SPEC 5.18: depth_attend" +
                                 (layer >= 0 ? " at application " + std::to_string(layer) : std::string()) + " has " +
                                 std::to_string(S) + " sources, and its pseudo-query '" + w + "' is not in the GGUF");
      }
    }
    if (has(n.attrs, "weight")) {
      std::string w = layer_name(n.attrs.at("weight").get<std::string>(), layer);
      if (!gguf.tensor_infos_map.count(w) && has(n.attrs, "tied_to"))
        w = layer_name(n.attrs.at("tied_to").get<std::string>(), layer);
      bind_tensor(w, gguf, n.op == "linear" || n.op == "embedding");
    }
    if (has(n.attrs, "bias")) bind_tensor(layer_name(n.attrs.at("bias").get<std::string>(), layer), gguf);
    if (n.op == "attention" && has(n.attrs, "position_bias") && n.attrs.at("position_bias").is_object() &&
        n.attrs.at("position_bias").contains("weight") && n.attrs.at("position_bias").at("weight").is_string())
      bind_tensor(layer_name(n.attrs.at("position_bias").at("weight").get<std::string>(), layer), gguf);
    if (n.op == "attention" && has(n.attrs, "relative_kv") && n.attrs.at("relative_kv").is_object())
      for (const char* member : {"key", "value"})
        if (n.attrs.at("relative_kv").contains(member) && n.attrs.at("relative_kv").at(member).is_string())
          bind_tensor(layer_name(n.attrs.at("relative_kv").at(member).get<std::string>(), layer), gguf);
    // rmsnorm weight_offset c: the weight is read as w + c, rounded to f32 once at load (Gemma's (1 + w)
    // for a file holding HF's raw w; llama.cpp's convert already stores 1 + w).
    if (n.op == "rmsnorm" && has(n.attrs, "weight_offset") && n.attrs.at("weight_offset").get<float>() != 0.0f) {
      const std::string w = layer_name(n.attrs.at("weight").get<std::string>(), layer);
      if (offset_applied_.insert(w).second) {
        const float off = n.attrs.at("weight_offset").get<float>();
        for (float& x : weights_.at(w).f32) x = x + off;
      }
    }
    if (n.op == "moe_ffn") {
      const std::string norm = n.attrs.value("gate_norm", "renorm"), act = n.attrs.value("activation", "silu");
      if (norm != "renorm" && norm != "raw") throw std::runtime_error("IR moe_ffn gate_norm '" + norm + "' is not implemented");
      if (act != "silu") throw std::runtime_error("IR moe_ffn activation '" + act + "' is not implemented");
      bind_tensor(layer_name(n.attrs.at("router").get<std::string>(), layer), gguf);
      for (const char* part : {"gate", "up", "down"})
        bind_tensor(layer_name(n.attrs.at("experts").at(part).get<std::string>(), layer), gguf);
      if (has(n.attrs, "selection_bias")) {  // ne {n_experts} (IR_SPEC 5.16, rule 44, revision 1.6)
        const std::string b = layer_name(n.attrs.at("selection_bias").get<std::string>(), layer);
        bind_tensor(b, gguf);
        const int E = at_layer(n.attrs.at("n_experts"), std::max(layer, 0)).get<int>();
        if (bound(b).numel != static_cast<size_t>(E) || shapes_.at(b).size() != 1)
          throw std::runtime_error("IR_SPEC 5.16 (rule 44): moe_ffn selection_bias '" + b + "' is not of ne {n_experts} = {" +
                                   std::to_string(E) + "}");
      }
    }
    if (n.op == "mla_attention") {  // k_b ne {dn, r, H}, v_b ne {r, dv, H} (IR_SPEC 5.20, rule 42)
      const auto at = [&](const char* k) { return static_cast<int64_t>(at_layer(n.attrs.at(k), std::max(layer, 0)).get<int>()); };
      const int64_t H = at("heads"), r = at("kv_rank"), dn = at("qk_nope"), dv = at("v_head_dim");
      const std::pair<const char*, std::vector<int64_t>> parts[] = {{"k_b", {dn, r, H}}, {"v_b", {r, dv, H}}};
      for (const auto& [key, want] : parts) {
        const std::string w = layer_name(n.attrs.at(key).get<std::string>(), layer);
        bind_tensor(w, gguf);
        std::vector<int64_t> got = shapes_.at(w);
        while (got.size() < 3) got.push_back(1);  // trailing ne of 1 (one head) may be left out
        if (got != want)
          throw std::runtime_error("IR_SPEC 5.20 (rule 42): mla_attention " + std::string(key) + " '" + w +
                                   "' is not of ne {" + std::to_string(want[0]) + ", " + std::to_string(want[1]) + ", " +
                                   std::to_string(want[2]) + "}");
      }
    }
    if (n.op == "ssm_scan") {
      const size_t sn = n.attrs.at("n").get<size_t>();
      const std::pair<const char*, size_t> parts[] = {{"a", sn * sn}, {"b", sn}, {"c", sn}, {"d", 1}};
      for (const auto& [part, numel] : parts) {
        const std::string w = layer_name(ssm_weight(n.attrs, part), layer);
        bind_tensor(w, gguf);
        if (bound(w).numel != numel)
          throw std::runtime_error("IR ssm_scan '" + w + "' has " + std::to_string(bound(w).numel) +
                                   " elements, expected " + std::to_string(numel));
      }
    }
    auto expect = [&](const std::string& w, size_t numel) {
      if (bound(w).numel != numel)
        throw std::runtime_error("IR " + n.op + " '" + w + "' has " + std::to_string(bound(w).numel) +
                                 " elements, expected " + std::to_string(numel));
    };
    if (n.op == "conv1d_causal") {  // weight [channels, kernel], bias [channels] (bound above)
      const size_t C = n.attrs.at("channels").get<size_t>(), K = n.attrs.at("kernel").get<size_t>();
      expect(layer_name(n.attrs.at("weight").get<std::string>(), layer), C * K);
      if (has(n.attrs, "bias")) expect(layer_name(n.attrs.at("bias").get<std::string>(), layer), C);
    }
    if (n.op == "selective_scan") {  // a [channels, d_state] as stored, d [channels] optional
      const size_t D = n.attrs.at("channels").get<size_t>(), N = n.attrs.at("d_state").get<size_t>();
      const std::string a = layer_name(n.attrs.at("a").get<std::string>(), layer);
      bind_tensor(a, gguf);
      expect(a, D * N);
      if (has(n.attrs, "d")) {
        const std::string d = layer_name(n.attrs.at("d").get<std::string>(), layer);
        bind_tensor(d, gguf);
        expect(d, D);
      }
    }
    if (n.op == "recurrence") {  // each weight as many elements as its declaration, read row-major (IR_SPEC 6.10.1)
      const auto it = recs_.find(&n);
      if (it == recs_.end()) throw std::runtime_error("IR recurrence: the node was not compiled");
      for (const auto& wd : it->second->weights()) {
        const std::string w = layer_name(wd.tensor, layer);
        bind_tensor(w, gguf);
        if (bound(w).numel != wd.numel)
          throw std::runtime_error("IR_SPEC 6.10 (rule 34): recurrence weight " + wd.decl + " reads '" + w + "', which has " +
                                   std::to_string(bound(w).numel) + " elements; its declaration has " +
                                   std::to_string(wd.numel));
      }
    }
  }
}

References a, j, k, n, numel, and off.

◆ bind_tensor()

void ember_ir::Interpreter::bind_tensor	(	const std::string &	name,
		const GGUFData &	gguf,
		bool	keep_quantised = `false`
	)

private

Definition at line 1882 of file graph_ir.cpp.

                                                                                              {
  if (auto bw = weights_.find(name); bw != weights_.end()) {
    Weight& w = bw->second;
    if (keep_quantised || !w.quantised()) return;
    // Bound quantised for a row-wise op, now needed whole (a norm, a bias): f32.
    w.f32.resize(w.numel);
    dequant_span_to_f32(w.type, w.qp, 0, w.numel, w.f32.data());
    std::vector<uint8_t>().swap(w.q);
    w.qp = nullptr;
    w.qbytes = 0;
    return;
  }
  auto it = gguf.tensor_infos_map.find(name);
  if (it == gguf.tensor_infos_map.end()) throw std::runtime_error("IR weight '" + name + "' is not in the GGUF");
  const GGUFTensorInfo& info = it->second;
  const uint8_t* base = gguf.mapped_tensor_data
                            ? static_cast<const uint8_t*>(gguf.mapped_tensor_data) + gguf.offset_diff_for_mmap
                            : gguf.tensor_data.data();
  const uint8_t* src = base + info.offset;
  std::vector<int64_t> shape(info.shape.begin(), info.shape.end());
  const GGMLType t = info.type;
  const bool k_quant = t == GGMLType::GGML_TYPE_Q4_K || t == GGMLType::GGML_TYPE_Q5_K || t == GGMLType::GGML_TYPE_Q6_K;
  const bool block_type = t == GGMLType::GGML_TYPE_Q8_0 || t == GGMLType::GGML_TYPE_Q4_0 ||
                          t == GGMLType::GGML_TYPE_Q4_1 || t == GGMLType::GGML_TYPE_Q5_0 || t == GGMLType::GGML_TYPE_Q5_1 ||
                          k_quant;
  const int64_t width = k_quant ? GGML_QK_K : 32;  // elements per block
  const bool half = t == GGMLType::GGML_TYPE_F16 || t == GGMLType::GGML_TYPE_BF16;
  if (keep_quantised && !dequant_all_ && !shape.empty() && (half || (block_type && shape[0] % width == 0))) {
    // Rows (ne0 elements) start on block boundaries, so a row is a whole
    // number of blocks and dequantises on its own.
    const size_t bytes = half ? info.num_elements * 2 : info.num_elements / width * ggml_type_size(t);
    if (info.size_in_bytes < bytes)
      throw std::runtime_error("IR weight '" + name + "' is shorter than its element count implies");
    Weight w;
    if (gguf.mapped_tensor_data) {
      w.qp = src;  // read in place: the mapping lives as long as the GGUFData
    } else {
      w.q.assign(src, src + bytes);
      w.qp = w.q.data();  // a vector's buffer survives the moves below
    }
    w.qbytes = bytes;
    w.type = t;
    w.numel = info.num_elements;
    shapes_[name] = shape;
    weights_[name] = std::move(w);
    return;
  }
  std::vector<float> out;
  if (info.type == GGMLType::GGML_TYPE_F32) {
    out.resize(info.num_elements);
    std::memcpy(out.data(), src, info.num_elements * sizeof(float));
  } else if (info.type == GGMLType::GGML_TYPE_Q8_0) {
    std::vector<block_q8_0> blocks(info.size_in_bytes / sizeof(block_q8_0));
    std::memcpy(blocks.data(), src, blocks.size() * sizeof(block_q8_0));
    dequantize_vector_q8_0_to_f32(blocks, out, info.num_elements);
  } else if (!dequant_legacy_to_f32(info, src, out)) {
    throw std::runtime_error("IR weight '" + name + "': tensor type " + ggml_type_name(info.type) +
                             " is not supported by ember's interpreter");
  }
  shapes_[name] = shape;
  Weight w;
  w.numel = out.size();
  w.f32 = std::move(out);
  weights_[name] = std::move(w);
}

References dequant_legacy_to_f32(), dequant_span_to_f32(), dequantize_vector_q8_0_to_f32(), ember_ir::Interpreter::Weight::f32, GGML_QK_K, GGML_TYPE_BF16, GGML_TYPE_F16, GGML_TYPE_F32, ggml_type_name(), GGML_TYPE_Q4_0, GGML_TYPE_Q4_1, GGML_TYPE_Q4_K, GGML_TYPE_Q5_0, GGML_TYPE_Q5_1, GGML_TYPE_Q5_K, GGML_TYPE_Q6_K, GGML_TYPE_Q8_0, ggml_type_size(), j, name, GGUFTensorInfo::num_elements, ember_ir::Interpreter::Weight::numel, GGUFTensorInfo::offset, out, ember_ir::Interpreter::Weight::qp, GGUFTensorInfo::shape, shape, GGUFTensorInfo::size_in_bytes, and GGUFTensorInfo::type.

◆ bound()

const Interpreter::Weight & ember_ir::Interpreter::bound ( const std::string & name ) const

private

Definition at line 2345 of file graph_ir.cpp.

                                                                       {
  auto it = weights_.find(name);
  if (it == weights_.end()) throw std::runtime_error("IR weight '" + name + "' was not bound");
  return it->second;
}

References j, and name.

◆ build_rope_tables()

void ember_ir::Interpreter::build_rope_tables ( const GGUFData & gguf )

private

Definition at line 1802 of file graph_ir.cpp.

                                                        {
  rope_tabs_.clear();
  rope_tab_short_.clear();
  rope_tab_long_.clear();
  rope_ff_ctx_ = 0;
  if (rope_theta_layer_.empty()) return;
  const int half = rope_dim_ / 2;
  std::string short_name, long_name;  // tensor names as written ({i} allowed)
  std::vector<float> values;           // freq_factors given by value (revision 1.6: YaRN's divisors), each as f32
  if (rope_ff_.is_array()) {           // n_rot / 2 positive numbers (rule 41, validate_ir)
    for (const auto& v : rope_ff_) values.push_back(v.get<float>());
    if (values.size() != static_cast<size_t>(half))
      throw std::runtime_error("IR_SPEC 5.14 (rule 41): rope freq_factors holds " + std::to_string(values.size()) +
                               " values, not n_rot / 2 = " + std::to_string(half));
    short_name = "=values";  // the table's key; no tensor is read
  } else if (rope_ff_.is_string()) {
    short_name = rope_ff_.get<std::string>();
  } else if (rope_ff_.is_object()) {  // {short, long, original_ctx}: rule 41 holds (validate_ir)
    short_name = rope_ff_.at("short").get<std::string>();
    long_name = rope_ff_.at("long").get<std::string>();
    rope_ff_ctx_ = rope_ff_.at("original_ctx").get<int>();
  }
  std::map<std::tuple<float, float, float, std::string>, int> index;
  auto table = [&](int l, const std::string& pattern) -> int {
    const float theta = rope_theta_layer_[l], fs = rope_scale_layer_[l], m = rope_attn_layer_.empty() ? 1.0f : rope_attn_layer_[l];
    if (!(theta > 0.0f)) return -1;
    const std::string name = pattern.empty() || pattern == "=values" ? pattern : layer_name(pattern, l);
    const auto key = std::make_tuple(theta, fs, m, name);
    if (const auto it = index.find(key); it != index.end()) return it->second;
    std::vector<float> f;
    if (pattern == "=values") {
      f = values;
    } else if (!name.empty()) {
      bind_tensor(name, gguf);
      f = weight(name);
      if (f.size() != static_cast<size_t>(half))
        throw std::runtime_error("IR_SPEC 5.14 (rule 41): rope freq_factors '" + name + "' holds " +
                                 std::to_string(f.size()) + " values, not n_rot / 2 = " + std::to_string(half));
    }
    std::vector<std::pair<float, float>> tab(static_cast<size_t>(max_len_) * half);
    for (int pos = 0; pos < max_len_; ++pos)
      for (int i = 0; i < rope_dim_; i += 2) {
        const float freq = 1.0f / std::pow(theta, float(i) / rope_dim_);
        float val = (static_cast<float>(pos) * fs) * freq;
        if (!f.empty()) val = val / f[i / 2];
        tab[static_cast<size_t>(pos) * half + i / 2] = {std::cos(val) * m, std::sin(val) * m};
      }
    rope_tabs_.push_back(std::move(tab));
    index[key] = static_cast<int>(rope_tabs_.size()) - 1;
    return index[key];
  };
  for (int l = 0; l < static_cast<int>(rope_theta_layer_.size()); ++l) {
    rope_tab_short_.push_back(table(l, short_name));
    if (!long_name.empty()) rope_tab_long_.push_back(table(l, long_name));
  }
  if (rope_ff_ctx_ > 0)
    Logger::info("[IR] LongRoPE: frequency factors '" + short_name + "' up to " + std::to_string(rope_ff_ctx_) +
                 " positions, '" + long_name + "' past them; the set follows the sequence's length (a call or a "
                 "truncate that changes it runs the sequence again)");
}

References index, Logger::info(), j, name, pos, and v.

◆ evict_declared()

bool ember_ir::Interpreter::evict_declared ( ) const

inline

Definition at line 181 of file graph_ir.h.

181{ return evict_declared_; }

ember_ir::Interpreter::evict_declared_

bool evict_declared_

Definition graph_ir.h:222

References evict_declared_.

Referenced by init().

◆ evicts()

bool ember_ir::Interpreter::evicts ( ) const

inline

Definition at line 169 of file graph_ir.h.

169{ return evict_; }

ember_ir::Interpreter::evict_

bool evict_

Definition graph_ir.h:302

References evict_.

◆ family_hint()

const std::string & ember_ir::Interpreter::family_hint ( ) const

inline

Definition at line 138 of file graph_ir.h.

138{ return family_hint_; }

ember_ir::Interpreter::family_hint_

std::string family_hint_

Definition graph_ir.h:306

References family_hint_.

◆ forward()

std::vector< float > ember_ir::Interpreter::forward	(	const std::vector< int > &	tokens,
		int	start_pos,
		bool	all_rows = `false`,
		int	tail_rows = `0`
	)

Definition at line 2957 of file graph_ir.cpp.

                                                                                                               {
  const int T = static_cast<int>(tokens.size());
  if (T == 0) return {};
  if (rope_ff_ctx_ <= 0) return forward_impl(tokens, start_pos, all_rows, tail_rows);
  // LongRoPE with two sets (IR_SPEC 5.14, 6.9; revision 1.5): n = start_pos + T selects the long set when
  // n > original_ctx, for this call's rows and every row the kv slot holds.
  if (start_pos < 0 || start_pos > static_cast<int>(ids_.size()))
    throw std::runtime_error("IR forward at position " + std::to_string(start_pos) + ": the sequence holds " +
                             std::to_string(ids_.size()) + " ids");
  const bool want = start_pos + T > rope_ff_ctx_;
  if (want == rope_long_ || start_pos == 0) {
    set_rope_long(want);
    std::vector<float> out = forward_impl(tokens, start_pos, all_rows, tail_rows);
    ids_.resize(start_pos);
    ids_.insert(ids_.end(), tokens.begin(), tokens.end());
    return out;
  }
  // The held rows were rotated with the other set: the whole sequence again from position 0, as Phi-3's own
  // generation recomputes its cache when the length first passes original_ctx; the same rows of logits come back.
  std::vector<int> all(ids_.begin(), ids_.begin() + start_pos);
  all.insert(all.end(), tokens.begin(), tokens.end());
  set_rope_long(want);
  Logger::info("[IR] LongRoPE: the sequence's length " + std::to_string(all.size()) + (want ? " passes " : " is within ") +
               std::to_string(rope_ff_ctx_) + ": the " + (want ? "long" : "short") +
               " set, every kv row recomputed from position 0");
  std::vector<float> out;
  if (all_rows) {
    out = forward_impl(all, 0, true, 0);
    const size_t width = out.size() / all.size();  // the head's row width
    out.erase(out.begin(), out.end() - static_cast<size_t>(T) * width);
  } else {
    out = forward_impl(all, 0, false, std::min(T, std::max(1, tail_rows)));
  }
  ids_ = std::move(all);
  return out;
}

References Logger::info(), j, and out.

◆ forward_impl()

std::vector< float > ember_ir::Interpreter::forward_impl	(	const std::vector< int > &	tokens,
		int	start_pos,
		bool	all_rows,
		int	tail_rows
	)

private

Definition at line 2994 of file graph_ir.cpp.

                                                            {
  const int T = static_cast<int>(tokens.size());
  if (T == 0) return {};
  if (rec_floats_ > 0 && (start_pos < 0 || start_pos + T > max_len_))
    throw std::runtime_error("IR recurrent state: positions " + std::to_string(start_pos) + ".." +
                             std::to_string(start_pos + T - 1) + " exceed model.ctx (" + std::to_string(max_len_) + ")");
#ifdef HAS_CUDA
  if (gpu_) {
    std::vector<float> out = forward_device(tokens, start_pos, all_rows, tail_rows);
    seq_len_ = start_pos + T;
    return out;
  }
#endif
  if (state_k_ > 0 || turn_state_) ts_begin(tokens, start_pos);
  if (rec_floats_ > 0) rec_begin(start_pos, T, all_rows);
  Env env;
  env["tokens"] = Buf{std::vector<float>(tokens.begin(), tokens.end()), 1};
  run_nodes(embed_, -1, env, T, start_pos, false);
  if (!mor_step_.empty()) {
    // Resuming at start_pos (a new sequence, prefix reuse, a rewind): each compact cache keeps
    // the rows of positions below it.
    for (auto& rows : mor_rows_)
      while (!rows.empty() && rows.back() >= start_pos) rows.pop_back();
    mor_depth_.clear();
  }
  for (int l = 0; l < n_layers_; ++l) {
    if (!mor_step_.empty() && mor_step_[l] >= 0)
      mor_step(l, env, T, start_pos);
    else
      run_nodes(block_, l, env, T, start_pos, false);
  }
  if (turn_state_) {
    // Every block has written this call's rows at kv_row0_ ..: they now hold these positions.
    for (int t = 0; t < T; ++t) kv_pos_.push_back(start_pos + t);
    // The text rows of segments before the last token's are invisible to every later query.
    if (evict_) ts_evict_below(ts_seg_.back());
  }
  if (rec_floats_ > 0) rec_end(start_pos, T);
  // Head ops are row-wise: run them on the last row only, or on the last tail_rows rows.
  const int rows = all_rows ? T : std::min(T, std::max(1, tail_rows));
  if (rows < T) {
    auto last_rows = [T, rows](Buf& b) {
      if (b.width > 0 && b.v.size() == static_cast<size_t>(T) * b.width)
        b.v.erase(b.v.begin(), b.v.end() - static_cast<size_t>(rows) * b.width);
    };
    for (auto& kv : env) {
      last_rows(kv.second);
      for (Buf& item : kv.second.items) last_rows(item);  // a list's tensors
    }
  }
  run_nodes(head_, -1, env, rows, start_pos, !all_rows);
  auto it = env.find("logits");
  if (it == env.end()) throw std::runtime_error("IR head produced no 'logits'");
  seq_len_ = start_pos + T;
  return std::move(it->second.v);
}

References j, and out.

◆ has_depth_route()

bool ember_ir::Interpreter::has_depth_route ( ) const

inline

Definition at line 183 of file graph_ir.h.

183{ return !mor_step_.empty(); }

References mor_step_.

◆ has_recurrent()

bool ember_ir::Interpreter::has_recurrent ( ) const

inline

Definition at line 127 of file graph_ir.h.

127{ return rec_floats_ > 0; }

References rec_floats_.

◆ held_positions()

std::vector< int > ember_ir::Interpreter::held_positions ( int application ) const

Definition at line 3077 of file graph_ir.cpp.

                                                                {
  if (application < 0 || application >= n_layers_)
    throw std::runtime_error("IR held_positions: application " + std::to_string(application) + " of " +
                             std::to_string(n_layers_));
  if (kv_heads_ <= 0) return {};
  if (turn_state_) return kv_pos_;
  if (!mor_step_.empty() && mor_step_[application] >= 0) return mor_rows_[application];
  std::vector<int> p(static_cast<size_t>(seq_len_));
  for (int i = 0; i < seq_len_; ++i) p[i] = i;
  return p;
}

References j, and p.

◆ init()

void ember_ir::Interpreter::init	(	const nlohmann::json &	j,
		const std::string &	ir_json,
		const GGUFData &	gguf,
		bool	use_gpu
	)

private

Definition at line 1111 of file graph_ir.cpp.

                                                                                                  {
  const int version = j.at("ir_version").get<int>();
  // 1 is spec/IR_SPEC.md's first normative text of the same document as 0 (paperkiln, 25 Oct 2026)
  if (version != 0 && version != 1)
    throw std::runtime_error("IR version " + std::to_string(version) + " is not supported (ember reads 0 and 1)");
  family_hint_ = j.value("family_hint", "");
  embed_ = parse_nodes(j.at("embed"), "embed");
  block_ = parse_nodes(j.at("block"), "block");
  head_ = parse_nodes(j.at("head"), "head");
  for (const auto& n : head_)
    if (!kRowWiseOps.count(n.op)) throw std::runtime_error("IR head op '" + n.op + "' is not row-wise");
  n_layers_ = j.at("block").at("repeat").get<int>();
  if (n_layers_ < 1) throw std::runtime_error("IR_SPEC 3.3: block.repeat " + std::to_string(n_layers_) + " is not a positive integer");
  // Tied-block recursion (paperkiln ir.hpp weight_index): application i reads
  // stored layer m_i's tensors; its state slots, kv cache and per-layer
  // attributes stay its own (indexed by i). Stored layers are 0 .. max m_i, each
  // used, and their count is the file's <arch>.block_count.
  if (j.at("block").contains("weight_index")) {
    const json& wi = j.at("block").at("weight_index");
    if (!is_per_layer(wi) || !wi.at("per_layer").is_array())
      throw std::runtime_error("IR block weight_index must be {\"per_layer\": [m_0, ..]}");
    for (const auto& e : wi.at("per_layer")) {
      if (!e.is_number_integer() || e.get<int>() < 0)
        throw std::runtime_error("IR block weight_index entries must be integers >= 0");
      weight_map_.push_back(e.get<int>());
    }
    if (static_cast<int>(weight_map_.size()) != n_layers_)
      throw std::runtime_error("IR block weight_index has " + std::to_string(weight_map_.size()) + " entries for " +
                               std::to_string(n_layers_) + " block applications");
    const int stored = *std::max_element(weight_map_.begin(), weight_map_.end()) + 1;
    std::vector<bool> used(static_cast<size_t>(stored), false);
    for (int m : weight_map_) used[m] = true;
    for (int m = 0; m < stored; ++m)
      if (!used[m]) throw std::runtime_error("IR block weight_index never uses stored layer " + std::to_string(m));
    const std::string arch = meta_string(gguf, "general.architecture").value_or("");
    if (const auto bc = meta_number(gguf, arch + ".block_count"); bc && static_cast<int>(*bc) != stored)
      throw std::runtime_error("IR block weight_index uses " + std::to_string(stored) + " stored layers, the file's " +
                               arch + ".block_count is " + std::to_string(static_cast<int>(*bc)));
    std::string m;
    for (int v : weight_map_) m += (m.empty() ? "" : ", ") + std::to_string(v);
    Logger::info("[IR] tied-block recursion: " + std::to_string(n_layers_) + " applications of " +
                 std::to_string(stored) + " stored layers, weight_index [" + m + "]");
  }
  vocab_ = j.at("model").at("vocab").get<int>();
  max_len_ = j.at("model").value("ctx", 0);
 
  bool has_kv = false, kv_rows_active = false, evict_declared = false;
  for (const auto& s : j.at("block").value("state", json::array())) {
    const std::string kind = s.at("kind").get<std::string>();
    if (kind == "recurrent") {
      RecSlot slot{s.at("name").get<std::string>(), s.at("op").get<std::string>(), {}};
      if (!kRecurrentOps.count(slot.op))
        throw std::runtime_error("IR recurrent state '" + slot.name + "': op '" + slot.op + "' is not a recurrent op");
      if (s.value("dtype", "f32") != "f32")
        throw std::runtime_error("IR recurrent state '" + slot.name + "': dtype must be f32");
      if (!s.contains("tensors") || s.at("tensors").empty())
        throw std::runtime_error("IR recurrent state '" + slot.name + "' declares no tensors");
      for (const auto& t : s.at("tensors")) {
        RecTensor rt;
        rt.name = t.at("name").get<std::string>();
        rt.shape = t.at("shape").get<std::vector<int64_t>>();
        rt.floats = 1;
        for (int64_t e : rt.shape) {
          if (e <= 0) throw std::runtime_error("IR recurrent state '" + slot.name + "': bad shape for " + rt.name);
          rt.floats *= static_cast<size_t>(e);
        }
        // init 0, or "body" for op recurrence: the node's init assigns it (IR_SPEC 6.10.1; set after binding)
        if (t.contains("init") && t.at("init").is_string()) {
          if (slot.op != "recurrence" || t.at("init") != "body")
            throw std::runtime_error("IR recurrent state '" + slot.name + "': init " + t.at("init").dump() + " for " +
                                     rt.name + " (0, or \"body\" for op recurrence; IR_SPEC 6.3, 10.1 rule 13)");
          rt.init = 0.0f;
        } else {
          rt.init = t.value("init", 0.0f);
        }
        rt.offset = rec_floats_;
        rec_floats_ += rt.floats;
        slot.tensors.push_back(std::move(rt));
      }
      for (const auto& other : rec_slots_)
        if (other.name == slot.name) throw std::runtime_error("IR recurrent state '" + slot.name + "' is declared twice");
      rec_slots_.push_back(std::move(slot));
      continue;
    }
    if (kind != "kv") throw std::runtime_error("IR state kind '" + kind + "' has no CPU implementation in ember");
    has_kv = true;
    kv_slot_ = s.value("name", "kv");
    if (s.contains("kv_from")) throw std::runtime_error("IR kv state: kv_from (recursive KV sharing) is not implemented in ember");
    if (s.contains("kv_rows")) {
      if (s.at("kv_rows") != "active")
        throw std::runtime_error("IR kv state: kv_rows '" + s.at("kv_rows").dump() + "' is not implemented in ember");
      kv_rows_active = true;
    }
    if (has(s, "evict")) {
      if (s.at("evict") != "turn_state")
        throw std::runtime_error("IR kv state: evict " + s.at("evict").dump() + " is not implemented in ember (turn_state only)");
      evict_declared = true;
    }
    for (const char* k : {"heads", "kv_heads", "head_dim", "max_len"})
      if (has(s, k) && is_per_layer(s.at(k)))
        throw std::runtime_error(std::string("IR kv state: a per-layer '") + k + "' is not implemented in ember");
    kv_heads_ = s.at("kv_heads").get<int>();
    head_dim_ = s.at("head_dim").get<int>();
    max_len_ = s.at("max_len").get<int>();
    // Sliding window + sinks (paperkiln swa_attention). The cache keeps every
    // position up to max_len, which holds the first `sinks` and the last
    // `window` rows the rule needs.
    if (has(s, "window") && is_per_layer(s.at("window"))) {
      window_layer_ = s.at("window").at("per_layer").get<std::vector<int>>();
      if (static_cast<int>(window_layer_.size()) != n_layers_)
        throw std::runtime_error("IR kv state: per-layer window has " + std::to_string(window_layer_.size()) +
                                 " entries for " + std::to_string(n_layers_) + " blocks");
      window_ = *std::max_element(window_layer_.begin(), window_layer_.end());
    } else {
      window_ = s.value("window", 0);
    }
    if (has(s, "sinks") && is_per_layer(s.at("sinks")))
      throw std::runtime_error("IR kv state: per-layer sinks are not implemented in ember");
    sinks_ = s.value("sinks", 0);
  }
  for (const auto& n : block_) {
    if (n.op != "attention" && n.op != "mla_attention") continue;
    if (n.op == "mla_attention" && !(n.attrs.contains("causal") && n.attrs.at("causal") == json(true)))
      throw std::runtime_error("IR_SPEC 5.20 (rule 14): mla_attention causal must be true");
    if (!n.attrs.value("causal", true))
      throw std::runtime_error("IR attention with causal=false is not implemented in ember");
    for (int l = 0; l < n_layers_; ++l) {
      int w = 0, s = 0;
      attention_window(n, l, w, s);
      if (w < 0 || s < 0) throw std::runtime_error("IR attention window/sinks must be non-negative");
      // A windowed kv state declares what its cache holds; an attention node
      // reading it with another window or sink count is an exporter error.
      const int kvw = window_layer_.empty() ? window_ : window_layer_[l];
      if (kvw > 0 && (w != kvw || s != sinks_))
        throw std::runtime_error("IR attention window/sinks (" + std::to_string(w) + ", " + std::to_string(s) +
                                 ") in block " + std::to_string(l) + " differ from its kv state's (" +
                                 std::to_string(kvw) + ", " + std::to_string(sinks_) + ")");
      if (w > 0 && (l == 0 || window_layer_.size() > 0))
        Logger::info("[IR] sliding-window attention: block " + std::to_string(l) + " window " + std::to_string(w) +
                     ", sinks " + std::to_string(s));
    }
  }
  for (const auto* sub : {&embed_, &block_, &head_})
    for (const auto& n : *sub)
      if (n.op == "gelu" && n.attrs.value("approx", "none") != "tanh" && n.attrs.value("approx", "none") != "none")
        throw std::runtime_error("IR gelu approx '" + n.attrs.value("approx", "") + "' is not implemented in ember");
  // The gating ops' attributes and arities (paperkiln ir.hpp): checked here so
  // a malformed node fails at load, not mid-generation.
  for (const auto* sub : {&embed_, &block_, &head_})
    for (const auto& n : *sub) {
      auto arity = [&](size_t k, bool at_least) {
        if (at_least ? n.in.size() < k : n.in.size() != k)
          throw std::runtime_error("IR " + n.op + " needs " + (at_least ? "at least " : "") + std::to_string(k) +
                                   " input(s), has " + std::to_string(n.in.size()));
        if (n.out.size() != 1) throw std::runtime_error("IR " + n.op + " needs one output");
      };
      if (n.op == "sub" || n.op == "mul_col" || n.op == "append") arity(2, false);
      if (n.op == "relu" || n.op == "sigmoid" || n.op == "affine" || n.op == "rms_row" || n.op == "scale") arity(1, false);
      if (n.op == "scale" && !has(n.attrs, "weight")) throw std::runtime_error("IR scale needs its weight [d]");
      if (n.op == "list" || n.op == "depth_attend") arity(1, true);
      if (n.op == "affine" && (!has(n.attrs, "scale") || !has(n.attrs, "shift")))
        throw std::runtime_error("IR affine needs float attributes scale and shift");
      if (n.op == "rms_row" && !has(n.attrs, "eps")) throw std::runtime_error("IR rms_row needs its eps attribute");
      for (const auto& kv : n.attrs.items()) {  // per-layer attributes: block nodes, the attributes ember reads per block
        if (!is_per_layer(kv.value())) continue;
        const bool ok = sub == &block_ && (((n.op == "attention" || n.op == "mla_attention") &&
                                            (kv.key() == "window" || kv.key() == "sinks")) ||
                                           (n.op == "attention" && kv.key() == "softcap") ||
                                           (n.op == "rope" && (kv.key() == "theta" || kv.key() == "freq_scale" ||
                                                               kv.key() == "attn_factor")) ||
                                           (n.op == "affine" && (kv.key() == "scale" || kv.key() == "shift")) ||
                                           (n.op == "depth_route" && kv.key() == "step") ||
                                           (n.op == "moe_ffn" && kv.key() == "scoring") ||
                                           kv.key() == "apply");  // revision 1.6 (IR_SPEC 3.8): any stateless node
        if (ok && kv.key() == "apply")  // its entries are true or false (IR_SPEC 3.8, rule 43, revision 1.6)
          for (const auto& e : kv.value().at("per_layer"))
            if (!e.is_boolean())
              throw std::runtime_error("IR_SPEC 3.8 (rule 43): block " + n.op + ": apply entry " + e.dump() +
                                       " is not true or false");
        if (!ok && n.op == "rope" && kv.key() == "freq_factors")  // structural (IR_SPEC 3.8, revision 1.5)
          throw std::runtime_error("IR_SPEC 3.8 (rule 6): rope freq_factors is structural and cannot be per-layer");
        if (!ok)
          throw std::runtime_error("IR " + n.op + ": a per-layer '" + kv.key() + "' is not implemented in ember");
        if (static_cast<int>(kv.value().at("per_layer").size()) != n_layers_)
          throw std::runtime_error("IR " + n.op + ": per-layer '" + kv.key() + "' has " +
                                   std::to_string(kv.value().at("per_layer").size()) + " entries for " +
                                   std::to_string(n_layers_) + " blocks");
      }
      if (n.op == "expr") {  // a runtime-defined op: parsed here, so a bad string fails at load naming its column
        arity(1, true);
        if (!n.attrs.contains("expr") || !n.attrs.at("expr").is_string())
          throw std::runtime_error("IR expr needs its string attribute expr");
        try {
          expr::cached(n.attrs.at("expr").get<std::string>(), n.in.size());
        } catch (const std::invalid_argument& e) {
          throw std::runtime_error(std::string("IR ") + e.what());
        }
      }
    }
#ifdef HAS_CUDA
  if (use_gpu && !device_supported(ir_json))
    throw std::runtime_error("IR uses an op or a shape that has no CUDA kernel (EMBER_IR_DEVICE=cpu runs it)");
#endif
  // Multi-head latent attention (IR_SPEC 5.20): its kv slot holds [c | kr] per position (kv_heads 1, head_dim
  // r + dr), and its rope rotates qr and kr, dr wide per head, not the slot's head_dim.
  bool mla = false;
  for (const auto& n : block_) mla = mla || n.op == "mla_attention";
  for (const auto& n : block_) {
    if (n.op == "mla_attention") {  // IR_SPEC 5.20: the slot holds c and kr, r + dr values per position (rule 12)
      for (const char* k : {"heads", "kv_rank", "qk_nope", "qk_rope", "v_head_dim"})
        if (!n.attrs.contains(k) || !n.attrs.at(k).is_number_integer() || n.attrs.at(k).get<int>() <= 0)
          throw std::runtime_error(std::string("IR_SPEC 5.20 (rule 8): mla_attention needs a positive integer ") + k);
      for (const char* k : {"k_b", "v_b"})
        if (!n.attrs.contains(k) || !n.attrs.at(k).is_string())
          throw std::runtime_error(std::string("IR_SPEC 5.20 (rule 8): mla_attention needs its tensor ") + k);
      const int r = n.attrs.at("kv_rank").get<int>(), dr = n.attrs.at("qk_rope").get<int>();
      if (kv_heads_ != 1 || head_dim_ != r + dr)
        throw std::runtime_error("IR_SPEC 5.20 (rule 12): mla_attention reads a kv slot of kv_heads " +
                                 std::to_string(kv_heads_) + " and head_dim " + std::to_string(head_dim_) +
                                 "; it holds c and kr: kv_heads 1, head_dim kv_rank + qk_rope = " + std::to_string(r + dr));
    }
  }
  for (const auto& n : block_) {
    if (n.op == "rope") {
      const std::string pairing = n.attrs.value("pairing", "interleaved");
      if (pairing != "interleaved" && pairing != "neox")
        throw std::runtime_error("IR rope pairing '" + pairing + "' is not implemented in ember");
      if (n.attrs.value("rope_heads", "all") != "all")
        throw std::runtime_error("IR rope_heads '" + n.attrs.value("rope_heads", "") + "' is not implemented in ember");
      if (kv_heads_ > 0 && has(n.attrs, "kv_heads") && n.attrs.at("kv_heads").get<int>() != kv_heads_)
        throw std::runtime_error("IR rope kv_heads differs from its kv state's");
      // The table is the kv state's head_dim wide, as before, unless the node
      // names its own n_rot (partial rotary) or NEOX pairing: then n_rot
      // (default head_dim) must fit the kv state's heads.
      const int hd = n.attrs.at("head_dim").get<int>();
      int n_rot = head_dim_;
      if (has(n.attrs, "n_rot") || pairing == "neox" || mla) {
        n_rot = has(n.attrs, "n_rot") ? n.attrs.at("n_rot").get<int>() : hd;
        if ((!mla && hd != head_dim_) || n_rot <= 0 || n_rot % 2 != 0 || n_rot > hd)
          throw std::runtime_error("IR rope: n_rot " + std::to_string(n_rot) + " / head_dim " + std::to_string(hd) +
                                   " (kv state head_dim " + std::to_string(head_dim_) + ") is not usable");
      }
      std::vector<float> thetas(static_cast<size_t>(n_layers_));
      std::vector<float> scales(static_cast<size_t>(n_layers_), 1.0f), attn(static_cast<size_t>(n_layers_), 1.0f);
      for (int l = 0; l < n_layers_; ++l) {
        thetas[l] = at_layer(n.attrs.at("theta"), l).get<float>();
        if (has(n.attrs, "freq_scale")) scales[l] = at_layer(n.attrs.at("freq_scale"), l).get<float>();
        if (!(scales[l] > 0.0f)) throw std::runtime_error("IR rope freq_scale must be positive");
        // LongRoPE's / YaRN's attention scaling (IR_SPEC 5.14, revision 1.5): the table holds cos * m and sin * m.
        if (has(n.attrs, "attn_factor") && at_layer(n.attrs.at("attn_factor"), l).is_number())
          attn[l] = at_layer(n.attrs.at("attn_factor"), l).get<float>();
      }
      const json ff = has(n.attrs, "freq_factors") ? n.attrs.at("freq_factors") : json();
      if (rope_dim_ != 0 && (rope_theta_layer_ != thetas || rope_scale_layer_ != scales || rope_attn_layer_ != attn ||
                             rope_dim_ != n_rot || rope_neox_ != (pairing == "neox") || rope_ff_ != ff))
        throw std::runtime_error(
            "IR rope nodes with different theta, freq_scale, attn_factor, freq_factors, n_rot or pairing are not "
            "implemented in ember");
      rope_theta_layer_ = thetas;
      rope_scale_layer_ = scales;
      rope_attn_layer_ = attn;
      rope_ff_ = ff;  // its form is rule 41's (validate_ir below); the tables are built after it
      rope_theta_ = thetas[0];
      rope_freq_scale_ = scales[0];
      rope_dim_ = n_rot;
      rope_neox_ = pairing == "neox";
    }
  }
  // The rmsnorm groups (per-head q/k norm) and the attention head counts.
  for (const auto& n : block_) {
    if (n.op == "rmsnorm" && has(n.attrs, "group") && n.attrs.at("group").get<int>() <= 0)
      throw std::runtime_error("IR rmsnorm group must be positive");
    if (n.op == "attention") {
      const int hq = n.attrs.at("heads").get<int>(), hkv = n.attrs.at("kv_heads").get<int>();
      if (hkv <= 0 || hq % hkv != 0)
        throw std::runtime_error("IR attention: heads " + std::to_string(hq) + " is not a multiple of kv_heads " +
                                 std::to_string(hkv));
    }
  }
  if (mla) {
    for (const auto& n : block_)
      if (n.op == "depth_route")
        throw std::runtime_error("IR mla_attention under depth_route is not implemented in ember");
  }
  // Mixture-of-Recursions (ir.hpp depth_route, llama_mor): one depth_route node, the block's
  // first, outside it nowhere; its steps one consecutive run of equal-length steps 0 .. R-1; causal
  // token choice; a kv state with kv_rows "active"; no recurrent state. A window and sinks count
  // positions over each application's rows (IR_SPEC 5.15 and 7.5, v1.4.2, Q16).
  for (const auto* sub : {&embed_, &head_})
    for (const auto& n : *sub)
      if (n.op == "depth_route") throw std::runtime_error("IR depth_route outside the block");
  for (size_t k = 0; k < block_.size(); ++k) {
    const Node& n = block_[k];
    if (n.op != "depth_route") continue;
    if (k != 0) throw std::runtime_error("IR depth_route must be the block's first node");
    if (n.in.size() != 1 || n.out.size() != 1) throw std::runtime_error("IR depth_route reads [x] and writes one output");
    if (n.attrs.value("route", "") != "token_choice")
      throw std::runtime_error("IR depth_route route '" + n.attrs.value("route", "") +
                               "' is not implemented in ember (causal token choice only)");
    const std::string gate = n.attrs.value("gate", "");
    if (gate != "score" && gate != "none") throw std::runtime_error("IR depth_route gate '" + gate + "' is not score or none");
    mor_R_ = n.attrs.at("n_rec").get<int>();
    if (mor_R_ < 1) throw std::runtime_error("IR depth_route n_rec must be at least 1");
    if (!is_per_layer(n.attrs.at("step"))) throw std::runtime_error("IR depth_route step must be {\"per_layer\": [...]}");
    mor_step_ = n.attrs.at("step").at("per_layer").get<std::vector<int>>();
    if (static_cast<int>(mor_step_.size()) != n_layers_)
      throw std::runtime_error("IR depth_route step has " + std::to_string(mor_step_.size()) + " entries for " +
                               std::to_string(n_layers_) + " block applications");
    // One run of applications with steps 0, 0, .., 1, 1, .., R-1, all steps the same length.
    int first = -1, last = -1;
    for (int i = 0; i < n_layers_; ++i)
      if (mor_step_[i] >= 0) {
        if (first < 0) first = i;
        last = i;
      }
    const int span = first < 0 ? 0 : last - first + 1;
    if (first < 0 || span % mor_R_ != 0) throw std::runtime_error("IR depth_route step: the recursion is not R equal steps");
    for (int i = 0; i < n_layers_; ++i) {
      const int want = i < first || i > last ? -1 : (i - first) / (span / mor_R_);
      if (mor_step_[i] != want)
        throw std::runtime_error("IR depth_route step: application " + std::to_string(i) + " has step " +
                                 std::to_string(mor_step_[i]) + ", expected " + std::to_string(want) +
                                 " (consecutive, ascending, equal-length steps)");
    }
    mor_gate_score_ = gate == "score";
    mor_x_ = n.in[0];
    mor_out_ = n.out[0];
    mor_router_ = n.attrs.at("weight").get<std::string>();
    if (mor_router_.find("{i}") != std::string::npos)
      throw std::runtime_error("IR depth_route weight names its tensor in full (no {i})");
  }
  if (!mor_step_.empty()) {
    if (!kv_rows_active) throw std::runtime_error("IR depth_route needs a kv state with kv_rows \"active\"");
    if (!rec_slots_.empty()) throw std::runtime_error("IR depth_route with recurrent state is not implemented in ember");
    for (const auto& n : block_)
      if (n.op == "attention" && has(n.attrs, "position_bias"))
        throw std::runtime_error("IR depth_route with attention position_bias is not implemented in ember");
      else if (n.op == "attention" && has(n.attrs, "relative_kv"))
        throw std::runtime_error("IR depth_route with attention relative_kv is not implemented in ember");
    mor_rows_.assign(n_layers_, {});
  }
  // Per-turn state tokens (paperkiln IR_SPEC_DRAFT.md 1.1, 1.3, 1.4): the extended embedding
  // (state_weight, one embedding in the embed subgraph, its name in full), attention "mask"
  // "turn_state" (window 0 and sinks 0 only), kv "evict" "turn_state" (every attention node
  // turn_state). Not with depth_route (MoR's compact caches) or, for eviction, recurrent state.
  for (const auto* sub : {&embed_, &block_, &head_})
    for (const auto& n : *sub) {
      if (n.op == "embedding" && has(n.attrs, "state_weight")) {
        if (sub != &embed_) throw std::runtime_error("IR embedding state_weight outside the embed subgraph");
        if (!state_weight_.empty()) throw std::runtime_error("IR: more than one embedding with a state_weight");
        state_weight_ = n.attrs.at("state_weight").get<std::string>();
        if (state_weight_.find("{i}") != std::string::npos)
          throw std::runtime_error("IR embedding state_weight names its tensor in full (no {i})");
      }
      if (n.op == "attention" && n.attrs.contains("mask")) {
        const std::string mask = n.attrs.at("mask").is_string() ? n.attrs.at("mask").get<std::string>() : "";
        if (mask != "causal" && mask != "turn_state")
          throw std::runtime_error("IR attention mask " + n.attrs.at("mask").dump() +
                                   " is not implemented in ember (causal or turn_state)");
        if (mask == "turn_state") turn_state_ = true;
      }
    }
  if (turn_state_) {
    for (const auto& n : block_) {
      if (n.op != "attention" || n.attrs.value("mask", "causal") != "turn_state") continue;
      for (int l = 0; l < n_layers_; ++l) {
        int w = 0, s = 0;
        attention_window(n, l, w, s);
        if (w != 0 || s != 0)
          throw std::runtime_error("IR attention mask turn_state needs window 0 and sinks 0 (block " + std::to_string(l) +
                                   " has window " + std::to_string(w) + ", sinks " + std::to_string(s) + ")");
      }
    }
    if (window_ > 0 || !window_layer_.empty() || sinks_ > 0)
      throw std::runtime_error("IR attention mask turn_state with a windowed kv state is not allowed");
    if (!mor_step_.empty() || kv_rows_active)
      throw std::runtime_error("IR attention mask turn_state with depth_route / kv_rows active is not implemented in ember");
  }
  {
    bool all_turn_state = true;
    for (const auto& n : block_)
      if (n.op == "attention" && n.attrs.value("mask", "causal") != "turn_state") all_turn_state = false;
    if (evict_declared && (!turn_state_ || !all_turn_state))
      throw std::runtime_error("IR kv state evict turn_state needs every attention node to have mask turn_state");
    evict_ = evict_declared;
    evict_declared_ = evict_declared;
    if (const char* e = std::getenv("EMBER_EVICT")) {
      if (std::string(e) == "0") evict_ = false;
      if (std::string(e) == "1" && turn_state_) {
        if (!all_turn_state)
          throw std::runtime_error("EMBER_EVICT=1: eviction needs every attention node to have mask turn_state");
        evict_ = true;
      }
    }
    if (evict_ && !rec_slots_.empty())
      throw std::runtime_error("IR kv state evict turn_state with recurrent state is not implemented in ember");
  }
  // EMBER_IR_CTX=N: hold at most N positions (a third-party IR declares the
  // model's whole trained context, e.g. 32768, which is ~0.8 GB of f32 KV for
  // Qwen2.5-0.5B).
  if (const char* c = std::getenv("EMBER_IR_CTX")) {
    const int cap = std::atoi(c);
    if (cap > 0 && max_len_ > cap) {
      Logger::info("[IR] EMBER_IR_CTX: the kv state holds " + std::to_string(cap) + " of the IR's " +
                   std::to_string(max_len_) + " positions");
      max_len_ = cap;
    }
  }
  // EMBER_KV=f16|q8|q4 (kv_quant.h, R2): the kv cache's storage format; unset or f32 changes nothing.
  kv_fmt_ = ember_kvq::parse_format(std::getenv("EMBER_KV"));
  if (kv_fmt_ != ember_kvq::kF32 && head_dim_ > 0) {
    const std::string why = ember_kvq::shape_error(kv_fmt_, head_dim_);
    if (!why.empty()) throw std::runtime_error(std::string("EMBER_KV=") + ember_kvq::format_name(kv_fmt_) + ": " + why);
    for (const auto& n : block_)  // its latent cache [c | kr] has its own write path, not quantised
      if (n.op == "mla_attention")
        throw std::runtime_error(std::string("EMBER_KV=") + ember_kvq::format_name(kv_fmt_) +
                                 ": mla_attention's latent cache is not quantised");
    Logger::info(std::string("[IR] EMBER_KV: the kv cache stores ") + ember_kvq::format_name(kv_fmt_));
  }
  bool has_attention = false;
  for (const auto& n : block_) has_attention = has_attention || n.op == "attention" || n.op == "mla_attention";
  if ((has_kv || has_attention) && (kv_heads_ <= 0 || head_dim_ <= 0 || max_len_ <= 0))
    throw std::runtime_error("IR block declares no usable kv state (kv_heads, head_dim, max_len)");
  if (max_len_ <= 0) throw std::runtime_error("IR model.ctx is required when no kv state bounds the sequence");
  for (const auto* sub : {&embed_, &head_})
    for (const auto& n : *sub)
      if (kRecurrentOps.count(n.op)) throw std::runtime_error("IR op '" + n.op + "' outside the block has no state");
  // Every recurrent node names a slot of its op whose tensors fit the node.
  for (const auto& n : block_) {
    if (n.op == "linear_attention") {
      const RecSlot& slot = rec_slot(n, "linear_attention");
      const int heads = n.attrs.at("heads").get<int>(), hd = n.attrs.at("head_dim").get<int>();
      if (n.attrs.value("feature_map", "elu_plus_one") != "elu_plus_one")
        throw std::runtime_error("IR linear_attention feature_map '" + n.attrs.value("feature_map", "") +
                                 "' is not implemented in ember");
      if (!n.attrs.value("causal", true))
        throw std::runtime_error("IR linear_attention with causal=false is not implemented in ember");
      if (n.in.size() != 3 || n.out.size() != 1) throw std::runtime_error("IR linear_attention reads [q, k, v] and writes one output");
      for (const char* t : {"num", "den"}) {
        bool found = false;
        for (const auto& rt : slot.tensors)
          if (rt.name == t) {
            found = true;
            if (rt.shape != std::vector<int64_t>{heads, hd})
              throw std::runtime_error("IR recurrent state '" + slot.name + "': " + t + " is not [heads, head_dim]");
          }
        if (!found) throw std::runtime_error("IR recurrent state '" + slot.name + "' has no tensor '" + t + "'");
      }
    } else if (n.op == "ssm_scan") {
      const RecSlot& slot = rec_slot(n, "ssm_scan");
      const int64_t sn = n.attrs.at("n").get<int64_t>();
      if (n.in.size() != 1 || n.out.size() != 1) throw std::runtime_error("IR ssm_scan reads [u] and writes one output");
      if (slot.tensors.size() != 1 || slot.tensors[0].name != "s" || slot.tensors[0].shape != std::vector<int64_t>{sn})
        throw std::runtime_error("IR recurrent state '" + slot.name + "' must hold one tensor s [n]");
    } else if (n.op == "conv1d_causal") {
      const RecSlot& slot = rec_slot(n, "conv1d_causal");
      const int64_t C = n.attrs.at("channels").get<int64_t>(), K = n.attrs.at("kernel").get<int64_t>();
      if (n.in.size() != 1 || n.out.size() != 1) throw std::runtime_error("IR conv1d_causal reads [x] and writes one output");
      if (K < 2 || C < 1) throw std::runtime_error("IR conv1d_causal needs channels >= 1 and kernel >= 2");
      if (slot.tensors.size() != 1 || slot.tensors[0].name != "x" ||
          slot.tensors[0].shape != std::vector<int64_t>{K - 1, C})
        throw std::runtime_error("IR recurrent state '" + slot.name + "' must hold one tensor x [kernel - 1, channels]");
    } else if (n.op == "selective_scan") {
      const RecSlot& slot = rec_slot(n, "selective_scan");
      const int64_t D = n.attrs.at("channels").get<int64_t>(), N = n.attrs.at("d_state").get<int64_t>();
      if (n.in.size() != 4 || n.out.size() != 1)
        throw std::runtime_error("IR selective_scan reads [x, dt, B, C] and writes one output");
      if (slot.tensors.size() != 1 || slot.tensors[0].name != "s" || slot.tensors[0].shape != std::vector<int64_t>{D, N})
        throw std::runtime_error("IR recurrent state '" + slot.name + "' must hold one tensor s [channels, d_state]");
    } else if (n.op == "recurrence") {  // IR_SPEC 6.10, 10.1 rules 31-33 (graph_ir_recurrence.h)
      auto prog = std::make_shared<rec::Recurrence>(rec::Recurrence::compile(n.attrs));
      if (n.out.size() != 1) throw std::runtime_error("IR recurrence writes one output");
      const std::string name = has(n.attrs, "state") ? n.attrs.at("state").get<std::string>() : "";
      const json* decl = nullptr;  // into the document (not a copy: it is read below)
      if (j.at("block").contains("state") && j.at("block").at("state").is_array())
        for (const auto& s : j.at("block").at("state"))
          if (s.is_object() && s.value("name", "") == name && s.value("kind", "") == "recurrent") decl = &s;
      if (!decl) throw std::runtime_error("IR_SPEC 6.10 (rule 33): recurrence names no recurrent slot '" + name + "'");
      const std::string why = prog->slot_errors(*decl);
      if (!why.empty()) throw std::runtime_error("IR_SPEC 6.10 (rule 33): " + why);
      rec_slot(n, "recurrence");
      for (const auto& other : block_)
        if (&other != &n && has(other.attrs, "state") && other.attrs.at("state") == name)
          throw std::runtime_error("IR_SPEC 6.10 (rule 33): the slot '" + name + "' of a recurrence is read by another node (" +
                                   other.op + ")");
      recs_[&n] = std::move(prog);
    }
  }
  if (rec_floats_ > 0) {
    // Every application's initial state: +0, or (op recurrence with "init" "body") the init body over the
    // application's weights, set after binding.
    rec_init_.assign(static_cast<size_t>(n_layers_) * rec_floats_, 0.0f);
    for (int l = 0; l < n_layers_; ++l)
      for (const auto& slot : rec_slots_)
        for (const auto& rt : slot.tensors)
          std::fill(rec_init_.begin() + l * rec_floats_ + rt.offset, rec_init_.begin() + l * rec_floats_ + rt.offset + rt.floats,
                    rt.init);
    rec_ = rec_init_;
    // Snapshots are a whole state each; at most 64, and at most 256 MiB.
    const size_t budget = (256u << 20) / sizeof(float);
    rec_max_snaps_ = std::max<size_t>(2, std::min<size_t>(64, budget / rec_.size()));
  }
 
  // Quantised weights stay in their blocks for large models on the CPU (the
  // memory saving pays for the per-row dequantisation); small models, like
  // the llama path's dequantize_all_if_small, and the CUDA path (f32 kernels)
  // dequantise at load. EMBER_IR_QUANT=1 / EMBER_IR_DEQUANT=1 force either.
  {
    size_t params = 0;
    for (const auto& kv : gguf.tensor_infos_map) params += kv.second.num_elements;
#ifdef HAS_CUDA
    const bool device = use_gpu;
#else
    const bool device = false;
#endif
    dequant_all_ = device || params <= 64u * 1000 * 1000;
    const char* q = std::getenv("EMBER_IR_QUANT");
    if (q && std::string(q) == "1" && !device) dequant_all_ = false;
    if (const char* dq = std::getenv("EMBER_IR_DEQUANT")) dequant_all_ = dequant_all_ || std::string(dq) == "1";
  }
  // depth_attend's source count S (IR_SPEC 5.18) is structural: the entries list and append put in its list
  // before the node, in run order (embed once, the block per application, the head), plus its other inputs.
  {
    std::unordered_map<std::string, int> entries;  // list values: their tensor count
    auto walk = [&](const std::vector<Node>& nodes, int layer) {
      for (const auto& n : nodes) {
        const auto held = [&](const std::string& v) {
          const auto it = entries.find(v);
          return it == entries.end() ? 0 : it->second;
        };
        if (n.op == "list") {
          entries[n.out.at(0)] = static_cast<int>(n.in.size());
        } else if (n.op == "append") {
          entries[n.out.at(0)] = held(n.in.at(0)) + 1;
        } else {
          if (n.op == "depth_attend") da_sources_[{&n, layer}] = held(n.in.at(0)) + static_cast<int>(n.in.size()) - 1;
          for (const auto& o : n.out) entries.erase(o);
        }
      }
    };
    walk(embed_, -1);
    for (int l = 0; l < n_layers_; ++l) walk(block_, l);
    walk(head_, -1);
  }
  bind(embed_, -1, gguf);
  for (int l = 0; l < n_layers_; ++l) bind(block_, l, gguf);
  bind(head_, -1, gguf);
  // A recurrence's initial state (IR_SPEC 6.10.3): its init bodies over each application's weights.
  for (const auto& [np, prog] : recs_) {
    const RecSlot& slot = rec_slot(*np, "recurrence");
    AppIndex ai(app_index_, *np);
    for (int l = 0; l < n_layers_; ++l) {
      std::vector<const float*> w;
      for (const auto& wd : prog->weights()) w.push_back(weight(layer_name(wd.tensor, l)).data());
      std::vector<float*> st;
      for (const auto& s : prog->state()) {
        const RecTensor* rt = nullptr;
        for (const auto& t : slot.tensors)
          if (t.name == s.name) rt = &t;
        st.push_back(rec_init_.data() + static_cast<size_t>(l) * rec_floats_ + rt->offset);
      }
      prog->init(w, st);
    }
  }
  if (!recs_.empty()) rec_ = rec_init_;
  // The state rows: state_weight [k, d] (ggml ne0 = d, the token embedding's width); k from its
  // size, equal to the file's paperkiln.state_tokens.k when the file names one.
  {
    const auto meta_k = meta_number(gguf, "paperkiln.state_tokens.k");
    if (!state_weight_.empty()) {
      bind_tensor(state_weight_, gguf, true);
      int d = 0;
      for (const auto& n : embed_)
        if (n.op == "embedding" && has(n.attrs, "state_weight")) {
          std::string w = layer_name(n.attrs.at("weight").get<std::string>(), -1);
          d = static_cast<int>(shapes_.at(w).at(0));
        }
      const auto& sh = shapes_.at(state_weight_);
      const size_t numel = bound(state_weight_).numel;
      if (sh.empty() || sh[0] != d || d <= 0 || numel % static_cast<size_t>(d) != 0 || numel == 0)
        throw std::runtime_error("IR embedding state_weight '" + state_weight_ + "' is not [k, d] with d " +
                                 std::to_string(d) + " (the token embedding's width)");
      state_k_ = static_cast<int>(numel / static_cast<size_t>(d));
      if (meta_k && static_cast<int>(*meta_k) != state_k_)
        throw std::runtime_error("paperkiln.state_tokens.k is " + std::to_string(static_cast<int>(*meta_k)) +
                                 ", the IR's state_weight '" + state_weight_ + "' has " + std::to_string(state_k_) +
                                 " rows");
      Logger::info("[IR] per-turn state tokens: k " + std::to_string(state_k_) + " (ids " + std::to_string(vocab_) +
                   " .. " + std::to_string(vocab_ + state_k_ - 1) + " read '" + state_weight_ + "')" +
                   (turn_state_ ? std::string(", attention mask turn_state") : std::string()) +
                   (evict_ ? std::string(", kv eviction at turn_state") : std::string()));
    } else if (meta_k && *meta_k > 0) {
      throw std::runtime_error("paperkiln.state_tokens.k is " + std::to_string(static_cast<int>(*meta_k)) +
                               " but the IR's embedding has no state_weight");
    }
  }
  // paperkiln spec/IR_SPEC.md section 10.1 (graph_ir_validate.h): after ember's own checks above, so a
  // document they refuse keeps their message; then what ember does not implement among valid
  // documents (section 10.3).
  validate_ir(j, gguf);
  {
    int kv_slots = 0;
    for (const auto& s : j.at("block").value("state", json::array())) kv_slots += s.value("kind", "") == "kv";
    if (kv_slots > 1)
      throw std::runtime_error("IR block declares " + std::to_string(kv_slots) +
                               " kv slots: more than one kv slot per block (IR_SPEC 6.1) is not implemented in ember");
  }
  size_t n_quant = 0, bytes_quant = 0, bytes_f32 = 0, n_mapped = 0;
  for (const auto& kv : weights_) {
    n_quant += kv.second.quantised();
    n_mapped += kv.second.quantised() && kv.second.q.empty();
    bytes_quant += kv.second.qbytes;
    bytes_f32 += kv.second.f32.size() * sizeof(float);
  }
 
  build_rope_tables(gguf);
#ifdef HAS_CUDA
  if (use_gpu) {
    init_device();
    gpu_ = true;
  }
#else
  (void)use_gpu;
#endif
  if (!gpu_ && kv_heads_ > 0) {
    kv_.resize(n_layers_);
    for (auto& layer : kv_) {
      layer.k.assign(static_cast<size_t>(max_len_) * kv_heads_ * head_dim_, 0.0f);
      layer.v.assign(layer.k.size(), 0.0f);
    }
  }
  Logger::info("[IR] interpreter ready: family_hint=" + family_hint_ + ", " + std::to_string(n_layers_) + " blocks of " +
               std::to_string(block_.size()) + " ops, " + std::to_string(weights_.size()) + " weights bound (" +
               std::to_string(n_quant) + " kept quantised, " + std::to_string(bytes_quant >> 20) + " MiB" +
               (n_mapped ? " (" + std::to_string(n_mapped) + " read in place from the memory map)" : std::string()) +
               "; f32 " +
               std::to_string(bytes_f32 >> 20) + " MiB" + (dequant_all_ ? "; dequantised at load" : "") + "), device " +
               (gpu_ ? "cuda" : "cpu"));
  {  // The document run, by its digest (IR_SPEC 3.10): a citation names it as sha256:<hex>.
    const DigestResult d = ir_digest(ir_json);
    Logger::info(d.ok ? "[IR] document sha256:" + d.sha256 : "[IR] the document has no digest (" + d.refused + ")");
  }
  if (rec_floats_ > 0)
    Logger::info("[IR] recurrent state: " + std::to_string(rec_slots_.size()) + " slot(s), " +
                 std::to_string(rec_floats_) + " floats per block; up to " + std::to_string(rec_max_snaps_) +
                 " snapshots for prefix reuse and speculative rewinds");
}

References arity, attention_window(), block_, c, ember_ir::expr::cached(), ember_ir::rec::Recurrence::compile(), ember_ir::device_supported(), embed_, evict_, evict_declared(), evict_declared_, family_hint_, ember_kvq::format_name(), head_, head_dim_, Logger::info(), j, k, ember_kvq::kF32, kind, kv_fmt_, kv_heads_, kv_slot_, max_len_, mor_gate_score_, mor_out_, mor_R_, mor_router_, mor_rows_, mor_step_, mor_x_, n, n_layers_, ember_ir::Interpreter::RecTensor::name, ember_ir::Interpreter::RecSlot::name, name, ember_kvq::parse_format(), rec_floats_, rec_slot(), rec_slots_, recs_, rope_attn_layer_, rope_dim_, rope_ff_, rope_freq_scale_, rope_neox_, rope_scale_layer_, rope_theta_, rope_theta_layer_, ember_kvq::shape_error(), sinks_, state_weight_, sub, ember_ir::Interpreter::RecSlot::tensors, turn_state_, v, vocab_, weight_map_, why, window_, and window_layer_.

Referenced by Interpreter().

◆ kv_cache_bytes()

size_t ember_ir::Interpreter::kv_cache_bytes ( ) const

Definition at line 1767 of file graph_ir.cpp.

                                         {
#ifdef HAS_CUDA
  if (gpu_)
    return kv_fmt_ != ember_kvq::kF32 ? kv_dev_bytes_
                                      : static_cast<size_t>(n_layers_) * 2 * max_len_ * kv_heads_ * head_dim_ * sizeof(float);
#endif
  return 0;
}

References j, and ember_kvq::kF32.

◆ kv_format()

const char * ember_ir::Interpreter::kv_format ( ) const

Definition at line 1765 of file graph_ir.cpp.

1765{ return ember_kvq::format_name(kv_fmt_); }

References ember_kvq::format_name().

◆ kv_positions()

std::vector< int > ember_ir::Interpreter::kv_positions ( int application ) const

Definition at line 3064 of file graph_ir.cpp.

                                                              {
  cpu_only(gpu_, "kv_positions");
  if (application < 0 || application >= n_layers_)
    throw std::runtime_error("IR kv_positions: application " + std::to_string(application) + " of " +
                             std::to_string(n_layers_));
  if (kv_.empty()) return {};
  if (turn_state_) return kv_pos_;
  if (!mor_step_.empty() && mor_step_[application] >= 0) return mor_rows_[application];
  std::vector<int> p(static_cast<size_t>(seq_len_));
  for (int i = 0; i < seq_len_; ++i) p[i] = i;
  return p;
}

References ember_ir::cpu_only(), j, and p.

◆ kv_rows_held()

int ember_ir::Interpreter::kv_rows_held ( ) const

inline

Definition at line 171 of file graph_ir.h.

171{ return turn_state_ ? static_cast<int>(kv_pos_.size()) : static_cast<int>(ts_seg_.size()); }

References kv_pos_, ts_seg_, and turn_state_.

◆ kv_rows_in_use()

int ember_ir::Interpreter::kv_rows_in_use ( int application ) const

private

Definition at line 3057 of file graph_ir.cpp.

                                                     {
  if (kv_.empty()) return 0;
  if (turn_state_) return static_cast<int>(kv_pos_.size());
  if (!mor_step_.empty() && mor_step_.at(application) >= 0) return static_cast<int>(mor_rows_.at(application).size());
  return seq_len_;
}

References j.

◆ kv_scratch_bytes()

size_t ember_ir::Interpreter::kv_scratch_bytes ( ) const

Definition at line 1776 of file graph_ir.cpp.

                                           {
#ifdef HAS_CUDA
  if (gpu_) return kv_scratch_dev_bytes_;
#endif
  return 0;
}

References j.

◆ kv_slot()

const std::string & ember_ir::Interpreter::kv_slot ( ) const

inline

Definition at line 180 of file graph_ir.h.

180{ return kv_slot_; }

References kv_slot_.

◆ last_depths()

const std::vector< int > & ember_ir::Interpreter::last_depths ( ) const

inline

Definition at line 141 of file graph_ir.h.

141{ return mor_depth_; }

References mor_depth_.

◆ layer_name()

std::string ember_ir::Interpreter::layer_name	(	const std::string &	pattern,
		int	layer
	)		const

private

Definition at line 1871 of file graph_ir.cpp.

                                                                           {
  std::string s = pattern;
  const auto p = s.find("{i}");
  if (p != std::string::npos) {
    if (layer < 0) throw std::runtime_error("IR weight '" + pattern + "' uses {i} outside the block");
    // m_i under weight_index, except in a node tagged "tensor_index": "application" (IR_SPEC 7.2): i.
    s.replace(p, 3, std::to_string(weight_map_.empty() || app_index_ ? layer : weight_map_.at(layer)));
  }
  return s;
}

References j, and p.

◆ load_state()

void ember_ir::Interpreter::load_state ( const State & s )

Definition at line 3180 of file graph_ir.cpp.

                                           {
  cpu_only(gpu_, "load_state");
  if (s.k.size() != kv_.size() || s.v.size() != kv_.size() || s.rec.size() != rec_.size() ||
      s.mor_rows.size() != mor_rows_.size())
    throw std::runtime_error("IR load_state: the state is not this interpreter's");
  for (size_t l = 0; l < kv_.size(); ++l) {
    if (s.k[l].size() > kv_[l].k.size() || s.v[l].size() != s.k[l].size())
      throw std::runtime_error("IR load_state: the kv rows do not fit the cache");
    std::copy(s.k[l].begin(), s.k[l].end(), kv_[l].k.begin());
    std::copy(s.v[l].begin(), s.v[l].end(), kv_[l].v.begin());
  }
  seq_len_ = s.seq_len;
  ts_seg_ = s.ts_seg;
  ts_sidx_ = s.ts_sidx;
  kv_pos_ = s.kv_pos;
  evicted_below_ = s.evicted_below;
  mor_rows_ = s.mor_rows;
  mor_depth_ = s.mor_depth;
  rec_ = s.rec;
  rec_pos_ = s.rec_pos;
  rec_snaps_ = s.rec_snaps;
  ids_ = s.ids;
  rope_long_ = s.rope_long;
}

References ember_ir::cpu_only(), and j.

◆ max_len()

int ember_ir::Interpreter::max_len ( ) const

inline

Definition at line 136 of file graph_ir.h.

136{ return max_len_; }

References max_len_.

◆ mor_step()

void ember_ir::Interpreter::mor_step	(	int	layer,
		Env &	env,
		int	T,
		int	start_pos
	)

private

Definition at line 2909 of file graph_ir.cpp.

                                                                    {
  const int s = mor_step_[layer];
  const bool first = layer == 0 || mor_step_[layer - 1] != s;
  const bool last = layer == n_layers_ - 1 || mor_step_[layer + 1] != s;
  auto it = env.find(mor_x_);
  if (it == env.end() || it->second.is_list) throw std::runtime_error("IR depth_route: input '" + mor_x_ + "' is not a tensor");
  Buf& x = it->second;
  const int d = x.width;
  if (s == 0 && first) {
    const auto& W = weight(mor_router_);
    if (shapes_.at(mor_router_).at(0) != d || W.size() != static_cast<size_t>(mor_R_) * d)
      throw std::runtime_error("IR depth_route '" + mor_router_ + "' is not [n_rec, d]");
    std::vector<float> logits(static_cast<size_t>(T) * mor_R_);
    matmul_f32_f32_batch_cpu(W, x.v, logits, T, mor_R_, d);  // x W^T, as linear
    mor_depth_.assign(T, 0);
    mor_g_.assign(T, 0.0f);
    for (int t = 0; t < T; ++t) cpu::mor_route(logits.data() + static_cast<size_t>(t) * mor_R_, mor_R_, mor_depth_[t], mor_g_[t]);
    env[mor_out_] = Buf{std::vector<float>(mor_depth_.begin(), mor_depth_.end()), 1};
  }
  if (first) mor_xin_ = x.v;
  std::vector<int> active;
  for (int t = 0; t < T; ++t)
    if (mor_depth_.at(t) > s) active.push_back(t);
  for (int t : active) {
    Env row;
    for (const auto& kv : env)
      if (!kv.second.is_list && kv.second.width > 0 && kv.second.v.size() == static_cast<size_t>(T) * kv.second.width) {
        const auto b = kv.second.v.begin() + static_cast<size_t>(t) * kv.second.width;
        row[kv.first] = Buf{std::vector<float>(b, b + kv.second.width), kv.second.width};
      }
    attn_row_ = static_cast<int>(mor_rows_[layer].size());
    try {
      run_nodes(block_, layer, row, 1, start_pos + t, false);
    } catch (...) {
      attn_row_ = -1;
      throw;
    }
    attn_row_ = -1;
    mor_rows_[layer].push_back(start_pos + t);
    const Buf& y = row.at(mor_x_);
    if (y.width != d) throw std::runtime_error("IR depth_route: the block changes the width of '" + mor_x_ + "'");
    std::copy(y.v.begin(), y.v.end(), x.v.begin() + static_cast<size_t>(t) * d);
  }
  if (last && mor_gate_score_)
    for (int t : active)
      cpu::mor_gate(mor_xin_.data() + static_cast<size_t>(t) * d, x.v.data() + static_cast<size_t>(t) * d, d, mor_g_[t]);
}

References j, matmul_f32_f32_batch_cpu(), ember_ir::Interpreter::Buf::v, and ember_ir::Interpreter::Buf::width.

◆ n_layers()

int ember_ir::Interpreter::n_layers ( ) const

inline

Definition at line 137 of file graph_ir.h.

137{ return n_layers_; }

References n_layers_.

◆ on_gpu()

bool ember_ir::Interpreter::on_gpu ( ) const

inline

Definition at line 147 of file graph_ir.h.

147{ return gpu_; }

References gpu_.

◆ operator=()

Interpreter & ember_ir::Interpreter::operator= ( const Interpreter & )

delete

◆ rec_begin()

void ember_ir::Interpreter::rec_begin	(	int	start_pos,
		int	T,
		bool	all_rows
	)

private

Definition at line 2242 of file graph_ir.cpp.

                                                               {
#ifdef HAS_CUDA
  const size_t bytes = rec_.size() * sizeof(float);
#endif
  if (start_pos == 0) {  // a new sequence
#ifdef HAS_CUDA
    if (gpu_) {
      gpuErrchk(cudaMemcpyAsync(drec_, drec_init_, bytes, cudaMemcpyDeviceToDevice, stream_));
      while (!dsnaps_.empty()) dsnap_drop(dsnaps_.begin());
    } else
#endif
    {
      rec_ = rec_init_;  // every application's initial state
      rec_snaps_.clear();
    }
  } else if (start_pos != rec_pos_) {
    auto missing = [&] {
      return std::runtime_error("IR recurrent state holds " + std::to_string(rec_pos_) + " tokens and has no snapshot at " +
                                std::to_string(start_pos) + " (resume_point says where a prefix can resume)");
    };
#ifdef HAS_CUDA
    if (gpu_) {
      auto it = dsnaps_.find(start_pos);
      if (it == dsnaps_.end()) throw missing();
      gpuErrchk(cudaMemcpyAsync(drec_, it->second, bytes, cudaMemcpyDeviceToDevice, stream_));
    } else
#endif
    {
      auto it = rec_snaps_.find(start_pos);
      if (it == rec_snaps_.end()) throw missing();
      rec_ = it->second;
    }
  }
  // Snapshots past start_pos belong to the continuation being replaced.
#ifdef HAS_CUDA
  if (gpu_)
    while (!dsnaps_.empty() && dsnaps_.rbegin()->first > start_pos) dsnap_drop(std::prev(dsnaps_.end()));
  else
#endif
    rec_snaps_.erase(rec_snaps_.upper_bound(start_pos), rec_snaps_.end());
  rec_pos_ = -1;  // until this call completes
  // Rows whose state is kept: every row of an all_rows call (speculative
  // verification), else the second-to-last row (the same prompt again).
  rec_snap_rows_.assign(T, -1);
  int kept = 0;
  for (int t = 0; t + 1 < T; ++t)
    if (all_rows || t == T - 2) rec_snap_rows_[t] = kept++;
  // Each kept row starts as the current state (a slot no node updates keeps it).
#ifdef HAS_CUDA
  if (gpu_) {
    if (kept > 0) grow(drec_rows_, static_cast<size_t>(kept) * rec_.size());
    for (int k = 0; k < kept; ++k)
      gpuErrchk(cudaMemcpyAsync(drec_rows_.alloc + k * rec_.size(), drec_, bytes, cudaMemcpyDeviceToDevice, stream_));
    return;
  }
#endif
  rec_rows_.resize(static_cast<size_t>(kept) * rec_.size());
  for (int k = 0; k < kept; ++k) std::copy(rec_.begin(), rec_.end(), rec_rows_.begin() + k * rec_.size());
}

References j, and k.

◆ rec_end()

void ember_ir::Interpreter::rec_end	(	int	start_pos,
		int	T
	)

private

Definition at line 2318 of file graph_ir.cpp.

                                              {
#ifdef HAS_CUDA
  if (gpu_) {
    const size_t bytes = rec_.size() * sizeof(float);
    for (int t = 0; t + 1 < T; ++t) {
      const int k = rec_snap_rows_[t];
      if (k < 0) continue;
      gpuErrchk(cudaMemcpyAsync(dsnap_take(start_pos + t + 1), drec_rows_.alloc + k * rec_.size(), bytes,
                                cudaMemcpyDeviceToDevice, stream_));
    }
    gpuErrchk(cudaMemcpyAsync(dsnap_take(start_pos + T), drec_, bytes, cudaMemcpyDeviceToDevice, stream_));
    while (dsnaps_.size() > rec_max_snaps_) dsnap_drop(dsnaps_.begin());  // the oldest positions go first
    rec_pos_ = start_pos + T;
    return;
  }
#endif
  for (int t = 0; t + 1 < T; ++t) {
    const int k = rec_snap_rows_[t];
    if (k < 0) continue;
    rec_snaps_[start_pos + t + 1].assign(rec_rows_.begin() + k * rec_.size(), rec_rows_.begin() + (k + 1) * rec_.size());
  }
  rec_snaps_[start_pos + T] = rec_;
  while (rec_snaps_.size() > rec_max_snaps_) rec_snaps_.erase(rec_snaps_.begin());  // the oldest positions go first
  rec_pos_ = start_pos + T;
  std::vector<float>().swap(rec_rows_);
}

References j, and k.

◆ rec_row_done()

void ember_ir::Interpreter::rec_row_done	(	const RecSlot &	slot,
		int	layer,
		int	t
	)

private

Definition at line 2302 of file graph_ir.cpp.

                                                                    {
  const int k = rec_snap_rows_[t];
  if (k < 0) return;
  // This slot only: another slot of the block may be at another row.
  const size_t off = static_cast<size_t>(layer) * rec_floats_ + slot.tensors.front().offset;
  const size_t floats = slot.tensors.back().offset + slot.tensors.back().floats - slot.tensors.front().offset;
#ifdef HAS_CUDA
  if (gpu_) {
    gpuErrchk(cudaMemcpyAsync(drec_rows_.alloc + k * rec_.size() + off, drec_ + off, floats * sizeof(float),
                              cudaMemcpyDeviceToDevice, stream_));
    return;
  }
#endif
  std::copy(rec_.begin() + off, rec_.begin() + off + floats, rec_rows_.begin() + k * rec_.size() + off);
}

References j, k, off, and ember_ir::Interpreter::RecSlot::tensors.

◆ rec_slot()

const Interpreter::RecSlot & ember_ir::Interpreter::rec_slot	(	const Node &	n,
		const char *	op
	)		const

private

Definition at line 2066 of file graph_ir.cpp.

                                                                                   {
  if (!has(n.attrs, "state")) throw std::runtime_error(std::string("IR ") + op + " names no state slot");
  const std::string name = n.attrs.at("state").get<std::string>();
  for (const auto& slot : rec_slots_)
    if (slot.name == name) {
      if (slot.op != op)
        throw std::runtime_error("IR recurrent state '" + name + "' belongs to op '" + slot.op + "', not " + op);
      return slot;
    }
  throw std::runtime_error(std::string("IR ") + op + ": no recurrent state named '" + name + "'");
}

References j, n, and name.

Referenced by init().

◆ rec_tensor()

float * ember_ir::Interpreter::rec_tensor	(	const RecSlot &	slot,
		const char *	tensor,
		int	layer,
		size_t	floats
	)

private

Definition at line 2078 of file graph_ir.cpp.

                                                                                                {
  for (const auto& rt : slot.tensors)
    if (rt.name == tensor) {
      if (rt.floats != floats)
        throw std::runtime_error("IR recurrent state '" + slot.name + "': " + tensor + " does not fit the op");
      return rec_.data() + static_cast<size_t>(layer) * rec_floats_ + rt.offset;
    }
  throw std::runtime_error("IR recurrent state '" + slot.name + "' has no tensor '" + tensor + "'");
}

References j, ember_ir::Interpreter::RecSlot::name, tensor, and ember_ir::Interpreter::RecSlot::tensors.

◆ recurrent_tensors()

std::vector< Interpreter::RecurrentTensorInfo > ember_ir::Interpreter::recurrent_tensors ( ) const

Definition at line 3089 of file graph_ir.cpp.

                                                                               {
  std::vector<RecurrentTensorInfo> out;
  for (const auto& slot : rec_slots_)
    for (const auto& rt : slot.tensors) out.push_back({slot.name, rt.name, rt.shape});
  return out;
}

References j, and out.

◆ recurrent_values()

std::vector< float > ember_ir::Interpreter::recurrent_values	(	const std::string &	slot,
		const std::string &	tensor,
		int	application
	)		const

Definition at line 3096 of file graph_ir.cpp.

                                                                        {
  cpu_only(gpu_, "recurrent_values");
  if (application < 0 || application >= n_layers_)
    throw std::runtime_error("IR recurrent_values: application " + std::to_string(application) + " of " +
                             std::to_string(n_layers_));
  for (const auto& s : rec_slots_)
    if (s.name == slot)
      for (const auto& rt : s.tensors)
        if (rt.name == tensor) {
          const auto b = rec_.begin() + static_cast<size_t>(application) * rec_floats_ + rt.offset;
          return std::vector<float>(b, b + rt.floats);
        }
  throw std::runtime_error("IR recurrent_values: no tensor '" + tensor + "' in a slot '" + slot + "'");
}

References ember_ir::cpu_only(), j, and tensor.

◆ resume_point()

int ember_ir::Interpreter::resume_point ( int want ) const

Definition at line 2222 of file graph_ir.cpp.

                                            {
  if (turn_state_ && evict_ && want > 0 && want <= static_cast<int>(ts_seg_.size())) {
    const int m = ts_seg_after(want - 1);
    if (m < evicted_below_) return ts_segment_start(m);
  }
  if (rec_floats_ == 0 || want <= 0) return std::max(want, 0);
  if (want == rec_pos_) return want;
  auto below = [want](const auto& snaps) {
    auto it = snaps.upper_bound(want);  // the first snapshot past want
    return it == snaps.begin() ? 0 : std::prev(it)->first;
  };
#ifdef HAS_CUDA
  if (gpu_) return below(dsnaps_);
#endif
  return below(rec_snaps_);
}

References j.

◆ rope_table()

const std::vector< std::pair< float, float > > & ember_ir::Interpreter::rope_table ( int layer ) const

private

Definition at line 1789 of file graph_ir.cpp.

                                                                               {
  const std::vector<int>& idx = rope_long_ && !rope_tab_long_.empty() ? rope_tab_long_ : rope_tab_short_;
  const int l = layer < 0 ? 0 : layer;
  if (l >= static_cast<int>(idx.size()) || idx[l] < 0)
    throw std::runtime_error("IR rope: block " + std::to_string(l) + " has no RoPE table (theta must be positive)");
  return rope_tabs_[idx[l]];
}

References j.

◆ run_nodes()

void ember_ir::Interpreter::run_nodes	(	const std::vector< Node > &	nodes,
		int	layer,
		Env &	env,
		int	T,
		int	start_pos,
		bool	head_last_only
	)

private

Definition at line 2357 of file graph_ir.cpp.

                                    {
  auto value = [&](const Node& n, size_t k, bool want_list) -> Buf& {
    auto it = env.find(n.in.at(k));
    if (it == env.end()) throw std::runtime_error("IR op '" + n.op + "': input '" + n.in.at(k) + "' is not defined");
    if (it->second.is_list != want_list)
      throw std::runtime_error("IR op '" + n.op + "': input '" + n.in.at(k) + "' is " +
                               (want_list ? "a tensor, not a list" : "a list, not a tensor"));
    return it->second;
  };
  auto in = [&](const Node& n, size_t k) -> Buf& { return value(n, k, false); };
  auto list_in = [&](const Node& n, size_t k) -> Buf& { return value(n, k, true); };
  auto wname = [&](const Node& n) {
    std::string w = layer_name(n.attrs.at("weight").get<std::string>(), layer);
    if (!weights_.count(w) && has(n.attrs, "tied_to")) w = layer_name(n.attrs.at("tied_to").get<std::string>(), layer);
    return w;
  };
  for (const auto& n : nodes) {
    if (!applies(n, layer)) continue;  // apply false in this application: reads, names and writes nothing
    AppIndex ti(app_index_, n);
    if (n.op == "embedding") {
      const auto& tokens = in(n, 0);  // token ids carried as floats (exact below 2^24)
      const std::string wn = wname(n);
      const Weight& W = bound(wn);
      const int d = static_cast<int>(shapes_.at(wn).at(0));
      Buf out{std::vector<float>(static_cast<size_t>(T) * d), d};
      // Extended embedding (state_weight): an id in [vocab, vocab + k) reads row id - vocab of the
      // state rows; without one, an id past the table is refused. EXACT.
      const bool ext = has(n.attrs, "state_weight");
      for (int t = 0; t < T; ++t) {
        const int id = static_cast<int>(tokens.v[t]);
        const Weight* src = &W;
        int row = id;
        if (ext && id >= vocab_) {
          if (id >= vocab_ + state_k_)
            throw std::runtime_error("IR embedding: token id " + std::to_string(id) + " is past the state slots (vocab " +
                                     std::to_string(vocab_) + ", k " + std::to_string(state_k_) + ")");
          src = &bound(state_weight_);
          row = id - vocab_;
        } else if (id < 0 || id >= vocab_ || static_cast<size_t>(id + 1) * d > W.numel) {  // IR_SPEC 5.1, 6.8.1
          throw std::runtime_error("IR embedding: token id " + std::to_string(id) + " out of range");
        }
        float* dst = out.v.data() + static_cast<size_t>(t) * d;
        if (src->quantised())
          dequant_span_to_f32(src->type, src->qp, static_cast<size_t>(row) * d, d, dst);
        else
          std::memcpy(dst, src->f32.data() + static_cast<size_t>(row) * d, d * sizeof(float));
      }
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "pos_embedding") {  // x + position_embd[pos], pos absolute. EXACT.
      Buf x = in(n, 0);
      const std::string wn = wname(n);
      const auto& P = weight(wn);
      const int d = x.width;
      if (shapes_.at(wn).at(0) != d) throw std::runtime_error("IR pos_embedding '" + wn + "': width differs from x");
      if (static_cast<size_t>(start_pos + T) * d > P.size())
        throw std::runtime_error("IR pos_embedding: position " + std::to_string(start_pos + T - 1) + " exceeds '" + wn + "'");
      for (int t = 0; t < T; ++t)
        for (int i = 0; i < d; ++i) x.v[static_cast<size_t>(t) * d + i] += P[static_cast<size_t>(start_pos + t) * d + i];
      env[n.out.at(0)] = std::move(x);
    } else if (n.op == "layernorm") {  // (x - mean) / sqrt(biased var + eps) * w + b. ROUNDING.
      const Buf& x = in(n, 0);
      const auto& w = weight(wname(n));
      const auto& b = weight(layer_name(n.attrs.at("bias").get<std::string>(), layer));
      const float eps = n.attrs.value("eps", 1e-5f);
      const int d = x.width;
      Buf out{std::vector<float>(x.v.size()), d};
      for (int t = 0; t < T; ++t) {
        const float* r = x.v.data() + static_cast<size_t>(t) * d;
        double mean = 0.0, var = 0.0;
        for (int i = 0; i < d; ++i) mean += r[i];
        mean /= d;
        for (int i = 0; i < d; ++i) var += (r[i] - mean) * (r[i] - mean);
        var /= d;
        const float inv = static_cast<float>(1.0 / std::sqrt(var + eps));
        for (int i = 0; i < d; ++i)
          out.v[static_cast<size_t>(t) * d + i] = (r[i] - static_cast<float>(mean)) * inv * w[i] + b[i];
      }
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "split") {  // columns of [T, sum(sizes)] into the outputs, in order. EXACT.
      const Buf& x = in(n, 0);
      const auto sizes = n.attrs.at("sizes").get<std::vector<int>>();
      int total = 0;
      for (int s : sizes) total += s;
      if (sizes.size() != n.out.size()) throw std::runtime_error("IR split: sizes do not match the outputs");
      // groups g (revision 1.5): x is g groups of sum(sizes) columns, output k the k-th part of every group
      const int groups = n.attrs.contains("groups") ? n.attrs.at("groups").get<int>() : 1;
      if (groups < 1) throw std::runtime_error("IR_SPEC 5.5 (rule 10): split groups " + std::to_string(groups) + " is not positive");
      if (static_cast<int64_t>(total) * groups != x.width)
        throw std::runtime_error("IR_SPEC 5.5 (rule 37): split sizes sum to " + std::to_string(total) +
                                 (groups > 1 ? " times groups " + std::to_string(groups) : std::string()) +
                                 ", not the input width " + std::to_string(x.width));
      if (groups > 1) {
        std::vector<Buf> parts;
        std::vector<float*> ptrs;
        for (int s : sizes) parts.push_back(Buf{std::vector<float>(static_cast<size_t>(T) * groups * s), groups * s});
        for (auto& b : parts) ptrs.push_back(b.v.data());
        cpu::split_groups(x.v.data(), T, sizes, groups, ptrs);
        for (size_t o = 0; o < sizes.size(); ++o) env[n.out[o]] = std::move(parts[o]);
        continue;
      }
      int col = 0;
      for (size_t o = 0; o < sizes.size(); ++o) {
        Buf part{std::vector<float>(static_cast<size_t>(T) * sizes[o]), sizes[o]};
        for (int t = 0; t < T; ++t)
          std::memcpy(part.v.data() + static_cast<size_t>(t) * sizes[o], x.v.data() + static_cast<size_t>(t) * x.width + col,
                      sizes[o] * sizeof(float));
        col += sizes[o];
        env[n.out[o]] = std::move(part);
      }
    } else if (n.op == "gelu") {  // tanh: 0.5x(1 + tanh(sqrt(2/pi)(x + 0.044715x^3))); none: 0.5x(1 + erf(x/sqrt2)).
      const Buf& x = in(n, 0);
      const bool tanh_form = n.attrs.value("approx", "none") == "tanh";
      Buf out{std::vector<float>(x.v.size()), x.width};
      const float c = std::sqrt(2.0f / 3.14159265358979323846f);
      for (size_t i = 0; i < x.v.size(); ++i) {
        const float v = x.v[i];
        out.v[i] = tanh_form ? 0.5f * v * (1.0f + std::tanh(c * (v + 0.044715f * v * v * v)))
                             : 0.5f * v * (1.0f + std::erf(v / std::sqrt(2.0f)));
      }
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "rmsnorm") {
      const Buf& x = in(n, 0);
      Buf out{{}, x.width};
      out.v.resize(x.v.size());
      // group g: each consecutive g columns normalised on their own, one
      // weight [g] shared (per-head q/k norm); a row of rows of width g.
      const int g = has(n.attrs, "group") ? n.attrs.at("group").get<int>() : x.width;
      const auto& w = weight(wname(n));
      if (g > 0 && x.width % g != 0)
        throw std::runtime_error("IR_SPEC 5.6 (rule 37): rmsnorm '" + wname(n) + "': group " + std::to_string(g) +
                                 " does not divide the input width " + std::to_string(x.width));
      if (g <= 0 || x.width % g != 0 || w.size() != static_cast<size_t>(g))
        throw std::runtime_error("IR rmsnorm '" + wname(n) + "': group " + std::to_string(g) + " does not fit width " +
                                 std::to_string(x.width) + " and weight [" + std::to_string(w.size()) + "]");
      rmsnorm_batch_cpu(x.v, w, out.v, T * (x.width / g), g, n.attrs.value("eps", 1e-5f));
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "linear") {
      const Buf& x = in(n, 0);
      const std::string wn = wname(n);
      const Weight& W = bound(wn);
      if (shapes_.at(wn).at(0) != x.width)
        throw std::runtime_error("IR linear '" + wn + "': input width " + std::to_string(x.width) +
                                 " does not match the weight's ne0 " + std::to_string(shapes_.at(wn).at(0)));
      const int in_dim = x.width, out_dim = static_cast<int>(W.numel / in_dim);
      Buf out{std::vector<float>(static_cast<size_t>(T) * out_dim), out_dim};
      if (W.quantised())  // one row of blocks at a time; equal to the f32 path (graph_ir.h)
        matmul_rows_f32_batch_cpu(
            [&](int64_t o, float* row) {
              dequant_span_to_f32(W.type, W.qp, static_cast<size_t>(o) * in_dim, in_dim, row);
            },
            x.v.data(), out.v.data(), T, out_dim, in_dim);
      else
        matmul_f32_f32_batch_cpu(W.f32, x.v, out.v, T, out_dim, in_dim);
      if (has(n.attrs, "bias")) {
        const auto& b = weight(layer_name(n.attrs.at("bias").get<std::string>(), layer));
        for (int t = 0; t < T; ++t)
          for (int o = 0; o < out_dim; ++o) out.v[static_cast<size_t>(t) * out_dim + o] += b[o];
      }
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "rope") {
      Buf q = in(n, 0), k = in(n, 1);
      const int hd = n.attrs.at("head_dim").get<int>();
      rope_widths(n, layer, q.width, k.width);
      const auto& freqs = rope_table(layer);
      if (!rope_neox_ && rope_dim_ == head_dim_) {  // full interleaved: the llama path's kernel
        apply_rope_batch_cpu(q.v, k.v, T, q.width / hd, k.width / hd, hd, start_pos, freqs, max_len_,
                             /*use_adjacent_pairing=*/true);
      } else {  // NEOX pairing or partial rotary
        if (start_pos < 0 || start_pos + T > max_len_)
          throw std::runtime_error("IR rope: position " + std::to_string(start_pos + T - 1) + " exceeds the rope table (" +
                                   std::to_string(max_len_) + ")");
        if (q.width % hd != 0 || k.width % hd != 0) throw std::runtime_error("IR rope: q/k widths are not whole heads");
        rope_rows_cpu(q.v, T, q.width / hd, hd, rope_dim_, rope_neox_, start_pos, freqs);
        rope_rows_cpu(k.v, T, k.width / hd, hd, rope_dim_, rope_neox_, start_pos, freqs);
      }
      env[n.out.at(0)] = std::move(q);
      env[n.out.at(1)] = std::move(k);
    } else if (n.op == "attention") {
      const Buf &q = in(n, 0), &k = in(n, 1), &v = in(n, 2);
      const int hd = n.attrs.at("head_dim").get<int>(), hq = n.attrs.at("heads").get<int>();
      const int hkv = n.attrs.at("kv_heads").get<int>();
      if (hd != head_dim_ || hkv != kv_heads_) throw std::runtime_error("IR attention shape differs from its kv state");
      // Under a recursion step (mor_step) the rows go to the application's compact cache, at
      // attn_row_, and the query reads the rows before and at it; under turn_state they go to the
      // compact cache at kv_row0_ and each query reads its visible rows (ts_vis_); elsewhere rows
      // sit at their positions.
      const bool ts = n.attrs.value("mask", "causal") == "turn_state";
      const int kpos = ts ? kv_row0_ : attn_row_ >= 0 ? attn_row_ : start_pos;
      if (kpos + T > max_len_)
        throw std::runtime_error("IR attention: position " + std::to_string(kpos + T) + " exceeds the kv state (" +
                                 std::to_string(max_len_) + ")");
      KVCacheLayer& cache = kv_.at(layer);
      const size_t row = static_cast<size_t>(hkv) * hd;
      std::memcpy(cache.k.data() + kpos * row, k.v.data(), static_cast<size_t>(T) * row * sizeof(float));
      std::memcpy(cache.v.data() + kpos * row, v.v.data(), static_cast<size_t>(T) * row * sizeof(float));
      if (kv_fmt_ != ember_kvq::kF32) {  // EMBER_KV: the rows as the quantised cache gives them back
        ember_kvq::roundtrip_rows(kv_fmt_, cache.k.data() + kpos * row, static_cast<size_t>(T), row);
        ember_kvq::roundtrip_rows(kv_fmt_, cache.v.data() + kpos * row, static_cast<size_t>(T), row);
      }
      const float scale = n.attrs.contains("scale") ? n.attrs.at("scale").get<float>() : 1.0f / std::sqrt(float(hd));
      Buf out{{}, hq * hd};
      int window = 0, sinks = 0;
      attention_window(n, layer, window, sinks);
      // Gemma 2's soft cap of the scaled scores (IR_SPEC 5.15, revision 1.5), before any position bias; 0: none.
      const float softcap = has(n.attrs, "softcap") ? at_layer(n.attrs.at("softcap"), std::max(layer, 0)).get<float>() : 0.0f;
      if (has(n.attrs, "position_bias")) {  // T5's relative position bias (IR_SPEC 5.15, revision 1.3)
        const json& pb = n.attrs.at("position_bias");
        const std::vector<int> buckets = pb.at("buckets").get<std::vector<int>>();
        const auto& W = weight(layer_name(pb.at("weight").get<std::string>(), layer));
        std::vector<int> qpos(static_cast<size_t>(T));
        for (int t = 0; t < T; ++t) qpos[t] = start_pos + t;
        // under turn_state the compact cache's row r holds position kv_pos_[r]; elsewhere row r is position r
        const std::vector<std::vector<int>> rows = ts ? ts_vis_ : attention_visible_rows(T, start_pos, window, sinks);
        attention_bias_batch_cpu(q.v, cache, out.v, T, rows, ts ? &kv_pos_ : nullptr, qpos, hq, hkv, hd, scale, W.data(),
                                 buckets, softcap);
      } else if (has(n.attrs, "relative_kv")) {  // Shaw's relative key and value vectors (IR_SPEC 5.15, revision 1.4)
        const json& rk = n.attrs.at("relative_kv");
        const int K = rk.at("max_distance").get<int>();
        // a table [K + 1, w], ne {w, K + 1}, w head_dim or heads * head_dim (rule 40, checked at load)
        auto table = [&](const char* member, int& w) -> const float* {
          w = 0;
          if (!rk.contains(member) || !rk.at(member).is_string()) return nullptr;
          const std::string name = layer_name(rk.at(member).get<std::string>(), layer);
          const auto& s = shapes_.at(name);
          if (s.size() != 2 || s[1] != K + 1 || (s[0] != hd && s[0] != static_cast<int64_t>(hq) * hd))
            throw std::runtime_error("IR_SPEC 5.15 (rule 40): relative_kv " + std::string(member) + " '" + name +
                                     "' is not of ne {head_dim or heads * head_dim, max_distance + 1}");
          w = static_cast<int>(s[0]);
          return weight(name).data();
        };
        int rk_w = 0, rv_w = 0;
        const float* RK = table("key", rk_w);
        const float* RV = table("value", rv_w);
        std::vector<int> qpos(static_cast<size_t>(T));
        for (int t = 0; t < T; ++t) qpos[t] = start_pos + t;
        const std::vector<std::vector<int>> rows = ts ? ts_vis_ : attention_visible_rows(T, start_pos, window, sinks);
        attention_relkv_batch_cpu(q.v, cache, out.v, T, rows, ts ? &kv_pos_ : nullptr, qpos, hq, hkv, hd, scale, RK, rk_w,
                                  RV, rv_w, K, softcap);
      } else if (ts)
        attention_rows_batch_cpu(q.v, cache, out.v, T, ts_vis_, hq, hkv, hd, scale, softcap);
      else if (window > 0 && attn_row_ >= 0) {
        // A window under depth_route (IR_SPEC 5.15 and 7.5, v1.4.2, Q16): the application's compact cache holds
        // positions mor_rows_[layer] at rows 0 .. attn_row_ - 1, and mor_step runs one token at a time, its row at
        // attn_row_ (position p = start_pos). A held row at position j is visible iff j < sink_hi or j >= win_lo
        // (swa_visible_ranges): the sinks and positions p - window + 1 .. p it holds, counted in positions, never the
        // last `window` rows; rows of tokens that exited are absent, not replaced.
        const std::vector<int>& held = mor_rows_.at(layer);
        if (T != 1 || static_cast<int>(held.size()) != attn_row_)
          throw std::runtime_error("IR depth_route: application " + std::to_string(layer) + " runs " +
                                   std::to_string(T) + " rows over " + std::to_string(held.size()) +
                                   " held positions for " + std::to_string(attn_row_) + " rows");
        int sink_hi = 0, win_lo = 0;
        swa_visible_ranges(start_pos, window, sinks, sink_hi, win_lo);
        std::vector<std::vector<int>> rows(1);
        for (int r = 0; r < attn_row_; ++r)
          if (held[r] < sink_hi || held[r] >= win_lo) rows[0].push_back(r);
        rows[0].push_back(attn_row_);  // its own row, position p
        attention_rows_batch_cpu(q.v, cache, out.v, 1, rows, hq, hkv, hd, scale, softcap);
      } else if (window > 0)
        attention_window_batch_cpu(q.v, cache, out.v, T, start_pos, hq, hkv, hd, scale, window, sinks, softcap);
      else
        attention_batch_cpu(q.v, cache, out.v, T, kpos, hq, hkv, hd, scale, softcap);
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "mla_attention") {  // multi-head latent attention (IR_SPEC 5.20). ROUNDING.
      const Buf &qn = in(n, 0), &qr = in(n, 1), &c = in(n, 2), &kr = in(n, 3);
      const int L = std::max(layer, 0);
      auto at = [&](const char* k) { return at_layer(n.attrs.at(k), L).get<int>(); };
      const int H = at("heads"), r = at("kv_rank"), dn = at("qk_nope"), dr = at("qk_rope"), dv = at("v_head_dim");
      if (qn.width != H * dn || qr.width != H * dr || c.width != r || kr.width != dr)
        throw std::runtime_error("IR_SPEC 5.20 (rule 37): mla_attention reads qn " + std::to_string(qn.width) + ", qr " +
                                 std::to_string(qr.width) + ", c " + std::to_string(c.width) + " and kr " +
                                 std::to_string(kr.width) + " wide, not heads x qk_nope " + std::to_string(H * dn) +
                                 ", heads x qk_rope " + std::to_string(H * dr) + ", kv_rank " + std::to_string(r) +
                                 " and qk_rope " + std::to_string(dr));
      if (start_pos + T > max_len_)
        throw std::runtime_error("IR mla_attention: position " + std::to_string(start_pos + T) + " exceeds the kv state (" +
                                 std::to_string(max_len_) + ")");
      // The slot's row of position p is [c_p | kr_p] (kv_heads 1, head_dim r + dr; rule 12 at load).
      KVCacheLayer& cache = kv_.at(layer);
      const size_t w = static_cast<size_t>(r) + dr;
      for (int t = 0; t < T; ++t) {
        float* dst = cache.k.data() + static_cast<size_t>(start_pos + t) * w;
        std::memcpy(dst, c.v.data() + static_cast<size_t>(t) * r, static_cast<size_t>(r) * sizeof(float));
        std::memcpy(dst + r, kr.v.data() + static_cast<size_t>(t) * dr, static_cast<size_t>(dr) * sizeof(float));
      }
      int window = 0, sinks = 0;
      attention_window(n, layer, window, sinks);
      const auto& Kb = weight(layer_name(n.attrs.at("k_b").get<std::string>(), layer));
      const auto& Vb = weight(layer_name(n.attrs.at("v_b").get<std::string>(), layer));
      Buf out{std::vector<float>(static_cast<size_t>(T) * H * dv), H * dv};
      cpu::mla_attention(qn.v.data(), qr.v.data(), cache.k.data(), out.v.data(), T,
                         attention_visible_rows(T, start_pos, window, sinks), H, r, dn, dr, dv, Kb.data(), Vb.data(),
                         n.attrs.at("scale").get<float>());
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "add" || n.op == "mul") {
      const Buf &a = in(n, 0), &b = in(n, 1);
      if (a.v.size() != b.v.size() || a.width != b.width) binary_shapes(n.op, a.width, b.width);
      Buf out{std::vector<float>(a.v.size()), a.width};
      if (n.op == "add")
        for (size_t i = 0; i < a.v.size(); ++i) out.v[i] = a.v[i] + b.v[i];
      else
        for (size_t i = 0; i < a.v.size(); ++i) out.v[i] = a.v[i] * b.v[i];
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "moe_ffn") {
      const Buf& x = in(n, 0);
      const auto& ex = n.attrs.at("experts");
      const std::string gn = layer_name(ex.at("gate").get<std::string>(), layer);
      const auto& R = weight(layer_name(n.attrs.at("router").get<std::string>(), layer));
      const auto& G = weight(gn);
      const auto& U = weight(layer_name(ex.at("up").get<std::string>(), layer));
      const auto& D = weight(layer_name(ex.at("down").get<std::string>(), layer));
      const int E = n.attrs.at("n_experts").get<int>(), k = n.attrs.at("top_k").get<int>();
      const auto& gs = shapes_.at(gn);  // ggml order {d, d_ff, E}
      if (gs.size() != 3 || gs[0] != x.width || gs[2] != E)
        throw std::runtime_error("IR moe_ffn '" + gn + "': expert stack shape does not match d and n_experts");
      const int d = x.width, d_ff = static_cast<int>(gs[1]);
      const bool renorm = n.attrs.value("gate_norm", "renorm") == "renorm";
      // DeepSeek-V3's router (IR_SPEC 5.16, revision 1.6): sigmoid scores, a selection bias, group-limited routing.
      const bool sigmoid = has(n.attrs, "scoring") && at_layer(n.attrs.at("scoring"), std::max(layer, 0)) == "sigmoid";
      const float* bias =
          has(n.attrs, "selection_bias") ? weight(layer_name(n.attrs.at("selection_bias").get<std::string>(), layer)).data()
                                         : nullptr;
      int groups_n = 1, groups_top = 1;
      if (has(n.attrs, "groups")) {
        groups_n = n.attrs.at("groups").at("n").get<int>();
        groups_top = n.attrs.at("groups").at("top").get<int>();
        if (groups_n < 1 || E % groups_n != 0 || groups_top < 1 || groups_top > groups_n || groups_top * (E / groups_n) < k)
          throw std::runtime_error("IR_SPEC 5.16 (rule 10): moe_ffn groups {n " + std::to_string(groups_n) + ", top " +
                                   std::to_string(groups_top) + "} do not fit n_experts " + std::to_string(E) + " and top_k " +
                                   std::to_string(k));
      }
      Buf out{std::vector<float>(static_cast<size_t>(T) * d), d};
#pragma omp parallel for schedule(static)
      for (int t = 0; t < T; ++t)
        moe_ffn_routed_f32_cpu(x.v.data() + static_cast<size_t>(t) * d, d, d_ff, E, k, R.data(), G.data(), U.data(),
                               D.data(), out.v.data() + static_cast<size_t>(t) * d, renorm, sigmoid, bias, groups_n, groups_top);
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "silu") {
      const Buf& x = in(n, 0);
      Buf out{std::vector<float>(x.v.size()), x.width};
      silu_cpu(x.v, out.v);
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "sub") {  // a - b. EXACT.
      const Buf &a = in(n, 0), &b = in(n, 1);
      if (a.v.size() != b.v.size() || a.width != b.width) binary_shapes(n.op, a.width, b.width);
      Buf out{std::vector<float>(a.v.size()), a.width};
      cpu::sub(a.v.data(), b.v.data(), out.v.data(), a.v.size());
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "relu" || n.op == "sigmoid") {  // relu EXACT; sigmoid ROUNDING.
      const Buf& x = in(n, 0);
      Buf out{std::vector<float>(x.v.size()), x.width};
      (n.op == "relu" ? cpu::relu : cpu::sigmoid)(x.v.data(), out.v.data(), x.v.size());
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "affine") {  // (x * scale) + shift. EXACT.
      const Buf& x = in(n, 0);
      Buf out{std::vector<float>(x.v.size()), x.width};
      cpu::affine(x.v.data(), out.v.data(), x.v.size(), at_layer(n.attrs.at("scale"), layer).get<float>(),
                  at_layer(n.attrs.at("shift"), layer).get<float>());
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "scale") {  // (x * g) + b per feature, g and b [d]. EXACT.
      const Buf& x = in(n, 0);
      const std::string wn = wname(n);
      const auto& g = weight(wn);
      const int d = x.width;
      if (d <= 0 || g.size() != static_cast<size_t>(d))
        throw std::runtime_error("IR_SPEC 5.4 (rule 37): scale '" + wn + "' has " + std::to_string(g.size()) +
                                 " values for width " + std::to_string(d));
      const float* b = nullptr;
      if (has(n.attrs, "bias")) {
        const std::string bn = layer_name(n.attrs.at("bias").get<std::string>(), layer);
        const auto& B = weight(bn);
        if (B.size() != static_cast<size_t>(d))
          throw std::runtime_error("IR_SPEC 5.4 (rule 37): scale bias '" + bn + "' has " + std::to_string(B.size()) +
                                   " values for width " + std::to_string(d));
        b = B.data();
      }
      Buf out{std::vector<float>(x.v.size()), d};
      cpu::scale(x.v.data(), g.data(), b, out.v.data(), static_cast<int>(x.v.size() / d), d);
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "mul_col") {  // x [T, w] * c [T, 1]. EXACT.
      const Buf &x = in(n, 0), &c = in(n, 1);
      if (c.width != 1 || c.v.size() != static_cast<size_t>(T) || x.v.size() != static_cast<size_t>(T) * x.width)
        throw std::runtime_error("IR_SPEC 5.12 (rule 37): mul_col: c must be [T, 1] beside x [T, w]");
      Buf out{std::vector<float>(x.v.size()), x.width};
      cpu::mul_col(x.v.data(), c.v.data(), out.v.data(), T, x.width);
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "rms_row") {  // [T, w] -> [T, 1]: sqrt(ss / w + eps). ROUNDING.
      const Buf& x = in(n, 0);
      if (x.width <= 0 || x.v.size() != static_cast<size_t>(T) * x.width)
        throw std::runtime_error("IR rms_row: input is not [T, w]");
      Buf out{std::vector<float>(static_cast<size_t>(T)), 1};
      cpu::rms_row(x.v.data(), out.v.data(), T, x.width, n.attrs.at("eps").get<float>());
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "expr") {  // runtime-defined elementwise op (graph_ir_expr.h). ROUNDING.
      std::vector<const float*> ptrs;
      const Buf& x0 = in(n, 0);
      for (size_t k = 0; k < n.in.size(); ++k) {
        const Buf& a = in(n, k);
        if (a.width != x0.width || a.v.size() != x0.v.size())
          throw std::runtime_error("IR_SPEC 5.19 (rule 37): expr inputs must share one shape [T, w]");
        ptrs.push_back(a.v.data());
      }
      if (x0.width <= 0 || x0.v.size() % x0.width) throw std::runtime_error("IR expr: input is not [T, w]");
      Buf out{std::vector<float>(x0.v.size()), x0.width};
      expr::eval(expr::cached(n.attrs.at("expr").get<std::string>(), n.in.size()), ptrs,
                 static_cast<int>(x0.v.size() / x0.width), x0.width, out.v.data());
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "list") {  // the inputs, in order, as one list value. EXACT.
      Buf out;
      out.is_list = true;
      for (size_t k = 0; k < n.in.size(); ++k) out.items.push_back(in(n, k));
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "append") {  // the list with x added at the end. EXACT.
      const Buf& x = in(n, 1);
      // The usual form appends in place (out names the list): no copy.
      Buf out;
      if (n.out.at(0) == n.in.at(0))
        out = std::move(list_in(n, 0));
      else
        out = list_in(n, 0);
      out.items.push_back(x);
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "depth_attend") {  // attention over depth (AttnRes block form). ROUNDING.
      const Buf& L = list_in(n, 0);
      std::vector<const float*> src;
      int d = 0;
      auto add_source = [&](const Buf& s) {
        if (d == 0) d = s.width;
        if (s.width != d || d <= 0 || s.v.size() != static_cast<size_t>(T) * d)
          throw std::runtime_error("IR_SPEC 5.18 (rule 37): depth_attend sources must all be one shape [T, d]");
        src.push_back(s.v.data());
      };
      for (const Buf& s : L.items) add_source(s);
      for (size_t k = 1; k < n.in.size(); ++k) add_source(in(n, k));
      if (src.empty()) throw std::runtime_error("IR depth_attend: no sources");
      const float* w = nullptr;
      if (src.size() > 1) {  // one source: out = s_0 and the pseudo-query is not read
        const std::string wn = layer_name(n.attrs.at("weight").get<std::string>(), layer);
        if (!weights_.count(wn)) throw std::runtime_error("IR depth_attend: pseudo-query '" + wn + "' is not in the GGUF");
        const auto& W = weight(wn);
        if (W.size() != static_cast<size_t>(d))
          throw std::runtime_error("IR depth_attend: pseudo-query '" + wn + "' is not [d]");
        w = W.data();
      }
      Buf out{std::vector<float>(static_cast<size_t>(T) * d), d};
      cpu::depth_attend(src, w, out.v.data(), T, d, n.attrs.value("eps", 1e-5f));
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "linear_attention") {  // elu+1 linear attention over the slot's num, den. ROUNDING.
      const Buf &q = in(n, 0), &k = in(n, 1), &v = in(n, 2);
      const int w = n.attrs.at("heads").get<int>() * n.attrs.at("head_dim").get<int>();
      if (q.width != w || k.width != w || v.width != w)
        throw std::runtime_error("IR linear_attention: q/k/v widths do not match heads * head_dim");
      const RecSlot& slot = rec_slot(n, "linear_attention");
      float* num = rec_tensor(slot, "num", layer, w);
      float* den = rec_tensor(slot, "den", layer, w);
      const float floor = n.attrs.value("den_floor", 1e-8f);
      Buf out{std::vector<float>(static_cast<size_t>(T) * w), w};
      for (int t = 0; t < T; ++t) {  // the decode update for t = 0..T-1, in order
        const size_t r = static_cast<size_t>(t) * w;
        linear_attention_row_cpu(q.v.data() + r, k.v.data() + r, v.v.data() + r, num, den, out.v.data() + r, w, floor);
        rec_row_done(slot, layer, t);
      }
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "ssm_scan") {  // the sequential state-space scan over the slot's s. ROUNDING.
      const Buf& u = in(n, 0);
      const int sn = n.attrs.at("n").get<int>();
      if (u.width != sn) throw std::runtime_error("IR ssm_scan: input width differs from n");
      const RecSlot& slot = rec_slot(n, "ssm_scan");
      float* s = rec_tensor(slot, "s", layer, sn);
      const auto& A = weight(layer_name(ssm_weight(n.attrs, "a"), layer));
      const auto& B = weight(layer_name(ssm_weight(n.attrs, "b"), layer));
      const auto& C = weight(layer_name(ssm_weight(n.attrs, "c"), layer));
      const float D = weight(layer_name(ssm_weight(n.attrs, "d"), layer)).at(0);
      std::vector<float> scratch(sn);
      Buf out{std::vector<float>(static_cast<size_t>(T) * sn), sn};
      for (int t = 0; t < T; ++t) {
        const size_t r = static_cast<size_t>(t) * sn;
        ssm_scan_row_cpu(u.v.data() + r, A.data(), B.data(), C.data(), D, s, scratch.data(), out.v.data() + r, sn);
        rec_row_done(slot, layer, t);
      }
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "conv1d_causal") {  // depthwise causal conv over the slot's previous inputs. ROUNDING.
      const Buf& x = in(n, 0);
      const int C = n.attrs.at("channels").get<int>(), K = n.attrs.at("kernel").get<int>();
      if (x.width != C) throw std::runtime_error("IR conv1d_causal: input width differs from channels");
      const RecSlot& slot = rec_slot(n, "conv1d_causal");
      float* st = rec_tensor(slot, "x", layer, static_cast<size_t>(K - 1) * C);
      const auto& w = weight(layer_name(n.attrs.at("weight").get<std::string>(), layer));
      const float* bias =
          has(n.attrs, "bias") ? weight(layer_name(n.attrs.at("bias").get<std::string>(), layer)).data() : nullptr;
      Buf out{std::vector<float>(static_cast<size_t>(T) * C), C};
      for (int t = 0; t < T; ++t) {
        const size_t r = static_cast<size_t>(t) * C;
        cpu::conv1d_causal_row(x.v.data() + r, w.data(), bias, st, out.v.data() + r, C, K);
        rec_row_done(slot, layer, t);
      }
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "selective_scan") {  // Mamba-1's scan over the slot's s. ROUNDING.
      const Buf &x = in(n, 0), &dt = in(n, 1), &bm = in(n, 2), &cm = in(n, 3);
      const int D = n.attrs.at("channels").get<int>(), N = n.attrs.at("d_state").get<int>();
      if (x.width != D || dt.width != D || bm.width != N || cm.width != N)
        throw std::runtime_error("IR selective_scan: x/dt must be [T, channels] and B/C [T, d_state]");
      const RecSlot& slot = rec_slot(n, "selective_scan");
      float* s = rec_tensor(slot, "s", layer, static_cast<size_t>(D) * N);
      const auto& A = weight(layer_name(n.attrs.at("a").get<std::string>(), layer));
      const float* d = has(n.attrs, "d") ? weight(layer_name(n.attrs.at("d").get<std::string>(), layer)).data() : nullptr;
      const bool softplus = n.attrs.value("dt_softplus", true);
      Buf out{std::vector<float>(static_cast<size_t>(T) * D), D};
      for (int t = 0; t < T; ++t) {
        cpu::selective_scan_row(x.v.data() + static_cast<size_t>(t) * D, dt.v.data() + static_cast<size_t>(t) * D,
                                bm.v.data() + static_cast<size_t>(t) * N, cm.v.data() + static_cast<size_t>(t) * N,
                                A.data(), d, s, out.v.data() + static_cast<size_t>(t) * D, D, N, softplus);
        rec_row_done(slot, layer, t);
      }
      env[n.out.at(0)] = std::move(out);
    } else if (n.op == "recurrence") {  // a recurrent op the document defines (IR_SPEC 6.10). EXACT or ROUNDING.
      const rec::Recurrence& prog = *recs_.at(&n);
      const RecSlot& slot = rec_slot(n, "recurrence");
      std::vector<const Buf*> ins;
      for (size_t k = 0; k < n.in.size(); ++k) {
        const Buf& b = in(n, k);
        if (b.is_list || static_cast<size_t>(b.width) != prog.input_widths()[k] ||
            b.v.size() != static_cast<size_t>(T) * b.width)
          throw std::runtime_error("IR_SPEC 6.10 (rule 34): recurrence input '" + n.in[k] + "' is " +
                                   std::to_string(b.width) + " wide; its declaration is " +
                                   std::to_string(prog.input_widths()[k]));
        ins.push_back(&b);
      }
      std::vector<const float*> w;
      for (const auto& wd : prog.weights()) w.push_back(weight(layer_name(wd.tensor, layer)).data());
      std::vector<float*> st;
      for (const auto& s : prog.state())
        st.push_back(rec_tensor(slot, s.name.c_str(), layer, s.numel));
      const size_t ow = prog.out_width();
      Buf out{std::vector<float>(static_cast<size_t>(T) * ow), static_cast<int>(ow)};
      std::vector<const float*> row(ins.size());
      for (int t = 0; t < T; ++t) {  // token by token, in order (prefill = decode)
        for (size_t k = 0; k < ins.size(); ++k) row[k] = ins[k]->v.data() + static_cast<size_t>(t) * ins[k]->width;
        prog.step(row, w, st, out.v.data() + static_cast<size_t>(t) * ow);
        rec_row_done(slot, layer, t);
      }
      env[n.out.at(0)] = std::move(out);
    }
  }
}

References a, apply_rope_batch_cpu(), attention_batch_cpu(), attention_bias_batch_cpu(), attention_relkv_batch_cpu(), attention_rows_batch_cpu(), attention_visible_rows(), attention_window_batch_cpu(), c, col, dequant_span_to_f32(), id, in, ember_ir::Interpreter::Buf::is_list, j, k, KVCacheLayer::k, ember_kvq::kF32, ember_ir::linear_attention_row_cpu(), matmul_f32_f32_batch_cpu(), matmul_rows_f32_batch_cpu(), moe_ffn_routed_f32_cpu(), n, ember_ir::Interpreter::RecSlot::name, name, out, rmsnorm_batch_cpu(), ember_kvq::roundtrip_rows(), scale, sigmoid, silu_cpu(), sizes, ember_ir::ssm_scan_row_cpu(), swa_visible_ranges(), ember_ir::Interpreter::Buf::v, v, value, and ember_ir::Interpreter::Buf::width.

◆ save_state()

Interpreter::State ember_ir::Interpreter::save_state ( ) const

Definition at line 3156 of file graph_ir.cpp.

                                               {
  cpu_only(gpu_, "save_state");
  State s;
  s.seq_len = seq_len_;
  const size_t row = static_cast<size_t>(kv_heads_) * head_dim_;
  for (int l = 0; l < static_cast<int>(kv_.size()); ++l) {
    const size_t n = static_cast<size_t>(kv_rows_in_use(l)) * row;
    s.k.emplace_back(kv_[l].k.begin(), kv_[l].k.begin() + n);
    s.v.emplace_back(kv_[l].v.begin(), kv_[l].v.begin() + n);
  }
  s.ts_seg = ts_seg_;
  s.ts_sidx = ts_sidx_;
  s.kv_pos = kv_pos_;
  s.evicted_below = evicted_below_;
  s.mor_rows = mor_rows_;
  s.mor_depth = mor_depth_;
  s.rec = rec_;
  s.rec_pos = rec_pos_;
  s.rec_snaps = rec_snaps_;
  s.ids = ids_;
  s.rope_long = rope_long_;
  return s;
}

References ember_ir::cpu_only(), j, k, n, ember_ir::Interpreter::State::seq_len, and v.

◆ seq_len()

int ember_ir::Interpreter::seq_len ( ) const

inline

Definition at line 182 of file graph_ir.h.

182{ return seq_len_; }

References seq_len_.

◆ set_rope_long()

void ember_ir::Interpreter::set_rope_long ( bool v )

private

Definition at line 1863 of file graph_ir.cpp.

                                      {
  if (v == rope_long_) return;
  rope_long_ = v;
#ifdef HAS_CUDA
  if (gpu_) ++alloc_epoch_;  // the captured decode step reads the other set's tables: it is captured again
#endif
}

References j, and v.

◆ state_k()

int ember_ir::Interpreter::state_k ( ) const

inline

Definition at line 167 of file graph_ir.h.

167{ return state_k_; }

References state_k_.

◆ truncate()

void ember_ir::Interpreter::truncate ( int n )

Definition at line 3112 of file graph_ir.cpp.

                                {
  cpu_only(gpu_, "truncate");
  if (n < 0 || n > seq_len_)
    throw std::runtime_error("IR truncate to " + std::to_string(n) + ": the state covers " + std::to_string(seq_len_) +
                             " positions");
  if (rec_floats_ > 0 && rec_pos_ < 0) throw std::runtime_error("IR truncate: the recurrent state is not usable (a call failed)");
  if (resume_point(n) != n)
    throw std::runtime_error("IR truncate to " + std::to_string(n) + ": the state can resume only at " +
                             std::to_string(resume_point(n)) + " (re-prefill from there)");
  if (n == seq_len_) return;
  if (state_k_ > 0 || turn_state_) {
    // n is its segment's start when that segment's text rows are gone (resume_point): every row
    // of segments >= m lies at or past n, so dropping positions >= n leaves segments < m only.
    const int m = n > 0 ? ts_seg_after(n - 1) : 0;
    if (turn_state_ && m < evicted_below_) evicted_below_ = m;
    ts_seg_.resize(n);
    ts_sidx_.resize(n);
    while (!kv_pos_.empty() && kv_pos_.back() >= n) kv_pos_.pop_back();
  }
  for (auto& rows : mor_rows_)
    while (!rows.empty() && rows.back() >= n) rows.pop_back();
  if (rec_floats_ > 0 && n != rec_pos_) {
    if (n == 0) {
      rec_ = rec_init_;  // every application's initial state
      rec_snaps_.clear();
    } else {
      rec_ = rec_snaps_.at(n);
      rec_snaps_.erase(rec_snaps_.upper_bound(n), rec_snaps_.end());
    }
    rec_pos_ = n;
  }
  seq_len_ = n;
  if (rope_ff_ctx_ > 0) {  // LongRoPE with two sets (IR_SPEC 6.9, revision 1.5)
    ids_.resize(static_cast<size_t>(n));
    if (n > 0 && (n > rope_ff_ctx_) != rope_long_) {
      // The kept rows were rotated with the other set: the first n ids again, with the set n selects.
      const std::vector<int> prefix = ids_;
      set_rope_long(n > rope_ff_ctx_);
      Logger::info("[IR] LongRoPE: a truncate to " + std::to_string(n) + " changes the set; the rows recomputed");
      forward_impl(prefix, 0, false, 1);
    }
  }
}

References ember_ir::cpu_only(), Logger::info(), j, and n.

◆ ts_begin()

void ember_ir::Interpreter::ts_begin	(	const std::vector< int > &	tokens,
		int	start_pos
	)

private

Definition at line 2146 of file graph_ir.cpp.

                                                                      {
  const int n = static_cast<int>(ts_seg_.size()), T = static_cast<int>(tokens.size());
  if (start_pos < 0 || start_pos > n)
    throw std::runtime_error("IR state tokens: a forward at position " + std::to_string(start_pos) +
                             ", but the sequence holds " + std::to_string(n) + " positions");
  // The segment the first token lands in, and how much of a state block is open before it.
  const int m = ts_seg_after(start_pos - 1);
  int open = 0;
  if (start_pos > 0 && ts_sidx_[start_pos - 1] >= 0 && ts_sidx_[start_pos - 1] < state_k_ - 1)
    open = ts_sidx_[start_pos - 1] + 1;
  // Ids alone fix s(p) and c(p): a state id must be the next slot of an open block (S_0 opens
  // one), a text id must not interrupt one. Checked before anything changes.
  std::vector<int> segs(T), sidx(T);
  int seg = m;
  for (int t = 0; t < T; ++t) {
    const int id = tokens[t];
    if (id >= vocab_ && state_k_ > 0) {
      const int j = id - vocab_;
      if (j >= state_k_)
        throw std::runtime_error("IR state tokens: id " + std::to_string(id) + " is past the " + std::to_string(state_k_) +
                                 " state slots (vocab " + std::to_string(vocab_) + ")");
      if (j != open)
        throw std::runtime_error("IR state tokens: state id S_" + std::to_string(j) + " at position " +
                                 std::to_string(start_pos + t) + ", expected " +
                                 (open == 0 ? std::string("S_0 to open a block") : "S_" + std::to_string(open)));
      segs[t] = seg;
      sidx[t] = j;
      if (++open == state_k_) {
        open = 0;
        ++seg;
      }
    } else {
      if (open != 0)
        throw std::runtime_error("IR state tokens: text id " + std::to_string(id) + " at position " +
                                 std::to_string(start_pos + t) + " inside an incomplete state block (S_0 .. S_" +
                                 std::to_string(open - 1) + " of " + std::to_string(state_k_) + ")");
      segs[t] = seg;
      sidx[t] = -1;
    }
  }
  // A rewind into a segment whose text rows are gone can resume only at that segment's start
  // (resume_point); the rows of segments >= m all lie at or past it and are dropped below.
  if (turn_state_ && m < evicted_below_) {
    const int s0 = ts_segment_start(m);
    if (start_pos != s0)
      throw std::runtime_error("IR state tokens: position " + std::to_string(start_pos) + " is inside turn " +
                               std::to_string(m) + ", whose text rows were evicted; resume from its start, position " +
                               std::to_string(s0) + " (resume_point)");
    evicted_below_ = m;
  }
  ts_seg_.resize(start_pos);
  ts_sidx_.resize(start_pos);
  ts_seg_.insert(ts_seg_.end(), segs.begin(), segs.end());
  ts_sidx_.insert(ts_sidx_.end(), sidx.begin(), sidx.end());
  if (!turn_state_) return;
  while (!kv_pos_.empty() && kv_pos_.back() >= start_pos) kv_pos_.pop_back();
  // Rows no query of this call can see: earlier segments' text.
  if (evict_ && T > 0) ts_evict_below(segs[0]);
  kv_row0_ = static_cast<int>(kv_pos_.size());
  if (kv_row0_ + T > max_len_)
    throw std::runtime_error("IR attention: " + std::to_string(kv_row0_ + T) + " kv rows exceed the kv state (" +
                             std::to_string(max_len_) + ")");
  // Visible rows per query (section 1.3), ascending: held rows, then this call's rows up to t. The
  // device path's kernel reads each row's code instead (ts_upload_codes).
  ts_vis_.assign(gpu_ ? 0 : T, {});
  for (int t = 0; t < T && !gpu_; ++t) {
    const int s = segs[t];
    auto visible = [&](int j) { return ts_seg_[j] == s || (ts_sidx_[j] >= 0 && ts_seg_[j] < s); };
    std::vector<int>& v = ts_vis_[t];
    for (int r = 0; r < kv_row0_; ++r)
      if (visible(kv_pos_[r])) v.push_back(r);
    for (int u = 0; u <= t; ++u)
      if (visible(start_pos + u)) v.push_back(kv_row0_ + u);
  }
}

References j, n, and v.

◆ ts_evict_below()

void ember_ir::Interpreter::ts_evict_below ( int seg )

private

Definition at line 2101 of file graph_ir.cpp.

                                        {
  if (!evict_ || seg <= evicted_below_) return;
  // Keep every state row and the text rows of segments >= seg; move them down in every block's
  // cache, in order (row w < row r, so the copies never overlap).
  const size_t row = static_cast<size_t>(kv_heads_) * head_dim_;
#ifdef HAS_CUDA
  if (gpu_) {  // the same rows kept, in the same order, compacted in each block's device cache
    std::vector<int> idx;
    std::vector<int> kept;
    for (size_t r = 0; r < kv_pos_.size(); ++r) {
      const int p = kv_pos_[r];
      if (ts_sidx_.at(p) < 0 && ts_seg_.at(p) < seg) continue;
      idx.push_back(static_cast<int>(r));
      kept.push_back(p);
    }
    int w0 = 0;
    while (w0 < static_cast<int>(idx.size()) && idx[w0] == w0) ++w0;
    if (w0 < static_cast<int>(idx.size())) {
      // Pageable source: the copy has read idx when cudaMemcpyAsync returns.
      gpuErrchk(cudaMemcpyAsync(dkv_idx_, idx.data(), idx.size() * sizeof(int), cudaMemcpyHostToDevice, stream_));
      for (int l = 0; l < n_layers_; ++l)
        ir_kv_compact_cuda(dk_[l], dv_[l], dkv_idx_, w0, static_cast<int>(idx.size()), static_cast<int>(row), stream_);
    }
    kv_pos_ = std::move(kept);
    evicted_below_ = seg;
    return;
  }
#endif
  size_t w = 0;
  for (size_t r = 0; r < kv_pos_.size(); ++r) {
    const int p = kv_pos_[r];
    if (ts_sidx_.at(p) < 0 && ts_seg_.at(p) < seg) continue;
    if (w != r) {
      for (auto& layer : kv_) {
        std::memcpy(layer.k.data() + w * row, layer.k.data() + r * row, row * sizeof(float));
        std::memcpy(layer.v.data() + w * row, layer.v.data() + r * row, row * sizeof(float));
      }
      kv_pos_[w] = p;
    }
    ++w;
  }
  kv_pos_.resize(w);
  evicted_below_ = seg;
}

References ember_ir::ir_kv_compact_cuda(), j, and p.

◆ ts_seg_after()

int ember_ir::Interpreter::ts_seg_after ( int p ) const

private

Definition at line 2088 of file graph_ir.cpp.

                                         {
  if (p < 0) return 0;
  return ts_seg_.at(p) + (state_k_ > 0 && ts_sidx_.at(p) == state_k_ - 1 ? 1 : 0);
}

References p.

◆ ts_segment_start()

int ember_ir::Interpreter::ts_segment_start ( int m ) const

private

Definition at line 2093 of file graph_ir.cpp.

                                             {
  if (m <= 0) return 0;
  // The position after the last state id of block m - 1.
  for (int p = static_cast<int>(ts_seg_.size()) - 1; p >= 0; --p)
    if (ts_seg_[p] == m - 1 && ts_sidx_[p] == state_k_ - 1) return p + 1;
  return static_cast<int>(ts_seg_.size());
}

References j, and p.

◆ turn_state()

bool ember_ir::Interpreter::turn_state ( ) const

inline

Definition at line 168 of file graph_ir.h.

168{ return turn_state_; }

References turn_state_.

◆ vocab()

int ember_ir::Interpreter::vocab ( ) const

inline

Definition at line 135 of file graph_ir.h.

135{ return vocab_; }

References vocab_.

◆ weight()

const std::vector< float > & ember_ir::Interpreter::weight ( const std::string & name ) const

private

Definition at line 2351 of file graph_ir.cpp.

                                                                       {
  const Weight& w = bound(name);
  if (w.quantised()) throw std::runtime_error("IR weight '" + name + "' is held quantised but read whole");
  return w.f32;
}

References j, and name.

Member Data Documentation

◆ app_index_

bool ember_ir::Interpreter::app_index_ = false

private

Definition at line 255 of file graph_ir.h.

◆ attn_row_

int ember_ir::Interpreter::attn_row_ = -1

private

Definition at line 290 of file graph_ir.h.

◆ block_

std::vector<Node> ember_ir::Interpreter::block_

private

Definition at line 266 of file graph_ir.h.

Referenced by init().

◆ da_sources_

std::map<std::pair<const Node*, int>, int> ember_ir::Interpreter::da_sources_

private

Definition at line 269 of file graph_ir.h.

◆ dequant_all_

bool ember_ir::Interpreter::dequant_all_ = false

private

Definition at line 307 of file graph_ir.h.

◆ embed_

std::vector<Node> ember_ir::Interpreter::embed_

private

Definition at line 266 of file graph_ir.h.

Referenced by init().

◆ evict_

bool ember_ir::Interpreter::evict_ = false

private

Definition at line 302 of file graph_ir.h.

Referenced by evicts(), and init().

◆ evict_declared_

bool ember_ir::Interpreter::evict_declared_ = false

private

Definition at line 222 of file graph_ir.h.

Referenced by evict_declared(), and init().

◆ evicted_below_

int ember_ir::Interpreter::evicted_below_ = 0

private

Definition at line 305 of file graph_ir.h.

◆ family_hint_

std::string ember_ir::Interpreter::family_hint_

private

Definition at line 306 of file graph_ir.h.

Referenced by family_hint(), and init().

◆ gpu_

bool ember_ir::Interpreter::gpu_ = false

private

Definition at line 340 of file graph_ir.h.

Referenced by on_gpu().

◆ head_

std::vector<Node> ember_ir::Interpreter::head_

private

Definition at line 266 of file graph_ir.h.

Referenced by init().

◆ head_dim_

int ember_ir::Interpreter::head_dim_ = 0

private

Definition at line 311 of file graph_ir.h.

Referenced by init().

◆ ids_

std::vector<int> ember_ir::Interpreter::ids_

private

Definition at line 335 of file graph_ir.h.

◆ kv_

std::vector<KVCacheLayer> ember_ir::Interpreter::kv_

private

Definition at line 310 of file graph_ir.h.

◆ kv_fmt_

int ember_ir::Interpreter::kv_fmt_ = 0

private

Definition at line 312 of file graph_ir.h.

Referenced by init().

◆ kv_heads_

int ember_ir::Interpreter::kv_heads_ = 0

private

Definition at line 311 of file graph_ir.h.

Referenced by init().

◆ kv_pos_

std::vector<int> ember_ir::Interpreter::kv_pos_

private

Definition at line 303 of file graph_ir.h.

Referenced by kv_rows_held().

◆ kv_row0_

int ember_ir::Interpreter::kv_row0_ = 0

private

Definition at line 305 of file graph_ir.h.

◆ kv_slot_

std::string ember_ir::Interpreter::kv_slot_

private

Definition at line 221 of file graph_ir.h.

Referenced by init(), and kv_slot().

◆ max_len_

int ember_ir::Interpreter::max_len_ = 0

private

Definition at line 270 of file graph_ir.h.

Referenced by init(), and max_len().

◆ mor_depth_

std::vector<int> ember_ir::Interpreter::mor_depth_

private

Definition at line 288 of file graph_ir.h.

Referenced by last_depths().

◆ mor_g_

std::vector<float> ember_ir::Interpreter::mor_g_

private

Definition at line 289 of file graph_ir.h.

◆ mor_gate_score_

bool ember_ir::Interpreter::mor_gate_score_ = false

private

Definition at line 286 of file graph_ir.h.

Referenced by init().

◆ mor_out_

std::string ember_ir::Interpreter::mor_out_

private

Definition at line 287 of file graph_ir.h.

Referenced by init().

◆ mor_R_

int ember_ir::Interpreter::mor_R_ = 0

private

Definition at line 285 of file graph_ir.h.

Referenced by init().

◆ mor_router_

std::string ember_ir::Interpreter::mor_router_

private

Definition at line 287 of file graph_ir.h.

Referenced by init().

◆ mor_rows_

std::vector<std::vector<int> > ember_ir::Interpreter::mor_rows_

private

Definition at line 284 of file graph_ir.h.

Referenced by init().

◆ mor_step_

std::vector<int> ember_ir::Interpreter::mor_step_

private

Definition at line 283 of file graph_ir.h.

Referenced by has_depth_route(), and init().

◆ mor_x_

std::string ember_ir::Interpreter::mor_x_

private

Definition at line 287 of file graph_ir.h.

Referenced by init().

◆ mor_xin_

std::vector<float> ember_ir::Interpreter::mor_xin_

private

Definition at line 289 of file graph_ir.h.

◆ n_layers_

int ember_ir::Interpreter::n_layers_ = 0

private

Definition at line 270 of file graph_ir.h.

Referenced by init(), and n_layers().

◆ offset_applied_

std::set<std::string> ember_ir::Interpreter::offset_applied_

private

Definition at line 339 of file graph_ir.h.

◆ rec_

std::vector<float> ember_ir::Interpreter::rec_

private

Definition at line 364 of file graph_ir.h.

◆ rec_floats_

size_t ember_ir::Interpreter::rec_floats_ = 0

private

Definition at line 360 of file graph_ir.h.

Referenced by has_recurrent(), and init().

◆ rec_init_

std::vector<float> ember_ir::Interpreter::rec_init_

private

Definition at line 361 of file graph_ir.h.

◆ rec_max_snaps_

size_t ember_ir::Interpreter::rec_max_snaps_ = 0

private

Definition at line 367 of file graph_ir.h.

◆ rec_pos_

int ember_ir::Interpreter::rec_pos_ = 0

private

Definition at line 365 of file graph_ir.h.

◆ rec_rows_

std::vector<float> ember_ir::Interpreter::rec_rows_

private

Definition at line 369 of file graph_ir.h.

◆ rec_slots_

std::vector<RecSlot> ember_ir::Interpreter::rec_slots_

private

Definition at line 359 of file graph_ir.h.

Referenced by init().

◆ rec_snap_rows_

std::vector<int> ember_ir::Interpreter::rec_snap_rows_

private

Definition at line 368 of file graph_ir.h.

◆ rec_snaps_

std::map<int, std::vector<float> > ember_ir::Interpreter::rec_snaps_

private

Definition at line 366 of file graph_ir.h.

◆ recs_

std::map<const Node*, std::shared_ptr<rec::Recurrence> > ember_ir::Interpreter::recs_

private

Definition at line 363 of file graph_ir.h.

Referenced by init().

◆ rope_attn_layer_

std::vector<float> ember_ir::Interpreter::rope_attn_layer_

private

Definition at line 322 of file graph_ir.h.

Referenced by init().

◆ rope_dim_

int ember_ir::Interpreter::rope_dim_ = 0

private

Definition at line 315 of file graph_ir.h.

Referenced by init().

◆ rope_ff_

nlohmann::json ember_ir::Interpreter::rope_ff_

private

Definition at line 332 of file graph_ir.h.

Referenced by init().

◆ rope_ff_ctx_

int ember_ir::Interpreter::rope_ff_ctx_ = 0

private

Definition at line 333 of file graph_ir.h.

◆ rope_freq_scale_

float ember_ir::Interpreter::rope_freq_scale_ = 1.0f

private

Definition at line 317 of file graph_ir.h.

Referenced by init().

◆ rope_long_

bool ember_ir::Interpreter::rope_long_ = false

private

Definition at line 334 of file graph_ir.h.

◆ rope_neox_

bool ember_ir::Interpreter::rope_neox_ = false

private

Definition at line 316 of file graph_ir.h.

Referenced by init().

◆ rope_scale_layer_

std::vector<float> ember_ir::Interpreter::rope_scale_layer_

private

Definition at line 322 of file graph_ir.h.

Referenced by init().

◆ rope_tab_long_

std::vector<int> ember_ir::Interpreter::rope_tab_long_

private

Definition at line 327 of file graph_ir.h.

◆ rope_tab_short_

std::vector<int> ember_ir::Interpreter::rope_tab_short_

private

Definition at line 327 of file graph_ir.h.

◆ rope_tabs_

std::vector<std::vector<std::pair<float, float> > > ember_ir::Interpreter::rope_tabs_

private

Definition at line 326 of file graph_ir.h.

◆ rope_theta_

float ember_ir::Interpreter::rope_theta_ = 0.0f

private

Definition at line 314 of file graph_ir.h.

Referenced by init().

◆ rope_theta_layer_

std::vector<float> ember_ir::Interpreter::rope_theta_layer_

private

Definition at line 322 of file graph_ir.h.

Referenced by init().

◆ seq_len_

int ember_ir::Interpreter::seq_len_ = 0

private

Definition at line 223 of file graph_ir.h.

Referenced by seq_len().

◆ shapes_

std::unordered_map<std::string, std::vector<int64_t> > ember_ir::Interpreter::shapes_

private

Definition at line 309 of file graph_ir.h.

◆ sinks_

int ember_ir::Interpreter::sinks_ = 0

private

Definition at line 313 of file graph_ir.h.

Referenced by init().

◆ state_k_

int ember_ir::Interpreter::state_k_ = 0

private

Definition at line 300 of file graph_ir.h.

Referenced by state_k().

◆ state_weight_

std::string ember_ir::Interpreter::state_weight_

private

Definition at line 301 of file graph_ir.h.

Referenced by init().

◆ ts_seg_

std::vector<int> ember_ir::Interpreter::ts_seg_

private

Definition at line 303 of file graph_ir.h.

Referenced by kv_rows_held().

◆ ts_sidx_

std::vector<int> ember_ir::Interpreter::ts_sidx_

private

Definition at line 303 of file graph_ir.h.

◆ ts_vis_

std::vector<std::vector<int> > ember_ir::Interpreter::ts_vis_

private

Definition at line 304 of file graph_ir.h.

◆ turn_state_

bool ember_ir::Interpreter::turn_state_ = false

private

Definition at line 302 of file graph_ir.h.

Referenced by init(), kv_rows_held(), and turn_state().

◆ vocab_

int ember_ir::Interpreter::vocab_ = 0

private

Definition at line 270 of file graph_ir.h.

Referenced by init(), and vocab().

◆ weight_map_

std::vector<int> ember_ir::Interpreter::weight_map_

private

Definition at line 273 of file graph_ir.h.

Referenced by init().

◆ weights_

std::unordered_map<std::string, Weight> ember_ir::Interpreter::weights_

private

Definition at line 308 of file graph_ir.h.

◆ window_

int ember_ir::Interpreter::window_ = 0

private

Definition at line 313 of file graph_ir.h.

Referenced by init().

◆ window_layer_

std::vector<int> ember_ir::Interpreter::window_layer_

private

Definition at line 321 of file graph_ir.h.

Referenced by init().

The documentation for this class was generated from the following files:

Classes

Public Member Functions

Private Types

Private Member Functions

Private Attributes

Detailed Description

Member Typedef Documentation

◆ Env

Constructor & Destructor Documentation

◆ Interpreter() [1/2]

◆ ~Interpreter()

◆ Interpreter() [2/2]

Member Function Documentation

◆ attention_window()

◆ bind()

◆ bind_tensor()

◆ bound()

◆ build_rope_tables()

◆ evict_declared()

◆ evicts()

◆ family_hint()

◆ forward()

◆ forward_impl()

◆ has_depth_route()

◆ has_recurrent()

◆ held_positions()

◆ init()

◆ kv_cache_bytes()

◆ kv_format()

◆ kv_positions()

◆ kv_rows_held()

◆ kv_rows_in_use()

◆ kv_scratch_bytes()

◆ kv_slot()

◆ last_depths()

◆ layer_name()

◆ load_state()

◆ max_len()

◆ mor_step()

◆ n_layers()

◆ on_gpu()

◆ operator=()

◆ rec_begin()

◆ rec_end()

◆ rec_row_done()

◆ rec_slot()

◆ rec_tensor()

◆ recurrent_tensors()

◆ recurrent_values()

◆ resume_point()

◆ rope_table()

◆ run_nodes()

◆ save_state()

◆ seq_len()

◆ set_rope_long()

◆ state_k()

◆ truncate()

◆ ts_begin()

◆ ts_evict_below()

◆ ts_seg_after()

◆ ts_segment_start()

◆ turn_state()

◆ vocab()

◆ weight()

Member Data Documentation

◆ app_index_

◆ attn_row_

◆ block_

◆ da_sources_

◆ dequant_all_

◆ embed_

◆ evict_

◆ evict_declared_

◆ evicted_below_

◆ family_hint_

◆ gpu_

◆ head_

◆ head_dim_

◆ ids_

◆ kv_