ronantakizawa/codereview-bench huggingface dataset:
Repository: ggerganov/llama.cpp
File: ggml/src/ggml-cpu/llamafile/sgemm.cpp (C++)

Code before:
        }
    }

    const TA *const A;
    const block_q8_0 *const B;
    float *C;
    const int64_t k;
    const int64_t lda;
    const int64_t ldb;
    const int64_t ldc;
    const int ith;
    const int nth;
};

class tinyBLAS_PPC {
  public:
    tinyBLAS_PPC(int64_t k,
                const float *A, int64_t lda,
                const float *B, int64_t ldb,
                float *C, int64_t ldc,
                int ith, int nth)
        : A(A), B(B), C(C), k(k), lda(lda), ldb(ldb), ldc(ldc), ith(ith), nth(nth) {
    }

    void matmul(int64_t m, int64_t n) {
       mnpack(0, m, 0, n);
    }

  private:

    void (tinyBLAS_PPC::*kernel)(int64_t, int64_t);

    inline void vector_permute_store_4(vector float *src, float *vecOffset) {
       vector float t1, t2, t3, t4, t5, t6, t7, t8;
           t1 = vec_mergeh(src[0], src[1]);
           t2 = vec_mergeh(src[2], src[3]);
           t3 = vec_mergel(src[0], src[1]);
           t4 = vec_mergel(src[2], src[3]);

           t5 = vec_xxpermdi(t1, t2, 0);
           t6 = vec_xxpermdi(t1, t2, 3);
           t7 = vec_xxpermdi(t3, t4, 0);
           t8 = vec_xxpermdi(t3, t4, 3);

           vec_xst(t5, 0, vecOffset);
           vec_xst(t6, 0, vecOffset + 4);
           vec_xst(t7, 0, vecOffset + 8);
           vec_xst(t8, 0, vecOffset + 12);
       }

    inline void vector_permute_store_8(vector float *src, float *vecOffset) {

Review comment (suggestion):
```suggestion
        if (m % mc == 0 && n % nc == 0 && k % kc == 0) {
```

Code after:
        }
    }

    const TA *const A;
    const block_q8_0 *const B;
    float *C;
    const int64_t k;
    const int64_t lda;
    const int64_t ldb;
    const int64_t ldc;
    const int ith;
    const int nth;
};

class tinyBLAS_PPC {
  public:
    tinyBLAS_PPC(int64_t k,
                const float * A, int64_t lda,
                const float * B, int64_t ldb,
                float * C, int64_t ldc,
                int ith, int nth)
        : A(A), B(B), C(C), k(k), lda(lda), ldb(ldb), ldc(ldc), ith(ith), nth(nth) {
    }

    void matmul(int64_t m, int64_t n) {
        int64_t mc = 256; int64_t nc = 256; int64_t kc = 256;
        if (m % mc == 0 && n % nc == 0 && k % kc == 0) {
            matmul_tiled(m, n, mc, nc, kc);
        } else {
            mnpack(0, m, 0, n);
        }
    }

  private:

    inline void save_acc(acc_t * ACC, int64_t ii, int64_t jj) {
        vec_t vec_C[4];
        __builtin_mma_disassemble_acc(vec_C, ACC);
        for (int I = 0; I < 4; I++) {
            for (int J = 0; J < 4; J++) {
                *((float *)(C+ii+((jj+J)*ldc)+I)) = *((float *)&vec_C[I]+J);
            }
        }
    }

