mirror of
https://github.com/LostRuins/koboldcpp.git
synced 2026-09-18 16:55:14 +02:00
Merge commit '52d42686560a9e8f441f9b9780c8890c37d2802d' into concedo_experimental
# Conflicts: # .github/actions/windows-setup-cuda/action.yml # .github/workflows/build-riscv.yml # .github/workflows/build-sanitize.yml # .github/workflows/build-self-hosted.yml # .github/workflows/release.yml # .github/workflows/server-sanitize.yml # docs/build-s390x.md # docs/build.md # ggml/CMakeLists.txt # ggml/cmake/common.cmake # ggml/src/ggml-cpu/CMakeLists.txt # ggml/src/ggml-cpu/repack.cpp # ggml/src/ggml-cuda/CMakeLists.txt # ggml/src/ggml-cuda/fattn.cu # ggml/src/ggml-hexagon/ggml-hexagon.cpp # ggml/src/ggml-hexagon/htp/hvx-sin-cos.h # ggml/src/ggml-hexagon/htp/rope-ops.c # ggml/src/ggml-hip/CMakeLists.txt # ggml/src/ggml-musa/CMakeLists.txt # ggml/src/ggml-vulkan/ggml-vulkan.cpp # pyproject.toml # requirements/requirements-convert_legacy_llama.txt # requirements/requirements-gguf_editor_gui.txt # requirements/requirements-server-bench.txt # requirements/requirements-tool_bench.txt # scripts/make-release-checks.sh # tests/test-backend-ops.cpp # tests/test-jinja.cpp # tools/cli/README.md # tools/completion/README.md # tools/llama-bench/llama-bench.cpp # tools/server/README.md
This commit is contained in:
Concedo
@@ -873,17 +873,6 @@ static bool common_params_parse_ex(int argc, char ** argv, common_params_context
|
||||
arg.c_str(), e.what(), opt.to_string().c_str()));
|
||||
}
|
||||
}
|
||||
|
||||
// TODO: remove this check after deprecating --mmap|mlock|dio
|
||||
auto has_arg = [&](std::initializer_list<const char *> names) {
|
||||
return std::any_of(names.begin(), names.end(), [&](const char * name) {
|
||||
return seen_args.count(name);
|
||||
});
|
||||
};
|
||||
if (has_arg({"-lm", "--load-mode"}) &&
|
||||
has_arg({"--mlock", "--mmap", "--no-mmap", "-dio", "--direct-io", "-ndio", "--no-direct-io"})) {
|
||||
LOG_WRN("DEPRECATED: `--load-mode` and `--mlock`/`--mmap`/`--direct-io` should not be combined; only the last flag on the command line will take effect\n");
|
||||
}
|
||||
};
|
||||
|
||||
// parse all CLI args now, so that -hf is available below for remote preset resolution
|
||||
@@ -2695,32 +2684,6 @@ common_params_context common_params_parser_init(common_params & params, llama_ex
|
||||
}
|
||||
).set_env("LLAMA_ARG_RPC"));
|
||||
}
|
||||
add_opt(common_arg(
|
||||
{"--mlock"},
|
||||
"DEPRECATED in favor of `--load-mode`: force system to keep model in RAM rather than swapping or compressing",
|
||||
[](common_params & params) {
|
||||
LOG_WRN("DEPRECATED: --mlock is deprecated. use --load-mode mlock instead\n");
|
||||
params.load_mode = LLAMA_LOAD_MODE_MLOCK;
|
||||
}
|
||||
).set_env("LLAMA_ARG_MLOCK"));
|
||||
add_opt(common_arg(
|
||||
{"--mmap"},
|
||||
{"--no-mmap"},
|
||||
"DEPRECATED in favor of `--load-mode`: whether to memory-map model. (if mmap disabled, slower load but may reduce pageouts if not using mlock)",
|
||||
[](common_params & params, bool value) {
|
||||
LOG_WRN("DEPRECATED: --mmap and --no-mmap are deprecated. use --load-mode mmap instead\n");
|
||||
params.load_mode = value ? LLAMA_LOAD_MODE_MMAP : LLAMA_LOAD_MODE_NONE;
|
||||
}
|
||||
).set_env("LLAMA_ARG_MMAP"));
|
||||
add_opt(common_arg(
|
||||
{"-dio", "--direct-io"},
|
||||
{"-ndio", "--no-direct-io"},
|
||||
"DEPRECATED in favor of `--load-mode`: use DirectIO if available",
|
||||
[](common_params & params, bool value) {
|
||||
LOG_WRN("DEPRECATED: --direct-io and --no-direct-io are deprecated. use --load-mode dio instead\n");
|
||||
params.load_mode = value ? LLAMA_LOAD_MODE_DIRECT_IO : LLAMA_LOAD_MODE_NONE;
|
||||
}
|
||||
).set_env("LLAMA_ARG_DIO"));
|
||||
add_opt(common_arg(
|
||||
{"-lm", "--load-mode"}, "MODE",
|
||||
"model loading mode (default: auto)\n"
|
||||
|
||||
@@ -170,6 +170,12 @@ value binary_expression::execute_impl(context & ctx) {
|
||||
}
|
||||
throw std::runtime_error("Cannot perform operation " + op.value + " on undefined values");
|
||||
} else if (is_val<value_none>(left_val) || is_val<value_none>(right_val)) {
|
||||
if (!is_val<value_none>(right_val) && (op.value == "in" || op.value == "not in")) {
|
||||
// case: none in {'low': 1}
|
||||
// A null left operand is looked up like any other value.
|
||||
bool member = test_is_in();
|
||||
return mk_val<value_bool>(op.value == "in" ? member : !member);
|
||||
}
|
||||
if (op.value == "+" || op.value == "~") {
|
||||
value res = mk_val<value_undefined>();
|
||||
if (workaround_concat_null_with_str(res)) {
|
||||
|
||||
@@ -1094,8 +1094,7 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
|
||||
|
||||
// Target prefill may contain token IDs or multimodal embeddings. Both
|
||||
// produce the target-layer features used to seed the draft KV cache, so
|
||||
// skipping the embedding batches leaves a hole in the draft's cache and
|
||||
// the next injection fails to initialize.
|
||||
// embeddings are injected too, except the pinned ones skipped below.
|
||||
// TODO: revisit after https://github.com/ggml-org/llama.cpp/pull/24669 is merged
|
||||
const bool has_tokens = batch_in.token != nullptr;
|
||||
const bool has_embeddings = batch_in.embd != nullptr;
|
||||
@@ -1131,6 +1130,13 @@ struct common_speculative_impl_draft_dflash : public common_speculative_impl {
|
||||
}
|
||||
const int32_t n_rows = i_batch_end[seq_id] - i_batch_beg[seq_id] + 1;
|
||||
|
||||
// an M-RoPE image pins all its rows to one position, so a windowed draft
|
||||
// cache cannot free cells for it - skip it, the draft can jump over the gap
|
||||
const bool pos_pinned = batch_in.pos[i_batch_beg[seq_id]] == batch_in.pos[i_batch_end[seq_id]];
|
||||
if (has_embeddings && n_rows > 1 && pos_pinned) {
|
||||
continue;
|
||||
}
|
||||
|
||||
for (int32_t offset = 0; offset < n_rows; offset += n_ubatch) {
|
||||
const int32_t n_chunk = std::min(n_ubatch, n_rows - offset);
|
||||
|
||||
|
||||
@@ -216,14 +216,14 @@ class NemotronHModel(GraniteHybridModel):
|
||||
hparams = kwargs.pop("hparams", None)
|
||||
if hparams is None:
|
||||
hparams = ModelBase.load_hparams(args[0], self.is_mistral_format)
|
||||
has_moe_params = (
|
||||
"num_experts_per_tok" in hparams
|
||||
or (isinstance(hparams.get("llm_config"), dict) and "num_experts_per_tok" in hparams["llm_config"])
|
||||
)
|
||||
llm_config = {**hparams, **(hparams.get("llm_config") or {})}
|
||||
|
||||
has_moe_params = "num_experts_per_tok" in llm_config
|
||||
layers_block_type = llm_config.get("layers_block_type")
|
||||
|
||||
if has_moe_params:
|
||||
self.model_arch = gguf.MODEL_ARCH.NEMOTRON_H_MOE
|
||||
self.is_moe = True
|
||||
layers_block_type = hparams.get("layers_block_type")
|
||||
if layers_block_type is not None:
|
||||
hparams["num_hidden_layers"] = len(layers_block_type)
|
||||
|
||||
|
||||
@@ -216,7 +216,6 @@
|
||||
// quants.c
|
||||
#define quantize_row_q8_K_generic quantize_row_q8_K
|
||||
#define ggml_vec_dot_nvfp4_q8_0_generic ggml_vec_dot_nvfp4_q8_0
|
||||
#define ggml_vec_dot_q1_0_q8_0_generic ggml_vec_dot_q1_0_q8_0
|
||||
#define ggml_vec_dot_q2_0_q8_0_generic ggml_vec_dot_q2_0_q8_0
|
||||
#define ggml_vec_dot_tq1_0_q8_K_generic ggml_vec_dot_tq1_0_q8_K
|
||||
#define ggml_vec_dot_tq2_0_q8_K_generic ggml_vec_dot_tq2_0_q8_K
|
||||
@@ -229,11 +228,9 @@
|
||||
#define ggml_vec_dot_iq1_s_q8_K_generic ggml_vec_dot_iq1_s_q8_K
|
||||
#define ggml_vec_dot_iq1_m_q8_K_generic ggml_vec_dot_iq1_m_q8_K
|
||||
// repack.cpp
|
||||
#define ggml_quantize_mat_q8_0_4x4_generic ggml_quantize_mat_q8_0_4x4
|
||||
#define ggml_quantize_mat_q8_0_4x8_generic ggml_quantize_mat_q8_0_4x8
|
||||
#define ggml_quantize_mat_q8_K_4x4_generic ggml_quantize_mat_q8_K_4x4
|
||||
#define ggml_quantize_mat_q8_K_4x8_generic ggml_quantize_mat_q8_K_4x8
|
||||
#define ggml_gemv_q4_0_4x4_q8_0_generic ggml_gemv_q4_0_4x4_q8_0
|
||||
#define ggml_gemv_q4_0_4x8_q8_0_generic ggml_gemv_q4_0_4x8_q8_0
|
||||
#define ggml_gemv_q4_0_8x8_q8_0_generic ggml_gemv_q4_0_8x8_q8_0
|
||||
#define ggml_gemv_q2_K_8x8_q8_K_generic ggml_gemv_q2_K_8x8_q8_K
|
||||
@@ -246,7 +243,6 @@
|
||||
#define ggml_gemv_iq4_nl_4x4_q8_0_generic ggml_gemv_iq4_nl_4x4_q8_0
|
||||
#define ggml_gemv_q8_0_4x4_q8_0_generic ggml_gemv_q8_0_4x4_q8_0
|
||||
#define ggml_gemv_q8_0_4x8_q8_0_generic ggml_gemv_q8_0_4x8_q8_0
|
||||
#define ggml_gemm_q4_0_4x4_q8_0_generic ggml_gemm_q4_0_4x4_q8_0
|
||||
#define ggml_gemm_q4_0_4x8_q8_0_generic ggml_gemm_q4_0_4x8_q8_0
|
||||
#define ggml_gemm_q4_0_8x8_q8_0_generic ggml_gemm_q4_0_8x8_q8_0
|
||||
#define ggml_gemm_q2_K_8x8_q8_K_generic ggml_gemm_q2_K_8x8_q8_K
|
||||
|
||||
@@ -146,6 +146,74 @@ void quantize_row_q8_1(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, i
|
||||
|
||||
//===================================== Dot products =================================
|
||||
|
||||
void ggml_vec_dot_q1_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) {
|
||||
const int qk = QK1_0; // 128
|
||||
const int nb = n / qk;
|
||||
|
||||
assert(n % qk == 0);
|
||||
assert(nrc == 1);
|
||||
UNUSED(nrc);
|
||||
UNUSED(bx);
|
||||
UNUSED(by);
|
||||
UNUSED(bs);
|
||||
|
||||
const block_q1_0 * GGML_RESTRICT x = vx;
|
||||
const block_q8_0 * GGML_RESTRICT y = vy;
|
||||
|
||||
#if defined(__VXE__) || defined(__VXE2__)
|
||||
float32x4_t v_sumf = vec_splats(0.0f);
|
||||
|
||||
const uint8x16_t v_zero = vec_splats((uint8_t)0x00); // zero
|
||||
const uint8x16_t v_bias = vec_splats((uint8_t)0x80); // bias from signed to unsigned
|
||||
// v ^ 0x80 == v + 128
|
||||
|
||||
const uint8x16_t v_idx = (const uint8x16_t){ 0, 0, 0, 0, 0, 0, 0, 0, 1, 1, 1, 1, 1, 1, 1, 1 };
|
||||
const uint8x16_t v_bit = (const uint8x16_t){ 1, 2, 4, 8, 16, 32, 64, 128, 1, 2, 4, 8, 16, 32, 64, 128 };
|
||||
|
||||
for (int i = 0; i < nb; ++i) {
|
||||
const uint8x16_t v_x = vec_xl(0, (const uint8_t *)x[i].qs);
|
||||
const float32x4_t v_xd = vec_splats(GGML_CPU_FP16_TO_FP32(x[i].d));
|
||||
|
||||
for (int k = 0; k < 4; ++k) {
|
||||
// sub-block k holds elements 32k .. 32k+31
|
||||
const block_q8_0 * GGML_RESTRICT yb = &y[i*4 + k];
|
||||
const float32x4_t v_yd = vec_splats(GGML_CPU_FP16_TO_FP32(yb->d));
|
||||
|
||||
const uint8x16_t v_xrl = vec_perm(v_x, v_x, vec_add(v_idx, vec_splats((uint8_t)(k*4 + 0))));
|
||||
const uint8x16_t v_xrh = vec_perm(v_x, v_x, vec_add(v_idx, vec_splats((uint8_t)(k*4 + 2))));
|
||||
|
||||
// isolate each lane's bit, then set all ones where that bit is clear, the -d case
|
||||
const int8x16_t v_ml = (int8x16_t)vec_cmpeq(vec_and(v_xrl, v_bit), v_zero);
|
||||
const int8x16_t v_mh = (int8x16_t)vec_cmpeq(vec_and(v_xrh, v_bit), v_zero);
|
||||
|
||||
const int8x16_t v_yl = vec_xl(0, (const int8_t *)yb->qs);
|
||||
const int8x16_t v_yh = vec_xl(QK8_0/2, (const int8_t *)yb->qs);
|
||||
|
||||
// weights are only +1 or -1, so negate y
|
||||
const int8x16_t v_ysl = vec_sub(vec_xor(v_yl, v_ml), v_ml);
|
||||
const int8x16_t v_ysh = vec_sub(vec_xor(v_yh, v_mh), v_mh);
|
||||
|
||||
// bias to unsigned, then vec_sum4 adds each group of 4 bytes into one word
|
||||
const uint32x4_t v_p = vec_add(vec_sum4(vec_xor((uint8x16_t)v_ysl, v_bias), v_zero),
|
||||
vec_sum4(vec_xor((uint8x16_t)v_ysh, v_bias), v_zero));
|
||||
|
||||
// each word summed 8 biased bytes, so take back 8 * 128
|
||||
const int32x4_t v_xy = vec_sub((int32x4_t)v_p, vec_splats((int32_t)1024));
|
||||
|
||||
// apply both block scales and add into the running total
|
||||
v_sumf = vec_madd(vec_float(v_xy), vec_mul(v_xd, v_yd), v_sumf);
|
||||
}
|
||||
}
|
||||
|
||||
*s = vec_hsum_f32x4(v_sumf);
|
||||
#else
|
||||
UNUSED(nb);
|
||||
UNUSED(x);
|
||||
UNUSED(y);
|
||||
ggml_vec_dot_q1_0_q8_0_generic(n, s, bs, vx, bx, vy, by, nrc);
|
||||
#endif
|
||||
}
|
||||
|
||||
void ggml_vec_dot_q4_0_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, size_t bx, const void * GGML_RESTRICT vy, size_t by, int nrc) {
|
||||
const int qk = QK8_0;
|
||||
const int nb = n / qk;
|
||||
|
||||
@@ -0,0 +1,223 @@
|
||||
#define GGML_COMMON_IMPL_CPP
|
||||
#define GGML_COMMON_DECL_CPP
|
||||
#include "ggml-common.h"
|
||||
#include "ggml-backend-impl.h"
|
||||
|
||||
#include "ggml-impl.h"
|
||||
#include "ggml-cpu.h"
|
||||
#include "ggml-cpu-impl.h"
|
||||
#include "simd-mappings.h"
|
||||
#include "traits.h"
|
||||
|
||||
#include <cmath>
|
||||
#include <cstring>
|
||||
#include <cassert>
|
||||
|
||||
#define GGML_CPU_CLANG_WORKAROUND
|
||||
#include "../../repack.h"
|
||||
|
||||
#define UNUSED GGML_UNUSED
|
||||
|
||||
void ggml_quantize_mat_q8_0_4x4(const float * GGML_RESTRICT x, void * GGML_RESTRICT vy, int64_t k) {
|
||||
assert(QK8_0 == 32);
|
||||
assert(k % QK8_0 == 0);
|
||||
const int nb = k / QK8_0;
|
||||
|
||||
block_q8_0x4 * GGML_RESTRICT y = (block_q8_0x4 *) vy;
|
||||
|
||||
#if defined(__VXE__) || defined(__VXE2__)
|
||||
float32x4_t v_src[4][8];
|
||||
float id[4];
|
||||
|
||||
for (int i = 0; i < nb; i++) {
|
||||
float32x4_t v_asrc[8];
|
||||
float32x4_t v_amax[8];
|
||||
|
||||
for (int row_iter = 0; row_iter < 4; row_iter++) {
|
||||
for (int j = 0; j < 8; j++) v_src[row_iter][j] = vec_xl(0, x + row_iter * k + i * 32 + 4 * j);
|
||||
for (int j = 0; j < 8; j++) v_asrc[j] = vec_abs(v_src[row_iter][j]);
|
||||
|
||||
for (int j = 0; j < 4; j++) v_amax[2 * j] = vec_max(v_asrc[2 * j], v_asrc[2 * j + 1]);
|
||||
for (int j = 0; j < 2; j++) v_amax[4 * j] = vec_max(v_amax[4 * j], v_amax[4 * j + 2]);
|
||||
for (int j = 0; j < 1; j++) v_amax[8 * j] = vec_max(v_amax[8 * j], v_amax[8 * j + 4]);
|
||||
|
||||
const float amax = MAX(MAX(vec_extract(v_amax[0], 0), vec_extract(v_amax[0], 1)),
|
||||
MAX(vec_extract(v_amax[0], 2), vec_extract(v_amax[0], 3)));
|
||||
|
||||
const float d = amax / ((1 << 7) - 1);
|
||||
id[row_iter] = d ? 1.0f / d : 0.0f;
|
||||
|
||||
y[i].d[row_iter] = GGML_CPU_FP32_TO_FP16(d);
|
||||
}
|
||||
|
||||
for (int j = 0; j < 8; j++) {
|
||||
/* Uses non-default rounding for vec_signed or vec_round */
|
||||
const int32x4_t v_qs0 = vec_signed(__builtin_s390_vfisb(vec_mul(v_src[0][j], id[0]), 4, 1));
|
||||
const int32x4_t v_qs1 = vec_signed(__builtin_s390_vfisb(vec_mul(v_src[1][j], id[1]), 4, 1));
|
||||
const int32x4_t v_qs2 = vec_signed(__builtin_s390_vfisb(vec_mul(v_src[2][j], id[2]), 4, 1));
|
||||
const int32x4_t v_qs3 = vec_signed(__builtin_s390_vfisb(vec_mul(v_src[3][j], id[3]), 4, 1));
|
||||
|
||||
const int16x8_t v_qs01 = vec_packs(v_qs0, v_qs1);
|
||||
const int16x8_t v_qs23 = vec_packs(v_qs2, v_qs3);
|
||||
|
||||
vec_xst(vec_packs(v_qs01, v_qs23), 0, y[i].qs + 16 * j);
|
||||
}
|
||||
}
|
||||
#else
|
||||
UNUSED(nb);
|
||||
UNUSED(y);
|
||||
ggml_quantize_mat_q8_0_4x4_generic(x, vy, k);
|
||||
#endif
|
||||
}
|
||||
|
||||
static inline int16x8_t vxe_dot_acc(const int8x16_t v_x, const int8x16_t v_y, const int16x8_t v_acc) {
|
||||
return vec_meadd(v_x, v_y, vec_moadd(v_x, v_y, v_acc));
|
||||
}
|
||||
|
||||
static inline int8x16_t vxe_splat_granule(const int8_t * qs) {
|
||||
uint32_t g;
|
||||
memcpy(&g, qs, sizeof(g));
|
||||
return (int8x16_t)vec_splats(g);
|
||||
}
|
||||
|
||||
static inline int32x4_t vxe_fold(const int16x8_t v_sumi) {
|
||||
const int16x8_t v_ones = vec_splats((int16_t)1);
|
||||
return vec_add(vec_mule(v_sumi, v_ones), vec_mulo(v_sumi, v_ones));
|
||||
}
|
||||
|
||||
void ggml_gemv_q4_0_4x4_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, const void * GGML_RESTRICT vy, int nr, int nc) {
|
||||
const int qk = QK8_0;
|
||||
const int nb = n / qk;
|
||||
const int ncols_interleaved = 4;
|
||||
|
||||
assert(nr == 1);
|
||||
assert(n % qk == 0);
|
||||
assert(nc % ncols_interleaved == 0);
|
||||
|
||||
UNUSED(bs);
|
||||
UNUSED(nr);
|
||||
|
||||
#if defined(__VXE__) || defined(__VXE2__)
|
||||
const block_q8_0 * a_ptr = (const block_q8_0 *) vy;
|
||||
float * res_ptr = s;
|
||||
|
||||
for (int x = 0; x < nc / ncols_interleaved; x++) {
|
||||
const block_q4_0x4 * b_ptr = (const block_q4_0x4 *) vx + (x * nb);
|
||||
|
||||
float32x4_t v_sumf = vec_splats(0.0f);
|
||||
|
||||
for (int l = 0; l < nb; l++) {
|
||||
const int8_t * x_qs = b_ptr[l].qs;
|
||||
|
||||
const int8x16_t v_x0 = vec_xl( 0, x_qs);
|
||||
const int8x16_t v_x1 = vec_xl(16, x_qs);
|
||||
const int8x16_t v_x2 = vec_xl(32, x_qs);
|
||||
const int8x16_t v_x3 = vec_xl(48, x_qs);
|
||||
|
||||
const int8x16_t v_x0l = vec_sra(vec_sl(v_x0, 4), 4);
|
||||
const int8x16_t v_x1l = vec_sra(vec_sl(v_x1, 4), 4);
|
||||
const int8x16_t v_x2l = vec_sra(vec_sl(v_x2, 4), 4);
|
||||
const int8x16_t v_x3l = vec_sra(vec_sl(v_x3, 4), 4);
|
||||
|
||||
const int8x16_t v_x0h = vec_sra(v_x0, 4);
|
||||
const int8x16_t v_x1h = vec_sra(v_x1, 4);
|
||||
const int8x16_t v_x2h = vec_sra(v_x2, 4);
|
||||
const int8x16_t v_x3h = vec_sra(v_x3, 4);
|
||||
|
||||
const int8_t * y_lo = a_ptr[l].qs;
|
||||
const int8_t * y_hi = y_lo + qk / 2;
|
||||
|
||||
int16x8_t v_sumi = vec_splats((int16_t)0);
|
||||
|
||||
v_sumi = vxe_dot_acc(v_x0l, vxe_splat_granule(y_lo + 0), v_sumi);
|
||||
v_sumi = vxe_dot_acc(v_x1l, vxe_splat_granule(y_lo + 4), v_sumi);
|
||||
v_sumi = vxe_dot_acc(v_x2l, vxe_splat_granule(y_lo + 8), v_sumi);
|
||||
v_sumi = vxe_dot_acc(v_x3l, vxe_splat_granule(y_lo + 12), v_sumi);
|
||||
|
||||
v_sumi = vxe_dot_acc(v_x0h, vxe_splat_granule(y_hi + 0), v_sumi);
|
||||
v_sumi = vxe_dot_acc(v_x1h, vxe_splat_granule(y_hi + 4), v_sumi);
|
||||
v_sumi = vxe_dot_acc(v_x2h, vxe_splat_granule(y_hi + 8), v_sumi);
|
||||
v_sumi = vxe_dot_acc(v_x3h, vxe_splat_granule(y_hi + 12), v_sumi);
|
||||
|
||||
const float32x4_t v_yd = vec_splats(GGML_CPU_FP16_TO_FP32(a_ptr[l].d));
|
||||
const float32x4_t v_xd = __lzs_f16cx4_load(b_ptr[l].d);
|
||||
const float32x4_t v_d = vec_mul(v_yd, v_xd);
|
||||
|
||||
v_sumf = vec_madd(vec_float(vxe_fold(v_sumi)), v_d, v_sumf);
|
||||
}
|
||||
|
||||
vec_xst(v_sumf, 0, res_ptr + x * ncols_interleaved);
|
||||
}
|
||||
#else
|
||||
UNUSED(nb);
|
||||
UNUSED(ncols_interleaved);
|
||||
ggml_gemv_q4_0_4x4_q8_0_generic(n, s, bs, vx, vy, nr, nc);
|
||||
#endif
|
||||
}
|
||||
|
||||
void ggml_gemm_q4_0_4x4_q8_0(int n, float * GGML_RESTRICT s, size_t bs, const void * GGML_RESTRICT vx, const void * GGML_RESTRICT vy, int nr, int nc) {
|
||||
const int qk = QK8_0;
|
||||
const int nb = n / qk;
|
||||
const int ncols_interleaved = 4;
|
||||
|
||||
assert(nr % 4 == 0);
|
||||
assert(n % qk == 0);
|
||||
assert(nc % ncols_interleaved == 0);
|
||||
|
||||
#if defined(__VXE__) || defined(__VXE2__)
|
||||
for (int y = 0; y < nr / 4; y++) {
|
||||
const block_q8_0x4 * a_ptr = (const block_q8_0x4 *) vy + (y * nb);
|
||||
|
||||
for (int x = 0; x < nc / ncols_interleaved; x++) {
|
||||
const block_q4_0x4 * b_ptr = (const block_q4_0x4 *) vx + (x * nb);
|
||||
|
||||
float32x4_t v_sumf[4];
|
||||
for (int m = 0; m < 4; m++) {
|
||||
v_sumf[m] = vec_splats(0.0f);
|
||||
}
|
||||
|
||||
for (int l = 0; l < nb; l++) {
|
||||
int16x8_t v_sumi0 = vec_splats((int16_t)0);
|
||||
int16x8_t v_sumi1 = vec_splats((int16_t)0);
|
||||
int16x8_t v_sumi2 = vec_splats((int16_t)0);
|
||||
int16x8_t v_sumi3 = vec_splats((int16_t)0);
|
||||
|
||||
for (int k = 0; k < 4; k++) {
|
||||
const int8x16_t v_x = vec_xl(0, b_ptr[l].qs + 16 * k);
|
||||
const int8x16_t v_xl = vec_sra(vec_sl(v_x, 4), 4);
|
||||
const int8x16_t v_xh = vec_sra(v_x, 4);
|
||||
|
||||
const int8_t * y_lo = a_ptr[l].qs + 16 * k;
|
||||
const int8_t * y_hi = y_lo + qk / 2 * 4;
|
||||
|
||||
v_sumi0 = vxe_dot_acc(v_xl, vxe_splat_granule(y_lo + 0), v_sumi0);
|
||||
v_sumi1 = vxe_dot_acc(v_xl, vxe_splat_granule(y_lo + 4), v_sumi1);
|
||||
v_sumi2 = vxe_dot_acc(v_xl, vxe_splat_granule(y_lo + 8), v_sumi2);
|
||||
v_sumi3 = vxe_dot_acc(v_xl, vxe_splat_granule(y_lo + 12), v_sumi3);
|
||||
|
||||
v_sumi0 = vxe_dot_acc(v_xh, vxe_splat_granule(y_hi + 0), v_sumi0);
|
||||
v_sumi1 = vxe_dot_acc(v_xh, vxe_splat_granule(y_hi + 4), v_sumi1);
|
||||
v_sumi2 = vxe_dot_acc(v_xh, vxe_splat_granule(y_hi + 8), v_sumi2);
|
||||
v_sumi3 = vxe_dot_acc(v_xh, vxe_splat_granule(y_hi + 12), v_sumi3);
|
||||
}
|
||||
|
||||
const float32x4_t v_yd = __lzs_f16cx4_load(a_ptr[l].d);
|
||||
const float32x4_t v_xd = __lzs_f16cx4_load(b_ptr[l].d);
|
||||
|
||||
v_sumf[0] = vec_madd(vec_float(vxe_fold(v_sumi0)), vec_mul(v_xd, vec_splat(v_yd, 0)), v_sumf[0]);
|
||||
v_sumf[1] = vec_madd(vec_float(vxe_fold(v_sumi1)), vec_mul(v_xd, vec_splat(v_yd, 1)), v_sumf[1]);
|
||||
v_sumf[2] = vec_madd(vec_float(vxe_fold(v_sumi2)), vec_mul(v_xd, vec_splat(v_yd, 2)), v_sumf[2]);
|
||||
v_sumf[3] = vec_madd(vec_float(vxe_fold(v_sumi3)), vec_mul(v_xd, vec_splat(v_yd, 3)), v_sumf[3]);
|
||||
}
|
||||
|
||||
for (int m = 0; m < 4; m++) {
|
||||
vec_xst(v_sumf[m], 0, s + (y * 4 + m) * bs + x * ncols_interleaved);
|
||||
}
|
||||
}
|
||||
}
|
||||
#else
|
||||
UNUSED(nb);
|
||||
UNUSED(ncols_interleaved);
|
||||
ggml_gemm_q4_0_4x4_q8_0_generic(n, s, bs, vx, vy, nr, nc);
|
||||
#endif
|
||||
}
|
||||
+169
-99
@@ -5,6 +5,64 @@
|
||||
#include "fattn-vec.cuh"
|
||||
#include "fattn.cuh"
|
||||
|
||||
// kcpp: mirror the CMake-generated FlashAttention defines for our fixed template instances.
|
||||
// Keep enabled pairs in sync with Makefile and CMakeLists.txt.
|
||||
#define GGML_CUDA_FA_F16_F16 1
|
||||
#define GGML_CUDA_FA_Q4_0_F16 0
|
||||
#define GGML_CUDA_FA_Q4_1_F16 0
|
||||
#define GGML_CUDA_FA_Q5_0_F16 0
|
||||
#define GGML_CUDA_FA_Q5_1_F16 0
|
||||
#define GGML_CUDA_FA_Q8_0_F16 0
|
||||
#define GGML_CUDA_FA_BF16_F16 0
|
||||
|
||||
#define GGML_CUDA_FA_F16_Q4_0 0
|
||||
#define GGML_CUDA_FA_Q4_0_Q4_0 1
|
||||
#define GGML_CUDA_FA_Q4_1_Q4_0 0
|
||||
#define GGML_CUDA_FA_Q5_0_Q4_0 0
|
||||
#define GGML_CUDA_FA_Q5_1_Q4_0 0
|
||||
#define GGML_CUDA_FA_Q8_0_Q4_0 0
|
||||
#define GGML_CUDA_FA_BF16_Q4_0 0
|
||||
|
||||
#define GGML_CUDA_FA_F16_Q4_1 0
|
||||
#define GGML_CUDA_FA_Q4_0_Q4_1 0
|
||||
#define GGML_CUDA_FA_Q4_1_Q4_1 0
|
||||
#define GGML_CUDA_FA_Q5_0_Q4_1 0
|
||||
#define GGML_CUDA_FA_Q5_1_Q4_1 0
|
||||
#define GGML_CUDA_FA_Q8_0_Q4_1 0
|
||||
#define GGML_CUDA_FA_BF16_Q4_1 0
|
||||
|
||||
#define GGML_CUDA_FA_F16_Q5_0 0
|
||||
#define GGML_CUDA_FA_Q4_0_Q5_0 0
|
||||
#define GGML_CUDA_FA_Q4_1_Q5_0 0
|
||||
#define GGML_CUDA_FA_Q5_0_Q5_0 0
|
||||
#define GGML_CUDA_FA_Q5_1_Q5_0 0
|
||||
#define GGML_CUDA_FA_Q8_0_Q5_0 0
|
||||
#define GGML_CUDA_FA_BF16_Q5_0 0
|
||||
|
||||
#define GGML_CUDA_FA_F16_Q5_1 0
|
||||
#define GGML_CUDA_FA_Q4_0_Q5_1 0
|
||||
#define GGML_CUDA_FA_Q4_1_Q5_1 0
|
||||
#define GGML_CUDA_FA_Q5_0_Q5_1 0
|
||||
#define GGML_CUDA_FA_Q5_1_Q5_1 1
|
||||
#define GGML_CUDA_FA_Q8_0_Q5_1 0
|
||||
#define GGML_CUDA_FA_BF16_Q5_1 0
|
||||
|
||||
#define GGML_CUDA_FA_F16_Q8_0 0
|
||||
#define GGML_CUDA_FA_Q4_0_Q8_0 0
|
||||
#define GGML_CUDA_FA_Q4_1_Q8_0 0
|
||||
#define GGML_CUDA_FA_Q5_0_Q8_0 0
|
||||
#define GGML_CUDA_FA_Q5_1_Q8_0 0
|
||||
#define GGML_CUDA_FA_Q8_0_Q8_0 1
|
||||
#define GGML_CUDA_FA_BF16_Q8_0 0
|
||||
|
||||
#define GGML_CUDA_FA_F16_BF16 0
|
||||
#define GGML_CUDA_FA_Q4_0_BF16 0
|
||||
#define GGML_CUDA_FA_Q4_1_BF16 0
|
||||
#define GGML_CUDA_FA_Q5_0_BF16 0
|
||||
#define GGML_CUDA_FA_Q5_1_BF16 0
|
||||
#define GGML_CUDA_FA_Q8_0_BF16 0
|
||||
#define GGML_CUDA_FA_BF16_BF16 1
|
||||
|
||||
#if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
|
||||
__launch_bounds__(256, 1)
|
||||
static __global__ void flash_attn_mask_to_sparse_indices(
|
||||
@@ -374,91 +432,101 @@ static void ggml_cuda_flash_attn_ext_mma_f16(ggml_backend_cuda_context & ctx, gg
|
||||
}
|
||||
}
|
||||
|
||||
#define FATTN_VEC_CASE(D, type_K, type_V) \
|
||||
{ \
|
||||
const bool type_K_okay = K->type == (type_K) || (K->type == GGML_TYPE_F32 && (type_K) == GGML_TYPE_F16); \
|
||||
const bool type_V_okay = V->type == (type_V) || (V->type == GGML_TYPE_F32 && (type_V) == GGML_TYPE_F16); \
|
||||
if (Q->ne[0] == (D) && type_K_okay && type_V_okay) { \
|
||||
ggml_cuda_flash_attn_ext_vec_case<D, type_K, type_V>(ctx, dst); \
|
||||
return; \
|
||||
} \
|
||||
} \
|
||||
#define FATTN_VEC_CASE(D, type_K_case, type_V_case) \
|
||||
if constexpr (GGML_CUDA_FA_##type_K_case##_##type_V_case) { \
|
||||
const bool type_K_okay = type_K == GGML_TYPE_##type_K_case || (type_K == GGML_TYPE_F32 && GGML_TYPE_##type_K_case == GGML_TYPE_F16); \
|
||||
const bool type_V_okay = type_V == GGML_TYPE_##type_V_case || (type_V == GGML_TYPE_F32 && GGML_TYPE_##type_V_case == GGML_TYPE_F16); \
|
||||
if (head_size == (D) && type_K_okay && type_V_okay) { \
|
||||
return ggml_cuda_flash_attn_ext_vec_case<D, GGML_TYPE_##type_K_case, GGML_TYPE_##type_V_case>; \
|
||||
} \
|
||||
} \
|
||||
|
||||
#define FATTN_VEC_CASES_ALL_D(type_K, type_V) \
|
||||
FATTN_VEC_CASE( 64, type_K, type_V) \
|
||||
FATTN_VEC_CASE(128, type_K, type_V) \
|
||||
FATTN_VEC_CASE(256, type_K, type_V) \
|
||||
#define FATTN_VEC_CASES_ALL_D(type_K_case, type_V_case) \
|
||||
FATTN_VEC_CASE( 64, type_K_case, type_V_case) \
|
||||
FATTN_VEC_CASE(128, type_K_case, type_V_case) \
|
||||
FATTN_VEC_CASE(256, type_K_case, type_V_case) \
|
||||
|
||||
typedef void (* fattn_vec_case_t)(ggml_backend_cuda_context & ctx, ggml_tensor * dst);
|
||||
|
||||
// Vector kernel for the given head size and K/V types, nullptr if its template instance was not compiled:
|
||||
static fattn_vec_case_t ggml_cuda_get_fattn_vec_case(const int64_t head_size, const ggml_type type_K, const ggml_type type_V) {
|
||||
FATTN_VEC_CASES_ALL_D(F16, F16)
|
||||
FATTN_VEC_CASES_ALL_D(Q4_0, F16)
|
||||
FATTN_VEC_CASES_ALL_D(Q4_1, F16)
|
||||
FATTN_VEC_CASES_ALL_D(Q5_0, F16)
|
||||
FATTN_VEC_CASES_ALL_D(Q5_1, F16)
|
||||
FATTN_VEC_CASES_ALL_D(Q8_0, F16)
|
||||
FATTN_VEC_CASES_ALL_D(BF16, F16)
|
||||
|
||||
FATTN_VEC_CASES_ALL_D(F16, Q4_0)
|
||||
FATTN_VEC_CASES_ALL_D(Q4_0, Q4_0)
|
||||
FATTN_VEC_CASES_ALL_D(Q4_1, Q4_0)
|
||||
FATTN_VEC_CASES_ALL_D(Q5_0, Q4_0)
|
||||
FATTN_VEC_CASES_ALL_D(Q5_1, Q4_0)
|
||||
FATTN_VEC_CASES_ALL_D(Q8_0, Q4_0)
|
||||
FATTN_VEC_CASES_ALL_D(BF16, Q4_0)
|
||||
|
||||
FATTN_VEC_CASES_ALL_D(F16, Q4_1)
|
||||
FATTN_VEC_CASES_ALL_D(Q4_0, Q4_1)
|
||||
FATTN_VEC_CASES_ALL_D(Q4_1, Q4_1)
|
||||
FATTN_VEC_CASES_ALL_D(Q5_0, Q4_1)
|
||||
FATTN_VEC_CASES_ALL_D(Q5_1, Q4_1)
|
||||
FATTN_VEC_CASES_ALL_D(Q8_0, Q4_1)
|
||||
FATTN_VEC_CASES_ALL_D(BF16, Q4_1)
|
||||
|
||||
FATTN_VEC_CASES_ALL_D(F16, Q5_0)
|
||||
FATTN_VEC_CASES_ALL_D(Q4_0, Q5_0)
|
||||
FATTN_VEC_CASES_ALL_D(Q4_1, Q5_0)
|
||||
FATTN_VEC_CASES_ALL_D(Q5_0, Q5_0)
|
||||
FATTN_VEC_CASES_ALL_D(Q5_1, Q5_0)
|
||||
FATTN_VEC_CASES_ALL_D(Q8_0, Q5_0)
|
||||
FATTN_VEC_CASES_ALL_D(BF16, Q5_0)
|
||||
|
||||
FATTN_VEC_CASES_ALL_D(F16, Q5_1)
|
||||
FATTN_VEC_CASES_ALL_D(Q4_0, Q5_1)
|
||||
FATTN_VEC_CASES_ALL_D(Q4_1, Q5_1)
|
||||
FATTN_VEC_CASES_ALL_D(Q5_0, Q5_1)
|
||||
FATTN_VEC_CASES_ALL_D(Q5_1, Q5_1)
|
||||
FATTN_VEC_CASES_ALL_D(Q8_0, Q5_1)
|
||||
FATTN_VEC_CASES_ALL_D(BF16, Q5_1)
|
||||
|
||||
FATTN_VEC_CASES_ALL_D(F16, Q8_0)
|
||||
FATTN_VEC_CASES_ALL_D(Q4_0, Q8_0)
|
||||
FATTN_VEC_CASES_ALL_D(Q4_1, Q8_0)
|
||||
FATTN_VEC_CASES_ALL_D(Q5_0, Q8_0)
|
||||
FATTN_VEC_CASES_ALL_D(Q5_1, Q8_0)
|
||||
FATTN_VEC_CASES_ALL_D(Q8_0, Q8_0)
|
||||
FATTN_VEC_CASES_ALL_D(BF16, Q8_0)
|
||||
|
||||
FATTN_VEC_CASES_ALL_D(F16, BF16)
|
||||
FATTN_VEC_CASES_ALL_D(Q4_0, BF16)
|
||||
FATTN_VEC_CASES_ALL_D(Q4_1, BF16)
|
||||
FATTN_VEC_CASES_ALL_D(Q5_0, BF16)
|
||||
FATTN_VEC_CASES_ALL_D(Q5_1, BF16)
|
||||
FATTN_VEC_CASES_ALL_D(Q8_0, BF16)
|
||||
FATTN_VEC_CASES_ALL_D(BF16, BF16)
|
||||
|
||||
return nullptr;
|
||||
}
|
||||
|
||||
static void ggml_cuda_flash_attn_ext_vec(ggml_backend_cuda_context & ctx, ggml_tensor * dst) {
|
||||
ggml_tensor * Q = dst->src[0];
|
||||
ggml_tensor * K = dst->src[1];
|
||||
ggml_tensor * V = dst->src[2];
|
||||
const ggml_tensor * Q = dst->src[0];
|
||||
const ggml_tensor * K = dst->src[1];
|
||||
const ggml_tensor * V = dst->src[2];
|
||||
|
||||
#ifdef GGML_CUDA_FA_ALL_QUANTS
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16, GGML_TYPE_F16)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_F16)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_1, GGML_TYPE_F16)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_0, GGML_TYPE_F16)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_F16)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_F16)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_F16)
|
||||
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16, GGML_TYPE_Q4_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q4_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_1, GGML_TYPE_Q4_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_0, GGML_TYPE_Q4_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_Q4_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q4_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_Q4_0)
|
||||
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16, GGML_TYPE_Q4_1)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q4_1)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_1, GGML_TYPE_Q4_1)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_0, GGML_TYPE_Q4_1)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_Q4_1)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q4_1)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_Q4_1)
|
||||
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16, GGML_TYPE_Q5_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q5_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_1, GGML_TYPE_Q5_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_0, GGML_TYPE_Q5_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_Q5_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q5_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_Q5_0)
|
||||
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16, GGML_TYPE_Q5_1)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q5_1)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_1, GGML_TYPE_Q5_1)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_0, GGML_TYPE_Q5_1)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_Q5_1)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q5_1)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_Q5_1)
|
||||
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16, GGML_TYPE_Q8_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q8_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_1, GGML_TYPE_Q8_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_0, GGML_TYPE_Q8_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_Q8_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q8_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_Q8_0)
|
||||
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16, GGML_TYPE_BF16)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_BF16)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_1, GGML_TYPE_BF16)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_0, GGML_TYPE_BF16)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_BF16)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_BF16)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_BF16)
|
||||
#else
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_F16, GGML_TYPE_F16)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q4_0, GGML_TYPE_Q4_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q5_1, GGML_TYPE_Q5_1)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_Q8_0, GGML_TYPE_Q8_0)
|
||||
FATTN_VEC_CASES_ALL_D(GGML_TYPE_BF16, GGML_TYPE_BF16)
|
||||
#endif // GGML_CUDA_FA_ALL_QUANTS
|
||||
|
||||
GGML_ABORT("fatal error");
|
||||
fattn_vec_case_t vec_case = ggml_cuda_get_fattn_vec_case(Q->ne[0], K->type, V->type);
|
||||
if (vec_case == nullptr) {
|
||||
static bool warned = false;
|
||||
if (!warned) {
|
||||
GGML_LOG_WARN("%s: no FlashAttention vector kernel compiled for K/V types %s-%s, converting K and V to f16 instead (slow). "
|
||||
"Add \"%s-%s\" to GGML_CUDA_FA_QUANTS to compile it.\n",
|
||||
__func__, ggml_type_name(K->type), ggml_type_name(V->type), ggml_type_name(K->type), ggml_type_name(V->type));
|
||||
warned = true;
|
||||
}
|
||||
vec_case = ggml_cuda_get_fattn_vec_case(Q->ne[0], GGML_TYPE_F16, GGML_TYPE_F16);
|
||||
}
|
||||
GGML_ASSERT(vec_case != nullptr);
|
||||
vec_case(ctx, dst);
|
||||
}
|
||||
|
||||
// Best FlashAttention kernel for a specific GPU:
|
||||
@@ -469,20 +537,17 @@ enum best_fattn_kernel {
|
||||
BEST_FATTN_KERNEL_MMA_F16 = 400,
|
||||
};
|
||||
|
||||
static bool ggml_cuda_fattn_kv_type_supported(ggml_type type) {
|
||||
// K/V types for which there is a vector kernel template instance, other kernels convert these to f16:
|
||||
static bool ggml_cuda_fattn_kv_type_supported(const ggml_type type) {
|
||||
switch (type) {
|
||||
case GGML_TYPE_F32:
|
||||
case GGML_TYPE_F16:
|
||||
return true;
|
||||
case GGML_TYPE_BF16:
|
||||
case GGML_TYPE_Q4_0:
|
||||
case GGML_TYPE_Q4_1:
|
||||
case GGML_TYPE_Q5_0:
|
||||
#ifndef GGML_CUDA_FA_ALL_QUANTS
|
||||
return false;
|
||||
#endif // GGML_CUDA_FA_ALL_QUANTS
|
||||
case GGML_TYPE_Q4_0:
|
||||
case GGML_TYPE_Q5_1: // kcpp: support q5_1 kv
|
||||
case GGML_TYPE_Q5_1:
|
||||
case GGML_TYPE_Q8_0:
|
||||
case GGML_TYPE_BF16:
|
||||
return true;
|
||||
default:
|
||||
return false;
|
||||
@@ -573,12 +638,6 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const
|
||||
return BEST_FATTN_KERNEL_NONE;
|
||||
}
|
||||
|
||||
#ifndef GGML_CUDA_FA_ALL_QUANTS
|
||||
if (K->type != V->type) {
|
||||
return BEST_FATTN_KERNEL_NONE;
|
||||
}
|
||||
#endif // GGML_CUDA_FA_ALL_QUANTS
|
||||
|
||||
if (!ggml_cuda_fattn_kv_type_supported(K->type) || !ggml_cuda_fattn_kv_type_supported(V->type)) {
|
||||
return BEST_FATTN_KERNEL_NONE;
|
||||
}
|
||||
@@ -622,6 +681,15 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const
|
||||
gqa_ratio_eff *= 2;
|
||||
}
|
||||
|
||||
if (volta_mma_available(cc) && Q->ne[0] != 40 && Q->ne[0] != 72) {
|
||||
if (can_use_vector_kernel && Q->ne[1] * gqa_ratio_eff <= 2) {
|
||||
return BEST_FATTN_KERNEL_VEC;
|
||||
}
|
||||
if (Q->ne[1] * gqa_ratio_eff <= 16) {
|
||||
return BEST_FATTN_KERNEL_TILE; // On Volta tensor cores are only faster for sufficiently large matrices.
|
||||
}
|
||||
return BEST_FATTN_KERNEL_MMA_F16;
|
||||
}
|
||||
|
||||
// AMD MFMA needs a certain minimum batch size to outscale the tile kernel for large head sizes.
|
||||
if ((amd_mfma_available(cc) && Q->ne[0] <= 256) && Q->ne[0] != 40 && Q->ne[0] != 72) {
|
||||
@@ -661,6 +729,7 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const
|
||||
size_t ggml_cuda_flash_attn_ext_get_alloc_size(int device, const ggml_tensor * dst) {
|
||||
GGML_ASSERT(dst->op == GGML_OP_FLASH_ATTN_EXT);
|
||||
|
||||
const ggml_tensor * Q = dst->src[0];
|
||||
const ggml_tensor * K = dst->src[1];
|
||||
const ggml_tensor * V = dst->src[2];
|
||||
|
||||
@@ -678,10 +747,11 @@ size_t ggml_cuda_flash_attn_ext_get_alloc_size(int device, const ggml_tensor * d
|
||||
need_f16_K = true;
|
||||
need_f16_V = true;
|
||||
break;
|
||||
case BEST_FATTN_KERNEL_VEC:
|
||||
need_f16_K = K->type == GGML_TYPE_F32;
|
||||
need_f16_V = V->type == GGML_TYPE_F32;
|
||||
break;
|
||||
case BEST_FATTN_KERNEL_VEC: {
|
||||
const bool f16_fallback = ggml_cuda_get_fattn_vec_case(Q->ne[0], K->type, V->type) == nullptr;
|
||||
need_f16_K = K->type == GGML_TYPE_F32 || f16_fallback;
|
||||
need_f16_V = V->type == GGML_TYPE_F32 || f16_fallback;
|
||||
} break;
|
||||
case BEST_FATTN_KERNEL_NONE:
|
||||
break;
|
||||
}
|
||||
|
||||
@@ -5654,8 +5654,8 @@ static ggml_backend_feature * ggml_backend_cuda_get_features(ggml_backend_reg_t
|
||||
features.push_back({ "USE_GRAPHS", "1" });
|
||||
#endif
|
||||
|
||||
#ifdef GGML_CUDA_FA_ALL_QUANTS
|
||||
features.push_back({ "FA_ALL_QUANTS", "1" });
|
||||
#ifdef GGML_CUDA_FA_QUANTS
|
||||
features.push_back({ "FA_QUANTS", GGML_CUDA_FA_QUANTS });
|
||||
#endif
|
||||
|
||||
{
|
||||
|
||||
@@ -171,7 +171,7 @@ void ggml_cuda_mul_mat_q(
|
||||
ne00, ne01, ne1, s01, ne11, s1,
|
||||
ne02, ne12, s02, s12, s2,
|
||||
ne03, ne13, s03, s13, s3,
|
||||
ne1};
|
||||
ne1, ne1};
|
||||
ggml_cuda_mul_mat_q_switch_type(ctx, args, stream);
|
||||
return;
|
||||
}
|
||||
@@ -244,6 +244,13 @@ void ggml_cuda_mul_mat_q(
|
||||
ne11 * ne10_padded * sizeof(block_q8_1) / (QK8_1 * sizeof(int));
|
||||
const int64_t s13 = ne12*s12;
|
||||
|
||||
// Each expert only sees ne12*n_expert_used/ne02 tokens on average.
|
||||
// On RDNA3 and RDNA4 it is faster to pick the tile size against this value instead of ne12.
|
||||
int64_t ncols_opt = ne12;
|
||||
if (GGML_CUDA_CC_IS_RDNA3_0(cc) || GGML_CUDA_CC_IS_RDNA4(cc)) {
|
||||
ncols_opt = (ne12*n_expert_used + ne02 - 1) / ne02;
|
||||
}
|
||||
|
||||
// Note that ne02 is used instead of ne12 because the number of y channels determines the z dimension of the CUDA grid.
|
||||
const mmq_args args = {
|
||||
src0_d, src0->type, (const int *) src1_q8_1.get(), ids_dst.get(), expert_bounds.get(), dst_d,
|
||||
@@ -251,7 +258,7 @@ void ggml_cuda_mul_mat_q(
|
||||
ne00, ne01, ne_get_rows, s01, ne_get_rows, s1,
|
||||
ne02, ne02, s02, s12, s2,
|
||||
ne03, ne13, s03, s13, s3,
|
||||
ne12};
|
||||
ne12, ncols_opt};
|
||||
|
||||
ggml_cuda_mul_mat_q_switch_type(ctx, args, stream);
|
||||
}
|
||||
|
||||
@@ -1377,6 +1377,7 @@ struct mmq_args {
|
||||
int64_t nchannels_x; int64_t nchannels_y; int64_t stride_channel_x; int64_t stride_channel_y; int64_t stride_channel_dst;
|
||||
int64_t nsamples_x; int64_t nsamples_y; int64_t stride_sample_x; int64_t stride_sample_y; int64_t stride_sample_dst;
|
||||
int64_t ncols_max;
|
||||
int64_t ncols_opt; // value to optimize the tile size against, launch grid still uses ncols_max
|
||||
};
|
||||
|
||||
static size_t mmq_get_nbytes_shared(const ggml_cuda_mmq_config & config, const int cc) {
|
||||
@@ -1487,7 +1488,7 @@ void mul_mat_q_switch_J(ggml_backend_cuda_context & ctx, const mmq_args & args,
|
||||
continue;
|
||||
}
|
||||
|
||||
const int ntiles_x = (args.ncols_max + config.J - 1) / config.J;
|
||||
const int ntiles_x = (args.ncols_opt + config.J - 1) / config.J;
|
||||
|
||||
if (ntiles_x < ntiles_J_best) {
|
||||
J_best = J;
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@@ -260,6 +260,20 @@ float16_t dequantFuncTQ1_0(const in decodeBufTQ1_0 bl, const in uint blockCoords
|
||||
return bl.block.d * (float16_t(int(xi)) - float16_t(1.0));
|
||||
}
|
||||
|
||||
f16vec4 dequantFuncTQ1_0_v(const in decodeBufTQ1_0 bl, const in uint blockCoords[2], const in uint coordInBlock[2])
|
||||
{
|
||||
const uint e = coordInBlock[1];
|
||||
f16vec4 v;
|
||||
[[unroll]] for (uint k = 0u; k < 4u; ++k) {
|
||||
const uint ee = e + k;
|
||||
const uint bidx = tq1_0_byte_of(ee);
|
||||
const uint qbyte = uint(bidx < 48u ? bl.block.qs[bidx] : bl.block.qh[bidx - 48u]);
|
||||
const uint xi = tq1_0_trit(qbyte, tq1_0_digit_of(ee));
|
||||
v[k] = bl.block.d * (float16_t(int(xi)) - float16_t(1.0));
|
||||
}
|
||||
return v;
|
||||
}
|
||||
|
||||
layout(buffer_reference, std430, buffer_reference_align = 2) buffer decodeBufTQ2_0 {
|
||||
block_tq2_0 block;
|
||||
};
|
||||
@@ -1054,7 +1068,7 @@ float16_t dequantFuncIQ2_S(const in decodeBufIQ2_S bl, const in uint blockCoords
|
||||
const uint scale = (bl.block.scales[ib32] >> ((idx & 0x10) >> 2)) & 0xf;
|
||||
const uint qs = bl.block.qs[ib8];
|
||||
const uint qh = bl.block.qh[ib32];
|
||||
const uint sign = bl.block.qs[QUANT_K / 8 + ib8] >> (idx & 0x6);
|
||||
const uint sign = bl.block.qs[QUANT_K_IQ2_S / 8 + ib8] >> (idx & 0x6);
|
||||
|
||||
const float d = float(bl.block.d);
|
||||
const float db = d * 0.25 * (0.5 + scale);
|
||||
@@ -1076,7 +1090,7 @@ f16vec4 dequantFuncIQ2_S_v(const in decodeBufIQ2_S bl, const in uint blockCoords
|
||||
const uint scale = (bl.block.scales[ib32] >> ((idx & 0x10) >> 2)) & 0xf;
|
||||
const uint qs = bl.block.qs[ib8];
|
||||
const uint qh = bl.block.qh[ib32];
|
||||
const uint sb = uint(bl.block.qs[QUANT_K / 8 + ib8]) >> (idx & 0x6u);
|
||||
const uint sb = uint(bl.block.qs[QUANT_K_IQ2_S / 8 + ib8]) >> (idx & 0x6u);
|
||||
|
||||
const float d = float(bl.block.d);
|
||||
const float db = d * 0.25 * (0.5 + scale);
|
||||
@@ -1107,7 +1121,7 @@ float16_t dequantFuncIQ3_XXS(const in decodeBufIQ3_XXS bl, const in uint blockCo
|
||||
uint idx = coordInBlock[1];
|
||||
|
||||
const uint iqs = (idx & 0xFC) >> 2; // 0..63
|
||||
const uint is = QUANT_K / 4 + ((idx & 0xE0) >> 3);// 8 values
|
||||
const uint is = QUANT_K_IQ3_XXS / 4 + ((idx & 0xE0) >> 3);// 8 values
|
||||
|
||||
const float d = float(bl.block.d);
|
||||
const uint qs = bl.block.qs[iqs];
|
||||
@@ -1130,7 +1144,7 @@ f16vec4 dequantFuncIQ3_XXS_v(const in decodeBufIQ3_XXS bl, const in uint blockCo
|
||||
const uint idx = coordInBlock[1];
|
||||
|
||||
const uint iqs = idx >> 2;
|
||||
const uint is = QUANT_K / 4 + ((idx & 0xE0) >> 3);
|
||||
const uint is = QUANT_K_IQ3_XXS / 4 + ((idx & 0xE0) >> 3);
|
||||
|
||||
const float d = float(bl.block.d);
|
||||
const uint qs = bl.block.qs[iqs];
|
||||
|
||||
@@ -1,32 +1,22 @@
|
||||
#if !defined(GGML_FA_TYPES_COMP)
|
||||
#define GGML_FA_TYPES_COMP
|
||||
|
||||
// FaTypeK / FaTypeV spec constant values. These mirror enum ggml_type so the
|
||||
// host can pass the type directly. Keep in sync with ggml.h.
|
||||
#define FA_TYPE_F32 0u
|
||||
#define FA_TYPE_F16 1u
|
||||
#define FA_TYPE_Q4_0 2u
|
||||
#define FA_TYPE_Q4_1 3u
|
||||
#define FA_TYPE_Q5_0 6u
|
||||
#define FA_TYPE_Q5_1 7u
|
||||
#define FA_TYPE_Q8_0 8u
|
||||
#define FA_TYPE_IQ4_NL 20u
|
||||
#define FA_TYPE_BF16 30u
|
||||
#include "ggml_type_ids.glsl"
|
||||
|
||||
// Number of matrix elements per buffer block, derived from the K/V type spec
|
||||
// constant. F32 is treated as a vec4 "block" of 4 floats. F16 uses block size 1
|
||||
// and bypasses the dequant path entirely. Quants follow their ggml block sizes.
|
||||
uint fa_block_elems(uint ty) {
|
||||
switch (ty) {
|
||||
case FA_TYPE_F32: return 4u;
|
||||
case FA_TYPE_F16: return 1u;
|
||||
case FA_TYPE_Q4_0: return uint(QUANT_K_Q4_0);
|
||||
case FA_TYPE_Q4_1: return uint(QUANT_K_Q4_1);
|
||||
case FA_TYPE_Q5_0: return uint(QUANT_K_Q5_0);
|
||||
case FA_TYPE_Q5_1: return uint(QUANT_K_Q5_1);
|
||||
case FA_TYPE_Q8_0: return uint(QUANT_K_Q8_0);
|
||||
case FA_TYPE_IQ4_NL: return uint(QUANT_K_IQ4_NL);
|
||||
case FA_TYPE_BF16: return 1u;
|
||||
case GGML_TYPE_F32: return 4u;
|
||||
case GGML_TYPE_F16: return 1u;
|
||||
case GGML_TYPE_Q4_0: return uint(QUANT_K_Q4_0);
|
||||
case GGML_TYPE_Q4_1: return uint(QUANT_K_Q4_1);
|
||||
case GGML_TYPE_Q5_0: return uint(QUANT_K_Q5_0);
|
||||
case GGML_TYPE_Q5_1: return uint(QUANT_K_Q5_1);
|
||||
case GGML_TYPE_Q8_0: return uint(QUANT_K_Q8_0);
|
||||
case GGML_TYPE_IQ4_NL: return uint(QUANT_K_IQ4_NL);
|
||||
case GGML_TYPE_BF16: return 1u;
|
||||
default: return 1u;
|
||||
}
|
||||
}
|
||||
@@ -36,18 +26,18 @@ uint fa_block_elems(uint ty) {
|
||||
// of int32s per 32-element block on the MMQ K path: ints_per_block == 8 / R.
|
||||
uint fa_quant_r_mmq(uint ty) {
|
||||
switch (ty) {
|
||||
case FA_TYPE_Q4_0: return uint(QUANT_R_Q4_0);
|
||||
case FA_TYPE_Q4_1: return uint(QUANT_R_Q4_1);
|
||||
case FA_TYPE_Q5_0: return uint(QUANT_R_Q5_0);
|
||||
case FA_TYPE_Q5_1: return uint(QUANT_R_Q5_1);
|
||||
case FA_TYPE_Q8_0: return uint(QUANT_R_Q8_0);
|
||||
case GGML_TYPE_Q4_0: return uint(QUANT_R_Q4_0);
|
||||
case GGML_TYPE_Q4_1: return uint(QUANT_R_Q4_1);
|
||||
case GGML_TYPE_Q5_0: return uint(QUANT_R_Q5_0);
|
||||
case GGML_TYPE_Q5_1: return uint(QUANT_R_Q5_1);
|
||||
case GGML_TYPE_Q8_0: return uint(QUANT_R_Q8_0);
|
||||
default: return 1u;
|
||||
}
|
||||
}
|
||||
|
||||
bool fa_type_needs_shmem(uint ty) {
|
||||
switch (ty) {
|
||||
case FA_TYPE_IQ4_NL: return true;
|
||||
case GGML_TYPE_IQ4_NL: return true;
|
||||
default: return false;
|
||||
}
|
||||
}
|
||||
|
||||
@@ -8,7 +8,9 @@ layout(local_size_x = 512, local_size_y = 1, local_size_z = 1) in;
|
||||
layout (binding = 0) writeonly buffer D {D_TYPE data_d[];};
|
||||
|
||||
void main() {
|
||||
const uint i = gl_GlobalInvocationID.x;
|
||||
// 2D grid flattening: each x workgroup covers gl_WorkGroupSize.x elements,
|
||||
// each y workgroup covers gl_NumWorkGroups.x * gl_WorkGroupSize.x elements.
|
||||
const uint i = (gl_GlobalInvocationID.y * gl_NumWorkGroups.x * gl_WorkGroupSize.x) + gl_GlobalInvocationID.x;
|
||||
|
||||
if (i >= p.KX) {
|
||||
return;
|
||||
|
||||
@@ -134,7 +134,7 @@ void main() {
|
||||
|
||||
// Q8_0 K only needs (qd, _); the asymmetric Q4_*/Q5_* family also stores
|
||||
// the row-sum scaled by qd, used in k_dot_correction.
|
||||
if (FaTypeK == FA_TYPE_Q8_0) {
|
||||
if (FaTypeK == GGML_TYPE_Q8_0) {
|
||||
if (buf_iqs == 0) {
|
||||
Qf[buf_ib].ds = FLOAT_TYPEV2(qd, 0.0f);
|
||||
}
|
||||
@@ -367,7 +367,7 @@ void main() {
|
||||
// Q4_*/Q5_* take the block-8 fast path when one step covers a full
|
||||
// block; Q8_0 always goes through the per-int get_k_qs* helpers
|
||||
// (its qs is byte-packed, not nibble-packed).
|
||||
const bool block8_fast = (d_per_step == 8) && (FaTypeK != FA_TYPE_Q8_0);
|
||||
const bool block8_fast = (d_per_step == 8) && (FaTypeK != GGML_TYPE_Q8_0);
|
||||
|
||||
if (SHMEM_STAGING != 0) {
|
||||
const uint k_block_idx = (d_tid * (HSK_per_thread / 4) + d_block) / 8;
|
||||
@@ -375,7 +375,7 @@ void main() {
|
||||
k_dm = ACC_TYPEV2(kblocksh[buf_ib].dm);
|
||||
|
||||
if (block8_fast) {
|
||||
const bool has_qh = (FaTypeK == FA_TYPE_Q5_0) || (FaTypeK == FA_TYPE_Q5_1);
|
||||
const bool has_qh = (FaTypeK == GGML_TYPE_Q5_0) || (FaTypeK == GGML_TYPE_Q5_1);
|
||||
[[unroll]] for (uint32_t d = 0; d < 4; d++) {
|
||||
uint vui = kblocksh[buf_ib].qs[d];
|
||||
k_quants[d ] = int32_t( vui & 0x0F0F0F0F);
|
||||
|
||||
@@ -105,8 +105,8 @@ layout (binding = 6) readonly buffer MO {uint32_t data_mask_opt[];};
|
||||
#define BLOCK_SIZE_V fa_block_elems(FaTypeV)
|
||||
// F16 reads f16 elements directly from the binding; everything else routes
|
||||
// through dequantize4 / the MMQ helpers to unpack from the packed block layout.
|
||||
#define USE_DECODE_K (FaTypeK != FA_TYPE_F16)
|
||||
#define USE_DECODE_V (FaTypeV != FA_TYPE_F16)
|
||||
#define USE_DECODE_K (FaTypeK != GGML_TYPE_F16)
|
||||
#define USE_DECODE_V (FaTypeV != GGML_TYPE_F16)
|
||||
|
||||
#define CEIL_DIV(a, b) (((a) + (b) - 1) / (b))
|
||||
|
||||
|
||||
@@ -40,26 +40,28 @@ layout(buffer_reference, std430, buffer_reference_align = 1) buffer decodeBufFA_
|
||||
#if !defined(BFLOAT16)
|
||||
float16_t faDecodeK(const decodeBufFA_K bl_in, const uint blockCoords[2], const uint coordInBlock[2]) {
|
||||
switch (FaTypeK) {
|
||||
case FA_TYPE_F32: return dequantFuncF32 (decodeBufF32 (bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q4_0: return dequantFuncQ4_0(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q4_1: return dequantFuncQ4_1(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q5_0: return dequantFuncQ5_0(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q5_1: return dequantFuncQ5_1(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q8_0: return dequantFuncQ8_0(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_IQ4_NL: return dequantFuncIQ4_NL(decodeBufIQ4_NL(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_F32: return dequantFuncF32 (decodeBufF32 (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q4_0: return dequantFuncQ4_0(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q4_1: return dequantFuncQ4_1(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q5_0: return dequantFuncQ5_0(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q5_1: return dequantFuncQ5_1(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q8_0: return dequantFuncQ8_0(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_IQ4_NL: return dequantFuncIQ4_NL(decodeBufIQ4_NL(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q1_0: return dequantFuncQ1_0(decodeBufQ1_0(bl_in), blockCoords, coordInBlock);
|
||||
default: return float16_t(0);
|
||||
}
|
||||
}
|
||||
|
||||
float16_t faDecodeV(const decodeBufFA_V bl_in, const uint blockCoords[2], const uint coordInBlock[2]) {
|
||||
switch (FaTypeV) {
|
||||
case FA_TYPE_F32: return dequantFuncF32 (decodeBufF32 (bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q4_0: return dequantFuncQ4_0(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q4_1: return dequantFuncQ4_1(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q5_0: return dequantFuncQ5_0(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q5_1: return dequantFuncQ5_1(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q8_0: return dequantFuncQ8_0(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_IQ4_NL: return dequantFuncIQ4_NL(decodeBufIQ4_NL(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_F32: return dequantFuncF32 (decodeBufF32 (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q4_0: return dequantFuncQ4_0(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q4_1: return dequantFuncQ4_1(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q5_0: return dequantFuncQ5_0(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q5_1: return dequantFuncQ5_1(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q8_0: return dequantFuncQ8_0(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_IQ4_NL: return dequantFuncIQ4_NL(decodeBufIQ4_NL(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q1_0: return dequantFuncQ1_0(decodeBufQ1_0(bl_in), blockCoords, coordInBlock);
|
||||
default: return float16_t(0);
|
||||
}
|
||||
}
|
||||
@@ -67,26 +69,26 @@ float16_t faDecodeV(const decodeBufFA_V bl_in, const uint blockCoords[2], const
|
||||
// V=4 vector decode for K/V; dispatches to per-format _v decoders.
|
||||
f16vec4 faDecodeKVector(const decodeBufFA_K bl_in, const uint blockCoords[2], const uint coordInBlock[2]) {
|
||||
switch (FaTypeK) {
|
||||
case FA_TYPE_F32: return f16vec4(decodeBufF32(bl_in).block);
|
||||
case FA_TYPE_Q4_0: return dequantFuncQ4_0_v(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q4_1: return dequantFuncQ4_1_v(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q5_0: return dequantFuncQ5_0_v(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q5_1: return dequantFuncQ5_1_v(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q8_0: return dequantFuncQ8_0_v(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_IQ4_NL: return dequantFuncIQ4_NL_v(decodeBufIQ4_NL(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_F32: return f16vec4(decodeBufF32(bl_in).block);
|
||||
case GGML_TYPE_Q4_0: return dequantFuncQ4_0_v(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q4_1: return dequantFuncQ4_1_v(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q5_0: return dequantFuncQ5_0_v(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q5_1: return dequantFuncQ5_1_v(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q8_0: return dequantFuncQ8_0_v(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_IQ4_NL: return dequantFuncIQ4_NL_v(decodeBufIQ4_NL(bl_in), blockCoords, coordInBlock);
|
||||
default: return f16vec4(0);
|
||||
}
|
||||
}
|
||||
|
||||
f16vec4 faDecodeVVector(const decodeBufFA_V bl_in, const uint blockCoords[2], const uint coordInBlock[2]) {
|
||||
switch (FaTypeV) {
|
||||
case FA_TYPE_F32: return f16vec4(decodeBufF32(bl_in).block);
|
||||
case FA_TYPE_Q4_0: return dequantFuncQ4_0_v(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q4_1: return dequantFuncQ4_1_v(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q5_0: return dequantFuncQ5_0_v(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q5_1: return dequantFuncQ5_1_v(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_Q8_0: return dequantFuncQ8_0_v(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
|
||||
case FA_TYPE_IQ4_NL: return dequantFuncIQ4_NL_v(decodeBufIQ4_NL(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_F32: return f16vec4(decodeBufF32(bl_in).block);
|
||||
case GGML_TYPE_Q4_0: return dequantFuncQ4_0_v(decodeBufQ4_0(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q4_1: return dequantFuncQ4_1_v(decodeBufQ4_1(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q5_0: return dequantFuncQ5_0_v(decodeBufQ5_0(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q5_1: return dequantFuncQ5_1_v(decodeBufQ5_1(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q8_0: return dequantFuncQ8_0_v(decodeBufQ8_0(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_IQ4_NL: return dequantFuncIQ4_NL_v(decodeBufIQ4_NL(bl_in), blockCoords, coordInBlock);
|
||||
default: return f16vec4(0);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -121,25 +121,25 @@ layout (binding = 1) readonly buffer K_PACKED_Q5_1_P32 { block_q5_1_packed32 dat
|
||||
FLOAT_TYPEV4 dequantize4(uint ib, uint iqs, uint a_offset, uint binding_idx) {
|
||||
if (binding_idx == BINDING_IDX_K) {
|
||||
switch (FaTypeK) {
|
||||
case FA_TYPE_F32: FA_DEQUANT4_F32 (k_packed_f32)
|
||||
case FA_TYPE_Q4_0: FA_DEQUANT4_Q4_0(k_packed_q4_0)
|
||||
case FA_TYPE_Q4_1: FA_DEQUANT4_Q4_1(k_packed_q4_1)
|
||||
case FA_TYPE_Q5_0: FA_DEQUANT4_Q5_0(k_packed_q5_0)
|
||||
case FA_TYPE_Q5_1: FA_DEQUANT4_Q5_1(k_packed_q5_1)
|
||||
case FA_TYPE_Q8_0: FA_DEQUANT4_Q8_0(k_packed_q8_0)
|
||||
case FA_TYPE_IQ4_NL: FA_DEQUANT4_IQ4_NL(k_packed_iq4_nl)
|
||||
case FA_TYPE_BF16: FA_DEQUANT4_BF16(k_packed_bf16)
|
||||
case GGML_TYPE_F32: FA_DEQUANT4_F32 (k_packed_f32)
|
||||
case GGML_TYPE_Q4_0: FA_DEQUANT4_Q4_0(k_packed_q4_0)
|
||||
case GGML_TYPE_Q4_1: FA_DEQUANT4_Q4_1(k_packed_q4_1)
|
||||
case GGML_TYPE_Q5_0: FA_DEQUANT4_Q5_0(k_packed_q5_0)
|
||||
case GGML_TYPE_Q5_1: FA_DEQUANT4_Q5_1(k_packed_q5_1)
|
||||
case GGML_TYPE_Q8_0: FA_DEQUANT4_Q8_0(k_packed_q8_0)
|
||||
case GGML_TYPE_IQ4_NL: FA_DEQUANT4_IQ4_NL(k_packed_iq4_nl)
|
||||
case GGML_TYPE_BF16: FA_DEQUANT4_BF16(k_packed_bf16)
|
||||
}
|
||||
} else {
|
||||
switch (FaTypeV) {
|
||||
case FA_TYPE_F32: FA_DEQUANT4_F32 (v_packed_f32)
|
||||
case FA_TYPE_Q4_0: FA_DEQUANT4_Q4_0(v_packed_q4_0)
|
||||
case FA_TYPE_Q4_1: FA_DEQUANT4_Q4_1(v_packed_q4_1)
|
||||
case FA_TYPE_Q5_0: FA_DEQUANT4_Q5_0(v_packed_q5_0)
|
||||
case FA_TYPE_Q5_1: FA_DEQUANT4_Q5_1(v_packed_q5_1)
|
||||
case FA_TYPE_Q8_0: FA_DEQUANT4_Q8_0(v_packed_q8_0)
|
||||
case FA_TYPE_IQ4_NL: FA_DEQUANT4_IQ4_NL(v_packed_iq4_nl)
|
||||
case FA_TYPE_BF16: FA_DEQUANT4_BF16(v_packed_bf16)
|
||||
case GGML_TYPE_F32: FA_DEQUANT4_F32 (v_packed_f32)
|
||||
case GGML_TYPE_Q4_0: FA_DEQUANT4_Q4_0(v_packed_q4_0)
|
||||
case GGML_TYPE_Q4_1: FA_DEQUANT4_Q4_1(v_packed_q4_1)
|
||||
case GGML_TYPE_Q5_0: FA_DEQUANT4_Q5_0(v_packed_q5_0)
|
||||
case GGML_TYPE_Q5_1: FA_DEQUANT4_Q5_1(v_packed_q5_1)
|
||||
case GGML_TYPE_Q8_0: FA_DEQUANT4_Q8_0(v_packed_q8_0)
|
||||
case GGML_TYPE_IQ4_NL: FA_DEQUANT4_IQ4_NL(v_packed_iq4_nl)
|
||||
case GGML_TYPE_BF16: FA_DEQUANT4_BF16(v_packed_bf16)
|
||||
}
|
||||
}
|
||||
return FLOAT_TYPEV4(0);
|
||||
|
||||
@@ -4,20 +4,20 @@
|
||||
|
||||
int32_t get_k_qs(uint ib, uint iqs, uint a_offset) {
|
||||
switch (FaTypeK) {
|
||||
case FA_TYPE_Q4_0: {
|
||||
case GGML_TYPE_Q4_0: {
|
||||
uint vui = pack32(u16vec2(k_packed_q4_0.data[a_offset + ib].qs[(iqs & 0xF) / 2 + 0],
|
||||
k_packed_q4_0.data[a_offset + ib].qs[(iqs & 0xF) / 2 + 1]));
|
||||
uint shift = (iqs & 0x10) >> 2;
|
||||
vui >>= shift;
|
||||
return int32_t(vui & 0x0F0F0F0F);
|
||||
}
|
||||
case FA_TYPE_Q4_1: { // uses packed32 alias
|
||||
case GGML_TYPE_Q4_1: { // uses packed32 alias
|
||||
uint vui = k_packed_q4_1_p32.data[a_offset + ib].qs[(iqs & 0xF) / 4];
|
||||
uint shift = (iqs & 0x10) >> 2;
|
||||
vui >>= shift;
|
||||
return int32_t(vui & 0x0F0F0F0F);
|
||||
}
|
||||
case FA_TYPE_Q5_0: {
|
||||
case GGML_TYPE_Q5_0: {
|
||||
uint vui = pack32(u16vec2(k_packed_q5_0.data[a_offset + ib].qs[(iqs & 0xF) / 2 + 0],
|
||||
k_packed_q5_0.data[a_offset + ib].qs[(iqs & 0xF) / 2 + 1]));
|
||||
uint qh = pack32(u16vec2(k_packed_q5_0.data[a_offset + ib].qh[0],
|
||||
@@ -27,7 +27,7 @@ int32_t get_k_qs(uint ib, uint iqs, uint a_offset) {
|
||||
uint qh_bits = (qh >> iqs) & 0xF;
|
||||
return int32_t(vui & 0x0F0F0F0F) | int32_t((qh_bits * 0x02040810u) & 0x10101010u);
|
||||
}
|
||||
case FA_TYPE_Q5_1: { // qs via packed32, qh via packed16
|
||||
case GGML_TYPE_Q5_1: { // qs via packed32, qh via packed16
|
||||
uint vui = k_packed_q5_1_p32.data[a_offset + ib].qs[(iqs & 0xF) / 4];
|
||||
uint qh = k_packed_q5_1.data[a_offset + ib].qh;
|
||||
uint shift = (iqs & 0x10) >> 2;
|
||||
@@ -35,7 +35,7 @@ int32_t get_k_qs(uint ib, uint iqs, uint a_offset) {
|
||||
uint qh_bits = (qh >> iqs) & 0xF;
|
||||
return int32_t(vui & 0x0F0F0F0F) | int32_t((qh_bits * 0x02040810u) & 0x10101010u);
|
||||
}
|
||||
case FA_TYPE_Q8_0: {
|
||||
case GGML_TYPE_Q8_0: {
|
||||
return pack32(i16vec2(k_packed_q8_0.data[a_offset + ib].qs[iqs / 2],
|
||||
k_packed_q8_0.data[a_offset + ib].qs[iqs / 2 + 1]));
|
||||
}
|
||||
@@ -47,11 +47,11 @@ int32_t get_k_qs(uint ib, uint iqs, uint a_offset) {
|
||||
// return (d, 0) so call sites always see the same shape.
|
||||
FLOAT_TYPEV2 get_k_scale(uint ib, uint a_offset) {
|
||||
switch (FaTypeK) {
|
||||
case FA_TYPE_Q4_0: return FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q4_0.data[a_offset + ib].d), 0.0);
|
||||
case FA_TYPE_Q4_1: return FLOAT_TYPEV2(k_packed_q4_1_p32.data[a_offset + ib].dm);
|
||||
case FA_TYPE_Q5_0: return FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q5_0.data[a_offset + ib].d), 0.0);
|
||||
case FA_TYPE_Q5_1: return FLOAT_TYPEV2(k_packed_q5_1_p32.data[a_offset + ib].dm);
|
||||
case FA_TYPE_Q8_0: return FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q8_0.data[a_offset + ib].d), 0.0);
|
||||
case GGML_TYPE_Q4_0: return FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q4_0.data[a_offset + ib].d), 0.0);
|
||||
case GGML_TYPE_Q4_1: return FLOAT_TYPEV2(k_packed_q4_1_p32.data[a_offset + ib].dm);
|
||||
case GGML_TYPE_Q5_0: return FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q5_0.data[a_offset + ib].d), 0.0);
|
||||
case GGML_TYPE_Q5_1: return FLOAT_TYPEV2(k_packed_q5_1_p32.data[a_offset + ib].dm);
|
||||
case GGML_TYPE_Q8_0: return FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q8_0.data[a_offset + ib].d), 0.0);
|
||||
default: return FLOAT_TYPEV2(0);
|
||||
}
|
||||
}
|
||||
@@ -61,16 +61,16 @@ void k_block_to_shmem(const uint buf_ib, const uint global_ib, const uint iqs, c
|
||||
// explicit casts. The bit pattern is what we care about here -- the actual
|
||||
// signed/unsigned interpretation happens downstream in the dot product.
|
||||
switch (FaTypeK) {
|
||||
case FA_TYPE_Q4_0: {
|
||||
case GGML_TYPE_Q4_0: {
|
||||
kblocksh[buf_ib].qs[iqs] = int32_t(pack32(u16vec2(k_packed_q4_0.data[a_offset + global_ib].qs[iqs * 2],
|
||||
k_packed_q4_0.data[a_offset + global_ib].qs[iqs * 2 + 1])));
|
||||
break;
|
||||
}
|
||||
case FA_TYPE_Q4_1: {
|
||||
case GGML_TYPE_Q4_1: {
|
||||
kblocksh[buf_ib].qs[iqs] = int32_t(k_packed_q4_1_p32.data[a_offset + global_ib].qs[iqs]);
|
||||
break;
|
||||
}
|
||||
case FA_TYPE_Q5_0: {
|
||||
case GGML_TYPE_Q5_0: {
|
||||
kblocksh[buf_ib].qs[iqs] = int32_t(pack32(u16vec2(k_packed_q5_0.data[a_offset + global_ib].qs[iqs * 2],
|
||||
k_packed_q5_0.data[a_offset + global_ib].qs[iqs * 2 + 1])));
|
||||
if (iqs == 0) {
|
||||
@@ -79,14 +79,14 @@ void k_block_to_shmem(const uint buf_ib, const uint global_ib, const uint iqs, c
|
||||
}
|
||||
break;
|
||||
}
|
||||
case FA_TYPE_Q5_1: {
|
||||
case GGML_TYPE_Q5_1: {
|
||||
kblocksh[buf_ib].qs[iqs] = int32_t(k_packed_q5_1_p32.data[a_offset + global_ib].qs[iqs]);
|
||||
if (iqs == 0) {
|
||||
kblocksh[buf_ib].qh = k_packed_q5_1.data[a_offset + global_ib].qh;
|
||||
}
|
||||
break;
|
||||
}
|
||||
case FA_TYPE_Q8_0: {
|
||||
case GGML_TYPE_Q8_0: {
|
||||
kblocksh[buf_ib].qs[iqs] = pack32(i16vec2(k_packed_q8_0.data[a_offset + global_ib].qs[iqs * 2],
|
||||
k_packed_q8_0.data[a_offset + global_ib].qs[iqs * 2 + 1]));
|
||||
break;
|
||||
@@ -96,11 +96,11 @@ void k_block_to_shmem(const uint buf_ib, const uint global_ib, const uint iqs, c
|
||||
if (iqs == 0) {
|
||||
// Q4_0/Q5_0/Q8_0 store dm.x = d; Q4_1/Q5_1 store dm = (d, m) pair.
|
||||
switch (FaTypeK) {
|
||||
case FA_TYPE_Q4_0: kblocksh[buf_ib].dm = FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q4_0.data[a_offset + global_ib].d), 0.0); break;
|
||||
case FA_TYPE_Q4_1: kblocksh[buf_ib].dm = FLOAT_TYPEV2(k_packed_q4_1_p32.data[a_offset + global_ib].dm); break;
|
||||
case FA_TYPE_Q5_0: kblocksh[buf_ib].dm = FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q5_0.data[a_offset + global_ib].d), 0.0); break;
|
||||
case FA_TYPE_Q5_1: kblocksh[buf_ib].dm = FLOAT_TYPEV2(k_packed_q5_1_p32.data[a_offset + global_ib].dm); break;
|
||||
case FA_TYPE_Q8_0: kblocksh[buf_ib].dm = FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q8_0.data[a_offset + global_ib].d), 0.0); break;
|
||||
case GGML_TYPE_Q4_0: kblocksh[buf_ib].dm = FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q4_0.data[a_offset + global_ib].d), 0.0); break;
|
||||
case GGML_TYPE_Q4_1: kblocksh[buf_ib].dm = FLOAT_TYPEV2(k_packed_q4_1_p32.data[a_offset + global_ib].dm); break;
|
||||
case GGML_TYPE_Q5_0: kblocksh[buf_ib].dm = FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q5_0.data[a_offset + global_ib].d), 0.0); break;
|
||||
case GGML_TYPE_Q5_1: kblocksh[buf_ib].dm = FLOAT_TYPEV2(k_packed_q5_1_p32.data[a_offset + global_ib].dm); break;
|
||||
case GGML_TYPE_Q8_0: kblocksh[buf_ib].dm = FLOAT_TYPEV2(FLOAT_TYPE(k_packed_q8_0.data[a_offset + global_ib].d), 0.0); break;
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -121,31 +121,31 @@ struct fa_k_qs_block8 {
|
||||
fa_k_qs_block8 get_k_qs_block8(uint ib, uint a_offset) {
|
||||
fa_k_qs_block8 r;
|
||||
uint qh = 0;
|
||||
if (FaTypeK == FA_TYPE_Q5_0) {
|
||||
if (FaTypeK == GGML_TYPE_Q5_0) {
|
||||
qh = pack32(u16vec2(k_packed_q5_0.data[a_offset + ib].qh[0],
|
||||
k_packed_q5_0.data[a_offset + ib].qh[1]));
|
||||
} else if (FaTypeK == FA_TYPE_Q5_1) {
|
||||
} else if (FaTypeK == GGML_TYPE_Q5_1) {
|
||||
qh = k_packed_q5_1.data[a_offset + ib].qh;
|
||||
}
|
||||
const bool has_qh = (FaTypeK == FA_TYPE_Q5_0) || (FaTypeK == FA_TYPE_Q5_1);
|
||||
const bool has_qh = (FaTypeK == GGML_TYPE_Q5_0) || (FaTypeK == GGML_TYPE_Q5_1);
|
||||
[[unroll]] for (uint32_t d = 0; d < 4; d++) {
|
||||
uint vui = 0;
|
||||
switch (FaTypeK) {
|
||||
case FA_TYPE_Q4_0: { // packed16
|
||||
case GGML_TYPE_Q4_0: { // packed16
|
||||
vui = pack32(u16vec2(k_packed_q4_0.data[a_offset + ib].qs[d * 2 + 0],
|
||||
k_packed_q4_0.data[a_offset + ib].qs[d * 2 + 1]));
|
||||
break;
|
||||
}
|
||||
case FA_TYPE_Q4_1: { // packed32 alias
|
||||
case GGML_TYPE_Q4_1: { // packed32 alias
|
||||
vui = k_packed_q4_1_p32.data[a_offset + ib].qs[d];
|
||||
break;
|
||||
}
|
||||
case FA_TYPE_Q5_0: { // packed16
|
||||
case GGML_TYPE_Q5_0: { // packed16
|
||||
vui = pack32(u16vec2(k_packed_q5_0.data[a_offset + ib].qs[d * 2 + 0],
|
||||
k_packed_q5_0.data[a_offset + ib].qs[d * 2 + 1]));
|
||||
break;
|
||||
}
|
||||
case FA_TYPE_Q5_1: { // packed32 alias
|
||||
case GGML_TYPE_Q5_1: { // packed32 alias
|
||||
vui = k_packed_q5_1_p32.data[a_offset + ib].qs[d];
|
||||
break;
|
||||
}
|
||||
@@ -164,21 +164,21 @@ fa_k_qs_block8 get_k_qs_block8(uint ib, uint a_offset) {
|
||||
|
||||
int32_t get_k_qs_shmem(const uint buf_ib, const uint pos) {
|
||||
switch (FaTypeK) {
|
||||
case FA_TYPE_Q4_0:
|
||||
case FA_TYPE_Q4_1: {
|
||||
case GGML_TYPE_Q4_0:
|
||||
case GGML_TYPE_Q4_1: {
|
||||
uint sub = pos % 4;
|
||||
uint shift = ((pos % 8) >= 4) ? 4u : 0u;
|
||||
return int32_t((uint(kblocksh[buf_ib].qs[sub]) >> shift) & 0x0F0F0F0Fu);
|
||||
}
|
||||
case FA_TYPE_Q5_0:
|
||||
case FA_TYPE_Q5_1: {
|
||||
case GGML_TYPE_Q5_0:
|
||||
case GGML_TYPE_Q5_1: {
|
||||
uint sub = pos % 4;
|
||||
uint shift = ((pos % 8) >= 4) ? 4u : 0u;
|
||||
int32_t result = int32_t((uint(kblocksh[buf_ib].qs[sub]) >> shift) & 0x0F0F0F0Fu);
|
||||
uint qh_bits = (kblocksh[buf_ib].qh >> (pos * 4u)) & 0xFu;
|
||||
return result | int32_t((qh_bits * 0x02040810u) & 0x10101010u);
|
||||
}
|
||||
case FA_TYPE_Q8_0: {
|
||||
case GGML_TYPE_Q8_0: {
|
||||
return kblocksh[buf_ib].qs[pos];
|
||||
}
|
||||
default: return 0;
|
||||
@@ -187,10 +187,10 @@ int32_t get_k_qs_shmem(const uint buf_ib, const uint pos) {
|
||||
|
||||
ACC_TYPE k_dot_correction(const uint qib, const ACC_TYPEV2 k_dm) {
|
||||
switch (FaTypeK) {
|
||||
case FA_TYPE_Q4_0: return -ACC_TYPE(8.0) * ACC_TYPE(Qf[qib].ds.y) * k_dm.x;
|
||||
case FA_TYPE_Q5_0: return -ACC_TYPE(16.0) * ACC_TYPE(Qf[qib].ds.y) * k_dm.x;
|
||||
case FA_TYPE_Q4_1:
|
||||
case FA_TYPE_Q5_1: return ACC_TYPE(Qf[qib].ds.y) * k_dm.y;
|
||||
case GGML_TYPE_Q4_0: return -ACC_TYPE(8.0) * ACC_TYPE(Qf[qib].ds.y) * k_dm.x;
|
||||
case GGML_TYPE_Q5_0: return -ACC_TYPE(16.0) * ACC_TYPE(Qf[qib].ds.y) * k_dm.x;
|
||||
case GGML_TYPE_Q4_1:
|
||||
case GGML_TYPE_Q5_1: return ACC_TYPE(Qf[qib].ds.y) * k_dm.y;
|
||||
default: return ACC_TYPE(0.0);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -0,0 +1,34 @@
|
||||
#if !defined(GGML_TYPE_IDS_COMP)
|
||||
#define GGML_TYPE_IDS_COMP
|
||||
|
||||
// ggml_type enum values — must match ggml.h
|
||||
#define GGML_TYPE_F32 0u
|
||||
#define GGML_TYPE_F16 1u
|
||||
#define GGML_TYPE_Q4_0 2u
|
||||
#define GGML_TYPE_Q4_1 3u
|
||||
#define GGML_TYPE_Q5_0 6u
|
||||
#define GGML_TYPE_Q5_1 7u
|
||||
#define GGML_TYPE_Q8_0 8u
|
||||
#define GGML_TYPE_Q2_K 10u
|
||||
#define GGML_TYPE_Q3_K 11u
|
||||
#define GGML_TYPE_Q4_K 12u
|
||||
#define GGML_TYPE_Q5_K 13u
|
||||
#define GGML_TYPE_Q6_K 14u
|
||||
#define GGML_TYPE_IQ2_XXS 16u
|
||||
#define GGML_TYPE_IQ2_XS 17u
|
||||
#define GGML_TYPE_IQ3_XXS 18u
|
||||
#define GGML_TYPE_IQ1_S 19u
|
||||
#define GGML_TYPE_IQ4_NL 20u
|
||||
#define GGML_TYPE_IQ3_S 21u
|
||||
#define GGML_TYPE_IQ2_S 22u
|
||||
#define GGML_TYPE_IQ4_XS 23u
|
||||
#define GGML_TYPE_IQ1_M 29u
|
||||
#define GGML_TYPE_BF16 30u
|
||||
#define GGML_TYPE_TQ1_0 34u
|
||||
#define GGML_TYPE_TQ2_0 35u
|
||||
#define GGML_TYPE_MXFP4 39u
|
||||
#define GGML_TYPE_NVFP4 40u
|
||||
#define GGML_TYPE_Q1_0 41u
|
||||
#define GGML_TYPE_Q2_0 42u
|
||||
|
||||
#endif // !defined(GGML_TYPE_IDS_COMP)
|
||||
@@ -0,0 +1,2 @@
|
||||
void init_iq_shmem(uvec3 wgsize) {
|
||||
}
|
||||
@@ -12,9 +12,9 @@
|
||||
|
||||
#include "types.glsl"
|
||||
#include "fa_types.glsl"
|
||||
#define FaTypeV FA_TYPE_F32
|
||||
#define FaTypeV GGML_TYPE_F32
|
||||
|
||||
layout(constant_id = 0) const uint FaTypeK = FA_TYPE_F32;
|
||||
layout(constant_id = 0) const uint FaTypeK = GGML_TYPE_F32;
|
||||
layout(constant_id = 1) const uint FaBlockBytesK = 4;
|
||||
layout(constant_id = 2) const uint SUBGROUP_SIZE = 32;
|
||||
|
||||
@@ -84,11 +84,11 @@ void main() {
|
||||
const uint k_block_elems = fa_block_elems(FaTypeK);
|
||||
const uint k_elem_bytes = FaBlockBytesK / k_block_elems;
|
||||
|
||||
if (FaTypeK == FA_TYPE_F16) {
|
||||
if (FaTypeK == GGML_TYPE_F16) {
|
||||
k_row[tid] = float(k_f16[k_offset / k_elem_bytes + tid]);
|
||||
} else if (FaTypeK == FA_TYPE_F32) {
|
||||
} else if (FaTypeK == GGML_TYPE_F32) {
|
||||
k_row[tid] = k_f32[k_offset / k_elem_bytes + tid];
|
||||
} else if (FaTypeK == FA_TYPE_BF16) {
|
||||
} else if (FaTypeK == GGML_TYPE_BF16) {
|
||||
k_row[tid] = bf16_to_fp32(uint(k_bf16[k_offset / k_elem_bytes + tid]));
|
||||
} else if (4 * tid < HEAD_SIZE) {
|
||||
const uint coord = 4 * tid;
|
||||
|
||||
@@ -0,0 +1,97 @@
|
||||
#version 450
|
||||
|
||||
#extension GL_EXT_shader_explicit_arithmetic_types_int32 : require
|
||||
|
||||
#include "mul_mat_vec_base.glsl"
|
||||
|
||||
layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in;
|
||||
|
||||
FLOAT_TYPE temp[NUM_COLS][NUM_ROWS];
|
||||
|
||||
// dedicated iq4_xs mat-vec, mirrors mul_mat_vec_iq3_s.comp
|
||||
// one packed32 word per l, so the 6-bit subblock scale is hoisted to a single fma after register accumulation
|
||||
|
||||
void calc_superblock(const uint a_offset, const uint b_offset, const uint ib32, const uint i, const uint num_blocks_per_row, const uint first_row, const uint num_rows) {
|
||||
const uint y_idx = i * QUANT_K + 32 * ib32;
|
||||
|
||||
uint ibi = a_offset + first_row * num_blocks_per_row + i;
|
||||
[[unroll]] for (uint n = 0; n < num_rows; ++n) {
|
||||
const float d = float(data_a[ibi].d);
|
||||
const uint sl = (data_a[ibi].scales_l[ib32/2] >> (4 * (ib32 & 1))) & 0xF;
|
||||
const uint sh = (data_a[ibi].scales_h >> (2 * ib32)) & 3;
|
||||
const float dscale = d * float(int(sl | (sh << 4)) - 32);
|
||||
|
||||
FLOAT_TYPE sum[NUM_COLS];
|
||||
[[unroll]] for (uint j = 0; j < NUM_COLS; ++j) {
|
||||
sum[j] = FLOAT_TYPE(0);
|
||||
}
|
||||
|
||||
[[unroll]] for (uint l = 0; l < 4; ++l) {
|
||||
const uint w = data_a_packed32[ibi].qs[4 * ib32 + l];
|
||||
const u8vec4 q0 = unpack8(w & 0x0F0F0F0F);
|
||||
const u8vec4 q1 = unpack8((w >> 4) & 0x0F0F0F0F);
|
||||
|
||||
[[unroll]] for (uint j = 0; j < NUM_COLS; ++j) {
|
||||
const vec4 b0 = vec4(data_b_v4[(j*p.batch_stride_b + b_offset + y_idx) / 4 + l]);
|
||||
const vec4 b1 = vec4(data_b_v4[(j*p.batch_stride_b + b_offset + y_idx) / 4 + 4 + l]);
|
||||
|
||||
sum[j] = fma(FLOAT_TYPE(b0.x), FLOAT_TYPE(kvalues_iq4nl[q0.x]),
|
||||
fma(FLOAT_TYPE(b0.y), FLOAT_TYPE(kvalues_iq4nl[q0.y]),
|
||||
fma(FLOAT_TYPE(b0.z), FLOAT_TYPE(kvalues_iq4nl[q0.z]),
|
||||
fma(FLOAT_TYPE(b0.w), FLOAT_TYPE(kvalues_iq4nl[q0.w]),
|
||||
fma(FLOAT_TYPE(b1.x), FLOAT_TYPE(kvalues_iq4nl[q1.x]),
|
||||
fma(FLOAT_TYPE(b1.y), FLOAT_TYPE(kvalues_iq4nl[q1.y]),
|
||||
fma(FLOAT_TYPE(b1.z), FLOAT_TYPE(kvalues_iq4nl[q1.z]),
|
||||
fma(FLOAT_TYPE(b1.w), FLOAT_TYPE(kvalues_iq4nl[q1.w]),
|
||||
sum[j]))))))));
|
||||
}
|
||||
}
|
||||
|
||||
[[unroll]] for (uint j = 0; j < NUM_COLS; ++j) {
|
||||
temp[j][n] = fma(dscale, sum[j], temp[j][n]);
|
||||
}
|
||||
|
||||
ibi += num_blocks_per_row;
|
||||
}
|
||||
}
|
||||
|
||||
void compute_outputs(const uint32_t first_row, const uint32_t num_rows) {
|
||||
uint a_offset, b_offset, d_offset;
|
||||
|
||||
get_offsets(a_offset, b_offset, d_offset);
|
||||
|
||||
const uint num_blocks_per_row = p.ncols / QUANT_K;
|
||||
|
||||
// 8 threads are used to process each block
|
||||
const uint blocks_per_wg = gl_WorkGroupSize.x/8;
|
||||
const uint tid = gl_LocalInvocationID.x;
|
||||
const uint itid = tid % 8; // 0...7
|
||||
const uint ix = tid / 8;
|
||||
|
||||
[[unroll]] for (uint j = 0; j < NUM_COLS; ++j) {
|
||||
[[unroll]] for (uint i = 0; i < NUM_ROWS; ++i) {
|
||||
temp[j][i] = FLOAT_TYPE(0);
|
||||
}
|
||||
}
|
||||
|
||||
[[unroll]] for (uint i = ix; i < num_blocks_per_row; i += blocks_per_wg)
|
||||
calc_superblock(a_offset, b_offset, itid, i, num_blocks_per_row, first_row, num_rows);
|
||||
|
||||
reduce_result(temp, d_offset, first_row, num_rows, tid);
|
||||
}
|
||||
|
||||
void main() {
|
||||
const uint first_row = NUM_ROWS * (gl_WorkGroupID.x + gl_NumWorkGroups.x * gl_WorkGroupID.z);
|
||||
|
||||
init_iq_shmem(gl_WorkGroupSize);
|
||||
|
||||
// do NUM_ROWS at a time, unless there aren't enough remaining rows
|
||||
if (first_row + NUM_ROWS <= p.stride_d) {
|
||||
compute_outputs(first_row, NUM_ROWS);
|
||||
} else {
|
||||
if (first_row >= p.stride_d) {
|
||||
return;
|
||||
}
|
||||
compute_outputs(first_row, p.stride_d - first_row);
|
||||
}
|
||||
}
|
||||
@@ -9,6 +9,9 @@
|
||||
#if defined(DATA_A_IQ1_M)
|
||||
#extension GL_EXT_shader_explicit_arithmetic_types_int16 : require
|
||||
#endif
|
||||
#if !defined(DATA_A_F32) && !defined(DATA_A_F16) && !defined(DATA_A_BF16)
|
||||
#extension GL_EXT_shader_explicit_arithmetic_types_int16 : require
|
||||
#endif
|
||||
|
||||
#if defined(DATA_A_BF16) && defined(COOPMAT)
|
||||
#extension GL_EXT_bfloat16 : enable
|
||||
@@ -28,24 +31,54 @@
|
||||
#extension GL_EXT_shader_explicit_arithmetic_types_int16 : require
|
||||
#endif
|
||||
|
||||
#ifdef MULMAT_QUANT
|
||||
#include "ggml_type_ids.glsl"
|
||||
layout (constant_id = 12) const uint MmTypeA = 0;
|
||||
#endif
|
||||
|
||||
#include "types.glsl"
|
||||
#include "dot_product_funcs.glsl"
|
||||
|
||||
#ifndef MULMAT_QUANT
|
||||
#ifndef LOAD_VEC_A
|
||||
#define LOAD_VEC_A 1
|
||||
#endif
|
||||
#endif
|
||||
#ifndef LOAD_VEC_B
|
||||
#define LOAD_VEC_B 1
|
||||
#endif
|
||||
|
||||
layout (constant_id = 11) const uint ALIGNED = 0;
|
||||
|
||||
#ifdef MULMAT_QUANT
|
||||
|
||||
uint mm_load_vec_a() {
|
||||
switch (MmTypeA) {
|
||||
case GGML_TYPE_Q1_0:
|
||||
case GGML_TYPE_Q4_0:
|
||||
case GGML_TYPE_Q4_1:
|
||||
case GGML_TYPE_Q5_1:
|
||||
return 8u;
|
||||
case GGML_TYPE_Q2_0:
|
||||
case GGML_TYPE_Q5_0:
|
||||
case GGML_TYPE_Q8_0:
|
||||
case GGML_TYPE_Q2_K:
|
||||
case GGML_TYPE_Q4_K:
|
||||
case GGML_TYPE_Q5_K:
|
||||
return 4u;
|
||||
default:
|
||||
return 2u;
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
#if !defined(TO_FLOAT_TYPE)
|
||||
#define TO_FLOAT_TYPE FLOAT_TYPE
|
||||
#endif
|
||||
|
||||
layout(local_size_x_id = 0, local_size_y = 1, local_size_z = 1) in;
|
||||
|
||||
#ifndef MULMAT_QUANT
|
||||
layout (binding = 0) readonly buffer A {A_TYPE data_a[];};
|
||||
#if defined(DATA_A_F32)
|
||||
layout (binding = 0) readonly buffer A_SCALAR {float data_a_scalar[];};
|
||||
@@ -60,6 +93,30 @@ layout (binding = 0) readonly buffer A_PACKED16 {A_TYPE_PACKED16 data_a_packed16
|
||||
#if defined(A_TYPE_PACKED32)
|
||||
layout (binding = 0) readonly buffer A_PACKED32 {A_TYPE_PACKED32 data_a_packed32[];};
|
||||
#endif
|
||||
#else
|
||||
// Unpacked struct aliases
|
||||
layout (binding = 0) readonly buffer BUF_Q1_0 { block_q1_0 data[]; } a_q1_0;
|
||||
layout (binding = 0) readonly buffer BUF_Q2_0 { block_q2_0 data[]; } a_q2_0;
|
||||
layout (binding = 0) readonly buffer BUF_Q2_K { block_q2_K data[]; } a_q2_k;
|
||||
layout (binding = 0) readonly buffer BUF_Q3_K { block_q3_K data[]; } a_q3_k;
|
||||
layout (binding = 0) readonly buffer BUF_Q4_K { block_q4_K data[]; } a_q4_k;
|
||||
layout (binding = 0) readonly buffer BUF_Q5_K { block_q5_K data[]; } a_q5_k;
|
||||
layout (binding = 0) readonly buffer BUF_Q6_K { block_q6_K data[]; } a_q6_k;
|
||||
layout (binding = 0) readonly buffer BUF_TQ1_0 { block_tq1_0 data[]; } a_tq1_0;
|
||||
layout (binding = 0) readonly buffer BUF_TQ2_0 { block_tq2_0 data[]; } a_tq2_0;
|
||||
// Packed16 aliases
|
||||
layout (binding = 0) readonly buffer BUF_Q4_0_P16 { block_q4_0_packed16 data[]; } a_q4_0_p16;
|
||||
layout (binding = 0) readonly buffer BUF_Q5_0_P16 { block_q5_0_packed16 data[]; } a_q5_0_p16;
|
||||
layout (binding = 0) readonly buffer BUF_Q8_0_P16 { block_q8_0_packed16 data[]; } a_q8_0_p16;
|
||||
layout (binding = 0) readonly buffer BUF_Q3_K_P16 { block_q3_K_packed16 data[]; } a_q3_k_p16;
|
||||
layout (binding = 0) readonly buffer BUF_Q6_K_P16 { block_q6_K_packed16 data[]; } a_q6_k_p16;
|
||||
// Packed32 aliases
|
||||
layout (binding = 0) readonly buffer BUF_Q4_1_P32 { block_q4_1_packed32 data[]; } a_q4_1_p32;
|
||||
layout (binding = 0) readonly buffer BUF_Q5_1_P32 { block_q5_1_packed32 data[]; } a_q5_1_p32;
|
||||
layout (binding = 0) readonly buffer BUF_Q2_K_P32 { block_q2_K_packed32 data[]; } a_q2_k_p32;
|
||||
layout (binding = 0) readonly buffer BUF_Q4_K_P32 { block_q4_K_packed32 data[]; } a_q4_k_p32;
|
||||
layout (binding = 0) readonly buffer BUF_Q5_K_P32 { block_q5_K_packed32 data[]; } a_q5_k_p32;
|
||||
#endif
|
||||
|
||||
layout (binding = 1) readonly buffer B {B_TYPE data_b[];};
|
||||
layout (binding = 1) readonly buffer B_SCALAR {B_TYPE_SCALAR data_b_scalar[];};
|
||||
@@ -121,8 +178,13 @@ layout (constant_id = 3) const uint BK = 16; // Assumed to be 32 if working wit
|
||||
#endif
|
||||
|
||||
#ifdef COOPMAT
|
||||
#ifdef MULMAT_QUANT
|
||||
layout(constant_id = 13) const uint SHMEM_STRIDE_PAD = 4;
|
||||
layout(constant_id = 14) const bool APPLY_SLM_A_RESHAPE = false;
|
||||
#else
|
||||
layout(constant_id = 12) const uint SHMEM_STRIDE_PAD = 4;
|
||||
layout(constant_id = 13) const bool APPLY_SLM_A_RESHAPE = false;
|
||||
#endif
|
||||
#else
|
||||
const uint SHMEM_STRIDE_PAD = 1;
|
||||
const bool APPLY_SLM_A_RESHAPE = false;
|
||||
@@ -141,6 +203,10 @@ shared ACC_TYPE coopmat_stage[TM * TN * NUM_WARPS];
|
||||
#include "mul_mm_id_funcs.glsl"
|
||||
#include "mul_mm_funcs.glsl"
|
||||
|
||||
#ifdef MULMAT_QUANT
|
||||
#include "iq_shmem_init.glsl"
|
||||
#endif
|
||||
|
||||
void main() {
|
||||
const uint ic = gl_WorkGroupID.y;
|
||||
|
||||
@@ -150,7 +216,7 @@ void main() {
|
||||
return;
|
||||
}
|
||||
#endif
|
||||
#ifdef NEEDS_INIT_IQ_SHMEM
|
||||
#if defined(NEEDS_INIT_IQ_SHMEM) || defined(MULMAT_QUANT)
|
||||
init_iq_shmem(gl_WorkGroupSize);
|
||||
#endif
|
||||
|
||||
@@ -200,9 +266,12 @@ void main() {
|
||||
#if defined(DATA_A_F32) || defined(DATA_A_F16) || defined(DATA_A_BF16)
|
||||
const uint LOAD_VEC_A_EFF = (ALIGNED != 0) ? LOAD_VEC_A : 1;
|
||||
const uint LOAD_VEC_BATCH_A = (ALIGNED != 0) ? 1 : 2;
|
||||
#else
|
||||
#elif !defined(MULMAT_QUANT)
|
||||
const uint LOAD_VEC_A_EFF = LOAD_VEC_A;
|
||||
const uint LOAD_VEC_BATCH_A = 1;
|
||||
#else
|
||||
const uint LOAD_VEC_A_EFF = mm_load_vec_a();
|
||||
const uint LOAD_VEC_BATCH_A = 1;
|
||||
#endif
|
||||
const uint LOAD_VEC_B_EFF = (ALIGNED != 0) ? LOAD_VEC_B : 1;
|
||||
const uint LOAD_VEC_BATCH_B = (ALIGNED != 0) ? 1 : 2;
|
||||
|
||||
@@ -21,6 +21,13 @@
|
||||
#extension GL_EXT_bfloat16 : enable
|
||||
#endif
|
||||
|
||||
#include "ggml_type_ids.glsl"
|
||||
|
||||
#ifdef MULMAT_QUANT
|
||||
layout (constant_id = 7) const uint MmTypeA = 0;
|
||||
layout (constant_id = 8) const uint MmABlockBytes = 2;
|
||||
#endif
|
||||
|
||||
#include "types.glsl"
|
||||
#include "utils.glsl"
|
||||
|
||||
@@ -37,6 +44,24 @@ layout (constant_id = 4) const bool enable_smaller_matrices = false;
|
||||
const uint BNover2 = enable_smaller_matrices ? (BN / 2) : BN;
|
||||
const uint BNover4 = enable_smaller_matrices ? (BN / 4) : BN;
|
||||
layout (constant_id = 5) const uint ALIGNED = 0;
|
||||
layout (constant_id = 6) const uint subgroup_size = 32;
|
||||
|
||||
#ifdef MULMAT_QUANT
|
||||
|
||||
uint mm_quant_k() {
|
||||
switch (MmTypeA) {
|
||||
case GGML_TYPE_Q4_0: case GGML_TYPE_Q4_1: case GGML_TYPE_Q5_0: case GGML_TYPE_Q5_1:
|
||||
case GGML_TYPE_Q8_0:
|
||||
return 32u;
|
||||
case GGML_TYPE_Q1_0:
|
||||
return 128u;
|
||||
case GGML_TYPE_Q2_0:
|
||||
return 64u;
|
||||
default:
|
||||
return 256u;
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
layout (push_constant) uniform parameter
|
||||
{
|
||||
@@ -72,21 +97,73 @@ layout (push_constant) uniform parameter
|
||||
} p;
|
||||
|
||||
|
||||
#ifndef MULMAT_QUANT
|
||||
layout (binding = 0) readonly buffer A {A_TYPE data_a[];};
|
||||
#else
|
||||
layout (binding = 0) readonly buffer A {uint8_t data_a[];};
|
||||
#endif
|
||||
layout (binding = 1) readonly buffer B {B_TYPE data_b[];};
|
||||
layout (binding = 2) writeonly buffer D {D_TYPE data_d[];};
|
||||
#if defined(MUL_MAT_ID) && defined(GGML_VULKAN_COOPMAT2_DECODE_VECTOR)
|
||||
layout (binding = 1) readonly buffer B4 {B_TYPEV4 data_b_v4[];};
|
||||
#endif
|
||||
|
||||
#if QUANT_K > 1
|
||||
#if defined(MULMAT_QUANT) || QUANT_K > 1
|
||||
#include "dequant_funcs_cm2.glsl"
|
||||
#ifndef MULMAT_QUANT
|
||||
// Per-type path: use the alias set by dequant_funcs_cm2.glsl
|
||||
#if defined(dequantFuncA_v) && defined(GGML_VULKAN_COOPMAT2_DECODE_VECTOR)
|
||||
#define DECODEFUNCA , dequantFuncA, dequantFuncA_v
|
||||
#else
|
||||
#define DECODEFUNCA , dequantFuncA
|
||||
#endif
|
||||
#else
|
||||
layout(buffer_reference, std430, buffer_reference_align = 1) buffer decodeBufA {
|
||||
uint8_t raw[MmABlockBytes];
|
||||
};
|
||||
|
||||
float16_t mmDecodeA(const in decodeBufA bl_in, const in uint blockCoords[2], const in uint coordInBlock[2]) {
|
||||
switch (MmTypeA) {
|
||||
case GGML_TYPE_Q1_0: return dequantFuncQ1_0 (decodeBufQ1_0 (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q2_0: return dequantFuncQ2_0 (decodeBufQ2_0 (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q4_0: return dequantFuncQ4_0 (decodeBufQ4_0 (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q4_1: return dequantFuncQ4_1 (decodeBufQ4_1 (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q5_0: return dequantFuncQ5_0 (decodeBufQ5_0 (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q5_1: return dequantFuncQ5_1 (decodeBufQ5_1 (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q8_0: return dequantFuncQ8_0 (decodeBufQ8_0 (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q2_K: return dequantFuncQ2_K (decodeBufQ2_K (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q3_K: return dequantFuncQ3_K (decodeBufQ3_K (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q6_K: return dequantFuncQ6_K (decodeBufQ6_K (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_TQ1_0: return dequantFuncTQ1_0(decodeBufTQ1_0(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_TQ2_0: return dequantFuncTQ2_0(decodeBufTQ2_0(bl_in), blockCoords, coordInBlock);
|
||||
default: return float16_t(0);
|
||||
}
|
||||
}
|
||||
|
||||
#ifdef GGML_VULKAN_COOPMAT2_DECODE_VECTOR
|
||||
f16vec4 mmDecodeA_v(const in decodeBufA bl_in, const in uint blockCoords[2], const in uint coordInBlock[2]) {
|
||||
switch (MmTypeA) {
|
||||
case GGML_TYPE_Q1_0: return dequantFuncQ1_0_v (decodeBufQ1_0 (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q2_0: return dequantFuncQ2_0_v (decodeBufQ2_0 (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q4_0: return dequantFuncQ4_0_v (decodeBufQ4_0 (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q4_1: return dequantFuncQ4_1_v (decodeBufQ4_1 (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q5_0: return dequantFuncQ5_0_v (decodeBufQ5_0 (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q5_1: return dequantFuncQ5_1_v (decodeBufQ5_1 (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q8_0: return dequantFuncQ8_0_v (decodeBufQ8_0 (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q2_K: return dequantFuncQ2_K_v (decodeBufQ2_K (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q3_K: return dequantFuncQ3_K_v (decodeBufQ3_K (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_Q6_K: return dequantFuncQ6_K_v (decodeBufQ6_K (bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_TQ1_0: return dequantFuncTQ1_0_v(decodeBufTQ1_0(bl_in), blockCoords, coordInBlock);
|
||||
case GGML_TYPE_TQ2_0: return dequantFuncTQ2_0_v(decodeBufTQ2_0(bl_in), blockCoords, coordInBlock);
|
||||
default: return f16vec4(0);
|
||||
}
|
||||
}
|
||||
#define DECODEFUNCA , mmDecodeA, mmDecodeA_v
|
||||
#else
|
||||
#define DECODEFUNCA , mmDecodeA
|
||||
#endif
|
||||
#endif
|
||||
#else
|
||||
#define DECODEFUNCA
|
||||
#endif
|
||||
|
||||
@@ -114,7 +191,6 @@ layout(buffer_reference, std430, buffer_reference_align = 2) buffer decodeBufB {
|
||||
};
|
||||
|
||||
uint _ne1;
|
||||
layout (constant_id = 6) const uint subgroup_size = 32;
|
||||
shared uvec4 ballots_sh[BLOCK_SIZE / subgroup_size];
|
||||
|
||||
B_TYPE decodeFuncB(const in decodeBufB bl, const in uint blockCoords[2], const in uint coordInBlock[2])
|
||||
@@ -246,6 +322,10 @@ void load_row_ids_hoisted(uint expert_idx, uint ic) {
|
||||
}
|
||||
#endif
|
||||
|
||||
#ifdef MULMAT_QUANT
|
||||
#include "iq_shmem_init.glsl"
|
||||
#endif
|
||||
|
||||
void main() {
|
||||
const uint tid = gl_LocalInvocationIndex;
|
||||
const uint ic = gl_WorkGroupID.y;
|
||||
@@ -264,7 +344,7 @@ void main() {
|
||||
#endif
|
||||
#endif
|
||||
|
||||
#ifdef NEEDS_INIT_IQ_SHMEM
|
||||
#if defined(NEEDS_INIT_IQ_SHMEM) || defined(MULMAT_QUANT)
|
||||
init_iq_shmem(gl_WorkGroupSize);
|
||||
#endif
|
||||
|
||||
@@ -305,22 +385,33 @@ void main() {
|
||||
const uint end_k = min(p.K, (ik + 1) * p.k_split);
|
||||
#endif
|
||||
|
||||
#ifdef MULMAT_QUANT
|
||||
const uint qk = mm_quant_k();
|
||||
#else
|
||||
const uint qk = QUANT_K;
|
||||
#endif
|
||||
|
||||
#ifdef MUL_MAT_ID
|
||||
uint pos_a = expert_idx * (p.batch_stride_a / QUANT_K);
|
||||
uint pos_a = expert_idx * (p.batch_stride_a / qk);
|
||||
uint pos_b = 0;
|
||||
#else
|
||||
uint pos_a = batch_idx_a * (p.batch_stride_a / QUANT_K);
|
||||
uint pos_a = batch_idx_a * (p.batch_stride_a / qk);
|
||||
uint pos_b = batch_idx * p.batch_stride_b;
|
||||
uint pos_d = batch_idx * p.batch_stride_d + ik * p.batch_stride_d * p.num_batches;
|
||||
#endif
|
||||
|
||||
uint stride_a = p.stride_a / QUANT_K;
|
||||
#ifdef MULMAT_QUANT
|
||||
// pos_a is a byte offset into the raw buffer; strides stay in block units
|
||||
pos_a *= MmABlockBytes;
|
||||
#endif
|
||||
|
||||
uint stride_a = p.stride_a / qk;
|
||||
uint stride_b = p.stride_b;
|
||||
|
||||
// Hint to the compiler that values are aligned (want 16B alignment).
|
||||
// Quants are always block-aligned, no alignment needed.
|
||||
if (ALIGNED != 0) {
|
||||
#if QUANT_K == 1
|
||||
#if !defined(MULMAT_QUANT) && QUANT_K == 1
|
||||
stride_a &= ~7;
|
||||
#endif
|
||||
stride_b &= ~7;
|
||||
@@ -335,10 +426,8 @@ void main() {
|
||||
#endif
|
||||
tensorLayoutNV<2, gl_CooperativeMatrixClampModeConstantNV> tensorLayoutD = createTensorLayoutNV(2, gl_CooperativeMatrixClampModeConstantNV);
|
||||
|
||||
#if QUANT_K > 1
|
||||
tensorLayoutA = setTensorLayoutBlockSizeNV(tensorLayoutA, 1, QUANT_K);
|
||||
tensorLayoutAClamp = setTensorLayoutBlockSizeNV(tensorLayoutAClamp, 1, QUANT_K);
|
||||
#endif
|
||||
tensorLayoutA = setTensorLayoutBlockSizeNV(tensorLayoutA, 1, qk);
|
||||
tensorLayoutAClamp = setTensorLayoutBlockSizeNV(tensorLayoutAClamp, 1, qk);
|
||||
#if defined(MUL_MAT_ID) && defined(GGML_VULKAN_COOPMAT2_DECODE_VECTOR)
|
||||
tensorLayoutB = setTensorLayoutBlockSizeNV(tensorLayoutB, 1, BK);
|
||||
#endif
|
||||
@@ -368,19 +457,19 @@ void main() {
|
||||
const uint START_ALIGN_K = 256;
|
||||
// For Qi_K (block size 256), unroll whole 256 element tiles.
|
||||
// For legacy quants (block size 32), unroll 8x.
|
||||
const uint UNROLL_K = (QUANT_K == 256) ? 256 : (BK * 8);
|
||||
const uint UNROLL_K = (qk == 256) ? 256 : (BK * 8);
|
||||
const uint unroll_count = UNROLL_K / BK;
|
||||
|
||||
// Detect a fast path where all loads are entirely in bounds and no clamping is required
|
||||
if ((ir + 1) * BM <= p.M && (ic + 1) * BN <= p.padded_N && (start_k % START_ALIGN_K) == 0 && (end_k % BK) == 0 &&
|
||||
#if QUANT_K == 1
|
||||
#if !defined(MULMAT_QUANT) && QUANT_K == 1
|
||||
(stride_a % 8) == 0 &&
|
||||
#endif
|
||||
(stride_b % 8) == 0) {
|
||||
// Hint to the compiler that values are aligned (want 16B alignment)
|
||||
start_k &= ~(START_ALIGN_K-1);
|
||||
stride_b &= ~7;
|
||||
#if QUANT_K == 1
|
||||
#if !defined(MULMAT_QUANT) && QUANT_K == 1
|
||||
stride_a &= ~7;
|
||||
#endif
|
||||
|
||||
@@ -551,10 +640,10 @@ void main() {
|
||||
[[dont_unroll]]
|
||||
for (uint block_k = start_k, i = 0; i < k_iters; block_k += BK, ++i) {
|
||||
|
||||
if ((block_k % QUANT_K) == 0) {
|
||||
if ((block_k % qk) == 0) {
|
||||
store_scales(tid);
|
||||
}
|
||||
if (block_k + BK < end_k && ((block_k + BK) % QUANT_K) == 0) {
|
||||
if (block_k + BK < end_k && ((block_k + BK) % qk) == 0) {
|
||||
fetch_scales(ir * BM, pos_a, stride_a, block_k + BK, tid, false);
|
||||
}
|
||||
|
||||
@@ -595,10 +684,10 @@ void main() {
|
||||
[[dont_unroll]]
|
||||
for (uint block_k = start_k, i = 0; i < k_iters; block_k += BK, ++i) {
|
||||
|
||||
if ((block_k % QUANT_K) == 0) {
|
||||
if ((block_k % qk) == 0) {
|
||||
store_scales(tid);
|
||||
}
|
||||
if (block_k + BK < end_k && ((block_k + BK) % QUANT_K) == 0) {
|
||||
if (block_k + BK < end_k && ((block_k + BK) % qk) == 0) {
|
||||
fetch_scales(ir * BM, pos_a, stride_a, block_k + BK, tid, false);
|
||||
}
|
||||
|
||||
@@ -639,10 +728,10 @@ void main() {
|
||||
[[dont_unroll]]
|
||||
for (uint block_k = start_k, i = 0; i < k_iters; block_k += BK, ++i) {
|
||||
|
||||
if ((block_k % QUANT_K) == 0) {
|
||||
if ((block_k % qk) == 0) {
|
||||
store_scales(tid);
|
||||
}
|
||||
if (block_k + BK < end_k && ((block_k + BK) % QUANT_K) == 0) {
|
||||
if (block_k + BK < end_k && ((block_k + BK) % qk) == 0) {
|
||||
fetch_scales(ir * BM, pos_a, stride_a, block_k + BK, tid, false);
|
||||
}
|
||||
|
||||
|
||||
File diff suppressed because it is too large
Load Diff
@@ -955,6 +955,7 @@ shared uint16_t iq1s_grid[2048];
|
||||
shared uint32_t iq1s_grid_gpu[2048];
|
||||
#endif
|
||||
|
||||
#if defined(DATA_A_IQ1_S) || defined(DATA_A_IQ1_M)
|
||||
#define NEEDS_INIT_IQ_SHMEM
|
||||
void init_iq_shmem(uvec3 wgsize)
|
||||
{
|
||||
@@ -978,6 +979,17 @@ void init_iq_shmem(uvec3 wgsize)
|
||||
barrier();
|
||||
}
|
||||
#endif
|
||||
#endif
|
||||
|
||||
#if defined(DATA_A_IQ2_XXS) || defined(DATA_A_IQ2_XS) || defined(DATA_A_IQ2_S)
|
||||
#if defined(DATA_A_IQ2_S)
|
||||
shared uvec2 iq2s_grid[1024];
|
||||
#elif defined(DATA_A_IQ2_XS)
|
||||
shared uvec2 iq2xs_grid[512];
|
||||
#else
|
||||
shared uvec2 iq2xxs_grid[256];
|
||||
#endif
|
||||
#endif
|
||||
|
||||
#define QUANT_K_IQ2_XXS 256
|
||||
#define QUANT_R_IQ2_XXS 1
|
||||
@@ -1063,8 +1075,7 @@ const uvec2[256] iq2xxs_grid_const = {
|
||||
uvec2(0x08080808, 0x2b2b082b), uvec2(0x08192b08, 0x2b2b1908), uvec2(0x19190808, 0x2b2b2b08), uvec2(0x08081908, 0x2b2b2b19)
|
||||
};
|
||||
|
||||
shared uvec2 iq2xxs_grid[256];
|
||||
|
||||
#if defined(DATA_A_IQ2_XXS)
|
||||
#define NEEDS_INIT_IQ_SHMEM
|
||||
void init_iq_shmem(uvec3 wgsize)
|
||||
{
|
||||
@@ -1076,12 +1087,15 @@ void init_iq_shmem(uvec3 wgsize)
|
||||
}
|
||||
barrier();
|
||||
}
|
||||
#endif
|
||||
|
||||
#if defined(DATA_A_IQ2_XXS)
|
||||
#define QUANT_K QUANT_K_IQ2_XXS
|
||||
#define QUANT_R QUANT_R_IQ2_XXS
|
||||
#define A_TYPE block_iq2_xxs
|
||||
#define A_TYPE_PACKED16 block_iq2_xxs_packed16
|
||||
#endif
|
||||
#endif
|
||||
|
||||
#define QUANT_K_IQ2_XS 256
|
||||
#define QUANT_R_IQ2_XS 1
|
||||
@@ -1233,8 +1247,7 @@ const uvec2 iq2xs_grid_const[512] = {
|
||||
uvec2(0x082b2b08, 0x2b2b2b2b), uvec2(0x082b2b2b, 0x2b2b2b2b), uvec2(0x2b190819, 0x2b2b2b2b), uvec2(0x2b2b2b2b, 0x2b2b2b2b),
|
||||
};
|
||||
|
||||
shared uvec2 iq2xs_grid[512];
|
||||
|
||||
#if defined(DATA_A_IQ2_XS)
|
||||
#define NEEDS_INIT_IQ_SHMEM
|
||||
void init_iq_shmem(uvec3 wgsize)
|
||||
{
|
||||
@@ -1246,12 +1259,15 @@ void init_iq_shmem(uvec3 wgsize)
|
||||
}
|
||||
barrier();
|
||||
}
|
||||
#endif
|
||||
|
||||
#if defined(DATA_A_IQ2_XS)
|
||||
#define QUANT_K QUANT_K_IQ2_XS
|
||||
#define QUANT_R QUANT_R_IQ2_XS
|
||||
#define A_TYPE block_iq2_xs
|
||||
#define A_TYPE_PACKED16 block_iq2_xs_packed16
|
||||
#endif
|
||||
#endif
|
||||
|
||||
#define QUANT_K_IQ2_S 256
|
||||
#define QUANT_R_IQ2_S 1
|
||||
@@ -1533,8 +1549,7 @@ const uvec2 iq2s_grid_const[1024] = {
|
||||
uvec2(0x082b082b, 0x2b2b2b2b), uvec2(0x082b2b08, 0x2b2b2b2b), uvec2(0x2b082b08, 0x2b2b2b2b), uvec2(0x2b2b2b2b, 0x2b2b2b2b)
|
||||
};
|
||||
|
||||
shared uvec2 iq2s_grid[1024];
|
||||
|
||||
#if defined(DATA_A_IQ2_S)
|
||||
#define NEEDS_INIT_IQ_SHMEM
|
||||
void init_iq_shmem(uvec3 wgsize)
|
||||
{
|
||||
@@ -1546,12 +1561,23 @@ void init_iq_shmem(uvec3 wgsize)
|
||||
}
|
||||
barrier();
|
||||
}
|
||||
#endif
|
||||
|
||||
#if defined(DATA_A_IQ2_S)
|
||||
#define QUANT_K QUANT_K_IQ2_S
|
||||
#define QUANT_R QUANT_R_IQ2_S
|
||||
#define A_TYPE block_iq2_s
|
||||
#define A_TYPE_PACKED16 block_iq2_s_packed16
|
||||
#endif
|
||||
#endif
|
||||
|
||||
#if defined(DATA_A_IQ3_XXS) || defined(DATA_A_IQ3_S)
|
||||
#if defined(DATA_A_IQ3_S)
|
||||
shared uint32_t iq3s_grid[512];
|
||||
#else
|
||||
shared uint32_t iq3xxs_grid[256];
|
||||
#endif
|
||||
#endif
|
||||
|
||||
#define QUANT_K_IQ3_XXS 256
|
||||
#define QUANT_R_IQ3_XXS 1
|
||||
@@ -1605,8 +1631,7 @@ const uint32_t iq3xxs_grid_const[256] = {
|
||||
0x3e1c1c1c, 0x3e1c3404, 0x3e24140c, 0x3e24240c, 0x3e2c0404, 0x3e2c0414, 0x3e2c1424, 0x3e341c04,
|
||||
};
|
||||
|
||||
shared uint32_t iq3xxs_grid[256];
|
||||
|
||||
#if defined(DATA_A_IQ3_XXS)
|
||||
#define NEEDS_INIT_IQ_SHMEM
|
||||
void init_iq_shmem(uvec3 wgsize)
|
||||
{
|
||||
@@ -1618,12 +1643,15 @@ void init_iq_shmem(uvec3 wgsize)
|
||||
}
|
||||
barrier();
|
||||
}
|
||||
#endif
|
||||
|
||||
#if defined(DATA_A_IQ3_XXS)
|
||||
#define QUANT_K QUANT_K_IQ3_XXS
|
||||
#define QUANT_R QUANT_R_IQ3_XXS
|
||||
#define A_TYPE block_iq3_xxs
|
||||
#define A_TYPE_PACKED16 block_iq3_xxs_packed16
|
||||
#endif
|
||||
#endif
|
||||
|
||||
#define QUANT_K_IQ3_S 256
|
||||
#define QUANT_R_IQ3_S 1
|
||||
@@ -1715,8 +1743,7 @@ const uint32_t iq3s_grid_const[512] = {
|
||||
0x0f090307, 0x0f090501, 0x0f090b01, 0x0f0b0505, 0x0f0b0905, 0x0f0d0105, 0x0f0d0703, 0x0f0f0101,
|
||||
};
|
||||
|
||||
shared uint32_t iq3s_grid[512];
|
||||
|
||||
#if defined(DATA_A_IQ3_S)
|
||||
#define NEEDS_INIT_IQ_SHMEM
|
||||
void init_iq_shmem(uvec3 wgsize)
|
||||
{
|
||||
@@ -1728,12 +1755,15 @@ void init_iq_shmem(uvec3 wgsize)
|
||||
}
|
||||
barrier();
|
||||
}
|
||||
#endif
|
||||
|
||||
#if defined(DATA_A_IQ3_S)
|
||||
#define QUANT_K QUANT_K_IQ3_S
|
||||
#define QUANT_R QUANT_R_IQ3_S
|
||||
#define A_TYPE block_iq3_s
|
||||
#define A_TYPE_PACKED16 block_iq3_s_packed16
|
||||
#endif
|
||||
#endif
|
||||
|
||||
#define QUANT_K_IQ4_XS 256
|
||||
#define QUANT_R_IQ4_XS 1
|
||||
@@ -1847,6 +1877,7 @@ const int8_t kvalues_iq4nl_const[16] = {
|
||||
|
||||
shared FLOAT_TYPE kvalues_iq4nl[16];
|
||||
|
||||
#if defined(DATA_A_IQ4_NL) || defined(DATA_A_IQ4_XS)
|
||||
#define NEEDS_INIT_IQ_SHMEM
|
||||
void init_iq_shmem(uvec3 wgsize)
|
||||
{
|
||||
@@ -1857,6 +1888,7 @@ void init_iq_shmem(uvec3 wgsize)
|
||||
barrier();
|
||||
}
|
||||
#endif
|
||||
#endif
|
||||
|
||||
#if defined(DATA_A_MXFP4) || defined(DATA_A_NVFP4)
|
||||
#if !defined(USE_OCP_FP4)
|
||||
@@ -1886,7 +1918,7 @@ float ue4m3_to_fp32_build(uint u) {
|
||||
}
|
||||
#endif
|
||||
|
||||
#if !defined(USE_OCP_FP4)
|
||||
#if (defined(DATA_A_MXFP4) || defined(DATA_A_NVFP4)) && !defined(USE_OCP_FP4)
|
||||
#define NEEDS_INIT_IQ_SHMEM
|
||||
void init_iq_shmem(uvec3 wgsize)
|
||||
{
|
||||
|
||||
@@ -260,6 +260,17 @@ bool is_iq_quant(const std::string& type_name) {
|
||||
return string_starts_with(type_name, "iq");
|
||||
}
|
||||
|
||||
bool is_lut_quant(const std::string& type_name) {
|
||||
return is_iq_quant(type_name) || type_name == "mxfp4" || type_name == "nvfp4";
|
||||
}
|
||||
|
||||
std::string lut_load_vec_a(const std::string& type_name) {
|
||||
if (type_name == "iq1_s" || type_name == "iq1_m" || type_name == "iq2_xxs" || type_name == "iq2_xs" || type_name == "iq2_s") {
|
||||
return "8";
|
||||
}
|
||||
return "4";
|
||||
}
|
||||
|
||||
static const char path_separator = '/';
|
||||
|
||||
std::string join_paths(const std::string& path1, const std::string& path2) {
|
||||
@@ -609,20 +620,28 @@ void matmul_shaders(bool fp16, MatMulIdType matmul_id_type, bool coopmat, bool c
|
||||
}
|
||||
|
||||
for (const auto& tname : type_names) {
|
||||
std::string load_vec_quant = "2";
|
||||
if ((tname == "q1_0") || (tname == "q4_0") || (tname == "q4_1") || (tname == "q5_1") || (tname == "iq1_s") || (tname == "iq1_m") || (tname == "iq2_xxs") || (tname == "iq2_xs") || (tname == "iq2_s"))
|
||||
load_vec_quant = "8";
|
||||
else if ((tname == "q2_0") || (tname == "q5_0") || (tname == "q8_0") || (tname == "q2_k") || (tname == "q4_k") || (tname == "q5_k") || (tname == "iq3_xxs") || (tname == "iq3_s") || (tname == "iq4_xs") || (tname == "iq4_nl") || (tname == "mxfp4") || (tname == "nvfp4"))
|
||||
load_vec_quant = "4";
|
||||
|
||||
if (tname == "bf16") {
|
||||
continue;
|
||||
}
|
||||
|
||||
std::string data_a_key = "DATA_A_" + to_uppercase(tname);
|
||||
// For aligned matmul loads
|
||||
std::string load_vec_a = (coopmat2 || tname == "f32" || tname == "f16" || tname == "bf16") ? load_vec : load_vec_quant;
|
||||
// Float types keep per-type compilation (different accumulation loop structure)
|
||||
if (tname == "f32" || tname == "f16") {
|
||||
std::string data_a_key = "DATA_A_" + to_uppercase(tname);
|
||||
|
||||
const std::map<std::string, std::string> float_type_dict = {
|
||||
{"FLOAT_TYPE", FLOAT_TYPE(1, tname)},
|
||||
{"FLOAT_TYPEV2", FLOAT_TYPE(2, tname)},
|
||||
{"FLOAT_TYPEV4", FLOAT_TYPE(4, tname)},
|
||||
{"FLOAT_TYPEV8", FLOAT_TYPE(8, tname)},
|
||||
};
|
||||
|
||||
if (!coopmat2) {
|
||||
string_to_spv(shader_name + "_" + tname + "_f32" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"LOAD_VEC_A", load_vec}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f32}, {"B_TYPE_SCALAR", "float"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
|
||||
}
|
||||
continue;
|
||||
}
|
||||
|
||||
std::string data_a_key = "DATA_A_" + to_uppercase(tname);
|
||||
const std::map<std::string, std::string> float_type_dict = {
|
||||
{"FLOAT_TYPE", FLOAT_TYPE(1, tname)},
|
||||
{"FLOAT_TYPEV2", FLOAT_TYPE(2, tname)},
|
||||
@@ -630,30 +649,52 @@ void matmul_shaders(bool fp16, MatMulIdType matmul_id_type, bool coopmat, bool c
|
||||
{"FLOAT_TYPEV8", FLOAT_TYPE(8, tname)},
|
||||
};
|
||||
|
||||
// don't generate f32 variants for coopmat2
|
||||
if (!coopmat2) {
|
||||
string_to_spv(shader_name + "_" + tname + "_f32" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"LOAD_VEC_A", load_vec_a}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f32}, {"B_TYPE_SCALAR", "float"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
|
||||
}
|
||||
|
||||
if (tname != "f16" && tname != "f32") {
|
||||
string_to_spv(shader_name + "_" + tname + "_f16" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"LOAD_VEC_A", load_vec_a}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f16}, {"B_TYPE_SCALAR", "float16_t"}, {"B_TYPEV4", "f16vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
|
||||
}
|
||||
|
||||
#if defined(GGML_VULKAN_FLOAT_E2M1_GLSLC_SUPPORT) && defined(GGML_VULKAN_FLOAT_E4M3_GLSLC_SUPPORT)
|
||||
if ((coopmat || coopmat2) && (tname == "mxfp4" || tname == "nvfp4")) {
|
||||
if (!coopmat2) {
|
||||
string_to_spv(shader_name + "_" + tname + "_f32_ocp" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"USE_OCP_FP4", "1"}, {"LOAD_VEC_A", load_vec_a}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f32}, {"B_TYPE_SCALAR", "float"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
|
||||
}
|
||||
string_to_spv(shader_name + "_" + tname + "_f16_ocp" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"USE_OCP_FP4", "1"}, {"LOAD_VEC_A", load_vec_a}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f16}, {"B_TYPE_SCALAR", "float16_t"}, {"B_TYPEV4", "f16vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
|
||||
}
|
||||
#endif
|
||||
|
||||
#if defined(GGML_VULKAN_INTEGER_DOT_GLSLC_SUPPORT)
|
||||
// Integer dot mmq performs better with f32 accumulators (different shader, skip for dot2)
|
||||
if (!f16acc && !coopmat && !coopmat2 && !dot2 && (is_legacy_quant(tname) || is_k_quant(tname) || tname == "mxfp4")) {
|
||||
string_to_spv(shader_name + "_" + tname + "_q8_1", "mul_mmq.comp", merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"D_TYPE", "float"},}), fp16, coopmat, coopmat2, f16acc);
|
||||
}
|
||||
#endif
|
||||
|
||||
if (is_lut_quant(tname)) {
|
||||
std::string lva = lut_load_vec_a(tname);
|
||||
|
||||
string_to_spv(shader_name + "_" + tname + "_f16" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"LOAD_VEC_A", lva}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f16}, {"B_TYPE_SCALAR", "float16_t"}, {"B_TYPEV4", "f16vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
|
||||
|
||||
if (!coopmat2) {
|
||||
string_to_spv(shader_name + "_" + tname + "_f32" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"LOAD_VEC_A", lva}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f32}, {"B_TYPE_SCALAR", "float"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
|
||||
}
|
||||
|
||||
#if defined(GGML_VULKAN_FLOAT_E2M1_GLSLC_SUPPORT) && defined(GGML_VULKAN_FLOAT_E4M3_GLSLC_SUPPORT)
|
||||
if ((tname == "mxfp4" || tname == "nvfp4") && (coopmat || coopmat2)) {
|
||||
string_to_spv(shader_name + "_" + tname + "_f16_ocp" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"USE_OCP_FP4", "1"}, {"LOAD_VEC_A", lva}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f16}, {"B_TYPE_SCALAR", "float16_t"}, {"B_TYPEV4", "f16vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
|
||||
if (!coopmat2) {
|
||||
string_to_spv(shader_name + "_" + tname + "_f32_ocp" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"USE_OCP_FP4", "1"}, {"LOAD_VEC_A", lva}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f32}, {"B_TYPE_SCALAR", "float"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
|
||||
}
|
||||
}
|
||||
#endif
|
||||
continue;
|
||||
}
|
||||
|
||||
// dedicated shader needed due to regression on Ampere
|
||||
if (coopmat2 && (tname == "q4_k" || tname == "q5_k")) {
|
||||
string_to_spv(shader_name + "_" + tname + "_f16" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, float_type_dict), {{data_a_key, "1"}, {"LOAD_VEC_A", load_vec}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f16}, {"B_TYPE_SCALAR", "float16_t"}, {"B_TYPEV4", "f16vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
|
||||
}
|
||||
}
|
||||
|
||||
// Quant shader: one SPIR-V for all quant types, selected via MmTypeA spec constant
|
||||
{
|
||||
const std::map<std::string, std::string> quant_float_type_dict = {
|
||||
{"FLOAT_TYPE", FLOAT_TYPE(1, "q4_0")},
|
||||
{"FLOAT_TYPEV2", FLOAT_TYPE(2, "q4_0")},
|
||||
{"FLOAT_TYPEV4", FLOAT_TYPE(4, "q4_0")},
|
||||
{"FLOAT_TYPEV8", FLOAT_TYPE(8, "q4_0")},
|
||||
};
|
||||
|
||||
string_to_spv(shader_name + "_quant_f16" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, quant_float_type_dict), {{"MULMAT_QUANT", "1"}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f16}, {"B_TYPE_SCALAR", "float16_t"}, {"B_TYPEV4", "f16vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
|
||||
|
||||
if (!coopmat2) {
|
||||
string_to_spv(shader_name + "_quant_f32" + dot2_sfx, source_name, merge_maps(merge_maps(base_dict, quant_float_type_dict), {{"MULMAT_QUANT", "1"}, {"LOAD_VEC_B", load_vec}, {"B_TYPE", aligned_b_type_f32}, {"B_TYPE_SCALAR", "float"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}}), fp16, coopmat, coopmat2, f16acc);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -761,7 +802,7 @@ void process_shaders() {
|
||||
for (const auto& tname : type_names) {
|
||||
// mul mat vec
|
||||
std::string data_a_key = "DATA_A_" + to_uppercase(tname);
|
||||
std::string shader = (string_ends_with(tname, "_k") || string_starts_with(tname, "iq1_") || string_starts_with(tname, "iq2_") || string_starts_with(tname, "iq3_") || tname == "tq2_0" || tname == "tq1_0") ? "mul_mat_vec_" + tname + ".comp" : "mul_mat_vec.comp";
|
||||
std::string shader = (string_ends_with(tname, "_k") || string_starts_with(tname, "iq1_") || string_starts_with(tname, "iq2_") || string_starts_with(tname, "iq3_") || tname == "iq4_xs" || tname == "tq2_0" || tname == "tq1_0") ? "mul_mat_vec_" + tname + ".comp" : "mul_mat_vec.comp";
|
||||
|
||||
string_to_spv("mul_mat_vec_" + tname + "_f32_f32", shader, merge_maps(base_dict, {{data_a_key, "1"}, {"B_TYPE", "float"}, {"B_TYPEV2", "vec2"}, {"B_TYPEV4", "vec4"}, {"D_TYPE", "float"}}));
|
||||
string_to_spv("mul_mat_vec_" + tname + "_f16_f32", shader, merge_maps(base_dict, {{data_a_key, "1"}, {"B_TYPE", "float16_t"}, {"B_TYPEV2", "f16vec2"}, {"B_TYPEV4", "f16vec4"}, {"D_TYPE", "float"}}));
|
||||
|
||||
@@ -7,7 +7,7 @@ dynamic = ["classifiers"]
|
||||
readme = "README.md"
|
||||
authors = [{name = "GGML", email = "ggml@ggml.ai"}]
|
||||
requires-python = '>=3.10'
|
||||
dependencies = ['numpy (>=1.17)', 'tqdm (>=4.27)', 'pyyaml (>=5.1)', 'requests (>=2.25)']
|
||||
dependencies = ['numpy (>=2.2.6)', 'tqdm (>=4.27)', 'pyyaml (>=5.1)', 'requests (>=2.25)']
|
||||
classifiers = [
|
||||
"Programming Language :: Python :: 3",
|
||||
"License :: OSI Approved :: MIT License",
|
||||
|
||||
+1
-1
@@ -1360,7 +1360,7 @@ extern "C" {
|
||||
LLAMA_API struct llama_sampler * llama_sampler_chain_get( struct llama_sampler * chain, int32_t i);
|
||||
|
||||
// the total number of samplers in the chain
|
||||
LLAMA_API int llama_sampler_chain_n (const struct llama_sampler * chain);
|
||||
LLAMA_API int32_t llama_sampler_chain_n (const struct llama_sampler * chain);
|
||||
|
||||
// after removing a sampler, the chain will no longer own it, and it will not be freed when the chain is freed
|
||||
LLAMA_API struct llama_sampler * llama_sampler_chain_remove( struct llama_sampler * chain, int32_t i);
|
||||
|
||||
@@ -55,6 +55,10 @@ llama_memory_hybrid_idx::llama_memory_hybrid_idx(
|
||||
// K-shift must not rotate them while the stream copies in the same update still apply
|
||||
hparams_idx.rope_type = LLAMA_ROPE_TYPE_NONE;
|
||||
|
||||
// fool llama_kv_cache into thinking this is a MLA cache, so it won't cache V tensors
|
||||
hparams_idx.n_embd_head_k_mla_impl = model.hparams.indexer_head_size;
|
||||
hparams_idx.n_embd_head_v_mla_impl = model.hparams.indexer_head_size;
|
||||
|
||||
LLAMA_LOG_INFO("%s: creating indexer KV cache, size = %u cells\n", __func__, kv_size);
|
||||
|
||||
return new llama_kv_cache(
|
||||
|
||||
@@ -1089,6 +1089,8 @@ const char * llm_type_name(llm_type type) {
|
||||
case LLM_TYPE_17B_16E: return "17Bx16E (Scout)";
|
||||
case LLM_TYPE_17B_128E: return "17Bx128E (Maverick)";
|
||||
case LLM_TYPE_A13B: return "A13B";
|
||||
case LLM_TYPE_1B_A400M: return "1B.A400M";
|
||||
case LLM_TYPE_3B_A800M: return "3B.A800M";
|
||||
case LLM_TYPE_7B_A1B: return "7B.A1B";
|
||||
case LLM_TYPE_8B_A1B: return "8B.A1B";
|
||||
case LLM_TYPE_7_9B_A1_3B: return "7.9B.A1.3B";
|
||||
@@ -1099,6 +1101,7 @@ const char * llm_type_name(llm_type type) {
|
||||
case LLM_TYPE_26B_A4B: return "26B.A4B";
|
||||
case LLM_TYPE_30B_A3B: return "30B.A3B";
|
||||
case LLM_TYPE_31B_A3_5B: return "31B.A3.5B";
|
||||
case LLM_TYPE_32B_A9B: return "32B.A9B";
|
||||
case LLM_TYPE_35B_A3B: return "35B.A3B";
|
||||
case LLM_TYPE_48B_A3B: return "48B.A3B";
|
||||
case LLM_TYPE_75B_A9B: return "75B.A9B";
|
||||
|
||||
@@ -116,6 +116,8 @@ enum llm_type {
|
||||
LLM_TYPE_17B_16E, // llama4 Scout
|
||||
LLM_TYPE_17B_128E, // llama4 Maverick
|
||||
LLM_TYPE_A13B,
|
||||
LLM_TYPE_1B_A400M, // Granite3 MoE
|
||||
LLM_TYPE_3B_A800M, // Granite3 MoE
|
||||
LLM_TYPE_7B_A1B,
|
||||
LLM_TYPE_8B_A1B, // lfm2moe
|
||||
LLM_TYPE_7_9B_A1_3B, // Ling-3.0-tiny
|
||||
@@ -126,6 +128,7 @@ enum llm_type {
|
||||
LLM_TYPE_26B_A4B, // Gemma4
|
||||
LLM_TYPE_30B_A3B,
|
||||
LLM_TYPE_31B_A3_5B,
|
||||
LLM_TYPE_32B_A9B, // Granite4 Hybrid
|
||||
LLM_TYPE_35B_A3B, // Qwen3.5
|
||||
LLM_TYPE_48B_A3B, // Kimi Linear
|
||||
LLM_TYPE_75B_A9B, // Nemotron 3 Puzzle
|
||||
|
||||
@@ -1006,7 +1006,7 @@ struct llama_sampler * llama_sampler_chain_remove(struct llama_sampler * chain,
|
||||
return result;
|
||||
}
|
||||
|
||||
int llama_sampler_chain_n(const struct llama_sampler * chain) {
|
||||
int32_t llama_sampler_chain_n(const struct llama_sampler * chain) {
|
||||
const auto * p = (const llama_sampler_chain *) chain->ctx;
|
||||
|
||||
return p->samplers.size();
|
||||
|
||||
+9
-21
@@ -29,15 +29,13 @@ void llama_model_bert::load_arch_tensors(llama_model_loader &) {
|
||||
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
|
||||
type_embd = create_tensor(tn(LLM_TENSOR_TOKEN_TYPES, "weight"), {n_embd, n_token_types}, TENSOR_NOT_REQUIRED);
|
||||
|
||||
if (arch == LLM_ARCH_BERT) {
|
||||
pos_embd = create_tensor(tn(LLM_TENSOR_POS_EMBD, "weight"), {n_embd, n_ctx_train}, 0);
|
||||
pos_embd = create_tensor(tn(LLM_TENSOR_POS_EMBD, "weight"), {n_embd, n_ctx_train}, 0);
|
||||
|
||||
cls = create_tensor(tn(LLM_TENSOR_CLS, "weight"), {n_embd, n_embd}, TENSOR_NOT_REQUIRED);
|
||||
cls_b = create_tensor(tn(LLM_TENSOR_CLS, "bias"), {n_embd}, TENSOR_NOT_REQUIRED);
|
||||
cls = create_tensor(tn(LLM_TENSOR_CLS, "weight"), {n_embd, n_embd}, TENSOR_NOT_REQUIRED);
|
||||
cls_b = create_tensor(tn(LLM_TENSOR_CLS, "bias"), {n_embd}, TENSOR_NOT_REQUIRED);
|
||||
|
||||
cls_out = create_tensor(tn(LLM_TENSOR_CLS_OUT, "weight"), {n_embd, hparams.n_cls_out}, TENSOR_NOT_REQUIRED);
|
||||
cls_out_b = create_tensor(tn(LLM_TENSOR_CLS_OUT, "bias"), {hparams.n_cls_out}, TENSOR_NOT_REQUIRED);
|
||||
}
|
||||
cls_out = create_tensor(tn(LLM_TENSOR_CLS_OUT, "weight"), {n_embd, hparams.n_cls_out}, TENSOR_NOT_REQUIRED);
|
||||
cls_out_b = create_tensor(tn(LLM_TENSOR_CLS_OUT, "bias"), {hparams.n_cls_out}, TENSOR_NOT_REQUIRED);
|
||||
|
||||
tok_norm = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD_NORM, "weight", 0), {n_embd}, 0);
|
||||
tok_norm_b = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD_NORM, "bias", 0), {n_embd}, 0);
|
||||
@@ -53,20 +51,10 @@ void llama_model_bert::load_arch_tensors(llama_model_loader &) {
|
||||
layer.attn_out_norm = create_tensor(tn(LLM_TENSOR_ATTN_OUT_NORM, "weight", i), {n_embd}, 0);
|
||||
layer.attn_out_norm_b = create_tensor(tn(LLM_TENSOR_ATTN_OUT_NORM, "bias", i), {n_embd}, 0);
|
||||
|
||||
if (hparams.moe_every_n_layers > 0 && i % hparams.moe_every_n_layers == 1) {
|
||||
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), { n_embd, n_ff, n_expert}, 0);
|
||||
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), { n_ff, n_embd, n_expert}, 0);
|
||||
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, 0);
|
||||
} else {
|
||||
layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), {n_embd, n_ff}, 0);
|
||||
layer.ffn_up_b = create_tensor(tn(LLM_TENSOR_FFN_UP, "bias", i), {n_ff}, TENSOR_NOT_REQUIRED);
|
||||
layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), {n_ff, n_embd}, 0);
|
||||
layer.ffn_down_b = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "bias", i), {n_embd}, TENSOR_NOT_REQUIRED);
|
||||
|
||||
if (arch == LLM_ARCH_NOMIC_BERT) {
|
||||
layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, 0);
|
||||
}
|
||||
}
|
||||
layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), {n_embd, n_ff}, 0);
|
||||
layer.ffn_up_b = create_tensor(tn(LLM_TENSOR_FFN_UP, "bias", i), {n_ff}, TENSOR_NOT_REQUIRED);
|
||||
layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), {n_ff, n_embd}, 0);
|
||||
layer.ffn_down_b = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "bias", i), {n_embd}, TENSOR_NOT_REQUIRED);
|
||||
|
||||
layer.layer_out_norm = create_tensor(tn(LLM_TENSOR_LAYER_OUT_NORM, "weight", i), {n_embd}, 0);
|
||||
layer.layer_out_norm_b = create_tensor(tn(LLM_TENSOR_LAYER_OUT_NORM, "bias", i), {n_embd}, 0);
|
||||
|
||||
@@ -30,7 +30,7 @@ void llama_model_granite_hybrid::load_arch_hparams(llama_model_loader & ml) {
|
||||
case 768: type = LLM_TYPE_350M; break;
|
||||
case 1536: type = (hparams.n_ff() == 512 ? LLM_TYPE_7B_A1B : LLM_TYPE_1B); break;
|
||||
case 2048: case 2560: type = LLM_TYPE_3B; break;
|
||||
case 4096: type = LLM_TYPE_32B; break;
|
||||
case 4096: type = LLM_TYPE_32B_A9B; break;
|
||||
default: type = LLM_TYPE_UNKNOWN;
|
||||
}
|
||||
|
||||
|
||||
@@ -8,8 +8,8 @@ void llama_model_granite_moe::load_arch_hparams(llama_model_loader & ml) {
|
||||
ml.get_key(LLM_KV_ATTENTION_SCALE, hparams.f_attention_scale, false);
|
||||
|
||||
switch (hparams.n_layer()) {
|
||||
case 32: type = LLM_TYPE_3B; break;
|
||||
case 40: type = LLM_TYPE_3B; break;
|
||||
case 24: type = LLM_TYPE_1B_A400M; break;
|
||||
case 32: type = LLM_TYPE_3B_A800M; break;
|
||||
// Add additional layer/vocab/etc checks here for other model sizes
|
||||
default: type = LLM_TYPE_UNKNOWN;
|
||||
}
|
||||
|
||||
+10
-1
@@ -38,7 +38,16 @@ void llama_model_granite::load_arch_hparams(llama_model_loader & ml) {
|
||||
|
||||
switch (hparams.n_layer()) {
|
||||
case 32: type = LLM_TYPE_3B; break;
|
||||
case 40: type = LLM_TYPE_3B; break;
|
||||
case 40: {
|
||||
switch (hparams.n_embd) {
|
||||
case 2048: type = LLM_TYPE_2B; break;
|
||||
case 2560: type = LLM_TYPE_3B; break;
|
||||
case 4096: type = LLM_TYPE_8B; break;
|
||||
default: type = LLM_TYPE_UNKNOWN;
|
||||
}
|
||||
break;
|
||||
}
|
||||
case 64: type = LLM_TYPE_30B; break;
|
||||
// Add additional layer/vocab/etc checks here for other model sizes
|
||||
default: type = LLM_TYPE_UNKNOWN;
|
||||
}
|
||||
|
||||
@@ -19,16 +19,6 @@ void llama_model_jina_bert_v3::load_arch_tensors(llama_model_loader &) {
|
||||
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
|
||||
type_embd = create_tensor(tn(LLM_TENSOR_TOKEN_TYPES, "weight"), {n_embd, n_token_types}, TENSOR_NOT_REQUIRED);
|
||||
|
||||
if (arch == LLM_ARCH_BERT) {
|
||||
pos_embd = create_tensor(tn(LLM_TENSOR_POS_EMBD, "weight"), {n_embd, n_ctx_train}, 0);
|
||||
|
||||
cls = create_tensor(tn(LLM_TENSOR_CLS, "weight"), {n_embd, n_embd}, TENSOR_NOT_REQUIRED);
|
||||
cls_b = create_tensor(tn(LLM_TENSOR_CLS, "bias"), {n_embd}, TENSOR_NOT_REQUIRED);
|
||||
|
||||
cls_out = create_tensor(tn(LLM_TENSOR_CLS_OUT, "weight"), {n_embd, hparams.n_cls_out}, TENSOR_NOT_REQUIRED);
|
||||
cls_out_b = create_tensor(tn(LLM_TENSOR_CLS_OUT, "bias"), {hparams.n_cls_out}, TENSOR_NOT_REQUIRED);
|
||||
}
|
||||
|
||||
tok_norm = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD_NORM, "weight", 0), {n_embd}, 0);
|
||||
tok_norm_b = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD_NORM, "bias", 0), {n_embd}, 0);
|
||||
|
||||
@@ -43,20 +33,10 @@ void llama_model_jina_bert_v3::load_arch_tensors(llama_model_loader &) {
|
||||
layer.attn_out_norm = create_tensor(tn(LLM_TENSOR_ATTN_OUT_NORM, "weight", i), {n_embd}, 0);
|
||||
layer.attn_out_norm_b = create_tensor(tn(LLM_TENSOR_ATTN_OUT_NORM, "bias", i), {n_embd}, 0);
|
||||
|
||||
if (hparams.moe_every_n_layers > 0 && i % hparams.moe_every_n_layers == 1) {
|
||||
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), { n_embd, n_ff, n_expert}, 0);
|
||||
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), { n_ff, n_embd, n_expert}, 0);
|
||||
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, 0);
|
||||
} else {
|
||||
layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), {n_embd, n_ff}, 0);
|
||||
layer.ffn_up_b = create_tensor(tn(LLM_TENSOR_FFN_UP, "bias", i), {n_ff}, TENSOR_NOT_REQUIRED);
|
||||
layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), {n_ff, n_embd}, 0);
|
||||
layer.ffn_down_b = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "bias", i), {n_embd}, TENSOR_NOT_REQUIRED);
|
||||
|
||||
if (arch == LLM_ARCH_NOMIC_BERT) {
|
||||
layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, 0);
|
||||
}
|
||||
}
|
||||
layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), {n_embd, n_ff}, 0);
|
||||
layer.ffn_up_b = create_tensor(tn(LLM_TENSOR_FFN_UP, "bias", i), {n_ff}, TENSOR_NOT_REQUIRED);
|
||||
layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), {n_ff, n_embd}, 0);
|
||||
layer.ffn_down_b = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "bias", i), {n_embd}, TENSOR_NOT_REQUIRED);
|
||||
|
||||
layer.layer_out_norm = create_tensor(tn(LLM_TENSOR_LAYER_OUT_NORM, "weight", i), {n_embd}, 0);
|
||||
layer.layer_out_norm_b = create_tensor(tn(LLM_TENSOR_LAYER_OUT_NORM, "bias", i), {n_embd}, 0);
|
||||
|
||||
@@ -4,12 +4,10 @@ void llama_model_nomic_bert_moe::load_arch_hparams(llama_model_loader & ml) {
|
||||
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_EPS, hparams.f_norm_eps);
|
||||
ml.get_key(LLM_KV_MOE_EVERY_N_LAYERS, hparams.moe_every_n_layers, 0);
|
||||
|
||||
if (hparams.n_layer() == 12 && hparams.n_embd == 768) {
|
||||
if (arch == LLM_ARCH_NOMIC_BERT) {
|
||||
type = LLM_TYPE_137M;
|
||||
} else if (arch == LLM_ARCH_NOMIC_BERT_MOE && hparams.moe_every_n_layers == 2) {
|
||||
type = LLM_TYPE_475M;
|
||||
}
|
||||
switch (hparams.n_layer()) {
|
||||
case 12:
|
||||
type = LLM_TYPE_475M; break;
|
||||
default: type = LLM_TYPE_UNKNOWN;
|
||||
}
|
||||
}
|
||||
|
||||
@@ -22,16 +20,6 @@ void llama_model_nomic_bert_moe::load_arch_tensors(llama_model_loader &) {
|
||||
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
|
||||
type_embd = create_tensor(tn(LLM_TENSOR_TOKEN_TYPES, "weight"), {n_embd, n_token_types}, TENSOR_NOT_REQUIRED);
|
||||
|
||||
if (arch == LLM_ARCH_BERT) {
|
||||
pos_embd = create_tensor(tn(LLM_TENSOR_POS_EMBD, "weight"), {n_embd, n_ctx_train}, 0);
|
||||
|
||||
cls = create_tensor(tn(LLM_TENSOR_CLS, "weight"), {n_embd, n_embd}, TENSOR_NOT_REQUIRED);
|
||||
cls_b = create_tensor(tn(LLM_TENSOR_CLS, "bias"), {n_embd}, TENSOR_NOT_REQUIRED);
|
||||
|
||||
cls_out = create_tensor(tn(LLM_TENSOR_CLS_OUT, "weight"), {n_embd, hparams.n_cls_out}, TENSOR_NOT_REQUIRED);
|
||||
cls_out_b = create_tensor(tn(LLM_TENSOR_CLS_OUT, "bias"), {hparams.n_cls_out}, TENSOR_NOT_REQUIRED);
|
||||
}
|
||||
|
||||
tok_norm = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD_NORM, "weight", 0), {n_embd}, 0);
|
||||
tok_norm_b = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD_NORM, "bias", 0), {n_embd}, 0);
|
||||
|
||||
@@ -55,10 +43,6 @@ void llama_model_nomic_bert_moe::load_arch_tensors(llama_model_loader &) {
|
||||
layer.ffn_up_b = create_tensor(tn(LLM_TENSOR_FFN_UP, "bias", i), {n_ff}, TENSOR_NOT_REQUIRED);
|
||||
layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), {n_ff, n_embd}, 0);
|
||||
layer.ffn_down_b = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "bias", i), {n_embd}, TENSOR_NOT_REQUIRED);
|
||||
|
||||
if (arch == LLM_ARCH_NOMIC_BERT) {
|
||||
layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, 0);
|
||||
}
|
||||
}
|
||||
|
||||
layer.layer_out_norm = create_tensor(tn(LLM_TENSOR_LAYER_OUT_NORM, "weight", i), {n_embd}, 0);
|
||||
|
||||
+10
-31
@@ -1,15 +1,12 @@
|
||||
#include "models.h"
|
||||
|
||||
void llama_model_nomic_bert::load_arch_hparams(llama_model_loader & ml) {
|
||||
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_EPS, hparams.f_norm_eps);
|
||||
ml.get_key(LLM_KV_MOE_EVERY_N_LAYERS, hparams.moe_every_n_layers, 0);
|
||||
ml.get_key(LLM_KV_ATTENTION_LAYERNORM_EPS, hparams.f_norm_eps);
|
||||
|
||||
if (hparams.n_layer() == 12 && hparams.n_embd == 768) {
|
||||
if (arch == LLM_ARCH_NOMIC_BERT) {
|
||||
type = LLM_TYPE_137M;
|
||||
} else if (arch == LLM_ARCH_NOMIC_BERT_MOE && hparams.moe_every_n_layers == 2) {
|
||||
type = LLM_TYPE_475M;
|
||||
}
|
||||
switch (hparams.n_layer()) {
|
||||
case 12:
|
||||
type = LLM_TYPE_137M; break;
|
||||
default: type = LLM_TYPE_UNKNOWN;
|
||||
}
|
||||
}
|
||||
|
||||
@@ -22,16 +19,6 @@ void llama_model_nomic_bert::load_arch_tensors(llama_model_loader &) {
|
||||
tok_embd = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD, "weight"), {n_embd, n_vocab}, 0);
|
||||
type_embd = create_tensor(tn(LLM_TENSOR_TOKEN_TYPES, "weight"), {n_embd, n_token_types}, TENSOR_NOT_REQUIRED);
|
||||
|
||||
if (arch == LLM_ARCH_BERT) {
|
||||
pos_embd = create_tensor(tn(LLM_TENSOR_POS_EMBD, "weight"), {n_embd, n_ctx_train}, 0);
|
||||
|
||||
cls = create_tensor(tn(LLM_TENSOR_CLS, "weight"), {n_embd, n_embd}, TENSOR_NOT_REQUIRED);
|
||||
cls_b = create_tensor(tn(LLM_TENSOR_CLS, "bias"), {n_embd}, TENSOR_NOT_REQUIRED);
|
||||
|
||||
cls_out = create_tensor(tn(LLM_TENSOR_CLS_OUT, "weight"), {n_embd, hparams.n_cls_out}, TENSOR_NOT_REQUIRED);
|
||||
cls_out_b = create_tensor(tn(LLM_TENSOR_CLS_OUT, "bias"), {hparams.n_cls_out}, TENSOR_NOT_REQUIRED);
|
||||
}
|
||||
|
||||
tok_norm = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD_NORM, "weight", 0), {n_embd}, 0);
|
||||
tok_norm_b = create_tensor(tn(LLM_TENSOR_TOKEN_EMBD_NORM, "bias", 0), {n_embd}, 0);
|
||||
|
||||
@@ -46,20 +33,12 @@ void llama_model_nomic_bert::load_arch_tensors(llama_model_loader &) {
|
||||
layer.attn_out_norm = create_tensor(tn(LLM_TENSOR_ATTN_OUT_NORM, "weight", i), {n_embd}, 0);
|
||||
layer.attn_out_norm_b = create_tensor(tn(LLM_TENSOR_ATTN_OUT_NORM, "bias", i), {n_embd}, 0);
|
||||
|
||||
if (hparams.moe_every_n_layers > 0 && i % hparams.moe_every_n_layers == 1) {
|
||||
layer.ffn_up_exps = create_tensor(tn(LLM_TENSOR_FFN_UP_EXPS, "weight", i), { n_embd, n_ff, n_expert}, 0);
|
||||
layer.ffn_down_exps = create_tensor(tn(LLM_TENSOR_FFN_DOWN_EXPS, "weight", i), { n_ff, n_embd, n_expert}, 0);
|
||||
layer.ffn_gate_inp = create_tensor(tn(LLM_TENSOR_FFN_GATE_INP, "weight", i), {n_embd, n_expert}, 0);
|
||||
} else {
|
||||
layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), {n_embd, n_ff}, 0);
|
||||
layer.ffn_up_b = create_tensor(tn(LLM_TENSOR_FFN_UP, "bias", i), {n_ff}, TENSOR_NOT_REQUIRED);
|
||||
layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), {n_ff, n_embd}, 0);
|
||||
layer.ffn_down_b = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "bias", i), {n_embd}, TENSOR_NOT_REQUIRED);
|
||||
layer.ffn_up = create_tensor(tn(LLM_TENSOR_FFN_UP, "weight", i), {n_embd, n_ff}, 0);
|
||||
layer.ffn_up_b = create_tensor(tn(LLM_TENSOR_FFN_UP, "bias", i), {n_ff}, TENSOR_NOT_REQUIRED);
|
||||
layer.ffn_down = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "weight", i), {n_ff, n_embd}, 0);
|
||||
layer.ffn_down_b = create_tensor(tn(LLM_TENSOR_FFN_DOWN, "bias", i), {n_embd}, TENSOR_NOT_REQUIRED);
|
||||
|
||||
if (arch == LLM_ARCH_NOMIC_BERT) {
|
||||
layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, 0);
|
||||
}
|
||||
}
|
||||
layer.ffn_gate = create_tensor(tn(LLM_TENSOR_FFN_GATE, "weight", i), {n_embd, n_ff}, 0);
|
||||
|
||||
layer.layer_out_norm = create_tensor(tn(LLM_TENSOR_LAYER_OUT_NORM, "weight", i), {n_embd}, 0);
|
||||
layer.layer_out_norm_b = create_tensor(tn(LLM_TENSOR_LAYER_OUT_NORM, "bias", i), {n_embd}, 0);
|
||||
|
||||
@@ -372,6 +372,7 @@ static bool is_webp_file(const unsigned char * buf, size_t len) {
|
||||
#ifdef MTMD_VIDEO
|
||||
static mtmd_bitmap * decode_webp_with_ffmpeg(const mtmd_context * mctx, const unsigned char * buf, size_t len, bool placeholder,
|
||||
const mtmd_helper_video_init_params & params);
|
||||
static void mtmd_helper_video_set_id(mtmd_helper_video * vctx, const std::string & id);
|
||||
#endif
|
||||
|
||||
mtmd_helper_bitmap_wrapper mtmd_helper_bitmap_init_from_buf(const mtmd_context * ctx, const unsigned char * buf, size_t len, bool placeholder,
|
||||
@@ -437,6 +438,7 @@ mtmd_helper_bitmap_wrapper mtmd_helper_bitmap_init_from_buf(const mtmd_context *
|
||||
LOG_ERR("%s: failed to decode buffer as either image/audio/video\n", __func__);
|
||||
return {nullptr, nullptr};
|
||||
}
|
||||
mtmd_helper_video_set_id(video_ctx, id); // propagate the hash to the frames
|
||||
result = mtmd_bitmap_init_lazy(ctx,
|
||||
id.empty() ? nullptr : id.c_str(),
|
||||
video_ctx,
|
||||
@@ -528,6 +530,7 @@ struct mtmd_helper_video {
|
||||
std::string ffprobe_bin;
|
||||
float fps_target = 0.0f;
|
||||
mtmd_helper_video_info info = {};
|
||||
std::string id; // hash of the input video
|
||||
|
||||
// RAII wrapper for managing subprocess
|
||||
struct subprocess_handle {
|
||||
@@ -786,9 +789,14 @@ struct mtmd_helper_video {
|
||||
}
|
||||
|
||||
LOG_DBG("%s: frame %d read OK\n", __func__, current_frame);
|
||||
current_frame++;
|
||||
mtmd_bitmap * frame = mtmd_bitmap_init(info.width, info.height, frame_buf.data());
|
||||
mtmd_bitmap_set_mergeable(frame, true);
|
||||
if (!id.empty()) {
|
||||
// each frame gets a unique id in the form of {hash}+{frame}, so that it can be identified in cache
|
||||
std::string frame_id = id + "+" + std::to_string(current_frame);
|
||||
mtmd_bitmap_set_id(frame, frame_id.c_str());
|
||||
}
|
||||
current_frame++;
|
||||
return frame;
|
||||
}
|
||||
|
||||
@@ -887,6 +895,10 @@ static std::string video_resolve_bin(const char * bin_dir, const char * name) {
|
||||
}
|
||||
|
||||
#ifdef MTMD_VIDEO
|
||||
static void mtmd_helper_video_set_id(mtmd_helper_video * vctx, const std::string & id) {
|
||||
vctx->id = id;
|
||||
}
|
||||
|
||||
static mtmd_bitmap * decode_webp_with_ffmpeg(const mtmd_context * mctx, const unsigned char * buf, size_t len, bool placeholder,
|
||||
const mtmd_helper_video_init_params & params) {
|
||||
mtmd_helper_video vctx;
|
||||
|
||||
@@ -2,7 +2,7 @@ aiohttp~=3.9.3
|
||||
pytest~=8.3.3
|
||||
pytest-xdist~=3.6
|
||||
filelock~=3.16
|
||||
numpy~=1.26.4
|
||||
numpy~=2.2.6
|
||||
openai~=2.14.0
|
||||
prometheus-client~=0.20.0
|
||||
requests~=2.32.3
|
||||
|
||||
Reference in New Issue
Block a user