CN114860341B - 线程配置方法、设备、装置、存储介质 - Google Patents

Abstract

本公开提供了一种线程配置方法、设备、装置、存储介质及程序产品，涉及大数据及人工智能技术领域，特别涉及深度学习技术领域。在本公开的一些实施例中，线程配置方法包括获取输入的第一维度和第二维度，以及需要计算的中间维度，其中，第二维度大于所述第一维度；根据第一维度和中间维度对线程块之中的线程进行配置，以生成线程块在x方向上的第一线程数和在y方向上的第二线程数；根据第二维度以及所述线程块在x方向上的第一线程数和在y方向上的第二线程数对线程栅格之中的线程进行配置，以生成线程栅格在x方向上的第三线程数和在y方向上的第四线程。通过本公开的实施例对线程块之中的线程进行配置，可以大幅提高访存效率。

Description

线程配置方法、设备、装置、存储介质

技术领域

本公开涉及计算机技术领域，具体涉及大数据及人工智能技术领域，特别涉及深度学习技术领域，尤其涉及一种线程配置方法、装置、设备、存储介质和程序产品。

背景技术

随着计算机技术的快速发展，深度学习框架的应用范围越来越广，例如图形处理器算子等等。

在深度学习框架中，底层算子的性能在很大程度上决定了该深度学习框架的性能，底层算子的性能由计算效率和访存效率决定，目前，在深度学习框架中，存在访问效率低的问题。

发明内容

本公开提供了一种线程配置方法、装置、设备、存储介质及程序产品。

根据本公开的一方面，提供了一种线程配置方法，包括：

获取输入的第一维度和第二维度，以及需要计算的中间维度，其中，所述第二维度大于所述第一维度；

根据所述第一维度和所述中间维度对线程块之中的线程进行配置，以生成所述线程块在x方向上的第一线程数和在y方向上的第二线程数；

根据所述第二维度和所述线程块在x方向上的第一线程数和在y方向上的第二线程数对线程栅格之中的线程进行配置，以生成所述线程栅格在x方向上的第三线程数和在y方向上的第四线程数。

根据本公开的另一方面，提供了一种线程配置装置，包括：

获取模块，用于获取输入的第一维度和第二维度，以及需要计算的中间维度，其中，所述第二维度大于所述第一维度；

第一生成模块，用于根据所述第一维度和所述中间维度对线程块之中的线程进行配置，以生成所述线程块在x方向上的第一线程数和在y方向上的第二线程数；

第二生成模块，用于根据所述第二维度和所述线程块在x方向上的第一线程数和在y方向上的第二线程数对线程栅格之中的线程进行配置，以生成所述线程栅格在x方向上的第三线程数和在y方向上的第四线程数。

根据本公开的另一方面，提供了一种电子设备，包括：

至少一个处理器；以及

与所述至少一个处理器通信连接的存储器；其中，

所述存储器存储有可被所述至少一个处理器执行的指令，所述指令被所述至少一个处理器执行，以使所述至少一个处理器能够执行上述的方法中的各步骤。

根据本公开的另一方面，提供一种存储有计算机指令的非瞬时计算机可读存储介质，其中，所述计算机指令用于使所述计算机执行上述的方法中的各步骤。

根据本公开的另一方面，提供一种计算机程序产品，包括计算机程序。所述计算机程序在被处理器执行时实现上述的方法中的各步骤。

在本公开的一些实施例中，通过为线程块block和线程栅格grid配置适合的线程数，从而可以提高访存效率。

应当理解，本部分所描述的内容并非旨在标识本公开的实施例的关键或重要特征，也不用于限制本公开的范围。本公开的其它特征将通过以下的说明书而变得容易理解。

附图说明

附图用于更好地理解本方案，不构成对本公开的限定。其中：

图1为本公开第一实施例提供的一种线程配置方法的流程示意图；

图2a为本公开示例性实施例提供的一种线程配置装置的结构框图；

图2b为本公开示例性实施例提供的另一种线程配置装置的结构框图；

图3示出了可以用来实施本公开的实施例的示例电子设备的示意性框图。

具体实施方式

以下结合附图对本公开的示范性实施例做出说明，其中包括本公开实施例的各种细节以助于理解，应当将它们认为仅仅是示范性的。因此，本领域普通技术人员应当认识到，可以对这里描述的实施例做出各种改变和修改，而不会背离本公开的范围和精神。同样，为了清楚和简明，以下的描述中省略了对公知功能和结构的描述。

本公开的技术方案中，所涉及的用户个人信息的收集、存储、使用、加工、传输、提供和公开等处理，均符合相关法律法规的规定，且不违背公序良俗。

本公开的技术方案中，所涉及的用户个人信息的获取、存储和应用等，均符合相关法律法规的规定，且不违背公序良俗。

目前，在深度学习框架中的线程配置中，存在访问效率低的问题。在深度学习框架中，底层算子的性能在很大程度上决定了该深度学习框架的性能，底层算子的性能由计算效率和访存效率决定，另外，IO密集型算子种类更多，占比更大，对于这类算子来说，其访存效率往往决定了其性能的优劣，IO密集型算子访存效率较低。

