寫(xiě)在前面：本文將對(duì) Nvidia BERT 推理解決方案 Faster Transformer 源碼進(jìn)行深度剖析，詳細(xì)分析作者的優(yōu)化意圖，并對(duì)源碼中的加速技巧進(jìn)行介紹，希望對(duì)讀者有所幫助。本文源碼解讀的內(nèi)容僅限 Faster Transformer v1.0 版本，更高版本的源碼將在后續(xù)文章中繼續(xù)解讀。

1 Faster Transformer

Transformer 模型最早在 2017 年由谷歌在論文中提出，拋棄了以往深度學(xué)習(xí)任務(wù)里面使用到的 CNN 和 RNN，取而代之的是一種 self-Attention 的結(jié)構(gòu)，將 Attention 思想發(fā)揮到了極致，一定程度上解決了 RNN 的長(zhǎng)序列信息丟失問(wèn)題，基本取代了 RNN 的地位。
Transformer 一經(jīng)面世便在 NLP 領(lǐng)域大殺四方，近兩年一些文章開(kāi)創(chuàng)性地將 Transformer 模型跨領(lǐng)域地引用到了 CV 任務(wù)中，并取得了不錯(cuò)地成果。這也被許多學(xué)者認(rèn)為是開(kāi)創(chuàng)了 CV 領(lǐng)域的新時(shí)代，甚至可能完全取代傳統(tǒng)的卷積操作。
雖然 Transformer在多種場(chǎng)景下都有優(yōu)秀的表現(xiàn)，但是在推理部署階段，其計(jì)算性能卻受到了巨大的挑戰(zhàn)：以 BERT 為原型的多層 Transformer 模型，其性能常常難以滿(mǎn)足在線業(yè)務(wù)對(duì)于低延遲和高吞吐的要求。以 BERT-BASE 為例，超過(guò) 90% 的計(jì)算時(shí)間消耗在 12 層 Transformer 的前向計(jì)算上。因此，一個(gè)高效的 Transformer 前向計(jì)算方案，既可以為在線業(yè)務(wù)帶來(lái)降本增效的作用，也有利于以 Transformer 結(jié)構(gòu)為核心的各類(lèi)網(wǎng)絡(luò)在更多實(shí)際工業(yè)場(chǎng)景中落地。
基于上述背景，NVIDIA GPU 計(jì)算專(zhuān)家團(tuán)隊(duì)針對(duì) Transformer 推理提出了的性能優(yōu)化方案：Faster Transformer。
Faster Transformer 是一個(gè) BERT Transformer 單層前向計(jì)算的高效實(shí)現(xiàn)，其代碼簡(jiǎn)潔明了，后續(xù)可以通過(guò)簡(jiǎn)單修改支持多種 Transformer 結(jié)構(gòu)。目前優(yōu)化集中在編碼器（encoder）的前向計(jì)算。底層由 CUDA 和 cuBLAS 實(shí)現(xiàn)，支持 FP16 和 FP32 兩種計(jì)算模式，其中 FP16 可以充分利用 Volta 和 Turing 架構(gòu) GPU 上的 Tensor Core 計(jì)算單元。

2 優(yōu)化原理

在深入了解 Faster Transformer 的優(yōu)化原理之前，我們先來(lái)了解一下主流深度學(xué)習(xí)框架 Tensorflow 中 Transformer 的實(shí)現(xiàn)情況，僅僅以一個(gè)基本的激活函數(shù) gelu 為例，這個(gè)函數(shù)在框架中是通過(guò) 8 個(gè)類(lèi)似 Pow、Add、和 Tanh 等基本 OP 來(lái)實(shí)現(xiàn)的。也就是說(shuō)每進(jìn)行一次 gelu 運(yùn)算要調(diào)用 8 次基本 OP，同時(shí)底層也對(duì)應(yīng) 8 次 GPU kernel 的調(diào)用，每一次調(diào)用也意味著一次顯存讀寫(xiě)，先不說(shuō) kernel 計(jì)算耗時(shí)，光是顯存讀寫(xiě)就已經(jīng)是大量的開(kāi)銷(xiāo)。如何降低這部分開(kāi)銷(xiāo)？最直觀的方法就是減少調(diào)用，讓數(shù)據(jù)一直留在顯存甚至寄存器里被訪問(wèn)，即 OP 融合，一次調(diào)用就實(shí)現(xiàn)整個(gè)計(jì)算邏輯。
出于性能最大化的考慮，在 Faster Transformer 內(nèi)部，Nividia 將除矩陣乘法以外的所有 kernel 都進(jìn)行了盡可能的融合，單層 Transformer 的計(jì)算流程如下圖所示:

如圖所示，基于 OP 融合的思想，F(xiàn)aster Transformer 只用了 14 個(gè) kernel 就完成了原來(lái)將近 60 個(gè) kernel 的計(jì)算邏輯。這其中，8 個(gè) kernel 是通過(guò)調(diào)用 cuBLAS 接口計(jì)算矩陣乘法（黃色框），其余 6 個(gè)是自定義 kernel（藍(lán)色框）。
接下來(lái)筆者將沿調(diào)用鏈逐步介紹每個(gè) kernel 的優(yōu)化邏輯。

3 調(diào)用鏈

Faster Transformer v1.0 版本源碼地址如下，有興趣的讀者可以前往閱讀。

https://github.com/NVIDIA/FasterTransformer/tree/v1.0/fastertransformer

通讀源碼后筆者對(duì)調(diào)用關(guān)系梳理如下。

