From e7e3f350904c087c7601f5d9d276c3c127f84345 Mon Sep 17 00:00:00 2001 From: Jassieluo <130133492+Jassieluo@users.noreply.github.com> Date: Fri, 26 Jun 2026 15:02:42 +0800 Subject: [PATCH] sycl : clamp softmax input to avoid underflow (#24941) --- ggml/src/ggml-sycl/softmax.cpp | 4 ++-- 1 file changed, 2 insertions(+), 2 deletions(-) diff --git a/ggml/src/ggml-sycl/softmax.cpp b/ggml/src/ggml-sycl/softmax.cpp index 18bf379bbe..67ea282b4b 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;