为了改一行代码，我花了10多天时间，让性能提升了40多倍---Pascal架构GPU在 ...

ChatGPT生成的文章摘要

这篇博客记录了作者在家中使用Pascal显卡运行大型模型时遇到的挑衅和解决方案。随着本地大型模型性能的提升，作者选择使用vllm库举行推理。然而，作者遇到了多个技术困难，必要自行编译vllm和PyTorch，以支持Pascal架构的显卡。编译过程中，作者深入研究了显卡不支持的问题，特殊是在量化矩阵乘法计算中发现性能瓶颈。终极，解决了性能问题，让性能提升了43倍。这次技术探索不仅解决了具体问题，还为作者提供了深入学习和扩展其他相干技术的机会，同时也展示了LLM在整个过程中提供的帮助。文章结尾，作者总结了经验并提出了进一步研究的方向。
背景

家里有张Pascal架构的显卡【划重点，背面要考】，近来发现本地大模型的性能在蹭蹭往上涨，于是开始研究下是否能在本地跑大模型。
之前我就了解vllm库，vllm的推理速度照旧很快的，并且我之前还给vllm提交过一个PR，对vllm比较认识，所以我选择了使用vllm来举行推理。
选择结束之后就开始了漫长的抗争之路，期间着实遇到了很多问题，也学到了很多知识，故写此文以记录。
第一关：下载安装

当时无知的我以为安装是一件很简单的事情，从前使用vllm，直接pip install vllm，不仅会帮忙安装好vllm，pytorch，还会帮忙下载对应的cuda库，本身啥都不消操心。
这次的安装也如从前一样顺利，
安装完后就是选择模型了，选择模型的话，对于消费级显卡来说，显存占用是一个重要的考量因素，你得先跑起来。获取模型的显存占用的方式有两种：

• 计算模型必要占用的显存巨细，好比一个7B的模型，它的参数目是7,000M个，一个float16的参数占2个字节，所以必要7,000M *2B=14GB的显存，除了参数外，还要考虑存储KV缓存，以及样本在中间传输时的值，量化元信息(假如涉及量化的话)，所以必要留一些buffer。
  
• 别的一个获取显存占用的方式是直接用这个工具[1]，输入模型在huggingface上的名称，然后选择精度，就可以看到模型占用的显存巨细了。
  
  
  
  必要留意的是，这里同样必要预留buffer，这上面的显存巨细是纯模型本身的巨细，量化的模型尤其要留意，必要考虑量化元数据带来的显存占用。
  

这样看下来，我这张12G显存的显卡，顶多只能跑一个7B-int8的模型，为了能跑稍微大一点的上下文，我终极选择了Qwen/Qwen2.5-7B-Instruct-GPTQ-Int4的模型（颠末项目标实测，Qwen模型现在在中文开源领域确实很不错）。
兴奋地下载完的模型后，噩梦在启动vllm server的时间开始了。
迎面而来的是第一个错误是：
RuntimeError: Error in model execution (input dumped to /tmp/err_execute_model_input_20241211-200011.pkl): CUDA error: no kernel image is available for execution on the device
这个问题去stackoverflow[2]了一下，大概率是vllm编译的时间没有支持对应的显卡架构，还记得重点么，没错，大概率就是不支持Pascal架构，我去官方文档[3]上看了一下，确实没有发现Pascal的显卡支持，支持矩阵长这样，没有Pascal架构呀：

没办法了，那就尝试本身编译vllm，看看能不能解决这个问题。
第二关：vllm编译

编译vllm

一开始编译的时间感觉还挺简单的，直接照着vllm的文档来，文档就只有一行命令pip install -e .，事情肯定没有这么简单，编译堕落了：

1. CMake Error at CMakeLists.txt:252 (cuda_archs_loose_intersection):
2.         cuda_archs_loose_intersection Function invoked with incorrect arguments for
3.         function named: cuda_archs_loose_intersection

复制代码

252行是这么写的：

1. cuda_archs_loose_intersection(MARLIN_ARCHS "8.0;8.6;8.7;8.9;9.0" ${CUDA_ARCHS})

复制代码

中间颠末了大量时间的定位，终极找到了问题所在，重要就是vllm设置了一个支持的显卡架构（实在它使用了算力来表示架构，算力和架构有对应关系[4]）：

1. set(CUDA_SUPPORTED_ARCHS "7.0;7.2;7.5;8.0;8.6;8.7;8.9;9.0")

复制代码