BertEncoderTransformer->forward()
    ->OpenMultiHeadAttention->forward()
        ->cublasGemmEx
        ->cublasGemmEx
        ->cublasGemmEx
        ->multiHeadAttr_nofuse_kernelLauncher
            ->add_QKV_bias  (kernel)
            ->cublasGemmStridedBatchedEx
            ->softmax_kernel    (kernel)
            ->cublasGemmStridedBatchedEx
            ->transpose (kernel)
    ->cublasGemmEx
    ->add_bias_input_layernorm_kernelLauncher   (kernel)
    ->cublasGemmEx
    ->add_bias_act_kernelLauncher   (kernel)
    ->cublasGemmEx
    ->add_bias_input_layernorm_kernelLauncher   (kernel)

從調(diào)用鏈也可以看出，總共 14 個(gè)步驟，與示意圖一一對(duì)應(yīng)。核心邏輯都在兩個(gè)類(lèi)中實(shí)現(xiàn)：BertEncoderTransformer 和 OpenMultiHeadAttention。

4 OpenMultiHeadAttention

OpenMultiHeadAttention 類(lèi)中有兩個(gè)重要的成員方法：構(gòu)造函數(shù)、forward 方法。其中構(gòu)造函數(shù)內(nèi)主要進(jìn)行一些參數(shù)初始化功能，設(shè)備內(nèi)存的申請(qǐng)和初始化也在該函數(shù)內(nèi)進(jìn)行。forward 方法內(nèi)主要是多頭注意力機(jī)制核心邏輯的具體實(shí)現(xiàn)。

4.1 cublasGemmEx for Q、K、V

forward 方法中首先就是對(duì)輸入的 3 個(gè) tensor 進(jìn)行線性變換，其實(shí)就是對(duì) 3 個(gè) tensor 分別進(jìn)行 Dense 層變換，我們知道 Dense 是包含一個(gè)矩陣乘法和一個(gè) add_bias 操作，這里只進(jìn)行矩陣乘法，add_bias 操作放在后面的 kernel 進(jìn)行。這里使用了 cuBLAS 接口計(jì)算矩陣乘法，具體代碼如下：

check_cuda_error(cublasGemmEx(param_.cublas_handle, 
    CUBLAS_OP_N, CUBLAS_OP_N, 
    n, m, k, 
    &alpha, 
    param_.attr_kernel_Q, AType_, n, 
    param_.from_tensor, BType_, k, 
    &beta, 
    query_buf_, CType_, n, 
    computeType_, 
    static_cast(cublasAlgo_[0])));

check_cuda_error(cublasGemmEx(param_.cublas_handle, 
    CUBLAS_OP_N, CUBLAS_OP_N,
    n, m, k, 
    &alpha, 
    param_.attr_kernel_K, AType_, n, 
    param_.to_tensor, BType_, k, 
    &beta, 
    key_buf_, CType_, n, 
    computeType_, 
    static_cast(cublasAlgo_[0])));

check_cuda_error(cublasGemmEx(param_.cublas_handle, 
    CUBLAS_OP_N, CUBLAS_OP_N, 
    n, m, k,
    &alpha,
    param_.attr_kernel_V, AType_, n, 
    param_.to_tensor, BType_, k, 
    &beta, 
    value_buf_, CType_, n, 
    computeType_, 
    static_cast(cublasAlgo_[0])));

這里僅僅是矩陣乘法 API 的調(diào)用，按文檔傳參即可，這里不展開(kāi)介紹，筆者計(jì)劃另開(kāi)一篇文章專(zhuān)門(mén)介紹這個(gè) API 的調(diào)用方法。

4.2 multiHeadAttr_nofuse_kernelLauncher

4.2.1 add_QKV_bias

上面說(shuō)過(guò) Dense 層包含矩陣乘法和 add_bias 操作，其中 add_bias 操作在核函數(shù) add_QKV_bias 中完成，源碼針對(duì)兩種數(shù)據(jù)類(lèi)型 fp16 和 fp32 分別提供了一個(gè) kernel，只是網(wǎng)絡(luò)結(jié)構(gòu)有所差異。
針對(duì) fp32，每個(gè) block 處理一個(gè) word，總共有 batch_size * seq_len * 3 個(gè) block，對(duì)于 Q、K、V 3 個(gè) tensor 而言，前 batch_size * seq_len 個(gè) block 處理 Q，中間 batch_size * seq_len 個(gè) block 處理 K，后 batch_size * seq_len 個(gè) block 處理 V。示意圖如下：

/**
 * @brief 
 * 
 * @tparam T                          OperationType
 * @param Q                           [batch_size, seq_len, head_num, size_per_head], query
 * @param bias_Q                      [head_num * size_per_head, ] length is the same as word's embedding dim
 * @param K                           [batch_size, seq_len, head_num, size_per_head], key
 * @param bias_K                      [head_num * size_per_head, ] length is the same as word's embedding dim
 * @param V                           [batch_size, seq_len, head_num, size_per_head], value
 * @param bias_V                      [head_num * size_per_head, ] length is the same as word's embedding dim
 * @param q_buf_                      [batch_size, head_num, seq_len, size_per_head], transpose query & add bias
 * @param k_buf_                      [batch_size, head_num, seq_len, size_per_head], transpose key & add bias
 * @param v_buf_                      [batch_size, head_num, seq_len, size_per_head], transpose value & add bias
 * @param batch_size
 * @param seq_len                     
 * @param head_num 
 * @param size_per_head 
 * @param word_per_block              1
 * @return __global__ 
 */
