用于高吞吐量和低开销内核启动的压缩命令分组的制作方法

用于高吞吐量和低开销内核启动的压缩命令分组
1.相关申请的交叉引用
2.本技术要求于2020年12月23日提交的名称为“condensed command packet for high throughput and low overhead kernel launch”的未决的美国非临时专利申请17/133,574号的权益，其全部内容据此以引用方式并入本文。


背景技术：

3.许多高性能计算(hpc)应用程序(例如，kripke)包括在循环中多次启动的一系列内核(例如，“任务图”)。随着gpu执行时间的改进，启动每个内核所需的时间成为影响应用程序的总体性能的重要因素。换句话说，随着内核启动开销与内核执行时间的比率增加，启动开销成为影响应用程序性能的关键路径的越来越重要的部分。
附图说明
4.可以从以下描述中获得更详细的理解，通过示例结合附图给出，其中：
5.图1是可实现本公开的一个或多个特征的示例设备的框图；
6.图2是图1的设备的框图，示出了附加的细节；
7.图3是示出用于内核分组启动和执行的示例性过程的流程图；
8.图4是示出用于在示例性应用程序中执行的示例性内核的任务图；
9.图5是示出与处理参照图4描述的内核中的每一者相关联的示例性处理时间和开销时间分量的框图；
10.图6是示出使用示例性压缩内核调度分组的内核分组启动和执行的示例性过程的流程图；并且
11.图7是示出根据参照图6示出和描述的过程的与处理参照图4描述的内核中的每一者相关联的示例性处理时间和开销时间分量的框图。
具体实施方式
12.一些具体实施提供被配置为调度计算内核以供执行的内核代理。内核代理包括被配置为接收参考内核调度分组的电路。内核代理还包括被配置为处理参考内核调度分组以确定内核调度信息的电路。内核代理还包括被配置为存储内核调度信息的电路。内核代理还包括被配置为基于内核调度信息来调度内核的电路。
13.在一些具体实施中，内核代理包括：被配置为接收压缩内核调度分组的电路；被配置为处理压缩内核调度分组以检索所存储的内核调度信息的电路；以及被配置为基于所检索的内核调度信息来调度内核的电路。在一些具体实施中，内核代理包括：被配置为接收压缩内核调度分组的电路；被配置为处理压缩内核调度分组以检索内核调度信息以及确定差异信息的电路；被配置为基于该差异信息来修改所检索的内核调度信息的电路；以及被配置为基于经修改的所检索的内核调度信息来调度内核的电路。
14.在一些具体实施中，内核代理包括：被配置为接收压缩内核调度分组的电路；被配
置为处理压缩内核调度分组以检索所存储的内核调度信息并且检索所存储的第二内核调度信息的电路；以及被配置为基于所检索的内核调度信息来调度内核以及基于所检索的第二内核信息来调度第二内核的电路。在一些具体实施中，内核代理包括：被配置为接收压缩内核调度分组的电路；被配置为处理压缩内核调度分组以检索所存储的内核调度信息、确定第一差异信息、检索所存储的第二内核调度信息以及确定第二差异信息的电路；被配置为基于第一差异信息来修改所检索的内核调度信息的电路；被配置为基于第二差异信息来修改所检索的第二内核调度信息的电路；以及被配置为基于经修改的内核调度信息来调度第一内核以及基于经修改的第二内核调度信息来调度第二内核的电路。
15.在一些具体实施中，内核代理包括参考状态缓冲器，并且内核调度信息被存储在该参考状态缓冲器中。在一些具体实施中，内核代理包括暂存(scratch)随机存取存储器(ram)，并且内核代理将内核调度信息存储在该暂存ram中。在一些具体实施中，内核代理是图形处理单元(gpu)或包括gpu。在一些具体实施中，内核代理包括被配置为从主机处理器接收参考内核调度分组的电路。在一些具体实施中，参考内核调度分组包括架构排队语言(architected queuing language，aql)字段。
16.一些具体实施提供了用于调度计算内核以供执行的方法。参考内核调度分组由内核代理接收。参考内核调度分组由内核代理处理以确定内核调度信息。内核调度信息由内核代理存储。内核代理基于内核调度信息来调度内核。
17.在一些具体实施中，由内核代理接收压缩内核调度分组，由内核代理处理压缩内核调度分组以检索所存储的内核调度信息，并且由内核代理基于所检索的内核调度信息来调度内核。在一些具体实施中，由内核代理接收压缩内核调度分组，由内核代理处理压缩内核调度分组以检索内核调度信息以及确定差异信息，由内核代理基于该差异信息来修改所检索的内核调度信息；以及由内核代理基于经修改的所检索的内核调度信息来调度内核。
18.在一些具体实施中，由内核代理接收压缩内核调度分组，由内核代理处理压缩内核调度分组以检索所存储的内核调度信息以及检索所存储的第二内核调度信息，由内核代理基于所检索的内核调度信息来调度内核，以及由内核代理基于所检索的第二内核调度信息来调度第二内核。
19.在一些具体实施中，由内核代理接收压缩内核调度分组，由内核代理处理压缩内核调度分组以检索所存储的内核调度信息、确定第一差异信息、检索所存储的第二内核调度信息以及确定第二差异信息，基于第一差异信息来修改所检索的内核调度信息，基于第二差异信息来修改所检索的第二内核调度信息，基于经修改的内核调度信息来调度第一内核以及基于经修改的第二内核调度信息来调度第二内核。
20.在一些具体实施中，内核代理将内核调度信息存储在参考状态缓冲器中。在一些具体实施中，内核代理将内核调度信息存储在内核代理上的暂存随机存取存储器(ram)中。在一些具体实施中，内核代理是图形处理单元(gpu)或包括gpu。在一些具体实施中，从主机处理器接收参考内核调度分组。在一些具体实施中，参考内核调度分组包括架构排队语言(aql)字段。
21.图1是可实现本公开的一个或多个特征的示例设备100的框图。设备100可包括例如计算机、游戏设备、手持设备、机顶盒、电视、移动电话或平板计算机。设备100包括处理器102、存储器104、存储装置106、一个或多个输入设备108以及一个或多个输出设备110。设备
100还可任选地包括输入驱动器112和输出驱动器114。应当理解，设备100可包括图1中未示出的另外部件。
22.在各种另选方案中，处理器102包括中央处理单元(cpu)、图形处理单元(gpu)、位于同一管芯上的cpu和gpu、或一个或多个处理器核心，其中每个处理器核心可为cpu或gpu。在各种另选方案中，存储器104位于与处理器102相同的管芯上，或与处理器102分开定位。存储器104包括易失性或非易失性存储器，例如随机存取存储器(ram)、动态ram或高速缓存。
23.存储装置106包括固定或可移动存储装置，例如硬盘驱动器、固态驱动器、光盘或闪存驱动器。输入设备108包括但不限于键盘、小键盘、触摸屏、触控板、检测器、麦克风、加速度计、陀螺仪、生物扫描仪或网络连接(例如，用于发射和/或接收无线ieee 802信号的无线局域网卡)。输出设备110包括但不限于显示器、扬声器、打印机、触觉反馈设备、一个或多个灯、天线或网络连接(例如，用于发射和/或接收无线ieee 802信号的无线局域网卡)。
24.输入驱动器112与处理器102和输入设备108通信，并允许处理器102从输入设备108接收输入。输出驱动器114与处理器102和输出设备110通信，并允许处理器102向输出设备110发送输出。应注意，输入驱动器112和输出驱动器114是任选的部件，并且如果输入驱动器112和输出驱动器114不存在，则设备100将以相同方式操作。输出驱动器116包括联接到显示设备118的加速处理设备(“apd”)116。apd从处理器102接受计算命令和图形渲染命令，处理这些计算命令和图形渲染命令，并将像素输出提供给显示设备118进行显示。如下文所详述，apd 116包括根据单指令多数据(“simd”)范式来执行计算的一个或多个并行处理单元。因此，尽管这里将各种功能描述为由apd 116执行或与其结合执行，但在各种另选方案中，被描述为由apd 116执行的功能另外地或另选地由具有类似能力的其他计算设备执行，该其他计算设备不由主机处理器(例如，处理器102)驱动并且向显示设备118提供图形输出。例如，可以设想根据simd范式执行处理任务的任何处理系统都可执行本文所述的功能。另选地，设想不根据simd范式执行处理任务的计算系统执行本文所述的功能。
25.图2是设备100的框图，示出了涉及在apd 116上执行处理任务的附加细节。处理器102在系统存储器104中保持一个或多个控制逻辑模块以供处理器102执行。控制逻辑模块包括操作系统120、内核模式驱动器122和应用程序126。这些控制逻辑模块控制处理器102和apd 116的操作的各种特征。例如，操作系统120直接与硬件通信并为在处理器102上执行的其他软件提供到硬件的接口。内核模式驱动器122通过例如向在处理器102上执行的软件(例如，应用程序126)提供应用编程接口(“api”)来控制apd 116的操作，以访问apd 116的各种功能。内核模式驱动器122还包括即时编译器，该即时编译器编译程序以供apd 116的处理部件(诸如下文所详述的simd单元138)执行。
26.apd 116执行用于所选功能的命令和程序，诸如可适于并行处理的图形操作和非图形操作。apd 116可用于执行图形流水线操作，诸如像素操作、几何计算和基于从处理器102接收的命令将图像呈现给显示设备118。apd 116还基于从处理器102接收的命令来执行与图形操作不直接相关的计算处理操作，诸如与视频、物理模拟、计算流体动力学或其他任务相关的操作。
27.apd 116包括计算单元132，该计算单元包括根据simd范式以并行方式在处理器102的请求下执行操作的一个或多个simd单元138。simd范式是这样一种范式，其中多个处
理元件共用单个程序控制流单元和程序计数器并由此执行相同的程序，但能够执行具有不同数据的该程序。在一个示例中，每个simd单元138包括十六个通道，其中每个通道与simd单元138中的其他通道同时执行相同的指令，但可执行具有不同数据的该指令。如果不是所有通道都需要执行给定指令，则可通过预测来关闭通道。还可使用预测来执行具有发散控制流的程序。更具体地，对于具有条件分支或其中控制流基于由单个通道执行的计算的其他指令的程序，预测对应于当前未被执行的控制流路径的通道，并且不同控制流路径的串行执行可实现任意控制流。
28.计算单元132中的基本执行单元是工作项。每个工作项表示要在特定通道中并行执行的程序的单个实例化。可在单个simd处理单元138上作为“波前”同时执行工作项。一个或多个波前包括在一个“工作组”中，该“工作组”包括被指定执行相同程序的工作项的集合。可通过执行构成工作组的波前中的每一者来执行工作组。在另选方案中，波前在单个simd单元138上顺序地执行，或在不同simd单元138上部分地或完全地并行执行。波前可被视为可在单个simd单元138上同时执行的工作项的最大集合。因此，如果从处理器102接收的命令指示特定程序要被并行化到该程序不能在单个simd单元138上同时执行的程度，则该程序被划分成在两个或多个simd单元138上并行化或在同一simd单元138上串行化(或根据需要并行化和串行化)的波前。调度器136执行涉及调度不同计算单元132和simd单元138上的各种波前的操作。
29.由计算单元132提供的并行性适合图形相关操作，诸如像素值计算、顶点变换和其他图形操作。因此，在一些实例中，接受来自处理器102的图形处理命令的图形流水线134将计算任务提供给计算单元132以供并行执行。
30.计算单元132还用于执行不涉及图形或不作为图形流水线134的“正常”操作(例如，所执行的用以补充针对图形流水线134的操作执行的处理的自定义操作)的一部分而执行的计算任务。在处理器102上执行的应用程序126或其他软件将定义此类计算任务的程序发送到apd 116以供执行。
31.在一些hpc和其他应用程序中，主机处理器(例如，cpu)启动一个或多个处理器内核以在gpu或其他处理器上执行。执行内核(例如，在gpu的情况下，gpu内核)的gpu或其他处理器在一些上下文中称为内核代理。
32.通常，主机处理器通过使特定类型的命令分组入队以供内核代理处理，来启动内核以在内核代理上执行。这种类型的命令分组可被称为内核调度分组。例如，异构系统架构(hsa)标准为此目的规定了架构排队语言(aql)内核调度分组(称为hsa_kernel_dispatch_packet)。表1示出了示例性hsa_kernel_dispatch_packet。
33.表1
34.hsa_kernel_dispatch_packet{
35.unit8_t header＝
36.hsa_packet_type_kernel_dispatch；
37.unit8_t synch_scopes；
38.unit16_t setup；
39.unit16_t workgroup_size_x；
40.unit16_t workgroup_size_y；
41.unit16_t workgroup_size_z；
42.unit16_t reserved0；
43.unit32_t grid_size_x；
44.unit32_t grid_size_y；
45.unit32_t grid_size_z；
46.unit16_t private_segment_size；
47.unit32_t group_segment_size；
48.unit64_t kernel_object；
49.void*kernarg_address；
50.unit64_t reserved2；
51.hsa_signal_t completion_signal；
52.}；
53.该示例性内核调度分组的格式和字段是示例性的。应注意，其他具体实施使用其他格式和/或字段，并且/或者不是特定于aql的。在一些情况下，主机使内核调度分组入队到为内核代理指定的特定队列中。内核代理的分组处理器处理内核调度分组以确定内核执行信息(例如，调度和“清理”信息)。
54.在一些具体实施中，调度信息包括用于调度内核以在内核代理(在该示例中为gpu)上执行的信息。在表1的示例hsa_kernel_dispatch_packet中，同步范围(synch_scopes)、设置、工作组大小、网格大小、专用段大小、组段大小、内核对象和kernarg地址是调度信息的一部分。这些字段提供关于在gpu上启动工作之前要执行的获取操作的范围的信息(synch_scopes字段)、指示gpu线程在该内核中如何组织的gpu内核维度(设置字段)、gpu内核中的线程的数量(工作组和网格大小字段)、由该内核的gpu线程消耗的暂存和片上本地存储器的量(分别为专用段和组段大小)、gpu内核代码本身(代码对象)和gpu内核的自变量(kernarg_address)。这些字段是示例，并且在一些具体实施中，例如，内核调度分组包括不同的调度信息(例如，不同的字段，或更多或更少数量的字段)，具体取决于内核代理具体实施。
55.在一些具体实施中，清理信息包括用于在内核代理上的内核执行完成之后执行动作的信息。在表1的示例hsa_kernel_dispatch_packet中，synch_scopes和完成信号是清理信息的一部分。synch_scopes字段提供关于在gpu上完成工作之后要执行的释放操作的范围的信息。完成信号用于通知主机(例如，cpu)和/或等待该完成信号的其他代理关于工作的完成。
56.应注意，在该示例中，synch_scope字段提供调度信息和清理信息两者。例如，在执行内核之前获取存储器栅栏的范围是调度信息，并且在执行内核之后释放存储器栅栏的范围是清理信息。在一些具体实施中，在分开的字段中提供调度信息和清理信息。
57.在一些具体实施中，调度信息和清理信息来源于内核调度分组的字段，并且来源于该字段的调度信息和清理信息的结构是特定于具体实施的。
58.内核代理基于内核调度信息来调度内核以供执行，并且在内核执行完成之后基于清理信息来执行清理。这些步骤是示例性的，并且在其他具体实施中可包括子步骤、不同步骤、更多步骤或更少步骤。
59.通常，内核调度分组被入队和处理，并且内核被调度以供执行并且针对在应用程序中运行的每个内核被清理。在该示例性内核处理方法中，入队、分组处理和清理操作通常由命令处理器或内核代理的其他合适的分组处理硬件来执行，而内核执行通常由计算单元(例如，simd设备)或内核代理的其他主处理单元来执行。不管执行每个操作的硬件是什么硬件，执行入队、分组处理和清理操作所花费的时间都被认为是内核执行的开销。
60.因此，对于执行若干处理器内核的应用程序，应用程序运行时间将包括处理器内核中的每一者的内核执行时间和内核开销时间。另外，许多应用程序包括在循环中多次执行的一系列内核(例如，短时运行内核)。随着内核执行时间改进(即，变得更短)，与启动内核以供执行相关联的开销在总内核处理时间中所占的比例变得越来越大，并且对于应用程序的总性能变得越来越重要。
61.图3是示出用于内核分组启动和执行的示例性过程300的流程图。
62.在步骤302中，内核调度分组被入队以供内核代理处理。内核调度分组是hsa_kernel_dispatch_packet、这种分组的经修改版本(例如，如本文所述)，或用于支持内核启动和执行的任何其他合适的分组或信息。在一些具体实施中，内核调度分组被入队在对应于内核代理的队列中。在一些具体实施中，内核调度分组被诸如cpu的主机处理器入队以供内核代理处理。在一些具体实施中，内核代理是gpu、dsp、cpu或任何其他合适的处理设备，或包括它们。
63.在步骤304中，内核代理处理内核调度分组。在一些具体实施中，内核调度代理的分组处理器或其他分组处理电路处理内核调度分组。在其他具体实施中，内核代理的通用处理电路处理分组。在一些具体实施中，内核调度分组被处理以确定用于在内核代理上执行内核的信息。在一些具体实施中，该信息包括调度信息和清理信息。
64.在步骤306中，内核代理基于从内核调度分组处理的信息来调度内核以在内核代理(例如，gpu)上执行，并且内核执行直到完成为止。在内核执行完成的条件308下，在步骤310中执行清理操作。在一些具体实施中，清理操作由内核代理基于从内核调度分组处理的信息来执行。在应用程序未完成的条件312下，过程300从步骤302重复，其中使下一内核的内核调度分组入队。否则，过程300结束。
65.如从图3的示例可以看出，每次在内核代理上启动内核时，由于内核调度分组的入队和处理以及由于清理操作而产生开销。
66.图4是示出用于在示例性应用程序中执行的示例性内核的任务图400。虽然任务图400举例示出了kripke应用程序的典型内核，但该概念对于任何应用程序和内核的集合都是通用的。任务图400包括ltimes内核410、散射内核420、源内核430、lplustimes内核440、扫描内核450和群内核460。应注意，所述特定内核仅是示例性的，并且它们的特定名称和功能对于该示例而言并不重要。为了执行该应用程序，以所示的顺序启动和执行每个内核。在一些具体实施中，在已经启动和执行所有内核之后，再次启动和执行这些内核。例如，在kripke中，在一些情况下，以任务图中所示的顺序再次启动和执行内核，具体取决于由任务图的先前迭代产生的数据的收敛分析。
67.图6是示出根据参照图3示出和描述的过程300的与处理参照图4示出和描述的内核410、420、430、440、450、460中的每一者相关联的示例性处理时间和开销时间分量的框图。如图所示，每个内核包括由于使内核调度分组入队以及处理内核调度分组而导致的开
销时间、用于在内核代理上调度和执行内核的处理时间、以及用于清理操作的开销时间。所示的框示出了促成内核410、420、430、440、450、460的开销时间、处理时间、调度时间、执行时间和清理时间的操作，并且不旨在按比例绘制或暗示内核必须并行运行，尽管一些或所有内核实际上可以并行运行或在一些具体实施中可以重叠。
68.为了减少在应用程序执行期间的开销时间(诸如内核入队、分组处理和/或清理开销)，一些具体实施包括被配置用于存储与内核相关的信息(诸如调度、执行和/或清理信息)的分组。此类分组在本文中被称为参考内核调度分组。
69.在一些具体实施中，参考分组包括指示参考分组信息或从参考分组处理的信息将被存储在存储器中以供将来访问的信息。在一些具体实施中，参考分组包括对信息要被存储的位置的索引。在一些具体实施中，参考分组是内核调度分组的经修改版本。例如，表2示出了示例性经修改hsa_kernel_dispatch_packet，其中unit16_t reserved0字段被重新指定以包括参考号(uint16_t ref_num)。
70.表2
71.hsa_kernel_dispatch_packet{
72.unit8_t header＝
73.hsa_packet_type_kernel_dispatch；
74.unit8_t synch_scopes；
75.unit16_t setup；
76.unit16_t workgroup_size_x；
77.unit16_t workgroup_size_y；
78.unit16_t workgroup_size_z；
79.unit16_t ref_num；//参考号
80.unit32_t grid_size_x；
81.unit32_t grid_size_y；
82.unit32_t grid_size_z；
83.unit16_t private_segment_size；
84.unit32_t group_segment_size；
85.unit64_t kernel_object；
86.void*kernarg_address；
87.unit64_t reserved2；
88.hsa_signal_t completion_signal；
89.}；
90.该示例性参考调度分组的格式和字段是示例性的。应注意，其他具体实施使用其他格式和/或字段，并且/或者不是特定于aql的。在一些具体实施中，信息被存储在缓冲器中，该缓冲器可被称为参考状态缓冲器(rsb)。该rsb是任何合适的缓冲器，诸如内核代理上的暂存ram、内核代理的gpu存储器的区域或任何其他合适的存储器位置。在一些具体实施中，信息被存储在rsb的参考状态表(rst)中，例如由来自参考分组的参考号(例如，表2的示例性分组中的ref_num)索引。表3示出了示例性rst，其包括用于存储来自参考分组的信息的8个条目。
91.表3
[0092][0093]
在一些具体实施中，使用参考分组(例如，表2的经修改hsa_kernel_dispatch_packet)而非普通内核调度分组(例如，表1的hsa_kernel_dispatch_packet)以使用参照图3示出和描述的过程300来启动参照图4示出和描述的内核410、420、430、440、450、460，使得从每个参考内核调度分组处理的信息存储在rfb的rst(例如，表3的示例性rst)中。
[0094]
为了利用存储在rfb中的信息来减少应用程序执行期间的内核开销(例如，入队、启动分组处理和/或清理时间)，一些具体实施包括被配置用于调度多个内核的分组。此类分组在本文中被称为压缩内核调度分组。
[0095]
在一些具体实施中，压缩内核调度分组包括指示用于调度的内核的数量的信息、每个内核的参考信息(例如，存储在rfb中)的索引和/或每个内核的差异信息(例如，差异向量)。
[0096]
在一些具体实施中，用于调度的内核的数量基于由压缩内核调度分组所参考的信息来指示要启动的内核的数量。在一些具体实施中，差异信息指示一种或多种方式，在该一种或多种方式中，由压缩内核调度分组所参考的信息(例如，存储在rfb中的信息)应当被修改以根据压缩内核调度分组来调度内核(在本文中称为差异信息或“diff”)，或者由压缩内核调度分组所参考的信息不应当被修改以根据压缩内核调度分组来调度内核。
[0097]
例如，表4示出了示例性压缩内核调度分组格式：
[0098]
表4
[0099]
hsa_condensed_dispatch_packet{
[0100]
unit8_t header＝
[0101]
hsa_packet_type_condensed_dispatch；
[0102]
unit8_t num_kernels；
[0103]
unit16_t diff_values[31]；//62字节的diff信息；
[0104]
}；
[0105]
标头字段指定分组是压缩调度分组，并且该分组携带与每个调度的参考分组的diff。num_kernels字段指定该单个压缩调度分组调度的内核的数量。diff_values指定每个内核与它们各自的参考分组相比的diff。该示例性压缩调度分组的格式和字段是示例性的。应注意，其他具体实施使用其他格式和/或字段，并且/或者不是特定于aql的。
[0106]
例如，表5示出了用于表示与存储在rfb中的信息的差异(例如，“diff”信息)的示例性标头：
[0107]
表5
[0108]
struct diff_params{
[0109]
unsignedref_num:3；//参考号
[0110]
unsigneddiff_vector:13；//diff向量
[0111]
}；
[0112]
diff标头是指示内核与其参考分组的diff的前导码。diff标头是diff的前导码，其指示哪个参考表条目被用作diff的基线(即，在该示例中为ref_num)以及哪些字段是不同的(即，在该示例中为diff_vector)。在前导码之后，发送diff本身。换句话说，diff标头中的ref_num指定修改(即，“diffed”)哪个唯一参考分组信息(例如，存储其的rst的索引)以调度该内核。diff_vector指定与对应的参考分组信息不同的该调度的字段。因此，在该示例中，diff_vector中的13个位对应于参考aql分组中的13个字段，并且diff_vector中设置的位指示与参考分组信息相比，该调度的对应字段是不同的。如果在diff_vector中没有设置位，则这意味着该调度与参考分组信息相同。应注意，在其他具体实施中，压缩分组可直接发送存储在参考表中的参考信息的diff。在这种情况下，diff_vector指定表中参考信息中的字段，而不是参考aql分组中的字段。
[0113]
该示例性diff标头的格式和字段是示例性的。应注意，其他具体实施使用其他格式和/或字段，并且/或者不是特定于aql的。
[0114]
例如，表6示出了根据以上示例的示例性压缩分组(其中为了便于参考而添加了行编号)：
[0115]
表6
[0116]
1.condensed_pkt.header＝
[0117]
hsa_packet_type_condensed_dispatch；
[0118]
2.condensed_pkt.num_kernels＝2；//2个内核被压缩
[0119]
3.//ref_num＝4；diff仅用于完成信号(第12位)
[0120]
4.structdiff_paramsparam1＝{0x4,0x1000}
[0121]
5.//ref_num＝6；diff仅用于kernarg(第11位)
[0122]
6.structdiff_paramsparam2＝{0x6,0x0800}
[0123]
7.hsa_condensed_dispatch_packetcondensed_pkt；
[0124]
8.//第一内核编码
[0125]
9.condensed_pkt.diff[0]＝param1；//diff标头
[0126]
10.//完成信号将采用64位＝4个diff[]条目
[0127]
11.condensed_pkt.diff[1]＝0xdead；
[0128]
12.condensed_pkt.diff[2]＝0xbeef；
[0129]
13.condensed_pkt.diff[3]＝0xfeed；
[0130]
14.condensed_pkt.diff[4]＝0x0bad；
[0131]
15.//第二内核编码
[0132]
16.condensed_pkt.diff[5]＝param2；//diff标头
[0133]
17.//kernarg将采用64位＝4个diff[]条目
[0134]
18.condensed_pkt.diff[6]＝0x1234；
[0135]
19.condensed_pkt.diff[7]＝0x5678；
[0136]
20.condensed_pkt.diff[8]＝0xdeed；
[0137]
21.condensed_pkt.diff[1]＝0xface；
[0138]
在该示例中，行1将分组标头设置为hsa_packet_type_condensed_dispatch，从而指示此为压缩调度分组。行2设置num_kernels＝2，指示该压缩调度分组包括调度两个内核的信息。行4为第一调度创建diff_header，并将其标记为param1。diff标头的第一字段具有值＝4(十六进制记数法中的0x4)，指示第一调度正在使用来自参考分组#4(例如，通过索引4存储在参考表中)的信息用于其调度。diff标头的第二字段，即diff_vector，设置了第12位，其指示来自参考分组#4的第12字段应当被修改(即，“diffed”)以用于第一调度。第12字段是完成信号字段。该示例性压缩调度分组的格式和字段是示例性的。应注意，其他具体实施使用其他格式和/或字段，并且/或者不是特定于aql的。
[0139]
换种说法来说明该示例，param1指示第一调度类似于参考分组#4，不同的是其使用了不同的完成信号。类似地，param2在行6中被初始化，并且指示第二调度类似于参考分组#6，在第11字段中(即，内核args)除外。行9使用第一分组的diff_header(即，param1)来填充压缩分组的第一diff字段(diff[0])。接下来的4个diff字段(diff[1]到diff[4])使用第一调度的完成信号进行填充(行11到14)。该调度的完成信号不同于对应的参考分组，如对应的diff_header所指示的。类似地，在diff[5]中填充对应于第二调度的diff_header(行16)，并且在diff[6]到diff[9]中填充不同于其参考分组的第二调度的内核arg地址(行18到21)。
[0140]
图6是示出使用示例性压缩内核调度分组的内核分组启动、执行和清理的示例性过程600的流程图。
[0141]
在步骤602中，压缩内核调度分组被入队以供内核代理处理从而调度一个或多个内核。假设用于调度一个或多个内核的信息已经存储在例如rfb或其他合适的存储器中。在一些具体实施中，先前通过处理一个或多个内核中的每一者的参考内核调度分组而将信息存储在rfb中。
[0142]
在步骤604中，内核代理处理压缩内核调度分组。在一些具体实施中，内核调度代理的分组处理器或其他分组处理电路处理压缩内核调度分组。在其他具体实施中，内核代理的通用处理电路处理压缩内核调度分组。在一些具体实施中，压缩内核调度分组被处理以确定用于在内核代理上执行一个或多个内核的信息。在一些具体实施中，该信息包括调度信息和清理信息。在一些具体实施中，该信息被存储在rfb或其他合适的存储器位置中，并且由每个内核的压缩内核调度分组中的参考号(例如，ref_num)索引。在一些具体实施中，基于一个或多个内核的压缩内核调度分组中的差异信息(例如，diff_vector)来修改该信息。
[0143]
在步骤606中，内核代理基于从内核调度分组处理的信息(例如，包括从rfb检索的diff信息)来调度一个或多个内核中的第一内核，并且内核执行直到完成为止。在内核执行完成的条件608下，基于所处理的信息(例如，包括基于其从rfb检索的diff信息)，调度和执行下一内核(如果有的话)直到完成为止。在所有内核完成的条件610下，在步骤612中执行清理操作。在一些具体实施中，清理操作由内核代理基于从内核调度分组处理的信息来执行。在应用程序未完成的条件614下，过程600从步骤602重复，其中使另一内核调度分组入队(或进入不同的过程，例如，参照图3示出和描述的过程300，其中使标准内核调度分组或参考内核调度分组入队)。否则，过程600结束。
[0144]
如从图6的示例可以看出，对于由压缩内核调度分组在内核代理上启动的所有内核，由于压缩内核调度分组的入队和处理以及由于清理操作而产生一次开销。
[0145]
图7是示出根据参照图6示出和描述的过程600的与处理参照图4示出和描述的内核410、420、430、440、450、460中的每一者相关联的示例性处理时间和开销时间分量的框图。
[0146]
如图所示，仅第一内核410包括由于使内核调度分组入队以及处理内核调度分组而产生的处理时间，而内核410、420、430、440、450、460中的每一者包括用于在内核代理上处理内核的处理时间。最终分组460包括用于清理操作的处理时间。分组410、420、430、440、450包括或不包括用于清理操作的处理时间，具体取决于清理信息(由图中的虚线指示)。因此，所示的框示出了基于压缩内核调度分组的所有内核410、420、430、440、450、460的总处理时间小于(或至少包括更少的元素)基于常规或参考内核调度分组的所有内核410、420、430、440、450、460的总处理时间(例如，如参照图5所示和所述)。所示的框示出了促成内核410、420、430、440、450、460的处理时间的操作，并且不旨在按比例绘制或暗示内核必须并行运行，尽管一些或所有内核实际上可以并行运行或在一些具体实施中可以重叠。
[0147]
应当理解，基于本文的公开内容，可能有许多变化。尽管上述特征和元素在特定组合中进行了描述，但每个特征或元素可以在没有其他特征和元素的情况下单独使用，或者在有或没有其他特征或元素的各种组合中使用。
[0148]
图中和/或本文所述的各种功能单元(包括但不限于处理器102、输入驱动器112、输入设备108、输出驱动器114、输出设备110、加速处理设备116、调度器136、图形处理流水线134、计算单元132、simd单元138)可被实现为通用计算机、处理器或处理器核，或者实现为存储在非暂态计算机可读介质或另一介质中的可由通用计算机、处理器或处理器核执行的程序、软件或固件。所提供的方法可以在通用计算机、处理器或处理器核心中实现。举例来说，合适的处理器包括通用处理器、专用处理器、常规处理器、数字信号处理器(dsp)、多个微处理器、与dsp核相关联的一个或多个微处理器、控制器、微控制器、专用集成电路(asic)、现场可编程门阵列(fpga)电路，任何其他类型的集成电路(ic)和/或状态机。可以通过使用处理的硬件描述语言(hdl)指令和包括网表的其他中间数据(能够存储在计算机可读介质上的此类指令)的结果来配置制造过程而制造此类处理器。这种处理的结果可以是掩码，然后在半导体制造过程中使用这些掩码来制造实现本公开的特征的处理器。
[0149]
本文提供的方法或流程图可以在并入非暂态计算机可读存储介质中的计算机程序、软件或固件中实现，以供通用计算机或处理器执行。非暂态计算机可读存储介质的示例包括只读存储器(rom)、随机存取存储器(ram)、寄存器、高速缓存存储器、半导体存储器设备、磁性介质(例如内部硬盘和可移动磁盘)、磁光介质和光学介质(例如cd-rom磁盘)以及数字多功能磁盘(dvd)。

