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

Merge branch 'upstream' into concedo_experimental

Browse Source 
# Conflicts:
#	.github/workflows/build.yml
#	ci/run.sh
#	ggml/src/ggml-metal/CMakeLists.txt
#	tools/mtmd/clip.cpp
#	tools/rpc/rpc-server.cpp
This commit is contained in:
Concedo
2025-09-14 10:43:59 +08:00
parent bf8fc4659b aa0c461efe
commit 994c13a167
10 changed files with 888 additions and 192 deletions
Download Patch File Download Diff File Expand all files Collapse all files
common/arg.cpp
+6 -15
View File
@@ -1306,7 +1306,7 @@ static std::vector<ggml_backend_dev_t> parse_device_list(const std::string & val
} else {
for (const auto & device : dev_names) {
auto * dev = ggml_backend_dev_by_name(device.c_str());
if (!dev || ggml_backend_dev_type(dev) != GGML_BACKEND_DEVICE_TYPE_GPU) {
if (!dev || ggml_backend_dev_type(dev) == GGML_BACKEND_DEVICE_TYPE_CPU) {
throw std::invalid_argument(string_format("invalid device: %s", device.c_str()));
}
devices.push_back(dev);
@@ -1316,7 +1316,7 @@ static std::vector<ggml_backend_dev_t> parse_device_list(const std::string & val
return devices;
}
static void add_rpc_devices(std::string servers) {
static void add_rpc_devices(const std::string & servers) {
auto rpc_servers = string_split<std::string>(servers, ',');
if (rpc_servers.empty()) {
throw std::invalid_argument("no RPC servers specified");
@@ -2518,24 +2518,15 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
{"--list-devices"},
"print list of available devices and exit",
[](common_params &) {
std::vector<ggml_backend_dev_t> rpc_devices;
std::vector<ggml_backend_dev_t> all_devices;
std::vector<ggml_backend_dev_t> devices;
for (size_t i = 0; i < ggml_backend_dev_count(); ++i) {
auto * dev = ggml_backend_dev_get(i);
if (ggml_backend_dev_type(dev) == GGML_BACKEND_DEVICE_TYPE_GPU) {
ggml_backend_reg_t reg = ggml_backend_dev_backend_reg(dev);
if (ggml_backend_reg_name(reg) == std::string("RPC")) {
rpc_devices.push_back(dev);
} else {
all_devices.push_back(dev);
}
if (ggml_backend_dev_type(dev) != GGML_BACKEND_DEVICE_TYPE_CPU) {
devices.push_back(dev);
}
}
// insert RPC devices in front
all_devices.insert(all_devices.begin(), rpc_devices.begin(), rpc_devices.end());
printf("Available devices:\n");
for (size_t i = 0; i < all_devices.size(); ++i) {
auto * dev = all_devices[i];
for (auto * dev : devices) {
size_t free, total;
ggml_backend_dev_memory(dev, &free, &total);
printf(" %s: %s (%zu MiB, %zu MiB free)\n", ggml_backend_dev_name(dev), ggml_backend_dev_description(dev), total / 1024 / 1024, free / 1024 / 1024);
ggml/src/ggml-metal/ggml-metal-common.cpp
+445
View File
@@ -0,0 +1,445 @@
#include "ggml-metal-common.h"
#include "ggml-impl.h"
#include <vector>
struct ggml_mem_range {
uint64_t pb; // buffer id
uint64_t p0; // begin
uint64_t p1; // end
ggml_mem_range_type pt;
};
struct ggml_mem_ranges {
std::vector<ggml_mem_range> ranges;
int debug = 0;
};
struct ggml_mem_ranges * ggml_mem_ranges_init(int debug) {
auto * res = new ggml_mem_ranges;
res->ranges.reserve(256);
res->debug = debug;
return res;
}
void ggml_mem_ranges_free(ggml_mem_ranges * mrs) {
delete mrs;
}
void ggml_mem_ranges_reset(ggml_mem_ranges * mrs) {
mrs->ranges.clear();
}
static bool ggml_mem_ranges_add(ggml_mem_ranges * mrs, ggml_mem_range mrp) {
mrs->ranges.push_back(mrp);
return true;
}
static ggml_mem_range ggml_mem_range_from_tensor(const ggml_tensor * tensor, ggml_mem_range_type pt) {
// always use the base tensor
tensor = tensor->view_src ? tensor->view_src : tensor;
GGML_ASSERT(!tensor->view_src);
ggml_mem_range mrp;
if (tensor->buffer) {
// when the tensor is allocated, use the actual memory address range of the buffer
mrp = {
/*.pb =*/ (uint64_t) tensor->buffer,
/*.p0 =*/ (uint64_t) tensor->data,
/*.p1 =*/ (uint64_t) tensor->data + ggml_nbytes(tensor),
/*.pt =*/ pt,
};
} else {
// otherwise, the tensor ptr is used as an unique id of the memory ranges
// that the tensor will be using when it is allocated
mrp = {
/*.pb =*/ (uint64_t) tensor,
/*.p0 =*/ 0, //
/*.p1 =*/ 1024, // [0, 1024) is a dummy range, not used
/*.pt =*/ pt,
};
};
return mrp;
}
static ggml_mem_range ggml_mem_range_from_tensor_src(const ggml_tensor * tensor) {
return ggml_mem_range_from_tensor(tensor, MEM_RANGE_TYPE_SRC);
}
static ggml_mem_range ggml_mem_range_from_tensor_dst(const ggml_tensor * tensor) {
return ggml_mem_range_from_tensor(tensor, MEM_RANGE_TYPE_DST);
}
static bool ggml_mem_ranges_add_src(ggml_mem_ranges * mrs, const ggml_tensor * tensor) {
GGML_ASSERT(tensor);
ggml_mem_range mrp = ggml_mem_range_from_tensor_src(tensor);
if (mrs->debug > 2) {
GGML_LOG_DEBUG("%s: add src range buf=%lld, [%lld, %lld)\n", __func__, mrp.pb, mrp.p0, mrp.p1);
}
return ggml_mem_ranges_add(mrs, mrp);
}
static bool ggml_mem_ranges_add_dst(ggml_mem_ranges * mrs, const ggml_tensor * tensor) {
GGML_ASSERT(tensor);
ggml_mem_range mrp = ggml_mem_range_from_tensor_dst(tensor);
if (mrs->debug > 2) {
GGML_LOG_DEBUG("%s: add dst range buf=%lld, [%lld, %lld)\n", __func__, mrp.pb, mrp.p0, mrp.p1);
}
return ggml_mem_ranges_add(mrs, mrp);
}
bool ggml_mem_ranges_add(ggml_mem_ranges * mrs, const ggml_tensor * tensor) {
for (int i = 0; i < GGML_MAX_DIMS; i++) {
if (tensor->src[i]) {
ggml_mem_ranges_add_src(mrs, tensor->src[i]);
}
}
return ggml_mem_ranges_add_dst(mrs, tensor);
}
static bool ggml_mem_ranges_check(const ggml_mem_ranges * mrs, ggml_mem_range mrp) {
for (size_t i = 0; i < mrs->ranges.size(); i++) {
const auto & cmp = mrs->ranges[i];
if (mrp.pb != cmp.pb) {
continue;
}
if (mrp.pt == MEM_RANGE_TYPE_SRC && cmp.pt == MEM_RANGE_TYPE_SRC) {
continue;
}
if (mrp.p0 < cmp.p1 && mrp.p1 >= cmp.p0) {
if (mrs->debug > 2) {
GGML_LOG_DEBUG("%s: the %s range buf=%lld, [%lld, %lld) overlaps with a previous %s range buf=%lld, [%lld, %lld)\n",
__func__,
mrp.pt == MEM_RANGE_TYPE_SRC ? "src" : "dst",
mrp.pb, mrp.p0, mrp.p1,
cmp.pt == MEM_RANGE_TYPE_SRC ? "src" : "dst",
cmp.pb, cmp.p0, cmp.p1);
}
return false;
}
}
return true;
}
static bool ggml_mem_ranges_check_src(const ggml_mem_ranges * mrs, const ggml_tensor * tensor) {
GGML_ASSERT(tensor);
ggml_mem_range mrp = ggml_mem_range_from_tensor_src(tensor);
const bool res = ggml_mem_ranges_check(mrs, mrp);
return res;
}
static bool ggml_mem_ranges_check_dst(const ggml_mem_ranges * mrs, const ggml_tensor * tensor) {
GGML_ASSERT(tensor);
ggml_mem_range mrp = ggml_mem_range_from_tensor_dst(tensor);
const bool res = ggml_mem_ranges_check(mrs, mrp);
return res;
}
bool ggml_mem_ranges_check(const ggml_mem_ranges * mrs, const ggml_tensor * tensor) {
for (int i = 0; i < GGML_MAX_DIMS; i++) {
if (tensor->src[i]) {
if (!ggml_mem_ranges_check_src(mrs, tensor->src[i])) {
return false;
}
}
}
return ggml_mem_ranges_check_dst(mrs, tensor);
}
// TODO: move to ggml.h?
static bool is_empty(ggml_op op) {
switch (op) {
case GGML_OP_NONE:
case GGML_OP_RESHAPE:
case GGML_OP_TRANSPOSE:
case GGML_OP_VIEW:
case GGML_OP_PERMUTE:
return true;
default:
return false;
}
}
struct node_info {
ggml_tensor * node;
std::vector<ggml_tensor *> fused;
ggml_op op() const {
return node->op;
}
const ggml_tensor * dst() const {
return fused.empty() ? node : fused.back();
}
bool is_empty() const {
return ::is_empty(node->op);
}
void add_fused(ggml_tensor * t) {
fused.push_back(t);
}
};
static std::vector<int> ggml_metal_graph_optimize_reorder(const std::vector<node_info> & nodes) {
// helper to add node src and dst ranges
const auto & h_add = [](ggml_mem_ranges * mrs, const node_info & node) {
for (int i = 0; i < GGML_MAX_SRC; i++) {
if (node.node->src[i]) {
if (!ggml_mem_ranges_add_src(mrs, node.node->src[i])) {
return false;
}
}
}
for (const auto * fused : node.fused) {
for (int i = 0; i < GGML_MAX_SRC; i++) {
if (fused->src[i]) {
if (!ggml_mem_ranges_add_src(mrs, fused->src[i])) {
return false;
}
}
}
}
return ggml_mem_ranges_add_dst(mrs, node.dst());
};
// helper to check if a node can run concurrently with the existing set of nodes
const auto & h_check = [](const ggml_mem_ranges * mrs, const node_info & node) {
for (int i = 0; i < GGML_MAX_SRC; i++) {
if (node.node->src[i]) {
if (!ggml_mem_ranges_check_src(mrs, node.node->src[i])) {
return false;
}
}
}
for (const auto * fused : node.fused) {
for (int i = 0; i < GGML_MAX_SRC; i++) {
if (fused->src[i]) {
if (!ggml_mem_ranges_check_src(mrs, fused->src[i])) {
return false;
}
}
}
}
return ggml_mem_ranges_check_dst(mrs, node.dst());
};
// perform reorders only across these types of ops
// can be expanded when needed
// IMPORTANT: do not add ops such as GGML_OP_CPY or GGML_OP_SET_ROWS
// the dependencies from such ops are not always represented in the graph
const auto & h_safe = [](ggml_op op) {
switch (op) {
case GGML_OP_MUL_MAT:
case GGML_OP_MUL_MAT_ID:
case GGML_OP_ROPE:
case GGML_OP_NORM:
case GGML_OP_RMS_NORM:
case GGML_OP_GROUP_NORM:
case GGML_OP_SUM_ROWS:
case GGML_OP_MUL:
case GGML_OP_ADD:
case GGML_OP_DIV:
case GGML_OP_GLU:
case GGML_OP_SCALE:
case GGML_OP_GET_ROWS:
return true;
default:
return is_empty(op);
}
};
const int n = nodes.size();
std::vector<int> res;
res.reserve(n);
std::vector<bool> used(n, false);
ggml_mem_ranges * mrs0 = ggml_mem_ranges_init(0);
ggml_mem_ranges * mrs1 = ggml_mem_ranges_init(0);
for (int i0 = 0; i0 < n; i0++) {
if (used[i0]) {
continue;
}
const auto & node0 = nodes[i0];
// the node is not concurrent with the existing concurrent set, so we have to "put a barrier" (i.e reset mrs0)
// but before we do that, look forward for some other nodes that can be added to the concurrent set mrs0
//
// note: we can always add empty nodes to the concurrent set as they don't read nor write anything
if (!node0.is_empty() && !h_check(mrs0, node0)) {
// this will hold the set of memory ranges from the nodes that haven't been processed yet
// if a node is not concurrent with this set, we cannot reorder it
ggml_mem_ranges_reset(mrs1);
// initialize it with the current node
h_add(mrs1, node0);
// that many nodes forward to search for a concurrent node
constexpr int N_FORWARD = 8;
for (int i1 = i0 + 1; i1 < i0 + N_FORWARD && i1 < n; i1++) {
if (used[i1]) {
continue;
}
const auto & node1 = nodes[i1];
// disallow reordering of certain ops
if (!h_safe(node1.op())) {
break;
}
const bool is_empty = node1.is_empty();
// to add a concurrent node, it has to be:
// + empty or concurrent with all nodes in the existing concurrent set (mrs0)
// + concurrent with all nodes prior to it that haven't been processed yet (mrs1)
if ((is_empty || h_check(mrs0, node1)) && h_check(mrs1, node1)) {
// add the node to the existing concurrent set (i.e. reorder it for early execution)
h_add(mrs0, node1);
res.push_back(i1);
// mark as used, so we skip re-processing it later
used[i1] = true;
} else {
// expand the set of nodes that haven't been processed yet
h_add(mrs1, node1);
}
}
// finalize the concurrent set and begin a new one
ggml_mem_ranges_reset(mrs0);
}
// expand the concurrent set with the current node
{
h_add(mrs0, node0);
res.push_back(i0);
}
}
ggml_mem_ranges_free(mrs0);
ggml_mem_ranges_free(mrs1);
return res;
}
void ggml_metal_graph_optimize(ggml_cgraph * gf) {
constexpr int MAX_FUSE = 16;
const int n = gf->n_nodes;
enum ggml_op ops[MAX_FUSE];
std::vector<node_info> nodes;
nodes.reserve(gf->n_nodes);
// fuse nodes:
// we don't want to make reorders that break fusing, so we first pack all fusable tensors
// and perform the reorder over the fused nodes. after the reorder is done, we unfuse
for (int i = 0; i < n; i++) {
node_info node = {
/*.node =*/ gf->nodes[i],
/*.fused =*/ {},
};
// fuse only ops that start with these operations
// can be expanded when needed
if (node.op() == GGML_OP_ADD ||
node.op() == GGML_OP_RMS_NORM) {
ops[0] = node.op();
int f = i + 1;
while (f < n && f < i + MAX_FUSE) {
// conservatively allow fusing only these ops
// can be expanded when needed
if (gf->nodes[f]->op != GGML_OP_ADD &&
gf->nodes[f]->op != GGML_OP_MUL &&
gf->nodes[f]->op != GGML_OP_RMS_NORM) {
break;
}
ops[f - i] = gf->nodes[f]->op;
f++;
}
f -= i;
for (; f > 1; f--) {
if (ggml_can_fuse(gf, i, ops, f)) {
break;
}
}
// add the fused tensors into the node info so we can unfuse them later
for (int k = 1; k < f; k++) {
++i;
// the .dst() becomes the last fused tensor
node.add_fused(gf->nodes[i]);
}
}
nodes.push_back(std::move(node));
}
// reorder to improve concurrency
#if 1
const auto order = ggml_metal_graph_optimize_reorder(nodes);
#else
std::vector<int> order(nodes.size());
for (size_t i = 0; i < nodes.size(); i++) {
order[i] = i;
}
#endif
// unfuse
{
int j = 0;
for (const auto i : order) {
const auto & node = nodes[i];
gf->nodes[j++] = node.node;
for (auto * fused : node.fused) {
gf->nodes[j++] = fused;
}
}
}
}
ggml/src/ggml-metal/ggml-metal-common.h
+52
View File
@@ -0,0 +1,52 @@
// helper functions for ggml-metal that are too difficult to implement in Objective-C
#pragma once
#include <stdbool.h>
#ifdef __cplusplus
extern "C" {
#endif
struct ggml_tensor;
struct ggml_cgraph;
enum ggml_mem_range_type {
MEM_RANGE_TYPE_SRC = 0,
MEM_RANGE_TYPE_DST = 1,
};
// a helper object that can be used for reordering operations to improve concurrency
//
// the fundamental idea is that a set of tasks (either ggml ops, or something else) can run concurrently if they
// don't write to a memory that is being read by another task or written to by another task in the set
//
// with this structure, we can add tasks to the set, setting memory constraints. we can also check if a new task
// can be added to the set without violating the constraints (i.e. if it can be executed concurrently with the
// tasks already in the set)
//
struct ggml_mem_ranges;
struct ggml_mem_ranges * ggml_mem_ranges_init(int debug);
void ggml_mem_ranges_free(struct ggml_mem_ranges * mrs);
// remove all ranges from the set
void ggml_mem_ranges_reset(struct ggml_mem_ranges * mrs);
// add src or dst ranges to track
bool ggml_mem_ranges_add(struct ggml_mem_ranges * mrs, const struct ggml_tensor * tensor);
// return false if:
// - new src range overlaps with any existing dst range
// - new dst range overlaps with any existing range (src or dst)
bool ggml_mem_ranges_check(const struct ggml_mem_ranges * mrs, const struct ggml_tensor * tensor);
// reorder the nodes in the graph to improve concurrency, while respecting fusion
//
// note: this implementation is generic and not specific to metal
// if it proves to work well, we can start using it for other backends in the future
void ggml_metal_graph_optimize(struct ggml_cgraph * gf);
#ifdef __cplusplus
}
#endif
ggml/src/ggml-metal/ggml-metal.m
+320 -153
View File
@@ -3,6 +3,7 @@
#import "ggml-impl.h"
#import "ggml-backend-impl.h"
#import "ggml-metal-impl.h"
#import "ggml-metal-common.h"
#import <Foundation/Foundation.h>
@@ -61,8 +62,11 @@ static struct ggml_backend_metal_device_context {
bool has_bfloat;
bool use_bfloat;
bool use_fusion;
bool use_concurrency;
bool use_shared_buffers;
bool use_graph_optimize;
int debug_graph;
int debug_fusion;
// how many times a given op was fused
@@ -83,7 +87,10 @@ static struct ggml_backend_metal_device_context {
/*.has_bfloat =*/ false,
/*.use_bfloat =*/ false,
/*.use_fusion =*/ true,
/*.use_concurrency =*/ true,
/*.use_shared_buffers =*/ true,
/*.use_graph_optimize =*/ true,
/*.debug_graph =*/ 0,
/*.debug_fusion =*/ 0,
/*.fuse_cnt =*/ { 0 },
/*.max_size =*/ 0,
@@ -124,7 +131,14 @@ static id<MTLDevice> ggml_backend_metal_device_acq(struct ggml_backend_metal_dev
#else
ctx->use_bfloat = false;
#endif
ctx->use_fusion = getenv("GGML_METAL_FUSION_DISABLE") == nil;
ctx->use_fusion = getenv("GGML_METAL_FUSION_DISABLE") == nil;
ctx->use_concurrency = getenv("GGML_METAL_CONCURRENCY_DISABLE") == nil;
{
const char * val = getenv("GGML_METAL_GRAPH_DEBUG");
ctx->debug_graph = val ? atoi(val) : 0;
}
{
const char * val = getenv("GGML_METAL_FUSION_DEBUG");
@@ -137,6 +151,12 @@ static id<MTLDevice> ggml_backend_metal_device_acq(struct ggml_backend_metal_dev
ctx->use_shared_buffers = false;
}
ctx->use_graph_optimize = true;
if (getenv("GGML_METAL_GRAPH_OPTIMIZE_DISABLE") != NULL) {
ctx->use_graph_optimize = false;
}
memset(ctx->fuse_cnt, 0, sizeof(ctx->fuse_cnt));
ctx->max_size = ctx->mtl_device.maxBufferLength;
@@ -212,28 +232,6 @@ struct ggml_metal_kernel {
@end
enum ggml_metal_kernel_type {
GGML_METAL_KERNEL_TYPE_ADD,
GGML_METAL_KERNEL_TYPE_ADD_FUSE_2,
GGML_METAL_KERNEL_TYPE_ADD_FUSE_3,
GGML_METAL_KERNEL_TYPE_ADD_FUSE_4,
GGML_METAL_KERNEL_TYPE_ADD_FUSE_5,
GGML_METAL_KERNEL_TYPE_ADD_FUSE_6,
GGML_METAL_KERNEL_TYPE_ADD_FUSE_7,
GGML_METAL_KERNEL_TYPE_ADD_FUSE_8,
GGML_METAL_KERNEL_TYPE_ADD_ROW_C4,
GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_2,
GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_3,
GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_4,
GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_5,
GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_6,
GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_7,
GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_8,
GGML_METAL_KERNEL_TYPE_SUB,
GGML_METAL_KERNEL_TYPE_SUB_ROW_C4,
GGML_METAL_KERNEL_TYPE_MUL,
GGML_METAL_KERNEL_TYPE_MUL_ROW_C4,
GGML_METAL_KERNEL_TYPE_DIV,
GGML_METAL_KERNEL_TYPE_DIV_ROW_C4,
GGML_METAL_KERNEL_TYPE_ADD_ID,
GGML_METAL_KERNEL_TYPE_REPEAT_F32,
GGML_METAL_KERNEL_TYPE_REPEAT_F16,
@@ -299,9 +297,6 @@ enum ggml_metal_kernel_type {
GGML_METAL_KERNEL_TYPE_SET_ROWS_Q5_0,
GGML_METAL_KERNEL_TYPE_SET_ROWS_Q5_1,
GGML_METAL_KERNEL_TYPE_SET_ROWS_IQ4_NL,
GGML_METAL_KERNEL_TYPE_RMS_NORM,
GGML_METAL_KERNEL_TYPE_RMS_NORM_MUL,
GGML_METAL_KERNEL_TYPE_RMS_NORM_MUL_ADD,
GGML_METAL_KERNEL_TYPE_L2_NORM,
GGML_METAL_KERNEL_TYPE_GROUP_NORM,
GGML_METAL_KERNEL_TYPE_NORM,
@@ -628,7 +623,7 @@ static void ggml_metal_heap_free(struct ggml_metal_heap * heap) {
@end
//
// ggml_metal_mem_pool
// ggml_metal_mem_pool [TAG_MEM_POOL_REMOVE]
//
struct ggml_metal_mem_pool {
@@ -791,6 +786,9 @@ struct ggml_metal_command_buffer {
// each command buffer has a memory pool from which it can allocate temporary buffers during the compute
struct ggml_metal_mem_pool * mem_pool;
// used to enable concurrent execution of ops in the command buffers
struct ggml_mem_ranges * mem_ranges;
};
struct ggml_backend_metal_context {
@@ -1091,7 +1089,9 @@ static struct ggml_backend_metal_context * ggml_metal_init(ggml_backend_dev_t de
GGML_LOG_INFO("%s: has bfloat = %s\n", __func__, ctx_dev->has_bfloat ? "true" : "false");
GGML_LOG_INFO("%s: use bfloat = %s\n", __func__, ctx_dev->use_bfloat ? "true" : "false");
GGML_LOG_INFO("%s: use fusion = %s\n", __func__, ctx_dev->use_fusion ? "true" : "false");
GGML_LOG_INFO("%s: use concurrency = %s\n", __func__, ctx_dev->use_concurrency ? "true" : "false");
GGML_LOG_INFO("%s: use shared buffers = %s\n", __func__, ctx_dev->use_shared_buffers ? "true" : "false");
GGML_LOG_INFO("%s: use graph optimize = %s\n", __func__, ctx_dev->use_graph_optimize ? "true" : "false");
GGML_LOG_INFO("%s: hasUnifiedMemory = %s\n", __func__, ctx_dev->mtl_device.hasUnifiedMemory ? "true" : "false");
ctx->capture_next_compute = false;
@@ -1105,6 +1105,10 @@ static struct ggml_backend_metal_context * ggml_metal_init(ggml_backend_dev_t de
ctx->cmd_bufs[i].mem_pool = ggml_metal_mem_pool_init();
ctx->cmd_bufs[i].mem_pool->device = device;
if (ctx_dev->use_concurrency) {
ctx->cmd_bufs[i].mem_ranges = ggml_mem_ranges_init(ctx_dev->debug_graph);
}
}
ctx->cmd_bufs_ext = [[NSMutableArray alloc] init];
@@ -1148,28 +1152,6 @@ static struct ggml_backend_metal_context * ggml_metal_init(ggml_backend_dev_t de
// simd_sum and simd_max requires MTLGPUFamilyApple7
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ADD, add, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ADD_FUSE_2, add_fuse_2, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ADD_FUSE_3, add_fuse_3, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ADD_FUSE_4, add_fuse_4, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ADD_FUSE_5, add_fuse_5, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ADD_FUSE_6, add_fuse_6, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ADD_FUSE_7, add_fuse_7, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ADD_FUSE_8, add_fuse_8, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ADD_ROW_C4, add_row_c4, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_2, add_row_c4_fuse_2, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_3, add_row_c4_fuse_3, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_4, add_row_c4_fuse_4, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_5, add_row_c4_fuse_5, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_6, add_row_c4_fuse_6, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_7, add_row_c4_fuse_7, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_8, add_row_c4_fuse_8, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_SUB, sub, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_SUB_ROW_C4, sub_row_c4, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_MUL, mul, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_MUL_ROW_C4, mul_row_c4, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_DIV, div, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_DIV_ROW_C4, div_row_c4, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_ADD_ID, add_id, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_REPEAT_F32, repeat_f32, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_REPEAT_F16, repeat_f16, true);
@@ -1235,9 +1217,6 @@ static struct ggml_backend_metal_context * ggml_metal_init(ggml_backend_dev_t de
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_SET_ROWS_Q5_0, set_rows_q5_0, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_SET_ROWS_Q5_1, set_rows_q5_1, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_SET_ROWS_IQ4_NL, set_rows_iq4_nl, true);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_RMS_NORM, rms_norm, has_simdgroup_reduction);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_RMS_NORM_MUL, rms_norm_mul, has_simdgroup_reduction);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_RMS_NORM_MUL_ADD, rms_norm_mul_add, has_simdgroup_reduction);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_L2_NORM, l2_norm, has_simdgroup_reduction);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_GROUP_NORM, group_norm, has_simdgroup_reduction);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_NORM, norm, true);
@@ -1361,7 +1340,6 @@ static struct ggml_backend_metal_context * ggml_metal_init(ggml_backend_dev_t de
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_MUL_MM_Q5_1_F32, mul_mm_q5_1_f32, has_simdgroup_mm);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_MUL_MM_Q8_0_F32, mul_mm_q8_0_f32, has_simdgroup_mm);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_MUL_MM_MXFP4_F32, mul_mm_mxfp4_f32, has_simdgroup_mm);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_MUL_MM_MXFP4_F32, mul_mm_mxfp4_f32, has_simdgroup_mm);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_MUL_MM_Q2_K_F32, mul_mm_q2_K_f32, has_simdgroup_mm);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_MUL_MM_Q3_K_F32, mul_mm_q3_K_f32, has_simdgroup_mm);
GGML_METAL_ADD_KERNEL(GGML_METAL_KERNEL_TYPE_MUL_MM_Q4_K_F32, mul_mm_q4_K_f32, has_simdgroup_mm);
@@ -1521,6 +1499,9 @@ static id<MTLComputePipelineState> ggml_metal_compile_kernel(ggml_backend_t back
NSString * key = [NSString stringWithUTF8String:name];
[ctx->kernels_ext setObject:obj forKey:key];
[metal_function release];
[obj release];
GGML_LOG_DEBUG("%s: loaded %-40s %16p | th_max = %4d | th_width = %4d\n", __func__, name, (void *) kernel.pipeline,
(int) kernel.pipeline.maxTotalThreadsPerThreadgroup,
(int) kernel.pipeline.threadExecutionWidth);
@@ -1542,8 +1523,6 @@ static id<MTLComputePipelineState> ggml_metal_get_pipeline_flash_attn_ext(
char name[256];
@autoreleasepool {
MTLFunctionConstantValues * cv = [[MTLFunctionConstantValues alloc] init];
const int32_t dk = (int32_t) op->src[1]->ne[0];
const int32_t dv = (int32_t) op->src[2]->ne[0];
@@ -1575,7 +1554,7 @@ static id<MTLComputePipelineState> ggml_metal_get_pipeline_flash_attn_ext(
return res;
}
cv = [[MTLFunctionConstantValues alloc] init];
MTLFunctionConstantValues * cv = [[MTLFunctionConstantValues alloc] init];
[cv setConstantValue:&has_mask type:MTLDataTypeBool atIndex:FC_FLASH_ATTN_EXT + 0];
[cv setConstantValue:&has_sinks type:MTLDataTypeBool atIndex:FC_FLASH_ATTN_EXT + 1];
@@ -1586,7 +1565,11 @@ static id<MTLComputePipelineState> ggml_metal_get_pipeline_flash_attn_ext(
[cv setConstantValue:&ns20 type:MTLDataTypeInt atIndex:FC_FLASH_ATTN_EXT + 21];
[cv setConstantValue:&nsg type:MTLDataTypeInt atIndex:FC_FLASH_ATTN_EXT + 22];
return ggml_metal_compile_kernel(backend, base, name, cv);
res = ggml_metal_compile_kernel(backend, base, name, cv);
[cv release];
return res;
}
}
@@ -1604,8 +1587,6 @@ static id<MTLComputePipelineState> ggml_metal_get_pipeline_flash_attn_ext_vec(
char name[256];
@autoreleasepool {
MTLFunctionConstantValues * cv = [[MTLFunctionConstantValues alloc] init];
const int32_t dk = (int32_t) op->src[1]->ne[0];
const int32_t dv = (int32_t) op->src[2]->ne[0];
@@ -1637,7 +1618,7 @@ static id<MTLComputePipelineState> ggml_metal_get_pipeline_flash_attn_ext_vec(
return res;
}
cv = [[MTLFunctionConstantValues alloc] init];
MTLFunctionConstantValues * cv = [[MTLFunctionConstantValues alloc] init];
[cv setConstantValue:&has_mask type:MTLDataTypeBool atIndex:FC_FLASH_ATTN_EXT_VEC + 0];
[cv setConstantValue:&has_sinks type:MTLDataTypeBool atIndex:FC_FLASH_ATTN_EXT_VEC + 1];
@@ -1649,7 +1630,11 @@ static id<MTLComputePipelineState> ggml_metal_get_pipeline_flash_attn_ext_vec(
[cv setConstantValue:&nsg type:MTLDataTypeInt atIndex:FC_FLASH_ATTN_EXT_VEC + 22];
[cv setConstantValue:&nwg type:MTLDataTypeInt atIndex:FC_FLASH_ATTN_EXT_VEC + 23];
return ggml_metal_compile_kernel(backend, base, name, cv);
res = ggml_metal_compile_kernel(backend, base, name, cv);
[cv release];
return res;
}
}
@@ -1663,8 +1648,6 @@ static id<MTLComputePipelineState> ggml_metal_get_pipeline_flash_attn_ext_vec_re
char name[256];
@autoreleasepool {
MTLFunctionConstantValues * cv = [[MTLFunctionConstantValues alloc] init];
snprintf(base, 256, "kernel_flash_attn_ext_vec_reduce");
snprintf(name, 256, "kernel_flash_attn_ext_vec_reduce_dv=%d_nwg=%d", dv, nwg);
@@ -1674,12 +1657,83 @@ static id<MTLComputePipelineState> ggml_metal_get_pipeline_flash_attn_ext_vec_re
return res;
}
cv = [[MTLFunctionConstantValues alloc] init];
MTLFunctionConstantValues * cv = [[MTLFunctionConstantValues alloc] init];
[cv setConstantValue:&dv type:MTLDataTypeInt atIndex:FC_FLASH_ATTN_EXT_VEC_REDUCE + 0];
[cv setConstantValue:&nwg type:MTLDataTypeInt atIndex:FC_FLASH_ATTN_EXT_VEC_REDUCE + 1];
return ggml_metal_compile_kernel(backend, base, name, cv);
res = ggml_metal_compile_kernel(backend, base, name, cv);
[cv release];
return res;
}
GGML_UNUSED(op);
}
static id<MTLComputePipelineState> ggml_metal_get_pipeline_bin(
ggml_backend_t backend, enum ggml_op op,
int32_t n_fuse,
bool row) {
struct ggml_backend_metal_context * ctx = backend->context;
char base[256];
char name[256];
@autoreleasepool {
const char * op_str = "undefined";
switch (op) {
case GGML_OP_ADD: op_str = "add"; break;
case GGML_OP_SUB: op_str = "sub"; break;
case GGML_OP_MUL: op_str = "mul"; break;
case GGML_OP_DIV: op_str = "div"; break;
default: GGML_ABORT("fatal error");
};
if (row) {
snprintf(base, 256, "kernel_%s_row_c4_fuse_%d", op_str, n_fuse);
} else {
snprintf(base, 256, "kernel_%s_fuse_%d", op_str, n_fuse);
}
snprintf(name, 256, "%s", base);
id<MTLComputePipelineState> res = ggml_metal_get_kernel(ctx, name);
if (res) {
// kernel found
return res;
}
return ggml_metal_compile_kernel(backend, base, name, nil);
}
}
static id<MTLComputePipelineState> ggml_metal_get_pipeline_rms_norm(
ggml_backend_t backend, struct ggml_tensor * op,
int32_t n_fuse) {
struct ggml_backend_metal_context * ctx = backend->context;
char base[256];
char name[256];
@autoreleasepool {
switch (n_fuse) {
case 1: snprintf(base, 256, "kernel_rms_norm"); break;
case 2: snprintf(base, 256, "kernel_rms_norm_mul"); break;
case 3: snprintf(base, 256, "kernel_rms_norm_mul_add"); break;
default: GGML_ABORT("fatal error");
}
snprintf(name, 256, "%s", base);
id<MTLComputePipelineState> res = ggml_metal_get_kernel(ctx, name);
if (res) {
// kernel found
return res;
}
return ggml_metal_compile_kernel(backend, base, name, nil);
}
GGML_UNUSED(op);
@@ -1707,6 +1761,10 @@ static void ggml_metal_free(struct ggml_backend_metal_context * ctx) {
}
ggml_metal_mem_pool_free(ctx->cmd_bufs[i].mem_pool);
if (ctx->cmd_bufs[i].mem_ranges) {
ggml_mem_ranges_free(ctx->cmd_bufs[i].mem_ranges);
}
}
[ctx->cmd_bufs_ext removeAllObjects];
@@ -2063,12 +2121,51 @@ static bool ggml_metal_supports_op(const struct ggml_backend_metal_device_contex
}
}
static int ggml_metal_encode_node(
ggml_backend_t backend,
int idx,
int idx_end,
id<MTLComputeCommandEncoder> encoder,
struct ggml_metal_mem_pool * mem_pool) {
struct ggml_metal_encode_context {
ggml_backend_t backend;
id<MTLComputeCommandEncoder> encoder;
struct ggml_metal_mem_pool * mem_pool;
struct ggml_mem_ranges * mem_ranges;
};
static bool ggml_metal_encode_concurrency_reset(struct ggml_metal_encode_context * ctx) {
if (!ctx->mem_ranges) {
return true;
}
[ctx->encoder memoryBarrierWithScope:MTLBarrierScopeBuffers];
ggml_mem_ranges_reset(ctx->mem_ranges);
return true;
}
static bool ggml_metal_encode_concurrency_check(struct ggml_metal_encode_context * ctx, const struct ggml_tensor * node) {
if (!ctx->mem_ranges) {
return false;
}
return ggml_mem_ranges_check(ctx->mem_ranges, node);
}
static bool ggml_metal_encode_concurrency_add(struct ggml_metal_encode_context * ctx, const struct ggml_tensor * node) {
if (!ctx->mem_ranges) {
return true;
}
return ggml_mem_ranges_add(ctx->mem_ranges, node);
}
static int ggml_metal_encode_node(struct ggml_metal_encode_context * ctx_enc, int idx, int idx_end) {
ggml_backend_t backend = ctx_enc->backend;
id<MTLComputeCommandEncoder> encoder = ctx_enc->encoder;
struct ggml_metal_mem_pool * mem_pool = ctx_enc->mem_pool;
struct ggml_backend_metal_context * ctx = backend->context;
struct ggml_backend_metal_device_context * ctx_dev = backend->device->context;
@@ -2151,38 +2248,71 @@ static int ggml_metal_encode_node(
const uint64_t nb2 = dst ? dst->nb[2] : 0;
const uint64_t nb3 = dst ? dst->nb[3] : 0;
size_t offs_src[GGML_MAX_SRC];
id<MTLBuffer> id_src[GGML_MAX_SRC];
enum ggml_type srct[GGML_MAX_SRC];
for (int i = 0; i < GGML_MAX_SRC; i++) {
offs_src[i] = 0;
id_src[i] = node->src[i] ? ggml_metal_get_buffer(node->src[i], &offs_src[i]) : nil;
srct[i] = node->src[i] ? node->src[i]->type : GGML_TYPE_COUNT;
}
// TODO: tmp shorthands - remove
size_t offs_src0 = offs_src[0];
size_t offs_src1 = offs_src[1];
size_t offs_src2 = offs_src[2];
id<MTLBuffer> id_src0 = id_src[0];
id<MTLBuffer> id_src1 = id_src[1];
id<MTLBuffer> id_src2 = id_src[2];
const enum ggml_type src0t = src0 ? src0->type : GGML_TYPE_COUNT;
const enum ggml_type src1t = src1 ? src1->type : GGML_TYPE_COUNT;
const enum ggml_type src2t = src2 ? src2->type : GGML_TYPE_COUNT;
const enum ggml_type dstt = dst ? dst->type : GGML_TYPE_COUNT;
size_t offs_src0 = 0;
size_t offs_src1 = 0;
size_t offs_src2 = 0;
size_t offs_dst = 0;
size_t offs_dst = 0;
id<MTLBuffer> id_src0 = src0 ? ggml_metal_get_buffer(src0, &offs_src0) : nil;
id<MTLBuffer> id_src1 = src1 ? ggml_metal_get_buffer(src1, &offs_src1) : nil;
id<MTLBuffer> id_src2 = src2 ? ggml_metal_get_buffer(src2, &offs_src2) : nil;
id<MTLBuffer> id_dst = dst ? ggml_metal_get_buffer(dst, &offs_dst) : nil;
id<MTLBuffer> id_dst = dst ? ggml_metal_get_buffer(dst, &offs_dst) : nil;
int n_fuse = 1;
#if 0
GGML_LOG_INFO("%s: op - %s\n", __func__, ggml_op_name(dst->op));
if (src0) {
GGML_LOG_INFO("%s: src0 - %4s [%5lld, %5lld, %5lld, %5lld] [%5lld, %5lld, %5lld, %5lld], %d, %s\n", __func__, ggml_type_name(src0t), ne00, ne01, ne02, ne03, nb00, nb01, nb02, nb03,
ggml_is_contiguous(src0), src0->name);
// check if the current node can run concurrently with other nodes before it
// the condition is that:
// - the current node cannot write to any previous src or dst ranges
// - the current node cannot read from any previous dst ranges
//
// if the condition is not satisfied, we put a memory barrier and clear all ranges
// otherwise, we add the new ranges to the encoding context and process the node concurrently
//
{
const bool is_concurrent = ggml_metal_encode_concurrency_check(ctx_enc, node);
if (!is_concurrent) {
ggml_metal_encode_concurrency_reset(ctx_enc);
}
if (ctx_dev->debug_graph > 0) {
GGML_LOG_DEBUG("%s: node[%5d] - %-12s %s\n", __func__, idx, ggml_op_name(dst->op), is_concurrent ? "(concurrent)" : "");
}
if (ctx_dev->debug_graph > 1) {
if (src0) {
GGML_LOG_DEBUG("%s: src0 - %4s [%5lld, %5lld, %5lld, %5lld] [%5lld, %5lld, %5lld, %5lld], %d, %s\n", __func__, ggml_type_name(src0t), ne00, ne01, ne02, ne03, nb00, nb01, nb02, nb03,
ggml_is_contiguous(src0), src0->name);
}
if (src1) {
GGML_LOG_DEBUG("%s: src1 - %4s [%5lld, %5lld, %5lld, %5lld] [%5lld, %5lld, %5lld, %5lld], %d, %s\n", __func__, ggml_type_name(src1t), ne10, ne11, ne12, ne13, nb10, nb11, nb12, nb13,
ggml_is_contiguous(src1), src1->name);
}
if (dst) {
GGML_LOG_DEBUG("%s: dst - %4s [%5lld, %5lld, %5lld, %5lld] [%5lld, %5lld, %5lld, %5lld], 1, %s\n", __func__, ggml_type_name(dstt), ne0, ne1, ne2, ne3, nb0, nb1, nb2, nb3,
dst->name);
}
}
}
if (src1) {
GGML_LOG_INFO("%s: src1 - %4s [%5lld, %5lld, %5lld, %5lld] [%5lld, %5lld, %5lld, %5lld], %d, %s\n", __func__, ggml_type_name(src1t), ne10, ne11, ne12, ne13, nb10, nb11, nb12, nb13,
ggml_is_contiguous(src1), src1->name);
}
if (dst) {
GGML_LOG_INFO("%s: dst - %4s [%5lld, %5lld, %5lld, %5lld] [%5lld, %5lld, %5lld, %5lld], 1, %s\n", __func__, ggml_type_name(dstt), ne0, ne1, ne2, ne3, nb0, nb1, nb2, nb3,
dst->name);
}
#endif
id<MTLDevice> device = ctx_dev->mtl_device;
@@ -2246,8 +2376,6 @@ static int ggml_metal_encode_node(
bool bcast_row = false;
id<MTLComputePipelineState> pipeline = nil;
ggml_metal_kargs_bin args = {
/*.ne00 =*/ ne00,
/*.ne01 =*/ ne01,
@@ -2328,59 +2456,31 @@ static int ggml_metal_encode_node(
}
}
id<MTLComputePipelineState> pipeline = nil;
if (ggml_nelements(src1) == ne10 && ggml_is_contiguous(src1) && ne00 % 4 == 0 && ne10 % 4 == 0) {
GGML_ASSERT(ggml_is_contiguous(src0));
// src1 is a row
GGML_ASSERT(ne11 == 1);
switch (dst->op) {
case GGML_OP_ADD:
{
switch (n_fuse) {
case 1: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD_ROW_C4 ].pipeline; break;
case 2: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_2].pipeline; break;
case 3: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_3].pipeline; break;
case 4: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_4].pipeline; break;
case 5: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_5].pipeline; break;
case 6: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_6].pipeline; break;
case 7: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_7].pipeline; break;
case 8: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD_ROW_C4_FUSE_8].pipeline; break;
default: GGML_ABORT("fatal error");
}
} break;
case GGML_OP_SUB: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_SUB_ROW_C4].pipeline; break;
case GGML_OP_MUL: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_MUL_ROW_C4].pipeline; break;
case GGML_OP_DIV: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_DIV_ROW_C4].pipeline; break;
default: GGML_ABORT("fatal error");
}
pipeline = ggml_metal_get_pipeline_bin(backend, dst->op, n_fuse, true);
bcast_row = true;
} else {
switch (dst->op) {
case GGML_OP_ADD:
{
switch (n_fuse) {
case 1: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD ].pipeline; break;
case 2: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD_FUSE_2].pipeline; break;
case 3: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD_FUSE_3].pipeline; break;
case 4: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD_FUSE_4].pipeline; break;
case 5: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD_FUSE_5].pipeline; break;
case 6: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD_FUSE_6].pipeline; break;
case 7: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD_FUSE_7].pipeline; break;
case 8: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD_FUSE_8].pipeline; break;
default: GGML_ABORT("fatal error");
}
} break;
case GGML_OP_SUB: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_SUB].pipeline; break;
case GGML_OP_MUL: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_MUL].pipeline; break;
case GGML_OP_DIV: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_DIV].pipeline; break;
default: GGML_ABORT("fatal error");
}
pipeline = ggml_metal_get_pipeline_bin(backend, dst->op, n_fuse, false);
}
if (n_fuse > 1) {
id_dst = ggml_metal_get_buffer(nodes[n_fuse - 1], &offs_dst);
for (int i = 1; i < n_fuse; ++i) {
if (!ggml_metal_encode_concurrency_check(ctx_enc, nodes[i])) {
ggml_metal_encode_concurrency_reset(ctx_enc);
break;
}
}
}
[encoder setComputePipelineState:pipeline];
@@ -2525,9 +2625,9 @@ static int ggml_metal_encode_node(
const int nth = MIN((int) pipeline.maxTotalThreadsPerThreadgroup, ne00);
[encoder dispatchThreadgroups:MTLSizeMake(ne01, ne02, ne03) threadsPerThreadgroup:MTLSizeMake(nth, 1, 1)];
}
const id<MTLComputePipelineState> pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD].pipeline;
ggml_metal_encode_concurrency_reset(ctx_enc);
}
ggml_metal_kargs_bin args = {
/*.ne00 =*/ ne00,
@@ -2558,6 +2658,9 @@ static int ggml_metal_encode_node(
/*.o1 =*/ { offs_src1},
};
//const id<MTLComputePipelineState> pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_ADD].pipeline;
const id<MTLComputePipelineState> pipeline = ggml_metal_get_pipeline_bin(backend, GGML_OP_ADD, 1, false);
[encoder setComputePipelineState:pipeline];
[encoder setBytes:&args length:sizeof(args) atIndex:0];
[encoder setBuffer:id_src0 offset:offs_src0 atIndex:1];
@@ -3989,6 +4092,12 @@ static int ggml_metal_encode_node(
default: break;
}
// TODO: using mem pool allocations with enabled concurrency is not safe because the mem pool
// reuses buffers. this can result in 2 concurrent MUL_MAT_ID ops using the same mem pool buffer.
// so we add this extra barrier to prevent the race.
// the correct solution is to remove mem pools and then remove this barrier [TAG_MEM_POOL_REMOVE]
ggml_metal_encode_concurrency_reset(ctx_enc);
// tokens per expert
const size_t s_tpe = ggml_type_size(GGML_TYPE_I32)*ne02;
id<MTLBuffer> h_tpe = ggml_metal_mem_pool_alloc(mem_pool, s_tpe);
@@ -4049,6 +4158,9 @@ static int ggml_metal_encode_node(
[encoder dispatchThreadgroups:MTLSizeMake(1, 1, 1) threadsPerThreadgroup:MTLSizeMake(ne02, 1, 1)];
}
// this barrier is always needed because the next kernel has to wait for the id maps to be computed
ggml_metal_encode_concurrency_reset(ctx_enc);
{
id<MTLComputePipelineState> pipeline = nil;
@@ -4517,16 +4629,17 @@ static int ggml_metal_encode_node(
if (n_fuse > 1) {
id_dst = ggml_metal_get_buffer(nodes[n_fuse - 1], &offs_dst);
for (int i = 1; i < n_fuse; ++i) {
if (!ggml_metal_encode_concurrency_check(ctx_enc, nodes[i])) {
ggml_metal_encode_concurrency_reset(ctx_enc);
break;
}
}
}
id<MTLComputePipelineState> pipeline;
switch (n_fuse) {
case 1: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_RMS_NORM ].pipeline; break;
case 2: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_RMS_NORM_MUL ].pipeline; break;
case 3: pipeline = ctx->kernels[GGML_METAL_KERNEL_TYPE_RMS_NORM_MUL_ADD].pipeline; break;
default: GGML_ABORT("unsupported n_fuse = %d\n", n_fuse);
}
const id<MTLComputePipelineState> pipeline = ggml_metal_get_pipeline_rms_norm(backend, node, n_fuse);
int nth = 32; // SIMD width
@@ -4660,7 +4773,6 @@ static int ggml_metal_encode_node(
} break;
case GGML_OP_ROPE:
{
// make sure we have one or more position id(ne10) per token(ne02)
GGML_ASSERT(ne10 % ne02 == 0);
GGML_ASSERT(ne10 >= ne02);
@@ -5419,6 +5531,10 @@ static int ggml_metal_encode_node(
GGML_ASSERT(ne01*ne02*ne03 == ne1*ne2*ne3);
GGML_ASSERT(ne1*ne2*ne3 <= (1u << 31));
// using mem pool allocations with enabled concurrency is not safe [TAG_MEM_POOL_REMOVE]
// still, we assume that concurrent FA won't happen before we do the refactor
//ggml_metal_encode_concurrency_reset(ctx_enc);
const int32_t nrows = ne1*ne2*ne3;
// temp buffer for writing the results from each workgroup
@@ -5439,6 +5555,8 @@ static int ggml_metal_encode_node(
[encoder setThreadgroupMemoryLength:smem atIndex:0];
[encoder dispatchThreadgroups:MTLSizeMake((ne01 + nqptg - 1)/nqptg, ne02, ne03*nwg) threadsPerThreadgroup:MTLSizeMake(32, nsg, 1)];
ggml_metal_encode_concurrency_reset(ctx_enc);
// reduce the results from the workgroups
{
ggml_metal_kargs_flash_attn_ext_vec_reduce args0 = {
@@ -5669,7 +5787,7 @@ static int ggml_metal_encode_node(
[encoder dispatchThreadgroups:MTLSizeMake(n_tg, 1, 1) threadsPerThreadgroup:MTLSizeMake(n_threads, 1, 1)];
} break;
case GGML_OP_ARGMAX:
case GGML_OP_ARGMAX:
{
GGML_ASSERT(src0->type == GGML_TYPE_F32);
GGML_ASSERT(ggml_is_contiguous_1(src0));
@@ -5701,6 +5819,19 @@ static int ggml_metal_encode_node(
}
}
if (ctx_dev->debug_graph > 0) {
if (n_fuse > 1) {
GGML_LOG_DEBUG("%s: fuse %d ops\n", __func__, n_fuse);
}
}
// update the mem ranges in the encoding context
for (int i = 0; i < n_fuse; ++i) {
if (!ggml_metal_encode_concurrency_add(ctx_enc, nodes[i])) {
ggml_metal_encode_concurrency_reset(ctx_enc);
}
}
return n_fuse;
}
@@ -5711,7 +5842,7 @@ static enum ggml_status ggml_metal_graph_compute(
struct ggml_backend_metal_device_context * ctx_dev = backend->device->context;
// number of nodes encoded by the main thread (empirically determined)
const int n_main = 128;
const int n_main = 64;
// number of threads in addition to the main thread
const int n_cb = ctx->n_cb;
@@ -5766,10 +5897,14 @@ static enum ggml_status ggml_metal_graph_compute(
// cannot use commandBufferWithUnretainedReferences because the buffers from the memory pool can get destroyed
// TODO: when the memory pools are removed, we can again use commandBufferWithUnretainedReferences
// https://github.com/ggml-org/llama.cpp/pull/15832#discussion_r2334215009
// [TAG_MEM_POOL_REMOVE]
//id<MTLCommandBuffer> cmd_buf = [ctx->queue commandBufferWithUnretainedReferences];
id<MTLCommandBuffer> cmd_buf = [ctx->queue commandBuffer];
[cmd_buf retain];
if (ctx->cmd_bufs[n_cb].obj) {
[ctx->cmd_bufs[n_cb].obj release];
}
ctx->cmd_bufs[n_cb].obj = cmd_buf;
[cmd_buf enqueue];
@@ -6536,6 +6671,18 @@ static enum ggml_status ggml_backend_metal_graph_compute(ggml_backend_t backend,
return ggml_metal_graph_compute(backend, cgraph);
}
static void ggml_backend_metal_graph_optimize(ggml_backend_t backend, struct ggml_cgraph * cgraph) {
struct ggml_backend_metal_device_context * ctx_dev = backend->device->context;
//const int64_t t_start = ggml_time_us();
if (ctx_dev->use_graph_optimize) {
ggml_metal_graph_optimize(cgraph);
}
//printf("%s: graph optimize took %.3f ms\n", __func__, (ggml_time_us() - t_start) / 1000.0);
}
static void ggml_backend_metal_set_n_cb(ggml_backend_t backend, int n_cb) {
GGML_ASSERT(ggml_backend_is_metal(backend));
@@ -6562,12 +6709,25 @@ static void ggml_backend_metal_set_n_cb(ggml_backend_t backend, int n_cb) {
const int n_nodes_per_cb = ctx->n_nodes_per_cb;
id<MTLCommandBuffer> cmd_buf = ctx->cmd_bufs[cb_idx].obj;
struct ggml_metal_mem_pool * mem_pool = ctx->cmd_bufs[cb_idx].mem_pool;
id<MTLCommandBuffer> cmd_buf = ctx->cmd_bufs[cb_idx].obj;
struct ggml_metal_mem_pool * mem_pool = ctx->cmd_bufs[cb_idx].mem_pool;
struct ggml_mem_ranges * mem_ranges = ctx->cmd_bufs[cb_idx].mem_ranges;
ggml_metal_mem_pool_reset(mem_pool);
id<MTLComputeCommandEncoder> encoder = [cmd_buf computeCommandEncoder];
if (mem_ranges) {
ggml_mem_ranges_reset(mem_ranges);
}
id<MTLComputeCommandEncoder> encoder;
struct ggml_backend_metal_device_context * ctx_dev = backend->device->context;
if (ctx_dev->use_concurrency) {
encoder = [cmd_buf computeCommandEncoderWithDispatchType: MTLDispatchTypeConcurrent];
} else {
encoder = [cmd_buf computeCommandEncoder];
}
int node_start = 0;
int node_end = n_nodes_0;
@@ -6579,12 +6739,19 @@ static void ggml_backend_metal_set_n_cb(ggml_backend_t backend, int n_cb) {
const bool should_capture = ctx->capture_next_compute;
struct ggml_metal_encode_context ctx_enc = {
/*.backend =*/ backend,
/*.encoder =*/ encoder,
/*.mem_pool =*/ mem_pool,
/*.mem_ranges =*/ mem_ranges,
};
for (int idx = node_start; idx < node_end;) {
if (should_capture) {
[encoder pushDebugGroup:[NSString stringWithCString:ggml_op_desc(ggml_graph_node(ctx->gf, idx)) encoding:NSUTF8StringEncoding]];
}
const int res = ggml_metal_encode_node(backend, idx, node_end, encoder, mem_pool);
const int res = ggml_metal_encode_node(&ctx_enc, idx, node_end);
if (idx + res > node_end) {
GGML_ABORT("fusion error: nodes spanning multiple encoders have been fused. this indicates a bug in the fusion logic %s",
"https://github.com/ggml-org/llama.cpp/pull/14849");
@@ -6627,7 +6794,7 @@ static struct ggml_backend_i ggml_backend_metal_i = {
// https://developer.apple.com/documentation/metal/mtlcommandbuffer#Synchronizing-Passes-with-Events
/* .event_record = */ NULL,
/* .event_wait = */ NULL,
/* .optimize_graph = */ NULL,
/* .optimize_graph = */ ggml_backend_metal_graph_optimize,
};
static ggml_guid_t ggml_backend_metal_guid(void) {
ggml/src/ggml-metal/ggml-metal.metal
+9 -15
View File
@@ -928,7 +928,7 @@ kernel void kernel_add_fuse_impl(
typedef decltype(kernel_add_fuse_impl<2>) kernel_add_fuse_t;
template [[host_name("kernel_add")]] kernel kernel_add_fuse_t kernel_add_fuse_impl<1>;
template [[host_name("kernel_add_fuse_1")]] kernel kernel_add_fuse_t kernel_add_fuse_impl<1>;
template [[host_name("kernel_add_fuse_2")]] kernel kernel_add_fuse_t kernel_add_fuse_impl<2>;
template [[host_name("kernel_add_fuse_3")]] kernel kernel_add_fuse_t kernel_add_fuse_impl<3>;
template [[host_name("kernel_add_fuse_4")]] kernel kernel_add_fuse_t kernel_add_fuse_impl<4>;
@@ -937,7 +937,7 @@ template [[host_name("kernel_add_fuse_6")]] kernel kernel_add_fuse_t kernel_add_
template [[host_name("kernel_add_fuse_7")]] kernel kernel_add_fuse_t kernel_add_fuse_impl<7>;
template [[host_name("kernel_add_fuse_8")]] kernel kernel_add_fuse_t kernel_add_fuse_impl<8>;
kernel void kernel_sub(
kernel void kernel_sub_fuse_1(
constant ggml_metal_kargs_bin & args,
device const char * src0,
device const char * src1,
@@ -963,7 +963,7 @@ kernel void kernel_sub(
}
}
kernel void kernel_mul(
kernel void kernel_mul_fuse_1(
constant ggml_metal_kargs_bin & args,
device const char * src0,
device const char * src1,
@@ -996,7 +996,7 @@ kernel void kernel_mul(
}
}
kernel void kernel_div(
kernel void kernel_div_fuse_1(
constant ggml_metal_kargs_bin & args,
device const char * src0,
device const char * src1,
@@ -1096,23 +1096,17 @@ kernel void kernel_add_row_c4_fuse_impl(
device const char * src1,
device char * dst,
uint tpig[[thread_position_in_grid]]) {
const uint nb = args.ne00/4;
const uint i = tpig % nb;
device const float4 * src0_row = (device const float4 *) (src0);
device float4 * dst_row = (device float4 *) (dst);
device const float4 * src1_row[F];
for (short j = 0; j < F; ++j) {
src1_row[j] = (device const float4 *) (src1 + args.o1[j]);
}
float4 res = src0_row[tpig];
#pragma unroll(F)
for (short j = 0; j < F; ++j) {
res += src1_row[j][i];
res += ((device const float4 *) (src1 + args.o1[j]))[i];
}
dst_row[tpig] = res;
@@ -1120,7 +1114,7 @@ kernel void kernel_add_row_c4_fuse_impl(
typedef decltype(kernel_add_row_c4_fuse_impl<1>) kernel_add_row_c4_fuse_t;
template [[host_name("kernel_add_row_c4")]] kernel kernel_add_row_c4_fuse_t kernel_add_row_c4_fuse_impl<1>;
template [[host_name("kernel_add_row_c4_fuse_1")]] kernel kernel_add_row_c4_fuse_t kernel_add_row_c4_fuse_impl<1>;
template [[host_name("kernel_add_row_c4_fuse_2")]] kernel kernel_add_row_c4_fuse_t kernel_add_row_c4_fuse_impl<2>;
template [[host_name("kernel_add_row_c4_fuse_3")]] kernel kernel_add_row_c4_fuse_t kernel_add_row_c4_fuse_impl<3>;
template [[host_name("kernel_add_row_c4_fuse_4")]] kernel kernel_add_row_c4_fuse_t kernel_add_row_c4_fuse_impl<4>;
@@ -1160,7 +1154,7 @@ kernel void kernel_sub_row_c4_fuse_impl(
typedef decltype(kernel_sub_row_c4_fuse_impl<1>) kernel_sub_row_c4_fuse_t;
template [[host_name("kernel_sub_row_c4")]] kernel kernel_sub_row_c4_fuse_t kernel_sub_row_c4_fuse_impl<1>;
template [[host_name("kernel_sub_row_c4_fuse_1")]] kernel kernel_sub_row_c4_fuse_t kernel_sub_row_c4_fuse_impl<1>;
template <short F>
kernel void kernel_mul_row_c4_fuse_impl(
@@ -1193,7 +1187,7 @@ kernel void kernel_mul_row_c4_fuse_impl(
typedef decltype(kernel_mul_row_c4_fuse_impl<1>) kernel_mul_row_c4_fuse_t;
template [[host_name("kernel_mul_row_c4")]] kernel kernel_mul_row_c4_fuse_t kernel_mul_row_c4_fuse_impl<1>;
template [[host_name("kernel_mul_row_c4_fuse_1")]] kernel kernel_mul_row_c4_fuse_t kernel_mul_row_c4_fuse_impl<1>;
template <short F>
kernel void kernel_div_row_c4_fuse_impl(
@@ -1226,7 +1220,7 @@ kernel void kernel_div_row_c4_fuse_impl(
typedef decltype(kernel_div_row_c4_fuse_impl<1>) kernel_div_row_c4_fuse_t;
template [[host_name("kernel_div_row_c4")]] kernel kernel_div_row_c4_fuse_t kernel_div_row_c4_fuse_impl<1>;
template [[host_name("kernel_div_row_c4_fuse_1")]] kernel kernel_div_row_c4_fuse_t kernel_div_row_c4_fuse_impl<1>;
kernel void kernel_scale(
device const float * src0,
ggml/src/ggml-vulkan/ggml-vulkan.cpp
+43
View File
@@ -17,8 +17,14 @@
#include "ggml-cpu.h"
#endif
// See https://github.com/KhronosGroup/Vulkan-Hpp?tab=readme-ov-file#extensions--per-device-function-pointers-
#define VULKAN_HPP_DISPATCH_LOADER_DYNAMIC 1
#include <vulkan/vulkan.hpp>
// See https://github.com/KhronosGroup/Vulkan-Hpp?tab=readme-ov-file#extensions--per-device-function-pointers-
VULKAN_HPP_DEFAULT_DISPATCH_LOADER_DYNAMIC_STORAGE
#include <algorithm>
#include <cmath>
#include <iomanip>
@@ -137,6 +143,8 @@ struct vk_pipeline_struct {
bool needed {};
// set to true when the shader has been compiled
bool compiled {};
// number of registers used, extracted from pipeline executable properties
uint32_t register_count {};
};
typedef std::shared_ptr<vk_pipeline_struct> vk_pipeline;
@@ -445,6 +453,8 @@ struct vk_device_struct {
bool coopmat2;
bool pipeline_executable_properties_support {};
size_t idx;
bool mul_mat_l[GGML_TYPE_COUNT];
@@ -1619,6 +1629,20 @@ static void ggml_vk_create_pipeline_func(vk_device& device, vk_pipeline& pipelin
vk_instance.pfn_vkSetDebugUtilsObjectNameEXT(device->device, &static_cast<VkDebugUtilsObjectNameInfoEXT &>(duoni));
}
if (device->pipeline_executable_properties_support) {
vk::PipelineExecutableInfoKHR executableInfo;
executableInfo.pipeline = pipeline->pipeline;
auto statistics = device->device.getPipelineExecutableStatisticsKHR(executableInfo);
for (auto & s : statistics) {
// "Register Count" is reported by NVIDIA drivers.
if (strcmp(s.name, "Register Count") == 0) {
VK_LOG_DEBUG(pipeline->name << " " << s.name << ": " << s.value.u64 << " registers");
pipeline->register_count = (uint32_t)s.value.u64;
}
}
}
{
std::lock_guard<std::recursive_mutex> guard(device->mutex);
device->all_pipelines.push_back(pipeline);
@@ -3632,6 +3656,7 @@ static vk_device ggml_vk_get_device(size_t idx) {
bool amd_shader_core_properties2 = false;
bool pipeline_robustness = false;
bool coopmat2_support = false;
bool pipeline_executable_properties_support = false;
device->coopmat_support = false;
device->integer_dot_product = false;
bool bfloat16_support = false;
@@ -3674,6 +3699,8 @@ static vk_device ggml_vk_get_device(size_t idx) {
!getenv("GGML_VK_DISABLE_BFLOAT16")) {
bfloat16_support = true;
#endif
} else if (strcmp("VK_KHR_pipeline_executable_properties", properties.extensionName) == 0) {
pipeline_executable_properties_support = true;
}
}
@@ -3908,8 +3935,18 @@ static vk_device ggml_vk_get_device(size_t idx) {
device_extensions.push_back("VK_KHR_shader_integer_dot_product");
}
VkPhysicalDevicePipelineExecutablePropertiesFeaturesKHR pep_features {};
pep_features.sType = VK_STRUCTURE_TYPE_PHYSICAL_DEVICE_PIPELINE_EXECUTABLE_PROPERTIES_FEATURES_KHR;
if (pipeline_executable_properties_support) {
last_struct->pNext = (VkBaseOutStructure *)&pep_features;
last_struct = (VkBaseOutStructure *)&pep_features;
device_extensions.push_back("VK_KHR_pipeline_executable_properties");
}
vkGetPhysicalDeviceFeatures2(device->physical_device, &device_features2);
device->pipeline_executable_properties_support = pipeline_executable_properties_support;
device->fp16 = device->fp16 && vk12_features.shaderFloat16;
#if defined(VK_KHR_shader_bfloat16)
@@ -4425,6 +4462,9 @@ static void ggml_vk_instance_init() {
}
VK_LOG_DEBUG("ggml_vk_instance_init()");
// See https://github.com/KhronosGroup/Vulkan-Hpp?tab=readme-ov-file#extensions--per-device-function-pointers-
VULKAN_HPP_DEFAULT_DISPATCHER.init(vkGetInstanceProcAddr);
uint32_t api_version = vk::enumerateInstanceVersion();
if (api_version < VK_API_VERSION_1_2) {
@@ -4492,6 +4532,9 @@ static void ggml_vk_instance_init() {
vk_perf_logger_enabled = getenv("GGML_VK_PERF_LOGGER") != nullptr;
// See https://github.com/KhronosGroup/Vulkan-Hpp?tab=readme-ov-file#extensions--per-device-function-pointers-
VULKAN_HPP_DEFAULT_DISPATCHER.init(vk_instance.instance);
std::vector<vk::PhysicalDevice> devices = vk_instance.instance.enumeratePhysicalDevices();
// Emulate behavior of CUDA_VISIBLE_DEVICES for Vulkan
ggml/src/ggml-vulkan/vulkan-shaders/dequant_iq2_s.comp
+1 -1
View File
@@ -29,7 +29,7 @@ void main() {
uint qs = data_a[ib].qs[4 * ib32 + l];
const uint8_t sign = data_a[ib].qs[QUANT_K / 8 + 4 * ib32 + l];
qs |= (qh << (8 - 2 * l)) & 0x300;
const uvec2 grid = iq2s_grid[qs & 511];
const uvec2 grid = iq2s_grid[qs];
const u8vec4 grid0 = unpack8(grid.x);
const u8vec4 grid1 = unpack8(grid.y);
data_b[b_idx + 8 * l + 0] = D_TYPE(db[l/2] * grid0.x * ((sign & 1) != 0 ? -1.0 : 1.0));
ggml/src/ggml-vulkan/vulkan-shaders/dequant_iq2_xxs.comp
+2 -1
View File
@@ -33,7 +33,8 @@ void main() {
[[unroll]] for (uint l = 0; l < 4; ++l) {
const uint sign7 = bitfieldExtract(signscale, 7 * int(l), 7);
const uint sign8 = sign7 | (bitCount(sign7) << 7); // parity bit
const uvec2 grid = iq2xxs_grid[data_a[ib].qs[8 * is + l]];
const uint qs = data_a[ib].qs[8 * is + l];
const uvec2 grid = iq2xxs_grid[qs];
const u8vec4 grid0 = unpack8(grid.x);
const u8vec4 grid1 = unpack8(grid.y);
data_b[b_idx + 8 * l + 0] = D_TYPE(db * grid0.x * ((sign8 & 1) != 0 ? -1.0 : 1.0));
ggml/src/ggml-vulkan/vulkan-shaders/dequant_iq3_s.comp
+6 -5
View File
@@ -22,15 +22,16 @@ void main() {
const uint b_idx = 256 * ib + 32 * is;
const float d = float(data_a[ib].d);
const float db = d * (1 + 2 * ((data_a[ib].scales[is] >> (4 * (is % 2))) & 0xf));
const float db = d * (1 + 2 * ((data_a[ib].scales[is / 2] >> (4 * (is % 2))) & 0xf));
// We must produce 32 values using 4 sign bytes, 1 qh byte, 8 qs bytes.
uint qh = data_a[ib].qh[is];
[[unroll]] for (uint l = 0; l < 8; ++l) {
uint qs = data_a[ib].qs[8 * is + l];
uint gidx = qs | ((qh << (8 - l)) & 256);
uint8_t signs = data_a[ib].signs[8 * is + l / 2] >> (4 * (l & 1));
u8vec4 grid = unpack8(iq3s_grid[gidx]);
const uint iqs = 8 * is + l;
const uint qs = data_a[ib].qs[iqs];
const uint gidx = qs | ((qh << (8 - l)) & 256);
const uint8_t signs = data_a[ib].signs[iqs / 2] >> (4 * (l & 1));
const u8vec4 grid = unpack8(iq3s_grid[gidx]);
data_b[b_idx + 4 * l + 0] = D_TYPE(db * grid.x * ((signs & 1) != 0 ? -1.0 : 1.0));
data_b[b_idx + 4 * l + 1] = D_TYPE(db * grid.y * ((signs & 2) != 0 ? -1.0 : 1.0));
data_b[b_idx + 4 * l + 2] = D_TYPE(db * grid.z * ((signs & 4) != 0 ? -1.0 : 1.0));
ggml/src/ggml-vulkan/vulkan-shaders/dequant_iq3_xxs.comp
+4 -2
View File
@@ -35,8 +35,10 @@ void main() {
const uint sign7 = bitfieldExtract(signscale, 7 * int(l), 7);
// Restore parity bit.
const uint sign8 = sign7 | (bitCount(sign7) << 7);
const u8vec4 grid0 = unpack8(iq3xxs_grid[data_a[ib].qs[8 * is + 2 * l]]);
const u8vec4 grid1 = unpack8(iq3xxs_grid[data_a[ib].qs[8 * is + 2 * l + 1]]);
const uint qs0 = data_a[ib].qs[8 * is + 2 * l];
const uint qs1 = data_a[ib].qs[8 * is + 2 * l + 1];
const u8vec4 grid0 = unpack8(iq3xxs_grid[qs0]);
const u8vec4 grid1 = unpack8(iq3xxs_grid[qs1]);
data_b[b_idx + 8 * l + 0] = D_TYPE(db * grid0.x * ((sign8 & 1) != 0 ? -1.0 : 1.0));
data_b[b_idx + 8 * l + 1] = D_TYPE(db * grid0.y * ((sign8 & 2) != 0 ? -1.0 : 1.0));
data_b[b_idx + 8 * l + 2] = D_TYPE(db * grid0.z * ((sign8 & 4) != 0 ? -1.0 : 1.0));