/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp

Bug Summary

File:	root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp
Warning:	line 2049, column 9 Access to field 'view_src' results in a dereference of a null pointer (loaded from variable 'dst')

Annotated Source Code

Press '?' to see keyboard shortcuts

Show analyzer invocation

clang -cc1 -cc1 -triple x86_64-pc-linux-gnu -O2 -analyze -disable-free -clear-ast-before-backend -disable-llvm-verifier -discard-value-names -main-file-name ggml-backend.cpp -analyzer-checker=core -analyzer-checker=apiModeling -analyzer-checker=unix -analyzer-checker=deadcode -analyzer-checker=cplusplus -analyzer-checker=security.insecureAPI.UncheckedReturn -analyzer-checker=security.insecureAPI.getpw -analyzer-checker=security.insecureAPI.gets -analyzer-checker=security.insecureAPI.mktemp -analyzer-checker=security.insecureAPI.mkstemp -analyzer-checker=security.insecureAPI.vfork -analyzer-checker=nullability.NullPassedToNonnull -analyzer-checker=nullability.NullReturnedFromNonnull -analyzer-output plist -w -setup-static-analyzer -analyzer-config-compatibility-mode=true -mrelocation-model pic -pic-level 2 -fhalf-no-semantic-interposition -mframe-pointer=all -relaxed-aliasing -ffp-contract=off -fno-rounding-math -mconstructor-aliases -funwind-tables=2 -target-cpu x86-64 -target-feature +avx -target-feature +avx2 -target-feature +bmi2 -target-feature +f16c -target-feature +fma -target-feature +sse4.2 -tune-cpu generic -debugger-tuning=gdb -fdebug-compilation-dir=/root/firefox-clang/obj-x86_64-pc-linux-gnu/third_party/llama.cpp -fcoverage-compilation-dir=/root/firefox-clang/obj-x86_64-pc-linux-gnu/third_party/llama.cpp -resource-dir /usr/lib/llvm-23/lib/clang/23 -include /root/firefox-clang/config/gcc_hidden.h -include /root/firefox-clang/obj-x86_64-pc-linux-gnu/mozilla-config.h -I /root/firefox-clang/obj-x86_64-pc-linux-gnu/dist/stl_wrappers -D _GLIBCXX_ASSERTIONS=1 -I /root/firefox-clang/obj-x86_64-pc-linux-gnu/dist/system_wrappers -U _FORTIFY_SOURCE -D _FORTIFY_SOURCE=2 -D DEBUG=1 -D _GNU_SOURCE=1 -D GGML_USE_CPU=1 -D GGML_VERSION="GGML_VERSION" -D GGML_COMMIT="GGML_COMMIT" -D GGML_SHARED=1 -D LLAMA_SHARED=1 -D GGML_BUILD=1 -D LLAMA_BUILD=1 -D GGML_BACKEND_SHARED=1 -D GGML_BACKEND_BUILD=1 -D MOZ_HAS_MOZGLUE -I /root/firefox-clang/third_party/llama.cpp -I /root/firefox-clang/obj-x86_64-pc-linux-gnu/third_party/llama.cpp -I /root/firefox-clang/third_party/llama.cpp/ggml -I /root/firefox-clang/third_party/llama.cpp/ggml/include -I /root/firefox-clang/third_party/llama.cpp/ggml/src -I /root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-cpu -I /root/firefox-clang/third_party/llama.cpp/include -I /root/firefox-clang/third_party/llama.cpp/src -I /root/firefox-clang/obj-x86_64-pc-linux-gnu/dist/include -I /root/firefox-clang/obj-x86_64-pc-linux-gnu/dist/include/nspr -I /root/firefox-clang/obj-x86_64-pc-linux-gnu/dist/include/nss -D MOZILLA_CLIENT -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/16/../../../../include/c++/16 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/16/../../../../include/x86_64-linux-gnu/c++/16 -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/16/../../../../include/c++/16/backward -internal-isystem /usr/lib/llvm-23/lib/clang/23/include -internal-isystem /usr/local/include -internal-isystem /usr/lib/gcc/x86_64-linux-gnu/16/../../../../x86_64-linux-gnu/include -internal-externc-isystem /usr/include/x86_64-linux-gnu -internal-externc-isystem /include -internal-externc-isystem /usr/include -Wno-error=pessimizing-move -Wno-error=large-by-value-copy=128 -Wno-error=implicit-int-float-conversion -Wno-error=thread-safety-analysis -Wno-error=tautological-type-limit-compare -Wno-invalid-offsetof -Wno-range-loop-analysis -Wno-deprecated-anon-enum-enum-conversion -Wno-deprecated-enum-enum-conversion -Wno-inline-new-delete -Wno-error=deprecated-declarations -Wno-error=array-bounds -Wno-error=free-nonheap-object -Wno-error=atomic-alignment -Wno-error=deprecated-builtins -Wno-psabi -Wno-error=builtin-macro-redefined -Wno-vla-cxx-extension -Wno-unknown-warning-option -Wno-character-conversion -Wno-sign-compare -Wno-unused-function -Wno-tautological-unsigned-enum-zero-compare -Wno-implicit-fallthrough -Wno-unreachable-code -std=gnu++20 -fdeprecated-macro -ferror-limit 19 -fstrict-flex-arrays=1 -stack-protector 2 -fstack-clash-protection -ftrivial-auto-var-init=pattern -fno-rtti -fgnuc-version=4.2.1 -fno-implicit-modules -fskip-odr-check-in-gmf -fno-sized-deallocation -fno-aligned-allocation -fdiagnostics-absolute-paths -vectorize-loops -vectorize-slp -analyzer-checker optin.performance.Padding -analyzer-output=html -analyzer-config stable-report-filename=true -mllvm -dwarf-linkage-names=Abstract -faddrsig -fdwarf2-cfi-asm -o /tmp/scan-build-2026-09-01-224014-2642839-1 -x c++ /root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp

/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp

→

1// Note: porting this file to C++ is a work in progress

3#ifdef _WIN32
4#define WIN32_LEAN_AND_MEAN
5#ifndef NOMINMAX
6#   define NOMINMAX
7#endif
8#include <windows.h>
9#endif

11#include "ggml-backend.h"
12#include "ggml-backend-impl.h"
13#include "ggml-alloc.h"
14#include "ggml-impl.h"

16#include <assert.h>
17#include <limits.h>
18#include <stdarg.h>
19#include <stdio.h>
20#include <stdlib.h>
21#include <string.h>
22#include <algorithm>
23#include <vector>

25#ifdef __APPLE__
26#include <sys/types.h>
27#include <sys/sysctl.h>
28#endif


31// backend buffer type

33const char * ggml_backend_buft_name(ggml_backend_buffer_type_t buft) {
  GGML_ASSERT(buft)if (!(buft)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 34, "GGML_ASSERT(%s) failed", "buft");
  return buft->iface.get_name(buft);
36}

38ggml_backend_buffer_t ggml_backend_buft_alloc_buffer(ggml_backend_buffer_type_t buft, size_t size) {
  GGML_ASSERT(buft)if (!(buft)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 39, "GGML_ASSERT(%s) failed", "buft");
  if (size == 0) {
      // return a dummy buffer for zero-sized allocations
      return ggml_backend_buffer_init(buft, {}, NULL__null, 0);
  }
  return buft->iface.alloc_buffer(buft, size);
45}

47size_t ggml_backend_buft_get_alignment(ggml_backend_buffer_type_t buft) {
  GGML_ASSERT(buft)if (!(buft)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 48, "GGML_ASSERT(%s) failed", "buft");
  return buft->iface.get_alignment(buft);
50}

52size_t ggml_backend_buft_get_max_size(ggml_backend_buffer_type_t buft) {
  GGML_ASSERT(buft)if (!(buft)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 53, "GGML_ASSERT(%s) failed", "buft");
  // get_max_size is optional, defaults to SIZE_MAX
  if (buft->iface.get_max_size) {
      return buft->iface.get_max_size(buft);
  }
  return SIZE_MAX(18446744073709551615UL);
59}

61size_t ggml_backend_buft_get_alloc_size(ggml_backend_buffer_type_t buft, const struct ggml_tensor * tensor) {
  GGML_ASSERT(buft)if (!(buft)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 62, "GGML_ASSERT(%s) failed", "buft");
  // get_alloc_size is optional, defaults to ggml_nbytes
  if (buft->iface.get_alloc_size) {
      size_t size = buft->iface.get_alloc_size(buft, tensor);
      assert(size >= ggml_nbytes(tensor))(static_cast <bool> (size >= ggml_nbytes(tensor)) ? void
 (0) : __assert_fail ("size >= ggml_nbytes(tensor)", __builtin_FILE
 (), __builtin_LINE (), __extension__ __PRETTY_FUNCTION__));
      return size;
  }
  return ggml_nbytes(tensor);
70}

72bool ggml_backend_buft_is_host(ggml_backend_buffer_type_t buft) {
  GGML_ASSERT(buft)if (!(buft)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 73, "GGML_ASSERT(%s) failed", "buft");
  if (buft->iface.is_host) {
      return buft->iface.is_host(buft);
  }
  return false;
78}

80ggml_backend_dev_t ggml_backend_buft_get_device(ggml_backend_buffer_type_t buft) {
  GGML_ASSERT(buft)if (!(buft)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 81, "GGML_ASSERT(%s) failed", "buft");
  return buft->device;
83}

85// backend buffer

87ggml_backend_buffer_t ggml_backend_buffer_init(
             ggml_backend_buffer_type_t buft,
      struct ggml_backend_buffer_i      iface,
             void *                     context,
             size_t                     size) {
  ggml_backend_buffer_t buffer = new ggml_backend_buffer {
      /* .interface = */ iface,
      /* .buft      = */ buft,
      /* .context   = */ context,
      /* .size      = */ size,
      /* .usage     = */ GGML_BACKEND_BUFFER_USAGE_ANY
  };

  return buffer;
101}

103const char * ggml_backend_buffer_name(ggml_backend_buffer_t buffer) {
  return ggml_backend_buft_name(ggml_backend_buffer_get_type(buffer));
105}

107void ggml_backend_buffer_free(ggml_backend_buffer_t buffer) {
  if (buffer == NULL__null) {
      return;
  }

  if (buffer->iface.free_buffer != NULL__null) {
      buffer->iface.free_buffer(buffer);
  }
  delete buffer;
116}

118size_t ggml_backend_buffer_get_size(ggml_backend_buffer_t buffer) {
  GGML_ASSERT(buffer)if (!(buffer)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 119, "GGML_ASSERT(%s) failed", "buffer");
  return buffer->size;
121}