一个操作系统内核在启动时，程序需要启动的线程块block和线程栅格grid可以为一维、二维或者三维，不同的线程配置方式，决定了通用并行计算架构操作系统内核不同的并行策略实现以及不同的并行执行效率，一维的线程块block和线程栅格grid最常见也最利于理解，深度学习框架中的算子在最初实现是往往会优先考虑这种实现方式，看i如对于某些规约操作，通常会采用一维的线程块block设置，每个线程块block完成一次完整的规约操作并得到结果。

另外，目前常见的线程块block和线程栅格grid的配置方式，虽然在逻辑上简单易懂，但是在一些场景中存在性能不理想的问题，比如无法实现内存合并访问，即一个线程束warp内的线程无法访问连续的内存块以及线程的并行度不高，无法充分利用多处理器资源的问题。

由于采用单指令多线程的架构，图形处理器对内存的访问指令是由线程束warp发起的，即线程束warp中每个线程同时执行内存操作指令，不过每个线程所访问的数据地址可以不同，图形处理器会根据这些不同的地址发起一次或多次动态随机存取存储器－＞缓存的传输(事务)，直到所有线程都拿到各自所需的数据(缓存－＞寄存器)。显然，我们可以通过减少动态随机存取存储器－＞缓存的传输事务来优化程序性能。内存合并访问是理想的全局内存访问方式，即一个图形处理器线程束warp内的线程访问连续的内存块。在进行带有跨度的非最低维数据读取时，一个图形处理器线程束warp内的线程无法实现内存合并访问，从而导致访存效率差的问题。

一个操作系统内核的并行效率并不完全取决于一个线程块block内的线程数量，其原因首先是能够同时并发的线程束warp数量会受到多处理器上的资源限制，因为每个多处理器上的资源如寄存器、共享内存等是有限的，因此线程束warp中每个线程所需的资源越多，能同时启动的线程束warp数就越少；其次，若不考虑实际多维数据的排布特点，仅使用一维线程块block和线程栅格grid配置中会导致负载不均衡，无法在多变的实际场景中都保证高效地并行效率。

针对上述存在的技术问题，在本公开的一些实施例中，线程配置装置获取输入的第一维度和第二维度，以及需要计算的中间维度，其中，所述第二维度大于所述第一维度；根据所述第一维度和所述中间维度对线程块block之中的线程进行配置，以生成所述线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y；根据所述第二维度和所述线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y对线程栅格grid之中的线程进行配置，以生成所述线程栅格grid在x方向上的第三线程数grid.x和在y方向上的第四线程数grid.y，对线程块block之中的线程进行配置，提高访存效率。以下结合附图，详细说明本公开各实施例提供的技术方案。

图1为本公开第一实施例提供的一种，线程配置方法的流程示意图。如图1所示，该方法包括：

S101：获取输入的第一维度和第二维度，以及需要计算的中间维度，其中，第二维度大于所述第一维度；

S102：根据第一维度和所述中间维度对线程块block之中的线程进行配置，以生成线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y；

S103：根据第二维度和所述线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y对线程栅格grid之中的线程进行配置，以生成所述线程栅格grid在x方向上的第三线程数grid.x和在y方向上的第四线程数grid.y。

在本实施例中，上述方法的执行主体线程配置装置可以为服务器。

当线程配置装置为服务器时，本实施例不限定服务器的实现形式，服务器可以是常规服务器、云服务器、云主机、虚拟中心等服务器设备。其中，服务器设备的构成主要包括处理器、硬盘、内存、系统总线等，和通用的计算机架构类似。

需要说明的是，第一维度是指最低维度，第二维度是指最高维度，线程是指最小的执行单位，多个线程并行执行一个图形处理器程序。

线程块block是由多个线程组成，同一个线程块block中的线程可以同步并且通过共享内存通信，其中，线程块block设置原则是保证线程间的负载均衡、避免每个线程的工作量过大，对于二维的线程块block，在保证线程块block线程数大于64的情况下，保证其形状接近正方形，例如所述线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y均为32，即线程块block线程数为1024＝32＊32。

线程栅格grid由多个线程块block组成。

需要注意的是，根据实际要处理的数据块排布特点，不局限于线程块block和线程栅格grid仅设置为一维，而是设置成二维或三维，充分利用了多处理器的资源，提高并行度。

因为线程束warp为32，因此在设置线程块block时，为了计算方便会先向上对齐，最低维度向上对齐到2的整数次幂后的值为low，需要计算的中间维度向上对齐到2的整数次幂后的值为mid，高维度的值为high。

在本实施例中，线程配置装置获取输入的第一维度和第二维度，以及需要计算的中间维度，其中，第二维度大于所述第一维度；线程配置装置根据所述第一维度和所述中间维度对线程块block之中的线程进行配置，以生成线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y；线程配置装置根据第二维度和线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y对线程栅格grid之中的线程进行配置，以生成线程栅格grid在x方向上的第三线程数grid.x和在y方向上的第四线程数grid.y。本公开实施例中，在对非最低微进行计算时，将现有实现方式中的一个线程块block内所有线程完成一个输出的计算改为一个线程块block内索引值的不同threadIdx.x的线程完成各自不同的输出，现内存合并访问，提升数据读取效率。

