CN114281554B - 用于3d图像处理的3d-cnn加速方法及装置、电子设备 - Google Patents

Abstract

本发明公开了一种用于3D图像处理的3D‑CNN加速方法及装置、电子设备，包括：将3D图像进行预处理，得到3D特征图数据，将所述3D特征图数据存放到全局内存；将所述全局内存中的3D特征图数据进行隐式数据转化，得到第一中间数据；将所述第一中间数据写入到共享内存中；在所述共享内存中分配出读取缓冲区与写入缓冲区，通过第一中间数据在读取缓冲区与写入缓冲区交替传输，实现所有第一中间数据都写入寄存器；对所述寄存器中的第一中间数据进行分块；将分块后的第一中间数据进行计算，得到第二中间数据；对所述第二中间数据进行所述隐式数据转化的逆过程，重新转化成新的3D特征图数据；将所述新的3D特征图数据存放到所述全局内存中。

Description

用于3D图像处理的3D-CNN加速方法及装置、电子设备

技术领域

本申请涉及数据处理技术领域，尤其涉及一种用于3D图像处理的3D-CNN加速方法及装置、电子设备。

背景技术

卷积神经网络（CNN）近年被广泛应用于计算机视觉中，包括分类、检测、分割等任务。这些任务一般都是针对图像进行的，使用的是二维卷积（即卷积核的维度为二维）。而对于基于视频分析的问题，2D convolution不能很好得捕获时序上的信息，因此3Dconvolution就被提出来了。

随着CNN识别精度的提升，CNN模型越来越大，网络结构越来越复杂，随之而来的是对计算能力的要求不断提升。由于传统的CPU处理器已无法应对CNN网络强大的并行计算需求，各种类型的加速器如GPU、华为昇腾AI加速卡，FPGA等。在这些加速平台中，GPU由于具备完善的开发工具链以及强大的计算能力，渐渐得到了研究人员的青睐。GPU提供商如Nvidia陆续推出和完善了GPU编程框架cuda以及众多官方库如cuDNN、cuBLAS，使得GPU编程难度得到有效地降低，GPU加速器开发周期大大缩短，更让GPU成为加速CNN的最佳选择之一。Nvidia推出的最新GPU装备了Tensor Core加速部件，大大提升了GPU的计算性能，如何充分利用GPU 上的Tensor Core进行计算是能否设计出高效的3DCNN以处理三维医学图像和科学数据的关键。

目前基于GPU Tensor core的CNN加速器都是面向2D-CNN的，并没有公开文献对基于GPU Tensor core的3D-CNN加速进行研究。与2D CNN相比，3D-CNN增加了一个维度之后，具备更高的计算和存储复杂度。GPU的显存空间有限，限制了3D-CNN可使用的通道数量。通过roofline模型分析，在通道数量受限时，3D-CNN的计算访存比与新一代的GPU的计算访存比有较大的差距，无法充分利用新一代GPU中Tensor-core的计算资源。

发明内容

本申请实施例的目的是提供一种用于3D图像处理的3D-CNN加速方法及装置、电子设备，以解决相关技术中存在的3D-CNN在GPU Tensor Cores上进行图像处理可能出现的访存瓶颈的技术问题。该方法加速深度神经网络卷积层的计算，提高计算资源的利用率。

根据本申请实施例的第一方面，提供一种用于3D图像处理的3D-CNN加速方法，包括：

将3D图像进行预处理，得到3D特征图数据，将所述3D特征图数据存放到全局内存；

将所述全局内存中的3D特征图数据进行隐式数据转化，所述隐式数据转化包括：将所述全局内存中的3D特征图数据进行指针起始位置和偏移量的预先计算，得到每个线程读取3D特征图数据的地址，根据所述地址，对所述3D特征图数据进行布局调整和去重处理，得到第一中间数据；

将所述第一中间数据写入到共享内存中；

在所述共享内存中分配出读取缓冲区与写入缓冲区；

向所述写入缓冲区写入共享内存中第n批第一中间数据，向所述读取缓冲区写入共享内存中第n+1批第一中间数据，同时，将所述写入缓冲区中的第n批第一中间数据读取到寄存器中，在第n批第一中间数据完成写入寄存器之后，交换所述读取缓冲区与写入缓冲区的内容，重复本步骤，直到所有第一中间数据都写入寄存器；

对所述寄存器中的第一中间数据进行分块；

将分块后的第一中间数据进行计算，得到第二中间数据；

对所述第二中间数据进行所述隐式数据转化的逆过程，重新转化成新的3D特征图数据；

将所述新的3D特征图数据存放到所述全局内存中。

进一步地，将3D图像进行预处理，得到3D特征图数据，将所述3D特征图数据存放到全局内存，包括：

将所述3D图像进行文件格式转换；

将转换后的3D图像依次进行以下处理得到3D特征图数据：缩放成统一的大小、通过直方图均衡增强图像对比对、去噪处理；

检查处理后的3D特征图数据的通道数是否是GPU一个线程束中线程数量的倍数，如果不是其倍数则填充空数据将通道数补足成线程束中线程数量的倍数，补足通道数后，将图像的通道维度进行分组，每个分组中通道的数量等于线程束中线程数量，一个分组中通道连续的存放在全局内存中；

