ronantakizawa/codereview-bench huggingface dataset:
Repository: ggerganov/llama.cpp
File: ggml/src/ggml-cuda/mmf.cuh (CUDA)

Code before:
#pragma once

#include "mma.cuh"
#include "common.cuh"
#include "convert.cuh"

using namespace ggml_cuda_mma;

#define MMF_ROWS_PER_BLOCK 32
#define MMF_ROWS_PER_BLOCK_CDNA 64

static __forceinline__ int64_t get_max_block_size(int cc) {
    if (GGML_CUDA_CC_IS_CDNA(cc)) {
        return 512;
    } else {
        return 256;
    }
}

static __forceinline__ int get_padded_size(int cc) {
    if (GGML_CUDA_CC_IS_CDNA(cc)) {
        return 2;
    } else {
        return 4;
    }
}

static constexpr __device__ int get_padded_size() {
#if defined(AMD_MFMA_AVAILABLE)
    return 2;
#else
    return 4;
#endif // defined(AMD_MFMA_AVAILABLE)
}

struct mmf_ids_data {
    const int32_t * ids_src_compact = nullptr;
    const int32_t * ids_dst_compact = nullptr;
    const int32_t * expert_bounds_dev = nullptr;
    int n_experts = 0;
    int sis1 = 0;
};

void ggml_cuda_mul_mat_f(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * ids, ggml_tensor * dst);


Review comment (suggestion):
```suggestion
static __forceinline__ int mmf_get_padding(int cc) {
```

Code after:
#pragma once

#include "mma.cuh"
#include "common.cuh"
#include "convert.cuh"

using namespace ggml_cuda_mma;

#define MMF_ROWS_PER_BLOCK 32
#define MMF_ROWS_PER_BLOCK_CDNA 64

static __forceinline__ int64_t mmf_get_max_block_size(int cc) {
    if (GGML_CUDA_CC_IS_CDNA(cc)) {
        return 512;
    } else {
        return 256;
    }
}

static __forceinline__ int mmf_get_padding(int cc) {
    if (GGML_CUDA_CC_IS_CDNA(cc)) {
        return 2;
    } else {
        return 4;
    }
}

static constexpr __device__ int mmf_get_padding() {
#if defined(AMD_MFMA_AVAILABLE)
    return 2;
#else
    return 4;
#endif // defined(AMD_MFMA_AVAILABLE)
}

struct mmf_ids_data {
    const int32_t * ids_src_compact = nullptr;
    const int32_t * ids_dst_compact = nullptr;
    const int32_t * expert_bounds_dev = nullptr;
    int n_experts = 0;
    int sis1 = 0;
};

void ggml_cuda_mul_mat_f(ggml_backend_cuda_context & ctx, const ggml_tensor * src0, const ggml_tensor * src1, const ggml_tensor * ids, ggml_tensor * dst);
Repository: ggerganov/llama.cpp
File: common/ngram-mod.h (C/C++)

Code before:
#pragma once

#include <cstdint>
#include <vector>
#include <memory>

#define COMMON_NGRAM_MOD_MAX_CHOICES 4

struct common_ngram_mod_entry {
    uint16_t head = 0;
    uint16_t n_choices = 0;

    int32_t choices[COMMON_NGRAM_MOD_MAX_CHOICES];
};

// basic n-gram hasher
struct common_ngram_mod {
    common_ngram_mod(uint16_t n, uint64_t size);

    uint64_t idx(const int32_t * tokens) const;

    void    add(const int32_t * tokens);
    int32_t get(const int32_t * tokens, int32_t offs) const; // return -1 if not found

    uint16_t n; // ngram size to hash

    std::vector<common_ngram_mod_entry> entries;
};

using common_ngram_mod_ptr = std::unique_ptr<common_ngram_mod>;


Review comment (suggestion):
in multiple places in the code, we need to cast this to `size_t`, so I think it's probably better to use `size_t`

```suggestion
    size_t n; // ngram size to hash
```

Code after:
#pragma once

#include <cstdint>
#include <vector>

//
// common_ngram_mod
// ref: https://github.com/ggml-org/llama.cpp/pull/19164
//

// basic n-gram hasher
struct common_ngram_mod {
    using entry_t = int32_t;

    static constexpr entry_t EMPTY = -1;

    common_ngram_mod(uint16_t n, size_t size);

    size_t  idx(const entry_t * tokens) const;
    void    add(const entry_t * tokens);
    entry_t get(const entry_t * tokens) const; // return -1 if not found