123void * ggml_backend_buffer_get_base(ggml_backend_buffer_t buffer) {
  GGML_ASSERT(buffer)if (!(buffer)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 124, "GGML_ASSERT(%s) failed", "buffer");
  // get_base is optional if the buffer is zero-sized
  if (!ggml_backend_buffer_is_meta(buffer) && buffer->size == 0) {
      return NULL__null;
  }

  // FIXME JG: a multi_buffer has a non-zero size, according to the above comment get_base is not optional,
  //     I don't know whether the above comment is correct
  if (!buffer->iface.get_base) {
      return NULL__null;
  }

  void * base = buffer->iface.get_base(buffer);

  GGML_ASSERT(base != NULL && "backend buffer base cannot be NULL")if (!(base != __null && "backend buffer base cannot be NULL"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 138, "GGML_ASSERT(%s) failed", "base != NULL && \"backend buffer base cannot be NULL\""
);

  return base;
141}

143enum ggml_status ggml_backend_buffer_init_tensor(ggml_backend_buffer_t buffer, struct ggml_tensor * tensor) {
  GGML_ASSERT(buffer)if (!(buffer)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 144, "GGML_ASSERT(%s) failed", "buffer");
  // init_tensor is optional
  if (buffer->iface.init_tensor) {
      return buffer->iface.init_tensor(buffer, tensor);
  }
  return GGML_STATUS_SUCCESS;
150}

152void ggml_backend_buffer_clear(ggml_backend_buffer_t buffer, uint8_t value) {
  GGML_ASSERT(buffer)if (!(buffer)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 153, "GGML_ASSERT(%s) failed", "buffer");
  // clear is optional if the buffer is zero-sized
  if (buffer->size == 0) {
      return;
  }

  buffer->iface.clear(buffer, value);
160}

162size_t ggml_backend_buffer_get_alignment(ggml_backend_buffer_t buffer) {
  return ggml_backend_buft_get_alignment(ggml_backend_buffer_get_type(buffer));
164}

166size_t ggml_backend_buffer_get_max_size(ggml_backend_buffer_t buffer) {
  return ggml_backend_buft_get_max_size(ggml_backend_buffer_get_type(buffer));
168}

170size_t ggml_backend_buffer_get_alloc_size(ggml_backend_buffer_t buffer, const struct ggml_tensor * tensor) {
  return ggml_backend_buft_get_alloc_size(ggml_backend_buffer_get_type(buffer), tensor);
172}

174bool ggml_backend_buffer_is_host(ggml_backend_buffer_t buffer) {
  return ggml_backend_buft_is_host(ggml_backend_buffer_get_type(buffer));
176}

178void ggml_backend_buffer_set_usage(ggml_backend_buffer_t buffer, enum ggml_backend_buffer_usage usage) {
  GGML_ASSERT(buffer)if (!(buffer)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 179, "GGML_ASSERT(%s) failed", "buffer");
  buffer->usage = usage;

  // FIXME: add a generic callback to the buffer interface
  if (ggml_backend_buffer_is_multi_buffer(buffer)) {
      ggml_backend_multi_buffer_set_usage(buffer, usage);
  }
186}

188enum ggml_backend_buffer_usage ggml_backend_buffer_get_usage(ggml_backend_buffer_t buffer) {
  GGML_ASSERT(buffer)if (!(buffer)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 189, "GGML_ASSERT(%s) failed", "buffer");
  return buffer->usage;
191}

193ggml_backend_buffer_type_t ggml_backend_buffer_get_type(ggml_backend_buffer_t buffer) {
  GGML_ASSERT(buffer)if (!(buffer)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 194, "GGML_ASSERT(%s) failed", "buffer");
  return buffer->buft;
196}

198void ggml_backend_buffer_reset(ggml_backend_buffer_t buffer) {
  GGML_ASSERT(buffer)if (!(buffer)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 199, "GGML_ASSERT(%s) failed", "buffer");
  if (buffer->iface.reset) {
      buffer->iface.reset(buffer);
  }
203}

205bool ggml_backend_buffer_copy_tensor(const struct ggml_tensor * src, struct ggml_tensor * dst) {
  ggml_backend_buffer_t dst_buf = dst->view_src ? dst->view_src->buffer : dst->buffer;
  if (dst_buf->iface.cpy_tensor) {
      return dst_buf->iface.cpy_tensor(dst_buf, src, dst);
  }
  return false;
211}

213// backend

215ggml_guid_t ggml_backend_guid(ggml_backend_t backend) {
  if (backend == NULL__null) {
      return NULL__null;
  }
  return backend->guid;
220}

222const char * ggml_backend_name(ggml_backend_t backend) {
  if (backend == NULL__null) {
      return "NULL";
  }
  return backend->iface.get_name(backend);
227}

229void ggml_backend_free(ggml_backend_t backend) {
  if (backend == NULL__null) {
      return;
  }

  backend->iface.free(backend);
235}

237ggml_backend_buffer_type_t ggml_backend_get_default_buffer_type(ggml_backend_t backend) {
  GGML_ASSERT(backend)if (!(backend)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 238, "GGML_ASSERT(%s) failed", "backend");
  return ggml_backend_dev_buffer_type(backend->device);
240}

242ggml_backend_buffer_t ggml_backend_alloc_buffer(ggml_backend_t backend, size_t size) {
  return ggml_backend_buft_alloc_buffer(ggml_backend_get_default_buffer_type(backend), size);
244}

246size_t ggml_backend_get_alignment(ggml_backend_t backend) {
  return ggml_backend_buft_get_alignment(ggml_backend_get_default_buffer_type(backend));
248}

250size_t ggml_backend_get_max_size(ggml_backend_t backend) {
  return ggml_backend_buft_get_max_size(ggml_backend_get_default_buffer_type(backend));
252}

254void ggml_backend_tensor_set_async(ggml_backend_t backend, struct ggml_tensor * tensor, const void * data, size_t offset, size_t size) {
  GGML_ASSERT(backend)if (!(backend)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 255, "GGML_ASSERT(%s) failed", "backend");
  GGML_ASSERT(tensor)if (!(tensor)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 256, "GGML_ASSERT(%s) failed", "tensor");
  GGML_ASSERT(tensor->data != NULL && "tensor not allocated")if (!(tensor->data != __null && "tensor not allocated"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 257, "GGML_ASSERT(%s) failed", "tensor->data != NULL && \"tensor not allocated\""
);
  GGML_ASSERT(offset + size <= ggml_nbytes(tensor) && "tensor write out of bounds")if (!(offset + size <= ggml_nbytes(tensor) && "tensor write out of bounds"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 258, "GGML_ASSERT(%s) failed", "offset + size <= ggml_nbytes(tensor) && \"tensor write out of bounds\""
);

  if (backend->iface.set_tensor_async == NULL__null) {
      ggml_backend_synchronize(backend);
      ggml_backend_tensor_set(tensor, data, offset, size);
  } else {
      backend->iface.set_tensor_async(backend, tensor, data, offset, size);
  }
266}

268void ggml_backend_tensor_get_async(ggml_backend_t backend, const struct ggml_tensor * tensor, void * data, size_t offset, size_t size) {
  GGML_ASSERT(backend)if (!(backend)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 269, "GGML_ASSERT(%s) failed", "backend");
  GGML_ASSERT(tensor)if (!(tensor)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 270, "GGML_ASSERT(%s) failed", "tensor");
  GGML_ASSERT(tensor->data != NULL && "tensor not allocated")if (!(tensor->data != __null && "tensor not allocated"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 271, "GGML_ASSERT(%s) failed", "tensor->data != NULL && \"tensor not allocated\""
);
  GGML_ASSERT(offset + size <= ggml_nbytes(tensor) && "tensor read out of bounds")if (!(offset + size <= ggml_nbytes(tensor) && "tensor read out of bounds"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 272, "GGML_ASSERT(%s) failed", "offset + size <= ggml_nbytes(tensor) && \"tensor read out of bounds\""
);

  if (backend->iface.get_tensor_async == NULL__null) {
      ggml_backend_synchronize(backend);
      ggml_backend_tensor_get(tensor, data, offset, size);
  } else {
      backend->iface.get_tensor_async(backend, tensor, data, offset, size);
  }
280}

282void ggml_backend_tensor_set_2d_async(ggml_backend_t backend, struct ggml_tensor * tensor, const void * data, size_t offset, size_t size,
          size_t n_copies, size_t stride_tensor, size_t stride_data) {
  GGML_ASSERT(backend)if (!(backend)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 284, "GGML_ASSERT(%s) failed", "backend");
  GGML_ASSERT(tensor)if (!(tensor)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 285, "GGML_ASSERT(%s) failed", "tensor");
  GGML_ASSERT(tensor->data != NULL && "tensor not allocated")if (!(tensor->data != __null && "tensor not allocated"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 286, "GGML_ASSERT(%s) failed", "tensor->data != NULL && \"tensor not allocated\""
);

  if (n_copies <= 1 || backend->iface.set_tensor_2d_async == NULL__null) {
      for (size_t i = 0; i < n_copies; i++) {
          ggml_backend_tensor_set_async(backend, tensor, (const char *) data + i*stride_data, offset + i*stride_tensor, size);
      }
      return;
  }
  if (size == 0) {
      return;
  }

  GGML_ASSERT(tensor->data != NULL && "tensor not allocated")if (!(tensor->data != __null && "tensor not allocated"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 298, "GGML_ASSERT(%s) failed", "tensor->data != NULL && \"tensor not allocated\""
);
  GGML_ASSERT(offset + (n_copies-1)*stride_tensor + size <= ggml_nbytes(tensor) && "tensor write out of bounds")if (!(offset + (n_copies-1)*stride_tensor + size <= ggml_nbytes
(tensor) && "tensor write out of bounds")) ggml_abort
("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 299, "GGML_ASSERT(%s) failed", "offset + (n_copies-1)*stride_tensor + size <= ggml_nbytes(tensor) && \"tensor write out of bounds\""
);
  backend->iface.set_tensor_2d_async(backend, tensor, data, offset, size, n_copies, stride_tensor, stride_data);
301}

303void ggml_backend_tensor_get_2d_async(ggml_backend_t backend, const struct ggml_tensor * tensor, void * data, size_t offset, size_t size,
          size_t n_copies, size_t stride_tensor, size_t stride_data) {
  GGML_ASSERT(backend)if (!(backend)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 305, "GGML_ASSERT(%s) failed", "backend");
  GGML_ASSERT(tensor)if (!(tensor)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 306, "GGML_ASSERT(%s) failed", "tensor");
  GGML_ASSERT(tensor->data != NULL && "tensor not allocated")if (!(tensor->data != __null && "tensor not allocated"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 307, "GGML_ASSERT(%s) failed", "tensor->data != NULL && \"tensor not allocated\""
);

  if (n_copies <= 1 || backend->iface.get_tensor_2d_async == NULL__null) {
      for (size_t i = 0; i < n_copies; i++) {
          ggml_backend_tensor_get_async(backend, tensor, (char *) data + i*stride_data, offset + i*stride_tensor, size);
      }
      return;
  }
  if (size == 0) {
      return;
  }

  GGML_ASSERT(tensor->data != NULL && "tensor not allocated")if (!(tensor->data != __null && "tensor not allocated"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 319, "GGML_ASSERT(%s) failed", "tensor->data != NULL && \"tensor not allocated\""
);
  GGML_ASSERT(offset + (n_copies-1)*stride_tensor + size <= ggml_nbytes(tensor) && "tensor read out of bounds")if (!(offset + (n_copies-1)*stride_tensor + size <= ggml_nbytes
(tensor) && "tensor read out of bounds")) ggml_abort(
"/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 320, "GGML_ASSERT(%s) failed", "offset + (n_copies-1)*stride_tensor + size <= ggml_nbytes(tensor) && \"tensor read out of bounds\""
);
  backend->iface.get_tensor_2d_async(backend, tensor, data, offset, size, n_copies, stride_tensor, stride_data);
322}

324void ggml_backend_tensor_set(struct ggml_tensor * tensor, const void * data, size_t offset, size_t size) {
  GGML_ASSERT(tensor)if (!(tensor)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 325, "GGML_ASSERT(%s) failed", "tensor");
  ggml_backend_buffer_t buf = tensor->view_src ? tensor->view_src->buffer : tensor->buffer;
  GGML_ASSERT(buf != NULL && "tensor buffer not set")if (!(buf != __null && "tensor buffer not set")) ggml_abort
("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 327, "GGML_ASSERT(%s) failed", "buf != NULL && \"tensor buffer not set\""
);

  if (size == 0) {
      return;
  }

  GGML_ASSERT(tensor->data != NULL && "tensor not allocated")if (!(tensor->data != __null && "tensor not allocated"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 333, "GGML_ASSERT(%s) failed", "tensor->data != NULL && \"tensor not allocated\""
);
  GGML_ASSERT(offset + size <= ggml_nbytes(tensor) && "tensor write out of bounds")if (!(offset + size <= ggml_nbytes(tensor) && "tensor write out of bounds"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 334, "GGML_ASSERT(%s) failed", "offset + size <= ggml_nbytes(tensor) && \"tensor write out of bounds\""
);

  buf->iface.set_tensor(buf, tensor, data, offset, size);
337}

339void ggml_backend_tensor_get(const struct ggml_tensor * tensor, void * data, size_t offset, size_t size) {
  GGML_ASSERT(tensor)if (!(tensor)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 340, "GGML_ASSERT(%s) failed", "tensor");
  ggml_backend_buffer_t buf = tensor->view_src ? tensor->view_src->buffer : tensor->buffer;
  GGML_ASSERT(buf != NULL && "tensor buffer not set")if (!(buf != __null && "tensor buffer not set")) ggml_abort
("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 342, "GGML_ASSERT(%s) failed", "buf != NULL && \"tensor buffer not set\""
);

  if (size == 0) {
      return;
  }

  GGML_ASSERT(tensor->data != NULL && "tensor not allocated")if (!(tensor->data != __null && "tensor not allocated"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 348, "GGML_ASSERT(%s) failed", "tensor->data != NULL && \"tensor not allocated\""
);
  GGML_ASSERT(offset + size <= ggml_nbytes(tensor) && "tensor read out of bounds")if (!(offset + size <= ggml_nbytes(tensor) && "tensor read out of bounds"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 349, "GGML_ASSERT(%s) failed", "offset + size <= ggml_nbytes(tensor) && \"tensor read out of bounds\""
);

  buf->iface.get_tensor(buf, tensor, data, offset, size);
352}

354void ggml_backend_tensor_set_2d(struct ggml_tensor * tensor, const void * data, size_t offset, size_t size,
          size_t n_copies, size_t stride_tensor, size_t stride_data) {
  GGML_ASSERT(tensor)if (!(tensor)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 356, "GGML_ASSERT(%s) failed", "tensor");
  ggml_backend_buffer_t buf = tensor->view_src ? tensor->view_src->buffer : tensor->buffer;
  GGML_ASSERT(buf != NULL && "tensor buffer not set")if (!(buf != __null && "tensor buffer not set")) ggml_abort
("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 358, "GGML_ASSERT(%s) failed", "buf != NULL && \"tensor buffer not set\""
);

  if (n_copies <= 1 || buf->iface.set_tensor_2d == NULL__null) {
      for (size_t i = 0; i < n_copies; i++) {
          ggml_backend_tensor_set(tensor, (const char *) data + i*stride_data, offset + i*stride_tensor, size);
      }
      return;
  }
  if (size == 0) {
      return;
  }

  GGML_ASSERT(tensor->data != NULL && "tensor not allocated")if (!(tensor->data != __null && "tensor not allocated"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 370, "GGML_ASSERT(%s) failed", "tensor->data != NULL && \"tensor not allocated\""
);
  GGML_ASSERT(offset + (n_copies-1)*stride_tensor + size <= ggml_nbytes(tensor) && "tensor write out of bounds")if (!(offset + (n_copies-1)*stride_tensor + size <= ggml_nbytes
(tensor) && "tensor write out of bounds")) ggml_abort
("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 371, "GGML_ASSERT(%s) failed", "offset + (n_copies-1)*stride_tensor + size <= ggml_nbytes(tensor) && \"tensor write out of bounds\""
);

  buf->iface.set_tensor_2d(buf, tensor, data, offset, size, n_copies, stride_tensor, stride_data);
374}

376void ggml_backend_tensor_get_2d(const struct ggml_tensor * tensor, void * data, size_t offset, size_t size,
          size_t n_copies, size_t stride_tensor, size_t stride_data) {
  GGML_ASSERT(tensor)if (!(tensor)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 378, "GGML_ASSERT(%s) failed", "tensor");
  ggml_backend_buffer_t buf = tensor->view_src ? tensor->view_src->buffer : tensor->buffer;
  GGML_ASSERT(buf != NULL && "tensor buffer not set")if (!(buf != __null && "tensor buffer not set")) ggml_abort
("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 380, "GGML_ASSERT(%s) failed", "buf != NULL && \"tensor buffer not set\""
);

  if (n_copies <= 1 || buf->iface.get_tensor_2d == NULL__null) {
      for (size_t i = 0; i < n_copies; i++) {
          ggml_backend_tensor_get(tensor, (char *) data + i*stride_data, offset + i*stride_tensor, size);
      }
      return;
  }
  if (size == 0) {
      return;
  }

  GGML_ASSERT(tensor->data != NULL && "tensor not allocated")if (!(tensor->data != __null && "tensor not allocated"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 392, "GGML_ASSERT(%s) failed", "tensor->data != NULL && \"tensor not allocated\""
);
  GGML_ASSERT(offset + (n_copies-1)*stride_tensor + size <= ggml_nbytes(tensor) && "tensor read out of bounds")if (!(offset + (n_copies-1)*stride_tensor + size <= ggml_nbytes
(tensor) && "tensor read out of bounds")) ggml_abort(
"/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 393, "GGML_ASSERT(%s) failed", "offset + (n_copies-1)*stride_tensor + size <= ggml_nbytes(tensor) && \"tensor read out of bounds\""
);

  buf->iface.get_tensor_2d(buf, tensor, data, offset, size, n_copies, stride_tensor, stride_data);
396}

398void ggml_backend_tensor_memset(struct ggml_tensor * tensor, uint8_t value, size_t offset, size_t size) {
  GGML_ASSERT(tensor)if (!(tensor)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 399, "GGML_ASSERT(%s) failed", "tensor");
  ggml_backend_buffer_t buf = tensor->view_src ? tensor->view_src->buffer : tensor->buffer;

  if (size == 0) {
      return;
  }

  GGML_ASSERT(buf != NULL && "tensor buffer not set")if (!(buf != __null && "tensor buffer not set")) ggml_abort
("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 406, "GGML_ASSERT(%s) failed", "buf != NULL && \"tensor buffer not set\""
);
  GGML_ASSERT(tensor->data != NULL && "tensor not allocated")if (!(tensor->data != __null && "tensor not allocated"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 407, "GGML_ASSERT(%s) failed", "tensor->data != NULL && \"tensor not allocated\""
);
  GGML_ASSERT(offset + size <= ggml_nbytes(tensor) && "tensor write out of bounds")if (!(offset + size <= ggml_nbytes(tensor) && "tensor write out of bounds"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 408, "GGML_ASSERT(%s) failed", "offset + size <= ggml_nbytes(tensor) && \"tensor write out of bounds\""
);
  GGML_ASSERT(buf->iface.memset_tensor != NULL && "memset not implemented by backend buffer")if (!(buf->iface.memset_tensor != __null && "memset not implemented by backend buffer"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 409, "GGML_ASSERT(%s) failed", "buf->iface.memset_tensor != NULL && \"memset not implemented by backend buffer\""
);

  buf->iface.memset_tensor(buf, tensor, value, offset, size);
412}

414void ggml_backend_synchronize(ggml_backend_t backend) {
  GGML_ASSERT(backend)if (!(backend)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 415, "GGML_ASSERT(%s) failed", "backend");
  if (backend->iface.synchronize == NULL__null) {
      return;
  }

  backend->iface.synchronize(backend);
421}

423ggml_backend_graph_plan_t ggml_backend_graph_plan_create(ggml_backend_t backend, struct ggml_cgraph * cgraph) {
  GGML_ASSERT(backend)if (!(backend)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 424, "GGML_ASSERT(%s) failed", "backend");
  GGML_ASSERT(backend->iface.graph_plan_create != NULL)if (!(backend->iface.graph_plan_create != __null)) ggml_abort
("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 425, "GGML_ASSERT(%s) failed", "backend->iface.graph_plan_create != NULL"
);

  return backend->iface.graph_plan_create(backend, cgraph);
428}

430void ggml_backend_graph_plan_free(ggml_backend_t backend, ggml_backend_graph_plan_t plan) {
  GGML_ASSERT(backend)if (!(backend)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 431, "GGML_ASSERT(%s) failed", "backend");
  GGML_ASSERT(backend->iface.graph_plan_free != NULL)if (!(backend->iface.graph_plan_free != __null)) ggml_abort
("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 432, "GGML_ASSERT(%s) failed", "backend->iface.graph_plan_free != NULL"
);

  backend->iface.graph_plan_free(backend, plan);
435}

437enum ggml_status ggml_backend_graph_plan_compute(ggml_backend_t backend, ggml_backend_graph_plan_t plan) {
  GGML_ASSERT(backend)if (!(backend)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 438, "GGML_ASSERT(%s) failed", "backend");
  GGML_ASSERT(backend->iface.graph_plan_compute != NULL)if (!(backend->iface.graph_plan_compute != __null)) ggml_abort
("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 439, "GGML_ASSERT(%s) failed", "backend->iface.graph_plan_compute != NULL"
);

  return backend->iface.graph_plan_compute(backend, plan);
442}

444enum ggml_status ggml_backend_graph_compute(ggml_backend_t backend, struct ggml_cgraph * cgraph) {
  enum ggml_status err = ggml_backend_graph_compute_async(backend, cgraph);
  ggml_backend_synchronize(backend);
  return err;
448}

450enum ggml_status ggml_backend_graph_compute_async(ggml_backend_t backend, struct ggml_cgraph * cgraph) {
  GGML_ASSERT(backend)if (!(backend)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 451, "GGML_ASSERT(%s) failed", "backend");
  return backend->iface.graph_compute(backend, cgraph);
453}

455bool ggml_backend_supports_op(ggml_backend_t backend, const struct ggml_tensor * op) {
  GGML_ASSERT(backend)if (!(backend)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 456, "GGML_ASSERT(%s) failed", "backend");
  return ggml_backend_dev_supports_op(backend->device, op);
458}

460bool ggml_backend_supports_buft(ggml_backend_t backend, ggml_backend_buffer_type_t buft) {
  GGML_ASSERT(backend)if (!(backend)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 461, "GGML_ASSERT(%s) failed", "backend");
  return ggml_backend_dev_supports_buft(backend->device, buft);
463}

465bool ggml_backend_offload_op(ggml_backend_t backend, const struct ggml_tensor * op) {
  GGML_ASSERT(backend)if (!(backend)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 466, "GGML_ASSERT(%s) failed", "backend");
  return ggml_backend_dev_offload_op(backend->device, op);
468}

470ggml_backend_dev_t ggml_backend_get_device(ggml_backend_t backend) {
  GGML_ASSERT(backend)if (!(backend)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 471, "GGML_ASSERT(%s) failed", "backend");
  return backend->device;
473}

475// backend copy

477void ggml_backend_tensor_copy(const struct ggml_tensor * src, struct ggml_tensor * dst) {
  GGML_ASSERT(ggml_are_same_layout(src, dst) && "cannot copy tensors with different layouts")if (!(ggml_are_same_layout(src, dst) && "cannot copy tensors with different layouts"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 478, "GGML_ASSERT(%s) failed", "ggml_are_same_layout(src, dst) && \"cannot copy tensors with different layouts\""
);

  if (src == dst) {
      return;
  }

  if (ggml_backend_buffer_is_host(src->buffer)) {
      ggml_backend_tensor_set(dst, src->data, 0, ggml_nbytes(src));
  } else if (ggml_backend_buffer_is_host(dst->buffer)) {
      ggml_backend_tensor_get(src, dst->data, 0, ggml_nbytes(src));
  } else if (!ggml_backend_buffer_copy_tensor(src, dst)) {
489#ifndef NDEBUG
      GGML_LOG_DEBUG("%s: warning: slow copy from %s to %s\n", __func__, ggml_backend_buffer_name(src->buffer), ggml_backend_buffer_name(dst->buffer))ggml_log_internal(GGML_LOG_LEVEL_DEBUG, "%s: warning: slow copy from %s to %s\n"
, __func__, ggml_backend_buffer_name(src->buffer), ggml_backend_buffer_name
(dst->buffer));
491#endif // NDEBUG
      size_t nbytes = ggml_nbytes(src);
      void * data = malloc(nbytes);
      ggml_backend_tensor_get(src, data, 0, nbytes);
      ggml_backend_tensor_set(dst, data, 0, nbytes);
      free(data);
  }
498}

500void ggml_backend_tensor_copy_async(ggml_backend_t backend_src, ggml_backend_t backend_dst, const struct ggml_tensor * src, struct ggml_tensor * dst) {
  GGML_ASSERT(ggml_are_same_layout(src, dst) && "cannot copy tensors with different layouts")if (!(ggml_are_same_layout(src, dst) && "cannot copy tensors with different layouts"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 501, "GGML_ASSERT(%s) failed", "ggml_are_same_layout(src, dst) && \"cannot copy tensors with different layouts\""
);

  if (src == dst) {
      return;
  }

  GGML_ASSERT(backend_dst)if (!(backend_dst)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 507, "GGML_ASSERT(%s) failed", "backend_dst");
  if (backend_dst->iface.cpy_tensor_async != NULL__null) {
      if (backend_dst->iface.cpy_tensor_async(backend_src, backend_dst, src, dst)) {
          return;
      }
  }

  // an async copy would normally happen after all the queued operations on both backends are completed
  // to simulate the same behavior, we need to synchronize both backends first, and do a blocking copy
  ggml_backend_synchronize(backend_src);
  ggml_backend_synchronize(backend_dst);
  ggml_backend_tensor_copy(src, dst);
519}

521// events

523ggml_backend_event_t ggml_backend_event_new(ggml_backend_dev_t device) {
  // null device is allowed for the transition period to the device interface
  if (device == NULL__null || device->iface.event_new == NULL__null) {
      return NULL__null;
  }
  return device->iface.event_new(device);
529}

531void ggml_backend_event_free(ggml_backend_event_t event) {
  if (event == NULL__null) {
      return;
  }
  event->device->iface.event_free(event->device, event);
536}

538void ggml_backend_event_record(ggml_backend_event_t event, ggml_backend_t backend) {
  GGML_ASSERT(backend)if (!(backend)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 539, "GGML_ASSERT(%s) failed", "backend");
  GGML_ASSERT(backend->iface.event_record != NULL)if (!(backend->iface.event_record != __null)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 540, "GGML_ASSERT(%s) failed", "backend->iface.event_record != NULL"
);

  backend->iface.event_record(backend, event);
543}

545void ggml_backend_event_synchronize(ggml_backend_event_t event) {
  GGML_ASSERT(event)if (!(event)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 546, "GGML_ASSERT(%s) failed", "event");
  GGML_ASSERT(event->device->iface.event_synchronize)if (!(event->device->iface.event_synchronize)) ggml_abort
("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 547, "GGML_ASSERT(%s) failed", "event->device->iface.event_synchronize"
);

  event->device->iface.event_synchronize(event->device, event);
550}

552void ggml_backend_event_wait(ggml_backend_t backend, ggml_backend_event_t event) {
  GGML_ASSERT(backend)if (!(backend)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 553, "GGML_ASSERT(%s) failed", "backend");
  GGML_ASSERT(backend->iface.event_wait != NULL)if (!(backend->iface.event_wait != __null)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 554, "GGML_ASSERT(%s) failed", "backend->iface.event_wait != NULL"
);

  backend->iface.event_wait(backend, event);
557}

559static void ggml_backend_graph_optimize(ggml_backend_t backend, struct ggml_cgraph * cgraph) {
  GGML_ASSERT(backend)if (!(backend)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 560, "GGML_ASSERT(%s) failed", "backend");
  if (backend->iface.graph_optimize != NULL__null) {
      backend->iface.graph_optimize(backend, cgraph);
  }
564}

566// Backend device

568const char * ggml_backend_dev_name(ggml_backend_dev_t device) {
  GGML_ASSERT(device)if (!(device)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 569, "GGML_ASSERT(%s) failed", "device");
  return device->iface.get_name(device);
571}

573const char * ggml_backend_dev_description(ggml_backend_dev_t device) {
  GGML_ASSERT(device)if (!(device)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 574, "GGML_ASSERT(%s) failed", "device");
  return device->iface.get_description(device);
576}

578void ggml_backend_dev_memory(ggml_backend_dev_t device, size_t * free, size_t * total) {
  GGML_ASSERT(device)if (!(device)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 579, "GGML_ASSERT(%s) failed", "device");
  device->iface.get_memory(device, free, total);
581}

583enum ggml_backend_dev_type ggml_backend_dev_type(ggml_backend_dev_t device) {
  GGML_ASSERT(device)if (!(device)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 584, "GGML_ASSERT(%s) failed", "device");
  return device->iface.get_type(device);
586}

588void ggml_backend_dev_get_props(ggml_backend_dev_t device, struct ggml_backend_dev_props * props) {
  GGML_ASSERT(device)if (!(device)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 589, "GGML_ASSERT(%s) failed", "device");
  memset(props, 0, sizeof(*props));
  device->iface.get_props(device, props);
592}

594ggml_backend_reg_t ggml_backend_dev_backend_reg(ggml_backend_dev_t device) {
  GGML_ASSERT(device)if (!(device)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 595, "GGML_ASSERT(%s) failed", "device");
  return device->reg;
597}

599ggml_backend_t ggml_backend_dev_init(ggml_backend_dev_t device, const char * params) {
  GGML_ASSERT(device)if (!(device)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 600, "GGML_ASSERT(%s) failed", "device");
  return device->iface.init_backend(device, params);
602}

604ggml_backend_buffer_type_t ggml_backend_dev_buffer_type(ggml_backend_dev_t device) {
  GGML_ASSERT(device)if (!(device)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 605, "GGML_ASSERT(%s) failed", "device");
  return device->iface.get_buffer_type(device);
607}

609ggml_backend_buffer_type_t ggml_backend_dev_host_buffer_type(ggml_backend_dev_t device) {
  GGML_ASSERT(device)if (!(device)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 610, "GGML_ASSERT(%s) failed", "device");
  if (device->iface.get_host_buffer_type == NULL__null) {
      return NULL__null;
  }

  return device->iface.get_host_buffer_type(device);
616}

618ggml_backend_buffer_t ggml_backend_dev_buffer_from_host_ptr(ggml_backend_dev_t device, void * ptr, size_t size, size_t max_tensor_size) {
  GGML_ASSERT(device)if (!(device)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 619, "GGML_ASSERT(%s) failed", "device");
  return device->iface.buffer_from_host_ptr(device, ptr, size, max_tensor_size);
621}

623bool ggml_backend_dev_supports_op(ggml_backend_dev_t device, const struct ggml_tensor * op) {
  GGML_ASSERT(device)if (!(device)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 624, "GGML_ASSERT(%s) failed", "device");
  return device->iface.supports_op(device, op);
626}

628bool ggml_backend_dev_supports_buft(ggml_backend_dev_t device, ggml_backend_buffer_type_t buft) {
  GGML_ASSERT(device)if (!(device)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 629, "GGML_ASSERT(%s) failed", "device");
  return device->iface.supports_buft(device, buft);
631}

633bool ggml_backend_dev_offload_op(ggml_backend_dev_t device, const struct ggml_tensor * op) {
  GGML_ASSERT(device)if (!(device)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 634, "GGML_ASSERT(%s) failed", "device");
  if (device->iface.offload_op != NULL__null) {
      return device->iface.offload_op(device, op);
  }

  return false;
640}

642// Backend (reg)

644const char * ggml_backend_reg_name(ggml_backend_reg_t reg) {
  GGML_ASSERT(reg)if (!(reg)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 645, "GGML_ASSERT(%s) failed", "reg");
  return reg->iface.get_name(reg);
647}

649size_t ggml_backend_reg_dev_count(ggml_backend_reg_t reg) {
  GGML_ASSERT(reg)if (!(reg)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 650, "GGML_ASSERT(%s) failed", "reg");
  return reg->iface.get_device_count(reg);
652}

654ggml_backend_dev_t ggml_backend_reg_dev_get(ggml_backend_reg_t reg, size_t index) {
  GGML_ASSERT(reg)if (!(reg)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 655, "GGML_ASSERT(%s) failed", "reg");
  return reg->iface.get_device(reg, index);
657}

659void * ggml_backend_reg_get_proc_address(ggml_backend_reg_t reg, const char * name) {
  GGML_ASSERT(reg)if (!(reg)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 660, "GGML_ASSERT(%s) failed", "reg");
  if (!reg->iface.get_proc_address) {
      return NULL__null;
  }
  return reg->iface.get_proc_address(reg, name);
665}

667// multi-buffer buffer

669struct ggml_backend_multi_buffer_context {
  ggml_backend_buffer_t * buffers;
  size_t n_buffers;
672};

674static void ggml_backend_multi_buffer_free_buffer(ggml_backend_buffer_t buffer) {
  GGML_ASSERT(buffer)if (!(buffer)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 675, "GGML_ASSERT(%s) failed", "buffer");
  ggml_backend_multi_buffer_context * ctx = (ggml_backend_multi_buffer_context *) buffer->context;
  for (size_t i = 0; i < ctx->n_buffers; i++) {
      ggml_backend_buffer_free(ctx->buffers[i]);
  }

  free(ctx->buffers);
  free(ctx);
683}

685static void ggml_backend_multi_buffer_clear(ggml_backend_buffer_t buffer, uint8_t value) {
  GGML_ASSERT(buffer)if (!(buffer)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 686, "GGML_ASSERT(%s) failed", "buffer");
  ggml_backend_multi_buffer_context * ctx = (ggml_backend_multi_buffer_context *) buffer->context;
  for (size_t i = 0; i < ctx->n_buffers; i++) {
      ggml_backend_buffer_clear(ctx->buffers[i], value);
  }
691}

693static const struct ggml_backend_buffer_i ggml_backend_multi_buffer_i = {
  /* .free_buffer     = */ ggml_backend_multi_buffer_free_buffer,
  /* .get_base        = */ NULL__null,
  /* .init_tensor     = */ NULL__null,
  /* .memset_tensor   = */ NULL__null,
  /* .set_tensor      = */ NULL__null,
  /* .get_tensor      = */ NULL__null,
  /* .set_tensor_2d   = */ NULL__null,
  /* .get_tensor_2d   = */ NULL__null,
  /* .cpy_tensor      = */ NULL__null,
  /* .clear           = */ ggml_backend_multi_buffer_clear,
  /* .reset           = */ NULL__null,
705};

707ggml_backend_buffer_t ggml_backend_multi_buffer_alloc_buffer(ggml_backend_buffer_t * buffers, size_t n_buffers) {
  ggml_backend_multi_buffer_context * ctx = (ggml_backend_multi_buffer_context *) malloc(sizeof(struct ggml_backend_multi_buffer_context));
  ctx->n_buffers = n_buffers;
  ctx->buffers = (ggml_backend_buffer_t *) malloc(n_buffers * sizeof(ggml_backend_buffer_t));

  GGML_ASSERT(ctx->buffers != NULL)if (!(ctx->buffers != __null)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 712, "GGML_ASSERT(%s) failed", "ctx->buffers != NULL");

  size_t total_size = 0;
  for (size_t i = 0; i < n_buffers; i++) {
      ctx->buffers[i] = buffers[i];
      total_size += ggml_backend_buffer_get_size(buffers[i]);
  }

  return ggml_backend_buffer_init(buffers[0]->buft, ggml_backend_multi_buffer_i, ctx, total_size);
721}

723bool ggml_backend_buffer_is_multi_buffer(ggml_backend_buffer_t buffer) {
  GGML_ASSERT(buffer)if (!(buffer)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 724, "GGML_ASSERT(%s) failed", "buffer");
  return buffer->iface.free_buffer == ggml_backend_multi_buffer_free_buffer;
726}

728void ggml_backend_multi_buffer_set_usage(ggml_backend_buffer_t buffer, enum ggml_backend_buffer_usage usage) {
  GGML_ASSERT(buffer)if (!(buffer)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 729, "GGML_ASSERT(%s) failed", "buffer");
  GGML_ASSERT(ggml_backend_buffer_is_multi_buffer(buffer))if (!(ggml_backend_buffer_is_multi_buffer(buffer))) ggml_abort
("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 730, "GGML_ASSERT(%s) failed", "ggml_backend_buffer_is_multi_buffer(buffer)"
);
  ggml_backend_multi_buffer_context * ctx = (ggml_backend_multi_buffer_context *) buffer->context;
  for (size_t i = 0; i < ctx->n_buffers; i++) {
      ggml_backend_buffer_set_usage(ctx->buffers[i], usage);
  }
735}

737// creates a copy of the tensor with the same memory layout
738static struct ggml_tensor * ggml_dup_tensor_layout(struct ggml_context * ctx, const struct ggml_tensor * tensor) {
  struct ggml_tensor * dup = ggml_dup_tensor(ctx, tensor);
  for (int i = 0; i < GGML_MAX_DIMS4; i++) {
      dup->nb[i] = tensor->nb[i];
  }
  return dup;
744}

746static bool ggml_is_view_op(enum ggml_op op) {
  return op == GGML_OP_VIEW || op == GGML_OP_RESHAPE || op == GGML_OP_PERMUTE || op == GGML_OP_TRANSPOSE;
748}

750// scheduler

752#ifndef GGML_SCHED_MAX_BACKENDS16
753#define GGML_SCHED_MAX_BACKENDS16 16
754#endif

756#ifndef GGML_SCHED_MAX_SPLIT_INPUTS30
757#define GGML_SCHED_MAX_SPLIT_INPUTS30 30
758#endif

760#ifndef GGML_SCHED_MAX_COPIES4
761#define GGML_SCHED_MAX_COPIES4 4
762#endif

764struct ggml_backend_sched_split {
  int backend_id;
  int i_start;
  int i_end;
  struct ggml_tensor * inputs[GGML_SCHED_MAX_SPLIT_INPUTS30];
  int n_inputs;
  // graph view of this split
  struct ggml_cgraph graph;
772};

774struct ggml_backend_sched {
  bool is_reset; // true if the scheduler has been reset since the last graph split
  bool is_alloc;

  int n_backends;

  ggml_backend_t backends[GGML_SCHED_MAX_BACKENDS16];
  ggml_backend_buffer_type_t bufts[GGML_SCHED_MAX_BACKENDS16];
  ggml_gallocr_t galloc;

  // hash map of the nodes in the graph
  struct ggml_hash_set  hash_set;
  int                 * hv_tensor_backend_ids; // [hash_set.size]
  struct ggml_tensor ** hv_tensor_copies;      // [hash_set.size][n_backends][n_copies]

  int * node_backend_ids; // [graph_size]
  int * leaf_backend_ids; // [graph_size]

  int * prev_node_backend_ids; // [graph_size]
  int * prev_leaf_backend_ids; // [graph_size]

  // copy of the graph with modified inputs
  struct ggml_cgraph graph;

  // graph splits
  struct ggml_backend_sched_split * splits;
  int n_splits;
  int splits_capacity;

  // pipeline parallelism support
  int n_copies;
  int cur_copy;
  int next_copy;
  ggml_backend_event_t events[GGML_SCHED_MAX_BACKENDS16][GGML_SCHED_MAX_COPIES4];
  struct ggml_tensor * graph_inputs[GGML_SCHED_MAX_SPLIT_INPUTS30];
  int n_graph_inputs;

  struct ggml_context * ctx;

  ggml_backend_sched_eval_callback callback_eval;
  void * callback_eval_user_data;

  char * context_buffer;
  size_t context_buffer_size;

  bool op_offload;

  int debug;

  // used for debugging graph reallocations [GGML_SCHED_DEBUG_REALLOC]
  // ref: https://github.com/ggml-org/llama.cpp/pull/17617
  int debug_realloc;
  int debug_graph_size;
  int debug_prev_graph_size;
828};

830#define hash_id(tensor)ggml_hash_find_or_insert(&sched->hash_set, tensor) ggml_hash_find_or_insert(&sched->hash_set, tensor)
831#define tensor_backend_id(tensor)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, tensor)] sched->hv_tensor_backend_ids[hash_id(tensor)ggml_hash_find_or_insert(&sched->hash_set, tensor)]
832#define tensor_id_copy(id, backend_id, copy_id)sched->hv_tensor_copies[(id) * sched->n_backends * sched
->n_copies + (backend_id) * sched->n_copies + (copy_id)
] sched->hv_tensor_copies[(id) * sched->n_backends * sched->n_copies + (backend_id) * sched->n_copies + (copy_id)]
833#define tensor_copy(tensor, backend_id, copy_id)sched->hv_tensor_copies[(ggml_hash_find_or_insert(&sched
->hash_set, tensor)) * sched->n_backends * sched->n_copies
 + (backend_id) * sched->n_copies + (copy_id)] tensor_id_copy(hash_id(tensor), backend_id, copy_id)sched->hv_tensor_copies[(ggml_hash_find_or_insert(&sched
->hash_set, tensor)) * sched->n_backends * sched->n_copies
 + (backend_id) * sched->n_copies + (copy_id)]

835// returns the priority of the backend, lower id is higher priority
836static int ggml_backend_sched_backend_id(ggml_backend_sched_t sched, ggml_backend_t backend) {
  for (int i = 0; i < sched->n_backends; i++) {
      if (sched->backends[i] == backend) {
          return i;
      }
  }
  return -1;
843}

845static int ggml_backend_sched_backend_from_buffer(ggml_backend_sched_t sched, const struct ggml_tensor * tensor, const struct ggml_tensor * op) {
  ggml_backend_buffer_t buffer = tensor->view_src ? tensor->view_src->buffer : tensor->buffer;
  if (buffer == NULL__null) {
      return -1;
  }

  // find highest prio backend that supports the buffer type and the op
  for (int i = 0; i < sched->n_backends; i++) {
      if (ggml_backend_supports_buft(sched->backends[i], buffer->buft) &&
          ggml_backend_supports_op(sched->backends[i], op)) {
          return i;
      }
  }

859#ifndef NDEBUG
  GGML_LOG_DEBUG("%s: warning: no backend supports op %s with a weight with buffer type %s used in tensor %s, the weight will need to be copied\n",ggml_log_internal(GGML_LOG_LEVEL_DEBUG, "%s: warning: no backend supports op %s with a weight with buffer type %s used in tensor %s, the weight will need to be copied\n"
, __func__, ggml_op_desc(tensor), ggml_backend_buffer_name(buffer
), tensor->name)
      __func__, ggml_op_desc(tensor), ggml_backend_buffer_name(buffer), tensor->name)ggml_log_internal(GGML_LOG_LEVEL_DEBUG, "%s: warning: no backend supports op %s with a weight with buffer type %s used in tensor %s, the weight will need to be copied\n"
, __func__, ggml_op_desc(tensor), ggml_backend_buffer_name(buffer
), tensor->name);
862#endif

  return -1;
865}

867#if 0
868#define GGML_SCHED_MAX_SPLITS_DEBUG 4096
869static char causes[GGML_DEFAULT_GRAPH_SIZE2048*16 + GGML_SCHED_MAX_SPLITS_DEBUG*GGML_SCHED_MAX_SPLIT_INPUTS30][128]; // debug only
870#define SET_CAUSE(node, ...) sprintf(causes[hash_id(node)ggml_hash_find_or_insert(&sched->hash_set, node)], __VA_ARGS__)
871#define GET_CAUSE(node)"" causes[hash_id(node)ggml_hash_find_or_insert(&sched->hash_set, node)]
872#else
873#define SET_CAUSE(node, ...)
874#define GET_CAUSE(node)"" ""
875#endif

877// returns the backend that should be used for the node based on the current locations
878static int ggml_backend_sched_backend_id_from_cur(ggml_backend_sched_t sched, struct ggml_tensor * tensor) {
  // assign pre-allocated nodes to their backend
  int cur_backend_id = ggml_backend_sched_backend_from_buffer(sched, tensor, tensor);
  if (cur_backend_id != -1) {
      SET_CAUSE(tensor, "1.dst");
      return cur_backend_id;
  }

  // view_src
  if (tensor->view_src != NULL__null) {
      cur_backend_id = ggml_backend_sched_backend_from_buffer(sched, tensor->view_src, tensor);
      if (cur_backend_id != -1) {
          SET_CAUSE(tensor, "1.vsrc");
          return cur_backend_id;
      }
  }

  if (tensor->buffer || (tensor->view_src && tensor->view_src->buffer)) {
      // since the tensor is pre-allocated, it cannot be moved to another backend
      ggml_backend_buffer_t buffer = tensor->view_src ? tensor->view_src->buffer : tensor->buffer;
      GGML_ABORT("pre-allocated tensor (%s) in a buffer (%s) that cannot run the operation (%s)", tensor->name, ggml_backend_buffer_name(buffer), ggml_op_name(tensor->op))ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 898, "pre-allocated tensor (%s) in a buffer (%s) that cannot run the operation (%s)"
, tensor->name, ggml_backend_buffer_name(buffer), ggml_op_name
(tensor->op));
  }

  // graph input
  if (tensor->flags & GGML_TENSOR_FLAG_INPUT) {
      cur_backend_id = sched->n_backends - 1; // last backend (assumed CPU)
      SET_CAUSE(tensor, "1.inp");
      return cur_backend_id;
  }

  // operations with weights are preferably run on the same backend as the weights
  for (int i = 0; i < GGML_MAX_SRC10; i++) {
      const struct ggml_tensor * src = tensor->src[i];
      if (src == NULL__null) {
          continue;
      }
      // skip ROPE since the rope freqs tensor is too small to choose a backend based on it
      // not an ideal solution
      if (tensor->op != GGML_OP_ROPE && src->buffer != NULL__null && src->buffer->usage == GGML_BACKEND_BUFFER_USAGE_WEIGHTS) {
          int src_backend_id = ggml_backend_sched_backend_from_buffer(sched, src, tensor);
          // check if a backend with higher prio wants to offload the op
          if (sched->op_offload && src_backend_id == sched->n_backends - 1 && ggml_backend_buffer_is_host(src->buffer)) {
              for (int b = 0; b < src_backend_id; b++) {
                  if (ggml_backend_supports_op(sched->backends[b], tensor) && ggml_backend_offload_op(sched->backends[b], tensor)) {
                      SET_CAUSE(tensor, "1.off");
                      return b;
                  }
              }
          }
          SET_CAUSE(tensor, "1.wgt%d", i);
          return src_backend_id;
      }
  }

  return -1;
933}

935static char * fmt_size(size_t size) {
  static char buffer[128];
  if (size >= 1024*1024) {
      snprintf(buffer, sizeof(buffer), "%zuM", size/1024/1024);
  } else {
      snprintf(buffer, sizeof(buffer), "%zuK", size/1024);
  }
  return buffer;
943}

945static void ggml_backend_sched_print_assignments(ggml_backend_sched_t sched, struct ggml_cgraph * graph) {
  int cur_split = 0;
  for (int i = 0; i < graph->n_nodes; i++) {
      if (cur_split < sched->n_splits && i == sched->splits[cur_split].i_start) {
          ggml_backend_t split_backend = sched->backends[sched->splits[cur_split].backend_id];
          GGML_LOG_DEBUG("\n## SPLIT #%d: %s # %d inputs", cur_split, ggml_backend_name(split_backend),ggml_log_internal(GGML_LOG_LEVEL_DEBUG, "\n## SPLIT #%d: %s # %d inputs"
, cur_split, ggml_backend_name(split_backend), sched->splits
[cur_split].n_inputs)
              sched->splits[cur_split].n_inputs)ggml_log_internal(GGML_LOG_LEVEL_DEBUG, "\n## SPLIT #%d: %s # %d inputs"
, cur_split, ggml_backend_name(split_backend), sched->splits
[cur_split].n_inputs);
          for (int j = 0; j < sched->splits[cur_split].n_inputs; j++) {
              if (j == 0) {
                  GGML_LOG_DEBUG(": ")ggml_log_internal(GGML_LOG_LEVEL_DEBUG, ": ");
              }
              GGML_LOG_DEBUG("[%s (%5.5s)] ", sched->splits[cur_split].inputs[j]->name,ggml_log_internal(GGML_LOG_LEVEL_DEBUG, "[%s (%5.5s)] ", sched
->splits[cur_split].inputs[j]->name, fmt_size(ggml_nbytes
(sched->splits[cur_split].inputs[j])))
                  fmt_size(ggml_nbytes(sched->splits[cur_split].inputs[j])))ggml_log_internal(GGML_LOG_LEVEL_DEBUG, "[%s (%5.5s)] ", sched
->splits[cur_split].inputs[j]->name, fmt_size(ggml_nbytes
(sched->splits[cur_split].inputs[j])));
          }
          GGML_LOG_DEBUG("\n")ggml_log_internal(GGML_LOG_LEVEL_DEBUG, "\n");
          cur_split++;
      }
      struct ggml_tensor * node = graph->nodes[i];
      if (ggml_is_view_op(node->op)) {
          continue;
      }
      if (sched->debug > 1) {
          ggml_backend_t tensor_backend = ggml_backend_sched_get_tensor_backend(sched, node);
          GGML_LOG_DEBUG("node #%3d (%10.10s): %20.20s (%5.5s) [%5.5s %8.8s] use=%d,c=%d:", i, ggml_op_desc(node), node->name,ggml_log_internal(GGML_LOG_LEVEL_DEBUG, "node #%3d (%10.10s): %20.20s (%5.5s) [%5.5s %8.8s] use=%d,c=%d:"
, i, ggml_op_desc(node), node->name, fmt_size(ggml_nbytes(
node)), tensor_backend ? ggml_backend_name(tensor_backend) : "NULL"
, "", graph->use_counts[ggml_hash_find(&graph->visited_hash_set
, node)], node->flags & GGML_TENSOR_FLAG_COMPUTE ? 1 :
 0)
              fmt_size(ggml_nbytes(node)), tensor_backend ? ggml_backend_name(tensor_backend) : "NULL", GET_CAUSE(node),ggml_log_internal(GGML_LOG_LEVEL_DEBUG, "node #%3d (%10.10s): %20.20s (%5.5s) [%5.5s %8.8s] use=%d,c=%d:"
, i, ggml_op_desc(node), node->name, fmt_size(ggml_nbytes(
node)), tensor_backend ? ggml_backend_name(tensor_backend) : "NULL"
, "", graph->use_counts[ggml_hash_find(&graph->visited_hash_set
, node)], node->flags & GGML_TENSOR_FLAG_COMPUTE ? 1 :
 0)
              graph->use_counts[ggml_hash_find(&graph->visited_hash_set, node)], node->flags & GGML_TENSOR_FLAG_COMPUTE ? 1 : 0)ggml_log_internal(GGML_LOG_LEVEL_DEBUG, "node #%3d (%10.10s): %20.20s (%5.5s) [%5.5s %8.8s] use=%d,c=%d:"
, i, ggml_op_desc(node), node->name, fmt_size(ggml_nbytes(
node)), tensor_backend ? ggml_backend_name(tensor_backend) : "NULL"
, "", graph->use_counts[ggml_hash_find(&graph->visited_hash_set
, node)], node->flags & GGML_TENSOR_FLAG_COMPUTE ? 1 :
 0);
          for (int j = 0; j < GGML_MAX_SRC10; j++) {
              struct ggml_tensor * src = node->src[j];
              if (src == NULL__null) {
                  continue;
              }
              ggml_backend_t src_backend = ggml_backend_sched_get_tensor_backend(sched, src);
              GGML_LOG_DEBUG(" %20.20s (%5.5s) [%5.5s %8.8s]", src->name,ggml_log_internal(GGML_LOG_LEVEL_DEBUG, " %20.20s (%5.5s) [%5.5s %8.8s]"
, src->name, fmt_size(ggml_nbytes(src)), src_backend ? ggml_backend_name
(src_backend) : "NULL", "")
                  fmt_size(ggml_nbytes(src)), src_backend ? ggml_backend_name(src_backend) : "NULL", GET_CAUSE(src))ggml_log_internal(GGML_LOG_LEVEL_DEBUG, " %20.20s (%5.5s) [%5.5s %8.8s]"
, src->name, fmt_size(ggml_nbytes(src)), src_backend ? ggml_backend_name
(src_backend) : "NULL", "");
          }
          GGML_LOG_DEBUG("\n")ggml_log_internal(GGML_LOG_LEVEL_DEBUG, "\n");
      }
  }
983}

985static bool ggml_backend_sched_buffer_supported(ggml_backend_sched_t sched, struct ggml_tensor * t, int backend_id) {
  ggml_backend_buffer_t buf = t->view_src ? t->view_src->buffer : t->buffer;
  ggml_backend_buffer_type_t buft = NULL__null;

  if (buf) {
      // the tensor is already allocated
      buft = buf->buft;
  } else {
      // see if the tensor already has a backend assigned, and use the buffer type of that backend
      int tensor_backend_id = tensor_backend_id(t)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, t)];
      if (tensor_backend_id == -1 && t->view_src) {
          tensor_backend_id = tensor_backend_id(t->view_src)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, t->view_src)];
      }
      if (tensor_backend_id != -1) {
          buft = sched->bufts[tensor_backend_id];
      }
  }

  return buft != NULL__null && ggml_backend_supports_buft(sched->backends[backend_id], buft);
1004}

1006static void ggml_backend_sched_set_if_supported(ggml_backend_sched_t sched, struct ggml_tensor * node, int cur_backend_id, int * node_backend_id) {
  if (ggml_backend_supports_op(sched->backends[cur_backend_id], node)) {
      *node_backend_id = cur_backend_id;
      SET_CAUSE(node, "2.sup");
  }
1011}

1013// assigns backends to ops and splits the graph into subgraphs that can be computed on the same backend
1014void ggml_backend_sched_split_graph(ggml_backend_sched_t sched, struct ggml_cgraph * graph) {
  // reset splits
  sched->n_splits = 0;
  sched->n_graph_inputs = 0;
  sched->is_reset = false;

  struct ggml_init_params params = {
      /* .mem_size =   */ sched->context_buffer_size,
      /* .mem_buffer = */ sched->context_buffer,
      /* .no_alloc =   */ true
  };

  ggml_free(sched->ctx);

  sched->ctx = ggml_init(params);
  if (sched->ctx == NULL__null) {
      GGML_ABORT("%s: failed to initialize context\n", __func__)ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1030, "%s: failed to initialize context\n", __func__);
  }

  graph->uid = ggml_graph_next_uid();

  // pass 1: assign backends to ops with pre-allocated inputs
  for (int i = 0; i < graph->n_leafs; i++) {
      struct ggml_tensor * leaf = graph->leafs[i];
      int * leaf_backend_id = &tensor_backend_id(leaf)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, leaf)];
      // do not overwrite user assignments
      if (*leaf_backend_id == -1) {
          *leaf_backend_id = ggml_backend_sched_backend_id_from_cur(sched, leaf);
      }
  }

  for (int i = 0; i < graph->n_nodes; i++) {
      struct ggml_tensor * node = graph->nodes[i];
      int * node_backend_id = &tensor_backend_id(node)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, node)];
      // do not overwrite user assignments
      if (*node_backend_id == -1) {
          *node_backend_id = ggml_backend_sched_backend_id_from_cur(sched, node);

1052#if 0
          // src
          if (node->op == GGML_OP_NONE) {
              continue;
          }

          for (int j = 0; j < GGML_MAX_SRC10; j++) {
              struct ggml_tensor * src = node->src[j];
              if (src == NULL__null) {
                  continue;
              }
              int * src_backend_id = &tensor_backend_id(src)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, src)];
              if (*src_backend_id == -1) {
                  *src_backend_id = ggml_backend_sched_backend_id_from_cur(sched, src);
              }
          }
1068#endif
      }
  }

  // pass 2: expand current backend assignments
  // assign the same backend to adjacent nodes
  // expand gpu backends (i.e. non last prio) up and down, ignoring cpu (the lowest priority backend)
  // thus, cpu will never be used unless weights are on cpu, or there are no gpu ops between cpu ops
  // ops unsupported by the backend being expanded will be left unassigned so that they can be assigned later when the locations of its inputs are known
  // expand gpu down
  {
      int cur_backend_id = -1;
      for (int i = 0; i < graph->n_nodes; i++) {
          struct ggml_tensor * node = graph->nodes[i];
          if (ggml_is_view_op(node->op)) {
              continue;
          }
          int * node_backend_id = &tensor_backend_id(node)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, node)];
          if (*node_backend_id != -1) {
              if (*node_backend_id == sched->n_backends - 1) {
                  // skip cpu (lowest prio backend)
                  cur_backend_id = -1;
              } else {
                  cur_backend_id = *node_backend_id;
              }
          } else if (cur_backend_id != -1) {
              ggml_backend_sched_set_if_supported(sched, node, cur_backend_id, node_backend_id);
          }
      }
  }
  // expand gpu up
  {
      int cur_backend_id = -1;
      for (int i = graph->n_nodes - 1; i >= 0; i--) {
          struct ggml_tensor * node = graph->nodes[i];
          if (ggml_is_view_op(node->op)) {
              continue;
          }
          int * node_backend_id = &tensor_backend_id(node)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, node)];
          if (*node_backend_id != -1) {
              if (*node_backend_id == sched->n_backends - 1) {
                  // skip cpu (lowest prio backend)
                  cur_backend_id = -1;
              } else {
                  cur_backend_id = *node_backend_id;
              }
          } else if (cur_backend_id != -1) {
              ggml_backend_sched_set_if_supported(sched, node, cur_backend_id, node_backend_id);
          }
      }
  }
  // expand rest down
  {
      int cur_backend_id = -1;
      for (int i = 0; i < graph->n_nodes; i++) {
          struct ggml_tensor * node = graph->nodes[i];
          if (ggml_is_view_op(node->op)) {
              continue;
          }
          int * node_backend_id = &tensor_backend_id(node)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, node)];
          if (*node_backend_id != -1) {
              cur_backend_id = *node_backend_id;
          } else if (cur_backend_id != -1) {
              ggml_backend_sched_set_if_supported(sched, node, cur_backend_id, node_backend_id);
          }
      }
  }
  // expand rest up
  {
      int cur_backend_id = -1;
      for (int i = graph->n_nodes - 1; i >= 0; i--) {
          struct ggml_tensor * node = graph->nodes[i];
          if (ggml_is_view_op(node->op)) {
              continue;
          }
          int * node_backend_id = &tensor_backend_id(node)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, node)];
          if (*node_backend_id != -1) {
              cur_backend_id = *node_backend_id;
          } else if (cur_backend_id != -1) {
              ggml_backend_sched_set_if_supported(sched, node, cur_backend_id, node_backend_id);
          }
      }
  }

  // pass 3: upgrade nodes to higher prio backends with compatible buffer types
  // if the tensor is already in the same buffer type (*) as another higher priority backend, we should move it there
  // however, we also need to verify that the sources are in compatible buffer types
  // (*) the actual requirement is more relaxed, the buffer type of the backend should be supported by all the users of this tensor further down the graph
  // however, this is slow to verify, so we have a more strict requirement that the buffer type is the same
  // this is not uncommon since multiple backends can use host memory, with the same buffer type (eg. BLAS and CPU)
  // additionally, set remaining unassigned nodes to the backend with the most supported inputs
  // only nodes that could not be assigned during expansion due to the backend not supporting the op should be unassigned at this point
  for (int i = 0; i < graph->n_nodes; i++) {
      struct ggml_tensor * node = graph->nodes[i];
      if (ggml_is_view_op(node->op)) {
          continue;
      }
      int * node_backend_id = &tensor_backend_id(node)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, node)];
      if (*node_backend_id == -1) {
          // unassigned node: find the backend with the most supported inputs
          int n_supported_best = -1;
          for (int b = 0; b < sched->n_backends; b++) {
              if (ggml_backend_supports_op(sched->backends[b], node)) {
                  int n_supported = 0;
                  for (int j = 0; j < GGML_MAX_SRC10; j++) {
                      struct ggml_tensor * src = node->src[j];
                      if (src == NULL__null) {
                          continue;
                      }
                      if ((tensor_backend_id(src)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, src)] != -1 || tensor_backend_id(src->view_src)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, src->view_src)] != -1) && ggml_backend_sched_buffer_supported(sched, src, b)) {
                          n_supported++;
                      }
                  }
                  if (n_supported > n_supported_best) {
                      n_supported_best = n_supported;
                      *node_backend_id = b;
                      SET_CAUSE(node, "3.best");
                  }
              }
          }
      } else {
          // assigned node: upgrade to higher prio backend if possible
          for (int b = 0; b < *node_backend_id; b++) {
              if (sched->bufts[b] == sched->bufts[*node_backend_id] && ggml_backend_supports_op(sched->backends[b], node)) {
                  bool supported = true;
                  for (int j = 0; j < GGML_MAX_SRC10; j++) {
                      struct ggml_tensor * src = node->src[j];
                      if (src == NULL__null) {
                          continue;
                      }
                      if (!ggml_backend_sched_buffer_supported(sched, src, b)) {
                          supported = false;
                          break;
                      }
                  }
                  if (supported) {
                      *node_backend_id = b;
                      SET_CAUSE(node, "3.upg");
                      break;
                  }
              }
          }
      }
  }

  // pass 4: assign backends to remaining src from dst and view_src
  for (int i = 0; i < graph->n_nodes; i++) {
      struct ggml_tensor * node = graph->nodes[i];
      int * cur_backend_id = &tensor_backend_id(node)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, node)];
      if (node->view_src != NULL__null && *cur_backend_id == -1) {
          *cur_backend_id = tensor_backend_id(node->view_src)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, node->view_src)];
          SET_CAUSE(node, "4.vsrc");
      }
      for (int j = 0; j < GGML_MAX_SRC10; j++) {
          struct ggml_tensor * src = node->src[j];
          if (src == NULL__null) {
              continue;
          }
          int * src_backend_id = &tensor_backend_id(src)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, src)];
          if (*src_backend_id == -1) {
              if (src->view_src != NULL__null) {
                  // views are always on the same backend as the source
                  *src_backend_id = tensor_backend_id(src->view_src)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, src->view_src)];
                  SET_CAUSE(src, "4.vsrc");
              } else {
                  *src_backend_id = *cur_backend_id;
                  SET_CAUSE(src, "4.cur");
              }
          }
      }
      // if the node is still unassigned, assign it to the first backend that supports it
      for (int b = 0; b < sched->n_backends && *cur_backend_id == -1; b++) {
          ggml_backend_sched_set_if_supported(sched, node, b, cur_backend_id);
      }
      GGML_ASSERT(*cur_backend_id != -1)if (!(*cur_backend_id != -1)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1242, "GGML_ASSERT(%s) failed", "*cur_backend_id != -1");
  }

  // pass 5: split graph, find tensors that need to be copied
  {
      int i_split = 0;
      struct ggml_backend_sched_split * split = &sched->splits[0];
      // find the backend of the first split, skipping view ops
      int i = 0;
      for (; i < graph->n_nodes; i++) {
          struct ggml_tensor * node = graph->nodes[i];
          if (!ggml_is_view_op(node->op)) {
              split->backend_id = tensor_backend_id(node)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, node)];
              break;
          }
      }
      split->i_start = 0;
      split->n_inputs = 0;
      int cur_backend_id = split->backend_id;
      for (; i < graph->n_nodes; i++) {
          struct ggml_tensor * node = graph->nodes[i];

          if (ggml_is_view_op(node->op)) {
              continue;
          }

          const int node_backend_id = tensor_backend_id(node)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, node)];

          GGML_ASSERT(node_backend_id != -1)if (!(node_backend_id != -1)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1270, "GGML_ASSERT(%s) failed", "node_backend_id != -1"); // all nodes should be assigned by now, this can happen if there is no CPU fallback

          // check if we should start a new split based on the sources of the current node
          bool need_new_split = false;
          if (node_backend_id == cur_backend_id && split->n_inputs > 0) {
              for (int j = 0; j < GGML_MAX_SRC10; j++) {
                  struct ggml_tensor * src = node->src[j];
                  if (src == NULL__null) {
                      continue;
                  }
                  // check if a weight is on a different and incompatible backend
                  // by starting a new split, the memory of the previously offloaded weights can be reused
                  if (src->buffer != NULL__null && src->buffer->usage == GGML_BACKEND_BUFFER_USAGE_WEIGHTS) {
                      int src_backend_id = tensor_backend_id(src)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, src)];
                      if (src_backend_id != cur_backend_id && !ggml_backend_sched_buffer_supported(sched, src, cur_backend_id)) {
                          need_new_split = true;
                          break;
                      }
                  }
                  // check if the split has too many inputs
                  // FIXME: count the number of inputs instead of only checking when full
                  if (split->n_inputs == GGML_SCHED_MAX_SPLIT_INPUTS30) {
                      const size_t id = hash_id(src)ggml_hash_find_or_insert(&sched->hash_set, src);
                      int src_backend_id = sched->hv_tensor_backend_ids[id];
                      bool supported = ggml_backend_sched_buffer_supported(sched, src, cur_backend_id);
                      if (src_backend_id != cur_backend_id && tensor_id_copy(id, cur_backend_id, 0)sched->hv_tensor_copies[(id) * sched->n_backends * sched
->n_copies + (cur_backend_id) * sched->n_copies + (0)] == NULL__null && !supported) {
                          need_new_split = true;
                          break;
                      }
                  }
              }
          }

          if (node_backend_id != cur_backend_id || need_new_split) {
              split->i_end = i;
              i_split++;
              if (i_split >= sched->splits_capacity) {
                  sched->splits_capacity *= 2;
                  sched->splits = (ggml_backend_sched_split *)
                      realloc(sched->splits, sched->splits_capacity * sizeof(struct ggml_backend_sched_split));
                  GGML_ASSERT(sched->splits != NULL)if (!(sched->splits != __null)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1310, "GGML_ASSERT(%s) failed", "sched->splits != NULL");
              }
              split = &sched->splits[i_split];
              split->backend_id = node_backend_id;
              split->i_start = i;
              split->n_inputs = 0;
              cur_backend_id = node_backend_id;
          }

          // find inputs that are not on the same backend
          for (int j = 0; j < GGML_MAX_SRC10; j++) {
              struct ggml_tensor * src = node->src[j];
              if (src == NULL__null) {
                  continue;
              }

              size_t src_id = hash_id(src)ggml_hash_find_or_insert(&sched->hash_set, src);
              const int src_backend_id = sched->hv_tensor_backend_ids[src_id];
              GGML_ASSERT(src_backend_id != -1)if (!(src_backend_id != -1)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1328, "GGML_ASSERT(%s) failed", "src_backend_id != -1"); // all inputs should be assigned by now

              if (src->flags & GGML_TENSOR_FLAG_INPUT && sched->n_copies > 1) {
                  if (tensor_id_copy(src_id, src_backend_id, 0)sched->hv_tensor_copies[(src_id) * sched->n_backends * sched
->n_copies + (src_backend_id) * sched->n_copies + (0)] == NULL__null) {
                      ggml_backend_t backend = sched->backends[src_backend_id];
                      for (int c = 0; c < sched->n_copies; c++) {
                          struct ggml_tensor * tensor_copy;
                          if (c == sched->cur_copy) {
                              tensor_copy = src; // use the original tensor as the current copy
                          } else {
                              tensor_copy = ggml_dup_tensor_layout(sched->ctx, src);
                              ggml_format_name(tensor_copy, "%s#%s#%d", ggml_backend_name(backend), src->name, c);
                          }
                          ggml_set_input(tensor_copy);
                          ggml_set_output(tensor_copy); // prevent ggml-alloc from overwriting the tensor
                          tensor_id_copy(src_id, src_backend_id, c)sched->hv_tensor_copies[(src_id) * sched->n_backends * sched
->n_copies + (src_backend_id) * sched->n_copies + (c)] = tensor_copy;
                          SET_CAUSE(tensor_copy, "4.cpy");
                      }
                      int n_graph_inputs = sched->n_graph_inputs++;
                      GGML_ASSERT(n_graph_inputs < GGML_SCHED_MAX_SPLIT_INPUTS)if (!(n_graph_inputs < 30)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1347, "GGML_ASSERT(%s) failed", "n_graph_inputs < GGML_SCHED_MAX_SPLIT_INPUTS"
);
                      sched->graph_inputs[n_graph_inputs] = src;
                  }
              }

              if (src_backend_id != cur_backend_id && !ggml_backend_sched_buffer_supported(sched, src, cur_backend_id)) {
                  // create a copy of the input in the split's backend
                  if (tensor_id_copy(src_id, cur_backend_id, 0)sched->hv_tensor_copies[(src_id) * sched->n_backends * sched
->n_copies + (cur_backend_id) * sched->n_copies + (0)] == NULL__null) {
                      ggml_backend_t backend = sched->backends[cur_backend_id];
                      for (int c = 0; c < sched->n_copies; c++) {
                          struct ggml_tensor * tensor_copy = ggml_dup_tensor_layout(sched->ctx, src);
                          ggml_format_name(tensor_copy, "%s#%s#%d", ggml_backend_name(backend), src->name, c);
                          if (sched->n_copies > 1) {
                              ggml_set_input(tensor_copy);
                              ggml_set_output(tensor_copy); // prevent ggml-alloc from overwriting the tensor
                          }
                          tensor_id_copy(src_id, cur_backend_id, c)sched->hv_tensor_copies[(src_id) * sched->n_backends * sched
->n_copies + (cur_backend_id) * sched->n_copies + (c)] = tensor_copy;
                          SET_CAUSE(tensor_copy, "4.cpy");
                      }
                      int n_inputs = split->n_inputs++;
                      GGML_ASSERT(n_inputs < GGML_SCHED_MAX_SPLIT_INPUTS)if (!(n_inputs < 30)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1367, "GGML_ASSERT(%s) failed", "n_inputs < GGML_SCHED_MAX_SPLIT_INPUTS"
);
                      split->inputs[n_inputs] = src;
                  }
                  node->src[j] = tensor_id_copy(src_id, cur_backend_id, sched->cur_copy)sched->hv_tensor_copies[(src_id) * sched->n_backends * sched
->n_copies + (cur_backend_id) * sched->n_copies + (sched
->cur_copy)];
              }
          }
      }
      split->i_end = graph->n_nodes;
      sched->n_splits = i_split + 1;
  }

  if (sched->debug) {
      ggml_backend_sched_print_assignments(sched, graph);
  }

  // swap node_backend_ids and leaf _backend_ids with prevs
  {
      int * tmp = sched->node_backend_ids;
      sched->node_backend_ids = sched->prev_node_backend_ids;
      sched->prev_node_backend_ids = tmp;

      tmp = sched->leaf_backend_ids;
      sched->leaf_backend_ids = sched->prev_leaf_backend_ids;
      sched->prev_leaf_backend_ids = tmp;
  }

  int graph_size = std::max(graph->n_nodes, graph->n_leafs) + sched->n_splits*GGML_SCHED_MAX_SPLIT_INPUTS30*2*sched->n_copies;

  // remember the actual graph_size for performing reallocation checks later [GGML_SCHED_DEBUG_REALLOC]
  sched->debug_prev_graph_size = sched->debug_graph_size;
  sched->debug_graph_size = graph_size;

  if (sched->graph.size < graph_size) {
      sched->graph.size = graph_size;
      sched->graph.nodes = (ggml_tensor **) realloc(sched->graph.nodes, graph_size * sizeof(struct ggml_tensor *));
      sched->graph.leafs = (ggml_tensor **) realloc(sched->graph.leafs, graph_size * sizeof(struct ggml_tensor *));
      GGML_ASSERT(sched->graph.nodes != NULL)if (!(sched->graph.nodes != __null)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1403, "GGML_ASSERT(%s) failed", "sched->graph.nodes != NULL"
);
      GGML_ASSERT(sched->graph.leafs != NULL)if (!(sched->graph.leafs != __null)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1404, "GGML_ASSERT(%s) failed", "sched->graph.leafs != NULL"
);
  }
  sched->graph.n_nodes = 0;
  sched->graph.n_leafs = 0;

  struct ggml_cgraph * graph_copy = &sched->graph;

  for (int i = 0; i < sched->n_splits; i++) {
      struct ggml_backend_sched_split * split = &sched->splits[i];
      split->graph = ggml_graph_view(graph, split->i_start, split->i_end);

      // Optimize this split of the graph. This needs to happen before we make graph_copy,
      // so they are in sync.
      ggml_backend_graph_optimize(sched->backends[split->backend_id], &split->graph);

      // add inputs to the graph copy so that they are allocated by ggml-alloc at the start of the split
      for (int j = 0; j < split->n_inputs; j++) {
          assert(graph_copy->size > (graph_copy->n_nodes + 1))(static_cast <bool> (graph_copy->size > (graph_copy
->n_nodes + 1)) ? void (0) : __assert_fail ("graph_copy->size > (graph_copy->n_nodes + 1)"
, __builtin_FILE (), __builtin_LINE (), __extension__ __PRETTY_FUNCTION__
));

          struct ggml_tensor * input = split->inputs[j];
          const size_t input_id = hash_id(input)ggml_hash_find_or_insert(&sched->hash_set, input);
          struct ggml_tensor * input_cpy = tensor_id_copy(input_id, split->backend_id, sched->cur_copy)sched->hv_tensor_copies[(input_id) * sched->n_backends *
 sched->n_copies + (split->backend_id) * sched->n_copies
 + (sched->cur_copy)];

          // add a dependency to the input source so that it is not freed before the copy is done
          struct ggml_tensor * input_dep = ggml_view_tensor(sched->ctx, input);
          input_dep->src[0] = input;
          sched->node_backend_ids[graph_copy->n_nodes] = sched->hv_tensor_backend_ids[input_id];
          graph_copy->nodes[graph_copy->n_nodes++] = input_dep;

          // add a dependency to the input copy so that it is allocated at the start of the split
          sched->node_backend_ids[graph_copy->n_nodes] = split->backend_id;
          graph_copy->nodes[graph_copy->n_nodes++] = input_cpy;
      }

      for (int j = split->i_start; j < split->i_end; j++) {
          assert(graph_copy->size > graph_copy->n_nodes)(static_cast <bool> (graph_copy->size > graph_copy
->n_nodes) ? void (0) : __assert_fail ("graph_copy->size > graph_copy->n_nodes"
, __builtin_FILE (), __builtin_LINE (), __extension__ __PRETTY_FUNCTION__
));
          sched->node_backend_ids[graph_copy->n_nodes] = tensor_backend_id(graph->nodes[j])sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, graph->nodes[j])];
          graph_copy->nodes[graph_copy->n_nodes++] = graph->nodes[j];
      }
  }

  if (sched->n_copies > 1) {
      // add input copies as leafs so that they are allocated first
      for (int i = 0; i < sched->n_graph_inputs; i++) {
          struct ggml_tensor * input = sched->graph_inputs[i];
          size_t id = hash_id(input)ggml_hash_find_or_insert(&sched->hash_set, input);
          int backend_id = tensor_backend_id(input)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, input)];
          for (int c = 0; c < sched->n_copies; c++) {
              struct ggml_tensor * input_cpy = tensor_id_copy(id, backend_id, c)sched->hv_tensor_copies[(id) * sched->n_backends * sched
->n_copies + (backend_id) * sched->n_copies + (c)];
              sched->leaf_backend_ids[graph_copy->n_leafs] = backend_id;
              assert(graph_copy->size > graph_copy->n_leafs)(static_cast <bool> (graph_copy->size > graph_copy
->n_leafs) ? void (0) : __assert_fail ("graph_copy->size > graph_copy->n_leafs"
, __builtin_FILE (), __builtin_LINE (), __extension__ __PRETTY_FUNCTION__
));
              graph_copy->leafs[graph_copy->n_leafs++] = input_cpy;
          }
      }

      for (int i = 0; i < sched->n_splits; i++) {
          struct ggml_backend_sched_split * split = &sched->splits[i];
          int backend_id = split->backend_id;
          for (int j = 0; j < split->n_inputs; j++) {
              struct ggml_tensor * input = split->inputs[j];
              size_t id = hash_id(input)ggml_hash_find_or_insert(&sched->hash_set, input);
              for (int c = 0; c < sched->n_copies; c++) {
                  struct ggml_tensor * input_cpy = tensor_id_copy(id, backend_id, c)sched->hv_tensor_copies[(id) * sched->n_backends * sched
->n_copies + (backend_id) * sched->n_copies + (c)];
                  sched->leaf_backend_ids[graph_copy->n_leafs] = backend_id;
                  assert(graph_copy->size > graph_copy->n_leafs)(static_cast <bool> (graph_copy->size > graph_copy
->n_leafs) ? void (0) : __assert_fail ("graph_copy->size > graph_copy->n_leafs"
, __builtin_FILE (), __builtin_LINE (), __extension__ __PRETTY_FUNCTION__
));
                  graph_copy->leafs[graph_copy->n_leafs++] = input_cpy;
              }
          }
      }
  }

  // add leafs from the original graph
  for (int i = 0; i < graph->n_leafs; i++) {
      struct ggml_tensor * leaf = graph->leafs[i];
      sched->leaf_backend_ids[graph_copy->n_leafs] = tensor_backend_id(leaf)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, leaf)];
      assert(graph_copy->size > graph_copy->n_leafs)(static_cast <bool> (graph_copy->size > graph_copy
->n_leafs) ? void (0) : __assert_fail ("graph_copy->size > graph_copy->n_leafs"
, __builtin_FILE (), __builtin_LINE (), __extension__ __PRETTY_FUNCTION__
));
      graph_copy->leafs[graph_copy->n_leafs++] = leaf;
  }

  // set ids for all splits
  for (int i = 0; i < sched->n_splits; ++i) {
      sched->splits[i].graph.uid = ggml_graph_next_uid();
  }
