CN111708511A

CN111708511A - 用于神经网络的数据压缩

Info

Publication number: CN111708511A
Application number: CN202010192465.7A
Authority: CN
Inventors: D·R·加西亚·加西亚
Original assignee: Nvidia Corp
Current assignee: Nvidia Corp
Priority date: 2019-03-18
Filing date: 2020-03-18
Publication date: 2020-09-25
Also published as: US20200302284A1; EP3713093A1

Abstract

本申请提供了一种用于神经网络的数据压缩。用于生成数据集的代表值的系统和方法，通过首先压缩数据集中的一部分值来确定第一公共值以及进一步压缩这一部分值的子集来确定第二公共值。代表值是通过取第一公共值和第二公共值之间的差生成的，其中代表值对应于第一和第二公共值与该一部分值的子集内的每个值之间的数学关系。相比于第一和第二公共值，代表值需要更少的存储。

Description

用于神经网络的数据压缩

背景技术

在存储压缩数据以实现所需存储度量时，能够压缩和解压缩数据以及进一步减少位数会涉及存在各种技术挑战。与确定压缩数据元素的压缩数值表示以及进一步存储这些压缩数值表示(通常由大量位表示)相关的算法可能会超过或以其他方式压倒与存储设备关联的存储容量。而且，由于这些数值表示的大小，与训练神经网络相关的计算成本也会受到影响。此外，在给定带宽限制和/或计算环境中的限制的情况下，需要大量位的数值表示通常会导致与在设备内或不同设备之间传输数据相关的复杂问题。

附图说明

参考附图来对各种技术进行描述，其中：

图1根据实施例说明了计算环境，在其中执行了块浮点(Block Floating Point，BFP)编码器以生成一个或更多个压缩数值表示；

图2根据实施例说明了包括由指数、符号和尾数表示的值的集合(例如，根据块浮点格式)的矩阵的图以及另一个包括由指数、压缩数值表示、符号和尾数表示的值的集合(根据具有压缩数值表示的块浮点格式)的矩阵的图；

图3根据实施例说明了示例性的矩阵、子矩阵(或块)、块的子矩阵(或四元组)和标量的图；

图4根据实施例说明了在其中确定块和四元组(quad)的公共指数值的图；

图5根据实施例说明了在其中确定与块中和四元组中的值之间的数学关系相对应的压缩数值表示(或指示符)的图；

图6根据实施例显示了用于生成与数据集中的指数值之间的数学关系相对应的压缩数值表示(或指示符)的过程的说明性示例；

图7根据实施例显示了用于生成与数据集中的指数值之间的数学关系相对应的压缩数值表示(或指示符)的过程的另一说明性示例；

图8根据实施例说明了并行处理单元(“PPU”)的示例；

图9根据实施例说明了通用处理集群(“GPC”)的示例；

图10根据实施例说明了存储器分区单元的示例；

图11根据实施例说明了流式多处理器的示例；以及

图12根据实施例说明了在其中可以实现各种实施例的计算机系统。

具体实施方式

在一个实施例中，根据本公开实现的系统和方法用于在数据集中生成压缩数值表示，该压缩数值表示同时非常适合表示深度神经网络中的激活、权值和梯度，并同时可以在硬件中有效地实现。特别地，在一个实施例中，压缩数值表示适合矩阵乘法的有效实现，因为这些乘法主要构成了训练神经网络的计算成本的大部分。在一个实施例中，压缩数值表示包括整数或其他数值，但在某些实施例中，它包括值的指示符、确定该值的指示符和/或确定值的数学关系的指示符。

也就是说，在一个实施例中，根据本公开实现的系统和方法用于在多步骤过程中压缩数据集，其中对于数据集中的每个元素，通过找到可以存储一次而不是多次的值来首先对整个数据集进行压缩。在一个实施例中，然后，以相同的方式对数据集的各个子集进行压缩，以找到(对于每个子集)第二值，然后不存储该第二值，而是存储如何从第一值到第二值获得的指示符。该指示符通常比第二值本身占用更少的空间(更少的位)，因此是一种更有效的存储数据集的值的方法。

换句话说，在一个实施例中，接收数据集中的多个值作为具有“m”行和“n”列的矩阵，其中矩阵的每个值(或标量)由普通的浮点数(或二进制值)表示。标量是用浮点格式近似一个实数值的计算机表示。在一个实施例中，接收数据集中的多个值作为多维张量(四维张量)。在一个实施例中，张量是标量的多维数组，包括但不限于R×C元素的矩阵。每个普通的浮点数都由位序列表示，其中位的一部分(通常是单个位)表示符号，位的另一部分表示指数，以及位的另一部分表示尾数。例如，在32位序列中，符号用单个位“1”表示，指数用8位“01111111”表示，以及尾数用23位“11000000000000000000000”表示，将生成“1.11.”的二进制值。除了浮点数(如对数表示)外，还可以实现针对值的不同格式的其他表示，浮点数只是本文描述的一个实施例。

在一个实施例中，矩阵被细分为N×M标量(例如，4×4标量)的非重叠图块(tile)的均匀网络。在一个实施例中，将矩阵划分为均匀网格的每个非重叠图块被称为子矩阵(或块)。这些图块可大小相同，也可不相同。矩阵中的元素通常被平铺以改善中间缓存存储中的局部性。更具体地说，该矩阵除以预定义大小的图块中的权重矩阵并在图块内共享指数。图块限制共享指数的值的数目。在一个实施例中，通过确定与该块的元素共有的信息而相互独立地压缩每个块，其使用更少的位的方式存储条目。这可以通过从块的每个元素的指数减去一个值来实现，例如块中所有元素的最大指数。还可以使用其他方法来减少存储每个元素的位或实现其他优点，例如提取公共位、因式分解或使用其他方式。在一个实施例中，最大或公共指数值称为“block_exp”。作为说明性示例，在一个实施例中，如果块中的每个指数至少为n(n是正整数)，则可以从块中的每个元素中减去n，以便可以使用更少的位来村存储块中的每个元素。

随后，在一个实施例中，每个块被进一步细分为较小的P×Q标量(例如，2×2标量)的非重叠的子图块的均匀的网格。在一个实施例中，这些非重叠的子图块的均匀的网格中的每一个被称为“四元组”。在一个实施例中，块中的每个四元组通过确定四元组中所有元素公共值(或指数值)来进行压缩，如前所述，例如在结合压缩矩阵的块时执行。在一个实施例中，四元组的公共值(或公共指数值)被称为“quad_exp”。在一个实施例中，设每个四元组属于同一块，quad_exp的值小于或等于block_exp。此外，在一个实施例中，只要N×M的块大小足够小，则很有可能block_exp和max_exp之间的差异足够小，从而可以进行额外的压缩。

在一个实施例中，不是让系统直接存储quad_exp，而是通过存储quad_exp和block_exp之间的差(或差的指示符)来获得更有效的编码。在某些情况下，该差可以称为每四指数增量，压缩数值表示或代理(替换quad_exp)，其提供信息以允许系统从block_exp中确定quad_exp。在一个实施例中，此差表示为具有预定位数的无符号整数。在一个实施例中，如果实际的差大于用预定位数所能表示的值，则该差被饱和为最大的代表值。可是，在一些实施例中，不存储差或差值本身，而且确定并存储数学关系的指示符，和/或，确定并存储用于确定quad_exp和block_exp与该四元组的所有标量中的值之间的数学关系的指示符。

本文所描述的技术用于生成提供信息的差值或指示符以指示标准块指数之间的数学关系。这些差值的存在增加了可以在子矩阵(或块)中支持的值的动态范围，同时与为每个标量具有单独的指数相比，仍然具有更低的存储器开销。更大的动态范围反过来意味着，子矩阵(或块)中的标量可以用比如果不存在差值时数值稳定性所需要的尾数位更少的尾数位来表示。因此，存储这些差值所需的存储空间小于存储每个标量的单个指数和单独指数所需的空间量。

除了提供了存储指数的更有效的方法外，本文所描述的技术还方便地允许计算矩阵的转置，由于子矩阵(或块)和四元组的规则结构意味着可以计算转置而不需要进一步的舍入或量化误差。此外，通过实施本文所描述的技术以及所节省的空间，透过广泛的各种神经网络和/或各种图形处理单元(GPU)它将大约提高两倍(2倍)的端到端训练软件配置。

在前面和下面的说明书中，描述了各种技术。以解释的目的，提出了具体的配置和细节以对实施这些技术的可能方法提供全面的理解。然而，很明显下面描述的技术可以在不同的配置中实践，而没有具体的细节。此外，公知的特征可以省略或简化，以避免模糊所描述的技术。