只支持到7.0算力，而Pascal架构是6.1算力，所以终极CUDA_ARCHS就为空，所以就报错了。
那简单呀，我直接给CUDA_SUPPORTED_ARCHS加上6.1就行了，然后重新编译…
这次编译很顺利，编译完成之后，我就继承兴奋地启动vllm了，不幸的是，又一次报了Cuda error: no kernel image is available for execution on the device错误。
于是我继承Google，找到了这么一个github的issue[5]，issue说这种环境是显卡不受支持了，必要本身编译（背面我本身测试了一下pytorch，实在我的pytorch是可以使用的，至于这里为什么报错，后续再研究研究吧），于是我就屁颠屁颠地去开始编译pytorch了。
编译pytorch

pytorch的编译就复杂很多了，不像vllm的编译命令，pytorch分了很多步。
先是要安装一堆前置工具：
- CuDNN
- cmake, ninja
- requirements.txt
- mkl-static, mkl-include
- magma-cuda121
- triton

1. 这些工具安装都还算顺利，要么照着说明安装，要么就是conda或pip安装，最后的triton就是一个make。

复制代码

这里有几个坑：

• pytorch要求先export CMAKE_PREFIX_PATH，并且给了个命令，检查一下执行完后的命令，有可能conda的路径没有找对，必要本身手动指定一下。
  
• cmake一开始会找不到cudnn，必要将cudnn-version.h（直接用find找一下本身安装的cudnn-version.h在哪）文件拷贝或link到cuda的include目录下。
  

编译完成！
下面重新编译一次vllm，由于我们必要使用本身编译的pytorch，所以必要执行一下python use_existing_torch.py，vllm会帮我们把pytorch从依赖里删除掉，然后执行pip install -r requirements-build.txt，安装一下依赖，最后执行pip install -e . --no-build-isolation，这样安装的时间，vllm就不会再去安装这部门依赖了。
中间假如出现version 'GLIBCXX_3.4.30' not found的错误，我是把我安装的gcc的libstdc++.so.6软链到conda的lib目录就行了。

1. strings /usr/lib/x86_64-linux-gnu/libstdc++.so.6 | grep GLIBCXX_3.4.30

复制代码

检查一下libstdc++.so.6是否包含GLIBCXX_3.4.30，假如包含，则软链到conda的lib目录下。

1. ln -s /usr/lib/x86_64-linux-gnu/libstdc++.so.6 ${CONDA_PREFIX_1}/lib/libstdc++.so.6

复制代码

编译完成！
再次满怀期待地启动vllm server，不出不测地又报错了，这次报错是没找到xformers，这个是由于vllm默认是不带留意力后端的，由于它也不知道你用什么留意力后端，所以必要本身安装一下。安装的时间发现它依赖了pytorch并且去下载了pytorch，那要不照旧本身编译一把吧。
xformers页面先容中支持Pascal架构，所以安装起来很丝滑，一行命令即可：

1. pip install -v -U git+https://github.com/facebookresearch/xformers.git@main#egg=xformers

复制代码

启动vllm server!

终于告一段落了，vllm server终于启动了，没有任何报错，我成功地看到了Loading model weights took 5.2035 GB。（这里可以印证我之前说的，量化的模型在考虑上量化元数据后，显存占用变大了很多，从计算得到的3.5GB，变成了5.2GB）
你以为故事到这就结束了？不不不，现在才是故事的开始。日记到了Loading model weights took 5.2035 GB就卡住了，我等了很久，发现它不停在卡在这。
第三关：定位性能问题的根因

初见端倪

出现这样的状况后，我是一点头绪都没有，只能像无头苍蝇一样，在vllm的Python代码里多打一些断点试试看了，在疯狂打了几十个断点之后，终于定位到卡哪了，vllm默认会先做一次profile run，来告诉你一些根本信息：

1. Memory profiling results: duration=11.82 seconds, total_gpu_memory=11.88GiB, initial_memory_usage=6.15GiB, peak_torch_memory=6.54GiB, memory_usage_post_profile=6.20GiB, non_torch_memory=1.05GiB, kv_cache_size=2.50GiB, gpu_memory_utilization=0.85.

复制代码

由于这里必要举行模型推理，所以卡住了，这时间我才意识到，看一下nvidia-smi看看显卡是否在工作实在就能知道它确实是在跑模型代码（虽然我一开始也有点意识到，却不停没往这个方面上想，毕竟再慢也不至于这么慢）。事实证明，卡住的时间，显卡确实在工作，所以问题很显着了，就是由于我的显卡推理速度“太慢”导致的。于是我就把max-model-len设置成了100，看看是否可以大概跑出结果来。等候了很长的时间后，服务真的启动了。
速度这么慢我是千万没有想到的，只能先换台呆板测一下看怎么样，用了一台A6000的呆板，发现人家一瞬间就启动了，那很显着了，问题就是只有我这边很慢。
初步定位问题