    void reset();

    size_t get_n()    const;
    size_t get_used() const;

    size_t size()       const;
    size_t size_bytes() const;

private:
    size_t n; // ngram size to hash

    size_t used;

    std::vector<entry_t> entries;
};
Repository: ggerganov/llama.cpp
File: src/models/jais2.cpp (C++)

Code before:

        // Residual connection
        ggml_tensor * ffn_inp = ggml_add(ctx0, cur, inpL);
        cb(ffn_inp, "ffn_inp", il);

        // Pre-FFN LayerNorm
        cur = build_norm(ffn_inp,
                model.layers[il].ffn_norm,
                model.layers[il].ffn_norm_b,
                LLM_NORM, il);
        cb(cur, "ffn_norm", il);

        // FFN with relu2 activation (ReLU squared) - no gate projection
        // up -> relu2 -> down
        cur = build_ffn(cur,
                model.layers[il].ffn_up,   model.layers[il].ffn_up_b,   NULL,
                NULL, NULL, NULL,  // no gate
                model.layers[il].ffn_down, model.layers[il].ffn_down_b, NULL,
                NULL,
                LLM_FFN_RELU_SQR, LLM_FFN_SEQ, il);
        cb(cur, "ffn_out", il);

        // Residual connection
        inpL = ggml_add(ctx0, cur, ffn_inp);
        cb(inpL, "l_out", il);
    }

    // Final LayerNorm
    cur = build_norm(inpL,
            model.output_norm,
            model.output_norm_b,
            LLM_NORM, -1);
    cb(cur, "result_norm", -1);

    res->t_embd = cur;

    // Output projection
    cur = build_lora_mm(model.output, cur);
    cb(cur, "result_output", -1);

    res->t_logits = cur;

    ggml_build_forward_expand(gf, cur);
}


Review comment (suggestion):
```suggestion
        inpL = build_cvec(inpL, il);
        cb(inpL, "l_out", il);
```

Code after:

        // Residual connection
        ggml_tensor * ffn_inp = ggml_add(ctx0, cur, inpL);
        cb(ffn_inp, "ffn_inp", il);

        // Pre-FFN LayerNorm
        cur = build_norm(ffn_inp,
                model.layers[il].ffn_norm,
                model.layers[il].ffn_norm_b,
                LLM_NORM, il);
        cb(cur, "ffn_norm", il);

        // FFN with relu2 activation (ReLU squared) - no gate projection
        // up -> relu2 -> down
        cur = build_ffn(cur,
                model.layers[il].ffn_up,   model.layers[il].ffn_up_b,   NULL,
                NULL, NULL, NULL,  // no gate
                model.layers[il].ffn_down, model.layers[il].ffn_down_b, NULL,
                NULL,
                LLM_FFN_RELU_SQR, LLM_FFN_SEQ, il);
        cb(cur, "ffn_out", il);

        // Residual connection
        inpL = ggml_add(ctx0, cur, ffn_inp);
        inpL = build_cvec(inpL, il);
        cb(inpL, "l_out", il);
    }

    // Final LayerNorm
    cur = build_norm(inpL,
            model.output_norm,
            model.output_norm_b,
            LLM_NORM, -1);
    cb(cur, "result_norm", -1);

    res->t_embd = cur;

    // Output projection
    cur = build_lora_mm(model.output, cur);
    cb(cur, "result_output", -1);

    res->t_logits = cur;

    ggml_build_forward_expand(gf, cur);
}
Repository: ggerganov/llama.cpp
File: tests/test-jinja.cpp (C++)

Code before:
        "None"
    );
}

