SOTA 2-bit quants

[ikawrakow]

TL;DR

This PR adds new "true" 2-bit quantization (but due to being implemented within the block-wise quantization approach of ggml/llama.cpp we end up using 2.0625 bpw, see below for more details). The approach achieves perplexities similar to QuIP# (see this discussion). Inference performance is reasonable, but not great. E.g., for a 7B LLaMA model for TG-128 I measure 155 t/s on an RTX-4080 using CUDA, 54 t/s on a 30-core M2 Max using Metal, and 24.3 t/s on a Ryzen 5975WX CPU. In comparison, for Q4_0, I have 130 t/s (CUDA), 63.5 t/s (Metal), and 15.4 t/s (Ryzen).

Caveats

• Quantization is not provided (but see heavily commented quantization source code below). For now I'll publish quantized models on HF. Adding the quantization functions is a fairly big change, so this may be added later.
• Not all back-ends are provided by this PR. I have implemented the necessary functions/kernels for CPU without SIMD, AVX2, ARM_NEON, Metal, and CUDA (but only tested with cuBLAS on RTX-4080).
• Quantized matrix multiplications on CUDA are not (yet) implemented, and neither is the LLAMA_CUDA_FORCE_DMMV option (is this still being used?)

Perplexities

The table gives some sample results for quantized model sizes and perplexities. Note that a context length of 4096 has been used for these results (and not 512, as typically done when reporting PPL in this repo). Model sizes are in GiB, not GB, as this is more relevant to see if a model will load into a device with limited RAM/VRAM.

Model	File size (GiB)	PPL
Mistral-7B	1.855	6.446
LLaMA-v2-7B	1.728	7.067
LLaMA-v2-13B	3.295	5.728
LLaMA-v2-70B	17.03	4.079

Some details

I have borrowed 2 ideas from QuIP#:

• They force an even number of positive (or negative) quant signs in a group of 8 quants. This allows to use 7 bits for 8 quants to record their sign. This looked strange to me at first sight (as one will obviously always have cases where the number of positive/negative weights in a group of 8 is odd), but after some more thought, this is actually quite brilliant as one can always flip the sign of the most "unimportant" quant in a group of 8 if necessary to ensure even number of positive/negative signs. As usual, the devil is in the detail. The QuIP# authors do not specify how they choose the quant to sign-flip. After some experimentation, I have settled on using the quant with minimum w * x^2, where x is the model weight and w is the importance of this weight derived from embedding statistics obtained in a calibration run (the "importance matrix").
• They use the E8 lattice (see https://en.wikipedia.org/wiki/E8_lattice) to encode the magnitude of the quants in a group of 8. For 3 allowed quant values (1/2, 3/2, 5/2) there are 3^8 = 6561 grid points. After taking into account the E8 requirement that the sum of the grid point coordinates is even, one is left with 3281 possible choices. For a 2-bit quantization one has at most 512 slots available (8 quants x 2 bits = 16 bits minus 7 bits spent on encoding signs, so 9 bits left), so one needs to select a subset of the possible E8 lattice points. The QuIP# authors use all points within an 8D-sphere with radius sqrt(10) (227 grid points) plus another 29 hand-picked points within radius sort(12), for a total of 256 grid points. They then use the remaining one bit to record a +/- 1/4 shift for the group of 8. Instead, I have used the following approach to pick the 256 grid points, which gives a completely different set of points from theirs (they call it a "codebook"):
  • Perform quantization using all possible E8 lattice points on a bunch of models
  • Count how many times each of the 3281 points occurred in the quantized models
  • Select the 256 points such that a) The total count of selected points is maximized and b) The maximum distance of not selected points to the closest selected point is minimized

I have not used any of the "fancy" QuiP# stuff (incoherent processing, Hadamard transforms, and such). Instead, I utilize the battle-tested block scaling that is used for all other ggml quants. Recall that we have spent 7 bits to record the signs and 8 bits to record the grid point index ("codebook") for a group of 8 quants, so we have one spare bit, which gives 4 spare bits in a block of 32. I use this for a 4-bit scale, ending up with exactly 64 bits per block of 32, so a "true" 2-bit quantization. I have verified in my private development repo that using only one additional floating point scale per tensor row is sufficient, this adds less than 0.01 bpw. But the change required in ggml for a row-wise quantization is quite significant, so for now I'm adding this new quantization type within the existing ggml block-wise framework. As with k-quants, there are super-blocks of 256 weights, and each such super-block has one fp16 super-block scale, so this adds 16/256 = 0.0625 bits per weight to what is achievable in a row-wise quantization.

For reference I'm also adding the quantization function from my private development repo that is being used to generate quantized models for this quantization type:

Quantization source code