有了方向之后，那要做的事情就比较简单了，由于我本身编译了pytorch、xformers以及vllm，所以我必要一个个地排查。
先在pytorch官网上找到了跑benchmark[6]的文档，分别在A6000呆板、我的呆板上本身编译的pytorch以及直接用pip install的pytorch上跑了一下，发现pytorch基础的性能是不差的。
然后使用xformers的benchmark[7]，同样测试了一下，发现xformers的性能也是ok的。
那问题多半就出在vllm了，由于我不确定到底问题出在什么地方，以及我大概率确定基础库是没啥问题的，所以我筹划把整个模型推理的各个步调都记录一下执行时间，来看看具体是什么地方出问题了，按照28原则，问题大概率出在20%的地方。
接下来就是想办法记录时间了，我本身没有特殊好的思路，所以就请教了一下LLM，LLM给了我一个思路，可以使用pytorch的register_forward_pre_hook和register_forward_hook来记录时间。它给的代码很粗糙直接使用time库来记录时间，而且只能记录一层模型。所以我就“稍”作修改，改成了递归地访问每一层模型，并且用cuda的Event（固然这个也是从LLM那问出来的）来记录时间。
时间记录的代码写完了，接下来就是运行一下，看看问题出在哪了。下面是我运行后跑出来的结果，各位来找找看以为那里有问题？

1. model: 134811.72338464856 ms
2.   model.embed_tokens: 37.62428665161133 ms
3.   model.layers: 134773.90933799744 ms
4.     model.layers.0: 4777.431374847889 ms
5.       model.layers.0.input_layernorm: 1.620192050933838 ms
6.       model.layers.0.self_attn: 673.7694255411625 ms
7.         model.layers.0.self_attn.qkv_proj: 411.43023681640625 ms
8.         model.layers.0.self_attn.rotary_emb: 0.1632319986820221 ms
9.         model.layers.0.self_attn.attn: 4.729087829589844 ms
10.         model.layers.0.self_attn.o_proj: 257.4468688964844 ms
11.       model.layers.0.post_attention_layernorm: 0.1900160014629364 ms
12.       model.layers.0.mlp: 4101.85174125433 ms
13.         model.layers.0.mlp.gate_up_proj: 2740.14697265625 ms
14.         model.layers.0.mlp.act_fn: 0.8391680121421814 ms
15.         model.layers.0.mlp.down_proj: 1360.8656005859375 ms

复制代码

不得不说134s才跑完profile真的是离谱，然后确实就是28原则，问题就出在了4个地方，分别是：

• model.layers.0.self_attn.qkv_proj
  
• model.layers.0.self_attn.o_proj
  
• model.layers.0.mlp.gate_up_proj
  
• model.layers.0.mlp.down_proj
  

这几个地方耗时都显着不正常，人家attention的计算才花了4ms，怎么这些操作要花几百乃至上千ms。
作为对比，我去查察了一下A6000呆板上的结果：

1. model: 7459.573736906052 ms
2.   model.embed_tokens: 265.0838928222656 ms
3.   model.layers: 7192.4459400177 ms
4.     model.layers.0: 259.46213555336 ms
5.       model.layers.0.input_layernorm: 1.3496320247650146 ms
6.       model.layers.0.self_attn: 145.4847927093506 ms
7.         model.layers.0.self_attn.qkv_proj: 129.69778442382812 ms
8.         model.layers.0.self_attn.rotary_emb: 1.3486080169677734 ms
9.         model.layers.0.self_attn.attn: 3.180543899536133 ms
10.         model.layers.0.self_attn.o_proj: 11.257856369018555 ms
11.       model.layers.0.post_attention_layernorm: 2.0490241050720215 ms
12.       model.layers.0.mlp: 110.57868671417236 ms
13.         model.layers.0.mlp.gate_up_proj: 69.62483215332031 ms
14.         model.layers.0.mlp.act_fn: 4.104191780090332 ms
15.         model.layers.0.mlp.down_proj: 36.84966278076172 ms

复制代码

结果很显着了，确实就是刚刚那几个地方的问题，其他地方的耗时根本上都差不多，有些乃至有领先（这个感觉应该属于偏差）。
ok，知道问题了就去看看代码吧。
通过Python源码定位问题