需要说明的是，线程块block和线程栅格grid可以设置为一维、二维或三维，本公开实施例以线程块block和线程栅格grid设置为二维进行说明。

例如，线程配置装置获取输入的第一维度值low＝48、第二维度值hig＝512以及计算得到的中间维度值mid＝896，线程配置装置根据所述第一维度值48和中间维度896对线程块block之中的线程进行配置，生成线程块block在x方向上的第一线程数block.x，第一线程数block.x＝32、线程块block在y方向上的第二线程数block.y，第二线程数block.y＝32；线程配置装置根据第二维度值512和线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y对线程栅格grid之中的线程进行配置，生成所述线程栅格grid在x方向上的第三线程数grid.x＝2、生成所述线程栅格grid在y方向上的第四线程数grid.y＝80。

在本实施例中，线程配置装置根据第一维度和中间维度对线程块block之中的线程进行配置，以生成线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y。一种可实现方式为，线程配置装置根据第一维度确定线程块block在x方向上的第五线程数block.x；线程配置装置根据中间维度和线程块block在x方向上的第五线程数block.x确定第二线程数block.y；线程配置装置根据第二线程数block.y和第一维度更新所述线程块block在x方向上的第五线程数block.x，以生成线程块block在x方向上的第一线程数block.x。

本公开实施例根据第一维度确定线程块block在x方向上的第五线程数block.x，根据中间维度和线程块block在x方向上的第五线程数block.x确定线程块block在y方向上的第二线程数block.y，根据线程块block在y方向上的第二线程数block.y和第一维度更新线程块block在x方向上的第五线程数block.x，以生成线程块block在x方向上的第一线程数block.x，防止出现线程块block在y方向上的第二线程数block.y过小以及线程块block整体线程过小的情况，其中，一般情况线程块block不小于64。

例如，线程配置装置根据第一维度值low＝38确定线程块block在x方向上的第五线程数block.x为32；线程配置装置根据中间维度值mid＝896和线程块block在x方向上的第五线程数block.x＝32确定第二线程数block.y＝32；线程配置装置根据第二线程数block.y为和第一维度更新线程块block在x方向上的第五线程数block.x，生成更新后的线程块block在x方向上第一线程数block.x＝32。

在本实施例中，线程配置装置根据第一维度确定线程块block在x方向上的第五线程数block.x。一种可实现方式为，线程配置获取线程束warp的大小；线程配置将线程束warp的大小和所述第一维度之中的较小值作为所述线程块block在x方向上的第五线程数block.x。

需要说明的是，线程束warp为是指图形处理器程序执行的调度单位，也是多处理器中最小执行单位，线线程束warp通常包括32个线程，其中，在同一个线程数warp中的线程，以不同的数据资源执行相同的指令，即单指令多线程。

在本实施例中，线程配置获取线程束warp的大小，本实施例中不对线程配置获取线程束warp的大小的方式进行限制；线程配置将线程束warp的大小和所述第一维度的值进行比较，将线程配置将线程束warp的大小和第一维度的值之中的较小值作为线程块block在x方向上的第五线程数block.x，得到需要的精确的线程块block在x方向上的第五线程数block.x的值。在进行带有跨度的非最低维数据读取时，一个图形处理器线程束warp内的线程无法实现内存合并访问，从而导致访存效率差，一个线程束warp内的线程读取最低维数据(没有stride)，进而实现连续内存块访问，即上述的内存合并访问，提升数据读取效率。

需要注意的是，所有线程数block.x最好不超过线程束warp的大小。

例如，线程配置获取线程束warp的大小为32，线程配置将线程束warp的大小和所述第一维度的值low＝48进行比较，此时线程束warp的大小为32小于第一维度的值low＝48，线程配置将线程束warp的大小32作为线程块block在x方向上的第五线程数block.x。

在本实施例中，线程配置装置根据所述中间维度和线程块block在x方向上的第五线程数block.x确定线程块block在y方向上的第二线程数block.y。一种可实现方式为，线程配置装置获取线程块block中最大可配置线程数；线程配置装置将中间维度与线程块block中最大可配置线程数之中的较小值作为第一值；线程配置装置将中间维度与线程块block在x方向上的第五线程数block.x之中的较小值作为第二值；线程配置装置将第一值与所述第二值之中的较小值作为线程块block在y方向上的第二线程数block.y。

需要说明的时，第一值是指中间维度与线程块block中最大可配置线程数之中的较小值，第二值是指中间维度与线程块block在x方向上的第五线程数block.x之中的较小值。

在本实施例中，线程配置装置获取线程块block中最大可配置线程数；线程配置装置将中间维度与线程块block中最大可配置线程数进行比较，将中间维度与线程块block中最大可配置线程数之中的较小值作为第一值；线程配置装置将中间维度与线程块block在x方向上的第五线程数block.x进行比较，将中间维度与线程块block在x方向上的第五线程数block.x之中的较小值作为第二值；线程配置装置将第一值与第二值进行比较，将第一值与第二值之中的较小值作为第二线程数block.y，由此取线程块block中最大可配置线程数以及中间维度得到第二线程数block.y，提升数据读取效率。