`

// * x points to the model weights to be quantized (input)
// * vy points to the buffer where the quantized weights will be stored (output)
// * n is the number of weights to be quantized
// * quant_weights is the importance matrix. In my case it is very simple, it contains just diagonal elements
//. stored consecutively
static void quantize_row_iq2_xxs_impl(const float * restrict x, void * restrict vy, int n, const float * restrict quant_weights) {

// kgrid_q2xs, kmap_q2xs, kneighbors_q2xs need to be initialized before the first call to this function
// Some are quite big (especially kneighbors_q2xs), so it does not make sense to have a separate
// copy for each thread in a multi-threaded quantization run.
//  * kgrid_q2xs contains the 256 selected E8-lattice points (so 256 x 8 uint8_t's)
//. * kmap_q2xs maps a 16-bit index for 8 quants, 2 bits per quant, to a corresponding grid point, or -1 if these 16
//    bits do not represent a point on the grid
// * kneighbors_q2xs contains the set of closest neighbors that are in the set of selected points for each of the 3281
//.   possible E8-lattice points 
GGML_ASSERT(kgrid_q2xs);   
GGML_ASSERT(kmap_q2xs);
GGML_ASSERT(kneighbors_q2xs);
GGML_ASSERT(n%QK_K == 0);

const int kMaxQ = 3;

const int nbl = n/256;  // number of super-blocks

block_iq2_xxs * y = vy;

float scales[QK_K/32];  // stores the  scales of 32-weight blocks in a super-block
float weight[32];            // stores importances
float xval[32];                //  model values after sign flips in a block of 32
int8_t L[32];                  //  quants in a block of 32
int8_t Laux[32];            //  tmp quants
float  waux[32];            //  auxiliary weight
bool   is_on_grid[4];         // flag for each of the 4 groups of 8 in a block of 32 indicating if the point was on the grid
bool   is_on_grid_aux[4]; // same as above for tmp usage
uint8_t block_signs[4];    //  sign bits for each group of 8 in a block of 32
uint32_t q2[2*(QK_K/32)]; // we will record the quantized data here before copying into the block_iq2_xss struct

for (int ibl = 0; ibl < nbl; ++ibl) { // for each super-block

    y[ibl].d = GGML_FP32_TO_FP16(0.f);
    memset(q2, 0, QK_K/4);

    float max_scale = 0;

    const float * xbl = x + QK_K*ibl;  // the model weights for this super block
    float sumx2 = 0;
    for (int i = 0; i < QK_K; ++i) sumx2 += xbl[i]*xbl[i];
    float sigma2 = sumx2/QK_K;  // the weight variance in this super-block

    for (int ib = 0; ib < QK_K/32; ++ib) { // for each block of 32 in this super-block 
        const float * xb = xbl + 32*ib;       // model weights for this block
        if (quant_weights) {
            // it is helpful to augment the importance matrix in the following way  
            const float * qw = quant_weights + QK_K*ibl + 32*ib;
            for (int i = 0; i < 32; ++i) weight[i] = qw[i] * sqrtf(sigma2 + xb[i]*xb[i]);
        } else {
            // OK, that's a version without an importance matrix, but don't even think using it.
            // For 2-bit quantization you will get total garbage without the importance matrix
            for (int i = 0; i < 32; ++i) weight[i] = xb[i]*xb[i];
        }
        // We use this augmented weight when searching for closest neighbors of points not on the grid
        for (int i = 0; i < 32; ++i) waux[i] = sqrtf(weight[i]);
        // Sign flips
        for (int k = 0; k < 4; ++k) {
            int nflip = 0;
            uint8_t s = 0;
            for (int i = 0; i < 8; ++i) {
                if (xb[8*k + i] >= 0) xval[8*k + i] = xb[8*k + i];
                else {
                    xval[8*k + i] = -xb[8*k + i]; ++nflip; s |= (1 << i);
                }
            }
            if (nflip%2) {
                // We have an odd number of negative weights. Need to flip one sign, so lets find the least important quant
                int imin = 0; float min = weight[8*k+imin]*xb[8*k+imin]*xb[8*k+imin];
                for (int i = 1; i < 8; ++i) {
                    float ax = weight[8*k+i]*xb[8*k+i]*xb[8*k+i];
                    if (ax < min) {
                        min = ax; imin = i;
                    }
                }
                // flip the sign
                xval[8*k+imin] = -xval[8*k+imin];
                s ^= (1 << imin);
            }
            block_signs[k] = s & 127;
        }
        // after the above loop, we have decided on sign flips and xval contains the "signless" model weights
        // for the block against which we will be comparing from here on.
        // Find max value in the block
        float max = xval[0];
        for (int i = 1; i < 32; ++i) max = MAX(max, xval[i]);
        if (!max) {
            // all zeros - nothing to do
            scales[ib] = 0;
            memset(L, 0, 32);
            continue;
        }
        // Now lets try to find the best scale for this block
        // We try a bunch of scales around max/5. Why 5 ? Because our possible quants are 1, 3, 5 
        float best = 0;
        float scale = max/(2*kMaxQ-1);
        for (int is = -9; is <= 9; ++is) {
            float id = (2*kMaxQ-1+is*0.1f)/max; // inverse scale
            float this_scale = 1/id;
            for (int k = 0; k < 4; ++k) { // for each group of 8
                // get the quants via RTN
                for (int i = 0; i < 8; ++i) {
                    int l = nearest_int(0.5f*(id*xval[8*k+i]-1));
                    Laux[8*k+i] = MAX(0, MIN(kMaxQ-1, l));
                }
                // convert the quants to a 16-bit integer and lookup the corresponding grid point 
                uint16_t u = 0;
                for (int i = 0; i < 8; ++i) u |= (Laux[8*k+i] << 2*i);
                int grid_index = kmap_q2xs[u];
                is_on_grid_aux[k] = true;
                if (grid_index < 0) {
                    // Not on the grid of selected points. Need to find a neigbouring grid point  
                    is_on_grid_aux[k] = false;
                    // the set of closest neigbhbors we will use.
                    // We could check all points on the grid, but that would be much too slow
                    // in a loop over 18 scale choices
                    const uint16_t * neighbours = kneighbors_q2xs - kmap_q2xs[u] - 1;
                    int num_neighbors = neighbours[0];
                    GGML_ASSERT(num_neighbors > 0);
                    // find the "closest" neighbor
                    float best_d2 = FLT_MAX;
                    for (int j = 1; j <= num_neighbors; ++j) {
                        const int8_t * pg = (const int8_t *)(kgrid_q2xs + neighbours[j]);
                        float d2 = 0;
                        for (int i = 0; i < 8; ++i) {
                            float q = pg[i];
                            float diff = this_scale*q - xval[8*k + i];
                            d2 += waux[8*k+i]*diff*diff;
                        }
                        if (d2 < best_d2) {
                            best_d2 = d2; grid_index = neighbours[j];
                        }
                    }
                    const int8_t * pg = (const int8_t *)(kgrid_q2xs + grid_index);
                    // store the closest neighbor quant values into Laux
                    for (int i = 0; i < 8; ++i) Laux[8*k+i] = (pg[i] - 1)/2;
                }
            }
            // OK, now we have in Laux quants that are all one of the 256 selected points from the E8 lattice
            // time to find the best scale via weighted RMSE minimization   
            float sumqx = 0, sumq2 = 0;
            for (int i = 0; i < 32; ++i) {
                float w = weight[i];
                float q = 2*Laux[i] + 1;
                sumqx += w*xval[i]*q;
                sumq2 += w*q*q;
            }
            if (sumq2 > 0 && sumqx*sumqx > best*sumq2) {
                // this set of quants is better than what we had so far => store the quants/scale
                scale = sumqx/sumq2; best = scale*sumqx;
                for (int i = 0; i < 32; ++i) L[i] = Laux[i];
                for (int k = 0; k <  4; ++k) is_on_grid[k] = is_on_grid_aux[k];
            }
        }
        int n_not_ongrid = 0;
        for (int k = 0; k < 4; ++k) if (!is_on_grid[k]) ++n_not_ongrid;
        if (n_not_ongrid > 0 && scale > 0) {
            // some of the points we selected weren't on the grid and we are using their closest grid neighbor
            // it is a good idea to re-quantize those with our final block scale 
            float id = 1/scale;
            for (int k = 0; k < 4; ++k) {
                if (is_on_grid[k]) continue;
                uint16_t u = 0;
                for (int i = 0; i < 8; ++i) {
                    int l = nearest_int(0.5f*(id*xval[8*k+i]-1));
                    l = MAX(0, MIN(kMaxQ-1, l));
                    u |= (l << 2*i);
                }
                int grid_index = kmap_q2xs[u];
                if (grid_index < 0) {
                    // still not on the grid, so find the closest neighbor (same as above)
                    const uint16_t * neighbours = kneighbors_q2xs - kmap_q2xs[u] - 1;
                    int num_neighbors = neighbours[0];
                    GGML_ASSERT(num_neighbors > 0);
                    float best_d2 = FLT_MAX;
                    for (int j = 1; j <= num_neighbors; ++j) {
                        const int8_t * pg = (const int8_t *)(kgrid_q2xs + neighbours[j]);
                        float d2 = 0;
                        for (int i = 0; i < 8; ++i) {
                            float q = pg[i];
                            float diff = scale*q - xval[8*k + i];
                            d2 += waux[8*k+i]*diff*diff;
                        }
                        if (d2 < best_d2) {
                            best_d2 = d2; grid_index = neighbours[j];
                        }
                    }
                }
                const int8_t * pg = (const int8_t *)(kgrid_q2xs + grid_index);
                // update this group of 8 quants
                for (int i = 0; i < 8; ++i) L[8*k+i] = (pg[i] - 1)/2;
            }
            // determine the best scale again via weighted RMSE minimization
            float sumqx = 0, sumq2 = 0;
            for (int i = 0; i < 32; ++i) {
                float w = weight[i];
                float q = 2*L[i] + 1;
                sumqx += w*xval[i]*q;
                sumq2 += w*q*q;
            }
            if (sumq2 > 0) scale = sumqx/sumq2;
        }
        // This should not actually happen, but just in case: flip the scale (so it is positive) and
        // correspondingly flip the quant signs in the block
        if (scale < 0) {
            scale = -scale;
            for (int k = 0; k < 4; ++k) block_signs[k] = (~block_signs[k]) & 127;
        }
        // encode the quants
        for (int k = 0; k < 4; ++k) {
            uint16_t u = 0;
            for (int i = 0; i < 8; ++i) u |= (L[8*k+i] << 2*i);
            int grid_index = kmap_q2xs[u];
            if (grid_index < 0) {
                printf("Oops: found point %u not on grid:", u);
                for (int i = 0; i < 8; ++i) printf(" %d", L[8*k+i]);
                printf("\n");
                GGML_ASSERT(false);
            }
            q2[2*ib+0] |= (grid_index << 8*k);
            q2[2*ib+1] |= (block_signs[k] << 7*k);
        }
        GGML_ASSERT(scale >= 0);
        scales[ib] = scale;
        max_scale = MAX(max_scale, scale);
    }

    if (!max_scale) {
        // all weights in the super-block were zero, so nothing to do
        memset(y[ibl].qs, 0, QK_K/4);
        continue;
    }

    // super-block scale
    float d = max_scale/31;
    y[ibl].d = GGML_FP32_TO_FP16(d);
    float id = 1/d;
    // Now quantize the block scales to 4 bits
    float sumqx = 0, sumq2 = 0;
    for (int ib = 0; ib < QK_K/32; ++ib) {
        int l = nearest_int(0.5f*(id*scales[ib]-1));
        l = MAX(0, MIN(15, l));
        // add the quantized scale to the encoded quant data
        q2[2*ib+1] |= ((uint32_t)l << 28);
        // In principle we are done here. We get a minor improvement by re-quantizing after
        // scaling with the quantized scales and re-optimizing the super-block scale
        // To do so, we need again the block importances
        const float * xb = xbl + 32*ib;
        if (quant_weights) {
            const float * qw = quant_weights + QK_K*ibl + 32*ib;
            for (int i = 0; i < 32; ++i) weight[i] = qw[i] * sqrtf(sigma2 + xb[i]*xb[i]);
        } else {
            for (int i = 0; i < 32; ++i) weight[i] = xb[i]*xb[i];
        }
        const uint8_t * aux8 = (const uint8_t *)(q2 + 2*ib);
        // The actual scale for this block
        const float db = d * (1 + 2*l);
        uint32_t u = 0;
        for (int k = 0; k < 4; ++k) {
            const int8_t * signs = keven_signs_q2xs + 8*((q2[2*ib+1] >> 7*k) & 127);
            const float * xk = xb + 8*k;
            const float * wk = weight + 8*k;
            const uint8_t * grid = (const uint8_t *)(kgrid_q2xs + aux8[k]);
            float best_mse = 0; int best_index = aux8[k];
            for (int j = 0; j < 8; ++j) {
                float diff = db * grid[j] * signs[j] - xk[j];
                best_mse += wk[j] * diff * diff;
            }
            // As we are doing this just once, we can afford to check all 256 points
            for (int idx = 0; idx < 256; ++idx) {
                grid = (const uint8_t *)(kgrid_q2xs + idx);
                float mse = 0;
                for (int j = 0; j < 8; ++j) {
                    float diff = db * grid[j] * signs[j] - xk[j];
                    mse += wk[j] * diff * diff;
                }
                if (mse < best_mse) {
                    best_mse = mse; best_index = idx;
                }
            }
            u |= (best_index << 8*k);
            grid = (const uint8_t *)(kgrid_q2xs + aux8[k]);
            for (int j = 0; j < 8; ++j) {
                float q = db * grid[j] * signs[j];
                sumqx += wk[j] * q * xk[j];
                sumq2 += wk[j] * q * q;
            }
        }
        q2[2*ib] = u;
        if (sumq2 > 0) y[ibl].d = GGML_FP32_TO_FP16(d*sumqx/sumq2);
    }
    memcpy(y[ibl].qs, q2, QK_K/4);
}

}
`

