0004-metal-add-missing-barriers-for-mul-mat-2699.patch 1.7 KB

1234567891011121314151617181920212223242526272829303132333435363738394041
  1. From 14b1d7e6f720dee41ce5a826376df738096d9033 Mon Sep 17 00:00:00 2001
  2. From: Shouzheng Liu <lshzh.hi@gmail.com>
  3. Date: Tue, 22 Aug 2023 02:18:40 -0400
  4. Subject: [PATCH] metal : add missing barriers for mul-mat (#2699)
  5. ---
  6. ggml-metal.metal | 5 +++--
  7. 1 file changed, 3 insertions(+), 2 deletions(-)
  8. diff --git a/ggml-metal.metal b/ggml-metal.metal
  9. index 88d48f6..ce3541f 100644
  10. --- a/ggml-metal.metal
  11. +++ b/ggml-metal.metal
  12. @@ -1850,6 +1850,7 @@ kernel void kernel_mul_mm(device const uchar * src0,
  13. //load data and store to threadgroup memory
  14. half4x4 temp_a;
  15. dequantize_func(x, il, temp_a);
  16. + threadgroup_barrier(mem_flags::mem_threadgroup);
  17. #pragma unroll(16)
  18. for (int i = 0; i < 16; i++) {
  19. *(sa + SG_MAT_SIZE * ((tiitg / THREAD_PER_ROW / 8) \
  20. @@ -1895,14 +1896,14 @@ kernel void kernel_mul_mm(device const uchar * src0,
  21. }
  22. } else {
  23. // block is smaller than 64x32, we should avoid writing data outside of the matrix
  24. + threadgroup_barrier(mem_flags::mem_threadgroup);
  25. threadgroup float *temp_str = ((threadgroup float *)shared_memory) \
  26. + 32 * (sgitg&1) + (16 * (sgitg>>1)) * BLOCK_SIZE_M;
  27. for (int i = 0; i < 8; i++) {
  28. - threadgroup_barrier(mem_flags::mem_device);
  29. simdgroup_store(c_res[i], temp_str + 8 * (i%4) + 8 * BLOCK_SIZE_M * (i/4), BLOCK_SIZE_M);
  30. }
  31. - threadgroup_barrier(mem_flags::mem_device);
  32. + threadgroup_barrier(mem_flags::mem_threadgroup);
  33. device float *C = dst + BLOCK_SIZE_M * r0 + (BLOCK_SIZE_N * r1) * ne0 + im*ne1*ne0;
  34. if (sgitg==0) {
  35. for (int i = 0; i < n_rows; i++) {
  36. --
  37. 2.41.0