template
__global__
void add_QKV_bias(T* Q, const T* bias_Q, T* K, const T* bias_K, T* V, const T* bias_V, T* q_buf_, T* k_buf_, T* v_buf_, 
  const int batch_size, const int seq_len, const int head_num, const int size_per_head, const int word_per_block)
{

  T* data_ptr;
  T* buf_ptr;
  const T* bias_ptr;
  // word counts per batch
  int m = batch_size * seq_len;
  // word embedding dim
  int n = head_num * size_per_head;
  // 總共有3m個(gè)block，第一部分處理q，第二部分處理k，第三部分處理v，這里使用qkv_id區(qū)分處理哪個(gè)矩陣
  int qkv_id = blockIdx.x * word_per_block / m;
  // 矩陣偏移量
  int row_offset = (blockIdx.x * word_per_block % m) * n;

  if(qkv_id == 0)
  {
    data_ptr = Q + row_offset;
    buf_ptr = q_buf_;
    bias_ptr = bias_Q;
  }
  else if(qkv_id == 1)
  {
    data_ptr = K + row_offset;
    buf_ptr = k_buf_;
    bias_ptr = bias_K;
  }
  else
  {
    data_ptr = V + row_offset;
    buf_ptr = v_buf_;
    bias_ptr = bias_V;
  }

  int batch_id = (blockIdx.x * word_per_block % m) / seq_len;
  int head_id = threadIdx.x / size_per_head;
  int id_in_head = threadIdx.x % size_per_head;
  // word_id in seq, not data_index
  int word_start_id = (blockIdx.x * word_per_block) % seq_len;

  T bias = __ldg(&bias_ptr[threadIdx.x]);

  for(int i = word_start_id; i < word_start_id + word_per_block; ++i)
  {
    // add bias
    T tmp = data_ptr[threadIdx.x] + bias;
    // buf's shape: [bacth_size, head_num, seq_len, size_per_head]
    int target_id = batch_id * (seq_len * head_num * size_per_head) + head_id * seq_len * size_per_head + 
      i * size_per_head + id_in_head;

    buf_ptr[target_id] = tmp;
    data_ptr += n;
  }
}

核函數(shù)第一部分先根據(jù) block_id 確定當(dāng)前處理的 tensor 具體是 Q、K、V 中的哪一個(gè)，從而拿到輸入輸出變量的內(nèi)存地址。
第二部分求出 tensor 中對(duì)應(yīng)元素的索引，首先我們知道輸入輸出 tensor 是一個(gè)四維的 array，所以應(yīng)該有四個(gè)索引，按維度順序依次是 batch_id、word_start_id、head_id、id_in_head，有讀者看到這里可能會(huì)有疑問(wèn)：為什么要計(jì)算這些索引，前面計(jì)算了矩陣偏移量 row_offset，完全可以在 block 內(nèi)按 thread_id 索引就可以拿到對(duì)應(yīng)元素。原因是在 add_QKV_bias 核函數(shù)中計(jì)算邏輯不僅僅是 add，還有 transpose，熟悉 multiHeadAttention 的讀者都知道對(duì) Q、K、V 線性映射之后，緊接著就是一個(gè) transpose 操作，目的是把 embedding_dim 這個(gè)維度劃分成多個(gè)獨(dú)立的 head，每個(gè) head 后面單獨(dú)進(jìn)行 attention，所以要把 head 維度移到 seq_len 維度前面。換句話(huà)說(shuō)這里的 transpose 解決的是“多頭”的問(wèn)題，和 attention 無(wú)關(guān)。
理解了前面的邏輯，第三部分就比較簡(jiǎn)單了，先進(jìn)行 add 操作，然后將結(jié)果按照 [bacth_size, head_num, seq_len, size_per_head] 的維度順序?qū)懺谳敵?tensor 中，這里隱含了一個(gè) transpose，需要注意的是這個(gè) transpose 操作是輸出的 tensor 中元素存儲(chǔ)順序相對(duì)于輸入 tensor 而言的，并不是對(duì)輸入 tensor 做了變換。

針對(duì) fp16，每個(gè) block 同時(shí)處理 Q、K、V 上的同一個(gè) word，同一個(gè)線程先后處理 3 個(gè) word 上對(duì)應(yīng)元素的計(jì)算邏輯，實(shí)際計(jì)算中把 half 都轉(zhuǎn)成了 half2，使用標(biāo)準(zhǔn)庫(kù)中的函數(shù) __hadd2 運(yùn)算。網(wǎng)絡(luò)結(jié)構(gòu)如下：

從圖中可以看出，block_size 為 embedding_dim 的一半，這是因?yàn)橛昧?half2 這個(gè)數(shù)據(jù)結(jié)構(gòu)，實(shí)際上每個(gè)線程處理了 2 個(gè)元素，所以線程數(shù)量縮減一半。核函數(shù)內(nèi)部邏輯分為 2 個(gè)部分：1、求出 tensor 中對(duì)應(yīng)元素的索引。2、一次對(duì) Q、K、V 進(jìn)行 add 和 transpose 操作。

template <>
__global__
void add_QKV_bias(__half* Q, const __half* bias_Q, __half* K, const __half* bias_K, __half* V, const __half* bias_V, 
  __half* q_buf_, __half* k_buf_, __half* v_buf_, 
  const int batch_size, const int seq_len, const int head_num, const int size_per_head, const int word_per_block)
{
  int tid = blockIdx.x * blockDim.x + threadIdx.x;
  int batch_id = tid / (head_num * seq_len * size_per_head);
  int seq_id = (tid % (head_num * seq_len * size_per_head)) / (head_num * size_per_head);
  int head_id = (tid % (head_num * size_per_head)) / size_per_head;
  // id in head
  int id = tid % size_per_head;
  // from [batch_size, seq_len, head_num, size_per_head] tanspose to [bacth_size, head_num, seq_len, size_per_head]
  int target_id = target_index(batch_id, seq_id, head_id, id, batch_size, seq_len, head_num, size_per_head);

  int bias_id = threadIdx.x;

  half2* src_ptr = (half2*)Q;
  half2* dst_ptr = (half2*)q_buf_;
  const half2* bias_ptr = (const half2*)bias_Q;
  dst_ptr[target_id] = __hadd2(src_ptr[tid],  __ldg(&bias_ptr[bias_id]));

  src_ptr = (half2*)K;
  dst_ptr = (half2*)k_buf_;
  bias_ptr = (const half2*)bias_K;
  dst_ptr[target_id] = __hadd2(src_ptr[tid],  __ldg(&bias_ptr[bias_id]));

  src_ptr = (half2*)V;
  dst_ptr = (half2*)v_buf_;
  bias_ptr = (const half2*)bias_V;
  dst_ptr[target_id] = __hadd2(src_ptr[tid],  __ldg(&bias_ptr[bias_id]));
}

