Skip to content

Commit 3cc7d80

Browse files
NexesenexNexesenex
committed
Merge branch 'concedo_experimental' into crokeso
2 parents ba4b822 + 8f34111 commit 3cc7d80

19 files changed

+255
-111
lines changed

‎convert_hf_to_gguf.py‎

Lines changed: 2 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -4121,7 +4121,7 @@ def set_gguf_parameters(self):
41214121
self.gguf_writer.add_block_count(block_count)
41224122
self.gguf_writer.add_head_count(hparams.get("num_attention_heads", 32))
41234123
self.gguf_writer.add_layer_norm_rms_eps(hparams.get("rms_norm_eps", 1e-06))
4124-
self.gguf_writer.add_rope_freq_base(hparams.get("rope_theta", 1000000.0))
4124+
self.gguf_writer.add_rope_freq_base(hparams.get("rope_theta", 10000))
41254125

41264126
# Mamba parameters
41274127
self.gguf_writer.add_ssm_state_size(hparams.get("mamba_d_state", 64))
@@ -4132,7 +4132,7 @@ def set_gguf_parameters(self):
41324132
self.gguf_writer.add_ssm_group_count(0)
41334133

41344134
# MLP feed forward parameters (for attention layers)
4135-
self.gguf_writer.add_feed_forward_length(hparams.get("intermediate_size", 16384))
4135+
self.gguf_writer.add_feed_forward_length(hparams.get("intermediate_size", 13312))
41364136
self.gguf_writer.add_file_type(self.ftype)
41374137

41384138
def modify_tensors(self, data_torch: Tensor, name: str, bid: int | None) -> Iterable[tuple[str, Tensor]]:

‎ggml/src/ggml-cuda/fattn-common.cuh‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -871,6 +871,7 @@ void launch_fattn(
871871
size_t nb23 = V ? V->nb[3] : nb13;
872872

873873
if (need_f16_K && K->type != GGML_TYPE_F16) {
874+
// GGML_ASSERT(ggml_is_contiguously_allocated(K));
874875
K_f16.alloc(ggml_nelements(K));
875876
to_fp16_cuda_t to_fp16 = ggml_get_to_fp16_cuda(K->type);
876877
to_fp16(K_data, K_f16.ptr, 1, ggml_nelements(K), main_stream);

‎ggml/src/ggml-cuda/fattn-mma-f16.cuh‎

Lines changed: 10 additions & 8 deletions
Original file line numberDiff line numberDiff line change
@@ -1365,14 +1365,16 @@ static __global__ void flash_attn_ext_f16(
13651365
ne01, ne02, stride_Q1, stride_Q2, stride_K, stride_V, stride_mask, jt, kb0_start_kernel, kb0_stop_kernel);
13661366
#else
13671367
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask);
1368-
GGML_UNUSED(dst); GGML_UNUSED(dst_meta); GGML_UNUSED(scale);
1369-
GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
1370-
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap); GGML_UNUSED(ne00);
1371-
GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03); GGML_UNUSED(ne10);
1372-
GGML_UNUSED(ne11); GGML_UNUSED(ne12); GGML_UNUSED(ne13); GGML_UNUSED(ne31); GGML_UNUSED(ne32);
1373-
GGML_UNUSED(nb31); GGML_UNUSED(nb32); GGML_UNUSED(nb01); GGML_UNUSED(nb02); GGML_UNUSED(nb03);
1374-
GGML_UNUSED(nb11); GGML_UNUSED(nb12); GGML_UNUSED(nb13); GGML_UNUSED(nb21);
1375-
GGML_UNUSED(nb22); GGML_UNUSED(nb23);
1368+
GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
1369+
GGML_UNUSED(scale); GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
1370+
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
1371+
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);
1372+
GGML_UNUSED(nb01); GGML_UNUSED(nb02); GGML_UNUSED(nb03);
1373+
GGML_UNUSED(ne10); GGML_UNUSED(ne11); GGML_UNUSED(ne12); GGML_UNUSED(ne13);
1374+
GGML_UNUSED(nb11); GGML_UNUSED(nb12); GGML_UNUSED(nb13);
1375+
GGML_UNUSED(nb21); GGML_UNUSED(nb22); GGML_UNUSED(nb23);
1376+
GGML_UNUSED(ne31); GGML_UNUSED(ne32); GGML_UNUSED(ne33);
1377+
GGML_UNUSED(nb31); GGML_UNUSED(nb32); GGML_UNUSED(nb33);
13761378
NO_DEVICE_CODE;
13771379
#endif // defined(FLASH_ATTN_AVAILABLE) && defined(NEW_MMA_AVAILABLE)
13781380
}

‎ggml/src/ggml-cuda/fattn-tile-f32.cu‎

