mirror of
https://github.com/ggml-org/llama.cpp.git
synced 2026-11-13 14:27:38 -06:00
Compare commits
| Author | SHA1 | Date | |
|---|---|---|---|
 AMIR
|
5193ef1d4f | ||
|
Max Krasnyansky
|
4df9b47976 | ||
|
thelittlefireman
|
b79696f1f5 | ||
|
Pascal
|
4f9dfa0e36 | ||
|
Max Krasnyansky
|
9daa0aa22b | ||
|
ynankani
|
fe14cad51b |
@@ -29,6 +29,7 @@ if (CUDAToolkit_FOUND)
|
||||
# 86 == RTX 3000, needs CUDA v11.1
|
||||
# 89 == RTX 4000, needs CUDA v11.8
|
||||
# 90 == Hopper H100/200, needs CUDA v11.8
|
||||
# 103 == Blackwell GB300, needs CUDA v12.9
|
||||
# 120 == Blackwell, needs CUDA v12.8, FP4 tensor cores
|
||||
#
|
||||
# XX-virtual == compile CUDA code as PTX, do JIT compilation to binary code on first run
|
||||
@@ -64,7 +65,7 @@ if (CUDAToolkit_FOUND)
|
||||
list(APPEND CMAKE_CUDA_ARCHITECTURES 120a-real)
|
||||
endif()
|
||||
if (CUDAToolkit_VERSION VERSION_GREATER_EQUAL "12.9")
|
||||
list(APPEND CMAKE_CUDA_ARCHITECTURES 121a-real)
|
||||
list(APPEND CMAKE_CUDA_ARCHITECTURES 103-real 121a-real)
|
||||
endif()
|
||||
endif()
|
||||
endif()
|
||||
|
||||
@@ -1,31 +1,36 @@
|
||||
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, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 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_Q1_0, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q1_0, 512, 2, 128, 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, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_0, 256, 2, 64, 24, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_0, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_0, 512, 2, 128, 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, 512, 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_Q2_0, 512, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q2_0, 512, 2, 128, 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, 24, 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, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_0, 256, 2, 64, 48, 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);
|
||||
@@ -37,8 +42,11 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
|
||||
|
||||
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, 24, 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, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_1, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_1, 512, 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);
|
||||
@@ -50,43 +58,45 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
|
||||
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, 48, 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, 16, 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, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_0, 512, 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, 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, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q5_1, 512, 2, 128, 48, 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, 16, 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, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q5_1, 512, 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, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q8_0, 512, 2, 128, 48, 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, 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);
|
||||
@@ -95,7 +105,10 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
|
||||
|
||||
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, 24, 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, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q2_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q2_K, 256, 2, 64, 48, 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);
|
||||
@@ -108,18 +121,20 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
|
||||
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, 48, 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, 256, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_Q3_K, 256, 2, 128, 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, 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, 512, 2, 128, 32, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_K, 512, 2, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q4_K, 256, 2, 128, 48, 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);
|
||||
@@ -131,6 +146,7 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
|
||||
|
||||
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, 24, 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);
|
||||
@@ -143,7 +159,11 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
|
||||
|
||||
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, 24, 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, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 48, GGML_CUDA_MMQ_SRAM_LAYOUT_Q6_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_Q6_K, 256, 2, 64, 56, 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);
|
||||
@@ -151,126 +171,151 @@ static constexpr __host__ __device__ ggml_cuda_mmq_config ggml_cuda_mmq_get_conf
|
||||
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, 56, 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, 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, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ1_S, 512, 2, 64, 48, 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, 512, 2, 128, 128, 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, 16, 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, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ1_S, 512, 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_IQ1_S, 512, 2, 128, 128, 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, 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, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 512, 2, 64, 48, 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, 512, 2, 128, 128, 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, 16, 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, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XXS, 512, 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_XXS, 512, 2, 128, 128, 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, 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, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_XS, 512, 2, 64, 48, 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, 512, 2, 128, 128, 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, 16, 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, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_XS, 512, 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_XS, 512, 2, 128, 128, 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, 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, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ2_S, 512, 2, 64, 48, 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, 512, 2, 128, 128, 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, 16, 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, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q3_K, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ2_S, 512, 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_IQ2_S, 512, 2, 128, 128, 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, 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, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 512, 2, 64, 48, 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, 512, 2, 128, 128, 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, 16, 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, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_XXS, 512, 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_XXS, 512, 2, 128, 128, 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, 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, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ3_S, 512, 2, 64, 48, 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, 512, 2, 128, 128, 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, 16, 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, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ3_S, 512, 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_IQ3_S, 512, 2, 128, 128, 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, 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, 64, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_XS, 512, 2, 64, 48, 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, 16, 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, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_XS, 512, 2, 128, 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, 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, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_IQ4_NL, 512, 2, 64, 48, 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, 16, 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, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_0, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_IQ4_NL, 512, 2, 128, 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, 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, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_MXFP4, 512, 2, 64, 48, 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, 16, 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, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_Q8_1, MMQ_ITER_K, false, false);
|
||||
CASE(GGML_TYPE_MXFP4, 512, 2, 128, 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, 512, 2, 128, 16, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 24, 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, 128, 40, GGML_CUDA_MMQ_SRAM_LAYOUT_NVFP4, MMQ_ITER_K, false, true);
|
||||
CASE(GGML_TYPE_NVFP4, 256, 2, 64, 48, 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, 512, 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);
|
||||
|
||||
@@ -316,6 +316,7 @@ struct ggml_hexagon_session {
|
||||
ggml_hexagon_opbatch* op_batch;
|
||||
ggml_hexagon_opqueue* op_queue;
|
||||
|
||||
std::unordered_set<ggml_hexagon_shared_buffer *> buffers;
|
||||
std::unordered_map<int, std::unique_ptr<ggml_hexagon_shared_buffer>> cloned_buffers;
|
||||
std::unordered_set<ggml_hexagon_session *> virt_peers;
|
||||
std::unordered_set<ggml_hexagon_session *> phys_peers;
|
||||
@@ -337,6 +338,7 @@ struct ggml_hexagon_session {
|
||||
ggml_hexagon_mdev_group mdev;
|
||||
ggml_backend_dev_t dev = nullptr;
|
||||
ggml_backend_hexagon_device_context * dev_ctx = nullptr;
|
||||
ggml_hexagon_shared_buffer * shm_buf = nullptr;
|
||||
ggml_hexagon_fence_buffer * fence_buf = nullptr;
|
||||
|
||||
ggml_hexagon_session(const ggml_hexagon_device_config & config, ggml_backend_dev_t dev = nullptr, uint32_t mdev_idx = 0, uint32_t mdev_count = 0) noexcept(false);
|
||||
@@ -373,8 +375,8 @@ struct ggml_hexagon_session {
|
||||
void flush_peers();
|
||||
void flush_pending(bool all = true);
|
||||
|
||||
ggml_hexagon_shared_buffer * mmap_tensor(const ggml_tensor * t);
|
||||
bool clone_buffer(const ggml_hexagon_shared_buffer*);
|
||||
void mmap_buffers();
|
||||
void clone_buffer(const ggml_hexagon_shared_buffer*);
|
||||
void release_buffer(const ggml_hexagon_shared_buffer*);
|
||||
void unclone_buffer(const ggml_hexagon_shared_buffer*);
|
||||
|
||||
@@ -434,7 +436,7 @@ struct ggml_hexagon_rpcmem_block {
|
||||
int fd = -1;
|
||||
size_t size = 0;
|
||||
|
||||
std::unordered_set<ggml_hexagon_session *> mapped_clones;
|
||||
std::unordered_set<ggml_hexagon_session *> clones;
|
||||
|
||||
ggml_hexagon_rpcmem_block(size_t size) {
|
||||
base = (uint8_t *) rpcmem_alloc2(RPCMEM_HEAP_ID_SYSTEM, RPCMEM_DEFAULT_FLAGS, size);
|
||||
@@ -468,16 +470,24 @@ struct ggml_hexagon_shared_buffer {
|
||||
size_t size() const { return mem ? mem->size : 0; }
|
||||
int fd() const { return mem ? mem->fd : -1; }
|
||||
|
||||
void mmap(bool extended = false) {
|
||||
void mmap() {
|
||||
if (!this->mem) return;
|
||||
if (this->mapped) return;
|
||||
|
||||
GGML_ASSERT(!this->pinned || !extended);
|
||||
GGML_ASSERT(!this->pinned || !this->extended);
|
||||
|
||||
this->extended = extended;
|
||||
fastrpc_map_flags flags = this->pinned ? FASTRPC_MAP_FD : (extended ? FASTRPC_MAP_FD_DELAYED_EXTENDED : FASTRPC_MAP_FD_DELAYED);
|
||||
fastrpc_map_flags flags = this->pinned ? FASTRPC_MAP_FD : (this->extended ? FASTRPC_MAP_FD_DELAYED_EXTENDED : FASTRPC_MAP_FD_DELAYED);
|
||||
|
||||
int err = fastrpc_mmap(sess->domain_id, fd(), (void *) base(), 0, size(), flags);
|
||||
if (err != 0) {
|
||||
if (this->extended) {
|
||||
flags = FASTRPC_MAP_FD_EXTENDED;
|
||||
}
|
||||
HEX_VERBOSE("ggml-hex: %s mmap failed (err 0x%08x), retrying in 100ms...\n", sess->c_name(), (unsigned) err);
|
||||
std::this_thread::sleep_for(std::chrono::milliseconds(100));
|
||||
err = fastrpc_mmap(sess->domain_id, fd(), (void *) base(), 0, size(), flags);
|
||||
}
|
||||
|
||||
if (err != 0) {
|
||||
GGML_LOG_ERROR("ggml-hex: %s buffer mapping failed : domain_id %d size %zu fd %d error 0x%08x\n", sess->c_name(),
|
||||
sess->domain_id, size(), fd(), (unsigned) err);
|
||||
@@ -550,10 +560,10 @@ struct ggml_hexagon_shared_buffer {
|
||||
|
||||
// Clone constructor for cross-session mapping
|
||||
ggml_hexagon_shared_buffer(ggml_hexagon_session * sess, const ggml_hexagon_shared_buffer & other) {
|
||||
this->sess = sess;
|
||||
this->mem = other.mem;
|
||||
this->mapped = false;
|
||||
this->pinned = other.pinned;
|
||||
this->sess = sess;
|
||||
this->mem = other.mem;
|
||||
this->mapped = false;
|
||||
this->pinned = false;
|
||||
this->extended = other.extended;
|
||||
}
|
||||
|
||||
@@ -563,19 +573,30 @@ struct ggml_hexagon_shared_buffer {
|
||||
};
|
||||
|
||||
struct ggml_hexagon_fence_buffer : public ggml_hexagon_shared_buffer {
|
||||
uint8_t * data = nullptr;
|
||||
uint32_t slot_count = 0;
|
||||
uint32_t slot_head = 0;
|
||||
std::vector<uint32_t> free_slots;
|
||||
ggml_backend_buffer backend_buffer{};
|
||||
|
||||
ggml_hexagon_fence_buffer(ggml_hexagon_session * sess, ggml_backend_buffer_type_t buft, size_t size)
|
||||
: ggml_hexagon_shared_buffer(sess, size, false /* pinned */),
|
||||
slot_count(size / HTP_FENCE_SLOT_SIZE),
|
||||
static size_t size_of(size_t n_slots) {
|
||||
return n_slots * HTP_FENCE_SLOT_SIZE;
|
||||
}
|
||||
|
||||
ggml_hexagon_fence_buffer(ggml_hexagon_shared_buffer * shm_buf, ggml_backend_buffer_type_t buft, size_t offset, size_t n_slots)
|
||||
: ggml_hexagon_shared_buffer(shm_buf->sess, *shm_buf),
|
||||
data(shm_buf->base() + offset),
|
||||
slot_count(n_slots),
|
||||
slot_head(0) {
|
||||
GGML_ASSERT((uintptr_t) data % HEX_L2_LINE_SIZE == 0);
|
||||
this->mapped = shm_buf->mapped;
|
||||
backend_buffer.buft = buft;
|
||||
backend_buffer.context = static_cast<ggml_hexagon_shared_buffer *>(this);
|
||||
backend_buffer.size = size;
|
||||
mmap(false);
|
||||
backend_buffer.size = size_of(n_slots);
|
||||
}
|
||||
|
||||
~ggml_hexagon_fence_buffer() {
|
||||
this->mapped = false;
|
||||
}
|
||||
|
||||
uint8_t * alloc_slot(uint32_t n_slots = 1) {
|
||||
@@ -583,11 +604,11 @@ struct ggml_hexagon_fence_buffer : public ggml_hexagon_shared_buffer {
|
||||
if (n_slots == 1 && !free_slots.empty()) {
|
||||
uint32_t slot = free_slots.back();
|
||||
free_slots.pop_back();
|
||||
ptr = base() + (size_t) slot * HTP_FENCE_SLOT_SIZE;
|
||||
ptr = data + (size_t) slot * HTP_FENCE_SLOT_SIZE;
|
||||
} else if (slot_head + n_slots <= slot_count) {
|
||||
uint32_t slot = slot_head;
|
||||
slot_head += n_slots;
|
||||
ptr = base() + (size_t) slot * HTP_FENCE_SLOT_SIZE;
|
||||
ptr = data + (size_t) slot * HTP_FENCE_SLOT_SIZE;
|
||||
}
|
||||
if (ptr) {
|
||||
memset(ptr, 0, (size_t) n_slots * HTP_FENCE_SLOT_SIZE);
|
||||
@@ -597,7 +618,7 @@ struct ggml_hexagon_fence_buffer : public ggml_hexagon_shared_buffer {
|
||||
|
||||
void free_slot(void * ptr, uint32_t n_slots = 1) {
|
||||
if (!ptr) return;
|
||||
uint32_t slot = ((uint8_t *) ptr - base()) / HTP_FENCE_SLOT_SIZE;
|
||||
uint32_t slot = ((uint8_t *) ptr - data) / HTP_FENCE_SLOT_SIZE;
|
||||
for (uint32_t i = 0; i < n_slots; i++) {
|
||||
free_slots.push_back(slot + i);
|
||||
}
|
||||
@@ -636,6 +657,7 @@ inline ggml_tensor ggml_hexagon_session::make_fence_tensor(const volatile void *
|
||||
|
||||
static void ggml_backend_hexagon_buffer_free_buffer(ggml_backend_buffer_t buffer) {
|
||||
auto sbuf = static_cast<ggml_hexagon_shared_buffer *>(buffer->context);
|
||||
sbuf->sess->buffers.erase(sbuf);
|
||||
sbuf->sess->unclone_buffer(sbuf);
|
||||
delete sbuf;
|
||||
}
|
||||
@@ -975,31 +997,44 @@ static bool ggml_backend_hexagon_host_buffer_type_is_host(ggml_backend_buffer_ty
|
||||
GGML_UNUSED(buft);
|
||||
}
|
||||
|
||||
struct ggml_backend_hexagon_alloc_buffer_n_plan_item {
|
||||
size_t size;
|
||||
int first;
|
||||
int last;
|
||||
};
|
||||
|
||||
using ggml_backend_hexagon_alloc_buffer_n_plan_t = std::vector<ggml_backend_hexagon_alloc_buffer_n_plan_item>;
|
||||
|
||||
static const char * ggml_hexagon_kv_layer_suffix(const struct ggml_tensor * t) {
|
||||
if (strncmp(t->name, "cache_", 6) != 0) {
|
||||
return NULL;
|
||||
}
|
||||
const char * p = strstr(t->name, "_l");
|
||||
if (!p || !isdigit((unsigned char)p[2])) {
|
||||
return NULL;
|
||||
}
|
||||
return p;
|
||||
}
|
||||
|
||||
struct ggml_backend_hexagon_alloc_unit {
|
||||
size_t size;
|
||||
int first;
|
||||
int last;
|
||||
};
|
||||
|
||||
struct ggml_backend_hexagon_alloc_buffer_n_plan_item {
|
||||
size_t size;
|
||||
std::vector<ggml_backend_hexagon_alloc_unit> units;
|
||||
};
|
||||
|
||||
using ggml_backend_hexagon_alloc_buffer_n_plan_t = std::vector<ggml_backend_hexagon_alloc_buffer_n_plan_item>;
|
||||
|
||||
static bool ggml_hexagon_tensor_layer(const struct ggml_tensor * t, int & kind, int & layer) {
|
||||
if (strncmp(t->name, "blk.", 4) == 0 && isdigit((unsigned char) t->name[4])) {
|
||||
char * end = NULL;
|
||||
long l = strtol(t->name + 4, &end, 10);
|
||||
if (end && *end == '.') {
|
||||
kind = 0;
|
||||
layer = (int) l;
|
||||
return true;
|
||||
}
|
||||
}
|
||||
if (strncmp(t->name, "cache_", 6) == 0) {
|
||||
const char * p = strstr(t->name, "_l");
|
||||
if (p && isdigit((unsigned char) p[2])) {
|
||||
char * end = NULL;
|
||||
long l = strtol(p + 2, &end, 10);
|
||||
if (end && *end == '\0') {
|
||||
kind = 1;
|
||||
layer = (int) l;
|
||||
return true;
|
||||
}
|
||||
}
|
||||
}
|
||||
return false;
|
||||
}
|
||||
|
||||
static ggml_backend_hexagon_alloc_buffer_n_plan_t ggml_backend_hexagon_alloc_buffer_n_plan(
|
||||
ggml_backend_buffer_type_t buft, struct ggml_tensor ** tensors, int n_tensors) {
|
||||
ggml_backend_hexagon_alloc_buffer_n_plan_t plan;
|
||||
@@ -1012,61 +1047,72 @@ static ggml_backend_hexagon_alloc_buffer_n_plan_t ggml_backend_hexagon_alloc_buf
|
||||
int i = 0;
|
||||
while (i < n_tensors) {
|
||||
struct ggml_tensor * t = tensors[i];
|
||||
size_t unit_size = 0;
|
||||
int unit_first = i;
|
||||
int unit_last = i + 1;
|
||||
|
||||
if (t->data == NULL && t->view_src == NULL) {
|
||||
unit_size += GGML_PAD(ggml_backend_buft_get_alloc_size(buft, t), alignment);
|
||||
if (t->data != NULL || t->view_src != NULL) {
|
||||
i++;
|
||||
continue;
|
||||
}
|
||||
|
||||
const char * layer_suffix = ggml_hexagon_kv_layer_suffix(t);
|
||||
size_t unit_size = GGML_PAD(ggml_backend_buft_get_alloc_size(buft, t), alignment);
|
||||
int unit_first = i;
|
||||
int unit_last = i + 1;
|
||||
|
||||
while (unit_last < n_tensors) {
|
||||
int kind = -1, layer = -1;
|
||||
const bool has_layer = ggml_hexagon_tensor_layer(t, kind, layer);
|
||||
|
||||
while (has_layer && unit_last < n_tensors) {
|
||||
struct ggml_tensor * next = tensors[unit_last];
|
||||
|
||||
if (next->view_src != NULL) {
|
||||
if (next->data != NULL || next->view_src != NULL) {
|
||||
unit_last++;
|
||||
continue;
|
||||
}
|
||||
|
||||
if (layer_suffix != NULL) {
|
||||
const char * next_suffix = ggml_hexagon_kv_layer_suffix(next);
|
||||
if (next_suffix != NULL && strcmp(layer_suffix, next_suffix) == 0) {
|
||||
if (next->data == NULL) {
|
||||
unit_size += GGML_PAD(ggml_backend_buft_get_alloc_size(buft, next), alignment);
|
||||
}
|
||||
unit_last++;
|
||||
continue;
|
||||
}
|
||||
int next_kind = -1, next_layer = -1;
|
||||
if (!ggml_hexagon_tensor_layer(next, next_kind, next_layer) ||
|
||||
next_kind != kind || next_layer != layer) {
|
||||
break;
|
||||
}
|
||||
|
||||
break;
|
||||
size_t next_size = GGML_PAD(ggml_backend_buft_get_alloc_size(buft, next), alignment);
|
||||
if (unit_size + next_size > max_size) {
|
||||
break;
|
||||
}
|
||||
|
||||
unit_size += next_size;
|
||||
unit_last++;
|
||||
}
|
||||
|
||||
units.push_back({ unit_size, unit_first, unit_last });
|
||||
i = unit_last;
|
||||
}
|
||||
|
||||
size_t cur_buf_size = 0;
|
||||
int cur_buf_first = 0;
|
||||
|
||||
for (const auto & unit : units) {
|
||||
if (unit.size == 0) {
|
||||
continue;
|
||||
}
|
||||
|
||||
if (cur_buf_size > 0 && (cur_buf_size + unit.size) > max_size) {
|
||||
plan.push_back({ cur_buf_size, cur_buf_first, unit.first });
|
||||
cur_buf_size = 0;
|
||||
cur_buf_first = unit.first;
|
||||
bool placed = false;
|
||||
for (auto & item : plan) {
|
||||
if (item.size + unit.size <= max_size) {
|
||||
item.size += unit.size;
|
||||
item.units.push_back(unit);
|
||||
placed = true;
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (!placed) {
|
||||
plan.push_back({ unit.size, { unit } });
|
||||
}
|
||||
|
||||
cur_buf_size += unit.size;
|
||||
}
|
||||
|
||||
if (cur_buf_size > 0) {
|
||||
plan.push_back({ cur_buf_size, cur_buf_first, n_tensors });
|
||||
std::stable_sort(plan.begin(), plan.end(), [](const auto & a, const auto & b) {
|
||||
return a.size > b.size;
|
||||
});
|
||||
|
||||
if (!plan.empty()) {
|
||||
for (const auto & unit : units) {
|
||||
if (unit.size == 0) {
|
||||
plan[0].units.push_back(unit);
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
return plan;
|
||||
@@ -1092,28 +1138,19 @@ static ggml_backend_buffer_t ggml_backend_hexagon_buffer_type_alloc_buffer_n(
|
||||
struct ggml_tallocr tallocr = ggml_tallocr_new(buffer);
|
||||
|
||||
struct ggml_tensor * t_failed = NULL;
|
||||
for (int j = item.first; j < item.last; j++) {
|
||||
struct ggml_tensor * t = tensors[j];
|
||||
if (t->data == NULL) {
|
||||
if (t->view_src == NULL) {
|
||||
for (const auto & unit : item.units) {
|
||||
for (int j = unit.first; j < unit.last; j++) {
|
||||
struct ggml_tensor * t = tensors[j];
|
||||
if (t->data == NULL && t->view_src == NULL) {
|
||||
if (ggml_tallocr_alloc(&tallocr, t) != GGML_STATUS_SUCCESS) {
|
||||
t_failed = t;
|
||||
break;
|
||||
}
|
||||
} else if (t->buffer == NULL) {
|
||||
if (ggml_backend_view_init(t) != GGML_STATUS_SUCCESS) {
|
||||
t_failed = t;
|
||||
break;
|
||||
}
|
||||
}
|
||||
} else {
|
||||
if (t->view_src != NULL && t->buffer == NULL) {
|
||||
if (ggml_backend_view_init(t) != GGML_STATUS_SUCCESS) {
|
||||
t_failed = t;
|
||||
break;
|
||||
}
|
||||
}
|
||||
}
|
||||
if (t_failed != NULL) {
|
||||
break;
|
||||
}
|
||||
}
|
||||
if (t_failed != NULL) {
|
||||
GGML_LOG_ERROR("%s: failed to initialize tensor %s\n", __func__, t_failed->name);
|
||||
@@ -1131,6 +1168,19 @@ static ggml_backend_buffer_t ggml_backend_hexagon_buffer_type_alloc_buffer_n(
|
||||
return NULL;
|
||||
}
|
||||
|
||||
for (int j = 0; j < n_tensors; j++) {
|
||||
struct ggml_tensor * t = tensors[j];
|
||||
if (t->view_src != NULL && t->buffer == NULL) {
|
||||
if (ggml_backend_view_init(t) != GGML_STATUS_SUCCESS) {
|
||||
GGML_LOG_ERROR("%s: failed to initialize tensor %s\n", __func__, t->name);
|
||||
for (ggml_backend_buffer_t b : buffers) {
|
||||
ggml_backend_buffer_free(b);
|
||||
}
|
||||
return NULL;
|
||||
}
|
||||
}
|
||||
}
|
||||
|
||||
if (buffers.size() == 1) {
|
||||
return buffers[0];
|
||||
}
|
||||
@@ -1200,7 +1250,6 @@ struct ggml_hexagon_opbatch {
|
||||
ggml_hexagon_session* sess;
|
||||
|
||||
std::vector<htp_opnode> ops; // htp_opnode of ops
|
||||
|
||||
std::vector<htp_buf_desc> h_bufs; // htp buffer descriptors
|
||||
std::vector<htp_tensor> h_tens; // htp tensor descriptors
|
||||
std::vector<htp_op_desc> h_ops; // htp op descriptors
|
||||
@@ -1529,45 +1578,45 @@ struct ggml_hexagon_registry {
|
||||
struct ggml_hexagon_opqueue {
|
||||
// Shared buffer for storing batches
|
||||
ggml_hexagon_shared_buffer *shm_buf;
|
||||
size_t shm_blk_size;
|
||||
uint8_t * data;
|
||||
size_t slot_size;
|
||||
size_t depth;
|
||||
|
||||
using opvec = std::vector<htp_opnode>;
|
||||
|
||||
std::vector<opvec> op_cache; // per batch op cache
|
||||
|
||||
ggml_hexagon_opqueue(ggml_hexagon_session *sess, size_t batch_size, size_t depth) : depth(depth) {
|
||||
static size_t size_of(size_t batch_size, size_t n_threads, size_t depth = 1) {
|
||||
size_t n_bufs = HTP_OP_MAX_BUFS;
|
||||
size_t n_ops = batch_size;
|
||||
size_t n_tensors = std::min<size_t>(n_ops * 4 + HTP_OP_MAX_INPUTS + HTP_OP_MAX_OUTPUTS, HTP_OP_MAX_TENSORS);
|
||||
|
||||
size_t tr_size = 0;
|
||||
if (opt_profile == 3) {
|
||||
tr_size = sess->n_threads * opt_optrace * sizeof(htp_trace_desc);
|
||||
tr_size = n_threads * opt_optrace * sizeof(htp_trace_desc);
|
||||
}
|
||||
|
||||
shm_blk_size = sizeof(htp_buf_desc) * n_bufs +
|
||||
sizeof(htp_tensor) * n_tensors +
|
||||
sizeof(htp_op_desc) * n_ops +
|
||||
sizeof(htp_prof_desc) * n_ops +
|
||||
tr_size;
|
||||
|
||||
shm_buf = new ggml_hexagon_shared_buffer(sess, shm_blk_size * depth, true /* pinned */);
|
||||
size_t slot_size = sizeof(htp_buf_desc) * n_bufs +
|
||||
sizeof(htp_tensor) * n_tensors +
|
||||
sizeof(htp_op_desc) * n_ops +
|
||||
sizeof(htp_prof_desc) * n_ops +
|
||||
tr_size;
|
||||
return slot_size * depth;
|
||||
}
|
||||
|
||||
ggml_hexagon_opqueue(ggml_hexagon_shared_buffer *shm_buf, size_t offset, size_t batch_size, size_t depth)
|
||||
: shm_buf(shm_buf),
|
||||
data(shm_buf->base() + offset),
|
||||
slot_size(size_of(batch_size, shm_buf->sess->n_threads)),
|
||||
depth(depth) {
|
||||
op_cache.resize(depth);
|
||||
|
||||
if (opt_verbose) {
|
||||
GGML_LOG_INFO("ggml-hex: %s allocated opqueue : batch-size %zu depth %zu shm-size %zu shm-block-size %zu\n",
|
||||
sess->c_name(), batch_size, depth, shm_buf->size(), shm_blk_size);
|
||||
GGML_LOG_INFO("ggml-hex: %s allocated opqueue : batch-size %zu depth %zu shm-size %zu slot-size %zu\n",
|
||||
shm_buf->c_name(), batch_size, depth, size_of(batch_size, shm_buf->sess->n_threads, depth), slot_size);
|
||||
}
|
||||
}
|
||||
|
||||
~ggml_hexagon_opqueue() {
|
||||
delete shm_buf;
|
||||
}
|
||||
|
||||
size_t shm_size() const { return shm_buf ? shm_buf->size() : 0; }
|
||||
|
||||
// push new batch
|
||||
bool push(htp_opbatch_req& req, dspqueue_buffer& dbuf, const ggml_hexagon_opbatch* op_batch, uint64_t seq) {
|
||||
static_assert(sizeof(htp_opbatch_req) % 8 == 0, "sizeof(htp_opbatch_req) must be multiple of 8");
|
||||
@@ -1603,13 +1652,13 @@ struct ggml_hexagon_opqueue {
|
||||
req.n_traces = 0;
|
||||
}
|
||||
|
||||
dbuf.ptr = shm_buf->base() + ((size_t) slot * shm_blk_size);
|
||||
dbuf.ptr = data + ((size_t) slot * slot_size);
|
||||
dbuf.fd = shm_buf->fd();
|
||||
dbuf.flags = DSPQUEUE_BUFFER_FLAG_FLUSH_SENDER | DSPQUEUE_BUFFER_FLAG_INVALIDATE_RECIPIENT;
|
||||
dbuf.offset = (uint8_t*) dbuf.ptr - (uint8_t*) shm_buf->base();
|
||||
dbuf.size = b_size + t_size + o_size + p_size + tr_size;
|
||||
|
||||
GGML_ASSERT(dbuf.size <= shm_blk_size);
|
||||
GGML_ASSERT(dbuf.size <= slot_size);
|
||||
|
||||
uint8_t * m_ptr = (uint8_t*) dbuf.ptr;
|
||||
uint8_t * b_ptr = m_ptr; m_ptr += b_size;
|
||||
@@ -1657,7 +1706,7 @@ struct ggml_hexagon_opqueue {
|
||||
}
|
||||
|
||||
const size_t m_size = b_size + t_size + o_size + p_size + tr_size;
|
||||
GGML_ASSERT(m_size <= shm_blk_size);
|
||||
GGML_ASSERT(m_size <= slot_size);
|
||||
|
||||
HEX_VERBOSE("ggml-hex: %s opqueue-pop batch #%llu : n-bufs %u n-tensors %u n-ops %u : m-size %zu b-size %zu t-size %zu o-size %zu\n",
|
||||
shm_buf->sess->c_name(), (unsigned long long) rsp.seq, rsp.n_bufs, rsp.n_tensors, rsp.n_ops,
|
||||
@@ -1775,6 +1824,7 @@ void ggml_hexagon_session::flush_batch(size_t min_ops) {
|
||||
if (op_batch->n_ops < min_ops) { return; }
|
||||
|
||||
op_batch->sort_buffers();
|
||||
this->mmap_buffers();
|
||||
|
||||
htp_opbatch_req req {};
|
||||
dspqueue_buffer dbuf{};
|
||||
@@ -1819,10 +1869,17 @@ void ggml_hexagon_session::flush_batch(size_t min_ops) {
|
||||
}
|
||||
|
||||
void ggml_hexagon_session::enqueue_op(const htp_opnode & node) {
|
||||
auto clone_tensor_buffer = [this](const ggml_tensor * t) {
|
||||
auto sbuf = this->mmap_tensor(t);
|
||||
if (!sbuf) return;
|
||||
if (sbuf->sess != this) {
|
||||
auto register_tensor = [&](const ggml_tensor * t) {
|
||||
if (!t) return;
|
||||
auto * sbuf = static_cast<ggml_hexagon_shared_buffer *>(t->buffer->context);
|
||||
if (ggml_backend_buffer_get_usage(t->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS) {
|
||||
sbuf->extended = opt_dma64;
|
||||
}
|
||||
if (sbuf->sess == this) {
|
||||
if (!sbuf->mapped) {
|
||||
this->buffers.insert(sbuf);
|
||||
}
|
||||
} else {
|
||||
this->clone_buffer(sbuf);
|
||||
}
|
||||
for (auto & sub : this->mdev.sessions) {
|
||||
@@ -1830,11 +1887,11 @@ void ggml_hexagon_session::enqueue_op(const htp_opnode & node) {
|
||||
}
|
||||
};
|
||||
|
||||
for (auto t : node.get_inputs()) {
|
||||
clone_tensor_buffer(t);
|
||||
for (const auto * src : node.get_inputs()) {
|
||||
register_tensor(src);
|
||||
}
|
||||
for (auto t : node.get_outputs()) {
|
||||
clone_tensor_buffer(t);
|
||||
for (const auto * dst : node.get_outputs()) {
|
||||
register_tensor(dst);
|
||||
}
|
||||
|
||||
const bool fits_tensors = op_batch->fit_tensors(node);
|
||||
@@ -1945,46 +2002,61 @@ void ggml_hexagon_session::enqueue_allreduce(
|
||||
this->enqueue_op(ar_node);
|
||||
}
|
||||
|
||||
ggml_hexagon_shared_buffer * ggml_hexagon_session::mmap_tensor(const ggml_tensor * t) {
|
||||
if (!t) return nullptr;
|
||||
void ggml_hexagon_session::mmap_buffers() {
|
||||
auto collect_session = [](ggml_hexagon_session * sess, std::unordered_set<ggml_hexagon_shared_buffer *> & unmapped) {
|
||||
for (auto * sbuf : sess->buffers) {
|
||||
if (!sbuf->mapped) {
|
||||
unmapped.insert(sbuf);
|
||||
}
|
||||
}
|
||||
for (auto & [fd, clone] : sess->cloned_buffers) {
|
||||
if (!clone->mapped) {
|
||||
unmapped.insert(clone.get());
|
||||
}
|
||||
}
|
||||
};
|
||||
|
||||
auto sbuf = static_cast<ggml_hexagon_shared_buffer *>(t->buffer->context);
|
||||
if (!sbuf->mapped) {
|
||||
const bool is_weight = ggml_backend_buffer_get_usage(t->buffer) == GGML_BACKEND_BUFFER_USAGE_WEIGHTS;
|
||||
const bool extended = opt_dma64 && is_weight;
|
||||
sbuf->mmap(extended);
|
||||
std::unordered_set<ggml_hexagon_shared_buffer *> unmapped;
|
||||
collect_session(this, unmapped);
|
||||
for (auto & sub : this->mdev.sessions) {
|
||||
collect_session(sub.get(), unmapped);
|
||||
}
|
||||
|
||||
if (!unmapped.empty()) {
|
||||
std::vector<ggml_hexagon_shared_buffer *> sorted(unmapped.begin(), unmapped.end());
|
||||
std::sort(sorted.begin(), sorted.end(), [](const auto * a, const auto * b) {
|
||||
if (a->size() != b->size()) {
|
||||
return a->size() > b->size();
|
||||
}
|
||||
if (a->fd() != b->fd()) {
|
||||
return a->fd() < b->fd();
|
||||
}
|
||||
return a->sess->domain_id < b->sess->domain_id;
|
||||
});
|
||||
for (auto * sbuf : sorted) {
|
||||
sbuf->mmap();
|
||||
}
|
||||
}
|
||||
return sbuf;
|
||||
}
|
||||
|
||||
bool ggml_hexagon_session::clone_buffer(const ggml_hexagon_shared_buffer * sbuf)
|
||||
void ggml_hexagon_session::clone_buffer(const ggml_hexagon_shared_buffer * sbuf)
|
||||
{
|
||||
GGML_ASSERT(sbuf && sbuf->mem);
|
||||
if (sbuf->sess == this) return true;
|
||||
if (sbuf->sess == this) return;
|
||||
|
||||
auto mem = sbuf->mem;
|
||||
int fd = mem->fd;
|
||||
|
||||
GGML_ASSERT(fd >= 0);
|
||||
|
||||
if (this->cloned_buffers.find(fd) != this->cloned_buffers.end()) return true;
|
||||
|
||||
GGML_ASSERT(sbuf->mapped);
|
||||
if (this->cloned_buffers.find(fd) != this->cloned_buffers.end()) return;
|
||||
|
||||
HEX_VERBOSE("ggml-hex: %s clone-buffer: %s base %p size %zu fd %d\n", this->name.c_str(),
|
||||
sbuf->c_name(), sbuf->base(), sbuf->size(), fd);
|
||||
|
||||
auto clone = std::make_unique<ggml_hexagon_shared_buffer>(this, *sbuf);
|
||||
try {
|
||||
clone->mmap(sbuf->extended);
|
||||
} catch (const std::exception & exc) {
|
||||
GGML_LOG_ERROR("ggml-hex: %s lazy mapping of buffer context failed: %s\n", this->c_name(), exc.what());
|
||||
return false;
|
||||
}
|
||||
|
||||
this->cloned_buffers[fd] = std::move(clone);
|
||||
mem->mapped_clones.insert(this);
|
||||
return true;
|
||||
mem->clones.insert(this);
|
||||
}
|
||||
|
||||
void ggml_hexagon_session::release_buffer(const ggml_hexagon_shared_buffer * sbuf) {
|
||||
@@ -2000,14 +2072,14 @@ void ggml_hexagon_session::release_buffer(const ggml_hexagon_shared_buffer * sbu
|
||||
auto clone = std::move(it->second);
|
||||
this->cloned_buffers.erase(it);
|
||||
}
|
||||
mem->mapped_clones.erase(this);
|
||||
mem->clones.erase(this);
|
||||
}
|
||||
|
||||
void ggml_hexagon_session::unclone_buffer(const ggml_hexagon_shared_buffer * sbuf) {
|
||||
GGML_ASSERT(sbuf && sbuf->mem);
|
||||
|
||||
auto mem = sbuf->mem;
|
||||
std::vector<ggml_hexagon_session *> sessions(mem->mapped_clones.begin(), mem->mapped_clones.end());
|
||||
std::vector<ggml_hexagon_session *> sessions(mem->clones.begin(), mem->clones.end());
|
||||
|
||||
for (auto * sess : sessions) {
|
||||
sess->release_buffer(sbuf);
|
||||
@@ -2276,9 +2348,14 @@ void ggml_hexagon_session::allocate(const ggml_hexagon_device_config & config) n
|
||||
}
|
||||
|
||||
// Allocate buffers and state for op batching
|
||||
this->op_queue = new ggml_hexagon_opqueue(this, opt_opbatch, opt_opqueue);
|
||||
const size_t n_fences = 512;
|
||||
const size_t opq_size = ggml_hexagon_opqueue::size_of(opt_opbatch, this->n_threads, opt_opqueue);
|
||||
const size_t fence_offset = GGML_PAD(opq_size, HEX_L2_LINE_SIZE);
|
||||
const size_t fence_size = ggml_hexagon_fence_buffer::size_of(n_fences);
|
||||
|
||||
this->fence_buf = new ggml_hexagon_fence_buffer(this, &dev_ctx->fence_buffer_type, 64 * 1024);
|
||||
this->shm_buf = new ggml_hexagon_shared_buffer(this, fence_offset + fence_size, true /* pinned */);
|
||||
this->op_queue = new ggml_hexagon_opqueue(this->shm_buf, 0, opt_opbatch, opt_opqueue);
|
||||
this->fence_buf = new ggml_hexagon_fence_buffer(this->shm_buf, &dev_ctx->fence_buffer_type, fence_offset, n_fences);
|
||||
if (this->mdev.count > 1) {
|
||||
this->mdev_fence_slot = this->alloc_fence(this->mdev.count);
|
||||
}
|
||||
@@ -2287,7 +2364,7 @@ void ggml_hexagon_session::allocate(const ggml_hexagon_device_config & config) n
|
||||
opt_vmem = ggml_hexagon_measure_max_vmem(this);
|
||||
GGML_LOG_INFO("ggml-hex: %s measured max vmem %zu\n", this->c_name(), opt_vmem);
|
||||
}
|
||||
const size_t shm_size = this->op_queue->shm_size();
|
||||
const size_t shm_size = this->shm_buf->size();
|
||||
this->max_vmem = (opt_vmem > shm_size) ? (opt_vmem - shm_size) : opt_vmem;
|
||||
|
||||
this->op_batch = new ggml_hexagon_opbatch(this, opt_opbatch, this->max_vmem);
|
||||
@@ -2328,8 +2405,6 @@ void ggml_hexagon_session::release() noexcept(true) {
|
||||
}
|
||||
}
|
||||
|
||||
delete this->op_batch;
|
||||
delete this->op_queue;
|
||||
for (auto & it : this->cpy_fence_slots) {
|
||||
free_fence((void *) it.second, 1);
|
||||
}
|
||||
@@ -2340,6 +2415,14 @@ void ggml_hexagon_session::release() noexcept(true) {
|
||||
delete this->fence_buf;
|
||||
this->fence_buf = nullptr;
|
||||
}
|
||||
|
||||
delete this->op_batch;
|
||||
delete this->op_queue;
|
||||
delete this->shm_buf;
|
||||
this->op_batch = nullptr;
|
||||
this->op_queue = nullptr;
|
||||
this->shm_buf = nullptr;
|
||||
|
||||
while (!this->cloned_buffers.empty()) {
|
||||
release_buffer(this->cloned_buffers.begin()->second.get());
|
||||
}
|
||||
@@ -2734,19 +2817,12 @@ static bool ggml_hexagon_cpy_tensor_async_phys(ggml_backend_t backend_src, ggml_
|
||||
auto sess_src = static_cast<ggml_hexagon_session *>(backend_src->context);
|
||||
auto sess_dst = static_cast<ggml_hexagon_session *>(backend_dst->context);
|
||||
|
||||
sess_src->mmap_tensor(src);
|
||||
auto sbuf_dst = sess_dst->mmap_tensor(dst);
|
||||
|
||||
if (!sess_src->clone_buffer(sbuf_dst)) { return false; }
|
||||
|
||||
const uint64_t src_key = ggml_hexagon_session_key(sess_src);
|
||||
auto & fence_slot = sess_dst->cpy_fence_slots[src_key];
|
||||
if (!fence_slot) {
|
||||
fence_slot = (volatile uint32_t *) sess_dst->alloc_fence(1);
|
||||
}
|
||||
|
||||
if (!sess_src->clone_buffer(sess_dst->fence_buf)) { return false; }
|
||||
|
||||
uint32_t fence_seq = sess_dst->next_fence_seq();
|
||||
|
||||
HEX_VERBOSE("ggml-hex: %s cpy-tensor-async %s -> %s size %zu : seq 0x%x\n",
|
||||
@@ -2766,11 +2842,6 @@ static bool ggml_hexagon_cpy_tensor_async_virt(ggml_backend_t backend_src, ggml_
|
||||
auto sess_src = static_cast<ggml_hexagon_session *>(backend_src->context);
|
||||
auto sess_dst = static_cast<ggml_hexagon_session *>(backend_dst->context);
|
||||
|
||||
auto sbuf_src = sess_src->mmap_tensor(src);
|
||||
sess_dst->mmap_tensor(dst);
|
||||
|
||||
if (!sess_dst->clone_buffer(sbuf_src)) { return false; }
|
||||
|
||||
HEX_VERBOSE("ggml-hex: %s cpy-tensor-async %s -> %s size %zu\n",
|
||||
sess_dst->name.c_str(), src->name, dst->name, ggml_nbytes(src));
|
||||
|
||||
@@ -2893,7 +2964,6 @@ static void ggml_backend_hexagon_event_wait(ggml_backend_t backend, ggml_backend
|
||||
return;
|
||||
}
|
||||
|
||||
sess->clone_buffer(hex_event->fence_sess->fence_buf);
|
||||
sess->add_peer(hex_event->sess);
|
||||
sess->enqueue_fence(&hex_event->fence_tensor, hex_event->seq, /* wait = */ true);
|
||||
}
|
||||
|
||||
@@ -82,6 +82,7 @@ struct htp_ops_context {
|
||||
struct htp_context {
|
||||
dspqueue_t dsp_queue;
|
||||
|
||||
struct htp_mmap shm;
|
||||
struct htp_mmap mmap[HTP_MAX_MMAPS];
|
||||
dma_queue_t dma[HTP_MAX_NTHREADS];
|
||||
struct htp_thread_trace trace[HTP_MAX_NTHREADS];
|
||||
|
||||
@@ -1089,7 +1089,12 @@ static int execute_op(struct htp_ops_context * octx) {
|
||||
}
|
||||
|
||||
static inline bool reuse_buf(struct htp_context *ctx, uint32_t *m_reuse, struct htp_buf_desc *b) {
|
||||
b->base = NULL;
|
||||
if (b->fd == ctx->shm.fd) {
|
||||
b->base = ctx->shm.base;
|
||||
return true;
|
||||
}
|
||||
|
||||
b->base = 0;
|
||||
|
||||
for (uint32_t i=0; i<HTP_MAX_MMAPS; i++) {
|
||||
struct htp_mmap *m = ctx->mmap + i;
|
||||
@@ -1203,7 +1208,7 @@ static void prep_op_bufs(struct htp_context *ctx, struct htp_buf_desc *bufs, uin
|
||||
|
||||
for (uint32_t i=0; i < n_bufs; i++) {
|
||||
struct htp_buf_desc *b = bufs + i;
|
||||
if (!(b->flags & HTP_BUF_EXTENDED)) {
|
||||
if (!(b->flags & HTP_BUF_EXTENDED) && b->fd != ctx->shm.fd) {
|
||||
b->base = 0;
|
||||
}
|
||||
if (!mmap_buf(ctx, b)) {
|
||||
@@ -1316,6 +1321,9 @@ static void process_opbatch(struct htp_context * ctx, const struct htp_opbatch_r
|
||||
n_bufs, n_tens, n_ops, req->n_traces, dbuf->size, b_size, t_size, o_size);
|
||||
|
||||
// Setup descriptor pointers
|
||||
ctx->shm.fd = dbuf->fd;
|
||||
ctx->shm.base = (uintptr_t) dbuf->ptr - dbuf->offset;
|
||||
|
||||
uint8_t * m_ptr = dbuf->ptr;
|
||||
struct htp_buf_desc* bufs = (struct htp_buf_desc*) m_ptr; m_ptr += b_size;
|
||||
struct htp_tensor* tens = (struct htp_tensor*) m_ptr; m_ptr += t_size;
|
||||
|
||||
+7
-4
@@ -6,7 +6,7 @@ version = "0.0.0"
|
||||
dynamic = ["classifiers"]
|
||||
readme = "README.md"
|
||||
authors = [{name = "GGML", email = "ggml@ggml.ai"}]
|
||||
requires-python = '>=3.10,<3.15'
|
||||
requires-python = '>=3.10,<3.16'
|
||||
dependencies = [
|
||||
'numpy (>=2.2.6,<3.0.0); python_version < "3.14"',
|
||||
'numpy (>=2.3.0,<3.0.0); python_version >= "3.14"',
|
||||
@@ -40,9 +40,12 @@ packages = [
|
||||
|
||||
[tool.poetry.dependencies]
|
||||
torch = [
|
||||
{ version = "==2.11.0", source = "pypi", markers = "sys_platform == 'darwin'" },
|
||||
{ version = "==2.11.0+cpu", source = "pytorch", markers = "sys_platform == 'linux'" },
|
||||
{ version = "==2.11.0", source = "pypi", markers = "sys_platform == 'win32'" },
|
||||
{ version = "==2.11.0", source = "pypi", markers = "sys_platform == 'darwin' and python_version < '3.15'" },
|
||||
{ version = "==2.11.0+cpu", source = "pytorch", markers = "sys_platform == 'linux' and python_version < '3.15'" },
|
||||
{ version = "==2.11.0", source = "pypi", markers = "sys_platform == 'win32' and python_version < '3.15'" },
|
||||
{ version = "==2.15.0", source = "pypi", markers = "sys_platform == 'darwin' and python_version >= '3.15'" },
|
||||
{ version = "==2.15.0+cpu", source = "pytorch", markers = "sys_platform == 'linux' and python_version >= '3.15'" },
|
||||
{ version = "==2.15.0", source = "pypi", markers = "sys_platform == 'win32' and python_version >= '3.15'" },
|
||||
]
|
||||
|
||||
[tool.poetry.group.dev.dependencies]
|
||||
|
||||
@@ -1,3 +1,4 @@
|
||||
tabulate~=0.9.0
|
||||
GitPython~=3.1.43
|
||||
matplotlib~=3.10.0
|
||||
matplotlib~=3.10.0; python_version < "3.11"
|
||||
matplotlib~=3.11.2; python_version >= "3.11"
|
||||
|
||||
@@ -2,4 +2,5 @@
|
||||
--extra-index-url https://download.pytorch.org/whl/cpu
|
||||
|
||||
## Embedding Gemma requires PyTorch 2.6.0 or later, bumped to 2.11.0 for compatibility
|
||||
torch==2.11.0
|
||||
torch==2.11.0; python_version < "3.15"
|
||||
torch==2.15.0; python_version >= "3.15"
|
||||
|
||||
@@ -1,5 +1,5 @@
|
||||
numpy~=2.2.6; python_version < "3.14"
|
||||
numpy~=2.4.6; python_version >= "3.14"
|
||||
numpy~=2.5.4; python_version >= "3.14"
|
||||
sentencepiece>=0.1.98,<0.3.0
|
||||
|
||||
transformers==4.57.6
|
||||
|
||||
@@ -1,4 +1,4 @@
|
||||
numpy~=2.2.6; python_version < "3.14"
|
||||
numpy~=2.4.6; python_version >= "3.14"
|
||||
numpy~=2.5.4; python_version >= "3.14"
|
||||
PySide6~=6.9.0
|
||||
gguf>=0.17.0
|
||||
|
||||
@@ -1,6 +1,7 @@
|
||||
datasets~=4.8.0
|
||||
matplotlib~=3.10.0
|
||||
matplotlib~=3.10.0; python_version < "3.11"
|
||||
matplotlib~=3.11.2; python_version >= "3.11"
|
||||
numpy~=2.2.6; python_version < "3.14"
|
||||
numpy~=2.4.6; python_version >= "3.14"
|
||||
numpy~=2.5.4; python_version >= "3.14"
|
||||
requests~=2.32.3
|
||||
tqdm~=4.67.1
|
||||
|
||||
@@ -1 +1 @@
|
||||
cffi~=1.16.0
|
||||
cffi~=2.1.1
|
||||
|
||||
@@ -1,9 +1,11 @@
|
||||
pytest~=8.3.3
|
||||
matplotlib~=3.10.0
|
||||
matplotlib~=3.10.0; python_version < "3.11"
|
||||
matplotlib~=3.11.2; python_version >= "3.11"
|
||||
numpy~=2.2.6; python_version < "3.14"
|
||||
numpy~=2.4.6; python_version >= "3.14"
|
||||
numpy~=2.5.4; python_version >= "3.14"
|
||||
openai~=2.14.0
|
||||
pandas~=2.3.3
|
||||
pandas~=2.3.3; python_version < "3.11"
|
||||
pandas~=3.0.6; python_version >= "3.11"
|
||||
prometheus-client~=0.20.0
|
||||
requests~=2.32.3
|
||||
wget~=3.2
|
||||
|
||||
@@ -68,8 +68,9 @@ log = logging.getLogger(__name__)
|
||||
|
||||
POLL_INTERVAL = 30
|
||||
JOB_TIMEOUT = 3600
|
||||
RUNNER_TIMEOUT = 300
|
||||
LOG_UPLOAD_TIMEOUT = 600
|
||||
CAPACITY_TIMEOUT = 1800
|
||||
CAPACITY_TIMEOUT = 300
|
||||
CAPACITY_POLL = 60
|
||||
MAX_CONCURRENT_JOBS = 5
|
||||
DEFAULT_RETRIES = 0
|
||||
@@ -82,6 +83,10 @@ class DeviceUnavailableError(Exception):
|
||||
"""Raised when the QDC device resource is not available (retryable)."""
|
||||
|
||||
|
||||
class RunnerUnavailableError(Exception):
|
||||
"""Raised when QDC runner or capacity is not available within timeout (non-retryable)."""
|
||||
|
||||
|
||||
_SCRIPTS_DIR = Path(__file__).parent
|
||||
_TESTS_DIR = _SCRIPTS_DIR / "tests"
|
||||
|
||||
@@ -253,8 +258,9 @@ DEVICE_PLATFORM: dict[str, Platform] = {
|
||||
# =============================================================================
|
||||
|
||||
|
||||
def wait_for_job(client, job_id: str, timeout: int) -> str:
|
||||
def wait_for_job(client, job_id: str, timeout: int, runner_timeout: int = RUNNER_TIMEOUT) -> str:
|
||||
elapsed = 0
|
||||
submitted_elapsed = 0
|
||||
last_state = None
|
||||
consecutive_errors = 0
|
||||
max_consecutive_errors = 5
|
||||
@@ -285,6 +291,26 @@ def wait_for_job(client, job_id: str, timeout: int) -> str:
|
||||
if raw != last_state:
|
||||
log.info("Job %s: %s", job_id, raw)
|
||||
last_state = raw
|
||||
|
||||
is_submitted = (status == JobState.SUBMITTED) or (
|
||||
isinstance(raw, str) and raw.lower() == "submitted"
|
||||
)
|
||||
if is_submitted:
|
||||
submitted_elapsed += POLL_INTERVAL
|
||||
if submitted_elapsed >= runner_timeout:
|
||||
try:
|
||||
qdc_api.abort_job(client, job_id)
|
||||
log.warning(
|
||||
"Aborted job %s: runner did not pick up job within %ds",
|
||||
job_id,
|
||||
runner_timeout,
|
||||
)
|
||||
except Exception as e:
|
||||
log.warning("Failed to abort job %s: %s", job_id, e)
|
||||
raise RunnerUnavailableError(
|
||||
f"Runner did not pick up job {job_id} within {runner_timeout}s"
|
||||
)
|
||||
|
||||
time.sleep(POLL_INTERVAL)
|
||||
elapsed += POLL_INTERVAL
|
||||
# Abort to free the QDC concurrency slot instead of leaking it.
|
||||
@@ -312,10 +338,10 @@ def wait_for_log_upload(client, job_id: str) -> None:
|
||||
log.warning("Timed out waiting for log upload after %ds", LOG_UPLOAD_TIMEOUT)
|
||||
|
||||
|
||||
def wait_for_capacity(client, max_jobs: int = MAX_CONCURRENT_JOBS) -> None:
|
||||
def wait_for_capacity(client, timeout: int = CAPACITY_TIMEOUT, max_jobs: int = MAX_CONCURRENT_JOBS) -> None:
|
||||
"""Block until the user's active (non-terminal) QDC job count is below max_jobs."""
|
||||
elapsed = 0
|
||||
while elapsed < CAPACITY_TIMEOUT:
|
||||
while elapsed < timeout:
|
||||
jobs_page = qdc_api.get_jobs_list(client, page_number=0, page_size=50)
|
||||
if jobs_page is None:
|
||||
log.warning(
|
||||
@@ -338,7 +364,7 @@ def wait_for_capacity(client, max_jobs: int = MAX_CONCURRENT_JOBS) -> None:
|
||||
time.sleep(CAPACITY_POLL)
|
||||
elapsed += CAPACITY_POLL
|
||||
raise TimeoutError(
|
||||
f"Capacity wait timed out after {CAPACITY_TIMEOUT}s"
|
||||
f"Capacity wait timed out after {timeout}s"
|
||||
)
|
||||
|
||||
|
||||
@@ -445,7 +471,11 @@ def fetch_logs_and_parse_tests(
|
||||
)
|
||||
|
||||
|
||||
def write_summary(result: JobResult, title: str = "QDC Test Results") -> None:
|
||||
def write_summary(
|
||||
result: JobResult,
|
||||
title: str = "QDC Test Results",
|
||||
note: str | None = None,
|
||||
) -> None:
|
||||
summary_path = os.environ.get("GITHUB_STEP_SUMMARY")
|
||||
if not summary_path:
|
||||
return
|
||||
@@ -454,8 +484,12 @@ def write_summary(result: JobResult, title: str = "QDC Test Results") -> None:
|
||||
|
||||
lines = [
|
||||
f"## {title}\n",
|
||||
f"Overall: {icon} {'PASSED' if result.passed else 'FAILED'}\n",
|
||||
]
|
||||
if note:
|
||||
lines.append(f"{note}\n")
|
||||
else:
|
||||
lines.append(f"Overall: {icon} {'PASSED' if result.passed else 'FAILED'}\n")
|
||||
|
||||
reportable = {n: ok for n, ok in result.tests.items() if "test_install" not in n}
|
||||
if reportable:
|
||||
lines += ["| Test | Result |", "| ---- | ------ |"]
|
||||
@@ -464,7 +498,7 @@ def write_summary(result: JobResult, title: str = "QDC Test Results") -> None:
|
||||
passed_n = sum(1 for v in reportable.values() if v)
|
||||
failed_n = sum(1 for v in reportable.values() if not v)
|
||||
lines += ["", f"**{passed_n} passed, {failed_n} failed**"]
|
||||
else:
|
||||
elif not note:
|
||||
lines.append("_No per-test data available._")
|
||||
|
||||
failed_names = [n for n, ok in reportable.items() if not ok]
|
||||
@@ -519,6 +553,8 @@ def parse_args() -> argparse.Namespace:
|
||||
help="Test suite to run (default: bench)")
|
||||
p.add_argument("--job-timeout", type=int, default=JOB_TIMEOUT, metavar="SECONDS",
|
||||
help=f"Max seconds to wait for job completion (default: {JOB_TIMEOUT})")
|
||||
p.add_argument("--runner-timeout", type=int, default=RUNNER_TIMEOUT, metavar="SECONDS",
|
||||
help=f"Max seconds to wait for runner to pick up job (default: {RUNNER_TIMEOUT})")
|
||||
p.add_argument("--retries", type=int, default=DEFAULT_RETRIES, metavar="N",
|
||||
help="Number of retries when device is unavailable (default: 0)")
|
||||
p.add_argument("--retry-delay", type=int, default=RETRY_DELAY, metavar="SECONDS",
|
||||
@@ -537,9 +573,11 @@ def _submit_and_run_job(client, args, spec, target_id, artifact_id) -> JobResult
|
||||
test failure).
|
||||
"""
|
||||
try:
|
||||
wait_for_capacity(client)
|
||||
wait_for_capacity(client, timeout=args.runner_timeout)
|
||||
except TimeoutError:
|
||||
raise DeviceUnavailableError("Capacity wait timed out — device busy")
|
||||
raise RunnerUnavailableError(
|
||||
f"Capacity wait timed out after {args.runner_timeout}s - device busy"
|
||||
)
|
||||
|
||||
job_name = spec.job_name_fmt.format(base="llama.cpp Hexagon tests")
|
||||
|
||||
@@ -559,11 +597,16 @@ def _submit_and_run_job(client, args, spec, target_id, artifact_id) -> JobResult
|
||||
job_parameters=[JobSubmissionParameter.WIFIENABLED],
|
||||
)
|
||||
if job_id is None:
|
||||
raise DeviceUnavailableError("Job submission failed — device may be unavailable")
|
||||
raise DeviceUnavailableError("Job submission failed - device may be unavailable")
|
||||
log.info("Job submitted: %s (device=%s)", job_id, args.device)
|
||||
|
||||
try:
|
||||
job_status = wait_for_job(client, job_id, timeout=args.job_timeout)
|
||||
job_status = wait_for_job(
|
||||
client,
|
||||
job_id,
|
||||
timeout=args.job_timeout,
|
||||
runner_timeout=args.runner_timeout,
|
||||
)
|
||||
except TimeoutError as e:
|
||||
raise DeviceUnavailableError(str(e))
|
||||
log.info("Job %s finished: %s", job_id, job_status)
|
||||
@@ -637,6 +680,14 @@ def main() -> int:
|
||||
try:
|
||||
result = _submit_and_run_job(client, args, spec, target_id, artifact_id)
|
||||
break
|
||||
except RunnerUnavailableError as e:
|
||||
log.warning("QDC runner unavailable for %s after %ds - job skipped: %s", args.device, args.runner_timeout, e)
|
||||
write_summary(
|
||||
JobResult(passed=True, tests={}),
|
||||
title=f"QDC Tests ({args.device})",
|
||||
note=f"SKIPPED: Runner did not pick up job within {args.runner_timeout}s. Skipped to avoid stalling CI.",
|
||||
)
|
||||
return 0
|
||||
except DeviceUnavailableError as e:
|
||||
if attempt < max_attempts:
|
||||
log.warning(
|
||||
|
||||
@@ -1,6 +1,8 @@
|
||||
-r ../../requirements/requirements-convert_legacy_llama.txt
|
||||
--extra-index-url https://download.pytorch.org/whl/cpu
|
||||
pillow~=11.3.0
|
||||
pillow~=12.3.0
|
||||
|
||||
torch==2.11.0 # check_requirements: ignore "=="
|
||||
torchvision==0.26.0; platform_machine != "s390x" # check_requirements: ignore "=="
|
||||
torch==2.11.0; python_version < "3.15" # check_requirements: ignore "=="
|
||||
torch==2.15.0; python_version >= "3.15" # check_requirements: ignore "=="
|
||||
torchvision==0.26.0; platform_machine != "s390x" and python_version < "3.15" # check_requirements: ignore "=="
|
||||
torchvision==0.29.2; platform_machine != "s390x" and python_version >= "3.15" # check_requirements: ignore "=="
|
||||
|
||||
@@ -2,7 +2,7 @@ pytest~=8.3.3
|
||||
pytest-xdist~=3.6
|
||||
filelock~=3.16
|
||||
numpy~=2.2.6; python_version < "3.14"
|
||||
numpy~=2.4.6; python_version >= "3.14"
|
||||
numpy~=2.5.4; python_version >= "3.14"
|
||||
openai~=2.14.0
|
||||
prometheus-client~=0.20.0
|
||||
requests~=2.32.3
|
||||
|
||||
@@ -1,6 +1,8 @@
|
||||
engine-strict=true
|
||||
ignore-scripts=true
|
||||
min-release-age=7
|
||||
# Allow root-declared remote tarball dependencies (llama-brand) on npm 12+
|
||||
allow-remote=root
|
||||
# The single-bundle production build needs more than the default heap on CI runners;
|
||||
# npm passes this to every script it runs, on every platform
|
||||
node-options=--max-old-space-size=6144
|
||||
|
||||
Reference in New Issue
Block a user