使用重要性重采样来减少光采样的内存不一致性的制作方法

使用重要性重采样来减少光采样的内存不一致性
1.优先权要求
2.本技术要求于2020年8月21日提交的题为“基于储层的照明的重采样技术(resampling technique for reservior-based lighting)”的美国临时申请第63/068,906号的权益，其全部内容通过引用并入本文。
技术领域
3.至少一个实施例涉及计算机图形。例如，至少一个实施例涉及用于使用本文描述的各种新颖技术来渲染图形图像的处理器或计算系统。


背景技术：

4.在计算机图形学中的光处理会消耗大量的时间、内存、处理能力和其他计算资源。这对于旨在产生良好视觉质量的技术(其可包括但不限于光线追踪)，以及在要渲染的场景中包括许多光的情况，尤其如此。因此，可以改进计算机图形学中处理光的技术。
附图说明
5.图1示出了根据至少一个实施例的采用光重采样来渲染虚拟场景的系统的示例；
6.图2示出了根据至少一个实施例的通过至少从场景光列表中采样并从存储器部分重采样来渲染虚拟场景的帧的过程的示例；
7.图3示出了根据至少一个实施例的用于通过从存储器部分重采样来渲染虚拟场景的帧的过程的示例；
8.图4示出了根据至少一个实施例的虚拟区域中光的示例；
9.图5描绘了根据至少一个实施例的从场景光列表中随机选择光的示例；
10.图6描绘了根据至少一个实施例的基于对一个或更多个预采样的光子集的选择来渲染图块的示例；
11.图7是根据至少一个实施例的使用聚集方法的基于储层的时空重要性重采样(“restir”)候选选择的图示；
12.图8是根据至少一个实施例的使用散射方法的restir候选选择的图示；
13.图9是根据至少一个实施例的子池重组的图示；
14.图10是根据至少一个实施例的使用多个开放图块渲染帧的图示；
15.图11是根据至少一个实施例的包括预随机化和渲染时阶段的示例过程的图示；
16.图12示出了根据至少一个实施例的示例性数据中心；
17.图13示出了根据至少一个实施例的处理系统；
18.图14示出了根据至少一个实施例的计算机系统；
19.图15示出了根据至少一个实施例的系统；
20.图16示出了根据至少一个实施例的示例性集成电路；
21.图17示出了根据至少一个实施例的计算系统；
22.图18示出了根据至少一个实施例的apu；
23.图19示出了根据至少一个实施例的cpu；
24.图20示出了根据至少一个实施例的示例性加速器集成切片；
25.图21a和图21b示出了根据至少一个实施例的示例性图形处理器；
26.图22a示出了根据至少一个实施例的图形核心；
27.图22b示出了根据至少一个实施例的gpgpu；
28.图23a示出了根据至少一个实施例的并行处理器；
29.图23b示出了根据至少一个实施例的处理集群；
30.图23c示出了根据至少一个实施例的图形多处理器；
31.图24示出了根据至少一个实施例的图形处理器；
32.图25示出了根据至少一个实施例的处理器；
33.图26示出了根据至少一个实施例的处理器；
34.图27示出了根据至少一个实施例的图形处理器核心；
35.图28示出了根据至少一个实施例的ppu；
36.图29示出了根据至少一个实施例的gpc；
37.图30示出了根据至少一个实施例的流式多处理器；
38.图31示出了根据至少一个实施例的编程平台的软件栈；
39.图32示出了根据至少一个实施例的图31的软件栈的cuda实现；
40.图33示出了根据至少一个实施例的图31的软件栈的rocm实现；
41.图34示出了根据至少一个实施例的图31的软件栈的opencl实现；
42.图35示出了根据至少一个实施例的由编程平台支持的软件；
43.图36示出了根据至少一个实施例的在图31-34的编程平台上执行的编译代码；
44.图37示出了根据至少一个实施例的在图31-34的编程平台上执行的更详细的编译代码；
45.图38示出了根据至少一个实施例的在编译源代码之前转换源代码；
46.图39a示出了根据至少一个实施例的被配置为使用不同类型的处理单元来编译和执行cuda源代码的系统；
47.图39b示出了根据至少一个实施例的被配置为使用cpu和启用cuda的gpu来编译和执行图39a的cuda源代码的系统；
48.图39c示出了根据至少一个实施例的被配置为使用cpu和未启用cuda的gpu来编译和执行图39a的cuda源代码的系统；
49.图40示出了根据至少一个实施例的由图39c的cuda到hip转换工具转换的示例性内核；
50.图41更详细地示出了根据至少一个实施例的图39c的未启用cuda的gpu；
51.图42示出了根据至少一个实施例的示例性cuda网格的线程如何被映射到图41的不同计算单元；以及
52.图43图示了根据至少一个实施例如何将现有cuda代码迁移到数据并行c 代码。
具体实施方式
53.在以下描述中，阐述了许多具体细节，以提供对至少一个实施例的更透彻的理解。然而，对于本领域的技术人员将显而易见的是，可以在没有一个或更多个这些具体细节的情况下实践本发明构思。
54.在至少一个实施例中，用于渲染计算图形的方法的实施例，结合了诸如基于储层时空重要性重采样(restir)的基于储层的照明(reservoir-based lighting)技术，包括优化计算机硬件和存储器使用的适配，包括诸如存储器访问延迟、高速缓存一致性、高速缓存利用率、抖动(thrashing)等问题。
55.本文公开的实施例可用于各种应用、装置和环境，包括本文所描述的那些。本文所描述的技术可用于渲染复杂的图形场景，诸如可在视频游戏、特效、计算机动画、计算机辅助设计等中生成的那些场景。
56.在各个实施例中，本文描述的技术用于在需要高渲染速度的情况下渲染图形场景。例如，本文描述的一些操作和技术能够在渲染计算机图形帧之前执行，或者以小于每帧的周期性执行，以提高渲染的每帧效率。
57.在各个实施例中，本文描述的技术对于共享类似于渲染或光线追踪特征的非图形应用和问题空间是有用的。例如，本文描述的实施例可适用于模拟涉及许多发射器的声学或电磁传输的效果。
58.图1示出了根据至少一个实施例的采用光重采样来渲染虚拟场景的系统的示例。在图1的示例100中，计算设备102生成图形输出，以驱动在屏幕108上的显示。
59.在至少一个实施例中，计算设备102使用图形管线104和图形卡106生成图形输出。在至少一个实施例中，图形卡106包括一个或更多个处理器，诸如图形处理单元。在至少一个实施例中，图形管线104包括软件、硬件或软件和硬件的组合，以生成图形输出。图形管线104可以根据多阶段过程生成图形输出，诸如包括图1中图形管线104内描绘的阶段110-118的过程。尽管阶段110-118在图1中被描绘为序列，但是实施例可以省略一些描绘的阶段110-118，以所描述的顺序以外的顺序(诸如并行地)来执行操作110-118中的一些，或者包括除了图1中所描述的那些阶段以外的阶段。因此，图1中所描述的顺序不应被解释为将潜在的实施例限制为仅符合所描述顺序的那些实施例。
60.在至少一个实施例中，图形管线104包括软件、硬件或硬件和软件的组合，以实现用于将应用程序数据转换为合适的(有或没有某些后管线步骤)由屏幕108显示的图形数据的多阶段过程。例如，图形管线104可以生成视频数据帧，该视频数据帧随后可转换为信号，以驱动该帧在屏幕108上的显示。在至少一个实施例中，这些阶段可以包括应用程序阶段110、几何阶段112、变换阶段114、照明和着色阶段116以及光栅化和纹理化阶段118。
61.在至少一个实施例中，阶段110-118中的一个或更多个利用光采样算法(包括但不一定限于restir)，以将光效果结合到虚拟场景的渲染中。
62.在至少一个实施例中，虚拟场景120包括模拟或计算机生成的环境，诸如景观、建筑物、运动场或其他区域。虚拟场景120有时可被称为或包括虚拟环境。虚拟环境可以与定义虚拟环境的内容和结构的数据结构、图形资产和其他数据相关联。例如，在至少一个实施例中，虚拟场景基于虚拟环境，所述虚拟环境包括景观的线框模型、驻留在场景内的各种纹理和对象等。虚拟环境还可以包括放置在场景内不同位置的光。
63.在某些情况下，可能存在大量这样的光，这在渲染虚拟场景120的描绘时会带来许多挑战。处理许多光是计算机图形学中的一个难题，特别是对于基于光线追踪的算法。例如，渲染虚拟场景的一种方法是为每个着色点评估场景中的所有光源。然而，增加光计数还可能增加要追踪的光线的数量和复杂度，从而也可能增加渲染过程的时间、计算资源和复杂度。
64.在至少一个实施例中，从虚拟场景120中所有光的列表122中选择光子集。基于至少部分随机过程来选择子集，并且所选子集被存储在存储器124的一部分中。在至少一个实施例中，所选子集存储在对应于虚拟场景的细分或单元的数据结构记录中。在其他实施例中，所选子集不与虚拟场景的任何特定区域绑定或关联。在其他实施例中，根据与光强度或场景的整体重要性成比例的概率随机选择光。在其他实施例中，根据与给定光对虚拟场景的细分的贡献的重要性成比例的概率随机选择光。随机(stochastic)技术或过程，有时可称为随意(random)技术或过程，通常指包括随机、伪随机或准随机因素的技术。随机(stochastic)因素或随意(random)因素的示例可以包括但不一定限于伪随机数生成器、蒙特卡罗序列和确定性散列。
65.一旦存储在存储器124的部分中，所选光子集就可用于渲染虚拟场景的描绘中的像素。然而，从光列表122中选择光的过程可能会对高效渲染造成各种障碍。如所指出的，光选择过程至少部分是随机的，因此对光列表122的这种访问也可以至少部分是随机的，并且涉及对存储器广泛分离区域的访问。列表122也可能非常大，可能包括数千、数万甚至数百万光。因此，访问列表中选定的光会导致各种低效率，诸如涉及存储器访问延迟、高速缓存一致性、高速缓存利用率、抖动等的低效率。
66.在至少一个实施例中，光是指虚拟照明源。在至少一个实施例中，这可以包括发射或反射光的源。光可以与包括光在虚拟场景中的位置和强度值在内的属性相关联。例如，光可以与表示光在虚拟场景中的位置的x、y、z值以及指示光的亮度的值相关联。光也可与附加的属性相关联，诸如描述强度、颜色、扩散模式等的参数。如本文所使用的，术语光通常是指描述光的数据，例如表示诸如这些之类的特性和参数的数据。
67.在至少一个实施例中，图形管线104的阶段110-118利用存储在存储器部分124中的光来渲染虚拟场景的各部分。在至少一个实施例中，计算设备102通过首先从光列表122中随机选择光子集，并将该子集存储在存储器部分124中，来渲染图形帧。在至少一个实施例中，这是在渲染帧之前完成的。在渲染帧期间，计算设备102通过从存储器部分124随机选择一个或更多个光来渲染像素。注意，光的随机采样是指使用一个或更多个随机过程从光池中选择光。在至少一个实施例中，随机过程包括用于从包含随机性、伪随机性或准随机性的至少一些元素的池中选择光的任何技术。在至少一个实施例中，随机过程基于与光的强度成比例的概率来选择光，使得较亮的光比较暗的光更有可能被选择。光池是指作为候选以供选择的光。
68.图2示出了根据至少一个实施例的通过从场景光的列表中采样并从存储器部分重采样来渲染虚拟场景的帧过程的示例。
69.尽管示例过程200被描绘为一系列操作，但是应当理解，在实施例中，所描绘的操作可以以各种方式改变，并且一些操作可以被省略、重新排序，或者与其他操作并行地执行，除非明确说明或逻辑暗示顺序，诸如当一个操作的输入取决于另一个操作的输出时。
70.图2所描绘的操作可以由诸如图1中描绘的系统100之类的系统来执行，所述系统包括至少一个处理器和存储有指令的存储器，所述指令响应于由所述至少一个处理器执行，使该系统执行所描述的操作。在至少一个实施例中，通过硬件和软件的组合来执行操作，其中所述硬件包括一个或更多个apu、cpu、gpu、ppu、gpgpu、并行处理器、处理集群、图形处理器、多处理器等，如由本文中各个图所描绘的。在至少一个实施例中，所述软件包括诸如cuda、opengl、openlc、rocm中的任何一个之类的库，并且还可以包括操作系统软件。
71.在202处，在至少一个实施例中，系统为光列表122中的一个或更多个光建立概率密度函数。在至少一个实施例中，概率密度函数指示选择给定光的可能性。在至少一个实施例中，概率密度函数结合了各种参数的效果，诸如光强度、颜色、到要渲染的点的距离等，使得一些光(例如，那些最有可能显著地有助于照明的光)比其他光更有可能被选中。
72.在至少一个实施例中，使用概率密度函数的替代方案。在至少一个实施例中，随机选择光，选择给定光的几率相等。在至少一个实施例中，对于这种情况，可以使用随机数来生成到光列表中的索引。各种结构(诸如树或数组)可用来存储该列表，并且可用来促进结合一个或更多个随机过程来选择光，在一些实施例中，所述随机过程可包括使用随机数生成器。
73.在204处，在至少一个实施例中，系统从光的列表中选择光子集。在至少一个实施例中，使用一个或更多个随机过程来选择光，诸如刚刚关于概率密度函数描述的那些随机过程。为子集选择的光的数量可以在实施例之间基于诸如要渲染的部分的大小、诸如处理器高速缓存之类的存储器结构的大小等因素而变化。在至少一个实施例中，选择光集合和那些光的子集，以便照明信息可以存储在一个或更多个级别的处理器高速缓存中。
74.在206处，在至少一个实施例中，系统将光子集存储在存储器部分中。所述存储器部分可以是计算机存储器的一部分，诸如由随机存取存储器(“ram”)设备实现的存储器的区域，或者虚拟计算机存储器的区域。在至少一个实施例中，所述存储器部分是物理或虚拟计算机存储器的连续区域。在至少一个实施例中，所述存储器部分是位于处理器高速缓存内的存储器。在至少一个实施例中，使用高级数据结构(例如用编程语言定义的数组或链表)读取或写入光信息，将导致光子集被存储在存储器部分中。例如，从ram的一部分读取光信息可能导致光信息也被存储在高速缓存存储器部分中。
75.在208处，在至少一个实施例中，系统选择并存储光的附加子集。选择和加载的子集的数量可以在实施例之间变化。在至少一个实施例中，子集被选择，并被加载到存储器部分中，使得在给定帧的渲染期间，只要需要，该子集就可以保持有效地可访问(例如，在高速缓存中)。合适数量的子集通常可以通过实验或通过考虑在其上执行渲染的特定系统的存储器特性来找到，并且还可以取决于用于渲染帧内像素的算法的变化。例如，具有多个高速缓存的多处理器系统可以加载足够数量的子集，使得每个高速缓存包括具有一光子集的存储器部分。
76.在210处，在至少一个实施例中，基于从存储在存储器部分中的光重采样的光来渲染图形帧。在至少一个实施例中，将帧细分为图块，并且使用从存储器部分中采样的一个或更多个光来渲染每个图块。在至少一个实施例中，在渲染这些图块之后，使用被加载到单独的存储器部分的不同子集来渲染其他图块。然后可以重复此过程，直到整个帧被渲染。
77.图3示出了根据至少一个实施例的用于通过从存储器部分重采样来渲染虚拟场景
的帧的过程的示例。尽管示例过程300被描绘为一系列操作，但是应当理解，在实施例中，所描绘的操作可以以各种方式改变，并且一些操作可以被省略、重新排序，或者与其他操作并行地执行，除非明确说明或逻辑暗示顺序，诸如当一个操作的输入取决于另一个操作的输出时。
78.图3所示的操作可以由诸如图1中描绘的系统100之类的系统来执行，所述系统包括至少一个处理器和存储有指令的存储器，所述指令响应于由所述至少一个处理器执行，使系统执行所描述的操作。在至少一个实施例中，通过硬件和软件的组合来执行操作，其中所述硬件包括一个或更多个apu、cpu、gpu、ppu、gpgpu、并行处理器、处理集群、图形处理器、多处理器等，如由本文中各个图所描绘的。在至少一个实施例中，所述软件包括诸如cuda、opengl、openlc、rocm中的任何一个之类的库，并且还可以包括操作系统软件。
79.在至少一个实施例中，在302处，系统执行预帧处理，包括从所有场景光的列表中采样光，并将采样的光存储在一个或更多个存储器部分中。在至少一个实施例中，所述采样是从少于所有场景光中完成的，而是从包括大量光的某个池中完成的，这样，由于诸如可用处理器高速缓存的低效使用等问题，从其采样会导致性能问题。例如，在至少一个实施例中，池可能足够小以容纳在l3高速缓存中，但太大而不能容纳在更高效的l2高速缓存中。
80.在至少一个实施例中，预帧处理是指在渲染计算机生成的场景描绘的个体帧之前完成的处理。如本文所使用的，帧通常是指为产生动画序列而生成的一系列帧中的一个，但在某些实施例中，也可用于指代计算机生成的图像的单个实例。
81.在至少一个实施例中，在304处，系统选择要渲染的图块。在至少一个实施例中，帧被细分为图块，每个图块表示该帧的一部分。在至少一个实施例中，图块的大小或数量至少部分地基于以下因素，这些因素可以包括从光列表中抽取的样本数量，存储这些样本的存储器部分的数量，可用于渲染的处理器、gpu等的数量，可用线程数量等。
82.在至少一个实施例中，在306处，系统会选择将从其重采样光的存储器部分。在至少一个实施例中，这通过分配用于渲染图块的处理器或执行线程来完成，所述处理器或执行线程被配置为使得所述处理器或线程从所选存储器部分抽取样本。
83.在至少一个实施例中，在308处，系统通过从所识别出的存储器部分采样来渲染落入图块内的像素。如本文更详细地描述的，在至少一个实施例中，可以通过从存储在存储器部分中的样本随机重采样一个或更多个光，并且使用存储在存储器部分中的光信息来确定如何渲染像素，来渲染图块内的像素。例如，在至少一个实施例中，可以使用重采样的光来执行光线追踪。
84.在至少一个实施例中，在310处，系统确定是否需要渲染任何附加图块。如果是，则可以再次执行关于元素304至308描述的操作。在至少一个实施例中，针对每个渲染的图块使用一新的存储器部分。在至少一个实施例中，给定的存储器部分在第一组图块之间被重用，选择新的存储器部分以供第二组图块使用，等等。在至少一个实施例中，多组图块被并行渲染，并且在至少一个实施例中，每组图块可以使用不同的存储器部分。
85.在至少一个实施例中，在312处，系统输出渲染帧。在至少一个实施例中，这包括将完成的帧的数据提供给系统内的另一组件，诸如驱动显示器的组件。
86.在至少一个实施例中，通过使用附加的随机因素(包括随机、伪随机、准随机因素和/或决定性因素)来调整图块边缘附近的渲染像素，以减少或防止渲染图像中的伪影。例
如，抖动处理可以与上面的元素308结合使用，以便对于靠近图块边界的像素，使得从中选择光的存储器部分的集合基于一些随机化元素或噪声而变化。
87.图4示出了根据至少一个实施例的在虚拟区域中的光的示例。在图4的示例400中，虚拟区域402是从顶视图和侧视图描绘的三维区域。地形408被包括在虚拟区域402的所描绘的示例中，但是虚拟区域402可以包括或省略各种特征，例如所描绘的地形408，以及图中未描绘的其他特征，诸如人物、障碍物、墙壁和其他对象。
88.在至少一个实施例中，光410也包括在虚拟区域502中，位于区域402内的各个位置。这些光410发出照明，该照明可以被考虑到基于虚拟区域402的计算机生成的图像的渲染中。可以有许多这样的光410，可能以几百、几千或者几百万编号。在某些方法中，处理这些数量的光可能具有挑战性或不切实际。
89.图5描绘了根据至少一个实施例的从场景光列表中随机选择光的示例。在示例500中，场景506包括多个场景光508。可能存在非常多的场景光508，例如数百、数千或数百万光。这些光可以作为场景光列表504存储在存储器或存储装置中。
90.在至少一个实施例中，场景光列表504包括一个或更多个存储器阵列，其中存储了描述场景光508的信息。为了描述目的，描述场景光的信息在本文中可以被称为光信息、光数据或光。在至少一个实施例中，场景光列表包括诸如数组、链表、树、b-树等之类的数据结构。场景光列表504可以存储在随机存取存储器(“ram”)、长期存储器(例如固态或机械磁盘驱动器)中或一些其他结构中。应当理解，这些示例旨在是说明性的，因此不应以将潜在实施例限制为仅包括所提供的特定示例的那些实施例的方式来解释。
91.在至少一个实施例中，基于一个或更多个至少部分随机的过程从场景光列表504中识别采样的光502，该过程可以包括各种随机、准随机或伪随机因素。例如，在至少一个实施例中，场景光列表504具有存储在具有n个存储位置的数组a中的n个光。在该示例性实施例中，随机过程可以生成0和n-1之间的随机数，并获得对存储在a[n]处的光的访问。如在示例500中所描绘的，每次访问可以是对场景光列表504的不同部分的访问。如示例500所描绘的，对列表504的每次访问可以在不同的随机位置。应当理解，这些示例旨在是说明性的，因此不应以将潜在实施例限制为仅包括所提供的特定示例的那些实施例的方式来解释。
[0092]
在至少一个实施例中，restir算法依赖于随机化来生成其中存在许多光的图像，但是在一些情况下，这种随机化可能导致差的性能。然而，如本文所述，样本的预随机化的使用可以解决这些性能问题。例如，本文描述的技术使用每帧预处理技术来充分地置换(permute)或扰动样本，以维持(无偏差的)图像收敛。在至少一个实施例中，这些预随机化的样本被存储在一种数据结构中，所述数据结构能够以避免高速缓存存储器的低效使用的方式被访问。例如，通过适当调整大小，可以将数组或其他结构存储在处理器高速缓存中。预随机化样本可以通过将非一致的存储器访问移动到预处理阶段，而不是在每像素渲染期间引起非一致的存储器访问，来提供效率增益。如果这个预处理阶段比其他每像素候选生成技术更短或使用更少的存储器，则可以获得附加的性能。
[0093]
诸如restir之类的渲染技术可以使用重采样的重要性重采样(ris)的迭代应用来解耦计算频率：
[0094][0095]
这将积分∫f0(x)f1(x)f2(x)dx分解为在不同频率下评估的各项的总和。其某些实现可能会导致不一致问题和性能差。例如，在采用诸如restir之类的技术的实施例中，不一致性问题可能是由于样本xk在可能很长的列表上稀疏分布造成的。然而，在至少一个实施例中，可以在选择样本集合{xk}之前再次应用重采样的重要性采样(“ris”)以减少不一致性。
[0096]
通过检查ris的两种退化形式，可以进一步理解本文所述的技术。假设一标准的ris估计器：
[0097][0098]
然后检查两个退化情况，(即,以及)：
[0099][0100][0101]
第一个看起来像随机层上的分层抽样。在第二个中，仍然通过首先选择m个样本，然后选择它们的子集n，来应用ris。然而，它给出了相同的估计量，就好像直接采样n个项一样。请注意，元素可以被划掉。
[0102]
同样，诸如restir之类的技术可以采用类似于以下形式的内容：
[0103][0104]
这使用了不一致的光列表{l}，并首先将其采样为大小为n2的较小子集{xk}，然后将此集合子采样为大小为n1的(偶数)较小子集{xj}，然后将其子采样为大小为n0的子集{xi}。
[0105]
在至少一个实施例中，预随机化仍然可以使用根据p(x)分布，但是来自于存储器中更小、更一致的集合的样本。
[0106][0107]
这将{l}子采样为四个子集：{l}
→
{x
l
}
→
{xk}
→
{xj}
→
{xi}。内和被退化：
[0108][0109]
在至少一个实施例中，该数值估计：
[0110][0111]
可以通过取域{l}来计算，其中在至少一个实施例中，{l}是对应于场景光的一组发射三角形。在至少一个实施例中，然后根据分布p(x)从{l}中抽取n3个样本以获得{x
l
}。由于插入{x
l
}是根据p(x)，因此{x
l
}中的样本已经根据p(x)分布。这意味着从{x
l
}均匀抽取n2个样本以获得{xk}。然后，实施例继续从该样本集合{xk}中采样{xj}和{xi}。
[0112]
重采样的形式可以分类为分层的和未分层的。如上所述，具有嵌套总和的形式是一种分层方法：
[0113][0114]
这使用m*n个候选x
ij
。
[0115]
非分层形式可以表述为：
[0116][0117]
这可能只需要m个候选xj并且可以重复使用这些候选来抽取所有n个样本。在这种预随机化形式的ris中：
[0118][0119]
如果放弃某种程度的分层是可以接受的，那么实际上可以从总和中抽出预随机化的集合：
[0120][0121]
由于这是在没有额外的具有n3个样本的集合{x
l
}的情况下不会出现的分层，因此在一些实施例中这可能不是重要因素。
[0122]
在至少一个实施例中，使用多个集合{x
l
}但使用少于完全分层所需的数量(在这种情况下，其将是n0n1n2个不同的集合{x
l
})，还可以存在部分分层。并且无论放在什么地方，退化项都抵消掉了，这表明了灵活性。各个实施例可以基于这种“部分分层”在各种预随机化算法中如何发生，如本文所述。
[0123]
在示例实施例中，执行预随机化阶段和渲染时(render-time)阶段。图11中描绘了此类算法的示例说明。
[0124]
预随机化阶段可以包括：1)输入光样本列表l、要生成的子集数量si＝{x
l
}以及每个子集的大小k；和2)对于每个si子集，根据“初始光候选”概率密度函数p(x)从l中抽取k个光。例如，p(x)
∝
le(x)，即每个光的发射功率。
[0125]
在渲染时阶段，对于每个像素，不是使用分布p(x)从光样本的总列表{l}中抽取m个样本，而是从预随机化的光子集si＝{x
l
}之一中均匀地抽取m个样本。这种均匀抽样可以以分层和交错的方式完成，以确保si的遍历是最大一致性的。
[0126]
图6描绘了根据至少一个实施例的基于对光的一个或更多个预采样的子集的选择来渲染图块的示例。如示例600所示，可以将帧602细分为图块，并且每个图块可以被单独渲染。在至少一个实施例中，在预帧处理期间生成多个光子集。在至少一个实施例中，每个子集都被存储在单独的数据结构或存储器部分中。渲染算法的一个实施例可以包括步骤，在该步骤中，选择光子集以用于渲染图块。在渲染期间，每个像素从原始光列表{l}的某个子集si中重采样。为确保扭曲一致性，共享处理器高速缓存或其他存储器资源的活动线程应从同一子集si中重采样。
[0127]
一些图像图块中的所有像素都可以从同一个si中采样。例如，在示例600中，来自第一图块604的所有像素都基于从子集s1重采样的像素进行渲染，来自第二图块606的像素使用从s2重采样的像素进行渲染，来自第三图块608的像素使用从s3重采样的像素进行渲染，并且来自第四图块610的像素使用从s4重采样的像素进行渲染。
[0128]
在至少一个实施例中，目标图块大小是基于硬件配置来配置的。实验表明，在一些实施例中，8x8和16x16图块可以产生良好的性能。在这些实施例中，较大的图块引入了伪影，而使用4x4图像图块增加了非一致性，从而降低了性能。
[0129]
图7是用于对诸如restir之类的采样技术进行候选选择的可能方法的图示。在示例700中，使用了聚集方法。如示例700中所示，帧702由各个像素706、708组成，每个像素都基于从光列表402中随机抽取的光子集进行渲染。在该示例中，假设在渲染帧之前识别出该子集，并将其重新用于帧702内的各个像素。为了渲染像素706，从列表704执行从#1到#32的32次读取，以从所识别出的子集中的光710获得信息，但是存储在包括所有场景光的光列表704中(在本示例中400万个场景光)。随后，当渲染另一像素708时，使用相同的光子集，系统可再次执行读取#1到#32以从光列表704访问相同的光子集。然而，因为其他像素可能已经使用不同的光(例如，来自另一个子集的光)进行了渲染，来自原始子集的光不再在高速缓存中。因此，这种方法可能具有有问题的性能特征。
[0130]
图8是根据至少一个实施例的使用散射方法的restir候选选择的图示。如本文所述，光可以被预随机化为“子池”si，如示例800中所示。此外，示例800还示出实施例可以改变子池如何在屏幕上分布。在至少一个实施例中，屏幕图块可以从一个子池或更多个子池中抓取。在至少一个实施例中，这由重用参数定义。例如，在至少一个实施例中，如重用参数所示，子池中的每个光被重用某预定次数。在至少一个实施例中，图块大小基于所指示的重用参数而动态变化。
[0131]
在至少一个实施例中，从所有场景光的列表生成光804的一个或更多个子集。所有场景光的列表可能非常大，例如，在一个示例中，四百万个光。光804的一个或更多个子集的
数量和大小可以根据配置而变化。在一个示例中，光的一个或更多个子集包括多个子池，每个子池包括1024个光。
[0132]
在至少一个实施例中，光804的一个或更多个子集是通过从所有场景光中采样或混洗(shuffle)光来生成的。在至少一个实施例中，在采样或混洗中使用一种或更多种随机过程，例如随机数生成。
[0133]
在至少一个实施例中，一个或更多个子集中的部分或全部被存储在存储器部分812中。例如，在至少一个实施例中，如图8所示，包括1024个光的子池被存储在存储器部分812中。然后，该子池可用于通过从存储器部分812中的子池中抽取预随机化的样本，来渲染帧802的各个图块806、808内的像素。
[0134]
在至少一个实施例中，这种预随机化方法提供了各种优点。例如，如果注意促进子池si中光样本的一致性统计分布，那么子池的创建方式就可以有灵活性。例如，子池可以以更便宜的方式创建，例如在光子池之间混洗，如图9所示。这避免了不一致的存储器读入全局光列表{l}(在这个例子中，400万个条目长)。
[0135]
图9是根据至少一个实施例的子池重组(reshuffle)的图示。在至少一个实施例中，多个子池904-908被加载到一个或更多个存储器部分，例如l1/l2/dram 910的一个或更多个部分。可以在步骤1a、1b和1c读取这些子池，并且用于渲染帧的图块。在步骤2a和2b，子池904、906、908中的光被混洗(shuffled)。例如，在至少一个实施例中，基于一个或更多个随机过程，子池904中的一些光与另一子池906中的一些光交换。如示例900中所描述的，混洗可以随时间发生。在至少一个实施例中，对l1/l2/dram 910的读写操作可以重叠，如读取步骤1a、1b、1c和写入步骤3a、3b、3c沿着时间线902的重叠所示。
[0136]
图10是根据至少一个实施例的使用多个打开的(opened)图块渲染帧的图示。打开的图块可以指针对其从大量场景光中采样至少一个光子集的图块。通过将该子集加载到与保存场景光的存储器或存储装置分离的存储器部分中，可以使该子集在渲染中可用。当多个图块打开时，它们可共享单个子池，或者依赖于已加载到存储器部分的多个子池。
[0137]
在至少一个实施例中，多个图块一次全部打开，多个子池之间存在循环。例如，如图10所示，可以使用每像素四个样本将帧1002划分为t＝16个图块，每图块2x2像素。然后可以将当前子池用于来自t＝16个图块中的每个图块的一个样本。在本示例中，需要将16个子池加载到存储器部分1012中，以完成所有16个图块，因为每个图块使用2x2x4＝16个样本。在至少一个实施例中，每个子池服务每像素两个样本，从而使用32个子池。在至少一个实施例中，图块的数量增加到t＝32，同样使用32个子池。应当理解，这些示例旨在是说明性的，因此不应以将潜在实施例的范围限制为仅结合所提供的特定示例的那些实施例的方式来解释。在各个实施例中，可以将这些参数调整为最佳利用可用硬件的值。通常，可以通过实验和考虑硬件特性来确定合适的值，可能包括处理器高速缓存的大小，例如l1、l2和l3处理器高速缓存。
[0138]
在至少一个实施例中，本文描述的预随机化技术提供了各种附加的优势。对光进行采样可能是一项昂贵的操作，尤其是对于一组异构的光类型，例如发射三角形、球体、网格、平面、圆柱体等。simd处理器上的控制流发散也可能是一个重要的原因或性能下降。通过预随机化，这种昂贵的发散被移到对性能敏感的内部渲染循环之外。相反，它发生在渲染之前，每帧的次数要少得多。不同的光类型可以分成不同的光池si，或者可以使用一致的
simd控制流对光进行采样，然后作为第二步，随机化到子池中。
[0139]
在至少一个实施例中，本文描述的预随机化技术提高了处理动态光的效率。当光在场景中移动并改变强度时，可能需要更新光的当前位置、更新光的强度，并可能更新用于选择光的采样分布。通过预随机化光，这些更新可以在较少数量的光上完成。例如，在至少一个实施例中，只有那些在帧的选定子池si中的光被完全更新。
[0140]
在至少一个实施例中，本文描述的预随机化技术增加允许更大的灵活性。例如，使用图块可以灵活地改变光采样的形状和域。
[0141]
图11是根据至少一个实施例的包括预随机化和渲染时阶段的示例过程的图示。尽管示例过程1100被描绘为一系列操作，但是应当理解，在实施例中，所描绘的操作可以以各种方式改变，并且一些操作可以被省略、重新排序或与其他操作并行执行，除非明确说明或逻辑隐含顺序，例如当一个操作的输入取决于另一个操作的输出时。
[0142]
图11所描绘的操作可以由诸如图1中描绘的系统100之类的系统来执行，所述系统包括至少一个处理器和存储有指令的存储器，所述指令响应于由所述至少一个处理器执行，使系统执行所描述的操作。在至少一个实施例中，通过硬件和软件的组合来执行操作，其中所述硬件包括一个或更多个apu、cpu、gpu、ppu、gpgpu、并行处理器、处理集群、图形处理器、多处理器等，如由本文中各个图所示。在至少一个实施例中，所述软件包括诸如cuda、opengl、openlc、rocm中的任何一个之类的库，并且还可以包括操作系统软件。
[0143]
在1102，在至少一个实施例中，系统接收光样本列表l、要生成的多个子集si＝{xi}以及每个子集的大小k。
[0144]
在1104，在至少一个实施例中，系统根据初始光候选概率函数p(x)为每个si子集从l中抽取k个光。
[0145]
在至少一个实施例中，在预随机化阶段1110期间执行先前的操作1102、1104，在该阶段中，从光列表l中抽取子集si。在包括操作1106和1108的第二渲染时阶段1112期间，子集用于渲染图形帧。
[0146]
在1106，在至少一个实施例中，系统从预随机化的光子集si＝{xi}之一中均匀地抽取m个样本。如在1108中所示，系统可以以分层和交错的方式执行均匀采样，以在遍历si期间提高高速缓存一致性。图9描绘了以分层和交错方式执行的采样的示例。
[0147]
光子池可以与特定组的像素、纹理像素或体素相关联。光子池可用于选择更小的子池。这可对应于不一致性的分层减少，本质上是从光的随机、伪随机或准随机选择中构建随机数据结构。在至少一个实施例中，该随机数据结构包括样本的层级，每个层级包括由至少部分随机的过程从其下的层级选择的样本。如果在某些情况下使用子池和屏幕图块会导致条带伪影，则光子池可以动态调整大小，以进行性能和质量权衡。子池的数量和大小可能因用户设备而异，以控制各种不同能力的硬件类型的性能、存储器利用率和质量。
[0148]
数据中心
[0149]
图12示出了根据至少一个实施例的示例数据中心1200。在至少一个实施例中，数据中心1200包括但不限于数据中心基础设施层1210、框架层1220、软件层1230和应用层1240。
[0150]
在至少一个实施例中，如图12所示，数据中心基础设施层1210可以包括资源协调器1212、分组的计算资源1214和节点计算资源(“节点c.r.”)1216(1)-1216(n)，其中“n”代
表任何完整的正整数。在至少一个实施例中，节点c.r.1216(1)-1216(n)可以包括但不限于任意数量的中央处理单元(“cpu”)或其他处理器(包括加速器、现场可编程门阵列(“fpga”)、图形处理器等)，存储器设备(例如动态只读存储器)，存储设备(例如固态硬盘或磁盘驱动器)，网络输入/输出(“nwi/o”)设备，网络交换机，虚拟机(“vm”)，电源模块和冷却模块等。在至少一个实施例中，节点c.r.1216(1)-1216(n)中的一个或更多个节点c.r.可以是具有一个或更多个上述计算资源的服务器。
[0151]
在至少一个实施例中，分组的计算资源1214可以包括容纳在一个或更多个机架内的节点c.r.的单独分组(未示出)，或者容纳在各个地理位置的数据中心内的许多机架(也未示出)。分组的计算资源1214内的节点c.r.的单独分组可以包括可以被配置或分配为支持一个或更多个工作负载的分组的计算、网络、存储器或存储资源。在至少一个实施例中，可以将包括cpu或处理器的几个节点c.r.分组在一个或更多个机架内，以提供计算资源来支持一个或更多个工作负载。在至少一个实施例中，一个或更多个机架还可以包括任意数量的电源模块、冷却模块和网络交换机，以任意组合。
[0152]
在至少一个实施例中，资源协调器1212可以配置或以其他方式控制一个或更多个节点c.r.1216(1)-1216(n)和/或分组的计算资源1214。在至少一个实施例中，资源协调器1212可以包括用于数据中心1200的软件设计基础结构(“sdi”)管理实体。在至少一个实施例中，资源协调器1212可以包括硬件、软件或其某种组合。
[0153]
在至少一个实施例中，如图12所示，框架层1220包括但不限于作业调度器1232、配置管理器1234、资源管理器1236和分布式文件系统1238。
[0154]
在至少一个实施例中，框架层1220可以包括支持软件层1230的软件1252和/或应用程序层1240的一个或更多个应用程序1242的框架。在至少一个实施例中，软件1252或应用程序1242可以分别包括基于web的服务软件或应用程序，例如由amazonwebservices，googlecloud和microsoftazure提供的服务或应用程序。在至少一个实施例中，框架层1220可以是但不限于一种免费和开放源软件网络应用框架，例如可以利用分布式文件系统1238来进行大范围数据处理(例如“大数据”)的apachesparktm(以下称为“spark”)。在至少一个实施例中，作业调度器1232可以包括spark驱动器，以促进对数据中心1200的各个层所支持的工作负载进行调度。在至少一个实施例中，配置管理器1234可以能够配置不同的层，例如软件层1230和包括spark和用于支持大规模数据处理的分布式文件系统1238的框架层1220。在至少一个实施例中，资源管理器1236能够管理映射到或分配用于支持分布式文件系统1238和作业调度器1232的集群或分组计算资源。在至少一个实施例中，集群或分组计算资源可以包括数据中心基础设施层1210上的分组的计算资源1214。在至少一个实施例中，资源管理器1236可以与资源协调器1212协调以管理这些映射的或分配的计算资源。
[0155]
在至少一个实施例中，包括在软件层1230中的软件1252可以包括由节点c.r.1216(1)-1216(n)的至少一部分，分组计算资源1214和/或框架层1220的分布式文件系统1238使用的软件。一种或更多种类型的软件可以包括但不限于internet网页搜索软件、电子邮件病毒扫描软件、数据库软件和流视频内容软件。
[0156]
在至少一个实施例中，应用层1240中包括的一个或更多个应用程序1242可以包括由节点c.r.1216(1)-1216(n)的至少一部分、分组的计算资源1214和/或框架层1220的分布式文件系统1238使用的一种或更多种类型的应用程序。一种或更多种类型的应用程序可以
包括但不限于cuda应用程序。
[0157]
在至少一个实施例中，配置管理器1234、资源管理器1236和资源协调器1212中的任何一个可以基于以任何技术上可行的方式获取的任意数量和类型的数据来实现任意数量和类型的自我修改动作。在至少一个实施例中，自我修改动作可以减轻数据中心1200的数据中心操作员做出可能不好的配置决定并且可以避免数据中心的未充分利用和/或执行差的部分。
[0158]
基于计算机的系统
[0159]
以下各图提出但不限于可用于实现至少一个实施例的示例性的基于计算机的系统。
[0160]
图13示出了根据至少一个实施例的处理系统1300。在至少一个实施例中，系统1300包括一个或更多个处理器1302和一个或更多个图形处理器1308，并且可以是单处理器台式机系统、多处理器工作站系统或具有大量处理器1302或处理器核心1307的服务器系统。在至少一个实施例中，处理系统1300是结合在片上系统(soc)集成电路内的处理平台，以用于移动、手持或嵌入式设备。
[0161]
在至少一个实施例中，处理系统1300可以包括或结合在基于服务器的游戏平台中，包括游戏和媒体控制台的游戏控制台、移动游戏控制台、手持游戏控制台或在线游戏控制台。在至少一个实施例中，处理系统1300是移动电话、智能电话、平板计算设备或移动互联网设备。在至少一个实施例中，处理系统1300还可包括与可穿戴设备耦合或集成在可穿戴设备中，例如智能手表可穿戴设备、智能眼镜设备、增强现实设备或虚拟现实设备。在至少一个实施例中，处理系统1300是电视或机顶盒设备，其具有一个或更多个处理器1302以及由一个或更多个图形处理器1308生成的图形界面。
[0162]
在至少一个实施例中，一个或更多个处理器1302每个包括一个或更多个处理器核心1307，以处理指令，该指令在被执行时执行针对系统和用户软件的操作。在至少一个实施例中，一个或更多个处理器核心1307中的每一个被配置为处理特定指令集1309。在至少一个实施例中，指令集1309可以促进复杂指令集计算(cisc)、精简指令集计算(risc)，或通过超长指令字(vliw)进行计算。在至少一个实施例中，多个处理器核心1307可以各自处理不同的指令集1309，该指令集1309可以包括有助于仿真其他指令集的指令。在至少一个实施例中，处理器核心1307还可以包括其他处理设备，例如数字信号处理器(dsp)。
[0163]
在至少一个实施例中，处理器1302包括高速缓存存储器(cache)1304。在至少一个实施例中，处理器1302可以具有单个内部高速缓存或多个级别的内部高速缓存。在至少一个实施例中，高速缓存存储器在处理器1302的各个组件之间共享。在至少一个实施例中，处理器1302还使用外部高速缓存(例如，三级(l3)高速缓存或最后一级高速缓存(llc))(未示出)，其可以使用已知的高速缓存一致性技术在处理器核心1307之间共享该逻辑。在至少一个实施例中，处理器1302中另外包括寄存器文件1306，处理器1302可以包括用于存储不同类型的数据的不同类型的寄存器(例如，整数寄存器、浮点寄存器、状态寄存器和指令指针寄存器)。在至少一个实施例中，寄存器文件1306可以包括通用寄存器或其他寄存器。
[0164]
在至少一个实施例中，一个或更多个处理器1302与一个或更多个接口总线1310耦合，以在处理器1302与系统1300中的其他组件之间传输通信信号，例如地址、数据或控制信号。在至少一个实施例中，接口总线1310在一个实施例中可以是处理器总线，例如直接媒体
接口(dmi)总线的版本。在至少一个实施例中，接口总线1310不限于dmi总线，并且可以包括一个或更多个外围组件互连总线(例如，pci，pciexpress)、存储器总线或其他类型的接口总线。在至少一个实施例中，处理器1302包括集成存储器控制器1316和平台控制器集线器1330。在至少一个实施例中，存储器控制器1316促进存储设备与处理系统1300的其他组件之间的通信，而平台控制器集线器(pch)1330通过本地i/o总线提供到输入/输出(i/o)设备的连接。
[0165]
在至少一个实施例中，存储设备1320可以是动态随机存取存储器(dram)设备、静态随机存取存储器(sram)设备、闪存设备、相变存储设备或具有适当的性能以用作处理器存储器。在至少一个实施例中，存储设备1320可以用作处理系统1300的系统存储器，以存储数据1322和指令1321，以在一个或更多个处理器1302执行应用或过程时使用。在至少一个实施例中，存储器控制器1316还与可选的外部图形处理器1312耦合，其可以与处理器1302中的一个或更多个图形处理器1308通信以执行图和媒体操作。在至少一个实施例中，显示设备1311可以连接至处理器1302。在至少一个实施例中，显示设备1311可以包括内部显示设备中的一个或更多个，例如在移动电子设备或便携式计算机设备或通过显示器接口(例如显示端口(displayport)等)连接的外部显示设备。在至少一个实施例中，显示设备1311可以包括头戴式显示器(hmd)，诸如用于虚拟现实(vr)应用或增强现实(ar)应用中的立体显示设备。
[0166]
在至少一个实施例中，平台控制器集线器1330使外围设备能够通过高速i/o总线连接到存储设备1320和处理器1302。在至少一个实施例中，i/o外围设备包括但不限于音频控制器1346、网络控制器1334、固件接口1328、无线收发器1326、触摸传感器1325、数据存储设备1324(例如，硬盘驱动器、闪存等)。在至少一个实施例中，数据存储设备1324可以经由存储器接口(例如，sata)或经由外围总线来连接，诸如外围组件互连总线(例如，pci、pcie)。在至少一个实施例中，触摸传感器1325可以包括触摸屏传感器、压力传感器或指纹传感器。在至少一个实施例中，无线收发器1326可以是wi-fi收发器、蓝牙收发器或移动网络收发器，诸如3g、4g或长期演进(lte)收发器。在至少一个实施例中，固件接口1328使能与系统固件的通信，并且可以是例如统一的可扩展固件接口(uefi)。在至少一个实施例中，网络控制器1334可以启用到有线网络的网络连接。在至少一个实施例中，高性能网络控制器(未示出)与接口总线1310耦合。在至少一个实施例中，音频控制器1346是多通道高清晰度音频控制器。在至少一个实施例中，处理系统1300包括可选的传统(legacy)i/o控制器1340，用于将遗留(例如，个人系统2(ps/2))设备耦合到处理系统1300。在至少一个实施例中，平台控制器集线器1330还可以连接到一个或更多个通用串行总线(usb)控制器1342，该控制器连接输入设备，诸如键盘和鼠标1343组合、相机1344或其他usb输入设备。
[0167]
在至少一个实施例中，存储器控制器1316和平台控制器集线器1330的实例可以集成到离散的外部图形处理器中，例如外部图形处理器1312。在至少一个实施例中，平台控制器集线器1330和/或存储控制器1316可以在一个或更多个处理器1302的外部。例如，在至少一个实施例中，处理系统1300可以包括外部存储控制器1316和平台控制器集线器1330，其可以配置成在与处理器1302通信的系统芯片组中的存储器控制器集线器和外围控制器集线器。
[0168]
图14示出了根据至少一个实施例的计算机系统1400。在至少一个实施例中，计算
机系统1400可以是具有互连的设备和组件，soc，或某种组合的系统。在至少一个实施例中，计算机系统1400由处理器1402形成，该处理器1402可以包括用于执行指令的执行单元。在至少一个实施例中，计算机系统1400可以包括但不限于组件，例如处理器1402，其采用包括逻辑的执行单元以执行用于过程数据的算法。在至少一个实施例中，计算机系统1400可以包括处理器，例如可从加利福尼亚圣塔克拉拉的英特尔公司(intelcorporationofsantaclara,california)获得的处理器家族、xeontm、xscaletm和/或strongarmtm，core
tm
或nervana
tm
微处理器，尽管也可以使用其他系统(包括具有其他微处理器的pc、工程工作站、机顶盒等)。在至少一个实施例中，计算机系统1400可以执行可从华盛顿州雷蒙德市的微软公司(microsoftcorporationofredmond,wash.)获得的windows操作系统版本，尽管其他操作系统(例如unix和linux)、嵌入式软件和/或图形用户界面也可以使用。
[0169]
在至少一个实施例中，计算机系统1400可以用在其他设备中，例如手持设备和嵌入式应用。手持设备的一些示例包括蜂窝电话、互联网协议(internetprotocol)设备、数码相机、个人数字助理(“pda”)和手持pc。在至少一个实施例中，嵌入式应用可以包括微控制器、数字信号处理器(“dsp”)、soc、网络计算机(“netpc”)、机顶盒、网络集线器、广域网(“wan”)交换机，或根据至少一个实施例可以执行一个或更多个指令的任何其他系统。
[0170]
在至少一个实施例中，计算机系统1400可包括但不限于处理器1402，该处理器1402可包括但不限于一个或更多个执行单元1408，其可以配置为执行计算统一设备架构(“cuda”)(由加利福尼亚州圣克拉拉的nvidiacorporation开发)程序。在至少一个实施例中，cuda程序是用cuda编程语言编写的软件应用程序的至少一部分。在至少一个实施例中，计算机系统1400是单处理器台式机或服务器系统。在至少一个实施例中，计算机系统1400可以是多处理器系统。在至少一个实施例中，处理器1402可以包括但不限于cisc微处理器、risc微处理器、vliw微处理器、实现指令集组合的处理器，或任何其他处理器设备，例如数字信号处理器。在至少一个实施例中，处理器1402可以耦合到处理器总线1410，该处理器总线1410可以在处理器1402与计算机系统1400中的其他组件之间传输数据信号。
[0171]
在至少一个实施例中，处理器1402可以包括但不限于1级(“l1”)内部高速缓存存储器(“cache”)1404。在至少一个实施例中，处理器1402可以具有单个内部高速缓存或多级内部缓存。在至少一个实施例中，高速缓存存储器可以驻留在处理器1402的外部。在至少一个实施例中，处理器1402可以包括内部和外部高速缓存的组合。在至少一个实施例中，寄存器文件1406可以在各种寄存器中存储不同类型的数据，包括但不限于整数寄存器、浮点寄存器、状态寄存器和指令指针寄存器。
[0172]
在至少一个实施例中，包括但不限于执行整数和浮点运算的逻辑的执行单元1408，其也位于处理器1402中。处理器1402还可以包括微码(“ucode”)只读存储器(“rom”)，用于存储某些宏指令的微代码。在至少一个实施例中，执行单元1408可以包括用于处理封装指令集1409的逻辑。在至少一个实施例中，通过将封装指令集1409包括在通用处理器1402的指令集中，以及要执行指令的相关电路，可以使用通用处理器1402中的封装数据来执行许多多媒体应用程序使用的操作。在至少一个实施例中，可以通过使用处理器的数据总线的全宽度来在封装的数据上执行操作来加速和更有效地执行许多多媒体应用程序，这
可能不需要在处理器的数据总线上传输较小的数据单元来一次对一个数据元素执行一个或更多个操作。
[0173]
在至少一个实施例中，执行单元1408也可以用在微控制器、嵌入式处理器、图形设备、dsp和其他类型的逻辑电路中。在至少一个实施例中，计算机系统1400可以包括但不限于存储器1420。在至少一个实施例中，存储器1420可以被实现为dram设备、sram设备、闪存设备或其他存储设备。存储器1420可以存储由处理器1402可以执行的由数据信号表示的指令1419和/或数据1421。
[0174]
在至少一个实施例中，系统逻辑芯片可以耦合到处理器总线1410和存储器1420。在至少一个实施例中，系统逻辑芯片可以包括但不限于存储器控制器集线器(“mch”)1416，并且处理器1402可以经由处理器总线1410与mch 1416通信。在至少一个实施例中，mch 1416可以提供到存储器1420的高带宽存储器路径1418以用于指令和数据存储以及用于图形命令、数据和纹理的存储。在至少一个实施例中，mch 1416可以在处理器1402、存储器1420和计算机系统1400中的其他组件之间启动数据信号，并且在处理器总线1410、存储器1420和系统i/o 1422之间桥接数据信号。在至少一个实施例中，系统逻辑芯片可以提供用于耦合到图形控制器的图形端口。在至少一个实施例中，mch 1416可以通过高带宽存储器路径1418耦合到存储器1420，并且图形/视频卡1412可以通过加速图形端口(acceleratedgraphicsport)(“agp”)互连1414耦合到mch 1416。
[0175]
在至少一个实施例中，计算机系统1400可以使用系统i/o 1422作为专有集线器接口总线来将mch 1416耦合到i/o控制器集线器(“ich”)1430。在至少一个实施例中，ich 1430可以通过本地i/o总线提供与某些i/o设备的直接连接。在至少一个实施例中，本地i/o总线可以包括但不限于用于将外围设备连接到存储器1420、芯片组和处理器1402的高速i/o总线。示例可以包括但不限于音频控制器1429、固件集线器(“flashbios”)1428、无线收发器1426、数据存储1424、包含用户输入1425的传统i/o控制器1423和键盘接口、串行扩展端口1427(例如usb)和网络控制器1434。数据存储1424可以包括硬盘驱动器、软盘驱动器、cd-rom设备、闪存设备或其他大容量存储设备。
[0176]
在至少一个实施例中，图14示出了包括互连的硬件设备或“芯片”的系统。在至少一个实施例中，图14可以示出示例性soc。在至少一个实施例中，图14中示出的设备可以与专有互连、标准化互连(例如，pcie)或其某种组合互连。在至少一个实施例中，系统1400的一个或更多个组件使用计算快速链路(cxl)互连来互连。
[0177]
图15示出了根据至少一个实施例的系统1500。在至少一个实施例中，系统1500是利用处理器1510的电子设备。在至少一个实施例中，系统1500可以是，例如但不限于，笔记本电脑、塔式服务器、机架服务器、刀片服务器、膝上型计算机、台式机、平板电脑、移动设备、电话、嵌入式计算机或任何其他合适的电子设备。
[0178]
在至少一个实施例中，系统1500可以包括但不限于通信地耦合到任何合适数量或种类的组件、外围设备、模块或设备的处理器1510。在至少一个实施例中，处理器1510使用总线或接口耦合，诸如i2c总线、系统管理总线(“smbus”)、低引脚数(lpc)总线、串行外围接口(“spi”)、高清音频(“hda”)总线、串行高级技术附件(“sata”)总线、usb(1、2、3版)或通用异步接收器/发送器(“uart”)总线。在至少一个实施例中，图15示出了系统，该系统包括互连的硬件设备或“芯片”。在至少一个实施例中，图15可以示出示例性soc。在至少一个实施
例中，图15中所示的设备可以与专有互连线、标准化互连(例如，pcie)或其某种组合互连。在至少一个实施例中，图15的一个或更多个组件使用计算快速链路(cxl)互连线来互连。
[0179]
在至少一个实施例中，图15可以包括显示器1524、触摸屏1525、触摸板1530、近场通信单元(“nfc”)1545、传感器集线器1540、热传感器1546、快速芯片组(“ec”)1535、可信平台模块(“tpm”)1538、bios/固件/闪存(“bios，fwflash”)1522、dsp 1560、固态磁盘(“ssd”)或硬盘驱动器(“hdd”)1520、无线局域网单元(“wlan”)1550、蓝牙单元1552、无线广域网单元(“wwan”)1556、全球定位系统(gps)1555、相机(“usb3.0相机”)1554(例如usb3.0相机)或以例如lpddr3标准实现的低功耗双倍数据速率(“lpddr”)存储器单元(“lpddr3”)1515。这些组件可以各自以任何合适的方式实现。
[0180]
在至少一个实施例中，其他组件可以通过以上讨论的组件通信地耦合到处理器1510。在至少一个实施例中，加速度计1541、环境光传感器(“als”)1542、罗盘1543和陀螺仪1544可以可通信地耦合到传感器集线器1540。在至少一个实施例中，热传感器1539、风扇1537、键盘1536和触摸板1530可以通信地耦合到ec 1535。在至少一个实施例中，扬声器1563、耳机1564和麦克风(“mic”)1565可以通信地耦合到音频单元(“音频编解码器和d类放大器”)1562，其又可以通信地耦合到dsp 1560。在至少一个实施例中，音频单元1562可以包括例如但不限于音频编码器/解码器(“编解码器”)和d类放大器。在至少一个实施例中，sim卡(“sim”)1557可以通信地耦合到wwan单元1556。在至少一个实施例中，组件(诸如wlan单元1550和蓝牙单元1552以及wwan单元1556)可以被实现为下一代形式因素(ngff)。
[0181]
图16示出了根据至少一个实施例的示例性集成电路1600。在至少一个实施例中，示例性集成电路1600是soc，其可使用一个或更多个ip核心制造。在至少一个实施例中，集成电路1600包括一个或更多个应用处理器1605(例如，cpu)、至少一个图形处理器1610，并且可以另外包括图像处理器1615和/或视频处理器1620，其中任意一个可能是模块化ip核心。在至少一个实施例中，集成电路1600包括外围或总线逻辑，其包括usb控制器1625、uart控制器1630、spi/sdio控制器1635和i2s/i2c控制器1640。在至少一个实施例中，集成电路1600可以包括显示设备1645耦合到高清多媒体接口(hdmi)控制器1650和移动工业处理器接口(mipi)显示接口1655中的一个或更多个。在至少一个实施例中，存储可以由闪存子系统1660提供，包括闪存和闪存控制器。在至少一个实施例中，可以经由存储器控制器1665提供存储器接口以用于访问sdram或sram存储器设备。在至少一个实施例中，一些集成电路还包括嵌入式安全引擎1670。
[0182]
图17示出了根据至少一个实施例的计算系统1700。在至少一个实施例中，计算系统1700包括处理子系统1701，其具有经由可以包括存储器集线器1705的互连路径通信的一个或更多个处理器1702和系统存储器1704。在至少一个实施例中，存储器集线器1705可以是芯片组组件内的单独组件，也可以集成在一个或更多个处理器1702内。在至少一个实施例中，存储器集线器1705通过通信链路1706与i/o子系统1711耦合。在至少一个实施例中，i/o子系统1711包括i/o集线器1707，其可以使计算系统1700能够接收来自一个或更多个输入设备1708的输入。在至少一个实施例中，i/o集线器1707可以使能显示控制器，其包括在一个或更多个处理器1702中，用于向一个或更多个显示设备1710a提供输出。在至少一个实施例中，与i/o集线器1707耦合的一个或更多个显示设备1710a可以包括本地、内部或嵌入式显示设备。
[0183]
在至少一个实施例中，处理子系统1701包括经由总线或其他通信链路1713耦合到存储器集线器1705的一个或更多个并行处理器1712。在至少一个实施例中，通信链路1713可以是许多基于标准的通信链路技术或协议中的一种，例如但不限于pcie，或者可以是针对供应商的通信接口或通信结构。在至少一个实施例中，一个或更多个并行处理器1712形成计算集中的并行或向量处理系统，该系统可以包括大量的处理核心和/或处理集群，例如多集成核心(mic)处理器。在至少一个实施例中，一个或更多个并行处理器1712形成可以将像素输出到经由i/o集线器1707耦合的一个或更多个显示设备1710a之一的图形处理子系统。在至少一个实施例中，一个或更多个并行处理器1712还可以包括显示控制器和显示接口(未示出)，以使得能够直接连接到一个或更多个显示设备1710b。
[0184]
在至少一个实施例中，系统存储单元1714可以连接到i/o集线器1707，以提供用于计算系统1700的存储机制。在至少一个实施例中，i/o交换机1716可以用于提供接口机制，以实现i/o集线器1707与其他组件之间的连接，例如可以集成到平台中的网络适配器1718和/或无线网络适配器1719，以及可以通过一个或更多个附加设备1720添加的各种其他设备。在至少一个实施例中，网络适配器1718可以是以太网适配器或另一有线网络适配器。在至少一个实施例中，无线网络适配器1719可以包括wi-fi、蓝牙、nfc的一个或更多个或其他包括一个或更多个无线电的网络设备。
[0185]
在至少一个实施例中，计算系统1700可以包括未明确示出的其他组件，包括usb或其他端口连接、光存储驱动器、视频捕获设备等，也可以连接到i/o集线器1707。在至少一个实施例中，对图17中的各个组件进行互连的通信路径可以使用任何合适的协议来实现，诸如基于pci(外围组件互连)的协议(例如，pcie)，或其他总线或点对点通信接口和/或协议(例如，nvlink高速互连或互连协议)。
[0186]
在至少一个实施例中，一个或更多个并行处理器1712包括针对图形和视频处理而优化的电路(包括例如视频输出电路)，并构成图形处理单元(gpu)。在至少一个实施例中，一个或更多个并行处理器1712包括针对通用处理而优化的电路。在至少一个实施例中，计算系统1700的组件可以与单个集成电路上的一个或更多个其他系统元件集成。例如，在至少一个实施例中，一个或更多个并行处理器1712、存储器集线器1705、处理器1702和i/o集线器1707可以被集成到片上系统(soc)集成电路中。在至少一个实施例中，计算系统1700的组件可以被集成到单个封装中以形成系统级封装(sip)配置。在至少一个实施例中，计算系统1700的组件的至少一部分可以被集成到多芯片模块(mcm)中，该多芯片模块可以与其他多芯片模块互连到模块化计算系统中。在至少一个实施例中，从计算系统1700中省略了i/o子系统1711和显示设备1710b。
[0187]
处理系统
[0188]
以下各图阐述了但不限于可用于实现至少一个实施例的示例性处理系统。
[0189]
图18示出了根据至少一个实施例的加速处理单元(“apu”)1800。在至少一个实施例中，apu 1800由加利福尼亚州圣克拉拉市的amd公司开发。在至少一个实施例中，apu 1800可以被配置为执行应用程序，诸如cuda程序。在至少一个实施例中，apu 1800包括但不限于核心复合体1810、图形复合体1840、结构1860、i/o接口1870、存储器控制器1880、显示控制器1892和多媒体引擎1894。在至少一个实施例中，apu 1800可以包括但不限于任意数量的核心复合体1810、任意数量的图形复合体1840、任意数量的显示控制器1892和任意数
量的多媒体引擎1894的任何组合。为了说明的目的，在本文中用附图标记表示相似对象的多个实例，其中附图标记标识该对象，并且括号中的数字标识所需要的实例。
[0190]
在至少一个实施例中，核心复合体1810是cpu，图形复合体1840是gpu，并且apu 1800是将不限于1810和1840集成到单个芯片上的处理单元。在至少一个实施例中，一些任务可以被分配给核心复合体1810，而其他任务可以被分配给图形复合体1840。在至少一个实施例中，核心复合体1810被配置为执行与apu 1800相关联的主控制软件，例如操作系统。在至少一个实施例中，核心复合体1810是apu 1800的主处理器，其控制和协调其他处理器的操作。在至少一个实施例中，核心复合体1810发出控制图形复合体1840的操作的命令。在至少一个实施例中，核心复合体1810可以被配置为执行从cuda源代码派生的主机可执行代码，并且图形复合体1840可以被配置为执行从cuda源代码派生的设备可执行代码。
[0191]
在至少一个实施例中，核心复合体1810包括但不限于核心1820(1)-1820(4)和l3高速缓存1830。在至少一个实施例中，核心复合体1810可以包括但不限于任意数量的核心1820以及任意数量和类型的高速缓存的任何组合。在至少一个实施例中，核心1820被配置为执行特定指令集架构(“isa”)的指令。在至少一个实施例中，每个核心1820是cpu核心。
[0192]
在至少一个实施例中，每个核心1820包括但不限于获取/解码单元1822，整数执行引擎1824，浮点执行引擎1826和l2高速缓存1828。在至少一个实施例中，获取/解码单元1822获取指令，对这些指令进行解码，生成微操作，并将单独的微指令分派给整数执行引擎1824和浮点执行引擎1826。在至少一个实施例中，获取/解码单元1822可以同时分派一个微指令到整数执行引擎1824和另一微指令到浮点执行引擎1826。在至少一个实施例中，整数执行引擎1824执行不限于整数和存储器操作。在至少一个实施例中，浮点引擎1826执行不限于浮点和向量运算。在至少一个实施例中，获取-解码单元1822将微指令分派给单个执行引擎，该执行引擎代替整数执行引擎1824和浮点执行引擎1826两者。
[0193]
在至少一个实施例中，每个核心1820(i)可以访问包括在核心1820(i)中的l2高速缓存1828(i)，其中i是表示核心1820的特定实例的整数。在至少一个实施例中，包括在核心复合体1810(j)中的每个核心1820经由包括在核心复合体1810(j)中的l3高速缓存1830(j)连接到包括在核心复合体1810(j)中的其他核心1820，其中j是表示核心复合体1810的特定实例的整数。在至少一个实施例中，包括在核心复合体1810(j)中的核心1820可以访问包括在核心复合体1810(j)中的所有l3高速缓存1830(j)，其中j是表示核心复合体1810的特定实例的整数。在至少一个实施例中，l3高速缓存1830可以包括但不限于任意数量的切片(slice)。
[0194]
在至少一个实施例中，图形复合体1840可以被配置为以高度并行的方式执行计算操作。在至少一个实施例中，图形复合体1840被配置为执行图形管线操作，诸如绘制命令、像素操作、几何计算以及与将图像渲染至显示器相关联的其他操作。在至少一个实施例中，图形复合体1840被配置为执行与图形无关的操作。在至少一个实施例中，图形复合体1840被配置为执行与图形有关的操作和与图形无关的操作。
[0195]
在至少一个实施例中，图形复合体1840包括但不限于任意数量的计算单元1850和l2高速缓存1842。在至少一个实施例中，计算单元1850共享l2高速缓存1842。在至少一个实施例中，l2高速缓存1842被分区。在至少一个实施例中，图形复合体1840包括但不限于任意数量的计算单元1850以及任意数量(包括零)和类型的高速缓存。在至少一个实施例中，图
形复合体1840包括但不限于任意数量的专用图形硬件。
[0196]
在至少一个实施例中，每个计算单元1850包括但不限于任意数量的simd单元1852和共享存储器1854。在至少一个实施例中，每个simd单元1852实现simd架构并且被配置为并行执行操作。在至少一个实施例中，每个计算单元1850可以执行任意数量的线程块，但是每个线程块在单个计算单元1850上执行。在至少一个实施例中，线程块包括但不限于任意数量的执行线程。在至少一个实施例中，工作组是线程块。在至少一个实施例中，每个simd单元1852执行不同的线程束(warp)。在至少一个实施例中，线程束是一组线程(例如16个线程)，其中线程束中的每个线程属于单个线程块，并且被配置为基于单个指令集来处理不同的数据集。在至少一个实施例中，可以使用预测(predication)来禁用线程束中的一个或更多个线程。在至少一个实施例中，通道是线程。在至少一个实施例中，工作项是线程。在至少一个实施例中，波前是线程束。在至少一个实施例中，线程块中的不同波前可一起同步并经由共享存储器1854进行通信。
[0197]
在至少一个实施例中，结构1860是系统互连，其促进跨核心复合体1810、图形复合体1840、i/o接口1870、存储器控制器1880、显示控制器1892和多媒体引擎1894的数据和控制传输。在至少一个实施例中，除了结构1860之外或代替结构1860，apu 1800还可以包括但不限于任意数量和类型的系统互连，该结构1860促进跨可以在apu 1800内部或外部的任意数量和类型的直接或间接链接的组件的数据和控制传输。在至少一个实施例中，i/o接口1870表示任意数量和类型的i/o接口(例如，pci，pci-extended(“pci-x”)，pcie，千兆以太网(“gbe”)，usb等)。在至少一个实施例中，各种类型的外围设备耦合到i/o接口1870。在至少一个实施例中，耦合到i/o接口1870的外围设备可以包括但不限于键盘，鼠标，打印机，扫描仪，操纵杆或其他类型的游戏控制器、媒体记录设备、外部存储设备、网络接口卡等。
[0198]
在至少一个实施例中，显示控制器1892在一个或更多个显示设备(例如液晶显示器(lcd)设备)上显示图像。在至少一个实施例中，多媒体引擎1894包括但不限于任意数量和类型的与多媒体相关的电路，例如视频解码器、视频编码器、图像信号处理器等。在至少一个实施例中，存储器控制器1880促进apu 1800与统一系统存储器1890之间的数据传输。在至少一个实施例中，核心复合体1810和图形复合体1840共享统一系统存储器1890。
[0199]
在至少一个实施例中，apu 1800实现种存储器子系统，其包括但不限于任意数量和类型的存储器控制器1880和可以专用于一个组件或在多个组件之间共享的存储器设备(例如，共享存储器1854)。在至少一个实施例中，apu 1800实现高速缓存子系统，其包括但不限于一个或更多个高速缓存存储器(例如，l2高速缓存1928，l3高速缓存1830和l2高速缓存1842)，每个高速缓存存储器可以是组件私有的或在任意数量的组件(例如，核心1820，核心复合体1810，simd单元1852，计算单元1850和图形复合体1840)之间共享。
[0200]
图19示出了根据至少一个实施例的cpu 1900。在至少一个实施例中，cpu 1900由加利福尼亚州圣克拉拉市的amd公司开发。在至少一个实施例中，cpu 1900可以被配置为执行应用程序。在至少一个实施例中，cpu 1900被配置为执行主控制软件，例如操作系统。在至少一个实施例中，cpu 1900发出控制外部gpu(未示出)的操作的命令。在至少一个实施例中，cpu 1900可以被配置为执行从cuda源代码派生的主机可执行代码，并且外部gpu可以被配置为执行从这种cuda源代码派生的设备可执行代码。在至少一个实施例中，cpu 1900包括但不限于任意数量的核心复合体1910，结构1960，i/o接口1970和存储器控制器1980。
[0201]
在至少一个实施例中，核心复合体1910包括但不限于核心1920(1)-1920(4)和l3高速缓存1930。在至少一个实施例中，核心复合体1910可以包括但不限于任意数量的核心1920以及任意数量和类型的高速缓存的任何组合。在至少一个实施例中，核心1920被配置为执行特定isa的指令。在至少一个实施例中，每个核心1920是cpu核心。
[0202]
在至少一个实施例中，每个核心1920包括但不限于获取/解码单元1922，整数执行引擎1924，浮点执行引擎1926和l2高速缓存1928。在至少一个实施例中，获取/解码单元1922获取指令，对这些指令进行解码，生成微操作，并将单独的微指令分派给整数执行引擎1924和浮点执行引擎1926。在至少一个实施例中，获取/解码单元1922可以同时分派一个微指令至整数执行引擎1924和另一微指令至浮点执行引擎1926。在至少一个实施例中，整数执行引擎1924执行不限于整数和存储器操作。在至少一个实施例中，浮点引擎1926执行不限于浮点和向量运算。在至少一个实施例中，获取-解码单元1922将微指令分派给单个执行引擎，该引擎代替整数执行引擎1924和浮点执行引擎1926两者。
[0203]
在至少一个实施例中，每个核心1920(i)可以访问包括在核心1920(i)中的l2高速缓存1928(i)，其中i是表示核心1920的特定实例的整数。在至少一个实施例中，包括在核心复合体1910(j)中的每个核心1920经由包括在核心复合体1910(j)中的l3高速缓存1930(j)连接到核心复合体1910(j)中的其他核心1920，其中j是表示核心复合体1910的特定实例的整数。在至少一个实施例中，包括在核心复合体1910(j)中的核心1920可以访问包括在核心复合体1910(j)中的所有l3高速缓存1930(j)，其中j是表示核心复合体1910的特定实例的整数。在至少一个实施例中，l3高速缓存1930可以包括但不限于任意数量的切片。
[0204]
在至少一个实施例中，结构1960是系统互连，其促进跨核心复合体1910(1)-1910(n)(其中n是大于零的整数)、i/o接口1970和存储器控制器1980的数据和控制传输。在至少一个实施例中，除了结构1960之外或代替结构1960，cpu 1900还可以包括但不限于任意数量和类型的系统互连，该结构1960促进跨可以在cpu 1900内部或外部的任意数量和类型的直接或间接链接的组件的数据和控制传输。在至少一个实施例中，i/o接口1970表示任意数量和类型的i/o接口(例如pci，pci-x，pcie，gbe，usb等)。在至少一个实施例中，各种类型的外围设备耦合到i/o接口1970。在至少一个实施例中，耦合到i/o接口1970的外围设备可以包括但不限于显示器，键盘，鼠标，打印机，扫描仪，操纵杆或其他类型的游戏控制器、媒体记录设备、外部存储设备、网络接口卡等。
[0205]
在至少一个实施例中，存储器控制器1980促进cpu 1900与系统存储器1990之间的数据传输。在至少一个实施例中，核心复合体1910和图形复合体1940共享系统存储器1990。在至少一个实施例中，cpu 1900实现存储器子系统，其包括但不限于任意数量和类型的存储器控制器1980和可以专用于一个组件或在多个组件之间共享的存储器设备。在至少一个实施例中，cpu 1900实现了高速缓存子系统，其包括但不限于一个或更多个高速缓存存储器(例如，l2高速缓存1928和l3高速缓存1930)，每个高速缓存存储器可以是组件私有的或在任意数量的组件(例如，核心1920和核心复合体1910)之间共享。
[0206]
图20示出了根据至少一个实施例的示例性加速器集成切片2090。如本文所使用的，“切片”包括加速器集成电路的处理资源的指定部分。在至少一个实施例中，加速器集成电路代表多个图形加速模块种的多个图形处理引擎提供高速缓存管理、存储器访问、环境管理和中断管理服务。图形处理引擎可以各自包括单独的gpu。可选地，图形处理引擎可包
括gpu内的不同类型的图形处理引擎，例如图形执行单元、媒体处理引擎(例如，视频编码器/解码器)、采样器和blit引擎。在至少一个实施例中，图形加速模块可以是具有多个图形处理引擎的gpu。在至少一个实施例中，图形处理引擎可以是集成在通用封装、线卡或芯片上的各个gpu。
[0207]
系统存储器2014内的应用程序有效地址空间2082存储进程元素2083。在一个实施例中，响应于来自处理器2007上执行的应用程序2080的gpu调用2081而存储进程元素2083。进程元素2083包含对应应用程序2080的处理状态。包含在进程元素2083中的工作描述符(wd)2084可以是应用程序请求的单个作业或可能包含指向作业队列的指针。在至少一个实施例中，wd 2084是指向应用程序有效地址空间2082中的作业请求队列的指针。
[0208]
图形加速模块2046和/或各个图形处理引擎可以由系统中的全部或部分进程共享。在至少一个实施例中，可以包括用于建立处理状态并将wd 2084发送到图形加速模块2046以在虚拟化环境中开始作业的基础设施。
[0209]
在至少一个实施例中，专用进程编程模型是针对实现的。在该模型中，单个进程拥有图形加速模块2046或个体图形处理引擎。由于图形加速模块2046由单个进程拥有，因此管理程序为拥有的分区初始化加速器集成电路，并且当分配图形加速模块2046时操作系统对加速器集成电路进行初始化以用于拥有的分区。
[0210]
在操作中，加速器集成切片2090中的wd获取单元2091获取下一个wd 2084，其中包括要由图形加速模块2046的一个或更多个图形处理引擎完成的工作的指示。来自wd 2084的数据可以存储在寄存器2045被存储器管理单元(mmu)2039、中断管理电路2047和/或环境管理电路2048使用，如图所示。例如，mmu 2039的一个实施例包括用于访问os虚拟地址空间2085内的段/页表2086的段/页面漫游电路。中断管理电路2047可以处理从图形加速模块2046接收到的中断事件(int)2092。当执行图操作时，由图形处理引擎产生的有效地址2093由mmu 2039转换为实际地址。
[0211]
在一个实施例中，为每个图形处理引擎和/或图形加速模块2046复制相同的寄存器组2045，并且可以由系统管理程序或操作系统来初始化。这些复制的寄存器中的每一个都可以包含在加速器集成切片2090中。表1中显示了可由管理程序初始化的示例性寄存器。
[0212]
表1
–
管理程序初始化的寄存器
[0213]
1切片控制寄存器2实地址(ra)计划的处理区域指针3授权掩码覆盖寄存器4中断向量表输入偏移5中断向量表入口限制6状态寄存器7逻辑分区id8实地址(ra)管理程序加速器利用率记录指针9存储描述寄存器
[0214]
表2中示出了可以由操作系统初始化的示例性寄存器。
[0215]
表2
–
操作系统初始化寄存器
[0216][0217][0218]
在一个实施例中，每个wd 2084特定于特定的图形加速模块2046和/或特定图形处理引擎。它包含图形处理引擎进行工作或工作所需的所有信息，或者它可以是指向存储器位置的指针，其中应用程序建立了要完成的工作的命令队列。
[0219]
图21a和图21b示出了根据本文至少一个实施例的示例性图形处理器。在至少一个实施例中，任何示例性图形处理器可以使用一个或更多个ip核心来制造。除了图示之外，在至少一个实施例中可以包括其他逻辑和电路，包括附加的图形处理器/核心、外围接口控制器或通用处理器核心。在至少一个实施例中，示例性图形处理器用于soc内。
[0220]
图21a示出了根据至少一个实施例的soc集成电路的示例性图形处理器2110，其可以使用一个或更多个ip核心来制造。图21b示出了根据至少一个实施例的soc集成电路的的附加示例性图形处理器2140，其可以使用一个或更多个ip核心来制造。在至少一个实施例中，图21a的图形处理器2110是低功耗图形处理器核心。在至少一个实施例中，图21b的图形处理器2140是更高性能的图形处理器核心。在至少一个实施例中，每个图形处理器2110、2140可以是图16的图形处理器1610的变体。
[0221]
在至少一个实施例中，图形处理器2110包括顶点处理器2105和一个或更多个片段处理器2115a-2115n(例如2115a、2115b、2115c、2115d至2115n-1和2115n)。在至少一个实施例中，图形处理器2110可以经由单独的逻辑来执行不同的着色器程序，使得顶点处理器2105被优化以执行针对顶点着色器程序的操作，而一个或更多个片段处理器2115a-2115n执行片段(例如，像素)着色操作用于片段或像素或着色器程序。在至少一个实施例中，顶点处理器2105执行3d图形管线的顶点处理阶段并生成图元和顶点数据。在至少一个实施例中，片段处理器2115a-2115n使用由顶点处理器2105生成的图元和顶点数据来生成在显示设备上显示的帧缓冲区。在至少一个实施例中，片段处理器2115a-2115n被优化以执行如在openglapi中所提供的片段着色器程序，其可以用于执行与在direct3dapi中所提供的像素着色器程序类似的操作。
[0222]
在至少一个实施例中，图形处理器2110附加地包括一个或更多个mmu 2120a-2120b、高速缓存2125a-2125b和电路互连2130a-2130b。在至少一个实施例中，一个或更多个mmu 2120a-2120b提供用于图形处理器2110的虚拟到物理地址的映射，包括用于顶点处
理器2105和/或片段处理器2115a-2115n，其可以引用存储在存储器中的顶点或图像/纹理数据，除了存储在一个或更多个高速缓存2125a-2125b中的顶点或图像/纹理数据之外。在至少一个实施例中，一个或更多个mmu 2120a-2120b可以与系统内的其他mmu同步，包括与图16的一个或更多个应用处理器1605、图像处理器1615和/或视频处理器1620相关联的一个或更多个mmu，使得每个处理器1605-1620可以参与共享或统一的虚拟存储器系统。在至少一个实施例中，一个或更多个电路互连2130a-2130b使图形处理器2110能够经由soc的内部总线或经由直接连接与soc内的其他ip核心相连接。
[0223]
在至少一个实施例中，图形处理器2140包括图21a的图形处理器2110的一个或更多个mmu 2120a-2120b、高速缓存2125a-2125b和电路互连2130a-2130b。在至少一个实施例中，图形处理器2140包括一个或更多个着色器核心2155a-2155n(例如，2155a、2155b、2155c、2155d、2155e、2155f、至2155n-1和2155n)，其提供了统一的着色器核心架构，其中单个核心或类型或核心可以执行所有类型的可编程着色器代码，包括用于实现顶点着色器、片段着色器和/或计算着色器的着色器程序代码。在至少一个实施例中，多个着色器核心可以变化。在至少一个实施例中，图形处理器2140包括核心间任务管理器2145，其充当线程分派器以将执行线程分派给一个或更多个着色器核心2155a-2155n和分块单元2158，以加速基于图块渲染的分块操作，其中在图像空间中细分了场景的渲染操作，例如，以利用场景内的局部空间一致性或优化内部缓存的使用。
[0224]
图22a示出了根据至少一个实施例的图形核心2200。在至少一个实施例中，图形核心2200可以包括在图16的图形处理器1610内。在至少一个实施例中，图形核心2200可以是图21b中统一的着色器核心2155a-2155n。在至少一个实施例中，图形核心2200包括共享指令高速缓存2202、纹理单元2218和高速缓存/共享存储器2220，它们是图形核心2200内的执行资源所共有的。在至少一个实施例中，图形核心2200可以包括多个切片(slice)2201a-2201n或每个核心的分区，图形处理器可以包括图形核心2200的多个实例。切片2201a-2201n可以包括支持逻辑，该支持逻辑包括本地指令高速缓存2204a-2204n、线程调度器2206a-2206n、线程分派器2208a-2208n和一组寄存器2210a-2210n。在至少一个实施例中，切片2201a-2201n可以包括一组附加功能单元(afu)2212a-2212n、浮点单元(fpu)2214a-2214n、整数算术逻辑单元(alu)2216a-2216n、地址计算单元(acu)2213a-2213n、双精度浮点单元(dpfpu)2215a-2215n和矩阵处理单元(mpu)2217a-2217n。
[0225]
在一个实施例中，fpu 2214a-2214n可以执行单精度(32位)和半精度(16位)浮点运算，而dpfpu 2215a-2215n可以执行双精度(64位)浮点运算点操作。在至少一个实施例中，alu 2216a-2216n可以以8位、16位和32位精度执行可变精度整数运算，并且可以被配置用于混合精度运算。在至少一个实施例中，mpu 2217a-2217n还可被配置用于混合精度矩阵运算，包括半精度浮点运算和8位整数运算。在至少一个实施例中，mpu 2217a-2217n可以执行各种矩阵操作以加速cuda程序，包括使得能够支持加速的通用矩阵到矩阵乘法(gemm)。在至少一个实施例中，afu 2212a-2212n可以执行浮点数或整数单元不支持的附加逻辑运算，包括三角运算(例如，sine、cosine等)。
[0226]
图22b示出了在至少一个实施例中的通用图形处理单元(gpgpu)2230。在至少一个实施例中，gpgpu 2230是高度并行的并且适合于部署在多芯片模块上。在至少一个实施例中，gpgpu 2230可以被配置为使得高度并行的计算操作能够由gpu阵列来执行。在至少一个
实施例中，gpgpu 2230可以直接链路到gpgpu 2230的其他实例，以创建多gpu集群以提高用于cuda程序的执行时间。在至少一个实施例中，gpgpu 2230包括主机接口2232以实现与主机处理器的连接。在至少一个实施例中，主机接口2232是pcie接口。在至少一个实施例中，主机接口2232可以是厂商专用的通信接口或通信结构。在至少一个实施例中，gpgpu 2230从主机处理器接收命令，并使用全局调度器2234将与那些命令相关联的执行线程分派给一组计算集群2236a-2236h。在至少一个实施例中，计算集群2236a-2236h共享高速缓存存储器2238。在至少一个实施例中，高速缓存存储器2238可以用作计算集群2236a-2236h内的高速缓存存储器的高级高速缓存。
[0227]
在至少一个实施例中，gpgpu 2230包括经由一组存储器控制器2242a-2242b与计算集群2236a-2236h耦合的存储器2244a-2244b。在至少一个实施例中，存储器2244a-2244b可以包括各种类型的存储器设备，包括动态随机存取存储器(dram)或图形随机存取存储器，例如同步图形随机存取存储器(sgram)，包括图形双倍数据速率(gddr)存储器。
[0228]
在至少一个实施例中，计算集群2236a-2236h各自包括一组图形核心，诸如图22a的图形核心2200，其可以包括多种类型的整数和浮点逻辑单元，可以以各种精度执行计算操作，包括适合与cuda程序相关的计算。例如，在至少一个实施例中，每个计算集群2236a-2236h中的浮点单元的至少一个子集可以配置为执行16位或32位浮点运算，而不同的浮点单元的子集可以配置为执行64位浮点运算。
[0229]
在至少一个实施例中，gpgpu 2230的多个实例可以被配置为操作为计算集群。计算集群2236a-2236h可以实现用于同步和数据交换的任何技术上可行的通信技术。在至少一个实施例中，gpgpu 2230的多个实例通过主机接口2232进行通信。在至少一个实施例中，gpgpu 2230包括i/o集线器2239，其将gpgpu 2230与gpu链路2240耦合，使得能够直接连接至gpgpu 2230的其他的实例。在至少一个实施例中，gpu链路2240耦合到专用gpu到gpu桥接器，其使得能够在gpgpu 2230的多个实例之间进行通信和同步。在至少一个实施例中，gpu链路2240与高速互连耦合，以向其他gpgpu或并行处理器发送和接收数据。在至少一个实施例中，gpgpu 2230的多个实例位于单独的数据处理系统中，并经由可经由主机接口2232访问的网络设备进行通信。在至少一个实施例中，gpu链路2240可被配置为能够连接到主机处理器，附加或替代主机接口2232。在至少一个实施例中，gpgpu 2230可以配置为执行cuda程序。
[0230]
图23a示出了根据至少一个实施例的并行处理器2300。在至少一个实施例中，并行处理器2300的各种组件可以使用一个或更多个集成电路设备来实现，例如可编程处理器、专用集成电路(asic)或fpga。
[0231]
在至少一个实施例中，并行处理器2300包括并行处理单元2302。在至少一个实施例中，并行处理单元2302包括i/o单元2304，其使得能够与其他设备进行通信，包括并行处理单元2302的其他实例。在至少一个实施例中，i/o单元2304可以直接连接到其他设备。在至少一个实施例中，i/o单元2304通过使用集线器或交换机接口(例如，存储器集线器2305)与其他设备连接。在至少一个实施例中，存储器集线器2305与i/o单元2304之间的连接形成通信链路。在至少一个实施例中，i/o单元2304与主机接口2306和存储器交叉开关2316连接，其中主机接口2306接收用于执行处理操作的命令，而存储器交叉开关2316接收用于执行存储器操作的命令。
[0232]
在至少一个实施例中，当主机接口2306经由i/o单元2304接收命令缓冲区时，主机接口2306可以引导工作操作以执行那些命令到前端2308。在至少一个实施例中，前端2308与调度器2310耦合，调度器2310配置成将命令或其他工作项分配给处理阵列2312。在至少一个实施例中，调度器2310确保在将任务分配给处理阵列2312之前，处理阵列2312被正确地配置并且处于有效状态。在至少一个实施例中，调度器2310通过在微控制器上执行的固件逻辑来实现。在至少一个实施例中，微控制器实现的调度器2310可配置成以粗粒度和细粒度执行复杂的调度和工作分配操作，从而实现对在处理阵列2312上执行的线程的快速抢占和环境切换。在至少一个实施例中，主机软件可以证明用于通过多个图形处理门铃之一在处理阵列2312上进行调度的工作负载。在至少一个实施例中，工作负载然后可以由包括调度器2310的微控制器内的调度器2310逻辑在处理阵列2312上自动分配。
[0233]
在至少一个实施例中，处理阵列2312可以包括多达“n”个处理集群(例如，集群2314a、集群2314b到集群2314n)。在至少一个实施例中，处理阵列2312的每个集群2314a-2314n可以执行大量并发线程。在至少一个实施例中，调度器2310可以使用各种调度和/或工作分配算法将工作分配给处理阵列2312的集群2314a-2314n，其可以根据每种程序或计算类型产生的工作负载而变化。在至少一个实施例中，调度可以由调度器2310动态地处理，或者可以在配置为由处理阵列2312执行的程序逻辑的编译期间部分地由编译器逻辑来辅助。在至少一个实施例中，可将处理阵列2312的不同的集群2314a-2314n分配用于处理不同类型的程序或用于执行不同类型的计算。
[0234]
在至少一个实施例中，处理阵列2312可以配置成执行各种类型的并行处理操作。在至少一个实施例中，处理阵列2312配置成执行通用并行计算操作。例如，在至少一个实施例中，处理阵列2312可以包括执行处理任务的逻辑，该处理任务包括对视频和/或音频数据的过滤，执行建模操作，包括物理操作以及执行数据转换。
[0235]
在至少一个实施例中，处理阵列2312配置成执行并行图形处理操作。在至少一个实施例中，处理阵列2312可以包括附加逻辑以支持这种图形处理操作的执行，包括但不限于执行纹理操作的纹理采样逻辑，以及镶嵌逻辑和其他顶点处理逻辑。在至少一个实施例中，处理阵列2312可以配置成执行与图形处理有关的着色器程序，例如但不限于顶点着色器、曲面细分着色器、几何着色器和像素着色器。在至少一个实施例中，并行处理单元2302可以经由i/o单元2304从系统存储器传送数据以进行处理。在至少一个实施例中，在处理期间，可以在处理期间将传送的数据存储到片上存储器(例如，并行处理器存储器2322)，然后将其写回到系统存储器。
[0236]
在至少一个实施例中，当并行处理单元2302用于执行图处理时，调度器2310可以配置成将处理工作负载划分为近似相等大小的任务，以更好地将图形处理操作分配给处理阵列2312的多个集群2314a-2314n。在至少一个实施例中，处理阵列2312的部分可以配置成执行不同类型的处理。例如，在至少一个实施例中，第一部分可以配置成执行顶点着色和拓扑生成，第二部分可以配置成执行镶嵌和几何着色，并且第三部分可以配置成执行像素着色或其他屏幕空间操作，以生成用于显示的渲染图像。在至少一个实施例中，可以将由集群2314a-2314n中的一个或更多个产生的中间数据存储在缓冲区中，以允许在集群2314a-2314n之间传输中间数据以进行进一步处理。
[0237]
在至少一个实施例中，处理阵列2312可以经由调度器2310接收要执行的处理任
务，该调度器2310从前端2308接收定义处理任务的命令。在至少一个实施例中，处理任务可以包括要被处理的数据的索引，例如可以包括表面(补丁)数据、原始数据、顶点数据和/或像素数据，以及状态参数和定义如何处理数据的命令(例如，要执行什么程序)。在至少一个实施例中，调度器2310可以配置成获取与任务相对应的索引，或者可以从前端2308接收索引。在至少一个实施例中，前端2308可以配置成确保在启动由传入命令缓冲区(例如，批缓冲区(batch-buffer)、推送缓冲区等)指定的工作负载之前，处理阵列2312配置成有效状态。
[0238]
在至少一个实施例中，并行处理单元2302的一个或更多个实例中的每一个可以与并行处理器存储器2322耦合。在至少一个实施例中，可以经由存储器交叉开关2316访问并行处理器存储器2322，所述存储器交叉开关2316可以接收来自处理阵列2312以及i/o单元2304的存储器请求。在至少一个实施例中，存储器交叉开关2316可以经由存储器接口2318访问并行处理器存储器2322。在至少一个实施例中，存储器接口2318可以包括多个分区单元(例如，分区单元2320a、分区单元2320b到分区单元2320n)，其可各自耦合至并行处理器存储器2322的一部分(例如，存储器单元)。在至少一个实施例中，多个分区单元2320a-2320n为配置为等于存储器单元的数量，使得第一分区单元2320a具有对应的第一存储器单元2324a，第二分区单元2320b具有对应的存储器单元2324b，第n分区单元2320n具有对应的第n存储器单元2324n。在至少一个实施例中，分区单元2320a-2320n的数量可以不等于存储器设备的数量。
[0239]
在至少一个实施例中，存储器单元2324a-2324n可以包括各种类型的存储器设备，包括动态随机存取存储器(dram)或图形随机存取存储器，例如同步图形随机存取存储器(sgram)，包括图形双倍数据速率(gddr)存储器。在至少一个实施例中，存储器单元2324a-2324n还可包括3d堆叠存储器，包括但不限于高带宽存储器(hbm)。在至少一个实施例中，可以跨存储器单元2324a-2324n来存储诸如帧缓冲区或纹理映射的渲染目标，从而允许分区单元2320a-2320n并行地写入每个渲染目标的部分，以有效地使用并行处理器存储器2322的可用带宽。在至少一个实施例中，可以排除并行处理器存储器2322的本地实例，以有利于利用系统存储器与本地高速缓存存储器结合的统一存储器设计。
[0240]
在至少一个实施例中，处理阵列2312的集群2314a-2314n中的任何一个都可以处理将被写入并行处理器存储器2322内的任何存储器单元2324a-2324n中的数据。在至少一个实施例中，存储器交叉开关2316可以配置为将每个集群2314a-2314n的输出传输到任何分区单元2320a-2320n或另一个集群2314a-2314n，集群2314a-2314n可以对输出执行其他处理操作。在至少一个实施例中，每个集群2314a-2314n可以通过存储器交叉开关2316与存储器接口2318通信，以从各种外部存储设备读取或写入各种外部存储设备。在至少一个实施例中，存储器交叉开关2316具有到存储器接口2318的连接以与i/o单元2304通信，以及到并行处理器存储器2322的本地实例的连接，从而使不同处理集群2314a-2314n内的处理单元与系统存储器或不是并行处理单元2302本地的其他存储器进行通信。在至少一个实施例中，存储器交叉开关2316可以使用虚拟通道来分离集群2314a-2314n和分区单元2320a-2320n之间的业务流。
[0241]
在至少一个实施例中，可以在单个插入卡上提供并行处理单元2302的多个实例，或者可以将多个插入卡互连。在至少一个实施例中，并行处理单元2302的不同实例可以配
置成相互操作，即使不同实例具有不同数量的处理核心，不同数量的本地并行处理器存储器和/或其他配置差异。例如，在至少一个实施例中，并行处理单元2302的一些实例可以包括相对于其他实例而言更高精度的浮点单元。在至少一个实施例中，结合并行处理单元2302或并行处理器2300的一个或更多个实例的系统可以以各种配置和形式因素来实现，包括但不限于台式机、膝上型计算机或手持式个人计算机、服务器、工作站、游戏机和/或嵌入式系统。
[0242]
图23b示出了根据至少一个实施例的处理集群2394。在至少一个实施例中，处理集群2394被包括在并行处理单元内。在至少一个实施例中，处理集群2394是图23a的处理集群2314a-2314n之一的实例。在至少一个实施例中，处理集群2394可以配置成并行执行许多线程，其中术语“线程”是指在特定的一组输入数据上执行的特定程序的实例。在至少一个实施例中，单指令多数据(simd)指令发布技术用于支持大量线程的并行执行而无需提供多个独立的指令单元。在至少一个实施例中，使用单指令多线程(simt)技术来支持并行执行大量一般同步的线程，这使用了公共指令单元，该公共指令单元配置成向每个处理集群2394内的一组处理引擎发出指令。
[0243]
在至少一个实施例中，可以通过将处理任务分配给simt并行处理器的管线管理器2332来控制处理集群2394的操作。在至少一个实施例中，管线管理器2332从图23a的调度器2310接收指令，通过图形多处理器2334和/或纹理单元2336管理这些指令的执行。在至少一个实施例中，图形多处理器2334是simt并行处理器的示例性实例。然而，在至少一个实施例中，处理集群2394内可以包括不同架构的各种类型的simt并行处理器。在至少一个实施例中，在处理集群2394内可以包括图形多处理器2334的一个或更多个实例。在至少一个实施例中，图形多处理器2334可以处理数据，并且数据交叉开关2340可以用于将处理后的数据分发到多个可能的目的(包括其他着色器单元)地之一。在至少一个实施例中，管线管理器2332可以通过指定要经由数据交叉开关2340分配的处理后的数据的目的地来促进处理后的数据的分配。
[0244]
在至少一个实施例中，处理集群2394内的每个图形多处理器2334可以包括相同的一组功能执行逻辑(例如，算术逻辑单元、加载存储单元(lsu)等)。在至少一个实施例中，可以以管线方式配置功能执行逻辑，其中可以在先前的指令完成之前发出新的指令。在至少一个实施例中，功能执行逻辑支持多种运算，包括整数和浮点算术、比较操作、布尔运算、移位和各种代数函数的计算。在至少一个实施例中，可以利用相同的功能单元硬件来执行不同的操作，并且可以存在功能单元的任何组合。
[0245]
在至少一个实施例中，传送到处理集群2394的指令构成线程。在至少一个实施例中，跨一组并行处理引擎执行的一组线程是线程组。在至少一个实施例中，线程组在不同的输入数据上执行程序。在至少一个实施例中，线程组内的每个线程可被分配给图形多处理器2334内的不同处理引擎。在至少一个实施例中，线程组可包括比图形多处理器2334内的多个处理引擎更少的线程。在至少一个实施例中，当线程组包括的线程数少于处理引擎的数量时，一个或更多个处理引擎在正在处理该线程组的循环期间可能是空闲的。在至少一个实施例中，线程组还可以包括比图形多处理器2334内的多个处理引擎更多的线程。在至少一个实施例中，当线程组包括比图形多处理器2334内的处理引擎的数量更多的线程时，可以在连续的时钟周期内执行处理。在至少一个实施例中，可以在图形多处理器2334上同
时执行多个线程组。
[0246]
在至少一个实施例中，图形多处理器2334包括内部高速缓存存储器，以执行加载和存储操作。在至少一个实施例中，图形多处理器2334可以放弃内部高速缓存并使用处理集群2394内的高速缓存存储器(例如，l1高速缓存2348)。在至少一个实施例中，每个图形多处理器2334还可以访问分区单元(例如，图23a的分区单元2320a-2320n)内的l2高速缓存，这些分区单元在所有处理集群2394之间共享并且可以用于在线程之间传输数据。在至少一个实施例中，图形多处理器2334还可以访问片外全局存储器，其可以包括本地并行处理器存储器和/或系统存储器中的一个或更多个。在至少一个实施例中，并行处理单元2302外部的任何存储器都可以用作全局存储器。在至少一个实施例中，处理集群2394包括图形多处理器2334的多个实例，它们可以共享可以存储在l1高速缓存2348中的公共指令和数据。
[0247]
在至少一个实施例中，每个处理集群2394可以包括配置成将虚拟地址映射为物理地址的mmu 2345。在至少一个实施例中，mmu 2345的一个或更多个实例可以驻留在图23a的存储器接口2318内。在至少一个实施例中，mmu 2345包括一组页表条目(pte)，其用于将虚拟地址映射到图块(谈论有关图块的更多信息)的物理地址以及可选地映射到高速缓存行索引。在至少一个实施例中，mmu 2345可以包括地址转换后备缓冲区(tlb)或可以驻留在图形多处理器2334或l1高速缓存2348或处理集群2394内的高速缓存。在至少一个实施例中，处理物理地址以分配表面数据访问局部性，以便在分区单元之间进行有效的请求交织。在至少一个实施例中，高速缓存行索引可以用于确定对高速缓存线的请求是命中还是未命中。
[0248]
在至少一个实施例中，可以配置处理集群2394，使得每个图形多处理器2334耦合到纹理单元2336，以执行纹理映射操作，例如，可以涉及确定纹理样本位置、读取纹理数据以及过滤纹理数据。在至少一个实施例中，根据需要从内部纹理l1高速缓存(未示出)或从图形多处理器2334内的l1高速缓存中读取纹理数据，并从l2高速缓存、本地并行处理器存储器或系统存储器中获取纹理数据。在至少一个实施例中，每个图形多处理器2334将处理后的任务输出到数据交叉开关2340，以将处理后的任务提供给另一处理集群2394以进行进一步处理或将处理后的任务存储在l2高速缓存、本地并行处理器存储器、或经由存储器交叉开关2316的系统存储器中。在至少一个实施例中，光栅前操作单元(prerop)2342配置成从图形多处理器2334接收数据，将数据引导至rop单元，该rop单元可以与本文所述的分区单元(例如，图23a的分区单元2320a-2320n)一起定位。在至少一个实施例中，prerop 2342单元可以执行用于颜色混合的优化、组织像素颜色数据以及执行地址转换。
[0249]
图23c示出了根据至少一个实施例的图形多处理器2396。在至少一个实施例中，图形多处理器2396是图23b的图形多处理器2334。在至少一个实施例中，图形多处理器2396与处理集群2394的管线管理器2332耦合。在至少一个实施例中，图形多处理器2396具有执行管线，该执行管线包括但不限于指令高速缓存2352、指令单元2354、地址映射单元2356、寄存器文件2358、一个或更多个gpgpu核心2362和一个或更多个lsu 2366。gpgpu核心2362和lsu 2366与高速缓存存储器2372和共享存储器2370通过存储器和高速缓存互连2368耦合。
[0250]
在至少一个实施例中，指令高速缓存2352从管线管理器2332接收要执行的指令流。在至少一个实施例中，将指令高速缓存在指令高速缓存2352中并将其分派以供指令单元2354执行。在一个实施例中，指令单元2354可以分派指令作为线程组(例如，线程束)，将
线程组的每个线程分配给gpgpu核心2362内的不同执行单元。在至少一个实施例中，指令可以通过在统一地址空间内指定地址来访问任何本地、共享或全局地址空间。在至少一个实施例中，地址映射单元2356可以用于将统一地址空间中的地址转换成可以由lsu 2366访问的不同的存储器地址。
[0251]
在至少一个实施例中，寄存器文件2358为图形多处理器2396的功能单元提供了一组寄存器。在至少一个实施例中，寄存器文件2358为连接到图形多处理器2396的功能单元(例如，gpgpu核心2362、lsu 2366)的数据路径的操作数提供了临时存储。在至少一个实施例中，在每个功能单元之间划分寄存器文件2358，使得为每个功能单元分配寄存器文件2358的专用部分。在至少一个实施例中，寄存器文件2358在图形多处理器2396正在执行的不同线程组之间划分。
[0252]
在至少一个实施例中，gpgpu核心2362可以各自包括用于执行图多处理器2396的指令的fpu和/或alu。gpgpu核心2362在架构上可以相似或架构可能有所不同。在至少一个实施例中，gpgpu核心2362的第一部分包括单精度fpu和整数alu，而gpgpu核心的第二部分包括双精度fpu。在至少一个实施例中，fpu可以实现用于浮点算法的ieee754-2308标准或启用可变精度浮点算法。在至少一个实施例中，图形多处理器2396可以另外包括一个或更多个固定功能或特殊功能单元，以执行特定功能，诸如复制矩形或像素混合操作。在至少一个实施例中，gpgpu核心2362中的一个或更多个也可以包括固定或特殊功能逻辑。
[0253]
在至少一个实施例中，gpgpu核心2362包括能够对多组数据执行单个指令的simd逻辑。在至少一个实施例中，gpgpu核心2362可以物理地执行simd4、simd8和simd9指令，并且在逻辑上执行simd1、simd2和simd32指令。在至少一个实施例中，用于gpgpu核心的simd指令可以在编译时由着色器编译器生成，或者在执行针对单程序多数据(spmd)或simt架构编写和编译的程序时自动生成。在至少一个实施例中，可以通过单个simd指令来执行为simt执行模型配置的程序的多个线程。例如，在至少一个实施例中，可以通过单个simd8逻辑单元并行执行执行相同或相似操作的八个simt线程。
[0254]
在至少一个实施例中，存储器和高速缓存互连2368是将图形多处理器2396的每个功能单元连接到寄存器文件2358和共享存储器2370的互连网络。在至少一个实施例中，存储器和高速缓存互连2368是交叉开关互连，其允许lsu 2366在共享存储器2370和寄存器文件2358之间实现加载和存储操作。在至少一个实施例中，寄存器文件2358可以以与gpgpu核心2362相同的频率操作，从而在gpgpu核心2362和寄存器文件2358之间进行数据传输的延迟非常低。在至少一个实施例中，共享存储器2370可以用于启用在图形多处理器2396内的功能单元上执行的线程之间的通信。在至少一个实施例中，高速缓存存储器2372可以用作例如数据高速缓存，以高速缓存在功能单元和纹理单元2336之间通信的纹理数据。在至少一个实施例中，共享存储器2370也可以用作程序管理的高速缓存。在至少一个实施例中，除了存储在高速缓存存储器2372中的自动高速缓存的数据之外，在gpgpu核心2362上执行的线程还可以以编程方式将数据存储在共享存储器中。
[0255]
在至少一个实施例中，如本文所述的并行处理器或gpgpu通信地耦合到主机/处理器核心，以加速图形操作、机器学习操作、图案分析操作以及各种通用gpu(gpgpu)功能。在至少一个实施例中，gpu可以通过总线或其他互连(例如，诸如pcie或nvlink的高速互连)通信地耦合到主机处理器/核心。在至少一个实施例中，gpu可以与核心集成在相同的封装或
芯片上，并通过内部处理器总线/互连(即，封装或芯片的内部)通信地耦合到核心。在至少一个实施例中，不管gpu连接的方式如何，处理器核心可以以wd包含的命令/指令序列的形式向gpu分配工作。在至少一个实施例中，gpu然后使用专用电路/逻辑来有效地处理这些命令/指令。
[0256]
图24示出了根据至少一个实施例的图形处理器2400。在至少一个实施例中，图形处理器2400包括环形互连2402、管线前端2404、媒体引擎2437和图形核心2480a-2480n。在至少一个实施例中，环形互连2402将图形处理器2400耦合到其他处理单元，包括其他图形处理器或一个或更多个通用处理器核心。在至少一个实施例中，图形处理器2400是集成在多核心处理系统内的许多处理器之一。
[0257]
在至少一个实施例中，图形处理器2400经由环形互连2402接收多批命令。在至少一个实施例中，输入命令由管线前端2404中的命令流转化器2403解释。在至少一个实施例中，图形处理器2400包括可缩放执行逻辑，以经由图形核心2480a-2480n执行3d几何处理和媒体处理。在至少一个实施例中，对于3d几何处理命令，命令流转化器2403将命令提供给几何管线2436。在至少一个实施例中，对于至少一些媒体处理命令，命令流转化器2403将命令提供给视频前端2434，其与媒体引擎2437耦合。在至少一个实施例中，媒体引擎2437包括用于视频和图像后处理的视频质量引擎(vqe)2430，以及用于提供硬件加速媒体数据编码和解码的多格式编码/解码(mfx)2433引擎。在至少一个实施例中，几何管线2436和媒体引擎2437各自生成用于由至少一个图形核心2480a提供的线程执行资源的执行线程。
[0258]
在至少一个实施例中，图形处理器2400包括以模块化图形核心2480a-2480n(有时称为核心切片)为特征的可缩放线程执行资源，每个模块核心具有多个子核心2450a-2450n、2460a-2460n(有时称为核心子切片)。在至少一个实施例中，图形处理器2400可以具有任意数量的图形核心2480a至2480n。在至少一个实施例中，图形处理器2400包括具有至少第一子核心2450a和第二子核心2460a的图形核心2480a。在至少一个实施例中，图形处理器2400是具有单个子核心(例如2450a)的低功率处理器。在至少一个实施例中，图形处理器2400包括多个图形核心2480a-2480n，每个图形核心包括一组第一子核心2450a-2450n和一组第二子核心2460a-2460n。在至少一个实施例中，第一子核心2450a-2450n中的每个子核心至少包括第一组执行单元(eu)2452a-2452n和媒体/纹理采样器2454a-2454n。在至少一个实施例中，第二子核心2460a-2460n中的每个子核心至少包括第二组执行单元2462a-2462n和采样器2464a-2464n。在至少一个实施例中，每个子核心2450a-2450n、2460a-2460n共享一组共享资源2470a-2470n。在至少一个实施例中，共享资源2470包括共享高速缓冲存储器和像素操作逻辑。
[0259]
图25示出了根据至少一个实施例的用于处理器2500。在至少一个实施例中，处理器2500可以包括但不限于执行指令的逻辑电路。在至少一个实施例中，处理器2500可以执行指令，包括x86指令、arm指令、用于asic的专用指令等。在至少一个实施例中，处理器2510可以包括用于存储封装数据的寄存器，例如作为加利福尼亚州圣克拉拉市英特尔公司采用mmx技术启用的微处理器中的64位宽mmxtm寄存器。在至少一个实施例中，整数和浮点数形式可用的mmx寄存器可以与封装的数据元素一起运行，所述封装的数据元素伴随simd和流式simd扩展(“sse”)指令。在至少一个实施例中，与sse2、sse3、sse4、avx或更高版本(一般称为“ssex”)技术有关的128位宽xmm寄存器可以保存此类封装数据操作数。在至少一个实
施例中，处理器2510可以执行指令以加速cuad程序。
[0260]
在至少一个实施例中，处理器2500包括有序前端(“前端”)2501，以提取要执行的指令并准备稍后在处理器管线中使用的指令。在至少一个实施例中，前端2501可以包括几个单元。在至少一个实施例中，指令预取器2526从存储器中获取指令并将指令提供给指令解码器2528，指令解码器2528又对指令进行解码或解释。例如，在至少一个实施例中，指令解码器2528将接收到的指令解码用于执行的所谓的“微指令”或“微操作”(也称为“微操作”或“微指令”)的一个或更多个操作。在至少一个实施例中，指令解码器2528将指令解析为操作码以及相应的数据和控制字段，其可以由微架构用来使用以执行操作。在至少一个实施例中，跟踪高速缓存2530可以将解码的微指令组装成微指令队列2534中的程序排序的序列或追踪以供执行。在至少一个实施例中，当追踪高速缓存2530遇到复杂指令时，微码rom 2532提供完成操作所需的微指令。
[0261]
在至少一个实施例中，可以将一些指令转换成单个微操作，而另一些指令则需要几个微操作来完成全部操作。在至少一个实施例中，如果需要多于四个的微指令来完成一条指令，则指令解码器2528可以访问微码rom 2532以执行指令。在至少一个实施例中，可以将指令解码为少量的微指令以在指令解码器2528处进行处理。在至少一个实施例中，如果需要多个微指令完成操作，则可以将指令存储在微码rom 2532中。在至少一个实施例中，追踪高速缓存器2530参考入口点可编程逻辑阵列(“pla”)以确定正确的微指令指针，用于根据至少一个实施例从微码rom 2532读取微码序列以完成一个或更多个指令。在至少一个实施例中，在微码rom 2532完成对指令的微操作排序之后，机器的前端2501可以恢复从追踪高速缓存2530获取微操作。
[0262]
在至少一个实施例中，乱序执行引擎(“乱序引擎”)2503可以准备用于执行的指令。在至少一个实施例中，乱序执行逻辑具有多个缓冲区，以使指令流平滑并重新排序，以在指令沿管线下降并被调度执行时优化性能。乱序执行引擎2503包括但不限于分配器/寄存器重命名器2540、存储器微指令队列2542、整数/浮点微指令队列2544、存储器调度器2546、快速调度器2502、慢速/通用浮点调度器(“慢速/通用fp调度器”)2504和简单浮点调度器(“简单fp调度器”)2506。在至少一个实施例中，快速调度器2502、慢速/通用浮点调度器2504和简单浮点调度器2506也统称为“微指令调度器2502、2504、2506”。分配器/寄存器重命名器2540分配每个微指令按顺序执行所需要的机器缓冲区和资源。在至少一个实施例中，分配器/寄存器重命名器2540将逻辑寄存器重命名为寄存器文件中的条目。在至少一个实施例中，分配器/寄存器重命名器2540还为两个微指令队列之一中的每个微指令分配条目，存储器微指令队列2542用于存储器操作和整数/浮点微指令队列2544用于非存储器操作，在存储器调度器2546和微指令调度器2502、2504、2506的前面。在至少一个实施例中，微指令调度器2502、2504、2506基于它们的从属输入寄存器操作数源的就绪性和需要完成的执行资源微指令的可用性来确定何时准备好执行微指令。在至少一个实施例中，至少一个实施例的快速调度器2502可以在主时钟周期的每个一半上调度，而慢速/通用浮点调度器2504和简单浮点调度器2506可以在每个主处理器时钟周期调度一次。在至少一个实施例中，微指令调度器2502、2504、2506对调度端口进行仲裁，以调度用于执行的微指令。
[0263]
在至少一个实施例中，执行块2511包括但不限于整数寄存器文件/支路网络2508、浮点寄存器文件/支路网络(“fp寄存器文件/支路网络”)2510、地址生成单元(“agu”)2512
和2514、快速算术逻辑单元(“快速alu”)2516和2518、慢速alu 2520、浮点alu(“fp”)2522和浮点移动单元(“fp移动”)2524。在至少一个实施例中，整数寄存器文件/支路网络2508和浮点寄存器文件/旁路网络2510在本文中也称为“寄存器文件2508、2510”。在至少一个实施例中，agus 2512和2514、快速alu 2516和2518、慢速alu 2520、浮点alu 2522和浮点移动单元2524在本文中也称为“执行单元2512、2514、2516、2518、2520、2522和2524”。在至少一个实施例中，执行框可以包括但不限于任意数量(包括零)和类型的寄存器文件、旁路网络、地址生成单元和执行单元(以任何组合)。
[0264]
在至少一个实施例中，寄存器文件2508、2510可以布置在微指令调度器2502、2504、2506与执行单元2512、2514、2516、2518、2520、2522和2524之间。在至少一个实施例中，整数寄存器文件/旁路网络2508执行整数运算。在至少一个实施例中，浮点寄存器文件/旁路网络2510执行浮点操作。在至少一个实施例中，寄存器文件2508、2510中的每一个可以包括但不限于旁路网络，该旁路网络可以绕过或转发尚未写入寄存器文件中的刚刚完成的结果到新的从属对象。在至少一个实施例中，寄存器文件2508、2510可以彼此通信数据。在至少一个实施例中，整数寄存器文件/旁路网络2508可以包括但不限于两个单独的寄存器文件、一个寄存器文件用于低阶32位数据，第二寄存器文件用于高阶32位数据。在至少一个实施例中，浮点寄存器文件/旁路网络2510可以包括但不限于128位宽的条目，因为浮点指令通常具有宽度为64至128位的操作数。
[0265]
在至少一个实施例中，执行单元2512、2514、2516、2518、2520、2522、2524可以执行指令。在至少一个实施例中，寄存器文件2508、2510存储微指令需要执行的整数和浮点数据操作数值。在至少一个实施例中，处理器2500可以包括但不限于任意数量的执行单元2512、2514、2516、2518、2520、2522、2524及其组合。在至少一个实施例中，浮点alu 2522和浮点移动单元2524，可以执行浮点、mmx、simd、avx和sse或其他操作，包括专门的机器学习指令。在至少一个实施例中，浮点alu 2522可以包括但不限于64位乘64位浮点除法器，以执行除法、平方根和余数微操作。在至少一个实施例中，可以用浮点硬件来处理涉及浮点值的指令。在至少一个实施例中，可以将alu操作传递给快速alu 2516、2518。在至少一个实施例中，快速alu 2516、2518可以以半个时钟周期的有效延迟执行快速操作。在至少一个实施例中，大多数复杂的整数运算进入慢速alu 2520，因为慢速alu 2520可以包括但不限于用于长延迟类型操作的整数执行硬件，例如乘法器、移位、标志逻辑和分支处理。在至少一个实施例中，存储器加载/存储操作可以由agu 2512、2514执行。在至少一个实施例中，快速alu 2516、快速alu 2518和慢速alu 2520可以对64位数据操作数执行整数运算。在至少一个实施例中，可以实现快速alu 2516、快速alu 2518和慢速alu 2520以支持包括16、32、128、256等的各种数据位大小。在至少一个实施例中，浮点alu 2522和浮点移动单元2524可以实现为支持具有各种宽度的位的一定范围的操作数。在至少一个实施例中，浮点alu 2522和浮点移动单元2524可以结合simd和多媒体指令对128位宽封装数据操作数进行操作。
[0266]
在至少一个实施例中，微指令调度器2502、2504、2506在父加载完成执行之前调度从属操作。在至少一个实施例中，由于可以在处理器2500中推测性地调度和执行微指令，处理器2500还可以包括用于处理存储器未命中的逻辑。在至少一个实施例中，如果数据高速缓存中的数据加载未命中，则可能存在在管线中正在运行的从属操作，其使调度器暂时没有正确的数据。在至少一个实施例中，一种重放机制追踪踪并重新执行使用不正确数据的
指令。在至少一个实施例中，可能需要重放从属操作并且可以允许完成独立操作。在至少一个实施例中，处理器的至少一个实施例的调度器和重放机制也可以设计为捕获用于文本串比较操作的指令序列。
[0267]
在至少一个实施例中，术语“寄存器”可以指代可以用作识别操作数的指令的一部分的机载处理器存储位置。在至少一个实施例中，寄存器可以是那些可以从处理器外部使用的寄存器(从程序员的角度来看)。在至少一个实施例中，寄存器可能不限于特定类型的电路。相反，在至少一个实施例中，寄存器可以存储数据、提供数据并执行本文描述的功能。在至少一个实施例中，本文描述的寄存器可以通过处理器内的电路使用多种不同技术来实现，例如专用物理寄存器、使用寄存器重命名动态分配的物理寄存器、专用和动态分配的物理寄存器的组合等。在至少一个实施例中，整数寄存器存储32位整数数据。至少一个实施例的寄存器文件还包含八个用于封装数据的多媒体simd寄存器。
[0268]
图26示出了根据至少一个实施例的处理器2600。在至少一个实施例中，处理器2600包括但不限于一个或更多个处理器核心(核心)2602a-2602n、集成存储器控制器2614和集成图形处理器2608。在至少一个实施例中，处理器2600可以包括直至并包括由虚线框表示的附加处理器核心2602n的附加核心。在至少一个实施例中，每个处理器核心2602a-2602n包括一个或更多个内部高速缓存单元2604a-2604n。在至少一个实施例中，每个处理器核心还可以访问一个或更多个共享高速缓存的单元2606。
[0269]
在至少一个实施例中，内部高速缓存单元2604a-2604n和共享高速缓存单元2606表示处理器2600内的高速缓存存储器层次结构。在至少一个实施例中，高速缓存存储器单元2604a-2604n可以包括每个处理器核心内的至少一级指令和数据以及共享中级缓存中的一级或更多级缓存，例如l2、l3、4级(l4)或其他级别的缓存，其中在外部存储器之前将最高级别的缓存归类为llc。在至少一个实施例中，高速缓存一致性逻辑维持各种高速缓存单元2606和2604a-2604n之间的一致性。
[0270]
在至少一个实施例中，处理器2600还可包括一组一个或更多个总线控制器单元2616和系统代理核心2610。在至少一个实施例中，一个或更多个总线控制器单元2616管理一组外围总线，例如一个或更多个pci或pciexpress总线。在至少一个实施例中，系统代理核心2610为各种处理器组件提供管理功能。在至少一个实施例中，系统代理核心2610包括一个或更多个集成存储器控制器2614，以管理对各种外部存储器设备(未示出)的访问。
[0271]
在至少一个实施例中，一个或更多个处理器核心2602a-2602n包括对多线程同时进行的支持。在至少一个实施例中，系统代理核心2610包括用于在多线程处理期间协调和操作处理器核心2602a-2602n的组件。在至少一个实施例中，系统代理核心2610可以另外包括电源控制单元(pcu)，该电源控制单元包括逻辑和组件以调节处理器核心2602a-2602n和图形处理器2608的一个或更多个电源状态。
[0272]
在至少一个实施例中，处理器2600另外包括图形处理器2608以执行图处理操作。在至少一个实施例中，图形处理器2608与共享高速缓存单元2606和包括一个或更多个集成存储器控制器2614的系统代理核心2610耦合。在至少一个实施例中，系统代理核心2610还包括用于驱动图形处理器输出到一个或更多个耦合的显示器的显示器控制器2611。在至少一个实施例中，显示器控制器2611也可以是经由至少一个互连与图形处理器2608耦合的独立模块，或者可以集成在图形处理器2608内。
[0273]
在至少一个实施例中，基于环的互连单元2612用于耦合处理器2600的内部组件。在至少一个实施例中，可以使用替代性互连单元，例如点对点互连、交换互连或其他技术。在至少一个实施例中，图形处理器2608经由i/o链路2613与环形互连2612耦合。
[0274]
在至少一个实施例中，i/o链路2613代表多种i/o互连中的至少一种，包括促进各种处理器组件与高性能嵌入式存储器模块2618(例如edram模块)之间的通信的封装i/o互连。在至少一个实施例中，处理器核心2602a-2602n和图形处理器2608中的每一个使用嵌入式存储器模块2618作为共享的llc。
[0275]
在至少一个实施例中，处理器核心2602a-2602n是执行公共指令集架构的同质核心。在至少一个实施例中，处理器核心2602a-2602n在isa方面是异构的，其中一个或更多个处理器核心2602a-2602n执行公共指令集，而一个或更多个其他处理器核心2602a-2602n执行公共指令集或不同指令集的子集。在至少一个实施例中，就微架构而言，处理器核心2602a-2602n是异构的，其中具有相对较高功耗的一个或更多个核心与具有较低功耗的一个或更多个功率核心耦合。在至少一个实施例中，处理器2600可以实现在一个或更多个芯片上或被实现为soc集成电路。
[0276]
图27示出了根据所描述的至少一个实施例的图形处理器核心2700。在至少一个实施例中，图形处理器核心2700被包括在图形核心阵列内。在至少一个实施例中，图形处理器核心2700(有时称为核心切片)可以是模块化图形处理器内的一个或更多个图形核心。在至少一个实施例中，图形处理器核心2700是一个图形核心切片的示例，并且本文所述的图形处理器可以基于目标功率和性能包络线包括多个图形核心切片。在至少一个实施例中，每个图形核心2700可以包括与多个子核心2701a-2701f耦合的固定功能块2730，也称为子切片，其包括通用和固定功能逻辑的模块块。
[0277]
在至少一个实施例中，固定功能块2730包括几何/固定功能管线2736，例如，在较低性能和/或较低功率的图形处理器实施方式中，该几何/固定功能管线2736可以由图形处理器2700中的所有子核心共享。在至少一个实施例中，几何/固定功能管线2736包括3d固定功能管线、视频前端单元，线程产生器和线程分派器以及管理统一返回缓冲区的统一返回缓冲区管理器。
[0278]
在至少一个实施例中，固定功能块2730还包括图形soc接口2737、图形微控制器2738和媒体管线2739。图形soc接口2737提供了图形核心2700以及soc集成电路系统中的其他处理器核心之间的接口。在至少一个实施例中，图形微控制器2738是可编程子处理器，其可配置为管理图形处理器2700的各种功能，包括线程分派、调度和抢占。在至少一个实施例中，媒体管线2739包括有助于对包括图像和视频数据的多媒体数据进行解码、编码、预处理和/或后处理的逻辑。在至少一个实施例中，媒体管线2739经由对子核心2701a-2701f内的计算或采样逻辑的请求来实现媒体操作。
[0279]
在至少一个实施例中，soc接口2737使图形核心2700能够与通用应用处理器核心(例如，cpu)和/或soc内的其他组件通信，包括存储器层次结构元素，诸如共享的llc存储器、系统ram和/或嵌入式片上或封装dram。在至少一个实施例中，soc接口2737还可以使得能够与soc内的固定功能设备(例如，相机成像管线)进行通信，并且使得能够使用和/或实现可以在图形核心2700和soc内部的cpu之间共享的全局存储器原子。在至少一个实施例中，soc接口2737还可以实现用于图形核心2700的电源管理控制，并且启用图形核心2700的
时钟域与soc内的其他时钟域之间的接口。在至少一个实施例中，soc接口2737使得能够从命令流转化器和全局线程分派器接收命令缓冲区，其配置为向图形处理器内的一个或更多个图形核心中的每一个提供命令和指令。在至少一个实施例中，当要执行媒体操作时，可以将命令和指令分派给媒体管线2739，或者当要执行图处理操作时，可以将其分配给几何形状和固定功能管线(例如，几何形状和固定功能管线2736、几何形状和固定功能管线2714)。
[0280]
在至少一个实施例中，图形微控制器2738可以配置为对图形核心2700执行各种调度和管理任务。在至少一个实施例中，图形微控制器2738可以在子核心2701a-2701f中的执行单元(eu)阵列2702a-2702f、2704a-2704f内的各种图形并行引擎上执行图和/或计算工作负载调度。在至少一个实施例中，在包括图形核心2700的soc的cpu核心上执行的主机软件可以提交多个图形处理器门铃之一的工作负载，其调用适当的图形引擎上的调度操作。在至少一个实施例中，调度操作包括确定接下来要运行哪个工作负载、将工作负载提交给命令流转化器、抢先在引擎上运行的现有工作负载、监控工作负载的进度以及在工作负载完成时通知主机软件。在至少一个实施例中，图形微控制器2738还可以促进图形核心2700的低功率或空闲状态，从而为图形核心2700提供在图形核心2700内独立于操作系统和/或系统上的图形驱动器软件的跨低功率状态转换的保存和恢复寄存器的能力。
[0281]
在至少一个实施例中，图形核心2700可以具有比所示的子核心2701a-2701f更多或更少的子核心，达n个模块化子核心。对于每组n个子核心，在至少一个实施例中，图形核心2700还可以包括共享功能逻辑2710、共享和/或高速缓存存储器2712、几何/固定功能管线2714以及附加的固定功能逻辑2716以加速各种图形和计算处理操作。在至少一个实施例中，共享功能逻辑2710可以包括可由图形核心2700内的每个n个子核心共享的逻辑单元(例如，采样器、数学和/或线程间通信逻辑)。共享和/或高速缓存存储器2712可以是图形核心2700内的n个子核心2701a-2701f的llc，并且还可以用作可由多个子核心访问的共享存储器。在至少一个实施例中，可以包括几何/固定功能管线2714来代替固定功能块2730内的几何/固定功能管线2736，并且可以包括相同或相似的逻辑单元。
[0282]
在至少一个实施例中，图形核心2700包括附加的固定功能逻辑2716，其可以包括供图形核心2700使用的各种固定功能加速逻辑。在至少一个实施例中，附加的固定功能逻辑2716包括用于仅位置着色中使用的附加的几何管线。在仅位置着色中，存在至少两个几何管线，而在几何/固定功能管线2716、2736内的完整几何管线和剔除管线中，其是可以包括在附加的固定功能逻辑2716中的附加几何管线。在至少一个实施例中，剔除管线是完整几何管线的修整版。在至少一个实施例中，完整管线和剔除管线可以执行应用程序的不同实例，每个实例具有单独的环境。在至少一个实施例中，仅位置着色可以隐藏被丢弃的三角形的长剔除运行，从而在某些情况下可以更早地完成着色。例如，在至少一个实施例中，附加固定功能逻辑2716中的剔除管线逻辑可以与主应用程序并行执行位置着色器，并且通常比完整管线更快地生成关键结果，因为剔除管线获取并遮蔽顶点的位置属性，无需执行光栅化和将像素渲染到帧缓冲区。在至少一个实施例中，剔除管线可以使用生成的临界结果来计算所有三角形的可见性信息，而与这些三角形是否被剔除无关。在至少一个实施例中，完整管线(在这种情况下可以称为重播管线)可以消耗可见性信息来跳过剔除的三角形以仅遮盖最终传递到光栅化阶段的可见三角形。
[0283]
在至少一个实施例中，附加的固定功能逻辑2716还可包括通用目标处理加速逻
辑，例如固定功能矩阵乘法逻辑，用于实现减速cuad程序。
[0284]
在至少一个实施例中，在每个图形子核心2701a-2701f内包括一组执行资源，其可用于响应于图形管线、媒体管线或着色器程序的请求来执行图、媒体和计算操作。在至少一个实施例中，图形子核心2701a-2701f包括多个eu阵列2702a-2702f、2704a-2704f，线程分派和线程间通信(td/ic)逻辑2703a-2703f，3d(例如，纹理)采样器2705a-2705f，媒体采样器2706a-2706f，着色器处理器2707a-2707f和共享本地存储器(slm)2708a-2708f。eu阵列2702a-2702f、2704a-2704f每个都包含多个执行单元，这些执行单元是gugpu，能够为图形、媒体或计算操作提供服务，执行浮点和整数/定点逻辑运算，包括图形、媒体或计算着色器程序。在至少一个实施例中，td/ic逻辑2703a-2703f为子核心内的执行单元执行本地线程分派和线程控制操作，并促进在子核心的执行单元上执行的线程之间的通信。在至少一个实施例中，3d采样器2705a-2705f可以将与纹理或其他3d图形相关的数据读取到存储器中。在至少一个实施例中，3d采样器可以基于与给定纹理相关联的配置的采样状态和纹理格式来不同地读取纹理数据。在至少一个实施例中，媒体采样器2706a-2706f可以基于与媒体数据相关联的类型和格式来执行类似的读取操作。在至少一个实施例中，每个图形子核心2701a-2701f可以可替代地包括统一的3d和媒体采样器。在至少一个实施例中，在每个子核心2701a-2701f内的执行单元上执行的线程可以利用每个子核心内的共享本地存储器2708a-2708f，以使在线程组内执行的线程能够使用片上存储器的公共池来执行。
[0285]
图28示出了根据至少一个实施例的并行处理单元(“ppu”)2800。在至少一个实施例中，ppu 2800配置有机器可读代码，该机器可读代码如果由ppu 2800执行，则使得ppu 2800执行贯穿本文描述的一些或全部过程和技术。在至少一个实施例中，ppu 2800是在一个或更多个集成电路设备上实现的多线程处理器，并且利用多线程作为被设计为处理在多个线程上并行执行的计算机可读指令(也称为机器可读指令或简单的指令)的延迟隐藏技术。在至少一个实施例中，线程是指执行线程，并且是被配置为由ppu 2800执行的一组指令的实例。在至少一个实施例中，ppu 2800是图形处理单元(“gpu”)，图形处理单元配置为实现用于处理三维(“3d”)图形数据的图形渲染管线，以便生成用于在显示设备(诸如lcd设备)上显示的二维(“2d”)图像数据。在至少一个实施例中，ppu 2800用于执行计算，诸如线性代数运算和机器学习运算。图28仅出于说明性目的示出了示例并行处理器，并且应被解释为在至少一个实施例中实现的处理器架构的非限制性示例。
[0286]
在至少一个实施例中，一个或更多个ppu 2800配置成加速高性能计算(“hpc”)、数据中心和机器学习应用程序。在至少一个实施例中，一个或更多个ppu 2800配置成加速cuda程序。在至少一个实施例中，ppu 2800包括但不限于i/o单元2806、前端单元2810、调度器单元2812、工作分配单元2814、集线器2816、交叉开关(“xbar”)2820、一个或更多个通用处理集群(“gpc”)2818和一个或更多个分区单元(“存储器分区单元”)2822。在至少一个实施例中，ppu 2800通过一个或更多个高速gpu互连(“gpu互连”)2808连接到主机处理器或其他ppu 2800。在至少一个实施例中，ppu 2800通过系统总线或互连2802连接到主机处理器或其他外围设备。在一实施例中，ppu 2800连接到包括一个或更多个存储器设备(“存储器”)2804的本地存储器。在至少一个实施例中，存储器设备2804包括但不限于一个或更多个动态随机存取存储器(“dram”)设备。在至少一个实施例中，一个或更多个dram设备配置和/或可配置为高带宽存储器(“hbm”)子系统，并且在每个设备内堆叠有多个dram管芯。
[0287]
在至少一个实施例中，高速gpu互连2808可以指代系统使用其来进行缩放的基于线的多通道通信链路，并包括与一个或更多个cpu结合的一个或更多个ppu 2800(“cpu”)，支持ppu 2800和cpu之间的高速缓存一致性以及cpu主控。在至少一个实施例中，高速gpu互连2808通过集线器2816将数据和/或命令传输到ppu 2800的其他单元，例如一个或更多个复制引擎、视频编码器、视频解码器、电源管理单元和/或在图28中可能未明确示出的其他组件。
[0288]
在至少一个实施例中，i/o单元2806配置为通过系统总线2802从主机处理器(图28中未示出)发送和接收通信(例如，命令、数据)。在至少一个实施例中，i/o单元2806直接通过系统总线2802或通过一个或更多个中间设备(例如内存桥)与主机处理器通信。在至少一个实施例中，i/o单元2806可以经由系统总线2802与一个或更多个其他处理器(例如一个或更多个ppu 2800)通信。在至少一个实施例中，i/o单元2806实现pcie接口，用于通过pcie总线进行通信。在至少一个实施例中，i/o单元2806实现用于与外部设备通信的接口。
[0289]
在至少一个实施例中，i/o单元2806对经由系统总线2802接收的分组进行解码。在至少一个实施例中，至少一些分组表示被配置为使ppu 2800执行各种操作的命令。在至少一个实施例中，i/o单元2806如命令所指定的那样将解码的命令发送到ppu 2800的各种其他单元。在至少一个实施例中，命令被发送到前端单元2810和/或被发送到集线器2816或ppu 2800的其他单元，例如一个或更多个复制引擎、视频编码器、视频解码器、电源管理单元等(图28中未明确示出)。在至少一个实施例中，i/o单元2806配置为在ppu 2800的各种逻辑单元之间路由通信。
[0290]
在至少一个实施例中，由主机处理器执行的程序在缓冲区中对命令流进行编码，该缓冲区将工作负载提供给ppu 2800以进行处理。在至少一个实施例中，工作负载包括指令和要由那些指令处理的数据。在至少一个实施例中，缓冲区是可由主机处理器和ppu 2800两者访问(例如，读/写)的存储器中的区域—主机接口单元可以配置为访问经由i/o单元2806通过系统总线2802传输的存储器请求连接到系统总线2802的系统存储器中的缓冲区。在至少一个实施例中，主机处理器将命令流写入缓冲区，然后将指示命令流开始的指针发送给ppu 2800，使得前端单元2810接收指向一个或更多个命令流指针并管理一个或更多个命令流，从命令流中读取命令并将命令转发到ppu 2800的各个单元。
[0291]
在至少一个实施例中，前端单元2810耦合到调度器单元2812，该调度器单元2812配置各种gpc 2818以处理由一个或更多个命令流定义的任务。在至少一个实施例中，调度器单元2812配置为跟踪与调度器单元2812管理的各种任务有关的状态信息，其中状态信息可以指示任务被分配给哪个gpc 2818，任务是活跃的还是非活跃的，与任务相关联的优先级等等。在至少一个实施例中，调度器单元2812管理在一个或更多个gpc 2818上执行的多个任务。
[0292]
在至少一个实施例中，调度器单元2812耦合到工作分配单元2814，该工作分配单元2814配置为分派任务以在gpc 2818上执行。在至少一个实施例中，工作分配单元2814跟踪从调度器单元2812接收到的多个调度任务并且工作分配单元2814管理每个gpc 2818的待处理任务池和活跃任务池。在至少一个实施例中，待处理任务池包括多个时隙(例如32个时隙)，这些时隙包含分配给要由特定的gpc 2818处理的任务；活跃任务池可包括用于由gpc 2818主动处理的任务的多个时隙(例如4个时隙)，以使随着gpc 2818中的一个完成任
务的执行，该任务将从gpc 2818的活动任务池中逐出，并且从待处理任务池中选择其他任务之一，并安排其在gpc 2818上执行。在至少一个实施例中，如果活跃任务在gpc 2818上处于空闲状态，例如在等待数据依赖性解决时，则活跃任务从gpc 2818中驱逐并返回到待处理任务池，同时选择了待处理任务池中的另一个任务并调度在gpc 2818上执行。
[0293]
在至少一个实施例中，工作分配单元2814经由xbar 2820与一个或更多个gpc 2818通信。在至少一个实施例中，xbar 2820是互连网络，其将ppu 2800的许多单元耦合到ppu 2800的其他单元，并且可以配置为将工作分配单元2814耦合到特定的gpc 2818。在至少一个实施例中，一个或更多个ppu 2800的其他单元也可以通过集线器2816连接到xbar 2820。
[0294]
在至少一个实施例中，任务由调度器单元2812管理，并由工作分配单元2814分配给gpc 2818之一。gpc 2818配置为处理任务并产生结果。在至少一个实施例中，结果可以由gpc 2818中的其他任务消耗，通过xbar2820路由到不同的gpc 2818或存储在存储器2804中。在至少一个实施例中，结果可以通过分区单元2822写到存储器2804中，其实现了用于向存储器2804写入数据或从存储器2804读取数据的存储器接口。在至少一个实施例中，结果可以经由高速gpu互连2808传输到另一ppu 2800或cpu。在至少一个实施例中，ppu 2800包括但不限于u个分区单元2822，其等于耦合到ppu 2800的分离且不同的存储器设备2804的数量。
[0295]
在至少一个实施例中，主机处理器执行驱动器核心，该驱动器核心实现应用程序编程接口(api)，该应用程序编程接口使在主机处理器上执行的一个或更多个应用程序能够调度操作以在ppu 2800上执行。在一个实施例中，多个计算应用由ppu 2800同时执行，并且ppu 2800为多个计算应用程序提供隔离、服务质量(“qos”)和独立的地址空间。在至少一个实施例中，应用程序生成指令(例如，以api调用的形式)，该指令使驱动器核心生成一个或更多个任务以供ppu 2800执行，并且驱动器核心将任务输出至由ppu 2800处理的一个或更多个流。在至少一个实施例中，每个任务包括一个或更多个相关线程组，其可以被称为线程束(warp)。在至少一个实施例中，线程束包括可以并行执行的多个相关线程(例如32个线程)。在至少一个实施例中，协作线程可以指代多个线程，包括用于执行任务并且通过共享存储器交换数据的指令。
[0296]
图29示出了根据至少一个实施例的gpc 2900。在至少一个实施例中，gpc 2900是图28的gpc 2818。在至少一个实施例中，每个gpc 2900包括但不限于用于处理任务的多个硬件单元，并且每个gpc 2900包括但不限于管线管理器2902、预光栅操作单元(“prop”)2904、光栅引擎2908、工作分配交叉开关(“wdx”)2916、存储器管理单元(“mmu”)2918、一个或更多个数据处理集群(“dpc”)2906，以及部件的任何合适组合。
[0297]
在至少一个实施例中，gpc 2900的操作由管线管理器2902控制。在至少一个实施例中，管线管理器2902管理一个或更多个dpc 2906的配置，以处理分配给gpc 2900的任务。在至少一个实施例中，管线管理器2902配置一个或更多个dpc 2906中的至少一个以实现图形渲染管线的至少一部分。在至少一个实施例中，dpc 2906配置为在可编程流式多处理器(“sm”)2914上执行顶点着色器程序。在至少一个实施例中，管线管理器2902配置为将从工作分配单元接收的数据包路由到gpc 2900内的适当逻辑单元，以及在至少一个实施例中，可以将一些数据包路由到prop 2904和/或光栅引擎2908中的固定功能硬件单元，而可以将
其他数据包路由到dpc 2906以由原始引擎2912或sm 2914进行处理。在至少一个实施例中，管线管理器2902配置dpc 2906中的至少一个以实现神经网络模型和/或计算管线。在至少一个实施例中，管线管理器2902配置dpc 2906中的至少一个以执行cuda程序的至少一部分。
[0298]
在至少一个实施例中，prop单元2904配置为将由光栅引擎2908和dpc 2906生成的数据路由到分区单元中的光栅操作(“rop”)单元，例如上面结合图28更详细描述的存储器分区单元2822等。在至少一个实施例中，prop单元2904配置为执行用于颜色混合的优化、组织像素数据、执行地址转换等等。在至少一个实施例中，光栅引擎2908包括但不限于配置为执行各种光栅操作的多个固定功能硬件单元，并且在至少一个实施例中，光栅引擎2908包括但不限于设置引擎、粗光栅引擎、剔除引擎、裁剪引擎、精细光栅引擎、图块聚合引擎及其任意合适的组合。在至少一个实施例中，设置引擎接收变换后的顶点并生成与由顶点定义的几何图元相关联的平面方程；平面方程式被传送到粗光栅引擎以生成基本图元的覆盖信息(例如，图块的x、y覆盖范围掩码)；粗光栅引擎的输出将传输到剔除引擎，在剔除引擎中与z测试失败的图元相关联的片段将被剔除，并传输到剪切引擎，在剪切引擎中剪切位于视锥范围之外的片段。在至少一个实施例中，将经过裁剪和剔除的片段传递给精细光栅引擎，以基于设置引擎生成的平面方程式生成像素片段的属性。在至少一个实施例中，光栅引擎2908的输出包括将由任何适当的实体(例如，由在dpc 2906内实现的片段着色器)处理的片段。
[0299]
在至少一个实施例中，包括在gpc 2900中的每个dpc 2906包括但不限于m管线控制器(“mpc”)2910；图元引擎2912；一个或更多个sm 2914；及其任何合适的组合。在至少一个实施例中，mpc 2910控制dpc 2906的操作，将从管线管理器2902接收的分组路由到dpc 2906中的适当单元。在至少一个实施例中，将与顶点相关联的分组路由到图元引擎2912，图元引擎2912配置为从存储器中获取与顶点关联的顶点属性；相反，可以将与着色器程序相关联的数据包发送到sm 2914。
[0300]
在至少一个实施例中，sm 2914包括但不限于可编程流式处理器，其配置为处理由多个线程表示的任务。在至少一个实施例中，sm 2914是多线程的并且配置为同时执行来自特定线程组的多个线程(例如32个线程)，并且实现单指令、多数据(“simd”)架构，其中将一组线程(例如，线程束)中的每个线程配置为基于相同的指令集来处理不同的数据集。在至少一个实施例中，线程组中的所有线程执行相同的指令。在至少一个实施例中，sm 2914实施单指令、多线程(“simt”)架构，其中一组线程中的每个线程配置为基于相同的指令集来处理不同的数据集，但是其中线程组中的各个线程允许在执行期间发散。在至少一个实施例中，为每个线程束维护程序计数器、调用栈和执行状态，从而当线程束中的线程发散时，实现线程束和线程束内的串行执行之间的并发性。在另一个实施例中，为每个单独的线程维护程序计数器、调用栈和执行状态，从而使得在线程束内和线程束之间的所有线程之间具有相等的并发性。在至少一个实施例中，为每个单独的线程维持执行状态，并且可以收敛并并行地执行执行相同指令的线程以提高效率。下面结合图30更详细地描述sm 2914的至少一个实施例。
[0301]
在至少一个实施例中，mmu 2918在gpc 2900和存储器分区单元(例如，图28的分区单元2822)之间提供接口，并且mmu 2918提供虚拟地址到物理地址的转换、存储器保护以及
存储器请求的仲裁。在至少一个实施例中，mmu 2918提供一个或更多个转换后备缓冲区(“tlb”)，用于执行虚拟地址到存储器中的物理地址的转换。
[0302]
图30示出了根据至少一个实施例的流式多处理器(“sm”)3000。在至少一个实施例中，sm 3000是图29的sm 2914。在至少一个实施例中，sm 3000包括但不限于指令高速缓存3002；一个或更多个调度器单元3004；寄存器文件3008；一个或更多个处理核心(“核心”)3010；一个或更多个特殊功能单元(“sfu”)3012；一个或更多个加载/存储单元(“lsu”)3014；互连网络3016；共享存储器/一级(“l1”)高速缓存3018；及其任何合适的组合。在至少一个实施例中，工作分配单元调度任务以在并行处理单元(“ppu”)的通用处理集群(“gpc”)上执行，并且每个任务被分配给gpc内部的特定数据处理集群(“dpc”)，并且如果任务与着色器程序相关联，则将任务分配给sm 3000之一。在至少一个实施例中，调度器单元3004从工作分配单元接收任务并管理分配给sm 3000的一个或更多个线程块的指令调度。在至少一个实施例中，调度器单元3004调度线程块以作为并行线程的线程束来执行，其中，每个线程块被分配至少一个线程束。在至少一个实施例中，每个线程束执行线程。在至少一个实施例中，调度器单元3004管理多个不同的线程块，将线程束分配给不同的线程块，然后在每个时钟周期内将来自多个不同的协作组的指令分派给各种功能单元(例如，处理核心3010、sfu 3012和lsu 3014)。
[0303]
在至少一个实施例中，“合作组”可以指用于组织通信线程组的编程模型，其允许开发人员表达线程正在通信的粒度，从而能够表达更丰富、更有效的并行分解。在至少一个实施例中，协作启动api支持线程块之间的同步以执行并行算法。在至少一个实施例中，常规编程模型的api提供了用于同步协作线程的单一、简单的构造：跨线程块的所有线程的屏障(例如，syncthreads()函数)。但是，在至少一个实施例中，程序员可以在小于线程块粒度的情形下来定义线程组，并在所定义的组内进行同步，以实现更高的性能、设计灵活性以及以集合组范围功能接口的形式实现软件重用。在至少一个实施例中，协作组使程序员能够以子块和多块粒度明确定义线程组，并执行集合操作，例如对协作组中的线程进行同步。在至少一个实施例中，子块粒度与单个线程一样小。在至少一个实施例中，编程模型支持跨软件边界的干净组合，从而库和实用程序功能可以在其本地环境中安全地同步，而不必进行关于收敛的假设。在至少一个实施例中，协作组图元使协作并行的新图案成为可能，包括但不限于生产者-消费者并行，机会主义并行以及整个线程块网格上的全局同步。
[0304]
在至少一个实施例中，分派单元3006配置为将指令发送到功能单元中的一个或更多个，并且调度器单元3004包括但不限于两个分派单元3006，该两个分派单元3006使得来自相同线程束的两个不同指令能够在每个时钟周期被分派。在至少一个实施例中，每个调度器单元3004包括单个分派单元3006或附加分派单元3006。
[0305]
在至少一个实施例中，每个sm 3000在至少一个实施例中包括但不限于寄存器文件3008，该寄存器文件3008为sm 3000的功能单元提供了一组寄存器。在至少一个实施例中，寄存器文件3008在每个功能单元之间划分，从而为每个功能单元分配寄存器文件3008的专用部分。在至少一个实施例中，寄存器文件3008在由sm 3000执行的不同线程束之间划分，并且寄存器文件3008为连接到功能单元的数据路径的操作数提供临时存储。在至少一个实施例中，每个sm 3000包括但不限于多个l个处理核心3010。在至少一个实施例中，sm 3000包括但不限于大量(例如128个或更多)不同的处理核心3010。在至少一个实施例中，每
个处理核心3010在至少一个实施例中包括但不限于全管线、单精度、双精度和/或混合精度处理单元，其包括但不限于浮点算术逻辑单元和整数算术逻辑单元。在至少一个实施例中，浮点算术逻辑单元实现用于浮点算术的ieee754-2008标准。在至少一个实施例中，处理核心3010包括但不限于64个单精度(32位)浮点核心、64个整数核心、32个双精度(64位)浮点核心和8个张量核心。
[0306]
在至少一个实施例中，张量核心配置为执行矩阵运算。在至少一个实施例中，一个或更多个张量核心包括在处理核心3010中。在至少一个实施例中，张量核心配置为执行深度学习矩阵算术，例如用于神经网络训练和推理的卷积运算。在至少一个实施例中，每个张量核心在4
×
4矩阵上操作并且执行矩阵乘法和累加运算d＝a
×
b c，其中a、b、c和d是4
×
4矩阵。
[0307]
在至少一个实施例中，矩阵乘法输入a和b是16位浮点矩阵，并且累加矩阵c和d是16位浮点或32位浮点矩阵。在至少一个实施例中，张量核心对16位浮点输入数据进行32位浮点累加运算。在至少一个实施例中，16位浮点乘法使用64个运算，并得到全精度乘积，然后使用32位浮点加法与其他中间乘积累加起来，以进行4x4x4矩阵乘法。在至少一个实施例中，张量核心用于执行由这些较小的元件构成的更大的二维或更高维度的矩阵运算。在至少一个实施例中，api(诸如cuda-c api)公开专门的矩阵加载、矩阵乘法和累加以及矩阵存储操作，以有效地使用来自cuda-c 程序的张量核心。在至少一个实施例中，在cuda级别，线程束级别接口假定跨越所有32个线程束线程的16
×
16大小的矩阵。
[0308]
在至少一个实施例中，每个sm 3000包括但不限于执行特殊功能(例如，属性评估、倒数平方根等)的m个sfu 3012。在至少一个实施例中，sfu 3012包括但不限于配置为遍历分层树数据结构的树遍历单元。在至少一个实施例中，sfu 3012包括但不限于配置为执行纹理映射过滤操作的纹理单元。在至少一个实施例中，纹理单元配置为从存储器中加载纹理映射(例如，纹理像素的2d阵列)和采样纹理映射，以产生采样的纹理值以供由sm 3000执行的着色器程序使用。在至少一个实施例中，将纹理映射存储在共享存储器/l1高速缓存3018中。在至少一个实施例中，纹理单元使用mip映射(mip-maps)(例如，细节级别不同的纹理映射)来实现纹理操作(诸如过滤操作)。在至少一个实施例中，每个sm 3000包括但不限于两个纹理单元。
[0309]
在至少一个实施例中，每个sm 3000包括但不限于实现共享存储器/l1高速缓存3018与寄存器文件3008之间的加载和存储操作的n个lsu 3014。在至少一个实施例中，每个sm 3000包括但不限于互连网络3016，互连网络3016将每个功能单元连接到寄存器文件3008，并且lsu 3014连接到寄存器文件3008和共享存储器/l1高速缓存3018。在至少一个实施例中，互连网络3016是交叉开关，其可以配置为将任何功能单元连接到寄存器文件3008中的任何寄存器，并且将lsu 3014连接到寄存器文件3008和共享存储器/l1高速缓存3018中的存储器位置。
[0310]
在至少一个实施例中，共享存储器/l1高速缓存3018是片上存储器的阵列，其在至少一个实施例中允许sm 3000与图元引擎之间以及sm 3000中的线程之间的数据存储和通信。在至少一个实施例中，共享存储器/l1高速缓存3018包括但不限于128kb的存储容量，并且位于从sm 3000到分区单元的路径中。在至少一个实施例中，共享存储器/l1高速缓存3018在至少一个实施例中用于高速缓存读取和写入。在至少一个实施例中，共享存储器/l1
高速缓存3018、l2高速缓存和存储器中的一个或更多个是后备存储。
[0311]
在至少一个实施例中，将数据高速缓存和共享存储器功能组合到单个存储器块中，为两种类型的存储器访问提供了改进的性能。在至少一个实施例中，容量由不使用共享存储器的程序使用或将其用作高速缓存，例如如果共享存储器配置为使用一半容量，则纹理和加载/存储操作可以使用剩余容量。根据至少一个实施例，在共享存储器/l1高速缓存3018内的集成使共享存储器/l1高速缓存3018能够用作用于流传输数据的高吞吐量管线，同时提供对频繁重用的数据的高带宽和低延迟访问。在至少一个实施例中，当配置用于通用并行计算时，与图形处理相比，可以使用更简单的配置。在至少一个实施例中，绕过固定功能gpu，从而创建了更加简单的编程模型。在至少一个实施例中，在通用并行计算配置中，工作分配单元直接将线程的块分配和分布给dpc。在至少一个实施例中，块中的线程执行相同的程序，在计算中使用唯一的线程id以确保每个线程生成唯一的结果，使用sm 3000执行程序并执行计算，使用共享存储器/l1高速缓存3018在线程之间进行通信，以及使用lsu 3014通过共享存储器/l1高速缓存3018和存储器分区单元来读写全局存储器。在至少一个实施例中，当被配置用于通用并行计算时，sm 3000向调度器单元3004写入可以用来在dpc上启动新工作的命令。
[0312]
在至少一个实施例中，ppu被包括在台式计算机、膝上型计算机、平板电脑、服务器、超级计算机、智能电话(例如，无线、手持设备)、pda、数码相机、车辆、头戴式显示器、手持式电子设备等中或与之耦合。在至少一个实施例中，ppu被实现在单个半导体衬底上。在至少一个实施例中，ppu与一个或更多个其他设备(例如附加的ppu、存储器、risc cpu，mmu、数模转换器(“dac”)等)一起被包括在片上系统(“soc”)中。
[0313]
在至少一个实施例中，ppu可以被包括在包括一个或更多个存储设备的图形卡上。图形卡可以配置为与台式计算机主板上的pcie插槽相连接。在至少一个实施例中，ppu可以是包括在主板的芯片组中的集成gpu(“igpu”)。
[0314]
通用计算的软件构造
[0315]
以下各图阐述但不限于用于实现至少一个实施例的示例性软件构造。
[0316]
图31示出了根据至少一个实施例的编程平台的软件栈。在至少一个实施例中，编程平台是用于利用计算系统上的硬件来加速计算任务的平台。在至少一个实施例中，软件开发人员可以通过库、编译器指令和/或对编程语言的扩展来访问编程平台。在至少一个实施例中，编程平台可以是但不限于cuda，radeon开放计算平台(“rocm”)，opencl(由khronosgroup开发的opencl
tm
)，sycl或inteloneapi。
[0317]
在至少一个实施例中，编程平台的软件栈3100为应用程序3101提供执行环境。在至少一个实施例中，应用程序3101可以包括能够在软件栈3100上启动的任何计算机软件。在至少一个实施例中，应用程序3101可以包括但不限于人工智能(“ai”)/机器学习(“ml”)应用程序，高性能计算(“hpc”)应用程序，虚拟桌面基础架构(“vdi”)或数据中心工作负载。
[0318]
在至少一个实施例中，应用程序3101和软件栈3100在硬件3107上运行。在至少一个实施例中，硬件3107可以包括一个或更多个gpu，cpu，fpga，ai引擎和/或支持编程平台的其他类型的计算设备。在至少一个实施例中，例如采用cuda，软件栈3100可以是厂商专用的，并且仅与来自特定厂商的设备兼容。在至少一个实施例中，例如在采用opencl中，软件栈3100可以与来自不同供应商的设备一起使用。在至少一个实施例中，硬件3107包括连接
到一个或更多个设备的主机，该设备可经由应用程序编程接口(api)调用被访问以执行计算任务。在至少一个实施例中，与硬件3107内的主机相比，其可以包括但不限于cpu(但还可以包括计算设备)及其存储器，硬件3107内的设备可以包括但不限于gpu，fpga，ai引擎或其他计算设备(但还可以包括cpu)及其存储器。
[0319]
在至少一个实施例中，编程平台的软件栈3100包括但不限于多个库3103，运行时(runtime)3105和设备内核驱动器3106。在至少一个实施例中，库3103中的每个库可以包括可以由计算机程序使用并在软件开发期间利用的数据和编程代码。在至少一个实施例中，库3103可以包括但不限于预写的代码和子例程，类，值，类型规范，配置数据，文档，帮助数据和/或消息模板。在至少一个实施例中，库3103包括被优化用于在一种或更多种类型的设备上执行的函数。在至少一个实施例中，库3103可以包括但不限于用于在设备上执行数学、深度学习和/或其他类型的运算的函数。在至少一个实施例中，库2903与对应的api 3102相关联，api 3102可包括一个或更多个api，其暴露在库3103中实现的函数。
[0320]
在至少一个实施例中，将应用程序3101编写为源代码，该源代码被编译成可执行代码，如下面结合图36-38更详细讨论的。在至少一个实施例中，应用程序3101的可执行代码可以至少部分地在由软件栈3100提供的执行环境上运行。在至少一个实施例中，在应用程序3101的执行期间，可以得到需要在设备(与主机相比)上运行的代码。在这种情况下，在至少一个实施例中，可以调用运行时3105以在设备上加载和启动必需的代码。在至少一个实施例中，运行时3105可以包括能够支持应用程序3101的执行的任何技术上可行的运行时系统。
[0321]
在至少一个实施例中，运行时3105被实现为与对应的api(其被示为api 3104)相关联的一个或更多个运行时库。在至少一个实施例中，一个或更多个这样的运行时库可以包括但不限于用于存储器管理，执行控制，设备管理，错误处理和/或同步等等的函数。在至少一个实施例中，存储器管理函数可以包括但不限于用于分配、解除分配和复制设备存储器以及在主机存储器和设备存储器之间传输数据的函数。在至少一个实施例中，执行控制函数可以包括但不限于在设备上启动函数(当函数是可从主机调用的全局函数时，有时称为“内核”)的函数，和用于在运行时库为要在设备上执行的给定函数维护的缓冲区中设置属性值的函数。
[0322]
在至少一个实施例中，可以任何技术上可行的方式来实现运行时库和相应的api 3104。在至少一个实施例中，一个(或任意数量的)api可以公开用于设备的细粒度控制的低级函数集，而另一(或任意数量的)api可以公开这样的较高级的函数集。在至少一个实施例中，可以在低级api之上构建高级运行时api。在至少一个实施例中，一个或更多个运行时api可以是在与语言无关的运行时api之上分层的特定于语言的api。
[0323]
在至少一个实施例中，设备内核驱动器3106被配置为促进与底层设备的通信。在至少一个实施例中，设备内核驱动器3106可以提供诸如api 3104之类的api和/或其他软件所依赖的低级函数。在至少一个实施例中，设备内核驱动器3106可以被配置为在运行时将中间表示(“ir”)代码编译成二进制代码。在至少一个实施例中，对于cuda，设备内核驱动器3106可以在运行时将非硬件专用的并行线程执行(“ptx”)ir代码编译为用于特定目标设备的二进制代码(高速缓存已编译的二进制代码)，其有时也称为“最终”代码。在至少一个实施例中，这样做可以允许最终代码在目标设备上运行，而当源代码最初被编译为ptx代码
时，该目标设备可能不存在。备选地，在至少一个实施例中，设备源代码可以离线地编译成二进制代码，而不需要设备内核驱动器3106在运行时编译ir代码。
[0324]
图32示出了根据至少一个实施例的图31的软件栈3100的cuda实现。在至少一个实施例中，可在其上启动应用程序3201的cuda软件栈3200包括cuda库3203，cuda运行时3205，cuda驱动器3207和设备内核驱动器3208。在至少一个实施例中，cuda软件栈3200在硬件3209上执行，该硬件3209可以包括支持cuda的gpu，其由加利福尼亚州圣克拉拉市的nvidia公司开发。
[0325]
在至少一个实施例中，应用程序3201、cuda运行时3205和设备内核驱动器3208可以分别执行与应用程序3101、运行时3105和设备内核驱动器3106类似的功能，以上结合图31对其进行了描述。在至少一个实施例中，cuda驱动器3207包括实现cuda驱动器api 3206的库(libcuda.so)。在至少一个实施例中，类似于由cuda运行时库(cudart)实现的cuda运行时api 3204，cuda驱动器api 3206可以公开但不限于用于存储器管理、执行控制、设备管理、错误处理、同步和/或图形互操作性等的函数。在至少一个实施例中，cuda驱动器api 3206与cuda运行时api 3204的不同之处在于，cuda运行时api 3204通过提供隐式初始化、上下文(类似于进程)管理和模块(类似于动态加载的库)管理来简化设备代码管理。与高级cuda运行时api 3204相反，在至少一个实施例中，cuda驱动器api 3206是提供对设备的更细粒度控制的低级api，特别是关于上下文和模块加载。在至少一个实施例中，cuda驱动器api 3206可以公开没有由cuda运行时api 3204公开的用于上下文管理的函数。在至少一个实施例中，cuda驱动器api 3206也与语言无关，并且除了支持cuda运行时api 3204之外，还支持例如opencl。此外，在至少一个实施例中，包括cuda运行时3205在内的开发库可被视为与驱动器组件分离，包括用户模式的cuda驱动器3207和内核模式的设备驱动器3208(有时也称为“显示”驱动器)。
[0326]
在至少一个实施例中，cuda库3203可以包括但不限于数学库，深度学习库，并行算法库和/或信号/图像/视频处理库，并行计算应用程序(例如应用程序3201)可以利用这些库。在至少一个实施例中，cuda库3203可包括数学库，例如cublas库，其是用于执行线性代数运算的基本线性代数子程序(“blas”)的实现；用于计算快速傅立叶变换(“fft”)的cufft库，以及用于生成随机数的curand库等。在至少一个实施例中，cuda库3203可以包括深度学习库，诸如用于深度神经网络的基元的cudnn库和用于高性能深度学习推理的tensorrt平台等等。
[0327]
图33示出了根据至少一个实施例的图31的软件栈3100的rocm实现。在至少一个实施例中，可在其上启动应用程序3301的rocm软件栈3300包括语言运行时3303、系统运行时3305、thunk 3307和rocm内核驱动器3308。在至少一个实施例中，rocm软件栈3300在硬件3309上执行，硬件3309可以包括支持rocm的gpu，其由加利福尼亚州圣克拉拉市的amd公司开发。
[0328]
在至少一个实施例中，应用程序3301可以执行与以上结合图31讨论的应用程序3101类似的功能。另外，在至少一个实施例中，语言运行时3303和系统运行时3305可以执行与以上结合图31讨论的运行时3105类似的功能。在至少一个实施例中，语言运行时3303和系统运行时3305的不同之处在于，系统运行时3305是实现rocr系统运行时api 3304并利用异构系统架构(“hsa”)运行时api的语言无关运行时。在至少一个实施例中，hsa运行时api
是一种瘦用户模式api，它公开接口以供访问和与amdgpu交互，包括用于存储器管理、通过架构分派内核的执行控制、错误处理、系统和代理信息以及运行时初始化和关闭等的函数。在至少一个实施例中，与系统运行时3305相比，语言运行时3303是rocr系统运行时api 3304之上分层的特定于语言的运行时api 3302的实现。在至少一个实施例中，语言运行时api可以包括但不限于可移植异构计算接口(“hip”)语言运行时api，异构计算编译器(“hcc”)语言运行时api或openclapi等等。特别是，hip语言是c 编程语言的扩展，具有cuda机制的功能相似版本，并且在至少一个实施例中，hip语言运行时api包括与以上结合图32讨论的cuda运行时api 3204相似的函数，例如用于存储器管理、执行控制、设备管理、错误处理和同步等的函数。
[0329]
在至少一个实施例中，thunk(roct)3307是可用于与底层rocm驱动器3308交互的接口。在至少一个实施例中，rocm驱动器3308是rock驱动器，其是amdgpu驱动器和hsa内核驱动器(amdkfd)的组合。在至少一个实施例中，amdgpu驱动器是由amd开发的用于gpu的设备内核驱动器，其执行与以上结合图31讨论的设备内核驱动器3106类似的功能。在至少一个实施例中，hsa内核驱动器是允许不同类型的处理器经由硬件特征更有效地共享系统资源的驱动器。
[0330]
在至少一个实施例中，各种库(未示出)可以被包括在语言运行时3303上方的rocm软件栈3300中，并且提供与以上结合图32讨论的cuda库3203相似的功能。在至少一个实施例中，各种库可以包括但不限于数学、深度学习和/或其他库，例如实现与cudacublas类似的函数的hipblas库，类似于cudacufft用于计算fft的rocfft库等。
[0331]
图34示出了根据至少一个实施例的图31的软件栈3100的opencl实现。在至少一个实施例中，可以在其上启动应用程序3401的opencl软件栈3400包括opencl框架3410，opencl运行时3406和驱动器3407。在至少一个实施例中，opencl软件栈3400在不是特定于供应商的硬件3209上执行。在至少一个实施例中，由于由不同厂商开发的设备支持opencl，因此可能需要特定的opencl驱动器才能与来自此类厂商的硬件进行互操作。
[0332]
在至少一个实施例中，应用程序3401，opencl运行时3406，设备内核驱动器3407和硬件3408可以分别执行与上面结合图31讨论的应用程序3101、运行时3105、设备内核驱动器3106和硬件3107类似的功能。在至少一个实施例中，应用程序3401还包括具有将在设备上执行的代码的opencl内核3402。
[0333]
在至少一个实施例中，opencl定义了一种“平台”，其允许主机控制连接到该主机的设备。在至少一个实施例中，opencl框架提供平台层api和运行时api，示出为平台api 3403和运行时api 3405。在至少一个实施例中，运行时api3405使用上下文来管理设备上内核的执行。在至少一个实施例中，每个标识的设备可以与各自的上下文相关联，运行时api 3405可以使用该上下文来管理该设备的命令队列、程序对象和内核对象、共享存储器对象等。在至少一个实施例中，平台api 3403公开了允许设备上下文用于选择和初始化设备，经由命令队列将工作提交给设备，以及使得能够进行来自和去往设备的数据传输等的函数。另外，在至少一个实施例中，opencl框架提供各种内置函数(未示出)，包括数学函数、关系函数和图像处理函数等。
[0334]
在至少一个实施例中，编译器3404也被包括在opencl框架3410中。在至少一个实施例中，源代码可以在执行应用程序之前被离线编译或者在执行应用程序期间被在线编
译。与cuda和rocm相反，至少一个实施例中的opencl应用程序可以由编译器3404在线编译，编译器3404被包括以代表可以用于将源代码和/或ir代码(例如标准可移植中间表示(“spir-v”)代码)编译为二进制代码的任意数量的编译器。可替代地，在至少一个实施例中，可以在执行这样的应用程序之前离线编译opencl应用程序。
[0335]
图35示出了根据至少一个实施例的由编程平台支持的软件。在至少一个实施例中，编程平台3504被配置为支持应用程序3500可以依赖的各种编程模型3503，中间件和/或库3502以及框架3501。在至少一个实施例中，应用程序3500可以是使用例如深度学习框架(例如，mxnet，pytorch或tensorflow)实现的ai/ml应用，其可以依赖于诸如cudnn，nvidiacollectivecommunicationslibrary(“nccl”)”和/或nvidia开发人员数据加载库(“dali”)cuda库之类的库，以在底层硬件上提供加速的计算。
[0336]
在至少一个实施例中，编程平台3504可以是以上分别结合图32、图33和图34描述的cuda、rocm或opencl平台之一。在至少一个实施例中，编程平台3504支持多个编程模型3503，其是底层计算系统的抽象，其允许算法和数据结构的表达。在至少一个实施例中，编程模型3503可以暴露底层硬件的特征以便改善性能。在至少一个实施例中，编程模型3503可以包括但不限于cuda，hip，opencl，c 加速大规模并行性(“c amp”)，开放多处理(“openmp”)，开放加速器(“openacc”)和/或vulcan计算(vulcancompute)。
[0337]
在至少一个实施例中，库和/或中间件3502提供编程模型3504的抽象的实现。在至少一个实施例中，这样的库包括可由计算机程序使用并在软件开发期间利用的数据和编程代码。在至少一个实施例中，除了可以从编程平台3504获得的那些之外，这样的中间件还包括向应用程序提供服务的软件。在至少一个实施例中，库和/或中间件3502可以包括但不限于cublas、cufft、curand和其他cuda库，或rocblas、rocfft、rocrand和其他rocm库。另外，在至少一个实施例中，库和/或中间件3502可以包括nccl和rocm通信集合库(“rccl”)库，其提供用于gpu的通信例程，用于深度学习加速的miopen库和/或用于线性代数、矩阵和向量运算、几何变换、数值求解器以及相关算法的本征库。
[0338]
在至少一个实施例中，应用程序框架3501依赖于库和/或中间件3502。在至少一个实施例中，每个应用程序框架3501是用于实现应用软件的标准结构的软件框架。回到上面讨论的ai/ml示例，在至少一个实施例中，可以使用框架(诸如caffe，caffe2，tensorflow，keras，pytorch或mxnet深度学习框架)来实现ai/ml应用。
[0339]
图36示出了根据至少一个实施例的编译代码以在图31-34的编程平台之一上执行。在至少一个实施例中，编译器3601接收源代码3600，其包括主机代码以及设备代码两者。在至少一个实施例中，编译器3601被配置为将源代码3600转换为用于在主机上执行的主机可执行代码3602以及用于在设备上执行的设备可执行代码3603。在至少一个实施例中，源代码3600可以在执行应用程序之前离线编译，或者在执行应用程序期间在线编译。
[0340]
在至少一个实施例中，源代码3600可以包括编译器3601支持的任何编程语言的代码，例如c 、c、fortran等。在至少一个实施例中，源代码3600可以包括在单源(single-source)文件中，其具有主机代码和设备代码的混合，并在其中指示了设备代码的位置。在至少一个实施例中，单源文件可以是包括cuda代码的.cu文件或包括hip代码的.hip.cpp文件。备选地，在至少一个实施例中，源代码3600可以包括多个源代码文件，而不是单源文件，在该单源文件中主机代码和设备代码是分开的。
[0341]
在至少一个实施例中，编译器3601被配置为将源代码3600编译成用于在主机上执行的主机可执行代码3602和用于在设备上执行的设备可执行代码3603。在至少一个实施例中，编译器3601执行操作，包括将源代码3600解析为抽象系统树(ast)，执行优化以及生成可执行代码。在源代码3600包括单源文件的至少一个实施例中，编译器3601可以将设备代码与主机代码在这种单源文件中分开，将设备代码和主机代码分别编译成设备可执行代码3603和主机可执行代码3602，以及将设备可执行代码3603和主机可执行代码3602在单个文件中链接到一起，如下面关于图37更详细讨论的。
[0342]
在至少一个实施例中，主机可执行代码3602和设备可执行代码3603可以是任何合适的格式，例如二进制代码和/或ir代码。在cuda的情况下，在至少一个实施例中，主机可执行代码3602可以包括本地对象代码，而设备可执行代码3603可以包括ptx中间表示的代码。在至少一个实施例中，在rocm的情况下，主机可执行代码3602和设备可执行代码3603都可以包括目标二进制代码。
[0343]
图37是根据至少一个实施例的编译代码以在图31-34的编程平台之一上执行的更详细图示。在至少一个实施例中，编译器3701被配置为接收源代码3700，编译源代码3700，并输出可执行文件3710。在至少一个实施例中，源代码3700是单源文件，例如.cu文件，.hip.cpp文件或其他格式的文件，其包括主机代码和设备代码两者。在至少一个实施例中，编译器3701可以是但不限于用于在.cu文件中编译cuda代码的nvidiacuda编译器(“nvcc”)，或用于在.hip.cpp文件中编译hip代码的hcc编译器。
[0344]
在至少一个实施例中，编译器3701包括编译器前端3702，主机编译器3705，设备编译器3706和链接器3709。在至少一个实施例中，编译器前端3702被配置为在源代码3700中将设备代码3704与主机代码3703分开。在至少一个实施例中，设备代码3704由设备编译器3706编译成设备可执行代码3708，如所描述的，其可以包括二进制代码或ir代码。在至少一个实施例中，主机代码3703由主机编译器3705单独地编译成主机可执行代码3707。在至少一个实施例中，对于nvcc，主机编译器3705可以是但不限于输出本机目标代码的通用c/c 编译器，而设备编译器3706可以是但不限于基于低级虚拟机(“llvm”)的编译器，其将llvm编译器基础架构分叉，并输出ptx代码或二进制代码。在至少一个实施例中，对于hcc，主机编译器3705和设备编译器3706两者可以是但不限于输出目标二进制代码的基于llvm的编译器。
[0345]
在至少一个实施例中，在将源代码3700编译成主机可执行代码3707和设备可执行代码3708之后，链接器3709将主机和设备可执行代码3707和3708在可执行文件3710中链接到一起。在至少一个实施例中，主机和ptx的本机目标代码或设备的二进制代码可以在可执行和可链接格式(“elf”)文件中链接在一起，该文件是用于存储目标代码的容器格式。
[0346]
图38示出了根据至少一个实施例的在编译源代码之前转换源代码。在至少一个实施例中，源代码3800通过转换工具3801传递，转换工具3801将源代码3800转换成转换后的源代码3802。在至少一个实施例中，编译器3803用于将转换后的源代码3802编译成主机可执行代码3804和设备可执行代码3805，其过程类似于由编译器3601将源代码3600编译成主机可执行代码3602和设备可执行代码3603的过程，如以上结合图36所讨论的。
[0347]
在至少一个实施例中，由转换工具3801执行的转换被用于移植(port)源代码3800，以在与最初打算在其上运行的不同的环境中执行。在至少一个实施例中，转换工具
3902定义的任意数量的函数的任意数量的调用。在至少一个实施例中，cuda运行时api 3902可以包括但不限于在主机上执行的任意数量的函数，用于分配和解除分配设备存储器，在主机存储器和设备存储器之间传输数据，管理具有多个设备的系统等。在至少一个实施例中，cuda源代码3910还可以包括对在任意数量的其他cudaapi中指定的任意数量的函数的任意数量的调用。在至少一个实施例中，cudaapi可以是被设计为由cuda代码使用的任何api。在至少一个实施例中，cudaapi包括但不限于cuda运行时api 3902，cuda驱动器api，用于任意数量的cuda库的api等。在至少一个实施例中并且相对于cuda运行时api 3902，cuda驱动器api是较低级别的api，但可以提供对设备的更细粒度的控制。在至少一个实施例中，cuda库的示例包括但不限于cublas，cufft，curand，cudnn等。
[0355]
在至少一个实施例中，cuda编译器3950编译输入的cuda代码(例如，cuda源代码3910)以生成主机可执行代码3970(1)和cuda设备可执行代码3984。在至少一个实施例中，cuda编译器3950是nvcc。在至少一个实施例中，主机可执行代码3970(1)是在cpu 3990上可执行的输入源代码中包括的主机代码的编译版本。在至少一个实施例中，cpu 3990可以是针对顺序指令处理而优化的任何处理器。
[0356]
在至少一个实施例中，cuda设备可执行代码3984是在启用cuda的gpu 3994上可执行的输入源代码中包括的设备代码的编译版本。在至少一个实施例中，cuda设备可执行代码3984包括但不限于二进制代码。在至少一个实施例中，cuda设备可执行代码3984包括但不限于ir代码，例如ptx代码，该ir代码在运行时被设备驱动器进一步编译为用于特定目标设备(例如，启用cuda的gpu 3994)的二进制代码。在至少一个实施例中，启用cuda的gpu 3994可以是针对并行指令处理而优化并且支持cuda的任何处理器。在至少一个实施例中，启用cuda的gpu 3994由加利福尼亚州圣克拉拉市的nvidia公司开发。
[0357]
在至少一个实施例中，cuda到hip转换工具3920被配置为将cuda源代码3910转换成功能上相似的hip源代码3930。在至少一个实施例中，hip源代码3930是hip编程语言的人类可读代码的集合。在至少一个实施例中，hip代码是hip编程语言的人类可读代码。在至少一个实施例中，hip编程语言是c 编程语言的扩展，其包括但不限于cuda机制的功能上相似的版本，用于定义设备代码并区分设备代码和主机代码。在至少一个实施例中，hip编程语言可以包括cuda编程语言的功能的子集。在至少一个实施例中，例如，hip编程语言包括但不限于定义全局函数3912的机制，但是这样的hip编程语言可能缺乏对动态并行性的支持，因此，在hip代码中定义的全局函数3912仅可从主机调用。
[0358]
在至少一个实施例中，hip源代码3930包括但不限于任意数量(包括零)的全局函数3912，任意数量(包括零)的设备函数3914，任意数量(包括零)的主机函数3916以及任意数量(包括零)的主机/设备函数3918。在至少一个实施例中，hip源代码3930还可以包括对在hip运行时api 3932中指定的任意数量的函数的任意数量的调用。在一个实施例中，hip运行时api 3932包括但不限于cuda运行时api 3902中包括的函数的子集的功能上相似的版本。在至少一个实施例中，hip源代码3930还可以包括对在任意数量的其他hipapi中指定的任意数量的函数的任意数量的调用。在至少一个实施例中，hipapi可以是被设计为供hip代码和/或rocm使用的任何api。在至少一个实施例中，hipapi包括但不限于hip运行时api 3932，hip驱动器api，用于任意数量的hip库的api，用于任意数量的rocm库的api等。
[0359]
在至少一个实施例中，cuda到hip转换工具3920将cuda代码中的每个内核调用从
cuda语法转换为hip语法，并将cuda代码中的任意数量的其他cuda调用转换为任意数量的其他功能上相似的hip调用。在至少一个实施例中，cuda调用是对在cudaapi中指定的函数的调用，并且hip调用是对在hipapi中指定的函数的调用。在至少一个实施例中，cuda到hip转换工具3920将对在cuda运行时api 3902中指定的函数的任意数量的调用转换为对在hip运行时api 3932中指定的函数的任意数量的调用。
[0360]
在至少一个实施例中，cuda到hip转换工具3920是被称为hipify-perl的工具，其执行基于文本的转换过程。在至少一个实施例中，cuda到hip转换工具3920是被称为hipify-clang的工具，相对于hipify-perl，其执行更复杂且更鲁棒的转换过程，该过程涉及使用clang(编译器前端)解析cuda代码，然后转换得到的符号。在至少一个实施例中，除了由cuda到hip转换工具3920执行的那些修改之外，将cuda代码正确地转换成hip代码可能还需要修改(例如，手动编辑)。
[0361]
在至少一个实施例中，hip编译器驱动器3940是确定目标设备3946，然后配置与目标设备3946兼容的编译器以编译hip源代码3930的前端。在至少一个实施例中，目标设备3946是针对并行指令处理而优化的处理器。在至少一个实施例中，hip编译器驱动器3940可以以任何技术上可行的方式确定目标设备3946。
[0362]
在至少一个实施例中，如果目标设备3946与cuda兼容(例如，启用cuda的gpu 3994)，则hip编译器驱动器3940生成hip/nvcc编译命令3942。在至少一个实施例中并且结合图39b更详细地描述的，hip/nvcc编译命令3942配置cuda编译器3950以使用但不限于hip到cuda转换头和cuda运行时库来编译hip源代码3930。在至少一个实施例中并且响应于hip/nvcc编译命令3942，cuda编译器3950生成主机可执行代码3970(1)和cuda设备可执行代码3984。
[0363]
在至少一个实施例中，如果目标设备3946与cuda不兼容，则hip编译器驱动器3940生成hip/hcc编译命令3944。在至少一个实施例中并且如结合图39c更详细地描述的，hip/hcc编译命令3944配置hcc 3960以使用hcc头和hip/hcc运行时库编译hip源代码3930。在至少一个实施例中并且响应于hip/hcc编译命令3944，hcc 3960生成主机可执行代码3970(2)和hcc设备可执行代码3982。在至少一个实施例中，hcc设备可执行代码3982是hip源代码3930中包含的可在gpu 3992上执行的设备代码的编译版本。在至少一个实施例中，gpu 3992可以是针对并行指令处理而优化的、与cuda不兼容且与hcc兼容的任何处理器。在至少一个实施例中，gpu 3992由加利福尼亚州圣克拉拉市的amd公司开发。在至少一个实施例中，gpu 3992是不启用cuda的gpu 3992。
[0364]
仅出于说明性目的，在图39a中描绘了在至少一个实施例中可以实现为编译cuda源代码3910以在cpu 3990和不同设备上执行的三个不同流程。在至少一个实施例中，直接cuda流程编译cuda源代码3910以在cpu 3990和启用cuda的gpu 3994上执行，而无需将cuda源代码3910转换为hip源代码3930。在至少一个实施例中，间接cuda流程将cuda源代码3910转换为hip源代码3930，然后编译hip源代码3930以在cpu 3990和启用cuda的gpu 3994上执行。在至少一个实施例中，cuda/hcc流程将cuda源代码3910转换为hip源代码3930，然后编译hip源代码3930以在cpu 3990和gpu 3992上执行。
[0365]
可以通过虚线和一系列气泡注释a1-a3描绘可以在至少一个实施例中实现的直接cuda流程。在至少一个实施例中，并且如气泡注释a1所示，cuda编译器3950接收cuda源代码
3910和配置cuda编译器3950以编译cuda源代码3910的cuda编译命令3948。在至少一个实施例中，直接cuda流程中使用的cuda源代码3910是用cuda编程语言编写的，该cuda编程语言基于除c 之外的其他编程语言(例如c，fortran，python，java等)。在至少一个实施例中，并且响应于cuda编译命令3948，cuda编译器3950生成主机可执行代码3970(1)和cuda设备可执行代码3984(用气泡注释a2表示)。在至少一个实施例中并且如用气泡注释a3所示，主机可执行代码3970(1)和cuda设备可执行代码3984可以分别在cpu 3990和启用cuda的gpu 3994上执行。在至少一个实施例中，cuda设备可执行代码3984包括但不限于二进制代码。在至少一个实施例中，cuda设备可执行代码3984包括但不限于ptx代码，并且在运行时被进一步编译成用于特定目标设备的二进制代码。
[0366]
可以通过虚线和一系列气泡注释b1-b6来描述可以在至少一个实施例中实现的间接cuda流程。在至少一个实施例中并且如气泡注释b1所示，cuda到hip转换工具3920接收cuda源代码3910。在至少一个实施例中并且如气泡注释b2所示，cuda到hip转换工具3920将cuda源代码3910转换为hip源代码3930。在至少一个实施例中并如气泡注释b3所示，hip编译器驱动器3940接收hip源代码3930，并确定目标设备3946是否启用了cuda。
[0367]
在至少一个实施例中并且如气泡注释b4所示，hip编译器驱动器3940生成hip/nvcc编译命令3942，并将hip/nvcc编译命令3942和hip源代码3930两者都发送到cuda编译器3950。在至少一个实施例中并且如结合图39b更详细地描述的，hip/nvcc编译命令3942配置cuda编译器3950以使用但不限于hip到cuda转换头和cuda运行时库来编译hip源代码3930。在至少一个实施例中并且响应于hip/nvcc编译命令3942，cuda编译器3950生成主机可执行代码3970(1)和cuda设备可执行代码3984(用气泡注释b5表示)。在至少一个实施例中并且如气泡注释b6所示，主机可执行代码3970(1)和cuda设备可执行代码3984可以分别在cpu 3990和启用cuda的gpu 3994上执行。在至少一个实施例中，cuda设备可执行代码3984包括但不限于二进制代码。在至少一个实施例中，cuda设备可执行代码3984包括但不限于ptx代码，并且在运行时被进一步编译成用于特定目标设备的二进制代码。
[0368]
可以通过实线和一系列气泡注释c1-c6来描述可以在至少一个实施例中实现的cuda/hcc流程。在至少一个实施例中并且如气泡注释c1所示，cuda到hip转换工具3920接收cuda源代码3910。在至少一个实施例中并且如气泡注释c2所示，cuda到hip转换工具3920将cuda源代码3910转换为hip源代码3930。在至少一个实施例中并且如气泡注释c3所示，hip编译器驱动器3940接收hip源代码3930，并确定目标设备3946未启用cuda。
[0369]
在至少一个实施例中，hip编译器驱动器3940生成hip/hcc编译命令3944，并且将hip/hcc编译命令3944和hip源代码3930两者发送到hcc 3960(用气泡注释c4表示)。在至少一个实施例中并且如结合图39c更详细地描述的，hip/hcc编译命令3964配置hcc 3960以使用但不限于hcc头和hip/hcc运行时库编译hip源代码3930。在至少一个实施例中并且响应于hip/hcc编译命令3944，hcc 3960生成主机可执行代码3970(2)和hcc设备可执行代码3982(用气泡注释c5表示)。在至少一个实施例中并且如气泡注释c6所示，主机可执行代码3970(2)和hcc设备可执行代码3982可以分别在cpu 3990和gpu 3992上执行。
[0370]
在至少一个实施例中，在将cuda源代码3910转换为hip源代码3930之后，hip编译器驱动器3940可随后用于生成用于启用cuda的gpu 3994或gpu 3992的可执行代码，而无需将cuda重新执行为hip转换工具3920。在至少一个实施例中，cuda到hip转换工具3920将
cuda源代码3910转换为hip源代码3930，然后将其存储在存储器中。在至少一个实施例中，hip编译器驱动器3940然后配置hcc 3960以基于hip源代码3930生成主机可执行代码3970(2)和hcc设备可执行代码3982。在至少一个实施例中，hip编译器驱动器3940随后配置cuda编译器3950以基于存储的hip源代码3930生成主机可执行代码3970(1)和cuda设备可执行代码3984。
[0371]
图39b示出了根据至少一个实施例的被配置为使用cpu 3990和启用cuda的gpu 3994来编译和执行图39a的cuda源代码3910的系统3904。在至少一个实施例中，系统3904包括但不限于cuda源代码3910，cuda到hip转换工具3920，hip源代码3930，hip编译器驱动器3940，cuda编译器3950，主机可执行代码3970(1)，cuda设备可执行代码3984，cpu 3990和启用cuda的gpu 3994。
[0372]
在至少一个实施例中并且如本文先前结合图39a所描述的，cuda源代码3910包括但不限于任意数量(包括零)的全局函数3912，任意数量(包括零)的设备函数3914，任意数量(包括零)的主机函数3916以及任意数量(包括零)的主机/设备函数3918。在至少一个实施例中，cuda源代码3910还包括但不限于对在任意数量的cudaapi中指定的任意数量的函数的任意数量的调用。
[0373]
在至少一个实施例中，cuda到hip转换工具3920将cuda源代码3910转换成hip源代码3930。在至少一个实施例中，cuda到hip转换工具3920将cuda源代码3910中的每个内核调用从cuda语法转换为hip语法，并将cuda源代码3910中任意数量的其他cuda调用转换为任意数量的其他功能上相似的hip调用。
[0374]
在至少一个实施例中，hip编译器驱动器3940确定目标设备3946是启用cuda的，并且生成hip/nvcc编译命令3942。在至少一个实施例中，然后hip编译器驱动器3940经由hip/nvcc编译命令3942配置cuda编译器3950以编译hip源代码3930。在至少一个实施例中，作为配置cuda编译器3950的一部分，hip编译器驱动器3940提供对hip到cuda转换头3952的访问。在至少一个实施例中，hip到cuda转换头3952将任意数量的hip api中指定的任意数量的机制(例如，函数)转换为任意数量的cuda api中指定的任意数量的机制。在至少一个实施例中，cuda编译器3950将hip到cuda转换头3952与对应于cuda运行时api 3902的cuda运行时库3954结合使用，以生成主机可执行代码3970(1)和cuda设备可执行代码3984。在至少一个实施例中，然后可以分别在cpu 3990和启用cuda的gpu 3994上执行主机可执行代码3970(1)和cuda设备可执行代码3984。在至少一个实施例中，cuda设备可执行代码3984包括但不限于二进制代码。在至少一个实施例中，cuda设备可执行代码3984包括但不限于ptx代码，并且在运行时被进一步编译成用于特定目标设备的二进制代码。
[0375]
图39c示出了根据至少一个实施例的系统3906，该系统3906被配置为使用cpu 3990和未启用cuda的gpu 3992来编译和执行图39a的cuda源代码3910。在至少一个实施例中，系统3906包括但不限于cuda源代码3910，cuda到hip转换工具3920，hip源代码3930，hip编译器驱动器3940，hcc 3960，主机可执行代码3970(2)，hcc设备可执行代码3982，cpu 3990和gpu 3992。
[0376]
在至少一个实施例中，并且如本文先前结合图39a所描述的，cuda源代码3910包括但不限于任意数量(包括零)的全局函数3912，任意数量(包括零)的设备函数3914，任意数量(包括零)的主机函数3916以及任意数量(包括零)的主机/设备函数3918。在至少一个实
施例中，cuda源代码3910还包括但不限于对在任意数量的cuda api中指定的任意数量的函数的任意数量的调用。
[0377]
在至少一个实施例中，cuda到hip转换工具3920将cuda源代码3910转换成hip源代码3930。在至少一个实施例中，cuda到hip转换工具3920将cuda源代码3910中的每个内核调用从cuda语法转换为hip语法，并将源代码3910中任意数量的其他cuda调用转换为任意数量的其他功能上相似的hip调用。
[0378]
在至少一个实施例中，hip编译器驱动器3940随后确定目标设备3946不是启用cuda的，并生成hip/hcc编译命令3944。在至少一个实施例中，然后hip编译器驱动器3940配置hcc 3960以执行hip/hcc编译命令3944，从而编译hip源代码3930。在至少一个实施例中，hip/hcc编译命令3944将hcc 3960配置为使用但不限于hip/hcc运行时库3958和hcc头3956来生成主机可执行代码3970(2)和hcc设备可执行代码3982。在至少一个实施例中，hip/hcc运行时库3958对应于hip运行时api 3932。在至少一个实施例中，hcc头3956包括但不限于用于hip和hcc的任意数量和类型的互操作性机制。在至少一个实施例中，主机可执行代码3970(2)和hcc设备可执行代码3982可以分别在cpu3990和gpu 3992上执行。
[0379]
图40示出了根据至少一个实施例的由图39c的cuda到hip转换工具3920转换的示例性内核。在至少一个实施例中，cuda源代码3910将给定内核被设计为解决的总体问题划分为可以使用线程块独立解决的相对粗糙的子问题。在至少一个实施例中，每个线程块包括但不限于任意数量的线程。在至少一个实施例中，每个子问题被划分为相对细小的部分(pieces)，这些部分可以由线程块中的线程协作并行地解决。在至少一个实施例中，线程块内的线程可以通过共享存储器共享数据并通过同步执行以协调存储器访问来协作。
[0380]
在至少一个实施例中，cuda源代码3910将与给定内核相关联的线程块组织成线程块的一维、二维或三维网格。在至少一个实施例中，每个线程块包括但不限于任意数量的线程，并且网格包括但不限于任意数量的线程块。
[0381]
在至少一个实施例中，内核是使用“__global__”声明说明符(specifier)定义的设备代码中的函数。在至少一个实施例中，使用cuda内核启动语法4010来指定针对给定内核调用执行内核的网格的尺寸以及相关联的流。在至少一个实施例中，cuda内核启动语法4010被指定为“kernelname《《《gridsize，blocksize，sharedmemorysize，stream》》》(kernelarguments)；”。在至少一个实施例中，执行配置语法是“《《《...》》》”构造，其被插入在内核名称(“kernelname”)和内核参数的括号列表(“kernelarguments”)之间。在至少一个实施例中，cuda内核启动语法4010包括但不限于cuda启动函数语法而不是执行配置语法。
[0382]
在至少一个实施例中，“gridsize”是dim3类型的，并且指定网格的尺寸和大小。在至少一个实施例中，类型dim3是cuda定义的结构，其包括但不限于无符号整数x，y和z。在至少一个实施例中，如果未指定z，则z默认为1。在至少一个实施例中，如果未指定y，则y默认为1。在至少一个实施例中，网格中的线程块的数量等于gridsize.x、gridsize.y和gridsize.z的乘积。在至少一个实施例中，“blocksize”是dim3类型的，并且指定每个线程块的尺寸和大小。在至少一个实施例中，每线程块的线程数等于blocksize.x、blocksize.y和blocksize.z的乘积。在至少一个实施例中，给定执行内核的每个线程唯一的线程id，该线程id可通过内置变量(例如“threadidx”)在内核内访问。
[0383]
在至少一个实施例中，关于cuda内核启动语法4010，“sharedmemorysize”是一可选参数，它指定共享存储器中除静态分配的存储器外，针对给定内核调用为每个线程块动态分配的字节数。在至少一个实施例中并且关于cuda内核启动语法4010，sharedmemorysize默认为零。在至少一个实施例中并且关于cuda内核启动语法4010，“流”是可选的参数，其指定相关联的流并且默认为零以指定默认流。在至少一个实施例中，流是按顺序执行的命令序列(其可能由不同的主机线程发出)。在至少一个实施例中，不同的流可以相对于彼此无序地或同时地执行命令。
[0384]
在至少一个实施例中，cuda源代码3910包括但不限于用于示例性内核“matadd”的内核定义和主函数。在至少一个实施例中，主函数是在主机上执行的主机代码，并且包括但不限于使内核matadd在设备上执行的内核调用。在至少一个实施例中，如图所示，内核matadd将大小为nxn的两个矩阵a和b相加，其中n为正整数，并将结果存储在矩阵c中。在至少一个实施例中，主函数将threadsperblock变量定义为16x16，numblocks变量为n/16xn/16。在至少一个实施例中，然后主函数指定内核调用“matadd《《《numblocks，threadsperblock》》》(a，b，c)；”。在至少一个实施例中，并且根据cuda内核启动语法4010，使用尺寸为n/16
×
n/16的线程块网格来执行内核matadd，其中每个线程块的尺寸为16
×
16。在至少一个实施例中，每个线程块包括256个线程，创建具有足够块的网格以使每个矩阵元素具有一个线程，并且该网格中的每个线程执行内核matadd以执行一个逐对的加法。
[0385]
在至少一个实施例中，在将cuda源代码3910转换成hip源代码3930的同时，cuda到hip转换工具3920将cuda源代码3910中的每个内核调用从cuda内核启动语法4010转换成hip内核启动语法4020，并将源代码3910中的任意数量的其他cuda调用转换为任意数量的其他功能上相似的hip调用。在至少一个实施例中，hip内核启动语法4020被指定为“hiplaunchkernelggl(kernelname,gridsize,blocksize,sharedmemorysize,stream,kernelarguments)；”。在至少一个实施例中，kernelname，gridsize，blocksize，sharememorysize，stream和kernelarguments中的每一个在hip内核启动语法4020中具有与在cuda内核启动语法4010中(本文先前描述)相同的含义。在至少一个实施例中，参数sharedmemorysize和stream在hip内核启动语法4020中是必需的，而在cuda内核启动语法4010中是可选的。
[0386]
在至少一个实施例中，除了使内核matadd在设备上执行的内核调用之外，图40中描绘的hip源代码3930的一部分与图40中描绘的cuda源代码3910的一部分相同。在至少一个实施例中，在hip源代码3930中定义内核matadd，具有与在cuda源代码3910中定义内核matadd相同的“__global__”声明说明符。在至少一个实施例中，在hip源代码3930中的内核调用是“hiplaunchkernelggl(matadd，numblocks，threadsperblock，0、0，a，b，c)；”，而cuda源代码3910中的相应内核调用是“matadd《《《numblocks，threadsperblock》》》(a，b，c)；”。
[0387]
图41更详细地示出了根据至少一个实施例的图39c的未启用cuda的gpu 3992。在至少一个实施例中，gpu 3992由圣塔克拉拉市的amd公司开发。在至少一个实施例中，gpu 3992可以被配置为以高度并行的方式执行计算操作。在至少一个实施例中，gpu 3992被配置为执行图形管线操作，诸如绘制命令、像素操作、几何计算以及与将图像渲染到显示器相关联的其他操作。在至少一个实施例中，gpu 3992被配置为执行与图形无关的操作。在至少
一个实施例中，gpu 3992被配置为执行与图形有关的操作和与图形无关的操作两者。在至少一个实施例中，gpu 3992可以被配置为执行hip源代码3930中包括的设备代码。
[0388]
在至少一个实施例中，gpu 3992包括但不限于任意数量的可编程处理单元4120，命令处理器4110，l2高速缓存4122，存储器控制器4170，dma引擎4180(1)，系统存储器控制器4182，dma引擎4180(2)和gpu控制器4184。在至少一个实施例中，每个可编程处理单元4120包括但不限于工作负载管理器4130和任意数量的计算单元4140。在至少一个实施例中，命令处理器4110读取来自一个或更多个命令队列(未示出)的命令，并将命令分发给工作负载管理器4130。在至少一个实施例中，对于每个可编程处理单元4120，相关的工作负载管理器4130将工作分发给包括在可编程处理单元4120中的计算单元4140。在至少一个实施例中，每个计算单元4140可以执行任意数量的线程块，但是每个线程块在单个计算单元4140上执行。在至少一个实施例中，工作组是线程块。
[0389]
在至少一个实施例中，每个计算单元4140包括但不限于任意数量的simd单元4150和共享存储器4160。在至少一个实施例中，每个simd单元4150实现simd架构并且被配置为并行执行操作。在至少一个实施例中，每个simd单元4150包括但不限于向量alu 4152和向量寄存器文件4154。在至少一个实施例中，每个simd单元4150执行不同的线程束。在至少一个实施例中，线程束是一组线程(例如16个线程)，其中线程束中的每个线程属于单个线程块，并且被配置为基于单个指令集来处理不同的数据集。在至少一个实施例中，可以使用预测来禁用线程束中的一个或更多个线程。在至少一个实施例中，通道是线程。在至少一个实施例中，工作项是线程。在至少一个实施例中，波前是线程束。在至少一个实施例中，线程块中的不同波前可一起同步并经由共享存储器4160进行通信。
[0390]
在至少一个实施例中，可编程处理单元4120被称为“着色引擎”。在至少一个实施例中，除了计算单元4140之外，每个可编程处理单元4120还包括但不限于任意数量的专用图形硬件。在至少一个实施例中，每个可编程处理单元4120包括但不限于任意数量(包括零)的几何处理器，任意数量(包括零)的光栅化器，任意数量(包括零)的渲染后端，工作负载管理器4130和任意数量的计算单元4140。
[0391]
在至少一个实施例中，计算单元4140共享l2高速缓存4122。在至少一个实施例中，l2高速缓存4122被分区。在至少一个实施例中，gpu3992中的所有计算单元4140可访问gpu存储器4190。在至少一个实施例中，存储器控制器4170和系统存储器控制器4182促进gpu 3992与主机之间的数据传输，并且dma引擎4180(1)使能gpu 3992与此主机之间的异步存储器传输。在至少一个实施例中，存储器控制器4170和gpu控制器4184促进gpu 3992与其他gpu 3992之间的数据传输，并且dma引擎4180(2)使能gpu 3992与其他gpu 3992之间的异步存储器传输。
[0392]
在至少一个实施例中，gpu 3992包括但不限于任意数量和类型的系统互连，该系统互连促进在gpu 3992内部或外部的任意数量和类型的直接或间接链接的组件之间的数据和控制传输。在至少一个实施例中，gpu 3992包括但不限于耦合到任意数量和类型的外围设备的任意数量和类型的i/o接口(例如，pcie)。在至少一个实施例中，gpu 3992可以包括但不限于任意数量(包括零)的显示引擎和任意数量(包括零)的多媒体引擎。在至少一个实施例中，gpu 3992实现了存储器子系统，该存储器子系统包括但不限于任意数量和类型的存储器控制器(例如，存储器控制器4170和系统存储器控制器4182)以及专用于一个组件
或在多个组件之间共享的存储器设备(例如，共享存储器4160)。在至少一个实施例中，gpu3992实现了高速缓存子系统，该高速缓存子系统包括但不限于一个或更多个高速缓存存储器(例如，l2高速缓存4122)，每个高速缓存存储器可以是私有的或在任意数量的组件(例如，simd单元4150，计算单元4140和可编程处理单元4120)之间共享。
[0393]
图42示出了根据至少一个实施例的示例性cuda网格4220的线程如何被映射到图41的不同计算单元4140。在至少一个实施例中，并且仅出于说明目的，网格4220具有bx乘以by乘以1的gridsize和tx乘以ty乘以1的blocksize。因此，在至少一个实施例中，网格4220包括但不限于(bx*by)线程块4230，每个线程块4230包括但不限于(tx*ty)线程4240。线程4240在图42中被描绘为弯曲箭头。
[0394]
在至少一个实施例中，网格4220被映射到可编程处理单元4120(1)，该可编程处理单元4120(1)包括但不限于计算单元4140(1)-4140(c)。在至少一个实施例中并且如图所示，将(bj*by)线程块4230映射到计算单元4140(1)，并且将其余线程块4230映射到计算单元4140(2)。在至少一个实施例中，每个线程块4230可以包括但不限于任意数量的线程束，并且每个线程束被映射到图41的不同的simd单元4150。
[0395]
在至少一个实施例中，给定线程块4230中的线程束可以一起同步并通过关联的计算单元4140中包括的共享存储器4160进行通信。例如并且在至少一个实施例中，线程块4230(bj，1)中的线程束可以一起同步并通过共享存储器4160(1)进行通信。例如并且在至少一个实施例中，线程块4230(bj 1，1)中的线程束可以一起同步并通过共享存储器4160(2)进行通信。
[0396]
图43示出了根据至少一个实施例的如何将现有cuda代码迁移到数据并行c 代码。数据并行c (dpc )可以指开放的、基于标准的单个体系结构的专有语言的替代方案，它允许开发人员跨硬件目标(cpu和加速器，诸如gpu和fpga)重用代码，并且还允许为特定加速器执行自定义调整。dpc 根据开发人员可能熟悉的iso c 使用类似的和/或相同的c和c 构造。dpc 结合了khronos group的标准sycl，以支持数据并行性和异构编程。sycl是指跨平台抽象层，它建立在opencl的底层概念、可移植性和效率之上，使异构处理器的代码能够使用标准c 以“单个源”风格编写。sycl可以实现单源开发，其中c 模板函数可以包含主机代码和设备代码，以构建使用opencl加速的复杂算法，然后在不同类型的数据在整个源代码中重用它们。
[0397]
在至少一个实施例中，使用dpc 编译器来编译可以跨不同硬件目标部署的dpc 源代码。在至少一个实施例中，dpc 编译器用于生成可跨不同硬件目标部署的dpc 应用程序，并且dpc 兼容性工具可用于将cuda应用程序迁移到dpc 中的多平台程序。在至少一个实施例中，dpc 基础工具包包括：dpc 编译器，用于跨不同硬件目标部署应用程序；dpc 库，用于提高cpu、gpu和fpga的生产效率和性能；dpc 兼容性工具，用于将cuda应用程序迁移到多平台应用程序；及其任何合适的组合。
[0398]
在至少一个实施例中，dpc 编程模型用于通过使用现代c 特征来用称为数据并行c 的编程语言表达并行性，来简化与对cpu和加速器编程相关的一个或更多个方面。dpc 编程语言可用于使用单个源语言对主机(例如cpu)和加速器(例如gpu或fpga)进行代码重用，并清楚地传达执行和存储器依赖关系。dpc 代码中的映射可用于将应用程序过渡为在硬件或一组硬件设备上运行，这最佳地加速工作负载。即使在没有可用加速器的平台上，
也可以使用主机来简化设备代码的开发和调试。
[0399]
在至少一个实施例中，cuda源代码4300作为输入提供给dpc 兼容性工具4302，以生成人类可读的dpc 4304。在至少一个实施例中，人类可读的dpc 4304包括由dpc 兼容性工具4302生成的内联注释，该内联注释指导开发人员如何和/或在何处修改dpc 代码，以完成编码和调节到所需性能4306，从而生成dpc 源代码4308。
[0400]
在至少一个实施例中，cuda源代码4300是或包括cuda编程语言中的人类可读源代码的集合。在至少一个实施例中，cuda源代码4300是采用cuda编程语言的人类可读源代码。在至少一个实施例中，cuda编程语言是c 编程语言的扩展，其包括但不限于定义设备代码和区分设备代码和主机代码的机制。在至少一个实施例中，设备代码是在编译之后可在设备(例如，gpu或fpga)上执行的源代码，并且可以包括可以在设备的一个或更多个处理器核心上执行的一个或更多个可并行的工作流。在至少一个实施例中，设备可以是针对并行指令处理进行了优化的处理器，诸如支持cuda的gpu、gpu或另一gpgpu等。在至少一个实施例中，主机代码是编译后可在主机上执行的源代码。在至少一个实施例中，主机代码和设备代码中的一些或全部可以跨cpu和gpu/fpga并行执行。在至少一个实施例中，主机是针对顺序指令处理而优化的处理器，例如cpu。结合图43描述的cuda源代码4300可以与本文档中其他地方讨论的那些一致。
[0401]
在至少一个实施例中，dpc 兼容性工具4302指的是用于促进cuda源代码4300到dpc 源代码4308的迁移的可执行工具、程序、应用程序或任何其他合适类型的工具。在至少一个实施例中，dpc 兼容性工具4302是可用作dpc 工具包的一部分的基于命令行的代码迁移工具，其用于将现有cuda源移植到dpc 。在至少一个实施例中，dpc 兼容性工具4302将cuda应用程序的一些或全部源代码从cuda转换为dpc ，并生成至少部分用dpc 编写的结果文件，称为人类可读的dpc 4304。在至少一个实施例中，人类可读的dpc 4304包括由dpc 兼容性工具4302生成的注释，用于指示可能必需用户干预的地方。在至少一个实施例中，当cuda源代码4300调用没有类似dpc api的cuda api时，用户干预是必须的；需要用户干预的其他示例将在后面更详细地讨论。
[0402]
在至少一个实施例中，用于迁移cuda源代码4300(例如，应用程序或其部分)的工作流包括创建一个或更多个编译数据库文件；使用dpc 兼容性工具4302将cuda迁移到dpc ；完成迁移并验证正确性，从而生成dpc 源代码4308；并用dpc 编译器编译dpc 源代码4308，以生成dpc 应用程序。在至少一个实施例中，兼容性工具提供了一种实用程序(utility)，该实用程序拦截makefile执行时使用的命令并将它们存储在编译数据库文件中。在至少一个实施例中，文件以json格式存储。在至少一个实施例中，拦截构建(intercept-build)命令将makefile命令转换为dpc兼容性命令。
[0403]
在至少一个实施例中，拦截构建是一种实用程序脚本，其拦截构建过程以捕获编译选项、宏定义和包括路径，并将该数据写入编译数据库文件。在至少一个实施例中，编译数据库文件是json文件。在至少一个实施例中，dpc 兼容性工具4302解析编译数据库，并在迁移输入源时应用选项。在至少一个实施例中，拦截构建的使用是可选的，但强烈建议用于基于make或cmake的环境。在至少一个实施例中，迁移数据库包括命令、目录和文件：命令可以包括必要的编译标志；目录可以包括到报头(header)文件的路径；文件可包括到cuda文件的路径。
[0404]
在至少一个实施例中，dpc 兼容性工具4302通过尽可能生成dpc 来将以cuda编写的cuda代码(例如，应用程序)迁移到dpc 。在至少一个实施例中，dpc 兼容性工具4302可用作工具包的一部分。在至少一个实施例中，dpc 工具包包括拦截构建工具。在至少一个实施例中，拦截构建工具创建捕获编译命令以迁移cuda文件的编译数据库。在至少一个实施例中，由拦截构建工具生成的编译数据库被dpc 兼容性工具4302用于将cuda代码迁移到dpc 。在至少一个实施例中，非cuda c 代码和文件按原样迁移。在至少一个实施例中，dpc 兼容性工具4302生成人类可读的dpc 4304，其可以是dpc 代码，如由dpc 兼容性工具4302生成的，其不能由dpc 编译器编译并且需要附加的管线来验证未正确迁移的代码部分，并且可能涉及人工干预，例如由开发人员进行干预。在至少一个实施例中，dpc 兼容性工具4302提供嵌入代码中的提示或工具，以帮助开发人员手动迁移无法自动迁移的附加代码。在至少一个实施例中，迁移是源文件、项目或应用程序的一次性活动。
[0405]
在至少一个实施例中，dpc 兼容性工具4302能够成功地将cuda代码的所有部分迁移到dpc ，并且可以简单地存在用于手动验证和调节所生成的dpc 源代码的性能的可选步骤。在至少一个实施例中，dpc 兼容性工具4302直接生成由dpc 编译器编译的dpc 源代码4308，而不需要或不利用人工干预来修改由dpc 兼容性工具4302生成的dpc 代码。在至少一个实施例中，dpc 兼容性工具生成可编译的dpc 代码，开发人员可以可选地对性能、可读性、可维护性和其他各种考虑因素或其任何组合进行调节。
[0406]
在至少一个实施例中，至少部分地使用dpc 兼容性工具4302将一个或更多个cuda源文件迁移到dpc 源文件。在至少一个实施例中，cuda源代码包括一个或更多个报头文件，这些报头文件可包括cuda报头文件。在至少一个实施例中，cuda源文件包括可用于打印文本的《cuda.h》报头文件和《stdio.h》报头文件。在至少一个实施例中，向量加法内核cuda源文件的一部分可以写成或相关于：
[0407]
[0408]
[0409][0410]
在至少一个实施例中并且结合上述cuda源文件，dpc 兼容性工具4302解析cuda源代码，并用适当的dpc 和sycl报头文件替换报头文件。在至少一个实施例中，dpc 报头文件包括助手声明。在cuda中，有线程id的概念，相应地，在dpc 或sycl中，对于每个元素都有一本地标识符。
[0411]
在至少一个实施例中，并结合上述cuda源文件，存在初始化的两个向量a和b，并且向量相加结果作为vectoraddkernel()的一部分被放入向量c中。在至少一个实施例中，作为将cuda代码迁移到dpc 代码的一部分，dpc 兼容性工具4302经由本地id将用于索引工作元素的cuda线程id转换为用于工作元素的sycl标准寻址。在至少一个实施例中，可以优化由dpc 兼容性工具4302生成的dpc 代码——例如，通过降低nd_item的维数，从而增加存储器和/或处理器利用率。
[0412]
在至少一个实施例中，并且结合上述cuda源文件，迁移存储器分配。在至少一个实施例中，cudamalloc()被迁移到统一共享存储器sycl调用malloc_device()，设备和上下文被传递到该统一共享存储器sycl调用，这依赖于诸如平台、设备、上下文和队列之类的sycl概念。在至少一个实施例中，sycl平台可以具有多个设备(例如，主机和gpu设备)；设备可能有可以向其提交作业的多个队列；每个设备可能有一上下文；并且上下文可以有多个设备并管理共享存储器对象。
[0413]
在至少一个实施例中并结合上述cuda源文件，main()函数调用或调用vectoraddkernel()，以将两个向量a和b相加，并将结果存储在向量c中。在至少一个实施例中，调用vectoraddkernel()的cuda代码被dpc 代码替换，以将内核提交到命令队列以供执行。在至少一个实施例中，命令组处理程序cgh传递提交到队列的数据、同步和计算，parallel_for被调用用于调用vectoraddkernel()的工作组中的多个全局元素和多个工作项。
[0414]
在至少一个实施例中，并结合上述cuda源文件，将复制设备存储器然后释放存储器用于向量a、b和c的cuda调用迁移到对应的dpc 调用。在至少一个实施例中，c 代码(例如，用于打印浮点变量的向量的标准iso c 代码)被原样迁移，而无需由dpc 兼容性工具4302修改。在至少一个实施例中，dpc 兼容性工具4302修改cuda api，以用于存储器设置和/或主机调用，以在加速设备上执行内核。在至少一个实施例中并结合上述cuda源文件，相应的人类可读dpc 4304(例如，可编译的)被编写为或相关于：
[0415]
[0416]
[0417][0418]
在至少一个实施例中，人类可读dpc 4304指的是由dpc 兼容性工具4302生成的输出，并且可以以一种或另一种方式进行优化。在至少一个实施例中，由dpc 兼容性工具4302生成的人类可读dpc 4304可以在迁移之后由开发人员手动编辑，以使其更易于维护、性能或其他考虑。在至少一个实施例中，由dpc 兼容性工具4302生成的dpc 代码，例如所公开的dpc ，可以通过为每个malloc_device()调用去除对get_current_device()和/或get_default_context()的重复调用来优化。在至少一个实施例中，上面生成的dpc 代码使用了3维的nd_range，它可以被重构(refactored)为只使用单个维度，从而减少了存储器的使用。在至少一个实施例中，开发人员可以手动编辑由dpc 兼容性工具4302生成的dpc 代码，从而用访问器替换统一共享存储器的使用。在至少一个实施例中，dpc 兼容性工具4302具有改变其如何将cuda代码迁移到dpc 代码的选项。在至少一个实施例中，dpc 兼容性工具4302是冗长的，因为它使用通用模板将cuda代码迁移到适用于大量情况的dpc 代码。
[0419]
在至少一个实施例中，cuda到dpc 的迁移工作流包括以下步骤：使用拦截-构建脚本准备迁移；使用dpc 兼容性工具4302执行cuda项目到dpc 的迁移；人工审查和编辑迁移的源文件，以确保完整性和正确性；以及编译最终的dpc 代码以生成dpc 应用程序。在至少一个实施例中，在一种或更多种场景中可能需要人工审查dpc 源代码，包括但不限于：迁移的api不返回错误代码(cuda代码可以返回错误代码，该错误代码随后可以被应用程序使用，但是sycl使用异常(exception)来报告错误，因此不会使用错误代码来显露错误)；dpc 不支持cuda计算能力相关逻辑；无法删除语句。在至少一个实施例中，dpc 代码需要人工干预的场景可以包括但不限于：错误代码逻辑替换为(*,0)代码或注释掉；等效的dpc api不可用；cuda计算依赖于能力的逻辑；依赖于硬件的api(clock())；缺少不支持api的特征；执行时间测量逻辑；处理内置向量类型冲突；cublas api的迁移；以及更多。
[0420]
其他变型在本公开的精神内。因此，尽管公开的技术易于进行各种修改和替代构造，但是某些示出的其实施例在附图中示出并且已经在上面进行了详细描述。然而，应理解，无意将公开内容限制为所公开的一种或更多种特定形式，而是相反，其意图是涵盖落入如所附权利要求书所定义的本公开内容的精神和范围内的所有修改、替代构造和等同物。
[0421]
除非另有说明或显然与上下文矛盾，否则在描述所公开的实施例的上下文中(特别是在所附权利要求的上下文中)，术语“一”和“一个”和“该”以及类似指代的使用应被解释为涵盖单数和复数，而不是作为术语的定义。除非另有说明，否则术语“包括”、“具有”、“包含”和“含有”应被解释为开放式术语(意味着“包括但不限于”)。术语“连接”(在未经修改时指的是物理连接)应解释为部分或全部包含在内、附接到或连接在一起，即使有某些介入。除非本文另外指出，否则本文中对数值范围的引用仅旨在用作分别指代落入该范围内
的每个单独值的简写方法，并且每个单独值都被并入说明书中，就如同其在本文中被单独叙述一样。除非另外指出或与上下文矛盾，否则术语“集”(例如“项目集”)或“子集”的使用应解释为包括一个或更多个成员的非空集合。此外，除非另外指出或与上下文矛盾，否则术语相应集的“子集”不一定表示对应集的适当子集，而是子集和对应集可以相等。
[0422]
除非以其他方式明确指出或与上下文明显矛盾，否则诸如“a，b和c中的至少一个”或“a，b与c中的至少一个”形式的短语之类的连接语在上下文中理解为通常用来表示项目、条款等，其可以是a或b或c，也可以是a和b和c集的任何非空子集。例如，在具有三个成员的集的说明性示例中，连接短语“a，b和c中的至少一个”和“a，b与c中的至少一个”是指以下任意集：{a}，{b}，{c}，{a，b}，{a，c}，{b，c}，{a，b，c}。因此，这种连接语言通常不旨在暗示某些实施例要求存在a中的至少一个，b中的至少一个和c中的至少一个。另外，除非另有说明或与上下文矛盾，否则术语“多个”表示复数的状态(例如，“多个项目”表示多个项目)。多个项目中项目的数量至少为两个，但如果明确指示或通过上下文指示，则可以更多。此外，除非另有说明或从上下文中可以清楚得知，否则短语“基于”是指“至少部分基于”而不是“仅基于”。
[0423]
除非本文另外指出或与上下文明显矛盾，否则本文描述的过程的操作可以任何合适的顺序执行。在至少一个实施例中，诸如本文所述的那些过程(或其变形和/或其组合)之类的过程在配置有可执行指令的一个或更多个计算机系统的控制下执行，并且被实现为代码(例如，可执行指令，一个或更多个计算机程序或一个或更多个应用程序)，该代码通过硬件或其组合在一个或更多个处理器上共同执行。在至少一个实施例中，代码以例如计算机程序的形式存储在计算机可读存储介质上，该计算机程序包括可由一个或更多个处理器执行的多个指令。在至少一个实施例中，计算机可读存储介质是非暂时性计算机可读存储介质，其排除了暂时性信号(例如，传播的瞬态电或电磁传输)，但包括非暂时性数据存储电路(例如，缓冲区、高速缓存和队列)。在至少一个实施例中，代码(例如，可执行代码或源代码)被存储在其上存储有可执行指令的一组一个或更多个非暂时性计算机可读存储介质(或用于存储可执行指令的其他存储器)上，该可执行指令在由计算机系统的一个或更多个处理器执行时(即，作为被执行的结果)，使得计算机系统执行本文所述的操作。在至少一个实施例中，一组非暂时性计算机可读存储介质包括多个非暂时性计算机可读存储介质，并且多个非暂时性计算机可读存储介质中的个体非暂时性存储介质中的一个或更多个缺少全部代码，而是多个非暂时性计算机可读存储介质共同存储全部代码。在至少一个实施例中，可执行指令被执行，以使得不同的指令由不同的处理器执行，例如，非暂时性计算机可读存储介质存储指令，并且主中央处理单元(“cpu”)执行一些指令，而图形处理单元(“gpu”)执行其他指令。在至少一个实施例中，计算机系统的不同组件具有单独的处理器，并且不同的处理器执行指令的不同子集。
[0424]
因此，在至少一个实施例中，计算机系统被配置为实现单独地或共同地执行本文所述的过程的操作的一个或更多个服务，并且这样的计算机系统被配置有使能实施操作的适用的硬件和/或软件。此外，实现本公开的至少一个实施例的计算机系统是单个设备，并且在另一实施例中是分布式计算机系统，其包括以不同方式操作的多个设备，使得分布式计算机系统执行本文所述的操作，并且使得单个设备不执行所有操作。
[0425]
本公开的至少一个实施例可以鉴于以下条款进行描述：
[0426]
1.一种系统，包括：
[0427]
至少一个处理器；
[0428]
至少一个存储器，包括指令，所述指令响应于由所述至少一个处理器执行，使所述系统至少：
[0429]
从与虚拟场景相关联的多个光中选择光集合，所述光集合至少部分地基于第一一个或更多个随机因素来选择；
[0430]
至少部分地基于第二一个或更多个随机因素，从所述光集合中选择光子集，所述光子集将用于生成图形帧，其中选择所述光子集以使得指示所述光子集的数据具有小于被预测为在处理器高速缓存中可用的存储器的总大小；和
[0431]
至少部分地基于所述光子集渲染所述图形帧的像素。
[0432]
2.根据条款1所述的系统，所述至少一个存储器还包括指令，所述指令响应于由所述至少一个处理器执行，使所述系统至少：
[0433]
选择附加光子集，以用于生成附加的图形帧，所述附加光子集是从与所述虚拟场景相关联的所述多个光中选择的。
[0434]
3.根据条款1或2所述的系统，所述至少一个存储器还包括指令，所述指令响应于由所述至少一个处理器执行，使所述系统至少：
[0435]
使用从所述光集合中选择的第二光子集来渲染所述图形帧的第二像素。
[0436]
4.根据条款1-3中任一项所述的系统，其中所述像素是所述图形帧的第一图块中的多个像素中的一个像素，并且其中使用不同的光子集来渲染第二图块中的像素。
[0437]
5.根据条款1-4中任一项所述的系统，所述至少一个存储器还包括指令，所述指令响应于由所述至少一个处理器执行，使所述系统至少：
[0438]
使用所述光子集来渲染所述图形帧的多个图块，所述多个图块在所述图形帧中不连续，所述多个图块被连续渲染以保持所述光子集驻留在所述处理器高速缓存中。
[0439]
6.根据条款1-5中任一项所述的系统，所述至少一个存储器还包括指令，所述指令响应于由所述至少一个处理器执行，使所述系统至少：
[0440]
将来自所述光子集的一个或更多个光混洗到附加光子集中；以及
[0441]
使用所述附加光子集来渲染所述图形帧的附加像素。
[0442]
7.根据条款1-6中任一项所述的系统，其中所述光子集被选择为具有小于在渲染所述图形帧期间可用的处理器高速缓存存储器的量的总大小。
[0443]
8.根据条款1-7中任一项所述的系统，其中加权所述第一一个或更多个随机因素，以有利于至少部分地基于所选光的强度来选择光。
[0444]
9.一种方法，包括：
[0445]
至少部分地基于第一一个或更多个随机因素，从与虚拟场景相关联的多个光中选择光集合；
[0446]
至少通过以下步骤使用来自所述光集合的光来生成图形帧：
[0447]
至少部分地基于第二一个或更多个随机因素，从所述光集合中选择光子集；以及
[0448]
至少部分地基于所述光子集来渲染所述图形帧的像素。
[0449]
10.根据条款9所述的方法，还包括：
[0450]
使用附加光集合来生成附加图形帧，所述附加光集合是从与所述虚拟场景相关联
的所述多个光中选择的。
[0451]
11.根据条款9或10所述的方法，还包括：
[0452]
使用从所述光集合中选择的第二光子集来渲染所述图形帧的第二像素。
[0453]
12.根据条款9-11中任一项所述的方法，还包括：
[0454]
使用不同的光子集来渲染所述图形帧的不同图块中的像素。
[0455]
13.根据条款9-12中任一项所述的方法，还包括：
[0456]
使用所述光子集来渲染所述图形帧的多个图块，所述多个图块在所述图形帧中不连续但连续渲染。
[0457]
14.根据条款9-13中任一项所述的方法，还包括：
[0458]
选择所述光子集，以使其大小小于处理器高速缓存大小。
[0459]
15.根据条款9-14中任一项所述的方法，还包括：
[0460]
至少部分地基于随机选择所述光子集的至少一部分以将其包括在附加光子集中，来生成所述附加光子集；以及
[0461]
使用所述附加光子集来生成所述图形帧的附加图块。
[0462]
16.根据条款9-15中任一项所述的方法，还包括：
[0463]
通过至少用从所述光集合中选择的一个或更多个附加光替换所述光子集的一部分，来生成附加光子集。
[0464]
17.一种非暂时性计算机可读存储介质，包括指令，所述指令响应于由计算设备的至少一个处理器执行，使所述计算设备至少：
[0465]
从与虚拟场景相关联的光中选择光集合，所述光集合至少部分地随机选择，其中至少部分地基于所述光集合来渲染图形帧；以及
[0466]
使用来自所述光集合的光子集来生成所述图形帧的一部分，所述光子集至少部分地从所述光集合中随机选择，其中所述图形帧的所述一部分的像素至少部分地基于所述光子集来渲染。
[0467]
18.根据条款17所述的非暂时性计算机可读存储介质，还包括指令，所述指令响应于由所述计算设备的至少一个处理器执行，使所述计算设备至少：
[0468]
使用附加光集合来生成附加图形帧，所述附加光集合从与所述虚拟场景相关联的所述光中选择。
[0469]
19.根据条款17或18所述的非暂时性计算机可读存储介质，还包括指令，所述指令响应于由所述计算设备的至少一个处理器执行，使所述计算设备至少：
[0470]
使用不同的光子集来渲染所述图形帧的不同部分中的像素。
[0471]
20.根据条款17-19中任一项所述的非暂时性计算机可读存储介质，还包括指令，所述指令响应于由所述计算设备的至少一个处理器执行，使所述计算设备至少：
[0472]
使用所述光子集来渲染所述图形帧的多个部分，所述多个部分被连续渲染。
[0473]
21.根据条款17-20中任一项所述的非暂时性计算机可读存储介质，还包括指令，所述指令响应于由所述计算设备的至少一个处理器执行，使所述计算设备至少：
[0474]
至少部分地基于从所述光子集中选择一个或更多个光来生成附加光子集。
[0475]
22.根据条款17-21中任一项所述的非暂时性计算机可读存储介质，其中在所述图形帧的一个或更多个部分的渲染期间，所述光子集中的大多数光仍然驻留在处理器高速缓
存中。
[0476]
23.根据条款17-22中任一项所述的非暂时性计算机可读存储介质，其中至少部分地基于所选光的强度或到所选光的距离中的至少一个来选择所述光集合和光子集。
[0477]
本文提供的任何和所有示例或示例性语言(例如，“诸如”)的使用仅旨在更好地阐明本公开的实施例，并且不对公开的范围构成限制，除非另有要求。说明书中的任何语言都不应被解释为表示任何未要求保护的要素对于实践公开内容是必不可少的。
[0478]
本文引用的所有参考文献，包括出版物、专利申请和专利，均通过引用并入本文，其程度就如同每个参考文献被单独且具体地指示为以引用的方式并入本文并且其全部内容在本文中阐述一样。
[0479]
在说明书和权利要求中，可以使用术语“耦合”和“连接”以及它们的派生词。应当理解，这些术语可能不旨在作为彼此的同义词。相反，在特定示例中，“连接”或“耦合”可用于指示两个或更多个元件彼此直接或间接物理或电接触。“耦合”也可能意味着两个或更多个元素彼此不直接接触，但仍彼此协作或交互。
[0480]
除非另有明确说明，否则可以理解，在整个说明书中，诸如“处理”、“计算”、“计算”、“确定”等之类的术语，是指计算机或计算系统或类似的电子计算设备的动作和/或过程，其将计算系统的寄存器和/或存储器中表示为物理量(例如电子)的数据处理和/或转换为类似表示为计算系统的存储器、寄存器或其他此类信息存储、传输或显示设备中的物理量的其他数据。
[0481]
以类似的方式，术语“处理器”可以指处理来自寄存器和/或存储器的电子数据并将该电子数据转换成可以存储在寄存器和/或存储器中的其他电子数据的任何设备或存储器的一部分。作为非限制性示例，“处理器”可以是cpu或gpu。“计算平台”可以包括一个或更多个处理器。如本文所使用的，“软件”进程可以包括例如随时间执行工作的软件和/或硬件实体，诸如任务、线程和智能代理。同样，每个过程可以指代多个过程，以连续地或间歇地顺序地或并行地执行指令。术语“系统”和“方法”在本文中可以互换使用，只要系统可以体现一种或更多种方法，并且方法可以被认为是系统。
[0482]
在本文件中，可以参考获得、获取、接收或将模拟或数字数据输入子系统、计算机系统或计算机实现的机器中。可以通过多种方式来完成获得、获取、接收或输入模拟和数字数据的过程，例如通过接收作为函数调用或对应用程序编程接口的调用的参数的数据。在一些实现方式中，可以通过经由串行或并行接口传输数据来完成获得、获取、接收或输入模拟或数字数据的过程。在另一实现方式中，可以通过经由计算机网络将数据从提供实体传输到获取实体来完成获得、获取、接收或输入模拟或数字数据的过程。也可以参考提供、输出、传送、发送或呈现模拟或数字数据。在各种示例中，提供、输出、传送、发送或呈现模拟或数字数据的过程可以通过将数据作为函数调用的输入或输出参数、应用程序编程接口或进程间通信机制的参数进行传输来实现。
[0483]
尽管上面的讨论阐述了所描述的技术的示例实现，但是其他架构可以用于实现所描述的功能，并且旨在落入本公开的范围内。此外，尽管出于讨论的目的在上面定义了具体的职责分配，但是根据情况，可以以不同的方式分配和划分各种功能和职责。
[0484]
此外，尽管已经用特定于结构特征和/或方法动作的语言描述了主题，但是应当理解，所附权利要求书所要求保护的主题不必限于所描述的特定特征或动作。而是，公开了特
定的特征和动作作为实现权利要求的示例性形式。