CN112232496A - 一种基于Tensorcore处理int4数据类型的方法、***、设备及介质 - Google Patents
一种基于Tensorcore处理int4数据类型的方法、***、设备及介质 Download PDFInfo
- Publication number
- CN112232496A CN112232496A CN202010980721.9A CN202010980721A CN112232496A CN 112232496 A CN112232496 A CN 112232496A CN 202010980721 A CN202010980721 A CN 202010980721A CN 112232496 A CN112232496 A CN 112232496A
- Authority
- CN
- China
- Prior art keywords
- data
- dimension
- shared memory
- input
- calculation result
- Prior art date
- Legal status (The legal status is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the status listed.)
- Withdrawn
Links
- 101100317378 Mus musculus Wnt3 gene Proteins 0.000 title claims abstract description 71
- 238000012545 processing Methods 0.000 title claims abstract description 66
- 238000000034 method Methods 0.000 title claims abstract description 45
- 238000004364 calculation method Methods 0.000 claims abstract description 88
- 230000004044 response Effects 0.000 claims abstract description 9
- 238000004590 computer program Methods 0.000 claims description 6
- 230000006872 improvement Effects 0.000 abstract description 5
- 239000011159 matrix material Substances 0.000 description 49
- 238000010586 diagram Methods 0.000 description 18
- 238000013528 artificial neural network Methods 0.000 description 5
- 238000009825 accumulation Methods 0.000 description 4
- 230000008569 process Effects 0.000 description 4
- 230000000694 effects Effects 0.000 description 3
- 230000006870 function Effects 0.000 description 3
- 238000005457 optimization Methods 0.000 description 3
- 238000005192 partition Methods 0.000 description 3
- 238000004422 calculation algorithm Methods 0.000 description 2
- 238000012986 modification Methods 0.000 description 2
- 230000004048 modification Effects 0.000 description 2
- 230000003287 optical effect Effects 0.000 description 2
- HPTJABJPZMULFH-UHFFFAOYSA-N 12-[(Cyclohexylcarbamoyl)amino]dodecanoic acid Chemical compound OC(=O)CCCCCCCCCCCNC(=O)NC1CCCCC1 HPTJABJPZMULFH-UHFFFAOYSA-N 0.000 description 1
- 102100039134 Integrator complex subunit 4 Human genes 0.000 description 1
- 101710092887 Integrator complex subunit 4 Proteins 0.000 description 1
- 238000007792 addition Methods 0.000 description 1
- 230000009286 beneficial effect Effects 0.000 description 1
- 238000013135 deep learning Methods 0.000 description 1
- 238000000605 extraction Methods 0.000 description 1
- 239000012634 fragment Substances 0.000 description 1
- 230000014509 gene expression Effects 0.000 description 1
- 239000002184 metal Substances 0.000 description 1
- 238000002156 mixing Methods 0.000 description 1
- 238000010295 mobile communication Methods 0.000 description 1
- 239000007787 solid Substances 0.000 description 1
- 238000006467 substitution reaction Methods 0.000 description 1
- 238000012549 training Methods 0.000 description 1
- 230000009466 transformation Effects 0.000 description 1
- 238000000844 transformation Methods 0.000 description 1
Images
Classifications
-
- G—PHYSICS
- G06—COMPUTING; CALCULATING OR COUNTING
- G06N—COMPUTING ARRANGEMENTS BASED ON SPECIFIC COMPUTATIONAL MODELS
- G06N3/00—Computing arrangements based on biological models
- G06N3/02—Neural networks
- G06N3/06—Physical realisation, i.e. hardware implementation of neural networks, neurons or parts of neurons
- G06N3/063—Physical realisation, i.e. hardware implementation of neural networks, neurons or parts of neurons using electronic means
-
- G—PHYSICS
- G06—COMPUTING; CALCULATING OR COUNTING
- G06F—ELECTRIC DIGITAL DATA PROCESSING
- G06F9/00—Arrangements for program control, e.g. control units
- G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
- G06F9/46—Multiprogramming arrangements
- G06F9/54—Interprogram communication
- G06F9/544—Buffers; Shared memory; Pipes
Landscapes
- Engineering & Computer Science (AREA)
- Physics & Mathematics (AREA)
- Theoretical Computer Science (AREA)
- Software Systems (AREA)
- General Engineering & Computer Science (AREA)
- General Physics & Mathematics (AREA)
- Health & Medical Sciences (AREA)
- Life Sciences & Earth Sciences (AREA)
- Biomedical Technology (AREA)
- Biophysics (AREA)
- Neurology (AREA)
- Artificial Intelligence (AREA)
- Computational Linguistics (AREA)
- Data Mining & Analysis (AREA)
- Evolutionary Computation (AREA)
- General Health & Medical Sciences (AREA)
- Molecular Biology (AREA)
- Computing Systems (AREA)
- Mathematical Physics (AREA)
- Image Processing (AREA)
Abstract
本发明公开了一种基于Tensorcore处理int4数据类型的方法、***、设备和存储介质，方法包括：响应于接收到数据类型为int4的数据，根据数据的输入数据维度、权重维度和偏置维度判断数据的批处理大小、输入维度数和输出维度数是否符合要求；响应于数据的批处理大小、输入维度数和输出维度数符合要求，将数据的输入数据从全局内存写入第一共享内存，将数据的权重数据从全局内存写入第二共享内存；将基于第一共享内存和第二共享内存得到的第一计算结果存入第三共享内存以与偏置数据相加得到第二计算结果；以及将第二计算结果返回全局内存。本发明实现了TVM全连接层对于int4数据类型的支持，并相比int8带来很大的性能提升。
Description
技术领域
本发明涉及神经网络领域，更具体地，特别是指一种基于Tensorcore处理int4数据类型的方法、***、计算机设备及可读介质。
背景技术
Dense是神经网络中的全连接层，在神经网络特征提取中占有很重要的地位。假设data是dense的输入数据，形状为(batch，in_dim)，weight为dense的权重，形状为(out_dim，in_dim)，bias为dense的偏置，形状为(out_dim，)，在这里(out_dim，)代表这是一个长度为out_dim的一维向量，如
其中batch代表神经网络训练的批处理大小，in_dim代表dense的输入维度，out_dim代表输出维度。目前的全连接层只能支持data以及weight的数据类型为int8，遇到数据类型为int4的数据则没有办法进行处理。
发明内容
有鉴于此，本发明实施例的目的在于提出一种基于Tensorcore处理int4数据类型的方法、***、计算机设备及计算机可读存储介质，通过将int4数据类型的数据分别写入共享内存中，在共享内存中进行计算，再将最后的结果写入全局内存中，实现了TVM全连接层对于int4数据类型的支持，并相比int8带来很大的性能提升。
基于上述目的，本发明实施例的一方面提供了一种基于Tensorcore处理int4数据类型的方法，包括如下步骤：响应于接收到数据类型为int4的数据，根据所述数据的输入数据维度、权重维度和偏置维度判断所述数据的批处理大小、输入维度数和输出维度数是否符合要求；响应于所述数据的批处理大小、输入维度数和输出维度数符合要求，将所述数据的输入数据从全局内存写入第一共享内存，将所述数据的权重数据从全局内存写入第二共享内存；将基于所述第一共享内存和所述第二共享内存得到的第一计算结果存入第三共享内存以与偏置数据相加得到第二计算结果；以及将所述第二计算结果返回所述全局内存。
在一些实施方式中，所述根据所述数据的输入数据维度、权重维度和偏置维度判断所述数据的批处理大小、输入维度数和输出维度数是否符合要求包括：根据所述数据的输入数据维度、权重维度和偏置维度确定所述数据的批处理大小、输入维度数和输出维度数，并基于与接口形状的比较判断所述批处理大小、输入维度数和输出维度数是否符合要求。
在一些实施方式中，所述根据所述数据的输入数据维度、权重维度和偏置维度判断所述数据的批处理大小、输入维度数和输出维度数是否符合要求包括：判断所述数据的批处理大小是否为接口形状第一分量的整数倍；判断所述数据的输入维度数是否为接口形状第二分量的整数倍；以及判断所述数据的输出维度数是否为接口形状第三分量的整数倍。
在一些实施方式中，所述将基于所述第一共享内存和所述第二共享内存得到的第一计算结果存入第三共享内存以与偏置数据相加得到第二计算结果包括：将所述第一计算结果与int32类型的偏置数据相加得到第二计算结果。
在一些实施方式中，所述将基于所述第一共享内存和所述第二共享内存得到的第一计算结果存入第三共享内存以与偏置数据相加得到第二计算结果包括：将int4类型的输入数据和int4类型的权重数据相乘得到第一计算结果。
在一些实施方式中，所述将所述第二计算结果返回所述全局内存包括：按照批处理大小轴和输出维度轴对所述第二计算结果进行划分。
在一些实施方式中，所述将所述第二计算结果返回所述全局内存包括：将所述批处理大小划分为第一参数和第三参数，将所述输出维度划分成第二参数和第四参数；将所述第三参数和所述第四参数进行融合并再次进行划分。
本发明实施例的另一方面，还提供了一种基于Tensorcore处理int4数据类型***，包括：判断模块，配置用于响应于接收到数据类型为int4的数据，根据所述数据的输入数据维度、权重维度和偏置维度判断所述数据的批处理大小、输入维度数和输出维度数是否符合要求；写入模块，配置用于响应于所述数据的批处理大小、输入维度数和输出维度数符合要求，将所述数据的输入数据从全局内存写入第一共享内存，将所述数据的权重数据从全局内存写入第二共享内存；计算模块，配置用于将基于所述第一共享内存和所述第二共享内存得到的第一计算结果存入第三共享内存以与偏置数据相加得到第二计算结果；返回模块，配置用于将所述第二计算结果返回所述全局内存。
本发明实施例的又一方面，还提供了一种计算机设备，包括：至少一个处理器；以及存储器，所述存储器存储有可在所述处理器上运行的计算机指令，所述指令由所述处理器执行时实现如上方法的步骤。
本发明实施例的再一方面，还提供了一种计算机可读存储介质，计算机可读存储介质存储有被处理器执行时实现如上方法步骤的计算机程序。
本发明具有以下有益技术效果：通过将int4数据类型的数据分别写入共享内存中，在共享内存中进行计算，再将最后的结果写入全局内存中，实现了TVM全连接层对于int4数据类型的支持，并相比int8带来很大的性能提升。
附图说明
为了更清楚地说明本发明实施例或现有技术中的技术方案，下面将对实施例或现有技术描述中所需要使用的附图作简单地介绍，显而易见地，下面描述中的附图仅仅是本发明的一些实施例，对于本领域普通技术人员来讲，在不付出创造性劳动的前提下，还可以根据这些附图获得其他的实施例。
图1为本发明提供的基于Tensorcore处理int4数据类型的方法的实施例的示意图；
图2为本发明提供的基于Tensorcore处理int4数据类型的方法的实施例的流程图；
图3为全连接层整体计算的调度示意图；
图4为单位的矩阵乘法及累加操作计算结果存储到共享内存的调度示意图；
图5为单位的矩阵乘法及累加操作计算的调度示意图；
图6为wmma_matrix_a加载的调度示意图；
图7为wmma_matrix_b加载的调度示意图；
图8为原始输入矩阵和权重矩阵从全局内存加载到共享内存的调度示意图；
图9为本发明提供的基于Tensorcore处理int4数据类型的计算机设备的实施例的硬件结构示意图。
具体实施方式
为使本发明的目的、技术方案和优点更加清楚明白，以下结合具体实施例，并参照附图，对本发明实施例进一步详细说明。
需要说明的是，本发明实施例中所有使用“第一”和“第二”的表述均是为了区分两个相同名称非相同的实体或者非相同的参量，可见“第一”“第二”仅为了表述的方便，不应理解为对本发明实施例的限定，后续实施例对此不再一一说明。
基于上述目的，本发明实施例的第一个方面，提出了一种基于Tensorcore处理int4数据类型的方法的实施例。图1示出的是本发明提供的基于Tensorcore处理int4数据类型的方法的实施例的示意图。如图1所示，本发明实施例包括如下步骤：
S1、响应于接收到数据类型为int4的数据，根据数据的输入数据维度、权重维度和偏置维度判断数据的批处理大小、输入维度数和输出维度数是否符合要求；
S2、响应于数据的批处理大小、输入维度数和输出维度数符合要求，将数据的输入数据从全局内存写入第一共享内存，将数据的权重数据从全局内存写入第二共享内存；
S3、将基于第一共享内存和第二共享内存得到的第一计算结果存入第三共享内存以与偏置数据相加得到第二计算结果；以及
S4、将第二计算结果返回全局内存。
TVM是一个AI编译器，其位于神经网络图(High-Level Differentiable IR)的下方，底层硬件(LLVM，CUDA，Metal)的上方。Tensorcore表示张量核，是一种矩阵乘累加的计算单元，可以进行混合精度以及加速矩阵乘法的计算。矩阵乘累加计算在深度学习网络层算法中，比如卷积层、全连接层等是最重要、最耗时的一部分。Tensorcore可以在一个时钟周期内实现两个4×4矩阵乘法以及与另一个4×4矩阵加法。譬如执行D＝A*B+C运算，其中A、B、C和D是4*4矩阵。矩阵乘法输入A和B是FP16(16位的float类型数据，属于半精度浮点数)，int8(8位整型数据，一个int8占1字节)，int4(4位整型数据，一个int4占半个字节，在cpu上无法直接定义)矩阵，而累加矩阵C和D可能是FP16或FP32(32位的float类型数据，是普通的float类型)，int32(32位整型数据，一个int32占4字节)矩阵。本发明的主要对TVM中的dense操作提供了int4 Tensorcore方案的支持和优化，填补了TVM对dense int4支持的空白和int4 tensorcore支持的空白。
图2示出的是本发明提供的基于Tensorcore处理int4数据类型的方法的实施例的流程图，根据图2对本发明实施例进行说明。
响应于接收到数据类型为int4的数据，根据数据的输入数据维度、权重维度和偏置维度判断数据的批处理大小、输入维度数和输出维度数是否符合要求。也即是对应于图2中判断Batch是否为wmma_m的倍数且输入维度为wmma_k的倍数且输出维度为wmma_n的倍数。dense的算法中涉及到三个输入：输入数据维度为(batch，in_dim)，权重维度为(out_dim，in_dim)，偏置维度为(out_dim，_)。其中batch表示批处理大小，in_dim表示输入维度数，out_dim表示输出维度数。满足输入形状的int4数据可以被处理，可以根据数据的输入数据维度、权重维度和偏置维度判断数据的批处理大小、输入维度数和输出维度数是否符合输入形状的要求。
在一些实施方式中，所述根据所述数据的输入数据维度、权重维度和偏置维度判断所述数据的批处理大小、输入维度数和输出维度数是否符合要求包括：根据所述数据的输入数据维度、权重维度和偏置维度确定所述数据的批处理大小、输入维度数和输出维度数，并基于与接口形状的比较判断所述批处理大小、输入维度数和输出维度数是否符合要求。例如，若dense的输入数据维度是(8，256)，权重维度是(128，256)，偏置维度则为(128，)。那么此时batch＝8，in_dim＝256，out_dim＝128。
在一些实施方式中，所述根据所述数据的输入数据维度、权重维度和偏置维度判断所述数据的批处理大小、输入维度数和输出维度数是否符合要求包括：判断所述数据的批处理大小是否为接口形状第一分量的整数倍；判断所述数据的输入维度数是否为接口形状第二分量的整数倍；以及判断所述数据的输出维度数是否为接口形状第三分量的整数倍。Cuda(Compute Unified Device Architecture，统一计算设备架构)的wmma(warpmatrix multiply-accumulate warp，单位的矩阵乘法及累加操作)接口是调用底层tensorcore的一种方式，本发明实施例采用的是直接调用wmma接口来调用底层tensorcore计算单元的方式。wmma接口的核心有三个，分别是wmma::load_matrix_sync接口(用于矩阵的加载)，wmma::mma_sync接口(用于矩阵乘法的计算)，wmma::store_matrix_sync接口(用于矩阵乘计算完成之后矩阵的存储)。这些接口涉及到三个重要的维度，即<wmma_m，wmma_n，wmma_k>，调用一次wmma接口会完成一个wmma_m*wmma_k的矩阵与wmma_n*wmma_k的矩阵的矩阵乘法。Wmma接口目前支持的数据类型的输入，分别是fp16、int8和int4。对于int4类型的输入，目前仅支持<8，8，32>形状的输入，也就是，调用一次int4的wmma接口将会完成一个8*32矩阵与8*32矩阵的矩阵乘法，最终将得到一个8*8的矩阵结果。wmma_m、wmma_n和wmma_k分别对应batch、out_dim和in_dim，如果想采用Tensorcore，<batch，out_dim，in_dim>必须满足能被wmma接口形状整除。继续上例，batch＝8(对应于wmma_m)，in_dim＝256(对应于wmma_k)，out_dim＝128(对应于wmma_n)，此时能够满足<wmma_n＝8，wmma_n＝8，wmma_k＝32>的调用情况。
响应于数据的批处理大小、输入维度数和输出维度数符合要求，将数据的输入数据从全局内存写入第一共享内存，将数据的权重数据从全局内存写入第二共享内存。也即是对应于图2中从全局内存A根据数据划分取输入数据到共享内存AS和从全局内存B根据数据划分取输入数据到共享内存BS。输入数据和权重数据原本存在全局内存中，但是全局内存的访问速度很慢，共享内存的数据读取与放回耗时很小，并且进行矩阵乘需要进行频繁的数据访问，因而定义输入数据对应的矩阵A和权重数据对应的矩阵B从全局内存读入共享内存“shared”，分别标识为AS和BS。
使用wmma接口需要将数据从共享内存加载到内存作用域“wmma::fragment”。有三种类型的内存作用域:“wmma.matrix_a”、“wmma.matrix_b”和“wmma.accumulator”。其中AS加载到wmma.matrix_a，BS加载到wmma.matrix_b，对应于图2中从共享内存AS取数据到wmma.matrix_a域和从共享内存BS取数据到wmma.matrix_b域。在从共享内存AS(BS)取数据到wmma.matrix_a域(wmma.matrix_b域)后，需要将上述定义有关wmma的数据划分为cuda可以识别的wmma接口。
将基于第一共享内存和第二共享内存得到的第一计算结果存入第三共享内存以与偏置数据相加得到第二计算结果。通过wmma接口计算得到的结果存储在wmma.accumulator，对应于CF的定义。对应于图中wmma.matrix_a与wmma.matrix_b计算结果放到wmma.accumulator。可以将wmma.accumulator中结果放入共享内存CS中，对应于CS的定义。本发明实施例采用的是wmma计算矩阵乘的方式，bias在与最终结果相加的时候必须是在共享内存中，而不能与wmma.accumulator直接相加或者在全局内存中相加，因而最终结果需从wmma.accumulator存储到shared中。
在一些实施方式中，所述将基于所述第一共享内存和所述第二共享内存得到的第一计算结果存入第三共享内存以与偏置数据相加得到第二计算结果包括：将所述第一计算结果与int32类型的偏置数据相加得到第二计算结果。
在一些实施方式中，所述将基于所述第一共享内存和所述第二共享内存得到的第一计算结果存入第三共享内存以与偏置数据相加得到第二计算结果包括：将int4类型的输入数据和int4类型的权重数据相乘得到第一计算结果。
dense的计算公式如下：
其中i代表batch，j代表dense的输出维度out_dim，k代表dense的输入维度in_dim，output表示输出，data表示输入，weight表示权重，bias表示偏置。输入数据和权重均为int4类型，调用tensocore接口之后，两个int4矩阵的相乘得到的结果是int32类型，因而偏置的类型是int32。Oputput为计算规则加上一个标识，例如，标识为dense_tensorcore_int4，之后TVM会根据这个标识进入相应的调度部分。
将第二计算结果返回全局内存。将处理完bias的数据从共享内存写回全局内存。对应于图2中将共享内存CS中的结果放入全局内存C中。
在一些实施方式中，所述将所述第二计算结果返回所述全局内存包括：按照批处理大小轴和输出维度轴对所述第二计算结果进行划分。
在一些实施方式中，所述将所述第二计算结果返回所述全局内存包括：将所述批处理大小划分为第一参数和第三参数，将所述输出维度划分成第二参数和第四参数；将所述第三参数和所述第四参数进行融合并再次进行划分。
图3示出的是全连接层整体计算的调度示意图。如图3所示，在TVM中可以分别对batch轴和输出维度轴进行数据划分。将batch划分为blockIdx.x和wmma_m*warp_row_tiles*block_row_warps；将输出维度划分为blockIdx.y和wmma_n*warp_col_tiles*block_col_warps，并进行重排序操作，变划分顺序为blockIdx.x，blockIdx.y，wmma_m*warp_row_tiles*block_row_warps，wmma_n*warp_col_tiles*block_col_warps，然后将wmma_m*warp_row_tiles*block_row_warps和wmma_n*warp_col_tiles*block_col_warps进行融合，继续划分，然后整个划分最终的形式就是blockIdx.x，blockIdx.y，threadIdx.z，threadIdx.y，threadIdx，向量化参数。需要注意的是，TVM中存在reorder(重排序)操作，为了表示这个操作，最终划分的前后顺序就是最后生成的前后顺序，越往后的越在整个计算循环的内维，比如，最终的threadIdx.x和向量化参数，代表着向量化参数在最内维，threadIdx.x在向量化参数的外层。举个例子：
for(int j＝0；j<threadIdx.x；j++)
for(int i＝0；i<向量化参数；i++)
图4示出的是单位的矩阵乘法及累加操作计算结果存储到共享内存的调度示意图。将Wmma计算结果存储到共享内存的调度方式定义为CS。如图4所示，Wmma计算结果存储到共享内存的调度方式与dense整体计算的调度类似，也是将batch划分为blockIdx.x和wmma_m*warp_row_tiles*block_row_warps；将输出维度划分为blockIdx.y和wmma_n*warp_col_tiles*block_col_warps，并进行重排序操作，变划分顺序为blockIdx.x，blockIdx.y，wmma_m*warp_row_tiles*block_row_warps和wmma_n*warp_col_tiles*block_col_warps。后续是将wmma_m*warp_row_tiles*block_row_warps划分成block_row_warps、warp_row_tiles和wmma_m，将wmma_n*warp_col_tiles*block_col_warps划分成block_col_warps、warp_col_tiles和wmma_n。然后进行Tensorize操作，把TVM生成的代码替换成wmma::store_matrix_sync接口代码
图5示出的是单位的矩阵乘法及累加操作计算的调度示意图。将Wmma计算的调度定义为CF。如图5所示，Wmma计算的调度主要是加载到wmma.matrix_a和wmma.matrix_b中的数据经过wmma::mma_sync接口计算得出的结果放到wmma.accumulator的过程。CF的计算主要是在CS的block_col_warps以后的部分，CF承接了之前CS其它部分的划分。在将wmma.matrix_a和wmma.matrix_b计算结果放到wmma.accumulator后，可以将上述定义有关wmma的数据划分转换为cuda可以识别的wmma接口。
图6示出的是wmma_matrix_a加载的调度示意图。将wmma_matrix_a加载的调度定义为AF。如图6所示，Wmma_matrix_a加载的调度主要是从共享内存中加载输入数据到wmma.matrix_a内存作用域，这部分的计算主要是在CF的chunk以后的部分，AF承接了之前CF其它部分的划分。
图7示出的是wmma_matrix_b加载的调度示意图。将wmma_matrix_b加载的调度定义为BF。如图7所示，Wmma_matrix_b加载的调度主要是从共享内存中加载去权重数据到wmma.matrix_b内存作用域，这部分的计算主要是在CF的chunk以后的部分，BF承接了之前CF其它部分的划分。
图8示出的是原始输入矩阵和权重矩阵从全局内存加载到共享内存的调度示意图。将原始输入矩阵和权重矩阵从全局内存加载到共享内存的调度定义为AS(或BS)。如图8所示，这部分包括了输入数据和权重数据的调度，这部分的计算主要是在CF的chunk个数以后的部分，AS(BS)承接了之前CF其它部分的划分。在这个过程中可以加入双缓冲优化，即提前取出下一次计算所需要的数据到共享内存中，在TVM中可以通过调用double_buffer()函数实现。
将完成wmma接口对普通代码的替换定义为tensorize。这部分需要提供准确的步长参数。
在输入维度＝2048，输出维度＝1024的条件下，对于输入数据(batch，2048)，weights(1024，2048)的dense计算，结果如下表所示，int4 tensorcore相比于int8tensorcore提高了1.5倍左右的性能。在下表中，speedup一列的参数代表着加速比，计算方法是用int8_tensorcore的计算时间除以int4_tensorcore的时间。
本发明实施例在resnet50 int4网络中，当计算到达dense模块时，会优先判断batch、out_dim、in_dim是否能被wmma接口所支持的形状整除，若能，则进入int4Tensosore解决方案，经过Tensor Core的计算得出int32的结果传递给下层网络。若不满足条件，则会根据条件，进入其他的int4方案(只不过现在还没有其他的int4方案，需后续提供支持)。
本发明实施例主要对TVM中的dense操作提供了Tensor Core int4方案的支持，填补了int4情况下dense使用tensorcre的空白，并带来很大的性能提升。根据TVM现有规则，设计了基于Tensorcore的dense的compute(计算)及schedule(计划)实现，根据int4数据特点，定义了相关的数据划分，wmma接口调用等，并添加了减小bank(堆积)冲突，从全局中向量化读取数据到共享内存，从共享内存向量化写回全局内存，double buffer(双倍缓存)优化，确保其能在resnet50等神经网络中应用并提升网络性能。经过测试，tensorcore int4方案的性能为tensorcore int8方案的1.5倍左右。
需要特别指出的是，上述基于Tensorcore处理int4数据类型的方法的各个实施例中的各个步骤均可以相互交叉、替换、增加、删减，因此，这些合理的排列组合变换之于基于Tensorcore处理int4数据类型的方法也应当属于本发明的保护范围，并且不应将本发明的保护范围局限在实施例之上。
基于上述目的，本发明实施例的第二个方面，提出了一种基于Tensorcore处理int4数据类型的***，包括：判断模块，配置用于响应于接收到数据类型为int4的数据，根据所述数据的输入数据维度、权重维度和偏置维度判断所述数据的批处理大小、输入维度数和输出维度数是否符合要求；写入模块，配置用于响应于所述数据的批处理大小、输入维度数和输出维度数符合要求，将所述数据的输入数据从全局内存写入第一共享内存，将所述数据的权重数据从全局内存写入第二共享内存；计算模块，配置用于将基于所述第一共享内存和所述第二共享内存得到的第一计算结果存入第三共享内存以与偏置数据相加得到第二计算结果；返回模块，配置用于将所述第二计算结果返回所述全局内存。
在一些实施方式中，所述判断模块配置用于：根据所述数据的输入数据维度、权重维度和偏置维度确定所述数据的批处理大小、输入维度数和输出维度数，并基于与接口形状的比较判断所述批处理大小、输入维度数和输出维度数是否符合要求。
在一些实施方式中，所述判断模块配置用于：判断所述数据的批处理大小是否为接口形状第一分量的整数倍；判断所述数据的输入维度数是否为接口形状第二分量的整数倍；以及判断所述数据的输出维度数是否为接口形状第三分量的整数倍。
在一些实施方式中，所述计算模块配置用于：将所述第一计算结果与int32类型的偏置数据相加得到第二计算结果。
在一些实施方式中，所述计算模块配置用于：将int4类型的输入数据和int4类型的权重数据相乘得到第一计算结果。
在一些实施方式中，所述返回模块配置用于：按照批处理大小轴和输出维度轴对所述第二计算结果进行划分。
在一些实施方式中，所述返回模块配置用于：将所述批处理大小划分为第一参数和第三参数，将所述输出维度划分成第二参数和第四参数；将所述第三参数和所述第四参数进行融合并再次进行划分。
基于上述目的，本发明实施例的第三个方面，提出了一种计算机设备，包括：至少一个处理器；以及存储器，存储器存储有可在处理器上运行的计算机指令，指令由处理器执行以实现如下步骤：S1、响应于接收到数据类型为int4的数据，根据数据的输入数据维度、权重维度和偏置维度判断数据的批处理大小、输入维度数和输出维度数是否符合要求；S2、响应于数据的批处理大小、输入维度数和输出维度数符合要求，将数据的输入数据从全局内存写入第一共享内存，将数据的权重数据从全局内存写入第二共享内存；S3、将基于第一共享内存和第二共享内存得到的第一计算结果存入第三共享内存以与偏置数据相加得到第二计算结果；以及S4、将第二计算结果返回全局内存。
在一些实施方式中，所述根据所述数据的输入数据维度、权重维度和偏置维度判断所述数据的批处理大小、输入维度数和输出维度数是否符合要求包括：根据所述数据的输入数据维度、权重维度和偏置维度确定所述数据的批处理大小、输入维度数和输出维度数，并基于与接口形状的比较判断所述批处理大小、输入维度数和输出维度数是否符合要求。
在一些实施方式中，所述根据所述数据的输入数据维度、权重维度和偏置维度判断所述数据的批处理大小、输入维度数和输出维度数是否符合要求包括：判断所述数据的批处理大小是否为接口形状第一分量的整数倍；判断所述数据的输入维度数是否为接口形状第二分量的整数倍；以及判断所述数据的输出维度数是否为接口形状第三分量的整数倍。
在一些实施方式中，所述将基于所述第一共享内存和所述第二共享内存得到的第一计算结果存入第三共享内存以与偏置数据相加得到第二计算结果包括：将所述第一计算结果与int32类型的偏置数据相加得到第二计算结果。
在一些实施方式中，所述将基于所述第一共享内存和所述第二共享内存得到的第一计算结果存入第三共享内存以与偏置数据相加得到第二计算结果包括：将int4类型的输入数据和int4类型的权重数据相乘得到第一计算结果。
在一些实施方式中，所述将所述第二计算结果返回所述全局内存包括：按照批处理大小轴和输出维度轴对所述第二计算结果进行划分。
在一些实施方式中，所述将所述第二计算结果返回所述全局内存包括：将所述批处理大小划分为第一参数和第三参数，将所述输出维度划分成第二参数和第四参数；将所述第三参数和所述第四参数进行融合并再次进行划分。
如图9所示，为本发明提供的上述基于Tensorcore处理int4数据类型的计算机设备的一个实施例的硬件结构示意图。
以如图9所示的装置为例，在该装置中包括一个处理器301以及一个存储器302，并还可以包括：输入装置303和输出装置304。
处理器301、存储器302、输入装置303和输出装置304可以通过总线或者其他方式连接，图9中以通过总线连接为例。
存储器302作为一种非易失性计算机可读存储介质，可用于存储非易失性软件程序、非易失性计算机可执行程序以及模块，如本申请实施例中的基于Tensorcore处理int4数据类型的方法对应的程序指令/模块。处理器301通过运行存储在存储器302中的非易失性软件程序、指令以及模块，从而执行服务器的各种功能应用以及数据处理，即实现上述方法实施例的基于Tensorcore处理int4数据类型的方法。
存储器302可以包括存储程序区和存储数据区，其中，存储程序区可存储操作***、至少一个功能所需要的应用程序；存储数据区可存储根据基于Tensorcore处理int4数据类型的方法的使用所创建的数据等。此外，存储器302可以包括高速随机存取存储器，还可以包括非易失性存储器，例如至少一个磁盘存储器件、闪存器件、或其他非易失性固态存储器件。在一些实施例中，存储器302可选包括相对于处理器301远程设置的存储器，这些远程存储器可以通过网络连接至本地模块。上述网络的实例包括但不限于互联网、企业内部网、局域网、移动通信网及其组合。
输入装置303可接收输入的用户名和密码等信息。输出装置304可包括显示屏等显示设备。
一个或者多个基于Tensorcore处理int4数据类型的方法对应的程序指令/模块存储在存储器302中，当被处理器301执行时，执行上述任意方法实施例中的基于Tensorcore处理int4数据类型的方法。
执行上述基于Tensorcore处理int4数据类型的方法的计算机设备的任何一个实施例，可以达到与之对应的前述任意方法实施例相同或者相类似的效果。
本发明还提供了一种计算机可读存储介质，计算机可读存储介质存储有被处理器执行时执行如上方法的计算机程序。
最后需要说明的是，本领域普通技术人员可以理解实现上述实施例方法中的全部或部分流程，可以通过计算机程序来指令相关硬件来完成，基于Tensorcore处理int4数据类型的方法的程序可存储于一计算机可读取存储介质中，该程序在执行时，可包括如上述各方法的实施例的流程。其中，程序的存储介质可为磁碟、光盘、只读存储记忆体(ROM)或随机存储记忆体(RAM)等。上述计算机程序的实施例，可以达到与之对应的前述任意方法实施例相同或者相类似的效果。
以上是本发明公开的示例性实施例，但是应当注意，在不背离权利要求限定的本发明实施例公开的范围的前提下，可以进行多种改变和修改。根据这里描述的公开实施例的方法权利要求的功能、步骤和/或动作不需以任何特定顺序执行。此外，尽管本发明实施例公开的元素可以以个体形式描述或要求，但除非明确限制为单数，也可以理解为多个。
应当理解的是，在本文中使用的，除非上下文清楚地支持例外情况，单数形式“一个”旨在也包括复数形式。还应当理解的是，在本文中使用的“和/或”是指包括一个或者一个以上相关联地列出的项目的任意和所有可能组合。
上述本发明实施例公开实施例序号仅仅为了描述，不代表实施例的优劣。
本领域普通技术人员可以理解实现上述实施例的全部或部分步骤可以通过硬件来完成，也可以通过程序来指令相关的硬件完成，程序可以存储于一种计算机可读存储介质中，上述提到的存储介质可以是只读存储器，磁盘或光盘等。
所属领域的普通技术人员应当理解：以上任何实施例的讨论仅为示例性的，并非旨在暗示本发明实施例公开的范围(包括权利要求)被限于这些例子；在本发明实施例的思路下，以上实施例或者不同实施例中的技术特征之间也可以进行组合，并存在如上的本发明实施例的不同方面的许多其它变化，为了简明它们没有在细节中提供。因此，凡在本发明实施例的精神和原则之内，所做的任何省略、修改、等同替换、改进等，均应包含在本发明实施例的保护范围之内。
Claims (10)
1.一种基于Tensorcore处理int4数据类型的方法，其特征在于，包括以下步骤：
响应于接收到数据类型为int4的数据，根据所述数据的输入数据维度、权重维度和偏置维度判断所述数据的批处理大小、输入维度数和输出维度数是否符合要求；
响应于所述数据的批处理大小、输入维度数和输出维度数符合要求，将所述数据的输入数据从全局内存写入第一共享内存，将所述数据的权重数据从全局内存写入第二共享内存；
将基于所述第一共享内存和所述第二共享内存得到的第一计算结果存入第三共享内存以与偏置数据相加得到第二计算结果；以及
将所述第二计算结果返回所述全局内存。
2.根据权利要求1所述的方法，其特征在于，所述根据所述数据的输入数据维度、权重维度和偏置维度判断所述数据的批处理大小、输入维度数和输出维度数是否符合要求包括：
根据所述数据的输入数据维度、权重维度和偏置维度确定所述数据的批处理大小、输入维度数和输出维度数，并基于与接口形状的比较判断所述批处理大小、输入维度数和输出维度数是否符合要求。
3.根据权利要求2所述的方法，其特征在于，所述根据所述数据的输入数据维度、权重维度和偏置维度判断所述数据的批处理大小、输入维度数和输出维度数是否符合要求包括：
判断所述数据的批处理大小是否为接口形状第一分量的整数倍；
判断所述数据的输入维度数是否为接口形状第二分量的整数倍；以及
判断所述数据的输出维度数是否为接口形状第三分量的整数倍。
4.根据权利要求1所述的方法，其特征在于，所述将基于所述第一共享内存和所述第二共享内存得到的第一计算结果存入第三共享内存以与偏置数据相加得到第二计算结果包括：
将所述第一计算结果与int32类型的偏置数据相加得到第二计算结果。
5.根据权利要求4所述的方法，其特征在于，所述将基于所述第一共享内存和所述第二共享内存得到的第一计算结果存入第三共享内存以与偏置数据相加得到第二计算结果包括：
将int4类型的输入数据和int4类型的权重数据相乘得到第一计算结果。
6.根据权利要求1所述的方法，其特征在于，所述将所述第二计算结果返回所述全局内存包括：
按照批处理大小轴和输出维度轴对所述第二计算结果进行划分。
7.根据权利要求6所述的方法，其特征在于，所述将所述第二计算结果返回所述全局内存包括：
将所述批处理大小划分为第一参数和第三参数，将所述输出维度划分成第二参数和第四参数；
将所述第三参数和所述第四参数进行融合并再次进行划分。
8.一种基于Tensorcore处理int4数据类型的***，其特征在于，包括：
判断模块，配置用于响应于接收到数据类型为int4的数据，根据所述数据的输入数据维度、权重维度和偏置维度判断所述数据的批处理大小、输入维度数和输出维度数是否符合要求；
写入模块，配置用于响应于所述数据的批处理大小、输入维度数和输出维度数符合要求，将所述数据的输入数据从全局内存写入第一共享内存，将所述数据的权重数据从全局内存写入第二共享内存；
计算模块，配置用于将基于所述第一共享内存和所述第二共享内存得到的第一计算结果存入第三共享内存以与偏置数据相加得到第二计算结果；以及
返回模块，配置用于将所述第二计算结果返回所述全局内存。
9.一种计算机设备，其特征在于，包括：
至少一个处理器；以及
存储器，所述存储器存储有可在所述处理器上运行的计算机指令，所述指令由所述处理器执行时实现权利要求1-7任意一项所述方法的步骤。
10.一种计算机可读存储介质，所述计算机可读存储介质存储有计算机程序，其特征在于，所述计算机程序被处理器执行时实现权利要求1-7任意一项所述方法的步骤。
Priority Applications (2)
| Application Number | Priority Date | Filing Date | Title |
|---|---|---|---|
| CN202010980721.9A CN112232496A (zh) | 2020-09-17 | 2020-09-17 | 一种基于Tensorcore处理int4数据类型的方法、***、设备及介质 |
| PCT/CN2021/109214 WO2022057459A1 (zh) | 2020-09-17 | 2021-07-29 | 一种基于Tensorcore处理int4数据类型的方法、***、设备及介质 |
Applications Claiming Priority (1)
| Application Number | Priority Date | Filing Date | Title |
|---|---|---|---|
| CN202010980721.9A CN112232496A (zh) | 2020-09-17 | 2020-09-17 | 一种基于Tensorcore处理int4数据类型的方法、***、设备及介质 |
Publications (1)
| Publication Number | Publication Date |
|---|---|
| CN112232496A true CN112232496A (zh) | 2021-01-15 |
Family
ID=74107931
Family Applications (1)
| Application Number | Title | Priority Date | Filing Date |
|---|---|---|---|
| CN202010980721.9A Withdrawn CN112232496A (zh) | 2020-09-17 | 2020-09-17 | 一种基于Tensorcore处理int4数据类型的方法、***、设备及介质 |
Country Status (2)
| Country | Link |
|---|---|
| CN (1) | CN112232496A (zh) |
| WO (1) | WO2022057459A1 (zh) |
Cited By (1)
| Publication number | Priority date | Publication date | Assignee | Title |
|---|---|---|---|---|
| WO2022057459A1 (zh) * | 2020-09-17 | 2022-03-24 | 苏州浪潮智能科技有限公司 | 一种基于Tensorcore处理int4数据类型的方法、***、设备及介质 |
Family Cites Families (5)
| Publication number | Priority date | Publication date | Assignee | Title |
|---|---|---|---|---|
| CN111124656B (zh) * | 2018-10-31 | 2023-09-15 | 伊姆西Ip控股有限责任公司 | 用于向专用计算资源分配任务的方法、设备和计算机可读存储介质 |
| CN111539526B (zh) * | 2020-04-24 | 2022-12-06 | 苏州浪潮智能科技有限公司 | 一种神经网络卷积的方法和设备 |
| CN111859270B (zh) * | 2020-07-14 | 2022-11-25 | 苏州浪潮智能科技有限公司 | 一种神经网络的卷积方法和装置 |
| CN111860838B (zh) * | 2020-07-24 | 2022-12-20 | 苏州浪潮智能科技有限公司 | 一种神经网络的全连接层计算方法和装置 |
| CN112232496A (zh) * | 2020-09-17 | 2021-01-15 | 苏州浪潮智能科技有限公司 | 一种基于Tensorcore处理int4数据类型的方法、***、设备及介质 |
-
2020
- 2020-09-17 CN CN202010980721.9A patent/CN112232496A/zh not_active Withdrawn
-
2021
- 2021-07-29 WO PCT/CN2021/109214 patent/WO2022057459A1/zh active Application Filing
Cited By (1)
| Publication number | Priority date | Publication date | Assignee | Title |
|---|---|---|---|---|
| WO2022057459A1 (zh) * | 2020-09-17 | 2022-03-24 | 苏州浪潮智能科技有限公司 | 一种基于Tensorcore处理int4数据类型的方法、***、设备及介质 |
Also Published As
| Publication number | Publication date |
|---|---|
| WO2022057459A1 (zh) | 2022-03-24 |
Similar Documents
| Publication | Publication Date | Title |
|---|---|---|
| US20210319284A1 (en) | System and architecture including processor and neural network accelerator | |
| CN107341541B (zh) | 一种用于执行全连接层神经网络训练的装置和方法 | |
| KR20200000480A (ko) | 처리 장치 및 처리 방법 | |
| CN111915001A (zh) | 卷积计算引擎、人工智能芯片以及数据处理方法 | |
| CN111915003A (zh) | 一种神经网络硬件加速器 | |
| CN115034402A (zh) | 模型推理性能的优化方法、装置及相关产品 | |
| CN113313247B (zh) | 基于数据流架构的稀疏神经网络的运算方法 | |
| CN110009048B (zh) | 一种神经网络模型的构建方法以及设备 | |
| CN111949932A (zh) | 一种在TVM中实现TensorCore卷积计算的方法及*** | |
| CN112232496A (zh) | 一种基于Tensorcore处理int4数据类型的方法、***、设备及介质 | |
| US7734456B2 (en) | Method and apparatus for priority based data processing | |
| CN111860838B (zh) | 一种神经网络的全连接层计算方法和装置 | |
| CN112966729A (zh) | 一种数据处理方法、装置、计算机设备及存储介质 | |
| CN115469931B (zh) | 一种循环程序的指令优化方法、装置、***、设备及介质 | |
| CN113469337B (zh) | 用于优化神经网络模型的编译方法及其相关产品 | |
| KR102372869B1 (ko) | 인공 신경망을 위한 행렬 연산기 및 행렬 연산 방법 | |
| CN113642722A (zh) | 用于卷积计算的芯片及其控制方法、电子装置 | |
| CN115437756A (zh) | 计算流图调度方案的生成方法、装置、电子设备及计算机可读存储介质 | |
| US20190073584A1 (en) | Apparatus and methods for forward propagation in neural networks supporting discrete data | |
| CN113570053A (zh) | 一种神经网络模型的训练方法、装置以及计算设备 | |
| CN111832714A (zh) | 运算方法及装置 | |
| CN114692847B (zh) | 数据处理电路、数据处理方法及相关产品 | |
| CN117114055B (zh) | 面向工业应用场景的fpga二值神经网络加速方法 | |
| CN118093143B (zh) | 大语言模型解码阶段的数据调度方法和装置 | |
| EP3073387A1 (en) | Controlling data flow between processors in a processing system |
Legal Events
| Date | Code | Title | Description |
|---|---|---|---|
| PB01 | Publication | ||
| PB01 | Publication | ||
| SE01 | Entry into force of request for substantive examination | ||
| SE01 | Entry into force of request for substantive examination | ||
| WW01 | Invention patent application withdrawn after publication |
Application publication date: 20210115 |
|
| WW01 | Invention patent application withdrawn after publication |