解决深度学习和LLM中的可复现性挑战：我们的旅程

ingonyama
发布于 2024-09-23 11:48
阅读 12

本文讲述了在深度学习和LLM中实现可复现性的挑战，特别是使用Nvidia GPU时，浮点运算的非结合性以及硬件差异会导致结果不一致。文章分享了通过重写GEMM CUDA kernel，避免使用Tensor Cores，并确保运算顺序的确定性来解决这一问题，最终在不同硬件上实现了模型输出的一致性。

## **解决深度学习和 LLM 中的可复现性挑战：我们的历程**

深度学习中的可复现性长期以来一直是一个关键问题。无论是在自动驾驶汽车，医疗系统，还是仅仅是复现一项科学实验的结果，确保一致性都是关键。随着大型语言模型（LLM）的兴起，这一挑战变得更加突出，特别是当我们希望我们的模型始终如一地遵循特定模式时，例如在安全关键型应用中。

这是一个众所周知的问题（例如，[深度学习中的确定性](https://developer.nvidia.com/gtc/2019/video/s9911)，[使用 PyTorch 实现可复现的深度学习](https://learnblockchain.cn/article/14860)），并且据我们所知，对于一般情况，仍然没有规范的解决方案——只有减轻其影响的最佳实践。引用 Pytorch 关于可复现性的[文档](https://pytorch.org/docs/stable/notes/randomness.html)：

_“不能保证在 PyTorch 版本、各个提交或不同平台之间完全可复现的结果。此外，即使使用相同的种子，CPU 和 GPU 执行之间的结果也可能无法复现。”_

![](https://img.learnblockchain.cn/2025/04/26/51972051_image.jpg)

“创建一张工作室灯光下的红色苹果图片”

## **我们的动机**

当将 [**零知识证明 (ZKP)**](https://www.ingonyama.com/ingopedia/zkbeginner) 纳入推理或训练中以用于问责、监管或隐私目的时，不确定性是一个主要障碍：

经典的 ZK 证明器依赖于完全确定的算术约束（例如，R1CS、Plonkish、GKR）。如果没有确定性，该证明将用于无法复现或验证的跟踪。

作为一种高度并行的计算，ZK 证明器广泛依赖于硬件加速器，其中 NVIDIA GPU（通常用于 AI）起着主导作用。不幸的是，证明非确定性计算的成本比已经昂贵的 ZKP 过程高出几个数量级。大多数 ZKP 用例都需要一个公共验证器，这意味着任何支持 CUDA 的设备都必须能够复现 AI 计算，而无需强制特定的硬件架构。有关更多详细信息，请参见 Tomer 的 [Encode 演讲](https://www.youtube.com/watch?v=ghU_-ADHBaw)。

在这篇博文中，我们将详细介绍我们为确保深度学习模型始终产生可复现的结果所做的努力。

## 根本原因：浮点运算的非结合性

我们可复现性问题的核心是**浮点数的非结合性**。浮点运算缺乏完美的精度，因为计算机以二进制形式表示十进制数。加法和乘法等运算会根据执行顺序和使用的硬件产生略有不同的结果。因此，结合律并不总是得到保证：(a + b) + c ≠ a + (b + c)

通过控制执行操作的顺序，我们可以确保计算在运行和硬件之间保持一致。

这种非结合性意味着，即使在相同的初始条件下，小的计算差异也会随着时间的推移而累积，从而导致模型输出中出现明显的差异。这在深度学习中尤其成问题，因为模型会执行数十亿次此类操作。

在我们的实验中，我们专注于使用 CUDA 及其衍生产品（如 [cuBLAS](https://developer.nvidia.com/cublas)）在 Nvidia 硬件上运行 DNN。尽管这种设置非常常见，但不能保证可复现性。

## 我们的设置

我们在三台不同的机器上进行了测试，每台机器都有不同的 GPU 和操作系统：

- 具有 NVIDIA RTX 3090 的 Ubuntu 20
- 具有 NVIDIA RTX 4080 的 Ubuntu 22.04
- 具有 NVIDIA L4  的 Centos

我们选择这些 GPU 是为了在来自不同世代和架构（Ampere 和 Ada Lovelace）的卡上，以及针对不同的用例（游戏与专业）测试我们的论点。所有机器都运行 CUDA Toolkit 12.0。

我们研究了多个框架中的问题，首先是 [PyTorch](https://pytorch.org/)，这是一个领先的深度学习开发平台，然后是 [llama.cpp](https://github.com/ggerganov/llama.cpp)，它作为 LLM 推理的流行框架而受到越来越多的关注，尤其是对于量化模型。

## 使用可复现性标志

我们的第一种方法是使用深度学习框架（如 PyTorch）提供的可复现性标志并控制随机数生成器。以下是我们使用的代码片段：

```
环境变量：
export CUBLAS_WORKSPACE_CONFIG=:4096:8
```

```
import random
import numpy as np
import torch

random.seed(0)  # 为 Python 内置的 random 模块设置种子
np.random.seed(0)  # 为 NumPy 的随机数生成器设置种子

torch.manual_seed(0)  # 为 PyTorch 的 CPU 随机数生成器设置种子
torch.cuda.manual_seed(0)  # 为当前 GPU 设备设置种子
torch.cuda.manual_seed_all(0)  # 为所有可用的 GPU 设备设置种子

torch.use_deterministic_algorithms(True)  # 确保仅使用确定性算法

torch.backends.cuda.matmul.allow_tf32 = False  # 禁用 matmul 运算上的 TensorFloat32 (TF32)
torch.backends.cudnn.allow_tf32 = False  # 禁用 cuDNN 上的 TF32
torch.backends.cudnn.benchmark = False  # 禁用 cuDNN 自动调谐器
torch.backends.cudnn.deterministic = True  # 强制 cuDNN 使用确定性算法
##在最坏的情况下，使用这个：
torch.backends.cudnn.enabled = False  # 完全禁用 cuDNN

```

## 每个标志的解释：

- **random.seed(0):** 为 Python 的内置随机数生成器设置种子，确保每次运行对 random 的任何使用都产生相同的结果。
- **np.random.seed(0):** 为 NumPy 的随机数生成器设置种子，以确保 NumPy 运算的一致性。
- **torch.manual\_seed(0):** 为 PyTorch 的 CPU 随机数生成器设置种子。
- **torch.cuda.manual\_seed(0) 和 torch.cuda.manual\_seed\_all(0):** 为 GPU 设置种子，确保 CUDA 运算是确定性的。
- **torch.use\_deterministic\_algorithms(True):** 强制 PyTorch 在可用时使用确定性算法；如果确定性算法不可用，它将引发错误。
- **torch.backends.cuda.matmul.allow\_tf32 = False 和 torch.backends.cudnn.allow\_tf32 = False:** 禁用 matmul 和 cuDNN 上的 TensorFloat32 (TF32) 运算，确保计算使用 FP32 精度，FP32 精度更高且更一致。
- **torch.backends.cudnn.benchmark = False:** 禁用 cuDNN 自动调谐器，该调谐器选择最快的Rollup算法；由于算法选择不同，这可能会引入不确定性。
- **torch.backends.cudnn.deterministic = True:** 强制 cuDNN 使用确定性算法，该算法可能较慢，但在运行中是一致的。
- **torch.backends.cudnn.enabled = False:** 完全禁用 cuDNN，确保仅使用纯 PyTorch 运算。这可能会显着影响性能，但会提高可复现性。

**结果：** 虽然这些标志有助于减少可变性，但它们并未完全消除具有不同硬件配置的机器之间的差异。这些模型仍然在我们的平台上产生略有不同的输出。

尽管设置了所有可能的种子并禁用了非确定性行为，但浮点运算固有的非结合性与硬件差异相结合，仍然导致不一致。

## 使用 Llama.cpp 运行量化模型

接下来，我们转向 [llama.cpp](https://github.com/ggerganov/llama.cpp)，这是一个允许使用量化权重运行 LLM 以显着减少计算负载的项目。Llama.cpp 定义了一种流行的量化格式，称为 **GGUF**，它支持各种级别的量化，包括 int8。

**预期：** 我们期望以量化格式（例如 int8）运行模型可以解决可复现性问题。我们的假设是，使用整数而不是浮点数将消除不一致。

**现实：** 但是，我们遇到了两个关键问题：

1. **不完整的量化：** 并非模型中的所有层都已量化。具体来说，**1D 归一化层**（例如 LayerNorm）仍然未量化，继续以浮点精度执行计算，这会重新引入不确定性。[来源](https://github.com/ggerganov/llama.cpp/issues/397#issuecomment-1510860290)
2. **运行时反量化：** 即使对于量化层，权重在 **推理期间也会被反量化**，以适应计算 logits 等运算。因此，尽管以量化格式存储权重，但模型仍然依赖于运行时浮点计算，这会受到非结合性和硬件可变性的影响。

因此，尽管使用了量化和 GGUF 格式，但我们仍无法在不同的机器上实现可复现性。量化和未量化计算的结合，以及运行时反量化，重新引入了我们试图消除的浮点运算！

**本次尝试的结论：** 仅靠量化不能确保可复现性。未量化层的存在以及在推理期间反量化权重的需求意味着浮点运算（及其固有的不确定性）仍然是计算管道的一部分。

## 突破：重写 GEMM CUDA 内核

意识到问题可能源于矩阵乘法运算，我们决定在内核级别解决问题。

在为不同的架构编译代码后，我们将不确定性的来源追溯到 PTX 文件（本质上是 CUDA 机器代码）。我们观察到，某些内核会根据目标架构生成不同的指令。因此，我们开始研究矩阵乘法内核，它们是我们 AI 计算的核心。

我们进行了一项隔离 GEMM 内核的实验。

CuBLAS 库提供的 GEMM 内核（[链接](https://docs.nvidia.com/deeplearning/performance/dl-performance-matrix-multiplication/index.html)）是深度学习和并行计算加速中使用最广泛和优化的内核之一。但是，这种高水平的优化有时会以可复现性为代价，导致各种硬件上的结果不同。

当我们在不同的平台（L4、3090、4080）上运行这些内核时，结果的差异为 1e-4。虽然这看起来是一个很小的差异，但在 LLM 的上下文中，其中每个 token 都取决于前一个 token，这种错误会迅速累积并导致输出发散。

## 实现确定性 GEMM CUDA 内核

我们重写了被识别为非确定性的 llama.cpp 使用的通用矩阵乘法 (GEMM) CUDA 内核。我们的目标是提供每个内核的更简单、更确定的版本，而不影响性能。我们的策略是：

- **避免使用 Tensor Cores 并且仅使用 CUDA Cores：** Tensor cores 允许 Nvidia 不断创新，并在每一代中更新其架构和功能，而 CUDA cores 在需要时可靠地提供向后兼容性。通过将计算限制为 CUDA cores，我们确保在不同的架构上以一致的方式执行运算，因为不同的架构可能具有不同代的 Tensor cores。更多阅读材料请参见[此处](https://www.digitalocean.com/community/tutorials/understanding-tensor-cores)。
- **确保运算的确定性顺序：** 我们仔细管理浮点运算的顺序，以消除非结合性影响。

**我们的内核示例：C/C++**

```
__global__ void ingo_mul_mat_fp16_fp16_kernel(const half * __restrict__ A, const half * __restrict__ B, half * __restrict__ C, int m, int n, int k, int lda, int ldb, int ldc, bool transpose_a) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row < m && col < n) {
        half sum = 0.0f;
        if (!transpose_a) {
            // 无转置：正常矩阵乘法 A * B
            // A 以列优先访问，B 也以列优先访问
            for (int i = 0; i < k; i++) {
                sum += A[i * lda + row] * B[col * ldb + i];
            }
        } else {
            // 转置 A：乘以 A^T * B
            // A^T 实际上被访问，就好像 A 是行优先一样
            for (int i = 0; i < k; i++) {
                sum += A[row * lda + i] * B[col * ldb + i];
            }
        }
        C[col * ldc + row] = sum;  // 以列优先顺序存储结果
    }
}

void ingo_mul_mat(const half * A, const half * B, half * C,
                           int m, int n, int k, int lda, int ldb, int ldc, bool transpose_a, cudaStream_t stream) {
    dim3 blockSize(16, 16);
    dim3 gridSize((n + blockSize.x - 1) / blockSize.x, (m + blockSize.y - 1) / blockSize.y);

ingo_mul_mat_fp16_fp16_kernel<<<gridSize, blockSize, 0, stream>>>(A, B, C, m, n, k, lda, ldb, ldc, transpose_a);
    checkCudaError("kernel launch");

cudaDeviceSynchronize();
    checkCudaError("synchronizing after kernel execution");
}
```

**结果：** 这种方法成功解决了我们的可复现性问题！在所有三台机器上运行的模型都产生了相同的输出，证实我们已经减轻了硬件引起的可变性。有关 GEMM 内核优化的更多见解，请查看 Anthropic 的 Simon Boehm 的[这篇](https://siboehm.com/articles/22/CUDA-MMM)精彩博文。

我们已经在以下模型中测试了我们的结果：

- llama2–7b-fp16
- llama2–7b-Q8
- mistral-7b-instruct-v0.2.Q8\_0
- mixtral-8x7b.Q2
- Mistrall-7b-instruct-v0.2.fp\_16
- Meta-Llama-3–8B-fp16-gguf
- Meta-Llama-3–8B.Q8\_0.gguf

此外，我们通过评估我们的确定性内核的性能来扩展我们的实验。经过一轮基本的优化，处理多个线程的批处理量，我们观察到性能主要在提示处理阶段受到影响（302 tps vs. 43 tps），而在文本生成阶段保持在 44-45 tps 之间。通过进一步优化，这可能会得到改善。

这些测量是使用在 RTX 4080 上运行的 llama.cpp 推理代码进行的。

原始 llama.cpp: C/C++

```
llama_print_timings:        load time =    1311.27 ms
llama_print_timings:      sample time =       5.33 ms /   300 runs   (    0.02 ms per token, 56285.18 tokens per second)
llama_print_timings: prompt eval time =      23.14 ms /     7 tokens (    3.31 ms per token,   302.48 tokens per second)
llama_print_timings:        eval time =    6627.57 ms /   299 runs   (   22.17 ms per token,    45.11 tokens per second)
llama_print_timings:       total time =    6702.29 ms /   306 tokens
```

```
llama_print_timings:        load time =    1322.69 ms
llama_print_timings:      sample time =       5.18 ms /   300 runs   (    0.02 ms per token, 57881.54 tokens per second)
llama_print_timings: prompt eval time =     161.55 ms /     7 tokens (   23.08 ms per token,    43.33 tokens per second)
llama_print_timings:        eval time =    6739.35 ms /   299 runs   (   22.54 ms per token,    44.37 tokens per second)
llama_print_timings:       total time =    6963.47 ms /   306 tokens
```

## 结论

- **硬件很重要：** 即使运行相同的代码，GPU 架构的差异也会影响可复现性。
- **量化不是万能的：** 虽然量化减少了模型大小和计算量，但它本身并不能解决可复现性问题。
- **深度控制是必要的：** 有时，确保可复现性需要更深入地研究堆栈，甚至深入到内核实现。

深度学习中的可复现性是一个多方面的挑战，需要关注软件和硬件细节。通过了解根本原因并愿意深入研究底层实现，我们可以克服这些挑战。

我们的历程强调了细致工程的重要性，并为面临类似问题的其他人提供了路线图。我们希望我们的经验能够帮助社区朝着更可靠和一致的深度学习实践迈进。

![](https://img.learnblockchain.cn/2025/04/26/56147942_image.jpg)

## 接下来的步骤

我们计划在以下方向扩展我们的研究：

- **与框架开发人员合作：** 与 PyTorch 和 TensorFlow 团队合作以集成确定性内核。
- **探索跨硬件供应商的可复现性：** 专注于在各种硬件平台上实现可复现性，包括 Nvidia、AMD、CPU 和专用硬件。
- **分布式系统中的可复现性：** 将我们的工作扩展到多节点设置，解决分布式环境中可复现性的额外挑战。
- **将研究结果应用于 ZKP 框架：** 与 ZKP 领域的开发人员合作，以确保我们的解决方案满足零知识证明的严格要求。

如果你有任何问题，或者你有兴趣合作改进深度学习和零知识证明中的可复现性，请随时与我们联系。

## 关注 Ingonyama

Twitter: [https://twitter.com/Ingo\_zk](https://twitter.com/Ingo_zk)

YouTube: [https://www.youtube.com/@ingo\_zk](https://www.youtube.com/@ingo_zk)

GitHub: [https://github.com/ingonyama-zk](https://github.com/ingonyama-zk)

LinkedIn: [https://www.linkedin.com/company/ingonyama](https://www.linkedin.com/company/ingonyama)

加入我们: [https://www.ingonyama.com/career](https://www.ingonyama.com/careers)

>- 原文链接： [medium.com/@ingonyama/so...](https://medium.com/@ingonyama/solving-reproducibility-challenges-in-deep-learning-and-llms-our-journey-5e9f1f111da2)
>- 登链社区 AI 助手，为大家转译优秀英文文章，如有翻译不通的地方，还请包涵～

解决深度学习和 LLM 中的可复现性挑战：我们的历程

这是一个众所周知的问题（例如，深度学习中的确定性，使用 PyTorch 实现可复现的深度学习），并且据我们所知，对于一般情况，仍然没有规范的解决方案——只有减轻其影响的最佳实践。引用 Pytorch 关于可复现性的文档：

“不能保证在 PyTorch 版本、各个提交或不同平台之间完全可复现的结果。此外，即使使用相同的种子，CPU 和 GPU 执行之间的结果也可能无法复现。”

“创建一张工作室灯光下的红色苹果图片”

我们的动机

当将 零知识证明 (ZKP) 纳入推理或训练中以用于问责、监管或隐私目的时，不确定性是一个主要障碍：

经典的 ZK 证明器依赖于完全确定的算术约束（例如，R1CS、Plonkish、GKR）。如果没有确定性，该证明将用于无法复现或验证的跟踪。

在这篇博文中，我们将详细介绍我们为确保深度学习模型始终产生可复现的结果所做的努力。

根本原因：浮点运算的非结合性

我们可复现性问题的核心是浮点数的非结合性。浮点运算缺乏完美的精度，因为计算机以二进制形式表示十进制数。加法和乘法等运算会根据执行顺序和使用的硬件产生略有不同的结果。因此，结合律并不总是得到保证：(a + b) + c ≠ a + (b + c)

通过控制执行操作的顺序，我们可以确保计算在运行和硬件之间保持一致。

在我们的实验中，我们专注于使用 CUDA 及其衍生产品（如 cuBLAS）在 Nvidia 硬件上运行 DNN。尽管这种设置非常常见，但不能保证可复现性。

我们的设置

我们在三台不同的机器上进行了测试，每台机器都有不同的 GPU 和操作系统：

具有 NVIDIA RTX 3090 的 Ubuntu 20
具有 NVIDIA RTX 4080 的 Ubuntu 22.04
具有 NVIDIA L4 的 Centos

我们研究了多个框架中的问题，首先是 PyTorch，这是一个领先的深度学习开发平台，然后是 llama.cpp，它作为 LLM 推理的流行框架而受到越来越多的关注，尤其是对于量化模型。

使用可复现性标志

我们的第一种方法是使用深度学习框架（如 PyTorch）提供的可复现性标志并控制随机数生成器。以下是我们使用的代码片段：

环境变量：
export CUBLAS_WORKSPACE_CONFIG=:4096:8

import random
import numpy as np
import torch

random.seed(0)  # 为 Python 内置的 random 模块设置种子
np.random.seed(0)  # 为 NumPy 的随机数生成器设置种子

torch.manual_seed(0)  # 为 PyTorch 的 CPU 随机数生成器设置种子
torch.cuda.manual_seed(0)  # 为当前 GPU 设备设置种子
torch.cuda.manual_seed_all(0)  # 为所有可用的 GPU 设备设置种子

torch.use_deterministic_algorithms(True)  # 确保仅使用确定性算法

torch.backends.cuda.matmul.allow_tf32 = False  # 禁用 matmul 运算上的 TensorFloat32 (TF32)
torch.backends.cudnn.allow_tf32 = False  # 禁用 cuDNN 上的 TF32
torch.backends.cudnn.benchmark = False  # 禁用 cuDNN 自动调谐器
torch.backends.cudnn.deterministic = True  # 强制 cuDNN 使用确定性算法
##在最坏的情况下，使用这个：
torch.backends.cudnn.enabled = False  # 完全禁用 cuDNN

每个标志的解释：

random.seed(0): 为 Python 的内置随机数生成器设置种子，确保每次运行对 random 的任何使用都产生相同的结果。
np.random.seed(0): 为 NumPy 的随机数生成器设置种子，以确保 NumPy 运算的一致性。
torch.manual_seed(0): 为 PyTorch 的 CPU 随机数生成器设置种子。
torch.cuda.manual_seed(0) 和 torch.cuda.manual_seed_all(0): 为 GPU 设置种子，确保 CUDA 运算是确定性的。
torch.use_deterministic_algorithms(True): 强制 PyTorch 在可用时使用确定性算法；如果确定性算法不可用，它将引发错误。
torch.backends.cuda.matmul.allow_tf32 = False 和 torch.backends.cudnn.allow_tf32 = False: 禁用 matmul 和 cuDNN 上的 TensorFloat32 (TF32) 运算，确保计算使用 FP32 精度，FP32 精度更高且更一致。
torch.backends.cudnn.benchmark = False: 禁用 cuDNN 自动调谐器，该调谐器选择最快的Rollup算法；由于算法选择不同，这可能会引入不确定性。
torch.backends.cudnn.deterministic = True: 强制 cuDNN 使用确定性算法，该算法可能较慢，但在运行中是一致的。
torch.backends.cudnn.enabled = False: 完全禁用 cuDNN，确保仅使用纯 PyTorch 运算。这可能会显着影响性能，但会提高可复现性。

结果： 虽然这些标志有助于减少可变性，但它们并未完全消除具有不同硬件配置的机器之间的差异。这些模型仍然在我们的平台上产生略有不同的输出。

尽管设置了所有可能的种子并禁用了非确定性行为，但浮点运算固有的非结合性与硬件差异相结合，仍然导致不一致。

使用 Llama.cpp 运行量化模型

接下来，我们转向 llama.cpp，这是一个允许使用量化权重运行 LLM 以显着减少计算负载的项目。Llama.cpp 定义了一种流行的量化格式，称为 GGUF，它支持各种级别的量化，包括 int8。

预期： 我们期望以量化格式（例如 int8）运行模型可以解决可复现性问题。我们的假设是，使用整数而不是浮点数将消除不一致。

现实： 但是，我们遇到了两个关键问题：

不完整的量化： 并非模型中的所有层都已量化。具体来说，1D 归一化层（例如 LayerNorm）仍然未量化，继续以浮点精度执行计算，这会重新引入不确定性。来源
运行时反量化： 即使对于量化层，权重在 推理期间也会被反量化，以适应计算 logits 等运算。因此，尽管以量化格式存储权重，但模型仍然依赖于运行时浮点计算，这会受到非结合性和硬件可变性的影响。

本次尝试的结论： 仅靠量化不能确保可复现性。未量化层的存在以及在推理期间反量化权重的需求意味着浮点运算（及其固有的不确定性）仍然是计算管道的一部分。

突破：重写 GEMM CUDA 内核

意识到问题可能源于矩阵乘法运算，我们决定在内核级别解决问题。

我们进行了一项隔离 GEMM 内核的实验。

CuBLAS 库提供的 GEMM 内核（链接）是深度学习和并行计算加速中使用最广泛和优化的内核之一。但是，这种高水平的优化有时会以可复现性为代价，导致各种硬件上的结果不同。

实现确定性 GEMM CUDA 内核

避免使用 Tensor Cores 并且仅使用 CUDA Cores： Tensor cores 允许 Nvidia 不断创新，并在每一代中更新其架构和功能，而 CUDA cores 在需要时可靠地提供向后兼容性。通过将计算限制为 CUDA cores，我们确保在不同的架构上以一致的方式执行运算，因为不同的架构可能具有不同代的 Tensor cores。更多阅读材料请参见此处。
确保运算的确定性顺序： 我们仔细管理浮点运算的顺序，以消除非结合性影响。

我们的内核示例：C/C++

__global__ void ingo_mul_mat_fp16_fp16_kernel(const half * __restrict__ A, const half * __restrict__ B, half * __restrict__ C, int m, int n, int k, int lda, int ldb, int ldc, bool transpose_a) {
    int row = blockIdx.y * blockDim.y + threadIdx.y;
    int col = blockIdx.x * blockDim.x + threadIdx.x;
    if (row &lt; m && col &lt; n) {
        half sum = 0.0f;
        if (!transpose_a) {
            // 无转置：正常矩阵乘法 A * B
            // A 以列优先访问，B 也以列优先访问
            for (int i = 0; i &lt; k; i++) {
                sum += A[i * lda + row] * B[col * ldb + i];
            }
        } else {
            // 转置 A：乘以 A^T * B
            // A^T 实际上被访问，就好像 A 是行优先一样
            for (int i = 0; i &lt; k; i++) {
                sum += A[row * lda + i] * B[col * ldb + i];
            }
        }
        C[col * ldc + row] = sum;  // 以列优先顺序存储结果
    }
}

void ingo_mul_mat(const half * A, const half * B, half * C,
                           int m, int n, int k, int lda, int ldb, int ldc, bool transpose_a, cudaStream_t stream) {
    dim3 blockSize(16, 16);
    dim3 gridSize((n + blockSize.x - 1) / blockSize.x, (m + blockSize.y - 1) / blockSize.y);

    ingo_mul_mat_fp16_fp16_kernel&lt;&lt;&lt;gridSize, blockSize, 0, stream>>>(A, B, C, m, n, k, lda, ldb, ldc, transpose_a);
    checkCudaError("kernel launch");

    cudaDeviceSynchronize();
    checkCudaError("synchronizing after kernel execution");
}

结果： 这种方法成功解决了我们的可复现性问题！在所有三台机器上运行的模型都产生了相同的输出，证实我们已经减轻了硬件引起的可变性。有关 GEMM 内核优化的更多见解，请查看 Anthropic 的 Simon Boehm 的这篇精彩博文。

我们已经在以下模型中测试了我们的结果：

llama2–7b-fp16
llama2–7b-Q8
mistral-7b-instruct-v0.2.Q8_0
mixtral-8x7b.Q2
Mistrall-7b-instruct-v0.2.fp_16
Meta-Llama-3–8B-fp16-gguf
Meta-Llama-3–8B.Q8_0.gguf

这些测量是使用在 RTX 4080 上运行的 llama.cpp 推理代码进行的。

原始 llama.cpp: C/C++

llama_print_timings:        load time =    1311.27 ms
llama_print_timings:      sample time =       5.33 ms /   300 runs   (    0.02 ms per token, 56285.18 tokens per second)
llama_print_timings: prompt eval time =      23.14 ms /     7 tokens (    3.31 ms per token,   302.48 tokens per second)
llama_print_timings:        eval time =    6627.57 ms /   299 runs   (   22.17 ms per token,    45.11 tokens per second)
llama_print_timings:       total time =    6702.29 ms /   306 tokens

llama_print_timings:        load time =    1322.69 ms
llama_print_timings:      sample time =       5.18 ms /   300 runs   (    0.02 ms per token, 57881.54 tokens per second)
llama_print_timings: prompt eval time =     161.55 ms /     7 tokens (   23.08 ms per token,    43.33 tokens per second)
llama_print_timings:        eval time =    6739.35 ms /   299 runs   (   22.54 ms per token,    44.37 tokens per second)
llama_print_timings:       total time =    6963.47 ms /   306 tokens