4.2.2 計(jì)算 attention scores

先來(lái)看一下 attention 的計(jì)算公式，定義如下：

其中，，也就是說(shuō)這一步要解決的是一個(gè)矩陣計(jì)算，用 tensorflow 代碼表示如下：

scores = tf.matmul(query, key, transpose_b=True)

針對(duì)矩陣運(yùn)算，源碼中直接調(diào)用了 cuBLAS API，具體代碼如下：

DataType_ alpha = (DataType_)1.0f, beta = (DataType_)0.0f;
// 計(jì)算 q * k^T
check_cuda_error(cublasGemmStridedBatchedEx(cublas_handle,
    CUBLAS_OP_T, CUBLAS_OP_N,
    seq_len, seq_len, size_per_head,
    &alpha,
    k_buf_, AType_, size_per_head, seq_len * size_per_head,
    q_buf_, BType_, size_per_head, seq_len * size_per_head,
    &beta,
    qk_buf_, CType_, seq_len, seq_len * seq_len,
    batch_size * head_num,
    computeType_,
    static_cast(cublasAlgo_[1])));

不熟悉 attention 的讀者可能會(huì)問(wèn) attention scores 的具體含義是什么，筆者在早期的文章中有過(guò)介紹，其實(shí)就是兩個(gè)矩陣的詞向量?jī)蓛上喑耍蛄肯喑擞惺裁春x？相似度，這個(gè)分?jǐn)?shù)就代表 Q、K 的相似度。有興趣的讀者可以移步筆者之前的文章詳細(xì)了解。（https://mp.weixin.qq.com/s/0zen3ItKmDLt5rTUbF37Mg）

4.2.3 softmax_kernel

拿到 Q、K 的相似度之后，直觀上只要右乘一個(gè) V 就可以得到 attention out，其含義就是一個(gè)加權(quán)平均的概念，既然要加權(quán)平均，必然要對(duì)權(quán)值進(jìn)行歸一化處理，這里的 softmax 就是這個(gè)作用。關(guān)于 softmax 核函數(shù)的實(shí)現(xiàn)方法筆者在前兩篇文章也有介紹，OneFlow 官方給出了更為高效的實(shí)現(xiàn)方式，其高效的原因主要在訪存帶寬處理上，有興趣的讀者可以移步。【CUDA編程】OneFlow Softmax 算子源碼解讀之WarpSoftmax，【CUDA編程】OneFlow Softmax算子源碼解讀之BlockSoftmax

// 計(jì)算softmax(qk)
if(seq_len <= 32)
    block.x = 32;
else if(seq_len > 32 && seq_len <= 64)
    block.x = 64;
else if(seq_len > 64 && seq_len <= 128)
    block.x = 128;
else if(seq_len > 128 && seq_len <= 256)
    block.x = 256;
else if(seq_len > 256 && seq_len <= 512)
    block.x = 512;
else
    block.x = 1024;

if(batch_size * head_num <= 120)
{
    grid.x = batch_size * head_num * seq_len;
    softmax_kernel_v2<<>>(qk_buf_, attr_mask, batch_size, head_num, seq_len, scaler); 
}
else
{
    grid.x = batch_size * head_num;
    softmax_kernel<<>>(qk_buf_, attr_mask, batch_size, head_num, seq_len, scaler); 
}

源碼中核函數(shù)的 block_size 是根據(jù) seq_len 確定的，取大于 seq_len 且為 32 的 的最小值。
另外在調(diào)用 softmax kernel 之前，會(huì)根據(jù) batch_size * head_num 選擇不同的 softmax kernel，主要是為了保證在大 batch 的情況下的計(jì)算效率，這里以 120 為閾值，應(yīng)該是作者的經(jīng)驗(yàn)數(shù)值。這里作者給出了 2 個(gè) softmax kernel 的實(shí)現(xiàn)。
當(dāng) batch_size * head_num > 120 時(shí)，此時(shí) batch 內(nèi)元素較多，grid_size 取 batch_size * head_num，這時(shí)一個(gè)線程內(nèi)處理一個(gè) seq_len 的數(shù)據(jù)。

/**
 * @brief 
 * 
 * @tparam T 
 * @param qk_buf_                 [batch_size, head_num, seq_len, seq_len]
 * @param attr_mask               [batch_size, seq_len, seq_len]
 * @param batch_size 
 * @param head_num 
 * @param seq_len 
 * @param scaler                  縮放因子
 * @return __global__ 
 */