static void test_literals(testing & t) {
    test_template(t, "integer",
        "{{ 42 }}",
        json::object(),
        "42"
    );

    test_template(t, "float",
        "{{ 3.14 }}",
        json::object(),
        "3.14"
    );

    test_template(t, "string",
        "{{ 'hello' }}",
        json::object(),
        "hello"
    );

    test_template(t, "boolean true",
        "{{ true }}",
        json::object(),
        "True"
    );

    test_template(t, "boolean false",
        "{{ false }}",
        json::object(),
        "False"
    );

    test_template(t, "none",
        "{% if x is none %}null{% endif %}",
        {{"x", nullptr}},
        "null"
    );

    test_template(t, "list literal",
        "{% for i in [1, 2, 3] %}{{ i }}{% endfor %}",
        json::object(),
        "123"
    );

    test_template(t, "dict literal",
        "{% set d = {'a': 1} %}{{ d.a }}",
        json::object(),
        "1"

Review comment (suggestion):
```suggestion

    test_template(t, "indent",
        "{{ data|indent }}",
        { "data", "foo\n\nbar" },
        "foo\n\n    bar"
    );

    test_template(t, "indent first",
        "{{ data|indent(2, first=true) }}",
        { "data", "\n\nbar" },
        "  \n\n  bar"
    );

    test_template(t, "indent blank",
        "{{ data|indent('  ', blank=true) }}",
        { "data", "foo\n\nbar" },
        "foo\n  \n  bar"
    );
```

Code after:
        "     foo\n     \n     bar"
    );

    test_template(t, "indent with default width",
        "{{ data|indent() }}",
        {{ "data", "foo\nbar" }},
        "foo\n    bar"
    );

    test_template(t, "indent with no newline",
        "{{ data|indent }}",
        {{ "data", "foo" }},
        "foo"
    );

    test_template(t, "indent with trailing newline",
        "{{ data|indent(blank=true) }}",
        {{ "data", "foo\n" }},
        "foo\n    "
    );

    test_template(t, "indent with string",
        "{{ data|indent(width='>>>>') }}",
        {{ "data", "foo\nbar" }},
        "foo\n>>>>bar"
    );

    test_template(t, "chained filters",
        "{{ '  HELLO  '|trim|lower }}",
        json::object(),
        "hello"
    );

    test_template(t, "none to string",
        "{{ x|string }}",
        {{"x", nullptr}},
        "None"
    );
}

static void test_literals(testing & t) {
    test_template(t, "integer",
        "{{ 42 }}",
        json::object(),
        "42"
    );

    test_template(t, "float",
        "{{ 3.14 }}",
        json::object(),
        "3.14"
Repository: ggerganov/llama.cpp
File: tests/test-jinja.cpp (C++)

Code before:
        "{\n  \"a\": 1,\n  \"b\": [\n    1,\n    2\n  ]\n}"
    );

    test_template(t, "indent",
        "{{ data|indent(2) }}",
        {{ "data", "foo\nbar" }},
        "foo\n  bar"
    );

    test_template(t, "indent first only",
        "{{ data|indent(width=3,first=true) }}",
        {{ "data", "foo\nbar" }},
        "   foo\n   bar"
    );

    test_template(t, "indent blank lines and first line",
        "{{ data|indent(width=5,blank=true,first=true) }}",
        {{ "data", "foo\n\nbar" }},
        "     foo\n     \n     bar"
    );

    test_template(t, "indent with default width",
        "{{ data|indent() }}",
        {{ "data", "foo\nbar" }},
        "foo\n    bar"
    );

    test_template(t, "indent with string",
        "{{ data|indent(width='>>>>') }}",
        {{ "data", "foo\nbar" }},
        "foo\n>>>>bar"
    );

    test_template(t, "chained filters",
        "{{ '  HELLO  '|trim|lower }}",
        json::object(),
        "hello"
    );

    test_template(t, "none to string",
        "{{ x|string }}",
        {{"x", nullptr}},
        "None"
    );
}

