diff --git a/docs/backend/SYCL.md b/docs/backend/SYCL.md
index c4dcb02ee..4a640e442 100644
--- a/docs/backend/SYCL.md
+++ b/docs/backend/SYCL.md
@@ -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`.
0 - Disable
1 - Basic memory info, including current and peak allocations, as well allocations from other sources, around 50 lines per model load.
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. |
diff --git a/ggml/src/ggml-sycl/common.cpp b/ggml/src/ggml-sycl/common.cpp
index e1b6db13e..894006949 100644
--- a/ggml/src/ggml-sycl/common.cpp
+++ b/ggml/src/ggml-sycl/common.cpp
@@ -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(q.get_context());
diff --git a/ggml/src/ggml-sycl/common.hpp b/ggml/src/ggml-sycl/common.hpp
index 9f2a27b18..355dd442b 100644
--- a/ggml/src/ggml-sycl/common.hpp
+++ b/ggml/src/ggml-sycl/common.hpp
@@ -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 streams={});
diff --git a/ggml/src/ggml-sycl/fattn-buffers.cpp b/ggml/src/ggml-sycl/fattn-buffers.cpp
index 46cf6d551..78a52d2ab 100644
--- a/ggml/src/ggml-sycl/fattn-buffers.cpp
+++ b/ggml/src/ggml-sycl/fattn-buffers.cpp
@@ -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(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)));
}
}
diff --git a/ggml/src/ggml-sycl/ggml-sycl.cpp b/ggml/src/ggml-sycl/ggml-sycl.cpp
index 27804e073..bfe6f1016 100644
--- a/ggml/src/ggml-sycl/ggml-sycl.cpp
+++ b/ggml/src/ggml-sycl/ggml-sycl.cpp
@@ -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 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) {
diff --git a/ggml/src/ggml-sycl/memtrace.cpp b/ggml/src/ggml-sycl/memtrace.cpp
new file mode 100644
index 000000000..9c4f88539
--- /dev/null
+++ b/ggml/src/ggml-sycl/memtrace.cpp
@@ -0,0 +1,194 @@
+#include "memtrace.hpp"
+
+#include "common.hpp"
+#include "ggml-impl.h"
+
+#include
+#include
+#include
+
+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> 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 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 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 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 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 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();
+}
diff --git a/ggml/src/ggml-sycl/memtrace.hpp b/ggml/src/ggml-sycl/memtrace.hpp
new file mode 100644
index 000000000..426d90963
--- /dev/null
+++ b/ggml/src/ggml-sycl/memtrace.hpp
@@ -0,0 +1,28 @@
+#ifndef GGML_SYCL_MEMTRACE_HPP
+#define GGML_SYCL_MEMTRACE_HPP
+
+#include
+
+#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