template 
__global__
void softmax_kernel(T* qk_buf_, const T* attr_mask, const int batch_size, const int head_num, const int seq_len, 
  const T scaler)
{
    // grid_size = batch_size * head_num
    int batch_id = blockIdx.x / head_num;
    // batch偏移量
    int qk_offset = blockIdx.x * seq_len * seq_len;
    int mask_offset = batch_id * seq_len * seq_len;

    __shared__ float s_sum, s_max;

    // 每次處理一個(gè)seq_len的數(shù)據(jù)
    for(int i = 0; i < seq_len; ++i)
    {
      float qk = threadIdx.x < seq_len ? (float)qk_buf_[threadIdx.x + qk_offset] : 0.0f;
      float mask_val = threadIdx.x < seq_len ? (float)attr_mask[threadIdx.x + mask_offset] : 0.0f;
      // 對(duì)于某些padded的word，給一個(gè)很小的值使其近似達(dá)到不參與運(yùn)算的目的
      mask_val = (1.0f - mask_val) * -10000.0f;

      float tmp = threadIdx.x < seq_len ? (float)(qk * (float)scaler + mask_val): -1e20f;

      float max_val = blockReduceMax(tmp);

      if(threadIdx.x == 0)
        s_max = max_val;
      __syncthreads();

      qk = threadIdx.x < seq_len ? __expf(tmp - s_max) : 0.0f;

      float sum_val = blockReduceSum(qk);

      if(threadIdx.x == 0)
      {
        s_sum = sum_val + 1e-6f;
      }
      __syncthreads();

      if(threadIdx.x < seq_len)
        qk_buf_[threadIdx.x + qk_offset] = (T)(qk / s_sum);

      qk_offset += seq_len;
      mask_offset += seq_len;
    }
}

核函數(shù)內(nèi)首先計(jì)算了每個(gè)元素偏移量，對(duì)于輸入 tensor 而言，每個(gè) block 處理 seq_len * seq_len 個(gè)數(shù)據(jù)，所以 block 內(nèi)元素偏移量為 blockIdx.x * seq_len * seq_len，而對(duì)于 mask 矩陣而言，其維度為 [batch_size, seq_len, seq_len]，跟 head_num 無(wú)關(guān)，所以其偏移量為 batch_id * seq_len * seq_len。
接下來(lái)是一層循環(huán)，對(duì)于 seq_len * seq_len 矩陣而言，每個(gè)線程處理當(dāng)前 thread_id 列的元素，每輪循環(huán)結(jié)束，處理該列下一行的元素。在每一輪循環(huán)中，所有的線程一起處理一行數(shù)據(jù)，首先拿到數(shù)據(jù) qk 以及 mask_val。如果 mask_val 為 0，則給 mask_val 賦一個(gè)很小的值最后加在 qk 上使 qk 值很小，以致最終這個(gè) softmax 分量趨于 0；如果 mask_val 為 1，則 mask 不干預(yù)后續(xù)計(jì)算。每個(gè)線程拿到處理后的 qk 值即 tmp 后，進(jìn)行一次塊內(nèi)規(guī)約，即可求出當(dāng)前行的最大值 max_val，然后為了避免指數(shù)運(yùn)算導(dǎo)致數(shù)值溢出，讓 tmp 減去 max_val 并求其指數(shù)值賦給 qk ，然后對(duì) qk 再一次塊內(nèi)規(guī)約求出當(dāng)前行的和 s_sum，最后讓 qk 除以 s_sum 即可得到 softmax 值。核函數(shù)內(nèi)要注意在兩次塊內(nèi)規(guī)約后一定要進(jìn)行一次塊內(nèi)同步，否則可能計(jì)算錯(cuò)誤。

當(dāng) batch_size * head_num <= 120 時(shí)，此時(shí) batch 較小，grid_size 取 batch_size * head_num * seq_len，這時(shí)一個(gè)線程塊內(nèi)處理一行數(shù)據(jù)，每個(gè)線程內(nèi)只處理一個(gè)的數(shù)據(jù)。

template 
__global__
void softmax_kernel_v2(T* qk_buf_, const T* attr_mask, const int batch_size, const int head_num, 
  const int seq_len, const float scaler)
{
    int batch_id = blockIdx.x / head_num / seq_len;
    int seq_id = blockIdx.x % seq_len;
    int qk_offset = blockIdx.x * seq_len;
    int mask_offset = batch_id * seq_len * seq_len + seq_id * seq_len;

    __shared__ float s_sum, s_max;

    float qk = threadIdx.x < seq_len ? (float)qk_buf_[threadIdx.x + qk_offset] : 0.0f;
    float mask_val = threadIdx.x < seq_len ? (float)attr_mask[threadIdx.x + mask_offset] : 0.0f;
      
    mask_val = (1.0f - mask_val) * -10000.0f;

    float tmp = threadIdx.x < seq_len ? (float)(qk * (float)scaler + mask_val) : -1e20f;
    float max_val = blockReduceMax(tmp);
    if(threadIdx.x == 0)
      s_max = max_val;
    __syncthreads();

    float qk_tmp = threadIdx.x < seq_len ? __expf((float)(tmp - s_max)) : 0.0f;
    float sum_val = blockReduceSum(qk_tmp);

    if(threadIdx.x == 0)
    {
      s_sum = sum_val + 1e-6f;
    }
    __syncthreads();

    if(threadIdx.x < seq_len)
      qk_buf_[threadIdx.x + qk_offset] = (T)(qk_tmp / s_sum);
}

這種情況下不涉及循環(huán)處理，計(jì)算邏輯與前面 softmax_kernel 循環(huán)體內(nèi)部計(jì)算邏輯相同，不再贅述。

4.2.4 計(jì)算多頭 attention out

這一步的意思就是使用 softmax 后的相似度矩陣右乘一個(gè) V，得到多頭注意力輸出，注意這時(shí)候輸出 tensor 的維度為 [batch_size, head_num, seq_len, size_per_head]。源碼中直接調(diào)用了 cuBLAS API，具體代碼如下：

