ronantakizawa/codereview-bench huggingface dataset:
Repository: getumbrel/llama-gpt
File: api/run.sh (Shell)

Code before:
#!/bin/bash

if [ -z "$MODEL" ]
then
    echo "Please set the MODEL_FILE environment variable"
    exit 1
fi

if [ -z "$MODEL_DOWNLOAD_URL" ]
then
    echo "Please set the MODEL_DOWNLOAD_URL environment variable"
    exit 1
fi


if [ ! -f $MODEL ]; then
    echo "Model file not found. Downloading..."
    curl -L -o $MODEL $MODEL_DOWNLOAD_URL
else
    echo "$MODEL model found."
fi

# Get the number of available threads on the system
n_threads=$(grep -c ^processor /proc/cpuinfo)

# Define context window
n_ctx=4096

# Offload everything to CPU
n_gpu_layers=0

# Define batch size
n_batch=2096
# If total RAM is less than 8GB, set batch size to 1024
total_ram=$(cat /proc/meminfo | grep MemTotal | awk '{print $2}')
if [ $total_ram -lt 8000000 ]; then
    n_batch=1024
fi

echo "Initializing server with:"
echo "Batch size: $n_batch"
echo "Number of CPU threads: $n_threads"
echo "Number of GPU layers: $n_gpu_layers"
echo "Context window: $n_ctx"

exec python3 -m llama_cpp.server --n_ctx $n_ctx --n_threads $n_threads --n_gpu_layers $n_gpu_layers --n_batch $n_batch


Review comment (suggestion):
```suggestion
    apt-get update --yes --quiet
    apt-get install --yes --quiet curl
    curl -L -o $MODEL $MODEL_DOWNLOAD_URL
```

We can install `curl` here if model doesn't exist.

Code after:
#!/bin/bash

if [ -z "$MODEL" ]
then
    echo "Please set the MODEL_FILE environment variable"
    exit 1
fi

if [ -z "$MODEL_DOWNLOAD_URL" ]
then
    echo "Please set the MODEL_DOWNLOAD_URL environment variable"
    exit 1
fi

# check if curl is installed

if ! [ -x "$(command -v curl)" ]; then
    echo "curl is not installed. Installing..."
    apt-get update --yes --quiet
    apt-get install --yes --quiet curl
fi


if [ ! -f $MODEL ]; then
    echo "Model file not found. Downloading..."
    curl -L -o $MODEL $MODEL_DOWNLOAD_URL
else
    echo "$MODEL model found."
fi

make build
# Get the number of available threads on the system
n_threads=$(grep -c ^processor /proc/cpuinfo)

# Define context window
n_ctx=4096

# Offload everything to CPU
n_gpu_layers=0

# Define batch size
n_batch=2096
# If total RAM is less than 8GB, set batch size to 1024
total_ram=$(cat /proc/meminfo | grep MemTotal | awk '{print $2}')
if [ $total_ram -lt 8000000 ]; then
    n_batch=1024
fi

echo "Initializing server with:"
echo "Batch size: $n_batch"
echo "Number of CPU threads: $n_threads"
Repository: ggerganov/llama.cpp
File: scripts/hf.sh (Shell)

Code before:
#!/bin/bash
#
# Shortcut for downloading HF models
#
# Usage:
#   ./main -m $(./scripts/hf.sh https://huggingface.co/TheBloke/Mixtral-8x7B-v0.1-GGUF/resolve/main/mixtral-8x7b-v0.1.Q4_K_M.gguf)
#   ./main -m $(./scripts/hf.sh --url https://huggingface.co/TheBloke/Mixtral-8x7B-v0.1-GGUF/blob/main/mixtral-8x7b-v0.1.Q4_K_M.gguf)
#   ./main -m $(./scripts/hf.sh --repo TheBloke/Mixtral-8x7B-v0.1-GGUF --file mixtral-8x7b-v0.1.Q4_K_M.gguf)
#

# all logs go to stderr
function log {
    echo "$@" 1>&2
}

function usage {
    log "Usage: $0 [[--url] <url>] [--repo <repo>] [--file <file>] [--outdir dir [-h|--help]"
    exit 1
}

# check for curl or wget
function has_cmd {
    if ! [ -x "$(command -v $1)" ]; then
        return 1
    fi
}

if has_cmd wget; then
    cmd="wget -q --show-progress -c -O %s/%s %s"
elif has_cmd curl; then
    cmd="curl -C - -f --output-dir %s -o %s -L %s"