图1说明了计算环境100，其中块浮点(BFP)编码器104与前端服务102、数据存储106、压缩器108以及解压器110一起在计算设备130中实现以生成和存储一个或更多个压缩数值表示。一个或更多个压缩数值表示可以称为代表值、共同指数值、参考指数或差值。在一个实施例中，根据实施例，一个或更多个压缩数值表示被简单地称为数据集的指示符140。在一个实施例中，计算环境100是与神经网络相关联的表示，其中神经网络使用该压缩数值表示，以减少在神经网络中执行的计算成本。

在一个实施例中，计算设备130包括并将本文提供的所有描述与图形处理单元(GPU)相关联。在一个实施例中，计算设备130被配置为接收或获取输入数据120并生成差值，或至少指示符140，该指示符包括表示输入数据120的数据集的块指数值之间的数学关系的信息。计算设备130可以通过网络180(无线网络或有线网络)从计算环境外部或内部的另一个计算设备、计算服务、虚拟机、虚拟服务(未在图1中描述)接收输入数据120。

在一个实施例中，计算设备130包括前端服务102。前端服务102可以是配置为接收或获取输入数据120的计算机系统或其组件。例如，前端服务102可以是模拟到数字(ADC)转换器。在一个实施例中，ADC是能够接收或获取模拟信号作为输入并将这些模拟信号转换为数字信号的计算系统。模拟信号可以是提供给计算设备130的、由麦克风拾取的信号或进入数码相机的光。在一个实施例中，然后将ADC配置为将模拟信号转换为数字信号，然后数字信号为以位表示的实数，以形成数据集或矩阵(由行和列组成)。如本文关于图1所描述的，ADC是前端接收服务/计算设备的说明性实施例，其可以在采取、接收或获取输入数据120的计算设备130中实现。在另一个实施例中，输入数据120已经是数字信号并且这些信号已经在到达前端服务102(例如ADC)之前转化为矩阵中的值。因此，ADC可能被绕过或从计算设备130中减去。在一个实施例中，前端服务102不是接收信号的ADC，而是计算服务/设备，例如服务器计算机系统、服务器集群、虚拟计算机系统、虚拟服务器、虚拟运行时环境、容器环境、无服务器执行环境、服务托管系统，或使用并与计算设备130相关联以接收输入数据120的任何合适的计算实体。

在一个实施例中，一旦信号已经被转换成数字信号或已经被验证为数字信号，该数字信号被表示为数据集中的值。在一个实施例中，这些值被存储为矩阵中的输入数据120。在一个实施例中，计算设备130包括BFP编码器104。在一个实施例中，BFP编码器104被称为BFP生成器。可以配置BFP编码器104以从前端服务102接收或获取值的输入数据120集。也就是说，输入数据120可以使用浮点表示来表示实数。因此，如本文所述，块浮点或BFP编码器104被用于提供算术算法，该算法将来自输入数据120集的每个数据块指派一个指数。这样，块中每个元素的每个单独指数都不必存储。

因此，在一个实施例中，BFP编码器104对输入数据120集执行第一步操作，以识别输入数据120集的块中的所有值中的公共(或最大)指数值。具体地，在一个实施例中，如果输入数据120集的块是(sign)a*2^n格式的数字(例如，2的n次幂的倍数，其可以通过存储表示符号、a和n的比特来表示)，并且每个元素具有至少“100”的“n”，则可以将“100”标识为公共。执行该操作是因为为每个元素存储2^(n-100)所需的空间较少。例如，先前解决方案的示例试图为数据集生成公共或最大指数值，例如是BF16格式浮点数(bfloat16)(也称为e8m7)，和各种形式的块指数表示。

如在另一实施例中，如果输入数据120集的块中的条目是a*2^70、b*2^59、c*2^65和d*2^73，则公共元素将是73(因为所有条目的指数最大为73)。这将导致通过存储73、a*2^-3、b*2^-14、c*2^-8和d*2^0来使用与条目相关联的更小的存储空间，因为存储公共指数73和根据每个指数差右移每个尾数相比于分别存储每个元素的指数来说占用更小的空间。

在又一实施例中，如果输入数据120集的块中的条目是a*2^70、a*2^59、a*2^65和a*2^73，则公共值可以是“a”，因为每个元素中都有“a”；公共值可以是a*2^59，因为a*2^59可以从每个元素中因式分解出来；公共值可以是向量(a，59)，因为每个元素具有系数a和至少为59的指数。实际上，通过分解数据集的条目之间的公共值，可能需要更少的存储空间。

此外，在一个实施例中，BFP编码器104进一步执行第二步操作，以在输入数据120集的块的值的子集150中识别第二公共指数值。也就是说，在一个实施例中，在根据如上所述的从数据集的值识别公共指数值之后的后续步骤中，BFP编码器104再次对块的各个子集执行与第一步中相同的操作。也就是说，遵循上面的示例，BFP编码器104可以将“95”标识为块的子集150的每个元素的公共(因为子集中的每个数的形式是a*2^n，n至多为“95”)。

由于识别两个公共指数值(一个来自输入数据120集的块，另一个来自输入数据120集的块的子集150)，可以确定或计算差值。差值可以存储为指示符140，其中指示符140将向系统提供如何从输入数据120集的块的公共指数值获得数据集的子集150的公共指数值的信息。因此，该差值可以简单地存储为“5”(因为它指示系统将需要从“100”减去“5”，即在第一步中标识的整个块的公共指数值)。因此，当向BFP编码器104提供指令或信息以使指示符140被确定并最终以值“5”存储时，可以实现更大的压缩，因为对于子集150，存储“5”仅一次占用的空间小于存储“95”。

在一个实施例中，压缩器108向BFP编码器104提供关于如何执行压缩输入数据120集的块的第一步骤以及进一步如何执行压缩输入数据120集的块的各个子集的第二步骤的信息或指令。压缩器可以提供指令使得BFP编码器104选择输入数据120集中的第一数量的元素作为要被首先压缩的块，并且进一步选择另一数量的元素(其小于要被再次压缩的块)。压缩器108可至少部分基于用户输入和/或系统策略来识别并向BFP编码器104提供指令，以确定如何压缩数据集及其子集以最终确定差值。在一个实施例中，如图1所示，压缩器108不是与BFP编码器104分离的单独设备或服务，而是与BFP编码器104合并以执行上述操作。

具体地，作为示例性实施例，BFP编码器104确定来自输入数据120集(例如，矩阵)的块可以具有4×4标量的大小。一旦块被压缩，就可以确定第一公共指数或最大指数。随后，从压缩器108到BFP编码器104的指令可以指示被压缩的块的子集。在一个实施例中，块的子集或四元组具有2×2标量的大小。压缩四元组后，将确定第二公共或最大指数值。计算并进一步存储第一公共指数值和第二公共指数值之间的差值。换句话说，差值本身可以是存储的指示符，这允许系统基于第一公共指数值确定第二公共指数值和块的子集中的值。然而，在一个实施例中，不存储该差值，而是存储该差值的几种指示符。也就是说，在确定差值之后，系统可以简单地存储第一公共指数值和第二公共指数值与子集的值之间的数学关系的指示符，而不必存储第二公共指数值本身。具体地，在一个实施例中，指示符是可用于导出第一公共指数值和第二公共指数值的函数。

在一个实施例中，计算设备130包括数据存储106。数据存储106被配置为接收和存储数据集的公共指数值的存储设备。具体地，在一个实施例中，数据存储106被配置成存储差值、对应于第一公共值和第二公共值与值的一个或更多个子集内的每个值之间的数学关系的指示符140、和/或指示如何确定第一和第二公共指数值与值一个或更多个子集内的每个值之间的数学关系的指示。用于存储差值、指示符和/或第一公共指数值和第二公共指数值之间的数学关系的指示的存储要求将小于用于单独存储矩阵的条目或元素的每个指数值的存储要求。

在一个实施例中，计算设备130包括解压器110。解压器110可被配置成接收具有公共指数值的压缩浮点数，并将每个浮点数分解成字节数组，形成字节矩阵。压缩浮点数可以是由解压器110从BFP编码器104接收的，和/或它可以来自数据存储106。然后，在一个实施例中，解压器110对矩阵进行转置从而产生新的矩阵，其中第一行显示低熵(指数)，最后一行具有高熵(最后尾数位)。换句话说，解压器110接收压缩数据集以及存储的差值和/或指示符，并且被配置成解压缩数据集以生成数据输出。也就是说，解压器110可以被配置为接收一组指令和/或算法，这些指令和/或算法在执行时使得解压器110将压缩数据集扩展回其原始形式(例如，在压缩值之前，例如，通过有损压缩技术或无损压缩技术重新生成矩阵120中的值)。在一个实施例中，解压器110在对压缩数据集执行附加的数学运算之后对压缩数据集进行解压缩，并且因此生成了与矩阵120中的原始值不同的数据输出。