[slaren]

Very nice! A 70b quant that fits in 24gb is awesome.

You can use test-backend-ops to compare the GPU implementations with the CPU implementation, by adding the new quant type to all_types:

llama.cpp/tests/test-backend-ops.cpp

Lines 1453 to 1461 in a919280

	const ggml_type all_types[] = {
	GGML_TYPE_F32, GGML_TYPE_F16,
	GGML_TYPE_Q4_0, GGML_TYPE_Q4_1,
	GGML_TYPE_Q5_0, GGML_TYPE_Q5_1,
	GGML_TYPE_Q8_0,
	GGML_TYPE_Q2_K, GGML_TYPE_Q3_K,
	GGML_TYPE_Q4_K, GGML_TYPE_Q5_K,
	GGML_TYPE_Q6_K
	};

[ikawrakow]

Thank you @slaren for pointing out the automation in testing. I cannot quite use it yet because of the missing quantization capability (unless the automation is so clever as to pull quantized models from somewhere else). The automated testing is also failing because there is no quantization function provided.

[JohannesGaessler]

Quantized matrix multiplications on CUDA are not (yet) implemented, and neither is the LLAMA_CUDA_FORCE_DMMV option (is this still being used?)

The dequantize_mul_mat_vec kernel is almost always not being used. On all devices with compute capability 6.1 or higher mul_mat_vec_q is used. But if you were to use e.g. a P100 or a Maxwell card it would be used. I think it would be fine if those cards are simply not supported though.