else
    log "[E] curl or wget not found"
    exit 1
fi

url=""
repo=""
file=""
outdir="."

# parse args

Review comment (suggestion):
```suggestion
    log "Usage: $0 [[--url] <url>] [--repo <repo>] [--file <file>] [--outdir <dir>] [-h|--help]"
```

Code after:
#!/bin/bash
#
# Shortcut for downloading HF models
#
# Usage:
#   ./main -m $(./scripts/hf.sh https://huggingface.co/TheBloke/Mixtral-8x7B-v0.1-GGUF/resolve/main/mixtral-8x7b-v0.1.Q4_K_M.gguf)
#   ./main -m $(./scripts/hf.sh --url https://huggingface.co/TheBloke/Mixtral-8x7B-v0.1-GGUF/blob/main/mixtral-8x7b-v0.1.Q4_K_M.gguf)
#   ./main -m $(./scripts/hf.sh --repo TheBloke/Mixtral-8x7B-v0.1-GGUF --file mixtral-8x7b-v0.1.Q4_K_M.gguf)
#

# all logs go to stderr
function log {
    echo "$@" 1>&2
}

function usage {
    log "Usage: $0 [[--url] <url>] [--repo <repo>] [--file <file>] [--outdir <dir> [-h|--help]"
    exit 1
}

# check for curl or wget
function has_cmd {
    if ! [ -x "$(command -v $1)" ]; then
        return 1
    fi
}

if has_cmd wget; then
    cmd="wget -q --show-progress -c -O %s/%s %s"
elif has_cmd curl; then
    cmd="curl -C - -f --output-dir %s -o %s -L %s"
else
    log "[E] curl or wget not found"
    exit 1
fi

url=""
repo=""
file=""
outdir="."

# parse args
Repository: ggerganov/llama.cpp
File: include/llama.h (C/C++)

Code before:
    // NOTE: Used by llama.cpp examples, avoid using in third-party apps. Instead, do your own performance measurements.
    //

    // performance timing information
    struct llama_perf_data {
        // llama_context
        double t_start_ms;
        double t_load_ms;
        double t_p_eval_ms;
        double t_eval_ms;

        int32_t n_p_eval;
        int32_t n_eval;

        // llama_sampler_chain
        double t_sample_ms;

        int32_t n_sample;
    };

    enum llama_perf_type {
        LLAMA_PERF_TYPE_CONTEXT       = 0,
        LLAMA_PERF_TYPE_SAMPLER_CHAIN = 1,
    };

    LLAMA_API struct llama_perf_data llama_perf_get(const void * ctx, enum llama_perf_type type);

    LLAMA_API void llama_perf_print(const void * ctx, enum llama_perf_type type);
    LLAMA_API void llama_perf_reset(      void * ctx, enum llama_perf_type type);

    LLAMA_API void llama_perf_dump_yaml(FILE * stream, const struct llama_context * ctx);

#ifdef __cplusplus
}
#endif

#endif // LLAMA_H


Review comment (bug):
I think it would be preferable to have two separate functions, just to remove the possibility of calling it with the wrong type of pointer.

Code after:

    //
    // Performance utils
    //
    // NOTE: Used by llama.cpp examples, avoid using in third-party apps. Instead, do your own performance measurements.
    //

    struct llama_perf_context_data {
        double t_start_ms;
        double t_load_ms;
        double t_p_eval_ms;
        double t_eval_ms;

        int32_t n_p_eval;
        int32_t n_eval;
    };

    struct llama_perf_sampler_data {
        double t_sample_ms;

        int32_t n_sample;
    };

    LLAMA_API struct llama_perf_context_data llama_perf_context      (const struct llama_context * ctx);
    LLAMA_API void                           llama_perf_context_print(const struct llama_context * ctx);
    LLAMA_API void                           llama_perf_context_reset(      struct llama_context * ctx);

    // NOTE: the following work only with samplers constructed via llama_sampler_chain_init
    LLAMA_API struct llama_perf_sampler_data llama_perf_sampler      (const struct llama_sampler * chain);
    LLAMA_API void                           llama_perf_sampler_print(const struct llama_sampler * chain);
    LLAMA_API void                           llama_perf_sampler_reset(      struct llama_sampler * chain);

    LLAMA_API void llama_perf_dump_yaml(FILE * stream, const struct llama_context * ctx);

#ifdef __cplusplus
}
#endif

