mirror of
https://github.com/ggml-org/llama.cpp.git
synced 2026-09-07 16:37:57 +02:00
Compare commits
9
Commits
| Author | SHA1 | Date | |
|---|---|---|---|
 Xuan Son Nguyen
|
4c31a6aac2 | ||
|
Xuan Son Nguyen
|
f288e69886 | ||
|
Niklas WenzelandGitHub
|
6a1a922d26 | ||
|
Jingxin (Philip) LiandGitHub
|
4d9176092d | ||
|
Nick FarrellandGitHub
|
cd8cdf397d | ||
|
IsaacandGitHub
|
427291b5b3 | ||
|
nachobhandGitHub
|
85d5703a3b | ||
|
Hongqiang WangandGitHub
|
1548a240e3 | ||
|
|
4acf4a4cb8 |
@@ -3901,6 +3901,14 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
|
||||
common_log_set_file(common_log_main(), value.c_str());
|
||||
}
|
||||
).set_env("LLAMA_ARG_LOG_FILE"));
|
||||
add_opt(common_arg(
|
||||
{"--log-jsonl"},
|
||||
{"--no-log-jsonl"},
|
||||
"Log as JSONL (one JSON object per line) to stdout, this also disables colored logging (default: disabled)",
|
||||
[](common_params &, bool value) {
|
||||
common_log_set_jsonl(common_log_main(), value);
|
||||
}
|
||||
).set_env("LLAMA_ARG_LOG_JSONL"));
|
||||
add_opt(common_arg(
|
||||
{"--log-prompts-dir"}, "PATH",
|
||||
"Log prompts to directory (auto-created if not present; only used for debugging, default: disabled)",
|
||||
|
||||
+39
-1
@@ -1,5 +1,6 @@
|
||||
#include "common.h"
|
||||
#include "log.h"
|
||||
#include "json.h"
|
||||
|
||||
#include <chrono>
|
||||
#include <condition_variable>
|
||||
@@ -66,6 +67,17 @@ static const char* g_col[] = {
|
||||
"",
|
||||
};
|
||||
|
||||
static const char * level_str(enum ggml_log_level level) {
|
||||
switch (level) {
|
||||
case GGML_LOG_LEVEL_DEBUG: return "debug";
|
||||
case GGML_LOG_LEVEL_INFO: return "info";
|
||||
case GGML_LOG_LEVEL_WARN: return "warn";
|
||||
case GGML_LOG_LEVEL_ERROR: return "error";
|
||||
case GGML_LOG_LEVEL_CONT: return "cont";
|
||||
default: return "none";
|
||||
}
|
||||
}
|
||||
|
||||
struct common_log_entry {
|
||||
enum ggml_log_level level {GGML_LOG_LEVEL_INFO};
|
||||
|
||||
@@ -74,6 +86,7 @@ struct common_log_entry {
|
||||
int64_t timestamp { 0 };
|
||||
bool is_end { false }; // signals the worker thread to stop
|
||||
bool prefix { false };
|
||||
bool jsonl { false };
|
||||
|
||||
common_log_entry(size_t size = 256) : msg(size) { }
|
||||
|
||||
@@ -88,11 +101,23 @@ struct common_log_entry {
|
||||
|
||||
fcur = stdout;
|
||||
|
||||
if (level != GGML_LOG_LEVEL_NONE) {
|
||||
if (level != GGML_LOG_LEVEL_NONE && !jsonl) {
|
||||
fcur = stderr;
|
||||
}
|
||||
}
|
||||
|
||||
if (jsonl) {
|
||||
common_json obj = {
|
||||
{"type", "log"},
|
||||
{"time", timestamp},
|
||||
{"level", level_str(level)},
|
||||
{"msg", msg.data()},
|
||||
};
|
||||
fprintf(fcur, "%s\n", obj.dump_safe().c_str());
|
||||
fflush(fcur);
|
||||
return;
|
||||
}
|
||||
|
||||
if (level != GGML_LOG_LEVEL_NONE && level != GGML_LOG_LEVEL_CONT && prefix) {
|
||||
if (timestamp) {
|
||||
// [M.s.ms.us]
|
||||
@@ -131,6 +156,7 @@ struct common_log {
|
||||
file = nullptr;
|
||||
prefix = false;
|
||||
timestamps = false;
|
||||
jsonl = false;
|
||||
running = false;
|
||||
t_start = t_us();
|
||||
|
||||
@@ -158,6 +184,7 @@ private:
|
||||
|
||||
bool prefix;
|
||||
bool timestamps;
|
||||
bool jsonl;
|
||||
bool running;
|
||||
|
||||
int64_t t_start;
|
||||
@@ -246,6 +273,7 @@ public:
|
||||
entry.is_end = false;
|
||||
entry.level = level;
|
||||
entry.prefix = prefix;
|
||||
entry.jsonl = jsonl;
|
||||
entry.timestamp = 0;
|
||||
if (timestamps) {
|
||||
entry.timestamp = t_us() - t_start;
|
||||
@@ -360,6 +388,12 @@ public:
|
||||
|
||||
this->timestamps = timestamps;
|
||||
}
|
||||
|
||||
void set_jsonl(bool jsonl) {
|
||||
std::lock_guard<std::mutex> lock(mtx);
|
||||
|
||||
this->jsonl = jsonl;
|
||||
}
|
||||
};
|
||||
|
||||
//
|
||||
@@ -433,6 +467,10 @@ void common_log_set_timestamps(struct common_log * log, bool timestamps) {
|
||||
log->set_timestamps(timestamps);
|
||||
}
|
||||
|
||||
void common_log_set_jsonl(struct common_log * log, bool jsonl) {
|
||||
log->set_jsonl(jsonl);
|
||||
}
|
||||
|
||||
void common_log_flush(struct common_log * log) {
|
||||
log->pause();
|
||||
log->resume();
|
||||
|
||||
@@ -91,6 +91,7 @@ void common_log_set_file (struct common_log * log, const char * file); // n
|
||||
void common_log_set_colors (struct common_log * log, log_colors colors); // not thread-safe
|
||||
void common_log_set_prefix (struct common_log * log, bool prefix); // whether to output prefix to each log
|
||||
void common_log_set_timestamps(struct common_log * log, bool timestamps); // whether to output timestamps in the prefix
|
||||
void common_log_set_jsonl (struct common_log * log, bool jsonl); // print each log as a JSON object on one line, not thread-safe
|
||||
void common_log_flush (struct common_log * log); // flush all pending log messages
|
||||
|
||||
// helper macros for logging
|
||||
|
||||
@@ -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. |
|
||||
|
||||
@@ -111,6 +111,7 @@ ggml_metal_t ggml_metal_init(ggml_metal_device_t dev) {
|
||||
id<MTLCommandQueue> queue = ggml_metal_device_get_queue(dev);
|
||||
if (queue == nil) {
|
||||
GGML_LOG_ERROR("%s: error: failed to create command queue\n", __func__);
|
||||
free(res);
|
||||
return NULL;
|
||||
}
|
||||
|
||||
|
||||
@@ -1525,6 +1525,178 @@ constexpr fa_vec_entry_t fa_vec_tuned_table[] = {
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 576, 512, 2, 1 }, { 4, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 576, 512, 2, 2 }, { 4, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_F16, 576, 512, 3, 1 }, { 4, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 32, 32, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 32, 32, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 32, 32, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 32, 32, 2, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 32, 32, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 64, 64, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 64, 64, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 64, 64, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 64, 64, 1, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 64, 64, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 64, 64, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 64, 64, 3, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 96, 96, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 96, 96, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 96, 96, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 128, 128, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 128, 128, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 128, 128, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 128, 128, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 128, 128, 3, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 192, 192, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 192, 192, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 192, 192, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 192, 192, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 192, 192, 3, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 192, 192, 3, 2 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 192, 192, 3, 3 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 192, 128, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 192, 128, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 192, 128, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 192, 128, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 256, 256, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 256, 256, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 320, 256, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 320, 256, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 320, 256, 3, 4 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 512, 512, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 512, 512, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 576, 512, 2, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 576, 512, 3, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 576, 512, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_0, 576, 512, 1, 4 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 32, 32, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 32, 32, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 32, 32, 1, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 32, 32, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 32, 32, 2, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 32, 32, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 64, 64, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 64, 64, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 64, 64, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 64, 64, 1, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 64, 64, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 64, 64, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 96, 96, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 96, 96, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 96, 96, 1, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 96, 96, 2, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 128, 128, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 128, 128, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 128, 128, 1, 2 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 128, 128, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 128, 128, 3, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 128, 128, 3, 2 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 128, 128, 3, 3 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 192, 192, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 192, 192, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 192, 192, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 192, 192, 2, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 192, 192, 3, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 192, 128, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 192, 128, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 256, 256, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 256, 256, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 320, 256, 2, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 320, 256, 3, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 320, 256, 1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 320, 256, 2, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 320, 256, 3, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 320, 256, 3, 3 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 512, 512, 2, 0 }, { 4, 1 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 512, 512, 2, 4 }, { 4, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 512, 512, 3, 1 }, { 4, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 512, 512, 3, 2 }, { 4, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 576, 512, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 576, 512, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q4_1, 576, 512, 1, 4 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 32, 32, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 32, 32, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 32, 32, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 32, 32, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 64, 64, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 64, 64, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 64, 64, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 64, 64, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 64, 64, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 96, 96, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 96, 96, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 96, 96, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 96, 96, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 128, 128, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 128, 128, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 128, 128, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 128, 128, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 128, 128, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 192, 192, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 192, 192, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 192, 128, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 192, 128, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 192, 128, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 192, 128, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 192, 128, 2, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 192, 128, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 256, 256, -1, 0 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 256, 256, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 256, 256, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 256, 256, 2, 2 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 256, 256, 3, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 320, 256, 1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 320, 256, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 320, 256, 2, 2 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 320, 256, 3, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 512, 512, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 512, 512, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 512, 512, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 576, 512, 2, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 576, 512, 3, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 576, 512, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 576, 512, 1, 1 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_0, 576, 512, 1, 4 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 32, 32, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 32, 32, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 32, 32, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 32, 32, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 64, 64, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 64, 64, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 64, 64, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 64, 64, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 64, 64, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 96, 96, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 96, 96, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 96, 96, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 96, 96, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 128, 128, 3, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 128, 128, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 128, 128, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 128, 128, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 128, 128, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 192, 192, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 192, 192, -1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 192, 128, 1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 192, 128, 2, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 192, 128, -1, 1 }, { 4, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 192, 128, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 192, 128, 1, 4 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 192, 128, 2, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 192, 128, 3, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 256, 256, -1, 0 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 256, 256, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 256, 256, 3, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 320, 256, 3, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 320, 256, -1, 1 }, { 2, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 320, 256, 1, 2 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 320, 256, 2, 2 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 320, 256, 3, 2 }, { 1, 2 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 512, 512, -1, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 512, 512, -1, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 576, 512, 2, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 576, 512, 3, 0 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 576, 512, 2, 4 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 576, 512, 3, 1 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q5_1, 576, 512, 3, 3 }, { 1, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 32, 32, -1, 1 }, { 4, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 32, 32, 1, 1 }, { 2, 4 } },
|
||||
{ { GGML_METAL_DEVICE_M3, GGML_TYPE_Q8_0, 32, 32, 2, 1 }, { 2, 4 } },
|
||||
|
||||
@@ -222,6 +222,7 @@ set(GGML_OPENCL_KERNELS
|
||||
exp
|
||||
expm1
|
||||
abs
|
||||
unary_ext
|
||||
softplus
|
||||
pad
|
||||
repeat
|
||||
@@ -238,7 +239,7 @@ set(GGML_OPENCL_KERNELS
|
||||
)
|
||||
|
||||
if (GGML_OPENCL_USE_ADRENO_KERNELS)
|
||||
list(APPEND GGML_OPENCL_KERNELS gemm_xmem_f16_f32_os8)
|
||||
list(APPEND GGML_OPENCL_KERNELS gemm_xmem_f16_f32_os8 sdpa_xmem_f32_f16_os8)
|
||||
endif ()
|
||||
|
||||
foreach (K ${GGML_OPENCL_KERNELS})
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@@ -1,56 +1,66 @@
|
||||
kernel void kernel_concat_f32(
|
||||
global const char * src0,
|
||||
ulong offset0,
|
||||
global const char * src1,
|
||||
ulong offset1,
|
||||
global char * dst,
|
||||
ulong offsetd,
|
||||
int ne00,
|
||||
int ne01,
|
||||
int ne02,
|
||||
int ne03,
|
||||
ulong nb00,
|
||||
ulong nb01,
|
||||
ulong nb02,
|
||||
ulong nb03,
|
||||
ulong nb10,
|
||||
ulong nb11,
|
||||
ulong nb12,
|
||||
ulong nb13,
|
||||
int ne0,
|
||||
ulong nb0,
|
||||
ulong nb1,
|
||||
ulong nb2,
|
||||
ulong nb3,
|
||||
int dim
|
||||
) {
|
||||
src0 = src0 + offset0;
|
||||
src1 = src1 + offset1;
|
||||
dst = dst + offsetd;
|
||||
// concat is a pure copy, so the kernels are keyed by element byte size
|
||||
// (1/2/4/8) rather than logical type, matching the CUDA backend.
|
||||
|
||||
const int i3 = get_group_id(2);
|
||||
const int i2 = get_group_id(1);
|
||||
const int i1 = get_group_id(0);
|
||||
|
||||
int o[4] = {0, 0, 0, 0};
|
||||
o[dim] = dim == 0 ? ne00 : (dim == 1 ? ne01 : (dim == 2 ? ne02 : ne03));
|
||||
|
||||
global const float * x;
|
||||
|
||||
for (int i0 = get_local_id(0); i0 < ne0; i0 += get_local_size(0)) {
|
||||
if (i0 < ne00 && i1 < ne01 && i2 < ne02 && i3 < ne03) {
|
||||
x = (global const float *)(src0 + (i3 )*nb03 + (i2 )*nb02 + (i1 )*nb01 + (i0 )*nb00);
|
||||
} else {
|
||||
x = (global const float *)(src1 + (i3 - o[3])*nb13 + (i2 - o[2])*nb12 + (i1 - o[1])*nb11 + (i0 - o[0])*nb10);
|
||||
}
|
||||
|
||||
global float * y = (global float *)(dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
|
||||
|
||||
*y = *x;
|
||||
}
|
||||
#define KERNEL_CONCAT(SUFFIX, T) \
|
||||
kernel void kernel_concat_##SUFFIX( \
|
||||
global const char * src0, \
|
||||
ulong offset0, \
|
||||
global const char * src1, \
|
||||
ulong offset1, \
|
||||
global char * dst, \
|
||||
ulong offsetd, \
|
||||
int ne00, \
|
||||
int ne01, \
|
||||
int ne02, \
|
||||
int ne03, \
|
||||
ulong nb00, \
|
||||
ulong nb01, \
|
||||
ulong nb02, \
|
||||
ulong nb03, \
|
||||
ulong nb10, \
|
||||
ulong nb11, \
|
||||
ulong nb12, \
|
||||
ulong nb13, \
|
||||
int ne0, \
|
||||
ulong nb0, \
|
||||
ulong nb1, \
|
||||
ulong nb2, \
|
||||
ulong nb3, \
|
||||
int dim \
|
||||
) { \
|
||||
src0 = src0 + offset0; \
|
||||
src1 = src1 + offset1; \
|
||||
dst = dst + offsetd; \
|
||||
\
|
||||
const int i3 = get_group_id(2); \
|
||||
const int i2 = get_group_id(1); \
|
||||
const int i1 = get_group_id(0); \
|
||||
\
|
||||
int o[4] = {0, 0, 0, 0}; \
|
||||
o[dim] = dim == 0 ? ne00 : (dim == 1 ? ne01 : (dim == 2 ? ne02 : ne03)); \
|
||||
\
|
||||
global const T * x; \
|
||||
\
|
||||
for (int i0 = get_local_id(0); i0 < ne0; i0 += get_local_size(0)) { \
|
||||
if (i0 < ne00 && i1 < ne01 && i2 < ne02 && i3 < ne03) { \
|
||||
x = (global const T *)(src0 + (i3 )*nb03 + (i2 )*nb02 + (i1 )*nb01 + (i0 )*nb00); \
|
||||
} else { \
|
||||
x = (global const T *)(src1 + (i3 - o[3])*nb13 + (i2 - o[2])*nb12 + (i1 - o[1])*nb11 + (i0 - o[0])*nb10); \
|
||||
} \
|
||||
\
|
||||
global T * y = (global T *)(dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0); \
|
||||
\
|
||||
*y = *x; \
|
||||
} \
|
||||
}
|
||||
|
||||
kernel void kernel_concat_f32_pack(
|
||||
KERNEL_CONCAT(b1, char)
|
||||
KERNEL_CONCAT(b2, short)
|
||||
KERNEL_CONCAT(b4, int)
|
||||
KERNEL_CONCAT(b8, long)
|
||||
|
||||
// packed variant for the common dim==0, small-ne0 case (4-byte elements only).
|
||||
kernel void kernel_concat_b4_pack(
|
||||
global const char * src0,
|
||||
ulong offset0,
|
||||
global const char * src1,
|
||||
@@ -104,14 +114,14 @@ kernel void kernel_concat_f32_pack(
|
||||
o[dim] = dim == 0 ? ne00 : (dim == 1 ? ne01 : (dim == 2 ? ne02 : ne03));
|
||||
|
||||
for (int i0 = lane; i0 < ne0; i0 += tpr) {
|
||||
global const float * x;
|
||||
global const int * x;
|
||||
if (i0 < ne00 && i1 < ne01 && i2 < ne02 && i3 < ne03) {
|
||||
x = (global const float *)(src0 + (i3 )*nb03 + (i2 )*nb02 + (i1 )*nb01 + (i0 )*nb00);
|
||||
x = (global const int *)(src0 + (i3 )*nb03 + (i2 )*nb02 + (i1 )*nb01 + (i0 )*nb00);
|
||||
} else {
|
||||
x = (global const float *)(src1 + (i3 - o[3])*nb13 + (i2 - o[2])*nb12 + (i1 - o[1])*nb11 + (i0 - o[0])*nb10);
|
||||
x = (global const int *)(src1 + (i3 - o[3])*nb13 + (i2 - o[2])*nb12 + (i1 - o[1])*nb11 + (i0 - o[0])*nb10);
|
||||
}
|
||||
|
||||
global float * y = (global float *)(dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
|
||||
global int * y = (global int *)(dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0);
|
||||
|
||||
*y = *x;
|
||||
}
|
||||
|
||||
@@ -286,3 +286,28 @@ kernel void kernel_cpy_i32_i32(
|
||||
dst_data[i00] = src[0];
|
||||
}
|
||||
}
|
||||
|
||||
// Contiguous f32 copy, one work item per float4 over the whole tensor. The kernels above map
|
||||
// one workgroup to each row, which leaves a tensor with few long rows on a single compute unit.
|
||||
// vload4/vstore4 rather than a float4 cast: these buffers carry an arbitrary 4-byte view offset.
|
||||
kernel void kernel_cpy_f32_f32_flat(
|
||||
global float * src0,
|
||||
ulong offset0,
|
||||
global float * dst,
|
||||
ulong offsetd,
|
||||
ulong ne,
|
||||
ulong n4
|
||||
) {
|
||||
src0 = (global float*)((global char*)src0 + offset0);
|
||||
dst = (global float*)((global char*)dst + offsetd);
|
||||
|
||||
const ulong i = get_global_id(0);
|
||||
|
||||
if (i < n4) {
|
||||
vstore4(vload4(i, src0), i, dst);
|
||||
} else if (i == n4) {
|
||||
for (ulong t = n4 * 4; t < ne; ++t) {
|
||||
dst[t] = src0[t];
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,871 @@
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
#pragma OPENCL EXTENSION cl_qcom_subgroup_uniform_load : enable
|
||||
#pragma OPENCL EXTENSION cl_qcom_subgroup_constant_load : enable
|
||||
|
||||
#define bool2 uchar2
|
||||
#define bool3 uchar3
|
||||
#define bool4 uchar4
|
||||
|
||||
__constant sampler_t smp_none = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_NONE | CLK_FILTER_NEAREST;
|
||||
__constant sampler_t smp_zero = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;
|
||||
|
||||
__kernel void adreno_xmem_attn_q_f32_to_img_scaled(const global void * src_void,
|
||||
ulong src_offset,
|
||||
write_only image2d_t dst_image2d,
|
||||
const float scale,
|
||||
const int d_head,
|
||||
const int n_q,
|
||||
const int n_head,
|
||||
const int n_head_kv,
|
||||
const int n_batch,
|
||||
const ulong src_nb1,
|
||||
const ulong src_nb2,
|
||||
const ulong src_nb3) {
|
||||
const int x = get_global_id(0);
|
||||
const int flat_h = get_global_id(1);
|
||||
const int d = get_global_id(2);
|
||||
|
||||
const int heads_total = n_head * n_batch;
|
||||
const int kpack = d_head / 4;
|
||||
|
||||
if (x >= n_q || flat_h >= heads_total || d >= kpack) {
|
||||
return;
|
||||
}
|
||||
|
||||
const int batch = flat_h / n_head;
|
||||
const int head = flat_h % n_head;
|
||||
const int gqa = n_head / n_head_kv;
|
||||
const int head_kv = head / gqa;
|
||||
const int head_group = head - head_kv * gqa;
|
||||
const int compact_h = batch * n_head_kv + head_kv;
|
||||
const int compact_x = head_group * n_q + x;
|
||||
const int c = d * 4;
|
||||
|
||||
const global char * src_base = (const global char *) src_void + src_offset;
|
||||
const global float * row_ptr = (const global float *) (src_base + batch * src_nb3 + head * src_nb2 + x * src_nb1);
|
||||
|
||||
half4 out = (half4) (0.0h);
|
||||
out.x = convert_half(row_ptr[c + 0] * scale);
|
||||
if (c + 1 < d_head) {
|
||||
out.y = convert_half(row_ptr[c + 1] * scale);
|
||||
}
|
||||
if (c + 2 < d_head) {
|
||||
out.z = convert_half(row_ptr[c + 2] * scale);
|
||||
}
|
||||
if (c + 3 < d_head) {
|
||||
out.w = convert_half(row_ptr[c + 3] * scale);
|
||||
}
|
||||
|
||||
write_imageh(dst_image2d, (int2) (compact_x, compact_h * kpack + d), out);
|
||||
}
|
||||
|
||||
__kernel void adreno_xmem_attn_kv_f32_to_img_gqa(const global void * src_void,
|
||||
ulong src_offset,
|
||||
write_only image2d_t dst_image2d,
|
||||
const int d_head,
|
||||
const int n_kv,
|
||||
const int n_kv_padded,
|
||||
const int n_head_kv,
|
||||
const int n_batch,
|
||||
const ulong src_nb1,
|
||||
const ulong src_nb2,
|
||||
const ulong src_nb3) {
|
||||
const int x = get_global_id(0);
|
||||
const int flat_h = get_global_id(1);
|
||||
const int d = get_global_id(2);
|
||||
|
||||
const int kv_heads_total = n_head_kv * n_batch;
|
||||
const int kpack = d_head / 4;
|
||||
|
||||
if (x >= n_kv_padded || flat_h >= kv_heads_total || d >= kpack) {
|
||||
return;
|
||||
}
|
||||
|
||||
const int batch = flat_h / n_head_kv;
|
||||
const int head_kv = flat_h % n_head_kv;
|
||||
const int c = d * 4;
|
||||
|
||||
half4 out = (half4) (0.0h);
|
||||
if (x < n_kv) {
|
||||
const global char * src_base = (const global char *) src_void + src_offset;
|
||||
const global float * row_ptr =
|
||||
(const global float *) (src_base + batch * src_nb3 + head_kv * src_nb2 + x * src_nb1);
|
||||
out.x = convert_half(row_ptr[c + 0]);
|
||||
if (c + 1 < d_head) {
|
||||
out.y = convert_half(row_ptr[c + 1]);
|
||||
}
|
||||
if (c + 2 < d_head) {
|
||||
out.z = convert_half(row_ptr[c + 2]);
|
||||
}
|
||||
if (c + 3 < d_head) {
|
||||
out.w = convert_half(row_ptr[c + 3]);
|
||||
}
|
||||
}
|
||||
|
||||
write_imageh(dst_image2d, (int2) (x, flat_h * kpack + d), out);
|
||||
}
|
||||
|
||||
__kernel void adreno_xmem_attn_kv_f16_to_img_gqa(const global void * src_void,
|
||||
ulong src_offset,
|
||||
write_only image2d_t dst_image2d,
|
||||
const int d_head,
|
||||
const int n_kv,
|
||||
const int n_kv_padded,
|
||||
const int n_head_kv,
|
||||
const int n_batch,
|
||||
const ulong src_nb1,
|
||||
const ulong src_nb2,
|
||||
const ulong src_nb3) {
|
||||
const int x = get_global_id(0);
|
||||
const int flat_h = get_global_id(1);
|
||||
const int d = get_global_id(2);
|
||||
|
||||
const int kv_heads_total = n_head_kv * n_batch;
|
||||
const int kpack = d_head / 4;
|
||||
|
||||
if (x >= n_kv_padded || flat_h >= kv_heads_total || d >= kpack) {
|
||||
return;
|
||||
}
|
||||
|
||||
const int batch = flat_h / n_head_kv;
|
||||
const int head_kv = flat_h % n_head_kv;
|
||||
const int c = d * 4;
|
||||
|
||||
half4 out = (half4) (0.0h);
|
||||
if (x < n_kv) {
|
||||
const global char * src_base = (const global char *) src_void + src_offset;
|
||||
const global half * row_ptr =
|
||||
(const global half *) (src_base + batch * src_nb3 + head_kv * src_nb2 + x * src_nb1);
|
||||
out.x = row_ptr[c + 0];
|
||||
if (c + 1 < d_head) {
|
||||
out.y = row_ptr[c + 1];
|
||||
}
|
||||
if (c + 2 < d_head) {
|
||||
out.z = row_ptr[c + 2];
|
||||
}
|
||||
if (c + 3 < d_head) {
|
||||
out.w = row_ptr[c + 3];
|
||||
}
|
||||
}
|
||||
|
||||
write_imageh(dst_image2d, (int2) (x, flat_h * kpack + d), out);
|
||||
}
|
||||
|
||||
__kernel void adreno_xmem_attn_img_to_f32(global void * dst_void,
|
||||
ulong dst_offset,
|
||||
read_only image2d_t src_image2d,
|
||||
const int d_head,
|
||||
const int n_q,
|
||||
const int n_head,
|
||||
const int n_head_kv,
|
||||
const int n_batch,
|
||||
const ulong dst_nb1,
|
||||
const ulong dst_nb2,
|
||||
const ulong dst_nb3) {
|
||||
const int x = get_global_id(0);
|
||||
const int flat_h = get_global_id(1);
|
||||
const int d = get_global_id(2);
|
||||
|
||||
const int heads_total = n_head * n_batch;
|
||||
const int kpack = d_head / 4;
|
||||
|
||||
if (x >= n_q || flat_h >= heads_total || d >= kpack) {
|
||||
return;
|
||||
}
|
||||
|
||||
const int batch = flat_h / n_head;
|
||||
const int head = flat_h % n_head;
|
||||
const int gqa = n_head / n_head_kv;
|
||||
const int head_kv = head / gqa;
|
||||
const int head_group = head - head_kv * gqa;
|
||||
const int compact_h = batch * n_head_kv + head_kv;
|
||||
const int compact_x = head_group * n_q + x;
|
||||
const int c = d * 4;
|
||||
|
||||
global char * dst_base = (global char *) dst_void + dst_offset;
|
||||
global float * row_ptr = (global float *) (dst_base + batch * dst_nb3 + x * dst_nb2 + head * dst_nb1);
|
||||
|
||||
const half4 in_value = read_imageh(src_image2d, smp_zero, (int2) (compact_x, compact_h * kpack + d));
|
||||
row_ptr[c + 0] = convert_float(in_value.x);
|
||||
if (c + 1 < d_head) {
|
||||
row_ptr[c + 1] = convert_float(in_value.y);
|
||||
}
|
||||
if (c + 2 < d_head) {
|
||||
row_ptr[c + 2] = convert_float(in_value.z);
|
||||
}
|
||||
if (c + 3 < d_head) {
|
||||
row_ptr[c + 3] = convert_float(in_value.w);
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void adreno_xmem_attn_k_gather(global half4 * dst_tensor_buffer,
|
||||
read_only image2d_t src_tensor_image2d,
|
||||
const int4 shared_int4_0,
|
||||
const int4 shared_int4_1) {
|
||||
int X = get_global_id(0);
|
||||
int Y = get_global_id(1);
|
||||
int S = get_global_id(2);
|
||||
if (X >= shared_int4_0.w || Y >= shared_int4_0.y || S >= shared_int4_0.z) {
|
||||
return;
|
||||
}
|
||||
half temps[4];
|
||||
temps[0] = (half) (0.f);
|
||||
temps[1] = (half) (0.f);
|
||||
temps[2] = (half) (0.f);
|
||||
temps[3] = (half) (0.f);
|
||||
for (int i = 0; i < 4; ++i) {
|
||||
int dst_channel = S * 4 + i;
|
||||
if (dst_channel < shared_int4_0.x) {
|
||||
int s_y = Y;
|
||||
int s_x = dst_channel;
|
||||
int s_c = X;
|
||||
{
|
||||
int slice_coord_TMP = (s_c) / 4;
|
||||
int sub_ch_coord_TMP = (s_c) % 4;
|
||||
half4 src_TMP = read_imageh(src_tensor_image2d, smp_zero,
|
||||
(int2) ((s_x), ((s_y) *shared_int4_1.x + (slice_coord_TMP))));
|
||||
temps[i] = (half[4]){ src_TMP.x, src_TMP.y, src_TMP.z, src_TMP.w }[sub_ch_coord_TMP];
|
||||
};
|
||||
}
|
||||
}
|
||||
half4 result;
|
||||
result.x = temps[0];
|
||||
result.y = temps[1];
|
||||
result.z = temps[2];
|
||||
result.w = temps[3];
|
||||
dst_tensor_buffer[(((S) *shared_int4_0.y + (Y)) * shared_int4_0.w + (X))] = result;
|
||||
}
|
||||
|
||||
__kernel void adreno_xmem_attn_pack_k(global half4 * dst_tensor_buffer,
|
||||
read_only image1d_buffer_t src_image_buffer,
|
||||
const int4 shared_int4_0,
|
||||
const int4 shared_int4_1,
|
||||
const int4 shared_int4_2) {
|
||||
int linear_index = get_global_id(0);
|
||||
if (linear_index >= shared_int4_0.y) {
|
||||
return;
|
||||
}
|
||||
if (get_global_id(1) != 0) {
|
||||
return;
|
||||
}
|
||||
if (get_global_id(2) != 0) {
|
||||
return;
|
||||
}
|
||||
int dst_o_sp_i_ogroup = linear_index;
|
||||
int dst_ogroup = dst_o_sp_i_ogroup % shared_int4_0.x;
|
||||
int dst_o_sp_i = dst_o_sp_i_ogroup / shared_int4_0.x;
|
||||
int dst_i = dst_o_sp_i % shared_int4_0.z;
|
||||
int dst_o_sp = dst_o_sp_i / shared_int4_0.z;
|
||||
int dst_sp = dst_o_sp % shared_int4_1.x;
|
||||
int dst_o = dst_o_sp / shared_int4_1.x;
|
||||
int i_slice = dst_i;
|
||||
int o_slice = dst_o * shared_int4_0.x + dst_ogroup;
|
||||
int spatial_linear = dst_sp;
|
||||
int W = spatial_linear % shared_int4_1.y;
|
||||
int H = spatial_linear / shared_int4_1.y;
|
||||
half4 w0 = (half4) (0);
|
||||
half4 w1 = (half4) (0);
|
||||
half4 w2 = (half4) (0);
|
||||
half4 w3 = (half4) (0);
|
||||
|
||||
if (i_slice * 4 < shared_int4_0.w && o_slice < shared_int4_1.w) {
|
||||
w0 = read_imageh(src_image_buffer, (((o_slice) *shared_int4_1.z + (W)) * shared_int4_2.x + (i_slice * 4)));
|
||||
}
|
||||
if (i_slice * 4 + 1 < shared_int4_0.w && o_slice < shared_int4_1.w) {
|
||||
w1 = read_imageh(src_image_buffer, (((o_slice) *shared_int4_1.z + (W)) * shared_int4_2.x + (i_slice * 4 + 1)));
|
||||
}
|
||||
if (i_slice * 4 + 2 < shared_int4_0.w && o_slice < shared_int4_1.w) {
|
||||
w2 = read_imageh(src_image_buffer, (((o_slice) *shared_int4_1.z + (W)) * shared_int4_2.x + (i_slice * 4 + 2)));
|
||||
}
|
||||
if (i_slice * 4 + 3 < shared_int4_0.w && o_slice < shared_int4_1.w) {
|
||||
w3 = read_imageh(src_image_buffer, (((o_slice) *shared_int4_1.z + (W)) * shared_int4_2.x + (i_slice * 4 + 3)));
|
||||
}
|
||||
half4 r0 = w0;
|
||||
half4 r1 = w1;
|
||||
half4 r2 = w2;
|
||||
half4 r3 = w3;
|
||||
dst_tensor_buffer[linear_index * 4 + 0] = r0;
|
||||
dst_tensor_buffer[linear_index * 4 + 1] = r1;
|
||||
dst_tensor_buffer[linear_index * 4 + 2] = r2;
|
||||
dst_tensor_buffer[linear_index * 4 + 3] = r3;
|
||||
}
|
||||
|
||||
__attribute__((qcom_max_concurrent_subgroups(12))) __kernel void adreno_xmem_attn_qk_gemm(
|
||||
global half4 * dst_tensor_buffer,
|
||||
constant half8 * weights_buffer __attribute__((sub_group_uniform)),
|
||||
constant half8 * xmem_buffer __attribute__((max_constant_size((6144)))),
|
||||
read_only image2d_t src_tensor_image2d,
|
||||
const int4 shared_int4_0,
|
||||
const int4 shared_int4_1,
|
||||
const int4 shared_int4_2) {
|
||||
int X = get_group_id(1) * get_local_size(0) + get_local_id(0);
|
||||
int Y = get_group_id(2) * get_local_size(1) + get_local_id(1);
|
||||
int Z = get_group_id(0) * get_local_size(2) + get_local_id(2);
|
||||
if (X >= shared_int4_0.z || Y >= shared_int4_0.x) {
|
||||
return;
|
||||
}
|
||||
if (Z * 8 >= shared_int4_0.y) {
|
||||
return;
|
||||
}
|
||||
|
||||
half4 r0 = (half4) (0.f);
|
||||
half4 r1 = (half4) (0.f);
|
||||
half4 r2 = (half4) (0.f);
|
||||
half4 r3 = (half4) (0.f);
|
||||
half4 r4 = (half4) (0.f);
|
||||
half4 r5 = (half4) (0.f);
|
||||
half4 r6 = (half4) (0.f);
|
||||
half4 r7 = (half4) (0.f);
|
||||
int x_coord = mad24(X, shared_int4_2.y, shared_int4_1.y);
|
||||
int y_coord = mad24(Y, shared_int4_2.z, shared_int4_1.z);
|
||||
int coord_x, coord_y, coord_s;
|
||||
int f_offset = (Z * shared_int4_1.w + Y) * shared_int4_1.x * 32;
|
||||
|
||||
int subgroup_id = (int) ((0x1F & qcom_get_physical_sub_group_id()));
|
||||
subgroup_id = subgroup_id % 12;
|
||||
int c_offset = mul24(subgroup_id, shared_int4_0.w);
|
||||
__constant half16 * weights_cache = (__constant half16 *) &xmem_buffer[c_offset];
|
||||
coord_y = Y;
|
||||
coord_x = X;
|
||||
coord_s = 0;
|
||||
do {
|
||||
half4 src0 =
|
||||
read_imageh(src_tensor_image2d, smp_zero, (int2) ((coord_x), ((coord_y) *shared_int4_2.x + (coord_s))));
|
||||
coord_s++;
|
||||
half4 src1 =
|
||||
read_imageh(src_tensor_image2d, smp_zero, (int2) ((coord_x), ((coord_y) *shared_int4_2.x + (coord_s))));
|
||||
coord_s++;
|
||||
qcom_sub_group_constant_load8(xmem_buffer, weights_buffer, c_offset, f_offset >> 1, 32);
|
||||
f_offset += 64;
|
||||
qcom_sub_group_sync(QCOM_CLK_CONST_LOAD_SYNC);
|
||||
r0 += src0.x * weights_cache[0].s0123;
|
||||
r0 += src0.y * weights_cache[0].s4567;
|
||||
r0 += src0.z * weights_cache[0].s89ab;
|
||||
r0 += src0.w * weights_cache[0].scdef;
|
||||
r1 += src0.x * weights_cache[1].s0123;
|
||||
r1 += src0.y * weights_cache[1].s4567;
|
||||
r1 += src0.z * weights_cache[1].s89ab;
|
||||
r1 += src0.w * weights_cache[1].scdef;
|
||||
r2 += src0.x * weights_cache[2].s0123;
|
||||
r2 += src0.y * weights_cache[2].s4567;
|
||||
r2 += src0.z * weights_cache[2].s89ab;
|
||||
r2 += src0.w * weights_cache[2].scdef;
|
||||
r3 += src0.x * weights_cache[3].s0123;
|
||||
r3 += src0.y * weights_cache[3].s4567;
|
||||
r3 += src0.z * weights_cache[3].s89ab;
|
||||
r3 += src0.w * weights_cache[3].scdef;
|
||||
r4 += src0.x * weights_cache[4].s0123;
|
||||
r4 += src0.y * weights_cache[4].s4567;
|
||||
r4 += src0.z * weights_cache[4].s89ab;
|
||||
r4 += src0.w * weights_cache[4].scdef;
|
||||
r5 += src0.x * weights_cache[5].s0123;
|
||||
r5 += src0.y * weights_cache[5].s4567;
|
||||
r5 += src0.z * weights_cache[5].s89ab;
|
||||
r5 += src0.w * weights_cache[5].scdef;
|
||||
r6 += src0.x * weights_cache[6].s0123;
|
||||
r6 += src0.y * weights_cache[6].s4567;
|
||||
r6 += src0.z * weights_cache[6].s89ab;
|
||||
r6 += src0.w * weights_cache[6].scdef;
|
||||
r7 += src0.x * weights_cache[7].s0123;
|
||||
r7 += src0.y * weights_cache[7].s4567;
|
||||
r7 += src0.z * weights_cache[7].s89ab;
|
||||
r7 += src0.w * weights_cache[7].scdef;
|
||||
r0 += src1.x * weights_cache[8].s0123;
|
||||
r0 += src1.y * weights_cache[8].s4567;
|
||||
r0 += src1.z * weights_cache[8].s89ab;
|
||||
r0 += src1.w * weights_cache[8].scdef;
|
||||
r1 += src1.x * weights_cache[9].s0123;
|
||||
r1 += src1.y * weights_cache[9].s4567;
|
||||
r1 += src1.z * weights_cache[9].s89ab;
|
||||
r1 += src1.w * weights_cache[9].scdef;
|
||||
r2 += src1.x * weights_cache[10].s0123;
|
||||
r2 += src1.y * weights_cache[10].s4567;
|
||||
r2 += src1.z * weights_cache[10].s89ab;
|
||||
r2 += src1.w * weights_cache[10].scdef;
|
||||
r3 += src1.x * weights_cache[11].s0123;
|
||||
r3 += src1.y * weights_cache[11].s4567;
|
||||
r3 += src1.z * weights_cache[11].s89ab;
|
||||
r3 += src1.w * weights_cache[11].scdef;
|
||||
r4 += src1.x * weights_cache[12].s0123;
|
||||
r4 += src1.y * weights_cache[12].s4567;
|
||||
r4 += src1.z * weights_cache[12].s89ab;
|
||||
r4 += src1.w * weights_cache[12].scdef;
|
||||
r5 += src1.x * weights_cache[13].s0123;
|
||||
r5 += src1.y * weights_cache[13].s4567;
|
||||
r5 += src1.z * weights_cache[13].s89ab;
|
||||
r5 += src1.w * weights_cache[13].scdef;
|
||||
r6 += src1.x * weights_cache[14].s0123;
|
||||
r6 += src1.y * weights_cache[14].s4567;
|
||||
r6 += src1.z * weights_cache[14].s89ab;
|
||||
r6 += src1.w * weights_cache[14].scdef;
|
||||
r7 += src1.x * weights_cache[15].s0123;
|
||||
r7 += src1.y * weights_cache[15].s4567;
|
||||
r7 += src1.z * weights_cache[15].s89ab;
|
||||
r7 += src1.w * weights_cache[15].scdef;
|
||||
} while (coord_s < shared_int4_2.x);
|
||||
|
||||
coord_s = mul24(Z, 8);
|
||||
coord_x = X;
|
||||
coord_y = Y;
|
||||
if (coord_s < shared_int4_0.y) {
|
||||
half4 res = convert_half4(r0);
|
||||
if (coord_s < 0) {
|
||||
res += read_imageh(src_tensor_image2d, smp_zero, (int2) ((0), ((0) * shared_int4_2.x + (0))));
|
||||
}
|
||||
dst_tensor_buffer[(((coord_s) *shared_int4_0.x + (coord_y)) * shared_int4_0.z + (coord_x))] = res;
|
||||
coord_s++;
|
||||
}
|
||||
if (coord_s < shared_int4_0.y) {
|
||||
half4 res = convert_half4(r1);
|
||||
if (coord_s < 0) {
|
||||
res += read_imageh(src_tensor_image2d, smp_zero, (int2) ((0), ((0) * shared_int4_2.x + (0))));
|
||||
}
|
||||
dst_tensor_buffer[(((coord_s) *shared_int4_0.x + (coord_y)) * shared_int4_0.z + (coord_x))] = res;
|
||||
coord_s++;
|
||||
}
|
||||
if (coord_s < shared_int4_0.y) {
|
||||
half4 res = convert_half4(r2);
|
||||
if (coord_s < 0) {
|
||||
res += read_imageh(src_tensor_image2d, smp_zero, (int2) ((0), ((0) * shared_int4_2.x + (0))));
|
||||
}
|
||||
dst_tensor_buffer[(((coord_s) *shared_int4_0.x + (coord_y)) * shared_int4_0.z + (coord_x))] = res;
|
||||
coord_s++;
|
||||
}
|
||||
if (coord_s < shared_int4_0.y) {
|
||||
half4 res = convert_half4(r3);
|
||||
if (coord_s < 0) {
|
||||
res += read_imageh(src_tensor_image2d, smp_zero, (int2) ((0), ((0) * shared_int4_2.x + (0))));
|
||||
}
|
||||
dst_tensor_buffer[(((coord_s) *shared_int4_0.x + (coord_y)) * shared_int4_0.z + (coord_x))] = res;
|
||||
coord_s++;
|
||||
}
|
||||
if (coord_s < shared_int4_0.y) {
|
||||
half4 res = convert_half4(r4);
|
||||
if (coord_s < 0) {
|
||||
res += read_imageh(src_tensor_image2d, smp_zero, (int2) ((0), ((0) * shared_int4_2.x + (0))));
|
||||
}
|
||||
dst_tensor_buffer[(((coord_s) *shared_int4_0.x + (coord_y)) * shared_int4_0.z + (coord_x))] = res;
|
||||
coord_s++;
|
||||
}
|
||||
if (coord_s < shared_int4_0.y) {
|
||||
half4 res = convert_half4(r5);
|
||||
if (coord_s < 0) {
|
||||
res += read_imageh(src_tensor_image2d, smp_zero, (int2) ((0), ((0) * shared_int4_2.x + (0))));
|
||||
}
|
||||
dst_tensor_buffer[(((coord_s) *shared_int4_0.x + (coord_y)) * shared_int4_0.z + (coord_x))] = res;
|
||||
coord_s++;
|
||||
}
|
||||
if (coord_s < shared_int4_0.y) {
|
||||
half4 res = convert_half4(r6);
|
||||
if (coord_s < 0) {
|
||||
res += read_imageh(src_tensor_image2d, smp_zero, (int2) ((0), ((0) * shared_int4_2.x + (0))));
|
||||
}
|
||||
dst_tensor_buffer[(((coord_s) *shared_int4_0.x + (coord_y)) * shared_int4_0.z + (coord_x))] = res;
|
||||
coord_s++;
|
||||
}
|
||||
if (coord_s < shared_int4_0.y) {
|
||||
half4 res = convert_half4(r7);
|
||||
if (coord_s < 0) {
|
||||
res += read_imageh(src_tensor_image2d, smp_zero, (int2) ((0), ((0) * shared_int4_2.x + (0))));
|
||||
}
|
||||
dst_tensor_buffer[(((coord_s) *shared_int4_0.x + (coord_y)) * shared_int4_0.z + (coord_x))] = res;
|
||||
coord_s++;
|
||||
}
|
||||
}
|
||||
|
||||
__kernel void adreno_xmem_attn_softmax_reduce_basic(read_only image1d_buffer_t src_tensor_image_buffer,
|
||||
write_only image2d_t dst_tensor_image2d,
|
||||
const int4 shared_int4_0,
|
||||
const int4 shared_int4_1) {
|
||||
int X = get_global_id(0);
|
||||
int Y = get_global_id(1);
|
||||
if (X >= shared_int4_0.z || Y >= shared_int4_0.x) {
|
||||
return;
|
||||
}
|
||||
float sum = 0.0f;
|
||||
int end_channel = shared_int4_0.w;
|
||||
int end_slice = (end_channel + 3) / 4;
|
||||
int start_channel = 0;
|
||||
int start_slice = start_channel / 4;
|
||||
bool need_per_channels_check = start_channel % 4 != 0 || end_channel % 4 != 0;
|
||||
float maximum;
|
||||
{
|
||||
int slice_coord_TMP = (start_channel) / 4;
|
||||
int sub_ch_coord_TMP = (start_channel) % 4;
|
||||
float4 src_TMP = convert_float4(
|
||||
read_imageh(src_tensor_image_buffer, ((slice_coord_TMP) *shared_int4_1.x + (Y)) * shared_int4_1.y + (X)));
|
||||
maximum = (float[4]){ src_TMP.x, src_TMP.y, src_TMP.z, src_TMP.w }[sub_ch_coord_TMP];
|
||||
};
|
||||
for (int d = start_slice; d < end_slice; d += 1) {
|
||||
float4 mask_dot = (float4) (1.f);
|
||||
float4 src =
|
||||
convert_float4(read_imageh(src_tensor_image_buffer, ((d) *shared_int4_1.x + (Y)) * shared_int4_1.y + (X)));
|
||||
if (need_per_channels_check && (d == start_slice || d == end_slice - 1)) {
|
||||
if (d * 4 + 0 < start_channel || d * 4 + 0 >= end_channel) {
|
||||
mask_dot.x = 0.f;
|
||||
src.x = maximum;
|
||||
}
|
||||
if (d * 4 + 1 < start_channel || d * 4 + 1 >= end_channel) {
|
||||
mask_dot.y = 0.f;
|
||||
src.y = maximum;
|
||||
}
|
||||
if (d * 4 + 2 < start_channel || d * 4 + 2 >= end_channel) {
|
||||
mask_dot.z = 0.f;
|
||||
src.z = maximum;
|
||||
}
|
||||
if (d * 4 + 3 < start_channel || d * 4 + 3 >= end_channel) {
|
||||
mask_dot.w = 0.f;
|
||||
src.w = maximum;
|
||||
}
|
||||
}
|
||||
float new_max = max(src.x, src.y);
|
||||
new_max = max(new_max, src.z);
|
||||
new_max = max(new_max, src.w);
|
||||
new_max = max(new_max, maximum);
|
||||
float scale = native_exp(maximum - new_max);
|
||||
maximum = new_max;
|
||||
sum *= scale;
|
||||
float4 exp_res = native_exp(src - maximum);
|
||||
sum += dot(mask_dot, exp_res);
|
||||
}
|
||||
if (!isfinite(maximum) || sum == 0.0f) {
|
||||
write_imageh(dst_tensor_image2d, (int2) (X, Y), (half4) (0.0h));
|
||||
return;
|
||||
}
|
||||
write_imageh(dst_tensor_image2d, (int2) (X, Y),
|
||||
(half4) (convert_half(1.0f / sum), convert_half(maximum), 0.0h, 0.0h));
|
||||
}
|
||||
|
||||
__kernel void adreno_xmem_attn_softmax_apply_basic(global half4 * dst_tensor_buffer,
|
||||
read_only image1d_buffer_t src_tensor_image_buffer,
|
||||
read_only image2d_t src_tensor_1_image2d,
|
||||
const int4 shared_int4_0,
|
||||
const int4 shared_int4_1) {
|
||||
int X = get_global_id(0);
|
||||
int Y = get_global_id(1);
|
||||
int Z = get_global_id(2);
|
||||
if (X >= shared_int4_0.z || Y >= shared_int4_0.x || Z >= shared_int4_0.y) {
|
||||
return;
|
||||
}
|
||||
half4 src = read_imageh(src_tensor_image_buffer, ((Z) *shared_int4_1.x + (Y)) * shared_int4_1.y + (X));
|
||||
{
|
||||
half4 src_final;
|
||||
{
|
||||
{
|
||||
half4 exp_val = read_imageh(src_tensor_1_image2d, smp_zero, (int2) (X, Y));
|
||||
src_final = exp(src - exp_val.y) * exp_val.x;
|
||||
const int k = Z * 4;
|
||||
const int n_kv = shared_int4_1.z;
|
||||
if (k + 0 >= n_kv) {
|
||||
src_final.x = 0.0h;
|
||||
}
|
||||
if (k + 1 >= n_kv) {
|
||||
src_final.y = 0.0h;
|
||||
}
|
||||
if (k + 2 >= n_kv) {
|
||||
src_final.z = 0.0h;
|
||||
}
|
||||
if (k + 3 >= n_kv) {
|
||||
src_final.w = 0.0h;
|
||||
}
|
||||
}
|
||||
}
|
||||
dst_tensor_buffer[(((Z) *shared_int4_0.x + (Y)) * shared_int4_0.z + (X))] = src_final;
|
||||
};
|
||||
}
|
||||
|
||||
__kernel void adreno_xmem_attn_mask_scores(global half4 * dst_score_tensor_buffer,
|
||||
read_only image1d_buffer_t src_score_image_buffer,
|
||||
const global half * mask,
|
||||
const ulong mask_offset,
|
||||
const int q_width,
|
||||
const int n_q,
|
||||
const int n_kv,
|
||||
const int n_kv_padded,
|
||||
const int kv_heads_total,
|
||||
const int n_head,
|
||||
const int n_head_kv,
|
||||
const ulong mask_nb1,
|
||||
const ulong mask_nb2,
|
||||
const ulong mask_nb3,
|
||||
const int mask_ne2,
|
||||
const int mask_ne3) {
|
||||
const int X = get_global_id(0);
|
||||
const int Y = get_global_id(1);
|
||||
const int Z = get_global_id(2);
|
||||
const int npack = n_kv_padded / 4;
|
||||
if (X >= q_width || Y >= kv_heads_total || Z >= npack) {
|
||||
return;
|
||||
}
|
||||
|
||||
const int gqa = n_head / n_head_kv;
|
||||
const int head_kv = Y % n_head_kv;
|
||||
const int batch = Y / n_head_kv;
|
||||
const int head_group = X / n_q;
|
||||
const int q = X - head_group * n_q;
|
||||
const int head = head_kv * gqa + head_group;
|
||||
const int mask_head_idx = head % mask_ne2;
|
||||
const int mask_batch_idx = batch % mask_ne3;
|
||||
const global char * mask_base = (const global char *) mask + mask_offset;
|
||||
const global half * mask_row = (const global half *) (mask_base + mask_batch_idx * mask_nb3 +
|
||||
mask_head_idx * mask_nb2 + q * mask_nb1);
|
||||
|
||||
const half4 score = read_imageh(src_score_image_buffer, ((Z * kv_heads_total + Y) * q_width + X));
|
||||
float vals[4] = {
|
||||
convert_float(score.x),
|
||||
convert_float(score.y),
|
||||
convert_float(score.z),
|
||||
convert_float(score.w),
|
||||
};
|
||||
|
||||
for (int lane = 0; lane < 4; ++lane) {
|
||||
const int k_idx = Z * 4 + lane;
|
||||
if (k_idx >= n_kv) {
|
||||
vals[lane] = -INFINITY;
|
||||
} else {
|
||||
vals[lane] += convert_float(mask_row[k_idx]);
|
||||
}
|
||||
}
|
||||
|
||||
dst_score_tensor_buffer[((Z * kv_heads_total + Y) * q_width + X)] =
|
||||
(half4) (convert_half(vals[0]), convert_half(vals[1]), convert_half(vals[2]), convert_half(vals[3]));
|
||||
}
|
||||
|
||||
__kernel void adreno_xmem_attn_pack_v(global half4 * dst_tensor_buffer,
|
||||
read_only image2d_t src_image2d,
|
||||
const int4 shared_int4_0,
|
||||
const int4 shared_int4_1) {
|
||||
int linear_index = get_global_id(0);
|
||||
if (linear_index >= shared_int4_0.y) {
|
||||
return;
|
||||
}
|
||||
if (get_global_id(1) != 0) {
|
||||
return;
|
||||
}
|
||||
if (get_global_id(2) != 0) {
|
||||
return;
|
||||
}
|
||||
int dst_o_sp_i_ogroup = linear_index;
|
||||
int dst_ogroup = dst_o_sp_i_ogroup % shared_int4_0.x;
|
||||
int dst_o_sp_i = dst_o_sp_i_ogroup / shared_int4_0.x;
|
||||
int dst_i = dst_o_sp_i % shared_int4_0.z;
|
||||
int dst_o_sp = dst_o_sp_i / shared_int4_0.z;
|
||||
int dst_sp = dst_o_sp % shared_int4_1.x;
|
||||
int dst_o = dst_o_sp / shared_int4_1.x;
|
||||
int i_slice = dst_i;
|
||||
int o_slice = dst_o * shared_int4_0.x + dst_ogroup;
|
||||
int spatial_linear = dst_sp;
|
||||
int W = spatial_linear % shared_int4_1.y;
|
||||
int H = spatial_linear / shared_int4_1.y;
|
||||
half4 w0 = (half4) (0);
|
||||
half4 w1 = (half4) (0);
|
||||
half4 w2 = (half4) (0);
|
||||
half4 w3 = (half4) (0);
|
||||
|
||||
if (i_slice * 4 < shared_int4_0.w && o_slice < shared_int4_1.z) {
|
||||
w0 = read_imageh(src_image2d, smp_zero, (int2) ((i_slice * 4), ((W) *shared_int4_1.z + (o_slice))));
|
||||
}
|
||||
if (i_slice * 4 + 1 < shared_int4_0.w && o_slice < shared_int4_1.z) {
|
||||
w1 = read_imageh(src_image2d, smp_zero, (int2) ((i_slice * 4 + 1), ((W) *shared_int4_1.z + (o_slice))));
|
||||
}
|
||||
if (i_slice * 4 + 2 < shared_int4_0.w && o_slice < shared_int4_1.z) {
|
||||
w2 = read_imageh(src_image2d, smp_zero, (int2) ((i_slice * 4 + 2), ((W) *shared_int4_1.z + (o_slice))));
|
||||
}
|
||||
if (i_slice * 4 + 3 < shared_int4_0.w && o_slice < shared_int4_1.z) {
|
||||
w3 = read_imageh(src_image2d, smp_zero, (int2) ((i_slice * 4 + 3), ((W) *shared_int4_1.z + (o_slice))));
|
||||
}
|
||||
half4 r0 = w0;
|
||||
half4 r1 = w1;
|
||||
half4 r2 = w2;
|
||||
half4 r3 = w3;
|
||||
dst_tensor_buffer[linear_index * 4 + 0] = r0;
|
||||
dst_tensor_buffer[linear_index * 4 + 1] = r1;
|
||||
dst_tensor_buffer[linear_index * 4 + 2] = r2;
|
||||
dst_tensor_buffer[linear_index * 4 + 3] = r3;
|
||||
}
|
||||
|
||||
__attribute__((qcom_max_concurrent_subgroups(12))) __kernel void adreno_xmem_attn_pv_gemm(
|
||||
constant half8 * weights_buffer __attribute__((sub_group_uniform)),
|
||||
constant half8 * xmem_buffer __attribute__((max_constant_size((6144)))),
|
||||
read_only image1d_buffer_t src_tensor_image_buffer,
|
||||
write_only image2d_t dst_tensor_image2d,
|
||||
const int4 shared_int4_0,
|
||||
const int4 shared_int4_1,
|
||||
const int4 shared_int4_2,
|
||||
const int4 shared_int4_3) {
|
||||
int X = get_group_id(1) * get_local_size(0) + get_local_id(0);
|
||||
int Y = get_group_id(2) * get_local_size(1) + get_local_id(1);
|
||||
int Z = get_group_id(0) * get_local_size(2) + get_local_id(2);
|
||||
if (X >= shared_int4_0.z || Y >= shared_int4_0.x) {
|
||||
return;
|
||||
}
|
||||
if (Z * 8 >= shared_int4_0.y) {
|
||||
return;
|
||||
}
|
||||
|
||||
half4 r0 = (half4) (0.f);
|
||||
half4 r1 = (half4) (0.f);
|
||||
half4 r2 = (half4) (0.f);
|
||||
half4 r3 = (half4) (0.f);
|
||||
half4 r4 = (half4) (0.f);
|
||||
half4 r5 = (half4) (0.f);
|
||||
half4 r6 = (half4) (0.f);
|
||||
half4 r7 = (half4) (0.f);
|
||||
int x_coord = mad24(X, shared_int4_2.w, shared_int4_1.y);
|
||||
int y_coord = mad24(Y, shared_int4_3.x, shared_int4_1.z);
|
||||
int coord_x, coord_y, coord_s;
|
||||
int f_offset = (Z * shared_int4_1.w + Y) * shared_int4_1.x * 32;
|
||||
|
||||
int subgroup_id = (int) ((0x1F & qcom_get_physical_sub_group_id()));
|
||||
subgroup_id = subgroup_id % 12;
|
||||
int c_offset = mul24(subgroup_id, shared_int4_0.w);
|
||||
__constant half16 * weights_cache = (__constant half16 *) &xmem_buffer[c_offset];
|
||||
coord_y = Y;
|
||||
coord_x = X;
|
||||
int addr = (((0) * shared_int4_1.w + (coord_y)) * shared_int4_2.z + (coord_x));
|
||||
int dz = shared_int4_2.x;
|
||||
coord_s = 0;
|
||||
do {
|
||||
half4 src0 = read_imageh(src_tensor_image_buffer, addr);
|
||||
addr += dz;
|
||||
coord_s++;
|
||||
half4 src1 = read_imageh(src_tensor_image_buffer, addr);
|
||||
addr += dz;
|
||||
coord_s++;
|
||||
qcom_sub_group_constant_load8(xmem_buffer, weights_buffer, c_offset, f_offset >> 1, 32);
|
||||
f_offset += 64;
|
||||
qcom_sub_group_sync(QCOM_CLK_CONST_LOAD_SYNC);
|
||||
r0 += src0.x * weights_cache[0].s0123;
|
||||
r0 += src0.y * weights_cache[0].s4567;
|
||||
r0 += src0.z * weights_cache[0].s89ab;
|
||||
r0 += src0.w * weights_cache[0].scdef;
|
||||
r1 += src0.x * weights_cache[1].s0123;
|
||||
r1 += src0.y * weights_cache[1].s4567;
|
||||
r1 += src0.z * weights_cache[1].s89ab;
|
||||
r1 += src0.w * weights_cache[1].scdef;
|
||||
r2 += src0.x * weights_cache[2].s0123;
|
||||
r2 += src0.y * weights_cache[2].s4567;
|
||||
r2 += src0.z * weights_cache[2].s89ab;
|
||||
r2 += src0.w * weights_cache[2].scdef;
|
||||
r3 += src0.x * weights_cache[3].s0123;
|
||||
r3 += src0.y * weights_cache[3].s4567;
|
||||
r3 += src0.z * weights_cache[3].s89ab;
|
||||
r3 += src0.w * weights_cache[3].scdef;
|
||||
r4 += src0.x * weights_cache[4].s0123;
|
||||
r4 += src0.y * weights_cache[4].s4567;
|
||||
r4 += src0.z * weights_cache[4].s89ab;
|
||||
r4 += src0.w * weights_cache[4].scdef;
|
||||
r5 += src0.x * weights_cache[5].s0123;
|
||||
r5 += src0.y * weights_cache[5].s4567;
|
||||
r5 += src0.z * weights_cache[5].s89ab;
|
||||
r5 += src0.w * weights_cache[5].scdef;
|
||||
r6 += src0.x * weights_cache[6].s0123;
|
||||
r6 += src0.y * weights_cache[6].s4567;
|
||||
r6 += src0.z * weights_cache[6].s89ab;
|
||||
r6 += src0.w * weights_cache[6].scdef;
|
||||
r7 += src0.x * weights_cache[7].s0123;
|
||||
r7 += src0.y * weights_cache[7].s4567;
|
||||
r7 += src0.z * weights_cache[7].s89ab;
|
||||
r7 += src0.w * weights_cache[7].scdef;
|
||||
r0 += src1.x * weights_cache[8].s0123;
|
||||
r0 += src1.y * weights_cache[8].s4567;
|
||||
r0 += src1.z * weights_cache[8].s89ab;
|
||||
r0 += src1.w * weights_cache[8].scdef;
|
||||
r1 += src1.x * weights_cache[9].s0123;
|
||||
r1 += src1.y * weights_cache[9].s4567;
|
||||
r1 += src1.z * weights_cache[9].s89ab;
|
||||
r1 += src1.w * weights_cache[9].scdef;
|
||||
r2 += src1.x * weights_cache[10].s0123;
|
||||
r2 += src1.y * weights_cache[10].s4567;
|
||||
r2 += src1.z * weights_cache[10].s89ab;
|
||||
r2 += src1.w * weights_cache[10].scdef;
|
||||
r3 += src1.x * weights_cache[11].s0123;
|
||||
r3 += src1.y * weights_cache[11].s4567;
|
||||
r3 += src1.z * weights_cache[11].s89ab;
|
||||
r3 += src1.w * weights_cache[11].scdef;
|
||||
r4 += src1.x * weights_cache[12].s0123;
|
||||
r4 += src1.y * weights_cache[12].s4567;
|
||||
r4 += src1.z * weights_cache[12].s89ab;
|
||||
r4 += src1.w * weights_cache[12].scdef;
|
||||
r5 += src1.x * weights_cache[13].s0123;
|
||||
r5 += src1.y * weights_cache[13].s4567;
|
||||
r5 += src1.z * weights_cache[13].s89ab;
|
||||
r5 += src1.w * weights_cache[13].scdef;
|
||||
r6 += src1.x * weights_cache[14].s0123;
|
||||
r6 += src1.y * weights_cache[14].s4567;
|
||||
r6 += src1.z * weights_cache[14].s89ab;
|
||||
r6 += src1.w * weights_cache[14].scdef;
|
||||
r7 += src1.x * weights_cache[15].s0123;
|
||||
r7 += src1.y * weights_cache[15].s4567;
|
||||
r7 += src1.z * weights_cache[15].s89ab;
|
||||
r7 += src1.w * weights_cache[15].scdef;
|
||||
} while (coord_s < shared_int4_2.y);
|
||||
|
||||
coord_s = mul24(Z, 8);
|
||||
coord_x = X;
|
||||
coord_y = Y;
|
||||
if (coord_s < shared_int4_0.y) {
|
||||
half4 res = convert_half4(r0);
|
||||
if (coord_s < 0) {
|
||||
res += read_imageh(src_tensor_image_buffer, ((0) * shared_int4_1.w + (0)) * shared_int4_2.z + (0));
|
||||
}
|
||||
write_imageh(dst_tensor_image2d, (int2) ((coord_x), ((coord_y) *shared_int4_0.y + (coord_s))), res);
|
||||
coord_s++;
|
||||
}
|
||||
if (coord_s < shared_int4_0.y) {
|
||||
half4 res = convert_half4(r1);
|
||||
if (coord_s < 0) {
|
||||
res += read_imageh(src_tensor_image_buffer, ((0) * shared_int4_1.w + (0)) * shared_int4_2.z + (0));
|
||||
}
|
||||
write_imageh(dst_tensor_image2d, (int2) ((coord_x), ((coord_y) *shared_int4_0.y + (coord_s))), res);
|
||||
coord_s++;
|
||||
}
|
||||
if (coord_s < shared_int4_0.y) {
|
||||
half4 res = convert_half4(r2);
|
||||
if (coord_s < 0) {
|
||||
res += read_imageh(src_tensor_image_buffer, ((0) * shared_int4_1.w + (0)) * shared_int4_2.z + (0));
|
||||
}
|
||||
write_imageh(dst_tensor_image2d, (int2) ((coord_x), ((coord_y) *shared_int4_0.y + (coord_s))), res);
|
||||
coord_s++;
|
||||
}
|
||||
if (coord_s < shared_int4_0.y) {
|
||||
half4 res = convert_half4(r3);
|
||||
if (coord_s < 0) {
|
||||
res += read_imageh(src_tensor_image_buffer, ((0) * shared_int4_1.w + (0)) * shared_int4_2.z + (0));
|
||||
}
|
||||
write_imageh(dst_tensor_image2d, (int2) ((coord_x), ((coord_y) *shared_int4_0.y + (coord_s))), res);
|
||||
coord_s++;
|
||||
}
|
||||
if (coord_s < shared_int4_0.y) {
|
||||
half4 res = convert_half4(r4);
|
||||
if (coord_s < 0) {
|
||||
res += read_imageh(src_tensor_image_buffer, ((0) * shared_int4_1.w + (0)) * shared_int4_2.z + (0));
|
||||
}
|
||||
write_imageh(dst_tensor_image2d, (int2) ((coord_x), ((coord_y) *shared_int4_0.y + (coord_s))), res);
|
||||
coord_s++;
|
||||
}
|
||||
if (coord_s < shared_int4_0.y) {
|
||||
half4 res = convert_half4(r5);
|
||||
if (coord_s < 0) {
|
||||
res += read_imageh(src_tensor_image_buffer, ((0) * shared_int4_1.w + (0)) * shared_int4_2.z + (0));
|
||||
}
|
||||
write_imageh(dst_tensor_image2d, (int2) ((coord_x), ((coord_y) *shared_int4_0.y + (coord_s))), res);
|
||||
coord_s++;
|
||||
}
|
||||
if (coord_s < shared_int4_0.y) {
|
||||
half4 res = convert_half4(r6);
|
||||
if (coord_s < 0) {
|
||||
res += read_imageh(src_tensor_image_buffer, ((0) * shared_int4_1.w + (0)) * shared_int4_2.z + (0));
|
||||
}
|
||||
write_imageh(dst_tensor_image2d, (int2) ((coord_x), ((coord_y) *shared_int4_0.y + (coord_s))), res);
|
||||
coord_s++;
|
||||
}
|
||||
if (coord_s < shared_int4_0.y) {
|
||||
half4 res = convert_half4(r7);
|
||||
if (coord_s < 0) {
|
||||
res += read_imageh(src_tensor_image_buffer, ((0) * shared_int4_1.w + (0)) * shared_int4_2.z + (0));
|
||||
}
|
||||
write_imageh(dst_tensor_image2d, (int2) ((coord_x), ((coord_y) *shared_int4_0.y + (coord_s))), res);
|
||||
coord_s++;
|
||||
}
|
||||
}
|
||||
@@ -0,0 +1,85 @@
|
||||
#pragma OPENCL EXTENSION cl_khr_fp16 : enable
|
||||
|
||||
//------------------------------------------------------------------------------
|
||||
// Extended elementwise unary ops, same variant shape as abs.cl:
|
||||
// f32, f32_4 (vec4), f16, f16_4 (vec4), f32_nc, f16_nc (stride-addressed).
|
||||
//
|
||||
// sgn, step, elu, hardswish, hardsigmoid, floor, ceil, round, trunc.
|
||||
//
|
||||
// Semantics match the ggml CPU reference (ggml.c). Values are computed in float
|
||||
// (the f16 variants read/write half and convert), so the conditional ops match
|
||||
// the CPU bit-for-bit within tolerance. SEXPR is the scalar form, VEXPR the
|
||||
// float4 form (vector ternaries need select()).
|
||||
//------------------------------------------------------------------------------
|
||||
|
||||
#define UNARY_EXT(NAME, SEXPR, VEXPR) \
|
||||
kernel void kernel_##NAME##_f32( \
|
||||
global const float * src0, ulong offset0, \
|
||||
global float * dst, ulong offsetd) { \
|
||||
src0 = (global float*)((global char*)src0 + offset0); \
|
||||
dst = (global float*)((global char*)dst + offsetd); \
|
||||
float x = src0[get_global_id(0)]; \
|
||||
dst[get_global_id(0)] = (SEXPR); \
|
||||
} \
|
||||
kernel void kernel_##NAME##_f32_4( \
|
||||
global const float4 * src0, ulong offset0, \
|
||||
global float4 * dst, ulong offsetd) { \
|
||||
src0 = (global float4*)((global char*)src0 + offset0); \
|
||||
dst = (global float4*)((global char*)dst + offsetd); \
|
||||
float4 x = src0[get_global_id(0)]; \
|
||||
dst[get_global_id(0)] = (VEXPR); \
|
||||
} \
|
||||
kernel void kernel_##NAME##_f16( \
|
||||
global const half * src0, ulong offset0, \
|
||||
global half * dst, ulong offsetd) { \
|
||||
src0 = (global half*)((global char*)src0 + offset0); \
|
||||
dst = (global half*)((global char*)dst + offsetd); \
|
||||
float x = src0[get_global_id(0)]; \
|
||||
dst[get_global_id(0)] = (SEXPR); \
|
||||
} \
|
||||
kernel void kernel_##NAME##_f16_4( \
|
||||
global const half4 * src0, ulong offset0, \
|
||||
global half4 * dst, ulong offsetd) { \
|
||||
src0 = (global half4*)((global char*)src0 + offset0); \
|
||||
dst = (global half4*)((global char*)dst + offsetd); \
|
||||
float4 x = convert_float4(src0[get_global_id(0)]); \
|
||||
dst[get_global_id(0)] = convert_half4(VEXPR); \
|
||||
} \
|
||||
kernel void kernel_##NAME##_f32_nc( \
|
||||
global const char * src0, ulong offset0, \
|
||||
global char * dst, ulong offsetd, \
|
||||
int ne00, ulong nb00, ulong nb01, ulong nb02, ulong nb03, \
|
||||
ulong nb0, ulong nb1, ulong nb2, ulong nb3) { \
|
||||
src0 = src0 + offset0; dst = dst + offsetd; \
|
||||
const int i3 = get_group_id(2); \
|
||||
const int i2 = get_group_id(1); \
|
||||
const int i1 = get_group_id(0); \
|
||||
for (int i0 = get_local_id(0); i0 < ne00; i0 += get_local_size(0)) { \
|
||||
float x = *(global const float *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00); \
|
||||
*(global float *)(dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0) = (SEXPR); \
|
||||
} \
|
||||
} \
|
||||
kernel void kernel_##NAME##_f16_nc( \
|
||||
global const char * src0, ulong offset0, \
|
||||
global char * dst, ulong offsetd, \
|
||||
int ne00, ulong nb00, ulong nb01, ulong nb02, ulong nb03, \
|
||||
ulong nb0, ulong nb1, ulong nb2, ulong nb3) { \
|
||||
src0 = src0 + offset0; dst = dst + offsetd; \
|
||||
const int i3 = get_group_id(2); \
|
||||
const int i2 = get_group_id(1); \
|
||||
const int i1 = get_group_id(0); \
|
||||
for (int i0 = get_local_id(0); i0 < ne00; i0 += get_local_size(0)) {\
|
||||
float x = *(global const half *)(src0 + i3*nb03 + i2*nb02 + i1*nb01 + i0*nb00); \
|
||||
*(global half *)(dst + i3*nb3 + i2*nb2 + i1*nb1 + i0*nb0) = (SEXPR); \
|
||||
} \
|
||||
}
|
||||
|
||||
UNARY_EXT(sgn, sign(x), sign(x))
|
||||
UNARY_EXT(step, x > 0.0f ? 1.0f : 0.0f, select((float4)0.0f, (float4)1.0f, x > 0.0f))
|
||||
UNARY_EXT(elu, x > 0.0f ? x : expm1(x), select(expm1(x), x, x > 0.0f))
|
||||
UNARY_EXT(hardswish, x * fmin(1.0f, fmax(0.0f, (x + 3.0f) / 6.0f)), x * fmin((float4)1.0f, fmax((float4)0.0f, (x + 3.0f) / 6.0f)))
|
||||
UNARY_EXT(hardsigmoid, fmin(1.0f, fmax(0.0f, (x + 3.0f) / 6.0f)), fmin((float4)1.0f, fmax((float4)0.0f, (x + 3.0f) / 6.0f)))
|
||||
UNARY_EXT(floor, floor(x), floor(x))
|
||||
UNARY_EXT(ceil, ceil(x), ceil(x))
|
||||
UNARY_EXT(round, round(x), round(x))
|
||||
UNARY_EXT(trunc, trunc(x), trunc(x))
|
||||
@@ -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());
|
||||
|
||||
@@ -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={});
|
||||
|
||||
@@ -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)));
|
||||
}
|
||||
}
|
||||
|
||||
@@ -1,6 +1,50 @@
|
||||
#include "fwht.hpp"
|
||||
|
||||
#include <cmath>
|
||||
#define P 1.0f
|
||||
#define N -1.0f
|
||||
|
||||
// constant Hadamard matrix via Paley I construction
|
||||
static constexpr float H12[12][12] = {
|
||||
{ P, P, P, P, P, P, P, P, P, P, P, P },
|
||||
{ P, N, P, N, P, P, P, N, N, N, P, N },
|
||||
{ P, N, N, P, N, P, P, P, N, N, N, P },
|
||||
{ P, P, N, N, P, N, P, P, P, N, N, N },
|
||||
{ P, N, P, N, N, P, N, P, P, P, N, N },
|
||||
{ P, N, N, P, N, N, P, N, P, P, P, N },
|
||||
{ P, N, N, N, P, N, N, P, N, P, P, P },
|
||||
{ P, P, N, N, N, P, N, N, P, N, P, P },
|
||||
{ P, P, P, N, N, N, P, N, N, P, N, P },
|
||||
{ P, P, P, P, N, N, N, P, N, N, P, N },
|
||||
{ P, N, P, P, P, N, N, N, P, N, N, P },
|
||||
{ P, P, N, P, P, P, N, N, N, P, N, N }
|
||||
};
|
||||
|
||||
static constexpr float H20[20][20] = {
|
||||
{ P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P },
|
||||
{ P, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N },
|
||||
{ P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P },
|
||||
{ P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P },
|
||||
{ P, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N },
|
||||
{ P, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N },
|
||||
{ P, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N },
|
||||
{ P, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N },
|
||||
{ P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P },
|
||||
{ P, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N },
|
||||
{ P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P },
|
||||
{ P, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N },
|
||||
{ P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P },
|
||||
{ P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P },
|
||||
{ P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P },
|
||||
{ P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P },
|
||||
{ P, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N },
|
||||
{ P, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N },
|
||||
{ P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P },
|
||||
{ P, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N }
|
||||
};
|
||||
|
||||
#undef P
|
||||
#undef N
|
||||
|
||||
template <int N>
|
||||
static void fwht_kernel(const float * __restrict__ src, float * __restrict__ dst, const int64_t n_rows,
|
||||
@@ -80,6 +124,122 @@ static void launch_fwht(const float * src, float * dst, const int64_t n_rows, co
|
||||
});
|
||||
}
|
||||
|
||||
template <int N, int m>
|
||||
static void kronecker_kernel(const float * __restrict__ src,
|
||||
float * __restrict__ dst,
|
||||
const int64_t n_rows,
|
||||
const float scale,
|
||||
const sycl::nd_item<2> & item) {
|
||||
static_assert(m == 12 || m == 20, "block size has to be 12 or 20.");
|
||||
|
||||
const sycl::sub_group sg = item.get_sub_group();
|
||||
|
||||
const int64_t r = item.get_global_id(0);
|
||||
if (r >= n_rows) {
|
||||
return;
|
||||
}
|
||||
|
||||
src += r * N;
|
||||
dst += r * N;
|
||||
|
||||
constexpr int blocks_per_group = N / m;
|
||||
constexpr int el_w = blocks_per_group / WARP_SIZE;
|
||||
static_assert(el_w >= 1 && blocks_per_group % WARP_SIZE == 0, "blocks_per_group must be a multiple of WARP_SIZE");
|
||||
float reg[el_w * m];
|
||||
const int lane = sg.get_local_linear_id();
|
||||
|
||||
#pragma unroll
|
||||
for (int i = 0; i < el_w; ++i) {
|
||||
const int b_idx = i * WARP_SIZE + lane;
|
||||
|
||||
#pragma unroll
|
||||
for (int j = 0; j < m; ++j) {
|
||||
reg[i * m + j] = src[b_idx * m + j] * scale;
|
||||
}
|
||||
}
|
||||
|
||||
#pragma unroll
|
||||
for (int b = 0; b < el_w; ++b) {
|
||||
float z[m] = { 0.0f };
|
||||
|
||||
#pragma unroll
|
||||
for (int i = 0; i < m; ++i) {
|
||||
#pragma unroll
|
||||
for (int j = 0; j < m; ++j) {
|
||||
const float h = (m == 12 ? H12[j][i] : H20[j][i]);
|
||||
z[i] += reg[b * m + j] * h;
|
||||
}
|
||||
}
|
||||
|
||||
#pragma unroll
|
||||
for (int i = 0; i < m; ++i) {
|
||||
reg[b * m + i] = z[i];
|
||||
}
|
||||
}
|
||||
|
||||
#pragma unroll
|
||||
for (int h = 1; h < WARP_SIZE; h *= 2) {
|
||||
#pragma unroll
|
||||
for (int j = 0; j < el_w; ++j) {
|
||||
#pragma unroll
|
||||
for (int k = 0; k < m; ++k) {
|
||||
const float val = reg[j * m + k];
|
||||
const float val2 = dpct::permute_sub_group_by_xor(sg, val, h, WARP_SIZE);
|
||||
|
||||
reg[j * m + k] = (lane & h) == 0 ? val + val2 : val2 - val;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
#pragma unroll
|
||||
for (int h = WARP_SIZE; h < blocks_per_group; h *= 2) {
|
||||
const int step = h / WARP_SIZE;
|
||||
#pragma unroll
|
||||
for (int j = 0; j < el_w; j += 2 * step) {
|
||||
#pragma unroll
|
||||
for (int s = 0; s < step; ++s) {
|
||||
#pragma unroll
|
||||
for (int k = 0; k < m; ++k) {
|
||||
const float x = reg[(j + s) * m + k];
|
||||
const float y = reg[(j + s + step) * m + k];
|
||||
|
||||
reg[(j + s) * m + k] = x + y;
|
||||
reg[(j + s + step) * m + k] = x - y;
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
#pragma unroll
|
||||
for (int i = 0; i < el_w; ++i) {
|
||||
const int b_idx = i * WARP_SIZE + lane;
|
||||
#pragma unroll
|
||||
for (int k = 0; k < m; ++k) {
|
||||
dst[b_idx * m + k] = reg[i * m + k];
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
template <int N, int m>
|
||||
static void launch_kronecker(const float * src,
|
||||
float * dst,
|
||||
const int64_t n_rows,
|
||||
const float scale,
|
||||
dpct::queue_ptr stream) {
|
||||
constexpr int rows_per_block = 4;
|
||||
|
||||
const int64_t num_blocks = (n_rows + rows_per_block - 1) / rows_per_block;
|
||||
|
||||
// dim 1 is the fastest-varying, so a sub-group is exactly one row's WARP_SIZE lanes.
|
||||
const sycl::range<2> global(num_blocks * rows_per_block, WARP_SIZE);
|
||||
const sycl::range<2> local(rows_per_block, WARP_SIZE);
|
||||
|
||||
stream->parallel_for(sycl::nd_range<2>(global, local),
|
||||
[=](sycl::nd_item<2> item) [[sycl::reqd_sub_group_size(WARP_SIZE)]] {
|
||||
kronecker_kernel<N, m>(src, dst, n_rows, scale, item);
|
||||
});
|
||||
}
|
||||
|
||||
bool ggml_sycl_op_fwht(ggml_backend_sycl_context & ctx, const ggml_tensor * src, ggml_tensor * dst) {
|
||||
if (src->type != GGML_TYPE_F32 || dst->type != GGML_TYPE_F32) {
|
||||
return false;
|
||||
@@ -113,6 +273,18 @@ bool ggml_sycl_op_fwht(ggml_backend_sycl_context & ctx, const ggml_tensor * src,
|
||||
case 512:
|
||||
launch_fwht<512>(src_d, dst_d, rows, scale, stream);
|
||||
return true;
|
||||
case 384:
|
||||
launch_kronecker<384, 12>(src_d, dst_d, rows, scale, stream);
|
||||
return true;
|
||||
case 768:
|
||||
launch_kronecker<768, 12>(src_d, dst_d, rows, scale, stream);
|
||||
return true;
|
||||
case 640:
|
||||
launch_kronecker<640, 20>(src_d, dst_d, rows, scale, stream);
|
||||
return true;
|
||||
case 1280:
|
||||
launch_kronecker<1280, 20>(src_d, dst_d, rows, scale, stream);
|
||||
return true;
|
||||
default:
|
||||
return false;
|
||||
}
|
||||
|
||||
@@ -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) {
|
||||
|
||||
@@ -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();
|
||||
}
|
||||
@@ -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
|
||||
+115
-14
@@ -4742,6 +4742,51 @@ struct test_mul_mat : public test_case {
|
||||
}
|
||||
};
|
||||
|
||||
#define P 1.0f
|
||||
#define N -1.0f
|
||||
|
||||
// constant Hadamard matrix via Paley I construction
|
||||
static constexpr float H12[12][12] = {
|
||||
{ P, P, P, P, P, P, P, P, P, P, P, P },
|
||||
{ P, N, P, N, P, P, P, N, N, N, P, N },
|
||||
{ P, N, N, P, N, P, P, P, N, N, N, P },
|
||||
{ P, P, N, N, P, N, P, P, P, N, N, N },
|
||||
{ P, N, P, N, N, P, N, P, P, P, N, N },
|
||||
{ P, N, N, P, N, N, P, N, P, P, P, N },
|
||||
{ P, N, N, N, P, N, N, P, N, P, P, P },
|
||||
{ P, P, N, N, N, P, N, N, P, N, P, P },
|
||||
{ P, P, P, N, N, N, P, N, N, P, N, P },
|
||||
{ P, P, P, P, N, N, N, P, N, N, P, N },
|
||||
{ P, N, P, P, P, N, N, N, P, N, N, P },
|
||||
{ P, P, N, P, P, P, N, N, N, P, N, N }
|
||||
};
|
||||
|
||||
static constexpr float H20[20][20] = {
|
||||
{ P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P, P },
|
||||
{ P, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N },
|
||||
{ P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P },
|
||||
{ P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P },
|
||||
{ P, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N },
|
||||
{ P, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N, N },
|
||||
{ P, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N, N },
|
||||
{ P, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P, N },
|
||||
{ P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N, P },
|
||||
{ P, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P, N },
|
||||
{ P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N, P },
|
||||
{ P, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P, N },
|
||||
{ P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P, P },
|
||||
{ P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P, P },
|
||||
{ P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P, P },
|
||||
{ P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N, P },
|
||||
{ P, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N, N },
|
||||
{ P, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P, N },
|
||||
{ P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N, P },
|
||||
{ P, P, N, N, P, P, P, P, N, P, N, P, N, N, N, N, P, P, N, N }
|
||||
};
|
||||
|
||||
#undef P
|
||||
#undef N
|
||||
|
||||
// GGML_HINT_SRC0_IS_HADAMARD
|
||||
struct test_mul_mat_hadamard : public test_mul_mat {
|
||||
test_mul_mat_hadamard(ggml_type type_a = GGML_TYPE_F32, ggml_type type_b = GGML_TYPE_F32,
|
||||
@@ -4766,20 +4811,58 @@ struct test_mul_mat_hadamard : public test_mul_mat {
|
||||
void initialize_tensors(ggml_context * ctx) override {
|
||||
for (ggml_tensor * t = ggml_get_first_tensor(ctx); t != NULL; t = ggml_get_next_tensor(ctx, t)) {
|
||||
if (strcmp(t->name, "a") == 0) {
|
||||
const int64_t n_cols = t->ne[0];
|
||||
const int64_t n_rows = ggml_nrows(t);
|
||||
const int64_t n_cols = t->ne[0];
|
||||
const int64_t n_rows = ggml_nrows(t);
|
||||
std::vector<float> data(n_cols * n_rows);
|
||||
float scale = 1.0f / sqrtf((float)n_cols);
|
||||
for (int64_t r = 0; r < n_rows; r++) {
|
||||
float * row_data = data.data() + r * n_cols;
|
||||
for (int64_t i = 0; i < n_cols; i++) {
|
||||
int pop = 0;
|
||||
int64_t val = r & i;
|
||||
while (val) {
|
||||
pop += (val & 1);
|
||||
val >>= 1;
|
||||
float scale = 1.0f / sqrtf((float) n_cols);
|
||||
|
||||
auto is_pow2 = [](const int64_t a) {
|
||||
return (a > 0) && ((a & (a - 1)) == 0);
|
||||
};
|
||||
#ifdef GGML_USE_SYCL
|
||||
const bool is_kronecker =
|
||||
((n_cols % 12 == 0) && is_pow2(n_cols / 12)) || ((n_cols % 20 == 0) && is_pow2(n_cols / 20));
|
||||
#else
|
||||
const bool is_kronecker = false;
|
||||
#endif
|
||||
if (is_kronecker) {
|
||||
const int64_t B = (n_cols % 12 == 0 && is_pow2(n_cols / 12)) ? 12 : 20;
|
||||
for (int64_t r = 0; r < n_rows; r++) {
|
||||
float * row_data = data.data() + r * n_cols;
|
||||
const int64_t r_mod = r % n_cols;
|
||||
const int64_t r_b = r_mod / B;
|
||||
const int64_t r_m = r_mod % B;
|
||||
|
||||
for (int64_t i = 0; i < n_cols; i++) {
|
||||
const int64_t c_b = i / B;
|
||||
const int64_t c_m = i % B;
|
||||
|
||||
int pop = 0;
|
||||
int64_t val = r_b & c_b;
|
||||
while (val) {
|
||||
pop += (val & 1);
|
||||
val >>= 1;
|
||||
}
|
||||
const float sign_m = (pop % 2 == 0) ? 1.0f : -1.0f;
|
||||
const float sign_b = (B == 12) ? H12[c_m][r_m] : H20[c_m][r_m];
|
||||
|
||||
row_data[i] = scale * sign_b * sign_m;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
else if (is_pow2(n_cols)) {
|
||||
for (int64_t r = 0; r < n_rows; r++) {
|
||||
float * row_data = data.data() + r * n_cols;
|
||||
for (int64_t i = 0; i < n_cols; i++) {
|
||||
int pop_cnt = 0;
|
||||
int64_t val = r & i;
|
||||
while (val) {
|
||||
pop_cnt += (val & 1);
|
||||
val >>= 1;
|
||||
}
|
||||
row_data[i] = (pop_cnt % 2 == 0) ? scale : -scale;
|
||||
}
|
||||
row_data[i] = (pop % 2 == 0) ? scale : -scale;
|
||||
}
|
||||
}
|
||||
ggml_backend_tensor_set(t, data.data(), 0, data.size() * sizeof(float));
|
||||
@@ -9469,7 +9552,16 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 256, 512, 256)); // many rows
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 32, 1, 32)); // too small (N<64)
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 1024, 1, 1024)); // too big (N>512)
|
||||
|
||||
#ifdef GGML_USE_SYCL
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 1, 384)); // m=12 (N=384)
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 32, 384)); // m=12 (batch)
|
||||
test_cases.emplace_back(
|
||||
new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 4, 384, { 2, 3 })); // m=12 (multi-dim)
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 768, 1, 768)); // m=12 (N=768)
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 640, 1, 640)); // m=20 (N=640)
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 640, 32, 640)); // m=20 (batch)
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 1280, 1, 1280)); // m=20 (N=1280)
|
||||
#endif
|
||||
#if 0
|
||||
// > 4GB A matrix. Too slow to be enabled by default.
|
||||
test_cases.emplace_back(new test_mul_mat(GGML_TYPE_F16, GGML_TYPE_F16, 900000, 3, 2592, {1, 1}, {1, 1}));
|
||||
@@ -10739,7 +10831,16 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_perf() {
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 128, 2048, 128));
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 256, 2048, 256));
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 512, 2048, 512));
|
||||
|
||||
#ifdef GGML_USE_SYCL
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 1, 384)); // m=12 (N=384)
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 32, 384)); // m=12 (batch)
|
||||
test_cases.emplace_back(
|
||||
new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 384, 4, 384, { 2, 3 })); // m=12 (multi-dim)
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 768, 1, 768)); // m=12 (N=768)
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 640, 1, 640)); // m=20 (N=640)
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 640, 32, 640)); // m=20 (batch)
|
||||
test_cases.emplace_back(new test_mul_mat_hadamard(GGML_TYPE_F32, GGML_TYPE_F32, 1280, 1, 1280)); // m=20 (N=1280)
|
||||
#endif
|
||||
test_cases.emplace_back(new test_solve_tri(GGML_TYPE_F32, { 64, 64, 4, 4 }, { 32, 64, 4, 4 }));
|
||||
test_cases.emplace_back(new test_solve_tri(GGML_TYPE_F32, { 128, 128, 4, 2 }, { 32, 128, 4, 2 }));
|
||||
// qwen3next with CHUNK_SIZE 64
|
||||
|
||||
+2
-1
@@ -90,6 +90,7 @@
|
||||
| `-hft, --hf-token TOKEN` | Hugging Face access token (default: value from HF_TOKEN environment variable)<br/>(env: HF_TOKEN) |
|
||||
| `--log-disable` | Log disable |
|
||||
| `--log-file FNAME` | Log to file<br/>(env: LLAMA_ARG_LOG_FILE) |
|
||||
| `--log-jsonl, --no-log-jsonl` | Log as JSONL (one JSON object per line) to stdout, this also disables colored logging (default: disabled)<br/>(env: LLAMA_ARG_LOG_JSONL) |
|
||||
| `--log-colors [on\|off\|auto]` | Set colored logging ('on', 'off', or 'auto', default: 'auto')<br/>'auto' enables colors when output is to a terminal<br/>(env: LLAMA_ARG_LOG_COLORS) |
|
||||
| `-v, --verbose, --log-verbose` | Set verbosity level to infinity (i.e. log all messages, useful for debugging) |
|
||||
| `--offline` | Offline mode: forces use of cache, prevents network access<br/>(env: LLAMA_ARG_OFFLINE) |
|
||||
@@ -178,7 +179,7 @@
|
||||
| `--reasoning-effort LEVEL` | reasoning effort level given to the chat template: 'default' to keep the template default,<br/>or a level such as 'minimal', 'low', 'medium', 'high', 'xhigh' or 'max' (default: default)<br/>(env: LLAMA_ARG_REASONING_EFFORT) |
|
||||
| `--reasoning-budget N` | token budget for thinking: -1 for unrestricted, 0 for immediate end, N>0 for token budget (default: -1)<br/>(env: LLAMA_ARG_THINK_BUDGET) |
|
||||
| `--reasoning-budget-message MESSAGE` | message injected before the end-of-thinking tag when reasoning budget is exhausted (default: none)<br/>(env: LLAMA_ARG_THINK_BUDGET_MESSAGE) |
|
||||
| `--reasoning-preserve, --no-reasoning-preserve` | preserve reasoning trace in the full history, not just the last assistant message (default: template default)<br/>compatible with certain templates having 'supports_preserve_reasoning' capability<br/>example: https://docs.z.ai/guides/capabilities/thinking-mode#preserved-thinking<br/>(env: LLAMA_ARG_REASONING_PRESERVE) |
|
||||
| `--reasoning-preserve, --no-reasoning-preserve` | preserve reasoning trace in the full history, not just the last assistant message (default: enabled)<br/>compatible with certain templates having 'supports_preserve_reasoning' capability<br/>example: https://docs.z.ai/guides/capabilities/thinking-mode#preserved-thinking<br/>(env: LLAMA_ARG_REASONING_PRESERVE) |
|
||||
| `--chat-template JINJA_TEMPLATE` | set custom jinja chat template (default: template taken from model's metadata)<br/>if suffix/prefix are specified, template will be disabled<br/>only commonly used templates are accepted (unless --jinja is set before this flag):<br/>list of built-in templates:<br/>bailing, bailing-think, bailing2, chatglm3, chatglm4, chatml, command-r, deepseek, deepseek-ocr, deepseek2, deepseek3, exaone-moe, exaone3, exaone4, falcon3, gemma, gigachat, glmedge, gpt-oss, granite, granite-4.0, granite-4.1, grok-2, hunyuan-dense, hunyuan-moe, hunyuan-vl, kimi-k2, llama2, llama2-sys, llama2-sys-bos, llama2-sys-strip, llama3, llama4, megrez, minicpm, mistral-v1, mistral-v3, mistral-v3-tekken, mistral-v7, mistral-v7-tekken, monarch, openchat, orion, pangu-embedded, phi3, phi4, rwkv-world, seed_oss, smolvlm, solar-open, vicuna, vicuna-orca, yandex, zephyr<br/>(env: LLAMA_ARG_CHAT_TEMPLATE) |
|
||||
| `--chat-template-file JINJA_TEMPLATE_FILE` | set custom jinja chat template file (default: template taken from model's metadata)<br/>if suffix/prefix are specified, template will be disabled<br/>only commonly used templates are accepted (unless --jinja is set before this flag):<br/>list of built-in templates:<br/>bailing, bailing-think, bailing2, chatglm3, chatglm4, chatml, command-r, deepseek, deepseek-ocr, deepseek2, deepseek3, exaone-moe, exaone3, exaone4, falcon3, gemma, gigachat, glmedge, gpt-oss, granite, granite-4.0, granite-4.1, grok-2, hunyuan-dense, hunyuan-moe, hunyuan-vl, kimi-k2, llama2, llama2-sys, llama2-sys-bos, llama2-sys-strip, llama3, llama4, megrez, minicpm, mistral-v1, mistral-v3, mistral-v3-tekken, mistral-v7, mistral-v7-tekken, monarch, openchat, orion, pangu-embedded, phi3, phi4, rwkv-world, seed_oss, smolvlm, solar-open, vicuna, vicuna-orca, yandex, zephyr<br/>(env: LLAMA_ARG_CHAT_TEMPLATE_FILE) |
|
||||
| `--skip-chat-parsing, --no-skip-chat-parsing` | force a pure content parser, even if a Jinja template is specified; model will output everything in the content section, including any reasoning and/or tool calls (default: disabled)<br/>(env: LLAMA_ARG_SKIP_CHAT_PARSING) |
|
||||
|
||||
@@ -173,6 +173,7 @@ llama-completion.exe -m models\gemma-1.1-7b-it.Q4_K_M.gguf --ignore-eos -n -1
|
||||
| `-hft, --hf-token TOKEN` | Hugging Face access token (default: value from HF_TOKEN environment variable)<br/>(env: HF_TOKEN) |
|
||||
| `--log-disable` | Log disable |
|
||||
| `--log-file FNAME` | Log to file<br/>(env: LLAMA_ARG_LOG_FILE) |
|
||||
| `--log-jsonl, --no-log-jsonl` | Log as JSONL (one JSON object per line) to stdout, this also disables colored logging (default: disabled)<br/>(env: LLAMA_ARG_LOG_JSONL) |
|
||||
| `--log-colors [on\|off\|auto]` | Set colored logging ('on', 'off', or 'auto', default: 'auto')<br/>'auto' enables colors when output is to a terminal<br/>(env: LLAMA_ARG_LOG_COLORS) |
|
||||
| `-v, --verbose, --log-verbose` | Set verbosity level to infinity (i.e. log all messages, useful for debugging) |
|
||||
| `--offline` | Offline mode: forces use of cache, prevents network access<br/>(env: LLAMA_ARG_OFFLINE) |
|
||||
@@ -256,7 +257,7 @@ llama-completion.exe -m models\gemma-1.1-7b-it.Q4_K_M.gguf --ignore-eos -n -1
|
||||
| `--reasoning-effort LEVEL` | reasoning effort level given to the chat template: 'default' to keep the template default,<br/>or a level such as 'minimal', 'low', 'medium', 'high', 'xhigh' or 'max' (default: default)<br/>(env: LLAMA_ARG_REASONING_EFFORT) |
|
||||
| `--reasoning-budget N` | token budget for thinking: -1 for unrestricted, 0 for immediate end, N>0 for token budget (default: -1)<br/>(env: LLAMA_ARG_THINK_BUDGET) |
|
||||
| `--reasoning-budget-message MESSAGE` | message injected before the end-of-thinking tag when reasoning budget is exhausted (default: none)<br/>(env: LLAMA_ARG_THINK_BUDGET_MESSAGE) |
|
||||
| `--reasoning-preserve, --no-reasoning-preserve` | preserve reasoning trace in the full history, not just the last assistant message (default: template default)<br/>compatible with certain templates having 'supports_preserve_reasoning' capability<br/>example: https://docs.z.ai/guides/capabilities/thinking-mode#preserved-thinking<br/>(env: LLAMA_ARG_REASONING_PRESERVE) |
|
||||
| `--reasoning-preserve, --no-reasoning-preserve` | preserve reasoning trace in the full history, not just the last assistant message (default: enabled)<br/>compatible with certain templates having 'supports_preserve_reasoning' capability<br/>example: https://docs.z.ai/guides/capabilities/thinking-mode#preserved-thinking<br/>(env: LLAMA_ARG_REASONING_PRESERVE) |
|
||||
| `--chat-template JINJA_TEMPLATE` | set custom jinja chat template (default: template taken from model's metadata)<br/>if suffix/prefix are specified, template will be disabled<br/>only commonly used templates are accepted (unless --jinja is set before this flag):<br/>list of built-in templates:<br/>bailing, bailing-think, bailing2, chatglm3, chatglm4, chatml, command-r, deepseek, deepseek-ocr, deepseek2, deepseek3, exaone-moe, exaone3, exaone4, falcon3, gemma, gigachat, glmedge, gpt-oss, granite, granite-4.0, granite-4.1, grok-2, hunyuan-dense, hunyuan-moe, hunyuan-vl, kimi-k2, llama2, llama2-sys, llama2-sys-bos, llama2-sys-strip, llama3, llama4, megrez, minicpm, mistral-v1, mistral-v3, mistral-v3-tekken, mistral-v7, mistral-v7-tekken, monarch, openchat, orion, pangu-embedded, phi3, phi4, rwkv-world, seed_oss, smolvlm, solar-open, vicuna, vicuna-orca, yandex, zephyr<br/>(env: LLAMA_ARG_CHAT_TEMPLATE) |
|
||||
| `--chat-template-file JINJA_TEMPLATE_FILE` | set custom jinja chat template file (default: template taken from model's metadata)<br/>if suffix/prefix are specified, template will be disabled<br/>only commonly used templates are accepted (unless --jinja is set before this flag):<br/>list of built-in templates:<br/>bailing, bailing-think, bailing2, chatglm3, chatglm4, chatml, command-r, deepseek, deepseek-ocr, deepseek2, deepseek3, exaone-moe, exaone3, exaone4, falcon3, gemma, gigachat, glmedge, gpt-oss, granite, granite-4.0, granite-4.1, grok-2, hunyuan-dense, hunyuan-moe, hunyuan-vl, kimi-k2, llama2, llama2-sys, llama2-sys-bos, llama2-sys-strip, llama3, llama4, megrez, minicpm, mistral-v1, mistral-v3, mistral-v3-tekken, mistral-v7, mistral-v7-tekken, monarch, openchat, orion, pangu-embedded, phi3, phi4, rwkv-world, seed_oss, smolvlm, solar-open, vicuna, vicuna-orca, yandex, zephyr<br/>(env: LLAMA_ARG_CHAT_TEMPLATE_FILE) |
|
||||
| `--skip-chat-parsing, --no-skip-chat-parsing` | force a pure content parser, even if a Jinja template is specified; model will output everything in the content section, including any reasoning and/or tool calls (default: disabled)<br/>(env: LLAMA_ARG_SKIP_CHAT_PARSING) |
|
||||
|
||||
@@ -107,6 +107,7 @@ For the full list of features, please refer to [server's changelog](https://gith
|
||||
| `-hft, --hf-token TOKEN` | Hugging Face access token (default: value from HF_TOKEN environment variable)<br/>(env: HF_TOKEN) |
|
||||
| `--log-disable` | Log disable |
|
||||
| `--log-file FNAME` | Log to file<br/>(env: LLAMA_ARG_LOG_FILE) |
|
||||
| `--log-jsonl, --no-log-jsonl` | Log as JSONL (one JSON object per line) to stdout, this also disables colored logging (default: disabled)<br/>(env: LLAMA_ARG_LOG_JSONL) |
|
||||
| `--log-colors [on\|off\|auto]` | Set colored logging ('on', 'off', or 'auto', default: 'auto')<br/>'auto' enables colors when output is to a terminal<br/>(env: LLAMA_ARG_LOG_COLORS) |
|
||||
| `-v, --verbose, --log-verbose` | Set verbosity level to infinity (i.e. log all messages, useful for debugging) |
|
||||
| `--offline` | Offline mode: forces use of cache, prevents network access<br/>(env: LLAMA_ARG_OFFLINE) |
|
||||
@@ -236,7 +237,7 @@ For the full list of features, please refer to [server's changelog](https://gith
|
||||
| `--reasoning-effort LEVEL` | reasoning effort level given to the chat template: 'default' to keep the template default,<br/>or a level such as 'minimal', 'low', 'medium', 'high', 'xhigh' or 'max' (default: default)<br/>(env: LLAMA_ARG_REASONING_EFFORT) |
|
||||
| `--reasoning-budget N` | token budget for thinking: -1 for unrestricted, 0 for immediate end, N>0 for token budget (default: -1)<br/>(env: LLAMA_ARG_THINK_BUDGET) |
|
||||
| `--reasoning-budget-message MESSAGE` | message injected before the end-of-thinking tag when reasoning budget is exhausted (default: none)<br/>(env: LLAMA_ARG_THINK_BUDGET_MESSAGE) |
|
||||
| `--reasoning-preserve, --no-reasoning-preserve` | preserve reasoning trace in the full history, not just the last assistant message (default: template default)<br/>compatible with certain templates having 'supports_preserve_reasoning' capability<br/>example: https://docs.z.ai/guides/capabilities/thinking-mode#preserved-thinking<br/>(env: LLAMA_ARG_REASONING_PRESERVE) |
|
||||
| `--reasoning-preserve, --no-reasoning-preserve` | preserve reasoning trace in the full history, not just the last assistant message (default: enabled)<br/>compatible with certain templates having 'supports_preserve_reasoning' capability<br/>example: https://docs.z.ai/guides/capabilities/thinking-mode#preserved-thinking<br/>(env: LLAMA_ARG_REASONING_PRESERVE) |
|
||||
| `--chat-template JINJA_TEMPLATE` | set custom jinja chat template (default: template taken from model's metadata)<br/>if suffix/prefix are specified, template will be disabled<br/>only commonly used templates are accepted (unless --jinja is set before this flag):<br/>list of built-in templates:<br/>bailing, bailing-think, bailing2, chatglm3, chatglm4, chatml, command-r, deepseek, deepseek-ocr, deepseek2, deepseek3, exaone-moe, exaone3, exaone4, falcon3, gemma, gigachat, glmedge, gpt-oss, granite, granite-4.0, granite-4.1, grok-2, hunyuan-dense, hunyuan-moe, hunyuan-vl, kimi-k2, llama2, llama2-sys, llama2-sys-bos, llama2-sys-strip, llama3, llama4, megrez, minicpm, mistral-v1, mistral-v3, mistral-v3-tekken, mistral-v7, mistral-v7-tekken, monarch, openchat, orion, pangu-embedded, phi3, phi4, rwkv-world, seed_oss, smolvlm, solar-open, vicuna, vicuna-orca, yandex, zephyr<br/>(env: LLAMA_ARG_CHAT_TEMPLATE) |
|
||||
| `--chat-template-file JINJA_TEMPLATE_FILE` | set custom jinja chat template file (default: template taken from model's metadata)<br/>if suffix/prefix are specified, template will be disabled<br/>only commonly used templates are accepted (unless --jinja is set before this flag):<br/>list of built-in templates:<br/>bailing, bailing-think, bailing2, chatglm3, chatglm4, chatml, command-r, deepseek, deepseek-ocr, deepseek2, deepseek3, exaone-moe, exaone3, exaone4, falcon3, gemma, gigachat, glmedge, gpt-oss, granite, granite-4.0, granite-4.1, grok-2, hunyuan-dense, hunyuan-moe, hunyuan-vl, kimi-k2, llama2, llama2-sys, llama2-sys-bos, llama2-sys-strip, llama3, llama4, megrez, minicpm, mistral-v1, mistral-v3, mistral-v3-tekken, mistral-v7, mistral-v7-tekken, monarch, openchat, orion, pangu-embedded, phi3, phi4, rwkv-world, seed_oss, smolvlm, solar-open, vicuna, vicuna-orca, yandex, zephyr<br/>(env: LLAMA_ARG_CHAT_TEMPLATE_FILE) |
|
||||
| `--skip-chat-parsing, --no-skip-chat-parsing` | force a pure content parser, even if a Jinja template is specified; model will output everything in the content section, including any reasoning and/or tool calls (default: disabled)<br/>(env: LLAMA_ARG_SKIP_CHAT_PARSING) |
|
||||
|
||||
@@ -194,7 +194,7 @@
|
||||
/>
|
||||
{:else if section.type === AgenticSectionType.TOOL_CALL || section.type === AgenticSectionType.TOOL_CALL_PENDING || section.type === AgenticSectionType.TOOL_CALL_STREAMING}
|
||||
<ChatMessageToolCallBlock
|
||||
attachments={message?.extra}
|
||||
attachments={section.toolResultExtras}
|
||||
isExecuting={section.toolCallId !== undefined &&
|
||||
section.toolCallId === currentlyExecutingToolCallId}
|
||||
{isStreaming}
|
||||
|
||||
@@ -2,8 +2,11 @@ import type { AgenticConfig } from '$lib/types/agentic';
|
||||
|
||||
export const ATTACHMENT_SAVED_REGEX = /\[Attachment saved: ([^\]]+)\]/;
|
||||
|
||||
// JSON detection: trimmed content opens with an object or array literal.
|
||||
export const TOOL_RESULT_JSON_OPEN_REGEX = /^[[{]/;
|
||||
// JSON detection: an attachment placeholder also starts with `[`, but is
|
||||
// plain text (`[Attachment saved: ...]`), not an array literal. Require the
|
||||
// first array value (or the closing bracket for an empty array) to look like
|
||||
// a valid JSON token before attempting JSON.parse.
|
||||
export const TOOL_RESULT_JSON_OPEN_REGEX = /^(?:\{|\[\s*(?:[[\]"{\-0-9]|true|false|null))/;
|
||||
|
||||
// Search-summary wire format used by file-glob and grep tools:
|
||||
// <matches>
|
||||
|
||||
@@ -37,6 +37,10 @@ describe('classifyToolResult', () => {
|
||||
expect(classifyToolResult('["a", "b", "c"]')).toBe('json');
|
||||
});
|
||||
|
||||
it('classifies a nested JSON array', () => {
|
||||
expect(classifyToolResult('[[1, 2], [3, 4]]')).toBe('json');
|
||||
});
|
||||
|
||||
it('classifies a pretty-printed JSON object', () => {
|
||||
expect(classifyToolResult('{\n "key": "value"\n}')).toBe('json');
|
||||
});
|
||||
|
||||
Reference in New Issue
Block a user