    inline void add_save_acc(acc_t * ACC, int64_t ii, int64_t jj) {
        vec_t vec_C[4];
        __builtin_mma_disassemble_acc(vec_C, ACC);
        for (int I = 0; I < 4; I++) {
            for (int J = 0; J < 4; J++) {
                float * c_ptr = (float *)(C+ii+((jj+J)*ldc)+I);
Repository: ggerganov/llama.cpp
File: gguf-py/gguf/gguf_writer.py (Python)

Code before:
from __future__ import annotations

import logging
import os
import shutil
import struct
import sys
import tempfile
from dataclasses import dataclass
from enum import Enum, auto
from math import prod
from pathlib import Path
from io import BufferedWriter
from typing import IO, Any, Sequence, Mapping, Union
from string import ascii_letters, digits

import numpy as np

from .constants import (
    GGUF_DEFAULT_ALIGNMENT,
    GGUF_MAGIC,
    GGUF_VERSION,
    GGMLQuantizationType,
    GGUFEndian,
    GGUFValueType,
    Keys,
    RopeScalingType,
    PoolingType,
    TokenType,
    ExpertGatingFuncType,
)

from .quants import quant_shape_from_byte_shape

logger = logging.getLogger(__name__)


SHARD_NAME_FORMAT = "{:s}-{:05d}-of-{:05d}.gguf"


Review comment (suggestion):
```suggestion
from typing import IO, Any, Sequence, Mapping
```

Code after:
from __future__ import annotations

import logging
import os
import shutil
import struct
import sys
import tempfile
from dataclasses import dataclass
from enum import Enum, auto
from math import prod
from pathlib import Path
from io import BufferedWriter
from typing import IO, Any, Sequence, Mapping
from string import ascii_letters, digits

import numpy as np

from .constants import (
    GGUF_DEFAULT_ALIGNMENT,
    GGUF_MAGIC,
    GGUF_VERSION,
    GGMLQuantizationType,
    GGUFEndian,
    GGUFValueType,
    Keys,
    RopeScalingType,
    PoolingType,
    TokenType,
    ExpertGatingFuncType,
)

from .quants import quant_shape_from_byte_shape

logger = logging.getLogger(__name__)


SHARD_NAME_FORMAT = "{:s}-{:05d}-of-{:05d}.gguf"
Repository: ggerganov/llama.cpp
File: src/models/modern-bert.cpp (C++)

Code before:
#include "models.h"

template <bool iswa>
llm_build_modern_bert<iswa>::llm_build_modern_bert(const llama_model & model, const llm_graph_params & params) : llm_graph_context(params) {
    const int64_t n_embd_head     = hparams.n_embd_head_v;
    const int64_t n_embd_gqa      = hparams.n_embd_v_gqa();

    GGML_ASSERT(n_embd_head == hparams.n_embd_head_k);

    ggml_tensor * cur = nullptr;
    ggml_tensor * inpL = nullptr;
    ggml_tensor * inp_pos = build_inp_pos();

    // construct input embeddings (token, type, position)
    inpL = build_inp_embd(model.tok_embd);
    cb(inpL, "inp_embd", -1);

    // embed layer norm
    inpL = build_norm(inpL, model.tok_norm, nullptr, LLM_NORM, -1);
    cb(inpL, "inp_norm", -1);

    ggml_tensor * inp_out_ids = build_inp_out_ids();

    auto * inp_attn = build_attn_inp_no_cache();

    for (int il = 0; il < n_layer; ++il) {
        float freq_base_l  = 0.0f;

        if constexpr (iswa) {
            freq_base_l = model.get_rope_freq_base(cparams, il);
        } else {

Review comment (suggestion):
```suggestion
    const int64_t n_embd_head = hparams.n_embd_head_v;
    const int64_t n_embd_gqa  = hparams.n_embd_v_gqa();
```

Code after:
#include "models.h"

template <bool iswa>
llm_build_modern_bert<iswa>::llm_build_modern_bert(const llama_model & model, const llm_graph_params & params) : llm_graph_context(params) {
    const int64_t n_embd_head = hparams.n_embd_head_v;
    const int64_t n_embd_gqa  = hparams.n_embd_v_gqa();

    GGML_ASSERT(n_embd_head == hparams.n_embd_head_k);

    ggml_tensor * cur;
    ggml_tensor * inpL;
    ggml_tensor * inp_pos = build_inp_pos();

    // construct input embeddings (token, type, position)
    inpL = build_inp_embd(model.tok_embd);
    cb(inpL, "inp_embd", -1);

    // embed layer norm
    inpL = build_norm(inpL, model.tok_norm, nullptr, LLM_NORM, -1);
    cb(inpL, "inp_norm", -1);

    ggml_tensor * inp_out_ids = build_inp_out_ids();

    auto * inp_attn = build_attn_inp_no_cache();

    for (int il = 0; il < n_layer; ++il) {
        float freq_base_l  = 0.0f;

        if constexpr (iswa) {
            freq_base_l = model.get_rope_freq_base(cparams, il);
        } else {
Repository: ggerganov/llama.cpp
File: src/models/modern-bert.cpp (C++)

Code before:
#include "models.h"

template <bool iswa>
llm_build_modern_bert<iswa>::llm_build_modern_bert(const llama_model & model, const llm_graph_params & params) : llm_graph_context(params) {
    const int64_t n_embd_head     = hparams.n_embd_head_v;
    const int64_t n_embd_gqa      = hparams.n_embd_v_gqa();

    GGML_ASSERT(n_embd_head == hparams.n_embd_head_k);

    ggml_tensor * cur = nullptr;
    ggml_tensor * inpL = nullptr;
    ggml_tensor * inp_pos = build_inp_pos();

    // construct input embeddings (token, type, position)
    inpL = build_inp_embd(model.tok_embd);
    cb(inpL, "inp_embd", -1);

    // embed layer norm
    inpL = build_norm(inpL, model.tok_norm, nullptr, LLM_NORM, -1);
    cb(inpL, "inp_norm", -1);

    ggml_tensor * inp_out_ids = build_inp_out_ids();

    auto * inp_attn = build_attn_inp_no_cache();

    for (int il = 0; il < n_layer; ++il) {
        float freq_base_l  = 0.0f;

        if constexpr (iswa) {
            freq_base_l = model.get_rope_freq_base(cparams, il);
        } else {
            freq_base_l = freq_base;
        }

        ggml_tensor * cur = inpL;


Review comment (suggestion):
```suggestion
    ggml_tensor * cur;
    ggml_tensor * inpL;
```

Code after:
#include "models.h"

template <bool iswa>
llm_build_modern_bert<iswa>::llm_build_modern_bert(const llama_model & model, const llm_graph_params & params) : llm_graph_context(params) {
    const int64_t n_embd_head = hparams.n_embd_head_v;
    const int64_t n_embd_gqa  = hparams.n_embd_v_gqa();

    GGML_ASSERT(n_embd_head == hparams.n_embd_head_k);

    ggml_tensor * cur;
    ggml_tensor * inpL;
    ggml_tensor * inp_pos = build_inp_pos();

    // construct input embeddings (token, type, position)
    inpL = build_inp_embd(model.tok_embd);
    cb(inpL, "inp_embd", -1);

    // embed layer norm
    inpL = build_norm(inpL, model.tok_norm, nullptr, LLM_NORM, -1);
    cb(inpL, "inp_norm", -1);

    ggml_tensor * inp_out_ids = build_inp_out_ids();

    auto * inp_attn = build_attn_inp_no_cache();

    for (int il = 0; il < n_layer; ++il) {
        float freq_base_l  = 0.0f;

        if constexpr (iswa) {
            freq_base_l = model.get_rope_freq_base(cparams, il);
        } else {
            freq_base_l = freq_base;
        }

        cur = inpL;
Repository: ggerganov/llama.cpp
File: src/llama-hparams.h (C/C++)

Code before:
    uint32_t n_embd_k_gqa_max() const;
    uint32_t n_embd_v_gqa_max() const;

    // dimension of the rolling state embeddings
    // corresponds to Mamba's conv_states size or RWKV's token_shift states size
    uint32_t n_embd_r() const;

    // dimension of the recurrent state embeddings
    uint32_t n_embd_s() const;

    // whether or not the given layer is recurrent (for hybrid models)
    bool is_recurrent(uint32_t il) const;

    uint32_t n_pos_per_embd() const;

    bool is_swa(uint32_t il) const;

    bool has_kv(uint32_t il) const;

    // number of layers for which has_kv() returns true
    uint32_t n_layer_kv() const;

    bool is_masked_swa(llama_pos p0, llama_pos p1) const;
};

static_assert(std::is_trivially_copyable<llama_hparams>::value, "llama_hparams must be trivially copyable");



Review comment (refactor):
Changed this to a static function.

Maybe it should become a member like this:

```suggestion
   bool is_masked_swa(uint32_t il, llama_pos p0, llama_pos p1) const;
```

But let's refactor this after the `master` stabilized.

Code after:
    uint32_t n_embd_k_gqa_max() const;
    uint32_t n_embd_v_gqa_max() const;

    // dimension of the rolling state embeddings
    // corresponds to Mamba's conv_states size or RWKV's token_shift states size
    uint32_t n_embd_r() const;

    // dimension of the recurrent state embeddings
    uint32_t n_embd_s() const;

    // whether or not the given layer is recurrent (for hybrid models)
    bool is_recurrent(uint32_t il) const;

    uint32_t n_pos_per_embd() const;

    bool is_swa(uint32_t il) const;

    bool has_kv(uint32_t il) const;

    // number of layers for which has_kv() returns true
    uint32_t n_layer_kv() const;

    // note that this function uses different SWA parameters from those in the hparams
    // TODO: think of a better place for this function
    // TODO: pack the SWA params in a struct?
    static bool is_masked_swa(uint32_t n_swa, llama_swa_type swa_type, llama_pos p0, llama_pos p1);
};

static_assert(std::is_trivially_copyable<llama_hparams>::value, "llama_hparams must be trivially copyable");
Repository: ggerganov/llama.cpp
File: common/common.h (C/C++)

Code before:

//
// Split utils
//

namespace {

const char * const LLM_KV_SPLIT_NO            = "split.no";
const char * const LLM_KV_SPLIT_COUNT         = "split.count";
const char * const LLM_KV_SPLIT_TENSORS_COUNT = "split.tensors.count";

}

//
// MoE utils
//

const char * const LLM_FFN_EXPS_REGEX = "\\.ffn_(up|down|gate)_exps";

inline std::string llm_ffn_exps_block_regex(int idx) {
    return string_format("blk\\.%d%s", idx, LLM_FFN_EXPS_REGEX);
}

inline llama_model_tensor_buft_override llm_ffn_exps_cpu_override() {
    return { LLM_FFN_EXPS_REGEX, ggml_backend_cpu_buffer_type() };
}

//
// training utils
//

ggml_opt_dataset_t common_opt_dataset_init(struct llama_context * ctx, const std::vector<llama_token> & tokens, int64_t stride);

// "adamw" or "sgd" (case insensitive)
enum ggml_opt_optimizer_type common_opt_get_optimizer(const char *);


Review comment (suggestion):
`static` instead of `inline` should be a bit more correct here.

```suggestion
static std::string llm_ffn_exps_block_regex(int idx) {
    return string_format("blk\\.%d%s", idx, LLM_FFN_EXPS_REGEX);
}

static llama_model_tensor_buft_override llm_ffn_exps_cpu_override() {
    return { LLM_FFN_EXPS_REGEX, ggml_backend_cpu_buffer_type() };
}
```

Code after:

//
// Split utils
//

namespace {

const char * const LLM_KV_SPLIT_NO            = "split.no";
const char * const LLM_KV_SPLIT_COUNT         = "split.count";
const char * const LLM_KV_SPLIT_TENSORS_COUNT = "split.tensors.count";

}

//
// MoE utils
//

const char * const LLM_FFN_EXPS_REGEX = "\\.ffn_(up|down|gate)_exps";

static std::string llm_ffn_exps_block_regex(int idx) {
    return string_format("blk\\.%d%s", idx, LLM_FFN_EXPS_REGEX);
}

static llama_model_tensor_buft_override llm_ffn_exps_cpu_override() {
    return { LLM_FFN_EXPS_REGEX, ggml_backend_cpu_buffer_type() };
}

//
// training utils
//

ggml_opt_dataset_t common_opt_dataset_init(struct llama_context * ctx, const std::vector<llama_token> & tokens, int64_t stride);

// "adamw" or "sgd" (case insensitive)
enum ggml_opt_optimizer_type common_opt_get_optimizer(const char *);
Repository: ggerganov/llama.cpp
File: ggml/src/ggml-cuda/pad_reflect_1d.cu (CUDA)

Code before:
#include "pad_reflect_1d.cuh"

static __global__ void pad_reflect_1d_kernel_f32(
    const void * __restrict__ src0,
    void * __restrict__ dst,
    const int64_t ne0,
    const int64_t ne00,
    const int64_t ne01,
    const uint3 ne01_packed,
    const int64_t ne02,
    const int64_t ne03,
    const int64_t nb00,
    const int64_t nb01,
    const int64_t nb02,
    const int64_t nb03,
    const int64_t nb0,
    const int64_t nb1,
    const int64_t nb2,
    const int64_t nb3,
    const int p0,
    const int p1) {

    const int64_t i3 = blockIdx.z;
    const int64_t i2 = blockIdx.y;

    const uint2 div_mod_packed = fast_div_modulo(blockIdx.x, ne01_packed);
    const int64_t tile1  = div_mod_packed.y;     // i1
    const int64_t tile0  = div_mod_packed.x;     // nth i0 tile
    const int64_t i1     = tile1;
    const int64_t i0     = threadIdx.x + tile0 * blockDim.x;
    if (i0 >= ne0 || i1 >= ne01 || i2 >= ne02 || i3 >= ne03) {
        return;
    }


Review comment (suggestion):
```suggestion
    const uint3   ne01,
```

Pass only the `fastdiv` values and use `ne01.z` to get the original values. Feel free to add an assert in host code for the numerical range but if it's insufficient the kernel cannot launch anyways.

Code after:
#include "pad_reflect_1d.cuh"

static __global__ __launch_bounds__(CUDA_PAD_REFLECT_1D_BLOCK_SIZE, 1) void
    pad_reflect_1d_kernel_f32(
        const void * __restrict__ src0,
        void * __restrict__       dst,
        const int64_t             ne0,
        const int64_t             ne00,
        const uint3               ne01,
        const int64_t             ne02,
        const int64_t             ne03,
        const int64_t             nb00,
        const int64_t             nb01,
        const int64_t             nb02,
        const int64_t             nb03,
        const int64_t             nb0,
        const int64_t             nb1,
        const int64_t             nb2,
        const int64_t             nb3,
        const int                 p0,
        const int                 p1) {
    const int64_t i3 = blockIdx.z;
    const int64_t i2 = blockIdx.y;

    const uint2   div_mod_packed = fast_div_modulo(blockIdx.x, ne01);
    const int64_t tile1          = div_mod_packed.y;  // i1
    const int64_t tile0          = div_mod_packed.x;  // nth i0 tile
    const int64_t i1             = tile1;
    const int64_t i0             = threadIdx.x + tile0 * blockDim.x;

    // ne01.z is original value of unpacked ne01 (see init_fastdiv_values in common.cuh)
    if (i0 >= ne0 || i1 >= ne01.z || i2 >= ne02 || i3 >= ne03) {
        return;
    }
Repository: ggerganov/llama.cpp
File: tools/server/webui/src/lib/components/app/chat/ChatMessages/ChatMessageUser.svelte (Svelte)

Code before:
		class: className = '',
		message,
		isEditing,
		editedContent,
		siblingInfo = null,
		showDeleteDialog,
		deletionInfo,
		onCancelEdit,
		onSaveEdit,
		onEditKeydown,
		onEditedContentChange,
		onCopy,
		onEdit,
		onDelete,
		onConfirmDelete,
		onNavigateToSibling,
		onShowDeleteDialogChange,
		textareaElement = $bindable()
	}: Props = $props();

	let innerWidth = $state(0);
	let isMultiline = $state(false);
	let messageElement: HTMLElement | undefined = $state();

	$effect(() => {
		if (messageElement && message.content.trim() && innerWidth > 0) {
			if (message.content.includes('\n')) {
				isMultiline = true;
				return;
			}

			const computedStyle = window.getComputedStyle(messageElement);
			const lineHeight = parseFloat(computedStyle.lineHeight);
			const actualHeight = messageElement.scrollHeight;

			isMultiline = actualHeight > lineHeight * 1.2;
		}
	});
</script>

<svelte:window bind:innerWidth />

<div
	aria-label="User message with actions"
	class="group flex flex-col items-end gap-3 md:gap-2 {className}"
	role="group"
>
	{#if isEditing}
		<div class="w-full max-w-[80%]">
			<textarea
				bind:this={textareaElement}

Review comment (performance):
will doing some processing each time input content changed impacts on performance? in the past, some users complained about input field being very laggy on certain low-end systems

especially when `getComputedStyle` is not really good for performance. should we do this with css trick instead?

Code after:
		class: className = '',
		message,
		isEditing,
		editedContent,
		siblingInfo = null,
		showDeleteDialog,
		deletionInfo,
		onCancelEdit,
		onSaveEdit,
		onEditKeydown,
		onEditedContentChange,
		onCopy,
		onEdit,
		onDelete,
		onConfirmDelete,
		onNavigateToSibling,
		onShowDeleteDialogChange,
		textareaElement = $bindable()
	}: Props = $props();

	let isMultiline = $state(false);
	let messageElement: HTMLElement | undefined = $state();

	$effect(() => {
		if (!messageElement || !message.content.trim()) return;

		if (message.content.includes('\n')) {
			isMultiline = true;
			return;
		}

		const resizeObserver = new ResizeObserver((entries) => {
			for (const entry of entries) {
				const element = entry.target as HTMLElement;
				const estimatedSingleLineHeight = 24; // Typical line height for text-md

				isMultiline = element.offsetHeight > estimatedSingleLineHeight * 1.5;
			}
		});

		resizeObserver.observe(messageElement);

		return () => {
			resizeObserver.disconnect();
		};
	});
</script>

<div
	aria-label="User message with actions"
	class="group flex flex-col items-end gap-3 md:gap-2 {className}"
Repository: ggerganov/llama.cpp
File: ggml/include/ggml.h (C/C++)

Code before:
        GGML_OP_COUNT,
    };

    enum ggml_unary_op {
        GGML_UNARY_OP_ABS,
        GGML_UNARY_OP_SGN,
        GGML_UNARY_OP_NEG,
        GGML_UNARY_OP_STEP,
        GGML_UNARY_OP_TANH,
        GGML_UNARY_OP_ELU,
        GGML_UNARY_OP_RELU,
        GGML_UNARY_OP_SIGMOID,
        GGML_UNARY_OP_GELU,
        GGML_UNARY_OP_GELU_QUICK,
        GGML_UNARY_OP_SILU,
        GGML_UNARY_OP_HARDSWISH,
        GGML_UNARY_OP_HARDSIGMOID,
        GGML_UNARY_OP_EXP,
        GGML_UNARY_OP_GELU_ERF,
        GGML_UNARY_OP_FLOOR,
        GGML_UNARY_OP_CEIL,
        GGML_UNARY_OP_ROUND,
        GGML_UNARY_OP_TRUNC,
        GGML_UNARY_OP_XIELU,


        GGML_UNARY_OP_COUNT,
    };

    enum ggml_glu_op {
        GGML_GLU_OP_REGLU,
        GGML_GLU_OP_GEGLU,
        GGML_GLU_OP_SWIGLU,
        GGML_GLU_OP_SWIGLU_OAI,
        GGML_GLU_OP_GEGLU_ERF,
        GGML_GLU_OP_GEGLU_QUICK,

        GGML_GLU_OP_COUNT,
    };

    enum ggml_object_type {
        GGML_OBJECT_TYPE_TENSOR,
        GGML_OBJECT_TYPE_GRAPH,
        GGML_OBJECT_TYPE_WORK_BUFFER
    };

    enum ggml_log_level {
        GGML_LOG_LEVEL_NONE  = 0,
        GGML_LOG_LEVEL_DEBUG = 1,
        GGML_LOG_LEVEL_INFO  = 2,
        GGML_LOG_LEVEL_WARN  = 3,

Review comment (suggestion):
```suggestion
        GGML_UNARY_OP_XIELU,
        GGML_UNARY_OP_FLOOR,
        GGML_UNARY_OP_CEIL,
        GGML_UNARY_OP_ROUND,
        GGML_UNARY_OP_TRUNC,

        GGML_UNARY_OP_COUNT,
```

Code after:
        GGML_OP_COUNT,
    };

    enum ggml_unary_op {
        GGML_UNARY_OP_ABS,
        GGML_UNARY_OP_SGN,
        GGML_UNARY_OP_NEG,
        GGML_UNARY_OP_STEP,
        GGML_UNARY_OP_TANH,
        GGML_UNARY_OP_ELU,
        GGML_UNARY_OP_RELU,
        GGML_UNARY_OP_SIGMOID,
        GGML_UNARY_OP_GELU,
        GGML_UNARY_OP_GELU_QUICK,
        GGML_UNARY_OP_SILU,
        GGML_UNARY_OP_HARDSWISH,
        GGML_UNARY_OP_HARDSIGMOID,
        GGML_UNARY_OP_EXP,
        GGML_UNARY_OP_GELU_ERF,
        GGML_UNARY_OP_XIELU,
        GGML_UNARY_OP_FLOOR,
        GGML_UNARY_OP_CEIL,
        GGML_UNARY_OP_ROUND,
        GGML_UNARY_OP_TRUNC,

        GGML_UNARY_OP_COUNT,
    };

    enum ggml_glu_op {
        GGML_GLU_OP_REGLU,
        GGML_GLU_OP_GEGLU,
        GGML_GLU_OP_SWIGLU,
        GGML_GLU_OP_SWIGLU_OAI,
        GGML_GLU_OP_GEGLU_ERF,
        GGML_GLU_OP_GEGLU_QUICK,

        GGML_GLU_OP_COUNT,
    };

    enum ggml_object_type {
        GGML_OBJECT_TYPE_TENSOR,
        GGML_OBJECT_TYPE_GRAPH,
        GGML_OBJECT_TYPE_WORK_BUFFER
    };

    enum ggml_log_level {
        GGML_LOG_LEVEL_NONE  = 0,
        GGML_LOG_LEVEL_DEBUG = 1,
        GGML_LOG_LEVEL_INFO  = 2,
        GGML_LOG_LEVEL_WARN  = 3,
        GGML_LOG_LEVEL_ERROR = 4,