图2示出了根据实施例的矩阵的图200，该矩阵在填充时包括由公共指数202、符号204和尾数206表示的值的集。在一个实施例中，图200是具有使用浮点表示来表示实数的数据集的矩阵的表示。在一个实施例中，浮点数以位表示，其中，例如，位的子序列表示数字“2”的指数。具体地，数字可以表示为(sign)*a*2^n，它们的浮点表示具有符号(sign)、n和a的位。在一个实施例中，图200中的一个值由一系列位表示，其中公共指数202为“1000001”、符号204为“0”、尾数206为“01100000000000000000000”。此值的结果是十进制数字1.0011x2^4＝22。

在一个实施例中，如果矩阵的4x4大小的块中的每个元素的指数最多为“k”，则系统执行如上的关于图1所述的第一操作并存储一次“k”，其将被存储为公共指数202。此外，在一个实施例中，该系统存储具有指数“k-n”的矩阵的每个值，其比具有指数“n”的存储占用更少的比特。

此外，图2根据实施例示出了矩阵的另一图210，该矩阵包括由指数202、压缩数值表示208、符号204和尾数206(或具有压缩数值表示的块浮点)表示的值的集。在一个实施例中，图2的第二图210被细分为2x2子矩阵，并且在执行如上所示的相同处理之后，为所细分的集确定第二公共指数。所细分的集的大小可以是预定的、任意的、随机的和/或由用户输入确定的。在一个实施例中，不存储每个2x2子矩阵(或子块)的第二公共指数。相反，为这些2x2子矩阵中的每一个子矩阵计算压缩数值表示208。压缩数值表示208可以通过从公共指数202减去第二公共指数来计算，并存储为差值(或增量值)。压缩数值表示208的大小可以是用数字比特Q表示的预定比特量，其小于为每个2x2子矩阵(或子块)计算的第二公共指数的大小。换句话说，预定比特量的大小将不大于第二公共指数的比特。

换言之，从这些2x2子矩阵中的每一个子矩阵中，执行从块中提取公共指数值202的相同过程，但不存储可被拉出的指数，而是存储已经提取的指数与来自第一步操作的第一公共指数值之间的压缩数值表示208。在一个实施例中，该压缩数值表示208不是代表2x2子矩阵中的每一个子矩阵的整数，而是确定压缩数值表示208的某种指示符。压缩数值表示208通常很小，因此需要较少的比特来存储。在一个实施例中，压缩数值表示208只是所存储的函数的指示符，以便能够导出第二公共指数值。

图3根据实施例示出了示例性矩阵302(类似于图1的输入数据120集)、子矩阵(或块)304、块的子矩阵(或四元组)306和标量308的图。在一个实施例中，使用浮点表示来表示数据集，以将实数表示为矩阵302的元素。也就是说，在一个实施例中，矩阵302具有24行和16列。在一个实施例中，矩阵302被配置成不同数量的行和列。

在一个实施例中，子矩阵(或块)304是矩阵302细分为16个元素的块(例如，4x4标量)。块大小可随机确定或由用户输入预先确定。此外，块大小还可以至少部分地基于用户输入或与矩阵相关联的元数据来确定。

在一个实施例中，大小为2x2标量的四元组306是子矩阵(或块)304的子部分。四元组大小可随机确定或由用户输入预先确定。此外，还可以至少部分地基于用户输入或与矩阵相关联的元数据来确定四元组大小。

在一个实施例中，标量308是矩阵302的子部分、一个元素或单个元素。每个标量308可以是以位表示的浮点数。

图4根据实施例示出了图400，其中为矩阵的子矩阵确定公共指数值。在一个实施例中，子矩阵(或块)402是从图3中的矩阵302提取的大小为4x4标量的值的数据集。

在一个实施例中，如上文关于图1所述，块402及块402的大小由BFP编码器确定。在一个实施例中，块402及其大小由用户输入和/或系统策略确定。可以在块402中的所有值中确定第一公共或最大指数值406。在一个实施例中，第一公共或最大指数值406具有10位的大小。第一公共指数值406可以表示至少共享该第一公共指数值或最大指数值的符号和尾数404。这样，只为块402存储一次公共指数值406，而不必单独存储块402中每个元素的每个指数值。也就是说，不存储块402中16个标量的16个指数值，例如，而是仅存储块402中所有元素中的单个共享指数值。

在一个实施例中，通过将块402细分为2x2标量来确定四元组408。如上所述，四元组408是块402的子集，四元组的大小可以预先确定或选择，使得四元组408的公共指数值412至少与块402的公共指数值412相同或更小的概率很高。一旦确定了四元组408，接下来的步骤将执行相同的压缩操作(与对块402执行的压缩操作相同)，以确定四元组408的公共指数值412。四元组408的公共指数值412将是四元组408的符号和尾数410的指数值。

如上所述，确定块402的公共指数值406和四元组408的公共指数值412之间的差。然后，该差值被用作两个公共指数值406、412之间的关系或数学关系的指示符。差值通常很小(小于块402的公共指数值406的大小和/或四元组408的公共指数值412的大小)，并且可以被指派为表示为“q位”整数。用户输入和/或系统策略可以预先确定“q位”的大小或数量。在一个实施例中，连续地执行确定差值的这个迭代，直到确定矩阵和子矩阵的所有差值为止。迭代次数可能取决于起始的矩阵的大小。

图5根据实施例示出了图500，其中确定了与块和四元值中的公共值506、512之间的数学关系相对应的差值(或指示符516)。也就是说，根据关于图4的描述，在块(矩阵的子矩阵)502提取或确定其符号和尾数504的公共指数值506(例如，第一公共值)之后，公共指数值506可以存储在数据存储中(图5中未示出，但如图1中所述)。随后，可以确定块的子矩阵(例如，四元组)508，使得为四元组508的符号和尾数510确定另一公共指数值512(例如，第二公共值)。在一个实施例中，另一个公共指数值512不存储在数据存储中。相反，确定公共指数值506和另一公共指数值512之间的差值(例如，第三值)，并将其存储在数据存储中。差值可以是整数，其通常比另一个公共指数值512小。在一个实施例中，差值(或指示符516)的大小为6位。差值(或指示符516)的大小可以变化，因为这里描述的6位大小只是示例，并且可以存储其他可能大小的差值(或指示符516)。在一个实施例中，该差值不是整数，而是如何确定该差值的指示。因此，差值(或指示符516)仅仅是系统在必要时通过至少部分基于先前生成的公共指数值506使用计算方法514或的数学计算来确定另一公共指数值512的代理。

图6根据实施例示出了用于生成差值(或指示符)的过程600的说明性示例。在一个实施例中，系统具有至少一个或更多个处理器(或一个或更多个算术逻辑单元(ALU)处理器)，被配置成包括接收或获取由值组成的数据集的逻辑602。此外，一个或更多个处理器还包括在第一集合值的每个值内确定第一公共值(例如，第一公共指数值或第一最大指数值)的逻辑604。在一个实施例中，一个或更多个处理器还被配置成包括在第一集合值的一个或更多个值子集的每个值内确定第二公共值(例如，第二公共指数值或第二最大指数值)的逻辑606。此外，一个或更多个处理器还被配置成包括存储对应于第一和第二公共值与一个或更多个值子集内的每个值之间的数学关系的第三值的逻辑，其中，第三值所需的存储少于第一和第二公共值608。在一个实施例中，第三值是预定大小的整数。在一个实施例中，第三值只是可用于导出第一公共值和第二公共值的指示符或函数。

图7根据实施例示出了用于生成差值或指示符的过程的另一示例。在一个实施例中，系统具有至少一个或更多个处理器，被配置为包括接收或获取数据集的逻辑。数据集可以被压缩以确定对数据集的所有值公共的第一值702。此外，处理器还可以包括压缩数据集的子集以确定与子集的值公共的第二值的逻辑704。一旦确定了第一值和第二值，就可以确定并存储两个值之间的差706。差值可以是第一值和第二值与子集的每个值之间的数学关系的指示符。换言之，差值可以是代理(例如，代替第二值)，当与第一值一起使用时，该代理为系统提供信息以确定第二值和子集的值。与存储第一和/或第二值相比，此差值需要的存储空间更少。

