LLM-Inference/ik_llama.cpp
1
0
Fork 0
mirror of https://github.com/ikawrakow/ik_llama.cpp.git synced 2026-08-12 22:29:39 +04:00
Code Issues Packages Projects Releases Wiki Activity

Fix CUDA silu kernel for merged up/gate with limit

Browse Source
This commit is contained in:
Kawrakow
2026-08-06 05:18:57 +00:00
parent 415e5278bc
commit 5c8db78a84
2 changed files with 7 additions and 10 deletions
Download Patch File Download Diff File Expand all files Collapse all files
ggml/src/ggml-cuda/unary.cu
+4 -3
View File
@@ -113,7 +113,7 @@ static __global__ void fused_mul_silu_f32(const float * x, float * dst, const in
int row = i / ne0;
int j = i % ne0;
auto x_row = x + 2*row*ne0;
//dst[i] = x_row[j] * x_row[j + ne0] / (1.0f + expf(-x_row[j + ne0]));
// Note: gate is at the beginning of each row, up is offset by ne0
dst[i] = x_row[j] * x_row[j + ne0] / (1.0f + expf(-x_row[j]));
}
@@ -126,9 +126,10 @@ static __global__ void fused_mul_silu_f32(const float * x, float * dst, const in
int row = i / ne0;
int j = i % ne0;
auto x_row = x + 2*row*ne0;
float g = x_row[j + ne0] / (1.0f + expf(-x_row[j + ne0]));
// Note: gate is at the beginning of each row, up is offset by ne0
float g = x_row[j] / (1.0f + expf(-x_row[j]));
g = min(g, limit);
dst[i] = max(-limit, min(limit, x_row[j])) * g;
dst[i] = max(-limit, min(limit, x_row[j + ne0])) * g;
}
static __global__ void fused_mul_relu_f32(const float * x, const float * y, float * dst, const int k) {
src/llama-load-tensors.cpp
+3 -7
View File
@@ -2913,13 +2913,9 @@ bool create_tensors_helper::create_deepseek4_tensors(const LLM_TN & tn) {
layer.ffn_gate_inp = create_tensor_from_meta(ctx_split, layer_weight_name(i, "ffn_gate_inp"));
layer.ffn_norm = create_tensor_from_meta(ctx_split, layer_weight_name(i, "ffn_norm"));
if (i >= hparams.dsv4_hash_layer_count) {
use_mmap_buffer &= !create_std_ffn_exps_from_meta(tn, i, 0);
} else {
layer.ffn_gate_exps = create_tensor_from_meta(ctx_split, layer_weight_name(i, "ffn_gate_exps"));
layer.ffn_down_exps = create_tensor_from_meta(ctx_split, layer_weight_name(i, "ffn_down_exps"));
layer.ffn_up_exps = create_tensor_from_meta(ctx_split, layer_weight_name(i, "ffn_up_exps"));
}
use_mmap_buffer &= !create_std_ffn_exps_from_meta(tn, i, 0);
layer.ffn_gate_shexp = create_tensor_from_meta(ctx_split, layer_weight_name(i, "ffn_gate_shexp"));
layer.ffn_down_shexp = create_tensor_from_meta(ctx_split, layer_weight_name(i, "ffn_down_shexp"));
layer.ffn_up_shexp = create_tensor_from_meta(ctx_split, layer_weight_name(i, "ffn_up_shexp"));