从FasterTransformer源码解读开始了解大模型(2.1)代码通读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的输出进行layernorm,然后使用cublas矩阵乘,将词表维度的logits计算出来。在此之后,在504行,则会使用invokeLogProbFromLogits这个kernel.

template<typename T>
void invokeLogProbFromLogits(float*       cum_log_probs,const T*     logits,const int*   input_ids,const int*   input_lengths,const size_t max_input_length,const size_t batch_size,const size_t vocab_size,const size_t vocab_size_padded,void*        workspace,const size_t workspace_size,cudaStream_t stream,const bool   batch_first)
{// A batched version of log prob computation.//// cum_log_probs: [batch_size]// logits: [max_input_length, batch_size, vocab_size] or [batch_size, max_input_length, vocab_size]// input_ids: [max_input_length, batch_size] or [max_input_length, batch_size]// input_lengths: [batch_size]// workspace: workspace buffer of size at least sizeof(float) * max_input_length * batch_size.FT_LOG_DEBUG(__PRETTY_FUNCTION__);// block_size should be multiple of 32 to use warpReduceMax.const int block_size = vocab_size < 1024 ? (vocab_size + 31) / 32 * 32 : 1024;assert(block_size % 32 == 0);assert(workspace != nullptr && workspace_size >= sizeof(float) * max_input_length * batch_size);assert(vocab_size <= vocab_size_padded);float* log_probs = reinterpret_cast<float*>(workspace);int    gx        = batch_first ? batch_size : max_input_length - 1;int    gy        = batch_first ? max_input_length - 1 : batch_size;dim3   grid(gx, gy);log_probs_kernel<T><<<grid, block_size, 0, stream>>>(log_probs,logits,input_ids,input_lengths,max_input_length,batch_size,vocab_size,vocab_size_padded,batch_first);accumulate_log_probs<<<batch_size, block_size, 0, stream>>>(cum_log_probs, log_probs, input_lengths, max_input_length, batch_size, batch_first);
}__global__ void accumulate_log_probs(float*       cum_log_probs,const float* log_probs,const int*   lengths,const size_t max_input_length,const size_t batch_size,const bool   batch_first)
{// Accumulate the log probability along with the sequence dimension.//   cum_log_probs[j] = sum_i log(softmax(logits))[ids[i,j]]//// cum_log_probs: [batch_size], cumulative log probability// log_probs: [max_length - 1, batch_size] or [batch_size, max_length - 1],//   log probability of each token// lengths: [batch_size], sequence lengths// batch_size: [1], batch_size. in case of beam > 1, batch x beam.int bidx = blockIdx.x;   // batch dimint tidx = threadIdx.x;  // step dimif (bidx < batch_size) {int length = lengths[bidx];// reposition logits to data for the current batch.log_probs += batch_first ? bidx * (max_input_length - 1) : bidx;int   stride      = batch_first ? 1 : batch_size;  // stride along with seq dim.float local_accum = 0.0f;for (int step = tidx; step < length - 1; step += blockDim.x) {local_accum += static_cast<float>(log_probs[step * stride]);}float accum = blockDim.x <= 32 ? warpReduceSum(local_accum) : blockReduceSum<float>(local_accum);if (tidx == 0) {cum_log_probs[bidx] = accum;}}
}

在任务划分阶段,grid的x维度是batch维度,而y维度则是输入长度,进入kernel后,则是按照batch维度对任务的log_probs进行了ReduceSum,并写入到cum_log_probs中。

一、跳过prompt和prefix阶段

回到653行,接下来是ft中处理其自身特殊的prompt和prefix的逻辑,但我们的源码解读可以先跳过这一段。之所以要跳过这一段是在当前主要的大模型处理逻辑中,并不会需要在推理引擎这一层过多地关注prompt和prefix的逻辑。在对话的大模型agent中,对于第n次对话的输入/输出,其真正要进行forward的输入,往往是由前n-1轮模型的输出和用户的输入共同组成的。从推理引擎的角度来看,prefill阶段在真正端到端时延的占比并不算高,所以专门设计一个处理prompt的逻辑(还需要常驻的显存空间)多少显得有些得不偿失了