分组存放之后，将每个分组的3D图像数据，按照宽度、高度、深度这几个不同维度顺序地址变化由快到慢的存放在全局内存中。

进一步地，将所述全局内存中的3D特征图数据进行指针起始位置和偏移量的预先计算，得到每个线程读取3D特征图数据的地址，包括：

将3D-CNN的卷积核的长、宽、高，通道四个维度相乘得到卷积核每次要处理元素的数量；

将每次卷积核处理元素的数量与GPU每个线程束中线程的数量相乘，得到每个线程束每次卷积的元素个数；

将每个线程的线程ID整除每个线程束每次卷积的元素个数，得到每个线程所属的卷积核处理3D特征图数据的起始位置；

将每个线程的线程ID对每个线程束每次卷积的元素个数进行求余操作，得到每个线程在所属的卷积核处理3D特征图数据的块内偏移量；

根据所述起始位置和块内偏移量，得到线每个线程读取3D特征图数据的地址。

进一步地，根据所述地址，对所述3D特征图数据进行布局调整和去重处理，得到第一中间数据，包括：

根据所述的地址，得到对应的内存；

根据所述内存对应的地址，从GPU的一级缓存和二级缓存中查看所述3D特征图数据是否在缓存中，3D特征图数据中的每个数据第一次加载都需要从全局内存中读取，第一次读取完成之后就会被缓存到GPU的一级缓存和二级缓存中，下一次如果3D特征图中还有相同的元素就会从一级缓存和二级缓存中读取，以实现去重处理；

根据所述内存对应的地址，每个线程每次从全局内存中读取GPU 缓存行 大小的3D特征图数据；

将所述3D特征图数据进行重新排列；

将重新排列后的3D特征图数据写入共享内存，得到第一中间数据；

其中，将所述3D特征图数据进行重新排列，包括：

将每个线程读取的一个缓存行大小的3D特征图数据按照一个单精度浮点数类型为单位交替相邻存放在共享内存中作为一组，一个线程束中的所有线程写完第一个单精度浮点数之后，依次存放一个线程束中的所有线程读取的下个单精度浮点数大小的元素；其中每组3D特征图数据的元素个数与共享内存中的存储体的数量相同。

进一步地，在共享内存中分配出读取缓冲区与写入缓冲区，包括：

根据张量处理单元中寄存器最大可容纳的数据空间的大小，计算共享内存中开辟缓冲区的大小，其中所述缓冲区的大小等于张量处理单元中寄存器空间大小乘以一个经验值系数；

在共享内存中分配出两倍缓存区大小的共享内存，并将这块共享内存切分为均等大小的两块缓存区，一块为读缓冲区，一块为写缓存区。

进一步地，向所述写入缓冲区写入共享内存中第n批第一中间数据，向所述读取缓冲区写入共享内存中第n+1批第一中间数据，同时，将所述写入缓冲区中的第n批第一中间数据读取到寄存器中，在第n批第一中间数据完成写入寄存器之后，交换所述读取缓冲区与写入缓冲区的内容，包括：

提前将第n批第一中间数据存到写入缓冲区中，同时将这批第一中间存到寄存器分块等待张量处理单元进行计算；

在等待张量处理单元对第n批第一中间数据进行计算的同时，将n+1批第一中间数据存到读取缓冲区中；

当对第n批第一中间数据完成计算返回结果时，将读取缓冲区中的n+1批第一中间数据与写入缓冲区中的n批第一中间数据交换，然后将写入缓冲区的n+1批第一中间数据存到寄存器等待张量处理单元进行计算；

在等待计算结果返回过程中，读取缓冲区继续读取第n+2批第一中间数据。

进一步地，对所述寄存器中的第一中间数据进行分块，包括：

将GPU一个线程束中的所有线程被分为若干组，每组中的线程负责读取写入缓冲中的第一中间数据到寄存器分块fragment的矩阵中的一行；

每一组中的一个线程读取每一行中相邻的若干个8位的有符号整型的数据，恰好填满一个单精度浮点数类型的寄存器；

一个线程束所有线程分组将一个寄存器块fragment中的所有矩阵用第一中间数据填满，填充完成后，线程之间会进行数据交换，将矩阵的数据重新分布到各个线程。

根据本申请实施例的第二方面，提供一种用于3D图像处理的3D-CNN加速装置，包括：

预处理模块，用于将3D图像进行预处理，得到3D特征图数据，将所述3D特征图数据存放到全局内存；

隐式数据转化模块，用于将所述全局内存中的3D特征图数据进行隐式数据转化，所述隐式数据转化包括：将所述全局内存中的3D特征图数据进行指针起始位置和偏移量的预先计算，得到每个线程读取3D特征图数据的地址，根据所述地址，对所述3D特征图数据进行布局调整和去重处理，得到第一中间数据；

第一写入模块，用于将所述第一中间数据写入到共享内存中；

分配模块，用于在所述共享内存中分配出读取缓冲区与写入缓冲区；

第二写入模块，用于向所述写入缓冲区写入共享内存中第n批第一中间数据，向所述读取缓冲区写入共享内存中第n+1批第一中间数据，同时，将所述写入缓冲区中的第n批第一中间数据读取到寄存器中，在第n批第一中间数据完成写入寄存器之后，交换所述读取缓冲区与写入缓冲区的内容，重复第二写入模块的执行过程，直到所有第一中间数据都写入寄存器；