1487}

1489static bool ggml_backend_sched_alloc_splits(ggml_backend_sched_t sched) {
  bool backend_ids_changed = false;
  for (int i = 0; i < sched->graph.n_nodes; i++) {
      if (sched->node_backend_ids[i] != sched->prev_node_backend_ids[i] &&
          sched->bufts[sched->node_backend_ids[i]] != sched->bufts[sched->prev_node_backend_ids[i]]) {
          backend_ids_changed = true;
          break;
      }
  }
  if (!backend_ids_changed) {
      for (int i = 0; i < sched->graph.n_leafs; i++) {
          if (sched->leaf_backend_ids[i] != sched->prev_leaf_backend_ids[i] &&
              sched->bufts[sched->leaf_backend_ids[i]] != sched->bufts[sched->prev_leaf_backend_ids[i]]) {
              backend_ids_changed = true;
              break;
          }
      }
  }

  // allocate graph
  if (backend_ids_changed || !ggml_gallocr_alloc_graph(sched->galloc, &sched->graph)) {
1510#ifndef NDEBUG
      GGML_LOG_DEBUG("%s: failed to allocate graph, reserving (backend_ids_changed = %d)\n", __func__, backend_ids_changed)ggml_log_internal(GGML_LOG_LEVEL_DEBUG, "%s: failed to allocate graph, reserving (backend_ids_changed = %d)\n"
, __func__, backend_ids_changed);
1512#endif

      if (sched->debug_realloc > 0) {
          // we are interested only in situations where the graph was reallocated even though its size remained the same [GGML_SCHED_DEBUG_REALLOC]
          // example: https://github.com/ggml-org/llama.cpp/pull/17143
          const bool unexpected = !backend_ids_changed && sched->debug_prev_graph_size == sched->debug_graph_size;

          if (unexpected || sched->debug_realloc > 1) {
              GGML_ABORT("%s: unexpected graph reallocation (graph size = %d, nodes = %d, leafs = %d), debug_realloc = %d\n", __func__,ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1521, "%s: unexpected graph reallocation (graph size = %d, nodes = %d, leafs = %d), debug_realloc = %d\n"
, __func__, sched->debug_graph_size, sched->graph.n_nodes
, sched->graph.n_leafs, sched->debug_realloc)
                      sched->debug_graph_size, sched->graph.n_nodes, sched->graph.n_leafs, sched->debug_realloc)ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1521, "%s: unexpected graph reallocation (graph size = %d, nodes = %d, leafs = %d), debug_realloc = %d\n"
, __func__, sched->debug_graph_size, sched->graph.n_nodes
, sched->graph.n_leafs, sched->debug_realloc);
          }
      }

      // the re-allocation may cause the split inputs to be moved to a different address
      // synchronize without ggml_backend_sched_synchronize to avoid changing cur_copy
      for (int i = 0; i < sched->n_backends; i++) {
          ggml_backend_synchronize(sched->backends[i]);
      }

      ggml_gallocr_reserve_n(sched->galloc, &sched->graph, sched->node_backend_ids, sched->leaf_backend_ids);
      if (!ggml_gallocr_alloc_graph(sched->galloc, &sched->graph)) {
          GGML_LOG_ERROR("%s: failed to allocate graph\n", __func__)ggml_log_internal(GGML_LOG_LEVEL_ERROR, "%s: failed to allocate graph\n"
, __func__);
          return false;
      }
  }

  return true;
