Revert "fix bug(breaking change) remove (zeors -= 1)"

auto_gptq/modeling/_base.py

@@ -39,7 +39,7 @@ class BaseQuantizeConfig(PushToHubMixin):
     damp_percent: float = field(default=0.01)
     desc_act: bool = field(default=True)
     static_groups: bool = field(default=False)
-    sym: bool = field(default=False)
+    sym: bool = field(default=True)
     true_sequential: bool = field(default=True)
     model_name_or_path: Optional[str] = field(default=None)
     model_file_base_name: Optional[str] = field(default=None)
@@ -967,27 +967,6 @@ class BaseGPTQForCausalLM(nn.Module, PushToHubMixin):
                 checkpoint
             )
             model.load_state_dict(checkpoint)
-        # Preprocessing for backward compatibility
-        if quantize_config.sym:
-            QuantLinear = dynamically_import_QuantLinear(use_triton=use_triton, disable_exllama=disable_exllama, use_qigen=use_qigen,
-                                                         desc_act=quantize_config.desc_act, group_size=quantize_config.group_size, bits=quantize_config.bits)
-            for name, submodule in model.named_modules():
-                if isinstance(submodule, QuantLinear):
-                    if use_qigen:
-                        submodule.zeros.data = torch.full_like(submodule.zeros.data, (torch.tensor(2 ** quantize_config.bits - 1) + 1) / 2)
-                    else:
-                        if quantize_config.bits == 2:
-                            submodule.qzeros.data = torch.full_like(submodule.qzeros.data, -1431655766)
-                        elif quantize_config.bits == 3:
-                            submodule.qzeros.data[:,range(0,submodule.qzeros.data.shape[1],3)] = 613566756
-                            submodule.qzeros.data[:,range(1,submodule.qzeros.data.shape[1],3)] = 1227133513
-                            submodule.qzeros.data[:,range(2,submodule.qzeros.data.shape[1],3)] = -1840700270
-                        elif quantize_config.bits == 4:
-                            submodule.qzeros.data = torch.full_like(submodule.qzeros.data, -2004318072)
-                        elif quantize_config.bits == 8:
-                            submodule.qzeros.data = torch.full_like(submodule.qzeros.data, -2139062144)
-                        else:
-                            raise NotImplementedError("Only 2,3,4,8 bits are supported.")
             # == step4: set seqlen == #
         model_config = model.config.to_dict()
         seq_len_keys = ["max_position_embeddings", "seq_length", "n_positions"]

auto_gptq/nn_modules/fused_gptj_attn.py

@@ -8,8 +8,6 @@ from transformers.models.gptj.modeling_gptj import GPTJAttention
 from ._fused_base import FusedBaseAttentionModule
 from ..utils.import_utils import compare_pytorch_version, dynamically_import_QuantLinear
 
-from logging import getLogger
-logger = getLogger(__name__)
 
 def fixed_pos_embedding(x, seq_dim=1, seq_len=None):
     dim = x.shape[-1]
@@ -242,12 +240,7 @@ class FusedGPTJAttentionForQuantizedModel(FusedBaseAttentionModule):
         **kwargs
     ):
         config = model.config
-        
         QuantLinear = dynamically_import_QuantLinear(use_triton=use_triton, desc_act=desc_act, group_size=group_size, bits=bits, disable_exllama=disable_exllama, disable_exllamav2=disable_exllamav2)
-        if QuantLinear.QUANT_TYPE in ["exllama", "exllamav2"] and desc_act:
-            # See fused_llama_attn.py comment
-            logger.warning(f"Exllama kernel does not support query/key/value fusion with act-order. Because of this, Fused attention is automatically disabled.")
-            return False
 
         for name, m in model.named_modules():
             if not isinstance(m, GPTJAttention):
@@ -264,6 +257,10 @@ class FusedGPTJAttentionForQuantizedModel(FusedBaseAttentionModule):
             scales = torch.cat([q_proj.scales, k_proj.scales, v_proj.scales], dim=1)
 
             if QuantLinear.QUANT_TYPE == "exllama":
+                if desc_act:
+                    # See fused_llama_attn.py comment
+                    raise ValueError("Exllama kernel does not support query/key/value fusion with act-order. Please either use inject_fused_attention=False or disable_exllama=True.")
+                else:
                     g_idx = None
             else:
                 g_idx = torch.cat([q_proj.g_idx, k_proj.g_idx, v_proj.g_idx], dim=0)
@@ -301,6 +298,6 @@ class FusedGPTJAttentionForQuantizedModel(FusedBaseAttentionModule):
 
             setattr(parent, child_name, attn)
             del m
-        return True
+
 
 __all__ = ["FusedGPTJAttentionForQuantizedModel"]

auto_gptq/nn_modules/fused_llama_attn.py

@@ -7,8 +7,6 @@ from transformers.models.llama.modeling_llama import LlamaAttention, apply_rotar
 from ._fused_base import FusedBaseAttentionModule
 from ..utils.import_utils import compare_pytorch_version, dynamically_import_QuantLinear
 
-from logging import getLogger
-logger = getLogger(__name__)
 
 class FusedLlamaAttentionForQuantizedModel(FusedBaseAttentionModule):
     """Multi-headed attention from 'Attention Is All You Need' paper"""
@@ -144,14 +142,7 @@ class FusedLlamaAttentionForQuantizedModel(FusedBaseAttentionModule):
         """
         Replace all LlamaAttention modules with QuantLlamaAttention modules, fusing the q, k, v projections.
         """
-
         QuantLinear = dynamically_import_QuantLinear(use_triton=use_triton, desc_act=desc_act, group_size=group_size, bits=bits, disable_exllama=disable_exllama, disable_exllamav2=disable_exllamav2)
-        if QuantLinear.QUANT_TYPE in ["exllama", "exllamav2"] and desc_act:
-            # TODO: support it. The issue lies maybe in the line:
-            # int groups = qzeros.size(0);
-            # in exllama_ext.cpp
-            logger.warning(f"Exllama kernel does not support query/key/value fusion with act-order. Because of this, Fused attention is automatically disabled.")
-            return False
 
         for name, m in model.named_modules():
             if not isinstance(m, LlamaAttention):