图8示出了根据一个实施例的并行处理单元(“PPU”)800。在一个实施例中，使用机器可读代码配置PPU 800，如果由PPU执行该代码，则使PPU执行贯穿本公开描述的部分或全部过程和技术。在一个实施例中，PPU 800是多线程处理器，它包括在一个或更多个集成电路设备上实现的一个或更多个算术逻辑单元(ALU)，并利用多线程作为延迟隐藏技术，被设计为并行地处理多线程上的计算机可读指令(也称为机器可读指令或简单地为指令)。在一个实施例中，线程指执行的线程，且是被配置为由PPU 800执行的一组指令的实例。在一个实施例中，PPU 800是图形处理单元(“GPU”)，被配置成实现用于处理三维(“3D”)图形数据的图形渲染管线，以便生成二维(“2D”)图像数据，用于显示在显示设备上，诸如液晶显示(LCD)设备。在一个实施例中，PPU 800用于执行诸如线性代数运算和机器学习运算的计算。图8示出了仅用于说明目的的示例性并行处理器，并且应被解释为本公开的范围内设想的处理器架构的非限制性示例，并且可以使用任何合适的处理器来补充和/或替代该处理器。

在一个实施例中，一个或更多个PPU 300被配置为加速高性能计算(HPC)、数据中心和机器学习应用程序。在一个实施例中，PPU 800被配置为加速深度学习系统和应用程序，包括以下非限制性示例：自动驾驶汽车平台、深度学习、高精度语音、图像、文本识别系统、智能视频分析、分子模拟、药物发现、疾病诊断，天气预报、大数据分析、天文学、分子动力学模拟、金融建模、机器人、工厂自动化、实时语言翻译、在线搜索优化和个性化用户推荐等。

在一个实施例中，PPU 800包括输入/输出(“I/O”)单元805、前端单元810、调度器单元812、工作分配单元814、集线器816、交叉开关(“Xbar”)820、一个或更多个通用处理集群(“GPC”)818和一个或更多个分区单元822。在一个实施例中，PPU 800通过一个或更多个高速GPU互连808连接到主机处理器或其他PPU 800。在一个实施例中，PPU 800通过互连802连接到主机处理器或其他外围设备。在一个实施例中，PPU 800连接到包括一个或更多个存储器设备804的本地存储器。在一个实施例中，本地存储器包括一个或更多个动态随机存取存储器(“DRAM”)设备。在一个实施例中，一个或更多个DRAM设备被配置为和/或是可配置的为高带宽存储器(“HBM”)子系统，其中多个DRAM裸晶(die)堆叠在每个设备内。

高速GPU互连808可以指基于线的多通道通信链路，其由系统使用以扩展并且包括与一个或更多个CPU结合的一个或更多个PPU 800，支持PPU 800与CPU之间的高速缓存一致性，以及CPU主控。在一个实施例中，数据和/或命令通过高速GPU互连808经由集线器816发送到PPU 800的其它单元/从PPU 800的其它单元发送，例如一个或更多个复制引擎、视频编码器、视频解码器、电源管理单元和其它在图8中可能未明确示出的组件。

在一个实施例中，I/O单元805被配置成通过系统总线802从主机处理器(图8中未示出)发送和接收通信(例如命令、数据)。在一个实施例中，I/O单元805直接经由系统总线802或通过一个或更多个中间设备(例如存储器网桥)与主机处理器通信。在一个实施例中，I/O单元805可经由系统总线802与一个或更多个其它处理器(例如一个或更多个PPU 800)通信。在一个实施例中，I/O单元805实现外围组件互连快速(“PCIe”)接口以用于通过PCIe总线通信。在一个实施例中，I/O单元805实现用于与外部设备通信的接口。

在一个实施例中，I/O单元805对通过系统总线802接收的分组进行解码。在一个实施例中，至少一些分组表示被配置成使PPU 800执行各种操作的命令。在一个实施例中，I/O单元805将解码的命令发送到由命令指定的PPU 800的各种其它单元。在一个实施例中，命令被发送到前端单元810和/或发送到集线器816或PPU 800的其他单元，例如一个或更多个复制引擎、视频编码器、视频解码器、电源管理单元等(图8中未明确示出)。在一个实施例中，I/O单元805被配置成在PPU 800的各个逻辑单元之间路由通信。

在一个实施例中，由主机处理器执行的程序在缓冲器中编码命令流，缓冲器将工作负载提供给PPU 800进行处理。在一个实施例中，工作负载包括由这些指令处理的指令和数据。在一个实施例中，缓冲器是存储器中可被主机处理器和PPU 800访问(例如，读/写)的区域-主机接口单元可被配置成通过I/O单元805经系统总线802发送的存储器请求来访问连接到系统总线802的系统存储器中的缓冲器。在一个实施例中，主机处理器将命令流写入缓冲器，然后将指向命令流开始的指针发送到PPU 800，使得前端单元810接收指向一个或更多个命令流的指针并管理一个或更多个命令流，从流读取命令并将命令转发到PPU 800的各个单元。

在一个实施例中，前端单元810耦合到调度器单元812，调度器单元812配置各种GPC 818以处理由一个或更多个流定义的任务。在一个实施例中，调度器单元812被配置成跟踪与调度器单元812管理的各种任务相关的状态信息，其中状态信息可以指示任务被分配给哪个GPC 818、任务是活动的还是非活动的、与任务相关联的优先级等等。在一个实施例中，调度器单元812管理在一个或更多个GPC 818上的多个任务的执行。

在一个实施例中，调度器单元812耦合到工作分配单元814，工作分配单元814被配置成分派任务以在GPC 818上执行。在一个实施例中，工作分配单元814跟踪从调度器单元812接收到的多个调度任务，并且工作分配单元814为每个GPC 818管理待处理(pending)任务池和活动任务池。在一个实施例中，待处理任务池可以包括若干时隙(例如，32个时隙)，其中包含分配给特定GPC 818处理的任务；活动任务池可以包括若干时隙(例如，4个时隙)，用于正在由GPC 818主动处理的任务，目的是当GPC 818完成任务的执行时，该任务从GPC818的活动任务池中逐出，并且来自待处理任务池的其他任务之一被选择和调度以在GPC818上执行。在一个实施例中，如果活动任务在GPC 818上处于空闲状态，例如在等待数据依赖性被解决时，则活动任务从GPC 818中逐出并返回到待处理任务池，而待处理任务池中的另一个任务被选择并调度以在GPC 818上执行。

在一个实施例中，工作分配单元814经由XBar 820与一个或更多个GPC 818通信。在一个实施例中，XBar 820是将PPU 800的许多单元耦合到PPU 800的其他单元的互连网络，并且可以被配置为将工作分配单元814耦合到特定GPC 818。尽管未明确示出，PPU 800的一个或更多个其它单元也可以经由集线器816连接到XBar 820。

这些任务由调度器单元812管理，并由工作分配单元814分派到GPC 818。GPC 818被配置为处理任务并生成结果。结果可被GPC 818内的其它任务消耗、经由XBar 820路由到不同的GPC 818或存储在存储器804中。结果可以经由分区单元822写入存储器804，分区单元822实现用于从/向存储器804读写数据的存储器接口。结果可经由高速GPU互连808发送到另一PPU 804或CPU。在一个实施例中，PPU 800包括数量U个分区单元822，其等于耦合到PPU 800的独立和不同的存储器设备804的数量。下面将结合图10更详细地描述分区单元822。

在一个实施例中，主机处理器执行实现应用程序编程接口(“API”)的驱动内核，该API使在主机处理器上执行的一个或更多个应用程序能够调度在PPU 800上执行的操作。在一个实施例中，PPU 800同时执行多个计算应用程序，以及PPU 800为多个计算应用程序提供隔离、服务质量(“QoS”)和独立的地址空间。在一个实施例中，应用程序生成指令(例如，以API调用的形式)，使得驱动内核生成一个或更多个任务供PPU 800执行，并且驱动内核将任务输出到PPU 800正在处理的一个或更多个流。在一个实施例中，每个任务包括一组或更多组相关线程，这些线程可以称为线程束。在一个实施例中，线程束包括可并行执行的多个相关线程(例如，32个线程)。在一个实施例中，协作线程可以指多个线程，包括执行任务的指令以及通过共享存储器交换数据的指令。根据实施例，结合图10更详细地描述了线程和协作线程。

图9根据一个实施例示出了GPC 900，例如图8的PPU 800示出的GPC。在一个实施例中，每个GPC 900包括用于处理任务的多个硬件单元，并且每个GPC 900包括管线管理器902、预光栅操作单元(“PROP”)904、光栅引擎908、工作分配交叉单元(“WDX”)916、存储器管理单元(“MMU”)918、一个或更多个数据处理集群(“DPC”)906以及任何合适的部分的组合。应当理解，图9的GPC 900可以包括代替或附加于图9所示的单元的其他硬件单元。