static void test_literals(testing & t) {
    test_template(t, "integer",
        "{{ 42 }}",
        json::object(),
        "42"

Review comment (suggestion):
```suggestion

    test_template(t, "indent with default width",
        "{{ data|indent() }}",
        {{ "data", "foo\nbar" }},
        "foo\n    bar"
    );

    test_template(t, "indent with no newline",
        "{{ data|indent }}",
        {{ "data", "foo" }},
        "foo"
    );

    test_template(t, "indent with trailing newline",
        "{{ data|indent(blank=true) }}",
        {{ "data", "foo\n" }},
        "foo\n    "
    );
```

Code after:
        "{\n  \"a\": 1,\n  \"b\": [\n    1,\n    2\n  ]\n}"
    );

    test_template(t, "indent",
        "{{ data|indent(2) }}",
        {{ "data", "foo\nbar" }},
        "foo\n  bar"
    );

    test_template(t, "indent first only",
        "{{ data|indent(width=3,first=true) }}",
        {{ "data", "foo\nbar" }},
        "   foo\n   bar"
    );

    test_template(t, "indent blank lines and first line",
        "{{ data|indent(width=5,blank=true,first=true) }}",
        {{ "data", "foo\n\nbar" }},
        "     foo\n     \n     bar"
    );

    test_template(t, "indent with default width",
        "{{ data|indent() }}",
        {{ "data", "foo\nbar" }},
        "foo\n    bar"
    );

    test_template(t, "indent with no newline",
        "{{ data|indent }}",
        {{ "data", "foo" }},
        "foo"
    );

    test_template(t, "indent with trailing newline",
        "{{ data|indent(blank=true) }}",
        {{ "data", "foo\n" }},
        "foo\n    "
    );

    test_template(t, "indent with string",
        "{{ data|indent(width='>>>>') }}",
        {{ "data", "foo\nbar" }},
        "foo\n>>>>bar"
    );

    test_template(t, "chained filters",
        "{{ '  HELLO  '|trim|lower }}",
        json::object(),
        "hello"
    );
Repository: ggerganov/llama.cpp
File: tools/mtmd/models/nemotron-v2-vl.cpp (C++)

Code before:
#include "models.h"

ggml_cgraph * clip_graph_nemotron_v2_vl::build() {
    GGML_ASSERT(model.class_embedding != nullptr);
    GGML_ASSERT(model.position_embeddings != nullptr);

    const int n_registers = model.class_embedding->ne[1];
    const int n_pos = n_patches + n_registers;

    ggml_tensor * inp = build_inp();

    ggml_tensor * pos_embd = resize_position_embeddings();
    inp = ggml_add(ctx0, inp, pos_embd);
    cb(inp, "inp_pos", -1);

    inp = ggml_concat(ctx0, model.class_embedding, inp, 1);

    ggml_tensor * cur = build_vit(inp, n_pos, NORM_TYPE_NORMAL, hparams.ffn_op, nullptr, nullptr);

    cur = ggml_view_2d(ctx0, cur,
        n_embd, n_patches,
        ggml_row_size(cur->type, n_embd),
        n_registers * ggml_row_size(cur->type, n_embd));

    cur = build_patch_merge_permute(cur, model.hparams.n_merge);

    {
        cur = build_norm(cur, model.mm_0_w, nullptr, NORM_TYPE_RMS, 1e-6, -1);
        cur = build_ffn(cur, model.mm_1_w, nullptr, nullptr, nullptr, model.mm_3_w, nullptr, FFN_RELU_SQR, -1);
    }

    ggml_build_forward_expand(gf, cur);

    return gf;
}


Review comment (suggestion):
this seems to be redundant, the model doesn't support dynamic resolution if I understand correctly

`resize_position_embeddings` is only used by model with dynamic resolution support

Code after:
#include "models.h"

ggml_cgraph * clip_graph_nemotron_v2_vl::build() {
    GGML_ASSERT(model.class_embedding != nullptr);
    GGML_ASSERT(model.position_embeddings != nullptr);

    const int n_registers = model.class_embedding->ne[1];
    const int n_pos = n_patches + n_registers;

    ggml_tensor * inp = build_inp();

    // add position embeddings (pre-downsampled during GGUF conversion for fixed 512x512 input)
    inp = ggml_add(ctx0, inp, model.position_embeddings);
    cb(inp, "inp_pos", -1);

    inp = ggml_concat(ctx0, model.class_embedding, inp, 1);

    ggml_tensor * cur = build_vit(inp, n_pos, NORM_TYPE_NORMAL, hparams.ffn_op, nullptr, nullptr);

    cur = ggml_view_2d(ctx0, cur,
        n_embd, n_patches,
        ggml_row_size(cur->type, n_embd),
        n_registers * ggml_row_size(cur->type, n_embd));

    cur = build_patch_merge_permute(cur, model.hparams.n_merge);

    {
        cur = build_norm(cur, model.mm_0_w, nullptr, NORM_TYPE_RMS, 1e-6, -1);
        cur = build_ffn(cur, model.mm_1_w, nullptr, nullptr, nullptr, model.mm_3_w, nullptr, FFN_RELU_SQR, -1);
    }

    ggml_build_forward_expand(gf, cur);

    return gf;
}
Repository: ggerganov/llama.cpp
File: scripts/sync_vendor.py (Python)

Code before:
    # "https://github.com/mackron/miniaudio/raw/refs/tags/0.11.23/miniaudio.h": "vendor/miniaudio/miniaudio.h",
    "https://github.com/mackron/miniaudio/raw/669ed3e844524fcd883231b13095baee9f6de304/miniaudio.h": "vendor/miniaudio/miniaudio.h",

    f"https://raw.githubusercontent.com/yhirose/cpp-httplib/{HTTPLIB_VERSION}/httplib.h": "httplib.h",
    f"https://raw.githubusercontent.com/yhirose/cpp-httplib/{HTTPLIB_VERSION}/split.py":  "split.py",
    f"https://raw.githubusercontent.com/yhirose/cpp-httplib/{HTTPLIB_VERSION}/LICENSE":   "vendor/cpp-httplib/LICENSE",

    "https://raw.githubusercontent.com/sheredom/subprocess.h/b49c56e9fe214488493021017bf3954b91c7c1f5/subprocess.h": "vendor/sheredom/subprocess.h",
}