#endif // LLAMA_H
Repository: ggerganov/llama.cpp
File: src/llama-sampling.cpp (C++)

Code before:
    };
}

// xtc

struct llama_sampler_xtc {
    const float    probability;
    const float    threshold;
    const float    threshold_max;
    const size_t   min_keep;

    const uint32_t seed;
    uint32_t       seed_cur;

    std::mt19937   rng;
};

static const char * llama_sampler_xtc_name(const struct llama_sampler * /*smpl*/) {
    return "xtc";
}

static void llama_sample_xtc_apply(struct llama_sampler * smpl, llama_token_data_array * cur_p) {
    auto * ctx = (llama_sampler_xtc *) smpl->ctx;

    if (ctx->probability <= 0.0f
        || ctx->threshold >= 1.0f
        || ctx->threshold_max <= 0.0f
        || ctx->threshold_max <= ctx->threshold
        || cur_p->size <= 2) {
        return;
    }

    std::uniform_real_distribution<float> distance(0.0f, 1.0f);
    float chance = distance(ctx->rng);
    if (chance > ctx->probability) return;

    // in case it's not sorted/recalculated yet
    llama_sampler_softmax_impl(cur_p);

    int pos_first = -1;
    int pos_last = 0;

    for (size_t i = 0; i < cur_p->size; ++i) {
        if (cur_p->data[i].p - ctx->threshold >= -1e-5) {
            if (cur_p->data[i].p - ctx->threshold_max > 1e-3) pos_first = i;
            pos_last = i;
        } else {
            break;
        }
    }


Review comment (suggestion):
```suggestion
        || ctx->threshold > 0.5f
```

Any threshold > 0.5 already disables XTC, since it is impossible for more than one token to have a probability greater than 1/2.


Code after:
    };
}

// xtc

struct llama_sampler_xtc {
    const float    probability;
    const float    threshold;
    const size_t   min_keep;

    const uint32_t seed;
    uint32_t       seed_cur;

    std::mt19937   rng;
};

static const char * llama_sampler_xtc_name(const struct llama_sampler * /*smpl*/) {
    return "xtc";
}

static void llama_sample_xtc_apply(struct llama_sampler * smpl, llama_token_data_array * cur_p) {
    auto * ctx = (llama_sampler_xtc *) smpl->ctx;

    if (ctx->probability <= 0.0f
        || ctx->threshold > 0.5f
        || cur_p->size < 2) {
        return;
    }

    std::uniform_real_distribution<float> distribution(0.0f, 1.0f);
    float chance = distribution(ctx->rng);
    if (chance > ctx->probability) return;

    // in case it's not sorted/recalculated yet
    llama_sampler_softmax_impl(cur_p);

    int pos_last = 0;

    for (size_t i = 0; i < cur_p->size; ++i) {
        if (cur_p->data[i].p >= ctx->threshold) {
            pos_last = i;
        } else break;
    }

    if (cur_p->size - pos_last >= ctx->min_keep && pos_last > 0) {
        cur_p->data += pos_last;
        cur_p->size -= pos_last;
    }
}

static struct llama_sampler * llama_sampler_xtc_clone(const struct llama_sampler * smpl) {
Repository: ggerganov/llama.cpp
File: src/llama-sampling.cpp (C++)

Code before:
    };
}

// xtc

struct llama_sampler_xtc {
    const float    probability;
    const float    threshold;
    const size_t   min_keep;

    const uint32_t seed;
    uint32_t       seed_cur;

    std::mt19937   rng;
};

static const char * llama_sampler_xtc_name(const struct llama_sampler * /*smpl*/) {
    return "xtc";
}

static void llama_sample_xtc_apply(struct llama_sampler * smpl, llama_token_data_array * cur_p) {
    auto * ctx = (llama_sampler_xtc *) smpl->ctx;

    if (ctx->probability <= 0.0f
        || ctx->threshold > 0.5f
        || cur_p->size <= 2) {
        return;
    }

    std::uniform_real_distribution<float> distance(0.0f, 1.0f);
    float chance = distance(ctx->rng);
    if (chance > ctx->probability) return;

    // in case it's not sorted/recalculated yet
    llama_sampler_softmax_impl(cur_p);

    int pos_last = 0;

    for (size_t i = 0; i < cur_p->size; ++i) {
        if (cur_p->data[i].p - ctx->threshold >= -1e-5) {
            pos_last = i;
        } else break;
    }

    if (cur_p->size - pos_last >= ctx->min_keep && pos_last > 0) {

        size_t last_idx = cur_p->size - pos_last;

        for (size_t i = 0; i <= last_idx; ++i) {
            cur_p->data[i] = cur_p->data[i + pos_last];
        }

Review comment (suggestion):
If there are exactly two tokens we can still remove one of them if they are both above the threshold. So this inequality should probably be a strict `<`.

Code after:
    };
}

