CUDA + ggml: add sparse-fa for DSV4/GLM by am17an · Pull Request #27970 · ggml-org/llama.cpp

am17an · 2026-08-29T15:56:31Z

Overview

Based on @fairydreaming's PR #25917, it is the same idea but instead just relies on a hint from the API about the max number of live kv-entries per token. Further sparse attention methods like QSA(qwen4) can be enrolled at a later stage as they use similar mechanisms. Performance of 2-bit DSV4 quant on DGX spark. More testing will be appreciated

CPU	Model	Test	t/s `dfc29b6`	t/s sparse-fa	Speedup
CPU	deepseek4 ?B IQ2_XXS - 2.0625 bpw	pp2048	563.31	563.53	1.00
CPU	deepseek4 ?B IQ2_XXS - 2.0625 bpw	pp2048@d4096	525.57	528.31	1.01
CPU	deepseek4 ?B IQ2_XXS - 2.0625 bpw	pp2048@d8192	502.37	520.61	1.04
CPU	deepseek4 ?B IQ2_XXS - 2.0625 bpw	pp2048@d32768	389.50	461.63	1.19
CPU	deepseek4 ?B IQ2_XXS - 2.0625 bpw	pp2048@d131072	204.63	315.26	1.54
CPU	deepseek4 ?B IQ2_XXS - 2.0625 bpw	pp2048@d262144	122.99	216.94	1.76
CPU	deepseek4 ?B IQ2_XXS - 2.0625 bpw	pp2048@d524288	66.09	134.30	2.03
CPU	deepseek4 ?B IQ2_XXS - 2.0625 bpw	pp2048@d1048576	33.99	74.90	2.20
CPU	deepseek4 ?B IQ2_XXS - 2.0625 bpw	tg32	19.00	19.07	1.00
CPU	deepseek4 ?B IQ2_XXS - 2.0625 bpw	tg32@d4096	18.16	18.28	1.01
CPU	deepseek4 ?B IQ2_XXS - 2.0625 bpw	tg32@d8192	18.18	18.26	1.00
CPU	deepseek4 ?B IQ2_XXS - 2.0625 bpw	tg32@d32768	17.01	17.52	1.03
CPU	deepseek4 ?B IQ2_XXS - 2.0625 bpw	tg32@d131072	14.20	15.73	1.11
CPU	deepseek4 ?B IQ2_XXS - 2.0625 bpw	tg32@d262144	11.68	13.94	1.19
CPU	deepseek4 ?B IQ2_XXS - 2.0625 bpw	tg32@d524288	8.66	11.32	1.31
CPU	deepseek4 ?B IQ2_XXS - 2.0625 bpw	tg32@d1048576	5.31	8.30	1.56

Additional information

Requirements

I have read and agree with the contributing guidelines
AI usage disclosure: YES, paired with codex to write the gather kernel

ggerganov

Should we also try instead of extracting indices from the mask, to directly construct a dense KV cache (i.e. get_rows-style) and run the existing FA kernels without modifications?

ggerganov · 2026-08-30T07:03:11Z

    const bool need_f16_V = type_V == GGML_TYPE_F16;
    constexpr size_t nbytes_shared = 0;
-    launch_fattn<D, cols_per_block, 1>(ctx, dst, fattn_kernel, nwarps, nbytes_shared, D, need_f16_K, need_f16_V, false);
+    launch_fattn<D, cols_per_block, 1>(ctx, dst, fattn_kernel, nwarps, nbytes_shared, D, need_f16_K, need_f16_V, false, false);


Where are the "vec" FA kernels located? Just from the filenames, it looks as if the sparse indices are not used in the "vec" case.

Yeah we use the mma kernel for the vec (bs=1) case when f16 kv cache is used. I think the vec kernel is only used for quantized kv cache - I may be wrong

ggerganov · 2026-08-30T07:06:18Z

+static constexpr bool ggml_cuda_flash_attn_ext_mma_f16_may_use_sparse(
+        const int DKQ, const int DV, const int ncols1, const int ncols2) {
+    return (DKQ == 512 && DV == 512 && ncols1 == 1 && ncols2 == 8) ||
+           (DKQ == 576 && DV == 512 && ncols1 == 1 && ncols2 == 16);
+}


Do we need to restrict to only these head sizes?

Could you remind me what ncols1 and ncols2 refer to in the CUDA backend?