二、正常的逻辑起始和输入展开

让我们直接来到865行,考虑以下一种场景来简化我们的源码解读的结构:beam_width=1,tp=1,pp=1,use_shared_contexts不启用,也就是说,没有beam search,在单机单卡上进行一次简单的,不共享contexts的推理服务的进行。

从866行开始,由于是计算起始,所以先对一些计算中的buff进行清空。在870行,由于不进行beam search,所以也不需要对beam search用到的buf进行清空。

在877到887行,如果词表大小不对齐的话,需要将词表计算的权重进行对齐拷贝。

在889到905行,不使用shared context的话,这一段的逻辑也不需要进行处理。在914行处理prompt的逻辑也可以跳过,直接进入955行的处理逻辑。在955行,由于我们的输入是带batch的,所以需要将batch进行展开(tiled),这里使用的是kernel invokeTileGptPromptInputs

void invokeTileGptPromptInputs(int*         tiled_input_ids,int*         tiled_input_lengths,int*         tiled_prompt_lengths,const int*   input_ids,const int*   input_lengths,const int*   prefix_prompt_lengths,const int    batch_size,const int    beam_width,const int    max_input_length,cudaStream_t stream)
{dim3 grid(batch_size, beam_width);dim3 block(min(1024, max_input_length));if (prefix_prompt_lengths != nullptr) {tileGptPromptInputs<true><<<grid, block, 0, stream>>>(tiled_input_ids,tiled_input_lengths,tiled_prompt_lengths,input_ids,input_lengths,prefix_prompt_lengths,max_input_length);}else {tileGptPromptInputs<false><<<grid, block, 0, stream>>>(tiled_input_ids,tiled_input_lengths,tiled_prompt_lengths,input_ids,input_lengths,prefix_prompt_lengths,max_input_length);}
}template<bool PREFIX_PROMPT>
__global__ void tileGptPromptInputs(int*       tiled_input_ids,int*       tiled_input_lengths,int*       tiled_prompt_lengths,const int* input_ids,const int* input_lengths,const int* prefix_prompt_lengths,const int  max_input_length)
{if (threadIdx.x == 0) {tiled_input_lengths[blockIdx.x * gridDim.y + blockIdx.y] = input_lengths[blockIdx.x];if (PREFIX_PROMPT) {tiled_prompt_lengths[blockIdx.x * gridDim.y + blockIdx.y] = prefix_prompt_lengths[blockIdx.x];}}for (int index = threadIdx.x; index < max_input_length; index += blockDim.x) {tiled_input_ids[(blockIdx.x * gridDim.y + blockIdx.y) * max_input_length + index] =input_ids[blockIdx.x * max_input_length + index];}
}

在任务划分阶段,按照batch维度进行划分,每个任务起了至少1024个线程来进行拷贝。由于没有beam width,那么gridDim.y就会是1。在kernel中,首先将input_length拷贝到tiled_input_lengths,之后再处理

在kernel中,由于没有beam width,那么gridDim.y就会是1。在kernel中,首先将input_length拷贝到tiled_input_lengths,之后再按照batch维度进行处理,将一个batch*max_input_length的数据进行展开,并拷贝到对应的buf,这有利于我们进行接下来的后续embedding和MHA计算
在这里插入图片描述

下一回预告

下一回继续讲解forward函数中的处理逻辑,会简单讲解embedding, pre_layernorm,以及进入attention_layer之后的初步讲解

本文来自互联网用户投稿,该文观点仅代表作者本人,不代表本站立场。本站仅提供信息存储空间服务,不拥有所有权,不承担相关法律责任。如若转载,请注明出处:http://www.mzph.cn/bicheng/42491.shtml

如若内容造成侵权/违法违规/事实不符,请联系多彩编程网进行投诉反馈email:809451989@qq.com,一经查实,立即删除!

相关文章

架构师学习理解和总结