分块模块，用于对所述寄存器中的第一中间数据进行分块；

计算模块，用于将分块后的第一中间数据进行计算，得到第二中间数据；

逆隐式数据转化模块，用于对所述第二中间数据进行所述隐式数据转化的逆过程，重新转化成新的3D特征图数据；

存放模块，用于将所述新的3D特征图数据存放到所述全局内存中。

根据本申请实施例的第三方面，提供一种电子设备，包括：

一个或多个处理器；

存储器，用于存储一个或多个程序；

当所述一个或多个程序被所述一个或多个处理器执行，使得所述一个或多个处理器实现如第一方面所述的方法。

根据本申请实施例的第四方面，提供一种计算机可读存储介质，其上存储有计算机指令，该指令被处理器执行时实现如第一方面所述方法的步骤。

本申请的实施例提供的技术方案可以包括以下有益效果：

由上述实施例可知，本申请通过去重处理减少了访问全局内存的次数，通过划分读写缓存区进行双缓冲流水线读写部分掩盖访存延迟， 通过寄存器分块计算减少访问共享内存的次数，最大可能地降低了GPU计算过程中受限于访存开销的问题。经测试3D图像推理效果取得了比直接使用英伟达官方库cuDNN进行推理加速1.4~1.7倍的效果。

应当理解的是，以上的一般描述和后文的细节描述仅是示例性和解释性的，并不能限制本申请。

附图说明

此处的附图被并入说明书中并构成本说明书的一部分，示出了符合本申请的实施例，并与说明书一起用于解释本申请的原理。

图1是根据一示例性实施例示出的一种用于3D图像处理的3D-CNN加速方法的流程图。

图2是根据一示例性实施例示出的数据转化实例图。

图3是根据一示例性实施例示出的数据去重图。

图4是根据一示例性实施例示出的数据分块图。

图5是根据一示例性实施例示出的流水线计算和数据预取图。

图6是根据一示例性实施例示出的共享内存数据重新排列和调整图。

图7是根据一示例性实施例示出的一种用于3D图像处理的3D-CNN加速装置的结构示意图。

图8是根据一示例性实施例示出的电子设备的结构示意图。

具体实施方式

这里将详细地对示例性实施例进行说明，其示例表示在附图中。下面的描述涉及附图时，除非另有表示，不同附图中的相同数字表示相同或相似的要素。以下示例性实施例中所描述的实施方式并不代表与本申请相一致的所有实施方式。相反，它们仅是与如所附权利要求书中所详述的、本申请的一些方面相一致的装置和方法的例子。

在本申请使用的术语是仅仅出于描述特定实施例的目的，而非旨在限制本申请。在本申请和所附权利要求书中所使用的单数形式的“一种”、“所述”和“该”也旨在包括多数形式，除非上下文清楚地表示其他含义。还应当理解，本文中使用的术语“和/或”是指并包含一个或多个相关联的列出项目的任何或所有可能组合。

应当理解，尽管在本申请可能采用术语第一、第二、第三等来描述各种信息，但这些信息不应限于这些术语。这些术语仅用来将同一类型的信息彼此区分开。例如，在不脱离本申请范围的情况下，第一信息也可以被称为第二信息，类似地，第二信息也可以被称为第一信息。取决于语境，如在此所使用的词语“如果”可以被解释成为“在……时”或“当……时”或“响应于确定”。

对视频使用3D卷积似乎优势并不大，而在3D图像处理领域的应用前景更大一些。3D图像，比如医学数据，通常都是3D的，比如CT扫描的数据，虽然我们看的片子是2D的，但其实那只是一个切片，真正的扫描数据是3D的。而如果要分割出一些病变组织，比如肿瘤，也必须是3D的。另外，在科学计算领域，生产许多空间数据也往往是3D的。

以下通过3D医学图像为例来进行优选实施例的描述。

图1是根据一示例性实施例示出的一种用于3D图像处理的3D-CNN加速方法的流程图，如图1所示，可以包括以下步骤：

步骤S11，将3D图像进行预处理，得到3D特征图数据，将所述3D特征图数据存放到全局内存；

步骤S12，将所述全局内存中的3D特征图数据进行隐式数据转化，所述隐式数据转化包括：将所述全局内存中的3D特征图数据进行指针起始位置和偏移量的预先计算，得到每个线程读取3D特征图数据的地址，根据所述地址，对所述3D特征图数据进行布局调整和去重处理，得到第一中间数据；

步骤S13，将所述第一中间数据写入到共享内存中；

步骤S14，在所述共享内存中分配出读取缓冲区与写入缓冲区；

步骤S15，向所述写入缓冲区写入共享内存中第n批第一中间数据，向所述读取缓冲区写入共享内存中第n+1批第一中间数据，同时，将所述写入缓冲区中的第n批第一中间数据读取到寄存器中，在第n批第一中间数据完成写入寄存器之后，交换所述读取缓冲区与写入缓冲区的内容，重复本步骤，直到所有第一中间数据都写入寄存器；

步骤S16，对所述寄存器中的第一中间数据进行分块；

步骤S17，将分块后的第一中间数据进行计算，得到第二中间数据；