颠末一番研究，终极我把问题锁定到了量化计算上面，由于所有出问题的点都执行了量化的矩阵乘法计算。从网上搜了一张Qwen的架构图[8]，我把耗时长的点都用红框标出来了。
![外链图片转存失败,源站可能有防盗链机制,发起将图片保存下来直接上传](https://img-home.csdnimg.cn/images/20230724024159.png?origin_url=https%3A%2F%2Fnextcloud.aboydfd.com%2

Fs%2FFys66gZRTckSG36%2Fpreview&pos_id=img-t4NYUYrb-1734152322353)
从中我们可以看到，这些地方都执行了没有量化的输入和量化后的weight之间的矩阵乘法计算。
vllm的代码里则对应了：

1. class ColumnParallelLinear(LinearBase):
2.     ...
3.     def forward(self, input_):
4.         bias = self.bias if not self.skip_bias_add else None
6.         # Matrix multiply.
7.         assert self.quant_method is not None
8.         output_parallel = self.quant_method.apply(self, input_, bias) ## 就是这行进行了量化矩阵乘法
9.         if self.gather_output:
10.             # All-gather across the partitions.
11.             output = tensor_model_parallel_all_gather(output_parallel)
12.         else:
13.             output = output_parallel
14.         output_bias = self.bias if self.skip_bias_add else None
15.         return output, output_bias

复制代码

和

1. class RowParallelLinear(LinearBase):
2.     ...
3.     def forward(self, input_):
4.         ...
5.         assert self.quant_method is not None
6.         bias_ = None if (self.tp_rank > 0 or self.skip_bias_add) else self.bias
7.         output_parallel = self.quant_method.apply(self,               ## 就是这行进行了量化矩阵乘法
8.                                                   input_parallel,
9.                                                   bias=bias_)
10.         ...
12.         return output, output_bias

复制代码

由于我使用的是GPTQ量化模型，所以继承跟进必要去找的quant_method是GPTQ相干的。
跟进到self.quant_method.apply:

1. class GPTQLinearMethod(LinearMethodBase):
2.     ...
3.     def apply(self,
4.               layer: torch.nn.Module,
5.               x: torch.Tensor,
6.               bias: Optional[torch.Tensor] = None) -> torch.Tensor:
7.         out_shape = x.shape[:-1] + (layer.qweight.shape[-1], )
8.         reshaped_x = x.reshape(-1, x.shape[-1])
10.         output = ops.gptq_gemm(reshaped_x, layer.qweight, layer.qzeros,
11.                                layer.scales, layer.g_idx,
12.                                layer.exllama_state == ExllamaState.READY,
13.                                self.quant_config.weight_bits)
14.         if bias is not None:
15.             output.add_(bias)
16.         return output.reshape(out_shape)

复制代码

这里很显着问题就是gptq_gemm的计算（GEMM表示General Matrix Multiplication，通用矩阵乘法），继承：

1. def gptq_gemm(a: torch.Tensor, b_q_weight: torch.Tensor,
2.               b_gptq_qzeros: torch.Tensor, b_gptq_scales: torch.Tensor,
3.               b_g_idx: torch.Tensor, use_exllama: bool,
4.               bit: int) -> torch.Tensor:
5.     return torch.ops._C.gptq_gemm(a, b_q_weight, b_gptq_qzeros, b_gptq_scales,
6.                                   b_g_idx, use_exllama, bit)

复制代码

哎，终极照旧得去看cuda代码么！！！
小插曲

这里想说一下GPTQ的名字，大家一看到可能会以为它和GPT有关系，实在不是的，它这算是蹭GPT的热度，GPTQ的全称是Post-Training Quantization for Generative pre-trained transformers，确实是硬蹭的。Post-Training Quantization，指的是训练后量化，所以它是一种在模型训练完之后，不再继承训练，单纯对权重和/或激活值举行量化的方法，而GPTQ是对PTQ的一种。
由于要去看cuda的源码，我对此没有很强的信心，我一没看过cuda源码，二不了解量化计算是什么样的，所以我就去告急补课了一下，在网上找了个量化计算的视频[9]来看，这个视频讲得很详细，对量化感爱好的同学可以去看一下。看完视频过后我还不外瘾，我想弄清晰GPTQ的量化数学原理（GPTQ有一套完善的数学推理），只看了它的前身OBS、OBC、OBQ，在看GPTQ本身的时间，想到，我已经了解得足够多了，再看下去有点浪费时间了，照旧回归主线先把。
感爱好的同学可以参考下面2个链接，OBC/OBQ的论文本身写得也挺友好的，也可以看看：

• https://readpaper.feishu.cn/docx/OPP2dTuXAoaO0oxWhQAcC05Wnpc
  
• https://zhuanlan.zhihu.com/p/646210009
  
