矩阵相乘在GPU上的终极优化：深度解析Maxas汇编器工作原理

机器之心释放

作者:王晓玉

九章超过10，000字，这篇文章可能是迄今为止对Maxas汇编程序工作原理最全面和详细的分析。

当致力于深度学习框架的实现时，我了解到Nervana有一个名为Maxas的汇编代码生成器项目，它可以生成性能超过nVidia的矩阵乘法官方版本的图形处理器机器码，因此对其工作原理产生了兴趣。

项目地址:https://github.com/NervanaSystems/maxas

它的作者斯科特·格雷(Scott Gray)在代码之外提供了详细的文档(https://github.com/NervanaSystems/maxas/wiki/SGEMM)，但是由于算法的复杂性，它的编写晦涩难懂，它的逻辑跳跃，特别是一些方法的动机没有得到解释，所以很容易迷失在细节中。

根据作者对文档的理解，这篇文章可以被视为重写，但只需要尽可能详细地解释代码的前因后果。然而，总体结构仍然是按照maxas文档来排列的，文章中的所有图片也是来自该文档。

值得注意的是，Maxas使用的算法完全依赖于Maxwell体系结构的某些特性。随着新一代GPU体系结构的发展，该项目本身已经完全过时，但其解决问题的方法仍值得学习。

背景

单精度矩阵乘法(SGEMM)应该是每个学习CUDA编程的程序员最熟悉的情况。这是自从英伟达推出CUDA的第一个版本以来，英伟达官方教程中唯一的例子。这不仅是因为SGEMM是几乎所有计算密集型软件中最重要的底层模块，而且还因为通过这个例子可以很好地演示图形处理器中的各种优化技术，尤其是图形处理器内存级别的利用。

对于两个NxN的矩阵a和b的乘法，最简单的并行方法之一是为其输出矩阵c的每个元素打开一个线程(具有相同的大小)。线程加载一行a和一列b，然后生成向量的内积。然而，问题是在GPU上访问视频存储器的延迟相当大(大约100个时钟周期)。如果一行a可以通过使用GPU的超大视频内存带宽来分割加载时间和缓存以减少延迟，从而一次加载多个元素，那么对于n千的大矩阵，b的一列中的元素的内存地址也应该是千分之一。这意味着，除了一次加载所需的列元素之外，大多数数据都是无用的。同时，这种访问模式几乎不可能有缓存行命中。总之，这种并行方法的内存访问效率极低。

为了优化它，需要共享内存。共享内存是位于GPU上的片上高速缓存，其速度可与一级高速缓存一样快，同一线程块中的线程可通过共享内存交换数据。唯一的缺点是容量有限。为了利用这一小块高速存储器，必须改进原来的并行算法:矩阵将在每个维度上分成小块，输出矩阵C可以写成:

a和B也进行同样的划分，其中不是元素而是面片矩阵，当然当面片大小为1时，面片矩阵退化为单个元素。显然，矩阵乘法的定义在这里仍然适用。

如果你把一小块看作一个元素，整个矩阵的大小就相当于缩小了两倍。每个图块的结果可以由一组线程处理，其中每个线程对应于图块中的一个元素。该线程组将A的行切片和B的列切片一个接一个地加载到共享内存中，将它们乘以共享内存中的矩阵，然后将它们叠加到原始结果上。这样，芯片中的元件在每次加载时都可以在高速共享存储器中被访问k次，从而获得远高于原始方法的存储器访问效率。

在线找到的分割算法图比CUDA官方教程中的更清晰。

以上只是对该算法的简单描述。经过这样的优化，整个算法达到了相当高的效率，是进一步优化的基础。为了理解下面这篇文章的内容，读者需要阅读CUDA的官方教程，以便非常熟悉这个切片算法，并对nVidia GPU的编程模型有更深的理解。

基本想法

如上所述，使用片上高速缓存后，切片算法不仅可以大大加快芯片矩阵的乘法速度，还可以利用计算芯片矩阵乘法的时间将下一个芯片从主存储器转移到片上共享存储器。换句话说，整个矩阵乘法的时间完全由码片矩阵乘法决定。如果要进一步提高性能，需要芯片矩阵乘法。

虽然共享内存中的矩阵乘法已经很快了，但是离硬件性能的极限还有一段距离，主要的瓶颈是两个。首先，共享内存的延迟毕竟不比寄存器好。当寄存器在Maxwell/Pascal上延迟时，需要6个时钟周期，在共享存储器上达到23个周期。此外，GPU的运算单元不能直接操作共享存储器中的数据，因此它需要一个传输指令来将其发送到寄存器，并且该mov指令将占用与实际计算指令几乎相同的时间，这相当昂贵。为了达到硬件的最高性能，我们希望GPU的每个操作单元的流水线都充满操作指令，并且计算每个时钟周期的结果。为了在现有基础上实现这一目标，maxas和以前的一些研究提出了以下方法:

1.新增加的矢量指令，一条指令可以传输四个连续的浮点数，大大减少了传输指令的数量，有利于用计算指令隐藏传输指令所消耗的时间；

2.通过交叉排列计算指令和传输指令实现预读和计算数据的流水线；

3.切片算法使用高速共享内存将需要多次访问的数据缓存在主内存中，然后在芯片矩阵中进一步切片，使用寄存器将数据缓存在共享内存中，从而进一步加速。然而，这种新的切片算法不同于以前的算法，并且带来了额外的困难。

为了实现这些方法，对GPU指令和寄存器的精确控制不再属于CUDA语言表达能力的范围，因此其实现必须由GPU本地汇编语言(而不是PTX等伪汇编语言)来完成，但这并不妨碍实现由具有更强表达能力的C等伪代码来说明。从伪代码到实际汇编代码有相当直接的转换方法，这种转换是用perl在maxas中实现的。

maxas算法概述

只考虑两个矩阵的乘法。在先前的直观算法中，根据矩阵乘法的定义，将A中的一行和B中的一列作为内积来计算C矩阵的元素。A中的一行和B中的一列被使用64次。如果想充分利用寄存器的优势，必须在寄存器中放置三个矩阵(每个矩阵占用16KB)，这对寄存器文件(每个SM 64K)来说是一个巨大的压力。更严重的问题是，与共享内存不同，寄存器不能在线程间共享。如果每个线程想把它负责的c矩阵部分保存在自己的寄存器中，它还必须保存寄存器中使用的a行和b列。结果，大量寄存器被用来存储重复数据，并且对共享存储器的访问很可能导致存储体冲突，进一步恶化延迟。

如果我们改变我们的想法，不是从输出矩阵的角度，而是从输入矩阵的角度，我们不难发现，第K列的A只用于乘以第K行的B的元素，也就是说，如果我们取第K列的A和第K行的B，我们将所有的元素乘以2，并将它们加到它们所贡献的输出矩阵元素上:

这些行和列完成了它们在矩阵乘法中的任务，并且可以被丢弃。该算法可以大大减少输入矩阵对寄存器的占用，加载后可以对数据进行加法和乘法运算，完全满足寄存器进一步缓存共享内存数据的要求。不难看出，当A的列大小和B的行大小不同时，只要它们的列索引和行索引相同，该方法仍然可以应用。

Maxas使用64个线程并行实现芯片的矩阵乘法，其中每个线程负责计算矩阵的乘积，64个线程按照布局排列，从而将芯片的大小确定为一个具有一个边和一个元素的矩阵(每个线程8个元素x8个线程)。这不同于原切片算法中每个线程的计算矩阵中的一个元素，也是充分利用寄存器超低延迟的关键。

图2。MAXAS计算两个64x64矩阵相乘的示意图。绿色4x4补丁是线程0负责的部分元素，黄色是其他线程负责的部分的左上角元素。在图中，只有左上角4x4矩阵的螺纹号被标记，其他15个只是重复。每个黑盒包含32条线，代表一条经线。在这个32x32矩阵的计算完成之后，线程0和其他线程保持它们的相对位置，并依次移动到由其他三个绿色补丁标记的位置，以完成这个64x64矩阵的计算。左边的向量是矩阵A的一列，而最上面的向量是矩阵B中相应的一行，其中标记为绿色的数据(每个8个浮点数)被线程0使用，并不难类比其他线程需要什么。

maxas的整个实现就是为这个算法服务的，在这个算法的实现过程中也遇到了本文要解决的问题。上述参数的选择，如选择64个线程的原因，是根据GPU硬件资源来确定的，以便在满足每个线程所需的寄存器资源的基础上，创建尽可能多的线程扭曲，从而调度器可以在一些线程扭曲等待数据的同时，切换其他线程扭曲进行计算。

将输入矩阵载入共享内存

为了实现上述算法，首先使用64个线程将两个输入矩阵的所需部分从主存储器加载到共享存储器中。有必要指出原文中没有提到的一个假设，即默认矩阵是按照maxas中的第一行存储的，即矩阵的每一列在内存中都是连续的，否则，下面的一系列算法就无法解释了。由于算法的加载方法不同，对原始方法进行了一些优化:

1.B矩阵使用行数据，而存储在默认行优先级矩阵中的行数据是不连续的，需要传输

设置，行到列，以便A和B加载方法可以完全相同，以减少代码的复杂性；

2.A和B矩阵被加载为一维纹理，因此不仅可以利用纹理数据的高速缓存，而且不需要耗时的阵列跨边界检查，因为纹理加载会自动将跨边界数据设置为0，这不会对矩阵乘法产生任何影响。

理解上面的预处理可以使理解下面的伪代码变得更容易。纹理创建和转换应该在GPU内核执行之前完成，而不影响内核执行性能。

纹理存储器中的数据也分段加载到共享存储器中。但是，按照原来的方法，每个片段应该一个接一个地装入小块中。为了充分利用寄存器资源，maxas采用了完全不同的计算方法。如果线程块计算出矩阵A首先被一个接一个地分成每64行的条带，则所需的输入数据都在条带上。当然，这段数据仍然非常大，需要在以后的阶段加载。maxas中的每个负载和消耗都是分阶段完成的。矩阵B的方法类似，只是条被分成列，转置加载方法与矩阵A完全相同。内存布局如下图所示:

图3。一个纹理在每个循环中被扭曲加载到共享内存中，可以被视为Bj或转置的Ai，因此矩阵实际上返回到传统的列优先级存储。这幅画被调换后实际上更直观。

图中的每个网格代表一个线程负责加载的数据单元，绿色网格是线程0将连续加载的四个单元，黄色网格是其余31个线程首次加载的数据部分。整个经线一次加载两条线，这要进行四次。GPU上的执行单元由32个线程组成，因此64个线程被分成两个线程来执行。一个经线(线0-31)加载A，另一个经线(线363)加载B。图中的一个难题是图中的矩阵形状不是，因为每个线将使用矢量指令一次加载4个浮点数，即每个网格本身是4个浮点数。在后面的代码中，您将看到当在纹理内存中使用矢量指令时，相对于元素的实际数量，偏移量被除以4。

这种装载方法显然不是唯一的方法。我的理解是，因为A和B的加载方法完全相同，但是使用的纹理不同，所以与同时加载A和B的线程相比，用于与计算无关的指令的代码量可以减少。

加载的数据暂时存储在寄存器中，等待存储在共享存储器中。共享内存中的数据排列如下:

图4。共享内存中矩阵a中加载的数据的排列。

由于共享存储器的偏移单位是1字节，它最终返回到可视表示，因此可以看出上述两个图的数据存储方法完全相同，所有列都先存储。唯一不同的是，在共享内存中，B的数据将直接跟随A。

图5。输入矩阵A和b的示意图。注意图中lda、ldb、ldc、bx、by、k的含义。这些变量将在下面的代码中使用。

下面的伪代码是对整个过程的描述。英文注释取自maxas文档，附加注释为中文。

tid = threadId.xbx = blockId.x。//可以看作是C_ij的iby = blockId.y//可以看作是C_ij中的j

blk = tid >。= 32？by:bx；ldx = tid >。= 32？ldb/4:LDA/4；// lda和ldb是A的行数或B的列数，ldx是它在向量加载中的表示(一次4个浮点数除以4)，它可以被视为每行的span tex = tid >。= 32？TeXb:TeXa；// 64根线被分成2根经纱，第一个加载织构A和第二个加载织构B

tid2 = (tid >。>。4)和1；//这里，32个线程被分成两组，tid2=0是加载到图3中第一行的第一组，tid2=1是加载到第二行tid15 = tid & 15//此线程在每行的列中的位置

//这些磁道变量表示线程要加载的数据偏移量，以tex为单位，乘以4，即偏移量track 0 = BLK * 64/4+TID 15+(LDX * TID 2)，分别以$A_i$或$B_j_T$为单位；//此线程加载的数据的起始位置解释如下:track 2 = track 0+ldx * 2；//对于此线程加载的数据的第二部分，第一部分有两行之后的track4 = track0+ldx*4。track 6 = track 0+ldx * 6；

end = track 0+(k-8)* ldx；//加载结束标志，其中k是a的列数和b的行数，即两个乘法矩阵的公共维数。对于NxN的矩阵，k = n。因为每次加载8行，所以减去8作为最后加载的标志。

WRITES = tid 15 * 4 * 4+tid 2 * 64 * 4；//此线程在共享内存中加载的数据偏移量是track0第二项的16倍，因为向量加载的偏移量1表示(4个数字x每个浮点数4个字节)writeS += tid > = 32？2048:0；//如果矩阵b的数据被加载，它被放置在矩阵a的数据之后，并且矩阵a占据(64列x8行×每个元素4字节)=2048字节

同时(track0 & ltend){ tex.1d.v4.f32.s32 loadX0，

图6。每个线程中使用的数据的寄存器号。每个寄存器所在的库用不同的颜色标记。如果c元素的颜色与其对应的a或b元素的颜色相同，则会发生存储体冲突。这个元素在图中用一个黑框标出。

显然，这是通过调整寄存器数量可以获得的最佳结果。无论如何调整c矩阵的数量，在图中用黑盒标记的存储体冲突都是不可避免的，因为它的来源是a和b使用相同的存储体，并且a和b中的操作数需要占据所有四个存储体(每个存储体两个),并且需要与另一个矩阵中的所有其他操作数配对。a的每个寄存器不可避免地会与B中的两个寄存器产生存储体冲突。事实上，如果C使用最简单的寄存器编号方法，例如在第一行中占据0~7，那么每个寄存器都将与相应的B操作数产生存储体冲突，这是一种非常糟糕的编号方法。

为了进一步消除寄存器分配无法消除的存储体冲突，需要Maxwell提供的操作数重用功能。当发出一条指令时，它的一些操作数可以被设置为重用，并且硬件将把操作数发送到重用缓存。如果后一条指令在相同的位置使用相同的操作数，则该指令可以直接从高速缓存中提取操作数，而无需通过存储体，从而避免存储体冲突。

FFMA·R2，R4 .复用，R5，R2；#这里指定的R4将被重用并放入缓存FFMA R0，R4 .重用，R5，R0；# R0和R4最初产生存储体冲突，但是因为先前的指令缓存了第二操作数R4，所以只需要从存储体中提取R0，从而避免了存储体冲突FFMA R0、R5、R4、R0；# R0和R4之间的存储体冲突仍然会发生，因为缓存的R4在操作数2中，而R4在此指令中属于操作数3。

如果图6中c矩阵的64个寄存器在一个线程中被简单地遍历一行或一列，并且a的寄存器被设置为重用，则a和b的16个寄存器组冲突中的14个可以被解决，并且寄存器R3和R35不能被解决，因为它们是使用a操作数的行的第一个指令，并且之前没有指令将它们发送到重用高速缓存中。当遍历每一行时，只要偶数行是从右向左(0行是从26到3)，奇数行是从左向右(1行是从7到30)，在知道原因之后，两个存储体之间的冲突也可以容易地解决。然而，马克斯仍然对此不满意。它在前面的遍历中添加了一个漩涡，并提出了一个更奇怪的遍历方法:

1、0、2、3、5、4、6、7、33、32、34、35、37、36、38、39、45、44、46、47、41、40、42、43、13、12、14、15、9、8、10、11、17、16、18、19、21、20、22、23、49、48、50、51、53、55

根据文档中的陈述，每个操作数8字节的重用缓存可以缓存两个4字节的寄存器，其中只有一个用于逐行遍历缓存寄存器a的数据，因此缓存利用率较低。我推测，考虑到存在B操作数和重用缓存，但它们没有被利用，逐行遍历的重用缓存利用率为4/8/2=25%。往返旅行利用率的估计并不那么直观。该文档直接给出了39%的利用率，而对于涡旋遍历的利用率可以达到49%。从maxas最终生成的汇编代码来看，有些指令确实对A和B都使用了重用缓存，并同时为每个操作数缓存了两个寄存器:

- :-:-:-:1 FFMA R37，R71 .再利用，R72 .再利用，R37；- :-:-:-:1 FFMA R36，R71 .再利用，R73 .再利用，R36；#前2个指令缓存R72和R73用于第3个操作数，它们由接下来的2个指令使用-:-:-:1 FFMA R38，R69 .重用，R73，R38；- :-:-:-:1 FFMA R39，R69 .再利用，R72，R39；#第二个操作数的前四个指令缓存R69和R71，它们被接下来的四个指令使用-:-:-:1 FFMA R45，R71 .重用，R74 .重用，R45；- :-:-:-:1 FFMA R44，R71，R75 .再利用，R44；- :-:-:-:1 FFMA R46，R69 .再利用，R75 .再利用，R46；- :-:-:-:1 FFMA R47，R69，R74 .再利用，R47；

然而，试图提高重用高速缓存的重用率的目的不是提高重用率，即使来回遍历可以完全解决存储体冲突，并且因为从共享存储器到寄存器的加载指令是在FFMA指令之间插入的。这样做的目的是加载指令的延迟可以由不依赖于其数据的计算指令来填充。遍历C矩阵寄存器的前八条指令和插入它们之间的加载指令如下:

01:-:-:-:0 FFMA R1，R66。重用，R72。重用，R1；- :-:-:-:1 LDS。U.128 R80，

图7。左图显示了寄存器写入共享内存的线程布局，而右图显示了从同一共享内存读取的线程布局。该图中的每一列都是图2中矩阵C的一列，并且相邻的2列被矩阵C中的4列分开..

该方法的实现代码如下。尽管这种方法需要在寄存器和共享内存之间重复传输数据，但共享内存的延迟可以通过在两个扭曲之间切换任务来屏蔽，毕竟，这比多次访问主内存要快得多。值得注意的是，尽管这种方法显然是一步一步完成的，但代码中没有同步指令，因为所有数据交换都是在同一经线的32个线程中直接完成的，经线之间的数据保持独立。GPU的执行模型能够严格保证同一扭曲中的同一条指令能够同时完成。省略同步指令的能力也是图2的并行方法的优点。

tid31 = tid & amp31；tid32 = tid & amp32；//只有两个值是可能的，0是第一个扭曲，32是第二个扭曲。

//如果高位存在于lastloop的xor中，则将其移除。//还删除了2048添加的本体读取bs。//以前，a和b在共享内存中被分配两倍的所需容量(4096字节)。一个块用于计算加载的数据，另一个块用于加载下一个段A和B。每个块的前2048个字节存储A，第二个2048个字节存储B//这种与操作相当于占用第一个块存储A的内存来存储CreadAs & = 0x7ff64x64矩阵readBs & = 0x7ff中该线程左上角数据的行坐标；64x64矩阵中该线程左上角数据的列坐标

//使用与以前几乎相同的共享映射写入共享，但将读操作折叠到一级。写操作=(读操作数/ 4) * 64 +读操作数；//此线程左上角的数据相对于64x64矩阵左上角的一维偏移量是根据行坐标和列坐标计算的。64/4是行优先存储矩阵的矢量传输跨度。对于线程0，这是图7左侧绿色网格中的顶部一个。

//用一个适合于合并的全局写的映射读取数据寄存器= ((tid32

LD C4 = LDC * 4；// ldc是矩阵C在行首存储中的列方向上的跨度，因子4表示其单位是字节而不是浮点数。

CX = bx * 64+tid 31；// cx可视为对应于要写入整个矩阵C的主存储器中的整个列的行数cy = by*64+(tid32 > );>。1)；// cy可以看作是对应于要写入主存储器的整个列的整个矩阵C的列数，显然是相同的扭曲列数。

// Cy00、Cy04、Cy08、Cy12是在整个矩阵C//中图7的右图中的顶部四个绿色网格点的偏移，尽管它们在共享存储器中被1列分开，但是在矩阵C中它们之间的间隔是4列，并且它们之间的偏移是LD C4 * 4C00 =(cy * LDC+CX)* 4+C；Cy04 = Cy00 + ldc4 * 4。Cy08 = Cy00 + ldc4 * 8。Cy12 = Cy00 + ldc4 * 12。

Foreach从in to . v 4 . f32 cs0和cs4 {//step 1复制8个寄存器的垂直线。从不连续的寄存器到共享内存//在存储到全局之前通过扭曲洗牌馈送8个寄存器//这里缺少一个步骤。每个线程首先从其两个4x4矩阵上下取出同一列的四个元素。此时，为了避免银行冲突//它们必须位于不连续的寄存器中。该步骤将它们复制到8个连续的附加寄存器cs0-cs7，上面的矩阵使用cs0-3，下面的使用cs7st . shared . v4 . f32[写cs+4 * 00]，cs0；//用矢量指令将连续寄存器cs0、cs1、cs2、cs3转移到与共享存储器中的4个数字对应的位置ST.Shared.V4.F32 [WRITECS+4 * 32]和CS4//与前一行的操作相同，因为上部和下部4x4矩阵由32个数字分隔，所以需要向写入位置添加4*32字节的偏移量

//步骤2。从共享内存中读取寄存器并重用cs寄存器，但此时不需要它的连续特性。//CS0、CS2、CS4和CS6位于同一行。行优先级存储相差64个元素，4*64字节。共享。F32 CS0，[READ CS+4 *(0 * 64+00)]；//CS1、CS3、CS5和CS7位于另一行。从图7的右图中可以看出，有32个元素LD。共享。F32CS1，[READCS+4 * (0 * 64+32)]与前一行不同。ld.shared.f32 cs2，[readCs+4 *(1 * 64+00)]；ld.shared.f32 cs3，[readCs+4 *(1 * 64+32)]；ld.shared.f32 cs4，[readCs+4 *(2 * 64+00)]；ld.shared.f32 cs5，[readCs+4 *(2 * 64+32)]；ld.shared.f32 cs6，[readCs+4 *(3 * 64+00)]；ld.shared.f32 cs7，[readCs+4 *(3 * 64+32)]；//步骤3。将cs寄存器的数量写入主存储器。对于整个扭曲，它相当于将一行32个连续的浮点数写入主内存。从逻辑上讲，它可以被视为步骤2的相反过程，除了列在共享存储器和主存储器中的位置不同。st.global.f32 [Cy00 + 4*00]，cs0st.global.f32 [Cy00 + 4*32]，cs1st.global.f32 [Cy04 + 4*00]，cs2st.global.f32 [Cy04 + 4*32]，cs3st.global.f32 [Cy08 + 4*00]，cs4st.global.f32 [Cy08 + 4*32]，cs5st.global.f32 [Cy12 + 4*00]，cs6st.global.f32 [Cy12 + 4*32]，cs7//输出该线程在下一个周期计算的4×4矩阵的下一列，对应于C矩阵的下一列。注意不要将其与图7中共享内存的下一列混淆。Cy00 += ldc4。Cy04 += ldc4。Cy08 += ldc4。Cy12 += ldc4。

//在处理完forthset之后，切换到strip32registers//add，在4个周期之后，4个4x4矩阵中的左两个矩阵已经被传送到主存储器，右两个矩阵接下来将被传送。//通过移动32列，可以重叠C矩阵中左右对4×4矩阵的对应位置。考虑到矩阵本身的宽度有4列(在前4个周期中通过+= ldc4 4次实现)//实际上需要额外移位的是4x4矩阵的左右对之间的间距34=28列，这是如果(第4次迭代){ Cy00+= LDC 4 * 28；Cy04 += ldc4 * 28。Cy08 += ldc4 * 28。Cy12 += ldc4 * 28。}}

在maxas文档中有另一个数字来表达这个过程，但是它可能不会在这里被引用，因为它没有被完全理解，它的意义也不重要，但是它很容易引起混淆。代码本身足以描述这个过程。

在向主存储器的传输完成后，由maxas生成的GEMM内核的任务就完成了

256线程实现

基于上面描述的每个线程块64个线程，它可以扩展4倍到256个线程，并且每个线程所做的工作将保持不变。这样，由每个线程块计算的补丁矩阵的两个维度分别被扩展2倍以达到。此时，输入矩阵的加载和结果的输出将会有相应的变化。然而，在实现了64个线程之后，理解这些变化是非常简单的，这里没有必要重复它们。对于较大的矩阵256线程实现，有一些性能优势。详细的测试结果见maxas文件。

结论

尽管本文尽可能多地注释了原始文档中的伪代码，但它仍然是一个相对高级的实现。具体来说，GPU机器代码还有一个重要的主题，即控制代码在本文中没有涉及。考虑到本文的目的仅仅是介绍一些GPU优化的思想和实现方法，涉及到控制代码的maxas文档的这一部分没有被解释。

一般来说，maxas使用的优化思想是相对清晰的。根据它的陈述，它在以前的文献中已经被提出。最困难的一点是，英伟达不愿意透露其硬件的实现细节，甚至需要作者的逆向工程猜测才能达到硬件性能的极限。也许作者已经建立了一个测试平台来快速验证由一些指令的细微差异引起的性能影响。无论如何，这是一项伟大的工作，值得任何对影响硬件性能极限感兴趣的工程师进行深入研究。

原始链接:https://www.jianshu.com/p/e01024892afb

本文由 在线网速测试 整理编辑，转载请注明出处，原文链接：https://www.wangsu123.cn/news/6347.html。