// 計(jì)算qk * v
check_cuda_error(cublasGemmStridedBatchedEx(cublas_handle,
    CUBLAS_OP_N, CUBLAS_OP_N,
    size_per_head, seq_len, seq_len,
    &alpha,
    v_buf_, AType_, size_per_head, seq_len * size_per_head,
    qk_buf_, BType_, seq_len, seq_len * seq_len,
    &beta,
    transpose_dst_, CType_, size_per_head, seq_len * size_per_head,
    batch_size * head_num,
    computeType_,
    static_cast(cublasAlgo_[2])));

4.2.5 transpose

前面說(shuō)過(guò)，多頭 attention out 的維度為 [batch_size, head_num, seq_len, size_per_head]，此時(shí)這些 head 已經(jīng)完成使命了，通過(guò)獨(dú)立的 head_num 組 attention 參數(shù)計(jì)算出了 attention out，最后需要做的就是把這 head_num 組 attention out 拼接起來(lái)，體現(xiàn)在 tensor 上就是做一次 transpose，將維度變?yōu)?[batch_size, seq_len, head_num, size_per_head]。源碼針對(duì) fp16 和 fp32 分別提供了一個(gè)核函數(shù) transpose，計(jì)算邏輯和 add_QKV_bias 中 transpose 計(jì)算邏輯相同，索引按順序乘即可。具體代碼如下：

template
__global__
void transpose(T* src, T* dst, const int batch_size, const int seq_len, const int head_num, const int size_per_head)
{
  int batch_id = blockIdx.x / (head_num * seq_len);
  int seq_id = blockIdx.x % seq_len;
  int head_id = (blockIdx.x % (head_num * seq_len))/ seq_len;
  dst[batch_id * (head_num * seq_len * size_per_head) + seq_id * head_num * size_per_head
    + head_id * size_per_head + threadIdx.x] = src[blockIdx.x * size_per_head + threadIdx.x];
}

template<>
  __global__
void transpose(__half* src, __half* dst,
    const int batch_size, const int seq_len, const int head_num, const int size_per_head)
{
  int tid = blockIdx.x * blockDim.x + threadIdx.x;

  int batch_id = tid / (head_num * seq_len * size_per_head);
  int head_id = (tid % (head_num * seq_len * size_per_head)) / (seq_len * size_per_head);
  int seq_id = (tid % (seq_len * size_per_head)) / size_per_head;
  int id = tid % size_per_head;

  int target_id = target_index(batch_id, head_id, seq_id, id, batch_size, head_num, seq_len, size_per_head);
  half2* src_ptr = (half2*)src;
  half2* dst_ptr = (half2*)dst;

  dst_ptr[target_id] = src_ptr[tid];
}

5 BertEncoderTransformer

BertEncoderTransformer 類(lèi)中有兩個(gè)重要的成員方法：構(gòu)造函數(shù)、forward 方法。其中構(gòu)造函數(shù)內(nèi)主要進(jìn)行一些參數(shù)初始化功能，設(shè)備內(nèi)存的申請(qǐng)和初始化也在該函數(shù)內(nèi)進(jìn)行。forward 方法內(nèi)主要是核心邏輯的實(shí)現(xiàn)。

5.1 attention forward

根據(jù)調(diào)用鏈可知，BertEncoderTransformer->forward() 中第一步就是 attention_->forward()，其中 attention_ 對(duì)象在構(gòu)造函數(shù)中被定義，attention_->forward() 執(zhí)行的就是第 4 節(jié)的內(nèi)容。

5.2 對(duì) attention out 做線性變換

根據(jù)流程圖和調(diào)用鏈可知，這一步是對(duì)多頭注意力的輸出 tensor 做一層線性變換，右乘一個(gè)參數(shù)矩陣，其實(shí)就是一個(gè)不加激活函數(shù)的 Dense 層，分為矩陣乘法和 add bias 兩個(gè)操作步驟，這里調(diào)用了 cuBLAS API 實(shí)現(xiàn)矩陣乘法。

DataType_ alpha = (DataType_)1.0f;
DataType_ beta = (DataType_)0.0f;
int m = batch_size_ * from_seq_len_;
int k = head_num_ * size_per_head_;
int n = k;

check_cuda_error(cublasGemmEx(param_.cublas_handle, 
    CUBLAS_OP_N, CUBLAS_OP_N,
    n, m, k, 
    &alpha, 
    param_.attr_output_kernel, AType_, n, 
    attr_out_buf_, BType_, k, 
    &beta, 
    attr_matmul_buf_, CType_, n, 
    computeType_, 
    static_cast(cublasAlgo_[0])));

5.3 add_bias_input_layernorm_kernel

從核函數(shù)名字可以看出，這個(gè)核函數(shù)實(shí)現(xiàn)了 3 個(gè)操作：add bias、add input、layernomalization。其中 add bias 是完成上一步線性變換未完成的加偏置工作，add input 是 transformer 模型中的殘差結(jié)構(gòu)，layernomalization 則是層歸一化操作。綜合起來(lái)這個(gè)核函數(shù)的作用是：對(duì)線性變換后的 attention out 加偏置，然后加上原始輸入 tensor 組成一個(gè)殘差結(jié)構(gòu)，最后進(jìn)行一次層歸一化變換。源碼中針對(duì) fp16 和 fp32 分別提供了一個(gè)核函數(shù)實(shí)現(xiàn)，計(jì)算邏輯都一樣，這里只以 fp32 為例介紹。