1.架构设计理念 2.架构方法论 2.1需求分析 2.1.1常见需求层次 2.1.2 常见需求结果 2.1.3 需求与架构关系 2.2 领域分析 2.3 关键需求 2.4 概念架构设计 2.5 细化架构设计 2.6 架构设计验证 3.架构设计工具 3.1 DDD领域建模 3.2 41视图分析法 3.3 UML设计工具 4.架构师知…

【PYG】pytorch中size和shape有什么不同

一般使用tensor.shape打印维度信息&#xff0c;因为简单直接 在 PyTorch 中&#xff0c;size 和 shape 都用于获取张量的维度信息&#xff0c;但它们之间有细微的区别。下面是它们的定义和用法&#xff1a; size&#xff1a; size 是一个方法&#xff08;size()&#xff09;和…

系统化学习 H264视频编码(01)基础概念

说明&#xff1a;我们参考黄金圈学习法&#xff08;什么是黄金圈法则?->模型 黄金圈法则&#xff0c;本文使用&#xff1a;why-what&#xff09;来学习音H264视频编码。本系列文章侧重于理解视频编码的知识体系和实践方法&#xff0c;理论方面会更多地讲清楚 音视频中概念的…

Swift 中的方法调用机制

Swift 方法调用详解&#xff1a;与 Objective-C 的对比、V-Table 机制、Witness Table 机制 在 iOS 开发中&#xff0c;Swift 和 Objective-C 是两种常用的编程语言。尽管它们都能用于开发应用程序&#xff0c;但在方法调用的底层机制上存在显著差异。本文将详细介绍 Swift 的…

Calicoctl工具学习 —— 筑梦之路

官方文档&#xff1a; Calico Documentation | Calico Documentation 插件方式安装 calicoctl 工具 curl -o kubectl-calico -O -L "https://github.com/projectcalico/calicoctl/releases/download/v3.20.0/calicoctl"cp kubectl-calico /usr/bin/kubectl-calic…

05:C语言函数

C语言函数 1、函数的三要素1.1、函数定义1.2、函数声明1.3、函数调用 函数是C语言代码的基本组成部分&#xff0c;它是一个小的模块&#xff0c;整个程序由很多个功能独立的模块&#xff08;函数&#xff09;组成。这就是程序设计的基本分化方法。 之前接触过的函数&#xff1a…

实验2 Aprori关联挖掘算法

目 录 一、实验目的... 1 二、实验环境... 1 三、实验内容... 1 3.1 connect_string()函数解析... 1 3.2 find_rule()函数解析纠错... 2 3.3 关联规则挖掘... 4 四、心得体会... 7 一、实验目的 &#xff08;1&#xff09;理解Aprori关联挖掘算法的程序编写&#xff1b; &…

PYTHON自学笔记(一)vscode配置

安装python 自行官网下载 安装vscode 自行官网下载 环境变量设置 把python和scripts的文件路径&#xff0c;添加到环境变量的path中&#xff0c;如图&#xff1a; 此项不弄&#xff0c;在命令行模式中系统不会认为你装了python和pip&#xff0c;你的输入相关命令shell不会…

# [0705] Task06 DDPG 算法、PPO 算法、SAC 算法【理论 only】

easy-rl PDF版本 笔记整理 P5、P10 - P12 joyrl 比对 补充 P11 - P13 OpenAI 文档整理 ⭐ https://spinningup.openai.com/en/latest/index.html 最新版PDF下载 地址&#xff1a;https://github.com/datawhalechina/easy-rl/releases 国内地址(推荐国内读者使用)&#xff1a; 链…

[机器学习]-4 Transformer介绍和ChatGPT本质

Transformer Transformer是由Vaswani等人在2017年提出的一种深度学习模型架构&#xff0c;最初用于自然语言处理&#xff08;NLP&#xff09;任务&#xff0c;特别是机器翻译。Transformer通过自注意机制和完全基于注意力的架构&#xff0c;核心思想是通过注意力来捕捉输入序列…

实现Windows经典游戏——扫雷

