(Alpha) Auto-GPTQ 浅析

动机、参考资料、涉及内容

动机

• Auto-GPTQ 源码浅析
  • 怎么写一个 python 包 (setup.py)
  • Github Action 做分发
  • torch cpp extension 的使用
  • OpenAI triton 的使用
  • 怎么基于 🤗 做二次开发的例子【待定】

参考资料

• 原始代码仓库: https://github.com/PanQiWei/AutoGPTQ.git

涉及内容

• 本博客只分析 commit id: 046c031 的代码

不涉及内容

• GPTQ 算法 (float16 模型量化为整数权重的过程)
• torch cpp extension 及 OpenAI triton 的深入介绍

项目目录结构

README.md
setup.py
docs/
examples/
.github/                            # 与 GitHub 与发布相关的
auto_gptq/                          # 安装的 python 包
    __init__.py
    eval_tasks/                     # 评估量化结果
    modeling/
        __init__.py
        _base.py                    # 核心类: BaseQuantizeConfig, BaseGPTQForCausalLM
        auto.py                     # 核心类: AutoGPTQForCausalLM
        llama.py                    # bloom.py, gptq2.py, ..., 继承自 BaseGPTQForCausalLM, 但只修改几个类属性
        ...
    nn_modules/                     # 待研究
        __init__.py
        _fused_base.py
        fused_gptj_attn.py
        fused_llama_attn.py
        fused_llama_mlp.py
        qlinear/
            __init__.py
            qlinear_cuda.py
            qlinear_cuda_old.py
            qlinear_triton.py
        triton_utils/
            __init__.py
            custom_autotune.py
            kernels.py
            mixin.py
    quantization/
        __init__.py
        gptq.py
        quantizer.py
    utils/
        ...
        peft_utils.py               # 🤗 Peft 的集成
autogptq_cuda/                      # Pytorch CUDAExtension 
    autogptq_cuda_256.cpp
    autogptq_cuda_64.cpp
    autogptq_cuda_kernel_256.cu
    autogptq_cuda_kernel_64.cu

安装与 setup.py

根据官方文档的描述, 使用 pip install 的方式进行安装有如下几种选项

# 确认之前的安装被删除
pip uninstall autogptq_cuda -y

# 安装时使用 pytorch 编译 extension
pip install autogptq  # 等价与 BUILD_CUDA_EXT=1 pip install auto-gptq

# 安装时不编译 extension
BUILD_CUDA_EXT=0 pip install auto-gptq

# 额外安装 triton
pip install auto-gptq[triton]

上面的注释实际上有些“含糊”, 因此这里直接对 setup.py 进行分析, 以得到准确理解

# 关键部分代码
common_setup_kwargs = {
    "python_requires": f">=3.8.0",                            # 强制检查, 否则安装报错
    # 以下只有在 BUILD_CUDA_EXT="1", 且 torch 存在时才有【待确认什么叫torch存在】
    "ext_modules": [
        cpp_extension.CUDAExtension(
            "autogptq_cuda_64",
            [
                "autogptq_cuda/autogptq_cuda_64.cpp",
                "autogptq_cuda/autogptq_cuda_kernel_64.cu"
            ]
        ),
        cpp_extension.CUDAExtension(
            "autogptq_cuda_256",
            [
                "autogptq_cuda/autogptq_cuda_256.cpp",
                "autogptq_cuda/autogptq_cuda_kernel_256.cu"
            ]
        )
    ],
    "cmdclass": {'build_ext': cpp_extension.BuildExtension}
}

include_dirs = [
    "autogptq_cuda"
    ""  # 【需要使用Pytorch extension时还会增加】
]

setup(
    packages=find_packages(),       # ["auto_gptq"]
    install_requires=requirements,  # 自动安装: ["accelerate>=0.19.0", "torch>=1.13.0", "transformers>=4.29.0", "peft", ...]
    extras_require=extras_require,  # 自动安装, 使用 pip install auto-gptq[triton] 才会触发 {"triton": ["triton>=2.0.0"]}
    include_dirs=include_dirs,      # 也放入 site-packages 目录内
    **common_setup_kwargs
)

