CN113641470A - 基于gpu的线程排布方法、装置、设备以及存储介质 - Google Patents

基于gpu的线程排布方法、装置、设备以及存储介质 Download PDF

Info

Publication number
CN113641470A
CN113641470A CN202110835793.9A CN202110835793A CN113641470A CN 113641470 A CN113641470 A CN 113641470A CN 202110835793 A CN202110835793 A CN 202110835793A CN 113641470 A CN113641470 A CN 113641470A
Authority
CN
China
Prior art keywords
threads
thread
array
mapping
elements
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.)
Pending
Application number
CN202110835793.9A
Other languages
English (en)
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.)
Dawning Information Industry Beijing Co Ltd
Original Assignee
Shuguang Cloud Computing Group Co ltd
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 Shuguang Cloud Computing Group Co ltd filed Critical Shuguang Cloud Computing Group Co ltd
Priority to CN202110835793.9A priority Critical patent/CN113641470A/zh
Publication of CN113641470A publication Critical patent/CN113641470A/zh
Pending legal-status Critical Current

Links

Images

Classifications

    • 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/48—Program initiating; Program switching, e.g. by interrupt
    • G06F9/4806—Task transfer initiation or dispatching
    • G06F9/4843—Task transfer initiation or dispatching by program, e.g. task dispatcher, supervisor, operating system
    • 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/5027—Allocation of resources, e.g. of the central processing unit [CPU] to service a request the resource being a machine, e.g. CPUs, Servers, Terminals

Landscapes

  • Engineering & Computer Science (AREA)
  • Software Systems (AREA)
  • Theoretical Computer Science (AREA)
  • Physics & Mathematics (AREA)
  • General Engineering & Computer Science (AREA)
  • General Physics & Mathematics (AREA)
  • Advance Control (AREA)

Abstract

本发明公开了一种基于GPU的线程排布方法、装置、设备以及存储介质，所述方法包括：获取所需要的线程；将所述线程分配到一维网格中；通过核函数将待计算的数组映射到所述线程上。本发明的技术方案，通过优化线程的排布，将线程映射到一维网格中，使得线程可以按需分配，提高线程的利用率，可以使得很多计算不是非常复杂的核函数的性能得到一定程度的提升，且不需要修改核函数中的计算代码，只需要修改元素的映射代码，既可以降低操作难度也可以减少工作量；同时，本发明提供的映射方法，不仅可以计算三维数组，还可以很轻松的实现维数大于3的数组元素的并行计算，应用范围广泛。

Description