// xtc

struct llama_sampler_xtc {
    const float    probability;
    const float    threshold;
    const size_t   min_keep;

    const uint32_t seed;
    uint32_t       seed_cur;

    std::mt19937   rng;
};

static const char * llama_sampler_xtc_name(const struct llama_sampler * /*smpl*/) {
    return "xtc";
}

static void llama_sample_xtc_apply(struct llama_sampler * smpl, llama_token_data_array * cur_p) {
    auto * ctx = (llama_sampler_xtc *) smpl->ctx;

    if (ctx->probability <= 0.0f
        || ctx->threshold > 0.5f
        || cur_p->size < 2) {
        return;
    }

    std::uniform_real_distribution<float> distribution(0.0f, 1.0f);
    float chance = distribution(ctx->rng);
    if (chance > ctx->probability) return;

    // in case it's not sorted/recalculated yet
    llama_sampler_softmax_impl(cur_p);

    int pos_last = 0;

    for (size_t i = 0; i < cur_p->size; ++i) {
        if (cur_p->data[i].p >= ctx->threshold) {
            pos_last = i;
        } else break;
    }

    if (cur_p->size - pos_last >= ctx->min_keep && pos_last > 0) {
        cur_p->data += pos_last;
        cur_p->size -= pos_last;
    }
}

static struct llama_sampler * llama_sampler_xtc_clone(const struct llama_sampler * smpl) {
Repository: ggerganov/llama.cpp
File: src/llama-sampling.cpp (C++)

Code before:
        || ctx->threshold > 0.5f
        || cur_p->size <= 2) {
        return;
    }

    std::uniform_real_distribution<float> distance(0.0f, 1.0f);
    float chance = distance(ctx->rng);
    if (chance > ctx->probability) return;

    // in case it's not sorted/recalculated yet
    llama_sampler_softmax_impl(cur_p);

    int pos_last = 0;

    for (size_t i = 0; i < cur_p->size; ++i) {
        if (cur_p->data[i].p - ctx->threshold >= -1e-5) {
            pos_last = i;
        } else break;
    }

    if (cur_p->size - pos_last >= ctx->min_keep && pos_last > 0) {

        size_t last_idx = cur_p->size - pos_last;

        for (size_t i = 0; i <= last_idx; ++i) {
            cur_p->data[i] = cur_p->data[i + pos_last];
        }

        cur_p->size = cur_p->size - pos_last;
    }
}

static struct llama_sampler * llama_sampler_xtc_clone(const struct llama_sampler * smpl) {
    const auto * ctx = (const llama_sampler_xtc *) smpl->ctx;
    auto * result = llama_sampler_init_xtc(ctx->probability, ctx->threshold, ctx->min_keep, ctx->seed);

    // copy the state
    {
        auto * result_ctx = (llama_sampler_xtc *) result->ctx;

        result_ctx->rng = ctx->rng;
    }

    return result;
}

static void llama_sampler_xtc_free(struct llama_sampler * smpl) {
    delete (llama_sampler_xtc *) smpl->ctx;
}

static void llama_sampler_xtc_reset(struct llama_sampler * smpl) {

Review comment (suggestion):
Can't the whole thing be accomplished by just advancing the `cur_p->data` pointer to skip the elements to be removed? That would avoid the need to shift potentially hundreds of thousands of elements.

Code after:
        || ctx->threshold > 0.5f
        || cur_p->size < 2) {
        return;
    }

    std::uniform_real_distribution<float> distribution(0.0f, 1.0f);
    float chance = distribution(ctx->rng);
    if (chance > ctx->probability) return;

    // in case it's not sorted/recalculated yet
    llama_sampler_softmax_impl(cur_p);

    int pos_last = 0;

    for (size_t i = 0; i < cur_p->size; ++i) {
        if (cur_p->data[i].p >= ctx->threshold) {
            pos_last = i;
        } else break;
    }

    if (cur_p->size - pos_last >= ctx->min_keep && pos_last > 0) {
        cur_p->data += pos_last;
        cur_p->size -= pos_last;
    }
}