需要注意的是，block.x＊block.y不能超过线程块block最大可配置的线程数。

例如，线程配置装置将中间维度与线程块block中最大可配置线程数进行比较，中间维度值mid＝896，线程块block中最大可配置线程数为1024，中间维度值mid＝896小于线程块block中最大可配置线程数，将中间维度作为第一值；线程配置装置将中间维度与线程块block在x方向上的第五线程数block.x进行比较，block.x＝min(48,32)＝32，初始线程数block.x为中间维度与线程块block在x方向上的第五线程数block.x之中的较小值且该值作为第二值，此时第二值为32；线程配置装置将第一值896与第二值32进行比较，将两者中较小的值32作为第二线程数block.y。

在本实施例中，线程配置装置根据第二线程数block.y和第一维度更新初始线程数block.x，以生成线程块block在x方向上的第一线程数block.x。一种可实现方式为，线程配置装置将第一维度与线程块block在x方向上的第五线程数block.x之中的较小值作为第三值；线程配置装置将第一维度与第二线程数block.y之中的较小值作为第四值；线程配置装置根据第三值和第四值中的较小值，以及第一维度更新所述线程块block在x方向上的第五线程数block.x。

需要说明的时，第三值是指第一维度与所述线程块block在x方向上的第五线程数block.x之中的较小值，第四值是指第一维度与所述第二线程数block.y之中的较小值。

在本实施例中，线程配置装置将第一维度与线程块block在x方向上的第五线程数block.x进行对比，将第一维度与初始线程数block.x之中的较小值作为第三值；线程配置装置将第一维度与第二线程数block.y进行对比，线程配置装置将第一维度与第二线程数block.y之中的较小值作为第四值；线程配置装置根据第三值和第四值中的较小值，以及所述第一维度更新所述程块block在x方向上的第五线程数block.x，防止出现线程数block.y过小，线程块block整体线程数过小的情况，一般线程块block的数量不小于64。

例如，线程配置装置将第一维度与程块block在x方向上的第五线程数block.x进行对比，第一维度值low＝48与初始线程数block.x＝32进行对比，将两者中的较小值32作为第三值；线程配置装置将第一维度与第二线程数block.y进行对比，block.y＝32，线程配置装置将两者之中的较小值32作为第四值；线程配置装置根据第三值和第四值中的较小值32，以及所述第一维度更新所述程块block在x方向上的第五线程数block.x，更新后的block.x＝32。

在本实施例中，线程配置装置根据第二维度以及线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y对线程栅格grid之中的线程进行配置，以生成线程栅格grid在x方向上的第三线程数grid.x和在y方向上的第四线程数grid.y，一种可实现方式为：线程配置装置根据第一维度和更新后线程数block.x确定在所述x方向上的第三线程数grid.x；线程配置装置根据所述第二维度和所述初始在x方向上的第三线程数grid.x确定在y方向上的第四线程数grid.y。

在本实施例中，线程配置装置根据第一维度和第一线程数block.x确定在所述x方向上的第三线程数grid.x；线程配置装置根据第二维度和初始在x方向上的第三线程数grid.x确定在y方向上的第四线程数grid.y，从而得到线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y，通过第一维度，准确得到线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y。

例如，第一维度值low＝48、第二维度值high＝512、线程数block.x＝32，线程配置装置根据第一维度和第一线程数block.x确定在所述x方向上的线程数grid.x＝2；线程配置装置根据所述第二维度和所述初始在x方向上的第三线程数grid.x确定在所述y方向上的第四线程数grid.y＝80。

在本实施例中，线程配置装置根据第一维度和更新后线程数block.x确定在x方向上的第三线程数grid.x，一种可实现方式为：线程配置装置获取线程栅格grid的可同时启动最大线程块数量；线程配置装置将第一维度向上取整后的值与可同时启动最大线程块数量之中的较小值作为第五值；线程配置装置将程块block在x方向上的第五线程数block.x向上取整后的值与可同时启动最大线程块数量之中的较小值作为第六值；线程配置装置将第五值和第六值中较小值作为所述在x方向上的第三线程数grid.x。

需要说明的是，第五值是指第一维度向上取整后的值与可同时启动最大线程块数量之中的较小值，第六值是指程块block在x方向上的第五线程数block.x向上取整后的值与可同时启动最大线程块数量之中的较小值；结合当前硬件资源配置与线程块block的配置，可以同时启动最大block数量。

在本实施例中，线程配置装置获取线程栅格grid的可同时启动最大线程块数量；线程配置装置将第一维度向上取整后的值与可同时启动最大线程块数量进行比较，第一维度向上取整后的值与可同时启动最大线程块数量之中的较小值作为第五值；线程配置装置将第一线程数block.x向上取整后的值与可同时启动最大线程块数量进行比较，将第一线程数block.x向上取整后的值与可同时启动最大线程块数量之中的较小值作为第六值；线程配置装置将第五值和第六值进行比较并将第五值和第六值中较小值作为所述在x方向上的第三线程数grid.x，实现不同的线程完成各自不同输出的计算，提高在各场景下的高效运行效率。

本实施例表示计算最低维度配置下第一维度值个输出可配置的线程块block数量，并决定了计算时循环的步长和次数。