步骤S18，对所述第二中间数据进行所述隐式数据转化的逆过程，重新转化成新的3D特征图数据；

步骤S19，将所述新的3D特征图数据存放到所述全局内存中。

由上述实施例可知，本申请通过去重处理减少了访问全局内存的次数，通过划分读写缓存区进行双缓冲流水线读写部分掩盖访存延迟， 通过寄存器分块计算减少访问共享内存的次数， 最大可能地降低了GPU计算过程中受限于访存开销的问题。经测试3D图像推理效果取得了比直接使用英伟达官方库cuDNN进行推理加速1.4~1.7倍的效果。

在步骤S11的具体实施中，将3D图像进行预处理，得到3D特征图数据，将所述3D特征图数据存放到全局内存；

具体地，将3D医学图像数据进行预处理转换成GPU能够处理的格式，即将不同3D图像的格式统一转换成dicom格式，再将转换后的3D图像依次进行以下处理得到3D特征图数据：缩放成统一的大小、通过直方图均衡增强图像对比对、去噪处理。这里可以通过插值、采样等方式将图像缩放成统一的大小；用中值滤波的方法对3D图像去噪。

检查处理后的3D特征图数据的通道数是否是GPU一个线程束中线程数量的倍数，如果不是其倍数则填充空数据将通道数补足成线程束中线程数量的倍数，补足通道数后，将图像的通道维度进行分组，每个分组中通道的数量等于线程束中线程数量，一个分组中通道连续的存放在全局内存中；

分组存放之后，将每个分组的3D图像数据，按照宽度、高度、深度这几个不同维度顺序地址变化由快到慢的存放在全局内存中。

本实例3D特征图数据的通道维度按照32个通道进行分组，每 32 个通道连续的存放在全局内存中。由于采用32 个通道对齐的存储格式，因此卷积层要求输入和输出特征图的通道数都是32的倍数。3D特征图数据的其余维度按照宽度、高度、深度的顺序地址变化由快到慢的存放在全局内存中。

将3D特征图数据通过cudaMemcpyAsync / cudaMemcpy2DAsync /cudaMemcpy3DAsync接口异步传输到GPU显存。数据传输是走PCIe总线的，计算和PCIe总线里的数据流通完全独立，那么某些情况下，我们可以让计算和传输异步进行，而不是等数据传输完再做计算。

举个例子：如果一次传入两张图像，做处理运算。常规操作是使用cudaMemcpy或者cudaMemcpy2D把两张图像都传输到显存，再启动kernel运算。传输和运算是串行的，运算必须等待传输完成。而cudaMemcpyAsync / cudaMemcpy2DAsync / cudaMemcpy3DAsync 可以让传输和运算之间异步并行。上面的例子，如果用cudaMemcpyAsync或cudaMemcpy2DAsync，可以先传输第一张影像到显存，然后启动第一张影像的运算kernel，同时启动第二张影像的传输，此时第一张影像的运算和第二张影像的传输就是异步进行的，互相独立，便可隐藏掉第二张影像的传输耗时。

在步骤S12的具体实施中，将所述全局内存中的3D特征图数据进行隐式数据转化，如图2所示，所述隐式数据转化包括：将所述全局内存中的3D特征图数据进行指针起始位置和偏移量的预先计算，得到每个线程读取3D特征图数据的地址，根据所述地址，对所述3D特征图数据进行布局调整和去重处理，得到第一中间数据；

具体地，将所述全局内存中的3D特征图数据进行指针起始位置和偏移量的预先计算，得到每个线程读取3D特征图数据的地址，可以包括：

将3D-CNN的卷积核的长、宽、高，通道四个维度相乘得到卷积核每次要处理元素的数量；

将每次卷积核处理元素的数量与GPU每个线程束中线程的数量相乘，得到每个线程束每次卷积的元素个数；

将每个线程的线程ID整除每个线程束每次卷积的元素个数，得到每个线程所属的卷积核处理3D特征图数据的起始位置；

将每个线程的线程ID对每个线程束每次卷积的元素个数进行求余操作，得到每个线程在所属的卷积核处理3D特征图数据的块内偏移量；

根据所述起始位置和块内偏移量，得到线每个线程读取3D特征图数据的地址。

具体地，每个线程对应的3D特征图数据指针起始位置等于(线程ID) /（卷积核宽度×卷积核高度×卷积核深度×32），偏移量等于(线程ID %(卷积核宽度×卷积核高度×卷积核深度×32) )。

根据所述地址，对所述3D特征图数据进行布局调整和去重处理，得到第一中间数据，可以包括：

根据所述的地址，得到对应的内存；

如图3所示，线程在从全局内存读取3D特征图数据之前，根据所述内存对应的地址，先从GPU的一级缓存和二级缓存中查看所述3D特征图数据是否在缓存中，3D特征图数据中的每个数据第一次加载都需要从全局内存中读取，第一次读取完成之后就会被缓存到GPU的一级缓存和二级缓存中，下一次如果3D特征图中还有相同的元素就会从一级缓存和二级缓存中读取，避免从全局内存中读取重复的数据，从而实现去重处理；

根据所述内存对应的地址，每个线程每次从全局内存中读取GPU 缓存行（cacheline）大小的3D特征图数据；具体地，每个线程从全局内存中读取粒度为128位，也即16个8位的有符号整型（INT8）类型的3D特征图数据的到共享内存，共享内存中的数据的布局以 16个数据为一组，组织成一个行优先的矩阵。