@@ -166,6 +157,12 @@ class FusedLlamaAttentionForQuantizedModel(FusedBaseAttentionModule):
             scales = torch.cat([q_proj.scales, k_proj.scales, v_proj.scales], dim=1)
 
             if QuantLinear.QUANT_TYPE == "exllama":
+                if desc_act:
+                    # TODO: support it. The issue lies maybe in the line:
+                    # int groups = qzeros.size(0);
+                    # in exllama_ext.cpp
+                    raise ValueError("Exllama kernel does not support query/key/value fusion with act-order. Please either use inject_fused_attention=False or disable_exllama=True.")
+                else:
                     g_idx = None
             else:
                 g_idx = torch.cat([q_proj.g_idx, k_proj.g_idx, v_proj.g_idx], dim=0)
@@ -201,7 +198,6 @@ class FusedLlamaAttentionForQuantizedModel(FusedBaseAttentionModule):
                 child_name = name
 
             setattr(parent, child_name, attn)
-        return True
 
 
 __all__ = ["FusedLlamaAttentionForQuantizedModel"]

auto_gptq/nn_modules/qlinear/qlinear_cuda.py

@@ -157,6 +157,7 @@ class QuantLinear(nn.Module):
         qweight = qweight.astype(np.int32)
         self.qweight = torch.from_numpy(qweight)
 
