llama.cpp gfx1151 optimizations for Qwen3.8 27B

A developer contributed optimizations to llama.cpp for AMD's gfx1151 GPU, specifically targeting Qwen3.8 27B's SSM convolution input pattern. The patch introduces a fast LDS-transpose path for dim-0 concatenation when the second source tensor is transposed-contiguous, staging a 64x64 tile through shared memory to achieve fully coalesced memory access. Benchmarks on the Qwen3.5 shape show throughput improved from 35 GB/s to 70 GB/s.

| diff --git a/ggml/src/ggml-cuda/concat.cu b/ggml/src/ggml-cuda/concat.cu | | | index 6df89013c..e5337ec30 100644 | | | --- a/ggml/src/ggml-cuda/concat.cu | | | +++ b/ggml/src/ggml-cuda/concat.cu | | | @@ -79,6 +79,76 @@ static void concat cont cuda const T x, | | | concat cont<T, 2 <<<num blocks, CUDA CONCAT BLOCK SIZE, 0, stream x, y, dst, ne00, ne01, ne02, ne0, ne1, ne2 ; | | | } | | | +// Fast LDS-transpose path for a dim-0 concat whose src1 is transposed-contiguous | | | +// nb 1 == element size < nb 0 . That is the qwen35 ssm conv input pattern: | | | +// concat conv state 3, d inner , transpose x n tokens, d inner along dim 0. | | | +// | | | +// The generic non-cont kernel walks i0 with consecutive threads, so for a | | | +// transposed src1 every lane of a wave reads a different row - one memory | | | +// transaction per lane on both the load and the store side. Staging a 32x32 | | | +// tile through LDS makes both sides fully coalesced. | | | + define CONCAT TRANSPOSE TILE 64 // 64x64 measured best on gfx1151 35 - 70 GB/s on the qwen35 shape | | | + define CONCAT TRANSPOSE NTX 64 | | | + define CONCAT TRANSPOSE NTY 4 | | | + | | | +template <typename T | | | +static global void launch bounds CONCAT TRANSPOSE NTX CONCAT TRANSPOSE NTY | | | + concat dim0 transposed src1 | | | + const char restrict src0, | | | + const char restrict src1, | | | + char restrict dst, | | | + int64 t ne00, | | | + int64 t ne0, | | | + int64 t ne1, | | | + int64 t ne2, | | | + uint64 t nb00, uint64 t nb01, uint64 t nb02, uint64 t nb03, | | | + uint64 t nb10, uint64 t nb11, uint64 t nb12, uint64 t nb13, | | | + uint64 t nb0, uint64 t nb1, uint64 t nb2, uint64 t nb3 { | | | + constexpr int TILE = CONCAT TRANSPOSE TILE; | | | + constexpr int NTX = CONCAT TRANSPOSE NTX; | | | + constexpr int NTY = CONCAT TRANSPOSE NTY; | | | + shared T tile TILE TILE + 1 ; // +1 to avoid LDS bank conflicts on the transposed read | | | + | | | + const int tx = threadIdx.x; | | | + const int ty = threadIdx.y; | | | + const int i23 = blockIdx.z; | | | + const int i2 = i23 % int ne2; | | | + const int i3 = i23 / int ne2; | | | + const int64 t i0 base = int64 t blockIdx.x TILE; | | | + const int64 t i1 base = int64 t blockIdx.y TILE; | | | + | | | + // load: consecutive tx - consecutive i1, i.e. consecutive addresses in src1 | | | + pragma unroll | | | + for int c = 0; c < TILE; c += NTX { | | | + const int64 t load i1 = i1 base + c + tx; | | | + for int r = ty; r < TILE; r += NTY { | | | + const int64 t load i0 = i0 base + r; | | | + T v = T 0 ; | | | + if load i0 < ne0 && load i1 < ne1 { | | | + if load i0 < ne00 { | | | + v = const T src0 + i3 nb03 + i2 nb02 + load i1 nb01 + load i0 nb00 ; | | | + } else { | | | + v = const T src1 + i3 nb13 + i2 nb12 + load i0 - ne00 nb10 + load i1 nb11 ; | | | + } | | | + } | | | + tile r c + tx = v; | | | + } | | | + } | | | + syncthreads ; | | | + | | | + // store: consecutive tx - consecutive i0, i.e. consecutive addresses in dst | | | + pragma unroll | | | + for int c = 0; c < TILE; c += NTX { | | | + const int64 t store i0 = i0 base + c + tx; | | | + for int r = ty; r < TILE; r += NTY { | | | + const int64 t store i1 = i1 base + r; | | | + if store i0 < ne0 && store i1 < ne1 { | | | + T dst + i3 nb3 + i2 nb2 + store i1 nb1 + store i0 nb0 = tile c + tx r ; | | | + } | | | + } | | | + } | | | +} | | | + | | | // non-contiguous kernel slow | | | template <typename T, int dim | | | static global void launch bounds CUDA CONCAT BLOCK SIZE | | | @@ -164,6 +234,25 @@ static void concat cuda const ggml tensor src0, const ggml tensor src1, ggml | | | GGML ASSERT ggml is quantized src0- type ; | | | dim3 grid dim dst- ne 1 , dst- ne 2 , dst- ne 3 ; | | | + // transposed-src1 dim-0 concat qwen35 ssm conv input : stage through LDS | | | + if dim == 0 | | | + && src1- nb 1 == sizeof T && src1- nb 0 src1- nb 1 | | | + && src0- nb 0 == sizeof T | | | + && src0- ne 1 == dst- ne 1 && src1- ne 1 == dst- ne 1 { | | | + constexpr int TILE = CONCAT TRANSPOSE TILE; | | | + const dim3 grid dst- ne 0 + TILE - 1 / TILE, | | | + dst- ne 1 + TILE - 1 / TILE, | | | + unsigned dst- ne 2 dst- ne 3 ; | | | + const dim3 block CONCAT TRANSPOSE NTX, CONCAT TRANSPOSE NTY, 1 ; | | | + concat dim0 transposed src1<T <<<grid, block, 0, stream | | | + const char src0- data, const char src1- data, char dst- data, | | | + src0- ne 0 , dst- ne 0 , dst- ne 1 , dst- ne 2 , | | | + src0- nb 0 , src0- nb 1 , src0- nb 2 , src0- nb 3 , | | | + src1- nb 0 , src1- nb 1 , src1- nb 2 , src1- nb 3 , | | | + dst- nb 0 , dst- nb 1 , dst- nb 2 , dst- nb 3 ; | | | + return; | | | + } | | | + | | | auto launch kernel = & auto dim { | | | concat non cont<T, dim <<<grid dim, CUDA CONCAT BLOCK SIZE, 0, stream | | | const char src0- data, const char src1- data, char dst- data, | | | diff --git a/ggml/src/ggml-cuda/fattn-tile.cuh b/ggml/src/ggml-cuda/fattn-tile.cuh | | | index d1164b852..6611f8096 100644 | | | --- a/ggml/src/ggml-cuda/fattn-tile.cuh | | | +++ b/ggml/src/ggml-cuda/fattn-tile.cuh | | | @@ -312,8 +312,25 @@ static constexpr host device uint32 t ggml cuda fattn tile get config am | | | return 0; | | | } | | | +// gfx1151, DKQ = DV = 256 the qwen35 attention head : nbatch fa 64 - 128. | | | +// | | | +// The generic RDNA table's 64 is the right answer only while the kernel is spilling: with | | | +// the staging loops fully unrolled this shape sits at 256 VGPRs with 5 spilled, and every | | | +// config that wants more registers is punished for that rather than on its merits. Once the | | | +// unroll is capped see GGML FA TILE UNROLL STAGE the optimum moves -- 128 was the worst | | | +// point of the original sweep and is the best one now. LDS goes 37888 - ~33800 B, still | | | +// three workgroups per WGP. Measured 8.65 - 9.40 TFLOPS at 512x4096 and 8.93 - 9.18 at | | | +// 512x33280. | | | +static constexpr host device uint32 t ggml cuda fattn tile get config amd rdna3 5 const int DKQ, const int DV, const int ncols { | | | + GGML CUDA FATTN TILE CONFIG CASE 256, 256, 32, 256, 2, 128, 64 | | | + return ggml cuda fattn tile get config amd rdna DKQ, DV, ncols ; | | | +} | | | + | | | static host uint32 t ggml cuda fattn tile get config const int DKQ, const int DV, const int ncols, const int cc { | | | if GGML CUDA CC IS AMD cc { | | | + if GGML CUDA CC IS RDNA3 5 cc { | | | + return ggml cuda fattn tile get config amd rdna3 5 DKQ, DV, ncols ; | | | + } | | | if GGML CUDA CC IS RDNA cc { | | | return ggml cuda fattn tile get config amd rdna DKQ, DV, ncols ; | | | } | | | @@ -327,7 +344,9 @@ static host uint32 t ggml cuda fattn tile get config const int DKQ, const in | | | static constexpr device uint32 t ggml cuda fattn tile get config const int DKQ, const int DV, const int ncols { | | | ifdef GGML USE HIP | | | - ifdef RDNA | | | + ifdef RDNA3 5 | | | + return ggml cuda fattn tile get config amd rdna3 5 DKQ, DV, ncols ; | | | + elif defined RDNA | | | return ggml cuda fattn tile get config amd rdna DKQ, DV, ncols ; | | | else | | | return ggml cuda fattn tile get config amd DKQ, DV, ncols ; | | | @@ -482,6 +501,19 @@ static device forceinline void flash attn tile load tile | | | // Function that performs a single iteration in for the KQ matrix multiplication: | | | template <int warp size, int nwarps, int ncols1, int ncols2, int DKQ, int nbatch fa, int nbatch K, | | | bool use logit softcap, bool oob check, typename T vec dot | | | +// gfx1151: these two loops stage 24 half2 of K/Q resp. V/KQ per iteration, and fully | | | +// unrolling them lets the scheduler keep several iterations of loads in flight. Combined | | | +// with this target's raised -amdgpu-unroll-threshold-local that pushes the DKQ=256 kernel | | | +// to 256 VGPRs with 5 spilled, and the spills then stop it keeping any further loads in | | | +// flight -- the kernel runs at 45% instruction-issue utilisation. Capping the unroll takes | | | +// it to 127 VGPRs with no spills and +23% at 32k context. Measured on RDNA3.5 only, so it | | | +// is scoped there; other targets keep the full unroll. | | | + if defined GGML USE HIP && defined RDNA3 5 | | | + define GGML FA TILE UNROLL STAGE Pragma "unroll 4" | | | + else | | | + define GGML FA TILE UNROLL STAGE Pragma "unroll" | | | + endif | | | + | | | static device forceinline void flash attn tile iter KQ | | | T vec dot const Q tmp, | | | const half2 const restrict K h2, | | | @@ -504,7 +536,7 @@ static device forceinline void flash attn tile iter KQ | | | ifdef FAST FP16 AVAILABLE | | | static assert nbatch K/2 % cpy ne == 0, "bad nbatch K" ; | | | - pragma unroll | | | +GGML FA TILE UNROLL STAGE | | | for int k KQ 1 = 0; k KQ 1 < nbatch K/2; k KQ 1 += cpy ne { | | | align 16 half2 K k nbatch fa/ np warp size cpy ne ; | | | align 16 half2 Q k cpw cpy ne ; | | | @@ -723,7 +755,7 @@ static device forceinline void flash attn tile iter | | | syncthreads ; | | | ifdef FAST FP16 AVAILABLE | | | - pragma unroll | | | +GGML FA TILE UNROLL STAGE | | | for int k1 = 0; k1 < nbatch V; k1 += np { | | | align 16 half2 V k DVp/2 /warp size ; | | | align 16 half2 KQ k cpw ; | | | @@ -755,7 +787,7 @@ static device forceinline void flash attn tile iter | | | } | | | } | | | else | | | - pragma unroll | | | +GGML FA TILE UNROLL STAGE | | | for int k1 = 0; k1 < nbatch V; k1 += np { | | | align 16 float2 V k DVp/2 /warp size ; | | | align 16 float KQ k cpw ; | | | diff --git a/ggml/src/ggml-cuda/fattn-wmma-d256.cu b/ggml/src/ggml-cuda/fattn-wmma-d256.cu | | | new file mode 100644 | | | index 000000000..908eb00bb | | | --- /dev/null | | | +++ b/ggml/src/ggml-cuda/fattn-wmma-d256.cu | | | @@ -0,0 +1,55 @@ | | | + include "fattn-wmma-d256.cuh" | | | + include <cstdlib | | | + | | | +template<int ncols1, int ncols2, ggml type type K, ggml type type V | | | +static void launch wmma d256 ggml backend cuda context & ctx, ggml tensor dst { | | | + constexpr int DKQ = 256; | | | + constexpr int DV = 256; | | | + launch fattn<DV, ncols1, ncols2 | | | + ctx, dst, flash attn wmma d256<DKQ, DV, ncols1, ncols2, type K, type V , | | | + / nwarps = / FA WMMA D256 NWARPS, | | | + / nbytes shared= / getenv "GGML FA LDS ABLATE" | | | + ? size t atoi getenv "GGML FA LDS ABLATE" : ggml cuda fattn wmma d256 nbytes shared , | | | + / nbatch fa = / FA WMMA D256 FA, | | | + // false: the kernel dequantises K/V itself, so launch fattn must not pre-convert | | | + // the cache to f16 -- that conversion is exactly the DRAM traffic we are avoiding. | | | + / need f16 K = / false, | | | + / need f16 V = / false, | | | + / stream k = / false, | | | + / warp size = / 32 ; | | | +} | | | + | | | +void ggml cuda flash attn ext wmma d256 ggml backend cuda context & ctx, ggml tensor dst { | | | + const ggml tensor Q = dst- src 0 ; | | | + const ggml tensor K = dst- src 1 ; | | | + | | | + GGML ASSERT Q- ne 0 == 256 ; | | | + | | | + const int gqa ratio = Q- ne 2 / K- ne 2 ; | | | + | | | + // Widest power-of-two head group that divides the GQA ratio, capped so that | | | + // ncols1 = 64/ncols2 still leaves at least one Q column per block. | | | + int ncols2 = 1; | | | + while gqa ratio % 2 ncols2 == 0 && ncols2 < 8 { | | | + ncols2 = 2; | | | + } | | | + | | | + constexpr int NC = FA WMMA D256 NC; | | | + const ggml type tK = K- type; | | | + const ggml type tV = dst- src 2 - type; | | | + | | | + define FA WMMA DISPATCH KV N1, N2 \ | | | + if tK == GGML TYPE F16 && tV == GGML TYPE F16 { launch wmma d256<N1, N2, GGML TYPE F16, GGML TYPE F16 ctx, dst ; return; } \ | | | + if tK == GGML TYPE Q8 0 && tV == GGML TYPE Q8 0 { launch wmma d256<N1, N2, GGML TYPE Q8 0, GGML TYPE Q8 0 ctx, dst ; return; } \ | | | + if tK == GGML TYPE Q4 0 && tV == GGML TYPE Q4 0 { launch wmma d256<N1, N2, GGML TYPE Q4 0, GGML TYPE Q4 0 ctx, dst ; return; } \ | | | + GGML ABORT "unsupported KV type combination" ; | | | + | | | + switch ncols2 { | | | + case 1: FA WMMA DISPATCH KV NC/1, 1 ; break; | | | + case 2: FA WMMA DISPATCH KV NC/2, 2 ; break; | | | + case 4: FA WMMA DISPATCH KV NC/4, 4 ; break; | | | + case 8: FA WMMA DISPATCH KV NC/8, 8 ; break; | | | + default: GGML ABORT "fatal error" ; | | | + } | | | + undef FA WMMA DISPATCH KV | | | +} | | | diff --git a/ggml/src/ggml-cuda/fattn-wmma-d256.cuh b/ggml/src/ggml-cuda/fattn-wmma-d256.cuh | | | new file mode 100644 | | | index 000000000..88187a552 | | | --- /dev/null | | | +++ b/ggml/src/ggml-cuda/fattn-wmma-d256.cuh | | | @@ -0,0 +1,426 @@ | | | + pragma once | | | + | | | + include "common.cuh" | | | + include "fattn-common.cuh" | | | + | | | +// WMMA flash attention for RDNA3/3.5 at head size 256. | | | +// | | | +// The generic tile kernel emits no WMMA at all, and fattn-mma-f16 gives every warp the | | | +// full DV in registers, which at DV=256 spills 1 kB/lane to scratch on a 256-VGPR wave. | | | +// This kernel splits DV across warps instead, so the O accumulator is 64 VGPRs/lane. | | | +// | | | +// Fixed shape: 8 warps, 64 columns per block, 32 KV rows per iteration. | | | +// NQB = 4 q blocks, NKB = 2 kv blocks, NQB NKB == 8 - one 16x16 S tile per warp. | | | +// O phase: warp w owns q block w/NKB and DV slice w%NKB 128, +128 = 8 tiles. | | | +// | | | +// WMMA fragment layout wave32, gfx11 : | | | +// A: lane t supplies row t % 16 of A, all 16 k | | | +// B: lane t supplies column t % 16 of B, all 16 k | | | +// D: lane t holds COLUMN t % 16 ; its 8 slots are rows 2 l + t/16 | | | + | | | + define FA WMMA D256 D 256 | | | + ifndef FA WMMA D256 NC | | | + define FA WMMA D256 NC 64 | | | + endif | | | + ifndef FA WMMA D256 FA | | | + define FA WMMA D256 FA 32 | | | + endif | | | + define FA WMMA D256 NQB FA WMMA D256 NC / 16 | | | + define FA WMMA D256 NKB FA WMMA D256 FA / 16 | | | +// KSPLIT warps cooperate on each S tile by splitting the head dimension. It buys waves per | | | +// workgroup without any extra LDS, which is what this kernel is short of: the Q tile alone | | | +// is 33 kB, so only 2 workgroups fit per WGP however the rest is arranged. | | | + define FA WMMA D256 KSPLIT 1 | | | + define FA WMMA D256 NWARPS FA WMMA D256 NQB FA WMMA D256 NKB FA WMMA D256 KSPLIT | | | + define FA WMMA D256 NSW FA WMMA D256 NKB FA WMMA D256 KSPLIT // warps sharing a q block | | | + define FA WMMA D256 DVW FA WMMA D256 D / FA WMMA D256 NSW | | | + define FA WMMA D256 OT FA WMMA D256 DVW / 16 | | | + define FA WMMA D256 KPAD 8 | | | + define FA WMMA D256 FPAD 8 | | | + | | | +static constexpr size t ggml cuda fattn wmma d256 nbytes shared { | | | + // Q is staged through LDS once and then held in registers, so its footprint overlaps | | | + // the K/V/P working set instead of adding to it. Keeping Q resident would put the | | | + // block at 60 kB and 2 workgroups per WGP; this way it is 3, which measured -20% on | | | + // the long-context shape. | | | + const size t nb Q = size t FA WMMA D256 NC FA WMMA D256 D + FA WMMA D256 KPAD 2; | | | + const size t nb K = size t FA WMMA D256 FA FA WMMA D256 D + FA WMMA D256 KPAD 2; | | | + const size t nb V = size t FA WMMA D256 D FA WMMA D256 FA + FA WMMA D256 FPAD 2; | | | + const size t nb S = size t FA WMMA D256 NQB FA WMMA D256 NKB 32 8 sizeof float ; | | | + size t nb KV = nb K nb V ? nb K : nb V; | | | + if nb S nb KV { nb KV = nb S; } | | | + const size t nb P = size t FA WMMA D256 NC FA WMMA D256 FA + FA WMMA D256 FPAD 2; | | | + const size t nb R = size t 2 FA WMMA D256 NQB FA WMMA D256 NKB 16 sizeof float ; | | | + return nb Q + nb KV + nb P + nb R; | | | +} | | | + | | | + if defined AMD WMMA AVAILABLE && defined RDNA3 | | | + | | | +typedef Float16 fa half16 attribute ext vector type 16 ; | | | +typedef float fa float8 attribute ext vector type 8 ; | | | + | | | +union fa frag16 { fa half16 h; int4 v 2 ; }; | | | + | | | +static device forceinline fa half16 fa load frag const Float16 p { | | | + fa frag16 f; | | | + f.v 0 = const int4 p + 0 ; | | | + f.v 1 = const int4 p + 8 ; | | | + return f.h; | | | +} | | | + | | | +// Read 8 consecutive head-dimension elements of one K/V row, dequantising on the way in. | | | +// The point of doing this here rather than letting launch fattn pre-convert the cache to | | | +// f16 is that the DRAM traffic is what bounds this kernel: at head size 256 the arithmetic | | | +// intensity is columns per block /2 MAC per byte, so halving the bytes doubles the roofline. | | | +template <ggml type type KV | | | +static device forceinline int4 fa load8 const char restrict row, int i0 { | | | + int4 out; | | | + Float16 o = Float16 &out; | | | + if constexpr type KV == GGML TYPE F16 { | | | + out = const int4 row + i0 2 ; | | | + } else if constexpr type KV == GGML TYPE Q8 0 { | | | + const block q8 0 b = const block q8 0 row + i0/QK8 0; | | | + const float d = half2float b- d ; | | | + const int j0 = i0 % QK8 0; | | | + pragma unroll | | | + for int j = 0; j < 8; ++j { o j = Float16 d float b- qs j0 + j ; } | | | + } else { | | | + static assert type KV == GGML TYPE Q4 0, "unsupported KV type" ; | | | + const block q4 0 b = const block q4 0 row + i0/QK4 0; | | | + const float d = half2float b- d ; | | | + const int j0 = i0 % QK4 0; | | | + const int sh = j0 / 16 4; // low nibbles hold the first half of the block | | | + const uint8 t q = b- qs + j0 % 16 ; | | | + pragma unroll | | | + for int j = 0; j < 8; ++j { o j = Float16 d float int q j sh & 0x0F - 8 ; } | | | + } | | | + return out; | | | +} | | | + | | | +// Single element, for the transposed V staging where each lane walks kv, not the head dim. | | | +template <ggml type type KV | | | +static device forceinline Float16 fa load1 const char restrict row, int i { | | | + if constexpr type KV == GGML TYPE F16 { | | | + return const Float16 row + i 2 ; | | | + } else if constexpr type KV == GGML TYPE Q8 0 { | | | + const block q8 0 b = const block q8 0 row + i/QK8 0; | | | + return Float16 half2float b- d float b- qs i % QK8 0 ; | | | + } else { | | | + static assert type KV == GGML TYPE Q4 0, "unsupported KV type" ; | | | + const block q4 0 b = const block q4 0 row + i/QK4 0; | | | + const int j = i % QK4 0; | | | + return Float16 half2float b- d float int b- qs j % 16 j/16 4 & 0x0F - 8 ; | | | + } | | | +} | | | + | | | +template<int DKQ, int DV, int ncols1, int ncols2, ggml type type K, ggml type type V | | | + launch bounds FA WMMA D256 NWARPS 32, 1 | | | +static global void flash attn wmma d256 | | | + const char Q ptr, | | | + const char K ptr, | | | + const char V ptr, | | | + const char mask ptr, | | | + const char sinks ptr, | | | + const int KV max ptr, | | | + float dst ptr, | | | + float2 dst meta ptr, | | | + const float scale, | | | + const float max bias, | | | + const float m0, | | | + const float m1, | | | + const uint32 t n head log2, | | | + const float logit softcap, | | | + const int32 t ne00, const uint3 ne01, const int32 t ne02, const int32 t ne03, | | | + const int32 t nb01, const int32 t nb02, const int32 t nb03, | | | + const int32 t ne10, const int32 t ne11, const int32 t ne12, const int32 t ne13, | | | + const int32 t nb11, const int32 t nb12, const int64 t nb13, | | | + const int32 t nb21, const int32 t nb22, const int64 t nb23, | | | + const int32 t ne31, const int32 t ne32, const int32 t ne33, | | | + const int32 t nb31, const int32 t nb32, const int64 t nb33 { | | | + constexpr int D = FA WMMA D256 D; | | | + constexpr int NC = FA WMMA D256 NC; | | | + constexpr int FAB = FA WMMA D256 FA; | | | + constexpr int NWARPS = FA WMMA D256 NWARPS; | | | + constexpr int NQB = FA WMMA D256 NQB; | | | + constexpr int NKB = FA WMMA D256 NKB; | | | + constexpr int KSPLIT = FA WMMA D256 KSPLIT; | | | + constexpr int NSW = FA WMMA D256 NSW; | | | + constexpr int DVW = FA WMMA D256 DVW; | | | + constexpr int OT = FA WMMA D256 OT; | | | + constexpr int KPAD = FA WMMA D256 KPAD; | | | + constexpr int FPAD = FA WMMA D256 FPAD; | | | + constexpr int WS = 32; | | | + constexpr int NTHR = NWARPS WS; | | | + | | | + static assert DKQ == D && DV == D, "this kernel is head size 256 only" ; | | | + static assert ncols1 ncols2 == NC, "bad column count" ; | | | + static assert NQB NKB KSPLIT == NWARPS, "bad warp count" ; | | | + static assert D % 16 KSPLIT == 0, "bad head split" ; | | | + | | | + GGML UNUSED VARS max bias, m0, m1, n head log2, logit softcap, ne00, ne10, ne12, ne13, | | | + ne31, ne32, nb32, nb33, sinks ptr ; | | | + | | | + const int tid = threadIdx.x + threadIdx.y WS; | | | + const int warp = threadIdx.y; | | | + const int lane = threadIdx.x; | | | + const int l16 = lane % 16; | | | + const int lhi = lane / 16; | | | + | | | + const int col Q 0 = blockIdx.x ncols1; | | | + | | | + const int sequence = blockIdx.z / ne02/ncols2 ; | | | + const int head0 = blockIdx.z ncols2 - sequence ne02; | | | + const int gqa ratio = ne02 / ne12; | | | + | | | + const float Q f = const float Q ptr + nb03 sequence + nb02 head0 ; | | | + const char K base = K ptr + nb13 sequence + nb12 head0 / gqa ratio ; | | | + const char V base = V ptr + nb23 sequence + nb22 head0 / gqa ratio ; | | | + const half maskh = mask ptr ? const half mask ptr + nb33 sequence % ne33 : nullptr; | | | + | | | + const int stride mask = nb31 / sizeof half ; | | | + | | | + extern shared char fa smem ; | | | + Float16 Qs = Float16 fa smem; // NC D+KPAD | | | + Float16 Ks = Qs + NC D + KPAD ; // FAB D+KPAD | | | + Float16 Vt = Ks; // D FAB+FPAD , aliases Ks | | | + constexpr size t kv halves = size t FAB D + KPAD size t D FAB + FPAD | | | + ? size t FAB D + KPAD : size t D FAB + FPAD ; | | | + Float16 Ps = Ks + kv halves; // NC FAB+FPAD | | | + float Rm = float Ps + NC FAB + FPAD ; // NQB NKB 16 | | | + float Rl = Rm + NQB NKB 16; // NQB NKB 16 | | | + | | | + ggml cuda pdl sync ; | | | + | | | + // ---- Q into LDS, pre-scaled and converted to f16 ---- | | | + // column jc - q position col Q 0 + jc/ncols2, head head0 + jc%ncols2 | | | + for int idx = tid 8; idx < NC D; idx += NTHR 8 { | | | + const int jc = idx / D; | | | + const int i0 = idx % D; | | | + const int j = jc / ncols2; | | | + const int c = jc % ncols2; | | | + const float src = Q f + c nb02/sizeof float + fastmodulo col Q 0 + j, ne01 nb01/sizeof float + i0; | | | + fa frag16 f; | | | + pragma unroll | | | + for int i = 0; i < 8; ++i { | | | + Float16 &f.v 0 i = Float16 src i scale ; | | | + } | | | + int4 &Qs jc D + KPAD + i0 = f.v 0 ; | | | + } | | | + | | | + const int s tile = warp % NQB NKB ; | | | + const int k half = warp / NQB NKB ; | | | + const int qb = s tile / NKB; | | | + const int s kb = s tile % NKB; | | | + // warps sharing this q block: NSW of them, each taking one DV slice in the O phase | | | + const int o slot = s tile % NKB + NKB k half; | | | + const int o d0 = o slot DVW; | | | + | | | + fa float8 Oacc OT ; | | | + pragma unroll | | | + for int i = 0; i < OT; ++i { Oacc i = fa float8 0.0f ; } | | | + | | | + float m run 8 , l run 8 ; | | | + pragma unroll | | | + for int l = 0; l < 8; ++l { m run l = -FLT MAX/2.0f; l run l = 0.0f; } | | | + | | | + syncthreads ; | | | + | | | + const int k VKQ max = KV max ptr ? KV max ptr sequence gridDim.x + blockIdx.x : ne11; | | | + | | | + for int k VKQ 0 = blockIdx.y FAB; k VKQ 0 < k VKQ max; k VKQ 0 += gridDim.y FAB { | | | + const int kv sup = k VKQ max - k VKQ 0 < FAB ? k VKQ max - k VKQ 0 : FAB; | | | + | | | + // ---- K into LDS ---- | | | + for int idx = tid 8; idx < FAB D; idx += NTHR 8 { | | | + const int kv = idx / D; | | | + const int i0 = idx % D; | | | + int4 v = make int4 0, 0, 0, 0 ; | | | + if kv < kv sup { | | | + v = fa load8<type K K base + size t k VKQ 0 + kv nb11, i0 ; | | | + } | | | + int4 &Ks kv D + KPAD + i0 = v; | | | + } | | | + syncthreads ; | | | + | | | + // ---- S = Q K^T, one tile per warp ---- | | | + fa float8 Sacc = fa float8 0.0f ; | | | + pragma unroll | | | + for int k0 = k half D/KSPLIT ; k0 < k half + 1 D/KSPLIT ; k0 += 16 { | | | + const fa half16 a = fa load frag &Qs qb 16 + l16 D + KPAD + k0 ; | | | + const fa half16 b = fa load frag &Ks s kb 16 + l16 D + KPAD + k0 ; | | | + Sacc = builtin amdgcn wmma f32 16x16x16 f16 w32 a, b, Sacc ; | | | + } | | | + // combine the KSPLIT partial head ranges; K is dead now so its buffer is scratch | | | + if constexpr KSPLIT 1 { | | | + float Sred = float Ks; | | | + syncthreads ; | | | + if k half = 0 { | | | + pragma unroll | | | + for int l = 0; l < 8; ++l { Sred s tile 32 + lane 8 + l = Sacc l ; } | | | + } | | | + syncthreads ; | | | + if k half == 0 { | | | + pragma unroll | | | + for int l = 0; l < 8; ++l { Sacc l += Sred s tile 32 + lane 8 + l ; } | | | + pragma unroll | | | + for int l = 0; l < 8; ++l { Sred s tile 32 + lane 8 + l = Sacc l ; } | | | + } | | | + syncthreads ; | | | + if k half = 0 { | | | + pragma unroll | | | + for int l = 0; l < 8; ++l { Sacc l = Sred s tile 32 + lane 8 + l ; } | | | + } | | | + } | | | + // slot l - block column jc = qb 16 + 2 l + lhi, kv = s kb 16 + l16 | | | + | | | + // ---- mask / out-of-bounds ---- | | | + const int kv loc = s kb 16 + l16; | | | + pragma unroll | | | + for int l = 0; l < 8; ++l { | | | + const int jc = qb 16 + 2 l + lhi; | | | + float s = Sacc l ; | | | + if kv loc = kv sup { | | | + s = -FLT MAX/2.0f; | | | + } else if maskh { | | | + s += half2float maskh size t col Q 0 + jc/ncols2 stride mask + k VKQ 0 + kv loc ; | | | + } | | | + Sacc l = s; | | | + } | | | + | | | + // ---- row max: reduce over the 16 lanes of the subgroup, then over NKB warps ---- | | | + float m t 8 ; | | | + pragma unroll | | | + for int l = 0; l < 8; ++l { m t l = Sacc l ; } | | | + pragma unroll | | | + for int off = 1; off < 16; off <<= 1 { | | | + pragma unroll | | | + for int l = 0; l < 8; ++l { m t l = fmaxf m t l , shfl xor sync 0xFFFFFFFF, m t l , off, WS ; } | | | + } | | | + if l16 == 0 && k half == 0 { | | | + pragma unroll | | | + for int l = 0; l < 8; ++l { Rm s tile 16 + 2 l + lhi = m t l ; } | | | + } | | | + syncthreads ; | | | + | | | + float m new 8 , rescale 8 ; | | | + pragma unroll | | | + for int l = 0; l < 8; ++l { | | | + float mx = Rm qb NKB + 0 16 + 2 l + lhi ; | | | + pragma unroll | | | + for int w = 1; w < NKB; ++w { mx = fmaxf mx, Rm qb NKB + w 16 + 2 l + lhi ; } | | | + m new l = fmaxf m run l , mx ; | | | + rescale l = expf m run l - m new l ; | | | + m run l = m new l ; | | | + } | | | + | | | + // ---- P = exp S - m , row sums, P into LDS ---- | | | + float l t 8 ; | | | + pragma unroll | | | + for int l = 0; l < 8; ++l { | | | + const float p = expf Sacc l - m new l ; | | | + l t l = p; | | | + if k half == 0 { | | | + Ps qb 16 + 2 l + lhi FAB + FPAD + s kb 16 + l16 = Float16 p; | | | + } | | | + } | | | + pragma unroll | | | + for int off = 1; off < 16; off <<= 1 { | | | + pragma unroll | | | + for int l = 0; l < 8; ++l { l t l += shfl xor sync 0xFFFFFFFF, l t l , off, WS ; } | | | + } | | | + if l16 == 0 && k half == 0 { | | | + pragma unroll | | | + for int l = 0; l < 8; ++l { Rl s tile 16 + 2 l + lhi = l t l ; } | | | + } | | | + | | | + pragma unroll | | | + for int t = 0; t < OT; ++t { | | | + pragma unroll | | | + for int l = 0; l < 8; ++l { Oacc t l = rescale l ; } | | | + } | | | + syncthreads ; // Rl ready and K no longer needed | | | + | | | + pragma unroll | | | + for int l = 0; l < 8; ++l { | | | + float sum = 0.0f; | | | + pragma unroll | | | + for int w = 0; w < NKB; ++w { sum += Rl qb NKB + w 16 + 2 l + lhi ; } | | | + l run l = l run l rescale l + sum; | | | + } | | | + | | | + // ---- V transposed into the K buffer: Vt dv kv ---- | | | + // one lane owns 8 consecutive kv of a single dv, so the LDS write is one int4 | | | + for int w = tid; w < FAB/8 D; w += NTHR { | | | + const int dv = w % D; | | | + const int kc = w / D; | | | + fa frag16 g; | | | + pragma unroll | | | + for int j = 0; j < 8; ++j { | | | + const int kv = kc 8 + j; | | | + Float16 &g.v 0 j = kv < kv sup | | | + ? fa load1<type V V base + size t k VKQ 0 + kv nb21, dv : Float16 0.0f; | | | + } | | | + int4 &Vt size t dv FAB + FPAD + kc 8 = g.v 0 ; | | | + } | | | + syncthreads ; | | | + | | | + // ---- O += P V ---- | | | + pragma unroll | | | + for int t = 0; t < OT; ++t { | | | + const int d0 = o d0 + t 16; | | | + pragma unroll | | | + for int k0 = 0; k0 < FAB; k0 += 16 { | | | + const fa half16 a = fa load frag &Ps qb 16 + l16 FAB + FPAD + k0 ; | | | + const fa half16 b = fa load frag &Vt size t d0 + l16 FAB + FPAD + k0 ; | | | + Oacc t = builtin amdgcn wmma f32 16x16x16 f16 w32 a, b, Oacc t ; | | | + } | | | + } | | | + syncthreads ; | | | + } | | | + | | | + // ---- write back ---- | | | + pragma unroll | | | + for int l = 0; l < 8; ++l { | | | + const int jc = qb 16 + 2 l + lhi; | | | + const int j = jc / ncols2; | | | + const int c = jc % ncols2; | | | + | | | + if ncols1 1 && col Q 0 + j = int ne01.z { | | | + continue; | | | + } | | | + | | | + const float out scale = gridDim.y == 1 ? 1.0f/l run l : 1.0f; | | | + const int j dst unrolled = sequence int ne01.z + col Q 0 + j ne02 + head0 + c gridDim.y + blockIdx.y; | | | + | | | + pragma unroll | | | + for int t = 0; t < OT; ++t { | | | + dst ptr size t j dst unrolled DV + o d0 + t 16 + l16 = Oacc t l out scale; | | | + } | | | + | | | + if gridDim.y = 1 && l16 == 0 && o d0 == 0 { | | | + dst meta ptr j dst unrolled = make float2 m run l , l run l ; | | | + } | | | + } | | | +} | | | + | | | + else // AMD WMMA AVAILABLE && RDNA3 | | | + | | | +template<int DKQ, int DV, int ncols1, int ncols2, ggml type type K, ggml type type V | | | + launch bounds FA WMMA D256 NWARPS 32, 1 | | | +static global void flash attn wmma d256 | | | + const char , const char , const char , const char , const char , | | | + const int , float , float2 , | | | + const float, const float, const float, const float, const uint32 t, const float, | | | + const int32 t, const uint3, const int32 t, const int32 t, | | | + const int32 t, const int32 t, const int32 t, | | | + const int32 t, const int32 t, const int32 t, const int32 t, | | | + const int32 t, const int32 t, const int64 t, | | | + const int32 t, const int32 t, const int64 t, | | | + const int32 t, const int32 t, const int32 t, | | | + const int32 t, const int32 t, const int64 t { | | | + NO DEVICE CODE; | | | +} | | | + | | | + endif // AMD WMMA AVAILABLE && RDNA3 | | | + | | | +void ggml cuda flash attn ext wmma d256 ggml backend cuda context & ctx, ggml tensor dst ; | | | diff --git a/ggml/src/ggml-cuda/fattn.cu b/ggml/src/ggml-cuda/fattn.cu | | | index ab7a3b297..f95c57b01 100644 | | | --- a/ggml/src/ggml-cuda/fattn.cu | | | +++ b/ggml/src/ggml-cuda/fattn.cu | | | @@ -3,7 +3,9 @@ | | | include "fattn-mma-f16.cuh" | | | include "fattn-tile.cuh" | | | include "fattn-vec.cuh" | | | + include "fattn-wmma-d256.cuh" | | | include "fattn.cuh" | | | + include <cstdlib | | | template <int DKQ, int DV, int ncols2 | | | static void ggml cuda flash attn ext mma f16 switch ncols1 ggml backend cuda context & ctx, ggml tensor dst { | | | @@ -333,6 +335,7 @@ enum best fattn kernel { | | | BEST FATTN KERNEL TILE = 200, | | | BEST FATTN KERNEL VEC = 100, | | | BEST FATTN KERNEL MMA F16 = 400, | | | + BEST FATTN KERNEL WMMA D256 = 500, | | | }; | | | static bool ggml cuda fattn kv type supported ggml type type { | | | @@ -511,6 +514,28 @@ static best fattn kernel ggml cuda get best fattn kernel const int device, const | | | } | | | } | | | + // gfx1151, head size 256: the tile kernel emits no WMMA at all there, so use the | | | + // dedicated DV-split WMMA kernel. It needs 64 columns per block to pay for its LDS | | | + // footprint, so only for batches that can fill them, and only for the plain f16 case. | | | + float fa logit softcap = 0.0f; | | | + memcpy &fa logit softcap, const float KQV- op params + 2, sizeof float ; | | | + // Opt-in: 1.3-1.4x the tile kernel on short/medium KV in isolation, but its 60 kB LDS | | | + // footprint fits only 2 workgroups per WGP against the tile kernel's 3, and at long | | | + // context that costs more than the matrix units gain. See docs/exp205. | | | + // | | | + // Quantised K/V are read natively no pre-conversion to f16 so that the DRAM traffic | | | + // actually drops. Measured: it does not pay. q8 0 and q4 0 are 1.7-1.8x slower than | | | + // f16 here, because this kernel is bound by instruction issue and memory latency rather | | | + // than by bytes, and per-element dequantisation costs more than the bytes it saves. | | | + if GGML CUDA CC IS RDNA3 5 cc && getenv "GGML CUDA FA WMMA256" | | | + && Q- ne 0 == 256 && V- ne 0 == 256 | | | + && K- type == V- type | | | + && K- type == GGML TYPE F16 || K- type == GGML TYPE Q8 0 || K- type == GGML TYPE Q4 0 | | | + && mask && dst- src 4 && max bias == 0.0f && fa logit softcap == 0.0f | | | + && gqa opt applies && Q- ne 1 gqa ratio eff = FA WMMA D256 NC { | | | + return BEST FATTN KERNEL WMMA D256; | | | + } | | | + | | | // 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 { | | | return BEST FATTN KERNEL MMA F16; | | | @@ -553,6 +578,8 @@ size t ggml cuda flash attn ext get alloc size int device, const ggml tensor d | | | need f16 K = true; | | | need f16 V = true; | | | break; | | | + case BEST FATTN KERNEL WMMA D256: | | | + break; // reads K/V natively, quantised or not | | | case BEST FATTN KERNEL VEC: | | | need f16 K = K- type == GGML TYPE F32; | | | need f16 V = V- type == GGML TYPE F32; | | | @@ -567,8 +594,35 @@ size t ggml cuda flash attn ext get alloc size int device, const ggml tensor d | | | return f16 extra.end - uintptr t dst- data; | | | } | | | -void ggml cuda flash attn ext ggml backend cuda context & ctx, ggml tensor dst { | | | - ggml cuda set device ctx.device ; | | | +// Max Q columns per flash-attention launch, 0 to disable splitting. | | | +// One launch with many Q columns puts many column tiles on the GPU at the same time. The tiles | | | +// then read different parts of K/V and stop sharing them in cache, which only hurts when the KV | | | +// cache is much larger than the last level cache. Splitting Q keeps the number of concurrent | | | +// tiles low. On gfx1151 with ubatch 2048 this is 1.4x faster at 32k context and 1.8x at 64k. | | | +static int ggml cuda fattn q chunk const ggml tensor dst { | | | + static const int chunk = getenv "GGML CUDA FA QCHUNK" ? atoi getenv "GGML CUDA FA QCHUNK" : 256; | | | + | | | + if chunk <= 0 || GGML CUDA CC IS RDNA3 5 ggml cuda info .devices ggml cuda get device .cc { | | | + return 0; | | | + } | | | + | | | + // Only worth it if K/V does not fit in cache anyway, else the extra launches just add overhead. | | | + static const int min kv = getenv "GGML CUDA FA QCHUNK MINKV" ? atoi getenv "GGML CUDA FA QCHUNK MINKV" : 8192; | | | + | | | + if dst- src 1 - ne 1 < min kv { | | | + return 0; | | | + } | | | + | | | + // Q rows are dst dimension 2. If dst has more than one dimension 3 then a row range of dst is | | | + // not one continuous block, but the kernels write dst as continuous memory. Do not split. | | | + if dst- ne 3 = 1 { | | | + return 0; | | | + } | | | + | | | + return chunk; | | | +} | | | + | | | +static void ggml cuda flash attn ext dispatch ggml backend cuda context & ctx, ggml tensor dst { | | | switch ggml cuda get best fattn kernel ggml cuda get device , dst { | | | case BEST FATTN KERNEL NONE: | | | GGML ABORT "fatal error" ; | | | @@ -581,6 +635,50 @@ void ggml cuda flash attn ext ggml backend cuda context & ctx, ggml tensor dst | | | case BEST FATTN KERNEL MMA F16: | | | ggml cuda flash attn ext mma f16 ctx, dst ; | | | break; | | | + case BEST FATTN KERNEL WMMA D256: | | | + ggml cuda flash attn ext wmma d256 ctx, dst ; | | | + break; | | | + } | | | +} | | | + | | | +void ggml cuda flash attn ext ggml backend cuda context & ctx, ggml tensor dst { | | | + ggml cuda set device ctx.device ; | | | + | | | + const ggml tensor Q = dst- src 0 ; | | | + const ggml tensor mask = dst- src 3 ; | | | + | | | + const int chunk = ggml cuda fattn q chunk dst ; | | | + if chunk == 0 || Q- ne 1 <= chunk { | | | + ggml cuda flash attn ext dispatch ctx, dst ; | | | + return; | | | + } | | | + | | | + // Q rows are dst rows in dimension 2, mask rows are in dimension 1. | | | + ggml tensor Q chunk = Q; | | | + ggml tensor mask chunk = mask ? mask : ggml tensor ; | | | + ggml tensor dst chunk = dst; | | | + | | | + dst chunk.src 0 = &Q chunk; | | | + if mask { | | | + dst chunk.src 3 = &mask chunk; | | | + } | | | + | | | + for int64 t i0 = 0; i0 < Q- ne 1 ; i0 += chunk { | | | + const int64 t ncols = std::min<int64 t chunk, Q- ne 1 - i0 ; | | | + | | | + Q chunk.ne 1 = ncols; | | | + Q chunk.data = char Q- data + i0 Q- nb 1 ; | | | + | | | + // Keep all remaining mask rows so that reads past ncols stay in bounds. | | | + if mask { | | | + mask chunk.ne 1 = mask- ne 1 - i0; | | | + mask chunk.data = char mask- data + i0 mask- nb 1 ; | | | + } | | | + | | | + dst chunk.ne 2 = ncols; | | | + dst chunk.data = char dst- data + i0 dst- nb 2 ; | | | + | | | + ggml cuda flash attn ext dispatch ctx, &dst chunk ; | | | } | | | } | | | diff --git a/ggml/src/ggml-cuda/gated delta net.cu b/ggml/src/ggml-cuda/gated delta net.cu | | | index 1b431a724..97ab34724 100644 | | | --- a/ggml/src/ggml-cuda/gated delta net.cu | | | +++ b/ggml/src/ggml-cuda/gated delta net.cu | | | @@ -1,6 +1,17 @@ | | | include "gated delta net.cuh" | | | + include <cstdlib | | | include "ggml-cuda/common.cuh" | | | +// Lanes cooperating on one state column in the column-parallel kernel. | | | +// 0 disables it and falls back to the warp-per-column kernel. Overridable for tuning. | | | +static int ggml cuda gdn split { | | | + static const int split = { | | | + const char s = getenv "GGML CUDA GDN SPLIT" ; | | | + return s ? atoi s : 2; | | | + } ; | | | + return split; | | | +} | | | + | | | template <int S v, bool KDA, bool keep rs t | | | global void launch bounds ggml cuda get physical warp size < S v ? ggml cuda get physical warp size : S v 4, 2 | | | gated delta net cuda const float q, | | | @@ -166,6 +177,152 @@ gated delta net cuda const float q, | | | } | | | } | | | +// Column-parallel variant of the kernel above. | | | +// | | | +// The warp-per-column kernel gives one warp one state column and splits the column's | | | +// rows across the warp's lanes. Every warp therefore reads the head's whole q and k | | | +// vector, and with the columns spread over S v/nwarps blocks that means q and k get | | | +// re-read S v times per head per token. That redundancy, not the arithmetic, is what | | | +// bounds the kernel on prefill. | | | +// | | | +// Here a block owns all S v columns of one head, so q/k/g are read exactly once per | | | +// head per token into LDS. Each column is handled by SPLIT adjacent lanes, so a thread | | | +// holds S v/SPLIT state values in registers and the row reduction costs log2 SPLIT | | | +// shuffles instead of log2 warp size . SPLIT trades registers and hence occupancy | | | +// against those shuffles. | | | +template <int S v, int SPLIT, bool KDA, bool keep rs t | | | + global void launch bounds S v SPLIT, 1 | | | +gated delta net col cuda const float restrict q, | | | + const float restrict k, | | | + const float restrict v, | | | + const float restrict g, | | | + const float restrict beta, | | | + const float restrict curr state, | | | + float restrict dst, | | | + float restrict state, | | | + int64 t H, | | | + int64 t n tokens, | | | + int64 t n seqs, | | | + int64 t sq1, int64 t sq2, int64 t sq3, | | | + int64 t sv1, int64 t sv2, int64 t sv3, | | | + int64 t sb1, int64 t sb2, int64 t sb3, | | | + const uint3 neqk1 magic, | | | + const uint3 rq3 magic, | | | + float scale, | | | + int64 t state slot stride, | | | + int K { | | | + GGML UNUSED n seqs ; | | | + constexpr int RPT = S v / SPLIT; // state rows held per thread | | | + constexpr int warp size = ggml cuda get physical warp size ; | | | + static assert S v % SPLIT == 0, "SPLIT must divide S v" ; | | | + static assert SPLIT <= warp size, "the row reduction must stay inside one wave" ; | | | + | | | + const uint32 t h idx = blockIdx.x; | | | + const uint32 t sequence = blockIdx.y; | | | + const int tid = threadIdx.x; | | | + const int col = tid / SPLIT; // SPLIT adjacent lanes share a column | | | + const int part = tid % SPLIT; | | | + const int row0 = part RPT; | | | + | | | + const uint32 t iq1 = fastmodulo h idx, neqk1 magic ; | | | + const uint32 t iq3 = fastdiv sequence, rq3 magic ; | | | + | | | + float attn data = dst + sequence n tokens H + h idx S v; | | | + | | | + const int64 t state in offset = sequence H S v S v + h idx S v S v; | | | + const int64 t state out offset = sequence H + h idx S v S v; | | | + state += state out offset; | | | + curr state += state in offset + col S v + row0; | | | + | | | + // state is stored transposed: M col i = S i col , so row col is contiguous | | | + float s col RPT ; | | | + | | | + // q/k and the per-row gate for KDA are identical for every column, so stage them in | | | + // LDS once per token. Double-buffered so one barrier per token is enough. | | | + shared float sk 2 S v ; | | | + shared float sq 2 S v ; | | | + shared float sg 2 KDA ? S v : 1 ; | | | + | | | + ggml cuda pdl sync ; | | | + pragma unroll | | | + for int j = 0; j < RPT; j++ { | | | + s col j = curr state j ; | | | + } | | | + | | | + for int t = 0; t < n tokens; t++ { | | | + const float restrict q t = q + iq3 sq3 + t sq2 + iq1 sq1; | | | + const float restrict k t = k + iq3 sq3 + t sq2 + iq1 sq1; | | | + const float restrict v t = v + sequence sv3 + t sv2 + h idx sv1; | | | + | | | + const int64 t gb offset = sequence sb3 + t sb2 + h idx sb1; | | | + const float beta val = beta gb offset ; | | | + const float restrict g t = g + gb offset KDA ? S v : 1 ; | | | + | | | + const int buf = t & 1; | | | + if tid < S v { | | | + sk buf tid = k t tid ; | | | + sq buf tid = q t tid ; | | | + if KDA { | | | + sg buf tid = expf g t tid ; | | | + } | | | + } | | | + syncthreads ; | | | + | | | + const float restrict lk = sk buf + row0; | | | + const float restrict lq = sq buf + row0; | | | + const float restrict lg = sg buf + KDA ? row0 : 0 ; | | | + | | | + const float g val = KDA ? 0.0f : expf g t ; | | | + | | | + float kv = 0.0f; | | | + pragma unroll | | | + for int j = 0; j < RPT; j++ { | | | + kv = KDA ? fmaf lg j s col j , lk j , kv : fmaf s col j , lk j , kv ; | | | + } | | | + pragma unroll | | | + for int o = 1; o < SPLIT; o <<= 1 { | | | + kv += shfl xor sync 0xFFFFFFFF, kv, o, warp size ; | | | + } | | | + | | | + const float delta = v t col - KDA ? kv : g val kv beta val; | | | + | | | + float attn = 0.0f; | | | + pragma unroll | | | + for int j = 0; j < RPT; j++ { | | | + s col j = fmaf KDA ? lg j : g val, s col j , lk j delta ; | | | + attn = fmaf s col j , lq j , attn ; | | | + } | | | + pragma unroll | | | + for int o = 1; o < SPLIT; o <<= 1 { | | | + attn += shfl xor sync 0xFFFFFFFF, attn, o, warp size ; | | | + } | | | + | | | + if part == 0 { | | | + attn data col = attn scale; | | | + } | | | + attn data += S v H; | | | + | | | + if constexpr keep rs t { | | | + const int target slot = int n tokens - 1 - t; | | | + if target slot = 0 && target slot < K { | | | + float slot = state + target slot state slot stride + col S v + row0; | | | + pragma unroll | | | + for int j = 0; j < RPT; j++ { | | | + slot j = s col j ; | | | + } | | | + } | | | + } | | | + } | | | + | | | + if constexpr keep rs t { | | | + float out = state + col S v + row0; | | | + pragma unroll | | | + for int j = 0; j < RPT; j++ { | | | + out j = s col j ; | | | + } | | | + } | | | +} | | | + | | | template <bool KDA, bool keep rs t | | | static void launch gated delta net | | | const float q d, const float k d, const float v d, | | | @@ -186,6 +343,42 @@ static void launch gated delta net | | | const uint3 neqk1 magic = init fastdiv values neqk1 ; | | | const uint3 rq3 magic = init fastdiv values rq3 ; | | | + // Column-parallel kernel: one block per head, sequence , all S v columns inside it. | | | + // Needs at least a full wave per block, so only used from S v = 64 up. | | | + { | | | + const int gdn split = ggml cuda gdn split ; | | | + // For a single token the warp-per-column kernel wins: there is no q/k reuse to | | | + // recover and its narrower blocks start faster. | | | + if gdn split 0 && n tokens = 16 && S v == 64 || S v == 128 { | | | + const dim3 grid col H, n seqs, 1 ; | | | + const dim3 block col S v gdn split, 1, 1 ; | | | + const ggml cuda kernel launch params lp = ggml cuda kernel launch params grid col, block col, 0, stream ; | | | + define GDN COL LAUNCH SV, SP \ | | | + ggml cuda kernel launch gated delta net col cuda<SV, SP, KDA, keep rs t , lp, \ | | | + q d, k d, v d, g d, b d, s d, dst d, state d, H, \ | | | + n tokens, n seqs, sq1, sq2, sq3, sv1, sv2, sv3, \ | | | + sb1, sb2, sb3, neqk1 magic, rq3 magic, scale, state slot stride, K | | | + if S v == 64 { | | | + switch gdn split { | | | + case 1: GDN COL LAUNCH 64, 1 ; return; | | | + case 2: GDN COL LAUNCH 64, 2 ; return; | | | + case 4: GDN COL LAUNCH 64, 4 ; return; | | | + case 8: GDN COL LAUNCH 64, 8 ; return; | | | + default: break; | | | + } | | | + } else { | | | + switch gdn split { | | | + case 1: GDN COL LAUNCH 128, 1 ; return; | | | + case 2: GDN COL LAUNCH 128, 2 ; return; | | | + case 4: GDN COL LAUNCH 128, 4 ; return; | | | + case 8: GDN COL LAUNCH 128, 8 ; return; | | | + default: break; | | | + } | | | + } | | | + undef GDN COL LAUNCH | | | + } | | | + } | | | + | | | const ggml cuda kernel launch params launch params = ggml cuda kernel launch params grid dims, block dims, 0, stream ; | | | switch S v { | | | case 16: | | | diff --git a/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh b/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh | | | index 180b2d937..9838fd175 100644 | | | --- a/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh | | | +++ b/ggml/src/ggml-cuda/mmq-config-rdna3-5.cuh | | | @@ -1,290 +1,290 @@ | | | static constexpr host device ggml cuda mmq config ggml cuda mmq get config rdna3 5 ggml type type, int J, bool fallback { | | | CASE GGML TYPE Q1 0, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q1 0, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q1 0, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q1 0, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q1 0, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q1 0, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q1 0, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q1 0, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q1 0, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q1 0, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q1 0, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q1 0, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q1 0, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q1 0, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q1 0, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q1 0, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q1 0, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q1 0, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q1 0, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q1 0, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q2 0, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q2 0, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q2 0, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q2 0, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q2 0, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q2 0, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q2 0, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q2 0, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q2 0, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q2 0, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q2 0, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q2 0, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q2 0, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q2 0, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q2 0, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q2 0, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q2 0, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q2 0, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q2 0, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q2 0, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q4 0, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q4 0, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q4 0, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q4 0, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q4 0, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q4 0, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q4 0, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q4 0, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q4 0, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q4 0, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q4 0, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q4 0, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q4 0, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q4 0, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 0, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 0, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 0, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 0, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 0, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 0, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q4 1, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q4 1, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q4 1, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q4 1, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q4 1, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q4 1, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q4 1, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q4 1, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q4 1, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q4 1, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q4 1, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q4 1, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q4 1, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q4 1, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 1, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 1, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 1, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 1, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 1, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 1, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q5 0, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q5 0, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q5 0, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q5 0, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q5 0, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q5 0, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q5 0, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q5 0, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 0, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 0, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 0, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 0, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 0, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 0, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 0, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 0, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 0, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 0, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 0, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 0, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q5 1, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q5 1, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q5 1, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q5 1, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q5 1, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q5 1, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q5 1, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q5 1, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 1, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 1, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 1, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 1, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 1, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 1, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 1, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 1, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 1, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 1, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 1, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 1, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q8 0, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q8 0, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q8 0, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q8 0, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q8 0, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q8 0, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q8 0, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q8 0, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q8 0, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q8 0, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q8 0, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q8 0, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q8 0, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q8 0, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q8 0, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q8 0, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q8 0, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q8 0, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q8 0, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q8 0, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | // --------------------------------------------------------------------------------------------- | | | CASE GGML TYPE Q2 K, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q2 K, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q2 K, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q2 K, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q2 K, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q2 K, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q2 K, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q2 K, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q2 K, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q2 K, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q2 K, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q2 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q2 K, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q2 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q2 K, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q2 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q2 K, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q2 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q2 K, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q2 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q2 K, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q2 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q2 K, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q2 K, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q3 K, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q3 K, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q3 K, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q3 K, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q3 K, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q3 K, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q3 K, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q3 K, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q3 K, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q3 K, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q3 K, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q3 K, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q3 K, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q3 K, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q3 K, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q3 K, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q3 K, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q3 K, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q3 K, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q3 K, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q4 K, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q4 K, 128, 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, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q4 K, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q4 K, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q4 K, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q4 K, 128, 2, 64, 32, 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 Q4 K, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q4 K, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q4 K, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q4 K, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 K, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 K, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 K, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 K, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 K, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q4 K, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q5 K, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q5 K, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q5 K, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q5 K, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q5 K, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q5 K, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q5 K, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q5 K, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 K, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 K, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 K, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 K, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 K, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q5 K, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 K, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 K, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 K, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 K, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 K, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q5 K, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q6 K, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q6 K, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q6 K, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, true ; | | | - CASE GGML TYPE Q6 K, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q6 K, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, true ; | | | + CASE GGML TYPE Q6 K, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, true ; | | | CASE GGML TYPE Q6 K, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, false ; | | | CASE GGML TYPE Q6 K, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q6 K, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q6 K, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q6 K, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q6 K, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q6 K, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE Q6 K, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q6 K, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q6 K, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q6 K, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q6 K, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q6 K, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE Q6 K, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q6 K, MMQ ITER K, false, false ; | | | // --------------------------------------------------------------------------------------------- | | | CASE GGML TYPE IQ1 S, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE IQ1 S, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE IQ1 S, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE IQ1 S, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE IQ1 S, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE IQ1 S, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE IQ1 S, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE IQ1 S, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ1 S, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ1 S, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ1 S, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ1 S, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ1 S, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ1 S, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ1 S, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ1 S, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ1 S, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ1 S, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ1 S, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ1 S, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE IQ2 XXS, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE IQ2 XXS, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE IQ2 XXS, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE IQ2 XXS, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE IQ2 XXS, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE IQ2 XXS, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE IQ2 XXS, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE IQ2 XXS, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 XXS, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 XXS, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 XXS, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 XXS, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 XXS, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 XXS, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 XXS, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 XXS, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 XXS, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 XXS, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 XXS, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 XXS, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE IQ2 XS, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | CASE GGML TYPE IQ2 XS, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | - CASE GGML TYPE IQ2 XS, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | - CASE GGML TYPE IQ2 XS, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | + CASE GGML TYPE IQ2 XS, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | + CASE GGML TYPE IQ2 XS, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | CASE GGML TYPE IQ2 XS, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | CASE GGML TYPE IQ2 XS, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 XS, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 XS, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 XS, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 XS, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 XS, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 XS, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 XS, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 XS, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 XS, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 XS, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 XS, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 XS, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | CASE GGML TYPE IQ2 S, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | CASE GGML TYPE IQ2 S, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | - CASE GGML TYPE IQ2 S, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | - CASE GGML TYPE IQ2 S, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | + CASE GGML TYPE IQ2 S, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | + CASE GGML TYPE IQ2 S, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, true ; | | | CASE GGML TYPE IQ2 S, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | CASE GGML TYPE IQ2 S, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 S, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 S, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 S, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 S, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 S, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ2 S, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 S, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 S, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 S, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 S, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 S, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ2 S, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q3 K, MMQ ITER K, false, false ; | | | CASE GGML TYPE IQ3 XXS, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE IQ3 XXS, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE IQ3 XXS, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE IQ3 XXS, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE IQ3 XXS, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE IQ3 XXS, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE IQ3 XXS, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE IQ3 XXS, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ3 XXS, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ3 XXS, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ3 XXS, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ3 XXS, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ3 XXS, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ3 XXS, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ3 XXS, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ3 XXS, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ3 XXS, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ3 XXS, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ3 XXS, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ3 XXS, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE IQ3 S, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE IQ3 S, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE IQ3 S, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE IQ3 S, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE IQ3 S, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE IQ3 S, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE IQ3 S, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE IQ3 S, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ3 S, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ3 S, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ3 S, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ3 S, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ3 S, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ3 S, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ3 S, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ3 S, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ3 S, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ3 S, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ3 S, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ3 S, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE IQ4 XS, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE IQ4 XS, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE IQ4 XS, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE IQ4 XS, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE IQ4 XS, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE IQ4 XS, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE IQ4 XS, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE IQ4 XS, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ4 XS, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ4 XS, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ4 XS, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ4 XS, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ4 XS, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ4 XS, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ4 XS, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ4 XS, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ4 XS, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ4 XS, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ4 XS, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ4 XS, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE IQ4 NL, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE IQ4 NL, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE IQ4 NL, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | - CASE GGML TYPE IQ4 NL, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE IQ4 NL, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | + CASE GGML TYPE IQ4 NL, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, true ; | | | CASE GGML TYPE IQ4 NL, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | CASE GGML TYPE IQ4 NL, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ4 NL, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ4 NL, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ4 NL, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ4 NL, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ4 NL, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | - CASE GGML TYPE IQ4 NL, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ4 NL, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ4 NL, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ4 NL, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ4 NL, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ4 NL, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | + CASE GGML TYPE IQ4 NL, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 0, MMQ ITER K, false, false ; | | | // --------------------------------------------------------------------------------------------- | | | CASE GGML TYPE MXFP4, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | CASE GGML TYPE MXFP4, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | - CASE GGML TYPE MXFP4, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | - CASE GGML TYPE MXFP4, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | + CASE GGML TYPE MXFP4, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | + CASE GGML TYPE MXFP4, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, true ; | | | CASE GGML TYPE MXFP4, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | CASE GGML TYPE MXFP4, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE MXFP4, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE MXFP4, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE MXFP4, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE MXFP4, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE MXFP4, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | - CASE GGML TYPE MXFP4, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE MXFP4, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE MXFP4, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE MXFP4, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE MXFP4, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE MXFP4, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | + CASE GGML TYPE MXFP4, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT Q8 1, MMQ ITER K, false, false ; | | | CASE GGML TYPE NVFP4, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, true ; | | | CASE GGML TYPE NVFP4, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, true ; | | | - CASE GGML TYPE NVFP4, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, true ; | | | - CASE GGML TYPE NVFP4, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, true ; | | | + CASE GGML TYPE NVFP4, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, true ; | | | + CASE GGML TYPE NVFP4, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, true ; | | | CASE GGML TYPE NVFP4, 128, 2, 64, 16, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, false ; | | | CASE GGML TYPE NVFP4, 128, 2, 64, 32, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, false ; | | | - CASE GGML TYPE NVFP4, 256, 2, 128, 48, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, false ; | | | - CASE GGML TYPE NVFP4, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, false ; | | | - CASE GGML TYPE NVFP4, 256, 2, 128, 80, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, false ; | | | - CASE GGML TYPE NVFP4, 256, 2, 128, 96, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, false ; | | | - CASE GGML TYPE NVFP4, 256, 2, 128, 112, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, false ; | | | - CASE GGML TYPE NVFP4, 256, 2, 128, 128, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, false ; | | | + CASE GGML TYPE NVFP4, 128, 2, 64, 48, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, false ; | | | + CASE GGML TYPE NVFP4, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, false ; | | | + CASE GGML TYPE NVFP4, 128, 2, 64, 80, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, false ; | | | + CASE GGML TYPE NVFP4, 128, 2, 64, 96, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, false ; | | | + CASE GGML TYPE NVFP4, 128, 2, 64, 112, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, false ; | | | + CASE GGML TYPE NVFP4, 128, 2, 64, 128, GGML CUDA MMQ SRAM LAYOUT NVFP4, MMQ ITER K, false, false ; | | | - return ggml cuda mmq config GGML TYPE COUNT, 256, 2, 128, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, 256, false, true ; | | | + return ggml cuda mmq config GGML TYPE COUNT, 128, 2, 64, 64, GGML CUDA MMQ SRAM LAYOUT Q8 0, 256, false, true ; | | | } | | | diff --git a/ggml/src/ggml-cuda/mmvq.cu b/ggml/src/ggml-cuda/mmvq.cu | | | index 0589e65bd..3a9d732ca 100644 | | | --- a/ggml/src/ggml-cuda/mmvq.cu | | | +++ b/ggml/src/ggml-cuda/mmvq.cu | | | @@ -63,12 +63,17 @@ static constexpr host device int get vdr mmvq ggml type type { | | | } | | | } | | | + ifndef GGML MMVQ RDNA35 NWARPS | | | + define GGML MMVQ RDNA35 NWARPS 2 | | | + endif | | | + | | | enum mmvq parameter table id { | | | MMVQ PARAMETERS GENERIC = 0, | | | MMVQ PARAMETERS TURING, | | | MMVQ PARAMETERS GCN, | | | MMVQ PARAMETERS RDNA2, | | | MMVQ PARAMETERS RDNA3 0, | | | + MMVQ PARAMETERS RDNA3 5, | | | MMVQ PARAMETERS RDNA4 | | | }; | | | @@ -77,7 +82,9 @@ static constexpr device mmvq parameter table id get device table id { | | | return MMVQ PARAMETERS RDNA4; | | | elif defined RDNA3 0 | | | return MMVQ PARAMETERS RDNA3 0; | | | - elif defined RDNA2 || defined RDNA3 5 | | | + elif defined RDNA3 5 | | | + return MMVQ PARAMETERS RDNA3 5; | | | + elif defined RDNA2 | | | return MMVQ PARAMETERS RDNA2; | | | elif defined GCN || defined CDNA | | | return MMVQ PARAMETERS GCN; | | | @@ -95,7 +102,10 @@ static host mmvq parameter table id get device table id int cc { | | | if GGML CUDA CC IS RDNA3 0 cc { | | | return MMVQ PARAMETERS RDNA3 0; | | | } | | | - if GGML CUDA CC IS RDNA2 cc || GGML CUDA CC IS RDNA3 5 cc { | | | + if GGML CUDA CC IS RDNA3 5 cc { | | | + return MMVQ PARAMETERS RDNA3 5; | | | + } | | | + if GGML CUDA CC IS RDNA2 cc { | | | return MMVQ PARAMETERS RDNA2; | | | } | | | if GGML CUDA CC IS GCN cc || GGML CUDA CC IS CDNA cc { | | | @@ -406,6 +416,31 @@ static constexpr host device int calc nwarps ggml type type, int ncols d | | | } | | | return 1; | | | } | | | + if table id == MMVQ PARAMETERS RDNA3 5 { | | | + // Strix Halo: mat-vec decode is bandwidth bound, and one warp per block does not | | | + // issue enough loads to keep the memory system busy. GGML MMVQ RDNA35 NWARPS | | | + // warps cooperate on a row instead. | | | + if ncols dst == 1 { | | | + switch type { | | | + case GGML TYPE Q4 0: | | | + case GGML TYPE Q4 1: | | | + case GGML TYPE Q5 0: | | | + case GGML TYPE Q5 1: | | | + case GGML TYPE Q8 0: | | | + case GGML TYPE Q2 K: | | | + case GGML TYPE Q3 K: | | | + case GGML TYPE Q4 K: | | | + case GGML TYPE Q5 K: | | | + case GGML TYPE Q6 K: | | | + case GGML TYPE IQ4 NL: | | | + case GGML TYPE IQ4 XS: | | | + return GGML MMVQ RDNA35 NWARPS; | | | + default: | | | + return 1; | | | + } | | | + } | | | + return 1; | | | + } | | | if table id == MMVQ PARAMETERS RDNA3 0 { | | | // RDNA3 W7900 : stricter whitelist than RDNA4. | | | // Q2 K / Q5 K / IQ4 XS regress in full quant sweeps. | | | diff --git a/ggml/src/ggml-cuda/quantize.cu b/ggml/src/ggml-cuda/quantize.cu | | | index bcc772395..dd42db0a0 100644 | | | --- a/ggml/src/ggml-cuda/quantize.cu | | | +++ b/ggml/src/ggml-cuda/quantize.cu | | | @@ -1,4 +1,5 @@ | | | include "quantize.cuh" | | | + include <cstdlib | | | include <cstdint | | | if defined BLACKWELL MMA AVAILABLE | | | @@ -458,19 +459,25 @@ template <mmq q8 1 ds layout ds layout, bool scatter | | | static global void quantize mmq q8 1 | | | const float restrict x, const int32 t restrict ids, void restrict vy, | | | const int64 t ne00, const int64 t s01, const int64 t s02, const int64 t s03, | | | - const int64 t ne0, const int ne1, const int ne2, const int n expert used { | | | + const int64 t ne0, const int ne1, const int ne2, const int n expert used, const int chunks { | | | constexpr int vals per scale = ds layout == MMQ Q8 1 DS LAYOUT D2S6 ? 64 : 32; | | | constexpr int vals per sum = ds layout == MMQ Q8 1 DS LAYOUT D2S6 ? 16 : 32; | | | - const int64 t i0 = int64 t blockDim.x blockIdx.y + threadIdx.x 4; | | | + ggml cuda pdl sync ; | | | + | | | + // Each block walks chunks consecutive column chunks of one row. One chunk is only | | | + // 4 blockDim.x floats, so a chunk per block leaves the reads scattered over ne1 | | | + // separate rows; walking several keeps a longer contiguous run per row. | | | + for int chunk = 0; chunk < chunks; ++chunk { | | | + | | | + const int64 t i0 = int64 t blockDim.x blockIdx.y int64 t chunks + chunk + threadIdx.x 4; | | | if i0 = ne0 { | | | - return; | | | + break; | | | } | | | const int64 t i00 = i0; | | | - ggml cuda pdl sync ; | | | int64 t base idx; | | | if constexpr scatter { | | | @@ -529,7 +536,7 @@ static global void quantize mmq q8 1 | | | const int64 t i = ids int64 t blockIdx.x n expert used + slot ; | | | ib = k block ne1 + i; | | | } else { | | | - const int64 t ib0 = blockIdx.z int64 t gridDim.x gridDim.y blockDim.x/QK8 1 ; // first block of channel | | | + const int64 t ib0 = blockIdx.z ne0/QK8 1 MMQ int64 t ne1; // first block of channel | | | ib = ib0 + k block ne1 + blockIdx.x; | | | } | | | @@ -552,6 +559,8 @@ static global void quantize mmq q8 1 | | | } | | | } | | | } | | | + | | | + } // chunk | | | GGML UNUSED n expert used ; | | | } | | | @@ -572,6 +581,14 @@ void quantize row q8 1 cuda | | | GGML UNUSED type src0 ; | | | } | | | +static int ggml cuda quantize mmq chunks { | | | + static const int chunks = { | | | + const char s = getenv "GGML CUDA QUANTIZE CHUNKS" ; | | | + return s ? atoi s : 2; | | | + } ; | | | + return chunks; | | | +} | | | + | | | void quantize mmq q8 1 cuda | | | const float x, const int32 t ids, void vy, const ggml type type src0, | | | const int64 t ne00, const int64 t s01, const int64 t s02, const int64 t s03, | | | @@ -580,21 +597,23 @@ void quantize mmq q8 1 cuda | | | GGML ASSERT ne0 % QK8 1 MMQ == 0 ; | | | // ne1 tends to assume the highest values, therefore use it as the "x" dimension of the CUDA grid: | | | - const int64 t block num y = ne0 + 4 CUDA QUANTIZE BLOCK SIZE MMQ - 1 / 4 CUDA QUANTIZE BLOCK SIZE MMQ ; | | | + const int64 t block num y full = ne0 + 4 CUDA QUANTIZE BLOCK SIZE MMQ - 1 / 4 CUDA QUANTIZE BLOCK SIZE MMQ ; | | | + const int chunks = ggml cuda quantize mmq chunks ; | | | + const int64 t block num y = block num y full + chunks - 1 / chunks; | | | const dim3 num blocks ne1, block num y, ne2 ne3 ; | | | const dim3 block size CUDA QUANTIZE BLOCK SIZE MMQ, 1, 1 ; | | | switch mmq get q8 1 ds layout type src0 { | | | case MMQ Q8 1 DS LAYOUT D4: | | | quantize mmq q8 1<MMQ Q8 1 DS LAYOUT D4, false | | | - <<<num blocks, block size, 0, stream x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, / n expert used= /0 ; | | | + <<<num blocks, block size, 0, stream x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, / n expert used= /0, chunks ; | | | break; | | | case MMQ Q8 1 DS LAYOUT DS4: | | | quantize mmq q8 1<MMQ Q8 1 DS LAYOUT DS4, false | | | - <<<num blocks, block size, 0, stream x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, / n expert used= /0 ; | | | + <<<num blocks, block size, 0, stream x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, / n expert used= /0, chunks ; | | | break; | | | case MMQ Q8 1 DS LAYOUT D2S6: | | | quantize mmq q8 1<MMQ Q8 1 DS LAYOUT D2S6, false | | | - <<<num blocks, block size, 0, stream x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, / n expert used= /0 ; | | | + <<<num blocks, block size, 0, stream x, ids, vy, ne00, s01, s02, s03, ne0, ne1, ne2, / n expert used= /0, chunks ; | | | break; | | | default: | | | GGML ABORT "fatal error" ; | | | @@ -611,20 +630,21 @@ void quantize scatter mmq q8 1 cuda | | | GGML ASSERT ne0 % QK8 1 MMQ == 0 ; | | | const int64 t block num y = ne0 + 4 CUDA QUANTIZE BLOCK SIZE MMQ - 1 / 4 CUDA QUANTIZE BLOCK SIZE MMQ ; | | | + const int chunks = 1; | | | const dim3 num blocks n tokens, block num y, 1 ; | | | const dim3 block size CUDA QUANTIZE BLOCK SIZE MMQ, 1, 1 ; | | | switch mmq get q8 1 ds layout type src0 { | | | case MMQ Q8 1 DS LAYOUT D4: | | | quantize mmq q8 1<MMQ Q8 1 DS LAYOUT D4, true <<<num blocks, block size, 0, stream | | | - x, ids src1 inv, vy, ne00, / s01= /0, / s02= /stride token, / s03= /0, ne0, / ne1= / int nrows dst, / ne2= /1, n expert used ; | | | + x, ids src1 inv, vy, ne00, / s01= /0, / s02= /stride token, / s03= /0, ne0, / ne1= / int nrows dst, / ne2= /1, n expert used, chunks ; | | | break; | | | case MMQ Q8 1 DS LAYOUT DS4: | | | quantize mmq q8 1<MMQ Q8 1 DS LAYOUT DS4, true <<<num blocks, block size, 0, stream | | | - x, ids src1 inv, vy, ne00, / s01= /0, / s02= /stride token, / s03= /0, ne0, / ne1= / int nrows dst, / ne2= /1, n expert used ; | | | + x, ids src1 inv, vy, ne00, / s01= /0, / s02= /stride token, / s03= /0, ne0, / ne1= / int nrows dst, / ne2= /1, n expert used, chunks ; | | | break; | | | case MMQ Q8 1 DS LAYOUT D2S6: | | | quantize mmq q8 1<MMQ Q8 1 DS LAYOUT D2S6, true <<<num blocks, block size, 0, stream | | | - x, ids src1 inv, vy, ne00, / s01= /0, / s02= /stride token, / s03= /0, ne0, / ne1= / int nrows dst, / ne2= /1, n expert used ; | | | + x, ids src1 inv, vy, ne00, / s01= /0, / s02= /stride token, / s03= /0, ne0, / ne1= / int nrows dst, / ne2= /1, n expert used, chunks ; | | | break; | | | default: | | | GGML ABORT "fatal error" ; |