如图6所示，将所述3D特征图数据进行重新排列；其中，将所述3D特征图数据进行重新排列，包括：将每个线程读取的一个cacheline 大小的3D特征图数据按照一个单精度浮点数（float）类型为单位交替相邻存放在共享内存中作为一组，一个线程束中的所有线程写完第一个float之后，依次存放一个线程束中的所有线程读取的下个float大小的元素；其中每组3D特征图数据的元素个数与共享内存中的存储体（bank）的数量相同；本实例中将全局内存数据中每4行的32个相邻元素的数据作为一组，连续存放在共享内存中，紧接着会存放这4行的下一组32个相邻元素的数据。每组数据占据了128个字节，恰好是共享内存中的32个不同bank。每组数据在排列是进行了交错，保证了ldmatrix时不会发生访问同一个bank导致冲突的问题。

将重新排列后的3D特征图数据写入共享内存，得到第一中间数据。

以上数据预处理、布局调整以及去重的好处是：通过数据的一系列预处理可以使得3D图像在用神经网络进行推理时可以有更高的精度，把通道数设置为和GPU线程束中线程的数量一致以及数据布局调整的是为了使得后续每个线程访问全局内存时可以按照cacheline的大小进行对齐访问，减少访问全局内存的次数。进行数据去重也可以减少计算过程中大量对全局内存进行的不必要访问。

对共享内存中的数据进行重新排列和调整的好处是，避免ldmatrix时发生访问同一个bank导致冲突的问题， 减少Tensor Core计算时访问共享内存的次数。

在步骤S14的具体实施中，在所述共享内存中分配出读取缓冲区与写入缓冲区；

具体地，根据张量处理单元 （TensorCore）中寄存器最大可容纳的数据空间的大小，计算共享内存中开辟缓冲区的大小，其中所述缓冲区的大小等于TensorCore 中寄存器空间大小乘以一个经验值系数，本实例的经验值系数取1.2，当然不局限于此；

在共享内存中分配出两倍缓存区大小的共享内存，并将这块共享内存切分为均等大小的两块缓存区，一块为读缓冲区，一块为写缓存区。

在步骤S15的具体实施中，如图5所示，向所述写入缓冲区写入共享内存中第n批第一中间数据，向所述读取缓冲区写入共享内存中第n+1批第一中间数据，同时，将所述写入缓冲区中的第n批第一中间数据读取到寄存器中，在第n批第一中间数据完成写入寄存器之后，交换所述读取缓冲区与写入缓冲区的内容，重复以上过程，直到所有第一中间数据都写入寄存器；

具体地，提前将第n批第一中间数据存到写入缓冲区中，同时将这批第一中间存到寄存器分块等待TensorCore进行计算。

由于将第一中间数据的存到寄存器之后进行计算所需的时间周期比较多。所以在的等待TensorCore对第n批第一中间数据进行计算的同时，将n+1批第一中间数据存到读取缓冲区中。

当对第n批第一中间数据完成计算返回结果时，将读取缓冲区中的n+1批第一中间数据与写入缓冲区中的n批第一中间数据交换，然后将写入缓冲区的n+1批第一中间数据存到寄存器等待TensorCore进行计算。 在等待计算结果返回过程中，读取缓冲区继续读取第n+2批第一中间数据。

重复以上过程，直到所有第一中间数据都写入寄存器。

通过这种方式可以使得访问全局内存的数据和TensorCore的计算同时进行， 掩盖部分访问全局内存的延迟。

在步骤S16的具体实施中，如图4所示对所述寄存器中的第一中间数据进行分块；

具体地，将GPU一个线程束中的所有线程被分为若干组，每组中的线程负责读取写入缓冲中的第一中间数据到寄存器分块fragment的矩阵中的一行；

每一组中的一个线程读取每一行中相邻的若干个INT8的数据，恰好填满一个float类型的寄存器；

一个线程束所有线程分组将一个寄存器块fragment中的所有矩阵用第一中间数据填满，填充完成后，线程之间会进行数据交换，将矩阵的数据重新分布到各个线程。

更具体地，通cuda提供的ldmatrix接口可以精细地控制线程从共享存储单元的写入缓冲区加载第一中间数据给到寄存器分块fragment。

GPU一个线程束中的 32 个线程被分为8组，每组4个线程，负责读取写入缓冲中的第一中间数据到寄存器分块fragment 8×16 的矩阵中的一行。每一组中的一个线程读取每一行中相邻的 4 个INT8的数据，恰好填满一个 32 位的寄存器。

一个线程束8个线程分组将一个寄存器块fragment中的所有8×16的矩阵用第一中间数据填满。填充完成后，线程之间会进行数据交换，将矩阵的数据重新分布到各个线程。

通过将数据从共享内存的写缓冲区加载到寄存器分块fragment中，可以使得TensorCore在计算过程中减少对共享内存的访问，降低对共享内存的访问开销。

在步骤S17的具体实施中，将分块后的第一中间数据进行计算，得到第二中间数据；

具体地，将寄存器中分块的第一中间数据通过cuda提供的接口调用TensorCore进行矩阵乘法计算，得到第二中间数据。