基于GPU的线程排布方法、装置、设备以及存储介质
技术领域
本发明属于计算机技术领域，尤其涉及一种基于GPU的线程排布方法、装置、设备以及存储介质。
背景技术
CUDA(Compute Unified Device Architecture)引入了线程(thread)，块(block)，以及网格(grid)等概念，且他们都是三维的。一般的，计算三维数组都会默认使用三维的线程，块来开发异构计算代码，因为这样映射足够清晰明了，三维数组中每一个数组维度的计算对应一个维度的线程。
但是，这种线程排布方式在很多时候会有以下缺点：
1)由于虽然在逻辑概念上线程，块以及网格都可以是三维的，但实际上对应到真正执行计算的SP单元，仍然是一维的。当大量三维线程访问全局内存时，很容易会由于调度器的某些策略问题导致访存变慢，这会严重影响异构程序的执行效率，大量的线程在等待读取数据的过程中闲置，效率低下。
2)根据不同算法，会浪费一部分开辟的线程；因为开辟线程的数量和核函数启动时间是呈正相关的，所以假设我们开辟了96*128个线程，但是只使用了96*96个，那么浪费的这些线程在启动时也是有开销的，影响异构程序性能。
3)由于三维的线程块数量在每个维度计算时都要向上取整，这会导致每个维度的线程数量分配无法做到按需分配，每个维度的线程块都会出现浪费的情况。
发明内容
本发明旨在至少在一定程度上解决相关技术中的技术问题之一。为此，本发明的一个目的在于提出一种可以优化线程排布，进而提高线程的利用率的基于GPU的线程排布方法、装置、设备以及存储介质。
为了解决上述技术问题，本发明的实施例提供如下技术方案：
一种基于GPU的线程排布方法，包括：
获取所需要的线程；
将所述线程分配到一维网格中；
通过核函数将待计算的数组映射到所述线程上。
可选的，所述获取所述需要的线程，包括：
确定所需要的所述的线程的数量；其中，所述线程的数量根据所述数组的元素的数量确定；
根据所述数量，生成所述线程。
可选的，所述获取所述需要的线程后，包括：
启动所述核函数。
可选的，所述将所述线程分配到一维网格中，包括：
确定每个所述线程在所述一维网格中的标号；
根据所述标号将所述线程分配到一维网格中。
可选的，通过核函数将待计算的数组映射到所述线程上，包括：
确定所述数组的元素的个数的总和；
根据所述元素的个数的总和，通过和函数将所述元素映射到所述线程上。
可选的，所述数组为N维数组，N为大于等于3的正整数。
本发明的实施例还提供一种基于GPU的线程排布装置，包括：
获取模块，用于获取所需要的线程；
分配模块，用于将所述线程分配到一维网格中；
映射模块，用于通过核函数将待计算的数组映射到所述线程上。
本发明的实施例还提供一种计算机设备，包括存储器、处理器及存储在存储器上并可在处理器上运行的计算机程序，所述处理器执行所述计算机程序时实现如上所述的方法。
本发明的实施例还提供一种计算机可读存储介质，其上存储有计算机程序，所述计算机程序被处理器执行时实现如上所述的方法。
本发明的实施例，具有如下技术效果：
本发明的上述技术方案，1)通过优化线程的排布，将线程映射到一维网格中，使得线程可以按需分配，提高线程的利用率，可以使得很多计算不是非常复杂的核函数的性能得到一定程度的提升，且不需要修改核函数中的计算代码，只需要修改元素的映射代码，既可以降低操作难度也可以减少工作量。
2)同时，本发明提供的映射方法，不仅可以计算三维数组，还可以很轻松的实现维数大于3的数组元素的并行计算，应用范围广泛。
本发明附加的方面和优点将在下面的描述中部分给出，部分将从下面的描述中变得明显，或通过本发明的实践了解到。
附图说明
图1是本发明实施例提供的的流程示意图；
图2是本发明实施例提供的三维网格结构示意图；
图3是本发明实施例提供的一维网格的结构示意图；
图4是本发明实施例提供的计算机设备的结构示意图。
具体实施方式
下面详细描述本发明的实施例，所述实施例的示例在附图中示出，其中自始至终相同或类似的标号表示相同或类似的元件或具有相同或类似功能的元件。下面通过参考附图描述的实施例是示例性的，旨在用于解释本发明，而不能理解为对本发明的限制。
本发明提到的，GPU(graphics processing unit)图形处理器；CUDA(ComputeUnified Device Architecture)一种运算平台；thread：线程；block：块；grid：网格；硬件概念中的流处理器(streaming processor，SP)，GPU大核(streaming multiprocessor，SM)；knernel：核函数。
如图1所示，本发明的实施例提供一种基于GPU的线程排布方法，包括：
步骤S1：获取所需要的线程；
步骤S2：将所述线程分配到一维网格中；
步骤S3：通过核函数将待计算的数组映射到所述线程上。
具体的，如图2所示，为xyz三维网格，表明了在异构计算程序移植开发的过程中，分析GPU加速卡在细粒度并行计算过程中，SP、SM和三维线程、块及网格之间的映射关系，并结合移植的气象预报程序中的热点计算代码，如图3所示，将所需要计算的三维数组映射到一维网格中。
本发明的该实施例，通过优化线程的排布，将线程映射到一维网格中，使得线程可以按需分配，提高线程的利用率，可以使得很多计算不是非常复杂的核函数的性能得到一定程度的提升，且不需要修改核函数中的计算代码，只需要修改元素的映射代码，既可以降低操作难度也可以减少工作量。
本发明一可选的实施例，步骤S1中，所述获取所述需要的线程，包括：
步骤S11：确定所需要的所述的线程的数量；其中，所述线程的数量根据所述数组的元素的数量确定；
步骤S12：根据所述数量，生成所述线程。
具体的，线程总数一般与核函数计算数组中的元素个数对应，避免线程闲置，防止线程资源的浪费，减少线程在启动时产生的不必要的开销，提升了异构程序的性能。
本发明一可选的实施例，步骤S1中，所述获取所述需要的线程后，包括：
步骤S13：启动所述核函数。
本发明一可选的实施例，步骤S2中，所述将所述线程分配到一维网格中，包括：
步骤S21：确定每个所述线程在所述一维网格中的标号；
具体的标号为tid。
步骤S22：根据所述标号将所述线程分配到一维网格中。
本发明一可选的实施例，步骤S3中，通过核函数将待计算的数组映射到所述线程上，包括：
步骤S31：确定所述数组的元素的个数的总和；
步骤S32：根据所述元素的个数的总和，通过和函数将所述元素映射到所述线程上。
本发明一可选的实施例，步骤S3中，所述数组为N维数组，N为大于等于3的正整数。
具体的，以N为3为例，算出整个三维数组(i,j,k)的所有数组元素个数的总和ijkdim，接着通过映射公式将数组的三个维度在线程上映射出来，也即将数组映射在了一维网格上。
以如下的代码为例：
int tid＝(blockIdx.x)*blockDim.x+threadldx.x；
int ijkdim＝(ied-isd+1)*(jed-jsd+1)*npz；
int k＝(tid)/((ied-isd+1)*(jed-jsd+1))；
Int j＝(tid-(k)*((ied-isd+1)*(jed-jsd+1))/(ied-isd+1)；
int i＝(tid)-(j)*(ied-isd+1)-(k)*((ied-isd+1)*(jed-jsd+1))；
本发明的该实施例，通过上述映射方法，不仅可以计算三维数组，还可以实现维数大于3的数组元素的并行计算，应用范围广泛。
如下表1列出的数据是经过异构计算天气预报软件中的大量核函数性能测试结果。
表1某些三维数组计算核函数性能试验结果
Figure BDA0003176920790000041
Figure BDA0003176920790000051
通过对比表1的数据可知：在该软件中参与测试的核函数中，最高可以获得26％的提速，所有核函数平均提速16.15％，可以看出在大部分的核函数中，该映射方法相当有效，核函数的性能提升效果显著。
本发明的实施例还提供一种基于GPU的线程排布装置，包括：
获取模块，用于获取所需要的线程；
分配模块，用于将所述线程分配到一维网格中；
映射模块，用于通过核函数将待计算的数组映射到所述线程上。
本发明的该实施例，数组使用GPU并行计算时，将元素映射到具体线程上的操作实现在核函数代码中，而不是靠编译器，优化了线程排布，提高了线程的利用率。
如图4所示，本发明的实施例还提供一种计算机设备40，包括存储器42、处理器41及存储在存储器42上并可在处理器41上运行的计算机程序，所述处理器41执行所述计算机程序时实现如上所述的方法。
另外，本发明实施例的线程排布的其他作用对本领域的技术人员来说是已知的，为减少冗余，此处不做赘述。
需要说明的是，在流程图中表示或在此以其他方式描述的逻辑和/或步骤，例如，可以被认为是用于实现逻辑功能的可执行指令的定序列表，可以具体实现在任何计算机可读介质中，以供指令执行系统、装置或设备(如基于计算机的系统、包括处理器的系统或其他可以从指令执行系统、装置或设备取指令并执行指令的系统)使用，或结合这些指令执行系统、装置或设备而使用。就本说明书而言，"计算机可读介质"可以是任何可以包含、存储、通信、传播或传输程序以供指令执行系统、装置或设备或结合这些指令执行系统、装置或设备而使用的装置。计算机可读介质的更具体的示例(非穷尽性列表)包括以下：具有一个或多个布线的电连接部(电子装置)，便携式计算机盘盒(磁装置)，随机存取存储器(RAM)，只读存储器(ROM)，可擦除可编辑只读存储器(EPROM或闪速存储器)，光纤装置，以及便携式光盘只读存储器(CDROM)。另外，计算机可读介质甚至可以是可在其上打印所述程序的纸或其他合适的介质，因为可以例如通过对纸或其他介质进行光学扫描，接着进行编辑、解译或必要时以其他合适方式进行处理来以电子方式获得所述程序，然后将其存储在计算机存储器中。
应当理解，本发明的各部分可以用硬件、软件、固件或它们的组合来实现。在上述实施方式中，多个步骤或方法可以用存储在存储器中且由合适的指令执行系统执行的软件或固件来实现。例如，如果用硬件来实现，和在另一实施方式中一样，可用本领域公知的下列技术中的任一项或他们的组合来实现：具有用于对数据信号实现逻辑功能的逻辑门电路的离散逻辑电路，具有合适的组合逻辑门电路的专用集成电路，可编程门阵列(PGA)，现场可编程门阵列(FPGA)等。
在本说明书的描述中，参考术语“一个实施例”、“一些实施例”、“示例”、“具体示例”、或“一些示例”等的描述意指结合该实施例或示例描述的具体特征、结构、材料或者特点包含于本发明的至少一个实施例或示例中。在本说明书中，对上述术语的示意性表述不一定指的是相同的实施例或示例。而且，描述的具体特征、结构、材料或者特点可以在任何的一个或多个实施例或示例中以合适的方式结合。
在本发明的描述中，需要理解的是，术语“中心”、“纵向”、“横向”、“长度”、“宽度”、“厚度”、“上”、“下”、“前”、“后”、“左”、“右”、“竖直”、“水平”、“顶”、“底”“内”、“外”、“顺时针”、“逆时针”、“轴向”、“径向”、“周向”等指示的方位或位置关系为基于附图所示的方位或位置关系，仅是为了便于描述本发明和简化描述，而不是指示或暗示所指的装置或元件必须具有特定的方位、以特定的方位构造和操作，因此不能理解为对本发明的限制。
此外，术语“第一”、“第二”仅用于描述目的，而不能理解为指示或暗示相对重要性或者隐含指明所指示的技术特征的数量。由此，限定有“第一”、“第二”的特征可以明示或者隐含地包括至少一个该特征。在本发明的描述中，“多个”的含义是至少两个，例如两个，三个等，除非另有明确具体的限定。
在本发明中，除非另有明确的规定和限定，术语“安装”、“相连”、“连接”、“固定”等术语应做广义理解，例如，可以是固定连接，也可以是可拆卸连接，或成一体；可以是机械连接，也可以是电连接；可以是直接相连，也可以通过中间媒介间接相连，可以是两个元件内部的连通或两个元件的相互作用关系，除非另有明确的限定。对于本领域的普通技术人员而言，可以根据具体情况理解上述术语在本发明中的具体含义。
在本发明中，除非另有明确的规定和限定，第一特征在第二特征“上”或“下”可以是第一和第二特征直接接触，或第一和第二特征通过中间媒介间接接触。而且，第一特征在第二特征“之上”、“上方”和“上面”可是第一特征在第二特征正上方或斜上方，或仅仅表示第一特征水平高度高于第二特征。第一特征在第二特征“之下”、“下方”和“下面”可以是第一特征在第二特征正下方或斜下方，或仅仅表示第一特征水平高度小于第二特征。
尽管上面已经示出和描述了本发明的实施例，可以理解的是，上述实施例是示例性的，不能理解为对本发明的限制，本领域的普通技术人员在本发明的范围内可以对上述实施例进行变化、修改、替换和变型。

Claims (9)

1.一种基于GPU的线程排布方法，其特征在于，包括：
获取所需要的线程；
将所述线程分配到一维网格中；
通过核函数将待计算的数组映射到所述线程上。
2.根据权利要求1所述的方法，其特征在于，所述获取所述需要的线程，包括：
确定所需要的所述的线程的数量；其中，所述线程的数量根据所述数组的元素的数量确定；
根据所述数量，生成所述线程。
3.根据权利要求2所述的方法，其特征在于，所述获取所述需要的线程后，包括：
启动所述核函数。
4.根据权利要求3所述的方法，其特征在于，所述将所述线程分配到一维网格中，包括：
确定每个所述线程在所述一维网格中的标号；
根据所述标号将所述线程分配到一维网格中。
5.根据权利要求4所述的方法，其特征在于，通过核函数将待计算的数组映射到所述线程上，包括：
确定所述数组的元素的个数的总和；
根据所述元素的个数的总和，通过和函数将所述元素映射到所述线程上。
6.根据权利要求5所述的方法，其特征在于，所述数组为N维数组，N为大于等于3的正整数。
7.一种基于GPU的线程排布装置，其特征在于，包括：
获取模块，用于获取所需要的线程；
分配模块，用于将所述线程分配到一维网格中；
映射模块，用于通过核函数将待计算的数组映射到所述线程上。
8.一种计算机设备，包括存储器、处理器及存储在存储器上并可在处理器上运行的计算机程序，其特征在于，所述处理器执行所述计算机程序时实现如权利要求1至6任一项所述的方法。
9.一种计算机可读存储介质，其上存储有计算机程序，其特征在于，所述计算机程序被处理器执行时实现如权利要求1至6任一项所述的方法。
CN202110835793.9A 2021-07-23 2021-07-23 基于gpu的线程排布方法、装置、设备以及存储介质 Pending CN113641470A (zh)

Priority Applications (1)

Application Number Priority Date Filing Date Title
CN202110835793.9A CN113641470A (zh) 2021-07-23 2021-07-23 基于gpu的线程排布方法、装置、设备以及存储介质

Applications Claiming Priority (1)

Application Number Priority Date Filing Date Title
CN202110835793.9A CN113641470A (zh) 2021-07-23 2021-07-23 基于gpu的线程排布方法、装置、设备以及存储介质

Publications (1)

Publication Number Publication Date
CN113641470A true CN113641470A (zh) 2021-11-12

Family

ID=78418185

Family Applications (1)

Application Number Title Priority Date Filing Date
CN202110835793.9A Pending CN113641470A (zh) 2021-07-23 2021-07-23 基于gpu的线程排布方法、装置、设备以及存储介质

Country Status (1)

Country Link
CN (1) CN113641470A (zh)

Cited By (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN114860341A (zh) * 2022-05-19 2022-08-05 北京百度网讯科技有限公司 线程配置方法、设备、装置、存储介质及程序产品

Citations (3)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN107229051A (zh) * 2017-05-26 2017-10-03 西安电子科技大学 基于gpu的视频sar回波仿真并行实现方法
CN109753726A (zh) * 2019-01-04 2019-05-14 东南大学 一种基于边界盒搜索方法和gpu的球磨机介质运动仿真方法
US20190278593A1 (en) * 2018-03-09 2019-09-12 Nvidia Corporation Accelerating linear algebra kernels for any processor architecture

Patent Citations (3)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN107229051A (zh) * 2017-05-26 2017-10-03 西安电子科技大学 基于gpu的视频sar回波仿真并行实现方法
US20190278593A1 (en) * 2018-03-09 2019-09-12 Nvidia Corporation Accelerating linear algebra kernels for any processor architecture
CN109753726A (zh) * 2019-01-04 2019-05-14 东南大学 一种基于边界盒搜索方法和gpu的球磨机介质运动仿真方法

Cited By (2)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN114860341A (zh) * 2022-05-19 2022-08-05 北京百度网讯科技有限公司 线程配置方法、设备、装置、存储介质及程序产品
CN114860341B (zh) * 2022-05-19 2023-09-22 北京百度网讯科技有限公司 线程配置方法、设备、装置、存储介质

Similar Documents

Publication Publication Date Title
CN105022670B (zh) 一种云计算平台中的异构分布式任务处理系统及其处理方法
CN110515739B (zh) 深度学习神经网络模型负载计算方法、装置、设备及介质
US8364739B2 (en) Sparse matrix-vector multiplication on graphics processor units
US8782645B2 (en) Automatic load balancing for heterogeneous cores
Xu et al. Graph processing on GPUs: Where are the bottlenecks?
EP2798499B1 (en) Methods and apparatus to manage workload memory allocation
Xu et al. Software bloat analysis: Finding, removing, and preventing performance problems in modern large-scale object-oriented applications
US9329867B2 (en) Register allocation for vectors
US9477465B2 (en) Arithmetic processing apparatus, control method of arithmetic processing apparatus, and a computer-readable storage medium storing a control program for controlling an arithmetic processing apparatus
US10332229B2 (en) System and method for high performance k-means clustering on GPU with smart kernels
Yang et al. Intermediate data caching optimization for multi-stage and parallel big data frameworks
Takizawa et al. SPRAT: Runtime processor selection for energy-aware computing
US10268519B2 (en) Scheduling method and processing device for thread groups execution in a computing system
Maggioni et al. AdELL: An adaptive warp-balancing ELL format for efficient sparse matrix-vector multiplication on GPUs
CN112130901A (zh) 基于risc-v的协处理器、数据处理方法及存储介质
Huang et al. GPU computing performance analysis on matrix multiplication
CN114924748A (zh) 编译方法、装置及设备
CN113641470A (zh) 基于gpu的线程排布方法、装置、设备以及存储介质
Wang et al. Gpu register packing: Dynamically exploiting narrow-width operands to improve performance
Fang et al. An auto-tuning solution to data streams clustering in opencl
CN108228351B (zh) Gpu的性能均衡调度方法、存储介质及电子终端
CN106598552A (zh) 基于Gridding模块的数据点转换方法及装置
US11188315B1 (en) Method and apparatus for reusable and relative indexed register resource allocation in function calls
CN112130977B (zh) 一种任务调度方法、装置、设备及介质
CN105718223B (zh) 管理工作负载存储器分配的方法和设备

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
TA01 Transfer of patent application right
TA01 Transfer of patent application right

Effective date of registration: 20211224

Address after: 100193 No. 36 Building, No. 8 Hospital, Wangxi Road, Haidian District, Beijing

Applicant after: Dawning Information Industry (Beijing) Co.,Ltd.

Address before: 100193 5 floor, 36 building, No. 8 Northeast Road, Haidian District, Beijing.

Applicant before: Shuguang Cloud Computing Group Co.,Ltd.