做网站的叫什么职位,如何引流客源最快的方法,克隆视厅网站怎么做,网站开发技术招聘我自己的原文哦~ https://blog.51cto.com/whaosoft/11603901
#CSWin-UNet
将自注意力机制集成到UNet中#xff01;CSWin-UNet#xff1a;U型分割方法#xff0c;显著提高计算效率和感受野交互#xff01;本文提出了CSWin-UNet#xff0c;这是一种新颖的U型分割方法CSWin-UNetU型分割方法显著提高计算效率和感受野交互本文提出了CSWin-UNet这是一种新颖的U型分割方法它将CSWin自注意力机制集成到UNet中以实现水平和垂直条纹的自注意力。这种方法显著提高了计算效率和感受野交互。 深度学习尤其是卷积神经网络CNNs和Transformer架构在医学图像分割领域受到了广泛关注并取得了令人瞩目的成果。然而CNN固有的归纳偏置限制了它们在更复杂、更多变的分割场景中的有效性。 相反尽管基于Transformer的方法擅长捕捉全局和长距离的语义细节但它们面临着计算成本高的挑战。 在本研究中作者提出了CSWin-UNet这是一种新颖的U型分割方法它将CSWin自注意力机制集成到UNet中以实现水平和垂直条纹的自注意力。这种方法显著提高了计算效率和感受野交互。 此外作者创新性的解码器采用了一种内容感知重组算子该算子根据预测的核策略性地重组特征以精确恢复图像分辨率。作者在包括突触多器官CT、心脏MRI和皮肤病变在内的多样化数据集上的广泛实证评估表明CSWin-UNet在保持低模型复杂性的同时提供了高分割精度。 I Introduction
医学图像分割是医学图像计算和计算机辅助干预领域的一个基本研究课题主要通过处理图像以获取有益信息例如病变器官或组织的形状、大小和结构从而提供更准确和详细的诊断和治疗建议[1, 2]。
基于深度学习的医学图像分割方法能够直接在像素 Level 对整张图像进行分类并已在多个医学领域得到广泛应用包括肺部计算机断层扫描CT图像分割、脑部磁共振成像MRI分割以及心脏超声图像分割等。这些方法不仅提高了分割的准确性还进一步推动了医学成像领域的发展。卷积神经网络CNN是计算机视觉领域中应用最广泛的深度学习技术之一。全卷积网络FCN[5]作为CNN的一种扩展促进了医学图像分割领域的发展。现有研究提出了扩展卷积和上下文学习方法以解决传统卷积操作感受野有限的问题。此外UNet以其创新的U形编码器-解码器设计和跳跃连接将编码器和解码器的特征图合并保留了浅层的关键空间细节。这种架构已成为图像分割领域的标准。UNet的增强衍生版本如UNet、AttentionUNet 和 ResUNet进一步细化了分割能力并在多种成像模态上提供了改进的性能。
尽管基于卷积神经网络CNN的方法在医学图像分割中取得了成功但它们在捕获全局和长距离语义信息方面的能力有限并且存在固有的归纳偏置问题。受到 Transformer 架构在自然语言处理NLP领域[16]的变革性影响研究行人开始将这项技术应用于计算机视觉任务旨在缓解CNN的一些局限性。Transformer 架构的核心是自注意力机制它并行处理输入序列中所有位置嵌入的信息而不是顺序处理。这种机制使得 Transformer 能够熟练地管理长距离信息依赖关系并适应不同的输入序列长度。一种针对图像处理的特定改编——视觉 Transformer Vision Transformer[20]通过将输入图像分割成一系列固定 Patch 每个 Patch 转换成一个向量然后由 Transformer 编码器处理体现了这一点。通过编码阶段自注意力建立 Patch 间的关系捕捉全面的上下文信息。这些编码后的特征随后被用于目标检测和图像分割等任务利用解码器或分类器。视觉 Transformer 的引入不仅为图像处理注入了新视角而且取得了与传统CNN相媲美或超越的结果[21, 22, 23, 24]。尽管 Transformer 架构在处理全局和长距离语义信息方面表现出色但由于其自注意力机制的广泛性其计算效率往往受到影响。针对这种低效问题Swin Transformer[25]创新性地采用了窗口自注意力机制将注意力限制在图像中的离散窗口内极大地降低了计算复杂性。然而这种方法在一定程度上限制了感受野之间的交互。为了克服这一点CSWin Transformer[26]提出了交叉形状窗口CSWin自注意力它可以水平垂直并行地计算自注意力以更低的计算成本取得更好的结果。此外CSWin Transformer还引入了局部增强位置编码LePE在每个 Transformer 块上施加位置信息。与之前的位置编码方法[27, 28]不同LePE直接操纵注意力权重的结果而不是添加到注意力计算的输入中。LePE使得CSWin Transformer在目标检测和图像分割方面更为有效。随着 Transformer 的发展许多研究将CNN与 Transformer 块结合起来。TransUNet[13]和LeViT-UNet[29]将UNet与 Transformer 结合在腹部多器官和心脏分割数据集上取得了竞争性结果。此外一些研究行人还开发了使用纯 Transformer 的分割模型。Swin-UNet[30]采用Swin Transformer块构建类似UNet架构的编码器和解码器与TransUNet[31]相比性能有所提升。然而基于Swin Transformer的这种分割方法在感受野交互方面仍有限制且计算成本也相对较高。
医学图像通常具有高分辨率并包含许多相互关联的精细结构。作者主要关注的问题是如何在消耗较少计算资源的情况下更好地处理医学图像中的长距离依赖关系。此外与语义分割相比医学图像中准确的边界分割对于诊断和治疗至关重要。因此作者研究的另一个重点是在分割过程中如何保留更多详细信息并提供更明确的边界。受到创新的CSWin Transformer [26] 的启发作者提出了一种新型的基于Transformer的医疗图像分割方法名为CSWin-UNet。该方法旨在降低计算成本的同时提高分割准确性。
与TransUNet [13] 这种CNN-Transformer混合架构不同CSWin-UNet类似于Swin-UNet [30]是一种纯Transformer基础的U形架构。CSWin-UNet与Swin-UNet的关键区别在于前者在编码器和解码器中配备了CSWin Transformer块并根据不同尺度设计了不同数量的块。此外作者在解码器中引入了CARAFE内容感知特征重组层 [31] 用于上采样。
最初输入的医学图像被转换为卷积标记嵌入然后由编码器处理以提取上下文特征。这些特征随后由CARAFE层上采样该层能够精确地重新组装特征。此外作者还使用了跳跃连接以持续融合高级语义信息与低级空间细节。这个过程最终将特征嵌入转化为与原始输入尺寸相匹配的分割 Mask 。
通过十字形窗口自注意力机制作者的方法可以在降低计算复杂性的同时保持对医学图像的高效特征提取能力。此外结合经典的UNet架构它能够有效地在编码器和解码器中整合不同尺度的特征从而提高分割准确性。最后引入CARAFE层进行上采样可以更有效地保留分割目标的边缘和详细特征。
对CSWin-UNet方法的综合实验评估表明与现有方法相比它在分割准确性和稳健泛化能力方面具有优势。此外它在降低医学图像分割任务的计算复杂性方面也显示出显著的优势。
本研究的主要贡献如下
作者开发了一种新型的U形编码器-解码器网络架构CSWin-UNet专门针对医疗图像分割采用了CSWin Transformer块。引入了CSWin自注意力机制来实现水平和垂直条纹自注意力学习。这一增强显著扩大了每个标记的关注区域促进了更全面的分析和上下文整合。在解码器中采用了CARAFE层替代传统的转置卷积或插值策略进行上采样。这种选择使得能够更精确地生成像素级分割 Mask 。综合实验结果验证了CSWin-UNet不仅轻量级而且在计算效率和分割准确性方面都超过了现有方法。
本文的结构安排如下第二部分回顾了医疗图像分割领域的近期工作和进展为本研究引入的创新技术提供了背景。第三部分详细描述了新提出的CSWin-UNet的方法论突出了其架构及其组件的创新之处。第四部分展示了实验结果证明了CSWin-UNet与现有方法相比的有效性和效率。第五部分总结了全文。
II Related works
Self-attention mechanisms in image segmentation_
在图像分割领域中对自注意力机制的应用已经得到了广泛的研究。中的研究显示为适合的场景设计不同的自注意力机制可以显著提高分割性能。在医学图像分割任务中常涉及到微妙但关键的结构自注意力机制能够更好地捕捉这些复杂结构之间的关系使得设计出有效且适当的自注意力机制尤为重要。然而许多现有的视觉Transformer仍然使用计算复杂度高的全局注意力机制如图1(a)所示。为了解决这个问题Swin Transformer [25]采用了移位版的局部自注意力机制如图1(b)所示通过滑动窗口机制实现了不同窗口之间的交互。此外轴向自注意力[35]和交错注意力[36]分别沿水平和垂直方向计算条带内的注意力如图1(c)和(d)所示。然而轴向自注意力受限于序列机制和窗口大小而交错注意力在特定应用中由于窗口重叠而表现不佳。CSWin Transformer [26]引入了十字形窗口CSWin自注意力机制它能并行地计算水平和垂直条带区域的自我注意力。与之前的注意力机制相比这种注意力机制在处理图像处理任务时更为通用且有效。 CNN-based medical image segmentation_
在医学图像分割领域卷积神经网络CNNs被广泛采用一些关键架构推动了该领域的发展。其中全卷积网络FCN[5]以其端到端的架构脱颖而出直接对像素进行分类将全连接层转换为卷积层以适应任意大小的图像。UNet[9]模型其特点是具有对称的U形编码器-解码器架构在医学图像的精确分割方面表现出色。在FCN和UNet的基础上已经提出了许多改进方法。例如SegNet[37]结合了FCN和UNet的思想使用最大池化操作符来提高分割 Mask 的准确性并且已有效地应用于各种医学分割任务[38, 39]。UNet[10]通过整合密集嵌套的跳跃连接扩展了原始UNet的设计最小化了编码器和解码器之间的信息丢失从而提高了分割性能。AttentionUNet[11]通过在UNet架构中增加注意力机制提高了准确性和鲁棒性。最后nnU-Net[40]提出了一种自适应的网络架构选择方法能够根据特定任务需求和数据集特性自动优化模型配置从而在各种分割挑战中增强了适应性。此外MRNet[41]提出了一种多评分者一致模型来校准分割结果而Pan等人[42]设计了一种混合监督学习策略来解决医学图像标签稀缺的问题。
Transformer-based medical image segmentation
鉴于医学影像的高分辨率和复杂性它们包含了大量的像素和复杂的局部特征传统的基于CNN的医学图像分割方法虽然在捕捉详细图像信息方面有效但在获取全局和长距离语义上下文方面往往力不从心。相比之下凭借其全局上下文建模能力Transformer在有效编码更大接受域并学习远距离像素间关系方面发挥关键作用从而提升分割性能。这一优势促使研究行人将Transformer融入到医学图像分割框架中。例如TransUNet [13] 使用Transformer作为编码器来从医学图像中提取上下文表示并结合基于UNet的解码器进行精确的像素级分割。这种组合展示了Transformer捕捉全局上下文信息的增强能力从而提高了分割的准确性。同样TransFuse [43] 在单一框架内整合了CNN和Transformer分支并使用专门模块合并两条路径的输出以产生最终的分割 Mask 。此外UNetR [44] 利用Transformer编码输入的3D图像配合CNN解码器完成分割过程而MT-UNet [45] 引入了一种混合Transformer架构学习样本内和样本间的关系。HiFormer [46] 则提出了另一种混合模型将两个CNN与Swin Transformer模块和双 Level 融合模块结合以整合并传递多尺度特征信息到解码器。在纯Transformer方法中SwinUNet [30] 使用Swin Transformer [25] 作为编码器来捕捉全局上下文嵌入然后由UNet解码器逐步上采样利用跳跃连接增强细节保留。此外DFQ [47] 在Vision TransformerViT框架内引入了解耦的特征 Query 使分割模型能更广泛地适应不同的任务。
受到多头自注意力机制尤其是CSWin Transformer [26] 的启发作者开发了CSWin-UNet这是一种基于CSWin自注意力的医学图像分割方法。该模型在节约计算资源的同时提升了分割的准确性代表着将Transformer应用于医学图像分割领域的一个重大进步。
III Methodology
CSWin-UNet的整体架构如图2所示, 它由编码器、解码器和跳跃连接组成, 基本单元是 CSWin Transformer块。对于输入尺寸为 的医学图像, 与CvT[34]类似, 作者使用卷积标记嵌入使用 的核和步长为 4 来获得 的 Patch 标记, 其通道数为 。编码器和解码器均由四个阶段组成。与UNet[5]一样, 跳跃连接被用于在编码器和解码器的每个阶段合并特征, 以更好地保留上下文信息。在编码器中, 使用卷积层 的核和步长为 2) 进行下采样, 将分辨率降低到输入大小的一半, 同时将通道数加倍。解码器中的上采样通过CARAFE层完成将分辨率增加到输入大小的两倍同时将通道数减半。最后执行 CARAFE上采样操作将分辨率恢复到输入分辨率 , 并使用线性层将特征图转换为分割 Mask 。
CSWin Transformer块 传统的Transformer架构凭借其自注意力机制在处理所有像素位置以建立全局语义依赖方面表现出色然而这在高分辨率医学成像中会导致计算成本高昂。Swin Transformer [25]通过移位窗口注意力机制减轻了这些成本该方法将图像划分为不同的、不重叠的窗口从而实现局部的自注意力。这种适应有助于管理图像的高分辨率同时控制计算复杂度。然而这种方法的有效性取决于窗口大小较小的窗口可能会遗漏一些全局信息而较大的窗口可能会不必要地提高计算需求和存储要求。与移位窗口注意力机制相比CSWin自注意力将注意力组织成水平和垂直的条纹增强了并行计算能力。这种结构不仅节约计算资源还拓宽了感受野内的交互作用。如图3所示基于这种创新的自注意力设计构建的CSWin Transformer块包括一个CSWin自注意力模块、一个LayerNormLN层、一个多层感知机MLP以及跳跃连接。这种配置在局部和全局信息处理之间达到了最优平衡显著提高了复杂医学图像分割任务的效率和有效性。
在多头自注意力机制中, 输入特征 首先经过一次变换, 在 个头之间进行线性映射, 通常选择为偶数。与传统的自注意力和基于移位窗口的多头自注意力不同, CSWin自注意力独特地促进了在划分的水平或垂直条纹内进行局部自注意力学习, 如图4所示。这种配置允许每个头在其指定的条纹内水平或垂直地计算自注意力。这些操作并行进行, 有效地拓宽了注意力计算区域的范围, 同时减少了整体计算复杂度。 图4CSWin自注意力机制的说明。首先, 将多个头 分为两组 和 , 分别在水平和垂直条纹上并行执行自注意力, 并连接输出。接下来, 可以调整条纹的宽度 以达到最佳性能。通常, 对于更高分辨率选择较小的 , 对于更低分辨率选择较大的 。
在CSWin Transformer的水平条纹自注意力配置中, 输入特征 被系统地划分为 个不重叠的水平条纹, 表示为 , 其中每个条纹的宽度为 由比例 确定。参数 是可调整的, 对于平衡计算复杂度与模型的学习能力至关重要。具体来说, 较大的 增强了模型在每个条纹内探索长距离像素相关性的能力, 有可能捕捉到更广泛的环境信息。考虑在一个特定 Head 的计算, 记为第 个 Head 。在这种情况下, Query (Q)、键K和值 (V) 的维度各为 , 其中 是通道数, 是总 Head 数。第 个水平条纹内第 个 Head 的自注意力输出 计算如下: 其中 是第 个水平条纹的特征图; , 表示第 个 Head 的 、 和 V 的权重矩阵。这个操作分别对每个条纹并行执行,以允许在特定的水平条纹内进行自注意力。 个水平条纹的自注意力被连接起来, 构建第 个 Head 的水平自注意力H-Attention o 。
类似于水平条纹自注意力机制, 输入特征 被均匀划分为 个不重叠的垂直条纹 以进行垂直自注意力处理。其中条纹的高度也是 , 且 。以第 个注意力头为例, 其中 、 和 V 的维度为 。第 个注意力头在第 个垂直条纹中的自注意力输出 可以按以下方式计算: 其中, 是第 个垂直条纹的特征图。 个垂直条纹的自注意力被连接起来, 构建第 个注意力头的垂直自注意力 -Attention : 作者将 个注意力头分为两组, 每组包含 个头。这些组中的每个头都生成其自注意力输出。第一组负责学习水平条纹自注意力, 而第二组学习垂直条纹自注意力。在分别计算自注意力之后, 这两个组的输出被连接起来。这种连接沿着通道维度进行 式中 表示第 个注意力头; 是一个权重矩阵, 用于将多注意力头自注意力机制的拼接输出线性转换以产生最终的注意力输出, 这种线性转换有助于学习不同头之间的关系并融合注意力信息。拼接输出有效地结合了水平和垂直的上下文信息, 全面学习输入图像内的空间关系。
基于上述自注意力机制CSWin Transformer块可以定义为 其中 表示第 个CSWin Transformer块的输出或每个阶段的先前卷积层的输出。
Encoder
在编码器中, 输入图像的尺寸为 , 然后它们进入四个阶段进行特征提取。前三个阶段伴随着下采样操作。四个阶段中CSWin Transformer块的数量各不相同, 关于块数量的设置细节将在后文讨论。下采样层通过一个 Kernel 大小为 、步长为 2 的卷积层实现,将分辨率降低到输入大小的一半, 同时通道数翻倍。条带宽度 在不同阶段相应变化。随着分辨率的持续降低和通道数的增加, 在较大分辨率的阶段选择较小的 , 在较小分辨率的阶段选择较大的 , 有效地扩大了在每个较小分辨率阶段的每个标记的注意力区域。此外, 输入图像的分辨率为 。为了确保输入图像的中等特征图大小可以被 整除, 作者将四个阶段的 设置为 、、 和 7 。
Decoder
与编码器相对应, 解码器同样包含四个阶段。在最后三个阶段中, 通过CARAFE层实现图像分辨率和通道数的增加。这四个阶段中用于注意力学习的CSWin Transformer块的数量和条带宽度 与编码器中设定的参数一致。常用的上采样方法包括线性插值和转置卷积。双线性插值仅考虑相邻像素, 可能会模糊图像边缘, 导致分割结果的边界不清; 而转置卷积的感受野通常受限于核大小和步长, 这不仅限制了其表示局部变化的能力, 还需要学习转置卷积核的权重和偏置。与这些方法不同, 作者使用CARAFE [31] 来实现上采样。
CARAFE层是一种先进的上采样机制, 它主要由两个核心组件构成一个核预测模块和一个内容感知重组模块。核预测模块首先通过一个卷积层从编码特征中预测重组核。它包括三个子模块通道压缩器、上下文编码器和核归一化器。通道压缩器降低了输入特征图 中通道空间的维度, 从而降低了计算复杂性, 并专注于重要的特征信息。通道压缩之后, 上下文编码器处理降维后的特征图以编码上下文信息, 这对于生成重组核至关重要。每个预测的重组核通过核归一化器中的Softmax函数进行归一化, 以确保权重的输出分布是概率性的, 总和为 1 , 这增强了上采样过程的稳定性和性能。具有上采样比 其中 为整数, CARAFE旨在生成一个扩展的特征图 。对于 中的每个像素 , 它对应于 中的特定像素 , 由 和 确定。核预测模块 根据邻域 为每个像素 预测一个唯一的重组核 , 这是一个以 上的像素 为中心的 区域。这个邻域提取局部特征, 预测的核使用这些特征有效地重组并上采样特征图。 其中 表示内容编码器的感受野。
第二步是内容感知重组, 输入特征通过卷积层进行重组, 而内容感知重组模块 使用重组核 重组 。 其中 是重装核的大小。对于每个重装核 , 内容感知重装模块在局部方形区域内重新组装特征。模块 执行加权求和。对于像素位置 及其中心邻域 , 重装过程如下: 其中 。 中的每个像素对上采样像素 的贡献各不相同。重新组装的特征图能够增强对局部区域内相关信息关注, 相较于原始特征图, 提供了更稳健的语义信息。此外, 与 UNet [9]类似, 作者使用跳跃连接将编码器和解码器输出的特征图进行合并, 从而提供了更丰富、更精确的空间信息, 有助于恢复图像细节。随后, 使用 卷积核在拼接后减少通道数, 确保与上采样过程中的特征通道数保持一致。
IV Experiments
Implementation details
CSWin-UNet是使用Python和PyTorch框架实现的。模型的训练和评估是在一块拥有24GB VRAM的NVIDIA(r) GeForce RTX(tm) 3090 GPU上进行的。作者使用从ImageNet [48]预训练的权重来初始化CSWin Transformer块, 以利用先验知识并加速收玫过程。在数据增强方面, 采用了翻转和旋转等方案, 以增强训练数据集的多样性, 从而帮助模型更好地泛化到未见过的数据。在训练阶段, 批量大小设置为 24 , 学习率设置为 0.05 。使用带有 0.9 动量和 权重衰减的随机梯度下降SGD方法进行优化。这种设置旨在平衡快速学习和收玫稳定性。此外, 为了有效地训练CSWin-UNet, 作者采用了一个组合损失函数, 该函数融合了Dice损失和交叉熵损失, 定义如下: 其中, 和 是两个超参数, 分别用于平衡 和Loss 对最终损失的影响。这个组合损失旨在同时关注像素 Level 的准确性和整体分割质量, 确保在各种医学图像分割任务中实现健壮的学习和性能提升。
Datasets and metrics
深度学习模型的性能在很大程度上依赖于用于训练的数据集的质量和规模。近年来大规模数据集的发展显著推动了深度学习技术在各个领域的进步。在本节中作者回顾了文献中常用的数据集以及用于评估深度学习模型性能的评价指标。
数据集
大多数深度学习研究依赖于大规模和高质量数据集的可用性。以下作者总结了文献中广泛使用的几个流行数据集。
ImageNet
ImageNet是一个视觉数据库用于视觉目标识别软件研究。它是计算机视觉领域最有影响力的数据集之一包含超过1400万张图片和超过2万个类别。
CIFAR-10和CIFAR-100
CIFAR数据集是一组常用于训练机器学习和计算机视觉算法的图像。CIFAR-10包含60,000张32x32彩色图像分为10个类别而CIFAR-100有100个类别每个类别包含600张图像。
MNIST
MNIST数据集是机器学习社区中的一个经典数据集由28x28灰度手写数字图像组成。它包括一个包含60,000个样本的训练集和一个包含10,000个样本的测试集。
评价指标
评估深度学习模型的性能需要使用适当的评价指标。作者讨论文献中一些最常用的评价指标。
准确度
准确度是最直观的评价指标之一通常被用作分类任务的标准默认度量。它定义为正确预测的数量除以总预测数量。
精确度、召回率和F1分数
在处理不平衡数据集时精确度、召回率和F1分数是更具有信息量的评价指标。精确度衡量正确识别为阳性的比例而召回率衡量正确识别出的实际阳性样本比例。
Top-k准确度
Top-k准确度是标准准确度指标的一个变体。它衡量正确标签在top k预测中的百分比。
Iv-B1 Synapse dataset
synapse多器官分割数据集包括来自2015年MICCAI多图谱腹部器官分割挑战赛的30个CT扫描, 总共包含 3779 张腹部CT图像。每个CT扫描由85到198个切片组成, 每个切片像素为 , 每个 Voxel 的大小为 毫米 3 。按照文献 113 5018]中的设置选择了用于训练的集合以及用于评估的12个集合。对八种腹部器官主动脉、胆囊、左肾、右肾、肝脏、胰腺、脾脏、胃) 的分割性能使用平均Dice相似系数 (DSC)和平均Hausdorff距离 (HD) 作为评价指标。
Iv-B2 ACDC dataset
自动心脏诊断挑战ACDC数据集在2017年的ACDC挑战期间发布提供了一个包含多个类别的心脏3D MRI数据集其中包括通过电影式MR 1.5T和3T扫描仪获取的100组短轴MR心脏图像。医学专家为三个心脏结构提供了标注右心室RV、心肌MYO和左心室LV[51]。作者随机选择了70组MR图像用于训练10组用于验证20组用于评估。ACDC数据集使用平均_dice相似性系数DSC作为评估指标以评价三个心脏结构的分割结果。
Iv-B3 Skin lesion segmentation datasets
作者在ISIC2017 [52]ISIC2018 [53]以及PH[54]数据集上进行了实验。ISIC数据集包含了大量的皮肤镜图像覆盖了各种皮肤病变。遵循HiFormer [46]中的设置作者在ISIC2017数据集中使用了1400张图像进行训练200张图像进行验证以及400张图像进行测试在ISIC2018数据集中使用了1815张图像进行训练259张图像进行验证以及520张图像进行测试在数据集中使用了80张图像进行训练20张图像进行验证以及100张图像进行测试。作者使用平均Dice相似系数DSC、敏感性SE、特异性SP和准确率ACC作为指标来评估皮肤病变分割任务。
Results on Synapse dataset
如下表1所示作者在Synapse数据集上提出的方法改善了平均DSC和HD。同时作者在图5中展示了平均DSC、平均HD以及每个器官DSC的误差条95%置信区间。与TransUNet [13]和Swin-UNet [30]相比作者的平均DSC分别提高了3.64%和1.99%平均HD分别改善了12.83%和2.69%。值得注意的是在胰腺分割方面CSWin-UNet的DSC显著高于其他分割方法。与其他器官不同胰腺具有模糊的边界和多变性作者的方法在胰腺分割上取得了更精确的结果表明作者的CSWin-UNet在复杂的分割环境中提供了更高的分割精度。 为了更直观地评估所提出的方法作者进行了分割结果的视觉分析。图6展示了在Synapse数据集上的比较结果。第一行显示在分割像胆囊绿色标签这样的小器官时Swin-UNet和HiFormer-B出现了明显的错误Swin-UNet [30]未能准确勾勒边界而HiFormer-B [46]错误地将其他区域识别为胆囊。第二行表明Swin-UNet、TransUNet [13]、HiFormer-B和UNet [9]在完全分割胃橙色标签方面均失败。第三行揭示Swin-UNet和HiFormer-B错误地将大片其他器官区域标记为胰腺黄色标签。考虑到定量指标和视觉结果作者提出的CSWin-UNet实现了对精细和复杂器官的准确分割产生了更精确的分割结果展示了在复杂背景下更强的鲁棒性并且在边缘结构处理方面表现更佳。 Results on ACDC dataset
表2展示了作者提出的CSWin-UNet在ACDC数据集上的实验结果并将其与其他先进方法进行了比较。图7表示了每个心脏结构平均DSC和DSC的错误条95%置信区间。在表中RV代表右心室MYO代表心肌LV代表左心室。结果显示提出的CSWin-UNet能更好地识别和分割这些器官准确率达到91.40%显示出良好的泛化能力和鲁棒性。 Results on skin lesion segmentation datasets
表3展示了实验结果图8显示了在三个皮肤病变分割数据集上DSC、SE、SP和ACC的误差条95%置信区间。实验结果表明在大多数评估指标上所提出的CSWin-UNet方法优于其他方法。特别是与Swin-UNet [30]相比CSWin-UNet在大多数指标上取得了更好的性能显示出满意的泛化能力。作者还将在图9中可视化的皮肤病变分割结果。与Swin-UNet [30]相比作者的CSWin-UNet在保留分割目标的边缘和详细特征方面具有一定的优势。然而在低对比度或遮挡的情况下如图9(d)所示分割产生了显著的错误。 Comparison of computational efficiency
神经网络模型设计的一个基本目标是尽可能减少参数数量和计算复杂度同时保持其性能。这种减少对于在计算资源有限的设备上实现更高效的模型训练和部署至关重要。因此在评估一个模型时不仅要考虑其准确性和泛化能力还要考虑其参数数量和计算复杂度。在这里作者使用浮点运算次数FLOPs和参数数量以百万计M来衡量计算复杂度。在Synapse数据集上的性能比较显示在表4中。结果表明所提出的CSWin-UNet在最低复杂度条件下实现了出色的分割性能。 Ablation studies
在本文的这一部分作者对CSWin-UNet在Synapse数据集上的性能进行了消融研究。具体来说作者探讨了解码器中不同的上采样策略、跳跃连接的数量、不同的网络架构以及组合损失函数中不同超参数对性能的影响。
V-B1 Upsampling strategy
在编码器中通过使用步长为2的卷积层进行下采样相应地在解码器中需要上采样以恢复特征图从而保留更多信息。在本文中作者引入了CARAFE层以实现上采样并增加特征通道数该层使用输入特征本身的内容来指导上采样过程从而更准确、高效地进行特征重组。为了验证CARAFE层的有效性作者在Synapse数据集上进行了实验比较了双线性插值、转置卷积以及CARAFE层在CSWin-UNet中的表现如表5所示。采用CARAFE层进行上采样获得了最高的分割准确度。此外与转置卷积相比CARAFE引入的计算开销非常小。实验结果表明结合了CARAFE层的CSWin-UNet能够达到最优性能。 V-B2 Skip connection
类似于UNet作者也引入了跳跃连接以增强细粒度的分割细节通过恢复低级空间信息。在CSWin-UNet中跳跃连接位于1/4、1/8和1/16的分辨率尺度上。作者依次减少了1/16、1/8和1/4尺度上的跳跃连接将跳跃连接的数量设置为3、2、1和0以探索不同数量的跳跃连接对分割精度的影响。如表6所示分割精度通常随着跳跃连接数量的增加而提高。值得注意的是相比于大器官如肝脏、脾脏和胃CSWin-UNet在小器官如主动脉、胆囊、肾脏和胰腺的分割精度上有更显著的提升。因此为了达到最佳性能作者将跳跃连接的数量设置为3。 V-B3 Network architecture
层数过少的神经网络可能导致特征表示丰富度和准确性不足难以理解图像上下文从而造成分割性能不佳。相反过多的层数会增加计算负担使网络难以收敛。因此在设计网络架构时在网络深度和模型性能之间取得了平衡使得模型能够在有限的计算资源下实现高分割精度。此外为防止因层数过多而导致的不收敛问题在最后阶段将模块数量设置为1。通过比较其他基于Transformer的医学图像分割方法的参数数量和计算成本作者将四个阶段的模块数量设置为、和编码器和解码器模块对称排列。如表7所示具有模块设置的网络架构实现了最佳性能。 Iv-C4 Combined loss function
作者探讨了组合损失函数的不同超参数对分割准确性的影响。在这里,作者将方程式10中的 和 分别设置为 、、、 和 。作者在Synapse数据集上进行了一项消融研究, 实验结果表明, 使用组合损失函数比单独使用Dice损失或交叉摘损失能获得更高的分割准确度尤其是在仅使用Dice损失而不结合交叉摘损失的情况下。表8显示当 和 设置为 时, 分割性能达到最优。 Discussions
作者在三个不同类型的医学图像分割数据集上的全面实验结果证明了作者提出的CSwin-UNet在多种模态的医学图像中比其他最先进的医学图像分割方法更为先进和适用。这些数据集包括CT、MRI和皮肤病变图像。
然而作者的方法在一些具有挑战性的情况下表现出一些不足例如在Synapse数据集中胆囊和肾脏区域不同样本的分割精度存在显著差异如图6所示。根据图9的可视化结果在处理皮肤病变分割数据集中的低对比度图像时分割性能还有很大的提升空间。
此外模型的预训练对其性能产生了显著影响。在本次研究中作者使用在ImageNet [48]上由CSwin Transformer [26]训练的权重来初始化编码器和解码器。因此探索端到端的医学图像分割方法是作者在未来努力追求的研究课题之一。
V Conclusion
在本论文中作者通过引入一种高效且轻量级的方法——CSWin-UNet解决了先前基于Transformer的医疗图像分割模型在感受野交互方面的局限性。利用来自CSWin Transformer的CSWin自注意力机制作者将这项技术融入了一种U形编码器-解码器架构中。
这种融合不仅降低了计算成本还提升了感受野的交互作用和分割精度。在解码器中采用了CARAFE层进行上采样这有助于保留复杂的细节并提高器官边缘分割的精确度。
在三个大规模医疗图像分割数据集上的全面评估表明CSWin-UNet在分割精度上超越了其他最先进的方法。此外CSWin-UNet在模型参数和计算负载方面更为轻量这表明在复杂的医疗图像分割任务中深度学习应用具有进一步的优化和增强的巨大潜力。 #GLARE
GLARE 利用外部正常光照先验实现逼真的低光照增强效果
本文提出一种新型的基于生成式隐层特征的码本检索的低光照增强网络GLARE大量实验证明GLARE在多个基准数据集和真实数据上的卓越性能以及GLARE在低光照目标检测任务中的有效性进一步验证了其在高层次视觉应用中的适用性。
论文链接https://arxiv.org/pdf/2407.12431
GitHub链接https://github.com/LowLevelAI/GLARE
亮点直击
首次采用外部正常光码本作为指导去自然地增强低光照图像。提出一种名为 GLARE 的新型低光照图像增强器利用潜在归一化流来学习低光照图像特征分布使其与自然光图像特征对齐。提出一种具有可调功能的自适应特征变换模块以在保证输出自然性的同时巩固保真度。大量实验证明GLARE在5个配对基准和4个真实世界数据集上的低光照图像增强任务中显著优于现有的最先进方法并且作为高层次目标检测任务的预处理方法时GLARE也具有很强的竞争力。
本文提出一种新型的基于生成式隐层特征的码本检索的低光照增强网络命名为GLARE。其中码本先验是使用矢量量化(VQ)策略从未受损的正常光图像中提取的。更重要的是我们开发了一种生成式可逆的隐层归一化流模块将低光照图像的特征分布对齐到正常光图像的隐层表征从而保证了在码本中正确的代码检索。为了保留由码本先验提供的真实细节的同时进一步提高保真度我们设计了一种新颖的自适应特征变换模块。该模块包括一个自适应特征混合环节和一个双解码器结构并具备用户调节功能。大量实验证明GLARE在多个基准数据集和真实数据上的卓越性能以及GLARE在低光照目标检测任务中的有效性进一步验证了其在高层次视觉应用中的适用性。
方法
除了引入外部自然光码本来指导低光到正常光的映射外本文工作的创新之处还在于独特的可逆潜在归一化流和自适应特征变换模块。这些模块旨在最大限度地发挥自然光码本先验的潜力并生成具有高保真度的逼真结果。
本文提出的方法的概览如图3所示其中我们方法的训练可以分为三个阶段。在第一阶段在数千张清晰的自然光图像上预训练VQGAN以构建一个全面的码书。在第二阶段利用低光照-正常光照图像对训练I-LNF模块实现低光照和正常光照特征之间的分布变换。在最后阶段提出AFT模块包含固定的自然光解码器、自适应混合块和多尺度融合解码器用于保留码本所提供的自然性的同时增强细粒度细节。 阶段I正常光码本学习
为了学习一个通用且全面的码本先验, 我们利用了一个结构类似于[15]的 VQGAN。具体来说, 一个自然光图像 首先被编码并重塑为隐层表征 , 其中 和 分别表示图像宽度、图像高度、隐层特征的维度和隐层特征的总数; 是自然光编码器的下采样因子。每个潜在向量 可以使用最近邻匹配量化为对应的代码 , 公式如下: 其中 表示包含 个离散码的可学习码本, 每个码由 表示。随后, 量化后的码 被传送到自然光解码器 (记作 ) 以生成重建的图像 。
为了更好地展示正常光码本先验的优势和局限性本文在低光照-正常光图像对上微调了预训练的VQGAN编码器。具体来说本文在图2b的第2列中展示了增强后的结果并使用t-SNE可视化了由微调后的自然光编码器生成的低光照图像特征见图2a这证明了外部自然光先验在低光照图像增强中的有效性。此外这些可视化结果启发作者设计额外的网络来对齐低光照特征和正常光隐层表征以进一步提高增强性能。 阶段II生成式隐层特征学习
为了充分利用外部码本先验的潜力本文从减少低光照和正常光照特征分布差异的角度设计了额外的机制。具体来说本文开发了一种可逆潜在归一化流以实现低光照和正常光照特征分布之间的转换从而实现更准确的码本检索。
如图 3 所示, 本文在第二阶段优化了两个关键组件条件编码器和可逆归一化流模块。1) 条件编码器 其结构与自然光编码器 相同输入一个低光图像并输出条件特征 。2本工作中的可逆归一化流模块通过一个可逆网络实现表示为 。该模块利用 作为条件将复杂的自然光特征分布 转换为一个潜在特征即 。第二阶段训练的重点是获得一个在 空间中的简化分布 , 例如高斯分布。因此, 条件分布 [45] 可以隐式表示为: 与传统的归一化流方法 不同, 本文在特征层面而不是图像空间独特地应用了归一化流, 并且本文的设计中没有集成任何压缩层。此外, 本文提出使用由卷积层基于 生成的低光照特征 作为 的均值, 而不是使用标准高斯分布作为 的分布。公式 (2) 中的条件分布使得可以通过最小化公式 (3) 中的负对数似然来训练条件编码器和可逆归一化流模块。此外, 可以从 中采样 并使用完全可逆的网络 得到低光照输入的高质量特征 。 第二阶段训练完成后本文在LOL数据集[65]上评估了提出的模型以验证可逆归一化流模块的有效性。如图2a所示由可逆归一化流模块生成的低光照特征分布与正常光照特征分布高度一致促进了准确的码本组装。此外令人满意的增强结果图2b第3列表明本文提出的方法在第二阶段后已经具备良好的低光照图像增强性能。然而这些结果仍然有很大的改进空间特别是在保真度方面。例如颜色图2b第1行或结构细节图2b第2行与真实值有显著差异。这一观察促使作者将输入信息融入解码过程中以提高保真度。
阶段III自适应特征融合
为了进一步增强纹理细节和保真度, 本文提出了一种自适应特征变换模块, 该模块灵活地将条件编码器生成的特征 融入解码过程中, 其中 表示分辨率级别。具体来说, 为了保持正常光解码器 (NLD) 的真实输出并避免受损的低光照特征的影响, 本文采用了双解码器架构, 并参考 开发了多尺度融合解码器 (MFD)。双解码器设计使本文如方程 4 所示, 能够利用可变形卷积 (dconv) 来对 NLD 特征 ( ) 进行变形处理, 并将变形后的特征 输入多尺度融合解码 得到最终的增强结果。 其中 和 分别表示分辨率级别和目标特征本文设计了一种新颖的特征融合网络该网络自适应地将低光照信息融入上述变形操作并在实际测试中为用户提供潜在的调整功能。
自适应混合块 多尺度融合解码器MFD在结构上与正常光解码器NLD相似旨在解码生成的低光照特征 并获得中间层表征 其中 表示分辨率级别。在每个分辨率级别来自条件编码器信息 被添加到相应的 中, 以引入更多的低光照信息。不同于常规的特征融合操作本文使用了一种自适应混合策略 如下所示 其中 表示可学习系数, 表示 Sigmoid 运算符, 用于实际测试中的调整, 训练时设置为 1 。
灵活调整功能 尽管公式5中 β 在训练阶段被设置为1但在测试真实世界图像时用户可以根据个人偏好灵活调整。这种设计的灵感来自于许多现有方法在处理真实世界数据时通常表现不佳造成这一现象的原因是因为真实世界数据中的光照情况通常与训练阶段使用的图像有所不同。
实验
与现有方法的比较
比较的方法包括 KinD, MIRNet, DRBN, SNR, Restormer, Retformer, MRQ, LLFlow, LL-SKF等数据集LOLLOL-v2-real, LOL-v2-synthetic, SDSD-indoor, SDSD-outdoor, 分别在这些数据集的训练数据上训练在相应测试集上评估。跨数据集验证在LOL训练集上训练在真实世界数据集: MEF, LIME, DICM 和NPE上测试。
定量比较 如表1所示 GLARE在五个基准数据集上优于当前最先进的方法。本文提出的GLARE在PSNR上表现出色在LOL和LOL-v2-synthetic数据集上分别超过LL-SKF 0.55 dB和0.74 dB。此外在SDSD-indoor和SDSD-outdoor数据集上它比Retformer分别提高了0.33 dB和1.01 dB。另外GLARE的LPIPS在表1中分别超过了第二名20.9%和12.6%这表明本文提出的方法的增强结果与人类视觉系统更为一致。表2展示了在非配对的真实世界数据集上的跨数据集评估结果。本文首先在LOL训练集上训练GLARE。然后将在LOL测试数据上表现最佳的模型应用在四个非配对真实世界数据集上。与当前最先进的方法相比GLARE在DICM和MEF数据集上表现优越并在平均表现上达到最佳。这些结果不仅证明了本文方法在生成高质量视觉结果方面的优越性也展示了其良好的泛化能力。 定性比较 图4、图5和图6展示了本文提出的 GLARE 与其他方法的视觉对比结果。显然以往的方法在噪声抑制方面表现较差。此外它们还往往会产生明显的色彩失真参见图4中的KinD、LLFlow、Retformer、LL-SKF的增强结果以及图5中的SNR和LLFlow。此外从定性比较来看可以看到LLFlow、LL-SKF和Retformer在其增强结果中可能会导致细节缺失参见图4和图5而图4中的KinD和图5中的SNR由于引入了不自然的伪影而表现不佳。相比之下GLARE能够有效增强低光照图像同时可靠地保留色彩和纹理细节且。图6中在非配对真实世界数据集上的视觉比较也展示了本文提出的 GLARE 在细节恢复和色彩保持方面的优势。 低光照物体检测
在ExDark数据集上进行实验该数据集收集了7,363张低光图像分为12个类别并附有边界框注释旨在研究低光图像增强LLIE方法作为预处理步骤在改善低光物体检测任务中的有效性。本文使用在LOL数据集上训练的不同的LLIE模型来增强ExDark数据集内图片然后对增强后的图像进行目标检测。本文采用的目标检测器是预先在COCO数据集上训练的YOLOv3并使用计算平均精度AP和平均的APmAP作为评估指标。
表3 提供了本文方法 GLARE 与其他方法的定量比较结果。与KinD、MBLLEN、LLFlow和LL-SKF相比本文提出的GLARE在mAP方面至少提高了0.8分。更重要的是GLARE也优于在低光照目标检测中表现突出的IAT方法。图9展示了各个LLIE方法低光目标检测上的视觉比较结果。可以看出尽管每种LLIE方法在一定程度上提高了图像的可见度但本文方法 GLARE 实现了最佳的视觉表现从而对下游的检测任务最有利。不出所料GLARE 的增强效果使得YOLO-v3检测器能够以更高的置信度识别更多的目标。 消融实验
为了验证GLARE每个组件的有效性并证明用于训练的优化目标的合理性本节在LOL数据集上进行了广泛的消融实验。具体来说本节讨论了自适应特征融合模块、可逆归一化流模块和正常光码本先验的重要性。
自适应特征融合模块 从GLARE中移除自适应特征融合模块可以得到一个简单的低光图像增强模型称为SimGLARE。基本上SimGLARE仅利用自然光码本先验中的信息而不进行特征变换。SimGLARE的定量结果如表4所示。在PSNR、SSIM和LPIPS方面SimGLARE在低光照图像增强任务中相当具有竞争力与表1中的SOTA方法相比。然而借助提出的自适应特征融合模块本文提出的GLARE在定量指标和视觉结果上均取得了进一步的改进如图7所示。此外表4中还检验了各种损失函数显示了本文在第三阶段选择的损失函数是合理的。本文还设计了两个变体分别命名为Variant 1和Variant 2以说明所提出的双解码器架构和自适应混合块的重要性。具体来说Variant 1直接将低光照特征通过自适应混合块引入到正常光解码器NLD而Variant 2采用并行解码器策略但用skip-connection操作[26]替换了自适应混合块。通过比较表4中的(4)与(2)和(3)可以发现PSNR和SSIM与LPIPS呈负相关这验证了本文提出的自适应混合块和双解码器设计的有效性。 可逆归一化流模块 为了展示I-LNF和采用的NLL损失的重要性本文基于SimGLARE进行了几种调整(1)使用L1损失训练SimGLARE以验证本文采用的NLL损失的有效性。(2)用一个结构上类似于[84]的Transformer模型取代可逆归一化流模块直接预测码本中的码索引。(3)在(1)的基础上移除可逆归一化流模块并在低光-正常光图像对上训练条件编码器。定量结果如表5所示。通过比较(4)和(1)可以验证NLL损失的优越性。此外比较图8中第二列和第三列的图像也表明与L1损失相比使用NLL损失可以产生更清晰的轮廓和边缘。此外与基于Transformer的码预测方法相比本文提出的可逆归一化流模块可以帮助生成更好地与正常光特征对齐的低光特征从而确保更准确的码匹配并实现更优的性能。更重要的是当从SimGLAREL1中移除可逆归一化流模块时PSNR降低1.16 dB和SSIM降低0.017显著下降这证明了本文提出的可逆归一化流模块的有效性。 自然光码本先验 结论
本文提出了一种名为GLARE的低光图像增强LLIE新方法。鉴于LLIE的病态性质引起的不确定性和模糊性本文利用从清晰、曝光良好的图像中使用VQGAN获得的正常光码本来指导低光到正常光的映射。为了更好地发挥码本先验的潜力本文采用了可逆潜在归一化流来生成与正常光隐层表征对齐的低光照特征从而最大限度地提高码向量在码本中正确匹配的概率。最后本文引入了具有双解码器架构的自适应特征变换模块以灵活地在解码过程中提供输入信息从而在保持感知质量的同时进一步提高增强结果的保真度。大量实验表明本文提出的GLARE在5个配对数据集和4个真实世界数据集上显著优于当前最先进的方法。GLARE在低光目标检测中的卓越性能使其成为高层次视觉任务中有效的预处理工具。 #SLAB
华为开源通过线性注意力和PRepBN提升Transformer效率
本文提出了一种渐进策略通过使用超参数控制两种归一化层的比例逐步将LayerNorm替换为BatchNorm。
论文地址https://arxiv.org/abs/2405.11582
论文代码https://github.com/xinghaochen/SLAB
Introduction
transformer架构最初引入用于自然语言处理任务迅速成为语言模型领域的杰出模型。随着Vision Transformer(ViT)的引入其影响力显著扩展展示了基于transformer的架构的有效性和多样性。这些架构在与卷积神经网络CNNs相比在各种视觉任务中表现出了竞争力的性能基准。由于其强大的性能transformer已成为深度学习中的主流架构。然而transformer架构的计算需求构成了一个重大挑战这主要是由于其注意力机制的二次计算复杂性和LayerNorm组件在线统计计算的必要性。
许多工作致力于提升transformer架构的效率。有的方法试图通过限制自注意机制中token交互的范围来减少计算复杂度例如降采样键和值矩阵、采用稀疏全局注意模式以及在较小的窗口内计算自注意力。与此同时线性注意力作为一种替代策略出现通过将注意力机制分解为线性计算成本来增强计算效率然而在效率和准确性之间取得良好平衡仍然是一个具有挑战性的任务。此外由于LayerNorm在推理过程中额外的计算开销一些探索尝试将BatchNormBN替代transformer中的LayerNormLN比如在前向网络的两个线性层之间添加一个BatchNorm层来稳定训练。然而LayerNorm和BatchNorm的transformer之间仍存在性能差距。
论文的重点是通过深入研究计算效率低下的模块即归一化层和注意力模块来获取高效的transformer架构。首先论文探索了用BatchNorm替换LayerNorm以加速transformer的推理过程。BatchNorm可以降低推理延迟但可能导致训练崩溃和性能下降而LayerNorm可以稳定训练但在推理过程中会增加额外的计算成本。因此论文提出了一种渐进策略通过使用超参数控制两种归一化层的比例逐步将LayerNorm替换为BatchNorm。最初transformer架构由LayerNorm主导随着训练的进行逐渐过渡到纯BatchNorm。这种策略有效地减轻了训练崩溃的风险并且在推理过程中不再需要计算统计信息。除了渐进策略外论文还提出了一种新的BatchNorm重新参数化公式RepBN以增强训练稳定性和整体性能。
此外注意力机制的计算成本对于高效的transformer架构至关重要之前的方法在效率和准确性之间难以取得良好的平衡。因此论文提出了一种简化的线性注意力SLA模块该模块利用ReLU作为核函数结合深度可分卷积来进行局部特征增强。这种注意力机制比之前的线性注意力更高效而且能达到可比较的性能水平。
论文在各种架构和多个基准测试上广泛评估了提出的方法。渐进重新参数化的BatchNorm在图像分类和物体检测任务中表现出强大的性能以更低的推理延迟获得类似的准确性。此外结合渐进RepBN和简化线性注意力模块的SLAB transformer在提高计算效率的同时与Flatten transformer相比达到了竞争性的准确性。例如SLAB-Swin-S在ImageNet-1K上达到了83.6%的Top-1准确率推理延迟为16.2毫秒比Flatten-Swin-S的准确率高出0.1%延迟则减少了2.4毫秒。论文还对提出的方法在语言建模任务上进行了评估获得了可比较的性能和更低的推理延迟。
Preliminaries
给定输入为 个令牌的特征 , 其中 是特征维度, Transformer 块的一般架构可以写成:
其中 计算注意力分数 表示多层感知机 是归一化函数。在 Transformer 块的默认配置中 通常是一个 LayerNorm 操作 是基于 sof tmax 的注意力机制
注意力在 Transformer 中扮演着重要角色。将查询、键和值矩阵表示为 , softmax 注意力首先计算查询和键之间的成对相似性。成对相似性计算导致与查询和键的数量 相关的二次计算复杂度 , 使得 Transformer 在处理具有长序列输入的任务时计算成本昂贵。线性注意力旨在解耦 softmax 函数, 通过适当的近似方法或者用其他核函数先计算 , 计算复杂度变为 , 与查询和键的数量 线性相关。
然而LayerNorm在推理过程中需要统计计算因此占据了不可忽视的延迟部分。因此论文探索利用BatchNorm来构建高效的Transformer模型BatchNorm仅在训练过程中存在并且可以与前置或顺序线性层合并。此外注意力模块对于Transformer至关重要而基于softmax的注意力机制由于其二次计算复杂度而在计算效率上存在问题。因此论文提出了一种简单而高效的注意力形式极大地减少了延迟同时在各种视觉任务上保持了良好的性能。
Methods
论文专注于构建高效的Transformer模型并提出了一系列策略包括逐步替换LayerNormLN为重新参数化的BatchNormBN以及简化的线性注意力SLA模块。所提出的SLAB Transformer模型在与先前方法相比表现出了强大的性能同时具备更高的计算效率。
Progressive Re-parameterized BatchNorm
LayerNorm在训练和推理过程中都需要进行统计量计算因此显著影响了Transformer的运行速度。相比之下BatchNorm在推理过程中可以简单地与线性层合并更适合于高效的架构设计。然而直接在Transformer中使用BatchNorm会导致性能表现不佳。为此论文提出在训练过程中逐步替换LayerNorm为BatchNorm并且还提出了一种受Repvgg启发的新的BatchNorm重新参数化公式以进一步提高性能如图2所示。
Re-parameterized BatchNorm
RepBN公式如下
其中是一个可学习的参数以端到端的方式联合训练。一旦训练完成RepBN可以重新参数化为BatchNorm的一种规范形式。
根据引理 4.1, RepBN 输出的分布由 和 控制, 分别对应于方差和均值。R epBN 可以借助 和 来恢复分布。
同时, 当 时, 相当于跳过了 BatchNorm 。当 时, RepBN 则退化为纯粹的 BatchNorm 。
Progressive LN RepBN
为了促进基于纯粹BN的Transformer模型的训练论文建议在训练过程中逐步过渡从LN到RepBN即
其中, 是一个超参数, 用于控制不同归一化层的输出。通常, 在训练初期 LN 主导架构时, ; 在训练结束时, 为了确保过渡到基于纯粹 BN 的 Transformer, 。我们采用了一个简单而有效的衰减策略来调整 的值:
其中, 表示使用 LayerNorm 进行训练的总步数, 表示当前的训练步数。这种渐进策略有助于减轻训练纯粹基于BN的Transformer的难度从而在各种任务上实现强大的性能表现。 还有一些其他衰减策略可以逐渐减小 的值例如余弦衰减和阶梯衰减。从实验来看线性策略是比较有效且简单的一种方法。
Simplified Linear Attention
注意力模块是Transformer网络中最重要的部分通常表述为
其中, 将输入的标记投影到查询 (query) 、键 (key) 和值 value张量。 表示相似性函数。对于注意力的原始形式, 相似性函数是
这种基于softmax的注意力导致了较高的计算复杂度。近年来有几种方法研究了使用线性注意力来避免softmax计算从而提高Transformer的效率。然而这些方法仍然存在相当复杂的设计并且计算效率不够高。因此论文提出了一种简化的线性注意力SLA
其中表示深度可分离卷积depth-wise convolution。这是一种简单而高效的线性注意力方法因为它通过先计算享受了解耦的计算顺序从而显著减少了复杂度。此外该方法只使用了ReLU函数和深度可分离卷积这两种操作在大多数硬件上都具有良好的计算效率。
这里的整体逻辑跟FLatten Transformer基本一样只是将其提出聚焦函数替换为ReLU函数。这里的效率提升通过摘除softmax计算从而达到先计算 实现的公式7做下乘法结合律ReLU也有保证内积为正数的作用和DWC是补充计算顺序改变带来的性能损失。
为了展示该方法仍然保持特征多样性论文通过可视化注意力图表明了应用了渐进重新参数化批归一化和简化线性注意力SLAB策略的DeiT-T的效果如图3所示。可以看出论文提出的方法仍然保持了较高的排名表明其在捕捉注意力信息方面具有良好的能力。
Experiments #小扎老黄亲密换衣炉边对谈
小扎竟破防爆粗老黄自曝第一批Blackwell已出炉
就在刚刚老黄在SIGGRAPH大会上透露Blackwell的工程样片已在本周正式向全世界发送随后老黄和小扎展开了炉边对话并且亲密换衣说到激动处小扎气得一度爆粗。
惊爆消息来了
刚刚老黄在SIGGRAPH计算机图形会议上透露就在本周英伟达已经开始向全世界发送Blackwell的工程样片了
紧接着主持人Lauren Goode便调侃道没错大家低个头凳子下面就有。
值得一提的是在这款当今最强的AI芯片背后同样也离不开AI——
没有AIHopper将无法问世没有AIBlackwell也无法成为可能。
在他和小扎上演的一场炉边对话中讲到情绪激动时小扎甚至一度忍不住出口爆粗。
由于两位大佬之前的换衣效果实在是一言难尽。
这次小扎专门给老黄送上了定制的「黑皮衣风」棉服。
上身之后的确效果拔群
当然小扎也穿上了老黄仅用了2小时的「二手」皮衣。这可比全新的值钱多了
英伟达的数字「副本」世界
在大会上老黄宣布英伟达构建了世界上首个能够理解基于 OpenUSD语言、几何、材料、物理和空间的生成性AI模型
什么是OpenUSD呢它指的就是Universal Scene Description可以理解为一种通用场景描述。
老黄表示比起AI对文本执行的操作更令人兴奋的是我们可以对图像执行同样的操作。
比如英伟达创建的Edify AI模型就是一个文本到2D的基础模型。
对于品牌来说它可以创造出可口可乐、汽车、奢侈品等等然而控制提示就是一件困难的事。
这是因为词语的纬度非常低它在内容上极其压缩但同时又非常不精确。
为此英伟达创造了一种方法——创造另一个模型实现控制和调整它与更多条件的对齐。
使用Omniverse可以组合所有这些多模态数据和内容无论是3DAI动画还是材质。
我们可以改变它的姿势、位置总之改变我们想要的任何东西。
使用Omniverse中的条件化提示就跟检索增强生成一样可以理解为一种3D增强生成。
这样我们就可以按照喜欢的方式生成图像了。
接下来WPP用Shutterstock与世界知名品牌完成的作品直接震撼了全场。
在一个空房间里给我建一张桌子周围摆着椅子在一个繁忙的餐馆里。
在晨光中给我建一张摆着玉米卷和一碗莎莎酱的桌子。
在一条空旷的道路上给我建一辆车周围是树木靠近一座现代房屋。
在一片空旷的田野里给我建一棵树。
在所有方向给我建数百棵这样的树。
让这片树林有灌木丛和藤蔓悬挂其间。
给我建一片巨大的热带雨林里面有奇异的鲜花和阳光射线。
Omniverse现在能够理解文本到USD的转换。它能够理解文本并拥有一个语义数据库因此可以搜索所有的3D对象。
因此小女孩可以描绘自己想以什么方式填充3D树完成之后3D场景就会进入生成式AI模型将其转化为照片级真实感的模型。
从此越来越多的生成式AI会出现在Omniverse中帮人们创建这些模拟或数字孪生。
比如下面这个数字AI将使每家公司都有客户服务。
在当下客户服务是由人类完成的但在未来AI将参与其中。
客户服务会连接到一个数字人前端也就是一个IO。这个IO可以说话还能与我们眼神交流。
各类AI都可以连接到这个数字人甚至数字人还可以连接到英伟达的检索增强生成客户服务AI上。
NIM服务
这次大会上英伟达推出了一套全新的NIM微服务。
NIM专为不同的工作流量身定制包括OpenUSD、3D 建模、物理、材料、机器人、工业数字孪生和物理AI。
在AI和图形领域英伟达推出了专为生成物理AI应用程序设计的全新OpenUSD NIM微服务。
这个工作流包括用于机器人模拟等的新NIM微服务来加速人形机器人的开发。
「三体」创造机器人
老黄预言下一波AI是物理AI。
如果机器人技术想进步就需要先进的AI和逼真的虚拟世界而在部署下一代人形机器人之前我们需要对AI进行训练。
机器人技术需要三台计算机一台用于训练AI一台在物理精确的模拟中测试AI另一台位于机器人本身内部可以学习如何优化机器人。
也就是说第三台AI是实际运行AI 的计算机。
为此英伟达创造了三台计算机。
没有AI就没有H100/H200和B100
从1990年代开始的英伟达历史中真正的DNA就在于计算机图形学。
计算机图形学也一路把英伟达带到了今天的位置。
这幅图中展示了计算机行业一些重要的里程碑包括IMB 360系统、犹他茶壶、光追、可编程着色等等
1993年英伟达成立。八年后他们发明了第一个可编程着色GPU很大程度上推动了英伟达的发展历程。
可以说英伟达所做的一切背后的核心就是加速计算。他们坚信如果创建一种计算模型来增强通用计算就可以解决普通计算机无法解决的问题。
首选的领域就是计算机图形学。他们赌对了。
将计算机图形学应用到当时非主流的领域——3D图形视频游戏直接推动了英伟达的飞轮。
此后他们花了很长时间让CUDA无处不在然后在2012年就仿佛《星际迷航》中一般英伟达第一次接触了AlexNet。
在2012年那是一个爆炸性的时刻AlexNet在计算机视觉上取得了惊人的突破。它的核心——深度学习如此深刻再也不需要工程师们提供输入后去想象输出的样子了。
2016年英伟达推出了第一台为深度学习打造的计算机DGX-1被马斯克看上随后产品被交付给当时名不见经传的OpenAI。
随后RTX、DLSS被发明出来。
然后就是ChatGPT的诞生。
未来人人都有AI助手
如今我们已经学会用AI学习一切不仅仅是单词还有图像、视频、3D、化学物质、蛋白质、物理学、热力学、流体动力学、粒子物理学等等。
我们理解了所有这些不同模态的意义。
在老黄看来基于视觉计算的生成式AI革命正在增强人类的创造力。
我们真正处于革命性的时刻迈向软件3.0的时代——没有哪个行业能逃过AI的影响
老黄预言每个人都会有一个AI助手每家公司、公司内的每一项工作都将得到AI的帮助
加速计算让能源问题有解
虽然生成式AI有望提高人类生产力但AI基础设施的耗能问题却是困扰整个地球的大问题。
ChatGPT的一次搜索相当于10次谷歌搜索的电量。
数据中心消耗了全球总体能源的1%到2%甚至可能在十年内达到6%。
怎么办老黄有解。
他表示加速计算技术有望使计算更节能。
「加速计算可以帮我们节省大量能源它能节省20倍、50倍并且执行相同的处理」老黄说。
「作为一个社会我们要做的第一件事就是加速我们能做到的每一项应用这减少了全世界的能源使用量。」
这也是为什么Blackwell备受期待因为它使用同样的能量却大大加速了应用程序。
而且它还会越来越便宜。
老黄强调要记住生成式AI的目标并不是训练而是推理。理想情况下推理可以为我们创建预测天气的新模型、预测新材料、优化供应链等等。
要知道数据中心并不是唯一消耗能源的地方。全球数据中心只占总计算量的40%60%的能源消耗在网上移动着电子、比特和字节。
因此生成式AI将减少网上的能源消耗因为不需要去检索信息我们可以在现场直接生成了。
而且就在刚刚英伟达在GCP中部署了GPU来运行Pandas。
这个世界上最领先的数据科学平台直接把速度从50提升到了100倍超过了通用计算。
在过去10到12年的时间里我们已经将深度学习的速度提升了100万倍成本和能耗降低了100万倍这就是为什么LLM能够诞生。
不过英伟达还会通过设计新的处理器、新系统、Tensor核心GPU和NVLink交换机结构给AI带来新的创新。
老黄和小扎的炉边对谈
今年SIGGRAPH上两位CEO的炉边对谈让很多人期待已久。用小扎本人的话来说「两个行业内最资深的创始人」究竟会碰撞出怎样的火花
下一波浪潮
不出意外的是这两位「青梅煮酒」的英雄都各自分享了自己的预判大谈未来的技术发展趋势从GenAI到Agent再到小扎始终念念不忘的「元宇宙」。
老黄表示GenAI的技术力量也让他感到震撼「我不记得有哪项技术以如此迅猛的速度影响了消费者、企业、工业和学界而且跨越了从气候技术到生物技术再到物理科学的所有不同领域。」
小扎也表示GenAI很可能会重塑Meta的各类社交媒体软件。
曾经这些产品的核心——推荐系统仅仅是将感兴趣的内容推送给用户。
但GenAI将不再局限于已有内容不仅能协助创作者还会为用户创建即时内容或综合现有内容进行生成。
关于Agent的发展两人似乎也有类似的观点。
在之前的演讲中老黄就明确表示「未来每个人都将有自己的AI助手」。
这场对谈中小扎也表述了类似的愿景。他正在为Meta规划AI助手和AI Studio产品让每个人都能为不同用途创建自己的Agent。
未来每个企业都将拥有自己的AI正如今天所有公司都有自己的社交媒体和邮箱账号一样。
他们口中的「AI助手」究竟要达到何种程度的「智能」
我们目前看到的Llama 3仅仅是一个「聊天机器人」般的语言模型只能对人类的提问做出响应。但小扎希望可以给AI赋予「意图」intent。
老黄则将其描述为「规划能力」能像人类一样在脑海中形成「决策树」进而指导行为。
他甚至更大胆地预测这种AI助手的成本仅有每小时10美元却能大大提高工程师们的工作表现。「如果你还没雇佣AI就赶紧行动起来」
对于Meta最核心也是最独特的AR/VR技术小扎的蓝图也相当精确充分体现了他的强迫症人格。
据老黄爆料小扎切西红柿都有毫米级的精度而且每片西红柿都不能相互接触。
去年9月Meta联合雷朋推出了新款智能眼镜配备音频设备和摄像头让用户直接从双眼视角拍照或将眼镜中看到的视野直接直播到Facebook或Instagram上并集成了对话助手Meta AI。
小扎表示基于雷朋眼镜现在的情况定价300美元的无显示屏AI眼镜将成为非常热门的产品。
根据他的预测智能眼镜在未来会成为一种类似手机的设备每个戴眼镜的人都将带上智能眼镜全世界超过10亿人。
接下来几年时间Meta还将推出有全息AR功能的眼镜虽然成本依旧很高但将会是可行的产品。
与智能眼镜不同混合现实头显更类似于工作站或游戏机不方便携带但算力更强能为用户提供更加沉浸式的体验。
而且随着全息AR技术的发展「虚拟会议」即将成为现实。
不同于Zoom平台上的头像或视频每个人都将有自己的全息影像即使身处不同的物理空间也能让全息图打造的「虚拟人」在同一空间内合作交互。
开源是前进之路
提到Meta他们一贯实行的「开源」策略也是不得不谈的要点。
老黄就十分赞赏这种策略他表示Llama 2可能是去年AI领域最重要的事件加上PyTorch和最新发布的Llama 3.1Meta已经构建了一整个生态系统。
但小扎表示他们走上开源之路也是一种「随机应变」。
在很多赛道上尤其是分布式计算系统和数据中心方面Meta的起跑线其实落后于其他公司因此团队想到了开源尤其是开放计算open compute。
没有想到的是这种权宜之计反而成为「弯道超车」的关键策略。
正是开源让Meta发布的产品成为行业标准整个供应链也围绕其建立。通过将项目开源Meta甚至还节省了数十亿美元。
比如Meta进入GPU领域的时间其实晚于大多数公司但他们目前运营的GPU超算集群规模几乎超过了所有竞争对手。
当然这背后少不了老黄的大力支持毕竟Meta的60万张GPU也都出自英伟达之手。
开源虽然能推动这个社区和行业的进步但小扎也很诚实地表示开源不是做慈善我们选择这种策略并不是因为有一颗无私奉献的心。
更主要的目的是希望让正在构建的产品登峰造极成为最好的东西。
PyTorch就是最典型的例子全世界的开发人员包括英伟达的两三百名工程师都在帮这个开源框架找bug、做优化构成了老黄口中的「PyTorch的工程之山」。
虽然小扎自己也承认开源是有私心的但谈到「封闭」平台时还是克制不住地情绪激动。全场唯一脏话就出自这个话题。
虽然Meta坐拥一众王者级社交软件但这些应用程序都需要通过竞争对手的平台进行分发尤其是苹果应用商店以及谷歌的安卓系统。
让小扎很气恼的是他曾经有过很多产品创意但受限于这些移动平台的各种限制最终都无法成行。
移动互联网时代的这种极度依赖平台的特点与曾经PC时代的开放完全不同这让小扎相当怀念网页端的Facebook。
因此他满怀信心地表示我们正在塑造下一代计算平台即混合现实技术其中开源软件将重新拥有更大的价值。
下一代平台和生态系统将更具有开放性、包容性类似之前的Windows或安卓生态而非完全封闭的苹果。
这种「让开源重新伟大」的雄心让人不禁想起Llama 3.1发布时他的比喻——Llama 3.1是这个时代的Linux。
CEO不好做
整个对谈过程中两人颇有惺惺相惜之感而且经常谈到CEO这份职业的艰难。
时年61岁、穿着皮夹克的Jensen甚至一脸严肃地自比娇花「我们是CEO像娇嫩的花朵需要很多支持。」
小扎甚至紧跟着接了一句「我们现在相当憔悴了」。
这种感慨也许来自于两位资深创始人曾经和公司共同经历的风雨。
在小扎看来老黄当年顶着不被看好的压力硬要把计算机做成「超级巨兽」让英伟达终成行业传奇
在老黄看来小扎带着Meta多次转型从PC端到移动端从社交媒体到VR/AR和AI的研究。
对谈的最后老黄直言两人的共同之处「我知道那样做转型有多难我们两个都曾经被狠狠打击过但这就是成为先锋和创新所必备的。」
参考资料
https://www.youtube.com/watch?vH0WxJ7caZQU
https://www.youtube.com/watch?vw-cmMcMZoZ4 #HourVideo
空间智能版ImageNet来了李飞飞吴佳俊团队出品,来自斯坦福李飞飞吴佳俊团队 HourVideo一个用于评估多模态模型对长达一小时视频理解能力的基准数据集包含多种任务。 通过与现有模型对比揭示当前模型在长视频理解上与人类水平的差距。
2009年李飞飞团队在CVPR上首次对外展示了图像识别数据集ImageNet它的出现极大推动计算机视觉算法的发展——懂CV的都是知道这里面的门道有多深。
现在随着多模态迅猛发展团队认为“现有的视频基准测试大多集中在特定领域或短视频上”并且“这些数据集的平均视频长度较短限制了对长视频理解能力的全面评估”。
于是空间智能版ImageNet应运而生。
HourVideo包含500个来自Ego4D数据集的第一人称视角视频时长在20到120分钟之间涉及77种日常活动。
评测结果表示人类专家水平显著优于目前长上下文多模态模型中最厉害的Gemini Pro 1.585.0%对37.3%。
在多模态能力上大模型们还任重而道远。
HourVideo如何炼成
之所以提出HourVideo是因为研究人员发现目前长视频理解越来越重要而现有评估benchmark存在不足。
多模态越来越卷人们期待AI被赋予autonomous agents的类似能力而从人类角度来看由于人类具备处理长时间视觉处理的能力因此能在现实视觉中感知、计划和行动。
因此长视频理解对实现这一目标至关重要。
而当前的多模态评估benchmark主要还是集中在评测单张图像或短视频片段几秒到三分钟对长视频理解的探索还有待开发。
不可否认的是AI评估长视频理解面临诸多挑战譬如要设计任务、避免通过先验知识或简短片断回答等。
因此团队提出HourVideo。
这是一个为长视频理解而设计的基准数据集。
为了设计出需要长期理解的任务团队首先提出了一个新的任务对应套件包含总结、感知回忆、跟踪、视觉推理空间、时间、预测、因果、反事实和导航房间到房间、对象检索任务共18个子任务。
其中总结任务要求模型对视频中的关键事件、主要交互等进行概括性描述例如总结出脖子上挂了个相机的人在超市中有什么关键交互行为。
感知任务由两部分构成
一个是回忆任务包括事实回忆比如脖子上挂了个相机的人在超市拿起的乳制品和序列回忆比如那个人在超市称完西红柿过后做了什么以及对时间距离的判断比如吃了多久的披萨才扔掉盒子。
还有一个是跟踪任务主要用来识别脖子上挂了个相机的人在特定场景比如超市、药店中互动的独特个体。
接下来是视觉推理任务分为空间推理和时间推理。
空间推理负责判断物体之间的空间关系、空间接近度如微波炉与冰箱或水槽相比是否更近以及空间布局如选择正确描绘脖子上挂相机的人的公寓的布局图。
时间推理则包括对活动持续时间的比较、事件发生频率的判断、活动的先决条件、预测如洗完衣服后最可能做的活动、因果关系如第二次离开车库的原因以及反事实推理如用烤箱做土豆泥会怎样。
导航任务包含了房间到房间的导航、对象检索导航。
以上每个任务有精心设计的问题原型以确保正确回答问题需要对长视频中的多个时间片段进行信息识别和综合从而有效测试模型的长期理解能力。
与此同时研究人员通过pipeline来生成了HourVideo数据集。
第一步视频筛选。
团队从Ego4D数据集中手动审核1470个20到120分钟的视频让5位人类专家选择了其中500个视频
至于为啥要从Ego4D中选呢一来是其以自我为中心的视角与autonomous agents和助手的典型视觉输入非常一致二来是它具有广泛的视觉叙述有助于创建多样化的题三来Ego4D的访问许可非常友好。
第二步候选MCQ生成。
这需要在长视频中跨多个时间片段进行信息分析和合成。
具体来说研究人员以20分钟为间隔分割了视频提取信息转化为结构化格式供大模型处理。最终一共开发了25个特定任务的prompts。
第三步LLM优化与人工反馈。
在这个阶段团队实现了一个人工反馈系统7名经验丰富的人员人工评估每个问题的有效性、答案准确性、错误选项合理性。最终收集了400多个小时的人工反馈然后设计prompt自动优化 MCQ₂得到 MCQ₃。
第四步盲选。
这一阶段的目标是消除可以通过大模型先验知识的问题或者消除那些可以在不用视频中任何信息就可以回答的问题。
团队用两个独立的大模型——GPT-4-turbo和GPT-4对MCQ₃进行盲筛确保剩余 MCQ₄高质量且专门测试长视频语言理解。
第五步也是最后一步专家优化。
这一步是用来提升MCQ₄质量将宽泛问题精确化经此阶段得到高质量 MCQ₅。
4个专家干的事be like把 “挂着相机的人把钥匙放在哪里了” 精确成“挂着相机的人购物回家后把自行车钥匙放在哪里了”
如上pipeline中研究图纳队使用了GPT-4来遵循复杂的多步骤指令同时还使用了CoT提示策略。
此外pipeline中涉及大模型的所有阶段的问题被设为0.1。
据统计HourVideo涵盖77种日常生活场景包含500个Ego4D视频视频时长共381个小时、平均时长45.7分钟其中113个视频时长超过1小时。
每个视频有约26个高质量五选一题共计12976个问题。
除因果、反事实和导航任务外问题在任务套件中均匀分布。
最好表现仍远低于人类专家水平
在实验评估方面HourVideo采用五选多任务问答MCQ 任务以准确率作为评估指标分别报告每个任务以及整个数据集的准确率。
由于防止信息泄露是评估长视频中的MCQ时的一个重要挑战——理想情况下每个MCQ应独立评估但这种方法计算成本巨高且十分耗时。
因此实际评估中按任务或子任务对问题进行分批评估对于预测任务提供精确的时间戳以便对视频进行有针对性的剪辑从而平衡计算成本和评估准确性。
研究团队比较了不同的多模态模型在零镜头设置下理解长视频的性能。
主要评估了三类模型所有这些模型都在一个通用函数下运行
盲LLM
指是指在评估过程中不考虑视频内容仅依靠自身预先训练的知识来回答问题的大型语言模型。
实验中以GPT-4为代表。它的存在可以揭示模型在多大程度上依赖于其预训练知识而不是对视频中实际视觉信息的理解。
苏格拉底模型
对于大多数当前的多模态模型直接处理非常长的视频存在困难。
因此采用Socratic模型方法将视频总时长为t分钟分割成1分钟的间隔每个间隔独立加字幕然后将这些字幕聚合形成一个全面的基于语言的视频表示并与通用任务无关的提示一起作为输入进行长视频问答。
实验中分别使用GPT-4和LLaVA- NEXT-34-DPO 为视频字幕生成器并最终使用GPT-4进行实际问题回答。
原生多模态模型
像Gemini 1.5 Pro这样的原生多模态模型在多模态数据包括音频、视频、图像和文本上联合训练能够处理非常长的上下文长度*2M 适合直接对HourVideo进行端到端评估。
为了与模型性能进行对比实验人员从基准数据集中选取了14个视频涵盖18种场景包括手工制作/绘画、烹饪、建筑/装修、园艺、清洁/洗衣和庭院工作等。
然后邀请了3位人类专家对上述总时长11.2小时的视频内容进行进行评估共涉及213个MCQ。
为确保评估的公正性参与评估的人类专家未参与过这些视频的早期注释工作。
最终人类专家在评估中的准确率达到了85.0% 。
而盲LLM的准确率为19.6%Socratic模型准确率略高原生多模态模型准确率最高达到了37.3%仍然远低于人类专家水平。
此外独立评估每个MCQ与按任务级别评估相比性能下降2.1%但成本增加3倍以上证明了任务级评估方法的效率和有效性。
最后团队表示未来计划扩展基准测试包括更多样化的视频来源如体育和YouTube视频纳入音频模态支持并探索其他感官模态。
同时强调在开发模型时需考虑隐私、伦理等问题。
团队成员
HourVideo项目来自斯坦福李飞飞和吴佳俊团队。
论文共同一作是Keshigeyan Chandrasegaran和Agrim Gupta。
Keshigeyan Chandrasegaran是斯坦福大学计算机科学博士二年级学生从事计算机视觉和机器学习研究导师是李飞飞和斯坦福视觉与学习实验室SVL联合主任胡安·卡洛斯·尼贝莱斯。
共同一作Agrim Gupta是斯坦福大学计算机科学专业的博士生2019年秋季入学同样是李飞飞的学生。
此前他曾在微软、DeepMind有Meta的全职经历也在Google做过兼职。2018年时他就跟随李飞飞一同在CVPR上发表了论文。
目前Agrim的Google Scholar论文被引用量接近6400次。
李飞飞是大家熟悉的AI教母AI领域内最具影响力的女性和华人之一。
她33岁成为斯坦福计算机系终身教授44岁成为美国国家工程院院士现任斯坦福以人为本人工智能研究院HAI院长。
计算机视觉领域标杆成果ImageNet亦是由她一手推动。
此前李飞飞也曾短暂进入工业界出任谷歌副总裁即谷歌云AI首席科学家。她一手推动了谷歌AI中国中心正式成立这是Google在亚洲设立的第一个AI研究中心。并带领谷歌云推出了一系列有影响力的产品包括AutoML、Contact Center AI、Dialogflow Enterprise等。
今年李飞飞宣布创办空间智能公司World Labs公司成立不到4个月时间估值突破10亿美元。
所谓空间智能即“视觉化为洞察看见成为理解理解导致行动”。
吴佳俊现任斯坦福大学助理教授隶属于斯坦福视觉与学习实验室SVL和斯坦福人工智能实验室SAIL。
他在麻省理工学院完成博士学位本科毕业于清华大学姚班曾被誉为“清华十大学神”之一。
同时他也是李飞飞创业公司World Labs的顾问。
参考链接
[1]https://arxiv.org/abs/2411.04998v1
[2]https://www.worldlabs.ai/team
[3]https://keshik6.github.io/ #梳理下Flash Attention的dispatch逻辑
本文分析了Flash Attention在不同场景下的内核调度逻辑特别关注了在解码阶段何时会使用split_kv实现并探讨了影响这一决策的因素如K序列的最大长度、注意力头数和GPU的流处理器数量。
1. 前言
这篇文章来源是当运行下面的对HuggingFace Qwen2.5-7B-Instruct模型使用Flash Attention的代码时使用Nsight System工具抓取的kernel trace会发现在prefill和decode阶段Flash Attention调用了不同的kernel并且decoding的Flash Attention kernel使用了split_kv的实现。然后如果把下面代码中max_new_tokens改成64我发现在Nsight System工具抓取的kernel trace中decode阶段的Flash Attention kernel又变成了和prefill阶段一样的kernel并没有使用split_kv的实现。这篇文章就尝试跟踪下Flash Attention的dispatch逻辑弄清楚什么情况下decode阶段的Flash Attention kernel会使用split_kv的实现(split_kv的实现也被叫作Flash Decoding专用大模型的Decoding阶段)。
# /opt/nvidia/nsight-systems/2024.5.1/bin/nsys profile --trace-fork-before-exectrue --cuda-graph-tracenode -o hf_qwen2.5_7b_flash_attn python3 debug.py
import os
os.environ[CUDA_VISIBLE_DEVICES] 0
os.environ[TOKENIZERS_PARALLELISM] false import nvtx
import torch from transformers import AutoModelForCausalLM, AutoTokenizer model_name /mnt/bbuf/Qwen2.5-7B-Instruct model AutoModelForCausalLM.from_pretrained( model_name, torch_dtypeauto, device_mapauto, trust_remote_codeTrue
)
tokenizer AutoTokenizer.from_pretrained(model_name, trust_remote_codeTrue) prompt 帮我计划一次去北京的旅行我想明年春天出发大概五天的行程。 model_inputs tokenizer(prompt, return_tensorspt).to(model.device) for i in range(1): with nvtx.annotate(fstep{i}, colorblue): generated_ids model.generate( **model_inputs, max_new_tokens512 ) generated_ids [ output_ids[len(input_ids):] for input_ids, output_ids in zip(model_inputs.input_ids, generated_ids)
] response tokenizer.batch_decode(generated_ids, skip_special_tokensTrue)[0]
print(response) 这张图是max_new_tokens512时prefill和decode阶段的Flash Attention kernel的trace。红色框表示prefill阶段调用的Flash Attention kernel绿色框表示decode阶段调用的Flash Attention kernel。可以看到prefill阶段调用了flash_fwd_kerneldecode阶段调用了flash_fwd_splitkv_kernel和flash_fwd_splitkv_combine_kernel两种kernel。 这张图是max_new_tokens64时prefill和decode阶段的Flash Attention kernel的trace。可以看到两个阶段都调用了同一个flash_fwd_kernel。
为什么产生了这种差别什么情况下decode阶段的Flash Attention kernel会使用split_kv的实现我们需要深入看一下Flash Attention的相关Dispatch逻辑。
2. Qwen2是如何访问Flash Attention API的
下面是 HuggingFace Qwen2 模型 Qwen2FlashAttention2 模块的实现我们可以从这个代码中看到 flash attention 的 API 是如何被调用的。这里调用的 _flash_attention_forward 实际上又是调用了 flash-attention 库(https://github.com/Dao-AILab/flash-attention)中的 flash_attn_varlen_func api这个api是flash attention库中用来处理Attention前向计算的核心函数并且可以从名字看出来这个api还支持可变长的多个序列的Attention计算。
class Qwen2FlashAttention2(Qwen2Attention): # ... def forward( self, hidden_states: torch.Tensor, # 输入隐藏状态 attention_mask: Optional[torch.Tensor] None, # 注意力mask position_ids: Optional[torch.LongTensor] None, # 位置编码id past_key_value: Optional[Cache] None, # KV缓存 output_attentions: bool False, # 是否输出注意力权重 use_cache: bool False, # 是否使用KV缓存 cache_position: Optional[torch.LongTensor] None, # 缓存位置 position_embeddings: Optional[Tuple[torch.Tensor, torch.Tensor]] None, # 位置编码,在v4.46中将成为必需 ): # 获取输入维度 bsz, q_len, _ hidden_states.size() # QKV投影 query_states self.q_proj(hidden_states) key_states self.k_proj(hidden_states) value_states self.v_proj(hidden_states) # 重塑维度以适应多头注意力 query_states query_states.view(-1, self.num_heads*self.head_dim) key_states key_states.view(-1, self.num_key_value_heads*self.head_dim) value_states value_states.view(-1, self.num_key_value_heads*self.head_dim) # 应用旋转位置编码(RoPE) query_states, key_states self.rotary_emb(position_ids, query_states, key_states) # 重塑维度为[batch_size, num_heads, seq_len, head_dim] query_states query_states.view(bsz, q_len, self.num_heads, self.head_dim).transpose(1, 2) key_states key_states.view(bsz, q_len, self.num_key_value_heads, self.head_dim).transpose(1, 2) value_states value_states.view(bsz, q_len, self.num_key_value_heads, self.head_dim).transpose(1, 2) # 处理KV缓存 if past_key_value is not None: cache_kwargs {cache_position: cache_position} # RoPE模型特有的参数 key_states, value_states past_key_value.update(key_states, value_states, self.layer_idx, cache_kwargs) # 如果KV头数小于注意力头数,需要重复KV key_states repeat_kv(key_states, self.num_key_value_groups) value_states repeat_kv(value_states, self.num_key_value_groups) dropout_rate 0.0 if not self.training else self.attention_dropout # 处理数据类型转换 input_dtype query_states.dtype if input_dtype torch.float32: if torch.is_autocast_enabled(): target_dtype torch.get_autocast_gpu_dtype() elif hasattr(self.config, _pre_quantization_dtype): target_dtype self.config._pre_quantization_dtype else: target_dtype self.q_proj.weight.dtype logger.warning_once( f输入隐藏状态似乎被静默转换为float32,这可能与embedding或layer norm层被上采样到float32有关。 f我们会将输入转回{target_dtype}。 ) query_states query_states.to(target_dtype) key_states key_states.to(target_dtype) value_states value_states.to(target_dtype) # 重塑维度以适应Flash Attention query_states query_states.transpose(1, 2) key_states key_states.transpose(1, 2) value_states value_states.transpose(1, 2) # 处理滑动窗口注意力 if ( self.config.use_sliding_window and getattr(self.config, sliding_window, None) is not None and self.layer_idx self.config.max_window_layers ): sliding_window self.config.sliding_window else: sliding_window None # 调用Flash Attention前向传播 attn_output _flash_attention_forward( query_states, key_states, value_states, attention_mask, q_len, position_idsposition_ids, dropoutdropout_rate, sliding_windowsliding_window, is_causalself.is_causal, use_top_left_maskself._flash_attn_uses_top_left_mask, ) # 重塑输出并应用输出投影 attn_output attn_output.reshape(bsz, q_len, self.hidden_size).contiguous() attn_output self.o_proj(attn_output) if not output_attentions: attn_weights None return attn_output, attn_weights, past_key_value
这里的代码省略掉了类的相关初始化在forward函数中涉及到ropekv cache更新reshape输入以适应Flash Attention的输入格式以及调用Flash Attention以及应用输出投影等等Attention计算的细节。
3. Flash Attention单独的调用例子
这里来关注一下使用 flash_attn_varlen_func 这个 api 的单独例子。由于它可以支持多个不同的序列所以这里我们用2个序列来调用一下我写了一个测试脚本如下
import torch
import math from flash_attn import flash_attn_varlen_func # 朴素实现的缩放点积注意力函数
# Efficient implementation equivalent to the following:
def scaled_dot_product_attention(query, key, value, attn_maskNone, dropout_p0.0, is_causalFalse, scaleNone) - torch.Tensor: # 调整输入张量的维度顺序 query query.transpose(0, 1) # [nheads, seqlen, headdim] key key.transpose(0, 1) # [nheads, seqlen, headdim] value value.transpose(0, 1) # [nheads, seqlen, headdim] L, S query.size(1), key.size(1) scale_factor 1 / math.sqrt(query.size(-1)) if scale is None else scale attn_bias torch.zeros(L, S, dtypequery.dtype, devicequery.device) if is_causal: assert attn_mask is None temp_mask torch.ones(L, S, dtypetorch.bool, devicequery.device).tril(diagnotallow0) attn_bias.masked_fill_(temp_mask.logical_not(), float(-inf)) attn_bias attn_bias.to(query.dtype) if attn_mask is not None: if attn_mask.dtype torch.bool: attn_bias.masked_fill_(attn_mask.logical_not(), float(-inf)) else: attn_bias attn_mask # 调整注意力计算以适应多头 attn_weight torch.matmul(query, key.transpose(-2, -1)) * scale_factor # [nheads, L, S] attn_weight attn_weight attn_bias.unsqueeze(0) # 广播 attn_bias 到所有头 attn_weight torch.softmax(attn_weight, dim-1) if dropout_p 0.0: attn_weight torch.nn.functional.dropout(attn_weight, pdropout_p, trainingTrue) output torch.matmul(attn_weight, value) # [nheads, L, headdim] return output.transpose(0, 1) # 返回 [L, nheads, headdim] # 设置随机种子以确保结果可复现
torch.manual_seed(0) # 参数设置
batch_size 2
seq_lengths [128, 256] # 两个序列的长度
nheads 16
headdim 32
dropout_p 0.0
causal True # 是否使用因果性掩码
scale None # 缩放因子默认为 1 / sqrt(headdim) # 为每个序列生成随机的 q, k, v 张量
qs []
ks []
vs []
for seqlen in seq_lengths: q torch.randn(seqlen, nheads, headdim, requires_gradTrue, dtypetorch.bfloat16, devicecuda) # (L, nheads, headdim) k torch.randn(seqlen, nheads, headdim, requires_gradTrue, dtypetorch.bfloat16, devicecuda) v torch.randn(seqlen, nheads, headdim, requires_gradTrue, dtypetorch.bfloat16, devicecuda) qs.append(q) ks.append(k) vs.append(v) # 将所有序列的 q, k, v 拼接起来
q_total torch.cat(qs, dim0) # (total_q, nheads, headdim)
k_total torch.cat(ks, dim0)
v_total torch.cat(vs, dim0) # 计算累积序列长度用于索引
cu_seqlens_q torch.zeros(batch_size 1, dtypetorch.int32, devicecuda)
cu_seqlens_q[1:] torch.cumsum(torch.tensor(seq_lengths, dtypetorch.int32), dim0)
cu_seqlens_k cu_seqlens_q.clone() print(cu_seqlens_q: , cu_seqlens_q) # 最大序列长度
max_seqlen_q max(seq_lengths)
max_seqlen_k max(seq_lengths) # 任意传入一个softmax_scale
softmax_scale 0.2 # 调用 flash_attn_varlen_func 函数
out_flash flash_attn_varlen_func( q_total, k_total, v_total, cu_seqlens_q, cu_seqlens_k, max_seqlen_q, max_seqlen_k, dropout_pdropout_p, softmax_scalesoftmax_scale, causalcausal,
) # 使用朴素实现对每个序列进行计算并将输出拼接起来
outputs_naive []
for i in range(batch_size): q qs[i] # (L_i, nheads, headdim) k ks[i] v vs[i] out scaled_dot_product_attention( q, k, v, attn_maskNone, dropout_pdropout_p, is_causalcausal, scalesoftmax_scale ) # 输出形状为 (L_i, nheads, headdim) outputs_naive.append(out) # 将朴素实现的输出拼接起来
out_naive torch.cat(outputs_naive, dim0) # (total_q, nheads, headdim) print(out_naive st: , out_naive.flatten()[:10])
print(out_flash st: , out_flash.flatten()[:10])
print(*20)
print(out_naive en: , out_naive.flatten()[-10:])
print(out_flash en: , out_flash.flatten()[-10:]) # 比较两个实现的输出是否一致
assert torch.allclose(out_flash, out_naive, atol1e-2), Outputs do not match! print(测试通过)
这个测试是可以通过的相信通过上面2个对上层接口调用的例子可以让我们对Flash Attention的接口调用有比较清晰的认识。下面我们可以关注一下Flash Attention这个借口的实现我们不需要深入到cuda实现中只需要把握一下整体的调用逻辑搞清楚文章开头抛出的问题即可。
4. flash_attn_interface.py中的上层接口
flash-attention 库中使用 cuda 实现了Flash Attention的计算然后通过 Torch Binding 将varlen_fwd这个接口暴露给Python而flash_attn_varlen_func则是对varlen_fwd的进一步封装我们可以在 https://github.com/Dao-AILab/flash-attention/blob/main/flash_attn/flash_attn_interface.py 中查看到flash_attn_varlen_func这个接口的实现。去掉了反向相关的逻辑如下所示
def _flash_attn_varlen_forward( q: torch.Tensor, k: torch.Tensor, v: torch.Tensor, cu_seqlens_q: torch.Tensor, # Q序列的累积长度 cu_seqlens_k: torch.Tensor, # K序列的累积长度 max_seqlen_q: int, # Q序列的最大长度 max_seqlen_k: int, # K序列的最大长度 dropout_p: float, # dropout概率 softmax_scale: float, # softmax缩放因子 causal: bool, # 是否使用因果掩码 window_size_left: int -1, # 滑动窗口左侧大小 window_size_right: int -1, # 滑动窗口右侧大小 softcap: float 0.0, # softmax的上限值 alibi_slopes: Optional[torch.Tensor] None, # ALiBi位置编码的斜率 return_softmax: bool False, # 是否返回softmax结果 block_table: Optional[torch.Tensor] None, # 分块表 leftpad_k: Optional[torch.Tensor] None, # K序列左侧填充 seqused_k: Optional[torch.Tensor] None, # K序列使用的长度
) - Tuple[torch.Tensor, torch.Tensor, torch.Tensor, torch.Tensor]: # 确保输入张量是连续的内存布局 q, k, v [maybe_contiguous(x) for x in (q, k, v)] # 调用CUDA实现的前向传播函数 out, softmax_lse, S_dmask, rng_state flash_attn_cuda.varlen_fwd( q, k, v, None, # 原始掩码矩阵(未使用) cu_seqlens_q, cu_seqlens_k, seqused_k, leftpad_k, block_table, alibi_slopes, max_seqlen_q, max_seqlen_k, dropout_p, softmax_scale, False, # 未使用的参数 causal, window_size_left, window_size_right, softcap, return_softmax, None, # 随机数生成器状态(未使用) ) return out, softmax_lse, S_dmask, rng_state # FlashAttnVarlenQKVPackedFunc类实现了PyTorch的自动微分接口
class FlashAttnVarlenQKVPackedFunc(torch.autograd.Function): staticmethod def forward( ctx, # 上下文对象,用于保存反向传播需要的信息 qkv, # 打包的QKV张量 cu_seqlens, # 累积序列长度 max_seqlen, # 最大序列长度 dropout_p, # dropout概率 softmax_scale, # softmax缩放因子 causal, # 是否使用因果掩码 window_size, # 滑动窗口大小 softcap, # softmax上限值 alibi_slopes, # ALiBi位置编码斜率 deterministic, # 是否确定性计算 return_softmax, # 是否返回softmax结果 ): # 如果未指定缩放因子,使用默认的1/sqrt(head_dim) if softmax_scale is None: softmax_scale qkv.shape[-1] ** (-0.5) # 分离Q、K、V并detach,避免建立反向图 q, k, v qkv[:, 0].detach(), qkv[:, 1].detach(), qkv[:, 2].detach() # 获取原始head size head_size_og q.size(2) # 如果head size不是8的倍数,进行padding if head_size_og % 8 ! 0: q torch.nn.functional.pad(q, [0, 8 - head_size_og % 8]) k torch.nn.functional.pad(k, [0, 8 - head_size_og % 8]) v torch.nn.functional.pad(v, [0, 8 - head_size_og % 8]) # 调用前向计算函数 out_padded, softmax_lse, S_dmask, rng_state _flash_attn_varlen_forward( q, k, v, cu_seqlens, cu_seqlens, max_seqlen, max_seqlen, dropout_p, softmax_scale, causalcausal, window_size_leftwindow_size[0], window_size_rightwindow_size[1], softcapsoftcap, alibi_slopesalibi_slopes, return_softmaxreturn_softmax and dropout_p 0, block_tableNone, ) # 移除padding,恢复原始head size out out_padded[..., :head_size_og] # 根据需要返回softmax结果 return out if not return_softmax else (out, softmax_lse, S_dmask) def flash_attn_varlen_func( q, k, v, cu_seqlens_q, cu_seqlens_k, max_seqlen_q, max_seqlen_k, dropout_p0.0, softmax_scaleNone, causalFalse, window_size(-1, -1), # -1 means infinite context window softcap0.0, # 0.0 means deactivated alibi_slopesNone, deterministicFalse, return_attn_probsFalse, block_tableNone,
): return FlashAttnVarlenFunc.apply( q, k, v, cu_seqlens_q, cu_seqlens_k, max_seqlen_q, max_seqlen_k, dropout_p, softmax_scale, causal, window_size, softcap, alibi_slopes, deterministic, return_attn_probs, block_table, )
上面这段代码清晰展示了 flash_attn_varlen_func 这个接口的调用逻辑接下来我们就可以去看一下flash_attn_cuda.varlen_fwd这个接口的具体dispatch逻辑了。
5. flash_attn_cuda.varlen_fwd的初步dispatch逻辑
首先来到这里https://github.com/Dao-AILab/flash-attention/blob/main/csrc/flash_attn/flash_api.cpp#L1518
PYBIND11_MODULE(TORCH_EXTENSION_NAME, m) { m.doc() FlashAttention; m.def(fwd, mha_fwd, Forward pass); m.def(varlen_fwd, mha_varlen_fwd, Forward pass (variable length)); m.def(bwd, mha_bwd, Backward pass); m.def(varlen_bwd, mha_varlen_bwd, Backward pass (variable length)); m.def(fwd_kvcache, mha_fwd_kvcache, Forward pass, with KV-cache);
}
可以发现flash_attn_cuda.varlen_fwd接口对应了mha_varlen_fwd这个c函数。从这里我们应该就可以看到flash attention forward的dispatch逻辑了。
std::vectorat::Tensor
mha_varlen_fwd(at::Tensor q, // total_q x num_heads x head_size, total_q为每个batch中序列长度的总和 const at::Tensor k, // total_k x num_heads_k x head_size, total_k为每个batch中序列长度的总和如果有block_table则为num_blocks x page_block_size x num_heads_k x head_size const at::Tensor v, // total_k x num_heads_k x head_size, total_k为每个batch中序列长度的总和如果有block_table则为num_blocks x page_block_size x num_heads_k x head_size c10::optionalat::Tensor out_, // total_q x num_heads x head_size, total_q为每个batch中序列长度的总和 const at::Tensor cu_seqlens_q, // b1 const at::Tensor cu_seqlens_k, // b1 c10::optionalat::Tensor seqused_k, // b。如果提供了该参数则每个batch元素只使用这么多个key c10::optionalconst at::Tensor leftpad_k_, // batch_size c10::optionalat::Tensor block_table_, // batch_size x max_num_blocks_per_seq c10::optionalat::Tensor alibi_slopes_, // num_heads或b x num_heads int max_seqlen_q, const int max_seqlen_k, const float p_dropout, const float softmax_scale, const bool zero_tensors, bool is_causal, int window_size_left, int window_size_right, const float softcap, const bool return_softmax, c10::optionalat::Generator gen_) { // 获取当前CUDA设备的属性 auto dprops at::cuda::getCurrentDeviceProperties(); // 检查GPU架构版本 // 判断是否为Ampere(SM8x)架构 bool is_sm8x dprops-major 8 dprops-minor 0; // 判断是否为Hopper(SM90)架构 bool is_sm90 dprops-major 9 dprops-minor 0; // 检查GPU架构要求 - 目前只支持Ampere或更新的架构 TORCH_CHECK(is_sm90 || is_sm8x, FlashAttention only supports Ampere GPUs or newer.); // 检查输入数据类型 auto q_dtype q.dtype(); // 只支持fp16和bf16数据类型 TORCH_CHECK(q_dtype torch::kFloat16 || q_dtype torch::kBFloat16, FlashAttention only support fp16 and bf16 data type); // bf16只在Ampere及以上架构支持 if (q_dtype torch::kBFloat16) { TORCH_CHECK(is_sm90 || is_sm8x, bfloat16 is only supported on Ampere GPUs or newer); } // 检查QKV的数据类型一致性 TORCH_CHECK(k.dtype() q_dtype, query and key must have the same dtype); TORCH_CHECK(v.dtype() q_dtype, query and value must have the same dtype); // 检查序列长度累加和的数据类型为int32 TORCH_CHECK(cu_seqlens_q.dtype() torch::kInt32, cu_seqlens_q must have dtype int32); TORCH_CHECK(cu_seqlens_k.dtype() torch::kInt32, cu_seqlens_k must have dtype int32); // 检查所有输入tensor是否在同一设备上 CHECK_DEVICE(q); CHECK_DEVICE(k); CHECK_DEVICE(v); CHECK_DEVICE(cu_seqlens_q); CHECK_DEVICE(cu_seqlens_k); // 检查分块表相关参数 at::Tensor block_table; const bool paged_KV block_table_.has_value(); // 是否使用分页KV缓存 if (paged_KV) { block_table block_table_.value(); CHECK_DEVICE(block_table); // 检查设备 TORCH_CHECK(block_table.dtype() torch::kInt32, block_table必须是int32类型); TORCH_CHECK(block_table.stride(-1) 1, block_table最后一维必须连续); } // 检查QKV张量的内存布局 TORCH_CHECK(q.stride(-1) 1, 输入张量最后一维必须连续); TORCH_CHECK(k.stride(-1) 1, 输入张量最后一维必须连续); TORCH_CHECK(v.stride(-1) 1, 输入张量最后一维必须连续); CHECK_CONTIGUOUS(cu_seqlens_q); // 检查序列长度累加和是否连续 CHECK_CONTIGUOUS(cu_seqlens_k); const auto sizes q.sizes(); // 获取Q的形状 // 获取基本参数 const int batch_size cu_seqlens_q.numel() - 1; // 批次大小 int num_heads sizes[1]; // Q的注意力头数 const int head_size sizes[2]; // 每个头的维度 const int num_heads_k paged_KV ? k.size(2) : k.size(1); // K的注意力头数 // softcap和dropout不能同时使用 if (softcap 0.f) { TORCH_CHECK(p_dropout 0.f, Softcapping暂不支持dropout); } // 分页KV缓存相关参数 const int max_num_blocks_per_seq !paged_KV ? 0 : block_table.size(1); // 每个序列最大块数 const int num_blocks !paged_KV ? 0 : k.size(0); // 总块数 const int page_block_size !paged_KV ? 1 : k.size(1); // 每块大小 TORCH_CHECK(!paged_KV || page_block_size % 256 0, 分页KV缓存块大小必须是256的倍数); // 因果掩码和窗口大小相关处理 if (max_seqlen_q 1 !alibi_slopes_.has_value()) { is_causal false; } if (is_causal) { window_size_right 0; } void *cu_seqlens_q_d cu_seqlens_q.data_ptr(); // 判断是否需要对Q进行重排 // 满足以下条件时需要重排: // 1. Q序列长度为1(即解码阶段) // 2. Q的注意力头数大于K的注意力头数(即MQA/GQA场景) // 3. 不使用滑动窗口(window_size_left和window_size_right都为-1) // 4. 不使用dropout // 5. head_size是8的倍数 // 6. 不使用ALiBi位置编码 const int seqlenq_ngroups_swapped max_seqlen_q 1 num_heads num_heads_k window_size_left 0 window_size_right 0 p_dropout 0.f head_size % 8 0 !alibi_slopes_.has_value(); // 计算每个K/V头对应多少个Q头 const int ngroups num_heads / num_heads_k; // 如果需要重排 if (seqlenq_ngroups_swapped) { // 将Q的形状从(batch_size, 1, num_heads_k * ngroups, head_size) // 重排为(batch_size * ngroups, num_heads_k, head_size) // 这样可以让同一个K/V头对应的Q头在内存上连续,提高访问效率 q q.reshape({batch_size, num_heads_k, ngroups, head_size}) .transpose(1, 2) .reshape({batch_size * ngroups, num_heads_k, head_size}); // 更新相关参数 max_seqlen_q ngroups; // Q序列长度变为ngroups num_heads num_heads_k; // Q的头数变为K的头数 cu_seqlens_q_d nullptr; // 不再需要Q的序列长度累加和 } const int total_q q.sizes()[0]; // Q的总token数 // 检查输入参数的合法性 // 1. batch_size必须为正数 TORCH_CHECK(batch_size 0, batch size must be positive); // 2. head_size必须小于等于256,这是Flash Attention的限制 TORCH_CHECK(head_size 256, FlashAttention forward only supports head dimension at most 256); // 3. head_size必须是8的倍数,这是为了内存对齐和CUDA优化 TORCH_CHECK(head_size % 8 0, query, key, value, and out_ must have a head_size that is a multiple of 8); // 4. Q的head数必须是K/V的head数的整数倍,这是为了支持MQA/GQA TORCH_CHECK(num_heads % num_heads_k 0, Number of heads in key/value must divide number of heads in query); // 如果滑动窗口大小超过了K序列的最大长度,则设置为-1表示不使用滑动窗口 if (window_size_left max_seqlen_k) { window_size_left -1; } if (window_size_right max_seqlen_k) { window_size_right -1; } // 检查Q张量的形状是否正确: [total_q, num_heads, head_size] CHECK_SHAPE(q, total_q, num_heads, head_size); // 根据是否使用分页KV缓存来检查K/V张量的形状 if (!paged_KV) { // 不使用分页KV缓存时,K/V的形状应为[total_k, num_heads_k, head_size] const int total_k k.size(0); CHECK_SHAPE(k, total_k, num_heads_k, head_size); CHECK_SHAPE(v, total_k, num_heads_k, head_size); } else { // 使用分页KV缓存时,K/V的形状应为[num_blocks, page_block_size, num_heads_k, head_size] // block_table的形状应为[batch_size, max_num_blocks_per_seq] CHECK_SHAPE(k, num_blocks, page_block_size, num_heads_k, head_size); CHECK_SHAPE(v, num_blocks, page_block_size, num_heads_k, head_size); CHECK_SHAPE(block_table, batch_size, max_num_blocks_per_seq); } // 检查序列长度累加和张量的形状,应为[batch_size 1] CHECK_SHAPE(cu_seqlens_q, batch_size 1); CHECK_SHAPE(cu_seqlens_k, batch_size 1); // 如果提供了K序列使用长度的信息,检查其属性 if (seqused_k.has_value()){ auto seqused_k_ seqused_k.value(); // 数据类型必须是int32 TORCH_CHECK(seqused_k_.dtype() torch::kInt32, seqused_k must have dtype int32); // 必须在CUDA设备上 TORCH_CHECK(seqused_k_.is_cuda(), seqused_k must be on CUDA device); // 必须是连续的内存布局 TORCH_CHECK(seqused_k_.is_contiguous(), seqused_k must be contiguous); // 形状必须是[batch_size] CHECK_SHAPE(seqused_k_, batch_size); } // 创建输出张量 at::Tensor out; // 如果提供了输出张量 if (out_.has_value()) { out out_.value(); // 检查输出张量的属性: // 1. 数据类型必须与输入相同 TORCH_CHECK(out.dtype() q_dtype, Output must have the same dtype as inputs); // 2. 必须在同一设备上 CHECK_DEVICE(out); // 3. 最后一维必须是连续的 TORCH_CHECK(out.stride(-1) 1, Output tensor must have contiguous last dimension); // 4. 形状必须正确 CHECK_SHAPE(out, sizes[0], sizes[1], head_size); // 如果序列长度和组数需要交换 if (seqlenq_ngroups_swapped) { // 重塑张量维度并转置,用于处理分组注意力 out out.reshape({batch_size, num_heads_k, ngroups, head_size}) .transpose(1, 2) .reshape({batch_size * ngroups, num_heads_k, head_size}); } } else { // 如果没有提供输出张量,创建一个与输入形状相同的空张量 out torch::empty_like(q); } // 定义一个lambda函数,用于将数字向上取整到m的倍数 auto round_multiple [](int x, int m) { return (x m - 1) / m * m; }; // 计算head_size的对齐值: // - 如果head_size 192,向上取整到32的倍数 // - 否则设为256 const int head_size_rounded head_size 192 ? round_multiple(head_size, 32) : 256; // 将Q序列长度向上取整到128的倍数 const int seqlen_q_rounded round_multiple(max_seqlen_q, 128); // 将K序列长度向上取整到128的倍数 const int seqlen_k_rounded round_multiple(max_seqlen_k, 128); // 设置CUDA设备,确保在正确的GPU上执行 at::cuda::CUDAGuard device_guard{(char)q.get_device()}; // 获取输入张量q的选项(设备、数据类型等) auto opts q.options(); // 创建softmax_lse张量,用于存储每个注意力头的log-sum-exp值 auto softmax_lse torch::empty({num_heads, total_q}, opts.dtype(at::kFloat)); at::Tensor p; // 只有在有dropout时才返回softmax结果,以减少编译时间 if (return_softmax) { // 确保dropout概率大于0 TORCH_CHECK(p_dropout 0.0f, return_softmax is only supported when p_dropout 0.0); // 创建p张量存储softmax结果 p torch::empty({ batch_size, num_heads, seqlen_q_rounded, seqlen_k_rounded }, opts); } else { // 如果不需要返回softmax,创建一个空张量 p torch::empty({ 0 }, opts); } // 如果需要将张量初始化为0 if (zero_tensors) { out.zero_(); // 输出张量置0 softmax_lse.fill_(-std::numeric_limitsfloat::infinity()); // softmax_lse填充负无穷 if (return_softmax) {p.zero_();} // softmax结果张量置0 } // 创建前向传播参数结构体 Flash_fwd_params params; // 设置前向传播的各项参数 set_params_fprop(params, batch_size, max_seqlen_q, max_seqlen_k, seqlen_q_rounded, seqlen_k_rounded, num_heads, num_heads_k, head_size, head_size_rounded, q, k, v, out, cu_seqlens_q_d, cu_seqlens_k.data_ptr(), seqused_k.has_value() ? seqused_k.value().data_ptr() : nullptr, return_softmax ? p.data_ptr() : nullptr, softmax_lse.data_ptr(), p_dropout, softmax_scale, window_size_left, window_size_right, softcap, seqlenq_ngroups_swapped, /*unpadded_lse*/true); params.total_q total_q; // 如果使用分页KV缓存 if (paged_KV) { params.block_table block_table.data_ptrint(); // 设置分块表指针 params.block_table_batch_stride block_table.stride(0); // 设置分块表的batch步长 params.k_batch_stride k.stride(0); // 设置K的batch步长 params.v_batch_stride v.stride(0); // 设置V的batch步长 } params.page_block_size page_block_size; // 设置页块大小 // 保持对这些张量的引用以延长其生命周期 at::Tensor softmax_lse_accum, out_accum; if (seqlenq_ngroups_swapped) { // 仅在解码时应用split-k std::tie(softmax_lse_accum, out_accum) set_params_splitkv(params, batch_size, num_heads, head_size, max_seqlen_k, max_seqlen_q, head_size_rounded, p_dropout, /*num_splits*/ 0, dprops, opts); } // 如果提供了K序列的左侧填充信息 if (leftpad_k_.has_value()) { auto leftpad_k leftpad_k_.value(); // 检查:不能同时使用分页KV和左侧填充 TORCH_CHECK(!paged_KV, We dont support Paged KV and leftpad_k running at the same time yet); // 检查数据类型必须是int32 TORCH_CHECK(leftpad_k.dtype() torch::kInt32, leftpad_k must have dtype int32); CHECK_DEVICE(leftpad_k); // 检查设备 CHECK_CONTIGUOUS(leftpad_k); // 检查连续性 CHECK_SHAPE(leftpad_k, batch_size); // 检查形状 params.leftpad_k static_castint *(leftpad_k.data_ptr()); // 设置左侧填充指针 } // 为每个线程生成随机数的次数,用于偏移THC随机状态中的philox计数器 // 我们使用自定义的RNG,将偏移量增加batch_size * num_heads * 32 int64_t counter_offset params.b * params.h * 32; // 创建一个CUDA上的float32类型的张量选项 auto options torch::TensorOptions().dtype(torch::kFloat32).device(torch::kCUDA); // 创建一个大小为2的int64类型的空张量,用于存储RNG状态 auto rng_state torch::empty({2}, options.dtype(torch::kInt64)); // 前向传播kernel将用种子和偏移量填充内存 params.rng_state reinterpret_castuint64_t*(rng_state.data_ptr()); // 如果设置了dropout if (p_dropout 0.0) { // 获取默认的CUDA生成器或使用提供的生成器 auto gen at::get_generator_or_defaultat::CUDAGeneratorImpl( gen_, at::cuda::detail::getDefaultCUDAGenerator()); // 使用互斥锁保护随机数生成器的访问 std::lock_guardstd::mutex lock(gen-mutex_); // 设置philox随机数生成器的状态 params.philox_args gen-philox_cuda_state(counter_offset); } // 设置ALiBi(Attention with Linear Biases)的参数 set_params_alibi(params, alibi_slopes_, batch_size, num_heads); // 如果K序列长度大于0,执行前向传播 if (max_seqlen_k 0) { // 获取当前CUDA流 auto stream at::cuda::getCurrentCUDAStream().stream(); // 运行前向传播kernel run_mha_fwd(params, stream, paged_KV); } else { // 如果K序列长度为0,说明是空张量,需要将输出置零 out.zero_(); // 将softmax的对数和填充为负无穷 softmax_lse.fill_(std::numeric_limitsfloat::infinity()); } // 如果进行了序列长度和组数的交换 if (seqlenq_ngroups_swapped) { // 定义reshape前后的维度大小 int64_t size_before[] {batch_size, max_seqlen_q, num_heads_k, head_size}; int64_t size_after[] {batch_size, num_heads_k * max_seqlen_q, head_size}; // 重新排列输出张量的维度 out out.reshape(size_before).transpose(1, 2).reshape(size_after); q q.reshape(size_before).transpose(1, 2).reshape(size_after); // 重新排列softmax对数和的维度 softmax_lse softmax_lse.reshape({num_heads * max_seqlen_q, batch_size}); } // 返回输出张量、softmax对数和、注意力分布(如果需要)和RNG状态 return {out, softmax_lse, p, rng_state};
}
由于Flash Attention的准备工作比较多上面的代码很长我们主要关注
if (seqlenq_ngroups_swapped) { // 仅在解码时应用split-k std::tie(softmax_lse_accum, out_accum) set_params_splitkv(params, batch_size, num_heads, head_size, max_seqlen_k, max_seqlen_q, head_size_rounded, p_dropout, /*num_splits*/ 0, dprops, opts); }
和
if (max_seqlen_k 0) { // 获取当前CUDA流 auto stream at::cuda::getCurrentCUDAStream().stream(); // 运行前向传播kernel run_mha_fwd(params, stream, paged_KV); }
这几行代码即可set_params_splitkv决定了是否使用split-k以及要在kv的序列纬度上切分多少次run_mha_fwd会根据set_params_splitkv的配置以及在上面的函数中其它部分设置的params的参数来dispatch不同的kernel。现在来看一下set_params_splitkv的实现
std::tupleat::Tensor, at::Tensor set_params_splitkv(Flash_fwd_params ¶ms, const int batch_size, const int num_heads, const int head_size, const int max_seqlen_k, const int max_seqlen_q, const int head_size_rounded, const float p_dropout, const int num_splits, cudaDeviceProp *dprops, struct c10::TensorOptions opts) { // 这里的block_n需要和run_mha_fwd_splitkv_dispatch中的配置匹配 // 根据head_size的大小选择不同的block_n: // - head_size 64: block_n 256 // - 64 head_size 128: block_n 128 // - head_size 128: block_n 64 const int block_n head_size 64 ? 256 : (head_size 128 ? 128 : 64); // 计算在K序列维度上需要多少个block const int num_n_blocks (max_seqlen_k block_n - 1) / block_n; // 对于splitKV kernel,kBlockM固定为64 // 一般在推理时Q序列长度不会超过64 const int num_m_blocks (max_seqlen_q 64 - 1) / 64; // 设置切分数量 params.num_splits num_splits; // 声明用于存储中间结果的tensor at::Tensor softmax_lse_accum; at::Tensor out_accum; // splitKV目前不支持dropout if (p_dropout 0.0f) { if (num_splits 1) { // 如果num_splits 1,则使用启发式方法计算切分数量 // 这里乘以2是因为每个block使用128个线程 params.num_splits num_splits_heuristic(batch_size * num_heads * num_m_blocks, dprops-multiProcessorCount * 2, num_n_blocks, 128); } // 如果需要切分(num_splits 1) if (params.num_splits 1) { // 分配存储中间结果的tensor softmax_lse_accum torch::empty({params.num_splits, batch_size, num_heads, max_seqlen_q}, opts.dtype(at::kFloat)); out_accum torch::empty({params.num_splits, batch_size, num_heads, max_seqlen_q, head_size_rounded}, opts.dtype(at::kFloat)); // 设置指向中间结果的指针 params.softmax_lseaccum_ptr softmax_lse_accum.data_ptr(); params.oaccum_ptr out_accum.data_ptr(); } // 切分数量不能超过128 TORCH_CHECK(params.num_splits 128, num_splits 128 not supported); } return std::make_tuple(softmax_lse_accum, out_accum);
}
由于调用set_params_splitkv时设置了num_splits0所以上面的代码会进入到启发式计算切分数量的逻辑中启发式计算切分数量的逻辑在num_splits_heuristic中我们来看一下这个函数的实现
// 这个函数用于找到最大化 GPU 占用率的切分数量。
// 例如,如果 batch * n_heads 48,且有 108 个 SM,那么:
// - 使用 2 个切分(效率 0.89)比使用 3 个切分(效率 0.67)更好
// 但是我们也不希望切分太多,因为这会导致更多的 HBM 读写。
// 所以我们先找到最佳效率,然后找到能达到最佳效率 85% 的最小切分数量。
inline int num_splits_heuristic(int batch_nheads_mblocks, int num_SMs, int num_n_blocks, int max_splits) { // 如果当前 batch_nheads_mblocks 已经能填充 80% 的 SM,就不需要切分了 if (batch_nheads_mblocks 0.8f * num_SMs) { return 1; } // 取 max_splits、SM数量和 n_blocks 三者的最小值作为最大切分数量 max_splits std::min({max_splits, num_SMs, num_n_blocks}); float max_efficiency 0.f; std::vectorfloat efficiency; efficiency.reserve(max_splits); // 向上取整除法 auto ceildiv [](int a, int b) { return (a b - 1) / b; }; // 有些切分数量是无效的。例如,如果我们有 64 个 blocks: // - 选择 11 个切分,我们会有 6 * 10 4 个 blocks // - 选择 12 个切分,我们会有 6 * 11 (-2) 个 blocks(实际上还是 11 个切分) // 所以我们需要检查每个切分的 block 数量是否与前一个切分数量相同 auto is_split_eligible [ceildiv, num_n_blocks](int num_splits) { return num_splits 1 || ceildiv(num_n_blocks, num_splits) ! ceildiv(num_n_blocks, num_splits - 1); }; // 第一轮循环:计算每个切分数量的效率,并找到最大效率 for (int num_splits 1; num_splits max_splits; num_splits) { if (!is_split_eligible(num_splits)) { efficiency.push_back(0.f); } else { // n_waves 表示每个 SM 平均需要处理多少波 blocks float n_waves float(batch_nheads_mblocks * num_splits) / num_SMs; // 效率 理论处理时间 / 实际处理时间 float eff n_waves / ceil(n_waves); if (eff max_efficiency) { max_efficiency eff; } efficiency.push_back(eff); } } // 第二轮循环:找到能达到最佳效率 85% 的最小切分数量 for (int num_splits 1; num_splits max_splits; num_splits) { if (!is_split_eligible(num_splits)) { continue; } if (efficiency[num_splits - 1] 0.85 * max_efficiency) { return num_splits; } } return 1;
}
从上面的代码我们就可以看出来影响splitkv的参数不仅有max_seqlen_khead_num还有SM个数等等。对于文章开头的例子head_num和SM个数是固定的但由于max_new_tokens从512变成64引起了max_seqlen_k的改变从而导致了num_splits的改变最终表现为我们在max_new_tokens为512的nsys中观察到了decoding时使用了splitkv的flash attention实现而在max_new_tokens为64的nsys中则没有使用splitkv的flash attention实现。run_mha_fwd的dispatch逻辑为
void run_mha_fwd(Flash_fwd_params ¶ms, cudaStream_t stream, bool force_split_kernelfalse) { FP16_SWITCH(!params.is_bf16, [] { HEADDIM_SWITCH(params.d, [] { BOOL_SWITCH(params.is_causal, Is_causal, [] { if (params.num_splits 1 !force_split_kernel) { // If we dont set it num_splits 0 run_mha_fwd_elem_type, kHeadDim, Is_causal(params, stream); } else { run_mha_fwd_splitkv_dispatchelem_type, kHeadDim, Is_causal(params, stream); } }); }); });
}
可以看到这里对num_splits进行判断如果num_splits 1且没有设置force_split_kernel则dispatch不使用splitkv的kernel否则dispatch使用splitkv的kernel。flash_attn_cuda.varlen_fwd的初步dispatch逻辑就梳理完了不过我们从文章开头的nsys可以看到调用splitkv实现的时候每个decoding step的每个Attenion计算都有2个kernel 在KV的seq纬度切分之后还需要把单独计算的结果组合成最终的计算结果这就是flash_fwd_splitkv_combine_kernel的作用。实际上这个也被叫作Flash Decoding你可以参考https://mp.weixin.qq.com/s/hvqPhNo3l0tL_-lf978euw 这里的介绍。
5. run_mha_fwd_splitkv_dispatch的上层实现逻辑
templatetypename Kernel_traits, bool Is_causal
void run_flash_splitkv_fwd(Flash_fwd_params ¶ms, cudaStream_t stream) { // 确保kernel特征不支持Q在寄存器中和Q/K共享共享内存 static_assert(!Kernel_traits::Is_Q_in_regs, SplitKV implementation does not support Is_Q_in_regs); static_assert(!Kernel_traits::Share_Q_K_smem, SplitKV implementation does not support Share_Q_K_smem); // 获取共享内存大小 constexpr size_t smem_size Kernel_traits::kSmemSize; // 计算M维度的block数量 const int num_m_block (params.seqlen_q Kernel_traits::kBlockM - 1) / Kernel_traits::kBlockM; // 设置grid维度: // x: M维度的block数 // y: 如果有splits则为splits数量,否则为batch size // z: 如果有splits则为batch*heads,否则为heads数量 dim3 grid(num_m_block, params.num_splits 1 ? params.num_splits : params.b, params.num_splits 1 ? params.b * params.h : params.h); // 判断序列长度是否能被block大小整除 const bool is_even_MN params.cu_seqlens_q nullptr params.cu_seqlens_k nullptr params.seqlen_k % Kernel_traits::kBlockN 0 params.seqlen_q % Kernel_traits::kBlockM 0; // 判断head维度是否匹配 const bool is_even_K params.d Kernel_traits::kHeadDim; // 使用一系列宏来根据不同条件选择不同的kernel实现 BOOL_SWITCH(is_even_MN, IsEvenMNConst, [] { EVENK_SWITCH(is_even_K, IsEvenKConst, [] { LOCAL_SWITCH((params.window_size_left 0 || params.window_size_right 0) !Is_causal, Is_local, [] { BOOL_SWITCH(params.num_splits 1, Split, [] { BOOL_SWITCH(params.knew_ptr ! nullptr, Append_KV, [] { ALIBI_SWITCH(params.alibi_slopes_ptr ! nullptr, Has_alibi, [] { SOFTCAP_SWITCH(params.softcap 0.0, Is_softcap, [] { // 选择合适的kernel实现 auto kernel flash_fwd_splitkv_kernelKernel_traits, Is_causal, Is_local !Is_causal, Has_alibi, IsEvenMNConst !Append_KV IsEvenKConst !Is_local Kernel_traits::kHeadDim 128, IsEvenKConst, Is_softcap, Split, Append_KV; // 如果共享内存超过48KB,需要设置属性 if (smem_size 48 * 1024) { C10_CUDA_CHECK(cudaFuncSetAttribute( kernel, cudaFuncAttributeMaxDynamicSharedMemorySize, smem_size)); } // 启动kernel kernelgrid, Kernel_traits::kNThreads, smem_size, stream(params); C10_CUDA_KERNEL_LAUNCH_CHECK(); }); }); }); }); }); }); }); // 如果有splits,需要启动combine kernel来合并结果 if (params.num_splits 1) { // 根据head维度选择合适的block大小 constexpr static int kBlockM Kernel_traits::kHeadDim % 128 0 ? 4 : (Kernel_traits::kHeadDim % 64 0 ? 8 : 16); dim3 grid_combine((params.b * params.h * params.seqlen_q kBlockM - 1) / kBlockM); // 根据splits数量选择合适的combine kernel EVENK_SWITCH(is_even_K, IsEvenKConst, [] { if (params.num_splits 2) { flash_fwd_splitkv_combine_kernelKernel_traits, kBlockM, 1, IsEvenKConstgrid_combine, Kernel_traits::kNThreads, 0, stream(params); } else if (params.num_splits 4) { flash_fwd_splitkv_combine_kernelKernel_traits, kBlockM, 2, IsEvenKConstgrid_combine, Kernel_traits::kNThreads, 0, stream(params); } else if (params.num_splits 8) { flash_fwd_splitkv_combine_kernelKernel_traits, kBlockM, 3, IsEvenKConstgrid_combine, Kernel_traits::kNThreads, 0, stream(params); } else if (params.num_splits 16) { flash_fwd_splitkv_combine_kernelKernel_traits, kBlockM, 4, IsEvenKConstgrid_combine, Kernel_traits::kNThreads, 0, stream(params); } else if (params.num_splits 32) { flash_fwd_splitkv_combine_kernelKernel_traits, kBlockM, 5, IsEvenKConstgrid_combine, Kernel_traits::kNThreads, 0, stream(params); } else if (params.num_splits 64) { flash_fwd_splitkv_combine_kernelKernel_traits, kBlockM, 6, IsEvenKConstgrid_combine, Kernel_traits::kNThreads, 0, stream(params); } else if (params.num_splits 128) { flash_fwd_splitkv_combine_kernelKernel_traits, kBlockM, 7, IsEvenKConstgrid_combine, Kernel_traits::kNThreads, 0, stream(params); } C10_CUDA_KERNEL_LAUNCH_CHECK(); }); }
} // 根据head维度选择合适的block大小并调用run_flash_splitkv_fwd
templatetypename T, int Headdim, bool Is_causal
void run_mha_fwd_splitkv_dispatch(Flash_fwd_params ¶ms, cudaStream_t stream) { constexpr static int kBlockM 64; // 固定M维度的block大小为64 // 根据head维度选择N维度的block大小: // head维度64: 256 // head维度128: 128 // 其他: 64 constexpr static int kBlockN Headdim 64 ? 256 : (Headdim 128 ? 128 : 64); run_flash_splitkv_fwdFlash_fwd_kernel_traitsHeaddim, kBlockM, kBlockN, 4, false, false, T, Is_causal(params, stream);
}
我们可以看到无论是在序列维度切分计算的flash_fwd_split_kv_kernel还是最后合并结果的flash_fwd_splitkv_combine_kernel他们都有非常多的模板来决定当前的输入下应该使用哪种kernel来获得最佳性能。如果你对这里的cuda实现感兴趣可以自行阅读源码学习或者修改。
6. 总结
本文主要探讨了Flash Attention在不同场景下的kernel dispatch逻辑特别关注了decode阶段使用split_kv实现的触发条件。通过分析源码发现Flash Attention的dispatch逻辑主要由max_seqlen_kK序列的最大长度、head_num注意力头数、SM数量GPU的流处理器数量等因素决定。这些因素会通过启发式函数num_splits_heuristic来计算num_splitsKV序列维度的切分数量该函数的目标是找到能最大化GPU利用率的切分数量。当计算得到num_splits 1时会使用split_kv实现这种实现会启动两个kernelflash_fwd_splitkv_kernel用于在KV序列维度上进行切分计算flash_fwd_splitkv_combine_kernel用于合并各个切分的计算结果。这就解释了文章开头的例子中当max_new_tokens512时由于序列长度较长导致num_splits 1而使用split_kv实现而max_new_tokens64时由于序列长度较短导致num_splits 1而使用普通实现的现象。这种灵活的dispatch机制设计使得Flash Attention能够在不同场景下都获得较好的性能表现在长序列场景下通过split_kv更好地利用GPU资源在短序列场景下避免不必要的开销。
