CN101937425B - 基于gpu众核平台的矩阵并行转置方法 - Google Patents

Abstract

本发明公开了一种基于GPU众核平台的矩阵转置并行方法，该方法是以线程块为单位进行的，因为同一个线程块内分配在共享存储器上面的临时数据是对本线程块内所有的线程共享的，所以硬件上的实现是先把一个线程块的输入数据全导入共享存储器，然后再针对这线程块内的所有数据以一个处理核心负责一个线程的模式完成转置，因此实现效率高。这样就实现了在含有成百上千个处理核心的GPU上高效、并行完成矩阵转置。

Description

基于GPU众核平台的矩阵并行转置方法

技术领域

本发明涉及一种基于GPU众核平台的矩阵并行转置方法以及在工程实践中的应用。

背景技术

在工程实践中如SAR成像、通信、微波等，矩阵转置是一种经常需求的运算，并且常常在整个程序的运行时间里占很大一部分，虽然目前已经有各种各样的算法，但是均是在CPU上进行串行处理，考虑到CPU的集成度比较低，目前一块芯片最多可以集成8块CPU处理核心，因此算法并行度不高，运行效率比较低。GPGPU是一种处理密集型数据和并行数据的可以内含成百上千个处理核心的处理器，因此更适合于大规模并行计算。但是，目前并没有在数百个处理核心上同时并行处理完成矩阵转置运算。

发明内容

本发明所要解决的技术问题是提供一种基于众核平台的矩阵并行转置方法，能够在包含成百上千个处理核心的GPU上并行、高效的实现矩阵转置运算以及在工程中应用。

该基于众核平台的矩阵转置并行方法包括以下步骤：

第一步：转置的矩阵为M行N列，将每一行和后一行的数据首尾相接，将M×N的二维数组转化为包含M×N个元素的一维数组，并标识为A；

第二步：计算A所占存储空间的大小，并标识为B；

第三步：根据计算出的B，在GPU的全局存储器上分配同样大小的存储空间标识为C，然后将内存上的数据A拷贝到GPU的全局存储器的C；

第四步：数据分割：GPU的组织包括线程网格层、线程块层、线程层，一个网格包含两个以上线程块，为一维组织或二、三维组织，每个块包含两个以上运行线程，同时，每个块有一个被该块内所有运行线程可见的共享存储器；线程块层设置为二维且每维的大小一致；线程网格层设置成二维，大小由原矩阵的大小M、N决定；

第五步：输入数据索引：以线程块为单位运算，根据原矩阵模型和GPU的内置变量计算出行方向对原始矩阵的索引和列方向的索引，从而计算出对C的索引；

第六步：根据第五步得到的对输入数据的索引，以线程块为单位，将C中对应于每个线程块的元素导入各自线程块的共享存储器中；

第七步：输出数据索引：以线程块为单位运算，首先根据转置后的矩阵模型，对掉行索引和列索引，借助GPU通用编程的内置变量计算出行方向对输出矩阵的索引以及列方向的索引，从而计算出输出矩阵的索引；

第八步：根据第七步得到的对输出数据的索引，以线程块为单位，将各个线程块的共享存储器中的缓存导出到一维数组C；

第九步：将C拷贝到内存A，然后对A每M个元素切割一次成为一行并按序组合成一个N行M列的二维数组，此即为转置后的矩阵。

通过以上步骤就完成了矩阵转置在GPU众核平台下的实现。

本发明的有益效果：

1、通过科学的统筹安排，成百上千个处理核心并行处理数据，高效完成了运算并提高了运算精度；

2、进行了数据分割，采用线程块处理的模式，用编程的复杂度换取了程序运行效率，并减小了矩阵转置尺寸的限制；

3、使用了高速缓存，并依据单次通信海量运算的原则，大大隐藏了通信消耗。

附图说明

图1为本发明基于众核平台的矩阵并行转置方法的流程图；

具体实施方式

以在SAR系统中4096×4096点目标成像中二维矩阵转置的应用为例，本专利的实现主要包括以下流程：

1.对于原始4096×4096二维矩阵，将每一行和后一行的数据首尾相接，转化为包含4096×4096个元素的一维数组，并标识为idata[16777216]；

2.计算idata[16777216]所占存储空间的大小：在本系统应用中，矩阵元素的数据类型为浮点型，那么idata[16777216]所占存储空间为：mem_size＝sizeof(float)*4096*4096＝67108864bytes；

3.根据计算出的mem_size，在显存上分配同样大小的存储空间标识为idata_gpu[16777216]，然后将内存上的数据idata拷贝到显存的idata_gpu；