例如，线程配置装置获取线程栅格grid的可同时启动最大线程块数量，可同时启动最大线程块数量为160；线程配置装置将第一维度向上取整后的值与可同时启动最大线程块数量进行比较，第一维度向上取整后的值与可同时启动最大线程块数量之中的较小值作为第五值；线程配置装置将第一线程数block.x向上取整后的值与可同时启动最大线程块数量进行比较，第一线程数block.x向上取整后的值与可同时启动最大线程块数量之中的较小值作为第六值；线程配置装置将第五值和第六值进行比较并将第五值和第六值中较小值作为所述在x方向上的第三线程数grid.x，其中grid.x＝min(ceil(48/32),160)＝2；本实施例表示计算最低维度配置下第一维度值个输出可配置的线程块block数量，并决定了计算时循环的步长和次数。

在本实施例中，线程配置装置根据第二维度和初始在x方向上的第三线程数grid.x确定在y方向上的第四线程数grid.y，一种可实现方式为：线程配置装置将可同时启动最大线程块数量向上取整后的值与所述第二维度与之中的较小值为第七值；线程配置装置将在x方向上的第三线程数grid.x向上取整后的值与所述第二维度与之中的较小值作为第八值；线程配置装置将所述第七值和第八值中的较小值作为在y方向上的第四线程数grid.y。

需要说明的是，第七值是指可同时启动最大线程块数量向上取整后的值与所述第二维度与之中的较小值，第八值是指在x方向上的第三线程数grid.x向上取整后的值与所述第二维度与之中的较小值作。

在本实施例中，线程配置装置将可同时启动最大线程块数量向上取整后的值与所述第二维度进行比较，将可同时启动最大线程块数量向上取整后的值与所述第二维度之中的较小值为第七值；线程配置装置将在x方向上的第三线程数grid.x向上取整后的值与所述第二维度进行比较，将在x方向上的第三线程数grid.x向上取整后的值与所述第二维度之中的较小值作为第八值；线程配置装置将所述第七值和第八值中的较小值作为在y方向上的第四线程数grid.y，实现不同的线程完成各自不同输出的计算，提高在各场景下的高效运行。

例如，线程配置装置将可同时启动最大线程块数量向上取整后的值与所述第二维度进行比较，第二维度值high＝512，将可同时启动最大线程块数量向上取整后的值与所述第二维度之中的较小值160为第七值；线程配置装置将在x方向上的第三线程数grid.x向上取整后的值与所述第二维度进行比较，将在x方向上的第三线程数grid.x向上取整后的值与所述第二维度之中的较小值作为第八值；线程配置装置将所述第七值和第八值中的较小值作为在y方向上的线第四程数grid.y，其中，grid.y＝min(ceil(160/2),512)＝80。

需要注意的是，每个线程块block中具有相同的threadIdx.x的线程，首先完成线程内的最大值和求和的规约计算，然后再完成线程间的规约计算，根据计算结果得到的做大致和结果，对每个线程单独对逐个输入数据进行后续的计算，输出结果。

图2为本公开示例性实施例提供的一种线程配置装置20的结构框图。如图2所示，该线程配置装置20包括：

获取模块21，第一生成模块22，第二生成模块23。

获取模块21，用于获取输入的第一维度和第二维度，以及需要计算的中间维度，其中，第二维度大于所述第一维度；

第一生成模块22，用于根据第一维度和所述中间维度对线程块Block之中的线程进行配置，以生成线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y；

第二生成模块23，用于根据第二维度以及线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y对线程栅格grid之中的线程进行配置，以生成线程栅格grid在x方向上的第三线程数grid.x和在y方向上的第四线程数grid.y。

可选地，获取模块21，获取输入的第一维度和第二维度，以及需要计算的中间维度，线第一生成模块22根据第一维度和所述中间维度对线程块block之中的线程进行配置，以生成线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y；第二生成模块23根据第二维度以及线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y对线程栅格grid之中的线程进行配置，以生成线程栅格grid在x方向上的第三线程数grid.x和在y方向上的第四线程数grid.y。

可选地，如图2b所示，第一生成模块22包括第一确定子模块221，第二确定子模块222和第一生成子模块223。

第一确定子模块221，用于根据第一维度确定线程块x方向上的第五线程数block.x；

第二确定子模块222，用于根据中间维度和线程块x方向上的第五线程数block.x确定线程块y方向上的第二线程数block.y；

第一生成子模块223，用于根据线程块y方向上的第二线程数block.y和第一维度更新线程块x方向上的第五线程数block.x，以生成线程块x方向上的第一线程数block.x。

可选地，第一确定子模块221用于获取线程束warp的大小，并将线程束warp的大小和第一维度之中的较小值作为所述线程块x方向上的第五线程数block.x。

可选地，第二确定子模块222，用于获取线程块block中最大可配置线程数，并将所述中间维度与所述线程块block中最大可配置线程数之中的较小值作为第一值，以及将所述中间维度与线程块x方向上的第五线程数block.x之中的较小值作为第二值，并将所述第一值与所述第二值之中的较小值作为第二线程数block.y。