关于 extras_require 与 pip install xx[yy,zz]: 【待确认并移入Notes中】

• 先判断 "yy" 是否是 extras_require 的键, 如果匹配上, 就安装 extras_require["yy"] 的包, 匹配不上就报一个警告, 然后进行下一步
• 然后判断 yy 是否为一个包名(PyPI), 如果匹配上, 就安装 yy 包
• 都匹配不上就不进行安装 yy

源码安装

pip install .[triton]

如果 Pytorch CUDA extension 被正确安装, 安装位置为

/path/to/site-packages/autogptq_cuda_256.cpython-38-x86_64-linux-gnu.so
/path/to/site-packages/autogptq_cuda_64.cpython-38-x86_64-linux-gnu.so

提示: 如果希望看出一些安装过程具体干了什么, 可以尝试使用 pip install -v .

通常不推荐使用: python setup.py install, 参考 stackoverflow

pip install -v -e .
# pip install . 相当于在 python setup.py install 之外还做了自动安装依赖包等工作
# -e 参数方便调试, 项目目录下的修改会自动生效
# -v 用于显示安装过程, 例如让 setup.py 文件中的 print 能正常生效

日志如下

# 以 # 号开头的是注释

Processing /abspath/to/AutoGPTQ
  Running command python setup.py egg_info
  ...  # 这一部分会打印 setup.py 中的 print 语句输出
  running egg_info
  # 创建一些临时目录, 并写文件
  writing /tmp/pip-pip-egg-info-eoiqxgek/auto_gptq.egg-info/PKG-INFO
  ...
  Prepareing metadata (setup.py) ... done
# 安装依赖包
Collecting accelerate>=0.19.0
...
Building wheels for collected packages: auto-gptq
  Running command python setup.py bdist_wheel
  running bdist_wheel
  running build
  running build_py
  # 拷贝文件到 build 目录
  create build/lib/auto_gptq
  copy auto_gptq/__init__.py -> build/lib/auto_gptq
  ...
  # 如果BUILD_CUDA_EXT=1(默认值), 且安装auto-gptq之前就已经安装了torch时会触发
  running build_ext
  ...
  running install
  running install_lib
  # 再拷贝一次
  copy build/lib/auto_gptq/__init__py -> build/bdist.linux-x86_64/wheel/auto_gptq
  ...
  running install_egg_info
  running egg_info
  # 疑似是写到 site-packages/auto_gptq.egg-info 中
  writing auto_gptq.egg-info/PKG-INFO
  ...
  running install_scripts
  creating build/bdist.linux-x86_64/wheel/auto_gptq-0.3.0.dev0.dist-info/WHEEL
  creating '/tmp/pip-wheel-nlkqzuc6/auto_gptq-0.3.0.dev0-py3-none-any.whl' and adding 'build/bdist.linux-x86_64/wheel' to it
  adding 'auto_gptq/__init__.py'
  ...
  removing build/bdist.linux-x86_64/wheel
  Building wheel for auto-gptq (setup.py) ... done
  Created wheel for auto-gptq: filename=auto_gptq-0.3.0.dev0-py3-none-any.whl size=63666 sha256=12345163663
  Store in directory: /tmp/pip-ephem-wheel-cache-5yjlc9ij/wheels/36/9a/b2/12355772277
Successfully build auto-gptq
Installing collected packages: accelerate, auto-gptq
  # 上一行运行之后可能会卡住一段时间, 应该是把前面的临时文件夹进行拷贝
  # /tmp/pip-wheel-nlkqzuc6, /tmp/pip-ephem-wheel-cache-5yjlc9ij
  changing model of /path/to/python/bin/accelerate to 755
  changing model of /path/to/python/bin/accelerate-config to 755
  changing model of /path/to/python/bin/accelerate-launch to 755