1539}

1541static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t sched) {
  GGML_ASSERT(sched)if (!(sched)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1542, "GGML_ASSERT(%s) failed", "sched");
  struct ggml_backend_sched_split * splits = sched->splits;

  ggml_tensor * prev_ids_tensor = nullptr;
  std::vector<int32_t> ids;
  std::vector<ggml_bitset_t> used_ids;

  for (int split_id = 0; split_id < sched->n_splits; split_id++) {
      struct ggml_backend_sched_split * split = &splits[split_id];
      int split_backend_id = split->backend_id;
      ggml_backend_t split_backend = sched->backends[split_backend_id];

      // copy the input tensors to the split backend
      for (int input_id = 0; input_id < split->n_inputs; input_id++) {
          ggml_backend_t input_backend = ggml_backend_sched_get_tensor_backend(sched, split->inputs[input_id]);
          struct ggml_tensor * input = split->inputs[input_id];
          struct ggml_tensor * input_cpy = tensor_copy(input, split_backend_id, sched->cur_copy)sched->hv_tensor_copies[(ggml_hash_find_or_insert(&sched
->hash_set, input)) * sched->n_backends * sched->n_copies
 + (split_backend_id) * sched->n_copies + (sched->cur_copy
)];

          if (input->flags & GGML_TENSOR_FLAG_INPUT) {
              // inputs from the user must be copied immediately to prevent the user overwriting the data before the copy is done
              if (sched->events[split_backend_id][sched->cur_copy] != NULL__null) {
                  ggml_backend_event_synchronize(sched->events[split_backend_id][sched->cur_copy]);
              } else {
                  ggml_backend_synchronize(split_backend);
              }
              ggml_backend_tensor_copy(input, input_cpy);
          } else {
              // wait for the split backend to finish using the input before overwriting it
              if (sched->events[split_backend_id][sched->cur_copy] != NULL__null) {
                  ggml_backend_event_wait(split_backend, sched->events[split_backend_id][sched->cur_copy]);
              } else {
                  ggml_backend_synchronize(split_backend);
              }

              // when offloading MoE weights, we can reduce the amount of data copied by copying only the experts that are used
              ggml_tensor * node = split->graph.nodes[0];
              if (split->graph.n_nodes > 0 &&
                  ggml_backend_buffer_get_usage(input->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS &&
                  ggml_backend_buffer_is_host(input->buffer) && (
                  (node->src[0] == input_cpy && node->op == GGML_OP_MUL_MAT_ID)
                  //|| (node->src[1] == input_cpy && node->op == GGML_OP_ADD_ID) /* GGML_OP_ADD_ID weights are small and not worth splitting */
                  )) {

                  const int64_t n_expert   = node->op == GGML_OP_MUL_MAT_ID ? input->ne[2] : input->ne[1];
                  const size_t expert_size = node->op == GGML_OP_MUL_MAT_ID ? input->nb[2] : input->nb[1];

                  ggml_backend_synchronize(input_backend);

                  // get the ids
                  ggml_tensor * ids_tensor = node->src[2];
                  ggml_backend_t ids_backend = split_backend;

                  // if the ids tensor is also an input of the split, it may not have been copied yet to the split backend
                  // in that case, we use the original ids tensor
                  for (int i = input_id + 1; i < split->n_inputs; i++) {
                      if (ids_tensor == tensor_copy(split->inputs[i], split_backend_id, sched->cur_copy)sched->hv_tensor_copies[(ggml_hash_find_or_insert(&sched
->hash_set, split->inputs[i])) * sched->n_backends *
 sched->n_copies + (split_backend_id) * sched->n_copies
 + (sched->cur_copy)]) {
                          ids_tensor = split->inputs[i];
                          ids_backend = ggml_backend_sched_get_tensor_backend(sched, split->inputs[i]);
                          break;
                      }
                  }

                  if (ids_tensor != prev_ids_tensor) {
                      ids.resize(ggml_nbytes(ids_tensor) / sizeof(int32_t));
                      ggml_backend_tensor_get_async(ids_backend, ids_tensor, ids.data(), 0, ggml_nbytes(ids_tensor));
                      ggml_backend_synchronize(ids_backend);

                      // find the used experts
                      used_ids.clear();
                      used_ids.resize(ggml_bitset_size(n_expert));
                      for (int64_t i1 = 0; i1 < ids_tensor->ne[1]; i1++) {
                          for (int64_t i0 = 0; i0 < ids_tensor->ne[0]; i0++) {
                              int32_t id = ids[i1 * ids_tensor->nb[1]/sizeof(int32_t) + i0 * ids_tensor->nb[0]/sizeof(int32_t)];
                              GGML_ASSERT(id >= 0 && id < n_expert)if (!(id >= 0 && id < n_expert)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1615, "GGML_ASSERT(%s) failed", "id >= 0 && id < n_expert"
);
                              ggml_bitset_set(used_ids.data(), id);
                          }
                      }

                      prev_ids_tensor = ids_tensor;
                  }

                  // group consecutive experts and copy them together
                  auto copy_experts = [&](int32_t first_id, int32_t last_id) {
                      const size_t expert_offset = first_id * expert_size;
                      const size_t expert_size_copy =  (last_id - first_id + 1) * expert_size;
                      const size_t padding = std::min<size_t>(expert_size, 512);
                      const size_t padding_end = last_id < n_expert - 1 ? padding : 0;

                      ggml_backend_tensor_set_async(split_backend,
                          input_cpy,
                          (const uint8_t *)input->data + expert_offset, expert_offset,
                          // copy a bit extra at the to ensure there are no NaNs in the padding of the last expert
                          // this is necessary for MMQ in the CUDA backend
                          expert_size_copy + padding_end);
                  };

                  int id = 0;
                  while (!ggml_bitset_get(used_ids.data(), id)) {
                      id++;
                  }
                  int32_t first_id = id;
                  int32_t last_id = first_id;

                  for (++id; id < n_expert; ++id) {
                      if (!ggml_bitset_get(used_ids.data(), id)) {
                          continue;
                      }

                      if (id == last_id + 1) {
                          last_id = id;
                          continue;
                      }

                      copy_experts(first_id, last_id);

                      first_id = id;
                      last_id = id;
                  }
                  copy_experts(first_id, last_id);
              } else {
                  // try async copy, but if not possible, we can still use a sync copy without synchronizing the dst backend, since we handle the synchronization here with multiple copies and events
                  // TODO: add public function to facilitate this, since applications do not have direct access to the backend interface
                  if (!split_backend->iface.cpy_tensor_async || !split_backend->iface.cpy_tensor_async(input_backend, split_backend, input, input_cpy)) {
                      ggml_backend_synchronize(input_backend);
                      if (sched->events[split_backend_id][sched->cur_copy] != NULL__null) {
                          ggml_backend_event_synchronize(sched->events[split_backend_id][sched->cur_copy]);
                      } else {
                          ggml_backend_synchronize(split_backend);
                      }
                      ggml_backend_tensor_copy(input, input_cpy);
                  }
              }
          }
      }

      if (!sched->callback_eval) {
          enum ggml_status ec = ggml_backend_graph_compute_async(split_backend, &split->graph);
          if (ec != GGML_STATUS_SUCCESS) {
              return ec;
          }
      } else {
          // similar to ggml_backend_compare_graph_backend
          for (int j0 = 0; j0 < split->graph.n_nodes; j0++) {
              struct ggml_tensor * t = split->graph.nodes[j0];

              // check if the user needs data from this node
              bool need = sched->callback_eval(t, true, sched->callback_eval_user_data);

              int j1 = j0;

              // determine the range [j0, j1] of nodes that can be computed together
              while (!need && j1 < split->graph.n_nodes - 1) {
                  t = split->graph.nodes[++j1];
                  need = sched->callback_eval(t, true, sched->callback_eval_user_data);
              }

              struct ggml_cgraph gv = ggml_graph_view(&split->graph, j0, j1 + 1);

              enum ggml_status ec = ggml_backend_graph_compute_async(split_backend, &gv);
              if (ec != GGML_STATUS_SUCCESS) {
                  return ec;
              }

              // TODO: pass backend to the callback, then the user can decide if they want to synchronize
              ggml_backend_synchronize(split_backend);

              if (need && !sched->callback_eval(t, false, sched->callback_eval_user_data)) {
                  break;
              }

              j0 = j1;
          }
      }

      // record the event of this copy
      if (split->n_inputs > 0) {
          if (sched->events[split_backend_id][sched->cur_copy] != NULL__null) {
              ggml_backend_event_record(sched->events[split_backend_id][sched->cur_copy], split_backend);
          }
      }
  }

  return GGML_STATUS_SUCCESS;
1725}

1727ggml_backend_sched_t ggml_backend_sched_new(
      ggml_backend_t * backends,
      ggml_backend_buffer_type_t * bufts,
      int n_backends,
      size_t graph_size,
      bool parallel,
      bool op_offload) {
  GGML_ASSERT(n_backends > 0)if (!(n_backends > 0)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1734, "GGML_ASSERT(%s) failed", "n_backends > 0");
  GGML_ASSERT(n_backends <= GGML_SCHED_MAX_BACKENDS)if (!(n_backends <= 16)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1735, "GGML_ASSERT(%s) failed", "n_backends <= GGML_SCHED_MAX_BACKENDS"
);
  GGML_ASSERT(ggml_backend_dev_type(ggml_backend_get_device(backends[n_backends - 1])) == GGML_BACKEND_DEVICE_TYPE_CPU)if (!(ggml_backend_dev_type(ggml_backend_get_device(backends[
n_backends - 1])) == GGML_BACKEND_DEVICE_TYPE_CPU)) ggml_abort
("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1736, "GGML_ASSERT(%s) failed", "ggml_backend_dev_type(ggml_backend_get_device(backends[n_backends - 1])) == GGML_BACKEND_DEVICE_TYPE_CPU"
);

  struct ggml_backend_sched * sched = (ggml_backend_sched *) calloc(1, sizeof(struct ggml_backend_sched));

  const char * GGML_SCHED_DEBUG = getenv("GGML_SCHED_DEBUG");
  sched->debug = GGML_SCHED_DEBUG ? atoi(GGML_SCHED_DEBUG) : 0;

  sched->debug_realloc = 0;
1744#ifdef GGML_SCHED_NO_REALLOC
  sched->debug_realloc = 1;
1746#endif
  const char * GGML_SCHED_DEBUG_REALLOC = getenv("GGML_SCHED_DEBUG_REALLOC");
  sched->debug_realloc = GGML_SCHED_DEBUG_REALLOC ? atoi(GGML_SCHED_DEBUG_REALLOC) : sched->debug_realloc;

  sched->n_backends = n_backends;
  sched->n_copies = parallel ? GGML_SCHED_MAX_COPIES4 : 1;

  // initialize hash table
  // FIXME: needs to be size*2 to account for leafs (do it in graph_split instead)
  sched->hash_set    = ggml_hash_set_new(graph_size);
  sched->hv_tensor_backend_ids = (int *) malloc(sched->hash_set.size * sizeof(sched->hv_tensor_backend_ids[0]));
  sched->hv_tensor_copies      = (ggml_tensor **) malloc(sched->hash_set.size * sched->n_backends * sched->n_copies * sizeof(struct ggml_tensor *));

  const size_t ggml_sched_max_splits = graph_size; // at most there is one split for each node in the graph
  const size_t nodes_size = graph_size + ggml_sched_max_splits*GGML_SCHED_MAX_SPLIT_INPUTS30*2;
  sched->node_backend_ids = (int *) calloc(nodes_size, sizeof(sched->node_backend_ids[0]));
  sched->leaf_backend_ids = (int *) calloc(nodes_size, sizeof(sched->leaf_backend_ids[0]));
  sched->prev_node_backend_ids = (int *) calloc(nodes_size, sizeof(sched->prev_node_backend_ids[0]));
  sched->prev_leaf_backend_ids = (int *) calloc(nodes_size, sizeof(sched->prev_leaf_backend_ids[0]));

  sched->debug_graph_size = 0;
  sched->debug_prev_graph_size = 0;

  sched->context_buffer_size = ggml_sched_max_splits*GGML_SCHED_MAX_SPLIT_INPUTS30*2*sizeof(struct ggml_tensor) + ggml_graph_overhead_custom(graph_size, false);
  sched->context_buffer = (char *) malloc(sched->context_buffer_size);

  const int initial_splits_capacity = 16;
  sched->splits = (ggml_backend_sched_split *) calloc(initial_splits_capacity, sizeof(sched->splits[0]));
  sched->splits_capacity = initial_splits_capacity;

  for (int b = 0; b < n_backends; b++) {
      sched->backends[b] = backends[b];
      sched->bufts[b] = bufts ? bufts[b] : ggml_backend_get_default_buffer_type(backends[b]);
      GGML_ASSERT(ggml_backend_supports_buft(backends[b], sched->bufts[b]))if (!(ggml_backend_supports_buft(backends[b], sched->bufts
[b]))) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1779, "GGML_ASSERT(%s) failed", "ggml_backend_supports_buft(backends[b], sched->bufts[b])"
);

      if (sched->n_copies > 1) {
          for (int c = 0; c < sched->n_copies; c++) {
              sched->events[b][c] = ggml_backend_event_new(backends[b]->device);
          }
      }
  }

  sched->galloc = ggml_gallocr_new_n(sched->bufts, n_backends);
  sched->op_offload = op_offload;

  ggml_backend_sched_reset(sched);

  return sched;
1794}