可选的，第二生成模块23，用于根据第一维度和第一线程数block.x确定在x方向上的线程块在x方向上的第三初始线程数grid.x，并根据第二维度和所述线程块在x方向上的第三线程数grid.x确定在y方向上的第四线程数grid.y。

可选的，第二生成模块23还包括第三确定子模块和第四确定子模块，第三确定模块用于获取所述线程栅格grid的可同时启动最大线程块数量，并将所述第一维度向上取整后的值与所述可同时启动最大线程块数量之中的较小值作为第五值，以及将所述第一线程数block.x向上取整后的值与所述可同时启动最大线程块数量之中的较小值作为第六值，最后将所述第五值和第六值中较小值作为线程栅格在x方向上的第三线程数grid.x。

可选的，所述第四确定子模块用于将可同时启动最大线程块数量向上取整后的值与所述第二维度与之中的较小值为第七值，并将所述在x方向上的第三线程数grid.x向上取整后的值与所述第二维度与之中的较小值作为第八值，以及将所述第七值和第八值中的较小值作为在线程栅格y方向上的第四线程数grid.y。

根据本公开的实施例，本公开还提供了一种电子设备、一种可读存储介质和一种计算机程序产品。

图3示出了可以用来实施本公开的实施例的示例电子设备300的示意性框图。电子设备旨在表示各种形式的数字计算机，诸如，膝上型计算机、台式计算机、工作台、个人数字助理、服务器、刀片式服务器、大型计算机、和其它适合的计算机。电子设备还可以表示各种形式的移动装置，诸如，个人数字处理、蜂窝电话、智能电话、可穿戴设备和其它类似的计算装置。本文所示的部件、它们的连接和关系、以及它们的功能仅仅作为示例，并且不意在限制本文中描述的和/或者要求的本公开的实现。

如图3所示，设备300包括计算单元301，其可以根据存储在只读存储器(ROM)302中的计算机程序或者从存储单元606加载到随机访问存储器(RAM)303中的计算机程序，来执行各种适当的动作和处理。在RAM 303中，还可存储设备300操作所需的各种程序和数据。计算单元301、ROM 302以及RAM 303通过总线304彼此相连。输入/输出(I/O)接口305也连接至总线304。

设备300中的多个部件连接至I/O接口305，包括：输入单元306，例如键盘、鼠标等；输出单元307，例如各种类型的显示器、扬声器等；存储单元306，例如磁盘、光盘等；以及通信单元309，例如网卡、调制解调器、无线通信收发机等。通信单元309允许设备300通过诸如因特网的计算机网络和/或各种电信网络与其他设备交换信息/数据。

计算单元301可以是各种具有处理和计算能力的通用和/或专用处理组件。计算单元301的一些示例包括但不限于中央处理单元(CPU)、图形处理单元(GPU)、各种专用的人工智能(AI)计算芯片、各种运行机器学习模型算法的计算单元、数字信号处理器(DSP)、以及任何适当的处理器、控制器、微控制器等。计算单元301执行上文所描述的各个方法和处理，例如线程配置处理方法。例如，在一些实施例中，线程配置处理方法可被实现为计算机软件程序，其被有形地包含于机器可读介质，例如存储单元306。在一些实施例中，计算机程序的部分或者全部可以经由ROM 302和/或通信单元309而被载入和/或安装到设备300上。当计算机程序加载到RAM 603并由计算单元301执行时，可以执行上文描述的线程配置方法的一个或多个步骤。备选地，在其他实施例中，计算单元301可以通过其他任何适当的方式(例如，借助于固件)而被配置为执行线程配置处理方法。

本文中以上描述的系统和技术的各种实施方式可以在数字电子电路系统、集成电路系统、场可编程门阵列(FPGA)、专用集成电路(ASIC)、专用标准产品(ASSP)、芯片上系统的系统(SOC)、负载可编程逻辑设备(CPLD)、计算机硬件、固件、软件、和/或它们的组合中实现。这些各种实施方式可以包括：实施在一个或者多个计算机程序中，该一个或者多个计算机程序可在包括至少一个可编程处理器的可编程系统上执行和/或解释，该可编程处理器可以是专用或者通用可编程处理器，可以从存储系统、至少一个输入装置、和至少一个输出装置接收数据和指令，并且将数据和指令传输至该存储系统、该至少一个输入装置、和该至少一个输出装置。

用于实施本公开的方法的程序代码可以采用一个或多个编程语言的任何组合来编写。这些程序代码可以提供给通用计算机、专用计算机或其他可编程数据处理装置的处理器或控制器，使得程序代码当由处理器或控制器执行时使流程图和/或框图中所规定的功能/操作被实施。程序代码可以完全在机器上执行、部分地在机器上执行，作为独立软件包部分地在机器上执行且部分地在远程机器上执行或完全在远程机器或服务器上执行。