/**
 * @brief                       grid_size = m, block_size = n
 * 
 * @tparam T 
 * @param out                   [batch_size, sql_len, latent_dim]
 * @param input                 [batch_size, sql_len, latent_dim]
 * @param bias                  [latent_dim,]
 * @param gamma 
 * @param beta 
 * @param m                     batch_size * seq_len
 * @param n                     latent_dim
 * @return __global__ 
 */
template 
__global__ 
void add_bias_input_layernorm(T* out, const T* input, const T* bias, const T* gamma, const T* beta, int m, int n)
{
  int tid = threadIdx.x;

  __shared__ float s_mean;
  __shared__ float s_variance;
  float mean =  0.0f;
  float variance = 0.0f;

  float local_out = 0.0f;
  // add，一個(gè)block處理一行
  for(int i = tid; i < n; i += blockDim.x)
    local_out += (float)(out[blockIdx.x * n + i] + input[blockIdx.x * n + i] + __ldg(&bias[i]));
  // mean_i = sum(x_i[j] for j in range(k)) / k
  mean = blockReduceSum(local_out);
  if(threadIdx.x == 0)
    s_mean = mean / n;
  __syncthreads();
  // var_i = sum((x_i[j] - mean_i) ** 2 for j in range(k)) / k + epsilon
  variance = blockReduceSum((local_out - s_mean) * (local_out - s_mean));
  if(threadIdx.x == 0)
    s_variance = variance / n + 1e-6f;
  __syncthreads();
  // x_i_normalized = (x_i - mean_i) / sqrt(var_i)
  // output_i = x_i_normalized * gamma + beta
  for(int i = tid; i < n; i += blockDim.x)
    out[blockIdx.x * n + i] = 
    (T)(((local_out - s_mean) * rsqrtf(s_variance)) * (float)(__ldg(&gamma[i])) + (float)(__ldg(&beta[i])));
}

如示意圖所示，核函數(shù)中每個(gè) block 處理一行數(shù)據(jù)，共 latent_dim = head_num * size_per_head 個(gè)元素，核函數(shù)中首先計(jì)算了 add bias、add input 兩個(gè)操作，并將計(jì)算結(jié)果存儲(chǔ)在寄存器變量 local_out 中。
接下來(lái)就是標(biāo)準(zhǔn)的 layerNormalization 操作，我們先來(lái)看一下 layerNormalization 的操作步驟，參照一下 tensorflow 框架 API 文檔。

For each sample x_i in inputs with k features, we compute the mean and variance of the sample:
mean_i = sum(x_i[j] for j in range(k)) / k
var_i = sum((x_i[j] - mean_i) ** 2 for j in range(k)) / k

and then compute a normalized x_i_normalized, including a small factor epsilon for numerical stability.
x_i_normalized = (x_i - mean_i) / sqrt(var_i + epsilon)

And finally x_i_normalized is linearly transformed by gamma and beta, which are learned parameters:
output_i = x_i_normalized * gamma + beta
gamma and beta will span the axes of inputs specified in axis, and this part of the inputs' shape must be fully defined.

具體地，第一步計(jì)算均值和方差，核函數(shù)中使用塊內(nèi)規(guī)約計(jì)算出均值 s_mean 存儲(chǔ)在共享內(nèi)存中，所有塊內(nèi)線程都可以訪問(wèn)。然后根據(jù) s_mean 和線程內(nèi)的 local_out 以及 epsilon 系數(shù)再進(jìn)行一次塊內(nèi)規(guī)約計(jì)算出方差 s_variance 存儲(chǔ)在共享內(nèi)存中。
第二步進(jìn)行歸一化和線性變換，對(duì)應(yīng) tensorflow API 的二、三步，直接計(jì)算即可，沒(méi)有其他技巧，公式如下：

5.4 FeedForward 結(jié)構(gòu)

根據(jù) Transformer 模型結(jié)構(gòu)，多頭注意力之后為了增強(qiáng)表達(dá)能力，加了一個(gè) FeedForward 層，該結(jié)構(gòu)內(nèi)部就是兩個(gè) Dense 層，第一層 Dense 中使用了激活函數(shù)，第二層沒(méi)有激活函數(shù)。所以 FeedForward 層中包含了 5 個(gè)操作：矩陣乘法、add bias、activation、矩陣乘法、add bias。

5.4.1 attention out * inter kernel

FeedForward 層第一次線性變換會(huì)擴(kuò)展 tensor 的最后一個(gè)維度的長(zhǎng)度，源碼中將 latent_dim（也就是 n）擴(kuò)展為原來(lái)的 4 倍，所以這里的 inter kernel 的形狀為 [latent_dim, 4 * latent_dim]，矩陣運(yùn)算后的輸出 tensor 形狀為 [batch_size, seq_len, 4 * latent_dim]。

n *= 4;
check_cuda_error(cublasGemmEx(param_.cublas_handle, 
    CUBLAS_OP_N, CUBLAS_OP_N,
    n, m, k, 
    &alpha, 
    param_.inter_kernel, AType_, n, 
    attr_matmul_buf_, BType_, k, 
    &beta, 
    inter_matmul_buf_, CType_, n, 
    computeType_, 
    static_cast(cublasAlgo_[1])));

5.4.2 add_bias_act_kernel

顧名思義，add_bias_act_kernel 核函數(shù)包含了 add bias 和 activation 兩個(gè)操作。源碼中 block_size = n / 4 實(shí)際就是 latent_dim，為什么不直接取上一次運(yùn)算后的矩陣寬度 n = 4 * latent_dim 呢？這里是希望一行元素（4 * latent_dim）能在一個(gè) block 內(nèi)處理，如果 block_size 直接取 n = 4 * latent_dim，可能超過(guò) 1024，因此還取 latent_dim，線程內(nèi)循環(huán) 4 次處理即可。同樣，源碼中針對(duì) grid_size 也取了 m / 4，在線程中通過(guò)循環(huán)每次跨 m / 4 步長(zhǎng)處理 4 行數(shù)據(jù)。

