基于图形处理器的矩阵转置优化方法
【专利摘要】本发明是基于图形处理器的矩阵转置优化方法，包括将R行S列的输入矩阵转化为一维数组并分配存储空间且拷贝数据；设置二维索引空间；计算工作项的全局标识符、工作组标识符和局部标识符；对矩阵分块并与工作组相对应；在工作组内申请本地内存且将数据复制到本地存储器中并同步等待数据复制完成；计算转置后数据在全局存储器中的行和列索引；计算输出数据在全局存储器及本地存储器中的位置；将本地存储器数据赋值给全局存储器中的一维数组，实现无冲突存储器合并访问；将一维数组拷贝到内存中形成S行R列的转置后的矩阵。本发明实现合并访问和矩阵转置的并行计算，提高程序的执行效率。
【专利说明】基于图形处理器的矩阵转置优化方法
【技术领域】
[0001]本发明属于通用计算图形处理器【技术领域】，主要涉及一种基于图形处理器的矩阵转置优化方法。
【背景技术】
[0002]通用计算图形处理器(GeneralPurpose Computing on Graphics ProcessingUnits-GPGPU)，是利用显卡图形处理器来处理通用计算任务的技术。图形处理器分担了中央处理器的计算任务，将计算机的处理速度提升成百上千倍，甚至更快。由此而产生了开放计算语言(Open Computing Language)。开放计算语言由各种处理器和软件制造商代表组成的计算工作组管理，它提供了一套标准的应用程序编程接口，让程序员对图形处理器的编程更便捷。
[0003]在信号处理，图像显示等工程领域，经常会用到矩阵转置运算，目前各类算法多是基于中央处理器进行串行计算，这种做法消耗了大量的时间，效率较低。图形处理器具有数以千计的处理器核心，同时具有远高于中央处理器的内存带宽，特别适合大规模数据的并行计算。
[0004]但是，存储器带宽是影响计算机性能的主要瓶颈之一。处理器的计算能力远远超过内存带宽。目前并没有在图形处理器核心上对矩阵并行化转置进行内存访问优化的运算。
【发明内容】
[0005](一 )本发明的要解决的技术问题:
[0006]本发明的目的是提供一种基于图形处理器的矩阵并行转置优化方法，能够在拥有数以千计个核心处理器的图形处理器上实现并行、快速、高效的矩阵转置以及在工程领域的应用。
[0007]( 二)本发明的技术方案
[0008]为达成所述目的，本发明提供基于图形处理器的矩阵转置优化方法，实现该方法的步骤包括:
[0009]步骤S1:将输入矩阵存储在中央处理器的磁盘上并按行读入，生成具有R行S列的二维数组，将R行S列的二维数组转化为含有RXS个元素的一维数组；
[0010]步骤S2:根据一维数组的大小和输入矩阵数据的类型，计算出一维数组所需要的存储空间；
[0011]步骤S3:在图形处理器的全局存储器上为一维数组分配大小相同的两片存储空间；将一维数组中的数据拷贝到图形处理器的全局存储器中的存储空间；
[0012]步骤S4:根据图形处理器具有的N维域模型、工作组和多个工作项，设置具有二维域的索引空间；
[0013]步骤S5:根据图形处理器内置函数分别对全局存储器中每个工作项的行维度和列维度进行计算，得到全局存储器中每个工作项的行维度和列维度的全局标识符、工作组标识符和局部标识符；
[0014]步骤S6:设定输入矩阵数据个数与工作项个数相同，将多个工作项形成多个工作组；
[0015]步骤S7:对输入矩阵进行分块并得到矩阵块，以工作组为单位，将每个矩阵块与工作组相对应；
[0016]步骤S8:在每个工作组内部的图形处理器的本地存储器中申请一片本地内存；
[0017]步骤S9:根据每个工作项的行维度和列维度的全局标识符，以工作组为单位，将图形处理器的全局存储器中输入矩阵的数据复制到图形处理器的本地存储器；
[0018]步骤SlO:执行图形处理器的本地存储器的同步等待指令，等待同一工作组中的数据全部复制到图形处理器的本地存储器中，在本地存储器中获得同一工作组中的数据；
[0019]步骤Sll:利用行维度上工作组标识符、工作组的大小和局部标识符计算得到转置后的输出数据在全局存储器中的行索引；利用列维度上工作组标识符、工作组的大小和局部标识符计算得到转置后的输出数据在全局存储器中的列索引；
[0020]步骤S12:利用所述输出数据在全局存储器中的行索引、列索引和R行的大小计算得到输出数据在全局存储器中的位置；利用输出数据的行维度的局部标识符、列维度的局部标识符和矩阵块的大小计算得到输出数据在本地存储器中的位置；
[0021]步骤S13:将输出数据在本地存储器位置上的数据赋值给输出数据在全局存储器中的一维数组，获得一维数组中的输出矩阵数据，用以实现无冲突的全局存储器合并访问；
[0022]步骤S14:将一维数组中的输出矩阵数据拷贝到内存中，对输出矩阵数据按照S行R列的方式进行切分，得到转置后的矩阵，实现矩阵转置在图形处理器平台上的并行、优化。
[0023](三)本发明的有益效果:
[0024]本发明的目的在于优化矩阵转置算法的性能，使中央处理器与图形处理器协同计算。利用图形处理器核心内部的共享内存进行线程间的通信，大大提高了算法的效率。同时本发明-开放计算语言(OpenCL:Open Computing Language)架构中将矩阵进行分块处理，在访问全局内存时避免了非合并访问和分区冲突问题。提高了图形处理器的访存带宽。
[0025]1.本发明将数据按照分块的方式分配在数以千计的图形处理器上，有效的实现了数据的并行化处理，提高了程序的运算速度和运算精度。
[0026]2.本发明将多个线程的内存访问合并到较少的内存请求命令中，改变了普通的图形处理器上线程映射的方式，借助了本地存储器作为中间元素，实现了全局内存的合并访问，这样可以大大提高带宽利用率。
[0027]3.本发明将连续的本地存储器访问地址映射到不同的本地存储器上，使得在一个线程调度单位上，没有本地存储器访问冲突，否则一个线程调度单位上的访问冲突会等待所有的本地内存访问完成，造成访问时延。
【专利附图】
【附图说明】
[0028]图1为本发明基于图形处理器平台的矩阵转置并行优化算法的流程图。【具体实施方式】
[0029]为使本发明的目的、技术方案和优点更加清楚明白，以下结合具体实施例，并参照附图，对本发明进一步详细说明。
[0030]图1示出了本发明基于图形处理器的矩阵转置并行优化方法的流程图，实施例所采用的硬件平台是:华硕主板，显卡；软件平台是:微软操作系统，微软公司开发套件，但不限于此。
[0031]本发明基于图形处理器的矩阵转置优化包括如下步骤:
[0032]步骤S1:转置的矩阵为R行S列，每个矩阵元素为复数，即包括数的实部和虚部。由于在通用的图形处理器编程标准语言中没有对二维数组的直接操作函数，因此，需要将矩阵转换为一维数组。将输入矩阵存储在中央处理器的磁盘上并按行读入，生成具有R行S列的二维数组，将R行S列的二维数组转化为含有RXS个元素的一维数组,标识为input ；所述输入矩阵为R行S列，每个矩阵数据为复数，每个矩阵元素包括实部和虚部。
[0033]步骤S2:根据一维数组的大小和输入矩阵数据的类型，计算出一维数组所需要的存储空间。
[0034]步骤S3:在图形处理器的全局存储器上为一维数组分配大小相同的两片存储空间；将一维数组中的数据拷贝到图形处理器的全局存储器中的存储空间。所述两片存储空间包括第一片存储空间和第二片存储空间，第一片存储空间存储从内存拷贝到图形处理器的全局存储器中的输入矩阵数据；第二片存储空间存储输出矩阵数据。
[0035]步骤S4:根据图形处理器可视为三层模型:N维域、工作组和工作项，设置具有二维域的索引空间；我们定义索引空间为二维域。所述二维域的每一维度的大小由输入矩阵的R和S决定。定义工作组的维度为二维维度。
[0036]步骤S5:根据图形处理器内置函数分别对全局存储器中每个工作项的行维度和列维度进行计算，得到全局存储器中每个工作项的行维度和列维度的全局标识符、工作组标识符和局部标识符。
[0037]步骤S6:设定输入矩阵数据个数与工作项个数相同，将多个工作项形成多个工作组；
[0038]步骤S7:对输入矩阵进行分块划分并得到矩阵块。定义每个子块的大小为mXn。以工作组为单位，将每个矩阵块与工作组相对应，以工作组为单位进行计算，将得到的工作组数据存储在全局存储器中，通过图形处理器的内置函数分别计算出矩阵的行和列方向的全局标识符。
[0039]步骤S8:在每个工作组内部的图形处理器的本地存储器中申请一片本地内存；借助本地存储器作为中间元素，实现了内存的合并访问。
[0040]步骤S9:根据每个工作项的行维度和列维度的全局标识符，以工作组为单位，将图形处理器的全局存储器中输入矩阵的数据复制到图形处理器的本地存储器；
[0041]步骤SlO:执行图形处理器的本地存储器的同步等待指令，等待同一工作组中的数据全部复制到图形处理器的本地存储器中，在本地存储器中获得同一工作组中的数据；
[0042]步骤Sll:利用行维度上工作组标识符、工作组的大小和局部标识符计算得到转置后的输出数据在全局存储器中的行索引；利用列维度上工作组.标识符、工作组的大小和局部标识符计算得到转置后的输出数据在全局存储器中的列索引；[0043]步骤S12:利用所述输出数据在全局存储器中的行索引、列索引和R行的大小计算得到输出数据在全局存储器中的位置；利用输出数据的行维度的局部标识符、列维度的局部标识符和矩阵块的大小计算得到输出数据在本地存储器中的位置；
[0044]步骤S13:将输出数据在本地存储器位置上的数据赋值给输出数据在全局存储器中的一维数组，获得一维数组中的输出矩阵数据，用以实现无冲突的全局存储器合并访问；改变工作项的映射方式，使得在一个线程调度单位上没有本地存储器访问冲突。
[0045]步骤S14:将一维数组中的输出矩阵数据拷贝到内存中，对输出矩阵数据按照S行R列的方式进行切分，得到转置后的矩阵，实现矩阵转置在图形处理器平台上的并行、优化。
[0046]实施实例是以在雷达信号处理系统中128X4096点脉冲压缩算法中矩阵转置的应用为例。本发明的实现流程如下:
[0047]1.将雷达的回波数据文件中的数读取128次，每次读取4096X2=8192个数，将8192个数组成复数，从文件的第O个数开始读取，偶数位置的数作为数的实部，奇数位置的数作为数的虚部。由此形成了含有128X4096=524288个复数的输入矩阵。由于开放计算语言(OpenCL)中没有对二维数组的直接操作函数，因此，需要将输入矩阵转换为一维数组。对于含有524288个复数的输入矩阵，实现按行读入，将输入矩阵的数据读入到含有524288个复数的一维数组中。将一维数组作为输入数据并标示为:input[524288];所述开放计算语言(OpenCL)是在超微半导体公司(AMD:Advanced Micro Devices, Inc.)的显卡产品上能够实现应用程序硬件加速的并行编程开放标准语言。
[0048]2.计算一维数组所占用的存储空间的大小:输入矩阵的大小为128X4096，输入矩阵的数据类型为复数，实部和虚部均为浮点数。输入数据所占据的存储空间为:128X4096X数据类型。所述数据类型是复数的数据类型,复数的实部和虚部均为浮点类型。
[0049]3.在图形处理器的全局内存中，为一维数组申请同样大小的两片存储空间，并分别标示为:input_gpu和output_gpu。调用开放计算语言(OpenCL)数据拷贝函数，将输入数据从中央处理器内存拷贝到图形处理器全局内存的存储空间input_gpu中。所述图形处理器的全局内存是指开放计算语言(OpenCL)的一种存储器，图形处理器中的任何数据对全局存储器有读，写权限。
[0050]4.数据的分块处理:将工作项和工作组的索引空间设为二维域，设置工作项的行维度大小等于矩阵的行数，列维度的大小等于矩阵的列数，即二维域的空间大小为128X4096。所述工作项是在开放计算语言(OpenCL)中，每个实际执行计算的线程称为一个工作项，开放计算语言(OpenCL)为所有工作项分配了一个索引空间。每个工作项在该索引空间中都有一个唯一的全局标识符。工作项被进一步组织成工作组。每个工作组被分配一个工作组标识符。位于同一个工作组的工作项在工作组内部都有一个唯一的局部标识符。
[0051]5.输入数据的索引号:首先，调用开放计算语言(OpenCL)内置函数get_global_id (O)和get_global_id(l),该函数的作用是:获得当前工作项的全局标识符。其次，调用开放计算语言(OpenCL)内置函数get_group_id(0)和get_group_id (I),该函数的作用是:获得当前工作项的工作组标识符。最后，调用开放计算语言(OpenCL)内置函数get_loCal_id (O)和get_local_id(l),该函数的作用是:获得当前工作项的局部标识符。[0052]6.设置工作项的个数等于输入矩阵数据的个数，即为524288。设置工作组的行维度大小为8，列维度大小为16。将8X16个工作项形成工作组。所述工作组的行和列维度的划分采用的是8X16的分割方式，这与显卡的硬件设备特性有关。根据本实施实例选用的显卡具有1408个处理器核心，共有22个计算单元数，因此每个计算单元内部拥有64个流处理器。因此，矩阵的分块方式应为64的整数倍，此实例选用8X 16=128，即为64的2倍。
[0053]7.对输入矩阵进行分块划分，每个矩阵块的大小等于工作组的大小，即为8X16。
[0054]8.在每个工作组内部的图形处理器的本地存储器中申请一片本地内存，本地内存的大小等于工作组的大小，即为8X16X数据类型。所述图形处理器的本地存储器是指它位于每一个多处理器的内部，它只能被工作组内部的单个线程访问。
[0055]9.根据第5步中得到的工作项的标识符，以工作组为单位，将以8X16分块的工作组中的输入数据存入每个工作组内部的本地存储器中。标示:block[get_local_id(I)X 8+get_local_id(O)]
[0056]=input_gpu[get_global_id(I)X128+get_global_id(O)]；
[0057]10.调用开放计算语言(OpenCL)同步函数，实现图形处理器的本地存储器的同步
等待;
[0058]11.转置后输出数据索引号:首先，根据当前工作项的列维度的工作组标识符get_group_id(l),工作项的列维度的的局部标识符get_local_id(I)和工作组列维度的大小16,来计算转置后输出矩阵的行索引:targetGIdx=get_group_id(l) X 16+get_local_id(l);其次,根据当前工作项的行维度的工作组标识符get_group_id(O),工作项的行维度的的局部标识符get_loCal_id(0)和工作组行维度的大小8，来计算转置后输出矩阵的列索引:targetGIdy=get_group_id(0) X 8+get_local_id (O);所述转置是指将 R 行S列的矩阵的行和列互换，得到的S行R列的矩阵即为转置后的矩阵。
[0059]12.首先，根据第11步中计算出的转置后输出矩阵的行索引targetGIdx，列索引targetGIdy和数据矩阵的列数R=4096计算出输出数据在全局存储器中的位置:targetldx=targetGIdyX4096+targetGIdx ;其次，根据第5步中计算出的工作项的行维度的局部标识符get_local_id (O),列维度的局部标识符get_local_id (I)和工作组的行维度大小8来计算出数据在本地存储器中的位置:sourceIdx=get_local_id(l) X8+get_local_id(O)。
[0060]13.将输出数据在本地存储器位置上的数据block[sourceIdx]赋值给输出数据在全局存储器中的一维数组output_gpu[targetIdx],以实现无冲突的全局存储器合并访问。所述全局存储器合并访问是将一个工作组内的工作项视为线程，每个线程对全局存储器的访问需要一次内存请求来完成。将多个线程对全局存储器的访问合并到一次内存请求中，即实现了全局存储器的合并访问。所述无冲突的全局存储器合并访问是通过改变线程的映射方式，用本地存储器作为中间元素，实现了输入矩阵和输出矩阵都是全局存储器合并访问。
[0061]14.调用开放计算语言(OpenCL)数据拷贝函数，将全局存储器中的一维数组out_gpu[targetldx]拷贝到内存output [524288]中，对拷贝到内存中的数据output [524288]按照S行R列的方式进行切分，得到转置后的矩阵。将转置后的矩阵输出到雷达信号处理系统的脉冲压缩算法中，最终得到脉冲压缩后矩阵的数值。
[0062]通过上述方法实现了雷达信号处理系统中的矩阵并行转置的优化。[0063]上述方法的实现过程所采用的是开放计算语言(OpenCL)，它具有平台无关性的特点以及真正意义上的异构计算资源解决方案。它能够支持不同的硬件和操作系统，对于操作系统的平台无关可以在软件上实现和维护。因此，本发明所提供的解决方案可以不受操作系统和平台的局限性限制，可实现跨平台运行。
[0064]此过程以工作组为单位，同时对矩阵进行了分块划分，将工作组与矩阵块进行一一对应。矩阵分块采用的是与显卡的硬件设备特性有关的分割方式，最终在显卡的多个处理器核心上实现了并行，高效的矩阵转置算法。
[0065]上述方法利用到了工作组内部的本地缓存实现了数据的合并访问以及减少了存储器访问冲突，因此提高了图形处理器内部数据的访问带宽以及减小了数据的访问时延。快速、高效的实现了在图形处理器上对矩阵进行并行、优化的转置算法。
[0066]以上所述，仅为本发明中的【具体实施方式】，但本发明的保护范围并不局限于此，任何熟悉该技术的人在本发明所揭露的技术范围内，可理解想到的变换或替换，都应涵盖在本发明的包含范围之内。
【权利要求】
1.一种基于图形处理器的矩阵转置优化方法，其特征在于，所述方法包括如下步骤: 步骤S1:将输入矩阵存储在中央处理器的磁盘上并按行读入，生成具有R行S列的二维数组，将R行S列的二维数组转化为含有RXS个元素的一维数组； 步骤S2:根据一维数组的大小和输入矩阵的数据类型，计算出一维数组所需要的存储空间； 步骤S3:在图形处理器的全局存储器上为一维数组分配大小相同的两片存储空间；将一维数组中的数据拷贝到图形处理器的全局存储器中的存储空间； 步骤S4:根据图形处理器具有的N维域模型、工作组和多个工作项，设置具有二维域的索引空间； 步骤S5:根据图形处理器内置函数分别对全局存储器中每个工作项的行维度和列维度进行计算，得到全局存储器中每个工作项的行维度和列维度的全局标识符、工作组标识符和局部标识符； 步骤S6:设定输入矩阵数据个数与工作项个数相同，将多个工作项形成多个工作组；步骤S7:对输入矩阵进行分块并得到矩阵块，以工作组为单位，将每个矩阵块与工作组相对应； 步骤S8:在每个工作组内部的图形处理器的本地存储器中申请一片本地内存； 步骤S9:根据每个工作项的行维度和列维度的全局标识符，以工作组为单位，将图形处理器的全局存储器中输入矩 阵的数据复制到图形处理器的本地存储器； 步骤SlO:执行图形处理器的本地存储器的同步等待指令，等待同一工作组中的数据全部复制到图形处理器的本地存储器中，在本地存储器中获得同一工作组中的数据； 步骤Sll:利用行维度上工作组标识符、工作组的大小和局部标识符计算得到转置后的输出数据在全局存储器中的行索引；利用列维度上工作组标识符、工作组的大小和局部标识符计算得到转置后的输出数据在全局存储器中的列索引； 步骤S12:利用所述输出数据在全局存储器中的行索引、列索引和R行的大小计算得到输出数据在全局存储器中的位置；利用输出数据的行维度的局部标识符、列维度的局部标识符和矩阵块的大小计算得到输出数据在本地存储器中的位置； 步骤S13:将输出数据在本地存储器位置上的数据赋值给输出数据在全局存储器中的一维数组，获得一维数组中的输出矩阵数据，用以实现无冲突的全局存储器合并访问； 步骤S14:将一维数组中的输出矩阵数据拷贝到内存中，对输出矩阵数据按照S行R列的方式进行切分，得到转置后的矩阵，实现矩阵转置在图形处理器平台上的并行、优化。
2.根据权利要求1所述矩阵转置优化方法，其特征在于，所述输入矩阵为R行S列，每个矩阵数据为复数，每个矩阵元素包括实部和虚部。
3.根据权利要求1所述矩阵转置优化方法，其特征在于，所述二维域的每一维度的大小由输入矩阵的R行S列决定。
4.根据权利要求1所述矩阵转置优化方法，其特征在于，所述本地内存的大小与工作组的大小一致。
5.根据权利要求1所述矩阵转置优化方法，其特征在于，所述两片存储空间包括第一片存储空间和第二片存储空间，第一片存储空间存储从内存拷贝到图形处理器的全局存储器中的输入矩阵数据；第二片存储空间存储输出矩阵数据。
6.根据权利要求1所述矩阵转置优化方法，其特征在于，所述工作组的维度为二维维度。
7.根据权利要求1所述矩阵转置优化方法，其特征在于，借助本地存储器作为中间元素，实现了内存的合并访问。
8.根据权利要求1所述矩阵转置优化方法，其特征在于，改变工作项的映射方式，使得在一个线程调度 单位上没有本地存储器访问冲突。
【文档编号】G06F9/38GK103761215SQ201410017822
【公开日】2014年4月30日 申请日期:2014年1月15日 优先权日:2014年1月15日
【发明者】田卓, 樊双丽 申请人:北京新松佳和电子系统股份有限公司