1796void ggml_backend_sched_free(ggml_backend_sched_t sched) {
  if (sched == NULL__null) {
      return;
  }
  for (int b = 0; b < sched->n_backends; b++) {
      for (int c = 0; c < sched->n_copies; c++) {
          ggml_backend_event_free(sched->events[b][c]);
      }
  }
  ggml_gallocr_free(sched->galloc);
  ggml_free(sched->ctx);
  ggml_hash_set_free(&sched->hash_set);
  free(sched->splits);
  free(sched->hv_tensor_backend_ids);
  free(sched->hv_tensor_copies);
  free(sched->node_backend_ids);
  free(sched->leaf_backend_ids);
  free(sched->prev_node_backend_ids);
  free(sched->prev_leaf_backend_ids);
  free(sched->context_buffer);
  free(sched->graph.nodes);
  free(sched->graph.leafs);
  free(sched);
1819}

1821void ggml_backend_sched_reset(ggml_backend_sched_t sched) {
  GGML_ASSERT(sched)if (!(sched)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1822, "GGML_ASSERT(%s) failed", "sched");
  // reset state for the next run
  if (!sched->is_reset) {
      ggml_hash_set_reset(&sched->hash_set);
      memset(sched->hv_tensor_backend_ids, -1, sched->hash_set.size * sizeof(sched->hv_tensor_backend_ids[0]));
      memset(sched->hv_tensor_copies,       0, sched->hash_set.size * sched->n_backends * sched->n_copies * sizeof(struct ggml_tensor *));
      sched->is_reset = true;
  }
  sched->is_alloc = false;
1831}

1833void ggml_backend_sched_reserve_size(ggml_backend_sched_t sched, struct ggml_cgraph * measure_graph, size_t * sizes) {
  GGML_ASSERT(sched)if (!(sched)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1834, "GGML_ASSERT(%s) failed", "sched");
  GGML_ASSERT((int)sched->hash_set.size >= measure_graph->n_nodes + measure_graph->n_leafs)if (!((int)sched->hash_set.size >= measure_graph->n_nodes
 + measure_graph->n_leafs)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1835, "GGML_ASSERT(%s) failed", "(int)sched->hash_set.size >= measure_graph->n_nodes + measure_graph->n_leafs"
);
  GGML_ASSERT(sizes)if (!(sizes)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1836, "GGML_ASSERT(%s) failed", "sizes");

  ggml_backend_sched_reset(sched);

  ggml_backend_sched_synchronize(sched);

  ggml_backend_sched_split_graph(sched, measure_graph);

  ggml_gallocr_reserve_n_size(sched->galloc, &sched->graph, sched->node_backend_ids, sched->leaf_backend_ids, sizes);
1845}

1847bool ggml_backend_sched_reserve(ggml_backend_sched_t sched, struct ggml_cgraph * measure_graph) {
  GGML_ASSERT(sched)if (!(sched)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1848, "GGML_ASSERT(%s) failed", "sched");
  GGML_ASSERT((int)sched->hash_set.size >= measure_graph->n_nodes + measure_graph->n_leafs)if (!((int)sched->hash_set.size >= measure_graph->n_nodes
 + measure_graph->n_leafs)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1849, "GGML_ASSERT(%s) failed", "(int)sched->hash_set.size >= measure_graph->n_nodes + measure_graph->n_leafs"
);

  ggml_backend_sched_synchronize(sched);

  ggml_backend_sched_split_graph(sched, measure_graph);

  if (!ggml_gallocr_reserve_n(sched->galloc, &sched->graph, sched->node_backend_ids, sched->leaf_backend_ids)) {
      return false;
  }

  ggml_backend_sched_reset(sched);

  return true;
1862}

1864bool ggml_backend_sched_alloc_graph(ggml_backend_sched_t sched, struct ggml_cgraph * graph) {
  GGML_ASSERT(sched)if (!(sched)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1865, "GGML_ASSERT(%s) failed", "sched");
  GGML_ASSERT((int)sched->hash_set.size >= graph->n_nodes + graph->n_leafs)if (!((int)sched->hash_set.size >= graph->n_nodes + graph
->n_leafs)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1866, "GGML_ASSERT(%s) failed", "(int)sched->hash_set.size >= graph->n_nodes + graph->n_leafs"
);
  GGML_ASSERT(!sched->is_alloc)if (!(!sched->is_alloc)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1867, "GGML_ASSERT(%s) failed", "!sched->is_alloc");

  sched->cur_copy = sched->next_copy;
  sched->next_copy = (sched->next_copy + 1) % sched->n_copies;

  ggml_backend_sched_split_graph(sched, graph);

  if (!ggml_backend_sched_alloc_splits(sched)) {
      return false;
  }

  sched->is_alloc = true;

  return true;
1881}

