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

基于gpu众核平台的矩阵并行转置方法 Download PDF

Info

Publication number
CN101937425A
CN101937425A CN2009100883723A CN200910088372A CN101937425A CN 101937425 A CN101937425 A CN 101937425A CN 2009100883723 A CN2009100883723 A CN 2009100883723A CN 200910088372 A CN200910088372 A CN 200910088372A CN 101937425 A CN101937425 A CN 101937425A
Authority
CN
China
Prior art keywords
thread block
index
matrix
gpu
dimension
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.)
Granted
Application number
CN2009100883723A
Other languages
English (en)
Other versions
CN101937425B (zh
Inventor
龙腾
刘海波
姚迪
靳星星
刘峰
Current Assignee (The listed assignees may be inaccurate. Google has not performed a legal analysis and makes no representation or warranty as to the accuracy of the list.)
Beijing Institute of Technology BIT
Original Assignee
Beijing Institute of Technology BIT
Priority date (The priority date 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 date listed.)
Filing date
Publication date
Application filed by Beijing Institute of Technology BIT filed Critical Beijing Institute of Technology BIT
Priority to CN2009100883723A priority Critical patent/CN101937425B/zh
Publication of CN101937425A publication Critical patent/CN101937425A/zh
Application granted granted Critical
Publication of CN101937425B publication Critical patent/CN101937425B/zh
Expired - Fee Related legal-status Critical Current
Anticipated expiration legal-status Critical

Links

Images

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 (3)

1.基于众核平台的矩阵转置并行方法，其特征在于：包括以下步骤：
第一步：转置的矩阵为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众核平台下的实现。
2.根据权利要求1所述的一种基于GPU众核平台的矩阵转置并行方法，其特征在于：在对数据进行分割时，对于线程块层设置为二维或二维以上且每维的大小一致。
3.根据权利要求1所述的一种基于GPU众核平台的矩阵转置并行方法，其特征在于：线层网格层大小设置为二维，且每维的大小由原矩阵的大小M、N决定。
CN2009100883723A 2009-07-02 2009-07-02 基于gpu众核平台的矩阵并行转置方法 Expired - Fee Related CN101937425B (zh)

Priority Applications (1)

Application Number Priority Date Filing Date Title
CN2009100883723A CN101937425B (zh) 2009-07-02 2009-07-02 基于gpu众核平台的矩阵并行转置方法

Applications Claiming Priority (1)

Application Number Priority Date Filing Date Title
CN2009100883723A CN101937425B (zh) 2009-07-02 2009-07-02 基于gpu众核平台的矩阵并行转置方法

Publications (2)

Publication Number Publication Date
CN101937425A true CN101937425A (zh) 2011-01-05
CN101937425B CN101937425B (zh) 2012-05-30

Family

ID=43390759

Family Applications (1)

Application Number Title Priority Date Filing Date
CN2009100883723A Expired - Fee Related CN101937425B (zh) 2009-07-02 2009-07-02 基于gpu众核平台的矩阵并行转置方法

Country Status (1)

Country Link
CN (1) CN101937425B (zh)

Cited By (9)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN102567283A (zh) * 2011-12-08 2012-07-11 清华大学 一种利用gpu对小矩阵求逆的方法
CN102708009A (zh) * 2012-04-19 2012-10-03 华为技术有限公司 一种基于cuda实现多任务共享gpu的方法
CN102881042A (zh) * 2012-09-05 2013-01-16 浪潮(北京)电子信息产业有限公司 电镜三维图像重构的方法及系统
WO2013159272A1 (en) * 2012-04-23 2013-10-31 Hewlett-Packard Development Company Statistical analysis using graphics processing unit
CN103761215A (zh) * 2014-01-15 2014-04-30 北京新松佳和电子系统股份有限公司 基于图形处理器的矩阵转置优化方法
CN106528054A (zh) * 2016-11-03 2017-03-22 东南大学 Gpu加速的稠密向量加法计算方法
CN106556873A (zh) * 2016-10-31 2017-04-05 华讯方舟科技有限公司 一种基于人体微波成像的安检方法及系统
CN106844022A (zh) * 2016-12-23 2017-06-13 中国石油天然气集团公司 一种数据处理的方法及系统
CN109471612A (zh) * 2018-09-18 2019-03-15 北京中科寒武纪科技有限公司 运算装置及方法

Family Cites Families (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN1889128A (zh) * 2006-07-17 2007-01-03 北京航空航天大学 基于gpu的预计算辐射度传递全频阴影的方法

Cited By (17)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN102567283A (zh) * 2011-12-08 2012-07-11 清华大学 一种利用gpu对小矩阵求逆的方法
CN102567283B (zh) * 2011-12-08 2014-12-31 清华大学 一种利用gpu对小矩阵求逆的方法
CN102708009B (zh) * 2012-04-19 2014-04-02 华为技术有限公司 一种基于cuda实现多任务共享gpu的方法
CN102708009A (zh) * 2012-04-19 2012-10-03 华为技术有限公司 一种基于cuda实现多任务共享gpu的方法
CN104662531A (zh) * 2012-04-23 2015-05-27 惠普发展公司，有限责任合伙企业 使用图形处理单元的统计分析
WO2013159272A1 (en) * 2012-04-23 2013-10-31 Hewlett-Packard Development Company Statistical analysis using graphics processing unit
GB2516192A (en) * 2012-04-23 2015-01-14 Hewlett Packard Development Co Statistical Analysis Using Graphics Processing Unit
CN102881042A (zh) * 2012-09-05 2013-01-16 浪潮(北京)电子信息产业有限公司 电镜三维图像重构的方法及系统
CN102881042B (zh) * 2012-09-05 2015-09-23 浪潮(北京)电子信息产业有限公司 电镜三维图像重构的方法及系统
CN103761215A (zh) * 2014-01-15 2014-04-30 北京新松佳和电子系统股份有限公司 基于图形处理器的矩阵转置优化方法
CN103761215B (zh) * 2014-01-15 2016-08-24 北京新松佳和电子系统股份有限公司 基于图形处理器的矩阵转置优化方法
CN106556873A (zh) * 2016-10-31 2017-04-05 华讯方舟科技有限公司 一种基于人体微波成像的安检方法及系统
WO2018076885A1 (zh) * 2016-10-31 2018-05-03 华讯方舟科技有限公司 一种基于人体微波成像的安检方法及系统
US11529067B2 (en) 2016-10-31 2022-12-20 China Communication Technology Co., Ltd. Security check method and system based on microwave imaging of human bodies
CN106528054A (zh) * 2016-11-03 2017-03-22 东南大学 Gpu加速的稠密向量加法计算方法
CN106844022A (zh) * 2016-12-23 2017-06-13 中国石油天然气集团公司 一种数据处理的方法及系统
CN109471612A (zh) * 2018-09-18 2019-03-15 北京中科寒武纪科技有限公司 运算装置及方法

Also Published As

Publication number Publication date
CN101937425B (zh) 2012-05-30

Similar Documents

Publication Publication Date Title
CN101937425B (zh) 基于gpu众核平台的矩阵并行转置方法
CN104636273B (zh) 一种带多级Cache的SIMD众核处理器上的稀疏矩阵存储方法
CN103049241B (zh) 一种提高cpu+gpu异构装置计算性能的方法
CN104572295B (zh) 匹配于高性能计算机体系结构的结构网格数据管理方法
CN103617150A (zh) 一种基于gpu的大规模电力系统潮流并行计算系统及其方法
CN103440163B (zh) 使用gpu并行实现的基于pic模型的加速器仿真方法
CN103761215B (zh) 基于图形处理器的矩阵转置优化方法
Yu et al. A parallel visualization pipeline for terascale earthquake simulations
CN101937555B (zh) 基于gpu众核平台的脉冲压缩参考矩阵并行生成方法
CN103345580A (zh) 基于格子Boltzmann方法的并行CFD方法
CN107341133A (zh) 基于任意维数矩阵lu分解的可重构计算结构的调度方法
CN109978171A (zh) 一种基于云计算的Grover量子仿真算法优化方法
Carr et al. Scalable contour tree computation by data parallel peak pruning
Zhang et al. Gpu-based implementation of finite element method for elasticity using cuda
Liu et al. Parallel reconstruction of neighbor-joining trees for large multiple sequence alignments using CUDA
CN112947870B (zh) 一种3D打印模型的G-code并行生成方法
Shi et al. Design of MILC lattice QCD application for GPU clusters
CN106484532B (zh) 面向sph流体模拟的gpgpu并行计算方法
Xu et al. Balancing cpu-gpu collaborative high-order cfd simulations on the tianhe-1a supercomputer
Deorowicz Solving longest common subsequence and related problems on graphical processing units
Capozzoli et al. The success of GPU computing in applied electromagnetics
Valero-Lara et al. LBM-HPC-an open-source tool for fluid simulations. case study: Unified parallel C (UPC-PGAS)
CN101937422B (zh) 基于gpu众核平台的fft并行方法
Zou et al. Direction-optimizing breadth-first search on CPU-GPU heterogeneous platforms
Gan et al. Million-core-scalable simulation of the elastic migration algorithm on Sunway TaihuLight supercomputer

Legal Events

Date Code Title Description
C06 Publication
PB01 Publication
C10 Entry into substantive examination
SE01 Entry into force of request for substantive examination
C14 Grant of patent or utility model
GR01 Patent grant
CF01 Termination of patent right due to non-payment of annual fee

Granted publication date: 20120530

Termination date: 20150702

EXPY Termination of patent right or utility model