Merge branch 'master' into concedo_experimental

# Conflicts:
#	README.md

CMakeLists.txt

@@ -167,11 +167,6 @@ if (${CMAKE_SYSTEM_PROCESSOR} MATCHES "arm" OR ${CMAKE_SYSTEM_PROCESSOR} MATCHES
  if (MSVC)
  # TODO: arm msvc?
  else()
- if (${CMAKE_SYSTEM_PROCESSOR} MATCHES "aarch64")
- # Apple M1, M2, etc.
- # Raspberry Pi 3, 4, Zero 2 (64-bit)
- add_compile_options(-mcpu=native)
- endif()
  if (${CMAKE_SYSTEM_PROCESSOR} MATCHES "armv6")
  # Raspberry Pi 1, Zero
  add_compile_options(-mfpu=neon-fp-armv8 -mfp16-format=ieee -mno-unaligned-access)

convert.py

@@ -136,7 +136,7 @@ def find_n_mult(n_ff: int, n_embd: int) -> int:
  calc_ff = (((8*n_embd) // 3 + n_mult - 1) // n_mult)*n_mult
  if calc_ff == n_ff:
  return n_mult
- return 1
+ raise Exception(f"failed to find n_mult for (n_ff={n_ff}, n_embd={n_embd}).")
 
 @dataclass
 class Params:
@@ -321,6 +321,10 @@ def astype(self, data_type: DataType) -> 'Tensor': ...
  @abstractmethod
  def permute(self, n_head: int) -> 'Tensor': ...
  @abstractmethod
+ def permute_part(self, n_part: int, n_head: int) -> 'UnquantizedTensor': ...
+ @abstractmethod
+ def part(self, n_part: int) -> 'UnquantizedTensor': ...
+ @abstractmethod
  def to_ggml(self) -> 'GGMLCompatibleTensor': ...
 
 
@@ -345,6 +349,14 @@ def astype(self, data_type: DataType) -> Tensor:
  def to_ggml(self) -> 'UnquantizedTensor':
  return self
 
+ def permute_part(self, n_part: int, n_head: int) -> 'UnquantizedTensor':
+ r = self.ndarray.shape[0] // 3
+ return UnquantizedTensor(permute(self.ndarray[r * n_part : r * n_part + r, ...], n_head))
+
+ def part(self, n_part: int) -> 'UnquantizedTensor':
+ r = self.ndarray.shape[0] // 3
+ return UnquantizedTensor(self.ndarray[r * n_part : r * n_part + r, ...])
+
  def permute(self, n_head: int) -> 'UnquantizedTensor':
  return UnquantizedTensor(permute(self.ndarray, n_head))
 
@@ -642,6 +654,19 @@ def load() -> Tensor:
  return lazy_tensor.load().permute(n_head)
  return LazyTensor(load, lazy_tensor.shape, lazy_tensor.data_type, f'permute({n_head}) ' + lazy_tensor.description)
 
+def permute_part_lazy(lazy_tensor: LazyTensor, n_part: int, n_head: int) -> LazyTensor:
+ def load() -> Tensor:
+ return lazy_tensor.load().permute_part(n_part, n_head)
+ s = lazy_tensor.shape.copy()
+ s[0] = s[0] // 3
+ return LazyTensor(load, s, lazy_tensor.data_type, f'permute({n_head}) ' + lazy_tensor.description)
+
+def part_lazy(lazy_tensor: LazyTensor, n_part: int) -> LazyTensor:
+ def load() -> Tensor:
+ return lazy_tensor.load().part(n_part)
+ s = lazy_tensor.shape.copy()
+ s[0] = s[0] // 3
+ return LazyTensor(load, s, lazy_tensor.data_type, 'part ' + lazy_tensor.description)
 
 def convert_transformers_to_orig(model: LazyModel, params: Params) -> LazyModel:
  out: LazyModel = {}
@@ -650,11 +675,17 @@ def convert_transformers_to_orig(model: LazyModel, params: Params) -> LazyModel:
  out["output.weight"] = model["lm_head.weight"]
 
  for i in itertools.count():
- if f"model.layers.{i}.self_attn.q_proj.weight" not in model:
+ if f"model.layers.{i}.self_attn.q_proj.weight" in model:
+ out[f"layers.{i}.attention.wq.weight"] = permute_lazy(model[f"model.layers.{i}.self_attn.q_proj.weight"], params.n_head)
+ out[f"layers.{i}.attention.wk.weight"] = permute_lazy(model[f"model.layers.{i}.self_attn.k_proj.weight"], params.n_head)
+ out[f"layers.{i}.attention.wv.weight"] = model[f"model.layers.{i}.self_attn.v_proj.weight"]
+ elif f"model.layers.{i}.self_attn.W_pack.weight" in model:
+ out[f"layers.{i}.attention.wq.weight"] = permute_part_lazy(model[f"model.layers.{i}.self_attn.W_pack.weight"], 0, params.n_head)
+ out[f"layers.{i}.attention.wk.weight"] = permute_part_lazy(model[f"model.layers.{i}.self_attn.W_pack.weight"], 1, params.n_head)
+ out[f"layers.{i}.attention.wv.weight"] = part_lazy(model[f"model.layers.{i}.self_attn.W_pack.weight"], 2)
+ else:
  break
- out[f"layers.{i}.attention.wq.weight"] = permute_lazy(model[f"model.layers.{i}.self_attn.q_proj.weight"], params.n_head)
- out[f"layers.{i}.attention.wk.weight"] = permute_lazy(model[f"model.layers.{i}.self_attn.k_proj.weight"], params.n_head)
- out[f"layers.{i}.attention.wv.weight"] = model[f"model.layers.{i}.self_attn.v_proj.weight"]
+
  out[f"layers.{i}.attention.wo.weight"] = model[f"model.layers.{i}.self_attn.o_proj.weight"]
 
  out[f"layers.{i}.feed_forward.w1.weight"] = model[f"model.layers.{i}.mlp.gate_proj.weight"]

examples/embd-input/embd-input-lib.cpp

@@ -210,9 +210,12 @@ llama_token sampling_id(struct MyModel* mymodel) {
 const char * sampling(struct MyModel * mymodel) {
  llama_context * ctx = mymodel->ctx;
  int id = sampling_id(mymodel);
- std::string ret;
- if (id == llama_token_eos()) ret = "</s>";
- else ret = llama_token_to_str(ctx, id);
+ static std::string ret;
+ if (id == llama_token_eos()) {
+ ret = "</s>";
+ } else {
+ ret = llama_token_to_str(ctx, id);
+ }
  eval_id(mymodel, id);
  return ret.c_str();
 }

examples/embd-input/embd-input.h

@@ -5,7 +5,6 @@
 #include "llama.h"
 #include "build-info.h"
 
-
 extern "C" {
 
 typedef struct MyModel {
@@ -14,14 +13,13 @@ typedef struct MyModel {
  int n_past = 0;
 } MyModel;
 
-
 struct MyModel* create_mymodel(int argc, char ** argv);
 
 bool eval_float(void* model, float* input, int N);
 bool eval_tokens(void* model, std::vector<llama_token> tokens);
 bool eval_id(struct MyModel* mymodel, int id);
 bool eval_string(struct MyModel* mymodel, const char* str);
-const char* sampling(struct MyModel* mymodel);
+const char * sampling(struct MyModel* mymodel);
 llama_token sampling_id(struct MyModel* mymodel);
 void free_mymodel(struct MyModel* mymodel);
 

examples/train-text-from-scratch/train-text-from-scratch.cpp

@@ -2671,7 +2671,8 @@ struct train_params {
  const char * fn_checkpoint_out;
  const char * fn_model_out;
 
- int seed;
+ uint32_t seed;
+
  int n_ctx;
  int n_embd;
  int n_mult;

ggml-cuda.cu

@@ -214,6 +214,11 @@ static_assert(sizeof(block_q6_K) == sizeof(ggml_fp16_t) + 13*QK_K/16, "wrong q6_
 static_assert(K_QUANTS_PER_ITERATION == 1 || K_QUANTS_PER_ITERATION == 2, "K_QUANTS_PER_ITERATION must be 1 or 2");
 #endif
 
+struct ggml_tensor_extra_gpu {
+ void * data_device[GGML_CUDA_MAX_DEVICES]; // 1 pointer for each device for split tensors
+ cudaEvent_t events[GGML_CUDA_MAX_DEVICES]; // events for synchronizing multiple GPUs
+};
+
 static __global__ void add_f32(const float * x, const float * y, float * dst, const int k) {
  const int i = blockDim.x*blockIdx.x + threadIdx.x;
 
@@ -1995,7 +2000,6 @@ inline void ggml_cuda_op_add(
  } else {
  GGML_ASSERT(false);
  }
- CUDA_CHECK(cudaGetLastError());
 
  (void) src1;
  (void) dst;
@@ -2027,7 +2031,6 @@ inline void ggml_cuda_op_mul(
 
  // compute
  mul_f32_cuda(src0_ddf_i01, src1_ddf_i01, dst_ddf_i01, ne00, ne10, cudaStream_main);
- CUDA_CHECK(cudaGetLastError());
  }
 
  (void) dst;
@@ -2048,7 +2051,6 @@ inline void ggml_cuda_op_silu(
 
  // compute
  silu_f32_cuda(src0_ddf_i, dst_ddf_i, ne00*i01_diff, cudaStream_main);
- CUDA_CHECK(cudaGetLastError());
 
  (void) src1;
  (void) dst;
@@ -2071,7 +2073,6 @@ inline void ggml_cuda_op_rms_norm(
 
  // compute
  rms_norm_f32_cuda(src0_ddf_i, dst_ddf_i, ne00, i01_diff, cudaStream_main);
- CUDA_CHECK(cudaGetLastError());
 
  (void) src1;
  (void) dst;
@@ -2150,7 +2151,6 @@ inline void ggml_cuda_op_dequantize_mul_mat_vec(
  GGML_ASSERT(false);
  break;
  }
- CUDA_CHECK(cudaGetLastError());
 
 #ifdef GGML_CUDA_DMMV_F16
  if (src1_convert_f16) {
@@ -2230,7 +2230,6 @@ inline void ggml_cuda_op_rope(
 
  // compute
  rope_f32_cuda(src0_ddf_i, dst_ddf_i, ne00, i01_diff, p, theta_scale, cudaStream_main);
- CUDA_CHECK(cudaGetLastError());
 
  (void) dst;
  (void) src0_ddq_i;
@@ -2254,7 +2253,6 @@ inline void ggml_cuda_op_diag_mask_inf(
 
  // compute
  diag_mask_inf_f32_cuda(src0_ddf_i, dst_ddf_i, ne00, i01_diff, ne01, n_past, cudaStream_main);
- CUDA_CHECK(cudaGetLastError());
 
  (void) dst;
  (void) src0_ddq_i;
@@ -2276,7 +2274,6 @@ inline void ggml_cuda_op_soft_max(
 
  // compute
  soft_max_f32_cuda(src0_ddf_i, dst_ddf_i, ne00, i01_diff, cudaStream_main);
- CUDA_CHECK(cudaGetLastError());
 
  (void) src1;
  (void) dst;
@@ -2372,10 +2369,11 @@ static void ggml_cuda_op(const ggml_tensor * src0, const ggml_tensor * src1, ggm
  size_t src1_asf[GGML_CUDA_MAX_DEVICES] = {0};
  size_t dst_asf[GGML_CUDA_MAX_DEVICES] = {0};
 
- // if multiple GPUs are used they need to wait for the main GPU to finish
+ // if multiple devices are used they need to wait for the main device
+ // here an event is recorded that signifies that the main device has finished calculating the input data
  if (split && g_device_count > 1) {
  CUDA_CHECK(cudaSetDevice(g_main_device));
- CUDA_CHECK(cudaDeviceSynchronize());
+ CUDA_CHECK(cudaEventRecord(src0_extra->events[g_main_device], g_cudaStreams_main[g_main_device]));
  }
 
  for (int id = 0; id < g_device_count; ++id) {
@@ -2401,6 +2399,12 @@ static void ggml_cuda_op(const ggml_tensor * src0, const ggml_tensor * src1, ggm
  int64_t row_diff = row_high - row_low;
 
  cudaSetDevice(id);
+ cudaStream_t cudaStream_main = g_cudaStreams_main[id];
+
+ // wait for main GPU data if necessary
+ if (split && id != g_main_device) {
+ CUDA_CHECK(cudaStreamWaitEvent(cudaStream_main, src0_extra->events[g_main_device]));
+ }
 
  if (src0_on_device && src0_is_contiguous) {
  if (src0_is_f32) {
@@ -2476,8 +2480,6 @@ static void ggml_cuda_op(const ggml_tensor * src0, const ggml_tensor * src1, ggm
  }
  const int64_t i11 = i13*ne12 + i12;
 
- cudaStream_t cudaStream_main = g_cudaStreams_main[id];
-
  // for split tensors the data begins at i0 == i0_offset_low
  char * src0_ddq_i = src0_ddq[id] + (i0 - i0_offset_low)*src0_stride*src0_ts/src0_bs;
  float * src0_ddf_i = src0_ddf[id] + (i0 - i0_offset_low)*src0_stride;
@@ -2537,6 +2539,7 @@ static void ggml_cuda_op(const ggml_tensor * src0, const ggml_tensor * src1, ggm
 
  // do the computation
  op(src0, src1, dst, src0_ddq_i, src0_ddf_i, src1_ddf_i, dst_ddf_i, i02, i01_low, i01_high, i11, cudaStream_main);
+ CUDA_CHECK(cudaGetLastError());
 
  // copy dst to host or other device if necessary
  if (!dst_on_device) {
@@ -2566,6 +2569,11 @@ static void ggml_cuda_op(const ggml_tensor * src0, const ggml_tensor * src1, ggm
  CUDA_CHECK(cudaMemcpyAsync(dhf_dst_i, dst_ddf_i, dst_stride*sizeof(float), kind, cudaStream_main));
  }
  }
+
+ // signify to main device that other device is done
+ if (split && g_device_count > 1 && id != g_main_device) {
+ CUDA_CHECK(cudaEventRecord(src0_extra->events[id], cudaStream_main));
+ }
  }
  }
  }
@@ -2577,7 +2585,6 @@ static void ggml_cuda_op(const ggml_tensor * src0, const ggml_tensor * src1, ggm
  }
 
  CUDA_CHECK(cudaSetDevice(id));
- CUDA_CHECK(cudaDeviceSynchronize());
 
  if (src0_asq[id] > 0) {
  ggml_cuda_pool_free(src0_ddq[id], src0_asq[id]);
@@ -2592,6 +2599,21 @@ static void ggml_cuda_op(const ggml_tensor * src0, const ggml_tensor * src1, ggm
  ggml_cuda_pool_free(dst_ddf[id], dst_asf[id]);
  }
  }
+
+ // main device waits for all other devices to be finished
+ if (split && g_device_count > 1) {
+ CUDA_CHECK(cudaSetDevice(g_main_device));
+ for (int id = 0; id < g_device_count; ++id) {
+ if (id != g_main_device) {
+ CUDA_CHECK(cudaStreamWaitEvent(g_cudaStreams_main[g_main_device], src0_extra->events[id]));
+ }
+ }
+ }
+
+ if (dst->backend == GGML_BACKEND_CPU) {
+ CUDA_CHECK(cudaSetDevice(g_main_device));
+ CUDA_CHECK(cudaDeviceSynchronize());
+ }
 }
 
 void ggml_cuda_add(const ggml_tensor * src0, const ggml_tensor * src1, ggml_tensor * dst) {
@@ -2831,6 +2853,10 @@ void ggml_cuda_transform_tensor(void * data, struct ggml_tensor * tensor) {
  cudaMemcpy(buf, buf_host, size, cudaMemcpyHostToDevice);
 
  extra->data_device[id] = buf;
+
+ if (backend == GGML_BACKEND_GPU_SPLIT) {
+ CUDA_CHECK(cudaEventCreateWithFlags(&extra->events[id], cudaEventDisableTiming));
+ }
  }
 
  tensor->extra = extra;
@@ -2844,12 +2870,15 @@ void ggml_cuda_free_data(struct ggml_tensor * tensor) {
  ggml_tensor_extra_gpu * extra = (ggml_tensor_extra_gpu *) tensor->extra;
 
  for (int id = 0; id < g_device_count; ++id) {
- if (extra->data_device[id] == nullptr) {
- continue;
+ if (extra->data_device[id] != nullptr) {
+ CUDA_CHECK(cudaSetDevice(id));
+ CUDA_CHECK(cudaFree(extra->data_device[id]));
  }
 
- CUDA_CHECK(cudaSetDevice(id));
- CUDA_CHECK(cudaFree(extra->data_device[id]));
+ if (extra->events[id] != nullptr) {
+ CUDA_CHECK(cudaSetDevice(id));
+ CUDA_CHECK(cudaEventDestroy(extra->events[id]));
+ }
  }
 
  delete extra;

ggml-cuda.h

@@ -8,10 +8,6 @@ extern "C" {
 
 #define GGML_CUDA_MAX_DEVICES 16
 
-struct ggml_tensor_extra_gpu {
- void * data_device[GGML_CUDA_MAX_DEVICES]; // 1 pointer for each device for split tensors
-};
-
 void ggml_init_cublas(void);
 void ggml_cuda_set_tensor_split(const float * tensor_split);
 

ggml-metal.m

@@ -202,7 +202,9 @@ @implementation GGMLMetalClass
 
 void ggml_metal_free(struct ggml_metal_context * ctx) {
  fprintf(stderr, "%s: deallocating\n", __func__);
-
+ for (int i = 0; i < ctx->n_buffers; ++i) {
+ [ctx->buffers[i].metal release];
+ }
  free(ctx);
 }