1883enum ggml_status ggml_backend_sched_graph_compute(ggml_backend_sched_t sched, struct ggml_cgraph * graph) {
  enum ggml_status err = ggml_backend_sched_graph_compute_async(sched, graph);
  ggml_backend_sched_synchronize(sched);
  return err;
1887}

1889enum ggml_status ggml_backend_sched_graph_compute_async(ggml_backend_sched_t sched, struct ggml_cgraph * graph) {
  GGML_ASSERT(sched)if (!(sched)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1890, "GGML_ASSERT(%s) failed", "sched");
  if (!sched->is_reset && !sched->is_alloc) {
      ggml_backend_sched_reset(sched);
  }

  if (!sched->is_alloc) {
      if (!ggml_backend_sched_alloc_graph(sched, graph)) {
          return GGML_STATUS_ALLOC_FAILED;
      }
  }

  return ggml_backend_sched_compute_splits(sched);
1902}

1904void ggml_backend_sched_synchronize(ggml_backend_sched_t sched) {
  GGML_ASSERT(sched)if (!(sched)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1905, "GGML_ASSERT(%s) failed", "sched");
  for (int i = 0; i < sched->n_backends; i++) {
      ggml_backend_synchronize(sched->backends[i]);
  }
  if (!sched->is_alloc) {
      // if the graph is not already allocated, always use copy 0 after a synchronization
      // this ensures that during generation the same copy is used every time,
      // which avoids changes in the graph that could cause CUDA or other graphs to be disabled
      sched->next_copy = 0;
  }
1915}

1917void ggml_backend_sched_set_eval_callback(ggml_backend_sched_t sched, ggml_backend_sched_eval_callback callback, void * user_data) {
  GGML_ASSERT(sched)if (!(sched)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1918, "GGML_ASSERT(%s) failed", "sched");
  sched->callback_eval = callback;
  sched->callback_eval_user_data = user_data;
1921}

1923int ggml_backend_sched_get_n_splits(ggml_backend_sched_t sched) {
  GGML_ASSERT(sched)if (!(sched)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1924, "GGML_ASSERT(%s) failed", "sched");
  return sched->n_splits;
1926}

1928int ggml_backend_sched_get_n_copies(ggml_backend_sched_t sched) {
  GGML_ASSERT(sched)if (!(sched)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1929, "GGML_ASSERT(%s) failed", "sched");
  return sched->n_copies;
1931}

1933int ggml_backend_sched_get_n_backends(ggml_backend_sched_t sched) {
  GGML_ASSERT(sched)if (!(sched)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1934, "GGML_ASSERT(%s) failed", "sched");
  return sched->n_backends;
1936}

1938ggml_backend_t ggml_backend_sched_get_backend(ggml_backend_sched_t sched, int i) {
  GGML_ASSERT(sched)if (!(sched)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1939, "GGML_ASSERT(%s) failed", "sched");
  GGML_ASSERT(i >= 0 && i < sched->n_backends)if (!(i >= 0 && i < sched->n_backends)) ggml_abort
("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1940, "GGML_ASSERT(%s) failed", "i >= 0 && i < sched->n_backends"
);
  return sched->backends[i];
1942}

1944ggml_backend_buffer_type_t ggml_backend_sched_get_buffer_type(ggml_backend_sched_t sched, ggml_backend_t backend) {
  GGML_ASSERT(sched)if (!(sched)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1945, "GGML_ASSERT(%s) failed", "sched");
  int backend_index = ggml_backend_sched_backend_id(sched, backend);
  GGML_ASSERT(backend_index >= 0 && backend_index < sched->n_backends)if (!(backend_index >= 0 && backend_index < sched
->n_backends)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1947, "GGML_ASSERT(%s) failed", "backend_index >= 0 && backend_index < sched->n_backends"
);

  return sched->bufts[backend_index];
1950}

1952size_t ggml_backend_sched_get_buffer_size(ggml_backend_sched_t sched, ggml_backend_t backend) {
  GGML_ASSERT(sched)if (!(sched)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1953, "GGML_ASSERT(%s) failed", "sched");
  int backend_index = ggml_backend_sched_backend_id(sched, backend);
  GGML_ASSERT(backend_index >= 0 && backend_index < sched->n_backends)if (!(backend_index >= 0 && backend_index < sched
->n_backends)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1955, "GGML_ASSERT(%s) failed", "backend_index >= 0 && backend_index < sched->n_backends"
);

  return ggml_gallocr_get_buffer_size(sched->galloc, backend_index);
1958}

1960void ggml_backend_sched_set_tensor_backend(ggml_backend_sched_t sched, struct ggml_tensor * node, ggml_backend_t backend) {
  GGML_ASSERT(sched)if (!(sched)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1961, "GGML_ASSERT(%s) failed", "sched");
  int backend_index = ggml_backend_sched_backend_id(sched, backend);
  GGML_ASSERT(backend_index >= 0 && backend_index < sched->n_backends)if (!(backend_index >= 0 && backend_index < sched
->n_backends)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1963, "GGML_ASSERT(%s) failed", "backend_index >= 0 && backend_index < sched->n_backends"
);
  tensor_backend_id(node)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, node)] = backend_index;
  SET_CAUSE(node, "usr");
  sched->is_reset = false;
1967}

1969ggml_backend_t ggml_backend_sched_get_tensor_backend(ggml_backend_sched_t sched, struct ggml_tensor * node) {
  GGML_ASSERT(sched)if (!(sched)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1970, "GGML_ASSERT(%s) failed", "sched");
  int backend_index = tensor_backend_id(node)sched->hv_tensor_backend_ids[ggml_hash_find_or_insert(&
sched->hash_set, node)];
  if (backend_index == -1) {
      return NULL__null;
  }
  return sched->backends[backend_index];
1976}

1978// utils

1980enum ggml_status ggml_backend_view_init(struct ggml_tensor * tensor) {
  GGML_ASSERT(tensor)if (!(tensor)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1981, "GGML_ASSERT(%s) failed", "tensor");
  GGML_ASSERT(tensor->buffer == NULL)if (!(tensor->buffer == __null)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1982, "GGML_ASSERT(%s) failed", "tensor->buffer == NULL"
);
  GGML_ASSERT(tensor->view_src != NULL)if (!(tensor->view_src != __null)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1983, "GGML_ASSERT(%s) failed", "tensor->view_src != NULL"
);
  GGML_ASSERT(tensor->view_src->buffer != NULL)if (!(tensor->view_src->buffer != __null)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1984, "GGML_ASSERT(%s) failed", "tensor->view_src->buffer != NULL"
);
  GGML_ASSERT(tensor->view_src->data != NULL)if (!(tensor->view_src->data != __null)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1985, "GGML_ASSERT(%s) failed", "tensor->view_src->data != NULL"
);

  tensor->buffer = tensor->view_src->buffer;
  tensor->data = (char *)tensor->view_src->data + tensor->view_offs;
  return ggml_backend_buffer_init_tensor(tensor->buffer, tensor);
1990}

1992enum ggml_status ggml_backend_tensor_alloc(ggml_backend_buffer_t buffer, struct ggml_tensor * tensor, void * addr) {
  GGML_ASSERT(tensor)if (!(tensor)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1993, "GGML_ASSERT(%s) failed", "tensor");
  GGML_ASSERT(tensor->buffer == NULL)if (!(tensor->buffer == __null)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1994, "GGML_ASSERT(%s) failed", "tensor->buffer == NULL"
);
  GGML_ASSERT(tensor->data == NULL)if (!(tensor->data == __null)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1995, "GGML_ASSERT(%s) failed", "tensor->data == NULL");
  GGML_ASSERT(tensor->view_src == NULL)if (!(tensor->view_src == __null)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1996, "GGML_ASSERT(%s) failed", "tensor->view_src == NULL"
);
  GGML_ASSERT(addr >= ggml_backend_buffer_get_base(buffer))if (!(addr >= ggml_backend_buffer_get_base(buffer))) ggml_abort
("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 1997, "GGML_ASSERT(%s) failed", "addr >= ggml_backend_buffer_get_base(buffer)"
);
  GGML_ASSERT(ggml_backend_buffer_is_meta(buffer) ||if (!(ggml_backend_buffer_is_meta(buffer) || (char *) addr + ggml_backend_buffer_get_alloc_size
(buffer, tensor) <= (char *) ggml_backend_buffer_get_base(
buffer) + ggml_backend_buffer_get_size(buffer))) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 2000, "GGML_ASSERT(%s) failed", "ggml_backend_buffer_is_meta(buffer) || (char *) addr + ggml_backend_buffer_get_alloc_size(buffer, tensor) <= (char *) ggml_backend_buffer_get_base(buffer) + ggml_backend_buffer_get_size(buffer)"
)
      (char *) addr + ggml_backend_buffer_get_alloc_size(buffer, tensor) <=if (!(ggml_backend_buffer_is_meta(buffer) || (char *) addr + ggml_backend_buffer_get_alloc_size
(buffer, tensor) <= (char *) ggml_backend_buffer_get_base(
buffer) + ggml_backend_buffer_get_size(buffer))) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 2000, "GGML_ASSERT(%s) failed", "ggml_backend_buffer_is_meta(buffer) || (char *) addr + ggml_backend_buffer_get_alloc_size(buffer, tensor) <= (char *) ggml_backend_buffer_get_base(buffer) + ggml_backend_buffer_get_size(buffer)"
)
      (char *) ggml_backend_buffer_get_base(buffer) + ggml_backend_buffer_get_size(buffer))if (!(ggml_backend_buffer_is_meta(buffer) || (char *) addr + ggml_backend_buffer_get_alloc_size
(buffer, tensor) <= (char *) ggml_backend_buffer_get_base(
buffer) + ggml_backend_buffer_get_size(buffer))) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 2000, "GGML_ASSERT(%s) failed", "ggml_backend_buffer_is_meta(buffer) || (char *) addr + ggml_backend_buffer_get_alloc_size(buffer, tensor) <= (char *) ggml_backend_buffer_get_base(buffer) + ggml_backend_buffer_get_size(buffer)"
);

  tensor->buffer = buffer;
  tensor->data = addr;
  return ggml_backend_buffer_init_tensor(buffer, tensor);
2005}

2007static struct ggml_tensor * graph_copy_dup_tensor(struct ggml_hash_set hash_set, struct ggml_tensor ** node_copies,
  struct ggml_context * ctx_allocated, struct ggml_context * ctx_unallocated, struct ggml_tensor * src) {

  GGML_ASSERT(src != NULL)if (!(src != __null)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 2010, "GGML_ASSERT(%s) failed", "src != NULL");
11
←
Assuming the condition is true→
  GGML_ASSERT(src->data && "graph must be allocated")if (!(src->data && "graph must be allocated")) ggml_abort
("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 2011, "GGML_ASSERT(%s) failed", "src->data && \"graph must be allocated\""
);
12
←
Taking false branch→
13
←
Assuming field 'data' is non-null→
14
←
Taking false branch→

  size_t id = ggml_hash_insert(&hash_set, src);
  if (id == GGML_HASHSET_ALREADY_EXISTS((size_t)-2)) {
15
←
Taking true branch→
      return node_copies[ggml_hash_find(&hash_set, src)];
  }

  struct ggml_tensor * dst = ggml_dup_tensor_layout(src->data && !src->view_src ? ctx_allocated : ctx_unallocated, src);
  if (src->view_src != NULL__null) {
      dst->view_src = graph_copy_dup_tensor(hash_set, node_copies, ctx_allocated, ctx_unallocated, src->view_src);
      dst->view_offs = src->view_offs;
  }
  dst->op = src->op;
  dst->flags = src->flags;
  memcpy(dst->op_params, src->op_params, sizeof(dst->op_params));
  ggml_set_name(dst, src->name);

  // copy src
  for (int i = 0; i < GGML_MAX_SRC10; i++) {
      struct ggml_tensor * s = src->src[i];
      if (s == NULL__null) {
          continue;
      }
      dst->src[i] = graph_copy_dup_tensor(hash_set, node_copies, ctx_allocated, ctx_unallocated, s);
  }

  node_copies[id] = dst;
  return dst;
2039}

2041static void graph_copy_init_tensor(struct ggml_hash_set * hash_set, struct ggml_tensor ** node_copies, bool * node_init, struct ggml_tensor * src) {
  size_t id = ggml_hash_find(hash_set, src);
23
←
Calling 'ggml_hash_find'→
30
←
Returning from 'ggml_hash_find'→
31
←
'id' initialized to 18446744073709551615→
  if (node_init[id]) {
32
←
Taking false branch→
      return;
  }
  node_init[id] = true;

  struct ggml_tensor * dst = node_copies[id];
33
←
'dst' initialized to a null pointer value→
  if (dst->view_src != NULL__null) {
34
←
Access to field 'view_src' results in a dereference of a null pointer (loaded from variable 'dst')
      graph_copy_init_tensor(hash_set, node_copies, node_init, src->view_src);
      enum ggml_status status = ggml_backend_view_init(dst);
      GGML_ASSERT(status == GGML_STATUS_SUCCESS)if (!(status == GGML_STATUS_SUCCESS)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 2052, "GGML_ASSERT(%s) failed", "status == GGML_STATUS_SUCCESS"
);
  }
  else {
      ggml_backend_tensor_copy(src, dst);
  }

  // init src
  for (int i = 0; i < GGML_MAX_SRC10; i++) {
      struct ggml_tensor * s = src->src[i];
      if (s == NULL__null) {
          continue;
      }
      graph_copy_init_tensor(hash_set, node_copies, node_init, s);
  }
2066}

2068struct ggml_backend_graph_copy ggml_backend_graph_copy(ggml_backend_t backend, struct ggml_cgraph * graph) {
  GGML_ASSERT(graph)if (!(graph)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 2069, "GGML_ASSERT(%s) failed", "graph");
2
←
Assuming 'graph' is non-null→
3
←
Taking false branch→
  struct ggml_hash_set hash_set = ggml_hash_set_new(graph->visited_hash_set.size);
  struct ggml_tensor ** node_copies = (ggml_tensor **) calloc(hash_set.size, sizeof(node_copies[0])); // NOLINT
4
←
Storing null pointer value→
  bool * node_init = (bool *) calloc(hash_set.size, sizeof(node_init[0]));

  struct ggml_init_params params = {
      /* .mem_size   = */ ggml_tensor_overhead()*hash_set.size + ggml_graph_overhead_custom(graph->size, false),
      /* .mem_buffer = */ NULL__null,
      /* .no_alloc   = */ true
  };

  struct ggml_context * ctx_allocated = ggml_init(params);
  struct ggml_context * ctx_unallocated = ggml_init(params);

  if (ctx_allocated == NULL__null || ctx_unallocated == NULL__null) {
5
←
Assuming 'ctx_allocated' is not equal to NULL→
6
←
Assuming 'ctx_unallocated' is not equal to NULL→
7
←
Taking false branch→
      GGML_LOG_ERROR("%s: failed to allocate context for graph copy\n", __func__)ggml_log_internal(GGML_LOG_LEVEL_ERROR, "%s: failed to allocate context for graph copy\n"
, __func__);
      ggml_hash_set_free(&hash_set);
      free(node_copies);
      free(node_init);
      ggml_free(ctx_allocated);
      ggml_free(ctx_unallocated);
      return {
          /* .buffer           = */ NULL__null,
          /* .ctx_allocated    = */ NULL__null,
          /* .ctx_unallocated  = */ NULL__null,
          /* .graph            = */ NULL__null,
      };
  }

  // dup nodes
  for (int i = 0; i < graph->n_nodes; i++) {
8
←
Assuming 'i' is < field 'n_nodes'→
9
←
Loop condition is true.  Entering loop body→
17
←
Assuming 'i' is >= field 'n_nodes'→
18
←
Loop condition is false. Execution continues on line 2105→
      struct ggml_tensor * node = graph->nodes[i];
      graph_copy_dup_tensor(hash_set, node_copies, ctx_allocated, ctx_unallocated, node);
10
←
Calling 'graph_copy_dup_tensor'→
16
←
Returning from 'graph_copy_dup_tensor'→
  }

  // allocate nodes
  ggml_backend_buffer_t buffer = ggml_backend_alloc_ctx_tensors(ctx_allocated, backend);
  if (buffer == NULL__null) {
19
←
Assuming 'buffer' is not equal to NULL→
20
←
Taking false branch→
      GGML_LOG_ERROR("%s: failed to allocate buffer for graph copy\n", __func__)ggml_log_internal(GGML_LOG_LEVEL_ERROR, "%s: failed to allocate buffer for graph copy\n"
, __func__);
      ggml_hash_set_free(&hash_set);
      free(node_copies);
      free(node_init);
      ggml_free(ctx_allocated);
      ggml_free(ctx_unallocated);
      return {
          /* .buffer           = */ NULL__null,
          /* .ctx_allocated    = */ NULL__null,
          /* .ctx_unallocated  = */ NULL__null,
          /* .graph            = */ NULL__null,
      };
  }

  //printf("copy buffer size: %zu MB\n", ggml_backend_buffer_get_size(buffer) / 1024 / 1024);

  // copy data and init views
  for (int i = 0; i < graph->n_nodes; i++) {
21
←
Loop condition is true.  Entering loop body→
      struct ggml_tensor * node = graph->nodes[i];
      graph_copy_init_tensor(&hash_set, node_copies, node_init, node);
22
←
Calling 'graph_copy_init_tensor'→
  }

  // build graph copy
  struct ggml_cgraph * graph_copy = ggml_new_graph_custom(ctx_allocated, graph->size, false);
  for (int i = 0; i < graph->n_nodes; i++) {
      struct ggml_tensor * node = graph->nodes[i];
      struct ggml_tensor * node_copy = node_copies[ggml_hash_find(&hash_set, node)];
      graph_copy->nodes[i] = node_copy;
  }
  graph_copy->n_nodes = graph->n_nodes;

  ggml_hash_set_free(&hash_set);
  free(node_copies);
  free(node_init);

  return {
      /* .buffer           = */ buffer,
      /* .ctx_allocated    = */ ctx_allocated,
      /* .ctx_unallocated  = */ ctx_unallocated,
      /* .graph            = */ graph_copy,
  };
2148}

2150void ggml_backend_graph_copy_free(struct ggml_backend_graph_copy copy) {
  ggml_backend_buffer_free(copy.buffer);
  ggml_free(copy.ctx_allocated);
  ggml_free(copy.ctx_unallocated);
2154}

2156bool ggml_backend_compare_graph_backend(ggml_backend_t backend1, ggml_backend_t backend2, struct ggml_cgraph * graph, ggml_backend_eval_callback callback, void * user_data, struct ggml_tensor const * const * test_nodes, size_t num_test_nodes) {
  struct ggml_backend_graph_copy copy = ggml_backend_graph_copy(backend2, graph);
1
Calling 'ggml_backend_graph_copy'→
  if (copy.buffer == NULL__null) {
      return false;
  }

  struct ggml_cgraph * g1 = graph;
  struct ggml_cgraph * g2 = copy.graph;

  assert(g1->n_nodes == g2->n_nodes)(static_cast <bool> (g1->n_nodes == g2->n_nodes) ?
 void (0) : __assert_fail ("g1->n_nodes == g2->n_nodes"
, __builtin_FILE (), __builtin_LINE (), __extension__ __PRETTY_FUNCTION__
));

  if (num_test_nodes != 0) {
      GGML_ASSERT(test_nodes)if (!(test_nodes)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 2168, "GGML_ASSERT(%s) failed", "test_nodes");
      // Compute the whole graph and only test the output for specific tensors
      ggml_backend_graph_compute(backend1, g1);
      ggml_backend_graph_compute(backend2, g2);

      bool verified = false;
      for (int i = 0; i < g1->n_nodes; i++) {
          for (size_t j = 0; j < num_test_nodes; ++j) {
              if (g1->nodes[i] == test_nodes[j]) {
                  callback(i, g1->nodes[i], g2->nodes[i], user_data);
                  verified = true;
              }
          }
      }
      GGML_ASSERT(verified)if (!(verified)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 2182, "GGML_ASSERT(%s) failed", "verified");
  } else {
      for (int i = 0; i < g1->n_nodes; i++) {
          struct ggml_tensor * t1 = g1->nodes[i];
          struct ggml_tensor * t2 = g2->nodes[i];

          assert(t1->op == t2->op && ggml_are_same_layout(t1, t2))(static_cast <bool> (t1->op == t2->op && ggml_are_same_layout
(t1, t2)) ? void (0) : __assert_fail ("t1->op == t2->op && ggml_are_same_layout(t1, t2)"
, __builtin_FILE (), __builtin_LINE (), __extension__ __PRETTY_FUNCTION__
));

          struct ggml_cgraph g1v = ggml_graph_view(g1, i, i + 1);
          struct ggml_cgraph g2v = ggml_graph_view(g2, i, i + 1);

          ggml_backend_graph_compute(backend1, &g1v);
          ggml_backend_graph_compute(backend2, &g2v);

          if (ggml_is_view_op(t1->op)) {
              continue;
          }

          // compare results, calculate rms etc
          if (!callback(i, t1, t2, user_data)) {
              break;
          }
      }
  }
  ggml_backend_graph_copy_free(copy);

  return true;
2209}

2211// CPU backend - buffer

2213static void * ggml_backend_cpu_buffer_get_base(ggml_backend_buffer_t buffer) {
  GGML_ASSERT(buffer)if (!(buffer)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 2214, "GGML_ASSERT(%s) failed", "buffer");
  uintptr_t data = (uintptr_t)buffer->context;

  // align the buffer
  if (data % TENSOR_ALIGNMENT32 != 0) {
      data = GGML_PAD(data, TENSOR_ALIGNMENT)(((data) + (32) - 1) & ~((32) - 1));
  }

  return (void *)data;
2223}

2225static void ggml_backend_cpu_buffer_free_buffer(ggml_backend_buffer_t buffer) {
  GGML_ASSERT(buffer)if (!(buffer)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 2226, "GGML_ASSERT(%s) failed", "buffer");
  ggml_aligned_free(buffer->context, buffer->size);
2228}