• https://arxiv.org/abs/2208.11580
  

通过cuda源码定位问题

接下来就是跟踪cuda源码了，通过搜索gptq_gemm找到对应的cuda源码：

1. torch::Tensor gptq_gemm(torch::Tensor a, torch::Tensor b_q_weight,
2.                         torch::Tensor b_gptq_qzeros,
3.                         torch::Tensor b_gptq_scales, torch::Tensor b_g_idx,
4.                         bool use_exllama, int64_t bit) {
5.   const at::cuda::OptionalCUDAGuard device_guard(device_of(a));
6.   auto options = torch::TensorOptions().dtype(a.dtype()).device(a.device());
7.   at::Tensor c = torch::empty({a.size(0), b_q_weight.size(1)}, options);
8.   at::Tensor temp_dq = torch::empty(
9.       {b_q_weight.size(0) * 32 / bit, b_q_weight.size(1)}, options);
11.   vllm::gptq::gemm_half_q_half_cuda(
12.       at::cuda::getCurrentCUDABlasHandle(), (const half*)a.data_ptr(),
13.       (const uint32_t*)b_q_weight.data_ptr(),
14.       (const uint32_t*)b_gptq_qzeros.data_ptr(),
15.       (const half*)b_gptq_scales.data_ptr(),
16.       b_g_idx.device().is_meta() ? NULL : (const int*)b_g_idx.data_ptr(),
17.       (half*)c.data_ptr(), (half*)temp_dq.data_ptr(),
18.       c.size(0),              // m
19.       c.size(1),              // n
20.       a.size(1),              // k
21.       b_gptq_qzeros.size(0),  // group number
22.       use_exllama, bit);
23.   return c;
24. }

复制代码

重要就是gemm_half_q_half_cuda这个函数，这个函数是GPTQ的量化矩阵乘法计算，a是输入，b_q_weight是量化后的权重，b_gptq_qzeros是公式里的Z，b_gptq_scales是公式里的S，然后use_exllama是是否使用exllama库。
由于use_exllama后续会影响到分支逻辑，所以先检查一下use_exllama是否为true。从这里的代码不停往上翻查，可以看到use_exllama是从config中读取的，qwen2.5的config中设置的是true。
继承跟进代码：

