动态N：M细粒度结构稀疏注意机制在GPU上的加速研究

369⇠⇥！动态N：M细粒度结构稀疏注意机制陈照东加州大学圣巴巴拉分校chenzd15thu@ucsb.edu加利福尼亚州刘柳美国加州圣巴巴拉伦斯勒理工学院liurpi.edu郑渠加州大学圣巴巴拉分校美国加利福尼亚州圣巴巴拉zhengqu@ucsb.edu丁宇飞加州大学圣巴巴拉分校美国加利福尼亚州圣巴巴拉yufeiding@ucsb.edu全玉英加州大学圣巴巴拉分校yuying_quan@ucsb.edu加利福尼亚州元勰ALIBABA DAMOAcademyyuanxie@gmail.com加州圣巴巴拉摘要Transformer正在成为NLP和计算机视觉等各种任务 尽管它们取得了成功，但注意力机制的高度复杂性阻碍了它们应用于延迟敏感的任务。加速注意机制的一个机会是杠杆化注意权重矩阵中的稀疏性然而，由于“动态”和“动态”之间的困境，研究未能在中等序列长度下实现GPU上的加速。他们还需要昂贵的再培训，以恢复准确性。 在本文中，我们提出了D，第一个GPU友好的动态机制，以解决这一困境。将全注意力得分矩阵D维修剪为N：M细粒度结构化稀疏模式。我们的关键见解是，在动态方面，N：M稀疏性对于GPU上的稀疏矩阵的修剪和编码是友好的。在稠密-稠密矩阵乘法内核上，这是同类中的第一个它将查询矩阵和关键矩阵相乘，修剪结果，并在没有开销的情况下编码压缩的稀疏矩阵。与以前的研究相比，D实现了在任意序列长度的加速只需要几新的调整时期，以达到与全atten标准精度作用机制我们提供了理论和经验证据来证明D是一个很好的近似的充分注意机制。 我们评估了di实现1下的1：2和2：4稀疏性。381. 在A100 GPU上获得86倍的加速比。在任务上，允许使数字或硬拷贝的全部或部分的工作，为个人或课堂使用的是免费提供的副本不制作或分发为亲或商业利益，副本承担本通知和第一页上的全部引文必须尊重作者以外的其他人拥有的本作品组件的版权允许使用学分进行摘要复制，或重新发布，张贴在服务器上或重新分发到 列 表 ， 需 要 事 先 的特 殊 许 可 和 / 或 费 用 。 请 求 权 限 请 发 邮 件 至permissions@acm.org。PPoPP©2023版权归所有者/作者所有。授权给ACM的出版权。序列长度从384到4096的各种域，其准确性与密集预训练模型的几个netuning epoch后的完全注意力相当CCS概念：·计算方法并行计算方法。关键词：Transformer，GPGPU，动态剪枝ACM参考格式：陈昭东、曲正、权玉英、刘柳、丁宇飞、谢元。2023年动态N：M 细 粒 度 结 构 化 稀 疏 注 意 机 制 。 ACM SIGPLAN AnnualSymphonyPrinciplesandPracticeofParallelProgramming（PPoPP'23）. ACM，纽约州纽约市，美国，20页。1介绍Transformers [23]在NLP [16]和计算机视觉[6]等各个领域都取得了有竞争力的。将它们与传统神经网络架构区分开来的关键特征是注意力机制[23]，它允许转换器以自适应和可学习的方式从输入序列中的元素嵌入中收集信息。然而，注意力机制带来的高计算成本和内存占用使其不能应用于延迟敏感的任务。 加速注意力机制的一种方式是利用注意力权重矩阵中的稀疏性，因为已经表明，只有少数条目对模型准确性至关重要。与传统的静态权重稀疏性不同，注意力的稀疏性是动态的、新粒度的。 它是动态的，因为注意力得分矩阵是从输入激活计算的，因此对于每个推理样本都是di。是n-grained，因为其中的任何条目都可以具有主导值。不幸的是，在GPU上存在e-dynamic和ne-grained之间的困境一方面，动态修剪、编码和解码y上的稀疏注意力得分矩阵需要规则的稀疏模式以在GPU上实现加速。另一方面，不规则的稀疏模式被期望捕获占主导地位的ne粒度分布。PPoPP陈赵东，郑，余英安，刘柳，丁宇飞，谢元370（）第二个⇠⇥/（/）⇥⇥p条目 先前的研究[1，11，19，20，22，27，28]使用粗粒度块稀疏性，其允许动态编码和解码，但不安全。因此，它们需要从头开始训练或重新训练，以迫使主要条目遵循启发式或学习的模式，即使对于小规模的Transformer模型，这也可能是耗时 其他方法，如低位近似或基于机器学习的预测，对GPU不友好或不能保证准确。2相关工作首先介绍了本文的相关工作2.1全注意力机制注意力机制正在成为复杂序列建模的一个组成部分全注意力机制对序列中每对条目之间的依赖关系进行建模。 给定输入序列^= ���1，.， =R= 0.3，完全注意力机制可以被定义为在本文中，我们直接解决了）namic”和U=（>5C0G（WQ/3）1，（1）使用以下公式动态地修剪完整的注意力得分矩阵：其中，W=^]，Q=^]，并且是query、key和N：M细粒度结构化稀疏模式。关键的洞察力@：E在动态方面，N：M稀疏性，最初是去-该算法只对稀疏矩阵的e译码进行符号化，使得动态剪枝和编码易于在GPU上并行实现。 通过保留每M个条目的局部最大值来对每行中的条目进行。我们提出了一种新的动态采样稠密-稠密矩阵乘法（SDDMM）内核，该内核将查询矩阵和密钥矩阵相乘，动态地将其修剪为具有幅度的N：M稀疏，并对压缩的非零和元数据进行编码。此外，压缩后的稀疏核可以被后续的核消耗，以实现加速和减少冗余占用。 据我们所知，深度学习软件栈中的rst内核可以在没有开销的情况下修剪和编码稀疏矩阵。 与以前的研究相比，我们的方法在任意序列长度的oGPU上实现了加速。 它只需要从BERT-large这样的预训练密集模型中进行几个netuning epoch（几个GPU小时），就可以达到同等精度。虽然我们专注于1：2和2：4结构稀疏性，它们的编码由o[15]支持，本文中的技术和观察可以在硬件支持下同等地应用于所有N和Ms。 我们的主要贡献概述如下：我们提出了D，一个动态的N：M稀疏注意机制，这是一个下拉式的完全注意机制的替代品，并正交于现有的e注意机制。 它的正确性有经验和理论两方面的证据。 据我们所知，它是第一个动态修剪技术下的N：M稀疏。我们提出了一个专门的CUDA内核设计，以完全消除修剪开销。修剪作为产生注意力得分矩阵的密集矩阵乘法的尾声来实现。 它是深度学习软件栈中的rst运算符，可应用于动态和非粒度稀疏矩阵，并生成零开销的压缩稀疏编码。我们在A100 GPU上对跨各种域和序列长度的任务进行D评估 它实现了1。381. 86加速超过完全注意力，没有准确性损失。值矩阵WQ）形成全二次邻接矩阵，其边权重是序列中所有元素这个邻接矩阵是用1p3标准化，以保持单位二阶矩然后用softmax进行归一化 最后，将\中的行特征向量按归一化邻接矩阵相乘进行聚合。其他地区其中G=（>5C0GWQ<）p3. WereferW（Q）作为注意力得分矩阵，G作为注意力权重矩阵。2.2有效注意机制完全注意机制中的高计算成本和内存占用来自G，其大小与序列长度=成二次方增长。为了解决这个问题，许多人提出了[22]。固 定 稀 疏 模 式 。 Beltagy 等 人 2020[28]， Zaheer etal.2020[1]在G上应用一组固定的稀疏注意模式，如全局注意和滑动窗口注意。这些模式是根据经验观察构建的，并设计为GPU友好，以实现墙时间加速。动态和新粒度的挑战是通过对主要条目进行昂贵的再培训来解决的。动态稀疏模式。Ham等人2021 [8]使用低成本二进制哈希在G上动态生成ne粒度的稀疏注意模式。然而，这种技术需要专门的硬件来实现加速，因此它不适用于通用硬件，如GPU。 Kitaev等人2020[11]，Tay等人2020 [22]，Roy et al. 2021 [19]应用各种聚类方法，仅计算每个聚类内的注意力。虽然在每个集群中计算全注意力对GPU更友好，由于新粒度的稀疏性，聚类方法包含几个GPU不友好的运算符，如top-k和排序，因此它们在中等序列长度下受益低秩/内核。Wang等人2020年[25]项目G从= =to=：使用线性投影。Choromanski等人2021[5]介绍了FAVOR+，它使用核方法近似软最大值。这使他们能够改变···动态N：M细粒度结构稀疏A机制PPoPP371YY#（！KT）KT02⇥⇥11(A) 注意事项（nn（n个（）n图1.A100 GPU 1：2和2：4细粒度结构化稀疏修剪。[第十五条](B) DFSSn（（n（nn/2 n/16）计算顺序和减少渐近复杂性线性。然而，低秩投影和内核构造也引入了相当大的开销。这使得这些方法只能在长序列长度下使用此外，低级投射极大地改变了注意力机制，成千上万的预先训练或训练，需要调整步骤以达到与原始的完全注意力机制相当的性能所以它们需要大量的工程技术来部署。2.3N：M细粒度结构稀疏度N：M流行趋势。 它保留了原始稠密矩阵中每1M个向量中的N个元素，并提供比粗粒度块稀疏多几个数量级的组合，同时保持GPU友好。NVIDIA推出1：2和2：4Ne-grainedstructuredpruningtoweightmatricesinneuralnetworks [15]. 如图1所示，密集权重矩阵被修剪，然后用cuSPARSELt库中的API压缩。在推理期间，压缩的权 重 矩阵与 输 入 激 活 相 乘 ， 其 中 稀 疏 - 密 集 矩 阵 乘 法（SpMM）由Ampere稀疏张量核心加速。 这使1。9比密集计数器部分加速。 后续研究[12，14，18，31，32]提出了提高准确性，减少训练时间并使用新硬件设计加速N：M权重稀疏性的SpMM的算法。然而，据我们所知，没有以前的研究图2. 我们的动态N：M细粒度结构稀疏注意机制的概述，其中N：M=1：2或2：4。f表示softmax，为简单起见省略了p33DFSS：D粒子N：M细粒度S结构S解析注意机制D解决了“本节我们首先给出我们的D方法的概述。然后，我们详细讨论了两种设计选择。此外，我们还在附录A.1提供了Dfss的更多理论分析，并在附录A.4中讨论了如何将我们的方法与现有的线性注意力3.1概述我们提出的D机制很简单，图2说明了1：2或2：4稀疏度下的情况。与完全二次注意机制相比，该方法动态地修剪注意力分数，而不会产生存储或计算开销，同时保持e重要的是，该方法可以在具有定制CUDA内核的现有GPU硬件上实现实际的注意力加速。 清单1显示了使用D的所有mod-i。3.2设计选择1：生成稀疏性的位置注意力机制的计算图如图3所示。 它可分为三个阶段：（WQ）、（>5C<0G）和（G1）。当我们能够将N：M稀疏性应用于动态矩阵，如注意力权重生成稀疏图案：从计算W（Q）开始。Transformer中的矩阵这部分是由于N：M稀疏性只用于权重修剪的刻板印象，以及修剪和压缩动态矩阵而计算WQ）。使用softmax在线，这可以o的硬件加速的SpMM的好处，我们的工作通过开发一个专用的SDDMM内核来打破这种刻板印象，该内核消除了修剪和压缩开销，允许在动态矩阵中使用N：M稀疏性。图3. 注意阶段。如果我们从0开始，W和Q之间的稠密矩阵乘法将被传统的静态采样稠密-稠密矩阵乘法（SDDMM）所取代稀疏张量核选择二比四X点积细粒度结构化剪枝压缩一比二1：2/2：4稀疏度非零元数据！KTV！KT*！V*！KT#（！KT）12）*S#$t a0 ！KTPPoPP陈赵东，郑，余英安，刘柳，丁宇飞，谢元372211⇥# Full attentionmechanismimporttorch.nn.functionalasFdeffull_attention（q，k，v）：attn_weight =火炬。bmm（q，k. transpose（1，2））attn_weight = F。softmax（attn_weight，-1）返回火炬。bmm（attn_weight，v）# DFSS注意机制M阻塞的GEMM线程块平铺翘曲瓷砖导入dspattdefdfss_attention（q，k，v）：attn_weight，metadata = dspattn. sddmm（q，k）attn_weight = F. softmax（attn_weight，-1）返回dspattn。spmm（attn_weight，metadata，v）清单1.使用D. “dspattn”是我们开发的包装。仅计算输入稀疏模式的条目IDENT1。 （>5C<0G仅对每行中的非零值进行运算。将G和I之间的原始稠密矩阵乘法替换为稀疏矩阵-矩阵乘法（SpMM），其将稀疏矩阵与稠密矩阵相乘。然而，不可能在计算W（ Q）之前准确地知道W（Q）中的哪个条目具有更高的幅度。 如果我们从0开始，则需要指定重要条目位置的其他组件。此外，用传统的SD-DMM代替稠密的WDMM，即使在高稀疏度下也会限制加速比。Chen等人2021[3]表明，即使采用一些结构化设计，传统SDDMM在80%稀疏度如果我们从1开始，则需要动态SDDMM内核，其基于W（Q）中的条目的幅度生成稀疏模式，并产生简化和压缩的注意力得分矩阵。 好处是我们可以显式地从W（Q）中选择重要条目而无需预测。softmax是一个单调递增的函数，从图4. 密集矩阵乘法分块设计。同时修剪、编码和消耗稀疏矩阵。其次，模式应该保留每行中的最大值。我们发现以前研究中使用的稀疏模式不符合这三个要求。不受约束的逐行top-k稀疏性违反了由GEMM内核中的di线程块生成的元素之间的比较。 因此，它的修剪和编码很难并行。此外，流行的基于压缩稀疏行（CSR）的SpMM需要超过95%的稀疏度才能与其密集对应物[3]相提并论，这太高而无法保持准确性。在以前的研究中广泛使用的块稀疏性违反了第二个要求，因为它不能保证所选择的块包含它所覆盖的所有行中的最大值。我们男人-第2.3节中提到的满足这两个要求。 对于需求1，在修剪和编码期间很容易并行。 由于N：M选择是在本地执行的，并且写入非零和元数据的地址是确定性的。当它被后续的内核消耗时，可以很容易地实现 由于行长度减少到原始注意力得分矩阵的N/M，softmax内核在没有修改器Core的情况下实现了加速，具有值矩阵的SpMM在N：M = 1：2的情况下也实现了1.7的加速。对于要求2，选择N从没有任何好处，但扔掉每个1M向量中的元素保证保留加速（>5C<0G. 在本文中，我们选择从基于两每一行中的最大条目虽然我们在本文中关注1：2和2：4注意事项。首先，从开始允许我们保持我们的机架GPU，也支持其他N：M比率，设计简单，使得它不引入额外的开销或超参数来调谐。第二，我们在下一节中开发的动态SD-DMM内核在50%稀疏性下比密集的运行更快。3.3设计选择2：稀疏结构为了在保持准确性的同时实现加速，我们需要对稀疏模式进行结构化约束，使其满足两个要求。首先，它应该对GPU友好表1.SQuADv1.1上的F1评分（不含Finetune）充分一比二二比四93. 17 ±0。2792. 86 ±0。2293. 00 ±0。16硬件支持N：M稀疏矩阵和密集矩阵之间的乘法。经验上，我们发现这种模式可以很好地近似于完全注意机制。 我们首先在SQuAD v1.1上对BERT大型模型进行了全面的关注。然后，我们直接用1：2和2：4注意力代替完全注意力，而不需要额外的网络调整。 F1评分总结见表1，≤95%。即使没有网络调整，精度损失也只有4动态SDDMM内核在这一节中，我们将介绍我们的动态新粒度SDDMM内核的设计我们的内核融合了密集的GEMM，1NNtileBfloat32bfloat1616K1632K姆蒂莱Warp3Warp2C一23Warp1Warp0一B动态N：M细粒度结构稀疏A机制PPoPP37301212⇥⇥⇥关于我们b（）/c⇥⇥⇥--关于我们在N：M稀疏性下基于幅度的修剪，并通过为CUTLASS [10]GEMM主循环开发新的动态修剪尾声来压缩到单个内核中。 平铺如图4所示。这三个阶段的融合节省了全局存储器和寄存器之间的两次往返。此外，我们提出了一种专门的线程和寄存器映射设计，确保所有的内存访问是128 B合并，并减少跨线程shu这些优化允许在GPU上实现N：M稀疏下的动态SDDMM。4.1SpMM核中稀疏矩阵的解码我们编码的稀疏矩阵应该很容易成功-block. 因此，只需要单个;30CA8G来加载所需线程映射下的元数据4.2动态修剪尾声SpMM核中稀疏矩阵的解码过程给我们的动态SDDMM核带来了两个约束首先，我们需要在编码元数据时执行两个逆转换。其次，由于逆变换发生在tf 32类型的32 16和f16/bf16类型的32 32的范围内，因此最小修剪单元为32 64字节。有了这两个约束，我们提出了动态prun-ing尾声如图6所示。有四个主要步骤：修剪每个连续8B数据的50%，生成非零值正在使用SpMM内核进行解码。图5（c）显示了线程和元数据;反转换（1）：交错两个带有tf32的mma.sp.m16n8k16元数据行乘以8;counter-transformation（2）：Switchtype或mma.sp.m16n8k32，在Tensor Core的硬件限制下使用f16，bf16类型每排有八个沿着次对角线的元数据。nonzero到全局内存。写入元数据，4位元数据。 我们使用）C_（10）A403I3G A46I3G来注释映射。每个A46I3G代表一个16位寄存器。 满足C ≤A403I3G<>3 32 2<>3 2 == 0的螺纹在其他螺纹中保持稳定。sp指令有一个参数用于在第一行和第二个16行之间进行选择。在这个线程映射下直接加载元数据将导致一个INE内存访问模式：每个线程详细地，基于它们的幅度选择4个2字节数据中的2个，并且将唯一的4比特元数据分配给0中的每个组合。 图6（b）中列举了选择模式的对应元数据。值得注意的是，随着oat32数据类型，每个32位数据占用两个连续的2字节时隙。因此，它只支持0x4和0xe下的模式。 在生成4比特元数据之后，它们中的连续四个被连接到2B元数据块。的每个指令仅加载16位，并且共享存储体元数据行以8英寸的间隔交错，骗局发生了。 为了解决这个问题，CUTLASS [10]使用了ldmatrix指令。如图5（a）所示，：2 [0， 7]加载128位数据，将其分为四个32位数据并将它们分发给线程4k<$4k +3，并带有特殊数据3BC_A>F= bA>F/32c=3 2+（A>F%8）=4+b（A>F%32）/8c。（二）在2中，右上方和左下方的元数据块每22格的开关。最后，在第三部分中，元数据硬件中的路径如图5（b）和（c）所示，为了匹配产生的数据被写入全局内存，对于具有MMA.SP的LDMatrix的线程映射，需要两个附加的变换：（1）在每2 × 2块中切换次对角条目（2）在每32行中将行交错4。CUTLASS通过在编码元数据时执行它们的逆变换来消除这两个变换o（1）在每32行中以8行交错行（2）在每2^2中切换次对角线条目以列为主的交错格式，跨距为4字节。这可以通过将两个连续的元数据解释为int对象，然后以列为主的方式将其写入DRAM来实现。非零值被写入row-major下的全局内存4.3线程和寄存器映射在修剪步骤中，将扭曲瓦片划分为（一）T0T1T2T7（b）第（1）款R\C01T0{0，1}T1{0，1}23T2{0，1}T3{0，1}45T4{0，1}T5{0，1}67T6{0，1}T7{0，1}89T8{0，1}T9{0，1}1011T10{0，1}T11{0，1}..2829T28{0，1}T29{0，1}3031T30{0，1}T31{0，1}（c）第（1）款..8x4位3264个由经线一次一个地穿过的小块。删除GEMM结果。为了减少跨线程的数据访问，我们在修剪过程中将所有比较都放在每个线程的本地。图7（a）中显示了32 6 4字节块中的线程和寄存器映射 ，threadId A46 I3分别表示具有oat32类型和b类型的线程C的A403I3的A46 I3 32位/16位寄存器。在oat32 type下，我们选择相邻两个条目中较大的一个。在b类型下，我们选择2较大人从相邻4条目而Ampere GPU中mma的输出映射匹配我们在oat32类型下，输出映射在b图5. CUTLASS中的元数据编码[10]。）C A403I3G A46I3G表示在线程C A403I3G C 16位寄存器A46I3G。图8（a）中的类型与我们的不一致因此，我们需要额外的经纱舒rst通过这4个条目到同一线程，然后比较它们，并获得2个较大的3R\C01T0{0}T0{1}T1{0}T1{1}23T2{0}T2{1}T3{0}T3{1}45T4{0}T4{1}T5{0}T5{1}67T6{0}T6{1}T7{0}T7{1}9T8{0}T8{1}T9{0}T9{1}11T10{0}T10{1}T11{0}T11{1}..29T28{0}T28{1}T29{0}T29{1}3031T30{0}T30{1}T31{0}T31{1}R\C01T0{0}T4{0}T1{0}T5{0}2..T8{0}T9{0}789T28{0}T0{1}T29{0}T1{1}10T4{1}T8{1}T5{1}T9{1}15T28{1}T29{1}1617T2{0}T6{0}T3{0}T7{0}18..T10{0}T11{0}2324T30{0}T2{1}T31{0}T3{1}2526..T6{1}T10{1}T7{1}T11{1}31T30{1}T31{1}PPoPP陈赵东，郑，余英安，刘柳，丁宇飞，谢元3740-001非零031..2423..1615..87..032BR123密集数据公司简介（一）非零：16B密集数据：32BLSB元数据：2BMSB01469BeF0123456789一 BC DeF0x40x80xd0xeR32B32B0..78..1516..2324..31R2B2B00-00-1..77-07-188-08-1..1515-015-11616-016-1..2323-023-12424-024-1..3131-031-1R2B2B00-00-118-08-1216-016-1324-024-1..287-07-12915-015-13023-023-13131-031-1R2B2B00-08-010-18-1216-024-1316-124-1..287-015-0297-115-13023-031-03123-131-1（b）第（1）款0 4 8 122B×44位元数据112116120124127...0xe0xd0x90xc0x80x4图6. 修剪密集数据并生成非零元数据。一个。这将引入额外的开销。 为了解决这个问题，当将矩阵B加载到共享内存时，我们只需在开始时操作指向全局内存的指针，就可以将列交错。 图8（b）中显示了到寄存器的映射结果，它相当于图7（a）b。交错后，连续的四个数据自然由同一个线程保存，我们从中选择两个较大的数据。 为了减少分支发散，通过比较任意两个数据的和来进行选择。生成元数据和非零值。 对于oat和b数据。在4.2节之后，我们需要通过按位OR将连续的四个元数据连接到一个16位元数据块中。 由于相邻的四个元数据由di个线程拥有，我们首先将线程4 t +k的4位元数据放置到int16寄存器中的[k 4：k 4 3]位，如图8（b）所示，其中我们将每个4位元数据表示为“T thread_id { register_id }[ bit_id ]"。 然后，我们与warpshu跨线程共享这些int16寄存器，并生成图7（c）。经纱舒的目的线是为了使整个经纱忙碌。图7表2. BERT-large SQuAD v1.1的F1评分（CI =95%）模型w/o netunew/netuneTransformer（Transformer（b93. 22 ±0。1593. 34 ±0。3193. 17 ±0。2793. 18 ±0。272016 - 02 - 22 01：02：02（D2：4（b）92. 86 ±0。2293. 00 ±0。1693. 07 ±0. 1793. 28 ± 0。295评价在本节中，我们首先评估我们的动态跨DI域任务的NE粒度结构化稀疏注意机制。然后，我们在NVIDIA A100 GPU上，在di序列长度从256到4096的情况下，证明了我们可以在任意序列长度下实现实际的加速比。 我们重点讨论了1：2和2：4稀疏度，因为它们的加速比可以直接在现成的硬件上进行评估，其他N：M稀疏度留待将来研究。5.1模型精度为了表明我们的方法在综合sce- narios中是有效的，我们首先评估了模型在任务上的准确性，(d) 以及（e）在图6之后和图6中示出了结果。D1结构域和序列长度。对于以下型号值得注意的是，这两个步骤仅改变元数据的逻辑映射，并且寄存器分配不是一个逻辑映射，因此这两个步骤不需要代码最后，我们需要将元数据和非零值写入全局内存，图6中如图7（e）所示，每行由同一线程的连续两个int16寄存器保存因此，我们可以简单地将其重新解释为= int 32对象，并将元数据写入column-major中的全局内存对于非零值，我们只需将它们合并到共享内存中，然后以row-major写入全局内存“b数据类型，我们rst netune他们从预训练模型下“提供了一个更精确的梯度，有助于收敛。后在netuning中，我们直接将模型中的所有参数为对于问题分类和掩蔽语言建模任务，我们报告了在di随机种子下8次运行的平均结果。提问提问。 我们在序列长度为384的SQuAD v1.1上评估BERT-large。 我们在Huggingface [26]中使用“bert- large-uncased-whole-word-masking”，320-08-00-18-116-0 24-1 16-1 24-07-015-0 7-115-1 23-0 31-0 23-1 31-1动态N：M细粒度结构稀疏A机制PPoPP375（一）float32bfloat16位浮点数R\C 01234 5678910 11 12 1314 150T0{0，1}T1{0，1}T2{0，1}T3{0，1}T0{4，5}T1{4，5}T2{4，5}T3{4，5}1T4{0，1}T5{0，1}T6{0，1}T7{0，1}T4{4，5}T5{4，5}T6{4，5}T7{4，5}2..7T28{0，1}T29{0，1}T30{0，1}T31{0，1}T28{4，5}T29{4，5}T30{4，5}T31{4，5}8T0{2，3}T1{2，3}T2{2，3}T3{2，3}T0{6，7}T1{6，7}T2{6，7}T3{6，7}9T4{2，3}T5{2，3}T6{2，3}T7{2，3}T4{6，7}T5{6，7}T6{6，7}T7{6，7}10..15T28{2，3}T29{2，3}T30{2，3}T31{2，3}T28{6，7}T29{6，7}T30{6，7}T31{6，7}16T0{0，1}T1{0，1}T2{0，1}T3{0，1}T0{4，5}T1{4，5}T2{4，5}T3{4，5}17T4{0，1}T5{0，1}T6{0，1}T7{0，1}T4{4，5}T5{4，5}T6{4，5}T7{4，5}18..23T28{0，1}T29{0，1}T30{0，1}T31{0，1}T28{4，5}T29{4，5}T30{4，5}T31{4，5}24T0{2，3}T1{2，3}T2{2，3}T3{2，3}T0{6，7}T1{6，7}T2{6，7}T3{6，7}R\C2019 - 05 -15 00：00：0016-1920-23 24-27 28-310T1 {0，1，4，5} T2{0，1，4，5} T3{0，1，4，5}T1 {0，1，4，5} T2{0，1，4，5} T3{0，1，4，5}1T5 {0，1，4，5} T6{0，1，4，5} T7{0，1，4，5}T5 {0，1，4，5} T6{0，1，4，5} T7{0，1，4，5}2..7T28{0，1，4，5} T29{0，1，4，5} T30{0，1，4，5} T31T28{0，1，4，5} T29{0，1，4，5} T30{0，1，4，5} T318T1 {2，3，6，7} T2{2，3，6，7} T3{2，3，6，7}T1 {2，3，6，7} T2{2，3，6，7} T3{2，3，6，7}9T4{2，3，6，7} T5{2，3，6，7} T6{2，3，6，7}T4{2，3，6，7} T5{2，3，6，7} T6{2，3，6，7}10..15T28{2，3，6，7} T29{2，3，6，7} T30{2，3，6，7} T31T28{2，3，6，7} T29{2，3，6，7} T30{2，3，6，7} T3116T1 {0，1，4，5} T2{0，1，4，5} T3{0，1，4，5}T1 {0，1，4，5} T2{0，1，4，5} T3{0，1，4，5}17T5 {0，1，4，5} T6{0，1，4，5} T7{0，1，4，5}T5 {0，1，4，5} T6{0，1，4，5} T7{0，1，4，5}18..23T28{0，1，4，5} T29{0，1，4，5} T30{0，1，4，5} T31T28{0，1，4，5} T29{0，1，4，5} T30{0，1，4，5} T3124T1 {2，3，6，7} T2{2，3，6，7} T3{2，3，6，7}T1 {2，3，6，7} T2{2，3，6，7} T3{2，3，6，7}25T4{2，3，6，7} T5{2，3，6，7} T6{2，3，6，7}T4{2，3，6，7} T5{2，3，6，7} T6{2，3，6，7}26..（b）（c）（d）（e）（f）]图7.线程、寄存器和数据、元数据之间的映射。预训练模型，并将其与默认con进行netune，HuggingFace1中的guration。利用“Transformer”的检查点，分别得到“1：2（”和“2：4（”的F1得分。“Transformer（”的F1得分是通过直接利用“Dfss1：2（”和“Dfss2：4（”的检查点，使用密集注意机制进行推理得到的。第1页https://github.com/huggingface/transformers/tree/master/示例/pytorch/问答如表2所示，通过网络调整，我们的1：2稀疏性只有0.1 F1分数损失，小于标准偏差。我们的2：4稀疏度甚至比密集基线实现了一点点性能改进。一个合理的解释是，虽然2：4稀疏度可以保留大多数重要边缘，但它偶尔也会丢弃一小部分重要边缘，其作用类似于注意力丢弃技术 [29] 。 此 外 ， 直 接 将 我 们 的 方 法 应 用 到 密 集 的Transformer没有netuning也取得了类似的结果。它只是很好地近似了密集注意机制。R\C0T0{0}[0]T1{0}[1]T2{0}[2]T3{0}[3]T0{2}[0]T1{2}[1]T2{2}[2]T3{2}[3]1T4{0}[0]T5{0}[1]T6{0}[2]T7{0}[3]T4{2}[0]T5{2}[1]T6{2}[2]T7{2}[3]2..7T28{0}[0]T29{0}[1]T30{0}[2] T31{0}[3]T28{2}[0]T29{2}[1]T30{2}[2] T31{2}[3]8T0{1}[0]T1{1}[1]T2{1}[2]T3{1}[3]T0{3}[0]T1{3}[1]T2{3}[2]T3{3}[3]9T4{1}[0]T5{1}[1]T6{1}[2]T7{1}[3]T4{3}[0]T5{3}[1]T6{3}[2]T7{3}[3]10..15T28{1}[0]T29{1}[1]T30{1}[2] T31{1}[3]T28{3}[0]T29{3}[1]T30{3}[2] T31{3}[3]16T0{4}[0]T1{4}[1]T2{4}[2]T3{4}[3]T0{6}[0]T1{6}[1]T2{6}[2]T3{6}[3]17T4{4}[0]T5{4}[1]T6{4}[2]T7{4}[3]T4{6}[0]T5{6}[1]T6{6}[2]T7{6}[3]18..23T28{4}[0]T29{4}[1]T30{4}[2] T31{4}[3]T28{6}[0]T29{6}[1]T30{6}[2] T31{6}[3]24T0{5}[0]T1{5}[1]T2{5}[2]T3{5}[3]T0{7}[0]T1{7}[1]T2{7}[2]T3{7}[3]25T4{5}[0]T5{5}[1]T6{5}[2]T7{5}[3]T4{7}[0]T5{7}[1]T6{7}[2]T7{7}[3]26..31T28{5}[0]T29{5}[1]T30{5}[2] T31{5}[3]T28{7}[0]T29{7}[1]T30{7}[2] T31{7}[3R\C0T0{0}T1{2}1T4{0}T5{2}2T8{0}T9{2}..7T28{0}T29{2}8T0{1}T1{3}9T4{1}T5{3}10T8{1}T9{3}..15T28{1}T29{3}16T2{4}T3{6}17T6{4}T7{6}18T10{4}T11{6}..23T30{4}T31{6}24T2{5}T3{7}25T6{5}T7{7}26T10{5}T11{7}..31T30{5}T31{7}R\C0T0{0}T1{2}1T0{1}T1{3}2T2{4}T3{6}3T2{5}T3{7}4T4{0}T5{2}5T4{1}T5{3}6T6{4}T7{6}7T6{5}T7{7}8T8{0}T9{2}9T8{1}T9{3}10T10{4}T11{6}11T10{5}T11{7}..28T28{0}T29{2}29T28{1}T29{3}30T30{4}T31{6}31T30{5}T31{7}R\C0T0{0，1}1T1{2，3}2T2{4，5}3T3{6，7}4T4{0，1}5T5{2，3}6T6{4，5}7T7{6，7}8T8{0，1}9T9{2，3}10T10{4，5}11T11{6，7}..28T28{0，1}29T29{2，3}30T30{4.5}31T31{6，7}R\C0T0{0，1}1T1{0，1}2T2{0，1}3T3{0，1}4T4{0，1}5T5{0，1}6T6{0，1}7T7{0，1}8T8{0，1}9T9{0，1}10T10{0，1}11T11{0，1}..28T28{0，1}29T29{0，1}30T30{0，1}31T31{0，1}PPoPP陈赵东，郑，余英安，刘柳，丁宇飞，谢元376⇠⇥⇠⇥⇠ ⇥ ⇠⇥（一）dst_c l =（b）第（1）款×）+c l * d）+×）+×/04. 我们的方法在所有四个长序列的基准上都达到了相当的精度5.2加速比在本节中，我们将展示我们的方法在di序列长度上实现的加速比，并将其与现有研究进行比较。内 核 加 速 。 我 们 首 先 将 我 们 的 动 态 SDDMM 内 核 与cuSPARSELt库进行比较，cuSPARSELt库提供了用于将输入矩阵修剪和编码为2：4稀疏度的API。 我们比 较 了cuDNN密集GEMM、cuDNN密集GEMM之后cuSPARSELt修剪以及表5中的我们的SD-DMM内核在di序列长度下的执行时间。批量大小、头数和嵌入大小为8，分别为16和64我们的内核加速了GEMM +图8. 交错矩阵B的列，以减少b的修剪期间的跨通道数据共享屏蔽语言建模。我们还评估了我们的模型，在Wikitext-2和Wikitext-103序列长度512下的掩码建模任务。与问答任务类似，我们选择“roberta-large” 作 为 预 训 练 模 型 ， 并 在 默 认 conHuggingface 2下对其进行netune。结果总结见表