通过cuda提供的mma接口使用 TensorCore来进行高速的矩阵乘加运算。

每个mma指令可以控制一个线程束中的而所有线程同步地完成若干个的矩阵块乘加运算。

参与矩阵乘法运算的分别是寄存器块fragment中的固定大小和形状的矩阵 A 和矩阵B，这两个输入矩阵的数据分布在同一线程束的所有线程。3.计算完成的结果存放在寄存器块fragment中的一个新的矩阵块中，这个计算结果就是第二中间数据。

更具体地，通过cuda提供的mma接口使用 TensorCore来进行高速的矩阵乘加运算。每个mma指令可以控制32 个线程同步地完成 8×8×16 的矩阵乘加运算。

参与矩阵乘法运算的分别是寄存器块fragment中的一个 8×16 的矩阵 A 和一个 16×8 的矩阵B，这两个输入矩阵的数据分布在同一线程束的 32 个线程。计算完成的结果存放在寄存器块fragment中的一个 8×8的矩阵中，这个计算结果就是第二中间数据。

在步骤S18的具体实施中，对所述第二中间数据进行所述隐式数据转化的逆过程，重新转化成新的3D特征图数据；

具体地，每个线程束中的所有线程将fragment中的第二中间数据写回全局内存，每个线程每次写回一个cacheline大小的数据。在写回过程中，需要查询所述指针和偏移，通过所述3D特征图布局调整的逆过程将第二中间数据重新转换为3D特征图数据。每个fragment块由一个线程束中的线程分成多次写回，多个线程束写回多个不同的fragment块中的数据。

更具体地，每32个线程为一组负责写回一个 64×64×64 的新3D特征图数据块的内容到全局内存。每个64×64×64的数据块分成8次写回，每次写回 32×8 的矩阵。 每次32个线程恰好写了全局内存上连续的 256 字节的数据，每个线程写回8个字节，这样可以使用64位宽的数据类型进行全局内存的写操作，尽可能提高带宽的利用率。

与前述的用于3D图像处理的3D-CNN加速方法的实施例相对应，本申请还提供了用于3D图像处理的3D-CNN加速装置的实施例。

图7是根据一示例性实施例示出的一种用于3D图像处理的3D-CNN加速装置框图。参照图7，该装置包括：

预处理模块21，用于将3D图像进行预处理，得到3D特征图数据，将所述3D特征图数据存放到全局内存；

隐式数据转化模块22，用于将所述全局内存中的3D特征图数据进行隐式数据转化，所述隐式数据转化包括：将所述全局内存中的3D特征图数据进行指针起始位置和偏移量的预先计算，得到每个线程读取3D特征图数据的地址，根据所述地址，对所述3D特征图数据进行布局调整和去重处理，得到第一中间数据；

第一写入模块23，用于将所述第一中间数据写入到共享内存中；

分配模块24，用于在所述共享内存中分配出读取缓冲区与写入缓冲区；

第二写入模块25，用于向所述写入缓冲区写入共享内存中第n批第一中间数据，向所述读取缓冲区写入共享内存中第n+1批第一中间数据，同时，将所述写入缓冲区中的第n批第一中间数据读取到寄存器中，在第n批第一中间数据完成写入寄存器之后，交换所述读取缓冲区与写入缓冲区的内容，重复第二写入模块的执行过程，直到所有第一中间数据都写入寄存器；

分块模块26，用于对所述寄存器中的第一中间数据进行分块；

计算模块27，用于将分块后的第一中间数据进行计算，得到第二中间数据；

逆隐式数据转化模块28，用于对所述第二中间数据进行所述隐式数据转化的逆过程，重新转化成新的3D特征图数据；

存放模块29，用于将所述新的3D特征图数据存放到所述全局内存中。

关于上述实施例中的装置，其中各个模块执行操作的具体方式已经在有关该方法的实施例中进行了详细描述，此处将不做详细阐述说明。

对于装置实施例而言，由于其基本对应于方法实施例，所以相关之处参见方法实施例的部分说明即可。以上所描述的装置实施例仅仅是示意性的，其中所述作为分离部件说明的单元可以是或者也可以不是物理上分开的，作为单元显示的部件可以是或者也可以不是物理单元，即可以位于一个地方，或者也可以分布到多个网络单元上。可以根据实际的需要选择其中的部分或者全部模块来实现本申请方案的目的。本领域普通技术人员在不付出创造性劳动的情况下，即可以理解并实施。

相应的，本申请还提供一种电子设备，包括：一个或多个处理器；存储器，用于存储一个或多个程序；当所述一个或多个程序被所述一个或多个处理器执行，使得所述一个或多个处理器实现如上述的用于3D图像处理的3D-CNN加速方法。如图8所示，为本发明实施例提供的一种用于3D图像处理的3D-CNN加速装置所在任意具备数据处理能力的设备的一种硬件结构图，除了图8所示的处理器、内存之外，实施例中装置所在的任意具备数据处理能力的设备通常根据该任意具备数据处理能力的设备的实际功能，还可以包括其他硬件，对此不再赘述。