在一个实施例中，GPC 900的操作由管线管理器902控制。管线管理器902管理一个或更多个DPC 906的配置以用于处理分配给GPC 900的任务。在一个实施例中，管线管理器902配置一个更或多个DPC 906中的至少一个以实现图形渲染管线的至少一部分。在一个实施例中，DPC 906被配置成在可编程流式多处理器(“SM”)914上执行顶点着色程序。在一个实施例中，管线管理器902被配置成将从工作分配接收的分组路由到GPC 900内的适当逻辑单元，并且一些分组可以被路由到PROP 904和/或光栅引擎908中的固定功能硬件单元，而其他分组可以被路由到DPC 906以供图元引擎912或SM 914处理。在一个实施例中，管线管理器902配置一个更或多个DPC 906中的至少一个以实现神经网络模型和/或计算管线。

在一个实施例中，PROP单元904被配置成将由光栅引擎908和DPC 906生成的数据路由到存储器分区单元中的光栅操作(“ROP”)单元，如上文所详细地描述。在一个实施例中，PROP单元904被配置成执行颜色混合的优化、组织像素数据、执行地址转换等。在一个实施例中，光栅引擎908包括被配置为执行各种光栅操作的若干固定功能硬件单元，并且光栅引擎908包括设置引擎、粗光栅引擎、剔除引擎、裁剪引擎、精细光栅引擎和图块聚合引擎及上述引擎的任意适当组合。在一个实施例中，设置引擎接收变换后的顶点并生成与由顶点定义的几何图元相关联的平面方程，该平面方程被发送到粗光栅引擎以生成图元的覆盖信息(例如，图块的x、y覆盖掩码)；粗光栅引擎的输出被发送到剔除引擎，其中与未通过z-测试的图元相关联的片段被剔除，并且被发送到裁剪引擎，其中位于视锥体之外的片段被裁剪掉。在一个实施例中，那些经过裁剪和剔除后留下来的片段可以被传递到精细光栅引擎，以基于由设置引擎生成的平面方程生成像素片段的属性。在一个实施例中，光栅引擎908的输出包括由任何合适实体处理的片段，例如由在DPC 906中实现的片段着色器处理的片段。

在一个实施例中，包括在GPC 900中的每个DPC 906包括M管线控制器(“MPC”)910、图元引擎912、一个或更多个SM 914以及它们的任何适当组合。在一个实施例中，MPC 910控制DPC 906的操作，将从管线管理器902接收到的分组路由到DPC 906中的适当单元。在一个实施例中，与顶点相关联的分组被路由到图元引擎912，图元引擎912被配置为从存储器中获取与顶点相关联的顶点属性；相反，与着色器程序相关联的分组可以被发送到SM 914。

在一个实施例中，SM 914包括被配置为处理由多个线程表示的任务的可编程流式处理器。在一个实施例中，SM 914是多线程的，并且被配置为同时执行来自特定线程组的多个线程(例如32个线程)并且实现SIMD(单指令，多数据)体系结构，其中一组线程(线程束)中的每个线程配置为基于同一组指令处理不同的数据集。在一个实施例中，线程组中的所有线程执行相同的指令。在一个实施例中，SM 914实现了SIMT(单指令，多线程)体系结构，其中一组线程中的每个线程被配置为基于同一组指令处理不同的数据集，但是在执行期间允许线程组中的单个线程发散。在一个实施例中，为每个线程束维护程序计数器、调用栈和执行状态，当线程束内的线程发散时，使线程束和线程束中的串行执行之间的并发成为可能。在另一个实施例中，为每个单独的线程维护程序计数器、调用栈和执行状态，从而在线程束内和线程束之间的所有线程之间实现相等的并发。在一个实施例中，为每个单独的线程维护执行状态，并且执行相同指令的线程可以被收敛和并行执行以获得最大效率。在一个实施例中，下面更详细地描述SM 914。

在一个实施例中，MMU 918提供GPC 900和存储器分区单元之间的接口，MMU 918提供虚拟地址到物理地址的转换、存储器保护和存储器请求的仲裁。在一个实施例中，MMU918提供一个或更多个转换后备缓冲器(“TLB”)，用于执行虚拟地址到存储器中的物理地址的翻译。

图10根据一个实施例示出了PPU的存储器分区单元。在一个实施例中，存储器分区单元1000包括光栅操作(“ROP”)单元1002、二级(“L2”)高速缓存1004、存储器接口1006以及它们的任何适当组合。存储器接口1006耦合到存储器。存储器接口1006可以实现用于高速数据传输的32、64、128、1024位数据总线等。在一个实施例中，PPU包括U个存储器接口1006，每对分区单元1000一个存储器接口1006，其中每对分区单元1000连接到相应的存储器设备。例如，PPU可以连接到多达Y个存储器设备，例如高带宽存储器堆叠或图形双数据速率、版本5、同步动态随机存取存储器(“GDDR5 SDRAM”)。

在一个实施例中，存储器接口1006实现HBM2存储器接口，并且Y等于一半U。在一个实施例中，HBM2存储器堆栈与PPU位于同一物理包上，与传统GDDR5 SDRAM系统相比提供了大量的功率和面积节省。在一个实施例中，每个HBM2堆叠包括四个存储器裸晶并且Y等于4，其中HBM2堆叠包括每个裸晶两个128位通道，总共8个通道和1024位的数据总线宽度。

在一个实施例中，存储器支持单纠错校正双纠错检测(“SECDED”)纠错码(“ECC”)以保护数据。对于数据损坏敏感的计算应用程序ECC提供了更高的可靠性。在大规模集群计算环境中，PPU处理非常大的数据集和/或长时间运行应用程序，可靠性尤其重要。

在一个实施例中，PPU实现了多级存储器层次。在一个实施例中，存储器分区单元1000支持统一存储器，以便为CPU和PPU存储器提供单个统一的虚拟地址空间，从而在虚拟存储器系统之间实现数据共享。在一个实施例中，跟踪PPU对位于其它处理器上的存储器的访问频率，以确保存储器页被移动到更频繁地访问页的PPU的物理存储器。在一个实施例中，高速GPU互连808支持地址转换服务，允许PPU直接访问CPU的页表，并提供由PPU对CPU存储器的完全访问。

在一个实施例中，复制引擎在多个PPU之间或PPU和CPU之间传输数据。在一个实施例中，复制引擎可以为未映射到页表中的地址生成页故障，然后存储器分区单元1000服务该页故障，将地址映射到页表中，然后复制引擎执行传输。在一个实施例中，针对多个处理器之间的多个复制引擎操作固定存储器(即，不可分页)，实质上减少了可用存储器。在一个实施例中，随着硬件页出错，地址可以被传递到复制引擎，而不考虑存储器页是否驻留，并且复制过程是透明的。

根据一个实施例，来自图8的存储器或其他系统存储器的数据由存储器分区单元1000获取并存储在位于芯片上并且在各个GPC之间共享的L2高速缓存1004中。在一个实施例中，每个存储器分区单元1000包括与相应的存储器设备相关联的L2高速缓存960的至少一部分。在一个实施例中，在GPC内的各个单元中实现低层缓存。在一个实施例中，SM 1040中的每一个可以实现一级(“L1”)高速缓存，其中L1高速缓存是专用于特定SM 1040的专用存储器，并且来自L2高速缓存1004的数据被提取并存储在每个L1高速缓存中以在SM 1040的功能单元中处理。在一个实施例中，L2高速缓存1004耦合到存储器接口1006和XBar 820。

在一个实施例中，ROP单元1002执行与像素颜色相关的图形光栅操作，例如颜色压缩、像素混合等。在一个实施例中，ROP单元$$50与光栅引擎1025一起实现深度测试，从光栅引擎1025的剔除引擎接收与像素片段相关联的样本位置的深度。在一个实施例中，测试与片段关联的样本位置相对于深度缓冲区中的对应深度的深度。在一个实施例中，如果片段通过了样本位置的深度测试，则ROP单元1002更新深度缓冲器并将深度测试的结果发送到光栅引擎1025。应当理解，分区单元1000的数目可以不同于GPC的数目，因此，在一个实施例中，每个ROP单元1002可以耦合到每个GPC。在一个实施例中，ROP单元1002跟踪从不同GPC接收的分组，并确定由ROP单元1002生成的结果经由Xbar路由到哪个GPC。