Regarding mul_mat_q: do you intend to do an implementation? In my experience MMQ is more difficult to work with than MMVQ so I would be willing to do the implementation instead (or just provide help).

But the change required in ggml for a row-wise quantization is quite significant, so for now I'm adding this new quantization type within the existing ggml block-wise framework.

I'm still prototyping but per-row/per-column scales may also be necessary for an efficient int8 tensor core implementation. The problem with the current block formats is that every time the scale changes you need to unload the tensor core accumulator which is very slow. In my case the necessary changes would be minor and isolated to ggml-cuda.cu though.

[JohannesGaessler]

Looks pretty good, would be amazing to have if NVIDIA releases another RTX ??90 with 24 GB VRAM.

Can the multiple instances of the hardcoded values be deduplicated? Perhaps by defining macros like IQ2XXS_GRID_VALUES?

Also consider #4755 if you haven't yet seen it. There seem to be issues with numerical precision when quantizing the hidden state for short context sizes. These issues will not show up in perplexity calculations because they exclude the first few tokens of a chunk. Though this does not necessarily mean that per-row scales for the weights are also problematic.

[JohannesGaessler]

Does this have the same effect as __constant__? In other words, does it actually put these values into constant memory? (Should be faster than if it is in global memory.)

[ikawrakow]

Just tried. Replacing

static const __device__ uint8_t ksigns_iq2xs[128]

with

static __constant__ __device__ uint8_t ksigns_iq2xs[128]

makes it massively slower (108 t/s vs 155 t/s)

[ikawrakow]

Can the multiple instances of the hardcoded values be deduplicated? Perhaps by defining macros like IQ2XXS_GRID_VALUES?

I thought about this, but did not do it (yet). Mainly because:

• The different files require different qualifiers (__device__ on CUDA, constexpr on Metal, etc.), so one either needs to work with pre-processor trickery, or needs to define the actual content as a macro. I did not like both options too much
• This would add an extra file, and my concept is that @ggerganov likes having as few files as possible.

But yes, absolutely, this is something one should consider.

[JohannesGaessler]

I see. The compiler is probably not copying the data from constant memory to registers. So for frequently used data it's slower as long as there is no register spilling.

[ikawrakow]

Also consider #4755 if you haven't yet seen it. There seem to be issues with numerical precision when quantizing the hidden state for short context sizes.

Yes, I know about this. It is context and model dependent. For some models (e.g. Falcon-7B) the difference between quantizing and not quantizing the hidden state can be quite dramatic. This is why the dequantize_mul_mat_vec kernels are actually useful, and I'm somewhat surprised they have fallen out of favor.

[ikawrakow]

Regarding mul_mat_q: do you intend to do an implementation? In my experience MMQ is more difficult to work with than MMVQ so I would be willing to do the implementation instead (or just provide help).