Successfully installed accelerate-0.19.0 auto-gptq-0.3.0-dev0

备注: /tmp 目录的使用实际上还有更多, 但最终会被清理掉

pip-build-tracker-xxxx
pip-ephem-wheel-cahce-xxxx
pip-install-xxxx
pip-pip-egg-info-xxxx
pip-unpack-xxxx
pip-wheel-xxxx

CUDA extension 安装可能报错的问题

Case 1

error identifier "__hfma2" is undefined

出现此问题的环境是: 驱动为最高支持 CUDA 11.5 版本, CUDA toolkit 版本号为 11.5, pytorch 版本号为 2.0.1 (安装的是 cuda11.7 的预编译包, 自动安装了一堆 Nvidia 的 CUDA 11.7 python 包)

疑似是显卡驱动/CUDA 版本的问题, __hfma2 定义在 CUDA_MATH_API 中, 这个函数至少在 CUDA 11.2 版本就已经支持, 故出现此问题的原因存疑.

Case 2

identifier "AT_DISPATCH_CASE" is undefined

出现此问题的环境是驱动: 为最高支持 CUDA 11.5 版本, CUDA toolkit 版本号为 11.5, pytorch 版本号为 1.11.0 (安装的是 cuda11.5 的预编译包), 并手动修改了 AutoGPTQ 的 setup.py 中 pytorch 的最低版本要求, 并注释掉了 auto_gptq/__init__.py 中 PEFT 的 import 语句.

这个问题应该是 Pytorch 版本太低的问题, AT_DISPATCH_CASE 这个宏定义在 pytorch ATEN 中: https://github.com/pytorch/pytorch/blob/main/aten/src/ATen/Dispatch.h, 而 Pytorch

一些探索【TODO: 后续整合到CUDA笔记中, 并删除此处的记录】 nvcc 怎么确定类似这种参数 -gencode=arch=compute_70,code=sm_70, 参考:

• 博客
• issue
• 问答

TORCH_CUDA_ARCH_LIST=7.0
/usr/local/cuda-11.7/bin/nvcc --list-gpu-arch

基本使用

基本的使用方式可以参考官方 示例代码 或者 README, 精简后如下:

import os

from transformers import AutoTokenizer, TextGenerationPipeline
from auto_gptq import AutoGPTQForCausalLM, BaseQuantizeConfig
import numpy as np
import torch
import torch.nn as nn

pretrained_model_dir = "/path/to/llama_7b"
quantized_model_dir = "llama_7b-4bit-128g"

def get_wikitext2(nsamples, seed, seqlen, model):
    from datasets import load_dataset
    traindata = load_dataset('wikitext', 'wikitext-2-raw-v1', split='train')
    testdata = load_dataset('wikitext', 'wikitext-2-raw-v1', split='test')

    from transformers import AutoTokenizer
    try:
        tokenizer = AutoTokenizer.from_pretrained(model, use_fast=False)
    except:
        tokenizer = AutoTokenizer.from_pretrained(model, use_fast=True)
    trainenc = tokenizer("\n\n".join(traindata['text']), return_tensors='pt')
    testenc = tokenizer("\n\n".join(testdata['text']), return_tensors='pt')

    import random
    random.seed(seed)
    np.random.seed(0)
    torch.random.manual_seed(0)
    
    traindataset = []
    for _ in range(nsamples):
        i = random.randint(0, trainenc.input_ids.shape[1] - seqlen - 1)
        j = i + seqlen
        inp = trainenc.input_ids[:, i:j]
        attention_mask = torch.ones_like(inp)
        traindataset.append({'input_ids':inp,'attention_mask': attention_mask})
    return traindataset, testenc