Lines changed: 14 additions & 14 deletions
Original file line numberDiff line numberDiff line change
@@ -37,16 +37,16 @@ static __global__ void flash_attn_tile_ext_f32(
3737
#endif // FP16_MMA_AVAILABLE
3838
if (use_logit_softcap && !(D == 128 || D == 256)) {
3939
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask);
40-
GGML_UNUSED(dst); GGML_UNUSED(dst_meta); GGML_UNUSED(scale);
41-
GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
40+
GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
41+
GGML_UNUSED(scale); GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
4242
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
43-
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02);
44-
GGML_UNUSED(ne03); GGML_UNUSED(ne10); GGML_UNUSED(ne11);
45-
GGML_UNUSED(ne12); GGML_UNUSED(ne13); GGML_UNUSED(ne31); GGML_UNUSED(ne32);
46-
GGML_UNUSED(nb31); GGML_UNUSED(nb32); GGML_UNUSED(nb01); GGML_UNUSED(nb02);
47-
GGML_UNUSED(nb03); GGML_UNUSED(nb11); GGML_UNUSED(nb12);
48-
GGML_UNUSED(nb13); GGML_UNUSED(nb21); GGML_UNUSED(nb22);
49-
GGML_UNUSED(nb23);
43+
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);
44+
GGML_UNUSED(nb01); GGML_UNUSED(nb02); GGML_UNUSED(nb03);
45+
GGML_UNUSED(ne10); GGML_UNUSED(ne11); GGML_UNUSED(ne12); GGML_UNUSED(ne13);
46+
GGML_UNUSED(nb11); GGML_UNUSED(nb12); GGML_UNUSED(nb13);
47+
GGML_UNUSED(nb21); GGML_UNUSED(nb22); GGML_UNUSED(nb23);
48+
GGML_UNUSED(ne31); GGML_UNUSED(ne32); GGML_UNUSED(ne33);
49+
GGML_UNUSED(nb31); GGML_UNUSED(nb32); GGML_UNUSED(nb33);
5050
NO_DEVICE_CODE;
5151
return;
5252
}
@@ -282,16 +282,16 @@ static __global__ void flash_attn_tile_ext_f32(
282282
}
283283
#else
284284
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask);
285-
GGML_UNUSED(dst); GGML_UNUSED(dst_meta); GGML_UNUSED(scale);
286-
GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
285+
GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
286+
GGML_UNUSED(scale); GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
287287
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
288288
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);
289-
GGML_UNUSED(ne10); GGML_UNUSED(ne11); GGML_UNUSED(ne12); GGML_UNUSED(ne13);
290-
GGML_UNUSED(ne31); GGML_UNUSED(ne32);
291-
GGML_UNUSED(nb31); GGML_UNUSED(nb32);
292289
GGML_UNUSED(nb01); GGML_UNUSED(nb02); GGML_UNUSED(nb03);
290+
GGML_UNUSED(ne10); GGML_UNUSED(ne11); GGML_UNUSED(ne12); GGML_UNUSED(ne13);
293291
GGML_UNUSED(nb11); GGML_UNUSED(nb12); GGML_UNUSED(nb13);
294292
GGML_UNUSED(nb21); GGML_UNUSED(nb22); GGML_UNUSED(nb23);
293+
GGML_UNUSED(ne31); GGML_UNUSED(ne32); GGML_UNUSED(ne33);
294+
GGML_UNUSED(nb31); GGML_UNUSED(nb32); GGML_UNUSED(nb33);
295295
NO_DEVICE_CODE;
296296
#endif // FLASH_ATTN_AVAILABLE
297297
}

‎ggml/src/ggml-cuda/fattn-vec-f16.cuh‎

