CN112837205B - 一种图形处理器上基于延迟修正的批量矩阵求逆方法 - Google Patents

一种图形处理器上基于延迟修正的批量矩阵求逆方法 Download PDF

Info

Publication number
CN112837205B
CN112837205B CN202110247100.4A CN202110247100A CN112837205B CN 112837205 B CN112837205 B CN 112837205B CN 202110247100 A CN202110247100 A CN 202110247100A CN 112837205 B CN112837205 B CN 112837205B
Authority
CN
China
Prior art keywords
matrix
inversion
graphics processor
block
batch
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.)
Active
Application number
CN202110247100.4A
Other languages
English (en)
Other versions
CN112837205A (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.)
Computer Network Information Center of CAS
Original Assignee
Computer Network Information Center of CAS
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 Computer Network Information Center of CAS filed Critical Computer Network Information Center of CAS
Priority to CN202110247100.4A priority Critical patent/CN112837205B/zh
Publication of CN112837205A publication Critical patent/CN112837205A/zh
Application granted granted Critical
Publication of CN112837205B publication Critical patent/CN112837205B/zh
Active legal-status Critical Current
Anticipated expiration legal-status Critical

Links

Images

Classifications

    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06T—IMAGE DATA PROCESSING OR GENERATION, IN GENERAL
    • G06T1/00—General purpose image data processing
    • G06T1/20—Processor architectures; Processor configuration, e.g. pipelining
    • 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/50—Allocation of resources, e.g. of the central processing unit [CPU]
    • G06F9/5005—Allocation of resources, e.g. of the central processing unit [CPU] to service a request
    • G06F9/5011—Allocation of resources, e.g. of the central processing unit [CPU] to service a request the resources being hardware resources other than CPUs, Servers and Terminals
    • G06F9/5016—Allocation of resources, e.g. of the central processing unit [CPU] to service a request the resources being hardware resources other than CPUs, Servers and Terminals the resource being the memory
    • 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)
  • Theoretical Computer Science (AREA)
  • Software Systems (AREA)
  • Physics & Mathematics (AREA)
  • General Physics & Mathematics (AREA)
  • General Engineering & Computer Science (AREA)
  • Image Processing (AREA)

Abstract

本发明涉及一种图形处理器上基于延迟修正的批量矩阵求逆方法，该方法包括以下步骤：产生需要批量求逆的矩阵数据；依次将所述矩阵数据从主机内传输到图形处理器全局内存；在所述图形处理器上建立Grid列条和Block块与矩阵的对应关系；所述图形处理器按照延迟修正方式完成求逆矩阵的批量求逆，并将批量求逆后，得到的逆矩阵从图形处理器的全局内存传会主机内存。本发明可以减少对全局内存的访问和加快数据的读写速度，减少访存次数的列交换方法，相比静态分配方法，性能得到明显提升。

Description