ncols1 is the number of query tokens and ncols2 is the k/v tokens per block. It's restricted here because I only ran test this for dsv4, for qwen4 it was slower for prefill and faster for decode but qwen4 has other issues w.r.t to the mask so I left for a future PR

am17an · 2026-08-30T07:38:57Z

Should we also try instead of extracting indices from the mask, to directly construct a dense KV cache (i.e. get_rows-style) and run the existing FA kernels without modifications?

At least for the CUDA backend this was slower than not doing anything at all for dsv4.

Mushoz · 2026-08-30T07:48:27Z

I have done a lot of attempts at optimizing sparse attention in DSv4 in a Vulkan focused llamacpp fork, and the biggest improvements was using dense attention for the continuous block of keys, sparse attention for the selected keys and then ultimately combining the result. The commit for this exact change can be found here: Nathanw1014@4bbe53e

Then I did some further memory usage improvements by changing the implementation to a tiled version in a follow-up commit here: Nathanw1014@80728f3

These changes were vibe coded with oversight from me, but the idea itself might be useful for optimizing sparse attention. My apologies if this PR already does something similar. I am on my phone so I haven't really look at the code yet, but I thought I would share my experiences with optimizing sparse attention when I saw this PR.

ggerganov · 2026-08-31T11:18:13Z

I have a Metal implementation that works with DSv4. But I also want to test it with Qwen4. Do you have a suggestion for a patch that enables sparse attention with the Qwen4 graph?

am17an · 2026-08-31T11:20:59Z

@ggerganov you can just pass n_kv_max (which should be top_k->ne[0]) to the build_mha part? (instead of 0)

ggerganov

I've been running DSv4 with sparse attention enabled (with #28098) and everything seems stable.

ServeurpersoCom · 2026-09-01T13:44:00Z

I run some tests on my container

ServeurpersoCom · 2026-09-01T15:05:18Z

Rebased onto current master and benchmarked against it on an RTX PRO 6000 Blackwell, DeepSeek-V4-Flash IQ2_XXS, f16 KV cache, single sequence. Clear win at long context, small cost below 32k, which matches the shape of your own table.

test	master `8887a48`	sparse-fa
prefill, 153806 token prompt	1050.2 / 1069.9 t/s	1262.8 / 1269.4 t/s
decode after that prompt	57.2 / 57.2 t/s	62.2 / 62.2 t/s
pp2048 @ d0	2056.1 t/s	1999.9 t/s
pp2048 @ d4096	1956.8 t/s	1900.7 t/s
pp2048 @ d16384	1739.6 t/s	1627.9 t/s
pp2048 @ d32768	1496.5 t/s	1503.4 t/s

Two passes per binary with the run order alternated, so the long context numbers are not a thermal artifact.

One correctness note: in flash_attn_mask_to_sparse_indices, ggml_cuda_pdl_lc() is called before the loop that pads the row with -1. flash_attn_ext_f16 calls ggml_cuda_pdl_sync() and PDL is on by default from Hopper up, so the FA kernel can start once launch completion is triggered and read indices[count..n_kv_max) before those -1 land. The buffer comes from the pool, so it would read leftover indices from a previous node, used unbounded as a row offset into K/V and into the mask. flash_attn_mask_to_KV_max never triggers launch completion, which is why it is not exposed. I could not reproduce it, the window is tiny, but moving the pdl_lc() below the padding loop costs nothing.

Patch :

diff --git a/ggml/src/ggml-cuda/fattn.cu b/ggml/src/ggml-cuda/fattn.cu
index d1e6bb4ee..4ed5e8e22 100644
--- a/ggml/src/ggml-cuda/fattn.cu
+++ b/ggml/src/ggml-cuda/fattn.cu
@@ -77,12 +77,15 @@ static __global__ void flash_attn_mask_to_sparse_indices(
         __syncthreads();
     }

-    ggml_cuda_pdl_lc();
-
     const int count = row_count;
     for (int i = count + tid; i < n_kv_max; i += blockDim.x) {
         indices[i] = -1;
     }
