correct cmd name

CONTRIBUTING.md

@@ -5,7 +5,6 @@
   - Execute [the full CI locally on your machine](ci/README.md) before publishing
 - Please rate the complexity of your PR (i.e. `Review Complexity : Low`, `Review Complexity : Medium`, `Review Complexity : High`). This makes it easier for maintainers to triage the PRs.
   - The PR template has a series of review complexity checkboxes `[ ]` that [you can mark as](https://docs.github.com/en/get-started/writing-on-github/working-with-advanced-formatting/about-task-lists) `[X]` for your convenience
-- Consider allowing write access to your branch for faster review
 - If your PR becomes stale, don't hesitate to ping the maintainers in the comments
 
 # Pull requests (for collaborators)

common/common.cpp

@@ -2039,8 +2039,8 @@ std::string fs_get_cache_file(const std::string & filename) {
 //
 // Model utils
 //
-struct llama_init_result llama_init_from_gpt_params(gpt_params & params) {
-    llama_init_result iparams;
+
+std::tuple<struct llama_model *, struct llama_context *> llama_init_from_gpt_params(gpt_params & params) {
     auto mparams = llama_model_params_from_gpt_params(params);
 
     llama_model * model = nullptr;
@@ -2055,7 +2055,7 @@ struct llama_init_result llama_init_from_gpt_params(gpt_params & params) {
 
     if (model == NULL) {
         fprintf(stderr, "%s: error: failed to load model '%s'\n", __func__, params.model.c_str());
-        return iparams;
+        return std::make_tuple(nullptr, nullptr);
     }
 
     auto cparams = llama_context_params_from_gpt_params(params);
@@ -2064,7 +2064,7 @@ struct llama_init_result llama_init_from_gpt_params(gpt_params & params) {
     if (lctx == NULL) {
         fprintf(stderr, "%s: error: failed to create context with model '%s'\n", __func__, params.model.c_str());
         llama_free_model(model);
-        return iparams;
+        return std::make_tuple(nullptr, nullptr);
     }
 
     if (!params.control_vectors.empty()) {
@@ -2075,7 +2075,7 @@ struct llama_init_result llama_init_from_gpt_params(gpt_params & params) {
         if (cvec.n_embd == -1) {
             llama_free(lctx);
             llama_free_model(model);
-            return iparams;
+            return std::make_tuple(nullptr, nullptr);
         }
 
         int err = llama_control_vector_apply(lctx,
@@ -2087,7 +2087,7 @@ struct llama_init_result llama_init_from_gpt_params(gpt_params & params) {
         if (err) {
             llama_free(lctx);
             llama_free_model(model);
-            return iparams;
+            return std::make_tuple(nullptr, nullptr);
         }
     }
 
@@ -2099,7 +2099,7 @@ struct llama_init_result llama_init_from_gpt_params(gpt_params & params) {
             fprintf(stderr, "%s: error: failed to apply lora adapter\n", __func__);
             llama_free(lctx);
             llama_free_model(model);
-            return iparams;
+            return std::make_tuple(nullptr, nullptr);
         }
         llama_lora_adapter_set(lctx, adapter, lora_scale);
     }
@@ -2135,9 +2135,7 @@ struct llama_init_result llama_init_from_gpt_params(gpt_params & params) {
         llama_reset_timings(lctx);
     }
 
-    iparams.model   = model;
-    iparams.context = lctx;
-    return iparams;
+    return std::make_tuple(model, lctx);
 }
 
 struct llama_model_params llama_model_params_from_gpt_params(const gpt_params & params) {

common/common.h

@@ -308,12 +308,8 @@ std::string fs_get_cache_file(const std::string & filename);
 // Model utils
 //
 
-struct llama_init_result {
-    struct llama_model * model = nullptr;
-    struct llama_context * context = nullptr;
-};
-
-struct llama_init_result    llama_init_from_gpt_params(gpt_params & params);
+// TODO: avoid tuplue, use struct
+std::tuple<struct llama_model *, struct llama_context *> llama_init_from_gpt_params(gpt_params & params);
 
 struct llama_model_params   llama_model_params_from_gpt_params  (const gpt_params & params);
 struct llama_context_params llama_context_params_from_gpt_params(const gpt_params & params);

convert_hf_to_gguf.py

@@ -2506,112 +2506,6 @@ class NomicBertModel(BertModel):
         self.gguf_writer.add_rope_freq_base(self.hparams["rotary_emb_base"])
 
 
-@Model.register("XLMRobertaModel")
-class XLMRobertaModel(BertModel):
-    model_arch = gguf.MODEL_ARCH.BERT
-
-    def __init__(self, *args, **kwargs):
-        super().__init__(*args, **kwargs)
-
-        # we need the pad_token_id to know how to chop down position_embd matrix
-        if (pad_token_id := self.hparams.get("pad_token_id")) is not None:
-            self._position_offset = 1 + pad_token_id
-            if "max_position_embeddings" in self.hparams:
-                self.hparams["max_position_embeddings"] -= self._position_offset
-        else:
-            self._position_offset = None
-
-    def set_vocab(self):
-        # to avoid TypeError: Descriptors cannot be created directly
-        # exception when importing sentencepiece_model_pb2
-        os.environ["PROTOCOL_BUFFERS_PYTHON_IMPLEMENTATION"] = "python"
-        from sentencepiece import SentencePieceProcessor
-        from sentencepiece import sentencepiece_model_pb2 as model
-
-        tokenizer_path = self.dir_model / 'sentencepiece.bpe.model'
-        if not tokenizer_path.is_file():
-            raise FileNotFoundError(f"File not found: {tokenizer_path}")
-
-        sentencepiece_model = model.ModelProto()  # pyright: ignore[reportAttributeAccessIssue]
-        sentencepiece_model.ParseFromString(open(tokenizer_path, "rb").read())
-        assert sentencepiece_model.trainer_spec.model_type == 1  # UNIGRAM
-
-        add_prefix = sentencepiece_model.normalizer_spec.add_dummy_prefix
-        remove_whitespaces = sentencepiece_model.normalizer_spec.remove_extra_whitespaces
-        precompiled_charsmap = sentencepiece_model.normalizer_spec.precompiled_charsmap
-
-        tokenizer = SentencePieceProcessor()
-        tokenizer.LoadFromFile(str(tokenizer_path))
-
-        vocab_size = self.hparams.get('vocab_size', tokenizer.vocab_size())
-
-        tokens: list[bytes] = [f"[PAD{i}]".encode("utf-8") for i in range(vocab_size)]
-        scores: list[float] = [-10000.0] * vocab_size
-        toktypes: list[int] = [SentencePieceTokenTypes.UNUSED] * vocab_size
-
-        for token_id in range(tokenizer.vocab_size()):
-            piece = tokenizer.IdToPiece(token_id)
-            text = piece.encode("utf-8")
-            score = tokenizer.GetScore(token_id)
-
-            toktype = SentencePieceTokenTypes.NORMAL
-            if tokenizer.IsUnknown(token_id):
-                toktype = SentencePieceTokenTypes.UNKNOWN
-            elif tokenizer.IsControl(token_id):
-                toktype = SentencePieceTokenTypes.CONTROL
-            elif tokenizer.IsUnused(token_id):
-                toktype = SentencePieceTokenTypes.UNUSED
-            elif tokenizer.IsByte(token_id):
-                toktype = SentencePieceTokenTypes.BYTE
-
-            tokens[token_id] = text
-            scores[token_id] = score
-            toktypes[token_id] = toktype
-
-        if vocab_size > len(tokens):
-            pad_count = vocab_size - len(tokens)
-            logger.debug(f"Padding vocab with {pad_count} token(s) - [PAD1] through [PAD{pad_count}]")
-            for i in range(1, pad_count + 1):
-                tokens.append(bytes(f"[PAD{i}]", encoding="utf-8"))
-                scores.append(-1000.0)
-                toktypes.append(SentencePieceTokenTypes.UNUSED)
-
-        # realign tokens (see HF tokenizer code)
-        tokens = [b'<s>', b'<pad>', b'</s>', b'<unk>'] + tokens[3:-1]
-        scores = [0.0, 0.0, 0.0, 0.0] + scores[3:-1]
-        toktypes = [
-            SentencePieceTokenTypes.CONTROL,
-            SentencePieceTokenTypes.CONTROL,
-            SentencePieceTokenTypes.CONTROL,
-            SentencePieceTokenTypes.UNKNOWN,
-        ] + toktypes[3:-1]
-
-        self.gguf_writer.add_tokenizer_model("t5")
-        self.gguf_writer.add_tokenizer_pre("default")
-        self.gguf_writer.add_token_list(tokens)
-        self.gguf_writer.add_token_scores(scores)
-        self.gguf_writer.add_token_types(toktypes)
-        self.gguf_writer.add_add_space_prefix(add_prefix)
-        self.gguf_writer.add_token_type_count(1)
-        self.gguf_writer.add_remove_extra_whitespaces(remove_whitespaces)
-        if precompiled_charsmap:
-            self.gguf_writer.add_precompiled_charsmap(precompiled_charsmap)
-
-        special_vocab = gguf.SpecialVocab(self.dir_model, n_vocab=len(tokens))
-        special_vocab.add_to_gguf(self.gguf_writer)
-
-        self.gguf_writer.add_add_bos_token(True)
-        self.gguf_writer.add_add_eos_token(True)
-
-    def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:
-        # position embeddings start at pad_token_id + 1, so just chop down the weight tensor
-        if name == "embeddings.position_embeddings.weight":
-            if self._position_offset is not None:
-                data_torch = data_torch[self._position_offset:,:]
-
-        return super().modify_tensors(data_torch, name, bid)
-
-
 @Model.register("GemmaForCausalLM")
 class GemmaModel(Model):
     model_arch = gguf.MODEL_ARCH.GEMMA

examples/cvector-generator/cvector-generator.cpp

@@ -414,10 +414,9 @@ int main(int argc, char ** argv) {
     llama_numa_init(params.numa);
 
     // load the model to get hparams
-    llama_init_result llama_init = llama_init_from_gpt_params(params);
-
-    llama_model * model = llama_init.model;
-    llama_context * ctx = llama_init.context;
+    llama_model * model;
+    llama_context * ctx;
+    std::tie(model, ctx) = llama_init_from_gpt_params(params);
 
     // int n_ctx = llama_n_ctx(ctx);
     int n_layers = llama_n_layer(model);

examples/embedding/embedding.cpp

@@ -79,11 +79,11 @@ int main(int argc, char ** argv) {
     llama_backend_init();
     llama_numa_init(params.numa);
 
-    // load the model
-    llama_init_result llama_init = llama_init_from_gpt_params(params);
+    llama_model * model;
+    llama_context * ctx;
 
-    llama_model * model = llama_init.model;
-    llama_context * ctx = llama_init.context;
+    // load the model
+    std::tie(model, ctx) = llama_init_from_gpt_params(params);
     if (model == NULL) {
         fprintf(stderr, "%s: error: unable to load model\n", __func__);
         return 1;

examples/eval-callback/eval-callback.cpp

@@ -163,10 +163,9 @@ int main(int argc, char ** argv) {
     params.warmup = false;
 
     // init
-    llama_init_result llama_init = llama_init_from_gpt_params(params);
-
-    llama_model * model = llama_init.model;
-    llama_context * ctx = llama_init.context;
+    llama_model * model;
+    llama_context * ctx;
+    std::tie(model, ctx) = llama_init_from_gpt_params(params);
     if (model == nullptr || ctx == nullptr) {
         fprintf(stderr, "%s : failed to init\n", __func__);
         return 1;

examples/imatrix/imatrix.cpp

@@ -611,10 +611,10 @@ int main(int argc, char ** argv) {
     params.warmup = false;
 
     // init
-    llama_init_result llama_init = llama_init_from_gpt_params(params);
+    llama_model * model;
+    llama_context * ctx;
 
-    llama_model * model = llama_init.model;
-    llama_context * ctx = llama_init.context;
+    std::tie(model, ctx) = llama_init_from_gpt_params(params);
     if (model == nullptr || ctx == nullptr) {
         fprintf(stderr, "%s : failed to init\n", __func__);
         return 1;

examples/infill/infill.cpp

@@ -179,10 +179,7 @@ int main(int argc, char ** argv) {
 
     // load the model and apply lora adapter, if any
     LOG("%s: load the model and apply lora adapter, if any\n", __func__);
-    llama_init_result llama_init = llama_init_from_gpt_params(params);
-
-    model = llama_init.model;
-    ctx = llama_init.context;
+    std::tie(model, ctx) = llama_init_from_gpt_params(params);
 
     if (model == NULL) {
         LOG_TEE("%s: error: unable to load model\n", __func__);

examples/lookahead/lookahead.cpp

@@ -58,11 +58,11 @@ int main(int argc, char ** argv) {
     llama_backend_init();
     llama_numa_init(params.numa);
 
-    // load the target model
-    llama_init_result llama_init = llama_init_from_gpt_params(params);
+    llama_model * model = NULL;
+    llama_context * ctx = NULL;
 
-    llama_model * model = llama_init.model;
-    llama_context * ctx = llama_init.context;
+    // load the target model
+    std::tie(model, ctx) = llama_init_from_gpt_params(params);
 
     // Tokenize the prompt
     std::vector<llama_token> inp;

examples/lookup/lookup-create.cpp

@@ -22,11 +22,11 @@ int main(int argc, char ** argv){
     llama_backend_init();
     llama_numa_init(params.numa);
 
-    // load the model
-    llama_init_result llama_init = llama_init_from_gpt_params(params);
+    llama_model * model = NULL;
+    llama_context * ctx = NULL;
 
-    llama_model * model = llama_init.model;
-    llama_context * ctx = llama_init.context;
+    // load the model
+    std::tie(model, ctx) = llama_init_from_gpt_params(params);
     GGML_ASSERT(model != nullptr);
 
     // tokenize the prompt

examples/lookup/lookup-stats.cpp

@@ -26,11 +26,11 @@ int main(int argc, char ** argv){
     llama_backend_init();
     llama_numa_init(params.numa);
 
-    // load the model
-    llama_init_result llama_init = llama_init_from_gpt_params(params);
+    llama_model * model = NULL;
+    llama_context * ctx = NULL;
 
-    llama_model * model = llama_init.model;
-    llama_context * ctx = llama_init.context;
+    // load the model
+    std::tie(model, ctx) = llama_init_from_gpt_params(params);
 
     // tokenize the prompt
     std::vector<llama_token> inp;

examples/lookup/lookup.cpp

@@ -34,11 +34,11 @@ int main(int argc, char ** argv){
     llama_backend_init();
     llama_numa_init(params.numa);
 
-    // load the model
-    llama_init_result llama_init = llama_init_from_gpt_params(params);
+    llama_model * model = NULL;
+    llama_context * ctx = NULL;
 
-    llama_model * model = llama_init.model;
-    llama_context * ctx = llama_init.context;
+    // load the model
+    std::tie(model, ctx) = llama_init_from_gpt_params(params);
 
     // tokenize the prompt
     std::vector<llama_token> inp;

examples/main/main.cpp

@@ -207,10 +207,7 @@ int main(int argc, char ** argv) {
 
     // load the model and apply lora adapter, if any
     LOG("%s: load the model and apply lora adapter, if any\n", __func__);
-    llama_init_result llama_init = llama_init_from_gpt_params(params);
-
-    model = llama_init.model;
-    ctx = llama_init.context;
+    std::tie(model, ctx) = llama_init_from_gpt_params(params);
     if (sparams.cfg_scale > 1.f) {
         struct llama_context_params lparams = llama_context_params_from_gpt_params(params);
         ctx_guidance = llama_new_context_with_model(model, lparams);

examples/parallel/parallel.cpp

@@ -129,11 +129,11 @@ int main(int argc, char ** argv) {
     llama_backend_init();
     llama_numa_init(params.numa);
 
-    // load the target model
-    llama_init_result llama_init = llama_init_from_gpt_params(params);
+    llama_model * model = NULL;
+    llama_context * ctx = NULL;
 
-    llama_model * model = llama_init.model;
-    llama_context * ctx = llama_init.context;
+    // load the target model
+    std::tie(model, ctx) = llama_init_from_gpt_params(params);
 
     // load the prompts from an external file if there are any
     if (params.prompt.empty()) {

examples/perplexity/perplexity.cpp

@@ -2018,11 +2018,11 @@ int main(int argc, char ** argv) {
     llama_backend_init();
     llama_numa_init(params.numa);
 
-    // load the model and apply lora adapter, if any
-    llama_init_result llama_init = llama_init_from_gpt_params(params);
+    llama_model * model;
+    llama_context * ctx;
 
-    llama_model * model = llama_init.model;
-    llama_context * ctx = llama_init.context;
+    // load the model and apply lora adapter, if any
+    std::tie(model, ctx) = llama_init_from_gpt_params(params);
     if (model == NULL) {
         fprintf(stderr, "%s: error: unable to load model\n", __func__);
         return 1;

examples/retrieval/retrieval.cpp

@@ -148,12 +148,11 @@ int main(int argc, char ** argv) {
     llama_backend_init();
     llama_numa_init(params.numa);
 
+    llama_model * model;
+    llama_context * ctx;
+
     // load the model
-    llama_init_result llama_init = llama_init_from_gpt_params(params);
-
-    llama_model * model = llama_init.model;
-    llama_context * ctx = llama_init.context;
-
+    std::tie(model, ctx) = llama_init_from_gpt_params(params);
     if (model == NULL) {
         fprintf(stderr, "%s: error: unable to load model\n", __func__);
         return 1;

examples/save-load-state/save-load-state.cpp

@@ -28,11 +28,10 @@ int main(int argc, char ** argv) {
     std::string result2;
 
     // init
-    llama_init_result llama_init = llama_init_from_gpt_params(params);
-
-    llama_model * model = llama_init.model;
-    llama_context * ctx = llama_init.context;
+    llama_model * model;
+    llama_context * ctx;
 
+    std::tie(model, ctx) = llama_init_from_gpt_params(params);
     if (model == nullptr || ctx == nullptr) {
         fprintf(stderr, "%s : failed to init\n", __func__);
         return 1;

examples/server/server.cpp

@@ -677,10 +677,7 @@ struct server_context {
         // dedicate one sequence to the system prompt
         params.n_parallel += 1;
 
-        llama_init_result llama_init = llama_init_from_gpt_params(params);
-
-        model = llama_init.model;
-        ctx = llama_init.context;
+        std::tie(model, ctx) = llama_init_from_gpt_params(params);
         params.n_parallel -= 1; // but be sneaky about it
         if (model == nullptr) {
             LOG_ERROR("unable to load model", {{"model", params.model}});

examples/simple/README.md

@@ -3,7 +3,7 @@
 The purpose of this example is to demonstrate a minimal usage of llama.cpp for generating text with a given prompt.
 
 ```bash
-./llama-simple -m ./models/llama-7b-v2/ggml-model-f16.gguf -p "Hello my name is"
+./simple -m ./models/llama-7b-v2/ggml-model-f16.gguf -p "Hello my name is"
 
 ...
 

examples/speculative/speculative.cpp

@@ -66,9 +66,7 @@ int main(int argc, char ** argv) {
     llama_context * ctx_dft = NULL;
 
     // load the target model
-    llama_init_result llama_init_tgt = llama_init_from_gpt_params(params);
-    model_tgt = llama_init_tgt.model;
-    ctx_tgt = llama_init_tgt.context;
+    std::tie(model_tgt, ctx_tgt) = llama_init_from_gpt_params(params);
 
     // load the draft model
     params.model = params.model_draft;
@@ -77,9 +75,7 @@ int main(int argc, char ** argv) {
         params.n_threads = params.n_threads_draft;
     }
     params.n_threads_batch = params.n_threads_batch_draft;
-    llama_init_result llama_init_dft = llama_init_from_gpt_params(params);
-    model_dft = llama_init_dft.model;
-    ctx_dft = llama_init_dft.context;
+    std::tie(model_dft, ctx_dft) = llama_init_from_gpt_params(params);
 
     const bool vocab_type_tgt = llama_vocab_type(model_tgt);
     LOG("vocab_type tgt: %d\n", vocab_type_tgt);

ggml/include/ggml.h

@@ -1140,17 +1140,16 @@ extern "C" {
 
     // group normalize along ne0*ne1*n_groups
     // used in stable-diffusion
+    // TODO: eps is hardcoded to 1e-6 for now
     GGML_API struct ggml_tensor * ggml_group_norm(
             struct ggml_context * ctx,
             struct ggml_tensor  * a,
-            int                   n_groups,
-            float                 eps);
+            int                   n_groups);
 
     GGML_API struct ggml_tensor * ggml_group_norm_inplace(
             struct ggml_context * ctx,
             struct ggml_tensor  * a,
-            int                   n_groups,
-            float                 eps);
+            int                   n_groups);
 
     // a - x
     // b - dy

ggml/src/ggml-cann.cpp

@@ -896,10 +896,11 @@ GGML_CALL static void ggml_backend_cann_buffer_init_tensor(
  * @param size Size of the data to be copied, in bytes.
  */
 GGML_CALL static void ggml_backend_cann_buffer_set_tensor(
-    ggml_backend_buffer_t buffer, ggml_tensor *tensor, const void *data,
+    ggml_backend_buffer_t buffer, ggml_tensor* tensor, const void* data,
     size_t offset, size_t size) {
-    ggml_backend_cann_buffer_context *ctx =
-        (ggml_backend_cann_buffer_context *)buffer->context;
+    // GGML_ASSERT(size == ggml_nbytes(tensor));
+    ggml_backend_cann_buffer_context* ctx =
+        (ggml_backend_cann_buffer_context*)buffer->context;
 
     ggml_cann_set_device(ctx->device);
     // TODO: refer to cann(#6017), it use thread's default stream.
@@ -907,21 +908,22 @@ GGML_CALL static void ggml_backend_cann_buffer_set_tensor(
     // Why aclrtSynchronizeDevice?
 
     if (!need_transform(tensor->type)) {
-        ACL_CHECK(aclrtMemcpy((char *)tensor->data + offset, size, data, size,
-                              ACL_MEMCPY_HOST_TO_DEVICE));
+        ACL_CHECK(aclrtMemcpy(tensor->data, size, (const char*)data + offset,
+                              size, ACL_MEMCPY_HOST_TO_DEVICE));
     } else {
-        void *transform_buffer = malloc(size);
-        ggml_backend_cann_transform(tensor, data, transform_buffer);
+        void* transform_buffer = malloc(size);
+        ggml_backend_cann_transform(tensor, (const char*)data + offset,
+                                    transform_buffer);
 
 #ifndef NDEBUG
-        void *check_buffer = malloc(size);
+        void* check_buffer = malloc(size);
         ggml_backend_cann_transform_back(tensor, transform_buffer,
                                          check_buffer);
-        GGML_ASSERT(memcmp(data, check_buffer, size) == 0);
+        GGML_ASSERT(memcmp((const char*)data + offset, check_buffer, size) ==
+                    0);
         free(check_buffer);
 #endif
-        ACL_CHECK(aclrtMemcpy((char *)tensor->data + offset, size,
-                              transform_buffer, size,
+        ACL_CHECK(aclrtMemcpy(tensor->data, size, transform_buffer, size,
                               ACL_MEMCPY_HOST_TO_DEVICE));
         free(transform_buffer);
     }
@@ -943,20 +945,21 @@ GGML_CALL static void ggml_backend_cann_buffer_set_tensor(
 GGML_CALL static void ggml_backend_cann_buffer_get_tensor(
     ggml_backend_buffer_t buffer, const ggml_tensor* tensor, void* data,
     size_t offset, size_t size) {
+    GGML_ASSERT(size == ggml_nbytes(tensor));
     ggml_backend_cann_buffer_context* ctx =
         (ggml_backend_cann_buffer_context*)buffer->context;
 
     ggml_cann_set_device(ctx->device);
 
     if (!need_transform(tensor->type)) {
-        ACL_CHECK(aclrtMemcpy(data, size, (char*)tensor->data + offset, size,
+        ACL_CHECK(aclrtMemcpy((char*)data + offset, size, tensor->data, size,
                               ACL_MEMCPY_DEVICE_TO_HOST));
     } else {
         void* transform_buffer = malloc(size);
-        ACL_CHECK(aclrtMemcpy(transform_buffer, size,
-                              (char*)tensor->data + offset, size,
+        ACL_CHECK(aclrtMemcpy(transform_buffer, size, tensor->data, size,
                               ACL_MEMCPY_DEVICE_TO_HOST));
-        ggml_backend_cann_transform_back(tensor, transform_buffer, data);
+        ggml_backend_cann_transform_back(tensor, transform_buffer,
+                                         (char*)data + offset);
         free(transform_buffer);
     }
 }
@@ -1445,41 +1448,42 @@ ggml_backend_cann_get_default_buffer_type(ggml_backend_t backend) {
  * @param size Size of the data to copy in bytes.
  */
 GGML_CALL static void ggml_backend_cann_set_tensor_async(ggml_backend_t backend,
-                                                         ggml_tensor *tensor,
-                                                         const void *data,
+                                                         ggml_tensor* tensor,
+                                                         const void* data,
                                                          size_t offset,
                                                          size_t size) {
-    ggml_backend_cann_context *cann_ctx =
-        (ggml_backend_cann_context *)backend->context;
+    ggml_backend_cann_context* cann_ctx =
+        (ggml_backend_cann_context*)backend->context;
 
     if (!need_transform(tensor->type)) {
-        ACL_CHECK(aclrtMemcpyAsync((char *)tensor->data + offset, size, data,
-                                   size, ACL_MEMCPY_HOST_TO_DEVICE,
-                                   cann_ctx->stream()));
+        ACL_CHECK(aclrtMemcpyAsync(
+            tensor->data, size, (const char*)data + offset, size,
+            ACL_MEMCPY_HOST_TO_DEVICE, cann_ctx->stream()));
     } else {
-        void *transform_buffer = malloc(size);
-        ggml_backend_cann_transform(tensor, data, transform_buffer);
+        void* transform_buffer = malloc(size);
+        ggml_backend_cann_transform(tensor, (const char*)data + offset,
+                                    transform_buffer);
 
 #ifndef NDEBUG
-        void *check_buffer = malloc(size);
+        void* check_buffer = malloc(size);
         ggml_backend_cann_transform_back(tensor, transform_buffer,
                                          check_buffer);
-        GGML_ASSERT(memcmp(data, check_buffer, size));
+        GGML_ASSERT(memcmp((const char*)data + offset, check_buffer, size));
         free(check_buffer);
 #endif
-        ACL_CHECK(aclrtMemcpyAsync(
-            (char *)tensor->data + offset, size, transform_buffer, size,
-            ACL_MEMCPY_HOST_TO_DEVICE, cann_ctx->stream()));
+        ACL_CHECK(aclrtMemcpyAsync(tensor->data, size, transform_buffer, size,
+                                   ACL_MEMCPY_HOST_TO_DEVICE,
+                                   cann_ctx->stream()));
         ACL_CHECK(aclrtSynchronizeStream(cann_ctx->stream()));
         free(transform_buffer);
     }
 }
 
 GGML_CALL static void ggml_backend_cann_get_tensor_async(
-    ggml_backend_t backend, const ggml_tensor *tensor, void *data,
+    ggml_backend_t backend, const ggml_tensor* tensor, void* data,
     size_t offset, size_t size) {
-    ggml_backend_cann_context *cann_ctx =
-        (ggml_backend_cann_context *)backend->context;
+    ggml_backend_cann_context* cann_ctx =
+        (ggml_backend_cann_context*)backend->context;
     ggml_backend_buffer_t buf =
         tensor->view_src ? tensor->view_src->buffer : tensor->buffer;
 
@@ -1487,16 +1491,17 @@ GGML_CALL static void ggml_backend_cann_get_tensor_async(
                 "unsupported buffer type");
 
     if (!need_transform(tensor->type)) {
-        ACL_CHECK(aclrtMemcpyAsync(data, size, (char *)tensor->data + offset,
+        ACL_CHECK(aclrtMemcpyAsync((char*)data + offset, size, tensor->data,
                                    size, ACL_MEMCPY_DEVICE_TO_HOST,
                                    cann_ctx->stream()));
     } else {
-        void *transform_buffer = malloc(size);
-        ACL_CHECK(aclrtMemcpyAsync(
-            transform_buffer, size, (char *)tensor->data + offset, size,
-            ACL_MEMCPY_DEVICE_TO_HOST, cann_ctx->stream()));
+        void* transform_buffer = malloc(size);
+        ACL_CHECK(aclrtMemcpyAsync(transform_buffer, size, tensor->data, size,
+                                   ACL_MEMCPY_DEVICE_TO_HOST,
+                                   cann_ctx->stream()));
         ACL_CHECK(aclrtSynchronizeStream(cann_ctx->stream()));
-        ggml_backend_cann_transform_back(tensor, transform_buffer, data);
+        ggml_backend_cann_transform_back(tensor, transform_buffer,
+                                         (char*)data + offset);
         free(transform_buffer);
     }
 }

ggml/src/ggml-cann/aclnn_ops.cpp

@@ -464,11 +464,9 @@ void ggml_cann_group_norm(ggml_backend_cann_context& ctx, ggml_tensor* dst) {
     aclTensor* acl_src = ggml_cann_create_tensor(src);
     aclTensor* acl_dst = ggml_cann_create_tensor(dst);
 
+    const float eps = 1e-6f;  // TODO: make this a parameter
     int n_groups = dst->op_params[0];
 
-    float eps;
-    memcpy(&eps, dst->op_params + 1, sizeof(float));
-
     uint64_t workspaceSize = 0;
     aclOpExecutor* executor;
     void* workspaceAddr = nullptr;

ggml/src/ggml-cuda.cu

@@ -1501,7 +1501,7 @@ static void ggml_cuda_op_mul_mat(
         }
 
         // If src0 is on a temporary compute buffers (partial offloading) there may be some padding that needs to be cleared:
-        if (ne00 % MATRIX_ROW_PADDING != 0 && ggml_is_quantized(src0->type) && ggml_backend_buffer_get_usage(src0->buffer) == GGML_BACKEND_BUFFER_USAGE_COMPUTE && src0->view_src == nullptr) {
+        if (ne00 % MATRIX_ROW_PADDING != 0 && ggml_backend_buffer_get_usage(src0->buffer) == GGML_BACKEND_BUFFER_USAGE_COMPUTE && src0->view_src == nullptr) {
             const int64_t nbytes_data    = ggml_row_size(src0->type, (dev[id].row_high - dev[id].row_low)*ne00);
             const int64_t nbytes_padding = ggml_row_size(src0->type, MATRIX_ROW_PADDING - ne00 % MATRIX_ROW_PADDING);
             CUDA_CHECK(cudaMemsetAsync(dev[id].src0_dd + nbytes_data , 0, nbytes_padding, stream));

ggml/src/ggml-cuda/norm.cu

@@ -142,7 +142,8 @@ static void norm_f32_cuda(const float * x, float * dst, const int ncols, const i
     }
 }
 
-static void group_norm_f32_cuda(const float * x, float * dst, const int num_groups, const float eps, const int group_size, const int ne_elements, cudaStream_t stream) {
+static void group_norm_f32_cuda(const float * x, float * dst, const int num_groups, const int group_size, const int ne_elements, cudaStream_t stream) {
+    static const float eps = 1e-6f;
     if (group_size < 1024) {
         const dim3 block_dims(WARP_SIZE, 1, 1);
         group_norm_f32<WARP_SIZE><<<num_groups, block_dims, 0, stream>>>(x, dst, group_size, ne_elements, eps);
@@ -195,12 +196,8 @@ void ggml_cuda_op_group_norm(ggml_backend_cuda_context & ctx, ggml_tensor * dst)
     GGML_ASSERT( dst->type == GGML_TYPE_F32);
 
     int num_groups = dst->op_params[0];
-
-    float eps;
-    memcpy(&eps, dst->op_params + 1, sizeof(float));
-
     int group_size = src0->ne[0] * src0->ne[1] * ((src0->ne[2] + num_groups - 1) / num_groups);
-    group_norm_f32_cuda(src0_d, dst_d, num_groups * src0->ne[3], eps, group_size, ggml_nelements(src0), stream);
+    group_norm_f32_cuda(src0_d, dst_d, num_groups * src0->ne[3], group_size, ggml_nelements(src0), stream);
 }
 
 void ggml_cuda_op_rms_norm(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {

ggml/src/ggml-metal.m

@@ -2229,8 +2229,10 @@ static enum ggml_status ggml_metal_graph_compute(
                         GGML_ASSERT(ne00 % 4 == 0);
                         GGML_ASSERT(ggml_is_contiguous(src0));
 
-                        float eps;
-                        memcpy(&eps, dst->op_params + 1, sizeof(float));
+                        //float eps;
+                        //memcpy(&eps, dst->op_params, sizeof(float));
+
+                        const float eps = 1e-6f; // TODO: temporarily hardcoded
 
                         const int32_t n_groups = ((int32_t *) dst->op_params)[0];
 

ggml/src/ggml-sycl/norm.cpp

@@ -225,8 +225,9 @@ static void norm_f32_sycl(const float* x, float* dst, const int ncols,
 }
 
 static void group_norm_f32_sycl(const float* x, float* dst,
-    const int num_groups, const float eps, const int group_size,
+    const int num_groups, const int group_size,
     const int ne_elements, queue_ptr stream, int device) {
+    static const float eps = 1e-6f;
     if (group_size < 1024) {
         const sycl::range<3> block_dims(1, 1, WARP_SIZE);
         stream->submit([&](sycl::handler& cgh) {
@@ -342,12 +343,8 @@ void ggml_sycl_op_group_norm(ggml_backend_sycl_context& ctx, const ggml_tensor*
     GGML_ASSERT(dst->type == GGML_TYPE_F32);
 
     int num_groups = dst->op_params[0];
-
-    float eps;
-    memcpy(&eps, dst->op_params + 1, sizeof(float));
-
     int group_size = src0->ne[0] * src0->ne[1] * ((src0->ne[2] + num_groups - 1) / num_groups);
-    group_norm_f32_sycl(src0_dd, dst_dd, num_groups, eps, group_size, src0->ne[0] * src0->ne[1] * src0->ne[2], main_stream, ctx.device);
+    group_norm_f32_sycl(src0_dd, dst_dd, num_groups, group_size, src0->ne[0] * src0->ne[1] * src0->ne[2], main_stream, ctx.device);
 
     (void)src1;
     (void)dst;

ggml/src/ggml.c

@@ -5374,7 +5374,6 @@ static struct ggml_tensor * ggml_group_norm_impl(
     struct ggml_context * ctx,
     struct ggml_tensor * a,
     int n_groups,
-    float eps,
     bool inplace) {
 
     bool is_node = false;
@@ -5385,8 +5384,7 @@ static struct ggml_tensor * ggml_group_norm_impl(
 
     struct ggml_tensor * result = inplace ? ggml_view_tensor(ctx, a) : ggml_dup_tensor(ctx, a);
 
-    ggml_set_op_params_i32(result, 0, n_groups);
-    ggml_set_op_params_f32(result, 1, eps);
+    result->op_params[0] = n_groups;
 
     result->op = GGML_OP_GROUP_NORM;
     result->grad = is_node ? ggml_dup_tensor(ctx, result) : NULL;
@@ -5398,17 +5396,15 @@ static struct ggml_tensor * ggml_group_norm_impl(
 struct ggml_tensor * ggml_group_norm(
     struct ggml_context * ctx,
     struct ggml_tensor * a,
-    int n_groups,
-    float eps) {
-    return ggml_group_norm_impl(ctx, a, n_groups, eps, false);
+    int n_groups) {
+    return ggml_group_norm_impl(ctx, a, n_groups, false);
 }
 
 struct ggml_tensor * ggml_group_norm_inplace(
     struct ggml_context * ctx,
     struct ggml_tensor * a,
-    int n_groups,
-    float eps) {
-    return ggml_group_norm_impl(ctx, a, n_groups, eps, true);
+    int n_groups) {
+    return ggml_group_norm_impl(ctx, a, n_groups, true);
 }
 
 // ggml_mul_mat
@@ -12099,10 +12095,9 @@ static void ggml_compute_forward_group_norm_f32(
 
     GGML_TENSOR_UNARY_OP_LOCALS
 
-    // TODO: optimize
+    const float eps = 1e-6f; // TODO: make this a parameter
 
-    float eps;
-    memcpy(&eps, dst->op_params + 1, sizeof(float));
+    // TODO: optimize
 
     int n_channels = src0->ne[2];
     int n_groups = dst->op_params[0];

ggml/src/vulkan-shaders/CMakeLists.txt

@@ -1,7 +1,5 @@
-find_package (Threads REQUIRED)
 
 set(TARGET vulkan-shaders-gen)
 add_executable(${TARGET} vulkan-shaders-gen.cpp)
 install(TARGETS ${TARGET} RUNTIME)
 target_compile_features(${TARGET} PRIVATE cxx_std_11)
-target_link_libraries(vulkan-shaders-gen PUBLIC Threads::Threads)

ggml/src/vulkan-shaders/vulkan-shaders-gen.cpp

@@ -22,7 +22,6 @@
 #ifdef _WIN32
     #include <windows.h>
     #include <direct.h> // For _mkdir on Windows
-    #include <algorithm> // For std::replace on w64devkit
 #else
     #include <unistd.h>
     #include <sys/wait.h>

src/llama-vocab.cpp

@@ -816,9 +816,6 @@ struct llm_tokenizer_ugm {
      * the best tokenization.
     */
     void tokenize(const std::string & text, std::vector<llama_vocab::id> & output) {
-        // get current size of output (for reversal later)
-        size_t output_size = output.size();
-
         // normalize the input first
         std::string normalized;
         normalize(text, &normalized);
@@ -898,7 +895,7 @@ struct llm_tokenizer_ugm {
         }
 
         // reverse the output since we added tokens starting from the end of the input
-        std::reverse(output.begin() + output_size, output.end());
+        std::reverse(output.begin(), output.end());
     }
 
 private:

tests/test-backend-ops.cpp

@@ -1511,7 +1511,6 @@ struct test_group_norm : public test_case {
     const ggml_type type;
     const std::array<int64_t, 4> ne;
     const int32_t num_groups;
-    const float eps;
 
     std::string vars() override {
         return VARS_TO_STR3(type, ne, num_groups);
@@ -1519,13 +1518,12 @@ struct test_group_norm : public test_case {
 
     test_group_norm(ggml_type type = GGML_TYPE_F32,
             std::array<int64_t, 4> ne = {64, 64, 320, 1},
-            int32_t num_groups = 32,
-            float eps = 1e-6f)
-        : type(type), ne(ne), num_groups(num_groups), eps(eps) {}
+            int32_t num_groups = 32)
+        : type(type), ne(ne), num_groups(num_groups) {}
 
     ggml_tensor * build_graph(ggml_context * ctx) override {
         ggml_tensor * a = ggml_new_tensor(ctx, type, 4, ne.data());
-        ggml_tensor * out = ggml_group_norm(ctx, a, num_groups, eps);
+        ggml_tensor * out = ggml_group_norm(ctx, a, num_groups);
         return out;
     }
 };