tokenizer = AutoTokenizer.from_pretrained(pretrained_model_dir, use_fast=True)
traindataset,testenc = get_wikitext2(128, 0, 2048, pretrained_model_dir)
quantize_config = BaseQuantizeConfig(
    bits=4,  # quantize model to 4-bit
    group_size=128,  # it is recommended to set the value to 128
    desc_act=False,  # desc_act and group size only works on triton
)
model = AutoGPTQForCausalLM.from_pretrained(pretrained_model_dir, quantize_config)
model.quantize(traindataset, use_triton=False)
model.save_quantized(quantized_model_dir)

model = AutoGPTQForCausalLM.from_quantized(quantized_model_dir, device="cuda:0")
print(tokenizer.decode(model.generate(**tokenizer("auto_gptq is", return_tensors="pt").to(model.device))[0]))

如果需要将 AutoGPTQ 用于更多的模型, 需要理解其核心类 auto_gptq.modeling._base.BaseGPTQForCausalLM

class BaseGPTQForCausalLM(nn.Module, PushToHubMixin):
    layer_type: str = None
    layers_block_name: str = None
    outside_layer_modules: List[str] = None
    inside_layer_modules: List[List[str]] = None
    lm_head_name: str = "lm_head"

    fused_attn_module_type: Optional[FusedBaseAttentionModule] = None
    fused_mlp_module_type: Optional[FusedBaseMLPModule] = None

    def quantize(...):
        ...

    def from_quantized(...):
        ...
    
    def save_quantized(...):
        ...
    
    def from_pretrained(...):
        ...
    
    def save_pretrained(...):
        ...

而每个具体的模型一般只需要定义类属性即可, 例如: auto_gptq.modeling.llama.LlamaGPTQForCausalLM

# LlamaGPTQForCausalLM 类的完整代码
class LlamaGPTQForCausalLM(BaseGPTQForCausalLM):
    layer_type = "LlamaDecoderLayer"
    layers_block_name = "model.layers"
    outside_layer_modules = ["model.embed_tokens", "model.norm"]
    inside_layer_modules = [
        ["self_attn.k_proj", "self_attn.v_proj", "self_attn.q_proj"],
        ["self_attn.o_proj"],
        ["mlp.up_proj", "mlp.gate_proj"],
        ["mlp.down_proj"]
    ]
    # 以下两个是非必须的
    fused_attn_module_type = FusedLlamaAttentionForQuantizedModel
    fused_mlp_module_type = FusedLlamaMLPForQuantizedModel

以 llama_7b 为例, 模型结构可以用这个方式进行打印:

from transformers.models.llama import LlamaForCausalLM, LlamaConfig
from accelerate import init_empty_weights
name = "./hf_download/llama_7B"
config = LlamaConfig.from_pretrained(name)
with init_empty_weights():
    model = LlamaForCausalLM(config)
print(model)

LlamaForCausalLM(
  (model): LlamaModel(
    (embed_tokens): Embedding(32000, 4096, padding_idx=0)
    (layers): ModuleList(
      (0-31): 32 x LlamaDecoderLayer(
        (self_attn): LlamaAttention(
          (q_proj): Linear(in_features=4096, out_features=4096, bias=False)
          (k_proj): Linear(in_features=4096, out_features=4096, bias=False)
          (v_proj): Linear(in_features=4096, out_features=4096, bias=False)
          (o_proj): Linear(in_features=4096, out_features=4096, bias=False)
          (rotary_emb): LlamaRotaryEmbedding()
        )
        (mlp): LlamaMLP(
          (gate_proj): Linear(in_features=4096, out_features=11008, bias=False)
          (down_proj): Linear(in_features=11008, out_features=4096, bias=False)
          (up_proj): Linear(in_features=4096, out_features=11008, bias=False)
          (act_fn): SiLUActivation()
        )
        (input_layernorm): LlamaRMSNorm()
        (post_attention_layernorm): LlamaRMSNorm()
      )
    )
    (norm): LlamaRMSNorm()
  )
  (lm_head): Linear(in_features=4096, out_features=32000, bias=False)
)