So, the MMQ kernels are mostly useful for contexts in the few - few tens of tokens range. I know this is an important use case for stuff such as speculative sampling, but in my private repo I have an MMQ implementation based on plain vector dot products that outperforms MMQ for, say, up to 16 tokens. If one could extend this up, and/or extend the dequantize/cuBLAS performance superiority down to fewer tokens, the MMQ kernels become unnecessary. This is the main reason I'm kind of reluctant with those, especially considering the amount of code and compilation time increase each new MMQ kernel adds.

[JohannesGaessler]

Originally the MMQ kernels were intended for large matrices. But it later turned out that the FP32 cuBLAS GEMM does not actually use tensor cores and that FP16 cuBLAS GEMM is still faster for Volta or newer. Georgi then repurposed the MMQ kernels for small batch sizes by changing the tile sizes. They were never intended or optimized by me for this use case and in my testing they still perform worse than FP16 cuBLAS even for small batch sizes:

However, because you do not need to dequantize the weight matrix MMQ should still be more efficient in terms of VRAM. Also on Pascal/RDNA2 or older there are no tensor cores so it is also faster than cuBLAS GEMM by a factor of ~2.

in my private repo I have an MMQ implementation based on plain vector dot products that outperforms MMQ for, say, up to 16 tokens.

I was thinking that you could extend MMVQ to allow for >1 y columns and probably get better performance than with MMQ/cuBLAS GEMM. Presumably this is very similar to what you have.

[ikawrakow]

But to follow up on your __constant__ comment, I did copy the data into shared memory on Metal. This boosted TG from about 48 t/s to ~54 t/s, so 12.5% speedup. There it was easy to implement without changing any other kernel. I was thinking that one could gain some performance on CUDA too by copying the grid/sign data to shared memory, but I didn't see how to do it without changing the MMVQ template and touching every single dot product kernel, so left it as is for now.

[JohannesGaessler]

I don't know how the shared memory on Metal works but with CUDA there are by default 32 shared memory banks with a size of 4 byte each. The values here are uint64 = 8 byte so unless you can ensure that the threads in a warp access different memory banks ((pointer % (32*4 bytes))/4 bytes needs to be different) I would expect you to get a lot of memory bank conflicts which drastically reduces the memory bandwidth.

From what I can tell you have not yet published any models using the new format so I cannot test this myself but with NVIDIA NSight Compute under the occupancy section you can see the number of registers needed per thread. If that number is not limiting occupancy I would not expect much of a performance gain from moving the values to shared memory (but doing that may still reduce cache evictions of other data so I could be wrong).

[JohannesGaessler]

I assume this is for testing different versions.

[ikawrakow]

Yes. The one currently active is slightly faster on my RTX-4080, but I left the other version behind (which was the initial implementation) just in case. You never know with all these different cards that are being supported.

[Dampfinchen]

On a side note, are Mixtral GGUFs "finalized" now to the point a breaking change is not to be expected? I remember that warning Slaren wrote in the first PR.

[ikawrakow]

I have posted some IQ2_XXS quantized models, including Mixtral-8x7B, on Huggingface. See https://huggingface.co/ikawrakow/various-2bit-sota-gguf/tree/main

[Dampfinchen]

I have posted some IQ2_XXS quantized models, including Mixtral-8x7B, on Huggingface. See https://huggingface.co/ikawrakow/various-2bit-sota-gguf/tree/main

Wow, only 12 GB for Mixtral? That is super impressive. Thank you for your amazing work. Do you plan to update the other quants as well?

[ikawrakow]

Do you plan to update the other quants as well?

I started playing with Mixtral this morning, so yes, I will post quantizations for the other ggml quants in the next days. I just finished a perplexity calculation for Q4_K_S, and I'm getting PPL = 4.1764 for context of 512. According to PR #4739, Q4_K_S perplexity is 4.5136 on current master. Is this possible? I have not yet seen such a massive difference between official k-quants and my activation aware quantization.

Update: I see that the perplexity of 4.5136 quoted in PR #4739 is simply wrong. I get PPL = 4.2523 for Q4_K_S with current llama.cpp master.

[JianbangZ]

Do you plan to update the other quants as well?

I started playing with Mixtral this morning, so yes, I will post quantizations for the other ggml quants in the next days. I just finished a perplexity calculation for Q4_K_S, and I'm getting PPL = 4.1764 for context of 512. According to PR #4739, Q4_K_S perplexity is 4.5136 on current master. Is this possible? I have not yet seen such a massive difference between official k-quants and my activation aware quantization.

I think you are onto something. I will take look at your repo and do some PPL and KL divergence calcualtions today

I have posted some IQ2_XXS quantized models, including Mixtral-8x7B, on Huggingface. See https://huggingface.co/ikawrakow/various-2bit-sota-gguf/tree/main

Can you talk more about how you quantize mixtral? follow the same principles as the others that leave gating part to fp16 or int8? are the code to quantize mixtral pushed to your repo? I will run some extensive test over the weekend for PPL and KL divergence

[ikawrakow]

Can you talk more about how you quantize mixtral? follow the same principles as the others that leave gating part to fp16 or int8?

Yes, the ffn_gate_inp.weight tensor is left as fp16 just like in mainline llama.cpp. Else the selection of quants is quantization type dependent. For Q4_K_S I have restored using Q5_K for ffn_down for all experts in the first 4 layers when quantizing to Q4_K_S (this has been lost in mainline llama.cpp), plus I use Q5_K for all attn_v tensors (these are just 4096 x 1024, so using 1 extra bit leads to negligible increase in size). I see someone has added a change to mainline llama.cpp to use Q8_0 for attn_k. I don't use this as in my experience quantization errors in the K and Q tensors have the least impact on quantization quality after token embeddings. For the "true" 2-bit model added via this PR (IQ2_XXS) I use Q2_K (so, 2.5625 bpw) for all attn_v, for token embeddings, and for all expert ffn_down tensors in the first 4 layers, everything else is IQ2_XXS.