for url, filename in vendor.items():
    print(f"downloading {url} to {filename}") # noqa: NP100
    urllib.request.urlretrieve(url, filename)

print("Splitting httplib.h...")
try:
    subprocess.check_call([
        sys.executable, "split.py",
        "--extension", "cpp",
        "--out", "vendor/cpp-httplib"
    ])
    os.remove("split.py")
    os.remove("httplib.h")
except subprocess.CalledProcessError as e:
    print(f"Error running split.py: {e}")
    sys.exit(1)


Review comment (suggestion):
Perhaps  this could be moved to a finally block to always be cleaned up:
```python:
finally:
    os.remove("split.py")
    os.remove("httplib.h")
```


Code after:
    # "https://github.com/mackron/miniaudio/raw/refs/tags/0.11.23/miniaudio.h": "vendor/miniaudio/miniaudio.h",
    "https://github.com/mackron/miniaudio/raw/669ed3e844524fcd883231b13095baee9f6de304/miniaudio.h": "vendor/miniaudio/miniaudio.h",

    f"https://raw.githubusercontent.com/yhirose/cpp-httplib/{HTTPLIB_VERSION}/httplib.h": "httplib.h",
    f"https://raw.githubusercontent.com/yhirose/cpp-httplib/{HTTPLIB_VERSION}/split.py":  "split.py",
    f"https://raw.githubusercontent.com/yhirose/cpp-httplib/{HTTPLIB_VERSION}/LICENSE":   "vendor/cpp-httplib/LICENSE",

    "https://raw.githubusercontent.com/sheredom/subprocess.h/b49c56e9fe214488493021017bf3954b91c7c1f5/subprocess.h": "vendor/sheredom/subprocess.h",
}

for url, filename in vendor.items():
    print(f"downloading {url} to {filename}") # noqa: NP100
    urllib.request.urlretrieve(url, filename)

print("Splitting httplib.h...") # noqa: NP100
try:
    subprocess.check_call([
        sys.executable, "split.py",
        "--extension", "cpp",
        "--out", "vendor/cpp-httplib"
    ])
except Exception as e:
    print(f"Error: {e}") # noqa: NP100
    sys.exit(1)
finally:
    os.remove("split.py")
    os.remove("httplib.h")
Repository: ggerganov/llama.cpp
File: tools/server/server-models.h (C/C++)

Code before:
    bool is_active() const {
        return status == SERVER_MODEL_STATUS_LOADED || status == SERVER_MODEL_STATUS_LOADING;
    }

    bool is_failed() const {
        return status == SERVER_MODEL_STATUS_UNLOADED && exit_code != 0;
    }

    void update_args(common_preset_context & ctx_presets, std::string bin_path);
};

struct subprocess_s;

struct server_models {
private:
    struct instance_t {
        std::shared_ptr<subprocess_s> subproc; // shared between main thread and monitoring thread
        std::thread th;
        server_model_meta meta;
        FILE * stdin_file = nullptr;
    };

    std::mutex mutex;
    std::condition_variable cv;
    std::map<std::string, instance_t> mapping;
    std::map<std::string, std::string> name_index; // alias/name -> canonical name

    // for stopping models
    std::condition_variable cv_stop;
    std::set<std::string> stopping_models;

    common_preset_context ctx_preset;

    common_params base_params;
    std::string bin_path;
    std::vector<std::string> base_env;
    common_preset base_preset; // base preset from llama-server CLI args

    void update_meta(const std::string & name, const server_model_meta & meta);

    // unload least recently used models if the limit is reached
    void unload_lru();

    // not thread-safe, caller must hold mutex
    void add_model(server_model_meta && meta);

public:
    server_models(const common_params & params, int argc, char ** argv);

    void load_models();


Review comment (suggestion):
I would prefer keeping just one flat map `std::map<std::string, instance_t>` for simplicity.

Also, we should throw an error or a warning if user accidentally define 2 models with the same alias

Code after:

    bool is_active() const {
        return status == SERVER_MODEL_STATUS_LOADED || status == SERVER_MODEL_STATUS_LOADING;
    }