2230static void ggml_backend_cpu_buffer_memset_tensor(ggml_backend_buffer_t buffer, struct ggml_tensor * tensor, uint8_t value, size_t offset, size_t size) {
  GGML_ASSERT(tensor)if (!(tensor)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 2231, "GGML_ASSERT(%s) failed", "tensor");
  memset((char *)tensor->data + offset, value, size);

  GGML_UNUSED(buffer)(void)(buffer);
2235}

2237static void ggml_backend_cpu_buffer_set_tensor(ggml_backend_buffer_t buffer, struct ggml_tensor * tensor, const void * data, size_t offset, size_t size) {
  GGML_ASSERT(tensor)if (!(tensor)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 2238, "GGML_ASSERT(%s) failed", "tensor");
  memcpy((char *)tensor->data + offset, data, size);

  GGML_UNUSED(buffer)(void)(buffer);
2242}

2244static void ggml_backend_cpu_buffer_get_tensor(ggml_backend_buffer_t buffer, const struct ggml_tensor * tensor, void * data, size_t offset, size_t size) {
  GGML_ASSERT(tensor)if (!(tensor)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 2245, "GGML_ASSERT(%s) failed", "tensor");
  memcpy(data, (const char *)tensor->data + offset, size);

  GGML_UNUSED(buffer)(void)(buffer);
2249}

2251static bool ggml_backend_cpu_buffer_cpy_tensor(ggml_backend_buffer_t buffer, const struct ggml_tensor * src, struct ggml_tensor * dst) {
  GGML_ASSERT(src)if (!(src)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 2252, "GGML_ASSERT(%s) failed", "src");
  if (ggml_backend_buffer_is_host(src->buffer)) {
      memcpy(dst->data, src->data, ggml_nbytes(src));
      return true;
  }
  return false;

  GGML_UNUSED(buffer)(void)(buffer);
2260}

2262static void ggml_backend_cpu_buffer_clear(ggml_backend_buffer_t buffer, uint8_t value) {
  GGML_ASSERT(buffer)if (!(buffer)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 2263, "GGML_ASSERT(%s) failed", "buffer");
  memset(buffer->context, value, buffer->size);
2265}

2267static const struct ggml_backend_buffer_i ggml_backend_cpu_buffer_i = {
  /* .free_buffer     = */ ggml_backend_cpu_buffer_free_buffer,
  /* .get_base        = */ ggml_backend_cpu_buffer_get_base,
  /* .init_tensor     = */ NULL__null, // no initialization required
  /* .memset_tensor   = */ ggml_backend_cpu_buffer_memset_tensor,
  /* .set_tensor      = */ ggml_backend_cpu_buffer_set_tensor,
  /* .get_tensor      = */ ggml_backend_cpu_buffer_get_tensor,
  /* .set_tensor_2d   = */ NULL__null,
  /* .get_tensor_2d   = */ NULL__null,
  /* .cpy_tensor      = */ ggml_backend_cpu_buffer_cpy_tensor,
  /* .clear           = */ ggml_backend_cpu_buffer_clear,
  /* .reset           = */ NULL__null,
2279};

2281static const struct ggml_backend_buffer_i ggml_backend_cpu_buffer_from_ptr_i = {
  /* .free_buffer     = */ NULL__null, // ptr is not owned by the buffer, so it does not need to be freed
  /* .get_base        = */ ggml_backend_cpu_buffer_get_base,
  /* .init_tensor     = */ NULL__null, // no initialization required
  /* .memset_tensor   = */ ggml_backend_cpu_buffer_memset_tensor,
  /* .set_tensor      = */ ggml_backend_cpu_buffer_set_tensor,
  /* .get_tensor      = */ ggml_backend_cpu_buffer_get_tensor,
  /* .set_tensor_2d   = */ NULL__null,
  /* .get_tensor_2d   = */ NULL__null,
  /* .cpy_tensor      = */ ggml_backend_cpu_buffer_cpy_tensor,
  /* .clear           = */ ggml_backend_cpu_buffer_clear,
  /* .reset           = */ NULL__null,
2293};

2295// CPU backend buffer type

2297// this buffer type is defined here to make it available to all backends

2299static const char * ggml_backend_cpu_buffer_type_get_name(ggml_backend_buffer_type_t buft) {
  return "CPU";

  GGML_UNUSED(buft)(void)(buft);
2303}

2305static ggml_backend_buffer_t ggml_backend_cpu_buffer_type_alloc_buffer(ggml_backend_buffer_type_t buft, size_t size) {
  void * data = ggml_aligned_malloc(size);

  if (data == NULL__null) {
      GGML_LOG_ERROR("%s: failed to allocate buffer of size %zu\n", __func__, size)ggml_log_internal(GGML_LOG_LEVEL_ERROR, "%s: failed to allocate buffer of size %zu\n"
, __func__, size);
      return NULL__null;
  }

  return ggml_backend_buffer_init(buft, ggml_backend_cpu_buffer_i, data, size);
2314}

2316static size_t ggml_backend_cpu_buffer_type_get_alignment(ggml_backend_buffer_type_t buft) {
  return TENSOR_ALIGNMENT32;

  GGML_UNUSED(buft)(void)(buft);
2320}

2322static bool ggml_backend_cpu_buffer_type_is_host(ggml_backend_buffer_type_t buft) {
  return true;

  GGML_UNUSED(buft)(void)(buft);
2326}

2328ggml_backend_buffer_type_t ggml_backend_cpu_buffer_type(void) {
  static struct ggml_backend_buffer_type ggml_backend_cpu_buffer_type = {
      /* .iface   = */ {
          /* .get_name         = */ ggml_backend_cpu_buffer_type_get_name,
          /* .alloc_buffer     = */ ggml_backend_cpu_buffer_type_alloc_buffer,
          /* .get_alignment    = */ ggml_backend_cpu_buffer_type_get_alignment,
          /* .get_max_size     = */ NULL__null, // defaults to SIZE_MAX
          /* .get_alloc_size   = */ NULL__null, // defaults to ggml_nbytes
          /* .is_host          = */ ggml_backend_cpu_buffer_type_is_host,
      },
      /* .device  = */ NULL__null, // FIXME ggml_backend_reg_dev_get(ggml_backend_cpu_reg(), 0),
      /* .context = */ NULL__null,
  };

  return &ggml_backend_cpu_buffer_type;
2343}

2345static const char * ggml_backend_cpu_buffer_from_ptr_type_get_name(ggml_backend_buffer_type_t buft) {
  return "CPU_Mapped";

  GGML_UNUSED(buft)(void)(buft);
2349}

2351static ggml_backend_buffer_type_t ggml_backend_cpu_buffer_from_ptr_type(void) {
  static struct ggml_backend_buffer_type ggml_backend_cpu_buffer_type = {
      /* .iface   = */ {
          /* .get_name         = */ ggml_backend_cpu_buffer_from_ptr_type_get_name,
          /* .alloc_buffer     = */ ggml_backend_cpu_buffer_type_alloc_buffer,
          /* .get_alignment    = */ ggml_backend_cpu_buffer_type_get_alignment,
          /* .get_max_size     = */ NULL__null, // defaults to SIZE_MAX
          /* .get_alloc_size   = */ NULL__null, // defaults to ggml_nbytes
          /* .is_host          = */ ggml_backend_cpu_buffer_type_is_host,
      },
      /* .device  = */ NULL__null, // FIXME ggml_backend_reg_dev_get(ggml_backend_cpu_reg(), 0),
      /* .context = */ NULL__null,
  };

  return &ggml_backend_cpu_buffer_type;
2366}

2368ggml_backend_buffer_t ggml_backend_cpu_buffer_from_ptr(void * ptr, size_t size) {
  GGML_ASSERT((uintptr_t)ptr % TENSOR_ALIGNMENT == 0 && "buffer pointer must be aligned")if (!((uintptr_t)ptr % 32 == 0 && "buffer pointer must be aligned"
)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-backend.cpp"
, 2369, "GGML_ASSERT(%s) failed", "(uintptr_t)ptr % TENSOR_ALIGNMENT == 0 && \"buffer pointer must be aligned\""
);
  return ggml_backend_buffer_init(ggml_backend_cpu_buffer_from_ptr_type(), ggml_backend_cpu_buffer_from_ptr_i, ptr, size);
2371}

←

/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-impl.h