static struct llama_sampler * llama_sampler_xtc_clone(const struct llama_sampler * smpl) {
    const auto * ctx = (const llama_sampler_xtc *) smpl->ctx;
    auto * result = llama_sampler_init_xtc(ctx->probability, ctx->threshold, ctx->min_keep, ctx->seed);

    // copy the state
    {
        auto * result_ctx = (llama_sampler_xtc *) result->ctx;

        result_ctx->rng = ctx->rng;
    }

    return result;
}

static void llama_sampler_xtc_free(struct llama_sampler * smpl) {
    delete (llama_sampler_xtc *) smpl->ctx;
}

static void llama_sampler_xtc_reset(struct llama_sampler * smpl) {
    auto * ctx = (llama_sampler_xtc *) smpl->ctx;
    ctx->seed_cur = get_rng_seed(ctx->seed);
    ctx->rng.seed(ctx->seed_cur);
}

static struct llama_sampler_i llama_sampler_xtc_i = {
Repository: ggerganov/llama.cpp
File: src/llama-sampling.cpp (C++)

Code before:
    auto * ctx = (llama_sampler_xtc *) smpl->ctx;

    if (ctx->probability <= 0.0f
        || ctx->threshold > 0.5f
        || cur_p->size < 2) {
        return;
    }

    std::uniform_real_distribution<float> distribution(0.0f, 1.0f);
    float chance = distribution(ctx->rng);
    if (chance > ctx->probability) return;

    // in case it's not sorted/recalculated yet
    llama_sampler_softmax_impl(cur_p);

    int pos_last = 0;

    for (size_t i = 0; i < cur_p->size; ++i) {
        if (cur_p->data[i].p >= ctx->threshold) {
            pos_last = i;
        } else break;
    }

    if (cur_p->size - pos_last >= ctx->min_keep && pos_last > 0) {
        cur_p->data += pos_last;
        cur_p->size = cur_p->size - pos_last;
    }
}

static struct llama_sampler * llama_sampler_xtc_clone(const struct llama_sampler * smpl) {
    const auto * ctx = (const llama_sampler_xtc *) smpl->ctx;
    auto * result = llama_sampler_init_xtc(ctx->probability, ctx->threshold, ctx->min_keep, ctx->seed);

    // copy the state
    {
        auto * result_ctx = (llama_sampler_xtc *) result->ctx;

        result_ctx->rng = ctx->rng;
    }

    return result;
}

static void llama_sampler_xtc_free(struct llama_sampler * smpl) {
    delete (llama_sampler_xtc *) smpl->ctx;
}

static void llama_sampler_xtc_reset(struct llama_sampler * smpl) {
    auto * ctx = (llama_sampler_xtc *) smpl->ctx;
    ctx->seed_cur = get_rng_seed(ctx->seed);
    ctx->rng.seed(ctx->seed_cur);

Review comment (suggestion):
```suggestion
        cur_p->size -= pos_last;
```

Code after:
    auto * ctx = (llama_sampler_xtc *) smpl->ctx;

    if (ctx->probability <= 0.0f
        || ctx->threshold > 0.5f
        || cur_p->size < 2) {
        return;
    }

    std::uniform_real_distribution<float> distribution(0.0f, 1.0f);
    float chance = distribution(ctx->rng);
    if (chance > ctx->probability) return;

    // in case it's not sorted/recalculated yet
    llama_sampler_softmax_impl(cur_p);

    int pos_last = 0;

    for (size_t i = 0; i < cur_p->size; ++i) {
        if (cur_p->data[i].p >= ctx->threshold) {
            pos_last = i;
        } else break;
    }

    if (cur_p->size - pos_last >= ctx->min_keep && pos_last > 0) {
        cur_p->data += pos_last;
        cur_p->size -= pos_last;
    }
}

static struct llama_sampler * llama_sampler_xtc_clone(const struct llama_sampler * smpl) {
    const auto * ctx = (const llama_sampler_xtc *) smpl->ctx;
    auto * result = llama_sampler_init_xtc(ctx->probability, ctx->threshold, ctx->min_keep, ctx->seed);

    // copy the state
    {
        auto * result_ctx = (llama_sampler_xtc *) result->ctx;

        result_ctx->rng = ctx->rng;
    }

    return result;
}

static void llama_sampler_xtc_free(struct llama_sampler * smpl) {
    delete (llama_sampler_xtc *) smpl->ctx;
}