在本公开的上下文中，机器可读介质可以是有形的介质，其可以包含或存储以供指令执行系统、装置或设备使用或与指令执行系统、装置或设备结合地使用的程序。机器可读介质可以是机器可读信号介质或机器可读储存介质。机器可读介质可以包括但不限于电子的、磁性的、光学的、电磁的、红外的、或半导体系统、装置或设备，或者上述内容的任何合适组合。机器可读存储介质的更具体示例会包括基于一个或多个线的电气连接、便携式计算机盘、硬盘、随机存取存储器(RAM)、只读存储器(ROM)、可擦除可编程只读存储器(EPROM或快闪存储器)、光纤、便捷式紧凑盘只读存储器(CD-ROM)、光学储存设备、磁储存设备、或上述内容的任何合适组合。

为了提供与用户的交互，可以在计算机上实施此处描述的系统和技术，该计算机具有：用于向用户显示信息的显示装置(例如，CRT(阴极射线管)或者LCD(液晶显示器)监视器)；以及键盘和指向装置(例如，鼠标或者轨迹球)，用户可以通过该键盘和该指向装置来将输入提供给计算机。其它种类的装置还可以用于提供与用户的交互；例如，提供给用户的反馈可以是任何形式的传感反馈(例如，视觉反馈、听觉反馈、或者触觉反馈)；并且可以用任何形式(包括声输入、语音输入或者、触觉输入)来接收来自用户的输入。

可以将此处描述的系统和技术实施在包括后台部件的计算系统(例如，作为数据服务器)、或者包括中间件部件的计算系统(例如，应用服务器)、或者包括前端部件的计算系统(例如，具有图形用户界面或者网络浏览器的用户计算机，用户可以通过该图形用户界面或者该网络浏览器来与此处描述的系统和技术的实施方式交互)、或者包括这种后台部件、中间件部件、或者前端部件的任何组合的计算系统中。可以通过任何形式或者介质的数字数据通信(例如，通信网络)来将系统的部件相互连接。通信网络的示例包括：局域网(LAN)、广域网(WAN)、互联网和区块链网络。

计算机系统可以包括客户端和服务器。客户端和服务器一般远离彼此并且通常通过通信网络进行交互。通过在相应的计算机上运行并且彼此具有客户端-服务器关系的计算机程序来产生客户端和服务器的关系。服务器可以是云服务器，又称为云计算服务器或云主机，是云计算服务体系中的一项主机产品，以解决了传统物理主机与VPS服务("Virtual Private Server"，或简称"VPS")中，存在的管理难度大，业务扩展性弱的缺陷。服务器也可以为分布式系统的服务器，或者是结合了区块链的服务器。

在本公开上述各装置、设备、存储介质以及计算机程序产品实施例中，线程配置装置获取输入的第一维度和第二维度，以及需要计算的中间维度，其中，第二维度大于第一维度；根据第一维度和中间维度对线程块block之中的线程进行配置，以生成线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y；根据第二维度和线程块block在x方向上的第一线程数block.x和在y方向上的第二线程数block.y对线程栅格grid之中的线程进行配置，以生成线程栅格grid在x方向上的第三线程数grid.x和在y方向上的第四线程数grid.y，对线程块block之中的线程进行配置，提高访存效率。

应该理解，可以使用上面所示的各种形式的流程，重新排序、增加或删除步骤。例如，本发公开中记载的各步骤可以并行地执行也可以顺序地执行也可以不同的次序执行，只要能够实现本公开公开的技术方案所期望的结果，本文在此不进行限制。

上述具体实施方式，并不构成对本公开保护范围的限制。本领域技术人员应该明白的是，根据设计要求和其他因素，可以进行各种修改、组合、子组合和替代。任何在本公开的精神和原则之内所作的修改、等同替换和改进等，均应包含在本公开保护范围之内。

Claims (13)

1.一种线程配置方法，包括：

获取输入的第一维度和第二维度，以及需要计算的中间维度，其中，所述第二维度大于所述第一维度，所述第一维度是指最低维度，所述第二维度是指最高维度；

根据所述第一维度和所述中间维度对线程块之中的线程进行配置，以生成所述线程块在x方向上的第一线程数和在y方向上的第二线程数，其中，包括：根据所述第一维度确定线程块在x方向上的第五线程数；

根据所述中间维度和所述线程块在x方向上的第五线程数确定所述线程块在y方向上的第二线程数；

根据所述线程块在y方向上的第二线程数和所述第一维度更新所述线程块在x方向上的第五线程数，以生成所述线程块在x方向上的第一线程数；

根据所述第二维度以及所述线程块在x方向上的第一线程数和在y方向上的第二线程数对线程栅格之中的线程进行配置，以生成所述线程栅格在x方向上的第三线程数和在y方向上的第四线程数，其中，包括：

根据所述第一维度和所述线程块在x方向上的第一线程数确定在所述线程栅格在x方向上的第三线程数；

根据所述第二维度和所述线程栅格在x方向上的第三线程数确定所述线程栅格在y方向上的第四线程数。

2.如权利要求1所述的方法，其中，所述根据所述第一维度确定所述线程块在x方向上的第五线程数，包括：

获取线程束的大小；

将所述线程束的大小和所述第一维度之中的较小值作为所述线程块在x方向上的第五线程数。

3.如权利要求1所述的方法，其中，所述根据所述中间维度和所述线程块在x方向上的第五线程数确定所述线程块在y方向上的第二线程数，包括：

获取线程块中最大可配置线程数；

将所述中间维度与所述线程块中最大可配置线程数之中的较小值作为第一值；