template 
__global__ 
void add_bias_act(T* out, const T* bias, int m, int n)
{
  T val, reg_bias;

  int row_id = blockIdx.x;
  int ite = n / blockDim.x;
  int tid = threadIdx.x;

  for(int i = 0; i < ite; ++i)
  {
    reg_bias = __ldg(&bias[i * blockDim.x + tid]);
    row_id = blockIdx.x;

    while(row_id < m){
      val = out[tid + i * blockDim.x + row_id * n]+ reg_bias;
      out[tid + i * blockDim.x + row_id * n] = gelu(val);
      row_id += gridDim.x;
    }
  }
}

核函數(shù)中先對(duì)列進(jìn)行循環(huán)，ite = 4，從全局內(nèi)存讀出當(dāng)前列的 bias，然后針對(duì)行進(jìn)行循環(huán)，步長(zhǎng)為 m / 4，循環(huán)體內(nèi)部對(duì)當(dāng)前行當(dāng)前列的元素進(jìn)行 add bias 和 gelu 操作，這里gelu 操作是一個(gè)簡(jiǎn)單的 element-wise 操作，比較簡(jiǎn)單不再介紹。
筆者點(diǎn)評(píng)：這里筆者私以為沒(méi)有必要 grid_size 也取 m / 4，cuda 本身對(duì)線程塊的數(shù)量沒(méi)有限制，完全可以直接取 m，每次每個(gè)線程只處理一行數(shù)據(jù)，一方面可以增加并行程度，另一方面代碼可閱讀性也更好。筆者給出代碼如下，親測(cè)可用。

dim3 grid(m);
dim3 block(n / 4);
assert(block.x <= 1024);
add_bias_act_v2<<>>(out, bias, m, n);

template 
__global__ 
void add_bias_act_v2(T* out, const T* bias, int m, int n) {
  T val, reg_bias;

  int row_id = blockIdx.x;
  int ite = n / blockDim.x;
  int tid = threadIdx.x;

  for(int i = 0; i < ite; ++i) {
    reg_bias = __ldg(&bias[i * blockDim.x + tid]);
    val = out[tid + i * blockDim.x + row_id * n]+ reg_bias;
    out[tid + i * blockDim.x + row_id * n] = gelu(val);
  }
}

5.4.3 inter out * out kernel

FeedForward 層第二次線性變換將 tensor 的最后一個(gè)維度的長(zhǎng)度轉(zhuǎn)換為原始大小，源碼中將 n 重新賦值為 latent_dim，所以這里的 out kernel 的形狀為 [4 * latent_dim, latent_dim]，矩陣運(yùn)算后的輸出 tensor 形狀為 [batch_size, seq_len, latent_dim]。

n = k;
k *= 4;
check_cuda_error(cublasGemmEx(param_.cublas_handle, 
    CUBLAS_OP_N, CUBLAS_OP_N,
    n, m, k, 
    &alpha, 
    param_.output_kernel, AType_, n, 
    inter_matmul_buf_, BType_, k, 
    &beta, 
    param_.transformer_out, CType_, n, 
    computeType_, 
    static_cast(cublasAlgo_[2])));

5.5 add_bias_input_layernorm_kernel

這個(gè)核函數(shù)的計(jì)算邏輯在 5.3 中已經(jīng)介紹過(guò)了，包含加偏置項(xiàng)、殘差結(jié)構(gòu)、層歸一化三個(gè)操作，不再贅述。

6 小結(jié)

至此，Transformer encoder 前向計(jì)算的 14 個(gè)操作優(yōu)化技巧已介紹完畢。總結(jié)如下：

針對(duì)半精度 fp16 的優(yōu)化方面。首先，在 kernel 的實(shí)現(xiàn)中，將輸入的 half 指針轉(zhuǎn)成 half2 類(lèi)型，并使用了 half2 相關(guān)的數(shù)學(xué)函數(shù)。這樣不僅僅可以達(dá)到 2 倍于 half 的訪存帶寬和計(jì)算吞吐，還可以極大地減少指令的發(fā)射數(shù)量。其次，在 softmax 以及 layerNormalization 的操作中，為防止求和溢出，將數(shù)據(jù)以 half2 的形式讀入后，會(huì)轉(zhuǎn)成 float2 類(lèi)型，來(lái)做求和計(jì)算，這里就非常細(xì)致了，盡可能地保障了較高精度，值得學(xué)習(xí)借鑒。

針對(duì)訪存帶寬方面，筆者以為除 fp16 以外其它數(shù)據(jù)類(lèi)型也可以進(jìn)一步優(yōu)化，比如可以自定義 pack 類(lèi)型進(jìn)行合并讀寫(xiě)，盡量把帶寬打滿(mǎn)。

針對(duì)線程網(wǎng)絡(luò)結(jié)構(gòu)方面，源碼中基本使用一個(gè) block 處理一行數(shù)據(jù)的模式進(jìn)行設(shè)計(jì)，這里筆者私以為針對(duì) seq_len 和 latent_dim 已知比較小的情況下（不超過(guò)1024），完全可以一個(gè)線程束處理一行數(shù)據(jù)，束內(nèi)同步的開(kāi)銷(xiāo)遠(yuǎn)小于塊內(nèi)同步。當(dāng)然，這個(gè)要求確實(shí)有些苛刻了。

源碼中提供了一個(gè)塊內(nèi)規(guī)約的代碼，思路非常好，值得初學(xué) cuda 的讀者品讀。

審核編輯：湯梓紅