相应的，本申请还提供一种计算机可读存储介质，其上存储有计算机指令，其特征在于，该指令被处理器执行时实现如上述的用于3D图像处理的3D-CNN加速方法。所述计算机可读存储介质可以是前述任一实施例所述的任意具备数据处理能力的设备的内部存储单元，例如硬盘或内存。所述计算机可读存储介质也可以是风力发电机的外部存储设备，例如所述设备上配备的插接式硬盘、智能存储卡（Smart Media Card，SMC）、SD卡、闪存卡（Flash Card）等。进一步的，所述计算机可读存储介还可以既包括任意具备数据处理能力的设备的内部存储单元也包括外部存储设备。所述计算机可读存储介质用于存储所述计算机程序以及所述任意具备数据处理能力的设备所需的其他程序和数据，还可以用于暂时地存储已经输出或者将要输出的数据。

本领域技术人员在考虑说明书及实践这里公开的内容后，将容易想到本申请的其它实施方案。本申请旨在涵盖本申请的任何变型、用途或者适应性变化，这些变型、用途或者适应性变化遵循本申请的一般性原理并包括本申请未公开的本技术领域中的公知常识或惯用技术手段。说明书和实施例仅被视为示例性的，本申请的真正范围和精神由下面的权利要求指出。

应当理解的是，本申请并不局限于上面已经描述并在附图中示出的精确结构，并且可以在不脱离其范围进行各种修改和改变。本申请的范围仅由所附的权利要求来限制。

Claims (9)

1.一种用于3D图像处理的3D-CNN加速方法，其特征在于，包括：

将3D图像进行预处理，得到3D特征图数据，将所述3D特征图数据存放到全局内存；

将所述全局内存中的3D特征图数据进行隐式数据转化，所述隐式数据转化包括：将所述全局内存中的3D特征图数据进行指针起始位置和偏移量的预先计算，得到每个线程读取3D特征图数据的地址，根据所述地址，对所述3D特征图数据进行布局调整和去重处理，得到第一中间数据；

将所述第一中间数据写入到共享内存中；

在所述共享内存中分配出读取缓冲区与写入缓冲区；

向所述写入缓冲区写入共享内存中第n批第一中间数据，向所述读取缓冲区写入共享内存中第n+1批第一中间数据，同时，将所述写入缓冲区中的第n批第一中间数据读取到寄存器中，在第n批第一中间数据完成写入寄存器之后，交换所述读取缓冲区与写入缓冲区的内容，重复本步骤，直到所有第一中间数据都写入寄存器；

对所述寄存器中的第一中间数据进行分块；

将分块后的第一中间数据进行计算，得到第二中间数据；

对所述第二中间数据进行所述隐式数据转化的逆过程，重新转化成新的3D特征图数据；

将所述新的3D特征图数据存放到所述全局内存中；

其中，向所述写入缓冲区写入共享内存中第n批第一中间数据，向所述读取缓冲区写入共享内存中第n+1批第一中间数据，同时，将所述写入缓冲区中的第n批第一中间数据读取到寄存器中，在第n批第一中间数据完成写入寄存器之后，交换所述读取缓冲区与写入缓冲区的内容，包括：

提前将第n批第一中间数据存到写入缓冲区中，同时将这批第一中间存到寄存器分块等待张量处理单元进行计算；

在等待张量处理单元对第n批第一中间数据进行计算的同时，将n+1批第一中间数据存到读取缓冲区中；

当对第n批第一中间数据完成计算返回结果时，将读取缓冲区中的n+1批第一中间数据与写入缓冲区中的n批第一中间数据交换，然后将写入缓冲区的n+1批第一中间数据存到寄存器等待张量处理单元进行计算；

在等待计算结果返回过程中，读取缓冲区继续读取第n+2批第一中间数据。

2.根据权利要求1所述的方法，其特征在于，将3D图像进行预处理，得到3D特征图数据，将所述3D特征图数据存放到全局内存，包括：

将所述3D图像进行文件格式转换；

将转换后的3D图像依次进行以下处理得到3D特征图数据：缩放成统一的大小、通过直方图均衡增强图像对比对、去噪处理；

检查处理后的3D特征图数据的通道数是否是GPU一个线程束中线程数量的倍数，如果不是其倍数则填充空数据将通道数补足成线程束中线程数量的倍数，补足通道数后，将图像的通道维度进行分组，每个分组中通道的数量等于线程束中线程数量，一个分组中通道连续的存放在全局内存中；

分组存放之后，将每个分组的3D图像数据，按照宽度、高度、深度这几个不同维度顺序地址变化由快到慢的存放在全局内存中。

3.根据权利要求1所述的方法，其特征在于，将所述全局内存中的3D特征图数据进行指针起始位置和偏移量的预先计算，得到每个线程读取3D特征图数据的地址，包括：

将3D-CNN的卷积核的长、宽、高，通道四个维度相乘得到卷积核每次要处理元素的数量；

将每次卷积核处理元素的数量与GPU每个线程束中线程的数量相乘，得到每个线程束每次卷积的元素个数；

将每个线程的线程ID整除每个线程束每次卷积的元素个数，得到每个线程所属的卷积核处理3D特征图数据的起始位置；

将每个线程的线程ID对每个线程束每次卷积的元素个数进行求余操作，得到每个线程在所属的卷积核处理3D特征图数据的块内偏移量；

根据所述起始位置和块内偏移量，得到线每个线程读取3D特征图数据的地址。