一种图形处理器上基于延迟修正的批量矩阵求逆方法
技术领域
本发明属于基于图形处理器的批量矩阵求逆技术，具体涉及一种图形处理器上基于延迟修正的批量矩阵求逆方法。
背景技术
将大量的复杂重复的任务交给图形处理器来处理，从而可以大大加快了数据处理的速度。相比于CPU处理器，图形处理器由于存在大量线程使得对批量任务的处理有着巨大的优势。批量矩阵求逆问题广泛出现于机器学习、数据挖掘、图像和信号处理中，由于需要处理的矩阵规模较小但数量巨大，这给图形处理器上批量处理技术带来巨大挑战。
目前的求逆技术采用了一种称为及时修正的技术，该技术将全局内存中需要处理的当前列条块数据读到寄存器，然后进行计算处理，处理完成后再写回到全局内存，及时修正技术存在着较大的全局内存的数据读写量，而对于图像处理器上批量小矩阵求逆问题，影响批处理性能的主要因素是数据的读写量，而不是数据处理时间。目前及时修正批处理求逆技术方案如下：
在采用及时修正技术时，如图1所示，首先从全局内存中将当前列条块
Figure BDA0002964485790000011
读到寄存器，然后进行当前列条块内的计算，计算完成后再将列条块写回到全局内存。这时对全局内存的数据读写量为2×(i×NB)×NB。然后与U12来修正列条块右边的部分
Figure BDA0002964485790000012
此时的对全局内存中矩阵W的数据读写量为2×(i×NB)×(n-i×NB)。整个修正尾部矩阵的过程共对矩阵U的上三角部分进行了一次访问，对全局内存中的矩阵U的数据读取量约
Figure BDA0002964485790000013
则整个及时修正的求U-1的块算法对全局内存的读写量约为：
Figure BDA0002964485790000014
为了减少批量处理过程中及时修正技术存在的全局内存的读写次数多和数据量大问题，我们发明了延迟修正技术。该技术相比及时修正的技术，将数据读写数量减少了50％，从而图形处理器的批量处理能力提升了1倍左右。同时，为了加快图形处理器对全局内存的数据读写速度，采用了更多利用寄存器和共享内存的优化方法和减少访存次数.并且为了避免线程的闲置和shared meomery等资源浪费，提出了按需申请资源的动态方法，相较于一次性分配资源的静态法性能得到明显提升
发明内容
本发明的目的在于，减少批量处理过程中及时修正技术存在的全局内存的读写次数多和数据量大问题。
为实现上述目的，本发明提供了一种图形处理器上基于延迟修正的批量矩阵求逆方法，该方法包括以下步骤：
产生需要批量求逆的矩阵数据；
依次将所述矩阵数据从主机内传输到图形处理器全局内存；
在所述图形处理器上建立Grid列条和Block块与矩阵的对应关系；
所述图形处理器按照延迟修正方式完成求逆矩阵的批量求逆，并将批量求逆后，得到的逆矩阵从图形处理器的全局内存传会主机内存。
优选地，延迟修正方式步骤，包括：
初始化寄存器中的当前列条块；
将上三角矩阵U的子矩阵U01和U11从全局内存中读到共享内存；
在求解当前第i个列条块W01之前，首先用当前列条块左边已经求解出前i-1个列条块W00的结果与上三角矩阵U的子矩阵U01来修正当前第i个列条块W01，修正完成后再求解当前列条块，然后再向右进行下一个列条块的求解，直到最后一个列条块，求解结果存入寄存器中；
将寄存器中的列条块的求解结果传回全局内存。
优选地，对于Block内线程的配置，采用1-D结构；1-D结构中一个线程对应矩阵的一行。
优选地，在图形处理器内实现LU分解之后的批量求逆算法时，采取了Thread-Block级并行块方法，由于一个线程对应矩阵的一行，并且列条块内每行的求解相互独立，因此可将当前列条块存储在线程的私有寄存器中。
优选地，求逆块算法包括列条块求逆和尾部矩阵修正两个阶段，求逆算法采用了更多利用寄存器和共享内存的优化方法，并设计了减少访存数据量的列交换方法。
优选地，采用一次性分配GPU资源造成的线程闲置和共享内存，设计了运行时按需分配线程、共享内存图形处理器资源的动态资源分配方法。
本发明相比及时修正的技术，可将数据读写数量减少了50％，从而图形处理器的批量处理能力提升了1倍左右。同时，为了加快图形处理器对全局内存的数据读写速度，采用了更多利用寄存器和共享内存的优化方法和减少访存次数。并且为了避免线程的闲置和shared meomery等资源浪费，提出了按需申请资源的动态方法，相较于一次性分配资源的静态法性能得到明显提升。
附图说明
图1为现有技术图形处理器上及时求逆流程示意图；
图2为图形处理器上延时修正批量求逆技术流程图；
图3为本发明实施例提供的一种图形处理器上基于延迟修正的批量矩阵求逆方法流程示意图；
图4为图形处理器上Grid和Block与批量矩阵的对应关系示意图；
图5为在图形处理器及时修正(a)和延时修正(b)求逆流程示意图。
具体实施方式
下面结合附图与具体实施方式对本发明作进一步详细描述。
图3为本发明实施例提供的一种图形处理器上基于延迟修正的批量矩阵求逆方法流程示意图。如图1所示，该方法包括步骤S101-S104:
步骤S101,产生需要批量求逆的矩阵数据。
步骤S102,依次将所述矩阵数据从主机内传输到图形处理器全局内存。
步骤S103,在所述图形处理器上建立Grid列条和Block块与矩阵的对应关系。
具体地，在图形处理器上实现批量矩阵求逆算法时，如图4所示一个Block对应一个矩阵。而对于Block内线程的配置，可以是1-D的也可以是2-D的。由于2-D线程结构需要使用n×n个线程对应矩阵的n×n个元素，限制了CUDA运行时每个SM(StreamingMultiprocessor)同时调度多个Block的能力。另外对于大多数矩阵规模而言，采用2-D结构的线程个数通常会超出每个Warp的线程数目(32)，这会导致在线程需要共享数据时会2-D比1-D结构更多必要的同步点。并且，当采用2-D结构对当前列条块进行处理时，由于一个线程对应矩阵的一个元素，这会造成条块外的线程处于空闲状态，导致线程利用率较低。另外，1-D结构中一个线程对应矩阵的一行，这使得每个线程被分配更多的任务，指令级并行性更好，因此本发明实施例采用了如图4所示的1-D结构。假设批量矩阵已经存储于图形处理器的全局内存上，并且矩阵元素是按照列主序存储，共有bs个矩阵。
步骤S104,所述图形处理器按照延迟修正方式完成求逆矩阵的批量求逆，并将批量求逆后，得到的逆矩阵从图形处理器的全局内存传会主机内存。
具体地，在图形处理器上实现批量LU分解算法时，采用了一个Block对应一个矩阵的Thread-Block级并行方法。通过使用块算法，只需将要处理的当前列条块放在共享内存中，而不是将整个矩阵放在共享内存中，这样处理的矩阵规模就不受共享内存大小的限制。在图形处理器内实现LU分解之后的批量求逆算法时，同样采取了Thread-Block级并行块方法.由于一个线程对应矩阵的一行，并且列条块内每行的求解相互独立，因此可将当前列条块存储在线程的私有寄存器中，加快数据的读写速度。
LU分解之后的求逆块算法包括列条块求逆和尾部矩阵修正两个阶段。为了减少修正过程对全局内存的读写数据量，本发明实施例在批量求逆算法的图形处理器实现中提出了延迟修正的矩阵求逆块算法。此外，为了减少对全局内存的访问和加快数据的读写速度，求逆算法采用了更多利用寄存器和共享内存的优化方法，并设计了减少访存数据量的列交换方法。另外，为了避免采用一次性分配GPU资源造成的线程闲置和共享内存浪费问题，本文设计了运行时按需分配线程、共享内存等GPU资源的动态资源分配方法，相比一次性分配的静态资源分配方法，性能得到了一定的提升。
图像处理器内延迟修正的求逆算法如下：
1for i＝1 to bn
2/*初始化寄存器中的当前列条块*/
3WR,01＝0,WR,11＝I；
4__syncthreads()；
5/*将U01和U11从全局内存读到共享内存*/
6
Figure BDA0002964485790000031
7__syncthreads()；
8/*修当前列条块*/
9WR，01-＝W00×US，01；
10__syncthreads()；
11/*求解列条块*/
12
Figure BDA0002964485790000041
13__syncthreads()；
14/*将列条块的求解结果放回到全局内存*/
15
Figure BDA0002964485790000042
16__syncthreads()；
17endfor
求逆块算法将矩阵分成一系列大小NB的列条块，共分成bn块(如图5所示)，然后逐次对这些列条块进行求解。延迟修正过程主要包括以下步骤：
初始化寄存器中的当前列条块，将0矩阵赋给寄存器中的变量WR,01,将单位矩阵赋给寄存器中的变量WR,11。
将上三角矩阵U的子矩阵U01和U11从全局内存中读到共享内存；
在求解当前第i个列条块W01之前，首先用当前列条块左边已经求解出前i-1个列条块W00的结果与上三角矩阵U的子矩阵U01来修正当前第i个列条块W01，修正完成后再求解当前列条块，然后再向右进行下一个列条块的求解，直到最后一个列条块，求解结果存入寄存器
Figure BDA0002964485790000043
中；
将寄存器中的列条块的求解结果传回全局内存
Figure BDA0002964485790000044
在条块内求解时，由于数据W中的当前条块中的数据需要多次读写，并且每个线程对应的每行数据的求解、修正相互独立，因此可以将其从全局内存读到线程私有的寄存器中，加快数据的读写速度，求解完成后再写回到全局内存即可。对于使用到的数据U和数据L中行/列条块内数据，由于对其进行了多次读取操作，因此可以将其从全局内存取到共享内存中，加快数据的读取速度。在算法中，后缀“R”表示存储在寄存器中的变量，后缀“S”表示存储在共享内存中的变量。
本发明实施例相比及时修正的技术，可将数据读写数量减少了50％，从而图形处理器的批量处理能力提升了1倍左右。同时，为了加快图形处理器对全局内存的数据读写速度，采用了更多利用寄存器和共享内存的优化方法和减少访存次数。并且为了避免线程的闲置和共享内存等资源浪费，提出了按需申请资源的动态方法，相较于一次性分配资源的静态法性能得到明显提升。
显而易见，在不偏离本发明的真实精神和范围的前提下，在此描述的本发明可以有许多变化。因此，所有对于本领域技术人员来说显而易见的改变，都应包括在本权利要求书所涵盖的范围之内。本发明所要求保护的范围仅由所述的权利要求书进行限定。

