ronantakizawa/codereview-bench huggingface dataset:
Repository: ggerganov/llama.cpp
File: src/llama-kv-cells.h (C/C++)

Code before:
#pragma once

#include "llama.h"
#include "llama-cparams.h"

#include <bitset>
#include <cassert>
#include <vector>
#include <set>
#include <map>

struct llama_kv_pos_mrope {
    llama_pos y = 0;
    llama_pos x = 0;
    // return true if this position is greater than the other position
    bool is_gt(const llama_kv_pos_mrope & other) const {
        return (y > other.y) || (y == other.y && x > other.x);
    }
};

// meta information about KV cells that can be part of multiple sequences at the same time
// TODO: add unit tests
class llama_kv_cells {
public:
    void reset() {
        for (uint32_t i = 0; i < pos.size(); ++i) {
            pos[i]   = -1;
            shift[i] =  0;
            seq[i].reset();
        }

        has_shift = false;

        used.clear();

        for (uint32_t s = 0; s < LLAMA_MAX_SEQ; ++s) {
            seq_pos[s].clear();
        }
    }

    void reset_shift() {
        has_shift = false;

        for (uint32_t i = 0; i < shift.size(); ++i) {
            shift[i] = 0;

Review comment (suggestion):
Similarly, I think we can decouple the concept of M-RoPE here by declaring this struct to be more generic:

```c++
struct llama_kv_cell_ext {
    // 2D spatial positions, typically used for M-RoPE
    llama_pos x = 0;
    llama_pos y = 0;

    // ... maybe more data in the future
};
```

Code after:
#pragma once

#include "llama.h"
#include "llama-cparams.h"

#include <bitset>
#include <cassert>
#include <cstring>
#include <map>
#include <set>
#include <vector>

struct llama_kv_cell_ext {
    // 2D spatial positions, typically used for M-RoPE
    llama_pos x = 0;
    llama_pos y = 0;

    // return true if the current 2D spatial position is greater than other
    bool is_2d_gt(llama_pos ox, llama_pos oy) const {
        return (y > oy) || (y == oy && x > ox);
    }

    void reset() {
        static_assert(std::is_trivially_copyable_v<llama_kv_cell_ext>);

        memset(this, 0, sizeof(*this));
    }
};

// meta information about KV cells that can be part of multiple sequences at the same time
// TODO: add unit tests
class llama_kv_cells {
public:
    void reset() {
        for (uint32_t i = 0; i < pos.size(); ++i) {
            pos[i]   = -1;
            ext[i].reset();
            shift[i] =  0;
            seq[i].reset();
        }

        has_shift = false;

        used.clear();
Repository: ggerganov/llama.cpp
File: src/llama-kv-cells.h (C/C++)

Code before:
        assert(i < pos.size());
        assert(pos[i] != -1);

        const llama_pos p_old = pos[i];

        seq_pos_rm(i);

        pos[i]   /= d;
        shift[i] += p_old - pos[i];

        seq_pos_add(i);

        has_shift = true;
    }

private:
    bool has_shift = false;

    // set of indices of used cells (i.e. pos[i] != -1, allowed to not have any seq_id)
    std::set<uint32_t> used;

    std::vector<llama_pos> pos;

    // stores addition info for M-RoPE positions
    std::vector<llama_kv_pos_mrope> pos_mrope;

    // this array accumulates any applied shifts to the pos array since the last reset_shift() call
    // this is used to queue multiple updates to the pos array, which in the end can be applied in one go:
    //
    //   cells.pos_add(x, shift_x);
    //   cells.pos_div(y, shift_y);
    //   ...
    //
    //   if (cells.has_shift()) {
    //      for (int i = 0; i < n; ++i) {
    //          auto shift_i = cells.get_shift(i);
    //          ...
    //      }
    //      cells.reset_shift();
    //   }
    //
    std::vector<llama_pos> shift;

    using seq_set_t = std::bitset<LLAMA_MAX_SEQ>;

    // the bitset seq[i] tells us which sequences are currently occupying the i-th cell
    std::vector<seq_set_t> seq;

    // the set seq_pos[s][p] tells us how many times the position p is currently present for sequence s
    // if the position p is not present, seq_pos[s][p] is not set
    // this way seq_pos[s].begin() and seq_pos[s].rbegin() give us the min/max positions currently in the cache

Review comment (suggestion):
```suggestion
    // stores extra optional cell info
    std::vector<llama_kv_cell_ext> ext;

```

Code after:

            used.erase(i);

            return true;
        }

        seq_pos_add(i);

        return false;
    }

    // pos[i] = pos[i] / d
    // sets "has_shift" to true
    // note: call only if the cell is not empty
    void pos_div(uint32_t i, int d) {
        assert(i < pos.size());
        assert(pos[i] != -1);

        const llama_pos p_old = pos[i];

        seq_pos_rm(i);

        pos[i]   /= d;
        shift[i] += p_old - pos[i];

        seq_pos_add(i);

        has_shift = true;
    }

private:
    bool has_shift = false;

    // set of indices of used cells (i.e. pos[i] != -1, allowed to not have any seq_id)
    std::set<uint32_t> used;

    std::vector<llama_pos> pos;

    // stores extra info per cell
    std::vector<llama_kv_cell_ext> ext;

    // this array accumulates any applied shifts to the pos array since the last reset_shift() call
    // this is used to queue multiple updates to the pos array, which in the end can be applied in one go:
    //
    //   cells.pos_add(x, shift_x);
    //   cells.pos_div(y, shift_y);
    //   ...
    //
    //   if (cells.has_shift()) {
    //      for (int i = 0; i < n; ++i) {
    //          auto shift_i = cells.get_shift(i);
Repository: ggerganov/llama.cpp
File: ggml/src/ggml-cuda/cpy.cu (CUDA)

Code before:
#include <algorithm>
#include "cpy.cuh"
#include "dequantize.cuh"
#include "cpy-utils.cuh"
#if defined(GGML_USE_MUSA) && defined(GGML_MUSA_MUDNN_COPY)
#include "ggml-musa/mudnn.cuh"
#endif // GGML_USE_MUSA && GGML_MUSA_MUDNN_COPY

typedef void (*cpy_kernel_t)(const char * cx, char * cdst);

const int CUDA_CPY_TILE_DIM_2D = 32; // 2D tile dimension for transposed blocks
const int CUDA_CPY_BLOCK_NM = 8;     // block size of 3rd dimension if available
const int CUDA_CPY_BLOCK_ROWS = 8;   // block dimension for marching through rows

template <cpy_kernel_t cpy_1>
static __global__ void cpy_flt(const char * cx, char * cdst, const int ne,
                               const int ne00, const int ne01, const int ne02, const int nb00, const int nb01, const int nb02,
                               const int nb03, const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
                               const int nb12, const int nb13) {
    const int64_t i = blockDim.x*blockIdx.x + threadIdx.x;

    if (i >= ne) {
        return;
    }

    // determine indices i03/i13, i02/i12, i01/i11, i00/i10 as a function of index i of flattened tensor

Review comment (suggestion):
```suggestion
```

I think this header is no longer needed.

Code after:
#include "cpy.cuh"
#include "dequantize.cuh"
#include "cpy-utils.cuh"
#if defined(GGML_USE_MUSA) && defined(GGML_MUSA_MUDNN_COPY)
#include "ggml-musa/mudnn.cuh"
#endif // GGML_USE_MUSA && GGML_MUSA_MUDNN_COPY

typedef void (*cpy_kernel_t)(const char * cx, char * cdst);

const int CUDA_CPY_TILE_DIM_2D = 32; // 2D tile dimension for transposed blocks
const int CUDA_CPY_BLOCK_NM = 8;     // block size of 3rd dimension if available
const int CUDA_CPY_BLOCK_ROWS = 8;   // block dimension for marching through rows

template <cpy_kernel_t cpy_1>
static __global__ void cpy_flt(const char * cx, char * cdst, const int ne,
                               const int ne00, const int ne01, const int ne02, const int nb00, const int nb01, const int nb02,
                               const int nb03, const int ne10, const int ne11, const int ne12, const int nb10, const int nb11,
                               const int nb12, const int nb13) {
    const int64_t i = blockDim.x*blockIdx.x + threadIdx.x;

    if (i >= ne) {
        return;
    }

    // determine indices i03/i13, i02/i12, i01/i11, i00/i10 as a function of index i of flattened tensor
    // then combine those indices with the corresponding byte offsets to get the total offsets
Repository: ggerganov/llama.cpp
File: ggml/src/ggml-cuda/moe-expert-reduce.cu (CUDA)

Code before:
#include "moe-expert-reduce.cuh"

// This kernel is fusion of the expert weight reduce, common in MoE models

template <int n_expert_used_template>
__global__ void moe_expert_reduce_cuda(const float * __restrict__ experts,
                                       const float * __restrict__ weights,
                                       float * __restrict__ dst,
                                       const int n_expert_used,
                                       const int n_cols) {
    const int row             = blockIdx.x;
    const int n_expert_used_t = n_expert_used_template == 0 ? n_expert_used : n_expert_used_template;

    const int col = blockIdx.y * blockDim.x + threadIdx.x;
    if (col >= n_cols) {
        return;
    }

    experts += row * n_cols * n_expert_used;
    weights += row * n_expert_used;
    dst += row * n_cols;

    float acc = 0.f;
    if constexpr (n_expert_used_template == 0) {
        for (int expert = 0; expert < n_expert_used_t; ++expert) {
            ggml_cuda_mad(acc, experts[col], weights[expert]);
            experts += n_cols;
        }

Review comment (suggestion):
```suggestion
// This kernel is a fusion of the expert weight reduce, common in MoE models.
```

Code after:
#include "moe-expert-reduce.cuh"

// This kernel is a fusion of the expert weight reduce, common in MoE models

template <int n_expert_used_template>
__global__ void moe_expert_reduce_cuda(const float * __restrict__ experts,
                                       const float * __restrict__ weights,
                                       float * __restrict__ dst,
                                       const int n_expert_used,
                                       const int n_cols) {
    const int row = blockIdx.x;
    const int col = blockIdx.y * blockDim.x + threadIdx.x;
    if (col >= n_cols) {
        return;
    }

    experts += row * n_cols * n_expert_used;
    weights += row * n_expert_used;
    dst += row * n_cols;

    float acc = 0.f;
    if constexpr (n_expert_used_template == 0) {
        for (int expert = 0; expert < n_expert_used; ++expert) {
            ggml_cuda_mad(acc, experts[col], weights[expert]);
            experts += n_cols;
        }
        dst[col] = acc;
    } else {
Repository: ggerganov/llama.cpp
File: ggml/src/ggml-cuda/moe-expert-reduce.cu (CUDA)

Code before:
#include "moe-expert-reduce.cuh"

// This kernel is fusion of the expert weight reduce, common in MoE models

template <int n_expert_used_template>
__global__ void moe_expert_reduce_cuda(const float * __restrict__ experts,
                                       const float * __restrict__ weights,
                                       float * __restrict__ dst,
                                       const int n_expert_used,
                                       const int n_cols) {
    const int row             = blockIdx.x;
    const int n_expert_used_t = n_expert_used_template == 0 ? n_expert_used : n_expert_used_template;

    const int col = blockIdx.y * blockDim.x + threadIdx.x;
    if (col >= n_cols) {
        return;
    }

    experts += row * n_cols * n_expert_used;
    weights += row * n_expert_used;
    dst += row * n_cols;

    float acc = 0.f;
    if constexpr (n_expert_used_template == 0) {
        for (int expert = 0; expert < n_expert_used_t; ++expert) {
            ggml_cuda_mad(acc, experts[col], weights[expert]);
            experts += n_cols;
        }
        dst[col] = acc;
    } else {
#pragma unroll
        for (int i = 0; i < n_expert_used_template; ++i) {
            ggml_cuda_mad(acc, experts[col], weights[i]);
            experts += n_cols;
        }
        dst[col] = acc;
    }

Review comment (suggestion):
I would suggest not using the `_t` suffix for variable names since that is how we designate types.

Code after:
#include "moe-expert-reduce.cuh"

// This kernel is a fusion of the expert weight reduce, common in MoE models

template <int n_expert_used_template>
__global__ void moe_expert_reduce_cuda(const float * __restrict__ experts,
                                       const float * __restrict__ weights,
                                       float * __restrict__ dst,
                                       const int n_expert_used,
                                       const int n_cols) {
    const int row = blockIdx.x;
    const int col = blockIdx.y * blockDim.x + threadIdx.x;
    if (col >= n_cols) {
        return;
    }

    experts += row * n_cols * n_expert_used;
    weights += row * n_expert_used;
    dst += row * n_cols;

    float acc = 0.f;
    if constexpr (n_expert_used_template == 0) {
        for (int expert = 0; expert < n_expert_used; ++expert) {
            ggml_cuda_mad(acc, experts[col], weights[expert]);
            experts += n_cols;
        }
        dst[col] = acc;
    } else {
#pragma unroll
        for (int i = 0; i < n_expert_used_template; ++i) {
            ggml_cuda_mad(acc, experts[col], weights[i]);
            experts += n_cols;
        }
        dst[col] = acc;
    }
}
Repository: ggerganov/llama.cpp
File: ggml/src/ggml-cuda/convert.cuh (CUDA)

Code before:

typedef to_t_nc_cuda_t<float> to_fp32_nc_cuda_t;
typedef to_t_nc_cuda_t<half> to_fp16_nc_cuda_t;
typedef to_t_nc_cuda_t<nv_bfloat16> to_bf16_nc_cuda_t;

to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type);
to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type);
to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type);

template<typename dst_t, typename src_t>
 __host__ __device__ inline dst_t ggml_cuda_cast(src_t x) {
    if constexpr (std::is_same_v<dst_t, src_t>) {
        return x;
    } else if constexpr(std::is_same_v<dst_t, nv_bfloat16>) {
        return __float2bfloat16(float(x));
    } else if constexpr(std::is_same_v<src_t, nv_bfloat16>) {
        return __bfloat162float(x);
    } else if constexpr(std::is_same_v<src_t, float2> && std::is_same_v<dst_t, half2>) {
        return __float22half2_rn(x);
    } else if constexpr(std::is_same_v<src_t, float2> && std::is_same_v<dst_t, nv_bfloat162>) {
        // bypass compile error on cuda 12.0.1
#if defined(__CUDACC__)
        return {x.x, x.y};
#else
        return __float22bfloat162_rn(x);
#endif
    } else if constexpr(std::is_same_v<dst_t, int32_t>) {
        return int32_t(x);
    } else {
        return float(x);
    }
}


Review comment (suggestion):
```suggestion
#ifdef GGML_USE_HIP
        return __float22bfloat162_rn(x);
#else
        return {x.x, x.y};
#endif // GGML_USE_HIP
```

Use this for consistency with the rest of the code, please.

Code after:

typedef to_t_nc_cuda_t<float> to_fp32_nc_cuda_t;
typedef to_t_nc_cuda_t<half> to_fp16_nc_cuda_t;
typedef to_t_nc_cuda_t<nv_bfloat16> to_bf16_nc_cuda_t;

to_fp32_nc_cuda_t ggml_get_to_fp32_nc_cuda(ggml_type type);
to_fp16_nc_cuda_t ggml_get_to_fp16_nc_cuda(ggml_type type);
to_bf16_nc_cuda_t ggml_get_to_bf16_nc_cuda(ggml_type type);

template<typename dst_t, typename src_t>
 __host__ __device__ inline dst_t ggml_cuda_cast(src_t x) {
    if constexpr (std::is_same_v<dst_t, src_t>) {
        return x;
    } else if constexpr(std::is_same_v<dst_t, nv_bfloat16>) {
        return __float2bfloat16(float(x));
    } else if constexpr(std::is_same_v<src_t, nv_bfloat16>) {
        return __bfloat162float(x);
    } else if constexpr(std::is_same_v<src_t, float2> && std::is_same_v<dst_t, half2>) {
        return __float22half2_rn(x);
    } else if constexpr(std::is_same_v<src_t, float2> && std::is_same_v<dst_t, nv_bfloat162>) {
        // bypass compile error on cuda 12.0.1
#ifdef GGML_USE_HIP
        return __float22bfloat162_rn(x);
#else
        return {x.x, x.y};
#endif // GGML_USE_HIP
    } else if constexpr(std::is_same_v<dst_t, int32_t>) {
        return int32_t(x);
    } else {
        return float(x);
    }
}
Repository: ggerganov/llama.cpp
File: tests/test-backend-ops.cpp (C++)

Code before:
// This file defines tests for various GGML ops and backends.
// For the forward pass it asserts that the results of multiple backends computing the same GGML ops are consistent.
// For the backward pass it asserts that the gradients from backpropagation are consistent
// with the gradients obtained via the method of finite differences ("grad" mode, this is optional).
// It is also possible to check the performance ("perf" mode).
//
// this file has three sections: Section 1 does general setup, section 2 defines the GGML ops to be tested,
// and section 3 defines which tests to run.
// Quick start for adding a new GGML op: Go to section 2 and create a struct that inherits from test_case,
// then go to section 3 and add an instantiation of your struct.


// ##############################
// ## Section 1: General Setup ##
// ##############################


#include <ggml.h>
#include <ggml-alloc.h>
#include <ggml-backend.h>
#include <ggml-cpp.h>
#include <ggml-cpu.h>

#include <algorithm>
#include <array>
#include <cfloat>
#include <cinttypes>
#include <cstdarg>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <ctime>
#include <future>
#include <memory>
#include <random>
#include <regex>
#include <set>
#include <string>
#include <string_view>
#include <thread>
#include <vector>

#ifdef __EMSCRIPTEN__
#   define N_THREADS 1
#else
#   define N_THREADS std::thread::hardware_concurrency()

Review comment (suggestion):
The functions in this header are not compatible with DL backends, so they cannot be used here. There is already code here to obtain the `set_threads` function dynamically, so it shouldn't be necessary.

Code after:
// This file defines tests for various GGML ops and backends.
// For the forward pass it asserts that the results of multiple backends computing the same GGML ops are consistent.
// For the backward pass it asserts that the gradients from backpropagation are consistent
// with the gradients obtained via the method of finite differences ("grad" mode, this is optional).
// It is also possible to check the performance ("perf" mode).
//
// this file has three sections: Section 1 does general setup, section 2 defines the GGML ops to be tested,
// and section 3 defines which tests to run.
// Quick start for adding a new GGML op: Go to section 2 and create a struct that inherits from test_case,
// then go to section 3 and add an instantiation of your struct.


// ##############################
// ## Section 1: General Setup ##
// ##############################


#include <ggml.h>
#include <ggml-alloc.h>
#include <ggml-backend.h>
#include <ggml-cpp.h>

#include <algorithm>
#include <array>
#include <cfloat>
#include <cinttypes>
#include <cstdarg>
#include <cstdint>
#include <cstdio>
#include <cstdlib>
#include <cstring>
#include <ctime>
#include <future>
#include <memory>
#include <random>
#include <regex>
#include <set>
#include <string>
#include <string_view>
#include <thread>
#include <vector>

#ifdef __EMSCRIPTEN__
#   define N_THREADS 1
#else
#   define N_THREADS std::thread::hardware_concurrency()
#endif
Repository: ggerganov/llama.cpp
File: src/models/plamo3.cpp (C++)

Code before:
#include "models.h"

llm_build_plamo3::llm_build_plamo3(const llama_model & model, const llm_graph_params & params) :
    llm_graph_context(params) {
    const int64_t head_dim_q = hparams.n_embd_head_k;
    const int64_t head_dim_v = hparams.n_embd_head_v;

    ggml_tensor * inpL = build_inp_embd(model.tok_embd);
    ggml_tensor * inp_pos = build_inp_pos();

    llm_graph_input_attn_kv_iswa * inp_attn_iswa = nullptr;
    llm_graph_input_attn_kv * inp_attn = nullptr;

    if (hparams.is_swa_any()) {
        inp_attn_iswa = build_attn_inp_kv_iswa();
    } else {
        inp_attn = build_attn_inp_kv();
    }

    ggml_tensor * inp_out_ids = build_inp_out_ids();

    for (int il = 0; il < n_layer; ++il) {
        ggml_tensor * residual = inpL;

        const float freq_base_l  = model.get_rope_freq_base (cparams, il);
        const float freq_scale_l = model.get_rope_freq_scale(cparams, il);
        ggml_tensor * rope_factors = model.get_rope_factors(cparams, il);

        ggml_tensor * cur = build_norm(inpL, model.layers[il].attn_norm, NULL, LLM_NORM_RMS, il);
        cb(cur, "attn_norm", il);

        ggml_tensor * qkv = build_lora_mm(model.layers[il].wqkv, cur);
        cb(cur, "wqkv", il);

        const int32_t n_head    = hparams.n_head(il);

Review comment (suggestion):
```suggestion
    ggml_tensor * cur;
    ggml_tensor * inpL = build_inp_embd(model.tok_embd);
```

Code after:
#include "models.h"

template <bool iswa>
llm_build_plamo3<iswa>::llm_build_plamo3(const llama_model & model, const llm_graph_params & params) :
    llm_graph_context(params) {
    const int64_t head_dim_q = hparams.n_embd_head_k;
    const int64_t head_dim_v = hparams.n_embd_head_v;

    ggml_tensor * cur;
    ggml_tensor * inpL = build_inp_embd(model.tok_embd);
    ggml_tensor * inp_pos = build_inp_pos();

    using inp_attn_type = std::conditional_t<iswa, llm_graph_input_attn_kv_iswa, llm_graph_input_attn_kv>;
    inp_attn_type * inp_attn = nullptr;

    if constexpr (iswa) {
        inp_attn = build_attn_inp_kv_iswa();
    } else {
        inp_attn = build_attn_inp_kv();
    }

    ggml_tensor * inp_out_ids = build_inp_out_ids();

    for (int il = 0; il < n_layer; ++il) {
        ggml_tensor * residual = inpL;

        float freq_base_l  = 0.0f;
        float freq_scale_l = 0.0f;
        if constexpr (iswa) {
            freq_base_l  = model.get_rope_freq_base (cparams, il);
            freq_scale_l = model.get_rope_freq_scale(cparams, il);
        } else {
            freq_base_l  = freq_base;
            freq_scale_l = freq_scale;
        }
Repository: ggerganov/llama.cpp
File: convert_lora_to_gguf.py (Python)

Code before:
                    # note: mergekit-extract-lora flip and transpose A and B
                    # here we only need to transpose token_embd.lora_a, see llm_build_inp_embd()
                    if "token_embd.weight" in dest_name:
                        lora_a = lora_a.T

                    yield (dest_name + ".lora_a", lora_a)
                    yield (dest_name + ".lora_b", lora_b)

        alpha: float = lparams["lora_alpha"]

        # Use local path if --base is provided,
        # otherwise use dummy path when loading from Hugging Face
        dir_model = dir_base_model if dir_base_model is not None else Path(".")

        model_instance = LoraModel(
            dir_model,
            ftype,
            fname_out,
            is_big_endian=args.bigendian,
            use_temp_file=False,
            eager=args.no_lazy,
            dry_run=args.dry_run,
            dir_lora_model=dir_lora,
            lora_alpha=alpha,
            hparams=hparams,
            remote_hf_model_id=remote_base_model_id,
        )

        logger.info("Exporting model...")
        model_instance.write()
        logger.info(f"Model successfully exported to {model_instance.fname_out}")


Review comment (suggestion):
```suggestion
            remote_hf_model_id=base_model_id,
```

Code after:
                    lora_a, lora_b = dest_data.get_lora_A_B()

                    # note: mergekit-extract-lora flip and transpose A and B
                    # here we only need to transpose token_embd.lora_a, see llm_build_inp_embd()
                    if "token_embd.weight" in dest_name:
                        lora_a = lora_a.T

                    yield (dest_name + ".lora_a", lora_a)
                    yield (dest_name + ".lora_b", lora_b)

        alpha: float = lparams["lora_alpha"]

        model_instance = LoraModel(
            dir_base_model,
            ftype,
            fname_out,
            is_big_endian=args.bigendian,
            use_temp_file=False,
            eager=args.no_lazy,
            dry_run=args.dry_run,
            dir_lora_model=dir_lora,
            lora_alpha=alpha,
            hparams=hparams,
            remote_hf_model_id=base_model_id,
        )

        logger.info("Exporting model...")
        model_instance.write()
        logger.info(f"Model successfully exported to {model_instance.fname_out}")