+    __syncthreads();
+
+    // the dependent grid reads indices, signal once the row is complete
+    ggml_cuda_pdl_lc();
+
     if (count > n_kv_max) {
         if (tid == 0) {
             printf("flash attention sparse mask row exceeds n_kv_max (%d > %d)\n", count, n_kv_max);

am17an · 2026-09-01T15:45:10Z

@ServeurpersoCom thanks, fixed in be029fd.

ServeurpersoCom · 2026-09-01T15:47:28Z

Optional nits :

Small cleanup patch on top of the PDL fix, nothing behavioural except one extra test case. It names the block size and the two activation thresholds, replaces the serial prefix sum with a warp scan over the eight warp counts, adds a static_assert for the ncols1 == 1 assumption the index addressing relies on, and drops a dead VARS_TO_STR18 plus a test case that was listed twice.

The one real change is in the mask generator: it now reduces one row in every thirty two to a single finite entry, which is the degenerate case of the first token of a sequence and was not covered before. Built and ran the full FLASH_ATTN_EXT suite on a Blackwell card, 2948 pass, one fewer than before because of the removed duplicate.

diff --git a/ggml/src/ggml-cuda/fattn-mma-f16.cuh b/ggml/src/ggml-cuda/fattn-mma-f16.cuh
index bf0d5412a..6ecd47744 100644
--- a/ggml/src/ggml-cuda/fattn-mma-f16.cuh
+++ b/ggml/src/ggml-cuda/fattn-mma-f16.cuh
@@ -1784,6 +1784,8 @@ static __global__ void flash_attn_ext_f16(
                             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) {
+    static_assert(!use_sparse || ncols1 == 1, "sparse indices are addressed with one mask row per query");
+
     ggml_cuda_pdl_sync(); // TODO optimize placement
 #if defined(FLASH_ATTN_AVAILABLE) && (defined(VOLTA_MMA_AVAILABLE) || defined(TURING_MMA_AVAILABLE) || defined(AMD_WMMA_AVAILABLE) || defined(AMD_MFMA_AVAILABLE))
     const char * GGML_CUDA_RESTRICT Q        = Q_ptr;
diff --git a/ggml/src/ggml-cuda/fattn.cu b/ggml/src/ggml-cuda/fattn.cu
index 4ed5e8e22..27218aa03 100644
--- a/ggml/src/ggml-cuda/fattn.cu
+++ b/ggml/src/ggml-cuda/fattn.cu
@@ -6,7 +6,10 @@
 #include "fattn.cuh"

 #if !defined(GGML_USE_HIP) && !defined(GGML_USE_MUSA)
-__launch_bounds__(256, 1)
+static constexpr int sparse_indices_nthreads = 256;
+static constexpr int sparse_indices_nwarps   = sparse_indices_nthreads / WARP_SIZE;
+
+__launch_bounds__(sparse_indices_nthreads, 1)
 static __global__ void flash_attn_mask_to_sparse_indices(
         const half * mask_ptr, int32_t * indices_ptr, const int ne30, const int n_kv_max,
         const int64_t s31, const int64_t s33) {
@@ -22,7 +25,7 @@ static __global__ void flash_attn_mask_to_sparse_indices(
     const half * mask = mask_ptr + sequence*s33 + query*s31;
     int32_t * indices = indices_ptr + (int64_t(sequence)*gridDim.x + query)*n_kv_max;

-    __shared__ int warp_offsets[256/WARP_SIZE];
+    __shared__ int warp_offsets[sparse_indices_nwarps];
     __shared__ int row_count;
     __shared__ int chunk_count;

@@ -47,14 +50,22 @@ static __global__ void flash_attn_mask_to_sparse_indices(
         }
         __syncthreads();

-        if (tid == 0) {
-            int offset = 0;
-            for (int iw = 0; iw < 256/WARP_SIZE; ++iw) {
-                const int count = warp_offsets[iw];
-                warp_offsets[iw] = offset;
-                offset += count;
+        if (warp == 0) {
+            const int count = lane < sparse_indices_nwarps ? warp_offsets[lane] : 0;
+            int prefix = count;
+#pragma unroll
+            for (int offset = 1; offset < sparse_indices_nwarps; offset *= 2) {
+                const int up = __shfl_up_sync(0xFFFFFFFF, prefix, offset);
+                if (lane >= offset) {
+                    prefix += up;
+                }
+            }
+            if (lane < sparse_indices_nwarps) {
+                warp_offsets[lane] = prefix - count;
+            }
+            if (lane == sparse_indices_nwarps - 1) {
+                chunk_count = prefix;
             }
-            chunk_count = offset;
         }
         __syncthreads();

@@ -104,7 +115,7 @@ void ggml_cuda_flash_attn_ext_compact_mask(
     const int64_t s31 = mask->nb[1] / sizeof(half);
     const int64_t s33 = mask->nb[3] / sizeof(half);
     const dim3 blocks_num(mask->ne[1], mask->ne[3], 1);
-    const dim3 block_dim(256, 1, 1);
+    const dim3 block_dim(sparse_indices_nthreads, 1, 1);
     const ggml_cuda_kernel_launch_params launch_params(blocks_num, block_dim, 0, stream);
     ggml_cuda_kernel_launch(flash_attn_mask_to_sparse_indices, launch_params,
         (const half *) mask->data, indices, int(mask->ne[0]), n_kv_max, s31, s33);
@@ -127,11 +138,15 @@ bool ggml_cuda_flash_attn_ext_mma_f16_shall_use_sparse(ggml_backend_cuda_context
     memcpy(&max_bias,      (const float *) dst->op_params + 1, sizeof(float));
     memcpy(&logit_softcap, (const float *) dst->op_params + 2, sizeof(float));

+    // the gather amortizes the mask scan only on a long KV that the sparse set covers sparsely
+    constexpr int64_t min_n_kv     = 4096;
+    constexpr int64_t min_kv_ratio = 2;
+
     const int32_t n_kv_max = ggml_get_op_params_i32(dst, 4);
     return GGML_CUDA_CC_IS_NVIDIA(cc) && turing_mma_available(cc) &&
         mask != nullptr && n_kv_max > 0 && max_bias == 0.0f && logit_softcap == 0.0f &&
         mask->ne[0] == K->ne[1] && mask->ne[1] >= Q->ne[1] && mask->ne[2] == 1 &&
-        K->ne[1] >= std::max<int64_t>(4096, 2LL*n_kv_max);
+        K->ne[1] >= std::max<int64_t>(min_n_kv, min_kv_ratio*n_kv_max);
 #endif
 }

diff --git a/tests/test-backend-ops.cpp b/tests/test-backend-ops.cpp
index ab11acafb..2897fc0c1 100644
--- a/tests/test-backend-ops.cpp
+++ b/tests/test-backend-ops.cpp
@@ -202,10 +202,15 @@ static void init_tensor_kq_mask_sparse(ggml_tensor * tensor, int64_t n_kv_max) {
         order[i] = i;
     }

+    // rows carry a spread of finite counts, and one row per group is reduced to a single entry
+    constexpr int64_t count_spread        = 17;
+    constexpr int64_t single_entry_stride = 32;
+
     std::mt19937 gen(0x5A17);
     for (int64_t row = 0; row < nrows; ++row) {
         std::shuffle(order.begin(), order.end(), gen);
-        const int64_t count = n_kv_max - row % std::min<int64_t>(n_kv_max, 17);
+        const int64_t count = row % single_entry_stride == 0 ?
+            1 : n_kv_max - row % std::min<int64_t>(n_kv_max, count_spread);
         std::sort(order.begin(), order.begin() + count);
         for (int64_t i = 0; i < count; ++i) {
             data_f32[row*ne0 + order[i]] = -0.03125f * (1 + (i + row) % 7);
@@ -461,7 +466,6 @@ static std::string var_to_str(ggml_scale_mode mode) {
 #define VARS_TO_STR15(a, b, c, d, e, f, g, h, i, j, k, l, m, n, o) VAR_TO_STR(a) + "," + VARS_TO_STR14(b, c, d, e, f, g, h, i, j, k, l, m, n, o)
 #define VARS_TO_STR16(a, b, c, d, e, f, g, h, i, j, k, l, m, n, o, p) VAR_TO_STR(a) + "," + VARS_TO_STR15(b, c, d, e, f, g, h, i, j, k, l, m, n, o, p)
 #define VARS_TO_STR17(a, b, c, d, e, f, g, h, i, j, k, l, m, n, o, p, q) VAR_TO_STR(a) + "," + VARS_TO_STR16(b, c, d, e, f, g, h, i, j, k, l, m, n, o, p, q)
-#define VARS_TO_STR18(a, b, c, d, e, f, g, h, i, j, k, l, m, n, o, p, q, r) VAR_TO_STR(a) + "," + VARS_TO_STR17(b, c, d, e, f, g, h, i, j, k, l, m, n, o, p, q, r)

 #ifdef GGML_USE_SYCL
 static bool inline _isinf(float f) {
@@ -10202,7 +10206,6 @@ static std::vector<std::unique_ptr<test_case>> make_test_cases_eval() {
     test_cases.emplace_back(new test_flash_attn_ext(512, 512, 1, { 8, 2}, 4096, 3, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16, {0, 1, 2, 3}, true, false,  768));
     test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {16, 1}, 4096, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16, {0, 1, 2, 3}, true, true,   512));
     test_cases.emplace_back(new test_flash_attn_ext(576, 512, 1, {16, 2}, 4096, 2, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16, {0, 1, 2, 3}, true, true,   768));
-    test_cases.emplace_back(new test_flash_attn_ext(512, 512, 1, { 8, 1}, 4096, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16, {0, 1, 2, 3}, true, false,  512 ));
     test_cases.emplace_back(new test_flash_attn_ext(512, 512, 1, { 8, 1}, 4096, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16, {0, 1, 2, 3}, true, false, 2304));
     test_cases.emplace_back(new test_flash_attn_ext(256, 256, 1, { 8, 1}, 4096, 1, true, false, 0, 0, GGML_PREC_F32, GGML_TYPE_F16, GGML_TYPE_F16, {0, 1, 2, 3}, true, false,  512));

Take or leave any of it, none of it is load bearing. Running a greedy check next and I will approve after that.

JohannesGaessler

Any comments regarding performance are only suggestions.

JohannesGaessler · 2026-09-01T14:55:36Z


+    // Use finite mask entries as a sparse K/V set. Set 0 to disable.
+    // n_kv_max must bound the number of finite entries in every mask row.
+    GGML_API void ggml_flash_attn_ext_set_sparse(


I don't feel strongly about this but wouldn't ggml_flash_attn_ext_set_n_kv_max be the more appropriate name?

JohannesGaessler · 2026-09-01T15:24:45Z

+                        const int32_t index = i < i_sup ? indices[i] : -1;
+                        src = index >= 0 ? KV + int64_t(index)*stride_KV + k*h2_per_chunk : zero;


It's not great to have a condition here. It would likely be preferable to pad indices with safe values that result in some redundant work.

JohannesGaessler · 2026-09-01T15:48:50Z

    // Skip unused kernel variants for faster compilation:
    if (use_logit_softcap && !(DKQ == 128 || DKQ == 256 || DKQ == 512)) {
        NO_DEVICE_CODE;
        return;
    }
    if (DKQ == 192 && ncols2 != 8 && ncols2 != 16) {
        NO_DEVICE_CODE;
        return;
    }
 #ifdef VOLTA_MMA_AVAILABLE
    if (ncols1*ncols2 < 32) {
        NO_DEVICE_CODE;
        return;
    }
 #endif // VOLTA_MMA_AVAILABLE

 #if __CUDA_ARCH__ == GGML_CUDA_CC_TURING
    if (ncols1*ncols2 > 32) {
        NO_DEVICE_CODE;
        return;
    }
 #endif // __CUDA_ARCH__ == GGML_CUDA_CC_TURING

 #if defined(AMD_WMMA_AVAILABLE)
    if (ncols1*ncols2 < 16 || ncols2 == 1 || DKQ > 128) {
        NO_DEVICE_CODE;
        return;
    }
 #endif // defined(AMD_WMMA_AVAILABLE)

 #if defined(AMD_MFMA_AVAILABLE)
    if (ncols1*ncols2 < 16 || DKQ > 256) {
        NO_DEVICE_CODE;
        return;
    }
 #endif // defined(AMD_MFMA_AVAILABLE)


Please skip unused kernel templates here, I would suggest you re-use ggml_cuda_flash_attn_ext_mma_f16_may_use_sparse.

does shall_use_sparse etc already do this?

No, because shall_use_sparse is a host function function intended for kernel selection logic.

In terms of program logic shall_use_sparse must be strictly a subset of may_use_sparse.

Okay it should be there

JohannesGaessler · 2026-09-01T16:00:08Z


        const half2 * V_h2 = V_is_K_view ? K_h2 : (const half2 *) (V + nb23*sequence + nb22*z_KV);
        const float * sinks_f = sinks ? (const float *) sinks + zt_Q : nullptr;
+        const int32_t * indices = use_sparse ? sparse_indices + (int64_t(sequence % ne33)*ne31 + jt)*ne11 : nullptr;


Suggested change

const int32_t * indices = use_sparse ? sparse_indices + (int64_t(sequence % ne33)*ne31 + jt)*ne11 : nullptr;

const int32_t * indices = use_sparse ? sparse_indices + (int64_t(sequence % ne33)*ne31 + jt*ncols1)*ne11 : nullptr;

I think the way you're calculating the indices may be incorrect here though this defect would as of right now not manifest as a bug due to ncols1 == 1.

Sorry, after thinking about it some more I'm not sure what I suggested here is correct. For > 1 tokens it's not clear to me what the best way to handle indices is.

ServeurpersoCom · 2026-09-01T16:21:46Z

Greedy output diverges from master, as expected since compacting the selected entries changes the accumulation order, so I checked perplexity instead. Same 153806 token corpus, run at c=32768, well past the 4096 KV length needed to arm the sparse path:

	chunk 1	chunk 2	chunk 3	chunk 4	final
master `8887a48`	1.1464	1.3313	1.2884	1.2421	1.2421 +/- 0.00451
sparse-fa	1.1455	1.3304	1.2876	1.2415	1.2415 +/- 0.00449

0.0006 apart against a 0.0045 confidence interval, same direction on all four chunks. No sign of anything dropped.

Good for me, I will leave the rest to @JohannesGaessler.

am17an requested review from a team, CISC and ggerganov as code owners August 29, 2026 15:56

github-actions Bot added model Model specific testing Everything test related ggml changes relating to the ggml tensor library for machine learning CUDA Related to the CUDA backend labels Aug 29, 2026

ggerganov reviewed Aug 30, 2026

View reviewed changes

Comment thread tests/test-backend-ops.cpp Outdated

ggerganov mentioned this pull request Aug 31, 2026

metal : add sparse attention (wip) #28098

Draft

0cc4m mentioned this pull request Aug 31, 2026

vulkan: support sparse Flash Attention #28105

Draft

am17an added 3 commits September 1, 2026 12:21

CUDA + ggml: add sparse flash attention

5577abb

fix PDL issues

33a4035

remove sparse hint

92b5b18

am17an force-pushed the sparse-fa branch from a9bed29 to 92b5b18 Compare September 1, 2026 09:48

ggerganov approved these changes Sep 1, 2026

View reviewed changes

cont: fix PDL placement

be029fd

JohannesGaessler reviewed Sep 1, 2026

View reviewed changes

am17an added 3 commits September 2, 2026 00:32

add comments to endif, add unroll, fix formatting

e6b567e

address rest of the review comments

999a81d

perf improvments suggested by Johannes

fb29470

		const int32_t index = i < i_sup ? indices[i] : -1;
		src = index >= 0 ? KV + int64_t(index)stride_KV + kh2_per_chunk : zero;

	const int32_t * indices = use_sparse ? sparse_indices + (int64_t(sequence % ne33)ne31 + jt)ne11 : nullptr;
	const int32_t * indices = use_sparse ? sparse_indices + (int64_t(sequence % ne33)ne31 + jtncols1)*ne11 : nullptr;

Uh oh!

Conversation

am17an commented Aug 29, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Overview

Additional information

Requirements

Uh oh!

ggerganov left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

am17an commented Aug 30, 2026

Uh oh!

Mushoz commented Aug 30, 2026

Uh oh!

Uh oh!

ggerganov commented Aug 31, 2026

Uh oh!

am17an commented Aug 31, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

ggerganov left a comment

Choose a reason for hiding this comment

Uh oh!

ServeurpersoCom commented Sep 1, 2026

Uh oh!

ServeurpersoCom commented Sep 1, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

am17an commented Sep 1, 2026

Uh oh!

ServeurpersoCom commented Sep 1, 2026

Uh oh!

JohannesGaessler left a comment

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Uh oh!

Choose a reason for hiding this comment

Uh oh!

Choose a reason for hiding this comment

Uh oh!

ServeurpersoCom commented Sep 1, 2026 • edited Loading Uh oh! There was an error while loading. Please reload this page.

Uh oh!

Uh oh!

Reviewers

Assignees

Labels

Projects

Milestone

Development

Uh oh!

5 participants

am17an commented Aug 29, 2026 •

edited

Loading

am17an commented Aug 31, 2026 •

edited

Loading

ServeurpersoCom commented Sep 1, 2026 •

edited

Loading

ServeurpersoCom commented Sep 1, 2026 •

edited

Loading