4.数据分割：先考虑线程块层，为了方便处理，我们设置为二维的而且每维的大小一致即每个线程块包含16×16个线程；考虑网格层，为方便处理，我们也设置成二维的，但是大小由原矩阵的大小4096*4096决定：沿行的方向gridx＝4096/16＝256，沿列的方向gridy＝4096/16＝256；

5.输入数据索引：以线程块为单位运算，首先根据原矩阵模型和GPU通用编程的内置变量blockIdx.x和threadIdx.x计算出行方向对原始矩阵的索引：xIndex＝blockIdx.x*16+threadIdx.x，同理计算出列方向的索引：yIndex＝blockIdx.y*16+threadIdx.y，根据xIndex和yIndex计算出对idata_gpu[16777216]的索引：index_in＝yIndex*4096+xIndex；

6.根据上面得到的对输入数据的索引，以线程块为单位，将idata_gpu[16777216]中对应于每个线程块的16×16个元素导入各自线程块的共享存储器中的缓存，记为block_temp[16][16]，导入时的对应关系为：block_temp[threadIdx.y][threadIdx.x]＝idata_gpu[index_in]；

7.输出数据索引：以线程块为单位运算，首先根据转置后的矩阵模型和GPU通用编程的内置变量blockIdx.x和threadIdx.x计算出行方向对输出矩阵的索引：xIndex＝blockIdx.y*16+threadIdx.x，同理计算出列方向的索引：yIndex＝blockIdx.x*16+threadIdx.y，根据xIndex和yIndex计算出对odata_gpu[16777216]的索引：index_out＝yIndex*4096+xIndex；

8.根据上面得到的对输出数据的索引，以线程块为单位，将各个线程块的共享存储器中的缓存block_temp[16][16]，导出到输出一维数组odata_gpu[16777216]，导出时的对应关系为：odata_gpu[index_out]＝block_temp[threadIdx.x][threadIdx.y]，注意与输入数据索引相比，block_temp的两个索引标号对调了；

9.将odata_gpu[16777216]通过函数cudamemcopy拷贝到内存odata[16777216]，然后对odata每4096个元素切割一次成为一行并按序组合成一个4096行4096列的二维数组，此即为转置后的矩阵。将得到的结果返回SAR成像算法中，并最终得到点目标的成像结果。

通过以上方法就实现了SAR成像系统中的矩阵并行转置。该过程是以线程块为单位进行的，因为同一个线程块内部分配在共享存储器上面的临时数据是对本线程块内所有的线程共享的，所以硬件上的实现是先把一个线程块的输入数据全导入共享存储器，然后再对这线程块内的所有数据每个处理核心负责一个线程完成转置，因此实现效率高。这样就实现了在含有成百上千个处理核心的GPU上高效、并行完成矩阵转置。

Claims (1)

1.一种基于GPU众核平台的矩阵并行转置方法，其特征在于：包括以下步骤：

第一步：转置的矩阵为M行N列，将每一行和后一行的数据首尾相接，将M×N的二维数组转化为包含M×N个元素的一维数组，并标识为A；

第二步：计算A所占存储空间的大小，并标识为B；

第三步：根据计算出的B，在GPU的全局存储器上分配同样大小的存储空间标识为C，然后将内存上的数据A拷贝到GPU的全局存储器的C；

第四步：数据分割：GPU的组织包括线程网格层、线程块层、线程层，一个网格包含两个以上线程块，为一维组织或二、三维组织，每个块包含两个以上运行线程，同时，每个块有一个被该块内所有运行线程可见的共享存储器；线程块层设置为二维且每维的大小一致；线程网格层设置成二维，大小由原矩阵的大小M、N决定；

第五步：输入数据索引：以线程块为单位运算，根据原矩阵模型和GPU的内置变量计算出行方向对原始矩阵的索引和列方向的索引，从而计算出对C的索引；

第六步：根据第五步得到的对输入数据的索引，以线程块为单位，将C中对应于每个线程块的元素导入各自线程块的共享存储器中；

第七步：输出数据索引：以线程块为单位运算，首先根据转置后的矩阵模型，对掉行索引和列索引，借助GPU通用编程的内置变量计算出行方向对输出矩阵的索引以及列方向的索引，从而计算出输出矩阵的索引；

第八步：根据第七步得到的对输出数据的索引，以线程块为单位，将各个线程块的共享存储器中的缓存导出到一维数组C；

第九步：将C拷贝到内存A，然后对A每M个元素切割一次成为一行并按序组合成一个N行M列的二维数组，此即为转置后的矩阵。