首先对类属性做简单介绍:

layer_type

用途主要是传递给 accelerate.utils.get_balanced_memory 的 no_split_module_classes 参数用

layers_block_name

这个是用来指定需要量化的层, 例如在 llama-7b 的例子里, 需要量化的参数都在 model.layers 中

outside_layer_modules

这个参数是用于指定不需要量化的层, 例如 llama-7b 的例子里, 取值为 ["model.embed_tokens", "model.norm"], 指定这个主要目的是 AutoGPTQ 在进行模型量化的过程里, 需要用校准数据获取 layers_block_name 里第一层的输入, 具体算法逻辑如下:

• 首先加载整个模型原始浮点数的权重至 CPU, 这一过程使用 BaseGPTQForCausalLM.from_pretrained 进行实现
• 然后将 outside_layer_modules 中的所有权重转移到 GPU 上, 然后将 layers_block_name 里第一层替换为 LayerHijacker, 而被替换的这个层只会记录输入参数, 然后触发 ValueError, 这样一来第一层的输入就被记录了下来
• 然后一层一层做量化

inside_layer_modules

这个参数用于指定 layers_block_name 中每一层需要量化的子模块, inside_layer_modules 是一个列表的列表, 列表中越靠前的组越先被优化, 这个顺序与模型在前向计算时的计算次序一致.

lm_head_name

这个参数对应模型的 head 部分, 但感觉这个参数在设计上可能可以被优化掉.

fused_attn_module_type

这个参数是可选的, 主要是在调用 from_quantized 时, 对一些层进行替换, 起到 kernel fused 的效果. 目前主要是使用 flashattention.

fused_mlp_module_type

这个参数是可选的, 类似 fused_attn_module_type

然后介绍 BaseGPTQForCausalLM 的几个主要方法【TODO待补充】

量化参数的保存格式

对于一个 nn.Linear 层, 量化前后的参数如下(注意与原始 gptq 实现上有一定的区别):

# 量化前
weight.dtype, weight.shape = float16, (out_features, in_features)
bias.dtype, weights.shape = float16, (out_features,)
# 量化后
qweight.dtype, qweight.shape = int32, (in_feature*bits/32, out_features)
qzeros.dtype, qzeros.shape = int32, (in_features/group_size, out_feature*bits/32)  # [0, 2**bits-2]. 保存时很奇怪地对量化过程中的 zeros 都进行了减 1 操作
scales.dtype, scales.shape = float16, (in_features/group_size, out_feature)
g_idx.dtype, g_idx.shape = int32, (in_features,)
bias.dtype, bias.shape = float16, (out_features,)

# 以 llama_7b 的 layers.0.mlp.down_proj 为例:
group_size, bits = 128, 4  # 量化至 4 位, 以 128 分组执行量化
# 量化前
weight.dtype, weight.shape = float16, (4096, 11008)
bias.dtype, weights.shape = float16, (4096,)
# 量化后
qweight.dtype, qweight.shape = int32, (1376, 4096)
qzeros.dtype, qzeros.shape = int32, (86, 512)
scales.dtype, scales.shape = float16, (86, 4096)
g_idx.dtype, g_idx.shape = int32, (11008,)
bias.dtype, bias.shape = float16, (4096,)

这里对 GPTQ 算法的量化函数稍加回顾: GPTQ 算法仅对 weight 进行量化, 量化的粒度为 weight 每一行的 group_size 个参数为一组, 也就是每 group_size 个参数确定一个只有 2**bits 的网格, 将浮点参数量化到这 2**bits 个浮点数上去, 因此会有 (in_features/group_size)*out_feature 个缩放系数 scale 和浮点数 0.0 量化的值 nbit_zeros. 由于 nbit_zeros 的取值位于 [0, 2**bits), 所以每个数字只需要 bits 位就可以存储, 但一般的硬件可能不支持 bits=2,3,4 的整数存储类型, 因此可以将多个整数打包在一起, 用 32 位整数来表示. 同样的道理, weight 被量化后的存储形式也会做类似的打包. 量化后的参数与量化前参数的转换规则如下:

\[w_{dequant-float} = (w_{int} - zero_{int}) * scale\]

这里具体对整数打包方式做个说明, qweight 和 qzeros 的具体排布方式进行如下说明:

bits, group_size = 4, 4
unpack_qweight.shape = (in_features, out_features) = (8, 8)
unpack_qweight = [
    [0, 1, 2, 0, 1, 2, 0, 1],
    [1, 2, 3, 1, 2, 3, 1, 2],
    [2, 3, 4, 2, 3, 4, 2, 3],
    [3, 4, 5, 3, 4, 5, 3, 4],
    [4, 5, 6, 4, 5, 6, 4, 5],
    [5, 6, 7, 5, 6, 7, 5, 6],
    [7, 8, 9, 7, 8, 9, 7, 8],
    [15,0,14,15, 0,14,15, 0]
]
qweight.shape = (in_features * 32 / bits, out_features) = (1, 8)
qweight = [
    [
        0 + 1*16 + 2*(16**2) + 3*(16**3) + 4*(16**4) + 5*(16**5) + 7*(16**6) + 15*(16**7),
        1 + 2*16 + 3*(16**2) + 4*(16**3) + 5*(16**4) + 6*(16**5) + 8*(16**6) + 0*(16**7),
        2 + 3*16 + 4*(16**2) + 5*(16**3) + 6*(16**4) + 7*(16**5) + 9*(16**6) + 14*(16**7),
        0 + 1*16 + 2*(16**2) + 3*(16**3) + 4*(16**4) + 5*(16**5) + 7*(16**6) + 15*(16**7),
        1 + 2*16 + 3*(16**2) + 4*(16**3) + 5*(16**4) + 6*(16**5) + 8*(16**6) + 0*(16**7),
        2 + 3*16 + 4*(16**2) + 5*(16**3) + 6*(16**4) + 7*(16**5) + 9*(16**6) + 14*(16**7),
        0 + 1*16 + 2*(16**2) + 3*(16**3) + 4*(16**4) + 5*(16**5) + 7*(16**6) + 15*(16**7),
        1 + 2*16 + 3*(16**2) + 4*(16**3) + 5*(16**4) + 6*(16**5) + 8*(16**6) + 0*(16**7),
    ]
]

unpack_qzeros.shape = (in_features/group_size, out_features) = (2, 8)
# unpack_zeros[0][0] = 1 表示的是在这个group里, 整数1代表浮点数 0.0
unpack_qzeros = [
    [1, 2, 3, 4, 15, 2, 3, 3],  # 注意按照GPTQ的量化函数, 0.0 对应的整数值总是位于 [1, 15]
    [2, 3, 4, 5, 4, 15, 1, 2]
]

qzeros.shape = (in_features/group_size, out_features * 32 / bits) = (2, 1)
qzeros = [
    [(1-1) + (2-1)*16 + (3-1)*(16**2) + (4-1)*(16**3) + (15-1)*(16**4) + (2-1)*(16**5) + (3-1)*(16**6) + (3-1)*(16**7)],
    [(2-1) + (3-1)*16 + (4-1)*(16**2) + (5-1)*(16**3) + (4-1)*(16**4) + (15-1)*(16**5) + (1-1)*(16**6) + (2-1)*(16**7)],
]

了解这一点后, 实际上可以从 AutoGPTQ 保存的量化后的权重参数文件 gptq_model-4bit-128g.bin 重新转为 float16 版本的权重文件(原始权重经过 AutoGPTQ 转换为量化后的权重相当于一次有损压缩, 通过这种方式恢复回来相当于是解压). 这里是一个简单的实现方法(脚本本身不具备通用性, 但可以用来做验证)【待补充】