图11根据一个实施例示出了流式多处理器，例如图9的流式多处理器。在一个实施例中，SM 1100包括：指令高速缓存1102；一个或更多个调度器单元1104；寄存器文件1108；一个或更多个处理核心1110；一个或更多个特殊功能单元(“SFU”)1112；一个或更多个加载/存储单元(“LSU”)1114；互连网络1116；共享存储器/L1高速缓存1118以及其中任何合适的组合。在一个实施例中，工作分配单元分派任务以在PPU的GPC上执行，并且每个任务被分配给GPC内的特定DPC，并且如果该任务与着色器程序相关联，则该任务被分配给SM 1100。在一个实施例中，调度器单元1104从工作分配单元接收任务并管理分配给SM 1100的一个或更多个线程块的指令调度。在一个实施例中，调度器单元1104将用于执行的线程块调度为并行线程的线程束，其中每个线程块被分配至少一个线程束。在一个实施例中，每个线程束执行线程。在一个实施例中，调度器单元1104管理多个不同的线程块，将线程束分配给不同的线程块，然后在每个时钟周期中将来自多个不同协作组的指令分派给各个功能单元(例如，核心1110、SFU 1112和LSU 1114)。

协作组可以指用于组织通信线程组的编程模型，其允许开发者表达线程正在进行通信所采用的粒度，使得能够表达更丰富、更高效的并行分解。在一个实施例中，协作启动API支持线程块之间的同步性，以执行并行算法。在一个实施例中，常规的编程模型的应用程序为同步协作线程提供了单一的简单结构：跨线程块的所有线程的栅栏(barrier)(例如，syncthreads()函数)。然而，程序员通常希望以小于线程块粒度的粒度定义线程组，并在所定义的组内同步，以集体的全组功能接口(collective group-wide functioninterface)的形式使能更高的性能、设计灵活性和软件重用。协作组使程序员能够在子块(即，小到单个线程)和多块粒度明确地定义线程组，并对协作组中的线程执行集体操作(如同步)。该编程模型支持跨软件边界的干净组合，因此库和实用程序函数可以在其本地上下文中安全地同步，而无需对收敛性进行假设。协作组图元支持协作并行新的模式，包括生产者-消费者并行、机会主义并行以及跨整个线程块网格的全局同步。

在一个实施例中，分派单元1106被配置成向一个或更多个功能单元发送指令，并且调度器单元1104包括两个分派单元1106，使得在每个时钟周期期间能够分派来自同一线程束的两个不同指令。在一个实施例中，每个调度器单元1104包括单个分派单元1106或附加分派单元1106。

在一个实施例中，每个SM 1100包括寄存器文件1108，该寄存器文件为SM 1100的功能单元提供一组寄存器。在一个实施例中，在每个功能单元之间划分寄存器文件1108，使得每个功能单元被分配到寄存器文件1108的专用部分。在一个实施例中，寄存器文件1108在SM 1100正在执行的不同线程束之间划分，并且寄存器文件1108为连接到功能单元的数据路径的操作数提供临时存储。在一个实施例中，每个SM 1100包括多个L个处理核心1110。在一个实施例中，SM 1100包括大量(例如，128个或更多)不同的处理核心1110。在一个实施例中，每个核心1110包括完全管线化的、单精度、双精度和/或混合精度处理单元，其包括浮点运算逻辑单元和整数运算逻辑单元。在一个实施例中，浮点运算逻辑单元实现用于浮点运算的IEEE 754-2008标准。在一个实施例中，核心1110包括64个单精度(32位)浮点核心、64个整数核心、32个双精度(64位)浮点核心和8个张量核心。

根据实施例，张量核心被配置成执行矩阵运算。在一个实施例中，一个或更多个张量核心包括在核心1110中。在一个实施例中，张量核心被配置成执行深度学习矩阵算法，例如用于神经网络训练和推断的卷积运算。在一个实施例中，每个张量核心对4x4矩阵进行运算，并执行矩阵乘法和累加运算D＝A×B+C，其中A、B、C和D是4x4矩阵。

在一个实施例中，矩阵乘法输入A和B是16位浮点矩阵，累积矩阵C和D是16位浮点或32位浮点矩阵。在一个实施例中，张量核心对16位浮点输入数据和32位浮点累积进行运算。在一个实施例中，16位浮点乘法需要64次运算，并产生全精度乘积，然后使用32位浮点加法与用于4x4x4矩阵乘法的其他中间乘积进行累积。在一个实施例中，张量核心用于执行由这些较小元素构建的更大的二维或更高维矩阵运算。在一个实施例中，API(如CUDA9C++API)公开了专门的矩阵加载、矩阵乘法和累加运算以及矩阵存储操作，以有效地使用来自CUDA-C++程序的张量核心。在一个实施例中，在CUDA层面，线程束级接口假定16×16尺寸矩阵跨越线程束的全部32个线程。

在一个实施例中，每个SM 1100包括执行特殊函数(例如，属性评估、倒数平方根等)的M个SFU 1112。在一个实施例中，SFU 1112包括被配置成遍历层次树数据结构的树遍历单元。在一个实施例中，SFU 1112包括被配置成执行纹理映射过滤操作的纹理单元。在一个实施例中，纹理单元被配置成从存储器中加载纹理贴图(例如，2D纹理数组)，并对纹理贴图进行采样以产生采样纹理值，以用于SM 1100执行的着色器程序中。在一个实施例中，纹理映射存储在共享存储器/L1高速缓存中。根据一个实施例，纹理单元实现纹理操作，例如使用mip映射(例如，不同细节级别的纹理映射)的过滤操作。在一个实施例中，每个SM 1100包括两个纹理单元。

在一个实施例中，每个SM 1100包括N个LSU 1054，其实现共享存储器/L1高速缓存1006和寄存器文件1108之间的加载操作和存储操作。在一个实施例中，每个SM 1100包括将每个功能单元连接到寄存器文件1108以及将LSU 1114连接到寄存器文件1108、共享存储器/L1高速缓存1006的互连网络1116。在一个实施例中，互连网络1116是交叉开关，其可以被配置为将任何功能单元连接到寄存器文件1108中的任何寄存器，以及将LSU 1114连接到寄存器文件和共享存储器/L1高速缓存1115中的存储器位置。

在一个实施例中，共享存储器/L1高速缓存1118是片上存储器阵列，其允许SM1100与图元引擎之间以及SM 1100中的线程之间的数据存储和通信。在一个实施例中，共享存储器/L1高速缓存1118包括128KB的存储容量，并且位于从SM 1100到分区单元的路径中。在一个实施例中，共享存储器/L1高速缓存1118用于高速缓存读和写。共享存储器/L1高速缓存1118、L2高速缓存和存储器中的一个或更多个正在备份存储。

在一个实施例中，将数据高速缓存和共享存储器功能结合到单个存储器块中，以为两种类型的存储器访问提供改进的性能。在一个实施例中，该容量可由程序用作或者是可用的不使用共享存储器的高速缓存，例如，如果共享存储器被配置为使用一半的容量，则纹理和加载/存储操作可以使用剩余的容量。根据实施例，在共享存储器/L1高速缓存1118内的集成使得共享存储器/L1高速缓存1118能够作为流数据的高吞吐量管线来工作，同时提供对频繁重复使用的数据的高带宽和低延迟访问。当配置为通用并行计算时，与图形处理相比，可以使用更简单的配置。在一个实施例中，绕过了固定功能的图形处理单元，创建了更简单的编程模型。在一个实施例中，在通用并行计算配置中，工作分配单元将线程块直接指派和分配给DPC。根据一个实施例，块中的线程执行相同的程序，在计算中使用唯一的线程ID以确保每个线程生成唯一的结果，使用SM 1100执行程序并执行计算，使用共享存储器/L1高速缓存1118以在线程之间通信，以及使用LSU 1114通过共享存储器/L1高速缓存1118和存储器分区单元读写全局存储器。在一个实施例中，当被配置为通用并行计算时，SM1100写入调度器单元可用于在DPC上启动新工作的命令。

在一个实施例中，PPU包括在台式计算机、膝上型计算机、平板计算机、服务器、超级计算机、智能电话(例如，无线、手持设备)、个人数字助理(“PDA”)、数字相机、车辆、头戴式显示器、手持电子设备等中或与之耦合。在一个实施例中，PPU设置在单个半导体衬底上。在一个实施例中，PPU与一个或更多个其它设备(诸如附加PPU、存储器、精简指令集计算机(“RISC”)CPU、存储器管理单元(“MMU”)、数模转换器(“DAC”)等)一起包括在片上系统(“SoC”)中。

在一个实施例中，PPU可以包括在包括一个或更多个存储器设备的图形卡上。图形卡可以配置为与台式计算机的主板上的PCIe插槽接口。在又一实施例中，PPU可以是包括在主板的芯片组中的集成图形处理单元(“iGPU”)。

图12根据一个实施例示出了可以在其上实现各种架构和/或功能的计算机系统1200。在一个实施例中，计算机系统1200被配置成实现贯穿本公开所描述的各种过程和方法。