    bool is_failed() const {
        return status == SERVER_MODEL_STATUS_UNLOADED && exit_code != 0;
    }

    void update_args(common_preset_context & ctx_presets, std::string bin_path);
};

struct subprocess_s;

struct server_models {
private:
    struct instance_t {
        std::shared_ptr<subprocess_s> subproc; // shared between main thread and monitoring thread
        std::thread th;
        server_model_meta meta;
        FILE * stdin_file = nullptr;
    };

    std::mutex mutex;
    std::condition_variable cv;
    std::map<std::string, instance_t> mapping;

    // for stopping models
    std::condition_variable cv_stop;
    std::set<std::string> stopping_models;

    common_preset_context ctx_preset;

    common_params base_params;
    std::string bin_path;
    std::vector<std::string> base_env;
    common_preset base_preset; // base preset from llama-server CLI args

    void update_meta(const std::string & name, const server_model_meta & meta);

    // unload least recently used models if the limit is reached
    void unload_lru();

    // not thread-safe, caller must hold mutex
    void add_model(server_model_meta && meta);

public:
    server_models(const common_params & params, int argc, char ** argv);

    void load_models();
Repository: ggerganov/whisper.cpp
File: whisper.cpp (C++)

Code before:
    whisper_vocab vocab;

    whisper_mel mel;

    std::vector<float> probs;
    std::vector<float> logits;

    std::vector<whisper_segment> result_all;

    std::vector<whisper_token> prompt_past;

    // [EXPERIMENTAL] token-level timestamps data
    int64_t t_beg;
    int64_t t_last;
    whisper_token tid_last;
    std::vector<float> energy; // PCM signal energy

    // [EXPERIMENTAL] speed-up techniques
    int32_t exp_n_audio_ctx; // 0 - use default
};

template<typename T>
static void read_safe(std::ifstream& fin, T& dest) {
    fin.read((char*)& dest, sizeof(T));
#if defined(GGML_BIG_ENDIAN)
    if constexpr (std::endian::native == std::endian::big) {
        dest = std::byteswap(dest);
    }
#endif
}

// load the model from a ggml file
//
// file format:
//
//   - hparams
//   - pre-computed mel filters
//   - vocab
//   - weights
//
// see the convert-pt-to-ggml.py script for details
//
static bool whisper_model_load(const std::string & fname, whisper_context & wctx) {
    fprintf(stderr, "%s: loading model from '%s'\n", __func__, fname.c_str());

    auto & model = wctx.model;
    auto & vocab = wctx.vocab;

    auto fin = std::ifstream(fname, std::ios::binary);
    if (!fin) {
        fprintf(stderr, "%s: failed to open '%s'\n", __func__, fname.c_str());

Review comment (bug):
I realise the reason for the guard is because `<bit>` is guarded by it, but this seems wrong; if that's defined, it will always be BE.

What you probably want is something like `#if __cplusplus >= 202110L` (the revision of C++2x that added `std::byteswap`).

Code after:
    struct ggml_tensor * cross_attn_k_w;

    // decoder.blocks.*.cross_attn.value
    struct ggml_tensor * cross_attn_v_w;
    struct ggml_tensor * cross_attn_v_b;

    // decoder.blocks.*.mlp_ln
    struct ggml_tensor * mlp_ln_w;
    struct ggml_tensor * mlp_ln_b;

    // decoder.blocks.*.mlp.0
    struct ggml_tensor * mlp_0_w;
    struct ggml_tensor * mlp_0_b;

    // decoder.blocks.*.mlp.2
    struct ggml_tensor * mlp_1_w;
    struct ggml_tensor * mlp_1_b;
};

struct whisper_kv_cache {
    struct ggml_tensor * k;
    struct ggml_tensor * v;

    struct ggml_context * ctx;

    std::vector<uint8_t> buf;

    int n; // number of tokens currently in the cache
};

struct whisper_model {
    e_model type = MODEL_UNKNOWN;

    whisper_hparams hparams;
    whisper_filters filters;

    // encoder.positional_embedding
    struct ggml_tensor * e_pe;

    // encoder.conv1
    struct ggml_tensor * e_conv_1_w;
    struct ggml_tensor * e_conv_1_b;

    // encoder.conv2
    struct ggml_tensor * e_conv_2_w;
    struct ggml_tensor * e_conv_2_b;

    // encoder.ln_post
    struct ggml_tensor * e_ln_w;
    struct ggml_tensor * e_ln_b;