Claims (3)

1.一种图形处理器上基于延迟修正的批量矩阵求逆方法，其特征在于，包括以下步骤：
产生需要批量求逆的矩阵数据；
依次将所述矩阵数据从主机内存传输到图形处理器全局内存；
在所述图形处理器上建立Grid列条和Block块与矩阵的对应关系；对于Block内线程的配置，采用1-D结构；1-D结构中一个线程对应矩阵的一行；
所述图形处理器按照延迟修正方式完成求逆矩阵的批量求逆，并将批量求逆后，得到的逆矩阵从图形处理器的全局内存传会主机内存；
所述延迟修正方式步骤，包括：
初始化寄存器中的当前列条块；
将上三角矩阵U的子矩阵U01和U11从全局内存中读到共享内存；
在求解当前第i个列条块W01之前，首先用当前列条块左边已经求解出前i-1个列条块W00的结果与上三角矩阵U的子矩阵U01来修正当前第i个列条块W01，修正完成后再求解当前列条块，然后再向右进行下一个列条块的求解，直到最后一个列条块，求解结果存入寄存器中；
将寄存器中的列条块的求解结果传回全局内存；
在图形处理器内实现LU分解之后的批量求逆算法时，采取Thread-Block级并行块方法，由于一个线程对应矩阵的一行，并且列条块内每行的求解相互独立，因此可将当前列条块存储在线程的私有寄存器中。
2.根据权利要求1所述的方法，其特征在于，求逆块算法包括列条块求逆和尾部矩阵修正两个阶段，求逆算法采用了更多利用寄存器和共享内存的优化方法，并设计了减少访存数据量的列交换方法。
3.根据权利要求1所述的方法，其特征在于，采用一次性分配GPU资源造成的线程闲置和共享内存，设计了运行时按需分配线程、共享内存图形处理器资源的动态资源分配方法。
CN202110247100.4A 2021-03-05 2021-03-05 一种图形处理器上基于延迟修正的批量矩阵求逆方法 Active CN112837205B (zh)