4.根据权利要求1所述的方法，其特征在于，根据所述地址，对所述3D特征图数据进行布局调整和去重处理，得到第一中间数据，包括：

根据所述的地址，得到对应的内存；

根据所述内存对应的地址，从GPU的一级缓存和二级缓存中查看所述3D特征图数据是否在缓存中，3D特征图数据中的每个数据第一次加载都需要从全局内存中读取，第一次读取完成之后就会被缓存到GPU的一级缓存和二级缓存中，下一次如果3D特征图中还有相同的元素就会从一级缓存和二级缓存中读取，以实现去重处理；

根据所述内存对应的地址，每个线程每次从全局内存中读取GPU 缓存行 大小的3D特征图数据；

将所述3D特征图数据进行重新排列；

将重新排列后的3D特征图数据写入共享内存，得到第一中间数据；

其中，将所述3D特征图数据进行重新排列，包括：

将每个线程读取的一个缓存行大小的3D特征图数据按照一个单精度浮点数类型为单位交替相邻存放在共享内存中作为一组，一个线程束中的所有线程写完第一个单精度浮点数之后，依次存放一个线程束中的所有线程读取的下个单精度浮点数大小的元素；其中每组3D特征图数据的元素个数与共享内存中的存储体的数量相同。

5.根据权利要求1所述的方法，其特征在于，在共享内存中分配出读取缓冲区与写入缓冲区，包括：

根据张量处理单元中寄存器最大可容纳的数据空间的大小，计算共享内存中开辟缓冲区的大小，其中所述缓冲区的大小等于张量处理单元中寄存器空间大小乘以一个经验值系数；

在共享内存中分配出两倍缓存区大小的共享内存，并将这块共享内存切分为均等大小的两块缓存区，一块为读缓冲区，一块为写缓存区。

6.根据权利要求1所述的方法，其特征在于，对所述寄存器中的第一中间数据进行分块，包括：

将GPU一个线程束中的所有线程被分为若干组，每组中的线程负责读取写入缓冲中的第一中间数据到寄存器分块fragment的矩阵中的一行；

每一组中的一个线程读取每一行中相邻的若干个8位的有符号整型的数据，恰好填满一个单精度浮点数类型的寄存器；

一个线程束所有线程分组将一个寄存器块fragment中的所有矩阵用第一中间数据填满，填充完成后，线程之间会进行数据交换，将矩阵的数据重新分布到各个线程。

7.一种用于3D图像处理的3D-CNN加速装置，其特征在于，包括：

预处理模块，用于将3D图像进行预处理，得到3D特征图数据，将所述3D特征图数据存放到全局内存；

隐式数据转化模块，用于将所述全局内存中的3D特征图数据进行隐式数据转化，所述隐式数据转化包括：将所述全局内存中的3D特征图数据进行指针起始位置和偏移量的预先计算，得到每个线程读取3D特征图数据的地址，根据所述地址，对所述3D特征图数据进行布局调整和去重处理，得到第一中间数据；

第一写入模块，用于将所述第一中间数据写入到共享内存中；

分配模块，用于在所述共享内存中分配出读取缓冲区与写入缓冲区；

第二写入模块，用于向所述写入缓冲区写入共享内存中第n批第一中间数据，向所述读取缓冲区写入共享内存中第n+1批第一中间数据，同时，将所述写入缓冲区中的第n批第一中间数据读取到寄存器中，在第n批第一中间数据完成写入寄存器之后，交换所述读取缓冲区与写入缓冲区的内容，重复第二写入模块的执行过程，直到所有第一中间数据都写入寄存器；

分块模块，用于对所述寄存器中的第一中间数据进行分块；

计算模块，用于将分块后的第一中间数据进行计算，得到第二中间数据；

逆隐式数据转化模块，用于对所述第二中间数据进行所述隐式数据转化的逆过程，重新转化成新的3D特征图数据；

存放模块，用于将所述新的3D特征图数据存放到所述全局内存中；

其中，向所述写入缓冲区写入共享内存中第n批第一中间数据，向所述读取缓冲区写入共享内存中第n+1批第一中间数据，同时，将所述写入缓冲区中的第n批第一中间数据读取到寄存器中，在第n批第一中间数据完成写入寄存器之后，交换所述读取缓冲区与写入缓冲区的内容，包括：

提前将第n批第一中间数据存到写入缓冲区中，同时将这批第一中间存到寄存器分块等待张量处理单元进行计算；

在等待张量处理单元对第n批第一中间数据进行计算的同时，将n+1批第一中间数据存到读取缓冲区中；

当对第n批第一中间数据完成计算返回结果时，将读取缓冲区中的n+1批第一中间数据与写入缓冲区中的n批第一中间数据交换，然后将写入缓冲区的n+1批第一中间数据存到寄存器等待张量处理单元进行计算；

在等待计算结果返回过程中，读取缓冲区继续读取第n+2批第一中间数据。

8.一种电子设备，其特征在于，包括：

一个或多个处理器；

存储器，用于存储一个或多个程序；

当所述一个或多个程序被所述一个或多个处理器执行，使得所述一个或多个处理器实现如权利要求1-6任一项所述的方法。

9.一种计算机可读存储介质，其上存储有计算机指令，其特征在于，该指令被处理器执行时实现如权利要求1-6中任一项所述方法的步骤。

Images