Skip to content
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

IQ3_S: a much better alternative to Q3_K #5676

Merged
merged 27 commits into from
Feb 24, 2024
Merged

IQ3_S: a much better alternative to Q3_K #5676

Changes from 1 commit
Commits
Show all changes
27 commits
Select commit Hold shift + click to select a range
10a47fa
iq4_nl: squash commits for easier rebase
Kawrakow Feb 19, 2024
5691fec
Resurrecting iq3_xs
Kawrakow Feb 20, 2024
76aff09
Minor PPL improvement via a block scale fudge factor
Kawrakow Feb 20, 2024
5be4e7a
Minor improvement via 3 neighbours
Kawrakow Feb 20, 2024
f1255c5
iq3_xs: working scalar and AVX2 dot products
Kawrakow Feb 20, 2024
76214ab
iq3_xs: ARM_NEON dot product - works but extremely slow (10 t/s)
Kawrakow Feb 20, 2024
38aa7b1
iq3_xs: working Metal implementation
Kawrakow Feb 20, 2024
2ec600b
Adding IQ3_M - IQ3_XS mix with mostly Q4_K
Kawrakow Feb 21, 2024
d83fdda
iiq3_xs: a 3.4375 bpw variant
Kawrakow Feb 22, 2024
eacff4a
iq3_xs: make CUDA work for new version
Kawrakow Feb 22, 2024
1fef4b8
iq3_xs: make scalar and AVX2 work for new version
Kawrakow Feb 22, 2024
1328331
iq3_s: make ARM_NEON work with new version
Kawrakow Feb 22, 2024
1777825
iq3_xs: make new version work on metal
Kawrakow Feb 22, 2024
87038fe
iq3_xs: tiny Metal speed improvement
Kawrakow Feb 22, 2024
4d5feeb
iq3_xs: tiny Metal speed improvement
Kawrakow Feb 22, 2024
b25f996
Fix stupid warning
Kawrakow Feb 22, 2024
272c7f7
Q3_K_XS now uses a mix of IQ3_XS and IQ3_XXS
Kawrakow Feb 22, 2024
2730225
iq3_xs: rename to iq3_s
Kawrakow Feb 22, 2024
47cf30b
iq3_s: make tests pass
Kawrakow Feb 22, 2024
cd6a0f0
Move Q3_K_XS mix to 3.25 bpw
Kawrakow Feb 23, 2024
436a146
Attempt to fix failing tests
Kawrakow Feb 23, 2024
303f3f3
Another attempt to fix the Windows builds
Kawrakow Feb 23, 2024
0d6d185
Attempt to fix ROCm
Kawrakow Feb 23, 2024
1d47de3
ROCm again
Kawrakow Feb 23, 2024
e6e61e3
iq3_s: partial fix for QK_K = 64
Kawrakow Feb 23, 2024
cbd950b
iq3_s: make it work on metal for QK_K = 64
Kawrakow Feb 23, 2024
e1b8efb
Will this fix ROCm?
Kawrakow Feb 23, 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
11 changes: 11 additions & 0 deletions ggml-cuda.cu
View file
Open in desktop
Original file line number Diff line number Diff line change
Expand Up @@ -196,6 +196,17 @@ static __device__ __forceinline__ int __vsub4(const int a, const int b) {
return __vsubss4(a, b);
}

static __device__ __forceinline__ unsigned int __vcmpeq4(unsigned int a, unsigned int b) {
const uint8x4_t& va = reinterpret_cast<const uint8x4_t&>(a);
Copy link
Collaborator

Choose a reason for hiding this comment

The reason will be displayed to describe this comment to others. Learn more. 

diff --git a/ggml-cuda.cu b/ggml-cuda.cu
index 92f9309b..9729ad73 100644
--- a/ggml-cuda.cu
+++ b/ggml-cuda.cu
@@ -172,6 +172,7 @@
 #endif
 
 typedef int8_t int8x4_t __attribute__((ext_vector_type(4)));
+typedef uint8_t uint8x4_t __attribute__((ext_vector_type(4)));
 static __device__ __forceinline__ int __vsubss4(const int a, const int b) {
     const int8x4_t va = reinterpret_cast<const int8x4_t&>(a);
     const int8x4_t vb = reinterpret_cast<const int8x4_t&>(b);

const uint8x4_t& vb = reinterpret_cast<const uint8x4_t&>(b);
unsigned int c;
uint8x4_t& vc = reinterpret_cast<uint8x4_t&>(c);
for (int i = 0; i < 4; ++i) {
vc[i] = va[i] == vb[i] ? 0xff : 0x00;
}
return c;
}

static __device__ __forceinline__ int __dp4a(const int a, const int b, int c) {
#if defined(__gfx906__) || defined(__gfx908__) || defined(__gfx90a__) || defined(__gfx1030__)
c = __builtin_amdgcn_sdot4(a, b, c, false);
Expand Down
Loading