量化模型的推理

AutoGPTQ 中包含了几种 kernel 实现 (commit id: 046c031)

这里的分发逻辑相对复杂, 因此在具体介绍每一种实现前需要理清分发逻辑

• triton
  • triton: 2,4,8 bit 量化
• cuda_old
  • cuda_faster_old: 2,3,4 bit 量化, 且入参的 batch_size (reshape 后) 小于阈值(默认为128) 且 use_cuda_fp16=True(默认值)
  • cuda_old: 2,3,4,8 bit 量化, 且入参的 batch_size (reshape 后) 小于阈值(默认为128) 且 use_cuda_fp16=False
  • torch 实现(兜底): 2,3,4,8 bit 量化, 且入参的 batch_size (reshape 后) 大于等于阈值(默认为128)
• cuda
  • cuda: 2,3,4,8 bit 量化, 且入参的 batch_size (reshape 后) 小于阈值(默认为128)
  • torch 实现(兜底): 2,3,4,8 bit 量化, 且入参 batch_size (reshape 后) 大于等于阈值(默认为128)

备注: cuda_old 和 cuda 的情形下, 如果入参维数与出参维数不能被 64 整除, 也会直接调用 torch 实现兜底

调用 from_quantized 方式时可以传入参数 use_triton=True 来使用 triton 的 kernel, 其余情况下, 如果 group_size!=-1 且 desc_act=True 时, 使用 cuda, 否则使用 cuda_old.

为了清晰起见, 这里列出分发到最底层的“kernel”:

• cuda 实现可直接参考 autogptq_cuda/autogptq_cuda_64.cpp 和 autogptq_cuda/autogptq_cuda_256.cpp 文件末尾的 pybind11 相关代码, 注意前者适用于 in_features 与 out_features 能被 64 整除时的情况, 后者适用于能被 256 整除的情况 (优先使用), 如果不满足条件, 则使用 torch 进行计算
• triton 实现可参考 auto_gptq/nn_modules/triton_utils/kernels.py 中的 quant_matmul_248_kernel 与 transpose_quant_matmul_248_kernel, 其中前者时前向计算的 kernel, 后者是反向求导的 kernel.
• torch 实现可参考 auto_gptq/nn_modules/qlinear_cuda_old.py 中 QuantLinear 的 forward 函数的实现

而最常见的情况是 4 bit 量化, 推理时是 batch_size=1, 输入输出形状一般能被 256 整除, 会触发 cuda_faster_old 或者 triton

triton 算子的实现【待定，本质上很朴素】

这部分实现与 triton 官方教程中的矩阵乘法类似, 具体逻辑是先将整数类型的权重重新转换回 float16, 然后与输入做乘积. 注意累加项采用的是 float32 类型 (应该是为了精度), 累加完毕后再将结果转换回 float16 作为输出.

AutoGPTQ 中对 triton.autotune 做了一些 hack【待补充】

CUDA 算子的实现【待定，本质上很朴素】

cuda_faster_old 与 cuda_old 的实现

.github 目录

相关链接

• 2023 年 6 月下旬至 7 月上旬, AutoGPTQ 项目似乎陷入了停滞 issue, 与此同时, huggingface 官方似乎有计划将 GPTQ 引入, 但 huggingface 的开发人员似乎计划使用 exllama 项目提供的算子 (exllama 的推理算子似乎更为高效). 截至至本文记录时间 2023/07/07, 许多其他项目对 GPTQ 的使用 (量化与推理) 采用的是 GPTQ-for-LLaMa. 总之, 可以期待一波 huggingface 对 GPTQ 的官方支持, 个人预计会出现在 huggingface optimum 仓库中. 注意 exllama 仅包含推理部分, 适用于 GPTQ-for-LLaMa 导出的量化后的模型 (应该也同样适用于 AutoGPTQ, 待验证).