在一个实施例中，计算机系统1200包括至少一个中央处理单元1202，所述中央处理单元1202连接到使用任何合适协议实现的通信总线1210，例如PCI(外围组件互连)、PCI-Express、AGP(加速图形端口)、超传输(HyperTransport)或任何其他总线或点对点通信协议。在一个实施例中，计算机系统1200包括主存储器1204，控制逻辑(例如，以硬件、软件或其组合实现)和存储在主存储器1204上的数据，主存储器1204可以采用随机存取存储器(“RAM”)的形式。在一个实施例中，网络接口子系统1222提供到其他计算设备和网络的接口，用于从计算机系统1200接收数据并将数据发送到其他系统。

在一个实施例中，计算机系统1200包括输入设备1208、并行处理系统1212和显示设备1206，其可以使用传统CRT(阴极射线管)、LCD(液晶显示器)、LED(发光二极管)、等离子显示器或其他合适的显示技术来实现。在一个实施例中，从输入设备1208(诸如键盘、鼠标、触摸板、麦克风等)接收用户输入。在一个实施例中，上述模块中的每一个可以位于单个半导体平台上以形成处理系统。

在本说明书中，单个半导体平台可指基于单半导体的集成电路或芯片。应当注意的是，术语单半导体平台还可以指具有增加的连接性的模拟片上操作的多芯片模块，并且相对于使用传统的中央处理单元(“CPU”)和总线实现作出实质性改进。当然，根据用户的需要，各个模块还可以分开放置或以半导体平台的各种组合来放置。

在一个实施例中，以机器可读可执行代码或计算机控制逻辑算法的形式的计算机程序存储在主存储器1204和/或辅助存储器中。根据一个实施例，如果由一个或更多个处理器执行计算机程序，则使得系统1200能够执行各种功能。存储器1204、存储器和/或任何其他存储器是计算机可读介质的可能性示例。辅助存储器可指任何合适的存储设备或系统，例如硬盘驱动器和/或可移动存储驱动器，表示软盘驱动器、磁带驱动器、光盘驱动器、数字多功能磁盘(“DVD”)驱动器、记录设备、通用串行总线(“USB”)闪存。

在一个实施例中，各种在先附图的架构和/或功能可在中央处理器1202、并行处理系统1212、能够同时具有中央处理器1202和并行处理系统1212的能力的至少一部分的集成电路、芯片组(例如，集成电路组，设计成作为执行相关功能的单元工作并出售等)以及任何合适的集成电路组合的上下文中实现。

在一个实施例中，各种在先附图的体系架构和/或功能可以在通用计算机系统、电路板系统、专用于娱乐目的的游戏控制台系统、专用系统等的上下文中实现。在一个实施例中，计算机系统1200可以采用如下的形式：台式计算机、膝上型计算机、平板计算机、服务器、超级计算机、智能电话(例如，无线、手持设备)、个人数字助理(“PDA”)、数字相机、车辆、头戴式显示器、手持电子设备、移动电话设备、电视、工作站、游戏机、嵌入式系统和/或任何其他类型的逻辑。

在一个实施例中，并行处理系统1212包括多个PPU 1214和相关联的存储器1216。在一个实施例中，PPU通过互连1218和交换机1220或多路复用器连接到主机处理器或其他外围设备。在一个实施例中，并行处理系统1212将计算任务分布在可并行化的PPU 1214上，例如，作为计算任务分布在多个GPU线程块上的一部分。在一个实施例中，存储器在PPU1214的部分或全部上是共享和可访问的(例如，用于读取和/或写入访问)，尽管这种共享存储器可能会因使用本地存储器和驻留在PPU上的寄存器而引起性能损失。在一个实施例中，PPU 1214的操作通过使用诸如__syncthreads()的命令来同步，该命令要求块中的所有线程(例如，跨多个PPU 1214执行)在继续之前达到代码的某个执行点。

因此，说明书和附图应被视为说明性的而不是限制性的。然而，显而易见的是，在不脱离权利要求中所述的本发明的更广泛的精神和范围的情况下，可以对其进行各种修改和改变。

其他变化在本公开的精神内。因此，虽然所公开的技术易受各种修改和替代构造的影响，但其某些示出的实施例在附图中示出并且已经在上文中详细描述。然而，应当理解的是，无意将本发明局限于所公开的一种或更多种特定形式，相反，其意图是涵盖如所附权利要求所定义的属于本发明的精神和范围的所有修改、替代结构和等同物。

在描述所公开的实施例的上下文中(特别是在以下权利要求的上下文中)，术语“一”和“一个”以及“所述”和类似的指代的使用应被解释为包括单数和复数，除非本文中另有说明或上下文明确反驳。除非另有说明，否则术语“包括”、“拥有”、“包括”和“包含”应解释为开放式术语(即，意思是“包括但不限于”)。术语“连接”，在未经修改和指代物理连接时，即使有东西介入，应被解释为部分或全部包含在、附着在或连接在一起。除非本文另有说明，否则本文中值的范围的列举仅仅是作为单独参考在范围内的每个单独值的简略表达，并且每个单独的值被并入说明书中，如同在本文中单独列举一样。除非上下文另有说明或反驳，否则术语“集”(如“项目集合”)或“子集”的使用应解释为包含一个或更多个成员的非空集合。此外，除非上下文另有说明或反驳，否则对应集合的术语“子集”不一定表示对应集合的适当子集，但该子集和对应集合可以相等。

连词(如“A、B和C中的至少一个”或“A、B、C中的至少一个”等形式的短语)，除非另有明确说明或与上下文有明显矛盾，否则应理解为通常用于表示项目、术语等的上下文可以是A、B或C，或A、B和C的集合的任何非空子集。例如，在具有三个成员的集合的示例中，连接短语“A、B和C中的至少一个”和“A、B和C中的至少一个”指下列集合中的任意一个：{A}、{B}、{C}、{A、B}、{A、C}、{B、C}、{A、B、C}。因此，这样的连词通常并不意味着某些实施例要求A中的至少一个、B中的至少一个和C中的至少一个存在。此外，除非上下文另有说明或反驳，否则术语“多个”表示复数状态(例如，“多个项目”表示多个项目)。多个项目中的项目数是至少两个，但当明确地或通过上下文这样表示时可以更多。此外，除非另有说明或上下文另有明确规定，短语“基于”是指“至少部分基于”而不是“仅基于”。

除非本文另有说明或上下文另有明确矛盾，否则本文所述过程的操作可以以任何合适的顺序执行。在一个实施例中，诸如本文描述的那些过程(或其变化和/或组合)的过程是在配置有可执行指令的一个或更多个计算机系统的控制下执行的，并且被实现为通过硬件或其组合在一个或更多个处理器上共同执行的代码(例如，可执行指令，一个或更多个计算机程序或一个或更多个应用程序)。在一个实施例中，代码例如以计算机程序的形式存储在计算机可读存储介质上，所述计算机程序包括可由一个或更多个处理器执行的多个指令。在一个实施例中，机器可读介质或计算机可读存储介质是非暂时性机器可读介质或计算机可读存储介质，其不包括暂时性信号(例如，传播的瞬时电或电磁传输)，但包括在临时信号的收发器内的非暂时性数据存储电路(例如，缓冲器、高速缓冲存储器和队列)。在一个实施例中，代码(例如，可执行代码或源代码)存储在一个或更多个非暂时性计算机可读存储介质的组合上，具有存储在其上的可执行指令(或存储可执行指令的其他存储器)，所述可执行指令在由计算机系统的一个或更多个处理器执行时(例如，作为被执行的结果)，使计算机系统执行本文所述的操作。在一个实施例中，一个或更多个非暂时性计算机可读存储介质或机器可读介质的组合包括信息，该信息如果由一个或更多个计算机指令使用，则配置计算机系统的一个或更多个处理器，并进一步使一个或更多个处理器执行本文描述的操作。在一个实施例中，非暂时性计算机可读存储介质的组合，包括多个非暂时性计算机可读存储介质和多个非暂时性计算机可读存储介质的一个或更多个单独的非暂时性存储介质缺少全部代码，而多个非暂时性计算机可读存储介质共同存储全部代码。在一个实施例中，可执行指令的执行使得不同的指令由不同的处理器执行，例如，非暂时性计算机可读存储介质存储指令和主CPU执行一些指令，而图形处理器单元执行其他指令。在一个实施例中，计算机系统的不同组件具有单独的处理器，并且不同的处理器执行指令的不同子集。

因此，在一个实施例中，计算机系统被配置成实现一个或更多个服务，其单独或共同地执行本文描述的过程的操作，并且这些计算机系统被配置有能够执行操作的适用硬件和/或软件。此外，实现本公开的实施例的计算机系统是单个设备，并且在另一实施例中是分布式计算机系统，其包括操作不同的多个设备，使得分布式计算机系统执行本文描述的操作并且使得单个设备不执行所有操作。

