mirror of
https://github.com/LostRuins/koboldcpp.git
synced 2026-09-20 01:31:42 +02:00
Merge commit 'ae9afff8d2c012ca760eb9c2adf41961cf6f6232' into concedo_experimental
# Conflicts: # .github/workflows/build-openvino.yml # .github/workflows/build-self-hosted.yml # .github/workflows/build-vulkan.yml # .github/workflows/build-wasm.yml # .github/workflows/build-webgpu.yml # .github/workflows/server-sanitize.yml # CMakeLists.txt # ci/run.sh # common/CMakeLists.txt # docs/backend/snapdragon/README.md # docs/backend/snapdragon/developer.md # docs/build.md # examples/speculative-simple/speculative-simple.cpp # ggml/src/ggml-cpu/CMakeLists.txt # ggml/src/ggml-hexagon/ggml-hexagon.cpp # ggml/src/ggml-hexagon/htp-opnode.h # ggml/src/ggml-hexagon/htp/act-ops.c # ggml/src/ggml-hexagon/htp/allreduce-ops.c # ggml/src/ggml-hexagon/htp/allreduce-ops.h # ggml/src/ggml-hexagon/htp/argsort-ops.c # ggml/src/ggml-hexagon/htp/binary-ops.c # ggml/src/ggml-hexagon/htp/concat-ops.c # ggml/src/ggml-hexagon/htp/cpy-ops.c # ggml/src/ggml-hexagon/htp/cumsum-ops.c # ggml/src/ggml-hexagon/htp/diag-ops.c # ggml/src/ggml-hexagon/htp/fill-ops.c # ggml/src/ggml-hexagon/htp/flash-attn-ops.c # ggml/src/ggml-hexagon/htp/flash-attn-ops.h # ggml/src/ggml-hexagon/htp/gated-delta-net-ops.c # ggml/src/ggml-hexagon/htp/get-rows-ops.c # ggml/src/ggml-hexagon/htp/hex-common.h # ggml/src/ggml-hexagon/htp/hex-utils.h # ggml/src/ggml-hexagon/htp/hmx-utils.h # ggml/src/ggml-hexagon/htp/htp-ctx.h # ggml/src/ggml-hexagon/htp/htp-ops.h # ggml/src/ggml-hexagon/htp/htp-tensor.c # ggml/src/ggml-hexagon/htp/htp-tensor.h # ggml/src/ggml-hexagon/htp/hvx-arith.h # ggml/src/ggml-hexagon/htp/hvx-div.h # ggml/src/ggml-hexagon/htp/hvx-inverse.h # ggml/src/ggml-hexagon/htp/hvx-scale.h # ggml/src/ggml-hexagon/htp/hvx-sigmoid.h # ggml/src/ggml-hexagon/htp/im2col-ops.c # ggml/src/ggml-hexagon/htp/main.c # ggml/src/ggml-hexagon/htp/matmul-ops.c # ggml/src/ggml-hexagon/htp/pad-ops.c # ggml/src/ggml-hexagon/htp/repeat-ops.c # ggml/src/ggml-hexagon/htp/rope-ops.c # ggml/src/ggml-hexagon/htp/set-rows-ops.c # ggml/src/ggml-hexagon/htp/softmax-ops.c # ggml/src/ggml-hexagon/htp/solve-tri-ops.c # ggml/src/ggml-hexagon/htp/ssm-conv.c # ggml/src/ggml-hexagon/htp/sum-rows-ops.c # ggml/src/ggml-hexagon/htp/unary-ops.c # ggml/src/ggml-metal/CMakeLists.txt # ggml/src/ggml-opencl/CMakeLists.txt # ggml/src/ggml-opencl/ggml-opencl.cpp # ggml/src/ggml-rpc/CMakeLists.txt # ggml/src/ggml-sycl/ggml-sycl.cpp # ggml/src/ggml-webgpu/CMakeLists.txt # ggml/src/ggml-webgpu/ggml-webgpu.cpp # scripts/snapdragon/run.py # scripts/sync_vendor.py # scripts/ui-assets.cmake # src/CMakeLists.txt # src/models/gemma4.cpp # tests/.gitignore # tests/CMakeLists.txt # tests/test-backend-ops.cpp # tests/test-jinja.cpp # tests/test-llama-archs.cpp # tests/test-quantize-fns.cpp # tests/test-save-load-state.cpp # tools/mtmd/CMakeLists.txt # tools/server/CMakeLists.txt
This commit is contained in:
Concedo
@@ -1712,6 +1712,10 @@ static enum ggml_status ggml_backend_sched_compute_splits(ggml_backend_sched_t s
|
||||
ggml_tensor * ids_tensor = node->src[2];
|
||||
ggml_backend_t ids_backend = split_backend;
|
||||
|
||||
if (ggml_nelements(ids_tensor) == 0) {
|
||||
continue;
|
||||
}
|
||||
|
||||
// if the ids tensor is also an input of the split, it may not have been copied yet to the split backend
|
||||
// in that case, we use the original ids tensor
|
||||
for (int i = input_id + 1; i < split->n_inputs; i++) {
|
||||
|
||||
@@ -18,7 +18,15 @@
|
||||
#endif
|
||||
#endif
|
||||
|
||||
// -Winterference-size was introduced in GCC 12
|
||||
#if defined(__cplusplus) && defined(__GNUC__) && !defined(__clang__) && __GNUC__ >= 12
|
||||
#pragma GCC diagnostic push
|
||||
#pragma GCC diagnostic ignored "-Winterference-size"
|
||||
#endif
|
||||
static const size_t CACHE_LINE_SIZE_F32 = CACHE_LINE_SIZE/sizeof(float);
|
||||
#if defined(__cplusplus) && defined(__GNUC__) && !defined(__clang__) && __GNUC__ >= 12
|
||||
#pragma GCC diagnostic pop
|
||||
#endif
|
||||
|
||||
// Work buffer size for im2col operations in CONV2D
|
||||
#define GGML_IM2COL_WORK_SIZE (16 * 1024 * 1024)
|
||||
|
||||
@@ -1129,12 +1129,21 @@ void launch_fattn(
|
||||
|
||||
dim3 blocks_num;
|
||||
if (stream_k) {
|
||||
// For short contexts it can be faster to have the SMs work on whole tiles because this lets us skip the fixup.
|
||||
const int max_blocks = max_blocks_per_sm*nsm;
|
||||
const int tiles_nwaves = (ntiles_dst + max_blocks - 1) / max_blocks;
|
||||
const int tiles_efficiency_percent = 100 * ntiles_dst / (max_blocks*tiles_nwaves);
|
||||
auto should_use_stream_k = [](const int cc, const int ntiles_dst, const int max_blocks, const int DKQ) {
|
||||
const int tiles_nwaves = (ntiles_dst + max_blocks - 1) / max_blocks;
|
||||
const int tiles_efficiency_percent = 100 * ntiles_dst / (max_blocks*tiles_nwaves);
|
||||
|
||||
const bool use_stream_k = cc >= GGML_CUDA_CC_ADA_LOVELACE || amd_wmma_available(cc) || tiles_efficiency_percent < 75;
|
||||
if (GGML_CUDA_CC_IS_NVIDIA(cc) && cc >= GGML_CUDA_CC_ADA_LOVELACE) {
|
||||
return true;
|
||||
}
|
||||
if (amd_wmma_available(cc) && DKQ == 64) {
|
||||
return true; // TODO better configuration
|
||||
}
|
||||
return tiles_efficiency_percent < 75;
|
||||
};
|
||||
|
||||
const int max_blocks = max_blocks_per_sm*nsm;
|
||||
const bool use_stream_k = should_use_stream_k(cc, ntiles_dst, max_blocks, Q->ne[0]);
|
||||
|
||||
blocks_num.x = ntiles_dst;
|
||||
blocks_num.y = 1;
|
||||
|
||||
@@ -158,8 +158,8 @@ static constexpr __host__ __device__ fattn_mma_config ggml_cuda_fattn_mma_get_co
|
||||
|
||||
GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 8, 64, 2, 32, 128, 128, 128, 1, true);
|
||||
GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 16, 64, 2, 32, 128, 128, 128, 1, true);
|
||||
GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 32, 128, 2, 64, 128, 128, 64, 1, true);
|
||||
GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 64, 128, 2, 64, 128, 128, 64, 1, true);
|
||||
GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 32, 256, 2, 64, 128, 128, 64, 1, true);
|
||||
GGML_CUDA_FATTN_MMA_CONFIG_CASE(256, 256, 64, 256, 2, 64, 128, 128, 64, 1, true);
|
||||
|
||||
GGML_CUDA_FATTN_MMA_CONFIG_CASE(320, 256, 32, 128, 2, 32, 160, 128, 128, 1, true);
|
||||
GGML_CUDA_FATTN_MMA_CONFIG_CASE(320, 256, 64, 128, 2, 32, 160, 128, 128, 1, true);
|
||||
@@ -1826,7 +1826,7 @@ static __global__ void flash_attn_ext_f16(
|
||||
#endif // __CUDA_ARCH__ == GGML_CUDA_CC_TURING
|
||||
|
||||
#if defined(AMD_WMMA_AVAILABLE)
|
||||
if (ncols1*ncols2 < 16 || ncols2 == 1 || DKQ > 128) {
|
||||
if (ncols1*ncols2 < 16 || ncols2 == 1 || DKQ > 256) {
|
||||
NO_DEVICE_CODE;
|
||||
return;
|
||||
}
|
||||
|
||||
@@ -279,6 +279,24 @@ static void ggml_cuda_flash_attn_ext_mma_f16_switch_ncols2(ggml_backend_cuda_con
|
||||
}
|
||||
}
|
||||
|
||||
// On RDNA it is preferable to minimize wasted compute vs. duplicate I/O for the mask.
|
||||
if (amd_wmma_available(cc)) {
|
||||
if (use_gqa_opt && gqa_ratio % 8 == 0) {
|
||||
ggml_cuda_flash_attn_ext_mma_f16_switch_ncols1<DKQ, DV, 8>(ctx, dst);
|
||||
return;
|
||||
}
|
||||
|
||||
if (use_gqa_opt && gqa_ratio % 4 == 0) {
|
||||
ggml_cuda_flash_attn_ext_mma_f16_switch_ncols1<DKQ, DV, 4>(ctx, dst);
|
||||
return;
|
||||
}
|
||||
|
||||
if (use_gqa_opt && gqa_ratio % 2 == 0) {
|
||||
ggml_cuda_flash_attn_ext_mma_f16_switch_ncols1<DKQ, DV, 2>(ctx, dst);
|
||||
return;
|
||||
}
|
||||
}
|
||||
|
||||
if (use_gqa_opt && gqa_ratio > 4) {
|
||||
ggml_cuda_flash_attn_ext_mma_f16_switch_ncols1<DKQ, DV, 8>(ctx, dst);
|
||||
return;
|
||||
@@ -704,8 +722,9 @@ static best_fattn_kernel ggml_cuda_get_best_fattn_kernel(const int device, const
|
||||
}
|
||||
}
|
||||
|
||||
// AMD WMMA is always faster than the tile kernel if the full tile width of 16 can be utilized.
|
||||
if ((amd_wmma_available(cc) && gqa_opt_applies && Q->ne[0] <= 128) && Q->ne[0] != 40 && Q->ne[0] != 72 && Q->ne[1] * gqa_ratio_eff > 8) {
|
||||
// AMD WMMA is faster than the tile kernel if the wide tiles with high arithmetic intensity can be utilized.
|
||||
if ((amd_wmma_available(cc) && gqa_opt_applies && Q->ne[0] <= 256) && Q->ne[0] != 40 && Q->ne[0] != 72 &&
|
||||
Q->ne[1] * gqa_ratio_eff > (Q->ne[0] <= 128 ? 8 : 16)) {
|
||||
return BEST_FATTN_KERNEL_MMA_F16;
|
||||
}
|
||||
|
||||
|
||||
@@ -0,0 +1,281 @@
|
||||
static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config_gcn(ggml_type type, int J, bool fallback) {
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_Q5_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_0, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_0, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_0, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_0, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_0, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_0, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_0, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_0, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_Q5_1, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_1, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_1, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_1, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_1, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_1, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_1, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_1, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_1, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_1, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_1, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 64, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 80, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 96, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 112, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 128, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
// ---------------------------------------------------------------------------------------------
|
||||
|
||||
CASE(GGML_TYPE_Q2_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_K, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_K, 512, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_K, 512, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_K, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_K, 512, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_K, 512, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_K, 512, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_K, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_K, 512, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_Q3_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q3_K, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q3_K, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q3_K, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q3_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q3_K, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q3_K, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q3_K, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q3_K, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q3_K, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q3_K, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_Q5_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_K, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_K, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_K, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_K, 256, 3, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_K, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_K, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_K, 256, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_K, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_K, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, false);
|
||||
|
||||
// ---------------------------------------------------------------------------------------------
|
||||
|
||||
CASE(GGML_TYPE_IQ1_S, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ1_S, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ1_S, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ1_S, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ1_S, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ1_S, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ1_S, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ1_S, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ1_S, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ1_S, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ1_S, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_IQ2_XXS, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_IQ2_XS, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XS, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XS, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XS, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XS, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XS, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XS, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XS, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XS, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XS, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XS, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_IQ2_S, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_S, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_S, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_S, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_S, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_S, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_S, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_S, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_S, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_S, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_S, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_IQ3_XXS, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_IQ3_S, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_S, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_S, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_S, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_S, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_S, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_S, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_S, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_S, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_S, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_S, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_IQ4_XS, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_XS, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_XS, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_XS, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_XS, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_XS, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_XS, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_XS, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_XS, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_XS, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_XS, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_IQ4_NL, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_NL, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_NL, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_NL, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_NL, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_NL, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_NL, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_NL, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_NL, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_NL, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_NL, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
|
||||
// ---------------------------------------------------------------------------------------------
|
||||
|
||||
CASE(GGML_TYPE_MXFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_MXFP4, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_MXFP4, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_MXFP4, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_MXFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_MXFP4, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_MXFP4, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_MXFP4, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_MXFP4, 256, 2, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_MXFP4, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_MXFP4, 512, 2, 128, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 8, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, false);
|
||||
|
||||
return ggml_cuda_mmq_config(GGML_TYPE_COUNT, 256, 2, 64, 64, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, 256, false, true);
|
||||
}
|
||||
@@ -219,6 +219,7 @@ struct ggml_cuda_mmq_config {
|
||||
#include "mmq-config-ampere.cuh"
|
||||
#include "mmq-config-blackwell.cuh"
|
||||
|
||||
#include "mmq-config-gcn.cuh"
|
||||
#include "mmq-config-cdna.cuh"
|
||||
#include "mmq-config-rdna2.cuh"
|
||||
#include "mmq-config-rdna3.cuh"
|
||||
@@ -229,6 +230,9 @@ struct ggml_cuda_mmq_config {
|
||||
|
||||
static __host__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(const ggml_type type, const int J, const bool fallback, const int cc) {
|
||||
if (GGML_CUDA_CC_IS_AMD(cc)) {
|
||||
if (GGML_CUDA_CC_IS_GCN(cc)) {
|
||||
return ggml_cuda_mmq_get_config_gcn(type, J, fallback);
|
||||
}
|
||||
if (GGML_CUDA_CC_IS_CDNA(cc)) {
|
||||
return ggml_cuda_mmq_get_config_cdna(type, J, fallback);
|
||||
}
|
||||
@@ -257,7 +261,9 @@ static __host__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(const ggml_type ty
|
||||
|
||||
static constexpr __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_config(ggml_type type, int J, bool fallback) {
|
||||
#ifdef GGML_USE_HIP
|
||||
#ifdef CDNA
|
||||
#ifdef GCN
|
||||
return ggml_cuda_mmq_get_config_gcn(type, J, fallback);
|
||||
#elif defined(CDNA)
|
||||
return ggml_cuda_mmq_get_config_cdna(type, J, fallback);
|
||||
#elif defined(RDNA4)
|
||||
return ggml_cuda_mmq_get_config_rdna4(type, J, fallback);
|
||||
|
||||
@@ -1,4 +1,5 @@
|
||||
#include "ggml-metal-common.h"
|
||||
#include "ggml-metal-fusion.h"
|
||||
|
||||
#include "ggml.h"
|
||||
#include "ggml-impl.h"
|
||||
@@ -390,59 +391,31 @@ static std::vector<int> ggml_metal_graph_optimize_reorder(const std::vector<node
|
||||
}
|
||||
|
||||
void ggml_graph_optimize(ggml_cgraph * gf) {
|
||||
constexpr int MAX_FUSE = 16;
|
||||
|
||||
const int n = gf->n_nodes;
|
||||
|
||||
enum ggml_op ops[MAX_FUSE];
|
||||
|
||||
std::vector<node_info> nodes;
|
||||
nodes.reserve(gf->n_nodes);
|
||||
|
||||
// fuse nodes:
|
||||
// we don't want to make reorders that break fusing, so we first pack all fusable tensors
|
||||
// and perform the reorder over the fused nodes. after the reorder is done, we unfuse
|
||||
//
|
||||
// the fusable sequences are declared in the fusion table (ggml-metal-fuse.cpp), so the
|
||||
// packing here is driven by the same patterns that the op encoders will later use
|
||||
for (int i = 0; i < n; i++) {
|
||||
node_info node = {
|
||||
/*.node =*/ gf->nodes[i],
|
||||
/*.fused =*/ {},
|
||||
};
|
||||
|
||||
// fuse only ops that start with these operations
|
||||
// can be expanded when needed
|
||||
if (node.op() == GGML_OP_ADD ||
|
||||
node.op() == GGML_OP_NORM ||
|
||||
node.op() == GGML_OP_RMS_NORM) {
|
||||
ops[0] = node.op();
|
||||
const int f = ggml_metal_fusion_max(gf, i);
|
||||
|
||||
int f = i + 1;
|
||||
while (f < n && f < i + MAX_FUSE) {
|
||||
// conservatively allow fusing only these ops
|
||||
// can be expanded when needed
|
||||
if (gf->nodes[f]->op != GGML_OP_ADD &&
|
||||
gf->nodes[f]->op != GGML_OP_MUL &&
|
||||
gf->nodes[f]->op != GGML_OP_NORM &&
|
||||
gf->nodes[f]->op != GGML_OP_RMS_NORM) {
|
||||
break;
|
||||
}
|
||||
ops[f - i] = gf->nodes[f]->op;
|
||||
f++;
|
||||
}
|
||||
// add the fused tensors into the node info so we can unfuse them later
|
||||
for (int k = 1; k < f; k++) {
|
||||
++i;
|
||||
|
||||
f -= i;
|
||||
for (; f > 1; f--) {
|
||||
if (ggml_can_fuse(gf, i, ops, f)) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
// add the fused tensors into the node info so we can unfuse them later
|
||||
for (int k = 1; k < f; k++) {
|
||||
++i;
|
||||
|
||||
// the .dst() becomes the last fused tensor
|
||||
node.add_fused(gf->nodes[i]);
|
||||
}
|
||||
// the .dst() becomes the last fused tensor
|
||||
node.add_fused(gf->nodes[i]);
|
||||
}
|
||||
|
||||
nodes.push_back(std::move(node));
|
||||
|
||||
@@ -33,6 +33,7 @@ ggml_metal_event_t ggml_metal_get_ev_cpy(ggml_metal_t ctx);
|
||||
|
||||
void ggml_metal_set_n_cb (ggml_metal_t ctx, int n_cb);
|
||||
void ggml_metal_set_abort_callback (ggml_metal_t ctx, ggml_abort_callback abort_callback, void * user_data);
|
||||
|
||||
bool ggml_metal_supports_family (ggml_metal_t ctx, int family);
|
||||
void ggml_metal_capture_next_compute(ggml_metal_t ctx);
|
||||
|
||||
|
||||
@@ -6,6 +6,7 @@
|
||||
#import "ggml-metal-impl.h"
|
||||
#import "ggml-metal-common.h"
|
||||
#import "ggml-metal-ops.h"
|
||||
#import "ggml-metal-fusion.h"
|
||||
|
||||
#import <Foundation/Foundation.h>
|
||||
|
||||
@@ -36,15 +37,12 @@ struct ggml_metal {
|
||||
// additional, inference-time compiled pipelines
|
||||
ggml_metal_pipelines_t pipelines_ext;
|
||||
|
||||
bool use_fusion;
|
||||
bool use_concurrency;
|
||||
bool use_graph_optimize;
|
||||
|
||||
int debug_graph;
|
||||
int debug_fusion;
|
||||
|
||||
// how many times a given op was fused
|
||||
uint64_t fuse_cnt[GGML_OP_COUNT];
|
||||
struct ggml_metal_fusion_info * finfo;
|
||||
|
||||
// capture state
|
||||
int capture_compute;
|
||||
@@ -139,7 +137,6 @@ ggml_metal_t ggml_metal_init(ggml_metal_device_t dev) {
|
||||
|
||||
res->d_queue = dispatch_queue_create("ggml-metal", DISPATCH_QUEUE_CONCURRENT);
|
||||
|
||||
res->use_fusion = getenv("GGML_METAL_FUSION_DISABLE") == nil;
|
||||
res->use_concurrency = getenv("GGML_METAL_CONCURRENCY_DISABLE") == nil;
|
||||
|
||||
{
|
||||
@@ -147,20 +144,19 @@ ggml_metal_t ggml_metal_init(ggml_metal_device_t dev) {
|
||||
res->debug_graph = val ? atoi(val) : 0;
|
||||
}
|
||||
|
||||
{
|
||||
const char * val = getenv("GGML_METAL_FUSION_DEBUG");
|
||||
res->debug_fusion = val ? atoi(val) : 0;
|
||||
}
|
||||
|
||||
res->use_graph_optimize = true;
|
||||
|
||||
if (getenv("GGML_METAL_GRAPH_OPTIMIZE_DISABLE") != NULL) {
|
||||
res->use_graph_optimize = false;
|
||||
}
|
||||
|
||||
memset(res->fuse_cnt, 0, sizeof(res->fuse_cnt));
|
||||
res->finfo = ggml_metal_device_get_fusion_info(dev);
|
||||
if (ggml_metal_fusion_info_stats(res->finfo)) {
|
||||
ggml_metal_fusion_info_labels_init(res->finfo);
|
||||
res->n_cb = 0;
|
||||
}
|
||||
|
||||
GGML_LOG_INFO("%s: use fusion = %s\n", __func__, res->use_fusion ? "true" : "false");
|
||||
GGML_LOG_INFO("%s: use fusion = %s\n", __func__, ggml_metal_fusion_info_enabled(res->finfo) ? "true" : "false");
|
||||
GGML_LOG_INFO("%s: use concurrency = %s\n", __func__, res->use_concurrency ? "true" : "false");
|
||||
GGML_LOG_INFO("%s: use graph optimize = %s\n", __func__, res->use_graph_optimize ? "true" : "false");
|
||||
|
||||
@@ -222,15 +218,18 @@ void ggml_metal_free(ggml_metal_t ctx) {
|
||||
ctx->pipelines_ext = nil;
|
||||
}
|
||||
|
||||
if (ctx->debug_fusion > 0) {
|
||||
if (ggml_metal_fusion_info_debug(ctx->finfo) > 0) {
|
||||
GGML_LOG_DEBUG("%s: fusion stats:\n", __func__);
|
||||
for (int i = 0; i < GGML_OP_COUNT; i++) {
|
||||
if (ctx->fuse_cnt[i] == 0) {
|
||||
|
||||
const int n_fusions = ggml_metal_fusion_info_n_fusions(ctx->finfo);
|
||||
for (int i = 0; i < n_fusions; i++) {
|
||||
const uint64_t count = ggml_metal_fusion_info_count(ctx->finfo, i);
|
||||
if (count == 0) {
|
||||
continue;
|
||||
}
|
||||
|
||||
// note: cannot use ggml_log here
|
||||
GGML_LOG_DEBUG("%s: - %s: %" PRIu64 "\n", __func__, ggml_op_name((enum ggml_op) i), ctx->fuse_cnt[i]);
|
||||
GGML_LOG_DEBUG("%s: - %s: %" PRIu64 "\n", __func__, ggml_metal_fusion_info_label(ctx->finfo, i), count);
|
||||
}
|
||||
}
|
||||
|
||||
@@ -481,10 +480,17 @@ enum ggml_status ggml_metal_graph_compute(ggml_metal_t ctx, struct ggml_cgraph *
|
||||
@autoreleasepool {
|
||||
ctx->gf = gf;
|
||||
|
||||
ctx->n_nodes_0 = MIN(n_main, gf->n_nodes);
|
||||
ctx->n_nodes_1 = gf->n_nodes - ctx->n_nodes_0;
|
||||
if (ctx->n_cb == 0) {
|
||||
// single-threaded encoding: the whole graph is encoded by one command buffer
|
||||
ctx->n_nodes_0 = gf->n_nodes;
|
||||
ctx->n_nodes_1 = 0;
|
||||
ctx->n_nodes_per_cb = 0;
|
||||
} else {
|
||||
ctx->n_nodes_0 = MIN(n_main, gf->n_nodes);
|
||||
ctx->n_nodes_1 = gf->n_nodes - ctx->n_nodes_0;
|
||||
|
||||
ctx->n_nodes_per_cb = (ctx->n_nodes_1 + ctx->n_cb - 1) / ctx->n_cb;
|
||||
ctx->n_nodes_per_cb = (ctx->n_nodes_1 + ctx->n_cb - 1) / ctx->n_cb;
|
||||
}
|
||||
|
||||
if (ctx->capture_compute >= 0) {
|
||||
ctx->capture_compute--;
|
||||
@@ -682,6 +688,12 @@ ggml_metal_event_t ggml_metal_get_ev_cpy(ggml_metal_t ctx) {
|
||||
}
|
||||
|
||||
void ggml_metal_set_n_cb(ggml_metal_t ctx, int n_cb) {
|
||||
// when fusion stats are collected the graph must be encoded by a single thread so the
|
||||
// counters are race-free; override whatever the caller requested
|
||||
if (ggml_metal_fusion_info_stats(ctx->finfo)) {
|
||||
n_cb = 0;
|
||||
}
|
||||
|
||||
if (ctx->n_cb != n_cb) {
|
||||
ctx->n_cb = MIN(n_cb, GGML_METAL_MAX_COMMAND_BUFFERS);
|
||||
|
||||
@@ -717,13 +729,12 @@ void ggml_metal_set_n_cb(ggml_metal_t ctx, int n_cb) {
|
||||
ctx->dev,
|
||||
cmd_buf,
|
||||
ctx->gf,
|
||||
ctx->finfo,
|
||||
idx_start,
|
||||
idx_end,
|
||||
ctx->use_fusion,
|
||||
ctx->use_concurrency,
|
||||
ctx->capture_compute,
|
||||
ctx->debug_graph,
|
||||
ctx->debug_fusion);
|
||||
ctx->debug_graph);
|
||||
|
||||
for (int idx = 0; idx < ggml_metal_op_n_nodes(ctx_op); ++idx) {
|
||||
const int res = ggml_metal_op_encode(ctx_op, idx);
|
||||
|
||||
@@ -932,12 +932,24 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta
|
||||
nsg = N_SG_IQ2_XXS;
|
||||
nr0 = N_R0_IQ2_XXS;
|
||||
smem = 256*8+128;
|
||||
|
||||
const int nb32 = ne00/32;
|
||||
if (nb32 < 32 && (32 % nb32) == 0) {
|
||||
nr0 = N_R0_IQ2_XXS_SPLIT;
|
||||
split = true;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_IQ2_XS:
|
||||
{
|
||||
nsg = N_SG_IQ2_XS;
|
||||
nr0 = N_R0_IQ2_XS;
|
||||
smem = 512*8+128;
|
||||
|
||||
const int nb32 = ne00/32;
|
||||
if (nb32 < 32 && (32 % nb32) == 0) {
|
||||
nr0 = N_R0_IQ2_XS_SPLIT;
|
||||
split = true;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_IQ3_XXS:
|
||||
{
|
||||
@@ -957,21 +969,45 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv(ggml_meta
|
||||
nsg = N_SG_IQ3_S;
|
||||
nr0 = N_R0_IQ3_S;
|
||||
smem = 512*4;
|
||||
|
||||
const int nb32 = ne00/32;
|
||||
if (nb32 < 32 && (32 % nb32) == 0) {
|
||||
nr0 = N_R0_IQ3_S_SPLIT;
|
||||
split = true;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_IQ2_S:
|
||||
{
|
||||
nsg = N_SG_IQ2_S;
|
||||
nr0 = N_R0_IQ2_S;
|
||||
|
||||
const int nb32 = ne00/32;
|
||||
if (nb32 < 32 && (32 % nb32) == 0) {
|
||||
nr0 = N_R0_IQ2_S_SPLIT;
|
||||
split = true;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_IQ1_S:
|
||||
{
|
||||
nsg = N_SG_IQ1_S;
|
||||
nr0 = N_R0_IQ1_S;
|
||||
|
||||
const int nb32 = ne00/32;
|
||||
if (nb32 < 32 && (32 % nb32) == 0) {
|
||||
nr0 = N_R0_IQ1_S_SPLIT;
|
||||
split = true;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_IQ1_M:
|
||||
{
|
||||
nsg = N_SG_IQ1_M;
|
||||
nr0 = N_R0_IQ1_M;
|
||||
|
||||
const int nb32 = ne00/32;
|
||||
if (nb32 < 32 && (32 % nb32) == 0) {
|
||||
nr0 = N_R0_IQ1_M_SPLIT;
|
||||
split = true;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_IQ4_NL:
|
||||
{
|
||||
@@ -1177,12 +1213,24 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id(ggml_m
|
||||
nsg = N_SG_IQ2_XXS;
|
||||
nr0 = N_R0_IQ2_XXS;
|
||||
smem = 256*8+128;
|
||||
|
||||
const int nb32 = ne00/32;
|
||||
if (nb32 < 32 && (32 % nb32) == 0) {
|
||||
nr0 = N_R0_IQ2_XXS_SPLIT;
|
||||
split = true;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_IQ2_XS:
|
||||
{
|
||||
nsg = N_SG_IQ2_XS;
|
||||
nr0 = N_R0_IQ2_XS;
|
||||
smem = 512*8+128;
|
||||
|
||||
const int nb32 = ne00/32;
|
||||
if (nb32 < 32 && (32 % nb32) == 0) {
|
||||
nr0 = N_R0_IQ2_XS_SPLIT;
|
||||
split = true;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_IQ3_XXS:
|
||||
{
|
||||
@@ -1202,21 +1250,45 @@ ggml_metal_pipeline_with_params ggml_metal_library_get_pipeline_mul_mv_id(ggml_m
|
||||
nsg = N_SG_IQ3_S;
|
||||
nr0 = N_R0_IQ3_S;
|
||||
smem = 512*4;
|
||||
|
||||
const int nb32 = ne00/32;
|
||||
if (nb32 < 32 && (32 % nb32) == 0) {
|
||||
nr0 = N_R0_IQ3_S_SPLIT;
|
||||
split = true;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_IQ2_S:
|
||||
{
|
||||
nsg = N_SG_IQ2_S;
|
||||
nr0 = N_R0_IQ2_S;
|
||||
|
||||
const int nb32 = ne00/32;
|
||||
if (nb32 < 32 && (32 % nb32) == 0) {
|
||||
nr0 = N_R0_IQ2_S_SPLIT;
|
||||
split = true;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_IQ1_S:
|
||||
{
|
||||
nsg = N_SG_IQ1_S;
|
||||
nr0 = N_R0_IQ1_S;
|
||||
|
||||
const int nb32 = ne00/32;
|
||||
if (nb32 < 32 && (32 % nb32) == 0) {
|
||||
nr0 = N_R0_IQ1_S_SPLIT;
|
||||
split = true;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_IQ1_M:
|
||||
{
|
||||
nsg = N_SG_IQ1_M;
|
||||
nr0 = N_R0_IQ1_M;
|
||||
|
||||
const int nb32 = ne00/32;
|
||||
if (nb32 < 32 && (32 % nb32) == 0) {
|
||||
nr0 = N_R0_IQ1_M_SPLIT;
|
||||
split = true;
|
||||
}
|
||||
} break;
|
||||
case GGML_TYPE_IQ4_NL:
|
||||
{
|
||||
|
||||
@@ -325,6 +325,11 @@ bool ggml_metal_device_supports_op(ggml_metal_device_t dev, const struct ggml_te
|
||||
|
||||
const struct ggml_metal_device_props * ggml_metal_device_get_props(ggml_metal_device_t dev);
|
||||
|
||||
struct ggml_metal_fusion_info;
|
||||
|
||||
// the device-owned fusion debugging context (NULL unless fusion debugging is enabled)
|
||||
struct ggml_metal_fusion_info * ggml_metal_device_get_fusion_info(ggml_metal_device_t dev);
|
||||
|
||||
//
|
||||
// device buffers
|
||||
//
|
||||
|
||||
@@ -1,4 +1,5 @@
|
||||
#import "ggml-metal-device.h"
|
||||
#import "ggml-metal-fusion.h"
|
||||
|
||||
#import "ggml-impl.h"
|
||||
#import "ggml-backend-impl.h"
|
||||
@@ -896,6 +897,9 @@ struct ggml_metal_device {
|
||||
|
||||
struct ggml_metal_device_props props;
|
||||
|
||||
// shared fusion debugging context
|
||||
struct ggml_metal_fusion_info * finfo;
|
||||
|
||||
// virtual address for GPU memory allocations
|
||||
atomic_uintptr_t addr_virt;
|
||||
};
|
||||
@@ -1280,6 +1284,13 @@ ggml_metal_device_t ggml_metal_device_init(int device, int n_devices) {
|
||||
dev->props.max_working_set_size = dev->mtl_device.maxBufferLength;
|
||||
}
|
||||
|
||||
{
|
||||
const char * val = getenv("GGML_METAL_FUSION_DEBUG");
|
||||
dev->finfo = ggml_metal_fusion_info_init(
|
||||
getenv("GGML_METAL_FUSION_DISABLE") == nil,
|
||||
val ? atoi(val) : 0);
|
||||
}
|
||||
|
||||
snprintf(dev->props.name, sizeof(dev->props.name), "%s%d", "MTL", device);
|
||||
const char * gpu_name = [[dev->mtl_device name] UTF8String];
|
||||
if (n_devices > 1) {
|
||||
@@ -1354,6 +1365,8 @@ void ggml_metal_device_free(ggml_metal_device_t dev) {
|
||||
assert(dev != NULL);
|
||||
|
||||
@autoreleasepool {
|
||||
ggml_metal_fusion_info_free(dev->finfo);
|
||||
|
||||
ggml_metal_rsets_free(dev->rsets);
|
||||
|
||||
ggml_metal_library_free(dev->library);
|
||||
@@ -1941,6 +1954,10 @@ static void ggml_metal_device_disable_tensor(ggml_metal_device_t dev) {
|
||||
dev->props.has_tensor = false;
|
||||
}
|
||||
|
||||
struct ggml_metal_fusion_info * ggml_metal_device_get_fusion_info(ggml_metal_device_t dev) {
|
||||
return dev->finfo;
|
||||
}
|
||||
|
||||
//
|
||||
// device buffers
|
||||
//
|
||||
|
||||
@@ -0,0 +1,502 @@
|
||||
#include "ggml-metal-fusion.h"
|
||||
|
||||
#include "ggml-backend-impl.h"
|
||||
#include "ggml-metal-device.h"
|
||||
|
||||
#include <algorithm>
|
||||
#include <string>
|
||||
#include <vector>
|
||||
|
||||
// ---- helpers -------------------------------------------------------------
|
||||
|
||||
// true if two tensors live in the same Metal buffer
|
||||
static bool ggml_metal_fusion_same_buffer(const ggml_tensor * a, const ggml_tensor * b) {
|
||||
if (!a || !b) {
|
||||
return false;
|
||||
}
|
||||
|
||||
ggml_backend_buffer_t ba = a->view_src ? a->view_src->buffer : a->buffer;
|
||||
ggml_backend_buffer_t bb = b->view_src ? b->view_src->buffer : b->buffer;
|
||||
|
||||
ggml_metal_buffer_t ca = (ggml_metal_buffer_t) ba->context;
|
||||
ggml_metal_buffer_t cb = (ggml_metal_buffer_t) bb->context;
|
||||
|
||||
return ggml_metal_buffer_get_id(ca, a).metal == ggml_metal_buffer_get_id(cb, b).metal;
|
||||
}
|
||||
|
||||
// ---- pattern checks ------------------------------------------------------
|
||||
|
||||
// NORM/RMS_NORM + MUL + ADD: the weight/bias of each fused step must match the norm input
|
||||
// width, be contiguous rows, and the fused outputs must stay F32
|
||||
static bool ggml_metal_fusion_check_norm(
|
||||
const ggml_metal_fusion * fusion,
|
||||
const ggml_tensor * const * nodes,
|
||||
ggml_metal_fusion_mode mode) {
|
||||
GGML_UNUSED(mode);
|
||||
|
||||
GGML_ASSERT(fusion->n_ops >= 2);
|
||||
|
||||
for (int j = 1; j < fusion->n_ops; j++) {
|
||||
// the fused MUL/ADD must read the previous node as src0
|
||||
if (nodes[j]->src[0] != nodes[j - 1]) {
|
||||
return false;
|
||||
}
|
||||
|
||||
// the weight/bias must have the same row width as the norm input
|
||||
if (nodes[j]->src[1]->ne[0] != nodes[0]->ne[0]) {
|
||||
return false;
|
||||
}
|
||||
|
||||
if (!ggml_is_contiguous_rows(nodes[j]->src[1])) {
|
||||
return false;
|
||||
}
|
||||
|
||||
if (nodes[j]->type != GGML_TYPE_F32) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
// ADD x N: each ADD reads the previous ADD as src0, and all addends must share layout
|
||||
// (and, in FULL mode, live in the same Metal buffer)
|
||||
static bool ggml_metal_fusion_check_add_chain(
|
||||
const ggml_metal_fusion * fusion,
|
||||
const ggml_tensor * const * nodes,
|
||||
ggml_metal_fusion_mode mode) {
|
||||
GGML_ASSERT(fusion->n_ops >= 2);
|
||||
|
||||
for (int j = 1; j < fusion->n_ops; j++) {
|
||||
if (nodes[j]->src[0] != nodes[j - 1]) {
|
||||
return false;
|
||||
}
|
||||
|
||||
if (!ggml_are_same_layout(nodes[j]->src[1], nodes[j - 1]->src[1])) {
|
||||
return false;
|
||||
}
|
||||
|
||||
if (mode == GGML_METAL_FUSION_FULL) {
|
||||
if (!ggml_metal_fusion_same_buffer(nodes[j]->src[1], nodes[0]->src[1])) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
// GATED_DELTA_NET + CPY: the trailing cpy scatters the gdn state snapshots into the recurrent
|
||||
// cache, so the gdn kernel writes them straight to the cache and the cpy is elided.
|
||||
// mirrors ggml_metal_op_can_fuse_gdn_cache (PR #25788). the gdn output has other consumers (the
|
||||
// attn scores view), so unlike the other patterns this is not an elision chain: the structural
|
||||
// checks live entirely in this callback (unsafe = true).
|
||||
static bool ggml_metal_fusion_check_gdn_cache(
|
||||
const ggml_metal_fusion * fusion,
|
||||
const ggml_tensor * const * nodes,
|
||||
ggml_metal_fusion_mode mode) {
|
||||
GGML_UNUSED(fusion);
|
||||
|
||||
const ggml_tensor * gdn = nodes[0];
|
||||
const ggml_tensor * cpy = nodes[1];
|
||||
|
||||
// the kernel skips the snapshot tail, so the gdn output must not be a graph output
|
||||
if (gdn->type != GGML_TYPE_F32 || (gdn->flags & GGML_TENSOR_FLAG_OUTPUT)) {
|
||||
return false;
|
||||
}
|
||||
|
||||
if (cpy->op != GGML_OP_CPY || (cpy->flags & GGML_TENSOR_FLAG_OUTPUT)) {
|
||||
return false;
|
||||
}
|
||||
|
||||
const int64_t S_v = gdn->src[2]->ne[0];
|
||||
const int64_t H = gdn->src[2]->ne[1];
|
||||
const int64_t n_tokens = gdn->src[2]->ne[2];
|
||||
const int64_t n_seqs = gdn->src[2]->ne[3];
|
||||
const int64_t K = ggml_get_op_params_i32(gdn, 0);
|
||||
const size_t tail_off = ggml_row_size(GGML_TYPE_F32, S_v * H * n_tokens * n_seqs);
|
||||
|
||||
const int64_t D = S_v * S_v * H;
|
||||
const int64_t n_written = std::min<int64_t>(n_tokens, K);
|
||||
|
||||
const ggml_tensor * src = cpy->src[0]; // gdn snapshot tail view
|
||||
const ggml_tensor * dst = cpy->src[1]; // cache view
|
||||
|
||||
// src must be this gdn's snapshot tail (contiguous, at the tail offset)
|
||||
if (src->op != GGML_OP_VIEW || src->view_src != gdn ||
|
||||
src->view_offs != tail_off || !ggml_is_contiguous(src)) {
|
||||
return false;
|
||||
}
|
||||
|
||||
const int64_t expected_ne[GGML_MAX_DIMS] = { D, n_seqs, n_written, 1 };
|
||||
if (dst->type != GGML_TYPE_F32 ||
|
||||
!std::equal(expected_ne, expected_ne + GGML_MAX_DIMS, dst->ne) ||
|
||||
dst->nb[0] != ggml_type_size(GGML_TYPE_F32) ||
|
||||
dst->nb[1] != ggml_row_size(GGML_TYPE_F32, D)) {
|
||||
return false;
|
||||
}
|
||||
|
||||
if (mode == GGML_METAL_FUSION_FULL) {
|
||||
// the cache must be allocated so the kernel can write straight to its buffer
|
||||
if (dst->data == nullptr) {
|
||||
return false;
|
||||
}
|
||||
}
|
||||
|
||||
return true;
|
||||
}
|
||||
|
||||
// MUL + SIN + SQR + MUL + ADD (snake activation)
|
||||
static bool ggml_metal_fusion_check_snake(
|
||||
const ggml_metal_fusion * fusion,
|
||||
const ggml_tensor * const * nodes,
|
||||
ggml_metal_fusion_mode mode) {
|
||||
GGML_UNUSED(fusion);
|
||||
GGML_UNUSED(mode);
|
||||
|
||||
const ggml_tensor * mul0 = nodes[0];
|
||||
const ggml_tensor * sin_node = nodes[1];
|
||||
const ggml_tensor * sqr = nodes[2];
|
||||
const ggml_tensor * mul1 = nodes[3];
|
||||
const ggml_tensor * add = nodes[4];
|
||||
|
||||
// x carries the full activation shape, a is the broadcast operand
|
||||
const ggml_tensor * x = ggml_are_same_shape(mul0, mul0->src[0]) ? mul0->src[0] : mul0->src[1];
|
||||
const ggml_tensor * a = (x == mul0->src[0]) ? mul0->src[1] : mul0->src[0];
|
||||
|
||||
// mul1 reads sqr and inv_b in either operand order
|
||||
const ggml_tensor * inv_b = (mul1->src[0] == sqr) ? mul1->src[1] : mul1->src[0];
|
||||
|
||||
// closure check: the trailing add reads the same x as the leading mul
|
||||
const ggml_tensor * x_in_add = (add->src[0] == mul1) ? add->src[1] : add->src[0];
|
||||
|
||||
// x is in the supported whitelist and every chain intermediate shares x's type.
|
||||
// a and inv_b bind as device const float * in the kernel, so they stay F32.
|
||||
const bool types_ok =
|
||||
(x->type == GGML_TYPE_F32 || x->type == GGML_TYPE_F16 || x->type == GGML_TYPE_BF16) &&
|
||||
(a->type == GGML_TYPE_F32) && (inv_b->type == GGML_TYPE_F32) &&
|
||||
(mul0->type == x->type) && (sin_node->type == x->type) &&
|
||||
(sqr->type == x->type) && (mul1->type == x->type) &&
|
||||
(add->type == x->type);
|
||||
|
||||
// a / inv_b collapse to [1, C, 1, 1], x and add stay 2D
|
||||
const bool shape_ok = ggml_are_same_shape(a, inv_b) && a->ne[0] == 1 && a->ne[1] == x->ne[1];
|
||||
const bool dim_ok =
|
||||
(x->ne[2] == 1) && (x->ne[3] == 1) &&
|
||||
(add->ne[2] == 1) && (add->ne[3] == 1) &&
|
||||
(a->ne[2] == 1) && (a->ne[3] == 1) &&
|
||||
(inv_b->ne[2] == 1) && (inv_b->ne[3] == 1);
|
||||
|
||||
// kernel reads x[idx] and a[c] / inv_b[c] linearly, so every operand is contiguous
|
||||
const bool contig_ok =
|
||||
ggml_is_contiguous(x) && ggml_is_contiguous(add) &&
|
||||
ggml_is_contiguous(a) && ggml_is_contiguous(inv_b);
|
||||
|
||||
return types_ok && shape_ok && dim_ok && contig_ok && x_in_add == x;
|
||||
}
|
||||
|
||||
// ---- patterns ------------------------------------------------------------
|
||||
|
||||
static const ggml_op ops_norm_mul[] = { GGML_OP_NORM, GGML_OP_MUL };
|
||||
static const ggml_op ops_norm_mul_add[] = { GGML_OP_NORM, GGML_OP_MUL, GGML_OP_ADD };
|
||||
static const ggml_op ops_rms_norm_mul[] = { GGML_OP_RMS_NORM, GGML_OP_MUL };
|
||||
static const ggml_op ops_rms_norm_mul_add[] = { GGML_OP_RMS_NORM, GGML_OP_MUL, GGML_OP_ADD };
|
||||
|
||||
static const ggml_op ops_add_2[] = { GGML_OP_ADD, GGML_OP_ADD };
|
||||
static const ggml_op ops_add_3[] = { GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD };
|
||||
static const ggml_op ops_add_4[] = { GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD };
|
||||
static const ggml_op ops_add_5[] = { GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD };
|
||||
static const ggml_op ops_add_6[] = { GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD };
|
||||
static const ggml_op ops_add_7[] = { GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD, GGML_OP_ADD };
|
||||
static const ggml_op ops_snake[] = { GGML_OP_MUL, GGML_OP_SIN, GGML_OP_SQR, GGML_OP_MUL, GGML_OP_ADD };
|
||||
|
||||
static const ggml_op ops_gdn_cache[] = { GGML_OP_GATED_DELTA_NET, GGML_OP_CPY };
|
||||
|
||||
static const ggml_metal_fusion ggml_metal_fusions[] = {
|
||||
{ GGML_METAL_FUSION_NORM_MUL, ops_norm_mul, 2, false, ggml_metal_fusion_check_norm },
|
||||
{ GGML_METAL_FUSION_NORM_MUL_ADD, ops_norm_mul_add, 3, false, ggml_metal_fusion_check_norm },
|
||||
{ GGML_METAL_FUSION_NORM_MUL, ops_rms_norm_mul, 2, false, ggml_metal_fusion_check_norm },
|
||||
{ GGML_METAL_FUSION_NORM_MUL_ADD, ops_rms_norm_mul_add, 3, false, ggml_metal_fusion_check_norm },
|
||||
{ GGML_METAL_FUSION_ADD_CHAIN, ops_add_2, 2, false, ggml_metal_fusion_check_add_chain },
|
||||
{ GGML_METAL_FUSION_ADD_CHAIN, ops_add_3, 3, false, ggml_metal_fusion_check_add_chain },
|
||||
{ GGML_METAL_FUSION_ADD_CHAIN, ops_add_4, 4, false, ggml_metal_fusion_check_add_chain },
|
||||
{ GGML_METAL_FUSION_ADD_CHAIN, ops_add_5, 5, false, ggml_metal_fusion_check_add_chain },
|
||||
{ GGML_METAL_FUSION_ADD_CHAIN, ops_add_6, 6, false, ggml_metal_fusion_check_add_chain },
|
||||
{ GGML_METAL_FUSION_ADD_CHAIN, ops_add_7, 7, false, ggml_metal_fusion_check_add_chain },
|
||||
{ GGML_METAL_FUSION_SNAKE, ops_snake, 5, false, ggml_metal_fusion_check_snake },
|
||||
{ GGML_METAL_FUSION_GDN_CACHE, ops_gdn_cache, 2, true, ggml_metal_fusion_check_gdn_cache },
|
||||
};
|
||||
|
||||
const ggml_metal_fusion * ggml_metal_fusion_all(int * n) {
|
||||
*n = (int) sizeof(ggml_metal_fusions) / sizeof(ggml_metal_fusions[0]);
|
||||
|
||||
return ggml_metal_fusions;
|
||||
}
|
||||
|
||||
// ---- shared fusion info ---------------------------------------------------
|
||||
|
||||
static std::string ggml_metal_fusion_label(const ggml_metal_fusion * fusion) {
|
||||
GGML_ASSERT(fusion != nullptr);
|
||||
|
||||
std::string label;
|
||||
for (int j = 0; j < fusion->n_ops; j++) {
|
||||
if (j > 0) {
|
||||
label += '+';
|
||||
}
|
||||
label += ggml_op_name(fusion->ops[j]);
|
||||
}
|
||||
return label;
|
||||
}
|
||||
|
||||
struct ggml_metal_fusion_info {
|
||||
std::vector<std::string> labels;
|
||||
std::vector<uint64_t> counts;
|
||||
bool enabled;
|
||||
bool stats;
|
||||
bool labels_set;
|
||||
int debug;
|
||||
};
|
||||
|
||||
struct ggml_metal_fusion_info * ggml_metal_fusion_info_init(bool enabled, int debug) {
|
||||
ggml_metal_fusion_info * finfo = new ggml_metal_fusion_info;
|
||||
finfo->enabled = enabled;
|
||||
finfo->stats = debug > 0;
|
||||
finfo->labels_set = false;
|
||||
finfo->debug = debug;
|
||||
|
||||
if (finfo->stats) {
|
||||
ggml_metal_fusion_info_labels_init(finfo);
|
||||
}
|
||||
|
||||
return finfo;
|
||||
}
|
||||
|
||||
void ggml_metal_fusion_info_free(struct ggml_metal_fusion_info * finfo) {
|
||||
delete finfo;
|
||||
}
|
||||
|
||||
bool ggml_metal_fusion_info_enabled(const struct ggml_metal_fusion_info * finfo) {
|
||||
return finfo->enabled;
|
||||
}
|
||||
|
||||
bool ggml_metal_fusion_info_stats(const struct ggml_metal_fusion_info * finfo) {
|
||||
return finfo->stats;
|
||||
}
|
||||
|
||||
int ggml_metal_fusion_info_debug(const struct ggml_metal_fusion_info * finfo) {
|
||||
return finfo->debug;
|
||||
}
|
||||
|
||||
int ggml_metal_fusion_info_n_fusions(const struct ggml_metal_fusion_info * finfo) {
|
||||
return (int) finfo->labels.size();
|
||||
}
|
||||
|
||||
const char * ggml_metal_fusion_info_label(const struct ggml_metal_fusion_info * finfo, int idx) {
|
||||
GGML_ASSERT(idx >= 0 && idx < (int) finfo->labels.size());
|
||||
return finfo->labels[idx].c_str();
|
||||
}
|
||||
|
||||
uint64_t ggml_metal_fusion_info_count(const struct ggml_metal_fusion_info * finfo, int idx) {
|
||||
GGML_ASSERT(idx >= 0 && idx < (int) finfo->counts.size());
|
||||
return finfo->counts[idx];
|
||||
}
|
||||
|
||||
void ggml_metal_fusion_info_count_fusion(struct ggml_metal_fusion_info * finfo, const struct ggml_metal_fusion * fusion) {
|
||||
if (!finfo->stats || fusion == nullptr) {
|
||||
return;
|
||||
}
|
||||
|
||||
int n = 0;
|
||||
const ggml_metal_fusion * all = ggml_metal_fusion_all(&n);
|
||||
|
||||
int idx = -1;
|
||||
for (int i = 0; i < n; i++) {
|
||||
if (&all[i] == fusion) {
|
||||
idx = i;
|
||||
break;
|
||||
}
|
||||
}
|
||||
|
||||
if (idx >= 0 && idx < (int) finfo->counts.size()) {
|
||||
finfo->counts[idx]++;
|
||||
}
|
||||
}
|
||||
|
||||
void ggml_metal_fusion_info_set_enabled(struct ggml_metal_fusion_info * finfo, bool enabled) {
|
||||
finfo->enabled = enabled;
|
||||
}
|
||||
|
||||
void ggml_metal_fusion_info_labels_init(struct ggml_metal_fusion_info * finfo) {
|
||||
if (finfo->labels_set) {
|
||||
return;
|
||||
}
|
||||
|
||||
int n = 0;
|
||||
const ggml_metal_fusion * all = ggml_metal_fusion_all(&n);
|
||||
|
||||
finfo->labels.clear();
|
||||
finfo->counts.assign(n, 0);
|
||||
finfo->labels.reserve(n);
|
||||
|
||||
for (int i = 0; i < n; i++) {
|
||||
finfo->labels.emplace_back(ggml_metal_fusion_label(&all[i]));
|
||||
}
|
||||
|
||||
finfo->labels_set = true;
|
||||
}
|
||||
|
||||
void ggml_metal_fusion_info_stats_init(struct ggml_metal_fusion_info * finfo) {
|
||||
finfo->stats = true;
|
||||
ggml_metal_fusion_info_labels_init(finfo);
|
||||
}
|
||||
|
||||
void ggml_metal_fusion_info_stats_reset(struct ggml_metal_fusion_info * finfo) {
|
||||
std::fill(finfo->counts.begin(), finfo->counts.end(), 0);
|
||||
}
|
||||
|
||||
int ggml_metal_fusion_info_stats_get(const struct ggml_metal_fusion_info * finfo, const char ** labels, uint64_t * counts, int n) {
|
||||
const int n_fusions = (int) finfo->labels.size();
|
||||
|
||||
if (labels == nullptr) {
|
||||
return n_fusions;
|
||||
}
|
||||
|
||||
const int n_fill = std::min(n, n_fusions);
|
||||
for (int i = 0; i < n_fill; i++) {
|
||||
labels[i] = finfo->labels[i].c_str();
|
||||
if (counts != nullptr) {
|
||||
counts[i] = finfo->counts[i];
|
||||
}
|
||||
}
|
||||
|
||||
return n_fill;
|
||||
}
|
||||
|
||||
// ---- queries -------------------------------------------------------------
|
||||
|
||||
// find the longest pattern matching the node sequence starting at idx
|
||||
// (idx is a position in node_idxs, which maps to graph node indices)
|
||||
const ggml_metal_fusion * ggml_metal_fusion_next(
|
||||
const ggml_cgraph * gf,
|
||||
const int * node_idxs,
|
||||
int n_idxs,
|
||||
int idx,
|
||||
ggml_metal_fusion_mode mode,
|
||||
int * n_out) {
|
||||
int n = 0;
|
||||
const ggml_metal_fusion * all = ggml_metal_fusion_all(&n);
|
||||
|
||||
const ggml_metal_fusion * res = nullptr;
|
||||
int best = 1;
|
||||
|
||||
for (int i = 0; i < n; i++) {
|
||||
const ggml_metal_fusion * fusion = &all[i];
|
||||
|
||||
// only look for a longer match than the current best
|
||||
if (fusion->n_ops <= best) {
|
||||
continue;
|
||||
}
|
||||
if (idx + fusion->n_ops > n_idxs) {
|
||||
continue;
|
||||
}
|
||||
|
||||
const ggml_tensor * nodes[GGML_METAL_FUSION_MAX];
|
||||
|
||||
// the op sequence must match exactly
|
||||
bool ok = true;
|
||||
for (int j = 0; j < fusion->n_ops; j++) {
|
||||
nodes[j] = gf->nodes[node_idxs[idx + j]];
|
||||
if (nodes[j]->op != fusion->ops[j]) {
|
||||
ok = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (!ok) {
|
||||
continue;
|
||||
}
|
||||
|
||||
if (!fusion->unsafe) {
|
||||
// common element-wise chain constraints: each node reads the previous one,
|
||||
// and all nodes have the same shape
|
||||
for (int j = 1; j < fusion->n_ops && ok; j++) {
|
||||
if (nodes[j]->src[0] != nodes[j - 1] && nodes[j]->src[1] != nodes[j - 1]) {
|
||||
ok = false;
|
||||
break;
|
||||
}
|
||||
if (!ggml_are_same_shape(nodes[j], nodes[j - 1])) {
|
||||
ok = false;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (!ok) {
|
||||
continue;
|
||||
}
|
||||
|
||||
// all current fusions are single-output elision chains, so the last node is the only output
|
||||
// TODO: multi-output fusions: store pattern-relative offsets in the table and translate them here
|
||||
int outputs_buf[1];
|
||||
outputs_buf[0] = node_idxs[idx + fusion->n_ops - 1];
|
||||
|
||||
// structural subgraph checks (op sequence, elidable uses, view containment)
|
||||
if (!ggml_can_fuse_subgraph_ext(gf, node_idxs + idx, fusion->n_ops, fusion->ops, outputs_buf, 1)) {
|
||||
continue;
|
||||
}
|
||||
}
|
||||
|
||||
// pattern-specific checks (the sole validator for unsafe patterns)
|
||||
if (fusion->check && !fusion->check(fusion, nodes, mode)) {
|
||||
continue;
|
||||
}
|
||||
|
||||
best = fusion->n_ops;
|
||||
res = fusion;
|
||||
}
|
||||
|
||||
*n_out = best;
|
||||
|
||||
return res;
|
||||
}
|
||||
|
||||
// optimize phase: maximum number of nodes starting at idx (a raw sequential graph index) that
|
||||
// could be fused, chaining patterns back-to-back. matching runs on the same filtered (view
|
||||
// transparent) node sequence that the compute phase uses, so the returned count is the raw index
|
||||
// span from idx to the last matched node (intermediate views are packed along).
|
||||
int ggml_metal_fusion_max(const ggml_cgraph * gf, int idx) {
|
||||
// an empty/view node cannot start a pattern - pack it alone
|
||||
if (ggml_op_is_empty(gf->nodes[idx]->op) || ggml_is_empty(gf->nodes[idx])) {
|
||||
return 1;
|
||||
}
|
||||
|
||||
// collect the non-empty node indices starting at idx
|
||||
int idxs[GGML_METAL_FUSION_MAX];
|
||||
int n_idxs = 0;
|
||||
for (int i = idx; i < gf->n_nodes && n_idxs < GGML_METAL_FUSION_MAX; i++) {
|
||||
if (!ggml_op_is_empty(gf->nodes[i]->op) && !ggml_is_empty(gf->nodes[i])) {
|
||||
idxs[n_idxs++] = i;
|
||||
}
|
||||
}
|
||||
if (n_idxs == 0) {
|
||||
return 1;
|
||||
}
|
||||
|
||||
int total = 0;
|
||||
int i_f = 0;
|
||||
|
||||
while (i_f < n_idxs && total < GGML_METAL_FUSION_MAX) {
|
||||
int len = 1;
|
||||
const ggml_metal_fusion * fusion = ggml_metal_fusion_next(gf, idxs, n_idxs, i_f, GGML_METAL_FUSION_STRUCTURAL, &len);
|
||||
if (!fusion || total + len > GGML_METAL_FUSION_MAX) {
|
||||
break;
|
||||
}
|
||||
|
||||
total += len;
|
||||
i_f += len;
|
||||
}
|
||||
|
||||
if (i_f == 0) {
|
||||
return 1;
|
||||
}
|
||||
|
||||
// map the matched non-empty nodes back to the raw index span (views are included)
|
||||
return std::min(GGML_METAL_FUSION_MAX, idxs[i_f - 1] - idx + 1);
|
||||
}
|
||||
@@ -0,0 +1,104 @@
|
||||
// single source of truth for the fusions supported by the Metal backend
|
||||
//
|
||||
// every fusable subgraph is declared exactly once as a ggml_metal_fusion entry in
|
||||
// the table in ggml-metal-fusion.cpp. both the graph optimizer (ggml_metal_fusion_max)
|
||||
// and the op encoders (ggml_metal_fusion_next) consult this same table, so the two
|
||||
// phases can never disagree about what can be fused.
|
||||
|
||||
#pragma once
|
||||
|
||||
#include "ggml-impl.h"
|
||||
|
||||
#include <stdbool.h>
|
||||
|
||||
#ifdef __cplusplus
|
||||
extern "C" {
|
||||
#endif
|
||||
|
||||
// the maximum number of nodes that can be fused in a single kernel
|
||||
// (also the maximum length of a packed fusion group during graph optimization)
|
||||
#define GGML_METAL_FUSION_MAX 16
|
||||
|
||||
typedef enum ggml_metal_fusion_mode {
|
||||
// structural checks only; used by the graph optimizer, at which point the graph
|
||||
// tensors are not allocated yet, so buffer placement cannot be verified
|
||||
GGML_METAL_FUSION_STRUCTURAL = 0,
|
||||
// full checks, including buffer placement; used by the op encoders
|
||||
GGML_METAL_FUSION_FULL,
|
||||
} ggml_metal_fusion_mode;
|
||||
|
||||
// identifier of each fusion pattern so the op encoders know which kernel to use
|
||||
typedef enum ggml_metal_fusion_id {
|
||||
GGML_METAL_FUSION_NONE = 0,
|
||||
GGML_METAL_FUSION_NORM_MUL, // NORM/RMS_NORM + MUL
|
||||
GGML_METAL_FUSION_NORM_MUL_ADD, // NORM/RMS_NORM + MUL + ADD
|
||||
GGML_METAL_FUSION_ADD_CHAIN, // ADD x N (N in [2, 7])
|
||||
GGML_METAL_FUSION_SNAKE, // MUL + SIN + SQR + MUL + ADD
|
||||
GGML_METAL_FUSION_GDN_CACHE, // GATED_DELTA_NET + CPY (write snapshots into the recurrent cache)
|
||||
} ggml_metal_fusion_id;
|
||||
|
||||
struct ggml_metal_fusion {
|
||||
ggml_metal_fusion_id id;
|
||||
|
||||
const enum ggml_op * ops; // op sequence (fixed length)
|
||||
int n_ops; // number of ops
|
||||
|
||||
// if unsafe: the generic chain/shape + ggml_can_fuse_subgraph checks are skipped and the
|
||||
// check callback below is the sole validator (used for patterns that are not elision chains,
|
||||
// e.g. the gdn + cache-cpy write-through fusion)
|
||||
bool unsafe;
|
||||
|
||||
// extra backend constraints on top of ggml_can_fuse_subgraph
|
||||
// nodes[j] is the j-th node of the pattern
|
||||
bool (*check)(const struct ggml_metal_fusion * fusion,
|
||||
const struct ggml_tensor * const * nodes,
|
||||
ggml_metal_fusion_mode mode);
|
||||
};
|
||||
|
||||
typedef struct ggml_metal_fusion ggml_metal_fusion;
|
||||
|
||||
// the single table of all fusions supported by the Metal backend
|
||||
const ggml_metal_fusion * ggml_metal_fusion_all(int * n);
|
||||
|
||||
// ---- shared fusion info ---------------------------------------------------
|
||||
|
||||
// shared fusion debugging context, owned by the device; newly created backend contexts for that
|
||||
// device register with it so the fusion counters are race-free and accumulate across contexts.
|
||||
struct ggml_metal_fusion_info; // defined in ggml-metal-fusion.cpp
|
||||
|
||||
struct ggml_metal_fusion_info * ggml_metal_fusion_info_init(bool enabled, int debug);
|
||||
void ggml_metal_fusion_info_free(struct ggml_metal_fusion_info * finfo);
|
||||
|
||||
bool ggml_metal_fusion_info_enabled(const struct ggml_metal_fusion_info * finfo);
|
||||
bool ggml_metal_fusion_info_stats (const struct ggml_metal_fusion_info * finfo);
|
||||
int ggml_metal_fusion_info_debug (const struct ggml_metal_fusion_info * finfo);
|
||||
|
||||
int ggml_metal_fusion_info_n_fusions(const struct ggml_metal_fusion_info * finfo);
|
||||
const char * ggml_metal_fusion_info_label (const struct ggml_metal_fusion_info * finfo, int idx);
|
||||
uint64_t ggml_metal_fusion_info_count (const struct ggml_metal_fusion_info * finfo, int idx);
|
||||
|
||||
void ggml_metal_fusion_info_count_fusion(struct ggml_metal_fusion_info * finfo, const struct ggml_metal_fusion * fusion);
|
||||
void ggml_metal_fusion_info_set_enabled (struct ggml_metal_fusion_info * finfo, bool enabled);
|
||||
|
||||
void ggml_metal_fusion_info_stats_init ( struct ggml_metal_fusion_info * finfo);
|
||||
void ggml_metal_fusion_info_stats_reset( struct ggml_metal_fusion_info * finfo);
|
||||
int ggml_metal_fusion_info_stats_get (const struct ggml_metal_fusion_info * finfo, const char ** labels, uint64_t * counts, int n);
|
||||
void ggml_metal_fusion_info_labels_init( struct ggml_metal_fusion_info * finfo);
|
||||
|
||||
// compute phase: longest fusion starting at idx (a position in node_idxs) that matches in `mode`.
|
||||
// returns the matching pattern (nullptr if no fusion) and sets *n_out to the number of nodes consumed.
|
||||
const ggml_metal_fusion * ggml_metal_fusion_next(
|
||||
const struct ggml_cgraph * gf,
|
||||
const int * node_idxs,
|
||||
int n_idxs,
|
||||
int idx,
|
||||
ggml_metal_fusion_mode mode,
|
||||
int * n_out);
|
||||
|
||||
// optimize phase: maximum number of nodes starting at idx (a raw sequential graph index) that
|
||||
// could be fused, chaining patterns back-to-back. returns at least 1.
|
||||
int ggml_metal_fusion_max(const struct ggml_cgraph * gf, int idx);
|
||||
|
||||
#ifdef __cplusplus
|
||||
}
|
||||
#endif
|
||||
@@ -62,18 +62,23 @@
|
||||
|
||||
#define N_R0_IQ1_S 4
|
||||
#define N_SG_IQ1_S 2
|
||||
#define N_R0_IQ1_S_SPLIT 8
|
||||
|
||||
#define N_R0_IQ1_M 4
|
||||
#define N_SG_IQ1_M 2
|
||||
#define N_R0_IQ1_M_SPLIT 8
|
||||
|
||||
#define N_R0_IQ2_XXS 4
|
||||
#define N_SG_IQ2_XXS 2
|
||||
#define N_R0_IQ2_XXS_SPLIT 8
|
||||
|
||||
#define N_R0_IQ2_XS 4
|
||||
#define N_SG_IQ2_XS 2
|
||||
#define N_R0_IQ2_XS_SPLIT 8
|
||||
|
||||
#define N_R0_IQ2_S 4
|
||||
#define N_SG_IQ2_S 2
|
||||
#define N_R0_IQ2_S_SPLIT 8
|
||||
|
||||
#define N_R0_IQ3_XXS 4
|
||||
#define N_SG_IQ3_XXS 2
|
||||
@@ -81,6 +86,7 @@
|
||||
|
||||
#define N_R0_IQ3_S 4
|
||||
#define N_SG_IQ3_S 2
|
||||
#define N_R0_IQ3_S_SPLIT 8
|
||||
|
||||
#define N_R0_IQ4_NL 2
|
||||
#define N_SG_IQ4_NL 2
|
||||
@@ -979,6 +985,7 @@ typedef struct {
|
||||
uint64_t nb1;
|
||||
uint64_t nb2;
|
||||
uint64_t nb3;
|
||||
uint64_t nb_out; // 0 => snapshots are appended after the attn scores (unfused)
|
||||
} ggml_metal_kargs_gated_delta_net;
|
||||
|
||||
typedef struct {
|
||||
|
||||
@@ -7,6 +7,7 @@
|
||||
#include "ggml-metal-impl.h"
|
||||
#include "ggml-metal-common.h"
|
||||
#include "ggml-metal-device.h"
|
||||
#include "ggml-metal-fusion.h"
|
||||
#include "ggml-metal-tuning.h"
|
||||
|
||||
#include <cassert>
|
||||
@@ -31,24 +32,22 @@ struct ggml_metal_op {
|
||||
ggml_metal_device_t dev,
|
||||
ggml_metal_cmd_buf_t cmd_buf,
|
||||
ggml_cgraph * gf,
|
||||
ggml_metal_fusion_info * finfo,
|
||||
int idx_start,
|
||||
int idx_end,
|
||||
bool use_fusion,
|
||||
bool use_concurrency,
|
||||
bool use_capture,
|
||||
int debug_graph,
|
||||
int debug_fusion) {
|
||||
int debug_graph) {
|
||||
this->dev = dev;
|
||||
this->lib = ggml_metal_device_get_library(dev);
|
||||
this->enc = ggml_metal_encoder_init(cmd_buf, use_concurrency);
|
||||
this->mem_ranges = ggml_mem_ranges_init(debug_graph);
|
||||
this->finfo = finfo;
|
||||
this->idx_start = idx_start;
|
||||
this->idx_end = idx_end;
|
||||
this->use_fusion = use_fusion;
|
||||
this->use_concurrency = use_concurrency;
|
||||
this->use_capture = use_capture;
|
||||
this->debug_graph = debug_graph;
|
||||
this->debug_fusion = debug_fusion;
|
||||
this->gf = gf;
|
||||
|
||||
idxs.reserve(gf->n_nodes);
|
||||
@@ -78,15 +77,24 @@ struct ggml_metal_op {
|
||||
return ggml_graph_node(gf, idxs[i]);
|
||||
}
|
||||
|
||||
bool can_fuse(int i0, const ggml_op * ops, int n_ops) const {
|
||||
assert(use_fusion);
|
||||
// consult the fusion table for the longest pattern starting at i0
|
||||
// returns the matching pattern (nullptr if no fusion) and sets *n_out to the number of nodes
|
||||
const ggml_metal_fusion * can_fuse(int i0, enum ggml_metal_fusion_mode mode, int * n_out) const {
|
||||
assert(use_fusion());
|
||||
assert(i0 >= 0 && i0 < n_nodes());
|
||||
|
||||
if (i0 + n_ops > n_nodes()) {
|
||||
return false;
|
||||
}
|
||||
return ggml_metal_fusion_next(gf, idxs.data(), (int) idxs.size(), i0, mode, n_out);
|
||||
}
|
||||
|
||||
return ggml_can_fuse_ext(gf, idxs.data() + i0, ops, n_ops);
|
||||
// whether to attempt fusion; the toggle lives in the shared fusion debugging context owned
|
||||
// by the device (initialized from GGML_METAL_FUSION_DISABLE, overridable by the test)
|
||||
bool use_fusion() const {
|
||||
return ggml_metal_fusion_info_enabled(finfo);
|
||||
}
|
||||
|
||||
// record that a fusion fired, indexed by the matching table entry
|
||||
void count_fusions(const ggml_metal_fusion * fusion) const {
|
||||
ggml_metal_fusion_info_count_fusion(finfo, fusion);
|
||||
}
|
||||
|
||||
ggml_metal_device_t dev;
|
||||
@@ -94,12 +102,13 @@ struct ggml_metal_op {
|
||||
ggml_metal_encoder_t enc;
|
||||
ggml_mem_ranges_t mem_ranges;
|
||||
|
||||
bool use_fusion;
|
||||
// shared fusion debugging context
|
||||
ggml_metal_fusion_info * finfo;
|
||||
|
||||
bool use_concurrency;
|
||||
bool use_capture;
|
||||
|
||||
int debug_graph;
|
||||
int debug_fusion;
|
||||
|
||||
private:
|
||||
ggml_cgraph * gf;
|
||||
@@ -115,24 +124,22 @@ ggml_metal_op_t ggml_metal_op_init(
|
||||
ggml_metal_device_t dev,
|
||||
ggml_metal_cmd_buf_t cmd_buf,
|
||||
ggml_cgraph * gf,
|
||||
ggml_metal_fusion_info * finfo,
|
||||
int idx_start,
|
||||
int idx_end,
|
||||
bool use_fusion,
|
||||
bool use_concurrency,
|
||||
bool use_capture,
|
||||
int debug_graph,
|
||||
int debug_fusion) {
|
||||
int debug_graph) {
|
||||
ggml_metal_op_t res = new ggml_metal_op(
|
||||
dev,
|
||||
cmd_buf,
|
||||
gf,
|
||||
finfo,
|
||||
idx_start,
|
||||
idx_end,
|
||||
use_fusion,
|
||||
use_concurrency,
|
||||
use_capture,
|
||||
debug_graph,
|
||||
debug_fusion);
|
||||
debug_graph);
|
||||
|
||||
return res;
|
||||
}
|
||||
@@ -1868,6 +1875,8 @@ int ggml_metal_op_gated_delta_net(ggml_metal_op_t ctx, int idx) {
|
||||
ggml_metal_library_t lib = ctx->lib;
|
||||
ggml_metal_encoder_t enc = ctx->enc;
|
||||
|
||||
const bool use_fusion = ctx->use_fusion();
|
||||
const int debug_fusion = ggml_metal_fusion_info_debug(ctx->finfo);
|
||||
|
||||
GGML_TENSOR_LOCALS( int32_t, ne0, op->src[0], ne);
|
||||
GGML_TENSOR_LOCALS(uint64_t, nb0, op->src[0], nb);
|
||||
@@ -1880,6 +1889,31 @@ int ggml_metal_op_gated_delta_net(ggml_metal_op_t ctx, int idx) {
|
||||
|
||||
auto pipeline = ggml_metal_library_get_pipeline_gated_delta_net(lib, op);
|
||||
|
||||
// when fused with the trailing cache cpy, the snapshots are written straight into the
|
||||
// recurrent cache and the cpy is skipped (see GGML_METAL_FUSION_GDN_CACHE)
|
||||
ggml_metal_buffer_id bid_out = ggml_metal_get_buffer_id(op);
|
||||
uint64_t nb_out = 0;
|
||||
int n_fuse = 1;
|
||||
|
||||
if (use_fusion) {
|
||||
int n = 1;
|
||||
const ggml_metal_fusion * fusion = ctx->can_fuse(idx, GGML_METAL_FUSION_FULL, &n);
|
||||
|
||||
if (fusion && fusion->id == GGML_METAL_FUSION_GDN_CACHE) {
|
||||
const ggml_tensor * dst_cache = ctx->node(idx + 1)->src[1]; // cache view
|
||||
|
||||
bid_out = ggml_metal_get_buffer_id(dst_cache);
|
||||
nb_out = dst_cache->nb[2]/sizeof(float);
|
||||
n_fuse = 2;
|
||||
|
||||
ctx->count_fusions(fusion);
|
||||
|
||||
if (debug_fusion > 1) {
|
||||
GGML_LOG_DEBUG("%s: fuse: GATED_DELTA_NET + CPY\n", __func__);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
int ida = 0;
|
||||
|
||||
ggml_metal_kargs_gated_delta_net args = {
|
||||
@@ -1918,23 +1952,25 @@ int ggml_metal_op_gated_delta_net(ggml_metal_op_t ctx, int idx) {
|
||||
/*.nb1 =*/ nb1,
|
||||
/*.nb2 =*/ nb2,
|
||||
/*.nb3 =*/ nb3,
|
||||
/*.nb_out =*/ nb_out,
|
||||
};
|
||||
|
||||
ggml_metal_encoder_set_pipeline(enc, pipeline);
|
||||
ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), ida++);
|
||||
ggml_metal_encoder_set_bytes (enc, &args, sizeof(args), ida++); // args
|
||||
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[0]), ida++); // q
|
||||
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[1]), ida++); // k
|
||||
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[2]), ida++); // v
|
||||
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[3]), ida++); // gate
|
||||
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[4]), ida++); // beta
|
||||
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op->src[5]), ida++); // state
|
||||
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op), ida++); // dst
|
||||
ggml_metal_encoder_set_buffer (enc, ggml_metal_get_buffer_id(op), ida++); // dst (attn)
|
||||
ggml_metal_encoder_set_buffer (enc, bid_out, ida++); // state_out
|
||||
|
||||
const int nsg = pipeline.nsg;
|
||||
|
||||
ggml_metal_encoder_dispatch_threadgroups(enc, op->src[2]->ne[0]/nsg, op->src[2]->ne[1], op->src[2]->ne[3], 32, nsg, 1);
|
||||
|
||||
return 1;
|
||||
return n_fuse;
|
||||
}
|
||||
|
||||
int ggml_metal_op_solve_tri(ggml_metal_op_t ctx, int idx) {
|
||||
@@ -3718,56 +3754,20 @@ int ggml_metal_op_flash_attn_ext(ggml_metal_op_t ctx, int idx) {
|
||||
return 1;
|
||||
}
|
||||
|
||||
// Snake activation autofuse: mul -> sin -> sqr -> mul -> add
|
||||
static bool ggml_metal_op_can_fuse_snake(ggml_metal_op_t ctx, int idx) {
|
||||
static constexpr ggml_op snake_ops[5] = { GGML_OP_MUL, GGML_OP_SIN, GGML_OP_SQR, GGML_OP_MUL, GGML_OP_ADD };
|
||||
|
||||
if (ctx->node(idx)->op != GGML_OP_MUL || !ctx->can_fuse(idx, snake_ops, 5)) {
|
||||
return false;
|
||||
}
|
||||
|
||||
const ggml_tensor * mul0 = ctx->node(idx + 0);
|
||||
const ggml_tensor * sin_node = ctx->node(idx + 1);
|
||||
const ggml_tensor * sqr = ctx->node(idx + 2);
|
||||
const ggml_tensor * mul1 = ctx->node(idx + 3);
|
||||
const ggml_tensor * add = ctx->node(idx + 4);
|
||||
|
||||
// x carries the full activation shape, a is the broadcast operand
|
||||
const ggml_tensor * x = ggml_are_same_shape(mul0, mul0->src[0]) ? mul0->src[0] : mul0->src[1];
|
||||
const ggml_tensor * a = (x == mul0->src[0]) ? mul0->src[1] : mul0->src[0];
|
||||
|
||||
// mul1 reads sqr and inv_b in either operand order
|
||||
const ggml_tensor * inv_b = (mul1->src[0] == sqr) ? mul1->src[1] : mul1->src[0];
|
||||
|
||||
// closure check: the trailing add reads the same x as the leading mul
|
||||
const ggml_tensor * x_in_add = (add->src[0] == mul1) ? add->src[1] : add->src[0];
|
||||
|
||||
// x is in the supported whitelist and every chain intermediate shares x's type.
|
||||
// a and inv_b bind as device const float * in the kernel, so they stay F32.
|
||||
const bool types_ok =
|
||||
(x->type == GGML_TYPE_F32 || x->type == GGML_TYPE_F16 || x->type == GGML_TYPE_BF16) &&
|
||||
(a->type == GGML_TYPE_F32) && (inv_b->type == GGML_TYPE_F32) &&
|
||||
(mul0->type == x->type) && (sin_node->type == x->type) &&
|
||||
(sqr->type == x->type) && (mul1->type == x->type) &&
|
||||
(add->type == x->type);
|
||||
// a / inv_b collapse to [1, C, 1, 1], x and add stay 2D
|
||||
const bool shape_ok = ggml_are_same_shape(a, inv_b) && a->ne[0] == 1 && a->ne[1] == x->ne[1];
|
||||
const bool dim_ok =
|
||||
(x->ne[2] == 1) && (x->ne[3] == 1) &&
|
||||
(add->ne[2] == 1) && (add->ne[3] == 1) &&
|
||||
(a->ne[2] == 1) && (a->ne[3] == 1) &&
|
||||
(inv_b->ne[2] == 1) && (inv_b->ne[3] == 1);
|
||||
// kernel reads x[idx] and a[c] / inv_b[c] linearly, so every operand is contiguous
|
||||
const bool contig_ok =
|
||||
ggml_is_contiguous(x) && ggml_is_contiguous(add) &&
|
||||
ggml_is_contiguous(a) && ggml_is_contiguous(inv_b);
|
||||
|
||||
return types_ok && shape_ok && dim_ok && contig_ok && x_in_add == x;
|
||||
}
|
||||
|
||||
int ggml_metal_op_bin(ggml_metal_op_t ctx, int idx) {
|
||||
if (ctx->use_fusion && ggml_metal_op_can_fuse_snake(ctx, idx)) {
|
||||
return ggml_metal_op_snake_fused(ctx, idx);
|
||||
int n_fuse = 1;
|
||||
const ggml_metal_fusion * fusion = nullptr;
|
||||
|
||||
if (ctx->use_fusion()) {
|
||||
int n = 1;
|
||||
fusion = ctx->can_fuse(idx, GGML_METAL_FUSION_FULL, &n);
|
||||
n_fuse = n;
|
||||
|
||||
// snake activation autofuse: mul -> sin -> sqr -> mul -> add
|
||||
if (fusion && fusion->id == GGML_METAL_FUSION_SNAKE) {
|
||||
ctx->count_fusions(fusion);
|
||||
return ggml_metal_op_snake_fused(ctx, idx);
|
||||
}
|
||||
}
|
||||
|
||||
ggml_tensor * op = ctx->node(idx);
|
||||
@@ -3775,9 +3775,9 @@ int ggml_metal_op_bin(ggml_metal_op_t ctx, int idx) {
|
||||
ggml_metal_library_t lib = ctx->lib;
|
||||
ggml_metal_encoder_t enc = ctx->enc;
|
||||
|
||||
const bool use_fusion = ctx->use_fusion;
|
||||
const bool use_fusion = ctx->use_fusion();
|
||||
|
||||
const int debug_fusion = ctx->debug_fusion;
|
||||
const int debug_fusion = ggml_metal_fusion_info_debug(ctx->finfo);
|
||||
|
||||
GGML_TENSOR_LOCALS( int32_t, ne0, op->src[0], ne);
|
||||
GGML_TENSOR_LOCALS(uint64_t, nb0, op->src[0], nb);
|
||||
@@ -3822,57 +3822,19 @@ int ggml_metal_op_bin(ggml_metal_op_t ctx, int idx) {
|
||||
/*.o1 =*/ { bid_src1.offs },
|
||||
};
|
||||
|
||||
ggml_op fops[8];
|
||||
|
||||
int n_fuse = 1;
|
||||
|
||||
// c[0] = add(a, b[0])
|
||||
// c[1] = add(c[0], b[1])
|
||||
// c[2] = add(c[1], b[2])
|
||||
// ...
|
||||
if (use_fusion) {
|
||||
fops[0] = GGML_OP_ADD;
|
||||
fops[1] = GGML_OP_ADD;
|
||||
fops[2] = GGML_OP_ADD;
|
||||
fops[3] = GGML_OP_ADD;
|
||||
fops[4] = GGML_OP_ADD;
|
||||
fops[5] = GGML_OP_ADD;
|
||||
fops[6] = GGML_OP_ADD;
|
||||
fops[7] = GGML_OP_ADD;
|
||||
|
||||
// note: in metal, we sometimes encode the graph in parallel so we have to avoid fusing ops
|
||||
// across splits. idx_end indicates the last node in the current split
|
||||
for (n_fuse = 0; n_fuse <= 6; ++n_fuse) {
|
||||
if (!ctx->can_fuse(idx + n_fuse, fops + n_fuse, 2)) {
|
||||
break;
|
||||
}
|
||||
|
||||
ggml_tensor * f0 = ctx->node(idx + n_fuse);
|
||||
ggml_tensor * f1 = ctx->node(idx + n_fuse + 1);
|
||||
|
||||
if (f0 != f1->src[0]) {
|
||||
break;
|
||||
}
|
||||
|
||||
// b[0] === b[1] === ...
|
||||
if (!ggml_are_same_layout(f0->src[1], f1->src[1])) {
|
||||
break;
|
||||
}
|
||||
|
||||
// only fuse ops if src1 is in the same Metal buffer
|
||||
ggml_metal_buffer_id bid_fuse = ggml_metal_get_buffer_id(f1->src[1]);
|
||||
if (bid_fuse.metal != bid_src1.metal) {
|
||||
break;
|
||||
}
|
||||
|
||||
//ctx->fuse_cnt[ops[n_fuse + 1]->op]++;
|
||||
|
||||
args.o1[n_fuse + 1] = bid_fuse.offs;
|
||||
if (use_fusion && fusion && fusion->id == GGML_METAL_FUSION_ADD_CHAIN) {
|
||||
// the offsets of the fused addends are relative to the start of the src1 buffer
|
||||
for (int i = 1; i < n_fuse; i++) {
|
||||
args.o1[i] = ggml_metal_get_buffer_id(ctx->node(idx + i)->src[1]).offs;
|
||||
}
|
||||
|
||||
++n_fuse;
|
||||
ctx->count_fusions(fusion);
|
||||
|
||||
if (debug_fusion > 1 && n_fuse > 1) {
|
||||
if (debug_fusion > 1) {
|
||||
GGML_LOG_DEBUG("%s: fuse: ADD x %d\n", __func__, n_fuse);
|
||||
}
|
||||
}
|
||||
@@ -4080,9 +4042,9 @@ int ggml_metal_op_norm(ggml_metal_op_t ctx, int idx) {
|
||||
ggml_metal_library_t lib = ctx->lib;
|
||||
ggml_metal_encoder_t enc = ctx->enc;
|
||||
|
||||
const bool use_fusion = ctx->use_fusion;
|
||||
const bool use_fusion = ctx->use_fusion();
|
||||
|
||||
const int debug_fusion = ctx->debug_fusion;
|
||||
const int debug_fusion = ggml_metal_fusion_info_debug(ctx->finfo);
|
||||
|
||||
GGML_TENSOR_LOCALS( int32_t, ne0, op->src[0], ne);
|
||||
GGML_TENSOR_LOCALS(uint64_t, nb0, op->src[0], nb);
|
||||
@@ -4110,8 +4072,6 @@ int ggml_metal_op_norm(ggml_metal_op_t ctx, int idx) {
|
||||
/*.nbf3 =*/ { nb03 },
|
||||
};
|
||||
|
||||
ggml_op fops[8];
|
||||
|
||||
int n_fuse = 1;
|
||||
|
||||
ggml_metal_buffer_id bid_fuse[2] = { bid_src0, bid_src0 };
|
||||
@@ -4120,55 +4080,35 @@ int ggml_metal_op_norm(ggml_metal_op_t ctx, int idx) {
|
||||
// d[1] = mul(d[0], b)
|
||||
// d[2] = add(d[1], c)
|
||||
if (use_fusion) {
|
||||
fops[0] = op->op;
|
||||
fops[1] = GGML_OP_MUL;
|
||||
fops[2] = GGML_OP_ADD;
|
||||
int n = 1;
|
||||
const ggml_metal_fusion * fusion = ctx->can_fuse(idx, GGML_METAL_FUSION_FULL, &n);
|
||||
|
||||
for (n_fuse = 0; n_fuse <= 1; ++n_fuse) {
|
||||
if (!ctx->can_fuse(idx + n_fuse, fops + n_fuse, 2)) {
|
||||
break;
|
||||
if (fusion && (fusion->id == GGML_METAL_FUSION_NORM_MUL || fusion->id == GGML_METAL_FUSION_NORM_MUL_ADD)) {
|
||||
n_fuse = n;
|
||||
|
||||
ctx->count_fusions(fusion);
|
||||
|
||||
for (int i = 1; i < n_fuse; i++) {
|
||||
const ggml_tensor * fn = ctx->node(idx + i);
|
||||
|
||||
bid_fuse[i - 1] = ggml_metal_get_buffer_id(fn->src[1]);
|
||||
|
||||
args.nef1[i] = fn->src[1]->ne[1];
|
||||
args.nef2[i] = fn->src[1]->ne[2];
|
||||
args.nef3[i] = fn->src[1]->ne[3];
|
||||
|
||||
args.nbf1[i] = fn->src[1]->nb[1];
|
||||
args.nbf2[i] = fn->src[1]->nb[2];
|
||||
args.nbf3[i] = fn->src[1]->nb[3];
|
||||
}
|
||||
|
||||
ggml_tensor * f0 = ctx->node(idx + n_fuse);
|
||||
ggml_tensor * f1 = ctx->node(idx + n_fuse + 1);
|
||||
|
||||
if (f0 != f1->src[0]) {
|
||||
break;
|
||||
}
|
||||
|
||||
if (f1->src[1]->ne[0] != op->ne[0]) {
|
||||
break;
|
||||
}
|
||||
|
||||
if (!ggml_is_contiguous_rows(f1->src[1])) {
|
||||
break;
|
||||
}
|
||||
|
||||
if (f1->type != GGML_TYPE_F32) {
|
||||
break;
|
||||
}
|
||||
|
||||
//ctx->fuse_cnt[f1->op]++;
|
||||
|
||||
bid_fuse[n_fuse] = ggml_metal_get_buffer_id(f1->src[1]);
|
||||
|
||||
args.nef1[n_fuse + 1] = f1->src[1]->ne[1];
|
||||
args.nef2[n_fuse + 1] = f1->src[1]->ne[2];
|
||||
args.nef3[n_fuse + 1] = f1->src[1]->ne[3];
|
||||
|
||||
args.nbf1[n_fuse + 1] = f1->src[1]->nb[1];
|
||||
args.nbf2[n_fuse + 1] = f1->src[1]->nb[2];
|
||||
args.nbf3[n_fuse + 1] = f1->src[1]->nb[3];
|
||||
}
|
||||
|
||||
++n_fuse;
|
||||
|
||||
if (debug_fusion > 1 && n_fuse > 1) {
|
||||
if (n_fuse == 2) {
|
||||
GGML_LOG_DEBUG("%s: fuse: %s + MUL\n", __func__, ggml_op_name(op->op));
|
||||
}
|
||||
if (n_fuse == 3) {
|
||||
GGML_LOG_DEBUG("%s: fuse: %s + MUL + ADD\n", __func__, ggml_op_name(op->op));
|
||||
if (debug_fusion > 1) {
|
||||
if (n_fuse == 2) {
|
||||
GGML_LOG_DEBUG("%s: fuse: %s + MUL\n", __func__, ggml_op_name(op->op));
|
||||
}
|
||||
if (n_fuse == 3) {
|
||||
GGML_LOG_DEBUG("%s: fuse: %s + MUL + ADD\n", __func__, ggml_op_name(op->op));
|
||||
}
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
@@ -8,17 +8,18 @@ extern "C" {
|
||||
|
||||
typedef struct ggml_metal_op * ggml_metal_op_t;
|
||||
|
||||
struct ggml_metal_fusion; // forward decl (ggml-metal-device.h)
|
||||
|
||||
ggml_metal_op_t ggml_metal_op_init(
|
||||
ggml_metal_device_t dev,
|
||||
ggml_metal_cmd_buf_t cmd_buf,
|
||||
struct ggml_cgraph * gf,
|
||||
struct ggml_metal_fusion_info * finfo,
|
||||
int idx_start,
|
||||
int idx_end,
|
||||
bool use_fusion,
|
||||
bool use_concurrency,
|
||||
bool use_capture,
|
||||
int debug_graph,
|
||||
int debug_fusion);
|
||||
int debug_graph);
|
||||
|
||||
void ggml_metal_op_free(ggml_metal_op_t ctx);
|
||||
|
||||
|
||||
@@ -4,6 +4,7 @@
|
||||
#include "ggml-backend-impl.h"
|
||||
|
||||
#include "ggml-metal-device.h"
|
||||
#include "ggml-metal-fusion.h"
|
||||
#include "ggml-metal-context.h"
|
||||
#include "ggml-metal-ops.h"
|
||||
#include "ggml-metal-tuning.h"
|
||||
@@ -906,6 +907,30 @@ static const char * ggml_backend_metal_tuning_device_token(ggml_backend_dev_t de
|
||||
return ggml_metal_device_id_token(ggml_metal_device_get_props(ctx_dev)->device_id);
|
||||
}
|
||||
|
||||
// generic fusion debugging API (ad-hoc proc-address mechanism): the test resolves the device
|
||||
// fusion context once and passes that opaque handle to the rest of the functions
|
||||
typedef void * ggml_backend_fusion_t;
|
||||
|
||||
static ggml_backend_fusion_t ggml_backend_metal_fusion_get(ggml_backend_dev_t dev) {
|
||||
return ggml_metal_device_get_fusion_info((ggml_metal_device_t)dev->context);
|
||||
}
|
||||
|
||||
static void ggml_backend_metal_fusion_stats_init(ggml_backend_fusion_t finfo) {
|
||||
ggml_metal_fusion_info_stats_init((struct ggml_metal_fusion_info *) finfo);
|
||||
}
|
||||
|
||||
static void ggml_backend_metal_fusion_stats_reset(ggml_backend_fusion_t finfo) {
|
||||
ggml_metal_fusion_info_stats_reset((struct ggml_metal_fusion_info *) finfo);
|
||||
}
|
||||
|
||||
static int ggml_backend_metal_fusion_stats_get(ggml_backend_fusion_t finfo, const char ** labels, uint64_t * counts, int n) {
|
||||
return ggml_metal_fusion_info_stats_get((struct ggml_metal_fusion_info *) finfo, labels, counts, n);
|
||||
}
|
||||
|
||||
static void ggml_backend_metal_fusion_set_enabled(ggml_backend_fusion_t finfo, bool enabled) {
|
||||
ggml_metal_fusion_info_set_enabled((struct ggml_metal_fusion_info *) finfo, enabled);
|
||||
}
|
||||
|
||||
static void * ggml_backend_metal_get_proc_address(ggml_backend_reg_t reg, const char * name) {
|
||||
if (strcmp(name, "ggml_backend_get_features") == 0) {
|
||||
return (void *)ggml_backend_metal_get_features;
|
||||
@@ -928,6 +953,23 @@ static void * ggml_backend_metal_get_proc_address(ggml_backend_reg_t reg, const
|
||||
if (strcmp(name, "ggml_backend_metal_tuning_device_token") == 0) {
|
||||
return (void *)ggml_backend_metal_tuning_device_token;
|
||||
}
|
||||
// generic fusion debugging API (ad-hoc proc-address mechanism, not part of the official
|
||||
// ggml backend interface yet; a backend that adopts it exports these exact names)
|
||||
if (strcmp(name, "ggml_backend_fusion_get") == 0) {
|
||||
return (void *)ggml_backend_metal_fusion_get;
|
||||
}
|
||||
if (strcmp(name, "ggml_backend_fusion_stats_init") == 0) {
|
||||
return (void *)ggml_backend_metal_fusion_stats_init;
|
||||
}
|
||||
if (strcmp(name, "ggml_backend_fusion_stats_reset") == 0) {
|
||||
return (void *)ggml_backend_metal_fusion_stats_reset;
|
||||
}
|
||||
if (strcmp(name, "ggml_backend_fusion_stats_get") == 0) {
|
||||
return (void *)ggml_backend_metal_fusion_stats_get;
|
||||
}
|
||||
if (strcmp(name, "ggml_backend_fusion_set_enabled") == 0) {
|
||||
return (void *)ggml_backend_metal_fusion_set_enabled;
|
||||
}
|
||||
|
||||
return NULL;
|
||||
|
||||
|
||||
@@ -15,6 +15,7 @@ kernel void kernel_gated_delta_net_impl(
|
||||
device const char * b,
|
||||
device const char * s,
|
||||
device char * dst,
|
||||
device char * dst_fuse,
|
||||
uint3 tgpig[[threadgroup_position_in_grid]],
|
||||
uint3 tpitg[[thread_position_in_threadgroup]],
|
||||
uint3 ntg[[threads_per_threadgroup]]) {
|
||||
@@ -65,6 +66,12 @@ kernel void kernel_gated_delta_net_impl(
|
||||
// per-(seq,head) offset within a slot
|
||||
const uint state_out_base = (i23*args.ne21 + i21)*S_v*S_v + i20*S_v;
|
||||
|
||||
// when fused with the cache cpy, write the snapshots straight into the cache buffer using
|
||||
// the slot stride; otherwise append them after the attn scores (nb_out == 0)
|
||||
const bool fused = args.nb_out > 0;
|
||||
const device float * state_out = fused ? (device float *)dst_fuse : (device float *)dst + attn_size;
|
||||
const uint slot_stride = fused ? (uint)args.nb_out : state_size_per_snap;
|
||||
|
||||
for (short t = 0; t < args.ne22; t++) {
|
||||
float s_k = 0.0f;
|
||||
|
||||
@@ -116,7 +123,7 @@ kernel void kernel_gated_delta_net_impl(
|
||||
if (K > 1) {
|
||||
const int target_slot = (int)args.ne22 - 1 - (int)t;
|
||||
if (target_slot >= 0 && target_slot < (int)K) {
|
||||
device float * dst_state = (device float *) (dst) + attn_size + (uint)target_slot * state_size_per_snap + state_out_base;
|
||||
device float * dst_state = (device float *)state_out + (uint)target_slot * slot_stride + state_out_base;
|
||||
FOR_UNROLL (short j = 0; j < NSG; j++) {
|
||||
const short is = tx*NSG + j;
|
||||
dst_state[is] = ls[j];
|
||||
@@ -126,7 +133,7 @@ kernel void kernel_gated_delta_net_impl(
|
||||
}
|
||||
|
||||
if (K == 1) {
|
||||
device float * dst_state = (device float *) (dst) + attn_size + state_out_base;
|
||||
device float * dst_state = (device float *)state_out + state_out_base;
|
||||
FOR_UNROLL (short j = 0; j < NSG; j++) {
|
||||
const short is = tx*NSG + j;
|
||||
dst_state[is] = ls[j];
|
||||
@@ -158,6 +165,7 @@ kernel void kernel_gated_delta_net_impl(
|
||||
device const char * b,
|
||||
device const char * s,
|
||||
device char * dst,
|
||||
device char * dst_fuse,
|
||||
uint3 tgpig[[threadgroup_position_in_grid]],
|
||||
uint3 tpitg[[thread_position_in_threadgroup]],
|
||||
uint3 ntg[[threads_per_threadgroup]]) {
|
||||
@@ -230,7 +238,13 @@ kernel void kernel_gated_delta_net_impl(
|
||||
dst_attn += args.ne21*S_v;
|
||||
}
|
||||
|
||||
device float * dst_state = (device float *) (dst) + args.ne23*args.ne22*args.ne21*S_v + (i23*args.ne21 + i21)*S_v*S_v + i20;
|
||||
// when fused with the cache cpy, write the snapshots straight into the cache buffer using
|
||||
// the slot stride; otherwise append them after the attn scores (nb_out == 0)
|
||||
const bool fused = args.nb_out > 0;
|
||||
const device float * state_out = fused ? (device float *)dst_fuse : (device float *)dst + args.ne23*args.ne22*args.ne21*S_v;
|
||||
const uint slot_stride = fused ? (uint)args.nb_out : S_v*S_v;
|
||||
|
||||
device float * dst_state = (device float *)state_out + (i23*args.ne21 + i21)*slot_stride + i20;
|
||||
device T * dstt_state = (device T *) (dst_state);
|
||||
|
||||
FOR_UNROLL (short j = 0; j < NSG; j++) {
|
||||
|
||||
@@ -496,6 +496,13 @@ kernel void kernel_mul_mm_id(
|
||||
+ args.nb11*i11
|
||||
+ args.nb10*iy);
|
||||
|
||||
// skip the upper half of the token tile when the expert did not fill it
|
||||
constexpr short NR1H = NR1/2;
|
||||
|
||||
const bool has_hi = nr1 > NR1H;
|
||||
|
||||
const short lb1 = (short) tiitg/NL1; // 0 .. NR1-1, this thread's row of the B tile
|
||||
|
||||
#ifndef GGML_METAL_HAS_TENSOR
|
||||
S0_8x8 ma[4];
|
||||
S1_8x8 mb[2];
|
||||
@@ -505,15 +512,22 @@ kernel void kernel_mul_mm_id(
|
||||
for (short i = 0; i < 8; i++){
|
||||
mc[i] = make_filled_simdgroup_matrix<float, 8>(0.f);
|
||||
}
|
||||
|
||||
// simdgroups 2,3 own rows NR1H..NR1-1
|
||||
const bool sg_active = has_hi || sgitg < 2;
|
||||
#else
|
||||
auto tA = tensor<threadgroup S0, dextents<int32_t, 2>, tensor_inline>(sa, dextents<int32_t, 2>(NK, NR0));
|
||||
auto tB = tensor<threadgroup S1, dextents<int32_t, 2>, tensor_inline>(sb, dextents<int32_t, 2>(NR1, NK ));
|
||||
auto tA = tensor<threadgroup S0, dextents<int32_t, 2>, tensor_inline>(sa, dextents<int32_t, 2>(NK, NR0));
|
||||
|
||||
// sb is [NR1][NK] row-major
|
||||
auto tB0 = tensor<threadgroup S1, dextents<int32_t, 2>, tensor_inline>(sb, dextents<int32_t, 2>(NK, NR1H));
|
||||
auto tB1 = tensor<threadgroup S1, dextents<int32_t, 2>, tensor_inline>(sb + NR1H*NK, dextents<int32_t, 2>(NK, NR1H));
|
||||
|
||||
mpp::tensor_ops::matmul2d<
|
||||
mpp::tensor_ops::matmul2d_descriptor(NR1, NR0, NK, false, true, false, mpp::tensor_ops::matmul2d_descriptor::mode::multiply_accumulate),
|
||||
mpp::tensor_ops::matmul2d_descriptor(NR1H, NR0, NK, false, true, false, mpp::tensor_ops::matmul2d_descriptor::mode::multiply_accumulate),
|
||||
execution_simdgroups<4>> mm;
|
||||
|
||||
auto cT = mm.get_destination_cooperative_tensor<decltype(tA), decltype(tB), float>();
|
||||
auto cT0 = mm.get_destination_cooperative_tensor<decltype(tA), decltype(tB0), float>();
|
||||
auto cT1 = mm.get_destination_cooperative_tensor<decltype(tA), decltype(tB1), float>();
|
||||
#endif
|
||||
|
||||
for (int loop_k = 0; loop_k < args.ne00; loop_k += NK) {
|
||||
@@ -656,37 +670,45 @@ kernel void kernel_mul_mm_id(
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
|
||||
#ifndef GGML_METAL_HAS_TENSOR
|
||||
// load matrices from threadgroup memory and conduct outer products
|
||||
threadgroup const S0 * lsma = (sa + 4*64*(sgitg%2));
|
||||
threadgroup const S1 * lsmb = (sb + 2*64*(sgitg/2));
|
||||
if (sg_active) {
|
||||
// load matrices from threadgroup memory and conduct outer products
|
||||
threadgroup const S0 * lsma = (sa + 4*64*(sgitg%2));
|
||||
threadgroup const S1 * lsmb = (sb + 2*64*(sgitg/2));
|
||||
|
||||
FOR_UNROLL (short ik = 0; ik < NK/8; ik++) {
|
||||
simdgroup_barrier(mem_flags::mem_none);
|
||||
FOR_UNROLL (short ik = 0; ik < NK/8; ik++) {
|
||||
simdgroup_barrier(mem_flags::mem_none);
|
||||
|
||||
FOR_UNROLL (short i = 0; i < 4; i++) {
|
||||
simdgroup_load(ma[i], lsma + 64*i, 8, 0, false);
|
||||
FOR_UNROLL (short i = 0; i < 4; i++) {
|
||||
simdgroup_load(ma[i], lsma + 64*i, 8, 0, false);
|
||||
}
|
||||
|
||||
simdgroup_barrier(mem_flags::mem_none);
|
||||
|
||||
FOR_UNROLL (short i = 0; i < 2; i++) {
|
||||
simdgroup_load(mb[i], lsmb + 64*i, 8, 0, false);
|
||||
}
|
||||
|
||||
simdgroup_barrier(mem_flags::mem_none);
|
||||
|
||||
FOR_UNROLL (short i = 0; i < 8; i++){
|
||||
simdgroup_multiply_accumulate(mc[i], mb[i/4], ma[i%4], mc[i]);
|
||||
}
|
||||
|
||||
lsma += 8*64;
|
||||
lsmb += 4*64;
|
||||
}
|
||||
|
||||
simdgroup_barrier(mem_flags::mem_none);
|
||||
|
||||
FOR_UNROLL (short i = 0; i < 2; i++) {
|
||||
simdgroup_load(mb[i], lsmb + 64*i, 8, 0, false);
|
||||
}
|
||||
|
||||
simdgroup_barrier(mem_flags::mem_none);
|
||||
|
||||
FOR_UNROLL (short i = 0; i < 8; i++){
|
||||
simdgroup_multiply_accumulate(mc[i], mb[i/4], ma[i%4], mc[i]);
|
||||
}
|
||||
|
||||
lsma += 8*64;
|
||||
lsmb += 4*64;
|
||||
}
|
||||
#else
|
||||
auto sA = tA.slice(0, 0);
|
||||
auto sB = tB.slice(0, 0);
|
||||
auto sA = tA.slice(0, 0);
|
||||
auto sB0 = tB0.slice(0, 0);
|
||||
|
||||
mm.run(sB, sA, cT);
|
||||
mm.run(sB0, sA, cT0);
|
||||
|
||||
if (has_hi) {
|
||||
auto sB1 = tB1.slice(0, 0);
|
||||
|
||||
mm.run(sB1, sA, cT1);
|
||||
}
|
||||
#endif
|
||||
}
|
||||
|
||||
@@ -694,13 +716,20 @@ kernel void kernel_mul_mm_id(
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
|
||||
#ifdef GGML_METAL_HAS_TENSOR
|
||||
auto tC = tensor<threadgroup float, dextents<int32_t, 2>, tensor_inline>(sc, dextents<int32_t, 2>(NR0, NR1));
|
||||
cT.store(tC);
|
||||
#else
|
||||
threadgroup float * temp_str = ((threadgroup float *) shmem) + 32*(sgitg&1) + (16*(sgitg >> 1))*NR0;
|
||||
auto tC0 = tensor<threadgroup float, dextents<int32_t, 2>, tensor_inline>(sc, dextents<int32_t, 2>(NR0, NR1H));
|
||||
cT0.store(tC0);
|
||||
|
||||
for (short i = 0; i < 8; i++) {
|
||||
simdgroup_store(mc[i], temp_str + 8*(i%4) + 8*NR0*(i/4), NR0, 0, false);
|
||||
if (has_hi) {
|
||||
auto tC1 = tensor<threadgroup float, dextents<int32_t, 2>, tensor_inline>(sc + NR1H*NR0, dextents<int32_t, 2>(NR0, NR1H));
|
||||
cT1.store(tC1);
|
||||
}
|
||||
#else
|
||||
if (sg_active) {
|
||||
threadgroup float * temp_str = ((threadgroup float *) shmem) + 32*(sgitg&1) + (16*(sgitg >> 1))*NR0;
|
||||
|
||||
for (short i = 0; i < 8; i++) {
|
||||
simdgroup_store(mc[i], temp_str + 8*(i%4) + 8*NR0*(i/4), NR0, 0, false);
|
||||
}
|
||||
}
|
||||
#endif
|
||||
|
||||
|
||||
@@ -1889,8 +1889,19 @@ void kernel_mul_mv_iq2_xxs_f32_impl(
|
||||
const uint i12 = im%FC_mul_mv_ne12;
|
||||
const uint i13 = im/FC_mul_mv_ne12;
|
||||
|
||||
const uint64_t offset0 = first_row*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03;
|
||||
const uint64_t offset1 = r1*args.nb11 + (i12 )*args.nb12 + (i13 )*args.nb13;
|
||||
const int nb32 = nb * (QK_K / 32);
|
||||
|
||||
const short ntx = FC_mul_mv_split ? nb32 : 32;
|
||||
const short nrep = 32 / ntx;
|
||||
|
||||
const short ix = tiisg % ntx;
|
||||
const short irep = tiisg / ntx;
|
||||
|
||||
const short row0 = (nr0 * irep ) / nrep;
|
||||
const short row1 = (nr0 * (irep + 1)) / nrep;
|
||||
|
||||
const uint64_t offset0 = (first_row + row0)*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03;
|
||||
const uint64_t offset1 = r1*args.nb11 + (i12 )*args.nb12 + (i13 )*args.nb13;
|
||||
|
||||
device const block_iq2_xxs * x = (device const block_iq2_xxs *) (src0 + offset0);
|
||||
device const float * y = (device const float *) (src1 + offset1);
|
||||
@@ -1898,8 +1909,6 @@ void kernel_mul_mv_iq2_xxs_f32_impl(
|
||||
float yl[32];
|
||||
float sumf[nr0]={0.f};
|
||||
|
||||
const int nb32 = nb * (QK_K / 32);
|
||||
|
||||
threadgroup uint64_t * svalues = (threadgroup uint64_t *)(shmem);
|
||||
threadgroup uint8_t * ssigns = (threadgroup uint8_t *)(svalues + 256);
|
||||
{
|
||||
@@ -1912,11 +1921,9 @@ void kernel_mul_mv_iq2_xxs_f32_impl(
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
}
|
||||
|
||||
const int ix = tiisg;
|
||||
|
||||
device const float * y4 = y + 32 * ix;
|
||||
|
||||
for (int ib32 = ix; ib32 < nb32; ib32 += 32) {
|
||||
for (int ib32 = ix; ib32 < nb32; ib32 += ntx) {
|
||||
for (short i = 0; i < 32; ++i) {
|
||||
yl[i] = y4[i];
|
||||
}
|
||||
@@ -1928,7 +1935,7 @@ void kernel_mul_mv_iq2_xxs_f32_impl(
|
||||
device const uint16_t * q2 = xr->qs + 4 * ib;
|
||||
device const half * dh = &xr->d;
|
||||
|
||||
for (short row = 0; row < nr0; row++) {
|
||||
for (short row = row0; row < row1; row++) {
|
||||
const float db = dh[0];
|
||||
device const uint8_t * aux8 = (device const uint8_t *)q2;
|
||||
const uint32_t aux32 = q2[2] | (q2[3] << 16);
|
||||
@@ -1948,7 +1955,7 @@ void kernel_mul_mv_iq2_xxs_f32_impl(
|
||||
q2 += args.nb01/2;
|
||||
}
|
||||
|
||||
y4 += 32 * 32;
|
||||
y4 += 32 * ntx;
|
||||
}
|
||||
|
||||
device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
|
||||
@@ -1961,6 +1968,23 @@ void kernel_mul_mv_iq2_xxs_f32_impl(
|
||||
}
|
||||
}
|
||||
|
||||
template<typename args_t>
|
||||
void kernel_mul_mv_iq2_xxs_f32_disp(
|
||||
args_t args,
|
||||
device const char * src0,
|
||||
device const char * src1,
|
||||
device char * dst,
|
||||
threadgroup char * shmem,
|
||||
uint3 tgpig,
|
||||
ushort tiisg,
|
||||
ushort sgitg) {
|
||||
if (FC_mul_mv_split) {
|
||||
kernel_mul_mv_iq2_xxs_f32_impl<N_R0_IQ2_XXS_SPLIT, args_t>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
} else {
|
||||
kernel_mul_mv_iq2_xxs_f32_impl<N_R0_IQ2_XXS, args_t>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
}
|
||||
}
|
||||
|
||||
[[host_name("kernel_mul_mv_iq2_xxs_f32")]]
|
||||
kernel void kernel_mul_mv_iq2_xxs_f32(
|
||||
constant ggml_metal_kargs_mul_mv & args,
|
||||
@@ -1971,7 +1995,7 @@ kernel void kernel_mul_mv_iq2_xxs_f32(
|
||||
uint3 tgpig[[threadgroup_position_in_grid]],
|
||||
ushort tiisg[[thread_index_in_simdgroup]],
|
||||
ushort sgitg[[simdgroup_index_in_threadgroup]]) {
|
||||
kernel_mul_mv_iq2_xxs_f32_impl<N_R0_IQ2_XXS, constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
kernel_mul_mv_iq2_xxs_f32_disp<constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
}
|
||||
|
||||
template<int nr0, typename args_t>
|
||||
@@ -1997,8 +2021,19 @@ void kernel_mul_mv_iq2_xs_f32_impl(
|
||||
const uint i12 = im%FC_mul_mv_ne12;
|
||||
const uint i13 = im/FC_mul_mv_ne12;
|
||||
|
||||
const uint64_t offset0 = first_row*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03;
|
||||
const uint64_t offset1 = r1*args.nb11 + (i12 )*args.nb12 + (i13 )*args.nb13;
|
||||
const int nb32 = nb * (QK_K / 32);
|
||||
|
||||
const short ntx = FC_mul_mv_split ? nb32 : 32;
|
||||
const short nrep = 32 / ntx;
|
||||
|
||||
const short ix = tiisg % ntx;
|
||||
const short irep = tiisg / ntx;
|
||||
|
||||
const short row0 = (nr0 * irep ) / nrep;
|
||||
const short row1 = (nr0 * (irep + 1)) / nrep;
|
||||
|
||||
const uint64_t offset0 = (first_row + row0)*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03;
|
||||
const uint64_t offset1 = r1*args.nb11 + (i12 )*args.nb12 + (i13 )*args.nb13;
|
||||
|
||||
device const block_iq2_xs * x = (device const block_iq2_xs *) (src0 + offset0);
|
||||
device const float * y = (device const float *) (src1 + offset1);
|
||||
@@ -2006,8 +2041,6 @@ void kernel_mul_mv_iq2_xs_f32_impl(
|
||||
float yl[32];
|
||||
float sumf[nr0]={0.f};
|
||||
|
||||
const int nb32 = nb * (QK_K / 32);
|
||||
|
||||
threadgroup uint64_t * svalues = (threadgroup uint64_t *)(shmem);
|
||||
threadgroup uint8_t * ssigns = (threadgroup uint8_t *)(svalues + 512);
|
||||
{
|
||||
@@ -2020,11 +2053,9 @@ void kernel_mul_mv_iq2_xs_f32_impl(
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
}
|
||||
|
||||
const int ix = tiisg;
|
||||
|
||||
device const float * y4 = y + 32 * ix;
|
||||
|
||||
for (int ib32 = ix; ib32 < nb32; ib32 += 32) {
|
||||
for (int ib32 = ix; ib32 < nb32; ib32 += ntx) {
|
||||
for (short i = 0; i < 32; ++i) {
|
||||
yl[i] = y4[i];
|
||||
}
|
||||
@@ -2037,7 +2068,7 @@ void kernel_mul_mv_iq2_xs_f32_impl(
|
||||
device const uint8_t * sc = xr->scales + ib;
|
||||
device const half * dh = &xr->d;
|
||||
|
||||
for (short row = 0; row < nr0; row++) {
|
||||
for (short row = row0; row < row1; row++) {
|
||||
const float db = dh[0];
|
||||
const uint8_t ls1 = sc[0] & 0xf;
|
||||
const uint8_t ls2 = sc[0] >> 4;
|
||||
@@ -2066,7 +2097,7 @@ void kernel_mul_mv_iq2_xs_f32_impl(
|
||||
sc += args.nb01;
|
||||
}
|
||||
|
||||
y4 += 32 * 32;
|
||||
y4 += 32 * ntx;
|
||||
}
|
||||
|
||||
device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
|
||||
@@ -2079,6 +2110,23 @@ void kernel_mul_mv_iq2_xs_f32_impl(
|
||||
}
|
||||
}
|
||||
|
||||
template<typename args_t>
|
||||
void kernel_mul_mv_iq2_xs_f32_disp(
|
||||
args_t args,
|
||||
device const char * src0,
|
||||
device const char * src1,
|
||||
device char * dst,
|
||||
threadgroup char * shmem,
|
||||
uint3 tgpig,
|
||||
ushort tiisg,
|
||||
ushort sgitg) {
|
||||
if (FC_mul_mv_split) {
|
||||
kernel_mul_mv_iq2_xs_f32_impl<N_R0_IQ2_XS_SPLIT, args_t>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
} else {
|
||||
kernel_mul_mv_iq2_xs_f32_impl<N_R0_IQ2_XS, args_t>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
}
|
||||
}
|
||||
|
||||
[[host_name("kernel_mul_mv_iq2_xs_f32")]]
|
||||
kernel void kernel_mul_mv_iq2_xs_f32(
|
||||
constant ggml_metal_kargs_mul_mv & args,
|
||||
@@ -2090,7 +2138,7 @@ kernel void kernel_mul_mv_iq2_xs_f32(
|
||||
ushort tiisg[[thread_index_in_simdgroup]],
|
||||
ushort sgitg[[simdgroup_index_in_threadgroup]]) {
|
||||
|
||||
kernel_mul_mv_iq2_xs_f32_impl<N_R0_IQ2_XS, constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
kernel_mul_mv_iq2_xs_f32_disp<constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
}
|
||||
|
||||
// FC_mul_mv_split: for nb32 < 32 (nb32 divides 32), 32/nb32 threads share each chunk and each takes a slice of the rows
|
||||
@@ -2117,8 +2165,19 @@ void kernel_mul_mv_iq3_xxs_f32_impl(
|
||||
const uint i12 = im%FC_mul_mv_ne12;
|
||||
const uint i13 = im/FC_mul_mv_ne12;
|
||||
|
||||
const uint64_t offset0 = first_row*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03;
|
||||
const uint64_t offset1 = r1*args.nb11 + (i12 )*args.nb12 + (i13 )*args.nb13;
|
||||
const int nb32 = nb * (QK_K / 32);
|
||||
|
||||
const short ntx = FC_mul_mv_split ? nb32 : 32;
|
||||
const short nrep = 32 / ntx;
|
||||
|
||||
const short ix = tiisg % ntx;
|
||||
const short irep = tiisg / ntx;
|
||||
|
||||
const short row0 = (nr0 * irep ) / nrep;
|
||||
const short row1 = (nr0 * (irep + 1)) / nrep;
|
||||
|
||||
const uint64_t offset0 = (first_row + row0)*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03;
|
||||
const uint64_t offset1 = r1*args.nb11 + (i12 )*args.nb12 + (i13 )*args.nb13;
|
||||
|
||||
device const block_iq3_xxs * x = (device const block_iq3_xxs *) (src0 + offset0);
|
||||
device const float * y = (device const float *) (src1 + offset1);
|
||||
@@ -2126,8 +2185,6 @@ void kernel_mul_mv_iq3_xxs_f32_impl(
|
||||
float yl[32];
|
||||
float sumf[nr0]={0.f};
|
||||
|
||||
const int nb32 = nb * (QK_K / 32);
|
||||
|
||||
threadgroup uint32_t * svalues = (threadgroup uint32_t *)(shmem);
|
||||
threadgroup uint8_t * ssigns = (threadgroup uint8_t *)(svalues + 256);
|
||||
{
|
||||
@@ -2140,15 +2197,6 @@ void kernel_mul_mv_iq3_xxs_f32_impl(
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
}
|
||||
|
||||
const short ntx = FC_mul_mv_split ? nb32 : 32;
|
||||
const short nrep = 32 / ntx;
|
||||
|
||||
const short ix = tiisg % ntx;
|
||||
const short irep = tiisg / ntx;
|
||||
|
||||
const short row0 = (nr0 * irep ) / nrep;
|
||||
const short row1 = (nr0 * (irep + 1)) / nrep;
|
||||
|
||||
device const float * y4 = y + 32 * ix;
|
||||
|
||||
for (int ib32 = ix; ib32 < nb32; ib32 += ntx) {
|
||||
@@ -2160,9 +2208,9 @@ void kernel_mul_mv_iq3_xxs_f32_impl(
|
||||
const int ib = ib32 % (QK_K / 32);
|
||||
|
||||
device const block_iq3_xxs * xr = x + ibl;
|
||||
device const uint8_t * q3 = xr->qs + 8 * ib + (uint64_t) row0*args.nb01;
|
||||
device const uint16_t * gas = (device const uint16_t *)(xr->qs + QK_K/4) + 2 * ib + (uint64_t) row0*args.nb01/2;
|
||||
device const half * dh = &xr->d + (uint64_t) row0*args.nb01/2;
|
||||
device const uint8_t * q3 = xr->qs + 8 * ib;
|
||||
device const uint16_t * gas = (device const uint16_t *)(xr->qs + QK_K/4) + 2 * ib;
|
||||
device const half * dh = &xr->d;
|
||||
|
||||
for (short row = row0; row < row1; row++) {
|
||||
const float db = dh[0];
|
||||
@@ -2253,8 +2301,19 @@ void kernel_mul_mv_iq3_s_f32_impl(
|
||||
const uint i12 = im%FC_mul_mv_ne12;
|
||||
const uint i13 = im/FC_mul_mv_ne12;
|
||||
|
||||
const uint64_t offset0 = first_row*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03;
|
||||
const uint64_t offset1 = r1*args.nb11 + (i12 )*args.nb12 + (i13 )*args.nb13;
|
||||
const int nb32 = nb * (QK_K / 32);
|
||||
|
||||
const short ntx = FC_mul_mv_split ? nb32 : 32;
|
||||
const short nrep = 32 / ntx;
|
||||
|
||||
const short ix = tiisg % ntx;
|
||||
const short irep = tiisg / ntx;
|
||||
|
||||
const short row0 = (nr0 * irep ) / nrep;
|
||||
const short row1 = (nr0 * (irep + 1)) / nrep;
|
||||
|
||||
const uint64_t offset0 = (first_row + row0)*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03;
|
||||
const uint64_t offset1 = r1*args.nb11 + (i12 )*args.nb12 + (i13 )*args.nb13;
|
||||
|
||||
device const block_iq3_s * x = (device const block_iq3_s *) (src0 + offset0);
|
||||
device const float * y = (device const float *) (src1 + offset1);
|
||||
@@ -2262,8 +2321,6 @@ void kernel_mul_mv_iq3_s_f32_impl(
|
||||
float yl[32];
|
||||
float sumf[nr0]={0.f};
|
||||
|
||||
const int nb32 = nb * (QK_K / 32);
|
||||
|
||||
threadgroup uint32_t * svalues = (threadgroup uint32_t *) shmem;
|
||||
{
|
||||
int nval = 8;
|
||||
@@ -2272,11 +2329,9 @@ void kernel_mul_mv_iq3_s_f32_impl(
|
||||
threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
}
|
||||
|
||||
const int ix = tiisg;
|
||||
|
||||
device const float * y4 = y + 32 * ix;
|
||||
|
||||
for (int ib32 = ix; ib32 < nb32; ib32 += 32) {
|
||||
for (int ib32 = ix; ib32 < nb32; ib32 += ntx) {
|
||||
for (short i = 0; i < 32; ++i) {
|
||||
yl[i] = y4[i];
|
||||
}
|
||||
@@ -2291,7 +2346,7 @@ void kernel_mul_mv_iq3_s_f32_impl(
|
||||
device const uint8_t * signs = xr->signs + 4 * ib;
|
||||
device const half * dh = &xr->d;
|
||||
|
||||
for (short row = 0; row < nr0; row++) {
|
||||
for (short row = row0; row < row1; row++) {
|
||||
const float db = dh[0];
|
||||
const float d = db * (1 + 2*((sc[0] >> 4*(ib%2)) & 0xf));
|
||||
|
||||
@@ -2315,7 +2370,7 @@ void kernel_mul_mv_iq3_s_f32_impl(
|
||||
signs += args.nb01;
|
||||
}
|
||||
|
||||
y4 += 32 * 32;
|
||||
y4 += 32 * ntx;
|
||||
}
|
||||
|
||||
device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
|
||||
@@ -2328,6 +2383,23 @@ void kernel_mul_mv_iq3_s_f32_impl(
|
||||
}
|
||||
}
|
||||
|
||||
template<typename args_t>
|
||||
void kernel_mul_mv_iq3_s_f32_disp(
|
||||
args_t args,
|
||||
device const char * src0,
|
||||
device const char * src1,
|
||||
device char * dst,
|
||||
threadgroup char * shmem,
|
||||
uint3 tgpig,
|
||||
ushort tiisg,
|
||||
ushort sgitg) {
|
||||
if (FC_mul_mv_split) {
|
||||
kernel_mul_mv_iq3_s_f32_impl<N_R0_IQ3_S_SPLIT, args_t>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
} else {
|
||||
kernel_mul_mv_iq3_s_f32_impl<N_R0_IQ3_S, args_t>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
}
|
||||
}
|
||||
|
||||
[[host_name("kernel_mul_mv_iq3_s_f32")]]
|
||||
kernel void kernel_mul_mv_iq3_s_f32(
|
||||
constant ggml_metal_kargs_mul_mv & args,
|
||||
@@ -2339,7 +2411,7 @@ kernel void kernel_mul_mv_iq3_s_f32(
|
||||
ushort tiisg[[thread_index_in_simdgroup]],
|
||||
ushort sgitg[[simdgroup_index_in_threadgroup]]) {
|
||||
|
||||
kernel_mul_mv_iq3_s_f32_impl<N_R0_IQ3_S, constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
kernel_mul_mv_iq3_s_f32_disp<constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
}
|
||||
|
||||
template<int nr0, typename args_t>
|
||||
@@ -2365,8 +2437,19 @@ void kernel_mul_mv_iq2_s_f32_impl(
|
||||
const uint i12 = im%FC_mul_mv_ne12;
|
||||
const uint i13 = im/FC_mul_mv_ne12;
|
||||
|
||||
const uint64_t offset0 = first_row*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03;
|
||||
const uint64_t offset1 = r1*args.nb11 + (i12 )*args.nb12 + (i13 )*args.nb13;
|
||||
const int nb32 = nb * (QK_K / 32);
|
||||
|
||||
const short ntx = FC_mul_mv_split ? nb32 : 32;
|
||||
const short nrep = 32 / ntx;
|
||||
|
||||
const short ix = tiisg % ntx;
|
||||
const short irep = tiisg / ntx;
|
||||
|
||||
const short row0 = (nr0 * irep ) / nrep;
|
||||
const short row1 = (nr0 * (irep + 1)) / nrep;
|
||||
|
||||
const uint64_t offset0 = (first_row + row0)*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03;
|
||||
const uint64_t offset1 = r1*args.nb11 + (i12 )*args.nb12 + (i13 )*args.nb13;
|
||||
|
||||
device const block_iq2_s * x = (device const block_iq2_s *) (src0 + offset0);
|
||||
device const float * y = (device const float *) (src1 + offset1);
|
||||
@@ -2374,8 +2457,6 @@ void kernel_mul_mv_iq2_s_f32_impl(
|
||||
float yl[32];
|
||||
float sumf[nr0]={0.f};
|
||||
|
||||
const int nb32 = nb * (QK_K / 32);
|
||||
|
||||
//threadgroup uint64_t * svalues = (threadgroup uint64_t *) shmem;
|
||||
//{
|
||||
// int nval = 32;
|
||||
@@ -2384,11 +2465,9 @@ void kernel_mul_mv_iq2_s_f32_impl(
|
||||
// threadgroup_barrier(mem_flags::mem_threadgroup);
|
||||
//}
|
||||
|
||||
const short ix = tiisg;
|
||||
|
||||
device const float * y4 = y + 32 * ix;
|
||||
|
||||
for (int ib32 = ix; ib32 < nb32; ib32 += 32) {
|
||||
for (int ib32 = ix; ib32 < nb32; ib32 += ntx) {
|
||||
for (short i = 0; i < 32; ++i) {
|
||||
yl[i] = y4[i];
|
||||
}
|
||||
@@ -2403,7 +2482,7 @@ void kernel_mul_mv_iq2_s_f32_impl(
|
||||
device const uint8_t * signs = qs + QK_K/8;
|
||||
device const half * dh = &xr->d;
|
||||
|
||||
for (short row = 0; row < nr0; row++) {
|
||||
for (short row = row0; row < row1; row++) {
|
||||
const float db = dh[0];
|
||||
const float d1 = db * (0.5f + (sc[0] & 0xf));
|
||||
const float d2 = db * (0.5f + (sc[0] >> 4));
|
||||
@@ -2428,7 +2507,7 @@ void kernel_mul_mv_iq2_s_f32_impl(
|
||||
signs += args.nb01;
|
||||
}
|
||||
|
||||
y4 += 32 * 32;
|
||||
y4 += 32 * ntx;
|
||||
}
|
||||
|
||||
device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
|
||||
@@ -2441,6 +2520,23 @@ void kernel_mul_mv_iq2_s_f32_impl(
|
||||
}
|
||||
}
|
||||
|
||||
template<typename args_t>
|
||||
void kernel_mul_mv_iq2_s_f32_disp(
|
||||
args_t args,
|
||||
device const char * src0,
|
||||
device const char * src1,
|
||||
device char * dst,
|
||||
threadgroup char * shmem,
|
||||
uint3 tgpig,
|
||||
ushort tiisg,
|
||||
ushort sgitg) {
|
||||
if (FC_mul_mv_split) {
|
||||
kernel_mul_mv_iq2_s_f32_impl<N_R0_IQ2_S_SPLIT, args_t>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
} else {
|
||||
kernel_mul_mv_iq2_s_f32_impl<N_R0_IQ2_S, args_t>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
}
|
||||
}
|
||||
|
||||
[[host_name("kernel_mul_mv_iq2_s_f32")]]
|
||||
kernel void kernel_mul_mv_iq2_s_f32(
|
||||
constant ggml_metal_kargs_mul_mv & args,
|
||||
@@ -2452,7 +2548,7 @@ kernel void kernel_mul_mv_iq2_s_f32(
|
||||
ushort tiisg[[thread_index_in_simdgroup]],
|
||||
ushort sgitg[[simdgroup_index_in_threadgroup]]) {
|
||||
|
||||
kernel_mul_mv_iq2_s_f32_impl<N_R0_IQ2_S, constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
kernel_mul_mv_iq2_s_f32_disp<constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
}
|
||||
|
||||
template<int nr0, typename args_t>
|
||||
@@ -2478,8 +2574,19 @@ void kernel_mul_mv_iq1_s_f32_impl(
|
||||
const uint i12 = im%FC_mul_mv_ne12;
|
||||
const uint i13 = im/FC_mul_mv_ne12;
|
||||
|
||||
const uint64_t offset0 = first_row*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03;
|
||||
const uint64_t offset1 = r1*args.nb11 + (i12 )*args.nb12 + (i13 )*args.nb13;
|
||||
const int nb32 = nb * (QK_K / 32);
|
||||
|
||||
const short ntx = FC_mul_mv_split ? nb32 : 32;
|
||||
const short nrep = 32 / ntx;
|
||||
|
||||
const short ix = tiisg % ntx;
|
||||
const short irep = tiisg / ntx;
|
||||
|
||||
const short row0 = (nr0 * irep ) / nrep;
|
||||
const short row1 = (nr0 * (irep + 1)) / nrep;
|
||||
|
||||
const uint64_t offset0 = (first_row + row0)*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03;
|
||||
const uint64_t offset1 = r1*args.nb11 + (i12 )*args.nb12 + (i13 )*args.nb13;
|
||||
|
||||
device const block_iq1_s * x = (device const block_iq1_s *) (src0 + offset0);
|
||||
device const float * y = (device const float *) (src1 + offset1);
|
||||
@@ -2487,13 +2594,9 @@ void kernel_mul_mv_iq1_s_f32_impl(
|
||||
float yl[32];
|
||||
float sumf[nr0]={0.f};
|
||||
|
||||
const int nb32 = nb * (QK_K / 32);
|
||||
|
||||
const short ix = tiisg;
|
||||
|
||||
device const float * y4 = y + 32 * ix;
|
||||
|
||||
for (int ib32 = ix; ib32 < nb32; ib32 += 32) {
|
||||
for (int ib32 = ix; ib32 < nb32; ib32 += ntx) {
|
||||
float sumy = 0;
|
||||
for (short i = 0; i < 32; ++i) {
|
||||
yl[i] = y4[i];
|
||||
@@ -2508,7 +2611,7 @@ void kernel_mul_mv_iq1_s_f32_impl(
|
||||
device const uint16_t * qh = xr->qh + ib;
|
||||
device const half * dh = &xr->d;
|
||||
|
||||
for (short row = 0; row < nr0; row++) {
|
||||
for (short row = row0; row < row1; row++) {
|
||||
constant uint8_t * grid1 = (constant uint8_t *)(iq1s_grid_gpu + (qs[0] | ((qh[0] << 8) & 0x700)));
|
||||
constant uint8_t * grid2 = (constant uint8_t *)(iq1s_grid_gpu + (qs[1] | ((qh[0] << 5) & 0x700)));
|
||||
constant uint8_t * grid3 = (constant uint8_t *)(iq1s_grid_gpu + (qs[2] | ((qh[0] << 2) & 0x700)));
|
||||
@@ -2528,7 +2631,7 @@ void kernel_mul_mv_iq1_s_f32_impl(
|
||||
qh += args.nb01/2;
|
||||
}
|
||||
|
||||
y4 += 32 * 32;
|
||||
y4 += 32 * ntx;
|
||||
}
|
||||
|
||||
device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
|
||||
@@ -2541,6 +2644,23 @@ void kernel_mul_mv_iq1_s_f32_impl(
|
||||
}
|
||||
}
|
||||
|
||||
template<typename args_t>
|
||||
void kernel_mul_mv_iq1_s_f32_disp(
|
||||
args_t args,
|
||||
device const char * src0,
|
||||
device const char * src1,
|
||||
device char * dst,
|
||||
threadgroup char * shmem,
|
||||
uint3 tgpig,
|
||||
ushort tiisg,
|
||||
ushort sgitg) {
|
||||
if (FC_mul_mv_split) {
|
||||
kernel_mul_mv_iq1_s_f32_impl<N_R0_IQ1_S_SPLIT, args_t>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
} else {
|
||||
kernel_mul_mv_iq1_s_f32_impl<N_R0_IQ1_S, args_t>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
}
|
||||
}
|
||||
|
||||
[[host_name("kernel_mul_mv_iq1_s_f32")]]
|
||||
kernel void kernel_mul_mv_iq1_s_f32(
|
||||
constant ggml_metal_kargs_mul_mv & args,
|
||||
@@ -2551,7 +2671,7 @@ kernel void kernel_mul_mv_iq1_s_f32(
|
||||
ushort tiisg[[thread_index_in_simdgroup]],
|
||||
ushort sgitg[[simdgroup_index_in_threadgroup]]) {
|
||||
|
||||
kernel_mul_mv_iq1_s_f32_impl<N_R0_IQ1_S, constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg);
|
||||
kernel_mul_mv_iq1_s_f32_disp<constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg);
|
||||
}
|
||||
|
||||
template<int nr0, typename args_t>
|
||||
@@ -2577,8 +2697,19 @@ void kernel_mul_mv_iq1_m_f32_impl(
|
||||
const uint i12 = im%FC_mul_mv_ne12;
|
||||
const uint i13 = im/FC_mul_mv_ne12;
|
||||
|
||||
const uint64_t offset0 = first_row*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03;
|
||||
const uint64_t offset1 = r1*args.nb11 + (i12 )*args.nb12 + (i13 )*args.nb13;
|
||||
const int nb32 = nb * (QK_K / 32);
|
||||
|
||||
const short ntx = FC_mul_mv_split ? nb32 : 32;
|
||||
const short nrep = 32 / ntx;
|
||||
|
||||
const short ix = tiisg % ntx;
|
||||
const short irep = tiisg / ntx;
|
||||
|
||||
const short row0 = (nr0 * irep ) / nrep;
|
||||
const short row1 = (nr0 * (irep + 1)) / nrep;
|
||||
|
||||
const uint64_t offset0 = (first_row + row0)*args.nb01 + (i12/FC_mul_mv_r2)*args.nb02 + (i13/FC_mul_mv_r3)*args.nb03;
|
||||
const uint64_t offset1 = r1*args.nb11 + (i12 )*args.nb12 + (i13 )*args.nb13;
|
||||
|
||||
device const block_iq1_m * x = (device const block_iq1_m *) (src0 + offset0);
|
||||
device const float * y = (device const float *) (src1 + offset1);
|
||||
@@ -2586,15 +2717,11 @@ void kernel_mul_mv_iq1_m_f32_impl(
|
||||
float yl[32];
|
||||
float sumf[nr0]={0.f};
|
||||
|
||||
const int nb32 = nb * (QK_K / 32);
|
||||
|
||||
const short ix = tiisg;
|
||||
|
||||
device const float * y4 = y + 32 * ix;
|
||||
|
||||
iq1m_scale_t scale;
|
||||
|
||||
for (int ib32 = ix; ib32 < nb32; ib32 += 32) {
|
||||
for (int ib32 = ix; ib32 < nb32; ib32 += ntx) {
|
||||
float4 sumy = {0.f};
|
||||
for (short i = 0; i < 8; ++i) {
|
||||
yl[i+ 0] = y4[i+ 0]; sumy[0] += yl[i+ 0];
|
||||
@@ -2611,7 +2738,7 @@ void kernel_mul_mv_iq1_m_f32_impl(
|
||||
device const uint8_t * qh = xr->qh + 2 * ib;
|
||||
device const uint16_t * sc = (device const uint16_t *)xr->scales;
|
||||
|
||||
for (short row = 0; row < nr0; row++) {
|
||||
for (short row = row0; row < row1; row++) {
|
||||
scale.u16 = (sc[0] >> 12) | ((sc[1] >> 8) & 0x00f0) | ((sc[2] >> 4) & 0x0f00) | (sc[3] & 0xf000);
|
||||
|
||||
constant uint8_t * grid1 = (constant uint8_t *)(iq1s_grid_gpu + (qs[0] | ((qh[0] << 8) & 0x700)));
|
||||
@@ -2637,7 +2764,7 @@ void kernel_mul_mv_iq1_m_f32_impl(
|
||||
qh += args.nb01;
|
||||
}
|
||||
|
||||
y4 += 32 * 32;
|
||||
y4 += 32 * ntx;
|
||||
}
|
||||
|
||||
device float * dst_f32 = (device float *) dst + (uint64_t)im*args.ne0*args.ne1 + (uint64_t)r1*args.ne0;
|
||||
@@ -2650,6 +2777,23 @@ void kernel_mul_mv_iq1_m_f32_impl(
|
||||
}
|
||||
}
|
||||
|
||||
template<typename args_t>
|
||||
void kernel_mul_mv_iq1_m_f32_disp(
|
||||
args_t args,
|
||||
device const char * src0,
|
||||
device const char * src1,
|
||||
device char * dst,
|
||||
threadgroup char * shmem,
|
||||
uint3 tgpig,
|
||||
ushort tiisg,
|
||||
ushort sgitg) {
|
||||
if (FC_mul_mv_split) {
|
||||
kernel_mul_mv_iq1_m_f32_impl<N_R0_IQ1_M_SPLIT, args_t>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
} else {
|
||||
kernel_mul_mv_iq1_m_f32_impl<N_R0_IQ1_M, args_t>(args, src0, src1, dst, shmem, tgpig, tiisg, sgitg);
|
||||
}
|
||||
}
|
||||
|
||||
[[host_name("kernel_mul_mv_iq1_m_f32")]]
|
||||
kernel void kernel_mul_mv_iq1_m_f32(
|
||||
constant ggml_metal_kargs_mul_mv & args,
|
||||
@@ -2660,7 +2804,7 @@ kernel void kernel_mul_mv_iq1_m_f32(
|
||||
ushort tiisg[[thread_index_in_simdgroup]],
|
||||
ushort sgitg[[simdgroup_index_in_threadgroup]]) {
|
||||
|
||||
kernel_mul_mv_iq1_m_f32_impl<N_R0_IQ1_M, constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg);
|
||||
kernel_mul_mv_iq1_m_f32_disp<constant ggml_metal_kargs_mul_mv &>(args, src0, src1, dst, nullptr, tgpig, tiisg, sgitg);
|
||||
}
|
||||
|
||||
template<int NR0, typename args_t>
|
||||
@@ -3239,13 +3383,13 @@ template [[host_name("kernel_mul_mv_id_q3_K_f32")]] kernel kernel_mul_mv_id_t
|
||||
template [[host_name("kernel_mul_mv_id_q4_K_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_q4_K_f32_impl <N_R0_Q4_K>>>;
|
||||
template [[host_name("kernel_mul_mv_id_q5_K_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_q5_K_f32_impl <N_R0_Q5_K>>>;
|
||||
template [[host_name("kernel_mul_mv_id_q6_K_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_q6_K_f32_impl <N_R0_Q6_K>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq1_s_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq1_s_f32_impl <N_R0_IQ1_S>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq1_m_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq1_m_f32_impl <N_R0_IQ1_M>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq2_xxs_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq2_xxs_f32_impl<N_R0_IQ2_XXS>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq2_xs_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq2_xs_f32_impl <N_R0_IQ2_XS>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq1_s_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq1_s_f32_disp<ggml_metal_kargs_mul_mv>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq1_m_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq1_m_f32_disp<ggml_metal_kargs_mul_mv>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq2_xxs_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq2_xxs_f32_disp<ggml_metal_kargs_mul_mv>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq2_xs_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq2_xs_f32_disp<ggml_metal_kargs_mul_mv>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq3_xxs_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq3_xxs_f32_disp<ggml_metal_kargs_mul_mv>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq3_s_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq3_s_f32_impl <N_R0_IQ3_S>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq2_s_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq2_s_f32_impl <N_R0_IQ2_S>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq3_s_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq3_s_f32_disp<ggml_metal_kargs_mul_mv>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq2_s_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq2_s_f32_disp<ggml_metal_kargs_mul_mv>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq4_nl_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq4_nl_f32_impl <N_R0_IQ4_NL>>>;
|
||||
template [[host_name("kernel_mul_mv_id_iq4_xs_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_iq4_xs_f32_impl <N_R0_IQ4_XS>>>;
|
||||
template [[host_name("kernel_mul_mv_id_tq2_0_f32")]] kernel kernel_mul_mv_id_t kernel_mul_mv_id<mmv_fn<kernel_mul_mv_tq2_0_f32_impl <N_R0_TQ2_0>>>;
|
||||
|
||||
@@ -5430,8 +5430,9 @@ static void ggml_vk_load_shaders(vk_device& device, vk_pipeline requested) {
|
||||
bool prefer_large = tiles_m > shader_core_count || tiles_l > shader_core_count ||
|
||||
(tiles_l <= shader_core_count / 3 && tiles_m > shader_core_count / 2);
|
||||
if (n > crossover_large && prefer_large) return last;
|
||||
uint32_t crossover_medium = configs[0].unaligned->wg_denoms[1];
|
||||
if (n > crossover_medium) return 1;
|
||||
uint32_t crossover_medium_m = configs[0].unaligned->wg_denoms[0];
|
||||
uint32_t crossover_medium_n = configs[0].unaligned->wg_denoms[1];
|
||||
if (m > crossover_medium_m && n > crossover_medium_n) return 1;
|
||||
return 0;
|
||||
};
|
||||
device->matmul_id_tile_selector = [](uint32_t /*m*/, uint32_t n, uint32_t /*k*/, uint32_t /*shader_core_count*/,
|
||||
@@ -9061,7 +9062,7 @@ static uint32_t ggml_vk_guess_split_k(ggml_backend_vk_context * ctx, uint32_t m,
|
||||
}
|
||||
|
||||
uint32_t split_k = 1;
|
||||
if (ctx->device->shader_core_count != 0 && m >= pipeline->wg_denoms[0] && n >= pipeline->wg_denoms[1]) {
|
||||
if (ctx->device->shader_core_count != 0 && n >= pipeline->wg_denoms[1]) {
|
||||
// If k is 'large' and the SMs will fill less than halfway, use split_k.
|
||||
uint32_t m_tiles = CEIL_DIV(m, pipeline->wg_denoms[0]);
|
||||
uint32_t n_tiles = CEIL_DIV(n, pipeline->wg_denoms[1]);
|
||||
@@ -9814,10 +9815,10 @@ static bool ggml_vk_should_use_mmvq(const vk_device& device, uint32_t m, uint32_
|
||||
GGML_UNUSED(m);
|
||||
}
|
||||
|
||||
static void ggml_vk_mul_mat_vec_q_f16(ggml_backend_vk_context * ctx, vk_context& subctx, const struct ggml_cgraph * cgraph, int node_idx) {
|
||||
static void ggml_vk_mul_mat_vec_q_f16(ggml_backend_vk_context * ctx, vk_context& subctx, const struct ggml_cgraph * cgraph, int node_idx, bool swap_inputs = false) {
|
||||
ggml_tensor * dst = cgraph->nodes[node_idx];
|
||||
const ggml_tensor * src0 = dst->src[0];
|
||||
const ggml_tensor * src1 = dst->src[1];
|
||||
const ggml_tensor * src0 = dst->src[swap_inputs ? 1 : 0];
|
||||
const ggml_tensor * src1 = dst->src[swap_inputs ? 0 : 1];
|
||||
|
||||
VK_LOG_DEBUG("ggml_vk_mul_mat_vec_q_f16((" << src0 << ", name=" << src0->name << ", type=" << src0->type << ", ne0=" << src0->ne[0] << ", ne1=" << src0->ne[1] << ", ne2=" << src0->ne[2] << ", ne3=" << src0->ne[3] << ", nb0=" << src0->nb[0] << ", nb1=" << src0->nb[1] << ", nb2=" << src0->nb[2] << ", nb3=" << src0->nb[3];
|
||||
std::cerr << "), (" << src1 << ", name=" << src1->name << ", type=" << src1->type << ", ne0=" << src1->ne[0] << ", ne1=" << src1->ne[1] << ", ne2=" << src1->ne[2] << ", ne3=" << src1->ne[3] << ", nb0=" << src1->nb[0] << ", nb1=" << src1->nb[1] << ", nb2=" << src1->nb[2] << ", nb3=" << src1->nb[3];
|
||||
@@ -9836,8 +9837,8 @@ static void ggml_vk_mul_mat_vec_q_f16(ggml_backend_vk_context * ctx, vk_context&
|
||||
const uint64_t ne12 = src1->ne[2];
|
||||
const uint64_t ne13 = src1->ne[3];
|
||||
|
||||
const uint64_t ne20 = dst->ne[0];
|
||||
const uint64_t ne21 = dst->ne[1];
|
||||
const uint64_t ne20 = dst->ne[swap_inputs ? 1 : 0];
|
||||
const uint64_t ne21 = dst->ne[swap_inputs ? 0 : 1];
|
||||
// const uint64_t ne22 = dst->ne[2];
|
||||
// const uint64_t ne23 = dst->ne[3];
|
||||
|
||||
@@ -10451,6 +10452,16 @@ static void ggml_vk_mul_mat(ggml_backend_vk_context * ctx, vk_context& subctx, c
|
||||
src0->ne[1] <= ctx->device->properties.limits.maxComputeWorkGroupCount[1] &&
|
||||
src1->ne[2] <= ctx->device->properties.limits.maxComputeWorkGroupCount[2]) {
|
||||
ggml_vk_mul_mat_vec_nc_f16_f32(ctx, subctx, cgraph, node_idx);
|
||||
// With one output row, B^T*A has the same flat output as A^T*B.
|
||||
} else if (ctx->num_additional_fused_ops == 0 &&
|
||||
(src0->type == GGML_TYPE_F32 || src0->type == GGML_TYPE_F16) &&
|
||||
(src1->type == GGML_TYPE_F32 || src1->type == GGML_TYPE_F16 || src1->type == GGML_TYPE_BF16 || ggml_is_quantized(src1->type)) &&
|
||||
dst->ne[0] == 1 && dst->ne[1] > mul_mat_vec_max_cols &&
|
||||
src0->ne[2] == 1 && src0->ne[3] == 1 &&
|
||||
src1->ne[2] == 1 && src1->ne[3] == 1 &&
|
||||
ggml_is_contiguous(src0) && ggml_is_contiguous(src1) && ggml_is_contiguous(dst) &&
|
||||
get_misalign_bytes(ctx, src0) == 0 && get_misalign_bytes(ctx, src1) == 0 && get_misalign_bytes(ctx, dst) == 0) {
|
||||
ggml_vk_mul_mat_vec_q_f16(ctx, subctx, cgraph, node_idx, true);
|
||||
// mul_mat_vec supports batching ne12*ne13 when ne11==1, or treating ne11 as the batch size (up to four)
|
||||
// when ne12 and ne13 are one.
|
||||
} else if ((dst->ne[1] == 1 || (dst->ne[1] <= mul_mat_vec_max_cols && src1->ne[2] * src1->ne[3] == 1)) &&
|
||||
@@ -17158,6 +17169,22 @@ static bool ggml_backend_vk_cpy_tensor_async(ggml_backend_t backend_src, ggml_ba
|
||||
return false;
|
||||
}
|
||||
|
||||
// If the backend is idle, use a CPU copy to avoid GPU synchronization overhead.
|
||||
static constexpr size_t max_cpu_copy_size = 128 * 1024;
|
||||
const bool src_backend_synchronous = backend_src->iface.synchronize == nullptr;
|
||||
const bool transfer_idle = !ctx->device->async_use_transfer_queue ||
|
||||
ctx->transfer_semaphore_last_submitted == ctx->transfer_semaphore.value;
|
||||
const bool backend_idle = ctx->compute_ctx.expired() && ctx->transfer_ctx.expired() &&
|
||||
!ctx->submit_pending && !ctx->almost_ready_fence_pending && transfer_idle;
|
||||
const bool dst_host_coherent =
|
||||
(dst_buf->memory_property_flags & (vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent)) ==
|
||||
(vk::MemoryPropertyFlagBits::eHostVisible | vk::MemoryPropertyFlagBits::eHostCoherent);
|
||||
|
||||
if ((backend_src == backend_dst || src_backend_synchronous) && backend_idle && dst_host_coherent && ggml_nbytes(src) <= max_cpu_copy_size) {
|
||||
ggml_vk_buffer_write(dst_buf, vk_tensor_offset(dst) + dst->view_offs, src->data, ggml_nbytes(src));
|
||||
return true;
|
||||
}
|
||||
|
||||
vk_context cpy_ctx;
|
||||
if (ctx->device->async_use_transfer_queue) {
|
||||
cpy_ctx = ggml_vk_get_transfer_ctx(ctx);
|
||||
@@ -17170,7 +17197,6 @@ static bool ggml_backend_vk_cpy_tensor_async(ggml_backend_t backend_src, ggml_ba
|
||||
src->data, ggml_nbytes(src));
|
||||
}
|
||||
|
||||
GGML_UNUSED(backend_src);
|
||||
return false;
|
||||
}
|
||||
|
||||
@@ -18351,38 +18377,30 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg
|
||||
|
||||
bool need_disable = false;
|
||||
|
||||
// topk_moe often overwrites the source, but for a given row all the src values are
|
||||
// loaded before anything is stored. If there's only one row, this is safe, so treat
|
||||
// this as a special case.
|
||||
bool is_topk_moe_single_row = ctx->fused_topk_moe_mode != TOPK_MOE_COUNT &&
|
||||
ggml_nrows(cgraph->nodes[i]->src[0]) == 1;
|
||||
|
||||
if (!is_topk_moe_single_row) {
|
||||
for (int j = 0; j < 2; ++j) {
|
||||
ggml_tensor *dst = output_nodes[j];
|
||||
if (!dst) {
|
||||
continue;
|
||||
}
|
||||
// Loop over all srcs of all nodes in the fusion. If the src overlaps
|
||||
// the destination and the src is not an intermediate node that's being
|
||||
// elided, then disable fusion.
|
||||
for (int k = 0; k <= ctx->num_additional_fused_ops; ++k) {
|
||||
for (uint32_t s = 0; s < GGML_MAX_SRC; ++s) {
|
||||
ggml_tensor *src = cgraph->nodes[i + k]->src[s];
|
||||
if (!src || src->op == GGML_OP_NONE) {
|
||||
continue;
|
||||
for (int j = 0; j < 2; ++j) {
|
||||
ggml_tensor *dst = output_nodes[j];
|
||||
if (!dst) {
|
||||
continue;
|
||||
}
|
||||
// Loop over all srcs of all nodes in the fusion. If the src overlaps
|
||||
// the destination and the src is not an intermediate node that's being
|
||||
// elided, then disable fusion.
|
||||
for (int k = 0; k <= ctx->num_additional_fused_ops; ++k) {
|
||||
for (uint32_t s = 0; s < GGML_MAX_SRC; ++s) {
|
||||
ggml_tensor *src = cgraph->nodes[i + k]->src[s];
|
||||
if (!src || src->op == GGML_OP_NONE) {
|
||||
continue;
|
||||
}
|
||||
if (ggml_vk_tensors_overlap(src, dst, op_srcs_fused_elementwise[k])) {
|
||||
bool found = false;
|
||||
for (int n = 0; n < k; ++n) {
|
||||
if (cgraph->nodes[i + n] == src) {
|
||||
found = true;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (ggml_vk_tensors_overlap(src, dst, op_srcs_fused_elementwise[k])) {
|
||||
bool found = false;
|
||||
for (int n = 0; n < k; ++n) {
|
||||
if (cgraph->nodes[i + n] == src) {
|
||||
found = true;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (!found) {
|
||||
need_disable = true;
|
||||
}
|
||||
if (!found) {
|
||||
need_disable = true;
|
||||
}
|
||||
}
|
||||
}
|
||||
@@ -18395,6 +18413,7 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg
|
||||
ctx->fused_topk_moe_scale = false;
|
||||
ctx->fused_topk_qsa = false;
|
||||
ctx->fused_rms_norm_mode = RMS_NORM_COUNT;
|
||||
fusion_string = nullptr;
|
||||
}
|
||||
}
|
||||
|
||||
@@ -18497,7 +18516,6 @@ static ggml_status ggml_backend_vk_graph_compute(ggml_backend_t backend, ggml_cg
|
||||
// Sort the graph for improved parallelism.
|
||||
static void ggml_vk_graph_optimize(ggml_backend_t backend, struct ggml_cgraph * graph, struct ggml_backend_graph_optimize_params * params)
|
||||
{
|
||||
GGML_UNUSED(params);
|
||||
VK_LOG_DEBUG("ggml_vk_graph_optimize(" << graph->n_nodes << " nodes)");
|
||||
ggml_backend_vk_context * ctx = (ggml_backend_vk_context *)backend->context;
|
||||
|
||||
@@ -18583,19 +18601,50 @@ static void ggml_vk_graph_optimize(ggml_backend_t backend, struct ggml_cgraph *
|
||||
return false;
|
||||
};
|
||||
|
||||
if (keep_pattern(topk_moe_early_softmax_norm)) {
|
||||
auto const &add_pattern_alloc_deps = [&](const std::initializer_list<ggml_op> &pattern, int last_node) {
|
||||
// Keep external inputs alive through the fused output.
|
||||
std::set<ggml_tensor *> seen;
|
||||
for (size_t j = 0; j < pattern.size(); ++j) {
|
||||
ggml_tensor * node = graph->nodes[first_unused + j];
|
||||
for (uint32_t s = 0; s < GGML_MAX_SRC; ++s) {
|
||||
ggml_tensor * src = node->src[s];
|
||||
if (src && seen.insert(src).second) {
|
||||
params->add_alloc_dep(params->user_data, src, graph->nodes[last_node]);
|
||||
}
|
||||
}
|
||||
seen.insert(node);
|
||||
}
|
||||
};
|
||||
|
||||
auto const &keep_topk_moe_pattern = [&](const std::initializer_list<ggml_op> &pattern) -> bool {
|
||||
if (!match_pattern(pattern, first_unused)) {
|
||||
return false;
|
||||
}
|
||||
|
||||
int last_node = first_unused + (int) pattern.size() - 1;
|
||||
// Some TOPK_MOE variants fuse a trailing scale.
|
||||
if (last_node + 1 < graph->n_nodes && graph->nodes[last_node + 1]->op == GGML_OP_SCALE) {
|
||||
last_node++;
|
||||
}
|
||||
|
||||
add_pattern_alloc_deps(pattern, last_node);
|
||||
|
||||
return keep_pattern(pattern);
|
||||
};
|
||||
|
||||
if (keep_topk_moe_pattern(topk_moe_early_softmax_norm)) {
|
||||
continue;
|
||||
}
|
||||
if (keep_pattern(topk_moe_sigmoid_norm_bias)) {
|
||||
if (keep_topk_moe_pattern(topk_moe_sigmoid_norm_bias)) {
|
||||
continue;
|
||||
}
|
||||
if (keep_pattern(topk_moe_sqrt_softplus_norm_bias)) {
|
||||
if (keep_topk_moe_pattern(topk_moe_sqrt_softplus_norm_bias)) {
|
||||
continue;
|
||||
}
|
||||
if (keep_pattern(topk_moe_early_softmax)) {
|
||||
if (keep_topk_moe_pattern(topk_moe_early_softmax)) {
|
||||
continue;
|
||||
}
|
||||
if (keep_pattern(topk_moe_late_softmax)) {
|
||||
if (keep_topk_moe_pattern(topk_moe_late_softmax)) {
|
||||
continue;
|
||||
}
|
||||
if (keep_pattern(snake_pattern)) {
|
||||
|
||||
@@ -33,7 +33,11 @@ void argsort(bool needs_bounds_check, const uint row) {
|
||||
const uint row_offset = row * p.ncols;
|
||||
|
||||
// initialize indices
|
||||
dst_row[col] = ivec2(col, floatBitsToInt(data_a[row_offset + col]));
|
||||
ivec2 value = ivec2(col, 0);
|
||||
if (!needs_bounds_check || col < p.ncols) {
|
||||
value.y = floatBitsToInt(data_a[row_offset + col]);
|
||||
}
|
||||
dst_row[col] = value;
|
||||
barrier();
|
||||
|
||||
uint num_outer_loop_iters = NCOLS_PADDED_LOG2;
|
||||
@@ -42,18 +46,20 @@ void argsort(bool needs_bounds_check, const uint row) {
|
||||
[[unroll]] for (uint j = k / 2, inner_idx = 0; inner_idx < num_inner_loop_iters; j /= 2, inner_idx++) {
|
||||
const int ixj = int(col ^ j);
|
||||
|
||||
int idx_0 = (col & k) == 0 ? col : ixj;
|
||||
int idx_1 = (col & k) == 0 ? ixj : col;
|
||||
if (ixj > col) {
|
||||
int idx_0 = (col & k) == 0 ? col : ixj;
|
||||
int idx_1 = (col & k) == 0 ? ixj : col;
|
||||
|
||||
ivec2 sh_idx_0 = dst_row[idx_0];
|
||||
ivec2 sh_idx_1 = dst_row[idx_1];
|
||||
bool idx_0_oob = needs_bounds_check ? sh_idx_0.x >= p.ncols : false;
|
||||
bool idx_1_oob = needs_bounds_check ? sh_idx_1.x >= p.ncols : false;
|
||||
ivec2 sh_idx_0 = dst_row[idx_0];
|
||||
ivec2 sh_idx_1 = dst_row[idx_1];
|
||||
bool idx_0_oob = needs_bounds_check ? sh_idx_0.x >= p.ncols : false;
|
||||
bool idx_1_oob = needs_bounds_check ? sh_idx_1.x >= p.ncols : false;
|
||||
|
||||
if ((idx_0_oob ||
|
||||
(!idx_1_oob && intBitsToFloat(sh_idx_0.y) > intBitsToFloat(sh_idx_1.y))) && (ixj > col)) {
|
||||
dst_row[idx_0] = sh_idx_1;
|
||||
dst_row[idx_1] = sh_idx_0;
|
||||
if (idx_0_oob ||
|
||||
(!idx_1_oob && intBitsToFloat(sh_idx_0.y) > intBitsToFloat(sh_idx_1.y))) {
|
||||
dst_row[idx_0] = sh_idx_1;
|
||||
dst_row[idx_1] = sh_idx_0;
|
||||
}
|
||||
}
|
||||
|
||||
barrier();
|
||||
|
||||
@@ -42,7 +42,10 @@ void argsort(bool needs_bounds_check, const uint row) {
|
||||
[[unroll]] for (int u = 0; u < WG_UNROLL_FACTOR; ++u) {
|
||||
uint c = u*BLOCK_SIZE + col;
|
||||
if (c < p.ncols_padded) {
|
||||
ivec2 v = ivec2(c, floatBitsToInt(data_a[row_offset + c]));
|
||||
ivec2 v = ivec2(c, 0);
|
||||
if (!needs_bounds_check || c < p.ncols) {
|
||||
v.y = floatBitsToInt(data_a[row_offset + c]);
|
||||
}
|
||||
tmp_idx[idx_offset + c] = v;
|
||||
}
|
||||
}
|
||||
|
||||
Reference in New Issue
Block a user