Public/koboldcpp
1
0
Fork 0
mirror of https://github.com/LostRuins/koboldcpp.git synced 2026-09-05 20:41:18 +02:00
Code Issues Packages Projects Releases Wiki Activity

sycl: attribute device allocations by site (GGML_SYCL_MEMTRACE) (#27631)

Browse Source 
define two new environment variables to better understand how much
memory is being allocated, and when. This has been invaluable in
inproving the --fit algorithm, and is likely to be useful when debugging
other memory-related issues.

`-lv 4` will be required to enable the following:

GGML_SYCL_MEMTRACE=1 will show per-site memory usage, updated whenever
it increases by more than 64MiB.
GGML_SYCL_MEMTRACE=2 will show every allocation and deallocation.

To change the default 64MiB threshold for reporting memory usage increases, use
GGML_SYCL_MEMTRACE_STEP.

A sample log line:
[SYCL-MEMTRACE] device memory query (dev): total 59493 MiB, free  4494, in use 54998; allocated     0 (buffers     0 + scratch     0), peak     0 MiB
This commit is contained in:
Nick Farrell
2026-09-05 12:36:02 +10:00
committed by GitHub
parent 427291b5b3
commit cd8cdf397d
7 changed files with 267 additions and 7 deletions
Download Patch File Download Diff File Expand all files Collapse all files
docs/backend/SYCL.md
+2
View File
@@ -805,6 +805,8 @@ User can use the device management in [docs/multi-gpu.md](https://github.com/ggm
| GGML_SYCL_ENABLE_VMM | 0 or 1 (default) | Enable the virtual-memory device pool. |
| GGML_SYCL_ENABLE_MKL_FA | 1 (default) or 0 | Enable oneMKL GEMM flash attention for XMX-accelerated prompt processing with quantized KV cache. Automatically activates during prefill (prompt processing) when all conditions are met: (1) flash-attn enabled (`-fa` or `--flash-attn on`), (2) KV cache quantized (`--cache-type-k q8_0 --cache-type-v q8_0` or other `*_0/*_1` types), (3) batch size ≥ 1024 (`--batch-size 1024`), (4) prompt length ≥ 1024 tokens. Set to 0 to force the TILE kernel for A/B testing. Example minimum command: `llama-cli -m model.gguf -fa -ngl 99 --cache-type-k q8_0 --cache-type-v q8_0 --batch-size 1024 -p "your prompt"` |
| GGML_SYCL_MKL_FA_DEBUG | 0 (default) or 1 | Enable per-call diagnostic logging for MKL flash attention: GEMM/softmax timings, interleaved-head detection, and buffer memory usage. |
| GGML_SYCL_MEMTRACE | 0 (default), 1, 2 | Enable record and output memory allocation diagnostics. Requires `-lv 4`. <br>0 - Disable<br>1 - Basic memory info, including current and peak allocations, as well allocations from other sources, around 50 lines per model load.<br>2 - More verbose, logging around 900 specific allocations and deallocations. |
| GGML_SYCL_MEMTRACE_STEP | 64 (default) or positive integer | With GGML_SYCL_MEMTRACE=1, the minimum growth in memory usage to trigger another log record. |
| GGML_SYCL_MKL_FA_DIAG | 0 (default) or 1 | Enable output fingerprinting for MKL flash attention. Dumps the first 64 float output values for the first 6 FA calls with n_kv ≥ 1024, labeled with kernel type (MKL/TILE/VEC) for cross-kernel comparison. |
| GGML_SYCL_ENABLE_FUSION | 0 or 1 (default) | Enable fused-kernel dispatch in graph compute. Unsupported types and layouts fall back to the standalone op kernels. See `ggml_sycl_can_fuse()`. |
| GGML_SYCL_ENABLE_ESIMD | 0 or 1 (default)| Enable ESIMD kernels when available. |
ggml/src/ggml-sycl/common.cpp
+11 -2
View File
@@ -94,7 +94,7 @@ static bool ggml_sycl_use_level_zero_device_alloc(sycl::queue &q) {
// Use Level Zero zeMemAllocDevice to avoid sycl::malloc_device triggering
// DMA-buf/TTM system RAM staging in the xe kernel driver during multi-GPU inference.
void * ggml_sycl_malloc_device(size_t size, sycl::queue &q) {
void * ggml_sycl_malloc_device(size_t size, sycl::queue &q, ggml_sycl_mem_type type) {
#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
if (ggml_sycl_use_level_zero_device_alloc(q)) {
void *ptr = nullptr;
@@ -117,16 +117,25 @@ void * ggml_sycl_malloc_device(size_t size, sycl::queue &q) {
#endif
ze_result_t r = zeMemAllocDevice(ze_ctx, &alloc_desc, size, 64, ze_dev, &ptr);
if (r == ZE_RESULT_SUCCESS && ptr) {
ggml_sycl_memtrace_add(type, ptr, size);
return ptr;
}
ggml_sycl_memtrace_fail(type, size);
return nullptr;
}
#endif
return sycl::malloc_device(size, q);
void * ptr = sycl::malloc_device(size, q);
if (ptr == nullptr) {
ggml_sycl_memtrace_fail(type, size);
return nullptr;
}
ggml_sycl_memtrace_add(type, ptr, size);
return ptr;
}
void ggml_sycl_free_device(void *ptr, sycl::queue &q) {
if (!ptr) return;
ggml_sycl_memtrace_del(ptr);
#ifdef GGML_SYCL_SUPPORT_LEVEL_ZERO_API
if (ggml_sycl_use_level_zero_device_alloc(q)) {
auto ze_ctx = sycl::get_native<sycl::backend::ext_oneapi_level_zero>(q.get_context());
ggml/src/ggml-sycl/common.hpp
+5 -1
View File
@@ -27,6 +27,7 @@
#include "type.hpp"
#include "sycl_hw.hpp"
#include "fattn-buffers.hpp"
#include "memtrace.hpp"
namespace syclexp = sycl::ext::oneapi::experimental;
@@ -69,6 +70,8 @@ extern int g_ggml_sycl_dev2dev_memcpy;
extern int g_ggml_sycl_fa_onednn;
extern int g_ggml_sycl_fa_onednn_max_kv;
extern int g_ggml_sycl_enable_mkl_fa;
extern int g_ggml_sycl_memtrace;
extern int g_ggml_sycl_memtrace_step;
#define CHECK_TRY_ERROR(expr) \
@@ -318,7 +321,8 @@ struct ggml_tensor_extra_gpu {
};
extern int g_ggml_sycl_use_level_zero_api;
void * ggml_sycl_malloc_device(size_t size, sycl::queue &q);
void * ggml_sycl_malloc_device(size_t size, sycl::queue &q,
ggml_sycl_mem_type type = GGML_SYCL_MEM_DIRECT);
void ggml_sycl_free_device(void *ptr, sycl::queue &q);
void release_extra_gpu(ggml_tensor_extra_gpu * extra, std::vector<queue_ptr> streams={});
ggml/src/ggml-sycl/fattn-buffers.cpp
+4
View File
@@ -21,6 +21,7 @@ sycl::half * ggml_sycl_fattn_kv_buffers::kv_buffer::ensure_half(size_t n_elems)
if (ptr) {
SYCL_CHECK(CHECK_TRY_ERROR(qptr->wait()));
ggml_sycl_memtrace_del(ptr);
SYCL_CHECK(CHECK_TRY_ERROR(sycl::free(ptr, *qptr)));
ptr = nullptr;
capacity = 0;
@@ -38,11 +39,13 @@ sycl::half * ggml_sycl_fattn_kv_buffers::kv_buffer::ensure_half(size_t n_elems)
if (!dev_ptr) {
GGML_LOG_ERROR("%s: can't allocate %lu Bytes of memory on device\n", __func__, cap);
ggml_sycl_memtrace_fail(GGML_SYCL_MEM_FATTN_KV, cap);
GGML_ABORT("fattn buffer alloc failed");
}
ptr = static_cast<sycl::half *>(dev_ptr);
capacity = cap;
ggml_sycl_memtrace_add(GGML_SYCL_MEM_FATTN_KV, ptr, cap);
return ptr;
}
@@ -51,6 +54,7 @@ ggml_sycl_fattn_kv_buffers::kv_buffer::~kv_buffer() {
GGML_LOG_INFO("ggml_sycl_fattn_kv_buffer[%d]: %.2f MiB\n", device, capacity / 1024.0 / 1024.0);
#endif
if (ptr) {
ggml_sycl_memtrace_del(ptr);
SYCL_CHECK(CHECK_TRY_ERROR(sycl::free(ptr, *qptr)));
}
}
ggml/src/ggml-sycl/ggml-sycl.cpp
+23 -4
View File
@@ -97,6 +97,8 @@ int g_ggml_sycl_enable_dnn = 1;
int g_ggml_sycl_fa_onednn = 1;
int g_ggml_sycl_fa_onednn_max_kv = 0;
int g_ggml_sycl_enable_mkl_fa = 1;
int g_ggml_sycl_memtrace = 0;
int g_ggml_sycl_memtrace_step = 64;
int g_ggml_sycl_enable_vmm = 1;
int g_ggml_sycl_enable_fusion = 1;
int g_ggml_sycl_enable_esimd = 1;
@@ -335,6 +337,8 @@ static void ggml_check_sycl() try {
g_ggml_sycl_fa_onednn = ggml_sycl_get_env("GGML_SYCL_FA_ONEDNN", 1);
g_ggml_sycl_fa_onednn_max_kv = ggml_sycl_get_env("GGML_SYCL_FA_ONEDNN_MAX_KV", 0);
g_ggml_sycl_enable_mkl_fa = ggml_sycl_get_env("GGML_SYCL_ENABLE_MKL_FA", 1);
g_ggml_sycl_memtrace = ggml_sycl_get_env("GGML_SYCL_MEMTRACE", 0);
g_ggml_sycl_memtrace_step = ggml_sycl_get_env("GGML_SYCL_MEMTRACE_STEP", 64);
g_ggml_sycl_enable_vmm = ggml_sycl_get_env("GGML_SYCL_ENABLE_VMM", 1);
g_ggml_sycl_enable_fusion = ggml_sycl_get_env("GGML_SYCL_ENABLE_FUSION", 1);
g_ggml_sycl_enable_esimd = ggml_sycl_get_env("GGML_SYCL_ENABLE_ESIMD", 1);
@@ -421,6 +425,8 @@ static void ggml_check_sycl() try {
#endif
GGML_LOG_INFO(" GGML_SYCL_FA_ONEDNN_MAX_KV: %d\n", g_ggml_sycl_fa_onednn_max_kv);
GGML_LOG_INFO(" GGML_SYCL_ENABLE_MKL_FA: %d\n", g_ggml_sycl_enable_mkl_fa);
GGML_LOG_INFO(" GGML_SYCL_MEMTRACE: %d\n", g_ggml_sycl_memtrace);
GGML_LOG_INFO(" GGML_SYCL_MEMTRACE_STEP: %d\n", g_ggml_sycl_memtrace_step);
#ifdef SYCL_FLASH_ATTN
GGML_LOG_INFO(" GGML_SYCL_ENABLE_FLASH_ATTN: %d\n", g_ggml_sycl_enable_flash_attention);
#else
@@ -964,7 +970,7 @@ ggml_backend_sycl_buffer_type_alloc_buffer(ggml_backend_buffer_type_t buft,
return nullptr;
}
} else {
SYCL_CHECK(CHECK_TRY_ERROR(dev_ptr = (void *)ggml_sycl_malloc_device(size, *stream)));
SYCL_CHECK(CHECK_TRY_ERROR(dev_ptr = (void *)ggml_sycl_malloc_device(size, *stream, GGML_SYCL_MEM_BUFFER)));
if (!dev_ptr) {
GGML_LOG_ERROR("%s: can't allocate %zu Bytes of memory on device\n", __func__, size);
return nullptr;
@@ -1217,7 +1223,7 @@ ggml_backend_sycl_split_buffer_init_tensor(ggml_backend_buffer_t buffer,
ggml_sycl_set_device(i);
const queue_ptr stream = ctx->streams[i];
char * buf;
SYCL_CHECK(CHECK_TRY_ERROR(buf = (char *)ggml_sycl_malloc_device(size, *stream)));
SYCL_CHECK(CHECK_TRY_ERROR(buf = (char *)ggml_sycl_malloc_device(size, *stream, GGML_SYCL_MEM_BUFFER)));
if (!buf) {
char err_buf[1024];
snprintf(err_buf, 1023, "%s: can't allocate %zu Bytes of memory on device\n", __func__, size);
@@ -1697,7 +1703,7 @@ struct ggml_sycl_pool_leg : public ggml_sycl_pool {
void * ptr;
size_t look_ahead_size = (size_t) (1.05 * size);
SYCL_CHECK(CHECK_TRY_ERROR(ptr = (void *)ggml_sycl_malloc_device(look_ahead_size, *qptr)));
SYCL_CHECK(CHECK_TRY_ERROR(ptr = (void *)ggml_sycl_malloc_device(look_ahead_size, *qptr, GGML_SYCL_MEM_POOL_LEG)));
if (!ptr) {
GGML_LOG_ERROR("%s: can't allocate %zu Bytes of memory on device/GPU\n", __func__, look_ahead_size);
return nullptr;
@@ -1786,6 +1792,13 @@ struct ggml_sycl_pool_vmm : public ggml_sycl_pool {
GGML_ASSERT(pool_size + reserve_size <= SYCL_POOL_VMM_MAX_SIZE);
if (ggml_sycl_memtrace_enabled()) {
GGML_LOG_INFO(GGML_SYCL_MEMTRACE_TAG " pool_vmm[%d] committing %5zu MiB (pool %5zu -> %5zu MiB)\n",
device, reserve_size / (1024 * 1024), pool_size / (1024 * 1024),
(pool_size + reserve_size) / (1024 * 1024));
ggml_sycl_memtrace_report("before pool_vmm commit");
}
// allocate more physical memory
std::optional<sycl::ext::oneapi::experimental::physical_mem> phys;
SYCL_CHECK(CHECK_TRY_ERROR(phys.emplace(dev, ctx, reserve_size)));
@@ -1811,6 +1824,7 @@ struct ggml_sycl_pool_vmm : public ggml_sycl_pool {
// add to the pool
pool_size += reserve_size;
ggml_sycl_memtrace_add(GGML_SYCL_MEM_POOL_VMM, map_ptr, reserve_size);
#ifdef DEBUG_SYCL_MALLOC
GGML_LOG_INFO("sycl pool[%d]: size increased to %llu MB (reserved %llu MB)\n",
@@ -4039,7 +4053,9 @@ static inline void * sycl_ext_malloc_device(dpct::queue_ptr stream, size_t size)
bool use_async = g_ggml_sycl_use_async_mem_op;
#if defined(GGML_SYCL_GRAPH) && SYCL_EXT_ONEAPI_ASYNC_MEMORY_ALLOC
if (use_async) {
return syclex::async_malloc(*stream, sycl::usm::alloc::device, size);
void * ptr = syclex::async_malloc(*stream, sycl::usm::alloc::device, size);
ggml_sycl_memtrace_add(GGML_SYCL_MEM_ASYNC, ptr, size);
return ptr;
}
#else
// If async allocation extension is not available, use_async should always be false.
@@ -4052,6 +4068,7 @@ static inline void sycl_ext_free(dpct::queue_ptr stream, void * ptr) {
bool use_async = g_ggml_sycl_use_async_mem_op;
#if defined(GGML_SYCL_GRAPH) && SYCL_EXT_ONEAPI_ASYNC_MEMORY_ALLOC
if (use_async) {
ggml_sycl_memtrace_del(ptr);
syclex::async_free(*stream, ptr);
return;
}
@@ -5643,6 +5660,7 @@ void ggml_backend_sycl_get_device_memory(int device, size_t * free, size_t * tot
if (!res) {
GGML_ABORT("[%s] failed to get device memory size", __func__);
}
ggml_sycl_memtrace_report_device("device memory query", device, *free, *total);
} catch (const sycl::exception & exc) {
std::cerr << exc.what() << "Exception caught at file:" << __FILE__ << ", line:" << __LINE__ << std::endl;
std::exit(1);
@@ -6082,6 +6100,7 @@ static void ggml_backend_sycl_device_get_memory(ggml_backend_dev_t dev, size_t *
if (!res) {
GGML_ABORT("[%s] failed to get device memory size", __func__);
}
ggml_sycl_memtrace_report_device("device memory query (dev)", ctx->device, *free, *total);
}
static enum ggml_backend_dev_type ggml_backend_sycl_device_get_type(ggml_backend_dev_t dev) {
ggml/src/ggml-sycl/memtrace.cpp
+194
View File
@@ -0,0 +1,194 @@
#include "memtrace.hpp"
#include "common.hpp"
#include "ggml-impl.h"
#include <cstdio>
#include <mutex>
#include <unordered_map>
constexpr size_t MIB = 1024 * 1024;
static const char * mem_type_name(ggml_sycl_mem_type type) {
switch (type) {
case GGML_SYCL_MEM_BUFFER: return "buffer";
case GGML_SYCL_MEM_POOL_LEG: return "pool_leg";
case GGML_SYCL_MEM_POOL_VMM: return "pool_vmm";
case GGML_SYCL_MEM_ASYNC: return "async";
case GGML_SYCL_MEM_FATTN_KV: return "fattn_kv";
case GGML_SYCL_MEM_DIRECT: return "direct";
default: GGML_ABORT("[%s] The type value %d is not supported\n", __func__, (int) type);
}
}
struct mem_tracker {
std::mutex mutex;
std::unordered_map<const void *, std::pair<ggml_sycl_mem_type, size_t>> live_by_ptr;
size_t live[GGML_SYCL_MEM_TYPE_COUNT] = {};
size_t peak[GGML_SYCL_MEM_TYPE_COUNT] = {};
size_t total_live = 0;
size_t total_peak = 0;
size_t last_logged_peak = 0;
};
static mem_tracker & get_tracker() {
static mem_tracker t;
return t;
}
static size_t step_bytes() {
const int mib = g_ggml_sycl_memtrace_step > 0 ? g_ggml_sycl_memtrace_step : 64;
return (size_t) mib * MIB;
}
static void report_sites_locked() {
mem_tracker & t = get_tracker();
for (int i = 0; i < GGML_SYCL_MEM_TYPE_COUNT; i++) {
if (t.peak[i] == 0) {
continue;
}
GGML_LOG_INFO(GGML_SYCL_MEMTRACE_TAG " %-9s allocated %5zu MiB, peak %5zu MiB\n",
mem_type_name((ggml_sycl_mem_type) i), t.live[i] / MIB, t.peak[i] / MIB);
}
}
static void report_locked(const char * tag) {
mem_tracker & t = get_tracker();
const size_t allocated = t.total_live / MIB;
const size_t buffers = t.live[GGML_SYCL_MEM_BUFFER] / MIB;
GGML_LOG_INFO(GGML_SYCL_MEMTRACE_TAG " %s: allocated %5zu MiB (buffers %5zu + scratch %5zu),"
" peak %5zu MiB\n",
tag, allocated, buffers, allocated - buffers, t.total_peak / MIB);
report_sites_locked();
}
static void log_event_locked(const char * op, ggml_sycl_mem_type type, const void * ptr, size_t bytes) {
GGML_LOG_INFO(GGML_SYCL_MEMTRACE_TAG " allocated %5zu MiB %-5s %-9s %9.3f MiB ptr=%p\n",
get_tracker().total_live / MIB, op, mem_type_name(type),
(double) bytes / MIB, ptr);
}
bool ggml_sycl_memtrace_enabled() {
return g_ggml_sycl_memtrace > 0;
}
void ggml_sycl_memtrace_add(ggml_sycl_mem_type type, const void * ptr, size_t bytes) {
if (!ggml_sycl_memtrace_enabled()) {
return;
}
GGML_ASSERT(ptr != nullptr);
GGML_ASSERT(bytes != 0);
mem_tracker & t = get_tracker();
std::lock_guard<std::mutex> lock(t.mutex);
auto it = t.live_by_ptr.find(ptr);
if (it != t.live_by_ptr.end()) {
t.live[it->second.first] -= it->second.second;
t.total_live -= it->second.second;
}
t.live_by_ptr[ptr] = { type, bytes };
t.live[type] += bytes;
t.total_live += bytes;
if (t.live[type] > t.peak[type]) {
t.peak[type] = t.live[type];
}
if (t.total_live > t.total_peak) {
t.total_peak = t.total_live;
}
if (g_ggml_sycl_memtrace >= 2) {
log_event_locked("alloc", type, ptr, bytes);
}
static const size_t step = step_bytes();
if (t.total_peak >= t.last_logged_peak + step) {
t.last_logged_peak = t.total_peak;
char tag[96];
std::snprintf(tag, sizeof(tag), "peak grew (+%zu MiB from %s)", bytes / MIB,
mem_type_name(type));
report_locked(tag);
}
}
void ggml_sycl_memtrace_del(const void * ptr) {
if (!ggml_sycl_memtrace_enabled() || ptr == nullptr) {
return;
}
mem_tracker & t = get_tracker();
std::lock_guard<std::mutex> lock(t.mutex);
auto it = t.live_by_ptr.find(ptr);
if (it == t.live_by_ptr.end()) {
return;
}
const ggml_sycl_mem_type type = it->second.first;
const size_t bytes = it->second.second;
t.live[type] -= bytes;
t.total_live -= bytes;
t.live_by_ptr.erase(it);
if (g_ggml_sycl_memtrace >= 2) {
log_event_locked("free", type, ptr, bytes);
}
}
void ggml_sycl_memtrace_fail(ggml_sycl_mem_type type, size_t bytes) {
GGML_LOG_ERROR(GGML_SYCL_MEMTRACE_TAG " alloc FAILED: %9.3f MiB %s\n",
(double) bytes / MIB, mem_type_name(type));
if (!ggml_sycl_memtrace_enabled()) {
return;
}
mem_tracker & t = get_tracker();
std::lock_guard<std::mutex> lock(t.mutex);
report_locked("at allocation failure");
}
void ggml_sycl_memtrace_report(const char * tag) {
if (!ggml_sycl_memtrace_enabled()) {
return;
}
mem_tracker & t = get_tracker();
std::lock_guard<std::mutex> lock(t.mutex);
report_locked(tag);
}
static bool device_memory_is_dedicated(int device) {
if (device < 0 || device >= ggml_sycl_info().device_count) {
return false;
}
const sycl_device_info & info = ggml_sycl_info().devices[device];
return info.l0_device_type_valid && info.l0_discrete_gpu;
}
void ggml_sycl_memtrace_report_device(const char * tag, int device, size_t dev_free, size_t dev_total) {
if (!ggml_sycl_memtrace_enabled()) {
return;
}
mem_tracker & t = get_tracker();
std::lock_guard<std::mutex> lock(t.mutex);
const size_t in_use = dev_total > dev_free ? dev_total - dev_free : 0;
const size_t total = dev_total / MIB;
const size_t freed = dev_free / MIB;
const size_t allocated = t.total_live / MIB;
const size_t buffers = t.live[GGML_SYCL_MEM_BUFFER] / MIB;
const size_t peak = t.total_peak / MIB;
if (in_use >= t.total_live && device_memory_is_dedicated(device) && total >= freed + allocated) {
GGML_LOG_INFO(GGML_SYCL_MEMTRACE_TAG " %s: total %5zu MiB = free %5zu + allocated %5zu"
" (buffers %5zu + scratch %5zu) + other %5zu, peak %5zu MiB\n",
tag, total, freed, allocated, buffers, allocated - buffers,
total - freed - allocated, peak);
} else {
GGML_LOG_INFO(GGML_SYCL_MEMTRACE_TAG " %s: total %5zu MiB, free %5zu, in use %5zu;"
" allocated %5zu (buffers %5zu + scratch %5zu), peak %5zu MiB\n",
tag, total, freed, in_use / MIB, allocated, buffers,
allocated - buffers, peak);
}
report_sites_locked();
}
ggml/src/ggml-sycl/memtrace.hpp
+28
View File
@@ -0,0 +1,28 @@
#ifndef GGML_SYCL_MEMTRACE_HPP
#define GGML_SYCL_MEMTRACE_HPP
#include <cstddef>
#define GGML_SYCL_MEMTRACE_TAG "[SYCL-MEMTRACE]"
enum ggml_sycl_mem_type {
GGML_SYCL_MEM_BUFFER = 0,
GGML_SYCL_MEM_POOL_LEG,
GGML_SYCL_MEM_POOL_VMM,
GGML_SYCL_MEM_ASYNC,
GGML_SYCL_MEM_FATTN_KV,
GGML_SYCL_MEM_DIRECT,
GGML_SYCL_MEM_TYPE_COUNT,
};
bool ggml_sycl_memtrace_enabled();
void ggml_sycl_memtrace_add(ggml_sycl_mem_type type, const void * ptr, size_t bytes);
void ggml_sycl_memtrace_del(const void * ptr);
void ggml_sycl_memtrace_report(const char * tag);
void ggml_sycl_memtrace_report_device(const char * tag, int device, size_t dev_free, size_t dev_total);
void ggml_sycl_memtrace_fail(ggml_sycl_mem_type type, size_t bytes);
#endif // GGML_SYCL_MEMTRACE_HPP