将所述中间维度与所述线程块在x方向上的第五线程数之中的较小值作为第二值；

将所述第一值与所述第二值之中的较小值作为所述线程块在y方向上的第二线程数。

4.如权利要求1所述的方法，其中，所述根据所述线程块在y方向上的第二线程数和所述第一维度更新所述线程块在x方向上的第五线程数，以生成所述线程块在x方向上的第一线程数，包括：

将所述第一维度与所述线程块在x方向上的第五线程数之中的较小值作为第三值；

将所述第一维度与所述线程块在y方向上的第二线程数之中的较小值作为第四值；

根据所述第三值和第四值中的较小值，以及所述第一维度更新所述线程块在x方向上的第一线程数。

5.如权利要求1所述的方法，其中，所述根据所述第一维度和所述线程块在x方向上的第一线程数确定所述线程栅格在x方向上的第三线程数，包括：

获取所述线程栅格的可同时启动最大线程块数量；

将所述第一维度向上取整后的值与所述可同时启动最大线程块数量之中的较小值作为第五值；

将所述线程块在x方向上的第一线程数向上取整后的值与所述可同时启动最大线程块数量之中的较小值作为第六值；

将所述第五值和第六值中较小值作为所述线程栅格在x方向上的第三线程数。

6.如权利要求5所述的方法，其中，所述根据所述第二维度和所述线程栅格在x方向上的第三线程数确定所述线程栅格在y方向上的第四线程数，包括：

将所述可同时启动最大线程块数量向上取整后的值与所述第二维度之中的较小值为第七值；

将所述线程栅格在x方向上的第三线程数向上取整后的值与所述第二维度之中的较小值作为第八值；

将所述第七值和第八值中的较小值作为所述线程栅格在y方向上的第四线程数。

7.一种线程配置的装置，包括：

获取模块，用于获取输入的第一维度和第二维度，以及需要计算的中间维度，其中，所述第二维度大于所述第一维度，所述第一维度是指最低维度，所述第二维度是指最高维度；

第一生成模块，用于根据所述第一维度和所述中间维度对线程块之中的线程进行配置，以生成所述线程块在x方向上的第一线程数和在y方向上的第二线程数；

第二生成模块，用于根据所述第二维度以及所述在x方向上的第一线程数和在y方向上的第二线程数对线程栅格之中的线程进行配置，以生成所述线程栅格在x方向上的第三线程数和在y方向上的第四线程数；

所述第一生成模块包括：

第一确定子模块，用于根据所述第一维度确定线程块在x方向上的第五线程数；

第二确定子模块，用于根据所述中间维度和所述线程块在x方向上的第五线程数确定所述线程块在y方向上的第二线程数；

第一生成子模块，用于根据所述线程块在y方向上的第二线程数和所述第一维度更新所述线程块在x方向上的第五线程数，以生成所述线程块在x方向上的第一线程数；

所述第二生成模块，用于根据所述第一维度和所述线程块在x方向上的第一线程数确定线程栅格在x方向上的第三线程数，并根据所述第二维度和所述线程栅格在x方向上的第三线程数确定所述线程栅格在y方向上的第四线程数。

8.如权利要求7所述的装置，其中，所述第一确定子模块，用于获取线程束的大小，并将所述线程束的大小和所述第一维度之中的较小值作为所述线程块在x方向上的第五线程数。

9.如权利要求7所述的装置，其中，所述第二确定子模块，用于获取线程块中最大可配置线程数，并将所述中间维度与所述线程块中最大可配置线程数之中的较小值作为第一值，以及将所述中间维度与所述线程块在x方向上的第五线程数之中的较小值作为第二值，并将所述第一值与所述第二值之中的较小值作为所述线程块在y方向上的第二线程数。

10.如权利要求7所述的装置，其中，所述第二生成模块包括第三确定子模块，所述第三确定子模块用于获取所述线程栅格的可同时启动最大线程块数量，并将所述第一维度向上取整后的值与所述可同时启动最大线程块数量之中的较小值作为第五值，以及将所述线程块在x方向上的第一线程数向上取整后的值与所述可同时启动最大线程块数量之中的较小值作为第六值，最后将所述第五值和第六值中较小值作为所述线程栅格在x方向上的第三线程数。

11.如权利要求10所述的装置，其中，所述第二生成模块还包括第四确定子模块，所述第四确定子模块，用于将所述可同时启动最大线程块数量向上取整后的值与所述第二维度与之中的较小值为第七值，并将所述在线程栅格x方向上的第三线程数向上取整后的值与所述第二维度与之中的较小值作为第八值，以及将所述第七值和第八值中的较小值作为所述线程栅格在y方向上的第四线程数。

12.一种电子设备，包括：

至少一个处理器；以及

与所述至少一个处理器通信连接的存储器；其中，

所述存储器存储有可被所述至少一个处理器执行的指令，所述指令被所述至少一个处理器执行，以使所述至少一个处理器能够执行权利要求1-6中任一项所述的方法。

13.一种存储有计算机指令的非瞬时计算机可读存储介质，其中，所述计算机指令用于使所述计算机执行根据权利要求1-6中任一项所述的方法。