扫雷游戏以前是windows的内置小游戏&#xff0c;相信不少人都玩过这么一款游戏。 1.扫雷游戏规则 以9*9的简单模式为例&#xff0c;游戏开始时系统会生成10颗雷随机分布在这9*9的格子当中。 [外链图片转存失败,源站可能有防盗链机制,建议将图片保存下来直接上传(img-Yv29OzCb…

中英双语介绍意大利(Italy):有哪些著名景点、出名品牌?

中文版 意大利概述 意大利&#xff0c;位于欧洲南部&#xff0c;是一个以其悠久的历史、丰富的文化遗产和美丽的自然风光而闻名的国家。意大利不仅是文艺复兴的发源地&#xff0c;还拥有众多世界著名的城市、景点和品牌。 著名城市 罗马&#xff08;Rome&#xff09;&#x…

DataWhale AI夏令营 对话分角色要素提取挑战赛 学习笔记

大模型技术概览 大模型技术是指利用庞大的数据集训练出的深度学习模型&#xff0c;这些模型在处理复杂自然语言处理&#xff08;NLP&#xff09;任务方面表现出色&#xff0c;例如文本分类、情感分析、机器翻译等。在对话系统中&#xff0c;大模型技术通过理解和生成更加自然和…

CosyVoice多语言、音色和情感控制模型,one-shot零样本语音克隆模型本地部署(Win/Mac),通义实验室开源

近日&#xff0c;阿里通义实验室开源了CosyVoice语音模型&#xff0c;它支持自然语音生成&#xff0c;支持多语言、音色和情感控制&#xff0c;在多语言语音生成、零样本语音生成、跨语言声音合成和指令执行能力方面表现卓越。 CosyVoice采用了总共超15万小时的数据训练&#…

【2024_CUMCM】时间序列算法ARMA

目录 2023-c-问题二 问题分析 介绍 单位根检验 白噪声检验 自相关和偏自相关图 利用信息准则定阶 构建AMIMA模型 D-W检验 预测 代码 2023-c-问题二 问题分析 ARMA适合多个领域的时间序列分析&#xff0c;不同时间的定价策略属于这类问题。 介绍 ARMA模型&…

C++入门7——string类详解

目录 1.什么是string类&#xff1f; 2.string类对象的常见构造 2.1 string(); 2.2 string (const char* s); 2.3 string (const string& str); 2.4 string (const string& str, size_t pos, size_t len npos); 2.5 string (const char* s, size_t n); 2.7 验证…

vue2项目迁移vue3与gogocode的使用

#背景 公司有个项目使用vue2jswebpack框架开发的&#xff0c;由于该项目内部需要安扫&#xff0c;导致很多框架出现了漏洞需要升级&#xff0c;其中主要需要从vue2升vue3,但是重新搭框架推翻重做成本太高&#xff0c;于是找到了gogocode。 #升级步骤踩坑 1. 安装 gogocode插…

牛客链表题:BM2 链表内指定区间反转

描述 将一个节点数为 size 链表 m 位置到 n 位置之间的区间反转&#xff0c;要求时间复杂度 &#x1d442;(&#x1d45b;)O(n)&#xff0c;空间复杂度 &#x1d442;(1)O(1)。 例如&#xff1a; 给出的链表为 1→2→3→4→5→&#x1d441;&#x1d448;&#x1d43f;&#x…

分表分库是一种数据库架构的优化策略,用于处理大规模数据和高并发请求,提高数据库的性能和可扩展性。

分表分库是一种数据库架构的优化策略&#xff0c;用于处理大规模数据和高并发请求&#xff0c;提高数据库的性能和可扩展性。以下是一些常见的分表分库技术方案&#xff1a; 1. **水平分表&#xff08;Horizontal Sharding&#xff09;**&#xff1a; - 将单表数据根据某个…

【机器学习】基于线性回归的医疗费用预测模型

文章目录 一、线性回归定义和工作原理假设表示 二、导入库和数据集矩阵表示可视化 三、成本函数向量的内积 四、正态方程五、探索性数据分析描述性统计检查缺失值数据分布图相关性热图保险费用分布保险费用与性别和吸烟情况的关系保险费用与子女数量的关系保险费用与地区和性别…