使用本文提供的任何和所有示例或示例性语言(例如，“例如”)仅仅是为了更好地说明本发明的实施例，并且不会对本发明的范围造成限制，除非另有要求。说明书中的任何语言都不应被解释为表示对本发明的实施至关重要的任何未声明的元素。

在本文描述的本公开的实施例，包括发明人已知的用于实施本发明的最佳模式。在阅读上述描述时，这些实施例的变化对于本领域的普通技术人员来说是显而易见的。发明人期望熟练的技术人员适当地采用这种变化，并且发明人意在以不同于本文特别描述的方式来实施本公开的实施例。因此，本公开的范围包括所附权利要求中所述的适用法律允许的主题的所有修改和等同。此外，除非本文中另有说明或上下文另有明确反驳，否则上述元素在其所有可能变化中的任何组合都包含在本公开的范围内。

本文引用的所有参考文献(包括出版物、专利申请和专利)在此通过引用合并，合并程度与单独明确指出通过引用合并的每个参考文献在本文完整列出的程度相同。

在说明书和权利要求中，可以使用术语“耦合”和“连接”及其衍生语。应当理解的是，这些术语可能不是彼此的同义词。相反，在具体的例子中，

“连接”或“耦合”可用于指示两个或更多个元件彼此直接或间接地物理或电接触。“耦合”也可能意味着两个或两个以上的元素彼此不直接接触，但仍然相互合作或交互。

除非另有特别说明，否则可以理解，在整个说明书中，例如“处理”、“计算”、“测算”、“确定”等，是指计算机或计算系统或类似电子计算设备的动作和/或过程，其将计算系统的寄存器和/或存储器中的物理量(如电子量)处理和/或转换成类似地表示为计算系统的存储器、寄存器或其它此类信息存储、传输或显示设备中的物理量的其它数据。

以类似的方式，术语“处理器”可指处理来自寄存器和/或存储器的电子数据并将该电子数据转换成可存储在寄存器和/或存储器中的其他电子数据的任何设备或设备的一部分。作为非限制性示例，“处理器”可以是中央处理单元(CPU)或图形处理单元(GPU)。“计算平台”可以包括一个或更多个处理器。如本文中所使用的，“软件”过程可以包括例如随时间执行工作的软件和/或硬件实体，例如任务、线程和智能代理。此外，每个过程可以指多个过程，用于顺序地或并行地、连续地或间歇地执行指令。这里互换地使用术语“系统”和“方法”，只要系统可以体现一个或更多个方法并且这些方法可以被认为是系统。

在本文件中，可参考获得、获取、接收或输入模拟或数字数据到子系统、计算机系统或计算机实现的机器。获得、获取、接收或输入模拟和数字数据的过程可以通过多种方式来完成，例如通过接收数据作为函数调用或对应用程序编程接口调用的参数。在一些实现方式中，通过串行或并行接口传输数据可以完成获得、获取、接收或输入模拟或数字数据的过程。在另一个实现方式中，可以通过经由计算机网络将数据从提供实体传送到获取实体来完成获得、获取、接收或输入模拟或数字数据的过程。也可以参考提供、输出、发送、发出或呈现模拟或数字数据。在各种示例中，提供、输出、发送、发出或呈现模拟或数字数据的过程可以通过将数据作为函数调用的输入或输出参数、应用程序编程接口或进程间通信机制的参数来实现。

尽管上述讨论阐述了所述技术的示例性实现方式，但是可以使用其他架构来实现所述功能，并且旨在本公开的范围内。此外，虽然为了讨论的目的在上文对责任的具体分配作了定义，但根据具体情况，各种功能和责任可能以不同的方式分配和划分。

此外，尽管主题已经用特定于结构特征和/或方法行为的语言描述，但是应当理解，所附权利要求中定义的主题不一定限于所描述的特定特征或行为。相反，所公开的具体特征和行为是实施权利要求的示例性形式。

Claims

1.一种处理器，包括：

一个或更多个算术逻辑单位(ALU)，配置为：

确定第一集合值的每个值内的第一公共值；

确定所述第一集合值的一个或更多个子集值的每个值内的第二公共值；以及

存储对应于所述第一和第二公共值与所述一个或更多个子集值中的每个值之间的数学关系的第三值，其中所述第三值比所述第一和第二公共值需要更少的存储。

2.根据权利要求1所述的处理器，其中所述第一集合值是来自浮点数的矩阵中的非重叠图块的均匀网格。

3.根据权利要求2所述的处理器，其中所述一个或更多个子集值中的每个子集是来自所述第一集合值中的非重叠图块的较小均匀网格。

4.根据权利要求1所述的处理器，其中所述第一公共值是从所述第一集合值的压缩集确定的，以及所述第二公共值是从所述一个或更多个子集值的第二压缩集确定的。

5.根据权利要求4所述的处理器，其中所述第三值包括用于所述处理器解压缩所述第一集合值的所述压缩集的信息。

6.根据权利要求1所述的处理器，其中对应于所述第一和第二公共值与所述一个或更多个子集值内的每个值之间的数学关系的所述第三值存储在存储设备中，而无需存储所述第二公共值。

7.根据权利要求6所述的处理器，其中所述第三值作为不超过预定的位数的无符号整数存储在所述存储设备中。

8.一种方法，包括：

确定第一集合值的每个值内的第一公共值；

存储对应于如何确定所述第一和第二公共值与所述一个或更多个子集值中的每个值之间的数学关系的指示的第三值，其中所述第三值比所述第一和第二公共值需要更少的存储。

9.根据权利要求8所述的方法，其中所述第一公共值是从所述第一集合值的压缩集确定的，以及所述第二公共值是从所述一个或更多个子集值的第二压缩集确定的，其中所述第三值与所述第一集合值的所述压缩集用于训练神经网络。

10.根据权利要求8所述的方法，其中所述指示是所述第二公共值与所述第一公共值之间的差值。

11.根据权利要求8所述的方法，其中所述第一集合值是浮点数矩阵的元素的数据集，其中所述浮点数中的每个浮点数的位表示指数值。

12.根据权利要求11所述的方法，其中确定所述第一公共值进一步包括确定所述第一集合值共有的第一最大指数值。

13.根据权利要求12所述的方法，其中在确定所述第一最大指数值之后，所述方法还包括：通过确定所述一个或更多个子集值共有的第二最大指数值来确定所述第二公共值。

14.一种机器可读介质，其上存储有指令集，所述指令集如果由一个或更多个处理器执行，使所述一个或更多个处理器至少：

计算第一集合值的每个值内的第一公共值；

计算所述第一集合值的一个或更多个子集值的每个值内的第二公共值；以及

存储对应于所述第一和第二公共值与所述一个或更多个子集值内的每个值之间的数学关系的第三值，其中所述第三值比所述第一和第二公共值需要更少的存储。

15.根据权利要求14所述的机器可读介质，其中存储所述第一公共值，而不存储所述第二公共值。

16.根据权利要求14所述的机器可读介质，其中所述第三值为所述第一公共值与所述第二公共值之间的差值。

17.根据权利要求16所述的机器可读介质，其中所述差值是用于从所述第一公共值确定所述第二公共值的代理。

18.根据权利要求16所述的机器可读介质，其中与所述第二公共值相比，所述差值至少具有相同或更少的位。

19.根据权利要求14所述的机器可读介质，其中所述指令集进一步包括指令，如果由所述一个或更多个处理器执行，则会使所述一个或更多个处理器压缩所述第一集合值以计算所述第一公共值。

20.根据权利要求14所述的机器可读介质，其中所述指令集进一步包括指令，如果由所述一个或更多个处理器执行，则会使所述一个或更多个处理器压缩所述一个或更多个子集值以计算所述第二公共值。

21.一种包含信息的机器可读介质，如果由一个或更多个计算机指令使用所述信息来配置一个或更多个处理器，使得所述一个或更多个处理器：首先通过压缩数据集来压缩所述数据集以确定第一值，然后通过用较小的元素替换子集的元素来压缩所述数据集的所述子集并存储第二值，其中所述第二值指示如何至少部分地基于所述较小的元素和所述第一值来计算被替换的元素。

22.根据权利要求21所述的机器可读介质，其中所述数据集包括浮点数矩阵。

23.根据权利要求21所述的机器可读介质，其中所述第一值是通过从第一集合值的每个值的指数中减去与所述数据集中所述第一集合值的所有值相关联的公共指数值来确定的。

24.根据权利要求21所述的机器可读介质，其中所述第二值是从所述子集的每个值的指数中减去与所述数据集的所述子集的所有值相关联的公共指数值来确定的。

25.根据权利要求21所述的机器可读介质，其中所述第二值比所述第一值需要更少的存储。