Priority Applications (1)

Application Number Priority Date Filing Date Title
CN202110247100.4A CN112837205B (zh) 2021-03-05 2021-03-05 一种图形处理器上基于延迟修正的批量矩阵求逆方法

Applications Claiming Priority (1)

Application Number Priority Date Filing Date Title
CN202110247100.4A CN112837205B (zh) 2021-03-05 2021-03-05 一种图形处理器上基于延迟修正的批量矩阵求逆方法

Publications (2)

Publication Number Publication Date
CN112837205A CN112837205A (zh) 2021-05-25
CN112837205B true CN112837205B (zh) 2022-07-26

Family

ID=75934673

Family Applications (1)

Application Number Title Priority Date Filing Date
CN202110247100.4A Active CN112837205B (zh) 2021-03-05 2021-03-05 一种图形处理器上基于延迟修正的批量矩阵求逆方法

Country Status (1)

Country Link
CN (1) CN112837205B (zh)

Citations (4)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN102567283A (zh) * 2011-12-08 2012-07-11 清华大学 一种利用gpu对小矩阵求逆的方法
CN104182209A (zh) * 2014-08-27 2014-12-03 中国科学院软件研究所 一种基于PETSc的GCRO-DR算法并行处理方法
CN104572588A (zh) * 2014-12-23 2015-04-29 中国电子科技集团公司第三十八研究所 矩阵求逆处理方法和装置
WO2020081543A1 (en) * 2018-10-16 2020-04-23 The Broad Institute, Inc. Methods of scaling computational genomics with specialized architectures for highly parallelized computations and uses thereof

