8000 metal : add quantized FA support by ggerganov · Pull Request #10149 · ggml-org/llama.cpp · GitHub
[go: up one dir, main page]
Skip to content

metal : add quantized FA support #10149

New issue

Have a question about this project? Sign up for a free GitHub account to open an issue and contact its maintainers and the community.

Sign up for GitHub

By clicking “Sign up for GitHub”, you agree to our terms of service and privacy statement. We’ll occasionally send you account related emails.

Already on GitHub? Sign in to your account

Merged
merged 8 commits into from
Nov 6, 2024
Merged

metal : add quantized FA support #10149

Changes from 1 commit
Commits
Show all changes
8 commits
Select commit Hold shift + click to select a range
6c484f3
metal : add quantized FA (vec) support
ggerganov Nov 3, 2024
e9565cc
metal : add quantized FA (non-vec) support
ggerganov Nov 4, 2024
13b87f2
metal : fix support check
ggerganov Nov 4, 2024
dd0d9ed
metal : clean-up
ggerganov Nov 4, 2024
1e12961
metal : clean-up (cont)
ggerganov Nov 4, 2024
d805404
metal : fix shared memory calc + reduce smem + comments
ggerganov Nov 5, 2024
73f378d
metal : float-correctness
ggerganov Nov 5, 2024
9c13f95
metal : minor [no ci]
ggerganov Nov 5, 2024
File filter

Filter by extension

Filter by extension
Conversations
Failed to load comments.
Loading
Jump to
Jump to file
Failed to load files.
Loading
Diff view
Diff view
Prev Previous commit
Next Next commit
metal : float-correctness
  • Loading branch information
@ggerganov
ggerganov committed Nov 5, 2024
commit 73f378df824ef9c79a1f16f68bcb892a646f12af
8 changes: 4 additions & 4 deletions ggml/src/ggml-metal.metal
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -2816,7 +2816,7 @@ kernel void kernel_flash_attn_ext(
threadgroup_barrier(mem_flags::mem_threadgroup);

{
float S[Q] = { [0 ... Q-1] = 0.0h };
float S[Q] = { [0 ... Q-1] = 0.0f };
float M[Q] = { [0 ... Q-1] = -FLT_MAX/2 };

// thread indices inside the simdgroup
Expand Down Expand Up @@ -3064,7 +3064,7 @@ kernel void kernel_flash_attn_ext(

// reduce the warps sequentially
for (short sg = 1; sg < nsg; ++sg) {
float S = { 0.0h };
float S = { 0.0f };
float M = { -FLT_MAX/2 };

threadgroup_barrier(mem_flags::mem_threadgroup);
Expand Down Expand Up @@ -3263,7 +3263,7 @@ kernel void kernel_flash_attn_ext_vec(

// zero out lo
for (short i = 0; i < D16/NW4; i += NW4) {
lo[i] = float4x4(0.0h);
lo[i] = float4x4(0.0f);
}

// zero out shared memory SH
Expand All @@ -3274,7 +3274,7 @@ kernel void kernel_flash_attn_ext_vec(
threadgroup_barrier(mem_flags::mem_threadgroup);

{
float S = 0.0h;
float S = 0.0f;
float M = -FLT_MAX/2;

// thread indices inside the simdgroup
Expand Down
Loading
0