From: luoyu-intel <redacted>
Date: Tue, 23 Jul 2024 07:43:28 +0000 (+0000)
Subject: fix scratch size of softmax (llama/8642)
X-Git-Tag: upstream/0.0.1642~503
X-Git-Url: https://git.djapps.eu/?a=commitdiff_plain;h=aad0ddec8f29dbdf85eb46815318171072dbac83;p=pkg%2Fggml%2Fsources%2Fggml

fix scratch size of softmax (llama/8642)
---

diff --git a/src/ggml-sycl/softmax.cpp b/src/ggml-sycl/softmax.cpp
index c5d9a837..17a542e4 100644
--- a/src/ggml-sycl/softmax.cpp
+++ b/src/ggml-sycl/softmax.cpp
@@ -152,7 +152,8 @@ static void soft_max_f32_sycl(const float * x, const float * mask,
 
     const sycl::range<3> block_dims(1, 1, nth);
     const sycl::range<3> block_nums(1, 1, nrows_x);
-    const size_t n_local_scratch = (GGML_PAD(ncols_x, WARP_SIZE) + WARP_SIZE);
+    const size_t n_val_tmp = nth / WARP_SIZE;
+    const size_t n_local_scratch = (GGML_PAD(ncols_x, WARP_SIZE) + n_val_tmp);
 
     const uint32_t n_head_kv   = nrows_x/nrows_y;
     const uint32_t n_head_log2 = 1u << (uint32_t) floorf(log2f((float) n_head_kv));