static void llama_sampler_xtc_reset(struct llama_sampler * smpl) {
    auto * ctx = (llama_sampler_xtc *) smpl->ctx;
    ctx->seed_cur = get_rng_seed(ctx->seed);
    ctx->rng.seed(ctx->seed_cur);
Repository: ggerganov/llama.cpp
File: examples/quantize/quantize.cpp (C++)

Code before:
                }
            }
        }
        imatrix_data = std::move(tmp);
    }
    if (!imatrix_data.empty()) {
        printf("%s: have %d importance matrix entries\n", __func__, int(imatrix_data.size()));
    }
    return m_last_call;
}

static ggml_type parse_ggml_type(const char * arg) {
    for (int i = 0; i < GGML_TYPE_COUNT; ++i) {
        auto type = (ggml_type)i;
        const auto * name = ggml_type_name(type);
        if (name && striequals(name, arg)) {
            return type;
        }
    }
    fprintf(stderr, "%s: invalid ggml_type '%s'\n", __func__, arg);
    return GGML_TYPE_COUNT;
}

// Allowed tensors for arbitrary quantization with --tensor-type option
static const std::vector<std::string> ALLOWED_TENSOR_TYPE = {
    "attn_k",
    "attn_kv_a_mqa",
    "attn_kv_b",
    "attn_out",
    "attn_q_a",
    "attn_q_b",
    "attn_q",
    "attn_qkv",
    "attn_v",
    "channel_mix_key",
    "channel_mix_receptance",
    "channel_mix_value",
    "cls_out",
    "cls",
    "dec_attn_k",
    "dec_attn_out",
    "dec_attn_q",
    "dec_attn_v",
    "dec_cross_attn_k",
    "dec_cross_attn_out",
    "dec_cross_attn_q",
    "dec_cross_attn_v",
    "ffn_act",
    "ffn_down_exp",
    "ffn_down_shexp",
    "ffn_down",

Review comment (refactor):
I'm rethinking about this, maybe we can simplify this functionality by adding just 2 flags:
- `--dump-mapping` to get the list of tensors and the target quantized type. User can then modify the target quant directly
- `--mapping FILE` user can then specify a custom mapping file given from step above

Code after:
                    tmp.emplace(std::move(e));
                }
            }
        }
        imatrix_data = std::move(tmp);
    }
    if (!imatrix_data.empty()) {
        printf("%s: have %d importance matrix entries\n", __func__, int(imatrix_data.size()));
    }
    return m_last_call;
}

static ggml_type parse_ggml_type(const char * arg) {
    for (int i = 0; i < GGML_TYPE_COUNT; ++i) {
        auto type = (ggml_type)i;
        const auto * name = ggml_type_name(type);
        if (name && striequals(name, arg)) {
            return type;
        }
    }
    fprintf(stderr, "%s: invalid ggml_type '%s'\n", __func__, arg);
    return GGML_TYPE_COUNT;
}

// Allowed tensors for arbitrary quantization with --tensor-type option
static const std::vector<std::string> ALLOWED_TENSOR_TYPE = {
    "attn_k",
    "attn_kv_a_mqa",
    "attn_kv_b",
    "attn_o",
    "attn_output",
    "attn_q",
    "attn_q_a",
    "attn_q_b",
    "attn_qkv",
    "attn_v",
    "channel_mix_key",
    "channel_mix_receptance",
    "channel_mix_value",
    "cls",
    "cls.output",
    "cross_attn_k",
    "cross_attn_o",
    "cross_attn_q",
    "cross_attn_v",
    "ffn_act",
    "ffn_down",
    "ffn_down_exps",
    "ffn_down_shexp",
    "ffn_gate",
    "ffn_gate_exps",
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 * src0,
    void * dst,
    const int64_t ne0,
    const int64_t ne00,
    const int64_t 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 int64_t i1 = blockIdx.x;

    if (i1 >= ne01 || i2 >= ne02 || i3 >= ne03) {
        return;
    }

    const char * src0_ptr = (const char *)src0 + i3*nb03 + i2*nb02 + i1*nb01;

Review comment (performance):
```suggestion
    const void * __restrict__ src0,
    void * __restrict__ dst,
```

On older GPUs especially you can gain some performance by telling the compiler that the pointers aren't aliased.

Code after:
#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 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 int64_t i1 = blockIdx.x;

    if (i1 >= ne01 || i2 >= ne02 || i3 >= ne03) {
        return;
    }

    const char * src0_ptr = (const char *)src0 + i3*nb03 + i2*nb02 + i1*nb01;