1#pragma once
2 
3// GGML internal header
4 
5#include "ggml.h"
6#include "gguf.h"
7 
8#include <assert.h>
9#include <math.h>
10#include <stdlib.h> // load `stdlib.h` before other headers to work around MinGW bug: https://sourceforge.net/p/mingw-w64/bugs/192/
11#include <stdbool.h>
12#include <stdint.h>
13#include <string.h>
14 
15#ifdef __ARM_FEATURE_SVE
16#include <arm_sve.h>
17#endif // __ARM_FEATURE_SVE
18 
19#if defined(__ARM_NEON) && !defined(__CUDACC__) && !defined(__MUSACC__)
20// if YCM cannot find <arm_neon.h>, make a symbolic link to it, for example:
21//
22//   $ ln -sfn /Library/Developer/CommandLineTools/usr/lib/clang/13.1.6/include/arm_neon.h ./src/
23//
24#include <arm_neon.h>
25#endif
26 
27#ifdef __cplusplus202002L
28extern "C" {
29#endif
30 
31void ggml_print_backtrace(void);
32 
33uint64_t ggml_graph_next_uid(void);
34 
35#ifndef MIN
36#    define MIN(a, b)((a) < (b) ? (a) : (b)) ((a) < (b) ? (a) : (b))
37#endif
38 
39#ifndef MAX
40#    define MAX(a, b)((a) > (b) ? (a) : (b)) ((a) > (b) ? (a) : (b))
41#endif
42 
43// required for mmap as gguf only guarantees 32-byte alignment
44#define TENSOR_ALIGNMENT32 32
45 
46// static_assert should be a #define, but if it's not,
47// fall back to the _Static_assert C11 keyword.
48// if C99 - static_assert is noop
49// ref: https://stackoverflow.com/a/53923785/4039976
50#ifndef __cplusplus202002L
51    #ifndef static_assert
52        #if defined(__STDC_VERSION__) && (__STDC_VERSION__ >= 201100L)
53            #define static_assert(cond, msg) _Static_assert(cond, msg)
54        #else
55            #define static_assert(cond, msg) struct global_scope_noop_trick
56        #endif
57    #endif
58#endif
59 
60static inline int ggml_up32(int n) {
61    return (n + 31) & ~31;
62}
63 
64//static inline int ggml_up64(int n) {
65//    return (n + 63) & ~63;
66//}
67 
68static inline int ggml_up(int n, int m) {
69    // assert m is a power of 2
70    GGML_ASSERT((m & (m - 1)) == 0)if (!((m & (m - 1)) == 0)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-impl.h"
, 70, "GGML_ASSERT(%s) failed", "(m & (m - 1)) == 0");
71    return (n + m - 1) & ~(m - 1);
72}
73 
74// TODO: move to ggml.h? (won't be able to inline)
75static bool ggml_are_same_layout(const struct ggml_tensor * a, const struct ggml_tensor * b) {
76    if (a->type != b->type) {
77        return false;
78    }
79    for (int i = 0; i < GGML_MAX_DIMS4; i++) {
80        if (a->ne[i] != b->ne[i]) {
81            return false;
82        }
83        if (a->nb[i] != b->nb[i]) {
84            return false;
85        }
86    }
87    return true;
88}
89 
90static bool ggml_op_is_empty(enum ggml_op op) {
91    switch (op) {
92        case GGML_OP_NONE:
93        case GGML_OP_RESHAPE:
94        case GGML_OP_TRANSPOSE:
95        case GGML_OP_VIEW:
96        case GGML_OP_PERMUTE:
97            return true;
98        default:
99            return false;
100    }
101}
102 
103static inline bool ggml_impl_is_view(const struct ggml_tensor * t) {
104    return t->view_src != NULL__null;
105}
106 
107static inline float ggml_compute_softplus_f32(float input) {
108    return (input > 20.0f) ? input : logf(1 + expf(input));
109}
110//
111// logging
112//
113 
114GGML_ATTRIBUTE_FORMAT(2, 3)__attribute__((format(printf, 2, 3)))
115GGML_API__attribute__ ((visibility ("default"))) extern void ggml_log_internal        (enum ggml_log_level level, const char * format, ...);
116GGML_API__attribute__ ((visibility ("default"))) extern void ggml_log_callback_default(enum ggml_log_level level, const char * text, void * user_data);
117 
118#define GGML_LOG(...)ggml_log_internal(GGML_LOG_LEVEL_NONE , ...)       ggml_log_internal(GGML_LOG_LEVEL_NONE , __VA_ARGS__)
119#define GGML_LOG_INFO(...)ggml_log_internal(GGML_LOG_LEVEL_INFO , ...)  ggml_log_internal(GGML_LOG_LEVEL_INFO , __VA_ARGS__)
120#define GGML_LOG_WARN(...)ggml_log_internal(GGML_LOG_LEVEL_WARN , ...)  ggml_log_internal(GGML_LOG_LEVEL_WARN , __VA_ARGS__)
121#define GGML_LOG_ERROR(...)ggml_log_internal(GGML_LOG_LEVEL_ERROR, ...) ggml_log_internal(GGML_LOG_LEVEL_ERROR, __VA_ARGS__)
122#define GGML_LOG_DEBUG(...)ggml_log_internal(GGML_LOG_LEVEL_DEBUG, ...) ggml_log_internal(GGML_LOG_LEVEL_DEBUG, __VA_ARGS__)
123#define GGML_LOG_CONT(...)ggml_log_internal(GGML_LOG_LEVEL_CONT , ...)  ggml_log_internal(GGML_LOG_LEVEL_CONT , __VA_ARGS__)
124 
125#define GGML_DEBUG0 0
126 
127#if (GGML_DEBUG0 >= 1)
128#define GGML_PRINT_DEBUG(...) GGML_LOG_DEBUG(__VA_ARGS__)ggml_log_internal(GGML_LOG_LEVEL_DEBUG, __VA_ARGS__)
129#else
130#define GGML_PRINT_DEBUG(...)
131#endif
132 
133#if (GGML_DEBUG0 >= 5)
134#define GGML_PRINT_DEBUG_5(...) GGML_LOG_DEBUG(__VA_ARGS__)ggml_log_internal(GGML_LOG_LEVEL_DEBUG, __VA_ARGS__)
135#else
136#define GGML_PRINT_DEBUG_5(...)
137#endif
138 
139#if (GGML_DEBUG0 >= 10)
140#define GGML_PRINT_DEBUG_10(...) GGML_LOG_DEBUG(__VA_ARGS__)ggml_log_internal(GGML_LOG_LEVEL_DEBUG, __VA_ARGS__)
141#else
142#define GGML_PRINT_DEBUG_10(...)
143#endif
144 
145// tensor params
146 
147static void ggml_set_op_params(struct ggml_tensor * tensor, const void * params, size_t params_size) {
148    GGML_ASSERT(tensor != NULL)if (!(tensor != __null)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-impl.h"
, 148, "GGML_ASSERT(%s) failed", "tensor != NULL"); // silence -Warray-bounds warnings
149    assert(params_size <= GGML_MAX_OP_PARAMS)(static_cast <bool> (params_size <= 64) ? void (0) :
 __assert_fail ("params_size <= GGML_MAX_OP_PARAMS", __builtin_FILE
 (), __builtin_LINE (), __extension__ __PRETTY_FUNCTION__));
150    memcpy(tensor->op_params, params, params_size);
151}
152 
153static int32_t ggml_get_op_params_i32(const struct ggml_tensor * tensor, uint32_t i) {
154    assert(i < GGML_MAX_OP_PARAMS / sizeof(int32_t))(static_cast <bool> (i < 64 / sizeof(int32_t)) ? void
 (0) : __assert_fail ("i < GGML_MAX_OP_PARAMS / sizeof(int32_t)"
, __builtin_FILE (), __builtin_LINE (), __extension__ __PRETTY_FUNCTION__
));
155    return ((const int32_t *)(tensor->op_params))[i];
156}
157 
158static float ggml_get_op_params_f32(const struct ggml_tensor * tensor, uint32_t i) {
159    assert(i < GGML_MAX_OP_PARAMS / sizeof(float))(static_cast <bool> (i < 64 / sizeof(float)) ? void (
0) : __assert_fail ("i < GGML_MAX_OP_PARAMS / sizeof(float)"
, __builtin_FILE (), __builtin_LINE (), __extension__ __PRETTY_FUNCTION__
));
160    return ((const float *)(tensor->op_params))[i];
161}
162 
163static void ggml_set_op_params_i32(struct ggml_tensor * tensor, uint32_t i, int32_t value) {
164    assert(i < GGML_MAX_OP_PARAMS / sizeof(int32_t))(static_cast <bool> (i < 64 / sizeof(int32_t)) ? void
 (0) : __assert_fail ("i < GGML_MAX_OP_PARAMS / sizeof(int32_t)"
, __builtin_FILE (), __builtin_LINE (), __extension__ __PRETTY_FUNCTION__
));
165    ((int32_t *)(tensor->op_params))[i] = value;
166}
167 
168static void ggml_set_op_params_f32(struct ggml_tensor * tensor, uint32_t i, float value) {
169    assert(i < GGML_MAX_OP_PARAMS / sizeof(float))(static_cast <bool> (i < 64 / sizeof(float)) ? void (
0) : __assert_fail ("i < GGML_MAX_OP_PARAMS / sizeof(float)"
, __builtin_FILE (), __builtin_LINE (), __extension__ __PRETTY_FUNCTION__
));
170    ((float *)(tensor->op_params))[i] = value;
171}
172 
173struct ggml_map_custom1_op_params {
174    ggml_custom1_op_t  fun;
175    int                n_tasks;
176    void             * userdata;
177};
178 
179struct ggml_map_custom2_op_params {
180    ggml_custom2_op_t   fun;
181    int                 n_tasks;
182    void              * userdata;
183};
184 
185struct ggml_map_custom3_op_params {
186    ggml_custom3_op_t fun;
187    int               n_tasks;
188    void            * userdata;
189};
190 
191struct ggml_custom_op_params {
192    ggml_custom_op_t fun;
193    int              n_tasks;
194    void           * userdata;
195};
196 
197// bitset
198 
199typedef uint32_t ggml_bitset_t;
200 
201static_assert(sizeof(ggml_bitset_t) == 4, "bitset_t constants must be updated");
202#define BITSET_SHR5 5 // log2(sizeof(ggml_bitset_t)*8)
203#define BITSET_MASK(sizeof(ggml_bitset_t)*8 - 1) (sizeof(ggml_bitset_t)*8 - 1)
204 
205static size_t ggml_bitset_size(size_t n) {
206    return (n + BITSET_MASK(sizeof(ggml_bitset_t)*8 - 1)) >> BITSET_SHR5;
207}
208 
209static inline bool ggml_bitset_get(const ggml_bitset_t * bitset, size_t i) {
210    return !!(bitset[i >> BITSET_SHR5] & (1u << (i & BITSET_MASK(sizeof(ggml_bitset_t)*8 - 1))));
211}
212 
213static inline void ggml_bitset_set(ggml_bitset_t * bitset, size_t i) {
214    bitset[i >> BITSET_SHR5] |= (1u << (i & BITSET_MASK(sizeof(ggml_bitset_t)*8 - 1)));
215}
216 
217static inline void ggml_bitset_clear(ggml_bitset_t * bitset, size_t i) {
218    bitset[i >> BITSET_SHR5] &= ~(1u << (i & BITSET_MASK(sizeof(ggml_bitset_t)*8 - 1)));
219}
220 
221// hash set
222 
223#define GGML_HASHSET_FULL((size_t)-1) ((size_t)-1)
224#define GGML_HASHSET_ALREADY_EXISTS((size_t)-2) ((size_t)-2)
225 
226struct ggml_hash_set {
227    size_t size;
228    ggml_bitset_t * used;       // whether or not the keys are in use i.e. set
229    struct ggml_tensor ** keys; // actual tensors in the set, keys[i] is only defined if ggml_bitset_get(used, i)
230};
231 
232struct ggml_hash_set ggml_hash_set_new(size_t size);
233void                 ggml_hash_set_free(struct ggml_hash_set * hash_set);
234 
235// returns the minimum size for a hash set that can hold min_sz elements
236size_t ggml_hash_size(size_t min_sz);
237 
238// remove all elements from the hash set
239void ggml_hash_set_reset(struct ggml_hash_set * hash_set);
240 
241// returns true if key is in the hash set
242static bool ggml_hash_contains(const struct ggml_hash_set * hash_set, struct ggml_tensor * key);
243 
244// returns GGML_HASHSET_FULL if table is full, otherwise the current index of the key or where it should be inserted
245static size_t ggml_hash_find(const struct ggml_hash_set * hash_set, const struct ggml_tensor * key);
246 
247// returns GGML_HASHSET_ALREADY_EXISTS if key already exists, index otherwise, asserts if table is full
248static size_t ggml_hash_insert(struct ggml_hash_set * hash_set, struct ggml_tensor * key);
249 
250// return index, asserts if table is full
251static size_t ggml_hash_find_or_insert(struct ggml_hash_set * hash_set, struct ggml_tensor * key);
252 
253// hash function for ggml_tensor
254static inline size_t ggml_hash(const struct ggml_tensor * p) {
255    // the last 4 bits are always zero due to alignment
256    return (size_t)(uintptr_t)p >> 4;
257}
258 
259static size_t ggml_hash_find(const struct ggml_hash_set * hash_set, const struct ggml_tensor * key) {
260    size_t h = ggml_hash(key) % hash_set->size;
261 
262    // linear probing
263    size_t i = h;
264    while (ggml_bitset_get(hash_set->used, i) && hash_set->keys[i] != key) {
24
←
Assuming the condition is true→
25
←
Loop condition is true.  Entering loop body→
265        i = (i + 1) % hash_set->size;
266        if (i == h) {
26
←
Assuming 'i' is equal to 'h'→
27
←
Taking true branch→
267            // visited all hash table entries -> not found
268            return GGML_HASHSET_FULL((size_t)-1);
28
←
Returning the value 18446744073709551615, which participates in a condition later→
29
←
Returning the value 18446744073709551615→
269        }
270    }
271    return i;
272}
273 
274static bool ggml_hash_contains(const struct ggml_hash_set * hash_set, struct ggml_tensor * key) {
275    size_t i = ggml_hash_find(hash_set, key);
276    return i != GGML_HASHSET_FULL((size_t)-1) && ggml_bitset_get(hash_set->used, i);
277}
278 
279static size_t ggml_hash_insert(struct ggml_hash_set * hash_set, struct ggml_tensor * key) {
280    size_t h = ggml_hash(key) % hash_set->size;
281 
282    // linear probing
283    size_t i = h;
284    do {
285        if (!ggml_bitset_get(hash_set->used, i)) {
286            ggml_bitset_set(hash_set->used, i);
287            hash_set->keys[i] = key;
288            return i;
289        }
290        if (hash_set->keys[i] == key) {
291            return GGML_HASHSET_ALREADY_EXISTS((size_t)-2);
292        }
293        i = (i + 1) % hash_set->size;
294    } while (i != h);
295 
296    // visited all hash table entries -> not found
297    GGML_ABORT("fatal error")ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-impl.h"
, 297, "fatal error");
298}
299 
300static size_t ggml_hash_find_or_insert(struct ggml_hash_set * hash_set, struct ggml_tensor * key) {
301    size_t h = ggml_hash(key) % hash_set->size;
302 
303    // linear probing
304    size_t i = h;
305    do {
306        if (!ggml_bitset_get(hash_set->used, i)) {
307            ggml_bitset_set(hash_set->used, i);
308            hash_set->keys[i] = key;
309            return i;
310        }
311        if (hash_set->keys[i] == key) {
312            return i;
313        }
314        i = (i + 1) % hash_set->size;
315    } while (i != h);
316 
317    // visited all hash table entries -> not found
318    GGML_ABORT("fatal error")ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-impl.h"
, 318, "fatal error");
319}
320 
321// computation graph
322 
323enum ggml_cgraph_eval_order {
324    GGML_CGRAPH_EVAL_ORDER_LEFT_TO_RIGHT = 0,
325    GGML_CGRAPH_EVAL_ORDER_RIGHT_TO_LEFT,
326    GGML_CGRAPH_EVAL_ORDER_COUNT
327};
328 
329struct ggml_cgraph {
330    int size;    // maximum number of nodes/leafs/grads/grad_accs
331    int n_nodes; // number of nodes currently in use
332    int n_leafs; // number of leafs currently in use
333 
334    struct ggml_tensor ** nodes;     // tensors with data that can change if the graph is evaluated
335    struct ggml_tensor ** grads;     // the outputs of these tensors are the gradients of the nodes
336    struct ggml_tensor ** grad_accs; // accumulators for node gradients
337    struct ggml_tensor ** leafs;     // tensors with constant data
338    int32_t             * use_counts;// number of uses of each tensor, indexed by hash table slot
339 
340    struct ggml_hash_set visited_hash_set;
341 
342    enum ggml_cgraph_eval_order order;
343 
344    // an optional identifier that can be utilized to recognize same graphs if two non-zero values match
345    // a value of 0 means it is not set and should be ignored
346    uint64_t uid;
347};
348 
349// returns a slice of cgraph with nodes [i0, i1)
350// the slice does not have leafs or gradients
351// if you need the gradients, get them from the original graph
352struct ggml_cgraph ggml_graph_view(struct ggml_cgraph * cgraph, int i0, int i1);
353 
354// ggml-alloc.c: true if the operation can reuse memory from its sources
355GGML_API__attribute__ ((visibility ("default"))) extern bool ggml_op_can_inplace(enum ggml_op op);
356 
357 
358// Memory allocation
359 
360GGML_API__attribute__ ((visibility ("default"))) extern void * ggml_aligned_malloc(size_t size);
361GGML_API__attribute__ ((visibility ("default"))) extern void ggml_aligned_free(void * ptr, size_t size);
362 
363// FP16 <-> FP32
364// ref: https://github.com/Maratyszcza/FP16
365 
366static inline float fp32_from_bits(uint32_t w) {
367    union {
368        uint32_t as_bits;
369        float as_value;
370    } fp32;
371    fp32.as_bits = w;
372    return fp32.as_value;
373}
374 
375static inline uint32_t fp32_to_bits(float f) {
376    union {
377        float as_value;
378        uint32_t as_bits;
379    } fp32;
380    fp32.as_value = f;
381    return fp32.as_bits;
382}
383 
384static inline float ggml_compute_fp16_to_fp32(ggml_fp16_t h) {
385    const uint32_t w = (uint32_t) h << 16;
386    const uint32_t sign = w & UINT32_C(0x80000000)0x80000000U;
387    const uint32_t two_w = w + w;
388 
389    const uint32_t exp_offset = UINT32_C(0xE0)0xE0U << 23;
390#if (defined(__STDC_VERSION__) && (__STDC_VERSION__ >= 199901L) || defined(__GNUC__4) && !defined(__STRICT_ANSI__)) && (!defined(__cplusplus202002L) || __cplusplus202002L >= 201703L)
391    const float exp_scale = 0x1.0p-112f;
392#else
393    const float exp_scale = fp32_from_bits(UINT32_C(0x7800000)0x7800000U);
394#endif
395    const float normalized_value = fp32_from_bits((two_w >> 4) + exp_offset) * exp_scale;
396 
397    const uint32_t magic_mask = UINT32_C(126)126U << 23;
398    const float magic_bias = 0.5f;
399    const float denormalized_value = fp32_from_bits((two_w >> 17) | magic_mask) - magic_bias;
400 
401    const uint32_t denormalized_cutoff = UINT32_C(1)1U << 27;
402    const uint32_t result = sign |
403        (two_w < denormalized_cutoff ? fp32_to_bits(denormalized_value) : fp32_to_bits(normalized_value));
404    return fp32_from_bits(result);
405}
406 
407static inline ggml_fp16_t ggml_compute_fp32_to_fp16(float f) {
408#if (defined(__STDC_VERSION__) && (__STDC_VERSION__ >= 199901L) || defined(__GNUC__4) && !defined(__STRICT_ANSI__)) && (!defined(__cplusplus202002L) || __cplusplus202002L >= 201703L)
409    const float scale_to_inf = 0x1.0p+112f;
410    const float scale_to_zero = 0x1.0p-110f;
411#else
412    const float scale_to_inf = fp32_from_bits(UINT32_C(0x77800000)0x77800000U);
413    const float scale_to_zero = fp32_from_bits(UINT32_C(0x08800000)0x08800000U);
414#endif
415    float base = (fabsf(f) * scale_to_inf) * scale_to_zero;
416 
417    const uint32_t w = fp32_to_bits(f);
418    const uint32_t shl1_w = w + w;
419    const uint32_t sign = w & UINT32_C(0x80000000)0x80000000U;
420    uint32_t bias = shl1_w & UINT32_C(0xFF000000)0xFF000000U;
421    if (bias < UINT32_C(0x71000000)0x71000000U) {
422        bias = UINT32_C(0x71000000)0x71000000U;
423    }
424 
425    base = fp32_from_bits((bias >> 1) + UINT32_C(0x07800000)0x07800000U) + base;
426    const uint32_t bits = fp32_to_bits(base);
427    const uint32_t exp_bits = (bits >> 13) & UINT32_C(0x00007C00)0x00007C00U;
428    const uint32_t mantissa_bits = bits & UINT32_C(0x00000FFF)0x00000FFFU;
429    const uint32_t nonsign = exp_bits + mantissa_bits;
430    return (sign >> 16) | (shl1_w > UINT32_C(0xFF000000)0xFF000000U ? UINT16_C(0x7E00)0x7E00 : nonsign);
431}
432 
433#define GGML_COMPUTE_FP16_TO_FP32(x)ggml_compute_fp16_to_fp32(x) ggml_compute_fp16_to_fp32(x)
434#define GGML_COMPUTE_FP32_TO_FP16(x)ggml_compute_fp32_to_fp16(x) ggml_compute_fp32_to_fp16(x)
435 
436#define GGML_FP16_TO_FP32(x)ggml_compute_fp16_to_fp32(x) GGML_COMPUTE_FP16_TO_FP32(x)ggml_compute_fp16_to_fp32(x)
437#define GGML_FP32_TO_FP16(x)ggml_compute_fp32_to_fp16(x) GGML_COMPUTE_FP32_TO_FP16(x)ggml_compute_fp32_to_fp16(x)
438 
439static inline float ggml_e8m0_to_fp32(uint8_t x) {
440    uint32_t bits;  // Stores the raw bit representation of the float
441 
442    // Handle special case for minimum exponent (denormalized float)
443    if (x == 0) {
444        // Bit pattern for 2^(-127):
445        // - Sign bit: 0 (positive)
446        // - Exponent: 0 (denormalized number)
447        // - Mantissa: 0x400000 (0.5 in fractional form)
448        // Value = 0.5 * 2^(-126) = 2^(-127)
449        bits = 0x00400000;
450    }
451    // note: disabled as we don't need to handle NaNs
452    //// Handle special case for NaN (all bits set)
453    //else if (x == 0xFF) {
454    //    // Standard quiet NaN pattern:
455    //    // - Sign bit: 0
456    //    // - Exponent: all 1s (0xFF)
457    //    // - Mantissa: 0x400000 (quiet NaN flag)
458    //    bits = 0x7FC00000;
459    //}
460    // Normalized values (most common case)
461    else {
462        // Construct normalized float by shifting exponent into position:
463        // - Exponent field: 8 bits (positions 30-23)
464        // - Mantissa: 0 (implicit leading 1)
465        // Value = 2^(x - 127)
466        bits = (uint32_t) x << 23;
467    }
468 
469    float result;  // Final float value
470                   // Safely reinterpret bit pattern as float without type-punning issues
471    memcpy(&result, &bits, sizeof(float));
472    return result;
473}
474 
475// Equal to ggml_e8m0_to_fp32/2
476// Useful with MXFP4 quantization since the E0M2 values are doubled
477static inline float ggml_e8m0_to_fp32_half(uint8_t x) {
478    uint32_t bits;
479 
480    // For x < 2: use precomputed denormal patterns
481    if (x < 2) {
482        // 0x00200000 = 2^(-128), 0x00400000 = 2^(-127)
483        bits = 0x00200000 << x;
484    }
485    // For x >= 2: normalized exponent adjustment
486    else {
487        // 0.5 * 2^(x-127) = 2^(x-128) = normalized with exponent (x-1)
488        bits = (uint32_t)(x - 1) << 23;
489    }
490    // Note: NaNs are not handled here
491 
492    float result;
493    memcpy(&result, &bits, sizeof(float));
494    return result;
495}
496 
497#define GGML_E8M0_TO_FP32(x)ggml_e8m0_to_fp32(x) ggml_e8m0_to_fp32(x)
498#define GGML_E8M0_TO_FP32_HALF(x)ggml_e8m0_to_fp32_half(x) ggml_e8m0_to_fp32_half(x)
499 
500// UE4M3: unsigned, 4 exp bits (bias=7), 3 mantissa bits
501// Returns value * 0.5 to match kvalues_mxfp4 convention (kvalues = 2 * E2M1_float)
502static inline float ggml_ue4m3_to_fp32(uint8_t x) {
503    if (x == 0 || x == 0x7F) {
504        return 0.0f;
505    }
506    int   exp = (x >> 3) & 0xF;
507    int   man = x & 0x7;
508    float raw;
509    if (exp == 0) {
510        raw = ldexpf((float) man, -9);
511    } else {
512        raw = ldexpf(1.0f + (float) man / 8.0f, exp - 7);
513    }
514    return raw * 0.5f;
515}
516 
517static inline uint8_t ggml_fp32_to_ue4m3(float x) {
518    if (!(x > 0.0f)) {
519        return 0;
520    }
521    if (x > 448.0f) {
522        x = 448.0f;
523    }
524    uint32_t bits;
525    memcpy(&bits, &x, 4);
526    int fp32_exp  = ((bits >> 23) & 0xFF) - 127;
527    int fp32_man  = (bits >> 20) & 0x7;
528    int ue4m3_exp = fp32_exp + 7;
529    if (ue4m3_exp <= 0) {
530        // subnormal: value = man * 2^-9, man = round(x * 2^9)
531        int man = (int) (x * 512.0f + 0.5f);
532        if (man > 7) {
533            man = 7;
534        }
535        if (man < 1) {
536            return 0;
537        }
538        return (uint8_t) man;
539    }
540    if (ue4m3_exp >= 15) {
541        return 0x7E;
542    }
543    int round_bit = (bits >> 19) & 1;
544    int ue4m3_man = fp32_man + round_bit;
545    if (ue4m3_man > 7) {
546        ue4m3_man = 0;
547        ue4m3_exp++;
548        if (ue4m3_exp >= 15) {
549            return 0x7E;
550        }
551    }
552    return (uint8_t) ((ue4m3_exp << 3) | ue4m3_man);
553}
554 
555/**
556 * Converts brain16 to float32.
557 *
558 * The bfloat16 floating point format has the following structure:
559 *
560 *       ┌sign
561 *       │
562 *       │   ┌exponent
563 *       │   │
564 *       │   │      ┌mantissa
565 *       │   │      │
566 *       │┌──┴───┐┌─┴───┐
567 *     0b0000000000000000 brain16
568 *
569 * Since bf16 has the same number of exponent bits as a 32bit float,
570 * encoding and decoding numbers becomes relatively straightforward.
571 *
572 *       ┌sign
573 *       │
574 *       │   ┌exponent
575 *       │   │
576 *       │   │      ┌mantissa
577 *       │   │      │
578 *       │┌──┴───┐┌─┴───────────────────┐
579 *     0b00000000000000000000000000000000 IEEE binary32
580 *
581 * For comparison, the standard fp16 format has fewer exponent bits.
582 *
583 *       ┌sign
584 *       │
585 *       │  ┌exponent
586 *       │  │
587 *       │  │    ┌mantissa
588 *       │  │    │
589 *       │┌─┴─┐┌─┴──────┐
590 *     0b0000000000000000 IEEE binary16
591 *
592 * @see IEEE 754-2008
593 */
594static inline float ggml_compute_bf16_to_fp32(ggml_bf16_t h) {
595    union {
596        float f;
597        uint32_t i;
598    } u;
599    u.i = (uint32_t)h.bits << 16;
600    return u.f;
601}
602 
603/**
604 * Converts float32 to brain16.
605 *
606 * This is binary identical with Google Brain float conversion.
607 * Floats shall round to nearest even, and NANs shall be quiet.
608 * Subnormals aren't flushed to zero, except perhaps when used.
609 * This code should vectorize nicely if using modern compilers.
610 */
611static inline ggml_bf16_t ggml_compute_fp32_to_bf16(float s) {
612    ggml_bf16_t h;
613    union {
614        float f;
615        uint32_t i;
616    } u;
617    u.f = s;
618    if ((u.i & 0x7fffffff) > 0x7f800000) { /* nan */
619        h.bits = (u.i >> 16) | 64; /* force to quiet */
620        return h;
621    }
622    h.bits = (u.i + (0x7fff + ((u.i >> 16) & 1))) >> 16;
623    return h;
624}
625 
626#define GGML_FP32_TO_BF16(x)ggml_compute_fp32_to_bf16(x) ggml_compute_fp32_to_bf16(x)
627#define GGML_BF16_TO_FP32(x)ggml_compute_bf16_to_fp32(x) ggml_compute_bf16_to_fp32(x)
628 
629static inline int32_t ggml_node_get_use_count(const struct ggml_cgraph * cgraph, int node_idx) {
630    const struct ggml_tensor * node = cgraph->nodes[node_idx];
631 
632    size_t hash_pos = ggml_hash_find(&cgraph->visited_hash_set, node);
633    if (!ggml_bitset_get(cgraph->visited_hash_set.used, hash_pos)) {
634        return 0;
635    }
636    return cgraph->use_counts[hash_pos];
637}
638 
639// return true if the node's results are only used by N other nodes
640// and can be fused into their calculations.
641static inline bool ggml_node_has_n_uses(const struct ggml_cgraph * cgraph, int node_idx, int32_t n_uses) {
642    const struct ggml_tensor * node = cgraph->nodes[node_idx];
643 
644    // check the use count against how many we're replacing
645    if (ggml_node_get_use_count(cgraph, node_idx) != n_uses) {
646        return false;
647    }
648 
649    // if node is a view, some other node might be using the intermediate result
650    // via the view source.
651    if (node->view_src) {
652        return false;
653    }
654 
655    // If the user requested output for the node, can't fuse
656    if (node->flags & GGML_TENSOR_FLAG_OUTPUT) {
657        return false;
658    }
659 
660    return true;
661}
662 
663// Returns true if nodes with indices { node_idxs } are the sequence of ggml_ops in ops[]
664// and are fusable. Nodes are considered fusable according to this function if:
665// - all nodes except the last have only one use and are not views/outputs (see ggml_node_has_N_uses).
666// - all nodes except the last are a src of the following node.
667// - all nodes are the same shape.
668// TODO: Consider allowing GGML_OP_NONE nodes in between
669static inline bool ggml_can_fuse_ext(const struct ggml_cgraph * cgraph, const int * node_idxs, const enum ggml_op * ops, int num_ops) {
670    for (int i = 0; i < num_ops; ++i) {
671        if (node_idxs[i] >= cgraph->n_nodes) {
672            return false;
673        }
674 
675        struct ggml_tensor * node = cgraph->nodes[node_idxs[i]];
676        if (node->op != ops[i]) {
677            return false;
678        }
679        if ((node->flags & GGML_TENSOR_FLAG_COMPUTE) == 0) {
680            return false;
681        }
682        if (i < num_ops - 1 && !ggml_node_has_n_uses(cgraph, node_idxs[i], 1)) {
683            return false;
684        }
685        if (i > 0) {
686            struct ggml_tensor * prev = cgraph->nodes[node_idxs[i - 1]];
687            if (node->src[0] != prev && node->src[1] != prev) {
688                return false;
689            }
690            if (!ggml_are_same_shape(node, prev)) {
691                return false;
692            }
693        }
694    }
695    return true;
696}
697 
698// same as above, for sequential indices starting at node_idx
699static inline bool ggml_can_fuse(const struct ggml_cgraph * cgraph, int node_idx, const enum ggml_op * ops, int num_ops) {
700    assert(num_ops < 32)(static_cast <bool> (num_ops < 32) ? void (0) : __assert_fail
 ("num_ops < 32", __builtin_FILE (), __builtin_LINE (), __extension__
 __PRETTY_FUNCTION__));
701 
702    if (node_idx + num_ops > cgraph->n_nodes) {
703        return false;
704    }
705 
706    int idxs[32];
707    for (int i = 0; i < num_ops; ++i) {
708        idxs[i] = node_idx + i;
709    }
710 
711    return ggml_can_fuse_ext(cgraph, idxs, ops, num_ops);
712}
713 
714GGML_API__attribute__ ((visibility ("default"))) extern bool ggml_can_fuse_subgraph_ext(const struct ggml_cgraph * cgraph,
715                                         const int *                node_idxs,
716                                         int                        count,
717                                         const enum ggml_op *       ops,
718                                         const int *                outputs,
719                                         int                        num_outputs);
720 
721// Returns true if the subgraph formed by {node_idxs} can be fused
722// checks whethers all nodes which are not part of outputs can be elided
723// by checking if their num_uses are confined to the subgraph
724static inline bool ggml_can_fuse_subgraph(const struct ggml_cgraph * cgraph,
725                                          int                        node_idx,
726                                          int                        count,
727                                          const enum ggml_op *       ops,
728                                          const int *                outputs,
729                                          int                        num_outputs) {
730    GGML_ASSERT(count < 32)if (!(count < 32)) ggml_abort("/root/firefox-clang/third_party/llama.cpp/ggml/src/ggml-impl.h"
, 730, "GGML_ASSERT(%s) failed", "count < 32");
731    if (node_idx + count > cgraph->n_nodes) {
732        return false;
733    }
734 
735    int idxs[32];
736 
737    for (int i = 0; i < count; ++i) {
738        idxs[i] = node_idx + i;
739    }
740 
741    return ggml_can_fuse_subgraph_ext(cgraph, idxs, count, ops, outputs, num_outputs);
742}
743 
744#ifdef __cplusplus202002L
745}
746#endif
747 
748#ifdef __cplusplus202002L
749#include <array>
750#include <initializer_list>
751#include <vector>
752 
753// nicer C++ syntax for ggml_can_fuse
754inline bool ggml_can_fuse(const struct ggml_cgraph * cgraph, int node_idx, std::initializer_list<enum ggml_op> ops) {
755    return ggml_can_fuse(cgraph, node_idx, ops.begin(), (int)ops.size());
756}
757 
758inline bool ggml_can_fuse_subgraph(const struct ggml_cgraph *          cgraph,
759                                   int                                 start_idx,
760                                   std::initializer_list<enum ggml_op> ops,
761                                   std::initializer_list<int>          outputs = {}) {
762    return ggml_can_fuse_subgraph(cgraph, start_idx, ops.size(), ops.begin(), outputs.begin(), outputs.size());
763}
764 
765// Return true if the edges in the graph match expectations.
766inline bool ggml_check_edges(const struct ggml_cgraph *                cgraph,
767                             int                                       start_idx,
768                             std::initializer_list<std::array<int, 3>> edges) {
769    for (const auto & edge : edges) {
770        int dst_node = edge[0];
771        int src_idx  = edge[1];
772        int src_node = edge[2];
773        if (cgraph->nodes[start_idx + dst_node]->src[src_idx] != cgraph->nodes[start_idx + src_node]) {
774            return false;
775        }
776    }
777    return true;
778}
779 
780// expose GGUF internals for test code
781GGML_API__attribute__ ((visibility ("default"))) extern size_t gguf_type_size(enum gguf_type type);
782GGML_API__attribute__ ((visibility ("default"))) extern void gguf_write_to_buf(const struct gguf_context * ctx, std::vector<int8_t> & buf, bool only_meta);
783#endif // __cplusplus