技术特征：
1.一种被配置为调度计算内核以供执行的内核代理，所述内核代理包括：被配置为接收参考内核调度分组的电路；被配置为处理所述参考内核调度分组以确定内核调度信息的电路；被配置为存储所述内核调度信息的电路；以及被配置为基于所述内核调度信息来调度内核的电路。2.根据权利要求1所述的内核代理，还包括：被配置为接收压缩内核调度分组的电路；被配置为处理所述压缩内核调度分组以检索所存储的内核调度信息的电路；以及被配置为基于所检索的内核调度信息来调度内核的电路。3.根据权利要求1所述的内核代理，还包括：被配置为接收压缩内核调度分组的电路；被配置为处理所述压缩内核调度分组以检索所述内核调度信息以及确定差异信息的电路；被配置为基于所述差异信息来修改所检索的内核调度信息的电路；以及被配置为基于经修改的所检索的内核调度信息来调度内核的电路。4.根据权利要求1所述的内核代理，还包括：被配置为接收压缩内核调度分组的电路；被配置为处理所述压缩内核调度分组以检索所存储的内核调度信息以及检索所存储的第二内核调度信息的电路；以及被配置为基于所检索的内核执行信息来调度内核以及基于所检索的第二内核信息来调度第二内核的电路。5.根据权利要求1所述的内核代理，还包括：被配置为接收压缩内核调度分组的电路；被配置为处理所述压缩内核调度分组以检索所存储的内核调度信息、确定第一差异信息、检索所存储的第二内核调度信息以及确定第二差异信息的电路；被配置为基于所述第一差异信息来修改所检索的内核调度信息的电路；被配置为基于所述第二差异信息来修改所检索的第二内核调度信息的电路；以及被配置为基于经修改的内核执行信息来调度第一内核以及基于经修改的第二内核信息来调度第二内核的电路。6.根据权利要求1所述的内核代理，还包括参考状态缓冲器，其中所述内核调度信息被存储在所述参考状态缓冲器中。7.根据权利要求1所述的内核代理，还包括暂存随机存取存储器(ram)，其中所述内核代理将所述内核调度信息存储在所述暂存ram中。8.根据权利要求1所述的内核代理，其中所述内核代理包括图形处理单元(gpu)。9.根据权利要求1所述的内核代理，还包括被配置为从主机处理器接收所述参考内核调度分组的电路。10.根据权利要求1所述的内核代理，其中所述参考内核调度分组包括架构排队语言(aql)字段。11.一种用于启动计算内核的方法，所述方法包括：
由内核代理接收参考内核调度分组；由所述内核代理处理所述参考内核调度分组以确定内核调度信息；由所述内核代理存储所述内核调度信息；以及基于所述内核调度信息来调度内核。12.根据权利要求11所述的方法，还包括：由所述内核代理接收压缩内核调度分组；由所述内核代理处理所述压缩内核调度分组以检索所存储的内核调度信息；以及基于所检索的内核调度信息来调度内核。13.根据权利要求11所述的方法，还包括：由所述内核代理接收压缩内核调度分组；由所述内核代理处理所述压缩内核调度分组以检索所述内核调度信息以及确定差异信息；基于所述差异信息来修改所检索的内核调度信息；以及基于经修改的所检索的内核调度信息来调度内核。14.根据权利要求11所述的方法，还包括：由所述内核代理接收压缩内核调度分组；由所述内核代理处理所述压缩内核调度分组，以检索所存储的内核调度信息以及检索所存储的第二内核调度信息；以及基于所检索的内核调度信息来调度内核以及基于所检索的第二调度信息来调度第二内核。15.根据权利要求11所述的方法，还包括：由所述内核代理接收压缩内核调度分组；由所述内核代理处理所述压缩内核调度分组以检索所存储的内核调度信息、确定第一差异信息、检索所存储的第二内核调度信息以及确定第二差异信息；基于所述第一差异信息来修改所检索的内核调度信息；基于所述第二差异信息来修改所检索的第二内核调度信息；以及基于经修改的内核执行信息来调度第一内核以及基于经修改的第二内核信息来调度第二内核。16.根据权利要求11所述的方法，其中所述内核代理将所述内核调度信息存储在参考状态缓冲器中。17.根据权利要求11所述的方法，其中所述内核代理将所述内核调度信息存储在所述内核代理上的暂存随机存取存储器(ram)中。18.根据权利要求11所述的方法，其中所述内核代理包括图形处理单元(gpu)。19.根据权利要求11所述的方法，其中所述内核代理从主机处理器接收所述参考内核调度分组。20.根据权利要求11所述的方法，其中所述参考内核调度分组包括架构排队语言(aql)字段。

技术总结
用于启动计算内核的方法、设备和系统。参考内核调度分组由内核代理接收。该参考内核调度分组由该内核代理处理以确定内核调度信息。该内核调度信息由该内核代理存储。该内核代理基于该内核调度信息来调度内核。在一些具体实施中，由该内核代理接收压缩内核调度分组，由该内核代理处理该压缩内核调度分组以检索所存储的内核调度信息，并且由该内核代理基于所检索的内核调度信息来调度内核。检索的内核调度信息来调度内核。检索的内核调度信息来调度内核。


技术研发人员：苏拉杰
受保护的技术使用者：超威半导体公司
技术研发日：2021.12.03
技术公布日：2023/8/24