[sakura-umi]

quant_weights is the importance matrix. In my case it is very simple, it contains just diagonal elements

can you explain more about how to get the importance matrix with a certain model (such as qwen)?

also see that the quantize_row_iq2_xxs_reference function not implement yet, how can I use the quantize_row_iq2_xxs_impl code above to make examples/quantize work?

[sorasoras]

server.exe don't seems to support SOTA2bit Quants yet.
No matter what I input. it would keep outputing
Llama: #���#�!������������������������!��$�$!����$#$������$������� #������� ����!�!�����������������$���#�����#$�$## �!����� ##��$ ��$#����� ����!�� #����!����$��#��� ������!

[Dampfinchen]

Can you talk more about how you quantize mixtral? follow the same principles as the others that leave gating part to fp16 or int8?

Yes, the ffn_gate_inp.weight tensor is left as fp16 just like in mainline llama.cpp. Else the selection of quants is quantization type dependent. For Q4_K_S I have restored using Q5_K for ffn_down for all experts in the first 4 layers when quantizing to Q4_K_S (this has been lost in mainline llama.cpp), plus I use Q5_K for all attn_v tensors (these are just 4096 x 1024, so using 1 extra bit leads to negligible increase in size). I see someone has added a change to mainline llama.cpp to use Q8_0 for attn_k. I don't use this as in my experience quantization errors in the K and Q tensors have the least impact on quantization quality after token embeddings. For the "true" 2-bit model added via this PR (IQ2_XXS) I use Q2_K (so, 2.5625 bpw) for all attn_v, for token embeddings, and for all expert ffn_down tensors in the first 4 layers, everything else is IQ2_XXS.

On anotther note, @TheBloke found out that while quantising Mixtral, Q4K_S and Q4K_M have the exact same size currently, so he decided to just upload the Q4K_M quants for now. Do you have any idea why that is and do you think your new quants could improve that?

[ggerganov]