Patent Citations (4)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN102567283A (zh) * 2011-12-08 2012-07-11 清华大学 一种利用gpu对小矩阵求逆的方法
CN104182209A (zh) * 2014-08-27 2014-12-03 中国科学院软件研究所 一种基于PETSc的GCRO-DR算法并行处理方法
CN104572588A (zh) * 2014-12-23 2015-04-29 中国电子科技集团公司第三十八研究所 矩阵求逆处理方法和装置
WO2020081543A1 (en) * 2018-10-16 2020-04-23 The Broad Institute, Inc. Methods of scaling computational genomics with specialized architectures for highly parallelized computations and uses thereof

Non-Patent Citations (2)

* Cited by examiner, † Cited by third party
Title
"并行对称矩阵三对角化算法在GPU集群上的有效实现";刘世芳， 赵永华， 于天禹， 黄荣锋;《计算机研究与发展》;20201231;第57卷(第12期);第2637页第2段、最后一段，第2638页第3段 *
基于GPU并行处理的大规模连续潮流批量计算;刘正元等;《电网技术》;20191111(第03期);全文 *

Also Published As

Publication number Publication date
CN112837205A (zh) 2021-05-25

Similar Documents

Publication Publication Date Title
US10255547B2 (en) Indirectly accessing sample data to perform multi-convolution operations in a parallel processing system
US7912889B1 (en) Mapping the threads of a CTA to the elements of a tile for efficient matrix multiplication
US7792895B1 (en) Efficient matrix multiplication on a parallel processing device
US9921847B2 (en) Tree-based thread management
CN108733415B (zh) 支持向量随机访存的方法及装置
US8751771B2 (en) Efficient implementation of arrays of structures on SIMT and SIMD architectures
US7506134B1 (en) Hardware resource based mapping of cooperative thread arrays (CTA) to result matrix tiles for efficient matrix multiplication in computing system comprising plurality of multiprocessors
US9262174B2 (en) Dynamic bank mode addressing for memory access
US20180121388A1 (en) Symmetric block sparse matrix-vector multiplication
US11061741B2 (en) Techniques for efficiently performing data reductions in parallel processing units
US7836118B1 (en) Hardware/software-based mapping of CTAs to matrix tiles for efficient matrix multiplication
CN108139994B (zh) 内存访问方法及内存控制器
CN110333827B (zh) 一种数据加载装置和数据加载方法
US9600852B2 (en) Hierarchical hash tables for SIMT processing and a method of establishing hierarchical hash tables
US20110072438A1 (en) Fast mapping table register file allocation algorithm for simt processors
JP2022508028A (ja) ３次元画像処理におけるデータの読み書き方法とシステム、記憶媒体及び端末
CN112837205B (zh) 一种图形处理器上基于延迟修正的批量矩阵求逆方法
US9928033B2 (en) Single-pass parallel prefix scan with dynamic look back
US10235208B2 (en) Technique for saving and restoring thread group operating state
JP2022512311A (ja) 行列数学命令セットのタイリングアルゴリズム
US9830161B2 (en) Tree-based thread management
US20220188380A1 (en) Data processing method and apparatus applied to graphics processing unit, and electronic device
CN115630013A (zh) 基于空间可重构阵列的便笺式缓存架构构建方法及系统
Wu et al. Optimizing dynamic programming on graphics processing units via data reuse and data prefetch with inter-block barrier synchronization
US11966999B2 (en) Real-time simulation using material point method on graphics processing units

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
GR01 Patent grant
GR01 Patent grant