+        zeros -= 1
         zeros = zeros.numpy().astype(np.uint32)
         qzeros = np.zeros((zeros.shape[0], zeros.shape[1] // 32 * self.bits), dtype=np.uint32)
         i = 0
@@ -220,6 +221,7 @@ class QuantLinear(nn.Module):
                 ).to(torch.int16 if self.bits == 8 else torch.int8)
                 torch.bitwise_and(zeros, (2 ** self.bits) - 1, out=zeros)
 
+                zeros = zeros + 1
                 zeros = zeros.reshape(self.scales.shape)
 
                 weight = torch.bitwise_right_shift(
@@ -237,6 +239,7 @@ class QuantLinear(nn.Module):
                 zeros = zeros & 0x7
                 zeros = torch.cat([zeros[:, :, 0, :11], zeros[:, :, 1, 1:12], zeros[:, :, 2, 1:11]], dim=2)
 
+                zeros = zeros + 1
                 zeros = zeros.reshape(self.scales.shape)
 
                 weight = self.qweight.reshape(

auto_gptq/nn_modules/qlinear/qlinear_cuda_old.py

@@ -157,6 +157,7 @@ class QuantLinear(nn.Module):
         qweight = qweight.astype(np.int32)
         self.qweight = torch.from_numpy(qweight)
 
+        zeros -= 1
         zeros = zeros.numpy().astype(np.uint32)
         qzeros = np.zeros((zeros.shape[0], zeros.shape[1] // 32 * self.bits), dtype=np.uint32)
         i = 0
@@ -230,6 +231,7 @@ class QuantLinear(nn.Module):
                zeros = torch.bitwise_right_shift(torch.unsqueeze(self.qzeros, 2).expand(-1, -1, 32 // self.bits), self.wf.unsqueeze(0)).to(torch.int16 if self.bits == 8 else torch.int8)
                torch.bitwise_and(zeros, (2 ** self.bits) - 1, out=zeros)
                    
+               zeros = zeros + 1
                zeros = zeros.reshape(-1, 1, zeros.shape[1] * zeros.shape[2])
    
                scales = self.scales
@@ -246,6 +248,7 @@ class QuantLinear(nn.Module):
                zeros = zeros & 0x7
                zeros = torch.cat([zeros[:,:,0,:11], zeros[:,:,1,1:12], zeros[:,:,2,1:11]], dim=2)
                 
+               zeros = zeros + 1
                zeros = zeros.reshape(-1, 1, zeros.shape[1] * zeros.shape[2]) 
                
                scales = self.scales

auto_gptq/nn_modules/qlinear/qlinear_exllama.py

@@ -146,6 +146,7 @@ class QuantLinear(nn.Module):
         qweight = qweight.astype(np.int32)
         self.qweight = torch.from_numpy(qweight)
 
+        zeros -= 1
         zeros = zeros.numpy().astype(np.uint32)
         qzeros = np.zeros((zeros.shape[0], zeros.shape[1] // 32 * self.bits), dtype=np.uint32)
         i = 0

auto_gptq/nn_modules/qlinear/qlinear_triton.py

@@ -114,6 +114,7 @@ class QuantLinear(nn.Module, TritonModuleMixin):
         qweight = qweight.astype(np.int32)
         self.qweight = torch.from_numpy(qweight)
 
+        zeros -= 1
         zeros = zeros.numpy().astype(np.uint32)
         qzeros = np.zeros((zeros.shape[0], zeros.shape[1] // 32 * self.bits), dtype=np.uint32)
         i = 0

auto_gptq/nn_modules/triton_utils/kernels.py

@@ -144,6 +144,7 @@ def quant_matmul_248_kernel(
         zeros = tl.load(zeros_ptrs + g_idx[:, None] * stride_zeros)  # (BLOCK_SIZE_K, BLOCK_SIZE_N,)
 
         zeros = (zeros >> zeros_shifter[None, :]) & maxq
+        zeros = (zeros + 1)
 
         a = tl.load(a_ptrs, mask=a_mask, other=0.)  # (BLOCK_SIZE_M, BLOCK_SIZE_K)
         b = tl.load(b_ptrs)  # (BLOCK_SIZE_K, BLOCK_SIZE_N), but repeated
@@ -289,6 +290,7 @@ def transpose_quant_matmul_248_kernel(
         zeros = tl.load(zeros_ptrs)  # (BLOCK_SIZE_K, BLOCK_SIZE_N,)
 
         zeros = (zeros >> zeros_shifter[None, :]) & maxq
+        zeros = (zeros + 1)
 
         a = tl.load(a_ptrs, mask=a_mask, other=0.)  # (BLOCK_SIZE_M, BLOCK_SIZE_N)
         b = tl.load(b_ptrs)  # (BLOCK_SIZE_K, BLOCK_SIZE_N), but repeated

autogptq_extension/cuda_256/autogptq_cuda_kernel_256.cu

@@ -30,9 +30,9 @@
 // }
 // #endif
 
-
 #if (defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 700) || defined(USE_ROCM)
 // adapted from https://github.com/torch/cutorch/blob/master/lib/THC/THCAtomics.cuh
+
 __device__ __forceinline__ void atomicAdd(c10::Half* address, c10::Half val) {
     unsigned int *address_as_ui = reinterpret_cast<unsigned int *>(reinterpret_cast<char *>(address) - (reinterpret_cast<size_t>(address) & 2));
     unsigned int old = *address_as_ui;
@@ -113,7 +113,6 @@ __global__ void VecQuant4MatMulKernel(
 	int zero_width
 );
 
-
 template <typename scalar_t>
 __global__ void VecQuant8MatMulKernel(
     const  scalar_t* __restrict__ vec,
@@ -313,7 +312,7 @@ __global__ void VecQuant2MatMulKernel(
 
     g = as_int(g_idx[g_h + k]);
     scalar_t scale = scales[g * width + w];
-    scalar_t zero = scalar_t(as_unsigned(zeros[g * zero_width + z_w]) >> z_mod & 0x3);
+    scalar_t zero = scalar_t((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod & 0x3) + 1);
 
     w_tmp = ((as_unsigned(mat[i + (k_w * width)]) >> k_bit) & 0x3);
 
@@ -447,12 +446,12 @@ __global__ void VecQuant3MatMulKernel(
     scalar_t zero;
     if (z_mod == 10) {
       z_tmp = (as_unsigned(zeros[g * zero_width + z_w]) >> 30) | ((as_unsigned(zeros[g * zero_width + (z_w + 1)]) << 2) & 0x4);
-      zero = scalar_t(z_tmp);
+      zero = scalar_t((z_tmp) + 1);
     } else if (z_mod == 21){
       z_tmp = (as_unsigned(zeros[g * zero_width + z_w]) >> 31) | ((as_unsigned(zeros[g * zero_width + (z_w + 1)]) << 1) & 0x6);
-      zero = scalar_t(z_tmp);
+      zero = scalar_t((z_tmp) + 1);
     } else {
-      zero = scalar_t((as_unsigned(zeros[g * zero_width + z_w]) >> z_bit) & 0x7);
+      zero = scalar_t(((as_unsigned(zeros[g * zero_width + z_w]) >> z_bit) & 0x7) + 1);
     }
 
     if (k_mod == 10) {
@@ -546,7 +545,7 @@ __global__ void VecQuant4MatMulKernel(
 
     g = as_int(g_idx[g_h + k]);
     scalar_t scale = scales[g * width + w];
-    scalar_t zero = scalar_t((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xF);
+    scalar_t zero = scalar_t(((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xF) + 1);
 
     w_tmp = ((as_unsigned(mat[i + (k_w * width)]) >> k_bit) & 0xF);
 
@@ -633,7 +632,7 @@ __global__ void VecQuant8MatMulKernel(
 
     g = as_int(g_idx[g_h + k]);
     scalar_t scale = scales[g * width + w];
-    scalar_t zero = scalar_t((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xFF);
+    scalar_t zero = scalar_t(((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xFF) + 1);
 
     w_tmp = ((as_unsigned(mat[i + (k_w * width)]) >> k_bit) & 0xFF);
 
@@ -724,7 +723,7 @@ __global__ void VecQuant2MatMulKernel_old(
 
     int g = (g_h + k) / groupsize;
     scalar_t scale = scales[g * width + w];
-    scalar_t zero = scale * scalar_t(as_unsigned(zeros[g * zero_width + z_w]) >> z_mod & 0x3);
+    scalar_t zero = scale * scalar_t((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod & 0x3) + 1);
 
     res += (scale * scalar_t((tmp >> 0) & 0x3) - zero) * blockvec[k + 0];
     res += (scale * scalar_t((tmp >> 2) & 0x3) - zero) * blockvec[k + 1];
@@ -847,12 +846,12 @@ __global__ void VecQuant3MatMulKernel_old(
     scalar_t zero;
     if (z_mod == 10) {
       z_tmp = (as_unsigned(zeros[g * zero_width + z_w]) >> 30) | ((as_unsigned(zeros[g * zero_width + (z_w + 1)]) << 2) & 0x4);
-      zero = scale * scalar_t(z_tmp);
+      zero = scale * scalar_t((z_tmp) + 1);
     } else if (z_mod == 21){
       z_tmp = (as_unsigned(zeros[g * zero_width + z_w]) >> 31) | ((as_unsigned(zeros[g * zero_width + (z_w + 1)]) << 1) & 0x6);
-      zero = scale * scalar_t(z_tmp);
+      zero = scale * scalar_t((z_tmp) + 1);
     } else {
-      zero = scale * scalar_t((as_unsigned(zeros[g * zero_width + z_w]) >> z_bit) & 0x7);
+      zero = scale * scalar_t(((as_unsigned(zeros[g * zero_width + z_w]) >> z_bit) & 0x7) + 1);
     }
 
     res += (scale * scalar_t((tmp1 >>  0) & 0x7) - zero) * blockvec[k + 0];
@@ -978,7 +977,7 @@ __global__ void VecQuant4MatMulKernel_old(
 
     int g = (g_h + k) / groupsize;
     scalar_t scale = scales[g * width + w];
-    scalar_t zero = scale * scalar_t((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xF);
+    scalar_t zero = scale * scalar_t(((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xF) + 1);
 
     res += (scale * scalar_t((tmp >> 0) & 0xF) - zero) * blockvec[k + 0];
     res += (scale * scalar_t((tmp >> 4) & 0xF) - zero) * blockvec[k + 1];
@@ -1065,7 +1064,7 @@ __global__ void VecQuant8MatMulKernel_old(
 
     int g = (g_h + k) / groupsize;
     scalar_t scale = scales[g * width + w];
-    scalar_t zero = scale * scalar_t((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xFF);
+    scalar_t zero = scale * scalar_t(((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xFF) + 1);
 
     res += (scale * scalar_t((tmp >> 0) & 0xFF) - zero) * blockvec[k + 0];
     res += (scale * scalar_t((tmp >> 8) & 0xFF) - zero) * blockvec[k + 1];
@@ -1160,7 +1159,7 @@ __global__ void VecQuant2MatMulKernelFaster_old(
     int g = (g_h + (k * 2)) / groupsize;
 	float scale_f = scales[g * width + w];
     half2 scale = __float2half2_rn(scale_f);
-    half2 zero = __float2half2_rn(-(scale_f * ((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0x3)));
+    half2 zero = __float2half2_rn(-(scale_f * (((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0x3) + 1)));
 
     std::memset(&res2, 0, sizeof(half2));
     tmp = as_unsigned(mat[i]);
@@ -1288,12 +1287,12 @@ __global__ void VecQuant3MatMulKernelFaster_old(
     half2 zero;
     if (z_mod == 10) {
       z_tmp = (as_unsigned(zeros[g * zero_width + z_w]) >> 30) | ((as_unsigned(zeros[g * zero_width + (z_w + 1)]) << 2) & 0x4);
-      zero = __float2half2_rn(-(scale_f * z_tmp));
+      zero = __float2half2_rn(-(scale_f * ((z_tmp) + 1)));
     } else if (z_mod == 21){
       z_tmp = (as_unsigned(zeros[g * zero_width + z_w]) >> 31) | ((as_unsigned(zeros[g * zero_width + (z_w + 1)]) << 1) & 0x6);
-      zero = __float2half2_rn(-(scale_f * z_tmp));
+      zero = __float2half2_rn(-(scale_f * ((z_tmp) + 1)));
     } else {
-      zero = __float2half2_rn(-(scale_f * ((as_unsigned(zeros[g * zero_width + z_w]) >> z_bit) & 0x7)));
+      zero = __float2half2_rn(-(scale_f * (((as_unsigned(zeros[g * zero_width + z_w]) >> z_bit) & 0x7) + 1)));
     }
 
     std::memset(&res2, 0, sizeof(half2));
@@ -1411,8 +1410,13 @@ __global__ void VecQuant4MatMulKernelFaster_old(
   while (k < blockwidth2) {
     int g = (g_h + (k * 2)) / groupsize;
 	float scale_f = scales[g * width + w];
+
     half2 scale = __float2half2_rn(scale_f);
-    half2 zero = __float2half2_rn(-(scale_f * ((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xF)));
+    half2 zero = __float2half2_rn(-(scale_f * (((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xF) + 1)));
+
+    //std::memset(&res2, 0, sizeof(half2));
+
+    //res2 = __float2half2_rn((float)0.);
 
     std::memset(&res2, 0, sizeof(half2));
     tmp = as_unsigned(mat[i]);
@@ -1422,7 +1426,9 @@ __global__ void VecQuant4MatMulKernelFaster_old(
     res2 = __hfma2(__hfma2(deq2[(tmp >> 24) & 0xff][off], scale, zero), blockvec[k + 3], res2);
 	i += width;
     k += 4;
+
     res += __low2float(res2) + __high2float(res2);
+
   }
 
   atomicAdd(&mul[b * width + w], res);

autogptq_extension/cuda_64/autogptq_cuda_kernel_64.cu

@@ -313,7 +313,7 @@ __global__ void VecQuant2MatMulKernel(
 	
     g = as_int(g_idx[g_h + k]);
     scalar_t scale = scales[g * width + w];
-    scalar_t zero = scalar_t(as_unsigned(zeros[g * zero_width + z_w]) >> z_mod & 0x3);
+    scalar_t zero = scalar_t((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod & 0x3) + 1);
 	
     w_tmp = ((as_unsigned(mat[i + (k_w * width)]) >> k_bit) & 0x3);
     
@@ -447,12 +447,12 @@ __global__ void VecQuant3MatMulKernel(
     scalar_t zero;
     if (z_mod == 10) {
       z_tmp = (as_unsigned(zeros[g * zero_width + z_w]) >> 30) | ((as_unsigned(zeros[g * zero_width + (z_w + 1)]) << 2) & 0x4);
-      zero = scalar_t(z_tmp);
+      zero = scalar_t((z_tmp) + 1);
     } else if (z_mod == 21){
       z_tmp = (as_unsigned(zeros[g * zero_width + z_w]) >> 31) | ((as_unsigned(zeros[g * zero_width + (z_w + 1)]) << 1) & 0x6);
-      zero = scalar_t(z_tmp);
+      zero = scalar_t((z_tmp) + 1);
     } else {
-      zero = scalar_t((as_unsigned(zeros[g * zero_width + z_w]) >> z_bit) & 0x7);
+      zero = scalar_t(((as_unsigned(zeros[g * zero_width + z_w]) >> z_bit) & 0x7) + 1);
     }
 	
     if (k_mod == 10) {
@@ -546,7 +546,7 @@ __global__ void VecQuant4MatMulKernel(
 	
     g = as_int(g_idx[g_h + k]);
     scalar_t scale = scales[g * width + w];
-    scalar_t zero = scalar_t((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xF);
+    scalar_t zero = scalar_t(((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xF) + 1);
 	
     w_tmp = ((as_unsigned(mat[i + (k_w * width)]) >> k_bit) & 0xF);
     
@@ -633,7 +633,7 @@ __global__ void VecQuant8MatMulKernel(
 	
     g = as_int(g_idx[g_h + k]);
     scalar_t scale = scales[g * width + w];
-    scalar_t zero = scalar_t((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xFF);
+    scalar_t zero = scalar_t(((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xFF) + 1);
 	
     w_tmp = ((as_unsigned(mat[i + (k_w * width)]) >> k_bit) & 0xFF);
     
@@ -724,7 +724,7 @@ __global__ void VecQuant2MatMulKernel_old(
 	
     int g = (g_h + k) / groupsize;
     scalar_t scale = scales[g * width + w];
-    scalar_t zero = scale * scalar_t(as_unsigned(zeros[g * zero_width + z_w]) >> z_mod & 0x3);
+    scalar_t zero = scale * scalar_t((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod & 0x3) + 1);
 	
     res += (scale * scalar_t((tmp >> 0) & 0x3) - zero) * blockvec[k + 0];
     res += (scale * scalar_t((tmp >> 2) & 0x3) - zero) * blockvec[k + 1];
@@ -847,12 +847,12 @@ __global__ void VecQuant3MatMulKernel_old(
     scalar_t zero;
     if (z_mod == 10) {
       z_tmp = (as_unsigned(zeros[g * zero_width + z_w]) >> 30) | ((as_unsigned(zeros[g * zero_width + (z_w + 1)]) << 2) & 0x4);
-      zero = scale * scalar_t(z_tmp);
+      zero = scale * scalar_t((z_tmp) + 1);
     } else if (z_mod == 21){
       z_tmp = (as_unsigned(zeros[g * zero_width + z_w]) >> 31) | ((as_unsigned(zeros[g * zero_width + (z_w + 1)]) << 1) & 0x6);
-      zero = scale * scalar_t(z_tmp);
+      zero = scale * scalar_t((z_tmp) + 1);
     } else {
-      zero = scale * scalar_t((as_unsigned(zeros[g * zero_width + z_w]) >> z_bit) & 0x7);
+      zero = scale * scalar_t(((as_unsigned(zeros[g * zero_width + z_w]) >> z_bit) & 0x7) + 1);
     }
 	
     res += (scale * scalar_t((tmp1 >>  0) & 0x7) - zero) * blockvec[k + 0];
@@ -978,7 +978,7 @@ __global__ void VecQuant4MatMulKernel_old(
 	
     int g = (g_h + k) / groupsize;
     scalar_t scale = scales[g * width + w];
-    scalar_t zero = scale * scalar_t((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xF);
+    scalar_t zero = scale * scalar_t(((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xF) + 1);
 	
     res += (scale * scalar_t((tmp >> 0) & 0xF) - zero) * blockvec[k + 0];
     res += (scale * scalar_t((tmp >> 4) & 0xF) - zero) * blockvec[k + 1];
@@ -1065,7 +1065,7 @@ __global__ void VecQuant8MatMulKernel_old(
 	
     int g = (g_h + k) / groupsize;
     scalar_t scale = scales[g * width + w];
-    scalar_t zero = scale * scalar_t((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xFF);
+    scalar_t zero = scale * scalar_t(((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xFF) + 1);
 	
     res += (scale * scalar_t((tmp >> 0) & 0xFF) - zero) * blockvec[k + 0];
     res += (scale * scalar_t((tmp >> 8) & 0xFF) - zero) * blockvec[k + 1];
@@ -1160,7 +1160,7 @@ __global__ void VecQuant2MatMulKernelFaster_old(
     int g = (g_h + (k * 2)) / groupsize;
 	float scale_f = scales[g * width + w];
     half2 scale = __float2half2_rn(scale_f);
-    half2 zero = __float2half2_rn(-(scale_f * ((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0x3)));
+    half2 zero = __float2half2_rn(-(scale_f * (((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0x3) + 1)));
 	
     std::memset(&res2, 0, sizeof(half2));
     tmp = as_unsigned(mat[i]);
@@ -1288,12 +1288,12 @@ __global__ void VecQuant3MatMulKernelFaster_old(
     half2 zero;
     if (z_mod == 10) {
       z_tmp = (as_unsigned(zeros[g * zero_width + z_w]) >> 30) | ((as_unsigned(zeros[g * zero_width + (z_w + 1)]) << 2) & 0x4);
-      zero = __float2half2_rn(-(scale_f * z_tmp));
+      zero = __float2half2_rn(-(scale_f * ((z_tmp) + 1)));
     } else if (z_mod == 21){
       z_tmp = (as_unsigned(zeros[g * zero_width + z_w]) >> 31) | ((as_unsigned(zeros[g * zero_width + (z_w + 1)]) << 1) & 0x6);
-      zero = __float2half2_rn(-(scale_f * z_tmp));
+      zero = __float2half2_rn(-(scale_f * ((z_tmp) + 1)));
     } else {
-      zero = __float2half2_rn(-(scale_f * ((as_unsigned(zeros[g * zero_width + z_w]) >> z_bit) & 0x7)));
+      zero = __float2half2_rn(-(scale_f * (((as_unsigned(zeros[g * zero_width + z_w]) >> z_bit) & 0x7) + 1)));
     }
 	
     std::memset(&res2, 0, sizeof(half2));
@@ -1412,7 +1412,7 @@ __global__ void VecQuant4MatMulKernelFaster_old(
     int g = (g_h + (k * 2)) / groupsize;
 	float scale_f = scales[g * width + w];
     half2 scale = __float2half2_rn(scale_f);
-    half2 zero = __float2half2_rn(-(scale_f * ((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xF)));
+    half2 zero = __float2half2_rn(-(scale_f * (((as_unsigned(zeros[g * zero_width + z_w]) >> z_mod) & 0xF) + 1)));
 	
     std::memset(&res2, 0, sizeof(half2));
     tmp = as_unsigned(mat[i]);

autogptq_extension/exllama/cuda_func/q4_matmul.cu

@@ -13,8 +13,6 @@
 const int THREADS_X = 32;       // Block size and thread count along columns in w and out
 const int THREADS_Y = 1;        // Block size and thread count along rows in x and out
 
-const int GROUP_STEP = 32;      // Assumed group size when block_size_z % groupsize != 0
-
 typedef void (*fp_q4_matmul_kernel)
 (
     const half*,
@@ -48,15 +46,12 @@ __global__ void q4_matmul_kernel
     bool no_zero
 )
 {
-    extern __shared__ half2 x_cache[];
-    half* x_cache_h = (half*)x_cache;
-
     // Start of block
 
     int x_column = block_size_z * blockIdx.z;
     int x_column_end = min(dim, block_size_z * (blockIdx.z + 1));
 
-    int w_column = THREADS_X * blockIdx.x + threadIdx.x;  // assume width of weight matrix divisible by THREADS_X
+    int w_column = THREADS_X * blockIdx.x + threadIdx.x;
     int x_row = THREADS_Y * blockIdx.y + threadIdx.y;
 
     int iterations = (x_column_end - x_column) / 8;
@@ -74,8 +69,8 @@ __global__ void q4_matmul_kernel
     if (!no_zero && blockIdx.z == 0 && (threadIdx.x & 1) == 0)
     {
         *((uint32_t*) out_.item_ptr(x_row, w_column)) = 0;
-    }
         __syncthreads();
+    }
 
     // Loop over part of x row (and w column)
 
@@ -89,56 +84,48 @@ __global__ void q4_matmul_kernel
 
         for (int k = x_column, group = x_column / groupsize; k < x_column + iterations * 8; group++, k += groupsize)
         {
-            for (int i = threadIdx.x; i < groupsize; i += THREADS_X)
-            {
-                if constexpr (use_x_map) x_cache_h[i] = *x_.item_ptr(x_row, x_map[k + i]);
-                else                     x_cache_h[i] = *x_.item_ptr(x_row, k + i);
-            }
-            __syncthreads();
-
             if constexpr (use_half2)
             {
                 half2 w_scale = w_scales_.item_half2half2(group, w_column);
-                uint32_t w_zero = w_zeros_.item(group, w_column);
+                uint32_t w_zero = w_zeros_.item(group, w_column) + 1;
-                acc = dot_product_8(acc, x_cache, w_, k, w_column, w_scale, w_zero, groupsize / 8);
+
+                if constexpr (use_x_map) acc = dot_product_8_x_map(acc, x_, x_row, k, w_, k, w_column, w_scale, w_zero, groupsize / 8, x_map);
+                else                     acc = dot_product_8      (acc, x_, x_row, k, w_, k, w_column, w_scale, w_zero, groupsize / 8);
             }
             else
             {
                 half w_scale = w_scales_.item(group, w_column);
-                uint32_t w_zero = w_zeros_.item(group, w_column);
+                uint32_t w_zero = w_zeros_.item(group, w_column) + 1;
-                acc_h = dot_product_8_h(acc_h, x_cache_h, w_, k, w_column, w_scale, w_zero, groupsize / 8);
+
+                if constexpr (use_x_map) acc_h = dot_product_8_x_map_h(acc_h, x_, x_row, k, w_, k, w_column, w_scale, w_zero, groupsize / 8, x_map);
+                else                     acc_h = dot_product_8_h      (acc_h, x_, x_row, k, w_, k, w_column, w_scale, w_zero, groupsize / 8);
             }
-            __syncthreads();
         }
     }
     else
     {
-        // Otherwise assume groupsize is a multiple of GROUP_STEP, do GROUP_STEP columns per iteration and trust the cache
+        // Otherwise assume groupsize is a multiple of 8, do 8 columns per iteration and trust the cache
 
-        for (int k = x_column; k < x_column + iterations * 8; k += GROUP_STEP)
+        for (int k = x_column; k < x_column + iterations * 8; k += 8)
         {
-            for (int i = threadIdx.x; i < GROUP_STEP; i += THREADS_X)
-            {
-                if constexpr (use_x_map) x_cache_h[i] = *x_.item_ptr(x_row, x_map[k + i]);
-                else                     x_cache_h[i] = *x_.item_ptr(x_row, k + i);
-            }
-            __syncthreads();
-
             if constexpr (use_half2)
             {
                 int group = k / groupsize;
                 half2 w_scale = w_scales_.item_half2half2(group, w_column);
-                uint32_t w_zero = w_zeros_.item(group, w_column);
+                uint32_t w_zero = w_zeros_.item(group, w_column) + 1;
-                acc = dot_product_8(acc, x_cache, w_, k, w_column, w_scale, w_zero, GROUP_STEP / 8);
+
+                if constexpr (use_x_map) acc = dot_product_8_x_map(acc, x_, x_row, k, w_, k, w_column, w_scale, w_zero, 1, x_map);
+                else                     acc = dot_product_8      (acc, x_, x_row, k, w_, k, w_column, w_scale, w_zero, 1);
             }
             else
             {
                 int group = k / groupsize;
                 half w_scale = w_scales_.item(group, w_column);
-                uint32_t w_zero = w_zeros_.item(group, w_column);
+                uint32_t w_zero = w_zeros_.item(group, w_column) + 1;
-                acc_h = dot_product_8_h(acc_h, x_cache_h, w_, k, w_column, w_scale, w_zero, GROUP_STEP / 8);
+
+                if constexpr (use_x_map) acc_h = dot_product_8_x_map_h(acc_h, x_, x_row, k, w_, k, w_column, w_scale, w_zero, 1, x_map);
+                else                     acc_h = dot_product_8_h      (acc_h, x_, x_row, k, w_, k, w_column, w_scale, w_zero, 1);
             }
-            __syncthreads();
         }
     }
 
@@ -146,7 +133,7 @@ __global__ void q4_matmul_kernel
 
     if constexpr (use_half2)
     {
-        half result = __hadd(acc.x, acc.y);
+        half result = __hadd(__low2half(acc), __high2half(acc));
         atomicAdd(out_.item_ptr(x_row, w_column), result);
     }
     else
@@ -228,8 +215,8 @@ void q4_matmul_cuda
     );
 
     fp_q4_matmul_kernel kernel = q4_matmul_kernel_pick(tuningParams, block_size_z, w->groupsize, x_map);
-    int shared_mem = (block_size_z % w->groupsize == 0 ? w->groupsize : GROUP_STEP) * sizeof(half);
+
-    kernel<<<blocks, threads, shared_mem, alt_stream>>>(x_mapped, w->cuda_qweight, out, w->cuda_scales, w->cuda_qzeros, height, dim, width, w->groupsize, block_size_z, x_map, no_zero);
+    kernel<<<blocks, threads, 0, alt_stream>>> (x_mapped, w->cuda_qweight, out, w->cuda_scales, w->cuda_qzeros, height, dim, width, w->groupsize, block_size_z, x_map, no_zero);
 }
 
 void q4_matmul_recons_cuda
@@ -253,7 +240,7 @@ void q4_matmul_recons_cuda
     const half* x_mapped = x;
     if (w->cuda_x_map)
     {
-        TORCH_CHECK(buffers->temp_state_size >= x_height * dim, "temp_state buffer is too small");
+        TORCH_CHECK(buffers->temp_state_size >= x_height * dim, "The temp_state buffer is too small in the exllama backend. Please call the exllama_set_max_input_length function to increase the buffer size. Example:\nfrom auto_gptq import exllama_set_max_input_length\nmodel = exllama_set_max_input_length(model, 4096)");
         column_remap_cuda(x, buffers->temp_state, x_height, dim, w->cuda_x_map);
         x_mapped = buffers->temp_state;
     }
@@ -261,18 +248,13 @@ void q4_matmul_recons_cuda
     w->reconstruct(buffers->temp_dq);
 
 #if defined(__CUDA_ARCH__) && __CUDA_ARCH__ < 700
-
     const float alpha = 1.0f;
     const float beta = no_zero ? 1.0f : 0.0f;
     cublasSgemmEx(handle, CUBLAS_OP_N, CUBLAS_OP_N, width, height, dim, &alpha, buffers->temp_dq, CUDA_R_16F, width,
                   x_mapped, CUDA_R_16F, dim, &beta, out, CUDA_R_16F, width);
-
 #else
-
     const half alpha = __float2half(1.0f);
     const half beta = no_zero ? __float2half(1.0f) : __float2half(0.0f);
     cublasHgemm(handle, CUBLAS_OP_N, CUBLAS_OP_N, width, height, dim, &alpha, buffers->temp_dq, width, x_mapped, dim, &beta, out, width);
-
 #endif
-
 }

autogptq_extension/exllama/cuda_func/q4_matrix.cu

@@ -197,7 +197,7 @@ __global__ void reconstruct_kernel
     int group = row / groupsize;
 
     half w_scale = w_scales_.item(group, column);
-    uint32_t w_zero = w_zeros_.item(group, column);
+    uint32_t w_zero = w_zeros_.item(group, column) + 1;
 
     uint32_t w_read = w_.item_uint32_t(row, column);
     half* out_ptr = out_.item_ptr(row, column);

autogptq_extension/exllama/matrix.cuh

@@ -87,15 +87,18 @@ public:
 __device__ __forceinline__ half2 dot_product_8
 (
     const half2 acc,
-    const half2* h_ptr,
+    MatrixView_half& h_,
+    const int h_row,
+    const int h_column,                 // divisible by 8
     MatrixView_q4_column& v_,
     const int v_row,                    // divisible by 8
     const int v_column,
     const half2 v_scale_2,
-    const uint32_t v_zero,
+    const uint32_t v_zero,              // + 1 (!!)
     const int count
 )
 {
+    const half2* h_ptr = (const half2*) h_.item_ptr(h_row, h_column);
     const uint32_t* v_ptr = (const uint32_t*) v_.item_uint32_ptr(v_row, v_column);
     half2 result = acc;
 
@@ -135,15 +138,18 @@ __device__ __forceinline__ half2 dot_product_8
 __device__ __forceinline__ half dot_product_8_h
 (
     const half acc,
-    const half* h_ptr,
+    MatrixView_half& h_,
+    const int h_row,
+    const int h_column,                 // divisible by 8
     MatrixView_q4_column& v_,
     const int v_row,                    // divisible by 8
     const int v_column,
     const half v_scale,
-    const uint32_t v_zero,
+    const uint32_t v_zero,              // + 1 (!!)
     const int count
 )
 {
+    const half* h_ptr = h_.item_ptr(h_row, h_column);
     const uint32_t* v_ptr = (const uint32_t*) v_.item_uint32_ptr(v_row, v_column);
     half result = acc;
 
@@ -174,4 +180,115 @@ __device__ __forceinline__ half dot_product_8_h
     return result;
 }
 
+// Accumulated dot product of 8-element row vectors in h and quantized column vectors in v, constant zero/scale, with x_map
+
+__device__ __forceinline__ half2 dot_product_8_x_map
+(
+    const half2 acc,
+    MatrixView_half& h_,
+    const int h_row,
+    const int h_column,                 // divisible by 8
+    MatrixView_q4_column& v_,
+    const int v_row,                    // divisible by 8
+    const int v_column,
+    const half2 v_scale_2,
+    const uint32_t v_zero,              // + 1 (!!)
+    const int count,
+    const uint32_t* x_map
+)
+{
+    const half* h_ptr = h_.item_ptr(h_row, 0);
+    const uint32_t* x_map_ptr = x_map + h_column;
+    const uint32_t* v_ptr = (const uint32_t*) v_.item_uint32_ptr(v_row, v_column);
+    half2 result = acc;
+
+    for (int i = 0; i < count; i++)
+    {
+        uint32_t v_read = *v_ptr; v_ptr += v_.width;
+
+        half v_0 = __int2half_rn((int)((v_read      ) & 0x0f) - v_zero);
+        half v_1 = __int2half_rn((int)((v_read >>  4) & 0x0f) - v_zero);
+        half v_2 = __int2half_rn((int)((v_read >>  8) & 0x0f) - v_zero);
+        half v_3 = __int2half_rn((int)((v_read >> 12) & 0x0f) - v_zero);
+        half v_4 = __int2half_rn((int)((v_read >> 16) & 0x0f) - v_zero);
+        half v_5 = __int2half_rn((int)((v_read >> 20) & 0x0f) - v_zero);
+        half v_6 = __int2half_rn((int)((v_read >> 24) & 0x0f) - v_zero);
+        half v_7 = __int2half_rn((int)((v_read >> 28)       ) - v_zero);
+
+        half2 v_01 = __halves2half2(v_0, v_1);
+        half2 v_23 = __halves2half2(v_2, v_3);
+        half2 v_45 = __halves2half2(v_4, v_5);
+        half2 v_67 = __halves2half2(v_6, v_7);
+
+        half h_0 = h_ptr[*x_map_ptr++];
+        half h_1 = h_ptr[*x_map_ptr++];
+        half h_2 = h_ptr[*x_map_ptr++];
+        half h_3 = h_ptr[*x_map_ptr++];
+        half h_4 = h_ptr[*x_map_ptr++];
+        half h_5 = h_ptr[*x_map_ptr++];
+        half h_6 = h_ptr[*x_map_ptr++];
+        half h_7 = h_ptr[*x_map_ptr++];
+
+        half2 h_01 = __halves2half2(h_0, h_1);
+        half2 h_23 = __halves2half2(h_2, h_3);
+        half2 h_45 = __halves2half2(h_4, h_5);
+        half2 h_67 = __halves2half2(h_6, h_7);
+
+        half2 tmp = __hmul2(h_01, v_01);
+        tmp = __hfma2(h_23, v_23, tmp);
+        tmp = __hfma2(h_45, v_45, tmp);
+        tmp = __hfma2(h_67, v_67, tmp);
+        result = __hfma2(v_scale_2, tmp, result);
+    }
+
+    return result;
+}
+
+__device__ __forceinline__ half dot_product_8_x_map_h
+(
+    const half acc,
+    MatrixView_half& h_,
+    const int h_row,
+    const int h_column,                 // divisible by 8
+    MatrixView_q4_column& v_,
+    const int v_row,                    // divisible by 8
+    const int v_column,
+    const half v_scale,
+    const uint32_t v_zero,              // + 1 (!!)
+    const int count,
+    const uint32_t* x_map
+)
+{
+    const half* h_ptr = h_.item_ptr(h_row, 0);
+    const uint32_t* x_map_ptr = x_map + h_column;
+    const uint32_t* v_ptr = (const uint32_t*) v_.item_uint32_ptr(v_row, v_column);
+    half result = acc;
+
+    for (int i = 0; i < count; i++)
+    {
+        uint32_t v_read = *v_ptr; v_ptr += v_.width;
+
+        half v_0 = __int2half_rn((int)((v_read      ) & 0x0f) - v_zero);
+        half v_1 = __int2half_rn((int)((v_read >>  4) & 0x0f) - v_zero);
+        half v_2 = __int2half_rn((int)((v_read >>  8) & 0x0f) - v_zero);
+        half v_3 = __int2half_rn((int)((v_read >> 12) & 0x0f) - v_zero);
+        half v_4 = __int2half_rn((int)((v_read >> 16) & 0x0f) - v_zero);
+        half v_5 = __int2half_rn((int)((v_read >> 20) & 0x0f) - v_zero);
+        half v_6 = __int2half_rn((int)((v_read >> 24) & 0x0f) - v_zero);
+        half v_7 = __int2half_rn((int)((v_read >> 28)       ) - v_zero);
+
+        half tmp = __hmul(h_ptr[*x_map_ptr++], v_0);
+        tmp = __hfma(h_ptr[*x_map_ptr++], v_1, tmp);
+        tmp = __hfma(h_ptr[*x_map_ptr++], v_2, tmp);
+        tmp = __hfma(h_ptr[*x_map_ptr++], v_3, tmp);
+        tmp = __hfma(h_ptr[*x_map_ptr++], v_4, tmp);
+        tmp = __hfma(h_ptr[*x_map_ptr++], v_5, tmp);
+        tmp = __hfma(h_ptr[*x_map_ptr++], v_6, tmp);
+        tmp = __hfma(h_ptr[*x_map_ptr++], v_7, tmp);
+        result = __hfma(v_scale, tmp, result);
+    }
+
+    return result;
+}
+
 #endif

autogptq_extension/qigen/generate.py

@@ -1162,7 +1162,7 @@ def unpack_zeros(bits):
     res += f"void unpack_zeros{bits}_cpu(const int* zv, float* ov, int n, int m)"
     packed = 32//bits
     mask = (2**bits)-1
-    res += "{\n"
+    res += "{\nconst __m256i ones = _mm256_set1_epi32(1);\n"
     res += f"const __m256i mask = _mm256_set1_epi32({mask});\n"
     if bits == 4:
         res += "const __m256i shift = _mm256_set_epi32(28,24,20,16,12,8,4,0);\n"
@@ -1179,14 +1179,15 @@ def unpack_zeros(bits):
         res += "__m256i z = _mm256_set1_epi32(zv[i*m/8 + j/8]);\n"
         res += "__m256i z0 = _mm256_srlv_epi32(z, shift);\n"
         res += "__m256i z1 = _mm256_and_si256(z0, mask);\n"
-        res += "__m256 z2 = _mm256_cvtepi32_ps(z1);\n"
+        res += "__m256i z2 = _mm256_add_epi32(z1, ones);\n"
-        res += "_mm256_storeu_ps(&ov[i*m +j], z2);\n"
+        res += "__m256 z3 = _mm256_cvtepi32_ps(z2);\n"
+        res += "_mm256_storeu_ps(&ov[i*m +j], z3);\n"
     elif bits == 2:
         res += f"for (int j = 0; j < m; j+={packed})"
         res += "{\n"
         res += f"for (int k = 0; k < {packed}; k++)"
         res += "{\n"
-        res += f"ov[i*m + j+k] = ((zv[j/{packed}] >> ({bits}*k)) & {mask});\n"
+        res += f"ov[i*m + j+k] = (((zv[j/{packed}] >> ({bits}*k)) & {mask})+1);\n"
         res += "}\n"
         # res += "for(int j = 0; j < m; j+=16){\n"
         # res += "__m256i z = _mm256_set1_epi32(zv[i*m/16 + j/16]);\n"