@Dampfinchen The Q4_K_S and Q4_K_M refer to the quantization mixtures used to quantize a model. A mixture can utilize various quantization types (see the log when loading the model), but the name of the mixture usually indicates the most prominent quantization type used - Q4_K in this case. The current quantization mixtures for Mixtral models are a first iteration (it was stated in the original PR #4406). As a first quick iteration, the current Q4_K_S and Q4_K_M quantization mixtures use the same quantization types as stated in #4406:

@ikawrakow

I see someone has added a change to mainline llama.cpp to use Q8_0 for attn_k. I don't use this as in my experience quantization errors in the K and Q tensors have the least impact on quantization quality after token embeddings.

The only consideration to bump both attn_k and attn_v to Q8_0 was that due to the GQA, these tensors are relatively small (4x smaller than attn_q) so using Q8_0 results in just extra ~200-300 MB added to the model, so I thought it wouldn't hurt. Haven't tested at all how this reflects to PPL - likely not much as you have observed. So it's something that can probably be improved - there are comments to remind us:

llama.cpp/llama.cpp

Lines 8917 to 8928 in c75ca5d

	if (qs.model.hparams.n_expert == 8) {
	// for the 8-expert model, bumping this to Q8_0 trades just ~128MB
	// TODO: explore better strategies
	new_type = GGML_TYPE_Q8_0;
	}
	++qs.i_attention_wv;
	} else if (name.find("attn_k.weight") != std::string::npos) {
	if (qs.model.hparams.n_expert == 8) {
	// for the 8-expert model, bumping this to Q8_0 trades just ~128MB
	// TODO: explore better strategies
	new_type = GGML_TYPE_Q8_0;
	}

[ikawrakow]

quant_weights is the importance matrix. In my case it is very simple, it contains just diagonal elements

can you explain more about how to get the importance matrix with a certain model (such as qwen)?

This will become available soon. I will either make my repo public, or I will make a PR to mainline llama.cpp to add the capability to compute the importance matrix used for these quantizations.

also see that the quantize_row_iq2_xxs_reference function not implement yet, how can I use the quantize_row_iq2_xxs_impl code above to make examples/quantize work?

You cannot use the existing implementation in the quantize example and llama.cpp. The importance matrix needs to be loaded, passed along to llama_model_quantize_internal, and from there propagated to the actual quantization functions. This requires a changes to the ggml quantization interface (ggml_quantize_chunk) and implementation. Once you make this change, you don't need to call quantize_row_iq2_xxs_reference but can call any function of your choice.

[he29-net]

On a RX 6800 with hipBLAS I'm getting:
GGML_ASSERT: ggml-cuda.cu:7556: false

I assumed it's related to the not yet implemented MMQ (mul_mat_q) so I tried to add -nommq to avoid using it, but it makes no difference, so maybe I'm misunderstanding what the parameter does. Just wanted to mention it in case it's not expected. CPU-only inference on an old CPU (no AVX2) worked fine.

[ikawrakow]

On a RX 6800 with hipBLAS I'm getting: GGML_ASSERT: ggml-cuda.cu:7556: false

I assumed it's related to the not yet implemented MMQ (mul_mat_q) so I tried to add -nommq to avoid using it, but it makes no difference, so maybe I'm misunderstanding what the parameter does. Just wanted to mention it in case it's not expected. CPU-only inference on an old CPU (no AVX2) worked fine.

The -nommq options has been lost. To work around that I have introduced ggml_supports_mmq, which returns true for all existing quants and false for the new IQ2_XXS. The problem was that the call to ggml_supports_mmq was in the wrong pre-processor nesting level, so did not work for hipBLAS. I have pushed a fix, can you you try now?

[he29-net]

Just tested it and it works perfectly, thanks. 👍

Mixtral-8x7b-2.10bpw now fully fits into the 16 GB VRAM with 4096 token context and gets about 29 tokens/s. Just a month ago it didn't even cross my mind this could be possible. :)

[jxy]

looks fine to me

./main -m models/mixtral-instruct-8x7b-2.10bpw.gguf --temp 0 --repeat-penalty 1.0 --no-penalize-nl -p '[INST] what is 12*8+7 [/INST]' --log-disable
 [INST] what is 12*8+7 [/INST] To calculate the expression 12*8+7, you need to follow the order of operations, which is often remembered by the acronym PEMDAS: Parentheses, Exponents, Multiplication and Division (from left to right), Addition and Subtraction (from left to right).

In this case, you should perform multiplication first:

12 * 8 = 96

Then, add 7:

96 + 7 = 103

So, the result of the expression 12*8+7 is 103.

[jxy]

@x4080 I thought you meant mixtral. But the mistral 7b from @ikawrakow's HF repo is not instruct tuned. So you need to give it a few shot prompt. This works.

./main -m models/mistral-7b-2.20bpw.gguf --temp 0 --repeat-penalty 1.0 --no-penalize-nl -p 'Q: What is 2*4?
A: 8

Q: What is 3+6?
A: 9

Q: What is 3*4+7?
A: 3 * 4 = 12
   12 + 7 = 19
   So 3 * 4 + 7 = 19

Q: What is 13*6+5?
A: 13 * 6 = 78
   78 + 5 = 83
   So 13 * 6 + 5 = 83

Q: What is 12*8+7?
A:' --log-disable -r '
Q:'

The output is

 Q: What is 2*4?
A: 8

Q: What is 3+6?
A: 9

Q: What is 3*4+7?
A: 3 * 4 = 12
   12 + 7 = 19
   So 3 * 4 + 7 = 19

Q: What is 13*6+5?
A: 13 * 6 = 78
   78 + 5 = 83
   So 13 * 6 + 5 = 83

Q: What is 12*8+7?
A: 12 * 8 = 96
   96 + 7 = 103
   So 12 * 8 + 7 = 103

Q:

Overall this 2bit quants perform quite good! Thanks @ikawrakow

[Ttl]

It looks like the output logit norm of llama-v2-7b-2.20bpw is slightly lower than fp16 model's output norm. After subtracting the mean, the output norm of quantized model is 1.035 times smaller. Multiplying that into the output weights and biases or just multiplying the output logits reduces the perplexity on wiki.test from 7.0681 to 7.0080. It's unlikely to matter much in practice as it just applies a very small temperature to output softmax, but just for optimizing numbers it does help a little.

Testing on wiki.test 2.8% of time the top token of llama-v2-7b-2.20bpw was not in top-10 of fp16 model output tokens. For Q2_K it's 0.85% and 0.10% for Q4_K_S.

There are 2100 instances in wiki.test where difference in probability of fp16 model's top-1 token is larger than 90% out of 280k tokens in the text. The same number is 0 for Q4_K_S so there is a definitely something lost in the 2-bit quantization. Looking at the differences most of them are dates, names and other facts that 2-bit model seems to have forgotten.

Some cherry picked examples with 0 temperature:

"American Beauty is a 1999 American drama film directed by"
fp16 model: " Sam Mendes and written by Alan Ball.". 99% probability for "Sam". (This is correct)
2.20 bpw: " Allan F.reviewed by the critics". 5% probability for "All".

"QuackShot , known in Japan as I Love Donald Duck : Georgia Ou no Hihou ( Japanese : アイラブドナルドダック グルジア王の秘"
fp16 model: "宝, ". 99% probability. (This is correct)
2.20 bpw: "増, ". 50% probability.

"Undertaker then returned and chokeslam"
fp16 model: "med Kane through the announce table.". 98% for "med".
2.20 bpw: "'d him, but he was able to escape." 5% for "'".

Last example generates reasonable text on CPU but garbage when run with CUDA: ./main -m llama-v2-7b-2.20bpw.gguf -p "Undertaker then returned and chokeslam" --top-k 1 -ngl 99. I'm not sure if this is a problem with this PR or #4755.

Still very good results for this extreme quantization.

[x4080]

@jxy Thanks for the tip. I'm glad that the 2bit works good, how many VRAM do you have to run the mixtral 2bit ?

[x4080]

Is the convert to 2bit available anywhere ?

Tried nous-hermes-2-34b-2.69bpw.gguf on my m2 16gb and it wont even start, maybe not enough memory ?

[JianbangZ]

Tested some mixtral 8x7 models. Generation speed is great, about 60% of single 7B model. But what is up with the prompt evaluation speed? Why it’s 1/7 of the 7B model? Some implementation inefficiency?

I have a RTX 6000 Ada, gen speed is about 80t/s, prompt evaluation speed is 480 t/s without 1k input (2k window). While 7B model is yielding 3300 t/s prompt evaluation speed and 120 t/s gen speed.

[joseph777111]

@ikawrakow Will the techniques used to make pure 2 quants be used to improve the larger quantizations as well? 🤔🤩

[ikawrakow]

Tested some mixtral 8x7 models. Generation speed is great, about 60% of single 7B model. But what is up with the prompt evaluation speed? Why it’s 1/7 of the 7B model? Some implementation inefficiency?

I have a RTX 6000 Ada, gen speed is about 80t/s, prompt evaluation speed is 480 t/s without 1k input (2k window). While 7B model is yielding 3300 t/s prompt evaluation speed and 120 t/s gen speed.

For prompt processing Mixtral is equivalent to a 46B model rather than a 13B model: each token from the prompt is given to a different set of 2 "experts", so one always ends up with all 8 "experts" being involved, so the model has effectively 46B parameters, so yes, about 1/7 the speed of a 7B model. I was confused about this myself as well due to the misleading use of "Mixture Of Experts". The 8 different feed forward networks are not really "experts", which are turned on and remain on based on the context being processed. Instead, this architecture allows to enlarge the feed forward network, use all of it when evaluating a prompt, but use 1/4 of it when generating tokens one-by-one.

[JianbangZ]

Tested some mixtral 8x7 models. Generation speed is great, about 60% of single 7B model. But what is up with the prompt evaluation speed? Why it’s 1/7 of the 7B model? Some implementation inefficiency?
I have a RTX 6000 Ada, gen speed is about 80t/s, prompt evaluation speed is 480 t/s without 1k input (2k window). While 7B model is yielding 3300 t/s prompt evaluation speed and 120 t/s gen speed.

For prompt processing Mixtral is equivalent to a 46B model rather than a 13B model: each token from the prompt is given to a different set of 2 "experts", so one always ends up with all 8 "experts" being involved, so the model has effectively 46B parameters, so yes, about 1/7 the speed of a 7B model. I was confused about this myself as well due to the misleading use of "Mixture Of Experts". The 8 different feed forward networks are not really "experts", which are turned on and remain on based on the context being processed. Instead, this architecture allows to enlarge the feed forward network, use all of it when evaluating a prompt, but use 1/4 of it when generating tokens one-by-one.

I get it now. Basically it's because of batch. All the experts were used for the entire batch of tokens (256, 512 etc.)。Might worth doing some experiment on "batched top-2 experts" to fix the top2/3 experts for the entire batch. Welcome other ideas to improve the prompt eval speed.

[x4080]

Hi, i just downloaded mistral-instruct-7b-2.43bpw.gguf but I got error when tried to running it, so I run it again with minimal config, but still the same error

./main -m ./models/mistral-instruct-7b-2.43bpw.gguf -p "Hello"

the error

...
llm_load_tensors: system memory used  = 6577.58 MiB
................................................................................................GGML_ASSERT: ggml-backend.c:1274: (char *)addr + ggml_backend_buffer_get_alloc_size(buffer, tensor) <= (char *)ggml_backend_buffer_get_base(buffer) + ggml_backend_buffer_get_size(buffer)
zsh: abort      ./main -m ./models/mistral-instruct-7b-2.43bpw.gguf -p "Hello"

system : m2 16gb

[tsengalb99]

@ikawrakow I just saw this - exciting to see some of the stuff from quip# being ported over! A few comments

• We don't actually "force" a number of sign flips during quantization. The shifted E8 lattice has the property that an even number of sign flips keeps you in it, and an odd number of sign flips moves you out of the lattice. All lattice points can be represented as an either an even number of sign flips from an all positive vector or an even number of sign flips from a vector with only one negative entry. Since the absv codebook is all positive, this means that every entry corresponds to either a even or odd number of flips (but not both) and so we can get away with storing 7/8 sign flips.
• You may want to try the fine tuning approach I posted about yesterday (more details in our arxiv). It seems to make a pretty big difference for smaller models. For ex, with our perplexity measuring method, 7B 2 bit goes from 8.2 to 6.2 ppl (fp16 is 5.x).

[ikawrakow]

@tsengalb99 Thanks for the update. The results of your latest paper are very impressive. Congratulations!

Once fine tuning/training gets involved in the quantization as it is the case in your latest paper or the AQLM paper, we are out of the competition in this repository due to the limited llama.cpp capabilities in that regard. But we did give you a good run for your money, I think (despite being consistently ignored in quantization related publications).

Just out of curiosity, how long does it take to quantize one of these models with QuIP#?

[tsengalb99]

With QuIP#, fine tuning doesn’t add any complexity to inference. This is probably the same with llama.cpp and E8 so you could finetune entirely offline and the final model should still work. My fine-tuning scripts are not very optimized and quantizing 70B with fine tuning takes about 60 GPU hours. Someone who actually knows how to write LLM training code could get that number down a lot. The hard part about citing llama.cpp is that, at least to my knowledge, 1) kquants is not well documented 2) the thing you’re doing is also not well documented 3) there isn’t a full set of eval numbers we can just pull off a paper where the fp16 matches ours. The first 2 things are important so people know exactly what a method does when comparing and can build upon your work if there’s important insights. The 3rd is useful since it allows Llama.cpp to be compared against without needing to rerun eval. There are already differences between papers using 2048 or 4096 context length for Llama 2 and having a 3rd method where the fp16 doesn’t match makes it even more annoying to compare methods. Almost all academic papers use the OPTQ perplexity sampling functions and LM eval to do zeroshot. These should be pretty easy to add into llama.cpp. From: Kawrakow ***@***.***> Sent: Tuesday, February 13, 2024 12:58 PM To: ggerganov/llama.cpp ***@***.***> Cc: Albert Tseng ***@***.***>; Mention ***@***.***> Subject: Re: [ggerganov/llama.cpp] SOTA 2-bit quants (PR #4773) @tsengalb99 <https://github.com/tsengalb99> Thanks for the update. The results of your latest paper are very impressive. Congratulations! Once fine tuning/training gets involved in the quantization as it is the case in your latest paper or the AQLM paper, we are out of the competition in this repository due to the limited llama.cpp capabilities in that regard. But we did give you a good run for your money, I think (despite being consistently ignored in quantization related publications). Just out of curiosity, how long does it take to quantize one of these models with QuIP#? — Reply to this email directly, view it on GitHub <#4773 (comment)> , or unsubscribe <https://github.com/notifications/unsubscribe-auth/AH6WZSGV7AZFKHL2QG2K3ZLYTOSRRAVCNFSM6AAAAABBM5O2LSVHI2DSMVQWIX3LMV43OSLTON2WKQ3PNVWWK3TUHMYTSNBSGEYDOMJYGI> . You are receiving this because you were mentioned.Message ID: ***@***.***>