从FasterTransformer源码解读开始了解大模型（2.2）代码通读03

从FasterTransformer源码解读开始了解大模型（2.2）代码解读03-forward函数

写在前面的话

本篇的内容继续解读forward函数，从650行开始进行解读

零、输出Context_embeddings和context_cum_log_probs的参数和逻辑

从653行开始，会从输入的请求tensors中读取一个配置，如果请求中配置了is_return_context_embeddings参数并设置为true时，则会在返回参数中增加一个context_embeddings的tensor，这个tensor中包含的数据是输入经过了ContextDecoder过程的所有的层之后的logits，并对其进行求和。可能有些类似于强化学习（RLHF）之类的场景会用到这里的输出，所以在这里做了一层准备。

跳转到1093行，可以看见，这里调用了invokeSumLengthDimension这个kennels，来将context_decoder_output_buf这块buf中的显存数据拷贝到输出tensor的context_embeddings中，可以简单看一下这个kernel的实现，在src/fastertransformer/kernels/gpt_kernels.cu中

template<typename T>
void invokeSumLengthDimension(float*       out_buf,
                              const T*     in_buf,
                              const size_t batch_size,
                              const size_t input_length,
                              const size_t hidden_dim,
                              cudaStream_t stream)
{
    dim3 gridSize(batch_size);
    dim3 blockSize(256);

    sum_length_dimension<<<gridSize, blockSize, 0, stream>>>(out_buf, in_buf, batch_size, input_length, hidden_dim);
}


template<typename T>
__global__ void sum_length_dimension(
    float* out_buf, const T* in_buf, const size_t batch_size, const size_t input_length, const size_t hidden_dim)
{
    const int bidx = blockIdx.x;

    for (int hidx = threadIdx.x; hidx < hidden_dim; hidx += blockDim.x) {
        float accum = 0.0f;
        for (int step = 0; step < input_length; step++) {
            accum += static_cast<float>(in_buf[(bidx * input_length + step) * hidden_dim + hidx]);
        }
        out_buf[bidx * hidden_dim + hidx] = accum;
    }
}

从kernel中可以看出，这个kernel任务划分为按照batch_size维度进行了grid task分配，并为每个任务划分了256个线程。在kernel内部则是将每个batch内所有输入的logits按照输入length维度进行了累加，并拷贝到输出buf中。

类似的一个设置和处理环节，在783行还处理了is_return_context_cum_log_probs，这里会将contextDecoder完成之后的输出进行cum log计算。其中主要的处理逻辑函数是403行的computeContextCumLogProbs函数，进入这个函数后，在425到502行的处理逻辑是，首先将context_decoder的输出进