2. void gemm_half_q_half_cuda(cublasHandle_t cublas_handle, const half* a,
3.                            const uint32_t* b_q_weight,
4.                            const uint32_t* b_gptq_qzeros,
5.                            const half* b_gptq_scales, const int* b_g_idx,
6.                            half* c, half* temp_dq, int size_m, int size_n,
7.                            int size_k, int groups, bool use_exllama, int bit) {
8.   bool use_reconstruct;
9.   if (use_exllama) {
10.     use_reconstruct = ((bit == 8 && size_m > MAX_Q_GEMM_ROWS_8BIT) ||
11.                        (bit != 8 && size_m > MAX_Q_GEMM_ROWS));
12.   } else {
13.     // The 2/3-bit kernels are somehow slower than dequant + gemm baseline, so
14.     // we disabled them for now.
15.     use_reconstruct = (bit < 4 || size_m > MAX_ALT_GEMM_ROWS);
16.   }
17.   if (use_reconstruct) {
18.     // Reconstruct FP16 matrix, then cuBLAS
19.     if (use_exllama) {
20.       reconstruct_exllama(b_q_weight, b_gptq_qzeros, b_gptq_scales, b_g_idx,
21.                           temp_dq, size_k, size_n, groups, bit);
22.     } else {
23.       reconstruct_gptq(b_q_weight, b_gptq_qzeros, b_gptq_scales, b_g_idx,
24.                        temp_dq, size_k, size_n, groups, bit);
25.     }
27.     const half alpha = __float2half(1.0f);
28.     const half beta = __float2half(0.0f);
29.     cublasHgemm(cublas_handle, CUBLAS_OP_N, CUBLAS_OP_N, size_n, size_m, size_k,
30.                 &alpha, temp_dq, size_n, a, size_k, &beta, c, size_n);
31. } else if (use_exllama) {
32.     // Quantized matmul
33.     int max_chunks = size_m / BLOCK_M_SIZE_MAX;
34.     int last_chunk = max_chunks * BLOCK_M_SIZE_MAX;
35.     int last_chunk_size = size_m - last_chunk;
37.     if (max_chunks) {
38.       gemm_half_q_half_cuda_part(a, b_q_weight, b_gptq_qzeros, b_gptq_scales,
39.                                  b_g_idx, c, last_chunk, size_n, size_k,
40.                                  BLOCK_M_SIZE_MAX, groups, bit);
41.     }
43.     if (last_chunk_size) {
44.       gemm_half_q_half_cuda_part(a + last_chunk * size_k, b_q_weight,
45.                                  b_gptq_qzeros, b_gptq_scales, b_g_idx,
46.                                  c + last_chunk * size_n, last_chunk_size,
47.                                  size_n, size_k, last_chunk_size, groups, bit);
48.     }
49.   } else {
50.     gemm_half_q_half_alt(a, b_q_weight, b_gptq_qzeros, b_gptq_scales, b_g_idx,
51.                          c, size_m, size_n, size_k, bit);
52.   }

复制代码

就是这部门代码，虽然现在看来比较明确它重要是走了use_reconstruct=True的分支，但是当时着实看了我很久的时间，要搞清晰走了哪个分支，得先知道这里的size_m代表着什么，它实在表示着输入a的行数，也就是seq_len*batch_size，而vllm在profile的时间，会使用到max_token_len大的seq_len。
大部门应该都是大于MAX_Q_GEMM_ROWS(=50)的，所以大部门是走了use_reconstruct=True的分支。这里我并没有深入研究reconstruct_exllama和reconstruct_gptq之间的差别点在哪，之后可以研究一下。
通过Nvidia的工具包定位问题

虽然代码大概看完了，但是我照旧不知道到底是什么函数出问题了呀，那就只能用老法子了，要么打印，要么用profile工具。所以我就问了问GPT，它给我保举了Nsight Compute，这是Nvidia出的一个工具，可以用来分析cuda程序的性能。吭哧吭哧学习了一下怎么用，然后现实给了我一顿暴击，Nsight Compute不支持Pascal架构，它的2019的版本才支持，但是2019的版本和现在的cuda版本又不兼容，尴尬。。。
不外幸运的是，在学习使用Nsight Compute的时间，我发现了Nsight System，这个也是Nvidia出的一个工具，可以用来分析cuda程序，看CPU和GPU联动的时间，问题出在哪，虽然不会像Nsight Compute那样详细地分析GPU的各个执行过程，但它能简单地分析cuda内核函数的耗时，这个恰好是我现在必要的。
上结果：

图中有两个关键信息：

• 大部门的耗时都在2个内核函数上，就是maxwell_hgemm_128x128和maxwell_hgemm_128x64。
  
• 在执行这俩函数前，都在执行reconstruct_exllama内核函数。
  

这样的话就比较轻易定位了，就是看reconstruct_exllama背面执行了什么，那不就是cublasHgemm么。
和cublasHgemm较劲

颠末一番搜索后，我了解了这个函数是啥，然后我就有点楞住了，啊？凭啥？这个是CuBLAS的函数，是Nvidia写的专门用来做向量和矩阵计算的，这怎么会有问题呢？这怎么能有问题呢？
为了验证它，我让GPT帮我写了个比较大的矩阵乘法并计算1000次来验证，结果确实是它的问题，执行起来很慢很慢，在A6000的呆板上结果险些是秒出，而我这边就会卡很久很久。
在这里我卡壳了很久，不知道这种环境下该咋办，感觉Pascal显卡就是该入土了，乃至想放弃了。背面想到，pytorch和xformers的性能不是没啥问题么，那肯定是有法子解决的。
于是我想了一个尝试的路子，我能不能换个库？我就去搜索了一下有没有CuBLAS的替换库。问了下GPT，还真就让我找到了，它就是CUTLASS，一个开源的CuBLAS库。
于是我就吭哧吭哧地又去编译了一下CUTLASS，3.0版本开始的CUTLASS就不支持PASCAL了，所以我只能用2.11版本。编译起来倒黑白常丝滑，没有任何问题，和最新的cuda也能兼容。
编译完成后，我照旧按照老思路，先找找看它的profile工具，确实有这个工具，于是我就举行了一次profile，就是CUTLASS的这次profile，帮我找到了问题的根因，官方的profile示例给的是用sgemm kernel:./tools/profiler/cutlass_profiler --kernels=sgemm --m=4352 --n=4096 --k=4096，我这边测试下来很快5s左右就执行完了，性能指标看着也还行：

1. Runtime: 15.7136  ms
2. Memory: 12.4296 GiB/s
3. Math: 9295.45 GFLOP/s

复制代码

当时我并不知道sgemm kernel的s表示什么，但我猜到了和精度相干，我一开始还猜是small（实在它表示单精度single-precision），就是精度很低，我就想，之前不是都是hgemm（半精度）么，我也来试试看它的profile是不是有这个kernel，这里纯属手贱，并不是想到了什么。但是就是这么一个不测，帮我找到了本次问题的根因。测试的结果是极其慢：

1. Runtime: 739.977  ms
2. Memory: 0.131972 GiB/s
3. Math: 197.391 GFLOP/s

复制代码

我当时就在想，这差距也太大了吧，就算是small，也不应该small得这么锋利，能差这么多呀。我就又测了一下dgemm（双精度），结果和hgemm根本类似。
然后我就去确认了一下，sgemm表示的是单精度的运算。到这，我根本上能知道怎么回事了，大概率是Pascal架构不支持半精度的运算，导致计算效率很低。为了验证我这个想法，趁便作为学习，我去翻了Nvidia的官网，找了各个时期的架构白皮书，看了一下里面重要的显卡性能先容。为了方便比较我先是让LLM帮我从各个白皮书里提取了性能信息，然后让它帮我输出json，我再用pandas将json转成了html方便我直观地对比。
这里给认识游戏显卡的同学稍微科普一下Nvidia的架构历史，从Maxwell开始：

• Maxwell 架构
  
  
  
  • 发布时间：2014年
    
  • 游戏卡定名：GTX 9xx 系列，如 GTX 970, GTX 980
    
  • 数据卡定名：Tesla Mxx 系列，如 Tesla M40, Tesla M60
    
  
  
• Pascal 架构
  
  
  
  • 发布时间：2016年
    
  • 游戏卡定名：GTX 10xx 系列，如 GTX 1070, GTX 1080, GTX 1080 Ti
    
  • 数据卡定名：Tesla Pxx 系列，如 Tesla P100
    
  
  
• Volta 架构
  
  
  
  • 发布时间：2017年
    
  • 游戏卡：N/A
    
  • 数据卡定名：Tesla Vxx系列，如 Tesla V100
    
  
  
• Turing 架构
  
  
  
  • 发布时间：2018年
    
  • 游戏卡定名：RTX 20xx 系列，如 RTX 2070, RTX 2080, RTX 2080 Ti; GTX 16xx 系列如 GTX 1660, GTX 1660 Ti（不包含RT核的变体）
    
  • 数据卡定名：Tesla Txx 系列，如 Tesla T4
    
  
  
• Ampere 架构
  
  
  
  • 发布时间：2020年
    
  • 游戏卡定名：RTX 30xx 系列，如 RTX 3070, RTX 3080, RTX 3090
    
  • 数据卡定名：A100, A30
    
  
  
• Ada Lovelace 架构
  
  
  
  • 发布时间：2022年
    
  • 游戏卡定名：RTX 40xx 系列，如 RTX 4070, RTX 4080, RTX 4090
    
  • 数据卡定名：L4
    
  
  
• Hopper 架构
  
  
  
  • 发布时间：2022年
    
  • 游戏卡定名：N/A
    
  • 数据卡定名：H100
    
  
  
• Blackwell 架构
  
  
  
  • 发布时间：2024年
    
  • 游戏卡定名：N/A
    
  • 数据卡定名：B100
    
  
  

可以看到，Pascal架构的P100并没有fp16的支持， 而要有fp16支持的条件也是tensor core，Pascal架构是没有tensor core，只有cuda core的。然后也能发现，为什么说4090的推理性能能强过A100，由于它的各个算力指标都好于A100，A100强的是它显存大，显存带宽大，有SXM的支持，显卡之间的互联带宽高，所以在训练上有巨大的上风。
这下百分百确定问题所在了，没有fp16的支持，计算本领自然就很弱了。
第四关：优化性能

接下来就是改代码了，我的第一个想法是直接改成fp32的计算，这样计算速度就有保障了。但我照旧决定去问一下LLM，看它有什么好的发起。它给我的发起是使用cublasGemmEx函数，这个函数也是CuBLAS的函数，它允许我们的输入输出矩阵都是fp16的，但是在计算的时间，转换成fp32来举行计算。
最后的改动就是这样：

1.     // cublasHgemm(cublas_handle, CUBLAS_OP_N, CUBLAS_OP_N, size_n, size_m, size_k,
2.     //             &alpha, temp_dq, size_n, a, size_k, &beta, c, size_n);
3.     cublasGemmEx(
4.       cublas_handle,                // Handle
5.       CUBLAS_OP_N,                  // transa
6.       CUBLAS_OP_N,                  // transb
7.       size_n,                       // m
8.       size_m,                       // n
9.       size_k,                       // k
10.       &alpha,                       // alpha
11.       temp_dq,                      // A
12.       CUDA_R_16F,                   // A type
13.       size_n,                       // lda
14.       a,                            // B
15.       CUDA_R_16F,                   // B type
16.       size_k,                       // ldb
17.       &beta,                        // beta
18.       c,                            // C
19.       CUDA_R_16F,                   // C type
20.       size_n,                       // ldc
21.       CUDA_R_32F,                   // computeType (FP32 for accumulation)
22.       CUBLAS_GEMM_DFALT_TENSOR_OP   // algo (default with potential Tensor Core usage)
23.     );

复制代码

结果就如标题所说，这一行代码的更改，让性能提升了43倍，现在再来看一下我之前的pytorch的耗时日记：

1. model: 3098.3325251191854 ms
2.   model.embed_tokens: 33.70710372924805 ms
3.   model.layers: 3064.419405385852 ms
4.     model.layers.0: 131.46515500545502 ms
5.       model.layers.0.input_layernorm: 0.6445760130882263 ms
6.       model.layers.0.self_attn: 30.52022334933281 ms
7.         model.layers.0.self_attn.qkv_proj: 20.16111946105957 ms
8.         model.layers.0.self_attn.rotary_emb: 0.16473600268363953 ms
9.         model.layers.0.self_attn.attn: 3.8500161170959473 ms
10.         model.layers.0.self_attn.o_proj: 6.344351768493652 ms
11.       model.layers.0.post_attention_layernorm: 0.22275200486183167 ms
12.       model.layers.0.mlp: 100.07760363817215 ms
13.         model.layers.0.mlp.gate_up_proj: 65.92633819580078 ms
14.         model.layers.0.mlp.act_fn: 0.9378560185432434 ms
15.         model.layers.0.mlp.down_proj: 33.213409423828125 ms
16.     model.layers.1: 115.83395344018936 ms
17.       model.layers.1.self_attn: 17.98700802028179 ms
18.         model.layers.1.self_attn.rotary_emb: 0.16617600619792938 ms

复制代码

可以看到，vllm的profile的耗时，从134s降到了3s，性能整整提升了43倍呀！！！
终于可以用我的Pascal显卡来推理了，爽！！
总结

第1点

对于一些程序员新人来说，希望这次的履历能给你一个参考，我们可以从一个问题点（一个好的问题从哪来确实也挺看运气的，我这次的问题刚好就是一个很深的问题，但是有时间我们可以刻意去创造一个问题，好比之前我看spark源码的时间，就是想搞清晰一个job的启动过程到底是怎么样的，这样也算是本身提出的一个好问题了）开始，然后不停深挖下去，这样你就认识了从外貌不停到内核的整个过程，然后你就可以选择在恣意感爱好的地方开枝散叶，就能认识一整个框架乃至领域了。
对于我本身来说，我接下来能研究的就有：

• 再去研究一下GPTQ的量化过程，把数据原理完全搞懂，有机会的话本身可以跑一遍模型量化
  
• 看看GGUF的量化是怎么做的
  
• 看看GEMM具体是怎么计算的，有哪些点可以做来举行优化
  
• 去看看xformers的留意力计算是怎么做的
  
• 去看看vllm的kv cache是怎么做的
  
• 也可以去学学cuda编程
  
• …
  

第2点

LLM在整个过程中起到了很大的作用，包括不限于：

• 表明一些源码
  
• 帮忙写部门测试用的代码
  
• 帮忙澄清一些概念
  
• 帮忙表明一些bug
  
• …
  

所以，赶紧用起来吧！
第3点

没事别瞎折腾别人不支持的东西，人家不支持是有原因的，除非你有折腾的觉悟和爱好。
参考资料

• https://stackoverflow.com/questions/75682385/runtimeerror-cuda-error-no-kernel-image-is-available-for-execution-on-the-devi
  
• https://docs.vllm.ai/en/latest/usage/compatibility_matrix.html
  
• https://developer.nvidia.com/cuda-gpus
  
• https://github.com/pytorch/pytorch/issues/31285
  
• https://pytorch.org/tutorials/recipes/recipes/benchmark.html
  
• https://github.com/facebookresearch/xformers/blob/main/BENCHMARKS.md
  
• https://blog.csdn.net/fan_fan_feng/article/details/138978901
  
• https://www.bilibili.com/video/BV17m411f7Cm?spm_id_from=333.788.videopod.sections&vd_source=68452628e4137592ea9efa4793a102a6
  

免责声明：如果侵犯了您的权益，请联系站长，我们会及时删除侵权内容，谢谢合作！更多信息从访问主页：qidao123.com:ToB企服之家，中国第一个企服评测及商务社交产业平台。