From: Jassieluo Date: Fri, 26 Jun 2026 07:02:42 +0000 (+0800) Subject: sycl : clamp softmax input to avoid underflow (#24941) X-Git-Tag: upstream/0.0.10438~631 X-Git-Url: https://git.djapps.eu/?a=commitdiff_plain;h=e7e3f350904c087c7601f5d9d276c3c127f84345;p=pkg%2Fggml%2Fsources%2Fllama.cpp sycl : clamp softmax input to avoid underflow (#24941) --- diff --git a/ggml/src/ggml-sycl/softmax.cpp b/ggml/src/ggml-sycl/softmax.cpp index 18bf379bb..67ea282b4 100644 --- a/ggml/src/ggml-sycl/softmax.cpp +++ b/ggml/src/ggml-sycl/softmax.cpp @@ -126,7 +126,7 @@ static void soft_max_f32(const float * x, break; } - const float val = sycl::native::exp(vals[col] - max_val); + const float val = sycl::native::exp(sycl::max(vals[col] - max_val, -80.0f)); tmp += val; vals[col] = val; } @@ -154,7 +154,7 @@ static void soft_max_f32(const float * x, tmp = warp_reduce_sum(tmp); } if (sinks) { - tmp += sycl::native::exp(sinks[i02] - max_val); + tmp += sycl::native::exp(sycl::max(sinks[i02] - max_val, -80.0f)); } const float inv_sum = 1.0f / tmp;