Lines changed: 9 additions & 9 deletions
Original file line numberDiff line numberDiff line change
@@ -329,16 +329,16 @@ static __global__ void flash_attn_vec_ext_f16(
329329
}
330330
#else
331331
GGML_UNUSED(Q); GGML_UNUSED(K); GGML_UNUSED(V); GGML_UNUSED(mask);
332-
GGML_UNUSED(dst); GGML_UNUSED(dst_meta); GGML_UNUSED(scale);
333-
GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
332+
GGML_UNUSED(dst); GGML_UNUSED(dst_meta);
333+
GGML_UNUSED(scale); GGML_UNUSED(max_bias); GGML_UNUSED(m0); GGML_UNUSED(m1);
334334
GGML_UNUSED(n_head_log2); GGML_UNUSED(logit_softcap);
335-
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02);
336-
GGML_UNUSED(ne03); GGML_UNUSED(ne10); GGML_UNUSED(ne11);
337-
GGML_UNUSED(ne12); GGML_UNUSED(ne13); GGML_UNUSED(ne31); GGML_UNUSED(ne32); GGML_UNUSED(ne32);
338-
GGML_UNUSED(nb31); GGML_UNUSED(nb32); GGML_UNUSED(nb33); GGML_UNUSED(nb01); GGML_UNUSED(nb02);
339-
GGML_UNUSED(nb03); GGML_UNUSED(nb11); GGML_UNUSED(nb12);
340-
GGML_UNUSED(nb13); GGML_UNUSED(nb21); GGML_UNUSED(nb22);
341-
GGML_UNUSED(nb23);
335+
GGML_UNUSED(ne00); GGML_UNUSED(ne01); GGML_UNUSED(ne02); GGML_UNUSED(ne03);
336+
GGML_UNUSED(nb01); GGML_UNUSED(nb02); GGML_UNUSED(nb03);
337+
GGML_UNUSED(ne10); GGML_UNUSED(ne11); GGML_UNUSED(ne12); GGML_UNUSED(ne13);
338+
GGML_UNUSED(nb11); GGML_UNUSED(nb12); GGML_UNUSED(nb13);
339+
GGML_UNUSED(nb21); GGML_UNUSED(nb22); GGML_UNUSED(nb23);
340+
GGML_UNUSED(ne31); GGML_UNUSED(ne32); GGML_UNUSED(ne33);
341+
GGML_UNUSED(nb31); GGML_UNUSED(nb32); GGML_UNUSED(nb33);
342342
NO_DEVICE_CODE;
343343
#endif // defined(FLASH_ATTN_AVAILABLE) && defined(FP16_AVAILABLE)
344344
}

‎ggml/src/ggml-cuda/set-rows.cu‎

Lines changed: 3 additions & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -44,6 +44,9 @@ static __global__ void k_set_rows_quant(
4444
block_type * dst_block = dst_row_ptr + i00 / qk;
4545

4646
quantize_func(src_block, dst_block);
47+
48+
GGML_UNUSED(ne10);
49+
GGML_UNUSED(ne13);
4750
}
4851

4952
// Template dispatch function for quantized set_rows

‎ggml/src/ggml-metal/ggml-metal-impl.h‎

Lines changed: 1 addition & 0 deletions
Original file line numberDiff line numberDiff line change
@@ -528,6 +528,7 @@ typedef struct {
528528
int64_t n_group;
529529
int64_t n_seq_tokens;
530530
int64_t n_seqs;
531+
int64_t s_off;
531532
uint64_t nb01;
532533
uint64_t nb02;
533534
uint64_t nb03;

‎ggml/src/ggml-metal/ggml-metal.m‎

Lines changed: 13 additions & 2 deletions
Original file line numberDiff line numberDiff line change
@@ -3141,6 +3141,7 @@ static int ggml_metal_encode_node(
31413141
/*.n_group =*/ n_group,
31423142
/*.n_seq_tokens =*/ n_seq_tokens,
31433143
/*.n_seqs =*/ n_seqs,
3144+
/*.s_off =*/ ggml_nelements(src1) * sizeof(float),
31443145
/*.nb01 =*/ nb01,
31453146
/*.nb02 =*/ nb02,
31463147
/*.nb03 =*/ nb03,
@@ -3169,12 +3170,22 @@ static int ggml_metal_encode_node(
31693170
[encoder setBuffer:id_dst offset:offs_dst atIndex:7];
31703171
[encoder setBytes:&args length:sizeof(args) atIndex:8];
31713172

3173+
// One shared memory bucket for each simd group in the threadgroup
3174+
// NOTE: Metal kernels require the buffer size to be multiple of 16 bytes
3175+
// https://developer.apple.com/documentation/metal/mtlcomputecommandencoder/1443142-setthreadgroupmemorylength
3176+
if (d_state >= 32) {
3177+
GGML_ASSERT((int64_t)(d_state / 32) <= 32);
3178+
const int64_t shmem_size = 32;
3179+
GGML_ASSERT(d_state <= (int64_t)pipeline.maxTotalThreadsPerThreadgroup);
3180+
[encoder setThreadgroupMemoryLength:(shmem_size)*sizeof(float) atIndex:0];
3181+
}
3182+
31723183
if (ne30 == 1) {
31733184
// Mamba-2
3174-
[encoder dispatchThreadgroups:MTLSizeMake(d_inner, n_head, n_seqs) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
3185+
[encoder dispatchThreadgroups:MTLSizeMake(d_inner, n_head, n_seqs) threadsPerThreadgroup:MTLSizeMake(d_state, 1, 1)];
31753186
} else {
31763187
GGML_ASSERT(d_inner == 1);
3177-
[encoder dispatchThreadgroups:MTLSizeMake(n_head, n_seqs, 1) threadsPerThreadgroup:MTLSizeMake(1, 1, 1)];
3188+
[encoder dispatchThreadgroups:MTLSizeMake(n_head, n_seqs, 1) threadsPerThreadgroup:MTLSizeMake(d_state, 1, 1)];
31783189
}
31793190
} break;
31803191
case GGML_OP_RWKV_WKV6:

0 commit comments

Comments
 (0)