CN115104119A - 生成插值视频帧的技术 - Google Patents

Abstract

用于生成插值视频帧的装置、系统和技术。在至少一个实施例中，至少部分地基于一个或更多个对象从第一视频帧到第二视频帧的多个可能运动中的一个来生成插值视频帧。

Description

生成插值视频帧的技术

相关申请的交叉引用

本申请要求于2020年7月30日提交的题为“生成插值视频帧的技术 (TECHNIQUESTO GENERATE INTERPOLATED VIDEO FRAMES)”的美国专利申请第16/944,066号的优先权，其全部内容通过引用整体并入本文并用于所有目的。

技术领域

至少一个实施例涉及用于生成视频帧的处理资源。例如，至少一个实施例涉及用于根据本文描述的各种新技术生成插值视频帧的处理器或计算系统。

背景技术

生成视频帧可能使用大量的存储器、时间或计算资源。用于生成视频帧的存储器、时间或计算资源的数量可以得到改善。

附图说明

图1是图示出根据至少一个实施例的用于生成插值视频帧的系统的框图；

图2是图示出根据至少一个实施例的帧插值的框图；

图3是图示出根据至少一个实施例的中间帧合成的图；

图4是图示出根据至少一个实施例的向前和向后扭曲的图；

图5是图示出根据至少一个实施例的用于生成插值视频帧的系统的框图；

图6是图示出根据至少一个实施例的帧插值的框图；

图7是图示出根据至少一个实施例的中间帧合成的图；

图8示出了根据至少一个实施例的生成中间视频帧的技术的流程图；

图9示出了根据至少一个实施例的对第一组像素数据进行采样的技术的流程图；

图10示出了根据至少一个实施例的对第二组像素数据进行采样的技术的流程图；

图11图示了根据至少一个实施例的生成中间视频帧的技术的流程图；

图12图示了根据至少一个实施例的第一输入视频帧和第二输入视频帧；

图13图示了根据至少一个实施例的覆盖中间视频帧的遮挡掩模和去遮挡掩模；

图14图示了根据至少一个实施例的生成的中间视频帧；

图15示出了根据至少一个实施例的示例性数据中心；

图16示出了根据至少一个实施例的处理系统；

图17示出了根据至少一个实施例的计算机系统；

图18示出了根据至少一个实施例的系统；

图19示出了根据至少一个实施例的示例性集成电路；

图20示出了根据至少一个实施例的计算系统；

图21示出了根据至少一个实施例的APU；

图22示出了根据至少一个实施例的CPU；

图23示出了根据至少一个实施例的示例性加速器集成切片；

图24A-24B示出了根据至少一个实施例的示例性图形处理器；

图25A示出了根据至少一个实施例的图形核心；

图25B示出了根据至少一个实施例的GPGPU；

图26A示出了根据至少一个实施例的并行处理器；

图26B示出了根据至少一个实施例的处理集群；

图26C图示了根据至少一个实施例的图形多处理器；

图27示出了根据至少一个实施例的图形处理器；

图28示出了根据至少一个实施例的处理器；

图29示出了根据至少一个实施例的处理器；

图30示出了根据至少一个实施例的图形处理器核；

图31示出了根据至少一个实施例的PPU；

图32示出了根据至少一个实施例的GPC；

图33示出了根据至少一个实施例的流式多处理器；

图34示出了根据至少一个实施例的编程平台的软件栈；

图35图示了根据至少一个实施例的图34的软件栈的CUDA实现；

图36示出了根据至少一个实施例的图34的软件栈的ROCm实现；

图37图示了根据至少一个实施例的图34的软件栈的OpenCL实现；

图38示出了根据至少一个实施例的由编程平台支持的软件；

图39示出了根据至少一个实施例的编译代码以在图34-37的编程平台上执行；

图40更详细地说明根据至少一个实施例编译代码以在图34-37的编程平台上执行；

图41示出了根据至少一个实施例在编译源代码之前翻译源代码；

图42A图示了根据至少一个实施例的被配置为使用不同类型的处理单元编译和执行CUDA源代码的系统；

图42B图示了根据至少一个实施例的被配置为使用CPU和启用CUDA 的GPU编译和执行图42A的CUDA源代码的系统；

图42C图示了根据至少一个实施例的被配置为使用CPU和非启用 CUDA的GPU编译和执行图42A的CUDA源代码的系统；

图43图示了根据至少一个实施例的由图42C的CUDA到HIP转换工具转换的示例性内核；

图44更详细地示出了根据至少一个实施例的图42C的非启用CUDA 的GPU；以及

图45图示了根据至少一个实施例的示例性CUDA网格的线程如何被映射到图44的不同计算单元。

具体实施方式

在以下描述中，阐述了许多具体细节以提供对至少一个实施例的更透彻的理解。然而，对于本领域技术人员来说显而易见的是，可以在没有这些特定细节中的一个或更多个的情况下实践本发明构思。

图1是根据至少一个实施例的生成插值视频帧的系统100的框图。在至少一个实施例中，系统100包括插值帧生成器102，其至少部分地基于帧数据106生成一组插值帧104。在至少一个实施例中，插值帧生成器102 是可通过应用编程接口(API)访问的计算机程序。在至少一个实施例中，插值帧生成器102在处理器(例如，中央处理单元(CPU)，为了清楚未示出)上运行。应当理解，当一个或更多个计算机程序和/或API(例如，与插值帧生成器102相关)被称为执行关于实施例的动作或技术的一个方面时，运行计算机程序和/或API的计算机系统的一个或更多个硬件组件(例如，CPU、GPU和/或其他硬件组件)执行动作或技术的方面。

在至少一个实施例中，帧数据106包括第一视频帧信息108和第二视频帧信息110。在至少一个实施例中，第一视频帧信息108包括第一视频帧 112和一组深度指示符114。在至少一个实施例中，第二视频帧信息110包括第二视频帧116和一组深度指示符118。在至少一个实施例中，第一视频帧112包括第一组像素值并且第二视频帧116包括第二组像素值。在至少一个实施例中，第一视频帧112和第二视频帧116的像素值集合包括像素 (例如，在由x和y值表示的二维网格中的位置处)的颜色信息(例如RGB值)。在至少一个实施例中，深度指示符组114是第一视频帧112的像素的一组深度值(例如，z缓冲器中的z值)。在至少一个实施例中，深度指示符组 118是第二视频帧116的像素的一组深度值。在至少一个实施例中，一些或全部帧数据106可在来自视频游戏引擎的一个或更多个缓冲器(例如，在诸如图形存储器之类的存储器中)中可用，并且可由插值帧生成器102 使用。

在至少一个实施例中，第一视频帧信息108包括一组前向运动向量120，而第二视频帧信息110包括一组后向运动向量122。在至少一个实施例中，一组前向运动向量120包括第一视频帧112中的像素的运动向量，其指向第二视频帧116中的像素位置。在至少一个实施例中，一组后向运动向量 122包括第二视频帧116中的像素的运动向量，其指向第一视频帧112中的像素位置。在至少一个实施例中，每个前向运动向量是投影的顶点运动(例如，到第二个视频帧中的像素位置)，并且每个后向运动向量是投影的顶点运动(例如，从第一个视频帧中的像素位置)。尽管关于帧数据106仅示出了两组视频帧信息，但应当理解，在至少一个实施例中，帧数据106包括附加的视频帧信息(例如，出现在帧112之前和/或帧116之后的视频帧的信息)。在至少一个实施例中，特定视频帧的视频帧信息包括前向和后向运动向量(例如，在至少一个实施例中，第二视频帧信息110还包括指向后续视频帧的像素的前向运动向量，为了清楚未示出)。

在至少一个实施例中，插值帧生成器102在第一视频帧112和第二视频帧116之间生成第一插值帧124、第二插值帧126和第三插值帧128。在至少一个实施例中，插值帧生成器102在第一视频帧112和第二视频帧116 之间生成与所示不同数量的帧(例如，单个帧或多于三个帧)。

图2是示出根据至少一个实施例的至少部分地基于前向运动向量(例如，前向运动向量120)和后向运动向量(例如，后向运动向量122)的帧插值200的图。在至少一个实施例中，插值帧生成器(例如插值帧生成器102) 执行帧插值200。在至少一个实施例中，帧插值200被称为双向运动向量 (BiMV)插值技术。在至少一个实施例中，插值帧生成器至少部分地基于在时间(t)、(t+1)和(t+2)所示的帧的视频帧信息生成插值帧(例如，使用像素颜色数据、深度和运动向量)。在至少一个实施例中，在给定帧对(例如，在时间(t)和(t+1))的情况下，插值帧生成器在预定时间(例如，在(t+0.25))合成插值帧。

在至少一个实施例中，插值帧生成器至少部分地基于输入RGB彩色图像(例如，第一视频帧112和第二视频帧116)、深度缓冲器(例如，深度指示符114和118)和几何运动向量(例如，前向运动向量120和后向运动向量122)来执行帧插值200。在至少一个实施例中，这些输入被表示为 (I₀,D₀,F_0→1)和(I₁,D₁,F_1→0),它们分别对应于时间T＝0和T＝1时的输入RGB 彩色图像、深度缓冲器和几何运动向量。在至少一个实施例中，输入的RGB 彩色图像是每秒低帧(FPS)图像。

在至少一个实施例中，帧插值200至少部分地基于以下假设：像素沿着与源像素和目的地像素相关的运动向量F_0→1和F_1→0从T＝0到T＝1(或T＝1 到T＝0)线性运动。在至少一个实施例中，当插值帧生成器在时间t＝0.5时生成中间插值帧时，插值帧生成器将中间帧估计为沿线性路径在t处(例如，t＝0.5，或中途)着陆的像素。在至少一个实施例中，帧插值200至少部分地基于加速度参数(例如，由插值帧生成器102至少部分地基于使用来自T＝0之前的一个或更多个附加帧的信息推断出的加速度)。

在至少一个实施例中，在大多数情况下，仅将像素从T＝0或T＝1移动到中间时间t就足以获得合理的中间像素。然而，在至少一个实施例中，像素的外观可能会从T＝0到T＝1急剧变化，例如当光线充足的场景变暗时。在至少一个实施例中，几何运动向量不能完全解释这种变化，并且仅从单个源移动像素将导致次优插值。在至少一个实施例中，为了说明这一点，中间时间t处的像素是由沿着由运动向量定义的线性线从两个源图像取出的像素的线性组合合成的。在至少一个实施例中，这在数学上表示为，

I_t＝(1-t)*Warp(I₀，F_t→0)+t*Warp(I₁，F_t→1)， (1)

其中F_t→0和F_t→1是由插值帧生成器(例如，如关于图4所述)生成的双向中间运动向量，而Warp是一种双线性采样操作，它从由运动向量引导的源RGB图像中拉出像素。

图3是图示出根据至少一个实施例的至少部分地基于前向运动向量 (例如，前向运动向量120)和后向运动向量(例如，后向运动向量122)的中间帧合成300的图。在至少一个实施例中，中间帧合成300是出于说明目的的简化表示，其仅示出了在不同时间的一维中的五个像素，而不是如将在至少一个实施例中使用的二维中的大量像素。在至少一个实施例中，插值帧生成器102执行中间帧合成300。在至少一个实施例中，中间帧合成 300被称为双向运动向量(BiMV)中间帧合成。在至少一个实施例中，中间帧合成300包括中间运动向量(F_t→1,F_t→0)，并且中间帧合成300由遮挡掩模和去遮挡掩模(occlusion and dis-occlusion masks)引导。在至少一个实施例中，圆圈302表示从T＝1开始扭曲帧的遮挡伪影，而圆圈304表示从T＝0 开始扭曲的去遮挡伪影。

在至少一个实施例中，源像素可能在针对所有中间像素的两个输入图像中都不可见，因为它们已被遮挡。例如，在至少一个实施例中，用F_t→0采样I₀将产生后沿伪影(也称为去遮挡伪影)，因为填充被遮挡区域的正确内容在I₀中是不可见的，但可能在I₁(第二张图像)中可见。在至少一个实施例中，类似地，用F_t→1对I₁进行采样将导致前沿伪影(也称为遮挡伪影)。然而，对于大多数中间像素，源像素在两个输入帧中都是可见的，因此等式(1) 将导致在这些区域中进行合理的插值。

在至少一个实施例中，为了考虑等式(1)未考虑的前沿和后沿伪影，等式(2)用于对被遮挡区域、被去遮挡区域和可见区域建模，如下所示：

其中M_0→t和M_1→t是遮挡掩模和去遮挡掩模。在至少一个实施例中，M_0→t允许对仅从输入I₀可见的像素进行采样，并且M_1→t允许对仅从输入I₁可见的像素进行采样。在至少一个实施例中，图3图示了使用等式(2)的帧合成的示例过程。在至少一个实施例中，插值帧生成器针对从输入I₀或I₁不可见的中间帧的像素(例如，同时被去遮挡和被遮挡)对附加输入帧(例如，在其中对应像素可见的I₀之前发生的第一帧)的像素进行采样。

在至少一个实施例中，插值帧生成器至少部分地基于是否应用遮挡和/ 或去遮挡掩模，和/或基于插值帧相对于第一视频帧和第二视频帧的相对时间值(例如，插值帧是否对应于时间t＝0.25、t＝0.5或t＝0.75)，将加权因子应用于从来自第一视频帧(例如，在T＝0)和第二视频帧(例如，在T＝1) 的像素中拉出的像素数据。在至少一个实施例中，对于受遮挡掩模(例如，圆圈302)影响的像素，插值帧生成器对从第一帧(例如，从T＝0到T＝t)拉出的像素应用1的加权因子，并且对从第二帧(例如，从T＝1到T＝t) 拉出的像素应用0的加权因子，如图3所示。在至少一个实施例中，对于受去遮挡掩模影响的像素(例如，圆圈304)，插值帧生成器对从第一帧(例如，从T＝0到T＝t)拉出的像素应用加权因子0，并且对从第二帧(例如，从T＝1到T＝t)拉出的像素应用加权因子1，如图3所示。在至少一个实施例中，对于不受遮挡掩模或去遮挡掩模影响的像素，插值帧生成器根据插值帧到第一视频帧和第二视频帧的时间相对距离来应用加权因子(例如，当插值帧介于两者之间时，将0.5应用于两者，如图3所示)。

图4是图示出根据至少一个实施例的向后和向前扭曲400的图。在至少一个实施例中，对于前向扭曲，源中的每个像素被迁移到由从源指向目的地的运动向量引导的目的地像素。在至少一个实施例中，对于前向扭曲，可能永远不会到达目的地像素并为其分配值(例如，在像素组402中)，或者多个源像素可能指向相同的目的地像素(例如，像素404)，从而在前向扭曲中产生模糊。

在至少一个实施例中，在图3的中间帧合成300中，插值帧生成器(例如，插值帧生成器102)通过投影通过中间像素位置的输入双向运动向量 (F_0→1,F_1→0)来近似中间运动向量(F_t→1，F_t→0)。在至少一个实施例中，插值帧生成器将输入运动矢量从T＝0或T＝1向前扭曲到T＝t，然后在T＝t中的目的地像素处，插值帧生成器在按(1-t)的因子缩放后，在源像素处分配运动向量。在至少一个实施例中，当t小时，比例因子更大，反之亦然。

在至少一个实施例中，插值帧生成器使用源像素处的深度(例如，来自深度指示符114和/或深度指示符118)来解决前向扭曲的模糊性。在至少一个实施例中，在到达同一目的地像素的多个像素中，插值帧生成器记录具有最小深度的源像素的运动向量，即没有被源遮挡的像素。在至少一个实施例中，这在数学上表示为，

F_t→1(x,y)＝(1-t)*F_0→1(x',y')，其中

(x',y')＝(x,y)-t*F_0→1(x,y) (3)

F_0→1(x,y)是一个运动向量，在从T＝0扭曲到T＝t时到达(x,y)的所有可能向量中，具有最小深度

在至少一个实施例中，插值帧生成器生成从T＝t指向T＝0的其他中间运动向量(例如，中间后向运动向量)，如

F_t→0(x，y)＝t*F_1→0(x′，y′)，其中

(x',y')＝(x,y)-(1-t)*F_1→0(x,y) (4)

而F_1→0(x,y)是一个运动向量，在从T＝1扭曲到T＝t时到达(x,y)的所有可能向量中，具有最小深度

在至少一个实施例中，中间帧合成300包括生成掩模以提供在关于等式(2)的后向图像扭曲中出现的遮挡和去遮挡伪影的掩蔽。在至少一个实施例中，中间帧合成300在时间T＝0或T＝1到T＝t时使用深度的前向投影，并将中间像素定义为被遮挡/被去遮挡的像素，如果源像素用其运动向量向前投影时没有源像素可以到达它。在至少一个实施例中，遮挡掩模在数学上表示为，

M_0→t(x,y)＝1，如果D_t(x,y)＝0,，即(x,y)在从1→t向前运动中不可达，否则

M_0→t(x,y)＝0,其中

D_t＝ForwardWarp(D₁,(1-t)F_1→0) (5)

在至少一个实施例中，类似地，去遮挡掩模由以下表示，

M_1→t(x,y)＝1，如果D_t(x,y)＝0,，即，(x,y)在从0→t向前运动中不可达，否则

M_1→t(x,y)＝0,其中

D_t＝ForwardWarp(D₀,tF_0→1) (6)

图5是根据至少一个实施例的生成插值视频帧的系统500的框图。在至少一个实施例中，系统500包括插值帧生成器502，其至少部分地基于帧数据506生成一组插值帧504。在至少一个实施例中，插值帧生成器502 是可通过应用编程接口(API)访问的计算机程序。在至少一个实施例中，插值帧生成器502在处理器(例如，中央处理单元(CPU)，为了清楚未示出)上运行。在至少一个实施例中，一些或全部帧数据506在来自视频游戏引擎的一个或更多个缓冲器(例如，在诸如图形存储器之类的存储器中)中可用，并且可供插值帧生成器502使用。

在至少一个实施例中，帧数据506包括第一视频帧信息508和第二视频帧信息510。在至少一个实施例中，第一视频帧信息508包括第一视频帧512和一组深度指示符514。在至少一个实施例中，前向运动向量在帧数据 506中不可用。在至少一个实施例中，第二视频帧信息510包括第二视频帧 516和一组深度指示符518。在至少一个实施例中，第一视频帧512包括第一组像素值并且第二视频帧516包括第二组像素值。在至少一个实施例中，第一视频帧512和第二视频帧516的像素值集合包括像素(例如，在由x和 y值表示的二维网格中的位置处)的颜色信息(例如RGB值)。在至少一个实施例中，一组深度指示符514是第一视频帧512的像素的一组深度值(例如，z缓冲器中的z值)。在至少一个实施例中，一组深度指示符518是第二视频帧516的像素的一组深度值。

在至少一个实施例中，第一视频帧信息508包括相机数据520并且第二视频帧信息510包括一组后向运动向量522和相机数据524。在至少一个实施例中，相机数据520和相机数据524作为矩阵存储在缓冲器中。在至少一个实施例中，一组向后向运动矢量522包括第二视频帧516中的像素的指向第一视频帧512中的像素位置的运动向量。在至少一个实施例中，每个后向运动向量是顶点运动的投影(例如，从第一视频帧中的像素位置)。尽管关于帧数据506仅示出了两组视频帧信息，但是应当理解，在至少一个实施例中，帧数据506包括附加的视频帧信息(例如，出现在帧512之前和/或帧516之后的视频帧的信息)。

在至少一个实施例中，插值帧生成器502在第一视频帧512和第二视频帧516之间生成第一插值帧526、第二插值帧528和第三插值帧530。在至少一个实施例中，插值帧生成器502在第一视频帧512和第二视频帧516 之间生成与所示不同数量的帧(例如，单个帧或多于三个帧)。

图6是图示出根据至少一个实施例的至少部分地基于后向运动向量(例如，后向运动向量522)的帧插值600的图。在至少一个实施例中，插值帧发生器(例如插值帧发生器502)执行帧插值600。在至少一个实施例中，帧插值600被称为单向运动向量(UniMV)插值技术。在至少一个实施例中，插值帧生成器至少部分地基于在时间(t)、(t+1)和(t+2)所示的帧的视频帧信息来生成插值帧(例如，使用像素颜色数据、深度和后向运动向量)。

在至少一个实施例中，插值帧生成器至少部分地基于输入RGB彩色图像(例如，第一视频帧512和第二视频帧516)、深度缓冲器(例如，深度指示符514和518)和几何运动向量(例如，后向运动向量522)来执行帧插值 600。在至少一个实施例中，这些输入被表示为(I₀,D₀)和(I₁,D₁)以及F_1→0，(I₀,D₀)和(I₁,D₁)分别对应于时间T＝0和T＝1时的输入RGB彩色图像和深度缓冲器，F_1→0是一个单向后向运动向量。在至少一个实施例中，输入RGB 彩色图像是每秒低帧(FPS)图像。

图7是图示出根据至少一个实施例的至少部分地基于后向运动向量(例如，后向运动向量522)的中间帧合成700的图。在至少一个实施例中，插值帧生成器502执行中间帧合成700。在至少一个实施例中，中间帧合成 700是出于说明目的的简化表示，其仅示出了在不同时间的一维中的五个像素，而不是如将在至少一个实施例中使用的二维中的大量像素。在至少一个实施例中，中间帧合成700被称为单向运动向量(UniMV)中间帧合成。在至少一个实施例中，中间帧合成700包括中间运动向量(F_t→1,F_t→0)，并且中间帧合成700由遮挡和去遮挡掩模引导。在至少一个实施例中，圆圈702 表示从T＝1开始扭曲帧的遮挡伪影，而圆圈704表示从T＝0开始扭曲的去遮挡伪影。

在至少一个实施例中，前向运动向量不可用于中间帧合成700，并且不能使用关于图3描述的中间帧合成300的技术。在至少一个实施例中，中间帧合成700包括经由前向扭曲操作从给定的后向运动向量估计丢失的前向运动向量。在至少一个实施例中，前向扭曲操作被称为散射操作。在至少一个实施例中，中间帧合成700包括使用估计的前向运动向量和给定的后向运动向量生成中间运动向量。

在至少一个实施例中，中间帧合成700包括通过散射然后翻转后向向量来估计前向运动向量

在至少一个实施例中，估计前向运动向量在数学上被描述为，

其中

(x′，y′)＝(x，y)-F_1→0(x，y) (7)

是一个运动向量，在从T＝0扭曲到T＝1时到达(x,y)的所有可能向量中，具有最小深度

在至少一个实施例中，这种具有最小深度的信息的散射和保留被中间帧合成700用来解决遮挡，因为具有相同目的地但深度更大的像素在从T＝1移动到T＝0时基本上被遮挡，因此它们被抑制。

在至少一个实施例中，如关于图4所讨论的，前向扭曲导致模糊性，例如目的地像素中的洞。在至少一个实施例中，中间帧合成700包括将这些洞标记为遮挡掩模(当从T＝1向T＝0向前扭曲时)，或去遮挡掩模(当从T＝0向T＝1向前扭曲时)。在至少一个实施例中，使用等式(7)的前向运动向量估计将导致目的地处(T＝0)的洞，在遮挡区域中创建具有未定义运动向量的像素。

在至少一个实施例中，插值帧生成器至少部分地基于是否应用遮挡和 /或去遮挡掩模，和/或基于插值帧相对于第一视频帧和第二视频帧的相对时间值(例如，插值帧是否对应于时间t＝0.25、t＝0.5或t＝0.75)，将加权因子应用于从来自第一视频帧(例如，在T＝0)和第二视频帧(例如，在T＝1) 的像素中拉出的像素数据。在至少一个实施例中，对于受遮挡掩模(例如，圆圈702)影响的像素，插值帧生成器对从第一帧拉出的像素(例如，从 T＝0到T＝t)应用1的加权因子，并且对从第二帧拉出的像素(例如，从 T＝1到T＝t)应用加权因子0，如图7所示。在至少一个实施例中，对于受去遮挡掩模(例如，圆圈704)影响的像素，插值帧生成器对从第一帧(例如，从T＝0到T＝t)拉出的像素应用加权因子0，并且对从第二帧拉出的像素(例如，从T＝1到T＝t)应用加权因子1，如图7所示。在至少一个实施例中，对于不受遮挡掩模或去遮挡掩模影响的像素，插值帧生成器根据插值帧与第一视频帧和第二视频帧的时间相对距离应用加权因子(例如，当插值帧处于中间时，对两者应用0.5，如图7所示)。在至少一个实施例中，插值帧生成器针对在T＝0或T＝1时从帧不可见的中间帧的像素，对附加输入帧(例如，在其中对应像素可见的T＝0之前发生的附加帧，，或在其中对应像素可见的T＝1之后发生的附加帧T＝2)的像素进行采样(例如，同时被去遮挡和被遮挡)。

在至少一个实施例中，插值帧生成器(例如，插值帧生成器502)恢复(例如，生成)T＝0处的遮挡区域(例如，等式(7)产生未定义向量的区域)的一个或更多个运动向量。在至少一个实施例中，插值帧生成器至少部分地基于Procrustes分析的实施来恢复运动向量。在至少一个实施例中，虽然使用这种近似通常导致可接受的中间帧合成，但它只能捕捉静态对象运动，也称为相机运动。在至少一个实施例中，插值帧生成器使用从T＝1 可见的在T＝t处的去遮挡区域。在至少一个实施例中，插值帧生成器在中间运动向量F_t→1的引导下，在去遮挡区域中的每个点与在T＝1时帧中的对应点之间建立一一对应关系。在至少一个实施例中，插值帧生成器以与关于图3和图4描述的方式类似的方式生成中间运动向量，但是当实际前向运动向量不能从缓冲器获得时(例如，不能从系统500中的缓冲器获得)，至少部分基于估计的前向运动向量。在至少一个实施例中，一旦建立对应关系，插值帧生成器为每个像素形成3D点集，将积分索引(x,y)和它们的深度 z视为3D中的一个点(x,y,z)。在至少一个实施例中，插值帧生成器计算登记那些3D点集的矩阵，给出全局或相机平移、旋转和缩放信息的近似估计。在至少一个实施例中，当可从缓冲器获得相机信息(例如，相机数据520和相机数据524)时，插值帧生成器至少部分地基于缓冲的相机信息(例如，在第一视频帧和第二视频帧的时间的相机矩阵)来确定相机位置和/或移动 (例如，平移、旋转和/或缩放)，而不是至少部分地基于3D点集来估计相机移动。

在至少一个实施例中，由于相机矩阵是适用于整个视频帧的全局信息，因此插值帧生成器将相同的相机矩阵外推到其他区域(例如，非去遮挡区域)，包括遮挡区域。在至少一个实施例中，在遮挡区域中未定义前向运动向量，并且将相机矩阵外推到遮挡区域会生成在前向运动向量未定义的情况下使用的近似运动向量。在至少一个实施例中，一旦生成了相机矩阵，插值帧生成器就在每个对应像素上(例如，在具有未定义向量的区域中)创建对应的运动向量，

在至少一个实施例中，从T＝t指向T＝1的最终估计中间向量在数学上由下式给出，

其中

是使用

以与针对图3和4描述的方式类似的方式生成的。在至少一个实施例中，用于使用等式(2)的中间帧合成的其他信息

是在用

替换F_0→1之后，如关于图3和图4所描述的那样生成的。

图8示出了根据至少一个实施例的生成中间视频帧的技术800的流程图。在至少一个实施例中，技术800由至少一个电路、至少一个系统(例如，系统100)、至少一个处理器、至少一个图形处理单元、至少一个并行处理器和/或在本文中描述和/或示出的至少一些其他处理器或其组件执行。在至少一个实施例中，插值帧生成器(例如插值帧生成器102)执行技术800。

在至少一个实施例中，在框802处，技术800包括对第一视频帧(例如，帧112)的第一组像素数据进行采样。在至少一个实施例中，在框804，技术800包括至少部分地基于从第一视频帧到第二视频帧的一组前向运动向量(例如，前向运动向量120)对第二视频帧(例如，帧116)的第二组像素数据进行采样。在至少一个实施例中，在框806处，技术800包括至少部分地基于第一组像素数据和第二组像素数据来生成中间视频帧(例如，插值帧124)。在至少一个实施例中，在框808处，技术800包括执行其他动作。

在至少一个实施例中，执行其他动作包括将生成的中间视频帧存储在缓冲器中。在至少一个实施例中，执行其他动作包括显示中间视频帧。在至少一个实施例中，执行其他动作包括通过网络传输中间视频帧。在至少一个实施例中，执行其他动作包括至少部分地基于第一视频帧、第二视频帧、一组前向运动向量、一组后向运动向量和一组深度指示符，在第一视频帧和第二视频帧之间生成一个或更多个附加中间视频帧。在至少一个实施例中，执行其他动作包括：识别中间视频帧的像素，该像素不具有使用前向运动向量(例如，沿着前向运动向量和/或对应于前向运动向量的函数) 识别的第二视频帧中的对应像素，或使用后向运动向量(例如，沿着后向运动向量和/或对应于后向运动向量的函数)识别的第一视频帧中的对应像素；以及至少部分地基于第一视频帧和第二视频帧中在识别像素的相同位置的像素的深度值，将中间视频帧的识别像素设置为在与第一视频帧或第二视频帧中识别像素相同位置处的像素之一的颜色。在至少一个实施例中，设置识别的像素颜色至少部分地基于将像素颜色设置为在具有最大深度的第一视频帧或第二视频帧中相同位置处的像素的颜色。在至少一个实施例中，执行其他动作包括将一个或更多个插值帧存储在缓冲器中和/或通知(例如，通过设置标志值或发送指示符)插值帧可用于显示的游戏引擎。

图9示出了根据至少一个实施例的对第一组像素数据进行采样的技术 900的流程图。在至少一个实施例中，技术900由至少一个电路、至少一个系统(例如，系统100)、至少一个处理器、至少一个图形处理单元、至少一个并行处理器和/或在此描述和/或示出的至少一些其他处理器或其组件来执行。在至少一个实施例中，插值帧生成器(例如，插值帧生成器102)执行技术900。在至少一个实施例中，技术900是图8的框802的更详细的表示。

在至少一个实施例中，在框902处，技术900包括生成一组中间后向运动向量(例如，包括图3中所示的F_t→0)。在至少一个实施例中，在框904 处，技术900包括生成遮挡掩模(例如，图3中所示的M_0→t)。在至少一个实施例中，在框906处，技术900包括至少部分地基于一组中间后向运动向量和生成的遮挡掩膜对第一视频帧的第一组像素数据进行采样。在至少一个实施例中，在框908处，技术900包括执行其他动作。

图10示出了根据至少一个实施例的对第二组像素数据进行采样的技术1000的流程图。在至少一个实施例中，技术1000由至少一个电路、至少一个系统(例如，系统100)、至少一个处理器、至少一个图形处理单元、至少一个并行处理器和/或在此描述和/或示出的至少一些其他处理器或其组件来执行。在至少一个实施例中，插值帧生成器(例如，插值帧生成器102) 执行技术1000。在至少一个实施例中，技术1000是图8的框804的更详细的表示。

在至少一个实施例中，在框1002处，技术1000包括生成一组中间前向运动向量(例如，包括图3中所示的F_t→1)。在至少一个实施例中，在框1004 处，技术1000包括生成去遮挡掩模(例如，图3中所示的M_1→t)。在至少一个实施例中，在框1006处，技术1000包括至少部分地基于一组中间前向运动向量和生成的去遮挡掩膜对第二组像素数据进行采样。在至少一个实施例中，在框1008处，技术1000包括执行其他动作。

图11示出了根据至少一个实施例的生成中间视频帧的技术1100的流程图。在至少一个实施例中，技术1100由至少一个电路、至少一个系统(例如，系统500)、至少一个处理器、至少一个图形处理单元、至少一个并行处理器和/或在此描述和/或示出的至少一些其他处理器或其组件来执行。在至少一个实施例中，插值帧生成器(例如插值帧生成器502)执行技术1100。

在至少一个实施例中，在框1102处，技术1100包括至少部分地基于一组后向运动向量(例如，后向运动向量522)生成一组估计的前向运动向量 (例如，包括图7所示的

)。在至少一个实施例中，在框1104，技术1100 包括生成一组中间前向运动向量(例如，包括图7中所示的F_t→1)和一组中间后向运动向量(例如，包括图7中所示的F_t→0)。在至少一个实施例中，插值帧生成器至少部分地基于所生成的一组估计的前向运动向量来生成一组中间前向运动向量。在至少一个实施例中，插值帧生成器至少部分地基于一组后向运动向量生成一组中间后向运动向量。

在至少一个实施例中，在框1106处，技术1100包括至少部分地基于所生成的一组中间后向运动向量对第一视频帧(例如，视频帧512)的第一组像素数据进行采样。在至少一个实施例中，对第一组像素数据进行采样包括生成遮挡掩模(例如图7所示的M_0→t)并至少部分地基于生成的遮挡掩模对第一组像素数据进行采样。在至少一个实施例中，在框1108处，技术 1100包括至少部分地基于所生成的一组中间前向运动向量对第二视频帧 (例如，视频帧516)的第二组像素数据进行采样。在至少一个实施例中，对第二组像素数据进行采样包括生成去遮挡掩模(例如图7所示的M_1→t)并且至少部分地基于生成的去遮挡掩模对第二组像素数据进行采样。

在至少一个实施例中，在框1110，技术1100包括至少部分地基于采样的第一组像素数据和采样的第二组像素数据来生成第三视频帧(例如，插值帧520)。在至少一个实施例中，第三视频帧是第一视频帧和第二视频帧之间的中间视频帧。在至少一个实施例中，生成第三视频帧被称为插值第三视频帧。在至少一个实施例中，插值第三视频帧至少部分地基于像素从第一视频帧到第二视频帧的多个可能运动中的一个，其中多个可能的运动对应于指向第一视频帧中的相同目的地像素的一组后向运动向量，并且插值帧生成器至少部分地基于与指向同一目的地像素的后向运动向量的其他源像素相比，与所选择的后向运动向量相关联的源像素的最短深度来选择一组后向运动向量中的一个。在至少一个实施例中，在框1112处，技术1100 包括执行其他动作。

在至少一个实施例中，执行其他动作包括将生成的第三视频帧存储在缓冲器中。在至少一个实施例中，执行其他动作包括显示中间视频帧。在至少一个实施例中，执行其他动作包括通过网络传输中间视频帧。在至少一个实施例中，执行其他动作包括至少部分地基于第一视频帧、第二视频帧、一组后向运动向量(例如，后向运动向量522)和一组深度指示符(例如，深度指示符514和/或深度指示符518)，在第一视频帧和第二视频帧之间生成一个或更多个附加中间视频帧(例如，插值帧522和/或插值帧524)。在至少一个实施例中，执行其他动作包括将一个或更多个插值帧存储在缓冲器中和/或通知(例如，通过设置标志值或发送指示符)插值帧可用于显示的游戏引擎。

图12图示了根据至少一个实施例的第一视频帧1200和第二视频帧 1202。在至少一个实施例中，第一视频帧1200对应于帧112并且第二视频帧1202对应于帧116。在至少一个实施例中，第一视频帧1200对应于帧 512并且第二视频帧1202对应于帧516。在至少一个实施例中，第一视频帧1200对应于技术800的第一视频帧并且第二视频帧1202对应于技术800 的第二视频帧。在至少一个实施例中，第一视频帧1200对应于技术1100 的第一视频帧并且第二视频帧1202对应于技术1100的第二视频帧。在至少一个实施例中，第一视频帧1200出现在视频帧序列中的第二视频帧1202 之前。在至少一个实施例中，像素数据(例如，RGB颜色值)、深度数据以及后向运动向量和前向运动向量中的一个或更多个由视频游戏引擎生成并在显示之前存储在一个或更多个缓冲器中。尽管第一视频帧1200和第二视频帧1202以灰度显示，但应当理解，第一视频帧1200和第二视频帧1202 通常以彩色显示。在至少一个实施例中，可以看出，在第一视频帧1200和第二视频帧1202中表示的角色在与第一视频帧1200相关联的第一时间和与第二视频帧1202相关联的第二更晚时间之间稍微向左移动并稍微抬起枪管。

图13图示了根据至少一个实施例的具有覆盖的去遮挡掩模的中间视频帧1300和具有覆盖的遮挡掩模的中间视频帧1302。在至少一个实施例中，去遮挡掩模包括区域1304，其示出了最初被枪隐藏的在被去遮挡的第一视频帧1200中的区域。在至少一个实施例中，遮挡掩模包括区域1306，其示出的区域对应于在第二视频帧1202中被枪管隐藏的区域。在至少一个实施例中，去遮挡掩模包括附加区域和/或遮挡掩模包括附加区域，然而，为了清楚起见，并非中间视频帧1300和中间视频帧1302中的所有去遮挡区域和遮挡区域都被具体标识。

图14图示了根据至少一个实施例的生成的中间视频帧1400。在至少一个实施例中，生成的中间视频帧1400对应于一组插值帧104中的插值帧。在至少一个实施例中，生成的中间视频帧1400对应于一组插值帧504中的插值帧。在至少一个实施例中，生成的中间视频帧1400对应于技术800的生成的中间视频帧。在至少一个实施例中，生成的中间视频帧1400对应于技术1100的生成的第三视频帧。在至少一个实施例中，对应于中间视频帧 1300的被去遮挡区域(例如，区域1304)的中间视频帧1400的区域包括从第二视频帧1202而非第一视频帧1200采样的像素数据。在至少一个实施例中，与中间视频帧1302的被遮挡区域(例如，区域1306)相对应的中间视频帧1400的区域包括从第一视频帧1200而非第二视频帧1202采样的像素数据。在至少一个实施例中，不对应于中间视频帧1300的被去遮挡区域或中间视频帧1302的被遮挡区域的中间视频帧1400的区域，至少部分地基于从第一视频帧1200和第二视频帧1202两者采样的像素数据进行插值。

在至少一个实施例中，中间视频帧1300、中间视频帧1302和/或中间视频帧1400将有所不同，取决于至少部分地基于由于在对于插值帧组104 中的插值帧的缓冲器中具有可用的前向运动向量而导致的插值差异，中间视频帧是对应于插值帧组104中的插值帧还是插值帧组504中的插值帧，但是使用插值帧组504中的插值帧的后向运动向量和深度信息来估计前向运动向量。在至少一个实施例中，尽管出于说明的目的示出了中间视频帧 1300和中间视频帧1302，但是插值帧生成器(例如，插值帧生成器102或插值帧生成器502)将至少部分地基于遮挡和去遮挡信息(例如，遮挡和去遮挡掩模)生成中间帧1400，而不生成和/或使用具有覆盖的去遮挡和遮挡区域的中间帧(例如，中间视频帧1300和/或中间视频帧1302)。在至少一个实施例中，插值帧生成器(例如，插值帧生成器102或插值帧生成器502)实时(例如，足够快以至于在玩游戏期间生成插值帧)生成插值帧(例如，中间视频帧1400)。在至少一个实施例中，插值帧生成器在视频游戏和/或光线追踪应用中生成插值帧。

在至少一个实施例中，插值帧生成器(例如，插值帧生成器102或插值帧生成器502)通过为低FPS视频合成每秒高帧数(FPS)视频来执行帧插值。在至少一个实施例中，插值帧生成器从给定的一对连续输入帧(例如，第一视频帧1200和第二视频帧1202)合成一个或更多个中间帧。在至少一个实施例中，使用插值帧生成器合成高FPS视频提供优于替代方法(例如使用高速相机，其具有高电力成本、大存储要求和降低的视频分辨率)的优点。在最后一个实施例中，使用插值帧生成器合成高FPS视频提供了优于通过图形处理单元(GPU)在视频游戏中使用初始高帧渲染速率的优势，其代价是降低了图像质量和/或功耗。在至少一个实施例中，使用插值帧生成器合成高FPS视频从低FPS帧序列中产生任意高FPS帧序列，具有最小的延迟或功率要求，这可以提高玩家在玩游戏过程中的游戏性能，并提供优于其他具有更高功率要求(例如，GPU最初的高FPS序列)或更高的延迟 (例如，某些机器学习方法)的方法的优势。在至少一个实施例中，使用插值帧生成器合成高FPS视频不使用来自机器学习(例如，深度学习)方法的推理操作。在至少一个实施例中，与一些深度学习方法相比，使用插值帧生成器合成高FPS视频提供了优势(例如，更好地处理被遮挡和被去遮挡区域)。在至少一个实施例中，使用插值帧生成器插值帧至少部分地基于以下假设：像素随时间线性运动，并沿由图形运动向量形成的线性路径对像素进行采样以在任意多(例如，预定数量，或参数控制的数量，例如通过API)个中间时间点创建完整帧。在至少一个实施例中，使用插值帧生成器插值帧至少部分地基于以下假设：像素以对应于图形运动向量的函数的方式运动，并且对使用图形运动向量识别的像素进行采样以创建完整帧(例如，沿图形运动向量采样像素，近似沿图形运动向量采样像素，例如通过基于图形运动向量和噪声值的采样，基于图形运动向量和预定加速度参数的采样，和/或一些其他合适的采样方法)。在至少一个实施例中，至少部分地基于两个连续帧的使用插值帧生成器来插值帧提供了优于不使用两个帧的方法的优点(例如，减少的伪影水平)。

数据中心

图15示出了根据至少一个实施例的示例数据中心1500。在至少一个实施例中，数据中心1500包括但不限于数据中心基础设施层1510、框架层 1520、软件层1530和应用层1540。

在至少一个实施例中，如图15所示，数据中心基础设施层1510可以包括资源协调器1512、分组的计算资源1514和节点计算资源(“节点C.R.”) 1516(1)-1516(N)，其中“N”代表任何完整的正整数。在至少一个实施例中，节点C.R.1516(1)-1516(N)可以包括但不限于任意数量的中央处理单元(“CPU”)或其他处理器(包括加速器、现场可编程门阵列(“FPGA”)、图形处理器等)，存储器设备(例如动态只读存储器)，存储设备(例如固态硬盘或磁盘驱动器)，网络输入/输出(“NW I/O”)设备，网络交换机，虚拟机(“VM”)，电源模块和冷却模块等。在至少一个实施例中，节点 C.R.1516(1)-1516(N)中的一个或更多个节点C.R.可以是具有一个或更多个上述计算资源的服务器。

在至少一个实施例中，分组的计算资源1514可以包括容纳在一个或更多个机架内的节点C.R.的单独分组(未示出)，或者容纳在各个地理位置的数据中心内的许多机架(也未示出)。分组的计算资源1514内的节点 C.R.的单独分组可以包括可以被配置或分配为支持一个或更多个工作负载的分组的计算、网络、内存或存储资源。在至少一个实施例中，可以将包括CPU或处理器的几个节点C.R.分组在一个或更多个机架内，以提供计算资源来支持一个或更多个工作负载。在至少一个实施例中，一个或更多个机架还可以包括任意数量的电源模块、冷却模块和网络交换机，以任意组合。

在至少一个实施例中，资源协调器1512可以配置或以其他方式控制一个或更多个节点C.R.1516(1)-1516(N)和/或分组的计算资源1514。在至少一个实施例中，资源协调器1512可以包括用于数据中心1500的软件设计基础结构(“SDI”)管理实体。在至少一个实施例中，资源协调器1512 可以包括硬件、软件或其某种组合。

在至少一个实施例中，如图15所示，框架层1520包括但不限于作业调度器1532、配置管理器1534、资源管理器1536和分布式文件系统1538。在至少一个实施例中，框架层1520可以包括支持软件层1530的软件1552 和/或应用程序层1540的一个或更多个应用程序1542的框架。在至少一个实施例中，软件1552或应用程序1542可以分别包括基于Web的服务软件或应用程序，例如由Amazon Web Services，Google Cloud和Microsoft Azure 提供的服务或应用程序。在至少一个实施例中，框架层1520可以是但不限于一种免费和开放源软件网络应用框架，例如可以利用分布式文件系统 1538来进行大范围数据处理(例如“大数据”)的Apache SparkTM(以下称为“Spark”)。在至少一个实施例中，作业调度器1532可以包括Spark驱动器，以促进对数据中心1500的各个层所支持的工作负载进行调度。在至少一个实施例中，配置管理器1534可以能够配置不同的层，例如软件层 1530和包括Spark和用于支持大规模数据处理的分布式文件系统1538的框架层1520。在至少一个实施例中，资源管理器1536能够管理映射到或分配用于支持分布式文件系统1538和作业调度器1532的集群或分组计算资源。在至少一个实施例中，集群或分组计算资源可以包括数据中心基础设施层 1510上的分组的计算资源1514。在至少一个实施例中，资源管理器1536 可以与资源协调器1512协调以管理这些映射的或分配的计算资源。

在至少一个实施例中，包括在软件层1530中的软件1552可以包括由节点C.R.1516(1)-1516(N)的至少一部分，分组计算资源1514和/或框架层1520的分布式文件系统1538使用的软件。一种或更多种类型的软件可以包括但不限于Internet网页搜索软件、电子邮件病毒扫描软件、数据库软件和流视频内容软件。

在至少一个实施例中，应用层1540中包括的一个或更多个应用程序 1542可以包括由节点C.R.1516(1)-1516(N)的至少一部分、分组的计算资源1514和/或框架层1520的分布式文件系统1538使用的一种或更多种类型的应用程序。一种或更多种类型的应用程序可以包括但不限于 CUDA应用程序。

在至少一个实施例中，配置管理器1534、资源管理器1536和资源协调器1512中的任何一个可以基于以任何技术上可行的方式获取的任意数量和类型的数据来实现任意数量和类型的自我修改动作。在至少一个实施例中，自我修改动作可以减轻数据中心1500的数据中心操作员做出可能不好的配置决定并且可以避免数据中心的未充分利用和/或执行差的部分。

基于计算机的系统

下列附图阐述了但不限于可用于实施至少一个实施例的示例性基于计算机的系统。在至少一个实施例中，下图的一个或更多个基于计算机的系统可以实现关于图1-14中的一个或更多个描述的实施例的一个或更多个方面，和/或关于图8-11描述的一种或更多个技术。

图16示出了根据至少一个实施例的处理系统1600。在至少一个实施例中，系统1600包括一个或更多个处理器1602和一个或更多个图形处理器1608，并且可以是单处理器台式机系统、多处理器工作站系统或具有大量处理器1602或处理器核心1607的服务器系统。在至少一个实施例中，处理系统1600是结合在片上系统(SoC)集成电路内的处理平台，以用于移动、手持或嵌入式设备。

在至少一个实施例中，处理系统1600可以包括或结合在基于服务器的游戏平台中，包括游戏和媒体控制台的游戏控制台、移动游戏控制台、手持游戏控制台或在线游戏控制台。在至少一个实施例中，处理系统1600 是移动电话、智能电话、平板计算设备或移动互联网设备。在至少一个实施例中，处理系统1600还可包括与可穿戴设备耦合或集成在可穿戴设备中，例如智能手表可穿戴设备、智能眼镜设备、增强现实设备或虚拟现实设备。在至少一个实施例中，处理系统1600是电视或机顶盒设备，其具有一个或更多个处理器1602以及由一个或更多个图形处理器1608生成的图形界面。

在至少一个实施例中，一个或更多个处理器1602每个包括一个或更多个处理器核心1607，以处理指令，该指令在被执行时执行针对系统和用户软件的操作。在至少一个实施例中，一个或更多个处理器核心1607中的每一个被配置为处理特定指令集1609。在至少一个实施例中，指令集1609 可以促进复杂指令集计算(CISC)、精简指令集计算(RISC)，或通过超长指令字(VLIW)进行计算。在至少一个实施例中，多个处理器核心1607 可以各自处理不同的指令集1609，该指令集1609可以包括有助于仿真其他指令集的指令。在至少一个实施例中，处理器核心1607还可以包括其他处理设备，例如数字信号处理器(DSP)。

在至少一个实施例中，处理器1602包括高速缓存存储器(cache)1604。在至少一个实施例中，处理器1602可以具有单个内部高速缓存或更多个级别的内部高速缓存。在至少一个实施例中，高速缓存存储器在处理器1602 的各个组件之间共享。在至少一个实施例中，处理器1602还使用外部高速缓存(例如，三级(L3)高速缓存或最后一级高速缓存(LLC))(未示出)，其可以使用已知的高速缓存一致性技术在处理器核心1607之间共享该逻辑。在至少一个实施例中，处理器1602中另外包括寄存器文件1606，处理器1602可以包括用于存储不同类型的数据的不同类型的寄存器(例如，整数寄存器、浮点寄存器、状态寄存器和指令指针寄存器)。在至少一个实施例中，寄存器文件1606可以包括通用寄存器或其他寄存器。

在至少一个实施例中，一个或更多个处理器1602与一个或更多个接口总线1610耦合，以在处理器1602与系统1600中的其他组件之间传输通信信号，例如地址、数据或控制信号。在至少一个实施例中，接口总线1610 在一个实施例中可以是处理器总线，例如直接媒体接口(DMI)总线的版本。在至少一个实施例中，接口总线1610不限于DMI总线，并且可以包括一个或更多个外围组件互连总线(例如，PCI，PCI Express)、存储器总线或其他类型的接口总线。在至少一个实施例中，处理器1602包括集成存储器控制器1610和平台控制器集线器1630。在至少一个实施例中，存储器控制器1610促进存储设备与处理系统1600的其他组件之间的通信，而平台控制器集线器(PCH)1630通过本地I/O总线提供到输入/输出(I/O)设备的连接。

在至少一个实施例中，存储设备1620可以是动态随机存取存储器 (DRAM)设备、静态随机存取存储器(SRAM)设备、闪存设备、相变存储设备或具有适当的性能以用作处理器存储器。在至少一个实施例中，存储设备1620可以用作处理系统1600的系统存储器，以存储数据1622和指令1621，以在一个或更多个处理器1602执行应用或过程时使用。在至少一个实施例中，存储器控制器1610还与可选的外部图形处理器1612耦合，其可以与处理器1602中的一个或更多个图形处理器1608通信以执行图和媒体操作。在至少一个实施例中，显示设备1611可以连接至处理器1602。在至少一个实施例中，显示设备1611可以包括内部显示设备中的一个或更多个，例如在移动电子设备或便携式计算机设备或通过显示器接口(例如显示端口(DisplayPort)等)连接的外部显示设备。在至少一个实施例中，显示设备1611可以包括头戴式显示器(HMD)，诸如用于虚拟现实(VR) 应用或增强现实(AR)应用中的立体显示设备。

在至少一个实施例中，平台控制器集线器1630使外围设备能够通过高速I/O总线连接到存储设备1620和处理器1602。在至少一个实施例中， I/O外围设备包括但不限于音频控制器1646、网络控制器1634、固件接口 1628、无线收发器1626、触摸传感器1625、数据存储设备1624(例如，硬盘驱动器、闪存等)。在至少一个实施例中，数据存储设备1624可以经由存储器接口(例如，SATA)或经由外围总线来连接，诸如外围组件互连总线(例如，PCI、PCIe)。在至少一个实施例中，触摸传感器1625可以包括触摸屏传感器、压力传感器或指纹传感器。在至少一个实施例中，无线收发器1626可以是Wi-Fi收发器、蓝牙收发器或移动网络收发器，诸如3G、 4G或长期演进(LTE)收发器。在至少一个实施例中，固件接口1628使能与系统固件的通信，并且可以是例如统一的可扩展固件接口(UEFI)。在至少一个实施例中，网络控制器1634可以启用到有线网络的网络连接。在至少一个实施例中，高性能网络控制器(未示出)与接口总线1610耦合。在至少一个实施例中，音频控制器1646是多通道高清晰度音频控制器。在至少一个实施例中，处理系统1600包括可选的传统(legacy)I/O控制器1640，用于将遗留(例如，个人系统2(PS/2))设备耦合到处理系统1600。在至少一个实施例中，平台控制器集线器1630还可以连接到一个或更多个通用串行总线(USB)控制器1642，该控制器连接输入设备，诸如键盘和鼠标1643组合、相机1644或其他USB输入设备。

在至少一个实施例中，存储器控制器1610和平台控制器集线器1630 的实例可以集成到离散的外部图形处理器中，例如外部图形处理器1612。在至少一个实施例中，平台控制器集线器1630和/或存储控制器1610可以在一个或更多个处理器1602的外部。例如，在至少一个实施例中，处理系统1600可以包括外部存储控制器1610和平台控制器集线器1630，其可以配置成在与处理器1602通信的系统芯片组中的存储器控制器集线器和外围控制器集线器。

图17示出了根据至少一个实施例的计算机系统1700。在至少一个实施例中，计算机系统1700可以是具有互连的设备和组件，SOC，或某种组合的系统。在至少一个实施例中，计算机系统1700由处理器1702形成，该处理器1702可以包括用于执行指令的执行单元。在至少一个实施例中，计算机系统1700可以包括但不限于组件，例如处理器1702，其采用包括逻辑的执行单元以执行用于过程数据的算法。在至少一个实施例中，计算机系统1700可以包括处理器，例如可从加利福尼亚圣塔克拉拉的英特尔公司 (Intel Corporation ofSanta Clara,California)获得的

处理器家族、XeonTM、

XScaleTM和/或StrongARMTM，

Core^TM或

Nervana^TM微处理器，尽管也可以使用其他系统(包括具有其他微处理器的PC、工程工作站、机顶盒等)。在至少一个实施例中，计算机系统1700可以执行可从华盛顿州雷蒙德市的微软公司(Microsoft Corporation ofRedmond,Wash.)获得的WINDOWS操作系统版本，尽管其他操作系统 (例如UNIX和Linux)、嵌入式软件和/或图形用户界面也可以使用。

在至少一个实施例中，计算机系统1700可以用在其他设备中，例如手持设备和嵌入式应用。手持设备的一些示例包括蜂窝电话、互联网协议 (Internet Protocol)设备、数码相机、个人数字助理(“PDA”)和手持PC。在至少一个实施例中，嵌入式应用可以包括微控制器、数字信号处理器 (“DSP”)、SoC、网络计算机(“NetPC”)、机顶盒、网络集线器、广域网(“WAN”)交换机，或根据至少一个实施例可以执行一个或更多个指令的任何其他系统。

在至少一个实施例中，计算机系统1700可包括但不限于处理器1702，该处理器1702可包括但不限于一个或更多个执行单元1708，其可以配置为执行计算统一设备架构(“CUDA”)(

由加利福尼亚州圣克拉拉的 NVIDIACorporation开发)程序。在至少一个实施例中，CUDA程序是用 CUDA编程语言编写的软件应用程序的至少一部分。在至少一个实施例中，计算机系统1700是单处理器台式机或服务器系统。在至少一个实施例中，计算机系统1700可以是多处理器系统。在至少一个实施例中，处理器1702 可以包括但不限于CISC微处理器、RISC微处理器、VLIW微处理器、实现指令集组合的处理器，或任何其他处理器设备，例如数字信号处理器。在至少一个实施例中，处理器1702可以耦合到处理器总线1710，该处理器总线1710可以在处理器1702与计算机系统1700中的其他组件之间传输数据信号。

在至少一个实施例中，处理器1702可以包括但不限于1级(“L1”) 内部高速缓存存储器(“cache”)1704。在至少一个实施例中，处理器1702 可以具有单个内部高速缓存或多级内部缓存。在至少一个实施例中，高速缓存存储器可以驻留在处理器1702的外部。在至少一个实施例中，处理器 1702可以包括内部和外部高速缓存的组合。在至少一个实施例中，寄存器文件1706可以在各种寄存器中存储不同类型的数据，包括但不限于整数寄存器、浮点寄存器、状态寄存器和指令指针寄存器。

在至少一个实施例中，包括但不限于执行整数和浮点运算的逻辑的执行单元1708，其也位于处理器1702中。处理器1702还可以包括微码 (“ucode”)只读存储器(“ROM”)，用于存储某些宏指令的微代码。在至少一个实施例中，执行单元1708可以包括用于处理封装指令集1709的逻辑。在至少一个实施例中，通过将封装指令集1709包括在通用处理器1702 的指令集中，以及要执行指令的相关电路，可以使用通用处理器1702中的封装数据来执行许多多媒体应用程序使用的操作。在至少一个实施例中，可以通过使用处理器的数据总线的全宽度来在封装的数据上执行操作来加速和更有效地执行许多多媒体应用程序，这可能不需要在处理器的数据总线上传输较小的数据单元来一次对一个数据元素执行一个或更多个操作。

在至少一个实施例中，执行单元1708也可以用在微控制器、嵌入式处理器、图形设备、DSP和其他类型的逻辑电路中。在至少一个实施例中，计算机系统1700可以包括但不限于存储器1720。在至少一个实施例中，存储器1720可以被实现为DRAM设备、SRAM设备、闪存设备或其他存储设备。存储器1720可以存储由处理器1702可以执行的由数据信号表示的指令1719和/或数据1721。

在至少一个实施例中，系统逻辑芯片可以耦合到处理器总线1710和存储器1720。在至少一个实施例中，系统逻辑芯片可以包括但不限于存储器控制器集线器(“MCH”)1716，并且处理器1702可以经由处理器总线 1710与MCH 1716通信。在至少一个实施例中，MCH1716可以提供到存储器1720的高带宽存储器路径1718以用于指令和数据存储以及用于图形命令、数据和纹理的存储。在至少一个实施例中，MCH 1716可以在处理器1702、存储器1720和计算机系统1700中的其他组件之间启动数据信号，并且在处理器总线1710、存储器1720和系统I/O 1722之间桥接数据信号。在至少一个实施例中，系统逻辑芯片可以提供用于耦合到图形控制器的图形端口。在至少一个实施例中，MCH 1716可以通过高带宽存储器路径1718 耦合到存储器1720，并且图形/视频卡1712可以通过加速图形端口(Accelerated Graphics Port)(“AGP”)互连1714耦合到MCH 1716。

在至少一个实施例中，计算机系统1700可以使用系统I/O 1722作为专有集线器接口总线来将MCH 1716耦合到I/O控制器集线器(“ICH”)1730。在至少一个实施例中，ICH1730可以通过本地I/O总线提供与某些I/O设备的直接连接。在至少一个实施例中，本地I/O总线可以包括但不限于用于将外围设备连接到存储器1720、芯片组和处理器1702的高速I/O总线。示例可以包括但不限于音频控制器1729、固件集线器(“Flash BIOS”)1728、无线收发器1726、数据存储1724、包含用户输入1725的传统I/O控制器 1723和键盘接口、串行扩展端口1727(例如USB)和网络控制器1734。数据存储1724可以包括硬盘驱动器、软盘驱动器、CD-ROM设备、闪存设备或其他大容量存储设备。

在至少一个实施例中，图17示出了包括互连的硬件设备或“芯片”的系统。在至少一个实施例中，图17可以示出示例性SoC。在至少一个实施例中，图17中示出的设备可以与专有互连、标准化互连(例如，PCIe)或其某种组合互连。在至少一个实施例中，系统1700的一个或更多个组件使用计算快速链路(CXL)互连来互连。

图18示出了根据至少一个实施例的系统1800。在至少一个实施例中，系统1800是利用处理器1810的电子设备。在至少一个实施例中，系统1800 可以是，例如但不限于，笔记本电脑、塔式服务器、机架服务器、刀片服务器、膝上型计算机、台式机、平板电脑、移动设备、电话、嵌入式计算机或任何其他合适的电子设备。

在至少一个实施例中，系统1800可以包括但不限于通信地耦合到任何合适数量或种类的组件、外围设备、模块或设备的处理器1810。在至少一个实施例中，处理器1810使用总线或接口耦合，诸如I²C总线、系统管理总线(“SMBus”)、低引脚数(LPC)总线、串行外围接口(“SPI”)、高清音频(“HDA”)总线、串行高级技术附件(“SATA”)总线、USB(1、2、 3版)或通用异步接收器/发送器(“UART”)总线。在至少一个实施例中，图18示出了系统，该系统包括互连的硬件设备或“芯片”。在至少一个实施例中，图18可以示出示例性SoC。在至少一个实施例中，图18中所示的设备可以与专有互连线、标准化互连(例如，PCIe)或其某种组合互连。在至少一个实施例中，图18的一个或更多个组件使用计算快速链路(CXL) 互连线来互连。

在至少一个实施例中，图18可以包括显示器1824、触摸屏1825、触摸板1830、近场通信单元(“NFC”)1845、传感器集线器1840、热传感器 1846、快速芯片组(“EC”)1835、可信平台模块(“TPM”)1838、BIOS/ 固件/闪存(“BIOS，FW Flash”)1822、DSP 1860、固态磁盘(“SSD”)或硬盘驱动器(“HDD”)1820、无线局域网单元(“WLAN”)1850、蓝牙单元1852、无线广域网单元(“WWAN”)1856、全球定位系统(GPS)1855、相机(“USB 3.0相机”)1854(例如USB3.0相机)或以例如LPDDR3标准实现的低功耗双倍数据速率(“LPDDR”)存储器单元(“LPDDR3”)1815。这些组件可以各自以任何合适的方式实现。

在至少一个实施例中，其他组件可以通过以上讨论的组件通信地耦合到处理器1810。在至少一个实施例中，加速度计1841、环境光传感器(“ALS”) 1842、罗盘1843和陀螺仪1844可以可通信地耦合到传感器集线器1840。在至少一个实施例中，热传感器1839、风扇1837、键盘1846和触摸板1830 可以通信地耦合到EC 1835。在至少一个实施例中，扬声器1863、耳机1864 和麦克风(“mic”)1865可以通信地耦合到音频单元(“音频编解码器和D 类放大器”)1864，其又可以通信地耦合到DSP 1860。在至少一个实施例中，音频单元1864可以包括例如但不限于音频编码器/解码器(“编解码器”)和 D类放大器。在至少一个实施例中，SIM卡(“SIM”)1857可以通信地耦合到WWAN单元1856。在至少一个实施例中，组件(诸如WLAN单元1850 和蓝牙单元1852以及WWAN单元1856)可以被实现为下一代形式因素 (NGFF)。

图19示出了根据至少一个实施例的示例性集成电路1900。在至少一个实施例中，示例性集成电路1900是SoC，其可使用一个或更多个IP核心制造。在至少一个实施例中，集成电路1900包括一个或更多个应用处理器1905(例如，CPU)、至少一个图形处理器1910，并且可以另外包括图像处理器1915和/或视频处理器1920，其中任意一个可能是模块化IP核心。在至少一个实施例中，集成电路1900包括外围或总线逻辑，其包括USB 控制器1925、UART控制器1930、SPI/SDIO控制器1935和I²S/I²C控制器1940。在至少一个实施例中，集成电路1900可以包括显示设备1945耦合到高清多媒体接口(HDMI)控制器1950和移动工业处理器接口(MIPI) 显示接口1955中的一个或更多个。在至少一个实施例中，存储可以由闪存子系统1960提供，包括闪存和闪存控制器。在至少一个实施例中，可以经由存储器控制器1965提供存储器接口以用于访问SDRAM或SRAM存储器设备。在至少一个实施例中，一些集成电路还包括嵌入式安全引擎1970。

图20示出了根据至少一个实施例的计算系统2000。在至少一个实施例中，计算系统2000包括处理子系统2001，其具有经由可以包括存储器集线器2005的互连路径通信的一个或更多个处理器2002和系统存储器2004。在至少一个实施例中，存储器集线器2005可以是芯片组组件内的单独组件，也可以集成在一个或更多个处理器2002内。在至少一个实施例中，存储器集线器2005通过通信链路2006与I/O子系统2011耦合。在至少一个实施例中，I/O子系统2011包括I/O集线器2007，其可以使计算系统2000能够接收来自一个或更多个输入设备2008的输入。在至少一个实施例中，I/O 集线器2007可以使能显示控制器，其包括在一个或更多个处理器2002中，用于向一个或更多个显示设备2010A提供输出。在至少一个实施例中，与 I/O集线器2007耦合的一个或更多个显示设备2010A可以包括本地、内部或嵌入式显示设备。

在至少一个实施例中，处理子系统2001包括经由总线或其他通信链路2013耦合到存储器集线器2005的一个或更多个并行处理器2012。在至少一个实施例中，通信链路2013可以是许多基于标准的通信链路技术或协议中的一种，例如但不限于PCIe，或者可以是针对供应商的通信接口或通信结构。在至少一个实施例中，一个或更多个并行处理器2012形成计算集中的并行或向量处理系统，该系统可以包括大量的处理核心和/或处理集群，例如多集成核心(MIC)处理器。在至少一个实施例中，一个或更多个并行处理器2012形成可以将像素输出到经由I/O集线器2007耦合的一个或更多个显示设备2010A之一的图形处理子系统。在至少一个实施例中，一个或更多个并行处理器2012还可以包括显示控制器和显示接口(未示出)，以使得能够直接连接到一个或更多个显示设备2010B。

在至少一个实施例中，系统存储单元2014可以连接到I/O集线器2007，以提供用于计算系统2000的存储机制。在至少一个实施例中，I/O交换机 2016可以用于提供接口机制，以实现I/O集线器2007与其他组件之间的连接，例如可以集成到平台中的网络适配器2018和/或无线网络适配器2019，以及可以通过一个或更多个附加设备2020添加的各种其他设备。在至少一个实施例中，网络适配器2018可以是以太网适配器或另一有线网络适配器。在至少一个实施例中，无线网络适配器2019可以包括Wi-Fi、蓝牙、NFC 的一个或更多个或其他包括一个或更多个无线电的网络设备。

在至少一个实施例中，计算系统2000可以包括未明确示出的其他组件，包括USB或其他端口连接、光存储驱动器、视频捕获设备等，也可以连接到I/O集线器2007。在至少一个实施例中，对图20中的各个组件进行互连的通信路径可以使用任何合适的协议来实现，诸如基于PCI(外围组件互连)的协议(例如，PCIe)，或其他总线或点对点通信接口和/或协议(例如，NVLink高速互连或互连协议)。

在至少一个实施例中，一个或更多个并行处理器2012包括针对图形和视频处理而优化的电路(包括例如视频输出电路)，并构成图形处理单元(GPU)。在至少一个实施例中，一个或更多个并行处理器2012包括针对通用处理而优化的电路。在至少一个实施例中，计算系统2000的组件可以与单个集成电路上的一个或更多个其他系统元件集成。例如，在至少一个实施例中，一个或更多个并行处理器2012、存储器集线器2005、处理器 2002和I/O集线器2007可以被集成到片上系统(SoC)集成电路中。在至少一个实施例中，计算系统2000的组件可以被集成到单个封装中以形成系统级封装(SIP)配置。在至少一个实施例中，计算系统2000的组件的至少一部分可以被集成到多芯片模块(MCM)中，该多芯片模块可以与其他多芯片模块互连到模块化计算系统中。在至少一个实施例中，从计算系统 2000中省略了I/O子系统2011和显示设备2010B。

处理系统

下列附图阐述了但不限于可用于实施至少一个实施例的示例性处理系统。在至少一个实施例中，下图的一个或更多个处理系统可以实现关于图1-14中的一个或更多个描述的实施例的一个或更多个方面，和/或关于图 8-11描述的一个或更多个技术。

图21示出了根据至少一个实施例的加速处理单元(“APU”)2100。在至少一个实施例中，APU 2100由加利福尼亚州圣克拉拉市的AMD公司开发。在至少一个实施例中，APU2100可以被配置为执行应用程序，诸如 CUDA程序。在至少一个实施例中，APU 2100包括但不限于核心复合体2110、图形复合体2140、结构2160、I/O接口2170、存储器控制器2180、显示控制器2192和多媒体引擎2194。在至少一个实施例中，APU 2100可以包括但不限于任意数量的核心复合体2110、任意数量的图形复合体2140、任意数量的显示控制器2192和任意数量的多媒体引擎2194的任何组合。为了说明的目的，在本文中用附图标记表示相似对象的多个实例，其中附图标记标识该对象，并且括号中的数字标识所需要的实例。

在至少一个实施例中，核心复合体2110是CPU，图形复合体2140 是GPU，并且APU2100是将不限于2110和2140集成到单个芯片上的处理单元。在至少一个实施例中，一些任务可以被分配给核心复合体2110，而其他任务可以被分配给图形复合体2140。在至少一个实施例中，核心复合体2110被配置为执行与APU 2100相关联的主控制软件，例如操作系统。在至少一个实施例中，核心复合体2110是APU 2100的主处理器，其控制和协调其他处理器的操作。在至少一个实施例中，核心复合体2110发出控制图形复合体2140的操作的命令。在至少一个实施例中，核心复合体2110 可以被配置为执行从CUDA源代码派生的主机可执行代码，并且图形复合体2140可以被配置为执行从CUDA源代码派生的设备可执行代码。

在至少一个实施例中，核心复合体2110包括但不限于核心2120(1) -2120(4)和L3高速缓存2130。在至少一个实施例中，核心复合体2110 可以包括但不限于任意数量的核心2120以及任意数量和类型的高速缓存的任何组合。在至少一个实施例中，核心2120被配置为执行特定指令集架构 (“ISA”)的指令。在至少一个实施例中，每个核心2120是CPU核心。

在至少一个实施例中，每个核心2120包括但不限于获取/解码单元 2122，整数执行引擎2124，浮点执行引擎2126和L2高速缓存2128。在至少一个实施例中，获取/解码单元2122获取指令，对这些指令进行解码，生成微操作，并将单独的微指令分派给整数执行引擎2124和浮点执行引擎 2126。在至少一个实施例中，获取/解码单元2122可以同时分派一个微指令到整数执行引擎2124和另一微指令到浮点执行引擎2126。在至少一个实施例中，整数执行引擎2124执行不限于整数和存储器操作。在至少一个实施例中，浮点引擎2126执行不限于浮点和向量运算。在至少一个实施例中，获取-解码单元2122将微指令分派给单个执行引擎，该执行引擎代替整数执行引擎2124和浮点执行引擎2126两者。

在至少一个实施例中，每个核心2120(i)可以访问包括在核心2120 (i)中的L2高速缓存2128(i)，其中i是表示核心2120的特定实例的整数。在至少一个实施例中，包括在核心复合体2110(j)中的每个核心2120 经由包括在核心复合体2110(j)中的L3高速缓存2130(j)连接到包括在核心复合体2110(j)中的其他核心2120，其中j是表示核心复合体2110 的特定实例的整数。在至少一个实施例中，包括在核心复合体2110(j)中的核心2120可以访问包括在核心复合体2110(j)中的所有L3高速缓存 2130(j)，其中j是表示核心复合体2110的特定实例的整数。在至少一个实施例中，L3高速缓存2130可以包括但不限于任意数量的切片(slice)。

在至少一个实施例中，图形复合体2140可以被配置为以高度并行的方式执行计算操作。在至少一个实施例中，图形复合体2140被配置为执行图形管线操作，诸如绘制命令、像素操作、几何计算以及与将图像渲染至显示器相关联的其他操作。在至少一个实施例中，图形复合体2140被配置为执行与图形无关的操作。在至少一个实施例中，图形复合体2140被配置为执行与图形有关的操作和与图形无关的操作。

在至少一个实施例中，图形复合体2140包括但不限于任意数量的计算单元2150和L2高速缓存2142。在至少一个实施例中，计算单元2150 共享L2高速缓存2142。在至少一个实施例中，L2高速缓存2142被分区。在至少一个实施例中，图形复合体2140包括但不限于任意数量的计算单元 2150以及任意数量(包括零)和类型的高速缓存。在至少一个实施例中，图形复合体2140包括但不限于任意数量的专用图形硬件。

在至少一个实施例中，每个计算单元2150包括但不限于任意数量的 SIMD单元2152和共享存储器2154。在至少一个实施例中，每个SIMD单元2152实现SIMD架构并且被配置为并行执行操作。在至少一个实施例中，每个计算单元2150可以执行任意数量的线程块，但是每个线程块在单个计算单元2150上执行。在至少一个实施例中，线程块包括但不限于任意数量的执行线程。在至少一个实施例中，工作组是线程块。在至少一个实施例中，每个SIMD单元2152执行不同的线程束(warp)。在至少一个实施例中，线程束是一组线程(例如16个线程)，其中线程束中的每个线程属于单个线程块，并且被配置为基于单个指令集来处理不同的数据集。在至少一个实施例中，可以使用预测(predication)来禁用线程束中的一个或更多个线程。在至少一个实施例中，通道是线程。在至少一个实施例中，工作项是线程。在至少一个实施例中，波前是线程束。在至少一个实施例中，线程块中的不同波前可一起同步并经由共享存储器2154进行通信。

在至少一个实施例中，结构2160是系统互连，其促进跨核心复合体 2110、图形复合体2140、I/O接口2170、存储器控制器2180、显示控制器 2192和多媒体引擎2194的数据和控制传输。在至少一个实施例中，除了结构2160之外或代替结构2160，APU 2100还可以包括但不限于任意数量和类型的系统互连，该结构2160促进跨可以在APU 2100内部或外部的任意数量和类型的直接或间接链接的组件的数据和控制传输。在至少一个实施例中，I/O接口2170表示任意数量和类型的I/O接口(例如，PCI， PCI-Extended(“PCI-X”)，PCIe，千兆以太网(“GBE”)，USB等)。在至少一个实施例中，各种类型的外围设备耦合到I/O接口2170。在至少一个实施例中，耦合到I/O接口2170的外围设备可以包括但不限于键盘，鼠标，打印机，扫描仪，操纵杆或其他类型的游戏控制器、媒体记录设备、外部存储设备、网络接口卡等。

在至少一个实施例中，显示控制器AMD92在一个或更多个显示设备 (例如液晶显示器(LCD)设备)上显示图像。在至少一个实施例中，多媒体引擎240包括但不限于任意数量和类型的与多媒体相关的电路，例如视频解码器、视频编码器、图像信号处理器等。在至少一个实施例中，存储器控制器2180促进APU 2100与统一系统存储器2190之间的数据传输。在至少一个实施例中，核心复合体2110和图形复合体2140共享统一系统存储器2190。

在至少一个实施例中，APU 2100实现种存储器子系统，其包括但不限于任意数量和类型的存储器控制器2180和可以专用于一个组件或在多个组件之间共享的存储器设备(例如，共享存储器2154)。组件。在至少一个实施例中，APU 2100实现高速缓存子系统，其包括但不限于一个或更多个高速缓存存储器(例如，L2高速缓存2228，L3高速缓存2130和L2高速缓存2142)，每个高速缓存存储器可以是组件私有的或在任意数量的组件(例如，核心2120，核心复合体2110，SIMD单元2152，计算单元2150 和图形复合体2140)之间共享。

图22示出了根据至少一个实施例的CPU 2200。在至少一个实施例中， CPU 2200由加利福尼亚州圣克拉拉市的AMD公司开发。在至少一个实施例中，CPU 2200可以被配置为执行应用程序。在至少一个实施例中，CPU 2200被配置为执行主控制软件，例如操作系统。在至少一个实施例中，CPU 2200发出控制外部GPU(未示出)的操作的命令。在至少一个实施例中， CPU 2200可以被配置为执行从CUDA源代码派生的主机可执行代码，并且外部GPU可以被配置为执行从这种CUDA源代码派生的设备可执行代码。在至少一个实施例中，CPU 2200包括但不限于任意数量的核心复合体 2210，结构2260，I/O接口2270和存储器控制器2280。

在至少一个实施例中，核心复合体2210包括但不限于核心2220(1) -2220(4)和L3高速缓存2230。在至少一个实施例中，核心复合体2210 可以包括但不限于任意数量的核心2220以及任意数量和类型的高速缓存的任何组合。在至少一个实施例中，核心2220被配置为执行特定ISA的指令。在至少一个实施例中，每个核心2220是CPU核心。

在至少一个实施例中，每个核心2220包括但不限于获取/解码单元 2222，整数执行引擎2224，浮点执行引擎2226和L2高速缓存2228。在至少一个实施例中，获取/解码单元2222获取指令，对这些指令进行解码，生成微操作，并将单独的微指令分派给整数执行引擎2224和浮点执行引擎 2226。在至少一个实施例中，获取/解码单元2222可以同时分派一个微指令至整数执行引擎2224和另一微指令至浮点执行引擎2226。在至少一个实施例中，整数执行引擎2224执行不限于整数和存储器操作。在至少一个实施例中，浮点引擎2226执行不限于浮点和向量运算。在至少一个实施例中，获取-解码单元2222将微指令分派给单个执行引擎，该引擎代替整数执行引擎2224和浮点执行引擎2226两者。

在至少一个实施例中，每个核心2220(i)可以访问包括在核心2220 (i)中的L2高速缓存2228(i)，其中i是表示核心2220的特定实例的整数。在至少一个实施例中，包括在核心复合体2210(j)中的每个核心2220 经由包括在核心复合体2210(j)中的L3高速缓存2230(j)连接到核心复合体2210(j)中的其他核心2220，其中j是表示核心复合体2210的特定实例的整数。在至少一个实施例中，包括在核心复合体2210(j)中的核心 2220可以访问包括在核心复合体2210(j)中的所有L3高速缓存2230(j)，其中j是表示核心复合体2210的特定实例的整数。在至少一个实施例中， L3高速缓存2230可以包括但不限于任意数量的切片。

在至少一个实施例中，结构2260是系统互连，其促进跨核心复合体 2210(1)-2210(N)(其中N是大于零的整数)、I/O接口2270和存储器控制器2280的数据和控制传输。在至少一个实施例中，除了结构2260 之外或代替结构2260，CPU 2200还可以包括但不限于任意数量和类型的系统互连，该结构2260促进跨可以在CPU 2200内部或外部的任意数量和类型的直接或间接链接的组件的数据和控制传输。在至少一个实施例中，I/O 接口2270表示任意数量和类型的I/O接口(例如PCI，PCI-X，PCIe，GBE， USB等)。在至少一个实施例中，各种类型的外围设备耦合到I/O接口2270。在至少一个实施例中，耦合到I/O接口2270的外围设备可以包括但不限于显示器，键盘，鼠标，打印机，扫描仪，操纵杆或其他类型的游戏控制器、媒体记录设备、外部存储设备、网络接口卡等。

在至少一个实施例中，存储器控制器2280促进CPU 2200与系统存储器2290之间的数据传输。在至少一个实施例中，核心复合体2210和图形复合体2240共享系统存储器2290。在至少一个实施例中，CPU 2200实现存储器子系统，其包括但不限于任意数量和类型的存储器控制器2280和可以专用于一个组件或在多个组件之间共享的存储器设备。在至少一个实施例中，CPU 2200实现了高速缓存子系统，其包括但不限于一个或更多个高速缓存存储器(例如，L2高速缓存2228和L3高速缓存2230)，每个高速缓存存储器可以是组件私有的或在任意数量的组件(例如，核心2220和核心复合体2210)之间共享。

图23示出了根据至少一个实施例的示例性加速器集成切片2390。如本文所使用的，“切片”包括加速器集成电路的处理资源的指定部分。在至少一个实施例中，加速器集成电路代表多个图形加速模块种的多个图形处理引擎提供高速缓存管理、存储器访问、环境管理和中断管理服务。图形处理引擎可以各自包括单独的GPU。可选地，图形处理引擎可包括GPU 内的不同类型的图形处理引擎，例如图形执行单元、媒体处理引擎(例如，视频编码器/解码器)、采样器和blit引擎。在至少一个实施例中，图形加速模块可以是具有多个图形处理引擎的GPU。在至少一个实施例中，图形处理引擎可以是集成在通用封装、线卡或芯片上的各个GPU。

系统存储器2314内的应用程序有效地址空间2382存储进程元素2383。在一个实施例中，响应于来自处理器2307上执行的应用程序2380的GPU 调用2381而存储进程元素2383。进程元素2383包含对应应用程序2380 的处理状态。包含在进程元素2383中的工作描述符(WD)2384可以是应用程序请求的单个作业或可能包含指向作业队列的指针。在至少一个实施例中，WD 2384是指向应用程序有效地址空间2382中的作业请求队列的指针。

图形加速模块2346和/或各个图形处理引擎可以由系统中的全部或部分进程共享。在至少一个实施例中，可以包括用于建立处理状态并将WD 2384发送到图形加速模块2346以在虚拟化环境中开始作业的基础设施。

在至少一个实施例中，专用进程编程模型是针对实现的。在该模型中，单个进程拥有图形加速模块2346或个体图形处理引擎。由于图形加速模块 2346由单个进程拥有，因此管理程序为拥有的分区初始化加速器集成电路，并且当分配图形加速模块2346时操作系统对加速器集成电路进行初始化以用于拥有的分区。

在操作中，加速器集成切片2390中的WD获取单元2391获取下一个 WD 2384，其中包括要由图形加速模块2346的一个或更多个图形处理引擎完成的工作的指示。来自WD2384的数据可以存储在寄存器2345被存储器管理单元(MMU)2339、中断管理电路2347和/或环境管理电路2348 使用，如图所示。例如，MMU 2339的一个实施例包括用于访问OS虚拟地址空间2385内的段/页表2386的段/页面漫游电路。中断管理电路2347 可以处理从图形加速模块2346接收到的中断事件(INT)2392。当执行图操作时，由图形处理引擎产生的有效地址2393由MMU 2339转换为实际地址。

在一个实施例中，为每个图形处理引擎和/或图形加速模块2346复制相同的寄存器组2345，并且可以由系统管理程序或操作系统来初始化。这些复制的寄存器中的每一个都可以包含在加速器集成切片2390中。表1中显示了可由管理程序初始化的示例性寄存器。

表1–管理程序初始化的寄存器

表2中示出了可以由操作系统初始化的示例性寄存器。

表2–操作系统初始化寄存器

1	进程和线程识别
		2	有效地址(EA)环境保存/还原指针
3	虚拟地址(VA)加速器利用率记录指针
		4	虚拟地址(VA)存储分段表指针
5	权限屏蔽
		6	工作描述符

在一个实施例中，每个WD 2384特定于特定的图形加速模块2346和 /或特定图形处理引擎。它包含图形处理引擎进行工作或工作所需的所有信息，或者它可以是指向存储器位置的指针，其中应用程序建立了要完成的工作的命令队列。

图24A-24B示出了根据本文至少一个实施例的示例性图形处理器。在至少一个实施例中，任何示例性图形处理器可以使用一个或更多个IP核心来制造。除了图示之外，在至少一个实施例中可以包括其他逻辑和电路，包括附加的图形处理器/核心、外围接口控制器或通用处理器核心。在至少一个实施例中，示例性图形处理器用于SoC内。

图24A示出了根据至少一个实施例的SoC集成电路的示例性图形处理器2410，其可以使用一个或更多个IP核心来制造。图24B示出了根据至少一个实施例的SoC集成电路的的附加示例性图形处理器2440，其可以使用一个或更多个IP核心来制造。在至少一个实施例中，图24A的图形处理器2410是低功耗图形处理器核心。在至少一个实施例中，图24B的图形处理器2440是更高性能的图形处理器核心。在至少一个实施例中，每个图形处理器2410、2440可以是图19的图形处理器1910的变体。

在至少一个实施例中，图形处理器2410包括顶点处理器2405和一个或更多个片段处理器2415A-2415N(例如2415A、2415B、2415C、2415D 至2415N-1和2415N)。在至少一个实施例中，图形处理器2410可以经由单独的逻辑来执行不同的着色器程序，使得顶点处理器2405被优化以执行针对顶点着色器程序的操作，而一个或更多个片段处理器2415A-2415N执行片段(例如，像素)着色操作用于片段或像素或着色器程序。在至少一个实施例中，顶点处理器2405执行3D图形管线的顶点处理阶段并生成图元和顶点数据。在至少一个实施例中，片段处理器2415A-2415N使用由顶点处理器2405生成的图元和顶点数据来生成在显示设备上显示的帧缓冲区。在至少一个实施例中，片段处理器2415A-2415N被优化以执行如在OpenGL API中所提供的片段着色器程序，其可以用于执行与在Direct 3D API中所提供的像素着色器程序类似的操作。

在至少一个实施例中，图形处理器2410附加地包括一个或更多个 MMU 2420A-2420B、高速缓存2425A-2425B和电路互连2430A-2430B。在至少一个实施例中，一个或更多个MMU 2420A-2420B提供用于图形处理器2410的虚拟到物理地址的映射，包括用于顶点处理器2405和/或片段处理器2415A-2415N，其可以引用存储在存储器中的顶点或图像/纹理数据，除了存储在一个或更多个高速缓存2425A-2425B中的顶点或图像/纹理数据之外。在至少一个实施例中，一个或更多个MMU 2420A-2420B可以与系统内的其他MMU同步，包括与图19的一个或更多个应用处理器1905、图像处理器1915和/或视频处理器1920相关联的一个或更多个MMU，使得每个处理器1905-1920可以参与共享或统一的虚拟存储器系统。在至少一个实施例中，一个或更多个电路互连2430A-2430B使图形处理器2410能够经由SoC的内部总线或经由直接连接与SoC内的其他IP核心相连接。

在至少一个实施例中，图形处理器2440包括图24A的图形处理器 2410的一个或更多个MMU 2420A-2420B、高速缓存2425A-2425B和电路互连2430A-2430B。在至少一个实施例中，图形处理器2440包括一个或更多个着色器核心2455A-2455N(例如，2455A、2455B、2455C、2455D、 2455E、2455F、至2455N-1和2455N)，其提供了统一的着色器核心架构，其中单个核心或类型或核心可以执行所有类型的可编程着色器代码，包括用于实现顶点着色器、片段着色器和/或计算着色器的着色器程序代码。在至少一个实施例中，多个着色器核心可以变化。在至少一个实施例中，图形处理器2440包括核心间任务管理器2445，其充当线程分派器以将执行线程分派给一个或更多个着色器核心2455A-2455N和分块单元2458，以加速基于图块渲染的分块操作，其中在图像空间中细分了场景的渲染操作，例如，以利用场景内的局部空间一致性或优化内部缓存的使用。

图25A示出了根据至少一个实施例的图形核心2500。在至少一个实施例中，图形核心2500可以包括在图19的图形处理器1910内。在至少一个实施例中，图形核心2500可以是图24B中统一的着色器核心 2455A-2455N。在至少一个实施例中，图形核心2500包括共享指令高速缓存2502、纹理单元2518和高速缓存/共享存储器2520，它们是图形核心2500 内的执行资源所共有的。在至少一个实施例中，图形核心2500可以包括多个切片(slice)2501A-2501N或每个核心的分区，图形处理器可以包括图形核心2500的多个实例。切片2501A-2501N可以包括支持逻辑，该支持逻辑包括本地指令高速缓存2504A-2504N、线程调度器2506A-2506N、线程分派器2508A-2508N和一组寄存器2510A-2510N。在至少一个实施例中，切片2501A-2501N可以包括一组附加功能单元(AFU)2512A-2512N、浮点单元(FPU)2514A-2514N、整数算术逻辑单元(ALU)2516A-2516N、地址计算单元(ACU)2513A-2513N、双精度浮点单元(DPFPU) 2515A-2515N和矩阵处理单元(MPU)2517A-2517N。

在一个实施例中，FPU 2514A-2514N可以执行单精度(32位)和半精度(16位)浮点运算，而DPFPU 2515A-2515N可以执行双精度(64位) 浮点运算点操作。在至少一个实施例中，ALU 2516A-2516N可以以8位、 16位和32位精度执行可变精度整数运算，并且可以被配置用于混合精度运算。在至少一个实施例中，MPU 2517A-2517N还可被配置用于混合精度矩阵运算，包括半精度浮点运算和8位整数运算。在至少一个实施例中，MPU 2517A-2517N可以执行各种矩阵操作以加速CUDA程序，包括使得能够支持加速的通用矩阵到矩阵乘法(GEMM)。在至少一个实施例中，AFU 2512A-2512N可以执行浮点数或整数单元不支持的附加逻辑运算，包括三角运算(例如，Sine、Cosine等)。

图25B示出了在至少一个实施例中的通用图形处理单元(GPGPU) 2530。在至少一个实施例中，GPGPU 2530是高度并行的并且适合于部署在多芯片模块上。在至少一个实施例中，GPGPU 2530可以被配置为使得高度并行的计算操作能够由GPU阵列来执行。在至少一个实施例中， GPGPU 2530可以直接链路到GPGPU 2530的其他实例，以创建多GPU集群以提高用于CUDA程序的执行时间。在至少一个实施例中，GPGPU 2530 包括主机接口2532以实现与主机处理器的连接。在至少一个实施例中，主机接口2532是PCIe接口。在至少一个实施例中，主机接口2532可以是厂商专用的通信接口或通信结构。在至少一个实施例中，GPGPU2530从主机处理器接收命令，并使用全局调度器2534将与那些命令相关联的执行线程分派给一组计算集群2536A-2536H。在至少一个实施例中，计算集群 2536A-2536H共享高速缓存存储器2538。在至少一个实施例中，高速缓存存储器2538可以用作计算集群2536A-2536H内的高速缓存存储器的高级高速缓存。

在至少一个实施例中，GPGPU 2530包括经由一组存储器控制器 2542A-2542B与计算集群2536A-2536H耦合的存储器2544A-2544B。在至少一个实施例中，存储器2544A-2544B可以包括各种类型的存储器设备，包括动态随机存取存储器(DRAM)或图形随机存取存储器，例如同步图形随机存取存储器(SGRAM)，包括图形双倍数据速率(GDDR)存储器。

在至少一个实施例中，计算集群2536A-2536H各自包括一组图形核心，诸如图25A的图形核心2500，其可以包括多种类型的整数和浮点逻辑单元，可以以各种精度执行计算操作，包括适合与CUDA程序相关的计算。例如，在至少一个实施例中，每个计算集群2536A-2536H中的浮点单元的至少一个子集可以配置为执行16位或32位浮点运算，而不同的浮点单元的子集可以配置为执行64位浮点运算。

在至少一个实施例中，GPGPU 2530的多个实例可以被配置为操作为计算集群。计算集群2536A-2536H可以实现用于同步和数据交换的任何技术上可行的通信技术。在至少一个实施例中，GPGPU 2530的多个实例通过主机接口2532进行通信。在至少一个实施例中，GPGPU 2530包括I/O 集线器2539，其将GPGPU 2530与GPU链路2540耦合，使得能够直接连接至GPGPU 2530的其他的实例。在至少一个实施例中，GPU链路2540 耦合到专用GPU到GPU桥接器，其使得能够在GPGPU 2530的多个实例之间进行通信和同步。在至少一个实施例中，GPU链路2540与高速互连耦合，以向其他GPGPU或并行处理器发送和接收数据。在至少一个实施例中，GPGPU 2530的多个实例位于单独的数据处理系统中，并经由可经由主机接口2532访问的网络设备进行通信。在至少一个实施例中，GPU链路 2540可被配置为能够连接到主机处理器，附加或替代主机接口2532。在至少一个实施例中，GPGPU 2530可以配置为执行CUDA程序。

图26A示出了根据至少一个实施例的并行处理器2600。在至少一个实施例中，并行处理器2600的各种组件可以使用一个或更多个集成电路设备来实现，例如可编程处理器、专用集成电路(ASIC)或FPGA。

在至少一个实施例中，并行处理器2600包括并行处理单元2602。在至少一个实施例中，并行处理单元2602包括I/O单元2604，其使得能够与其他设备进行通信，包括并行处理单元2602的其他实例。在至少一个实施例中，I/O单元2604可以直接连接到其他设备。在至少一个实施例中，I/O 单元2604通过使用集线器或交换机接口(例如，存储器集线器1405)与其他设备连接。在至少一个实施例中，存储器集线器1405与I/O单元2604 之间的连接形成通信链路。在至少一个实施例中，I/O单元2604与主机接口2606和存储器交叉开关2616连接，其中主机接口2606接收用于执行处理操作的命令，而存储器交叉开关2616接收用于执行存储器操作的命令。

在至少一个实施例中，当主机接口2606经由I/O单元2604接收命令缓冲区时，主机接口2606可以引导工作操作以执行那些命令到前端2608。在至少一个实施例中，前端2608与调度器2610耦合，调度器2610配置成将命令或其他工作项分配给处理阵列2612。在至少一个实施例中，调度器 2610确保在将任务分配给处理阵列2612中的处理阵列2612之前，处理阵列2612被正确地配置并且处于有效状态。在至少一个实施例中，调度器 2610通过在微控制器上执行的固件逻辑来实现。在至少一个实施例中，微控制器实现的调度器2610可配置成以粗粒度和细粒度执行复杂的调度和工作分配操作，从而实现对在处理阵列2612上执行的线程的快速抢占和环境切换。在至少一个实施例中，主机软件可以证明用于通过多个图形处理门铃之一在处理阵列2612上进行调度的工作负载。在至少一个实施例中，工作负载然后可以由包括调度器2610的微控制器内的调度器2610逻辑在处理阵列2612上自动分配。

在至少一个实施例中，处理阵列2612可以包括多达“N”个处理集群 (例如，集群2614A、集群2614B到集群2614N)。在至少一个实施例中，处理阵列2612的每个集群2614A-2614N可以执行大量并发线程。在至少一个实施例中，调度器2610可以使用各种调度和/或工作分配算法将工作分配给处理阵列2612的集群2614A-2614N，其可以根据每种程序或计算类型产生的工作负载而变化。在至少一个实施例中，调度可以由调度器2610 动态地处理，或者可以在配置为由处理阵列2612执行的程序逻辑的编译期间部分地由编译器逻辑来辅助。在至少一个实施例中，可将处理阵列2612 的不同的集群2614A-2614N分配用于处理不同类型的程序或用于执行不同类型的计算。

在至少一个实施例中，处理阵列2612可以配置成执行各种类型的并行处理操作。在至少一个实施例中，处理阵列2612配置成执行通用并行计算操作。例如，在至少一个实施例中，处理阵列2612可以包括执行处理任务的逻辑，该处理任务包括对视频和/或音频数据的过滤，执行建模操作，包括物理操作以及执行数据转换。

在至少一个实施例中，处理阵列2612配置成执行并行图形处理操作。在至少一个实施例中，处理阵列2612可以包括附加逻辑以支持这种图形处理操作的执行，包括但不限于执行纹理操作的纹理采样逻辑，以及镶嵌逻辑和其他顶点处理逻辑。在至少一个实施例中，处理阵列2612可以配置成执行与图形处理有关的着色器程序，例如但不限于顶点着色器、曲面细分着色器、几何着色器和像素着色器。在至少一个实施例中，并行处理单元 2602可以经由I/O单元2604从系统存储器传送数据以进行处理。在至少一个实施例中，在处理期间，可以在处理期间将传送的数据存储到片上存储器(例如，并行处理器存储器2622)，然后将其写回到系统存储器。

在至少一个实施例中，当并行处理单元2602用于执行图处理时，调度器2610可以配置成将处理工作负载划分为近似相等大小的任务，以更好地将图形处理操作分配给处理阵列2612的多个集群2614A-2614N。在至少一个实施例中，处理阵列2612的部分可以配置成执行不同类型的处理。例如，在至少一个实施例中，第一部分可以配置成执行顶点着色和拓扑生成，第二部分可以配置成执行镶嵌和几何着色，并且第三部分可以配置成执行像素着色或其他屏幕空间操作，以生成用于显示的渲染图像。在至少一个实施例中，可以将由集群2614A-2614N中的一个或更多个产生的中间数据存储在缓冲区中，以允许在集群2614A-2614N之间传输中间数据以进行进一步处理。

在至少一个实施例中，处理阵列2612可以经由调度器2610接收要执行的处理任务，该调度器2610从前端2608接收定义处理任务的命令。在至少一个实施例中，处理任务可以包括要被处理的数据的索引，例如可以包括表面(补丁)数据、原始数据、顶点数据和/或像素数据，以及状态参数和定义如何处理数据的命令(例如，要执行什么程序)。在至少一个实施例中，调度器2610可以配置成获取与任务相对应的索引，或者可以从前端 2608接收索引。在至少一个实施例中，前端2608可以配置成确保在启动由传入命令缓冲区(例如，批缓冲区(batch-buffer)、推送缓冲区等)指定的工作负载之前，处理阵列2612配置成有效状态。

在至少一个实施例中，并行处理单元2602的一个或更多个实例中的每一个可以与并行处理器存储器2622耦合。在至少一个实施例中，可以经由存储器交叉开关2616访问并行处理器存储器2622，所述存储器交叉开关 2616可以接收来自处理阵列2612以及I/O单元2604的存储器请求。在至少一个实施例中，存储器交叉开关2616可以经由存储器接口2618访问并行处理器存储器2622。在至少一个实施例中，存储器接口2618可以包括多个分区单元(例如，分区单元2620A、分区单元2620B到分区单元2620N)，其可各自耦合至并行处理器存储器2622的一部分(例如，存储器单元)。在至少一个实施例中，多个分区单元2620A-2620N为配置为等于存储器单元的数量，使得第一分区单元2620A具有对应的第一存储器单元2624A，第二分区单元2620B具有对应的存储器单元2624B，第N分区单元2620N 具有对应的第N存储器单元2624N。在至少一个实施例中，分区单元 2620A-2620N的数量可以不等于存储器设备的数量。

在至少一个实施例中，存储器单元2624A-2624N可以包括各种类型的存储器设备，包括动态随机存取存储器(DRAM)或图形随机存取存储器，例如同步图形随机存取存储器(SGRAM)，包括图形双倍数据速率(GDDR)存储器。在至少一个实施例中，存储器单元2624A-2624N还可包括3D堆叠存储器，包括但不限于高带宽存储器(HBM)。在至少一个实施例中，可以跨存储器单元2624A-2624N来存储诸如帧缓冲区或纹理映射的渲染目标，从而允许分区单元2620A-2620N并行地写入每个渲染目标的部分，以有效地使用并行处理器存储器2622的可用带宽。在至少一个实施例中，可以排除并行处理器存储器2622的本地实例，以有利于利用系统存储器与本地高速缓存存储器结合的统一存储器设计。

在至少一个实施例中，处理阵列2612的集群2614A-2614N中的任何一个都可以处理将被写入并行处理器存储器2622内的任何存储器单元 2624A-2624N中的数据。在至少一个实施例中，存储器交叉开关2616可以配置为将每个集群2614A-2614N的输出传输到任何分区单元2620A-2620N 或另一个集群2614A-2614N，集群2614A-2614N可以对输出执行其他处理操作。在至少一个实施例中，每个集群2614A-2614N可以通过存储器交叉开关2616与存储器接口2618通信，以从各种外部存储设备读取或写入各种外部存储设备。在至少一个实施例中，存储器交叉开关2616具有到存储器接口2618的连接以与I/O单元2604通信，以及到并行处理器存储器2622 的本地实例的连接，从而使不同处理集群2614A-2614N内的处理单元与系统存储器或不是并行处理单元2602本地的其他存储器进行通信。在至少一个实施例中，存储器交叉开关2616可以使用虚拟通道来分离集群 2614A-2614N和分区单元2620A-2620N之间的业务流。

在至少一个实施例中，可以在单个插入卡上提供并行处理单元2602 的多个实例，或者可以将多个插入卡互连。在至少一个实施例中，并行处理单元2602的不同实例可以配置成相互操作，即使不同实例具有不同数量的处理核心，不同数量的本地并行处理器存储器和/或其他配置差异。例如，在至少一个实施例中，并行处理单元2602的一些实例可以包括相对于其他实例而言更高精度的浮点单元。在至少一个实施例中，结合并行处理单元2602或并行处理器2600的一个或更多个实例的系统可以以各种配置和形式因素来实现，包括但不限于台式机、膝上型计算机或手持式个人计算机、服务器、工作站、游戏机和/或嵌入式系统。

图26B示出了根据至少一个实施例的处理集群2694。在至少一个实施例中，处理集群2694被包括在并行处理单元内。在至少一个实施例中，处理集群2694是图26的处理集群2614A-2614N之一的实例。在至少一个实施例中，处理集群2694可以配置成并行执行许多线程，其中术语“线程”是指在特定的一组输入数据上执行的特定程序的实例。在至少一个实施例中，单指令多数据(SIMD)指令发布技术用于支持大量线程的并行执行而无需提供多个独立的指令单元。在至少一个实施例中，使用单指令多线程 (SIMT)技术来支持并行执行大量一般同步的线程，这使用了公共指令单元，该公共指令单元配置成向每个处理集群2694内的一组处理引擎发出指令。

在至少一个实施例中，可以通过将处理任务分配给SIMT并行处理器的管线管理器2632来控制处理集群2694的操作。在至少一个实施例中，管线管理器2632从图26的调度器2610接收指令，通过图形多处理器2634 和/或纹理单元2636管理这些指令的执行。在至少一个实施例中，图形多处理器2634是SIMT并行处理器的示例性实例。然而，在至少一个实施例中，处理集群2694内可以包括不同架构的各种类型的SIMT并行处理器。在至少一个实施例中，在处理集群2694内可以包括图形多处理器2634的一个或更多个实例。在至少一个实施例中，图形多处理器2634可以处理数据，并且数据交叉开关2640可以用于将处理后的数据分发到多个可能的目的(包括其他着色器单元)地之一。在至少一个实施例中，管线管理器2632 可以通过指定要经由数据交叉开关2640分配的处理后的数据的目的地来促进处理后的数据的分配。

在至少一个实施例中，处理集群2694内的每个图形多处理器2634可以包括相同的一组功能执行逻辑(例如，算术逻辑单元、加载存储单元(LSU) 等)。在至少一个实施例中，可以以管线方式配置功能执行逻辑，其中可以在先前的指令完成之前发出新的指令。在至少一个实施例中，功能执行逻辑支持多种运算，包括整数和浮点算术、比较操作、布尔运算、移位和各种代数函数的计算。在至少一个实施例中，可以利用相同的功能单元硬件来执行不同的操作，并且可以存在功能单元的任何组合。

在至少一个实施例中，传送到处理集群2694的指令构成线程。在至少一个实施例中，跨一组并行处理引擎执行的一组线程是线程组。在至少一个实施例中，线程组在不同的输入数据上执行程序。在至少一个实施例中，线程组内的每个线程可被分配给图形多处理器2634内的不同处理引擎。在至少一个实施例中，线程组可包括比图形多处理器2634内的多个处理引擎更少的线程。在至少一个实施例中，当线程组包括的线程数少于处理引擎的数量时，一个或更多个处理引擎在正在处理该线程组的循环期间可能是空闲的。在至少一个实施例中，线程组还可以包括比图形多处理器2634 内的多个处理引擎更多的线程。在至少一个实施例中，当线程组包括比图形多处理器2634内的处理引擎的数量更多的线程时，可以在连续的时钟周期内执行处理。在至少一个实施例中，可以在图形多处理器2634上同时执行多个线程组。

在至少一个实施例中，图形多处理器2634包括内部高速缓存存储器，以执行加载和存储操作。在至少一个实施例中，图形多处理器2634可以放弃内部高速缓存并使用处理集群2694内的高速缓存存储器(例如，L1高速缓存2648)。在至少一个实施例中，每个图形多处理器2634还可以访问分区单元(例如，图26A的分区单元2620A-2620N)内的L2高速缓存，这些分区单元在所有处理集群2694之间共享并且可以用于在线程之间传输数据。在至少一个实施例中，图形多处理器2634还可以访问片外全局存储器，其可以包括本地并行处理器存储器和/或系统存储器中的一个或更多个。在至少一个实施例中，并行处理单元2602外部的任何存储器都可以用作全局存储器。在至少一个实施例中，处理集群2694包括图形多处理器2634 的多个实例，它们可以共享可以存储在L1高速缓存2648中的公共指令和数据。

在至少一个实施例中，每个处理集群2694可以包括配置成将虚拟地址映射为物理地址的MMU 2645。在至少一个实施例中，MMU 2645的一个或更多个实例可以驻留在图26的存储器接口2618内。在至少一个实施例中，MMU 2645包括一组页表条目(PTE)，其用于将虚拟地址映射到图块(谈论有关图块的更多信息)的物理地址以及可选地映射到高速缓存行索引。在至少一个实施例中，MMU 2645可以包括地址转换后备缓冲区 (TLB)或可以驻留在图形多处理器2634或L1高速缓存2648或处理集群 2694内的高速缓存。在至少一个实施例中，处理物理地址以分配表面数据访问局部性，以便在分区单元之间进行有效的请求交织。在至少一个实施例中，高速缓存行索引可以用于确定对高速缓存线的请求是命中还是未命中。

在至少一个实施例中，可以配置处理集群2694，使得每个图形多处理器2634耦合到纹理单元2636，以执行纹理映射操作，例如，可以涉及确定纹理样本位置、读取纹理数据以及过滤纹理数据。在至少一个实施例中，根据需要从内部纹理L1高速缓存(未示出)或从图形多处理器2634内的 L1高速缓存中读取纹理数据，并从L2高速缓存、本地并行处理器存储器或系统存储器中获取纹理数据。在至少一个实施例中，每个图形多处理器 2634将处理后的任务输出到数据交叉开关2640，以将处理后的任务提供给另一处理集群2694以进行进一步处理或将处理后的任务存储在L2高速缓存、本地并行处理器存储器、或经由存储器交叉开关2616的系统存储器中。在至少一个实施例中，光栅前操作单元(preROP)2642配置成从图形多处理器2634接收数据，将数据引导至ROP单元，该ROP单元可以与本文所述的分区单元(例如，图26的分区单元2620A-2620N)一起定位。在至少一个实施例中，PreROP 2642单元可以执行用于颜色混合的优化、组织像素颜色数据以及执行地址转换。

图26C示出了根据至少一个实施例的图形多处理器2696。在至少一个实施例中，图形多处理器2696是图26B的图形多处理器2634。在至少一个实施例中，图形多处理器2696与处理集群2694的管线管理器2632耦合。在至少一个实施例中，图形多处理器2696具有执行管线，该执行管线包括但不限于指令高速缓存2652、指令单元2654、地址映射单元2656、寄存器文件2658、一个或更多个GPGPU核心2662和一个或更多个LSU 2666。 GPGPU核心2662和LSU 2666与高速缓存存储器2672和共享存储器2670 通过存储器和高速缓存互连2668耦合。

在至少一个实施例中，指令高速缓存2652从管线管理器2632接收要执行的指令流。在至少一个实施例中，将指令高速缓存在指令高速缓存2652 中并将其分派以供指令单元2654执行。在一个实施例中，指令单元2654 可以分派指令作为线程组(例如，线程束)，将线程组的每个线程分配给 GPGPU核心2662内的不同执行单元。在至少一个实施例中，指令可以通过在统一地址空间内指定地址来访问任何本地、共享或全局地址空间。在至少一个实施例中，地址映射单元2656可以用于将统一地址空间中的地址转换成可以由LSU 2666访问的不同的存储器地址。

在至少一个实施例中，寄存器文件2658为图形多处理器2696的功能单元提供了一组寄存器。在至少一个实施例中，寄存器文件2658为连接到图形多处理器2696的功能单元(例如，GPGPU核心2662、LSU 2666)的数据路径的操作数提供了临时存储。在至少一个实施例中，在每个功能单元之间划分寄存器文件2658，使得为每个功能单元分配寄存器文件2658 的专用部分。在至少一个实施例中，寄存器文件2658在图形多处理器2696 正在执行的不同线程组之间划分。

在至少一个实施例中，GPGPU核心2662可以各自包括用于执行图多处理器2696的指令的FPU和/或ALU。GPGPU核心2662在架构上可以相似或架构可能有所不同。在至少一个实施例中，GPGPU核心2662的第一部分包括单精度FPU和整数ALU，而GPGPU核心的第二部分包括双精度 FPU。在至少一个实施例中，FPU可以实现用于浮点算法的IEEE 754-2608 标准或启用可变精度浮点算法。在至少一个实施例中，图形多处理器2696 可以另外包括一个或更多个固定功能或特殊功能单元，以执行特定功能，诸如复制矩形或像素混合操作。在至少一个实施例中，GPGPU核心2662 中的一个或更多个也可以包括固定或特殊功能逻辑。

在至少一个实施例中，GPGPU核心2662包括能够对多组数据执行单个指令的SIMD逻辑。在至少一个实施例中，GPGPU核心2662可以物理地执行SIMD4、SIMD8和SIMD9指令，并且在逻辑上执行SIMD1、SIMD2 和SIMD32指令。在至少一个实施例中，用于GPGPU核心的SIMD指令可以在编译时由着色器编译器生成，或者在执行针对单程序多数据(SPMD) 或SIMT架构编写和编译的程序时自动生成。在至少一个实施例中，可以通过单个SIMD指令来执行为SIMT执行模型配置的程序的多个线程。例如，在至少一个实施例中，可以通过单个SIMD8逻辑单元并行执行执行相同或相似操作的八个SIMT线程。

在至少一个实施例中，存储器和高速缓存互连2668是将图形多处理器2696的每个功能单元连接到寄存器文件2658和共享存储器2670的互连网络。在至少一个实施例中，存储器和高速缓存互连2668是交叉开关互连，其允许LSU 2666在共享存储器2670和寄存器文件2658之间实现加载和存储操作。在至少一个实施例中，寄存器文件2658可以以与GPGPU核心2662 相同的频率操作，从而在GPGPU核心2662和寄存器文件2658之间进行数据传输的延迟非常低。在至少一个实施例中，共享存储器2670可以用于启用在图形多处理器2696内的功能单元上执行的线程之间的通信。在至少一个实施例中，高速缓存存储器2672可以用作例如数据高速缓存，以高速缓存在功能单元和纹理单元2636之间通信的纹理数据。在至少一个实施例中，共享存储器2670也可以用作程序管理的高速缓存。在至少一个实施例中，除了存储在高速缓存存储器2672中的自动高速缓存的数据之外，在 GPGPU核心2662上执行的线程还可以以编程方式将数据存储在共享存储器中。

在至少一个实施例中，如本文所述的并行处理器或GPGPU通信地耦合到主机/处理器核心，以加速图形操作、机器学习操作、图案分析操作以及各种通用GPU(GPGPU)功能。在至少一个实施例中，GPU可以通过总线或其他互连(例如，诸如PCIe或NVLink的高速互连)通信地耦合到主机处理器/核心。在至少一个实施例中，GPU可以与核心集成在相同的封装或芯片上，并通过内部处理器总线/互连(即，封装或芯片的内部)通信地耦合到核心。在至少一个实施例中，不管GPU连接的方式如何，处理器核心可以以WD包含的命令/指令序列的形式向GPU分配工作。在至少一个实施例中，GPU然后使用专用电路/逻辑来有效地处理这些命令/指令。

图27示出了根据至少一个实施例的图形处理器2700。在至少一个实施例中，图形处理器2700包括环形互连2702、管线前端2704、媒体引擎 2737和图形核心2780A-2780N。在至少一个实施例中，环形互连2702将图形处理器2700耦合到其他处理单元，包括其他图形处理器或一个或更多个通用处理器核心。在至少一个实施例中，图形处理器2700是集成在多核心处理系统内的许多处理器之一。

在至少一个实施例中，图形处理器2700经由环形互连2702接收多批命令。在至少一个实施例中，输入命令由管线前端2704中的命令流转化器 2703解释。在至少一个实施例中，图形处理器2700包括可缩放执行逻辑，以经由图形核心2780A-2780N执行3D几何处理和媒体处理。在至少一个实施例中，对于3D几何处理命令，命令流转化器2703将命令提供给几何管线2736。在至少一个实施例中，对于至少一些媒体处理命令，命令流转化器2703将命令提供给视频前端2734，其与媒体引擎2737耦合。在至少一个实施例中，媒体引擎2737包括用于视频和图像后处理的视频质量引擎 (VQE)2730，以及用于提供硬件加速媒体数据编码和解码的多格式编码/ 解码(MFX)2733引擎。在至少一个实施例中，几何管线2736和媒体引擎2737各自生成用于由至少一个图形核心2780A提供的线程执行资源的执行线程。

在至少一个实施例中，图形处理器2700包括以模块化图形核心 2780A-2780N(有时称为核心切片)为特征的可缩放线程执行资源，每个模块核心具有多个子核心2750A-2750N、2760A-2760N(有时称为核心子切片)。在至少一个实施例中，图形处理器2700可以具有任意数量的图形核心2780A至2780N。在至少一个实施例中，图形处理器2700包括具有至少第一子核心2750A和第二子核心2760A的图形核心2780A。在至少一个实施例中，图形处理器2700是具有单个子核心(例如2750A)的低功率处理器。在至少一个实施例中，图形处理器2700包括多个图形核心 2780A-2780N，每个图形核心包括一组第一子核心2750A-2750N和一组第二子核心2760A-2760N。在至少一个实施例中，第一子核心2750A-2750N 中的每个子核心至少包括第一组执行单元(EU)2752A-2752N和媒体/纹理采样器2754A-2754N。在至少一个实施例中，第二子核心2760A-2760N中的每个子核心至少包括第二组执行单元2762A-2762N和采样器 2764A-2764N。在至少一个实施例中，每个子核心2750A-2750N、 2760A-2760N共享一组共享资源2770A-2770N。在至少一个实施例中，共享资源包括共享高速缓冲存储器和像素操作逻辑。

图28示出了根据至少一个实施例的用于处理器2800。在至少一个实施例中，处理器2800可以包括但不限于执行指令的逻辑电路。在至少一个实施例中，处理器2800可以执行指令，包括x86指令、ARM指令、用于 ASIC的专用指令等。在至少一个实施例中，处理器2810可以包括用于存储封装数据的寄存器，例如作为加利福尼亚州圣克拉拉市英特尔公司采用 MMX技术启用的微处理器中的64位宽MMXTM寄存器。在至少一个实施例中，整数和浮点数形式可用的MMX寄存器可以与封装的数据元素一起运行，所述封装的数据元素伴随SIMD和流式SIMD扩展(“SSE”)指令。在至少一个实施例中，与SSE2、SSE3、SSE4、AVX或更高版本(一般称为“SSEx”)技术有关的128位宽XMM寄存器可以保存此类封装数据操作数。在至少一个实施例中，处理器2810可以执行指令以加速CUAD程序。

在至少一个实施例中，处理器2800包括有序前端(“前端”)2801，以提取要执行的指令并准备稍后在处理器管线中使用的指令。在至少一个实施例中，前端2801可以包括几个单元。在至少一个实施例中，指令预取器2826从存储器中获取指令并将指令提供给指令解码器2828，指令解码器 2828又对指令进行解码或解释。例如，在至少一个实施例中，指令解码器2828将接收到的指令解码用于执行的所谓的“微指令”或“微操作”(也称为“微操作”或“微指令”)的一个或更多个操作。在至少一个实施例中，指令解码器2828将指令解析为操作码以及相应的数据和控制字段，其可以由微架构用来使用以执行操作。在至少一个实施例中，跟踪高速缓存2830可以将解码的微指令组装成微指令队列2834中的程序排序的序列或追踪以供执行。在至少一个实施例中，当追踪高速缓存2830遇到复杂指令时，微码 ROM2832提供完成操作所需的微指令。

在至少一个实施例中，可以将一些指令转换成单个微操作，而另一些指令则需要几个微操作来完成全部操作。在至少一个实施例中，如果需要多于四个的微指令来完成一条指令，则指令解码器2828可以访问微码ROM 2832以执行指令。在至少一个实施例中，可以将指令解码为少量的微指令以在指令解码器2828处进行处理。在至少一个实施例中，如果需要多个微指令完成操作，则可以将指令存储在微码ROM 2832中。在至少一个实施例中，追踪高速缓存器2830参考入口点可编程逻辑阵列(“PLA”)以确定正确的微指令指针，用于根据至少一个实施例从微码ROM 2832读取微码序列以完成一个或更多个指令。在至少一个实施例中，在微码ROM 2832 完成对指令的微操作排序之后，机器的前端2801可以恢复从追踪高速缓存 2830获取微操作。

在至少一个实施例中，乱序执行引擎(“乱序引擎”)2803可以准备用于执行的指令。在至少一个实施例中，乱序执行逻辑具有多个缓冲区，以使指令流平滑并重新排序，以在指令沿管线下降并被调度执行时优化性能。乱序执行引擎2803包括但不限于分配器/寄存器重命名器2840、存储器微指令队列2842、整数/浮点微指令队列2844、存储器调度器2846、快速调度器2802、慢速/通用浮点调度器(“慢速/通用FP调度器”)2804和简单浮点调度器(“简单FP调度器”)2806。在至少一个实施例中，快速调度器2802、慢速/通用浮点调度器2804和简单浮点调度器2806也统称为“微指令调度器 2802、2804、2806”。分配器/寄存器重命名器2840分配每个微指令按顺序执行所需要的机器缓冲区和资源。在至少一个实施例中，分配器/寄存器重命名器2840将逻辑寄存器重命名为寄存器文件中的条目。在至少一个实施例中，分配器/寄存器重命名器2840还为两个微指令队列之一中的每个微指令分配条目，存储器微指令队列2842用于存储器操作和整数/浮点微指令队列2844用于非存储器操作，在存储器调度器2846和微指令调度器2802、2804、2806的前面。在至少一个实施例中，微指令调度器2802、2804、2806 基于它们的从属输入寄存器操作数源的就绪性和需要完成的执行资源微指令的可用性来确定何时准备好执行微指令。在至少一个实施例中，至少一个实施例的快速调度器2802可以在主时钟周期的每个一半上调度，而慢速 /通用浮点调度器2804和简单浮点调度器2806可以在每个主处理器时钟周期调度一次。在至少一个实施例中，微指令调度器2802、2804、2806对调度端口进行仲裁，以调度用于执行的微指令。

在至少一个实施例中，执行块2811包括但不限于整数寄存器文件/支路网络2808、浮点寄存器文件/支路网络(“FP寄存器文件/支路网络”)2810、地址生成单元(“AGU”)2812和2814、快速算术逻辑单元(“快速ALU”) 2816和2818、慢速ALU 2820、浮点ALU(“FP”)2822和浮点移动单元 (“FP移动”)2824。在至少一个实施例中，整数寄存器文件/支路网络2808和浮点寄存器文件/旁路网络2810在本文中也称为“寄存器文件2808、2810”。在至少一个实施例中，AGUS 2812和2814、快速ALU 2816和2818、慢速 ALU 2820、浮点ALU 2822和浮点移动单元2824在本文中也称为“执行单元2812、2814、2816、2818、2820、2822和2824”。在至少一个实施例中，执行框可以包括但不限于任意数量(包括零)和类型的寄存器文件、支路网络、地址生成单元和执行单元(以任何组合)。

在至少一个实施例中，寄存器文件2808、2810可以布置在微指令调度器2802、2804、2806与执行单元2812、2814、2816、2818、2820、2822 和2824之间。在至少一个实施例中，整数寄存器文件/支路网络2808执行整数运算。在至少一个实施例中，浮点寄存器文件/支路网络2810执行浮点操作。在至少一个实施例中，寄存器文件2808、2810中的每一个可以包括但不限于支路网络，该支路网络可以绕过或转发尚未写入寄存器文件中的刚刚完成的结果到新的从属对象。在至少一个实施例中，寄存器文件2808、 2810可以彼此通信数据。在至少一个实施例中，整数寄存器文件/支路网络 2808可以包括但不限于两个单独的寄存器文件、一个寄存器文件用于低阶 32位数据，第二寄存器文件用于高阶32位数据。在至少一个实施例中，浮点寄存器文件/支路网络2810可以包括但不限于128位宽的条目，因为浮点指令通常具有宽度为64至128位的操作数。

在至少一个实施例中，执行单元2812、2814、2816、2818、2820、2822、 2824可以执行指令。在至少一个实施例中，寄存器文件2808、2810存储微指令需要执行的整数和浮点数据操作数值。在至少一个实施例中，处理器 2800可以包括但不限于任意数量的执行单元2812、2814、2816、2818、2820、 2822、2824及其组合。在至少一个实施例中，浮点ALU 2822和浮点移动单元2824，可以执行浮点、MMX、SIMD、AVX和SSE或其他操作，包括专门的机器学习指令。在至少一个实施例中，浮点ALU 2822可以包括但不限于64位乘64位浮点除法器，以执行除法、平方根和余数微操作。在至少一个实施例中，可以用浮点硬件来处理涉及浮点值的指令。在至少一个实施例中，可以将ALU操作传递给快速ALU 2816、2818。在至少一个实施例中，快速ALUS 2816、2818可以以半个时钟周期的有效延迟执行快速操作。在至少一个实施例中，大多数复杂的整数运算进入慢速ALU 2820，因为慢速ALU 2820可以包括但不限于用于长延迟类型操作的整数执行硬件，例如乘法器、移位、标志逻辑和分支处理。在至少一个实施例中，存储器加载/存储操作可以由AGUS 2812、2814执行。在至少一个实施例中，快速ALU 2816、快速ALU 2818和慢速ALU 2820可以对64位数据操作数执行整数运算。在至少一个实施例中，可以实现快速ALU 2816、快速ALU 2818和慢速ALU 2820以支持包括16、32、128、256等的各种数据位大小。在至少一个实施例中，浮点ALU 2822和浮点移动单元2824可以实现为支持具有各种宽度的位的一定范围的操作数。在至少一个实施例中，浮点ALU2822和浮点移动单元2824可以结合SIMD和多媒体指令对 128位宽封装数据操作数进行操作。

在至少一个实施例中，微指令调度器2802、2804、2806在父加载完成执行之前调度从属操作。在至少一个实施例中，由于可以在处理器2800 中推测性地调度和执行微指令，处理器2800还可以包括用于处理存储器未命中的逻辑。在至少一个实施例中，如果数据高速缓存中的数据加载未命中，则可能存在在管线中正在运行的从属操作，其使调度器暂时没有正确的数据。在至少一个实施例中，一种重放机制追踪踪并重新执行使用不正确数据的指令。在至少一个实施例中，可能需要重放从属操作并且可以允许完成独立操作。在至少一个实施例中，处理器的至少一个实施例的调度器和重放机制也可以设计为捕获用于文本串比较操作的指令序列。

在至少一个实施例中，术语“寄存器”可以指代可以用作识别操作数的指令的一部分的机载处理器存储位置。在至少一个实施例中，寄存器可以是那些可以从处理器外部使用的寄存器(从程序员的角度来看)。在至少一个实施例中，寄存器可能不限于特定类型的电路。相反，在至少一个实施例中，寄存器可以存储数据、提供数据并执行本文描述的功能。在至少一个实施例中，本文描述的寄存器可以通过处理器内的电路使用多种不同技术来实现，例如专用物理寄存器、使用寄存器重命名动态分配的物理寄存器、专用和动态分配的物理寄存器的组合等。在至少一个实施例中，整数寄存器存储32位整数数据。至少一个实施例的寄存器文件还包含八个用于封装数据的多媒体SIMD寄存器。

图29示出了根据至少一个实施例的处理器2900。在至少一个实施例中，处理器2900包括但不限于一个或更多个处理器核心(核心) 2902A-2902N、集成存储器控制器2914和集成图形处理器2908。在至少一个实施例中，处理器2900可以包括直至并包括由虚线框表示的附加处理器核心2902N的附加核心。在至少一个实施例中，每个处理器核心 2902A-2902N包括一个或更多个内部高速缓存单元2904A-2904N。在至少一个实施例中，每个处理器核心还可以访问一个或更多个共享高速缓存的单元2906。

在至少一个实施例中，内部高速缓存单元2904A-2904N和共享高速缓存单元2906表示处理器2900内的高速缓存存储器层次结构。在至少一个实施例中，高速缓存存储器单元2904A-2904N可以包括每个处理器核心内的至少一级指令和数据以及共享中级缓存中的一级或更多级缓存，例如 L2、L3、4级(L4)或其他级别的缓存，其中在外部存储器之前将最高级别的缓存归类为LLC。在至少一个实施例中，高速缓存一致性逻辑维持各种高速缓存单元2906和2904A-2904N之间的一致性。

在至少一个实施例中，处理器2900还可包括一组一个或更多个总线控制器单元2916和系统代理核心2910。在至少一个实施例中，一个或更多个总线控制器单元2916管理一组外围总线，例如一个或更多个PCI或PCI Express总线。在至少一个实施例中，系统代理核心2910为各种处理器组件提供管理功能。在至少一个实施例中，系统代理核心2910包括一个或更多个集成存储器控制器2914，以管理对各种外部存储器设备(未示出)的访问。

在至少一个实施例中，一个或更多个处理器核心2902A-2902N包括对多线程同时进行的支持。在至少一个实施例中，系统代理核心2910包括用于在多线程处理期间协调和操作处理器核心2902A-2902N的组件。在至少一个实施例中，系统代理核心2910可以另外包括电源控制单元(PCU)，该电源控制单元包括逻辑和组件以调节处理器核心2902A-2902N和图形处理器2908的一个或更多个电源状态。

在至少一个实施例中，处理器2900另外包括图形处理器2908以执行图处理操作。在至少一个实施例中，图形处理器2908与共享高速缓存单元 2906和包括一个或更多个集成存储器控制器2914的系统代理核心2910耦合。在至少一个实施例中，系统代理核心2910还包括用于驱动图形处理器输出到一个或更多个耦合的显示器的显示器控制器2911。在至少一个实施例中，显示器控制器2911也可以是经由至少一个互连与图形处理器2908 耦合的独立模块，或者可以集成在图形处理器2908内。

在至少一个实施例中，基于环的互连单元2912用于耦合处理器2900 的内部组件。在至少一个实施例中，可以使用替代性互连单元，例如点对点互连、交换互连或其他技术。在至少一个实施例中，图形处理器2908经由I/O链路2913与环形互连2912耦合。

在至少一个实施例中，I/O链路2913代表多种I/O互连中的至少一种，包括促进各种处理器组件与高性能嵌入式存储器模块2918(例如eDRAM 模块)之间的通信的封装I/O互连。在至少一个实施例中，处理器核心 2902A-2902N和图形处理器2908中的每一个使用嵌入式存储器模块2918 作为共享的LLC。

在至少一个实施例中，处理器核心2902A-2902N是执行公共指令集架构的同质核心。在至少一个实施例中，处理器核心2902A-2902N在ISA 方面是异构的，其中一个或更多个处理器核心2902A-2902N执行公共指令集，而一个或更多个其他处理器核心2902A-2902N执行公共指令集或不同指令集的子集。在至少一个实施例中，就微架构而言，处理器核心2902A-2902N是异构的，其中具有相对较高功耗的一个或更多个核心与具有较低功耗的一个或更多个功率核心耦合。在至少一个实施例中，处理器 2900可以实现在一个或更多个芯片上或被实现为SoC集成电路。

图30示出了根据所描述的至少一个实施例的图形处理器核心3000。在至少一个实施例中，图形处理器核心3000被包括在图形核心阵列内。在至少一个实施例中，图形处理器核心3000(有时称为核心切片)可以是模块化图形处理器内的一个或更多个图形核心。在至少一个实施例中，图形处理器核心3000是一个图形核心切片的示例，并且本文所述的图形处理器可以基于目标功率和性能包络线包括多个图形核心切片。在至少一个实施例中，每个图形核心3000可以包括与多个子核心3001A-3001F耦合的固定功能块3030，也称为子切片，其包括通用和固定功能逻辑的模块块。

在至少一个实施例中，固定功能块3030包括几何/固定功能管线3036，例如，在较低性能和/或较低功率的图形处理器实施方式中，该几何/固定功能管线3036可以由图形处理器3000中的所有子核心共享。在至少一个实施例中，几何/固定功能管线3036包括3D固定功能管线、视频前端单元，线程产生器和线程分派器以及管理统一返回缓冲区的统一返回缓冲区管理器。

在至少一个实施例中，固定功能块3030还包括图形SoC接口3037、图形微控制器3038和媒体管线3039。图形SoC接口3037提供了图形核心 3000以及SoC集成电路系统中的其他处理器核心之间的接口。在至少一个实施例中，图形微控制器3038是可编程子处理器，其可配置为管理图形处理器3000的各种功能，包括线程分派、调度和抢占。在至少一个实施例中，媒体管线3039包括有助于对包括图像和视频数据的多媒体数据进行解码、编码、预处理和/或后处理的逻辑。在至少一个实施例中，媒体管线3039 经由对子核心3001-3001F内的计算或采样逻辑的请求来实现媒体操作。

在至少一个实施例中，SoC接口3037使图形核心3000能够与通用应用处理器核心(例如，CPU)和/或SoC内的其他组件通信，包括存储器层次结构元素，诸如共享的LLC存储器、系统RAM和/或嵌入式片上或封装 DRAM。在至少一个实施例中，SoC接口3037还可以使得能够与SoC内的固定功能设备(例如，相机成像管线)进行通信，并且使得能够使用和/或实现可以在图形核心3000和SoC内部的CPU之间共享的全局存储器原子。在至少一个实施例中，SoC接口3037还可以实现用于图形核心3000的电源管理控制，并且启用图形核心3000的时钟域与SoC内的其他时钟域之间的接口。在至少一个实施例中，SoC接口3037使得能够从命令流转化器和全局线程分派器接收命令缓冲区，其配置为向图形处理器内的一个或更多个图形核心中的每一个提供命令和指令。在至少一个实施例中，当要执行媒体操作时，可以将命令和指令分派给媒体管线3039，或者当要执行图处理操作时，可以将其分配给几何形状和固定功能管线(例如，几何形状和固定功能管线3036、几何形状和固定功能管线3014)。

在至少一个实施例中，图形微控制器3038可以配置为对图形核心3000执行各种调度和管理任务。在至少一个实施例中，图形微控制器3038 可以在子核心3001A-3001F中的执行单元(EU)阵列3002A-3002F、 3004A-3004F内的各种图形并行引擎上执行图和/或计算工作负载调度。在至少一个实施例中，在包括图形核心3000的SoC的CPU核心上执行的主机软件可以提交多个图形处理器门铃之一的工作负载，其调用适当的图形引擎上的调度操作。在至少一个实施例中，调度操作包括确定接下来要运行哪个工作负载、将工作负载提交给命令流转化器、抢先在引擎上运行的现有工作负载、监控工作负载的进度以及在工作负载完成时通知主机软件。在至少一个实施例中，图形微控制器3038还可以促进图形核心3000的低功率或空闲状态，从而为图形核心3000提供在图形核心3000内独立于操作系统和/或系统上的图形驱动器软件的跨低功率状态转换的保存和恢复寄存器的能力。

在至少一个实施例中，图形核心3000可以具有比所示的子核心 3001A-3001F更多或更少的子核心，达N个模块化子核心。对于每组N个子核心，在至少一个实施例中，图形核心3000还可以包括共享功能逻辑 3010、共享和/或高速缓存存储器3012、几何/固定功能管线3014以及附加的固定功能逻辑3016以加速各种图形和计算处理操作。在至少一个实施例中，共享功能逻辑3010可以包括可由图形核心3000内的每个N个子核心共享的逻辑单元(例如，采样器、数学和/或线程间通信逻辑)。共享和/ 或高速缓存存储器3012可以是图形核心3000内的N个子核心3001A-3001F 的LLC，并且还可以用作可由多个子核心访问的共享存储器。在至少一个实施例中，可以包括几何/固定功能管线3014来代替固定功能块3030内的几何/固定功能管线3036，并且可以包括相同或相似的逻辑单元。

在至少一个实施例中，图形核心3000包括附加的固定功能逻辑3016，其可以包括供图形核心3000使用的各种固定功能加速逻辑。在至少一个实施例中，附加的固定功能逻辑3016包括用于仅位置着色中使用的附加的几何管线。在仅位置着色中，存在至少两个几何管线，而在几何/固定功能管线3016、3036内的完整几何管线和剔除管线中，其是可以包括在附加的固定功能逻辑3016中的附加几何管线。在至少一个实施例中，剔除管线是完整几何管线的修整版。在至少一个实施例中，完整管线和剔除管线可以执行应用程序的不同实例，每个实例具有单独的环境。在至少一个实施例中，仅位置着色可以隐藏被丢弃的三角形的长剔除运行，从而在某些情况下可以更早地完成着色。例如，在至少一个实施例中，附加固定功能逻辑3016 中的剔除管线逻辑可以与主应用程序并行执行位置着色器，并且通常比完整管线更快地生成关键结果，因为剔除管线获取并遮蔽顶点的位置属性，无需执行光栅化和将像素渲染到帧缓冲区。在至少一个实施例中，剔除管线可以使用生成的临界结果来计算所有三角形的可见性信息，而与这些三角形是否被剔除无关。在至少一个实施例中，完整管线(在这种情况下可以称为重播管线)可以消耗可见性信息来跳过剔除的三角形以仅遮盖最终传递到光栅化阶段的可见三角形。

在至少一个实施例中，附加的固定功能逻辑3016还可包括通用目标处理加速逻辑，例如固定功能矩阵乘法逻辑，用于实现减速CUAD程序。

在至少一个实施例中，在每个图形子核心3001A-3001F内包括一组执行资源，其可用于响应于图形管线、媒体管线或着色器程序的请求来执行图、媒体和计算操作。在至少一个实施例中，图形子核心3001A-3001F包括多个EU阵列3002A-3002F、3004A-3004F，线程分派和线程间通信(TD/IC) 逻辑3003A-3003F，3D(例如，纹理)采样器3005A-3005F，媒体采样器 3006A-3006F，着色器处理器3007A-3007F和共享本地存储器(SLM) 3008A-3008F。EU阵列3002A-3002F、3004A-3004F每个都包含多个执行单元，这些执行单元是GUGPU，能够为图形、媒体或计算操作提供服务，执行浮点和整数/定点逻辑运算，包括图形、媒体或计算着色器程序。在至少一个实施例中，TD/IC逻辑3003A-3003F为子核心内的执行单元执行本地线程分派和线程控制操作，并促进在子核心的执行单元上执行的线程之间的通信。在至少一个实施例中，3D采样器3005A-3005F可以将与纹理或其他3D图形相关的数据读取到存储器中。在至少一个实施例中，3D采样器可以基于与给定纹理相关联的配置的采样状态和纹理格式来不同地读取纹理数据。在至少一个实施例中，媒体采样器3006A-3006F可以基于与媒体数据相关联的类型和格式来执行类似的读取操作。在至少一个实施例中，每个图形子核心3001A-3001F可以可替代地包括统一的3D和媒体采样器。在至少一个实施例中，在每个子核心3001A-3001F内的执行单元上执行的线程可以利用每个子核心内的共享本地存储器3008A-3008F，以使在线程组内执行的线程能够使用片上存储器的公共池来执行。

图31示出了根据至少一个实施例的并行处理单元(“PPU”)3100。在至少一个实施例中，PPU 3100配置有机器可读代码，该机器可读代码如果由PPU 3100执行，则使得PPU3100执行贯穿本文描述的一些或全部过程和技术。在至少一个实施例中，PPU 3100是在一个或更多个集成电路设备上实现的多线程处理器，并且利用多线程作为被设计为处理在多个线程上并行执行的计算机可读指令(也称为机器可读指令或简单的指令)的延迟隐藏技术。在至少一个实施例中，线程是指执行线程，并且是被配置为由 PPU 3100执行的一组指令的实例。在至少一个实施例中，PPU 3100是图形处理单元(“GPU”)，图形处理单元配置为实现用于处理三维(“3D”)图形数据的图形渲染管线，以便生成用于在显示设备(诸如LCD设备)上显示的二维(“2D”)图像数据。在至少一个实施例中，PPU 3100用于执行计算，诸如线性代数运算和机器学习运算。图31仅出于说明性目的示出了示例并行处理器，并且应被解释为在至少一个实施例中实现的处理器架构的非限制性示例。

在至少一个实施例中，一个或更多个PPU 3100配置成加速高性能计算(“HPC”)、数据中心和机器学习应用程序。在至少一个实施例中，一个或更多个PPU 3100配置成加速CUDA程序。在至少一个实施例中，PPU 3100包括但不限于I/O单元3106、前端单元3110、调度器单元3112、工作分配单元3114、集线器3116、交叉开关(“Xbar”)3120、一个或更多个通用处理集群(“GPC”)3118和一个或更多个分区单元(“存储器分区单元”) 3122。在至少一个实施例中，PPU 3100通过一个或更多个高速GPU互连 (“GPU互连”)3108连接到主机处理器或其他PPU 3100。在至少一个实施例中，PPU 3100通过系统总线或互连3102连接到主机处理器或其他外围设备。在一实施例中，PPU 3100连接到包括一个或更多个存储器设备(“存储器”)3104的本地存储器。在至少一个实施例中，存储器设备3104包括但不限于一个或更多个动态随机存取存储器(“DRAM”)设备。在至少一个实施例中，一个或更多个DRAM设备配置和/或可配置为高带宽存储器 (“HBM”)子系统，并且在每个设备内堆叠有多个DRAM管芯。

在至少一个实施例中，高速GPU互连3108可以指代系统使用其来进行缩放的基于线的多通道通信链路，并包括与一个或更多个CPU结合的一个或更多个PPU 3100(“CPU”)，支持PPU 3100和CPU之间的高速缓存一致性以及CPU主控。在至少一个实施例中，高速GPU互连3108通过集线器3116将数据和/或命令传输到PPU 3100的其他单元，例如一个或更多个复制引擎、视频编码器、视频解码器、电源管理单元和/或在图31中可能未明确示出的其他组件。

在至少一个实施例中，I/O单元3106配置为通过系统总线3102从主机处理器(图31中未示出)发送和接收通信(例如，命令、数据)。在至少一个实施例中，I/O单元3106直接通过系统总线3102或通过一个或更多个中间设备(例如内存桥)与主机处理器通信。在至少一个实施例中，I/O 单元3106可以经由系统总线3102与一个或更多个其他处理器(例如一个或更多个PPU 3100)通信。在至少一个实施例中，I/O单元3106实现PCIe 接口，用于通过PCIe总线进行通信。在至少一个实施例中，I/O单元3106 实现用于与外部设备通信的接口。

在至少一个实施例中，I/O单元3106对经由系统总线3102接收的分组进行解码。在至少一个实施例中，至少一些分组表示被配置为使PPU 3100 执行各种操作的命令。在至少一个实施例中，I/O单元3106如命令所指定的那样将解码的命令发送到PPU 3100的各种其他单元。在至少一个实施例中，命令被发送到前端单元3110和/或被发送到集线器3116或PPU 3100 的其他单元，例如一个或更多个复制引擎、视频编码器、视频解码器、电源管理单元等(图31中未明确示出)。在至少一个实施例中，I/O单元3106 配置为在PPU 3100的各种逻辑单元之间路由通信。

在至少一个实施例中，由主机处理器执行的程序在缓冲区中对命令流进行编码，该缓冲区将工作负载提供给PPU 3100以进行处理。在至少一个实施例中，工作负载包括指令和要由那些指令处理的数据。在至少一个实施例中，缓冲区是可由主机处理器和PPU3100两者访问(例如，读/写) 的存储器中的区域—主机接口单元可以配置为访问经由I/O单元3106通过系统总线3102传输的存储器请求连接到系统总线3102的系统存储器中的缓冲区。在至少一个实施例中，主机处理器将命令流写入缓冲区，然后将指示命令流开始的指针发送给PPU 3100，使得前端单元3110接收指向一个或更多个命令流指针并管理一个或更多个命令流，从命令流中读取命令并将命令转发到PPU 3100的各个单元。

在至少一个实施例中，前端单元3110耦合到调度器单元3112，该调度器单元3112配置各种GPC 3118以处理由一个或更多个命令流定义的任务。在至少一个实施例中，调度器单元3112配置为跟踪与调度器单元3112 管理的各种任务有关的状态信息，其中状态信息可以指示任务被分配给哪个GPC 3118，任务是活跃的还是非活跃的，与任务相关联的优先级等等。在至少一个实施例中，调度器单元3112管理在一个或更多个GPC 3118上执行的多个任务。

在至少一个实施例中，调度器单元3112耦合到工作分配单元3114，该工作分配单元3114配置为分派任务以在GPC 3118上执行。在至少一个实施例中，工作分配单元3114跟踪从调度器单元3112接收到的多个调度任务并且工作分配单元3114管理每个GPC 3118的待处理任务池和活跃任务池。在至少一个实施例中，待处理任务池包括多个时隙(例如32个时隙)，这些时隙包含分配给要由特定的GPC 3118处理的任务；活跃任务池可包括用于由GPC 3118主动处理的任务的多个时隙(例如4个时隙)，以使随着 GPC 3118中的一个完成任务的执行，该任务将从GPC 3118的活动任务池中逐出，并且从待处理任务池中选择其他任务之一，并安排其在GPC 3118 上执行。在至少一个实施例中，如果活跃任务在GPC 3118上处于空闲状态，例如在等待数据依赖性解决时，则活跃任务从GPC 3118中驱逐并返回到待处理任务池，同时选择了待处理任务池中的另一个任务并调度在GPC 3118 上执行。

在至少一个实施例中，工作分配单元3114经由XBar 3120与一个或更多个GPC3118通信。在至少一个实施例中，XBar 3120是互连网络，其将PPU 3100的许多单元耦合到PPU 3100的其他单元，并且可以配置为将工作分配单元3114耦合到特定的GPC 3118。在至少一个实施例中，一个或更多个PPU 3100的其他单元也可以通过集线器3116连接到XBar3120。

在至少一个实施例中，任务由调度器单元3112管理，并由工作分配单元3114分配给GPC 3118之一。GPC 3118配置为处理任务并产生结果。在至少一个实施例中，结果可以由GPC 3118中的其他任务消耗，通过XBar 3120路由到不同的GPC 3118或存储在存储器3104中。在至少一个实施例中，结果可以通过分区单元3122写到存储器3104中，其实现了用于向存储器3104写入数据或从存储器3104读取数据的存储器接口。在至少一个实施例中，结果可以经由高速GPU互连3108传输到另一PPU 3100或CPU。在至少一个实施例中，PPU 3100包括但不限于U个分区单元3122，其等于耦合到PPU 3100的分离且不同的存储器设备3104的数量。

在至少一个实施例中，主机处理器执行驱动器核心，该驱动器核心实现应用程序编程接口(API)，该应用程序编程接口使在主机处理器上执行的一个或更多个应用程序能够调度操作以在PPU 3100上执行。在一个实施例中，多个计算应用由PPU 3100同时执行，并且PPU 3100为多个计算应用程序提供隔离、服务质量(“QoS”)和独立的地址空间。在至少一个实施例中，应用程序生成指令(例如，以API调用的形式)，该指令使驱动器核心生成一个或更多个任务以供PPU 3100执行，并且驱动器核心将任务输出至由PPU 3100处理的一个或更多个流。在至少一个实施例中，每个任务包括一个或更多个相关线程组，其可以被称为线程束(warp)。在至少一个实施例中，线程束包括可以并行执行的多个相关线程(例如32个线程)。在至少一个实施例中，协作线程可以指代多个线程，包括用于执行任务并且通过共享存储器交换数据的指令。

图32示出了根据至少一个实施例的GPC 3200。在至少一个实施例中， GPC 3200是图31的GPC 3118。在至少一个实施例中，每个GPC 3200包括但不限于用于处理任务的多个硬件单元，并且每个GPC 3200包括但不限于管线管理器3202、预光栅操作单元(“PROP”)3204、光栅引擎3208、工作分配交叉开关(“WDX”)3216、存储器管理单元(“MMU”)3218、一个或更多个数据处理集群(“DPC”)3206，以及部件的任何合适组合。

在至少一个实施例中，GPC 3200的操作由管线管理器3202控制。在至少一个实施例中，管线管理器3202管理一个或更多个DPC 3206的配置，以处理分配给GPC 3200的任务。在至少一个实施例中，管线管理器3202 配置一个或更多个DPC 3206中的至少一个以实现图形渲染管线的至少一部分。在至少一个实施例中，DPC 3206配置为在可编程流式多处理器(“SM”) 3214上执行顶点着色器程序。在至少一个实施例中，管线管理器3202配置为将从工作分配单元接收的数据包路由到GPC 3200内的适当逻辑单元，以及在至少一个实施例中，可以将一些数据包路由到PROP 3204和/或光栅引擎3208中的固定功能硬件单元，而可以将其他数据包路由到DPC 3206以由原始引擎3212或SM 3214进行处理。在至少一个实施例中，管线管理器 3202配置DPC 3206中的至少一个以实现神经网络模型和/或计算管线。在至少一个实施例中，管线管理器3202配置DPC 3206中的至少一个以执行 CUDA程序的至少一部分。

在至少一个实施例中，PROP单元3204配置为将由光栅引擎3208和 DPC 3206生成的数据路由到分区单元中的光栅操作(“ROP”)单元，例如上面结合图31更详细描述的存储器分区单元2522等。在至少一个实施例中，PROP单元3204配置为执行用于颜色混合的优化、组织像素数据、执行地址转换等等。在至少一个实施例中，光栅引擎3208包括但不限于配置为执行各种光栅操作的多个固定功能硬件单元，并且在至少一个实施例中，光栅引擎3208包括但不限于设置引擎、粗光栅引擎、剔除引擎、裁剪引擎、精细光栅引擎、图块聚合引擎及其任意合适的组合。在至少一个实施例中，设置引擎接收变换后的顶点并生成与由顶点定义的几何图元相关联的平面方程；平面方程式被传送到粗光栅引擎以生成基本图元的覆盖信息(例如，图块的x、y覆盖范围掩模)；粗光栅引擎的输出将传输到剔除引擎，在剔除引擎中与z测试失败的图元相关联的片段将被剔除，并传输到剪切引擎，在剪切引擎中剪切位于视锥范围之外的片段。在至少一个实施例中，将经过裁剪和剔除的片段传递给精细光栅引擎，以基于设置引擎生成的平面方程式生成像素片段的属性。在至少一个实施例中，光栅引擎3208的输出包括将由任何适当的实体(例如，由在DPC 3206内实现的片段着色器)处理的片段。

在至少一个实施例中，包括在GPC 3200中的每个DPC 3206包括但不限于M管线控制器(“MPC”)3210；图元引擎3212；一个或更多个SM 3214；及其任何合适的组合。在至少一个实施例中，MPC 3210控制DPC 3206的操作，将从管线管理器3202接收的分组路由到DPC3206中的适当单元。在至少一个实施例中，将与顶点相关联的分组路由到图元引擎3212，图元引擎3212配置为从存储器中获取与顶点关联的顶点属性；相反，可以将与着色器程序相关联的数据包发送到SM 3214。

在至少一个实施例中，SM 3214包括但不限于可编程流式处理器，其配置为处理由多个线程表示的任务。在至少一个实施例中，SM 3214是多线程的并且配置为同时执行来自特定线程组的多个线程(例如32个线程)，并且实现单指令、多数据(“SIMD”)架构，其中将一组线程(例如，线程束)中的每个线程配置为基于相同的指令集来处理不同的数据集。在至少一个实施例中，线程组中的所有线程执行相同的指令。在至少一个实施例中，SM 3214实施单指令、多线程(“SIMT”)架构，其中一组线程中的每个线程配置为基于相同的指令集来处理不同的数据集，但是其中线程组中的各个线程允许在执行期间发散。在至少一个实施例中，为每个线程束维护程序计数器、调用栈和执行状态，从而当线程束中的线程发散时，实现线程束和线程束内的串行执行之间的并发性。在另一个实施例中，为每个单独的线程维护程序计数器、调用栈和执行状态，从而使得在线程束内和线程束之间的所有线程之间具有相等的并发性。在至少一个实施例中，为每个单独的线程维持执行状态，并且可以收敛并并行地执行执行相同指令的线程以提高效率。下面结合图33更详细地描述SM 3214的至少一个实施例。

在至少一个实施例中，MMU 3218在GPC 3200和存储器分区单元(例如，图31的分区单元3122)之间提供接口，并且MMU 3218提供虚拟地址到物理地址的转换、存储器保护以及存储器请求的仲裁。在至少一个实施例中，MMU 3218提供一个或更多个转换后备缓冲区(“TLB”)，用于执行虚拟地址到存储器中的物理地址的转换。

图33示出了根据至少一个实施例的流式多处理器(“SM”)3300。在至少一个实施例中，SM 3300是图32的SM 3214。在至少一个实施例中， SM 3300包括但不限于指令高速缓存3302；一个或更多个调度器单元3304；寄存器文件3308；一个或更多个处理核心(“核心”)3310；一个或更多个特殊功能单元(“SFU”)3312；一个或更多个加载/存储单元(“LSU”)3314；互连网络3316；共享存储器/一级(“L1”)高速缓存3318；及其任何合适的组合。在至少一个实施例中，工作分配单元调度任务以在并行处理单元 (“PPU”)的通用处理集群(“GPC”)上执行，并且每个任务被分配给GPC 内部的特定数据处理集群(“DPC”)，并且如果任务与着色器程序相关联，则将任务分配给SM 3300之一。在至少一个实施例中，调度器单元3304 从工作分配单元接收任务并管理分配给SM 3300的一个或更多个线程块的指令调度。在至少一个实施例中，调度器单元3304调度线程块以作为并行线程的线程束来执行，其中，每个线程块被分配至少一个线程束。在至少一个实施例中，每个线程束执行线程。在至少一个实施例中，调度器单元 3304管理多个不同的线程块，将线程束分配给不同的线程块，然后在每个时钟周期内将来自多个不同的协作组的指令分派给各种功能单元(例如，处理核心3310、SFU 3312和LSU 3314)。

在至少一个实施例中，“合作组”可以指用于组织通信线程组的编程模型，其允许开发人员表达线程正在通信的粒度，从而能够表达更丰富、更有效的并行分解。在至少一个实施例中，协作启动API支持线程块之间的同步以执行并行算法。在至少一个实施例中，常规编程模型的API提供了用于同步协作线程的单一、简单的构造：跨线程块的所有线程的屏障(例如，syncthreads()函数)。但是，在至少一个实施例中，程序员可以在小于线程块粒度的情形下来定义线程组，并在所定义的组内进行同步，以实现更高的性能、设计灵活性以及以集合组范围功能接口的形式实现软件重用。在至少一个实施例中，协作组使程序员能够以子块和多块粒度明确定义线程组，并执行集合操作，例如对协作组中的线程进行同步。在至少一个实施例中，子块粒度与单个线程一样小。在至少一个实施例中，编程模型支持跨软件边界的干净组合，从而库和实用程序功能可以在其本地环境中安全地同步，而不必进行关于收敛的假设。在至少一个实施例中，协作组图元使协作并行的新图案成为可能，包括但不限于生产者-消费者并行，机会主义并行以及整个线程块网格上的全局同步。

在至少一个实施例中，分派单元3306配置为将指令发送到功能单元中的一个或更多个，并且调度器单元3304包括但不限于两个分派单元3306，该两个分派单元3306使得来自相同线程束的两个不同指令能够在每个时钟周期被分派。在至少一个实施例中，每个调度器单元3304包括单个分派单元3306或附加分派单元3306。

在至少一个实施例中，每个SM 3300在至少一个实施例中包括但不限于寄存器文件3308，该寄存器文件3308为SM 3300的功能单元提供了一组寄存器。在至少一个实施例中，寄存器文件3308在每个功能单元之间划分，从而为每个功能单元分配寄存器文件3308的专用部分。在至少一个实施例中，寄存器文件3308在由SM 3300执行的不同线程束之间划分，并且寄存器文件3308为连接到功能单元的数据路径的操作数提供临时存储。在至少一个实施例中，每个SM 3300包括但不限于多个L个处理核心3310。在至少一个实施例中，SM3300包括但不限于大量(例如128个或更多) 不同的处理核心3310。在至少一个实施例中，每个处理核心3310在至少一个实施例中包括但不限于全管线、单精度、双精度和/或混合精度处理单元，其包括但不限于浮点算术逻辑单元和整数算术逻辑单元。在至少一个实施例中，浮点算术逻辑单元实现用于浮点算术的IEEE 754-2008标准。在至少一个实施例中，处理核心3310包括但不限于64个单精度(32位)浮点核心、64个整数核心、32个双精度(64位)浮点核心和8个张量核心。

在至少一个实施例中，张量核心配置为执行矩阵运算。在至少一个实施例中，一个或更多个张量核心包括在处理核心3310中。在至少一个实施例中，张量核心配置为执行深度学习矩阵算术，例如用于神经网络训练和推理的卷积运算。在至少一个实施例中，每个张量核心在4×4矩阵上操作并且执行矩阵乘法和累加运算D＝A×B+C，其中A、B、C和D是4×4矩阵。

在至少一个实施例中，矩阵乘法输入A和B是16位浮点矩阵，并且累加矩阵C和D是16位浮点或32位浮点矩阵。在至少一个实施例中，张量核心对16位浮点输入数据进行32位浮点累加运算。在至少一个实施例中，16位浮点乘法使用64个运算，并得到全精度乘积，然后使用32位浮点加法与其他中间乘积累加起来，以进行4x4x4矩阵乘法。在至少一个实施例中，张量核心用于执行由这些较小的元件构成的更大的二维或更高维度的矩阵运算。在至少一个实施例中，API(诸如CUDA-C++API)公开专门的矩阵加载、矩阵乘法和累加以及矩阵存储操作，以有效地使用来自 CUDA-C++程序的张量核心。在至少一个实施例中，在CUDA级别，线程束级别接口假定跨越所有32个线程束线程的16×16大小的矩阵。

在至少一个实施例中，每个SM 3300包括但不限于执行特殊功能(例如，属性评估、倒数平方根等)的M个SFU 3312。在至少一个实施例中， SFU 3312包括但不限于配置为遍历分层树数据结构的树遍历单元。在至少一个实施例中，SFU 3312包括但不限于配置为执行纹理映射过滤操作的纹理单元。在至少一个实施例中，纹理单元配置为从存储器中加载纹理映射 (例如，纹理像素的2D阵列)和采样纹理映射，以产生采样的纹理值以供由SM 3300执行的着色器程序使用。在至少一个实施例中，将纹理映射存储在共享存储器/L1高速缓存3318中。在至少一个实施例中，纹理单元使用mip映射(mip-maps)(例如，细节级别不同的纹理映射)来实现纹理操作(诸如过滤操作)。在至少一个实施例中，每个SM 3300包括但不限于两个纹理单元。

在至少一个实施例中，每个SM 3300包括但不限于实现共享存储器/ L1高速缓存3318与寄存器文件3308之间的加载和存储操作的N个LSU 3314。在至少一个实施例中，每个SM 3300包括但不限于互连网络3316，互连网络3316将每个功能单元连接到寄存器文件3308，并且LSU 3314连接到寄存器文件3308和共享存储器/L1高速缓存3318。在至少一个实施例中，互连网络3316是交叉开关，其可以配置为将任何功能单元连接到寄存器文件3308中的任何寄存器，并且将LSU 3314连接到寄存器文件3308 和共享存储器/L1高速缓存3318中的存储器位置。

在至少一个实施例中，共享存储器/L1高速缓存3318是片上存储器的阵列，其在至少一个实施例中允许SM 3300与图元引擎之间以及SM 3300中的线程之间的数据存储和通信。在至少一个实施例中，共享存储器/ L1高速缓存3318包括但不限于128KB的存储容量，并且位于从SM 3300 到分区单元的路径中。在至少一个实施例中，共享存储器/L1高速缓存3318 在至少一个实施例中用于高速缓存读取和写入。在至少一个实施例中，共享存储器/L1高速缓存3318、L2高速缓存和存储器中的一个或更多个是后备存储。

在至少一个实施例中，将数据高速缓存和共享存储器功能组合到单个存储器块中，为两种类型的存储器访问提供了改进的性能。在至少一个实施例中，容量由不使用共享存储器的程序使用或将其用作高速缓存，例如如果共享存储器配置为使用一半容量，则纹理和加载/存储操作可以使用剩余容量。根据至少一个实施例，在共享存储器/L1高速缓存3318内的集成使共享存储器/L1高速缓存3318能够用作用于流传输数据的高吞吐量管线，同时提供对频繁重用的数据的高带宽和低延迟访问。在至少一个实施例中，当配置用于通用并行计算时，与图形处理相比，可以使用更简单的配置。在至少一个实施例中，绕过固定功能GPU，从而创建了更加简单的编程模型。在至少一个实施例中，在通用并行计算配置中，工作分配单元直接将线程的块分配和分布给DPC。在至少一个实施例中，块中的线程执行相同的程序，在计算中使用唯一的线程ID以确保每个线程生成唯一的结果，使用SM 3300执行程序并执行计算，使用共享存储器/L1高速缓存3318在线程之间进行通信，以及使用LSU3314通过共享存储器/L1高速缓存3318 和存储器分区单元来读写全局存储器。在至少一个实施例中，当被配置用于通用并行计算时，SM 3300向调度器单元3304写入可以用来在DPC上启动新工作的命令。

在至少一个实施例中，PPU被包括在台式计算机、膝上型计算机、平板电脑、服务器、超级计算机、智能电话(例如，无线、手持设备)、PDA、数码相机、车辆、头戴式显示器、手持式电子设备等中或与之耦合。在至少一个实施例中，PPU被实现在单个半导体衬底上。在至少一个实施例中， PPU与一个或更多个其他设备(例如附加的PPU、存储器、RISCCPU，MMU、数模转换器(“DAC”)等)一起被包括在片上系统(“SoC”)中。

在至少一个实施例中，PPU可以被包括在包括一个或更多个存储设备的图形卡上。图形卡可以配置为与台式计算机主板上的PCIe插槽相连接。在至少一个实施例中，PPU可以是包括在主板的芯片组中的集成GPU (“iGPU”)。

通用计算的软件构造

下图阐述了但不限于用于实现至少一个实施例的示例性软件构造。在至少一个实施例中，下图的一个或更多个软件构造可以实现关于图1-14中的一个或更多个描述的实施例的一个或更多个方面，和/或关于图8-11描述的一种或更多个技术。

图34示出了根据至少一个实施例的编程平台的软件栈。在至少一个实施例中，编程平台是用于利用计算系统上的硬件来加速计算任务的平台。在至少一个实施例中，软件开发人员可以通过库、编译器指令和/或对编程语言的扩展来访问编程平台。在至少一个实施例中，编程平台可以是但不限于CUDA，Radeon开放计算平台(“ROCm”)，OpenCL(由Khronosgroup开发的OpenCL^TM)，SYCL或Intel One API。

在至少一个实施例中，编程平台的软件栈3400为应用程序3401提供执行环境。在至少一个实施例中，应用程序3401可以包括能够在软件栈 3400上启动的任何计算机软件。在至少一个实施例中，应用程序3401可以包括但不限于人工智能(“AI”)/机器学习(“ML”)应用程序，高性能计算(“HPC”)应用程序，虚拟桌面基础架构(“VDI”)或数据中心工作负载。

在至少一个实施例中，应用程序3401和软件栈3400在硬件3407上运行。在至少一个实施例中，硬件3407可以包括一个或更多个GPU，CPU， FPGA，AI引擎和/或支持编程平台的其他类型的计算设备。在至少一个实施例中，例如采用CUDA，软件栈3400可以是厂商专用的，并且仅与来自特定厂商的设备兼容。在至少一个实施例中，例如在采用OpenCL中，软件栈3400可以与来自不同供应商的设备一起使用。在至少一个实施例中，硬件3407包括连接到一个或更多个设备的主机，该设备可经由应用程序编程接口(API)调用被访问以执行计算任务。在至少一个实施例中，与硬件 3407内的主机相比，其可以包括但不限于CPU(但还可以包括计算设备) 及其存储器，硬件3407内的设备可以包括但不限于GPU，FPGA，AI引擎或其他计算设备(但还可以包括CPU)及其存储器。

在至少一个实施例中，编程平台的软件栈3400包括但不限于多个库 3403，运行时(runtime)3405和设备内核驱动器3406。在至少一个实施例中，库3403中的每个库可以包括可以由计算机程序使用并在软件开发期间利用的数据和编程代码。在至少一个实施例中，库3403可以包括但不限于预写的代码和子例程，类，值，类型规范，配置数据，文档，帮助数据和/ 或消息模板。在至少一个实施例中，库3403包括被优化用于在一种或更多种类型的设备上执行的函数。在至少一个实施例中，库3403可以包括但不限于用于在设备上执行数学、深度学习和/或其他类型的运算的函数。在至少一个实施例中，库3403与对应的API3402相关联，API 3402可包括一个或更多个API，其暴露在库3403中实现的函数。

在至少一个实施例中，将应用程序3401编写为源代码，该源代码被编译成可执行代码，如下面结合图39-41更详细讨论的。在至少一个实施例中，应用程序3401的可执行代码可以至少部分地在由软件栈3400提供的执行环境上运行。在至少一个实施例中，在应用程序3401的执行期间，可以得到需要在设备(与主机相比)上运行的代码。在这种情况下，在至少一个实施例中，可以调用运行时3405以在设备上加载和启动必需的代码。在至少一个实施例中，运行时3405可以包括能够支持应用程序3401的执行的任何技术上可行的运行时系统。

在至少一个实施例中，运行时3405被实现为与对应的API(其被示为API 3404)相关联的一个或更多个运行时库。在至少一个实施例中，一个或更多个这样的运行时库可以包括但不限于用于存储器管理，执行控制，设备管理，错误处理和/或同步等等的函数。在至少一个实施例中，存储器管理函数可以包括但不限于用于分配、解除分配和复制设备存储器以及在主机存储器和设备存储器之间传输数据的函数。在至少一个实施例中，执行控制函数可以包括但不限于在设备上启动函数(当函数是可从主机调用的全局函数时，有时称为“内核”)的函数，和用于在运行时库为要在设备上执行的给定函数维护的缓冲区中设置属性值的函数。

在至少一个实施例中，可以任何技术上可行的方式来实现运行时库和相应的API3404。在至少一个实施例中，一个(或任意数量的)API可以公开用于设备的细粒度控制的低级函数集，而另一(或任意数量的)API 可以公开这样的较高级的函数集。在至少一个实施例中，可以在低级API 之上构建高级运行时API。在至少一个实施例中，一个或更多个运行时API 可以是在与语言无关的运行时API之上分层的特定于语言的API。

在至少一个实施例中，设备内核驱动器3406被配置为促进与底层设备的通信。在至少一个实施例中，设备内核驱动器3406可以提供诸如API 3404之类的API和/或其他软件所依赖的低级函数。在至少一个实施例中，设备内核驱动器3406可以被配置为在运行时将中间表示(“IR”)代码编译成二进制代码。在至少一个实施例中，对于CUDA，设备内核驱动器3406 可以在运行时将非硬件专用的并行线程执行(“PTX”)IR代码编译为用于特定目标设备的二进制代码(高速缓存已编译的二进制代码)，其有时也称为“最终”代码。在至少一个实施例中，这样做可以允许最终代码在目标设备上运行，而当源代码最初被编译为PTX代码时，该目标设备可能不存在。备选地，在至少一个实施例中，设备源代码可以离线地编译成二进制代码，而不需要设备内核驱动器3406在运行时编译IR代码。

图35示出了根据至少一个实施例的图34的软件栈3400的CUDA实现。在至少一个实施例中，可在其上启动应用程序3501的CUDA软件栈 3500包括CUDA库3503，CUDA运行时3505，CUDA驱动器3507和设备内核驱动器3508。在至少一个实施例中，CUDA软件栈3500在硬件3509 上执行，该硬件3509可以包括支持CUDA的GPU，其由加利福尼亚州圣克拉拉市的NVIDIA公司开发。

在至少一个实施例中，应用程序3501、CUDA运行时3505和设备内核驱动器3508可以分别执行与应用程序3401、运行时3405和设备内核驱动器3406类似的功能，以上结合图34对其进行了描述。在至少一个实施例中，CUDA驱动器3507包括实现CUDA驱动器API 3506的库(libcuda.so)。在至少一个实施例中，类似于由CUDA运行时库(cudart)实现的CUDA 运行时API 3504，CUDA驱动器API 3506可以公开但不限于用于存储器管理、执行控制、设备管理、错误处理、同步和/或图形互操作性等的函数。在至少一个实施例中，CUDA驱动器API3506与CUDA运行时API 3504 的不同之处在于，CUDA运行时API 3504通过提供隐式初始化、上下文(类似于进程)管理和模块(类似于动态加载的库)管理来简化设备代码管理。与高级CUDA运行时API 3504相反，在至少一个实施例中，CUDA驱动器API 3506是提供对设备的更细粒度控制的低级API，特别是关于上下文和模块加载。在至少一个实施例中，CUDA驱动器API 3506可以公开没有由CUDA运行时API 3504公开的用于上下文管理的函数。在至少一个实施例中，CUDA驱动器API 3506也与语言无关，并且除了支持CUDA运行时API3504之外，还支持例如OpenCL。此外，在至少一个实施例中，包括CUDA运行时3505在内的开发库可被视为与驱动器组件分离，包括用户模式的CUDA驱动器3507和内核模式的设备驱动器3508(有时也称为“显示”驱动器)。

在至少一个实施例中，CUDA库3503可以包括但不限于数学库，深度学习库，并行算法库和/或信号/图像/视频处理库，并行计算应用程序(例如应用程序3501)可以利用这些库。在至少一个实施例中，CUDA库3503 可包括数学库，例如cuBLAS库，其是用于执行线性代数运算的基本线性代数子程序(“BLAS”)的实现；用于计算快速傅立叶变换(“FFT”) 的cuFFT库，以及用于生成随机数的cuRAND库等。在至少一个实施例中， CUDA库3503可以包括深度学习库，诸如用于深度神经网络的基元的 cuDNN库和用于高性能深度学习推理的TensorRT平台等等。

图36示出了根据至少一个实施例的图34的软件栈3400的ROCm实现。在至少一个实施例中，可在其上启动应用程序3601的ROCm软件栈 3600包括语言运行时3603，系统运行时3605，thunk 3607，ROCm内核驱动器3608和设备内核驱动器。在至少一个实施例中，ROCm软件栈3600 在硬件3609上执行，硬件3609可以包括支持ROCm的GPU，其由加利福尼亚州圣克拉拉市的AMD公司开发。

在至少一个实施例中，应用程序3601可以执行与以上结合图34讨论的应用程序3401类似的功能。另外，在至少一个实施例中，语言运行时3603 和系统运行时3605可以执行与以上结合图34讨论的运行时3405类似的功能。在至少一个实施例中，语言运行时3603和系统运行时3605的不同之处在于，系统运行时3605是实现ROCr系统运行时API 3604并利用异构系统架构(“HSA”)运行时API的语言无关运行时。在至少一个实施例中，HSA运行时API是一种瘦用户模式API，它公开接口以供访问和与AMD GPU交互，包括用于存储器管理、通过架构分派内核的执行控制、错误处理、系统和代理信息以及运行时初始化和关闭等的函数。在至少一个实施例中，与系统运行时3605相比，语言运行时3603是ROCr系统运行时API 3604之上分层的特定于语言的运行时API 3602的实现。在至少一个实施例中，语言运行时API可以包括但不限于可移植异构计算接口(“HIP”)语言运行时API，异构计算编译器(“HCC”)语言运行时API或OpenCL API 等等。特别是，HIP语言是C++编程语言的扩展，具有CUDA机制的功能相似版本，并且在至少一个实施例中，HIP语言运行时API包括与以上结合图35讨论的CUDA运行时API 3504相似的函数，例如用于存储器管理、执行控制、设备管理、错误处理和同步等的函数。

在至少一个实施例中，thunk(ROCt)3607是可用于与底层ROCm驱动器3608交互的接口。在至少一个实施例中，ROCm驱动器3608是ROCk 驱动器，其是AMDGPU驱动器和HSA内核驱动器(amdkfd)的组合。在至少一个实施例中，AMDGPU驱动器是由AMD开发的用于GPU的设备内核驱动器，其执行与以上结合图34讨论的设备内核驱动器3406类似的功能。在至少一个实施例中，HSA内核驱动器是允许不同类型的处理器经由硬件特征更有效地共享系统资源的驱动器。

在至少一个实施例中，各种库(未示出)可以被包括在语言运行时3603 上方的ROCm软件栈3600中，并且提供与以上结合图35讨论的CUDA库3503相似的功能。在至少一个实施例中，各种库可以包括但不限于数学、深度学习和/或其他库，例如实现与CUDA cuBLAS类似的函数的hipBLAS 库，类似于CUDA cuFFT用于计算FFT的rocFFT库等。

图37示出了根据至少一个实施例的图34的软件栈3400的OpenCL 实现。在至少一个实施例中，可以在其上启动应用程序3701的OpenCL软件栈3700包括OpenCL框架3705，OpenCL运行时3706和驱动器3707。在至少一个实施例中，OpenCL软件栈3700在不是特定于供应商的硬件 3509上执行。在至少一个实施例中，由于由不同厂商开发的设备支持OpenCL，因此可能需要特定的OpenCL驱动器才能与来自此类厂商的硬件进行互操作。

在至少一个实施例中，应用程序3701，OpenCL运行时3706，设备内核驱动器3707和硬件3708可以分别执行与上面结合图34讨论的应用程序3401、运行时3405、设备内核驱动器3406和硬件3407类似的功能。在至少一个实施例中，应用程序3701还包括具有将在设备上执行的代码的 OpenCL内核3702。

在至少一个实施例中，OpenCL定义了一种“平台”，其允许主机控制连接到该主机的设备。在至少一个实施例中，OpenCL框架提供平台层 API和运行时API，示出为平台API3703和运行时API 3705。在至少一个实施例中，运行时API 3705使用上下文来管理设备上内核的执行。在至少一个实施例中，每个标识的设备可以与各自的上下文相关联，运行时API 3705可以使用该上下文来管理该设备的命令队列、程序对象和内核对象、共享存储器对象等。在至少一个实施例中，平台API 3703公开了允许设备上下文用于选择和初始化设备，经由命令队列将工作提交给设备，以及使得能够进行来自和去往设备的数据传输等的函数。另外，在至少一个实施例中，OpenCL框架提供各种内置函数(未示出)，包括数学函数、关系函数和图像处理函数等。

在至少一个实施例中，编译器3704也被包括在OpenCL框架3710中。在至少一个实施例中，源代码可以在执行应用程序之前被离线编译或者在执行应用程序期间被在线编译。与CUDA和ROCm相反，至少一个实施例中的OpenCL应用程序可以由编译器3704在线编译，编译器3704被包括以代表可以用于将源代码和/或IR代码(例如标准可移植中间表示(“SPIR-V”)代码)编译为二进制代码的任意数量的编译器。可替代地，在至少一个实施例中，可以在执行这样的应用程序之前离线编译OpenCL 应用程序。

图38示出了根据至少一个实施例的由编程平台支持的软件。在至少一个实施例中，编程平台3804被配置为支持应用程序3800可以依赖的各种编程模型3803，中间件和/或库3802以及框架3801。在至少一个实施例中，应用程序3800可以是使用例如深度学习框架(例如，MXNet，PyTorch 或TensorFlow)实现的AI/ML应用，其可以依赖于诸如cuDNN，NVIDIACollective Communications Library(“NCCL”)”和/或NVIDIA开发人员数据加载库(“DALI”)CUDA库之类的库，以在底层硬件上提供加速的计算。

在至少一个实施例中，编程平台3804可以是以上分别结合图35、图 36和图37描述的CUDA、ROCm或OpenCL平台之一。在至少一个实施例中，编程平台3804支持多个编程模型3803，其是底层计算系统的抽象，其允许算法和数据结构的表达。在至少一个实施例中，编程模型3803可以暴露底层硬件的特征以便改善性能。在至少一个实施例中，编程模型3803可以包括但不限于CUDA，HIP，OpenCL，C++加速大规模并行性(“C++ AMP”)，开放多处理(“OpenMP”)，开放加速器(“OpenACC”) 和/或Vulcan计算(Vulcan Compute)。

在至少一个实施例中，库和/或中间件3802提供编程模型3804的抽象的实现。在至少一个实施例中，这样的库包括可由计算机程序使用并在软件开发期间利用的数据和编程代码。在至少一个实施例中，除了可以从编程平台3804获得的那些之外，这样的中间件还包括向应用程序提供服务的软件。在至少一个实施例中，库和/或中间件3802可以包括但不限于 cuBLAS、cuFFT、cuRAND和其他CUDA库，或rocBLAS、rocFFT、rocRAND 和其他ROCm库。另外，在至少一个实施例中，库和/或中间件3802可以包括NCCL和ROCm通信集合库(“RCCL”)库，其提供用于GPU的通信例程，用于深度学习加速的MIOpen库和/或用于线性代数、矩阵和向量运算、几何变换、数值求解器以及相关算法的本征库。

在至少一个实施例中，应用程序框架3801依赖于库和/或中间件3802。在至少一个实施例中，每个应用程序框架3801是用于实现应用软件的标准结构的软件框架。回到上面讨论的AI/ML示例，在至少一个实施例中，可以使用框架(诸如Caffe，Caffe2，TensorFlow，Keras，PyTorch或MxNet 深度学习框架)来实现AI/ML应用。

图39示出了根据至少一个实施例的编译代码以在图34-37的编程平台之一上执行。在至少一个实施例中，编译器3901接收源代码3900，其包括主机代码以及设备代码两者。在至少一个实施例中，编译器3901被配置为将源代码3900转换为用于在主机上执行的主机可执行代码3902以及用于在设备上执行的设备可执行代码3903。在至少一个实施例中，源代码3900 可以在执行应用程序之前离线编译，或者在执行应用程序期间在线编译。

在至少一个实施例中，源代码3900可以包括编译器3901支持的任何编程语言的代码，例如C++、C、Fortran等。在至少一个实施例中，源代码3900可以包括在单源(single-source)文件中，其具有主机代码和设备代码的混合，并在其中指示了设备代码的位置。在至少一个实施例中，单源文件可以是包括CUDA代码的.cu文件或包括HIP代码的.hip.cpp文件。备选地，在至少一个实施例中，源代码3900可以包括多个源代码文件，而不是单源文件，在该单源文件中主机代码和设备代码是分开的。

在至少一个实施例中，编译器3901被配置为将源代码3900编译成用于在主机上执行的主机可执行代码3902和用于在设备上执行的设备可执行代码3903。在至少一个实施例中，编译器3901执行操作，包括将源代码 3900解析为抽象系统树(AST)，执行优化以及生成可执行代码。在源代码3900包括单源文件的至少一个实施例中，编译器3901可以将设备代码与主机代码在这种单源文件中分开，将设备代码和主机代码分别编译成设备可执行代码3903和主机可执行代码3902，以及将设备可执行代码3903 和主机可执行代码3902在单个文件中链接到一起，如下面关于图40更详细讨论的。

在至少一个实施例中，主机可执行代码3902和设备可执行代码3903 可以是任何合适的格式，例如二进制代码和/或IR代码。在CUDA的情况下，在至少一个实施例中，主机可执行代码3902可以包括本地对象代码，而设备可执行代码3903可以包括PTX中间表示的代码。在至少一个实施例中，在ROCm的情况下，主机可执行代码3902和设备可执行代码3903都可以包括目标二进制代码。

图40是根据至少一个实施例的编译代码以在图34-37的编程平台之一上执行的更详细图示。在至少一个实施例中，编译器4001被配置为接收源代码4000，编译源代码4000，并输出可执行文件4010。在至少一个实施例中，源代码4000是单源文件，例如.cu文件，.hip.cpp文件或其他格式的文件，其包括主机代码和设备代码两者。在至少一个实施例中，编译器4001 可以是但不限于用于在.cu文件中编译CUDA代码的NVIDIA CUDA编译器(“NVCC”)，或用于在.hip.cpp文件中编译HIP代码的HCC编译器。

在至少一个实施例中，编译器4001包括编译器前端4002，主机编译器4005，设备编译器4006和链接器4009。在至少一个实施例中，编译器前端4002被配置为在源代码4000中将设备代码4004与主机代码4003分开。在至少一个实施例中，设备代码4004由设备编译器4006编译成设备可执行代码4008，如所描述的，其可以包括二进制代码或IR代码。在至少一个实施例中，主机代码4003由主机编译器4005单独地编译成主机可执行代码4007。在至少一个实施例中，对于NVCC，主机编译器4005可以是但不限于输出本机目标代码的通用C/C++编译器，而设备编译器4006 可以是但不限于基于低级虚拟机(“LLVM”)的编译器，其将LLVM编译器基础架构分叉，并输出PTX代码或二进制代码。在至少一个实施例中，对于HCC，主机编译器4005和设备编译器4006两者可以是但不限于输出目标二进制代码的基于LLVM的编译器。

在至少一个实施例中，在将源代码4000编译成主机可执行代码4007 和设备可执行代码4008之后，链接器4009将主机和设备可执行代码4007 和4008在可执行文件4010中链接到一起。在至少一个实施例中，主机和 PTX的本机目标代码或设备的二进制代码可以在可执行和可链接格式 (“ELF”)文件中链接在一起，该文件是用于存储目标代码的容器格式。

图41示出了根据至少一个实施例的在编译源代码之前转换源代码。在至少一个实施例中，源代码4100通过转换工具4101传递，转换工具4101 将源代码4100转换成转换后的源代码4102。在至少一个实施例中，编译器 4103用于将转换后的源代码4102编译成主机可执行代码4104和设备可执行代码3405，其过程类似于由编译器3901将源代码3900编译成主机可执行代码3902和设备可执行代码3903的过程，如以上结合图39所讨论的。

在至少一个实施例中，由转换工具4101执行的转换被用于移植(port) 源代码4100，以在与最初打算在其上运行的不同的环境中执行。在至少一个实施例中，转换工具4101可以包括但不限于HIP转换器，其用于将用于 CUDA平台的CUDA代码“移植(hipify)”为可以在ROCm平台上编译和执行的HIP代码。在至少一个实施例中，源代码4100的转换可以包括：解析源代码4100，并将对由一个编程模型(例如，CUDA)提供的API的调用转换为对由另一编程模型(例如，例如，HIP)提供的API的相应调用，如下面结合图42A-43更详细地讨论的。返回到移植CUDA代码的示例，在至少一个实施例中，对CUDA运行时API、CUDA驱动器API和/或CUDA 库的调用可以被转换为对应的HIP API调用。在至少一个实施例中，由转换工具4101执行的自动转换有时可能是不完整的，需要额外的人工来完全移植源代码4100。

配置用于通用计算的GPU

下列附图阐述了但不限于用于根据至少一个实施例编译和执行计算源代码的示例性架构。在至少一个实施例中，下图的一个或更多个架构可以实现关于图1-14中的一个或更多个描述的实施例的一个或更多个方面，和/或关于图8-11描述的一个或更多个技术。

图42A示出了根据至少一个实施例的被配置为使用不同类型的处理单元来编译和执行CUDA源代码4210的系统4200。在至少一个实施例中，系统4200包括但不限于CUDA源代码4210，CUDA编译器4250，主机可执行代码4270(1)，主机可执行代码4270(2)，CUDA设备可执行代码 4284，CPU 4290，启用CUDA的GPU 4294，GPU 4292，CUDA到HIP转换工具4220，HIP源代码4230，HIP编译器驱动器4240，HCC 4260和HCC 设备可执行代码4282。

在至少一个实施例中，CUDA源代码4210是CUDA编程语言的人类可读代码的集合。在至少一个实施例中，CUDA代码是CUDA编程语言的人类可读代码。在至少一个实施例中，CUDA编程语言是C++编程语言的扩展，其包括但不限于定义设备代码以及区分设备代码和主机代码的机制。在至少一个实施例中，设备代码是在编译之后可在设备上并行执行的源代码。在至少一个实施例中，设备可以是针对并行指令处理而优化的处理器，例如启用CUDA的GPU 4290、GPU 4292或另一GPGPU等。在至少一个实施例中，主机代码是在编译后可以在主机上执行的源代码。在至少一个实施例中，主机是针对顺序指令处理而优化的处理器，例如CPU 4290。

在至少一个实施例中，CUDA源代码4210包括但不限于，任意数量 (包括零)的全局函数4212，任意数量(包括零)的设备函数4214，任意数量(包括零)的主机函数4216，以及任意数量(包括零)的主机/设备函数4218。在至少一个实施例中，全局函数4212，设备函数4214，主机函数 4216和主机/设备函数4218在CUDA源代码4210中可以混合。在至少一个实施例中，每个全局函数4212可在设备上执行并且可从主机调用。因此，在至少一个实施例中，全局函数4212中的一个或更多个可以充当设备的入口点。在至少一个实施例中，每个全局函数4212是内核。在至少一个实施例中以及在一种称为动态并行性的技术中，一个或更多个全局函数4212定义了一内核，该内核可以在设备上执行并且可以从这样的设备调用。在至少一个实施例中，内核在执行期间由设备上的N个不同线程并行执行N次 (其中N为任何正整数)。

在至少一个实施例中，每个设备函数4214在设备上执行并且只能从这样的设备调用。在至少一个实施例中，每个主机函数4216在主机上执行并且只能从这样的主机调用。在至少一个实施例中，每个主机/设备函数 4216既定义了在主机上可执行并且只能从这样的主机调用的函数的主机版本，也定义了在设备上可执行并且只能从这样的设备调用的函数的设备版本。

在至少一个实施例中，CUDA源代码4210还可包括但不限于对通过 CUDA运行时API4202定义的任意数量的函数的任意数量的调用。在至少一个实施例中，CUDA运行时API4202可以包括但不限于在主机上执行的任意数量的函数，用于分配和解除分配设备存储器，在主机存储器和设备存储器之间传输数据，管理具有多个设备的系统等。在至少一个实施例中， CUDA源代码4210还可以包括对在任意数量的其他CUDA API中指定的任意数量的函数的任意数量的调用。在至少一个实施例中，CUDA API可以是被设计为由CUDA代码使用的任何API。在至少一个实施例中，CUDA API包括但不限于CUDA运行时API 4202，CUDA驱动器API，用于任意数量的CUDA库的API等。在至少一个实施例中并且相对于CUDA运行时API4202，CUDA驱动器API是较低级别的API，但可以提供对设备的更细粒度的控制。在至少一个实施例中，CUDA库的示例包括但不限于 cuBLAS，cuFFT，cuRAND，cuDNN等。

在至少一个实施例中，CUDA编译器4250编译输入的CUDA代码(例如，CUDA源代码4210)以生成主机可执行代码4270(1)和CUDA设备可执行代码4284。在至少一个实施例中，CUDA编译器4250是NVCC。在至少一个实施例中，主机可执行代码4270(1)是在CPU 4290上可执行的输入源代码中包括的主机代码的编译版本。在至少一个实施例中，CPU 4290可以是针对顺序指令处理而优化的任何处理器。

在至少一个实施例中，CUDA设备可执行代码4284是在启用CUDA 的GPU 4294上可执行的输入源代码中包括的设备代码的编译版本。在至少一个实施例中，CUDA设备可执行代码4284包括但不限于二进制代码。在至少一个实施例中，CUDA设备可执行代码4284包括但不限于IR代码，例如PTX代码，该IR代码在运行时被设备驱动器进一步编译为用于特定目标设备(例如，启用CUDA的GPU 4294)的二进制代码。在至少一个实施例中，启用CUDA的GPU4294可以是针对并行指令处理而优化并且支持CUDA的任何处理器。在至少一个实施例中，启用CUDA的GPU 4294 由加利福尼亚州圣克拉拉市的NVIDIA公司开发。

在至少一个实施例中，CUDA到HIP转换工具4220被配置为将CUDA 源代码4210转换成功能上相似的HIP源代码4230。在至少一个实施例中，HIP源代码4230是HIP编程语言的人类可读代码的集合。在至少一个实施例中，HIP代码是HIP编程语言的人类可读代码。在至少一个实施例中， HIP编程语言是C++编程语言的扩展，其包括但不限于CUDA机制的功能上相似的版本，用于定义设备代码并区分设备代码和主机代码。在至少一个实施例中，HIP编程语言可以包括CUDA编程语言的功能的子集。在至少一个实施例中，例如，HIP编程语言包括但不限于定义全局函数4212的机制，但是这样的HIP编程语言可能缺乏对动态并行性的支持，因此，在 HIP代码中定义的全局函数4212仅可从主机调用。

在至少一个实施例中，HIP源代码4230包括但不限于任意数量(包括零)的全局函数4212，任意数量(包括零)的设备函数4214，任意数量 (包括零)的主机函数4216以及任意数量(包括零)的主机/设备函数4218。在至少一个实施例中，HIP源代码4230还可以包括对在HIP运行时API 4232中指定的任意数量的函数的任意数量的调用。在一个实施例中，HIP运行时API 4232包括但不限于CUDA运行时API 4202中包括的函数的子集的功能上相似的版本。在至少一个实施例中，HIP源代码4230还可以包括对在任意数量的其他HIP API中指定的任意数量的函数的任意数量的调用。在至少一个实施例中，HIP API可以是被设计为供HIP代码和/或ROCm 使用的任何API。在至少一个实施例中，HIP API包括但不限于HIP运行时 API 4232，HIP驱动器API，用于任意数量的HIP库的API，用于任意数量的ROCm库的API等。

在至少一个实施例中，CUDA到HIP转换工具4220将CUDA代码中的每个内核调用从CUDA语法转换为HIP语法，并将CUDA代码中的任意数量的其他CUDA调用转换为任意数量的其他功能上相似的HIP调用。在至少一个实施例中，CUDA调用是对在CUDA API中指定的函数的调用，并且HIP调用是对在HIP API中指定的函数的调用。在至少一个实施例中， CUDA到HIP转换工具4220将对在CUDA运行时API 4202中指定的函数的任意数量的调用转换为对在HIP运行时API 4232中指定的函数的任意数量的调用。

在至少一个实施例中，CUDA到HIP转换工具4220是被称为 hipify-perl的工具，其执行基于文本的转换过程。在至少一个实施例中， CUDA到HIP转换工具4220是被称为hipify-clang的工具，相对于hipify-perl，其执行更复杂且更鲁棒的转换过程，该过程涉及使用clang(编译器前端) 解析CUDA代码，然后转换得到的符号。在至少一个实施例中，除了由 CUDA到HIP转换工具4220执行的那些修改之外，将CUDA代码正确地转换成HIP代码可能还需要修改(例如，手动编辑)。

在至少一个实施例中，HIP编译器驱动器4240是确定目标设备4246，然后配置与目标设备4246兼容的编译器以编译HIP源代码4230的前端。在至少一个实施例中，目标设备4246是针对并行指令处理而优化的处理器。在至少一个实施例中，HIP编译器驱动器4240可以以任何技术上可行的方式确定目标设备4246。

在至少一个实施例中，如果目标设备4246与CUDA兼容(例如，启用CUDA的GPU4294)，则HIP编译器驱动器4240生成HIP/NVCC编译命令4242。在至少一个实施例中并且结合图42B更详细地描述的， HIP/NVCC编译命令4242配置CUDA编译器4250以使用但不限于HIP到 CUDA转换头和CUDA运行时库来编译HIP源代码4230。在至少一个实施例中并且响应于HIP/NVCC编译命令4242，CUDA编译器4250生成主机可执行代码4270(1)和CUDA设备可执行代码4284。

在至少一个实施例中，如果目标设备4246与CUDA不兼容，则HIP 编译器驱动器4240生成HIP/HCC编译命令4244。在至少一个实施例中并且如结合图42C更详细地描述的，HIP/HCC编译命令4244配置HCC 4260 以使用HCC头和HIP/HCC运行时库编译HIP源代码4230。在至少一个实施例中并且响应于HIP/HCC编译命令4244，HCC 4260生成主机可执行代码4270(2)和HCC设备可执行代码4282。在至少一个实施例中，HCC 设备可执行代码4282是HIP源代码4230中包含的可在GPU 4292上执行的设备代码的编译版本。在至少一个实施例中，GPU 4292可以是针对并行指令处理而优化的、与CUDA不兼容且与HCC兼容的任何处理器。在至少一个实施例中，GPU 4292由加利福尼亚州圣克拉拉市的AMD公司开发。在至少一个实施例中，GPU 4292是不启用CUDA的GPU 4292。

仅出于说明性目的，在图42A中描绘了在至少一个实施例中可以实现为编译CUDA源代码4210以在CPU 4290和不同设备上执行的三个不同流程。在至少一个实施例中，直接CUDA流程编译CUDA源代码4210以在CPU 4290和启用CUDA的GPU 4294上执行，而无需将CUDA源代码4210 转换为HIP源代码4230。在至少一个实施例中，间接CUDA流程将CUDA 源代码4210转换为HIP源代码4230，然后编译HIP源代码4230以在CPU 4290和启用CUDA的GPU 4294上执行。在至少一个实施例中，CUDA/HCC 流程将CUDA源代码4210转换为HIP源代码4230，然后编译HIP源代码 4230以在CPU 4290和GPU 4292上执行。

可以通过虚线和一系列气泡注释A1-A3描绘可以在至少一个实施例中实现的直接CUDA流程。在至少一个实施例中，并且如气泡注释A1所示，CUDA编译器4250接收CUDA源代码4210和配置CUDA编译器4250 以编译CUDA源代码4210的CUDA编译命令4248。在至少一个实施例中，直接CUDA流程中使用的CUDA源代码4210是用CUDA编程语言编写的，该CUDA编程语言基于除C++之外的其他编程语言(例如C，Fortran， Python，Java等)。在至少一个实施例中，并且响应于CUDA编译命令4248， CUDA编译器4250生成主机可执行代码4270(1)和CUDA设备可执行代码4284(用气泡注释A2表示)。在至少一个实施例中并且如用气泡注释 A3所示，主机可执行代码4270(1)和CUDA设备可执行代码4284可以分别在CPU 4290和启用CUDA的GPU 4294上执行。在至少一个实施例中，CUDA设备可执行代码4284包括但不限于二进制代码。在至少一个实施例中，CUDA设备可执行代码4284包括但不限于PTX代码，并且在运行时被进一步编译成用于特定目标设备的二进制代码。

可以通过虚线和一系列气泡注释B1-B6来描述可以在至少一个实施例中实现的间接CUDA流程。在至少一个实施例中并且如气泡注释B1所示，CUDA到HIP转换工具4220接收CUDA源代码4210。在至少一个实施例中并且如气泡注释B2所示，CUDA到HIP转换工具4220将CUDA源代码4210转换为HIP源代码4230。在至少一个实施例中并如气泡注释B3 所示，HIP编译器驱动器4240接收HIP源代码4230，并确定目标设备4246 是否启用了CUDA。

在至少一个实施例中并且如气泡注释B4所示，HIP编译器驱动器 4240生成HIP/NVCC编译命令4242，并将HIP/NVCC编译命令4242和 HIP源代码4230两者都发送到CUDA编译器4250。在至少一个实施例中并且如结合图42B更详细地描述的，HIP/NVCC编译命令4242配置CUDA 编译器4250以使用但不限于HIP到CUDA转换头和CUDA运行时库来编译HIP源代码4230。在至少一个实施例中并且响应于HIP/NVCC编译命令 4242，CUDA编译器4250生成主机可执行代码4270(1)和CUDA设备可执行代码4284(用气泡注释B5表示)。在至少一个实施例中并且如气泡注释B6所示，主机可执行代码4270(1)和CUDA设备可执行代码4284 可以分别在CPU 4290和启用CUDA的GPU 4294上执行。在至少一个实施例中，CUDA设备可执行代码4284包括但不限于二进制代码。在至少一个实施例中，CUDA设备可执行代码4284包括但不限于PTX代码，并且在运行时被进一步编译成用于特定目标设备的二进制代码。

可以通过实线和一系列气泡注释C1-C6来描述可以在至少一个实施例中实现的CUDA/HCC流程。在至少一个实施例中并且如气泡注释C1所示，CUDA到HIP转换工具4220接收CUDA源代码4210。在至少一个实施例中并且如气泡注释C2所示，CUDA到HIP转换工具4220将CUDA源代码4210转换为HIP源代码4230。在至少一个实施例中并且如气泡注释 C3所示，HIP编译器驱动器4240接收HIP源代码4230，并确定目标设备 4246未启用CUDA。

在至少一个实施例中，HIP编译器驱动器4240生成HIP/HCC编译命令4244，并且将HIP/HCC编译命令4244和HIP源代码4230两者发送到 HCC 4260(用气泡注释C4表示)。在至少一个实施例中并且如结合图42C 更详细地描述的，HIP/HCC编译命令4244配置HCC 4260以使用但不限于 HCC头和HIP/HCC运行时库编译HIP源代码4230。在至少一个实施例中并且响应于HIP/HCC编译命令4244，HCC 4260生成主机可执行代码4270 (2)和HCC设备可执行代码4282(用气泡注释C5表示)。在至少一个实施例中并且如气泡注释C6所示，主机可执行代码4270(2)和HCC设备可执行代码4282可以分别在CPU 4290和GPU 4292上执行。

在至少一个实施例中，在将CUDA源代码4210转换为HIP源代码 4230之后，HIP编译器驱动器4240可随后用于生成用于启用CUDA的GPU 4294或GPU 4292的可执行代码，而无需将CUDA重新执行为HIP转换工具4220。在至少一个实施例中，CUDA到HIP转换工具4220将CUDA源代码4210转换为HIP源代码4230，然后将其存储在存储器中。在至少一个实施例中，HIP编译器驱动器4240然后配置HCC 4260以基于HIP源代码4230生成主机可执行代码4270(2)和HCC设备可执行代码4282。在至少一个实施例中，HIP编译器驱动器4240随后配置CUDA编译器4250 以基于存储的HIP源代码4230生成主机可执行代码4270(1)和CUDA设备可执行代码4284。

图42B示出了根据至少一个实施例的被配置为使用CPU 4290和启用CUDA的GPU4294来编译和执行图42A的CUDA源代码4210的系统4204。在至少一个实施例中，系统4204包括但不限于CUDA源代码4210，CUDA 到HIP转换工具4220，HIP源代码4230，HIP编译器驱动器4240，CUDA 编译器4250，主机可执行代码4270(1)，CUDA设备可执行代码4284， CPU 4290和启用CUDA的GPU 4294。

在至少一个实施例中并且如本文先前结合图42A所描述的，CUDA 源代码4210包括但不限于任意数量(包括零)的全局函数4212，任意数量 (包括零)的设备函数4214，任意数量(包括零)的主机函数4216以及任意数量(包括零)的主机/设备函数4218。在至少一个实施例中，CUDA源代码4210还包括但不限于对在任意数量的CUDA API中指定的任意数量的函数的任意数量的调用。

在至少一个实施例中，CUDA到HIP转换工具4220将CUDA源代码 4210转换成HIP源代码4230。在至少一个实施例中，CUDA到HIP转换工具4220将CUDA源代码4210中的每个内核调用从CUDA语法转换为 HIP语法，并将CUDA源代码4210中任意数量的其他CUDA调用转换为任意数量的其他功能上相似的HIP调用。

在至少一个实施例中，HIP编译器驱动器4240确定目标设备4246是启用CUDA的，并且生成HIP/NVCC编译命令4242。在至少一个实施例中，然后HIP编译器驱动器4240经由HIP/NVCC编译命令4242配置CUDA编译器4250以编译HIP源代码4230。在至少一个实施例中，作为配置CUDA 编译器4250的一部分，HIP编译器驱动器4240提供对HIP到CUDA转换头4252的访问。在至少一个实施例中，HIP到CUDA转换头4252将任意数量的HIP API中指定的任意数量的机制(例如，函数)转换为任意数量的CUDA API中指定的任意数量的机制。在至少一个实施例中，CUDA编译器4250将HIP到CUDA转换头4252与对应于CUDA运行时API 4202 的CUDA运行时库4254结合使用，以生成主机可执行代码4270(1)和 CUDA设备可执行代码4284。在至少一个实施例中，然后可以分别在CPU 4290和启用CUDA的GPU 4294上执行主机可执行代码4270(1)和CUDA 设备可执行代码4284。在至少一个实施例中，CUDA设备可执行代码4284 包括但不限于二进制代码。在至少一个实施例中，CUDA设备可执行代码 4284包括但不限于PTX代码，并且在运行时被进一步编译成用于特定目标设备的二进制代码。

图42C示出了根据至少一个实施例的系统4206，该系统4206被配置为使用CPU4290和未启用CUDA的GPU 4292来编译和执行图42A的 CUDA源代码4210。在至少一个实施例中，系统4206包括但不限于CUDA 源代码4210，CUDA到HIP转换工具4220，HIP源代码4230，HIP编译器驱动器4240，HCC 4260，主机可执行代码4270(2)，HCC设备可执行代码4282，CPU4290和GPU 4292。

在至少一个实施例中，并且如本文先前结合图42A所描述的，CUDA 源代码4210包括但不限于任意数量(包括零)的全局函数4212，任意数量 (包括零)的设备函数4214，任意数量(包括零)的主机函数4216以及任意数量(包括零)的主机/设备函数4218。在至少一个实施例中，CUDA源代码4210还包括但不限于对在任意数量的CUDA API中指定的任意数量的函数的任意数量的调用。

在至少一个实施例中，CUDA到HIP转换工具4220将CUDA源代码 4210转换成HIP源代码4230。在至少一个实施例中，CUDA到HIP转换工具4220将CUDA源代码4210中的每个内核调用从CUDA语法转换为 HIP语法，并将源代码4210中任意数量的其他CUDA调用转换为任意数量的其他功能上相似的HIP调用。

在至少一个实施例中，HIP编译器驱动器4240随后确定目标设备4246 不是启用CUDA的，并生成HIP/HCC编译命令4244。在至少一个实施例中，然后HIP编译器驱动器4240配置HCC 4260以执行HIP/HCC编译命令4244，从而编译HIP源代码4230。在至少一个实施例中，HIP/HCC编译命令4244将HCC 4260配置为使用但不限于HIP/HCC运行时库4258和 HCC头4256来生成主机可执行代码4270(2)和HCC设备可执行代码4282。在至少一个实施例中，HIP/HCC运行时库4258对应于HIP运行时API 4232。在至少一个实施例中，HCC头4256包括但不限于用于HIP和HCC的任意数量和类型的互操作性机制。在至少一个实施例中，主机可执行代码4270 (2)和HCC设备可执行代码4282可以分别在CPU 4290和GPU 4292上执行。

图43示出了根据至少一个实施例的由图42C的CUDA到HIP转换工具4220转换的示例性内核。在至少一个实施例中，CUDA源代码4210将给定内核被设计为解决的总体问题划分为可以使用线程块独立解决的相对粗糙的子问题。在至少一个实施例中，每个线程块包括但不限于任意数量的线程。在至少一个实施例中，每个子问题被划分为相对细小的部分(pieces)，这些部分可以由线程块中的线程协作并行地解决。在至少一个实施例中，线程块内的线程可以通过共享存储器共享数据并通过同步执行以协调存储器访问来协作。

在至少一个实施例中，CUDA源代码4210将与给定内核相关联的线程块组织成线程块的一维、二维或三维网格。在至少一个实施例中，每个线程块包括但不限于任意数量的线程，并且网格包括但不限于任意数量的线程块。

在至少一个实施例中，内核是使用“__global__”声明说明符(specifier) 定义的设备代码中的函数。在至少一个实施例中，使用CUDA内核启动语法4310来指定针对给定内核调用执行内核的网格的尺寸以及相关联的流。在至少一个实施例中，CUDA内核启动语法4310被指定为“KernelName< <<GridSize，BlockSize，SharedMemorySize，Stream>>>(KernelArguments)；”。在至少一个实施例中，执行配置语法是“<<<...>>>”构造，其被插入在内核名称(“KernelName”)和内核参数的括号列表(“KernelArguments”) 之间。在至少一个实施例中，CUDA内核启动语法4310包括但不限于CUDA 启动函数语法而不是执行配置语法。

在至少一个实施例中，“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”)在内核内访问。

在至少一个实施例中，关于CUDA内核启动语法4310，“SharedMemorySize”是一可选参数，它指定共享存储器中除静态分配的存储器外，针对给定内核调用为每个线程块动态分配的字节数。在至少一个实施例中并且关于CUDA内核启动语法4310，SharedMemorySize默认为零。在至少一个实施例中并且关于CUDA内核启动语法4310，“流”是可选的参数，其指定相关联的流并且默认为零以指定默认流。在至少一个实施例中，流是按顺序执行的命令序列(其可能由不同的主机线程发出)。在至少一个实施例中，不同的流可以相对于彼此无序地或同时地执行命令。

在至少一个实施例中，CUDA源代码4210包括但不限于用于示例性内核“MatAdd”的内核定义和主函数。在至少一个实施例中，主函数是在主机上执行的主机代码，并且包括但不限于使内核MatAdd在设备上执行的内核调用。在至少一个实施例中，如图所示，内核MatAdd将大小为NxN 的两个矩阵A和B相加，其中N为正整数，并将结果存储在矩阵C中。在至少一个实施例中，主函数将threadsPerBlock变量定义为16x16， numBlocks变量为N/16xN/16。在至少一个实施例中，然后主函数指定内核调用“MatAdd<<<numBlocks，threadsPerBlock>>>(A，B，C)；”。在至少一个实施例中，并且根据CUDA内核启动语法4310，使用尺寸为 N/16×N/16的线程块网格来执行内核MatAdd，其中每个线程块的尺寸为 16×16。在至少一个实施例中，每个线程块包括256个线程，创建具有足够块的网格以使每个矩阵元素具有一个线程，并且该网格中的每个线程执行内核MatAdd以执行一个逐对的加法。

在至少一个实施例中，在将CUDA源代码4210转换成HIP源代码 4230的同时，CUDA到HIP转换工具4220将CUDA源代码4210中的每个内核调用从CUDA内核启动语法4310转换成HIP内核启动语法4320，并将源代码4210中的任意数量的其他CUDA调用转换为任意数量的其他功能上相似的HIP调用。在至少一个实施例中，HIP内核启动语法4320被指定为“hipLaunchKernelGGL(KernelName,GridSize,BlockSize, SharedMemorySize,Stream,KernelArguments)；”。在至少一个实施例中， KernelName，GridSize，BlockSize，ShareMemorySize，Stream和 KernelArguments中的每一个在HIP内核启动语法4320中具有与在CUDA 内核启动语法4310中(本文先前描述)相同的含义。在至少一个实施例中，参数SharedMemorySize和Stream在HIP内核启动语法4320中是必需的，而在CUDA内核启动语法4310中是可选的。

在至少一个实施例中，除了使内核MatAdd在设备上执行的内核调用之外，图43中描绘的HIP源代码4230的一部分与图43中描绘的CUDA 源代码4210的一部分相同。在至少一个实施例中，在HIP源代码4230中定义内核MatAdd，具有与在CUDA源代码4210中定义内核MatAdd相同的“__global__”声明说明符。在至少一个实施例中，在HIP源代码4230 中的内核调用是“hipLaunchKernelGGL(MatAdd，numBlocks， threadsPerBlock，0、0，A，B，C)；”，而CUDA源代码4210中的相应内核调用是“MatAdd<<<numBlocks，threadsPerBlock>>>(A，B，C)；”。

图44更详细地示出了根据至少一个实施例的图42C的未启用CUDA 的GPU 4292。在至少一个实施例中，GPU 4292由圣塔克拉拉市的AMD 公司开发。在至少一个实施例中，GPU4292可以被配置为以高度并行的方式执行计算操作。在至少一个实施例中，GPU 4292被配置为执行图形管线操作，诸如绘制命令、像素操作、几何计算以及与将图像渲染到显示器相关联的其他操作。在至少一个实施例中，GPU 4292被配置为执行与图形无关的操作。在至少一个实施例中，GPU 4292被配置为执行与图形有关的操作和与图形无关的操作两者。在至少一个实施例中，GPU 4292可以被配置为执行HIP源代码4230中包括的设备代码。

在至少一个实施例中，GPU 4292包括但不限于任意数量的可编程处理单元4420，命令处理器4410，L2高速缓存4422，存储器控制器4470， DMA引擎4480(1)，系统存储器控制器4482，DMA引擎4480(2)和 GPU控制器4484。在至少一个实施例中，每个可编程处理单元4420包括但不限于工作负载管理器4430和任意数量的计算单元4440。在至少一个实施例中，命令处理器4410读取来自一个或更多个命令队列(未示出)的命令，并将命令分发给工作负载管理器4430。在至少一个实施例中，对于每个可编程处理单元4420，相关的工作负载管理器4430将工作分发给包括在可编程处理单元4420中的计算单元4440。在至少一个实施例中，每个计算单元4440可以执行任意数量的线程块，但是每个线程块在单个计算单元4440上执行。在至少一个实施例中，工作组是线程块。

在至少一个实施例中，每个计算单元4440包括但不限于任意数量的 SIMD单元4450和共享存储器4460。在至少一个实施例中，每个SIMD单元4450实现SIMD架构并且被配置为并行执行操作。在至少一个实施例中，每个SIMD单元4450包括但不限于向量ALU 4452和向量寄存器文件4454。在至少一个实施例中，每个SIMD单元4450执行不同的线程束。在至少一个实施例中，线程束是一组线程(例如16个线程)，其中线程束中的每个线程属于单个线程块，并且被配置为基于单个指令集来处理不同的数据集。在至少一个实施例中，可以使用预测来禁用线程束中的一个或更多个线程。在至少一个实施例中，通道是线程。在至少一个实施例中，工作项是线程。在至少一个实施例中，波前是线程束。在至少一个实施例中，线程块中的不同波前可一起同步并经由共享存储器4460进行通信。

在至少一个实施例中，可编程处理单元4420被称为“着色引擎”。在至少一个实施例中，除了计算单元4440之外，每个可编程处理单元4420 还包括但不限于任意数量的专用图形硬件。在至少一个实施例中，每个可编程处理单元4420包括但不限于任意数量(包括零)的几何处理器，任意数量(包括零)的光栅化器，任意数量(包括零)的渲染后端，工作负载管理器4430和任意数量的计算单元4440。

在至少一个实施例中，计算单元4440共享L2高速缓存4422。在至少一个实施例中，L2高速缓存4422被分区。在至少一个实施例中，GPU 4292 中的所有计算单元4440可访问GPU存储器4490。在至少一个实施例中，存储器控制器4470和系统存储器控制器4482促进GPU 4292与主机之间的数据传输，并且DMA引擎4480(1)使能GPU 4292与此主机之间的异步存储器传输。在至少一个实施例中，存储器控制器4470和GPU控制器4484 促进GPU 4292与其他GPU 4292之间的数据传输，并且DMA引擎4480 (2)使能GPU 4292与其他GPU 4292之间的异步存储器传输。

在至少一个实施例中，GPU 4292包括但不限于任意数量和类型的系统互连，该系统互连促进在GPU 4292内部或外部的任意数量和类型的直接或间接链接的组件之间的数据和控制传输。在至少一个实施例中，GPU 4292 包括但不限于耦合到任意数量和类型的外围设备的任意数量和类型的I/O 接口(例如，PCIe)。在至少一个实施例中，GPU 4292可以包括但不限于任意数量(包括零)的显示引擎和任意数量(包括零)的多媒体引擎。在至少一个实施例中，GPU 4292实现了存储器子系统，该存储器子系统包括但不限于任意数量和类型的存储器控制器(例如，存储器控制器4470和系统存储器控制器4482)以及专用于一个组件或在多个组件之间共享的存储器设备(例如，共享存储器4460)。在至少一个实施例中，GPU4292实现了高速缓存子系统，该高速缓存子系统包括但不限于一个或更多个高速缓存存储器(例如，L2高速缓存4422)，每个高速缓存存储器可以是私有的或在任意数量的组件(例如，SIMD单元4450，计算单元4440和可编程处理单元4420)之间共享。

图45示出了根据至少一个实施例的示例性CUDA网格4520的线程如何被映射到图44的不同计算单元4440。在至少一个实施例中，并且仅出于说明目的，网格4520具有BX乘以BY乘以1的GridSize和TX乘以TY 乘以1的BlockSize。因此，在至少一个实施例中，网格4520包括但不限于(BX*BY)线程块4530，每个线程块4530包括但不限于(TX*TY) 线程4540。线程4540在图45中被描绘为弯曲箭头。

在至少一个实施例中，网格4520被映射到可编程处理单元4420(1)，该可编程处理单元4420(1)包括但不限于计算单元4440(1)-4440(C)。在至少一个实施例中并且如图所示，将(BJ*BY)线程块4530映射到计算单元4440(1)，并且将其余线程块4530映射到计算单元4440(2)。在至少一个实施例中，每个线程块4530可以包括但不限于任意数量的线程束，并且每个线程束被映射到图44的不同的SIMD单元4450。

在至少一个实施例中，给定线程块4530中的线程束可以一起同步并通过关联的计算单元4440中包括的共享存储器4460进行通信。例如并且在至少一个实施例中，线程块4530(BJ，1)中的线程束可以一起同步并通过共享存储器4460(1)进行通信。例如并且在至少一个实施例中，线程块4530(BJ+1，1)中的线程束可以一起同步并通过共享存储器4460(2) 进行通信。

至少一个实施例可以根据以下条款中的至少一项进行描述:

1.一种处理器，包括：

一个或更多个电路，用于至少部分地基于一个或更多个对象从第一视频帧到第二视频帧的多个可能运动中的一个来生成第三视频帧。

2.根据条款1所述的处理器，其中所述一个或更多个电路进一步用于至少部分地基于与所述第二视频帧的像素相关联的后向运动向量来确定所述多个可能运动。

3.根据条款1或2所述的处理器，其中所述一个或更多个对象是像素。

4.根据条款1-3中的任一项所述的处理器，其中所述一个或更多个电路用于进一步至少部分地基于相机视点的一个或更多个运动来生成所述第三视频帧。

5.根据条款1-4中的任一项所述的处理器，其中所述一个或更多个电路用于至少部分地基于深度信息来选择所述多个可能运动中的所述一个。

6.根据条款1-5中的任一项所述的处理器，其中所述一个或更多个电路用于至少部分地基于与第二视频帧的像素相关联的后向运动向量和相机视点的一个或更多个运动来生成一个或更多个附加视频帧。

7.根据条款1-6中的任一项所述的处理器，其中所述一个或更多个电路用于至少部分地基于从一个或更多个缓冲器接收到第一视频帧、第二视频帧、深度信息和后向运动向量来生成第三视频帧。

8.一种机器可读介质，其上存储有一组指令，如果由一个或更多个处理器执行所述一组指令，则使该一个或更多个处理器至少：

至少部分地基于一个或更多个对象从第一视频帧到第二视频帧的多个可能运动中的一个来生成第三视频帧。

9.根据条款8所述的机器可读介质，其中所述多个可能运动中的每一个与和第二视频帧的像素深度值相关联的一组后向运动向量中的从第二视频帧到第一视频帧的后向运动向量对应，并且如果由所述一个或更多个处理器执行所述指令，则进一步使所述一个或更多个处理器：

至少部分地基于与所述一组后向运动向量中的一个相关联的深度值来识别所述多个可能运动中的一个；以及

至少部分地基于所识别的运动生成第三视频帧。

10.根据条款8或9所述的机器可读介质，其中所述一个或更多个对象是像素。

11.根据条款8-10中的任一项所述的机器可读介质，其中如果由所述一个或更多个处理器执行所述指令，则进一步使所述一个或更多个处理器：

确定第一视频帧和第二视频帧之间的相机视点矩阵的变化；以及

至少部分地基于所确定的相机视点矩阵的变化来生成第三视频帧。

12.根据条款8-11中任一项所述的机器可读介质，其中如果由所述一个或更多个处理器执行所述指令，则进一步使所述一个或更多个处理器：

确定第三视频帧中的一组被遮挡像素位置；

确定第三视频帧中的一组被去遮挡像素位置；以及

至少部分地基于所述一组被遮挡像素位置和所述一组被去遮挡像素位置来生成第三视频帧。

13.根据条款8-12中的任一项所述的机器可读介质，其中如果由所述一个或更多个处理器执行所述指令，则进一步使所述一个或更多个处理器：

至少部分地基于一组后向运动向量生成从第一视频帧到第二视频帧的一组估计的前向运动向量，其中所述一组后向运动向量是从第二视频帧到第一视频帧；以及

至少部分地基于所生成的一组估计的前向运动向量和一组后向运动向量来生成第三视频帧。

14.根据条款8-13中的任一项所述的机器可读介质，其中如果由所述一个或更多个处理器执行所述指令，则进一步使所述一个或更多个处理器：

至少部分地基于所生成的一组估计的前向运动向量生成从第三视频帧到第二视频帧的一组中间前向运动向量；

生成从第三视频帧到第一个视频帧的一组中间后向运动向量；以及

至少部分地基于所述一组中间前向运动向量和所述一组中间后向运动向量来生成第三视频帧。

15.根据条款8-14中的任一项所述的机器可读介质，其中如果由所述一个或更多个处理器执行所述指令，则进一步使所述一个或更多个处理器：

至少部分地基于从一个或更多个缓冲器接收到第一视频帧、第二视频帧、深度信息和后向运动向量来生成第三视频帧。

16.一种方法，包括：

至少部分地基于一个或更多个对象从第一视频帧到第二视频帧的多个可能运动中的一个来生成第三视频帧。

17.根据条款16所述的方法，还包括至少部分地基于与第二视频帧的像素相关联的后向运动向量来确定所述多个可能运动，其中所述多个可能运动中的每一个对应于指向第一视频帧中相同像素位置的后向运动向量。

18.根据条款16或17所述的方法，其中所述一个或更多个对象是像素。

19.根据条款16-18中任一项的方法，进一步包括：

确定第一视频帧和第二视频帧之间的相机视点的一个或更多个运动；以及

至少部分地基于所确定的相机视点的一个或更多个运动来生成第三视频帧。

20.根据条款16-19中任一项所述的方法，其中选择所述多个可能运动中的所述一个是至少部分地基于所述第二视频帧中的像素的深度信息。

21.根据条款16-20中任一项所述的方法，还包括：

生成遮挡掩模；

生成去遮挡掩模；以及

至少部分地基于所述遮挡掩模和所述去遮挡掩模生成第三视频帧。

22.根据条款16-21中任一项所述的方法，其中生成第三视频帧是至少部分地基于从一个或更多个缓冲器接收到第一视频帧、第二视频帧、深度信息和后向运动向量。

23.一种系统，包括：

一个或更多个处理器，用于至少部分地基于一个或更多个对象从第一视频帧到第二视频帧的多个可能运动中的一个来生成第三视频帧；以及

一个或更多个存储器，用于存储第三视频帧。

24.根据条款23所述的系统，其中所述一个或更多个处理器进一步用于：

至少部分地基于与第二视频帧的像素相关联的后向运动向量来确定所述多个可能运动；以及

至少部分地基于深度信息选择所述多个可能运动之一。

25.根据条款23或24所述的系统，其中所述一个或更多个对象是像素。

26.根据条款23-25中的任一项所述的系统，其中所述一个或更多个处理器进一步用于：

确定第一视频帧和第二视频帧之间的相机视点变化；以及

至少部分地基于所确定的相机视点变化来生成第三视频帧。

27.根据条款23-26中的任一项所述的系统，其中所述一个或更多个处理器进一步用于：

识别第三视频帧的像素位置，该像素位置具有使用仅在第一视频帧和第二视频帧中的一个中的中间运动向量识别的对应像素；以及仅对具有所识别像素位置的对应像素的视频帧的像素数据进行采样。

28.根据条款23-27中的任一项所述的系统，其中一个或更多个处理器用于至少部分地基于从一个或更多个缓冲器接收到第一视频帧、第二视频帧、深度信息和后向运动向量来生成第三视频帧。

29.一种处理器，包括：

一个或更多个电路，用于：

采样第一视频帧的第一组像素数据；

至少部分地基于从第一视频帧到第二视频帧的一组前向运动向量对第二视频帧的第二组像素数据进行采样；以及

至少部分地基于所述第一组像素数据和所述第二组像素数据来生成第一视频帧和第二视频帧之间的中间视频帧。

30.根据条款29所述的处理器，其中所述一个或更多个电路用于至少部分地基于从所述第二视频帧到所述第一视频帧的一组后向运动向量来对所述第一组像素数据进行采样。

31.根据条款29或30所述的处理器，其中所述一个或更多个电路用于：识别中间视频帧的像素，该像素具有使用仅在第一视频帧和第二视频帧中的一个中的运动向量识别的对应像素；以及仅对具有所识别像素的对应像素的视频帧的像素数据进行采样。

32.根据条款29-31中的任一项所述的处理器，其中所述一个或更多个电路还用于至少部分地基于所述一组前向运动向量来生成从中间视频帧到第二视频帧的一组中间前向运动向量，并且其中所述一个或更多个电路用于至少部分地基于所述一组中间前向运动向量对第二组像素数据进行采样。

33.根据条款29-32中的任一项所述的处理器，其中所述一个或更多个电路进一步用于：

生成一组从中间视频帧到第二视频帧的中间前向运动向量；

生成一组从中间视频帧到第一视频帧的中间后向运动向量；以及

至少部分地基于所述一组中间前向运动向量和所述一组中间后向运动向量来生成中间视频帧。

34.根据条款29-33中的任一项所述的处理器，其中所述一个或更多个电路进一步用于：

至少部分地基于第一视频帧中的像素的深度值生成一组中间前向运动向量；以及

至少部分地基于第二视频帧中的像素的深度值生成一组中间后向运动向量。

35.根据条款29-34中任一项所述的处理器，其中所述一个或更多个电路进一步用于：

生成遮挡掩模；

生成去遮挡掩模；以及

至少部分地基于所述遮挡掩模和所述去遮挡掩模生成中间视频帧。

36.根据条款29-35中任一项所述的处理器，其中所述一个或更多个电路用于至少部分地基于从一个或更多个缓冲器接收第一视频帧、第二视频帧、深度信息、前向运动向量和后向运动向量来生成中间帧。

37.根据条款29-36中任一项所述的处理器，其中所述一组前向运动向量中的每个前向运动向量包括投影的顶点运动。

38.一种其上存储有一组指令的机器可读介质，如果由一个或更多个处理器执行所述一组指令，则使所述一个或更多个处理器至少：

对第一视频帧的第一组像素数据进行采样；

至少部分地基于从第一视频帧到第二视频帧的一组前向运动向量对第二视频帧的第二组像素数据进行采样；以及

至少部分地基于第一组像素数据和第二组像素数据生成第一视频帧和第二视频帧之间的中间视频帧。

39.根据条款38所述的机器可读介质，其中如果由所述一个或更多个处理器执行所述指令，则进一步使所述一个或更多个处理器：

至少部分地基于从第二视频帧到第一视频帧的一组后向运动向量对第一组像素数据进行采样。

40.根据条款38或39所述的机器可读介质，其中所述中间视频帧是第一中间视频帧，并且如果由所述一个或更多个处理器执行所述指令，则进一步使所述一个或更多个处理器：

至少部分地基于第一组像素数据和第二组像素数据在第一视频帧和第二视频帧之间生成第二中间视频帧。

41.根据条款38-40中的任一项所述的机器可读介质，其中如果由所述一个或更多个处理器执行所述指令，则进一步使所述一个或更多个处理器：

至少部分地基于来自第一组像素数据和第二组像素数据的、针对中间视频帧的像素的像素颜色信息的加权平均来生成中间视频帧，该像素在使用第一运动向量识别的第一视频帧中和在使用第二运动向量识别的第二视频帧中具有对应的像素。

42.根据条款38-41中任一项所述的机器可读介质，其中像素颜色信息是来自第一组像素数据的第一组RGB值和来自第二组像素数据的第二组RGB 值。

43.根据条款38-42中任一项所述的机器可读介质，其中如果由所述一个或更多个处理器执行所述指令，则进一步使所述一个或更多个处理器：

确定一组被遮挡的像素；

确定一组被去遮挡的像素；以及

至少部分地基于所述一组被遮挡的像素集和所述一组被去遮挡的像素来生成中间视频帧。

44.根据条款38-43中任一项所述的机器可读介质，其中所述一组前向运动向量中的每个前向运动向量包括投影的顶点运动。

45.一种方法，包括：

对第一视频帧的第一组像素数据进行采样；

至少部分地基于从第一视频帧到第二视频帧的一组前向运动向量对第二视频帧的第二组像素数据进行采样；以及

至少部分地基于第一组像素数据和第二组像素数据在第一视频帧和第二视频帧之间生成中间视频帧。

46.根据条款45所述的方法，其中对所述第一组像素数据进行采样是至少部分地基于从所述第二视频帧到所述第一视频帧的一组后向运动向量。

47.根据条款45或46所述的方法，还包括：

生成从中间视频帧到第二视频帧的一组中间前向运动向量；以及

生成从中间视频帧到第一视频帧的一组中间后向运动向量，其中对第一组像素数据进行采样是至少部分地基于所述一组中间后向运动向量。

48.根据条款45-47中的任一项所述的方法，其中生成所述一组中间前向运动向量是至少部分地基于所述第一视频帧中的像素的深度值。

49.根据条款45-48中任一项的方法，进一步包括：

识别中间视频帧中的在使用前向运动向量识别的第二视频帧中不具有对应像素，或在使用后向运动向量识别的第一视频帧中不具有对应像素的像素；

至少部分地基于在与第一视频帧和第二视频帧中的识别像素相同位置的像素的深度值，将中间视频帧的识别像素设置为在与第一视频帧或第二视频帧中的识别像素相同位置处的像素之一的颜色。

50.根据条款45-49中任一项的方法，进一步包括：

生成遮挡掩模；

生成去遮挡掩模；

至少部分地基于所述遮挡掩模对第一组像素数据进行采样；以及

至少部分地基于所述去遮挡掩模对第二组像素数据进行采样。

51.根据条款45-50中任一项所述的方法，其中所述一组前向运动向量中的每个前向运动向量包括投影的顶点运动。

52.一种系统，包括：

一个或更多个处理器，用于：

对第一视频帧的第一组像素数据进行采样；

至少部分地基于从第一视频帧到第二视频帧的一组前向运动向量对第二视频帧的第二组像素数据进行采样；以及

至少部分地基于第一组像素数据和第二组像素数据在第一视频帧和第二视频帧之间生成中间视频帧；以及

一个或更多个存储器，用于存储中间视频帧。

53.根据条款52所述的系统，其中所述一个或更多个处理器用于：

从一个或更多个缓冲器接收一组前向运动向量和一组后向运动向量；以及

至少部分地基于一组后向运动向量对第一组像素数据进行采样。

54.根据条款52或53所述的系统，其中所述一个或更多个处理器用于：

从缓冲器接收第一组深度值；

至少部分地基于所述一组前向运动向量和所述第一组深度值，生成从中间视频帧到第二视频帧的一组中间前向运动向量；以及

至少部分地基于所述一组中间前向运动向量生成中间视频帧。

55.根据条款52-54中的任一项所述的系统，其中所述一个或更多个处理器用于：

接收第二组深度值；

至少部分地基于一组后向运动向量和所述第二组深度值，生成从中间视频帧到第一视频帧的一组中间后向运动向量；以及

至少部分地基于所述一组中间后向运动向量生成中间视频帧。

56.根据条款52-55中的任一项所述的系统，其中所述一个或更多个处理器用于：

生成遮挡掩模；

生成去遮挡掩模；以及

至少部分地基于所述遮挡掩模和所述去遮挡掩模生成中间视频帧。

57.根据条款52-56中任一项所述的系统，其中所述一个或更多个处理器进一步用于：

至少部分地基于第一视频帧、第二视频帧、一组前向运动向量集、一组后向运动向量和一组深度指示符，在第一视频帧和第二视频帧之间生成一个或更多个附加的中间视频帧。

58.根据条款52-57中任一项所述的系统，其中所述一组前向运动向量中的每个前向运动向量包括投影的顶点运动。

其他变型在本公开的精神内。因此，尽管公开的技术易于进行各种修改和替代构造，但是某些示出的其实施例在附图中示出并且已经在上面进行了详细描述。然而，应理解，无意将公开内容限制为所公开的一种或更多种特定形式，而是相反，其意图是涵盖落入如所附权利要求书所定义的本公开内容的精神和范围内的所有修改、替代构造和等同物。

除非另有说明或显然与上下文矛盾，否则在描述所公开的实施例的上下文中(特别是在所附权利要求的上下文中)，术语“一”和“一个”和“该”以及类似指代的使用应被解释为涵盖单数和复数，而不是作为术语的定义。除非另有说明，否则术语“包括”、“具有”、“包含”和“含有”应被解释为开放式术语(意味着“包括但不限于”)。术语“连接” (在未经修改时指的是物理连接)应解释为部分或全部包含在内、附接到或连接在一起，即使有某些介入。除非本文另外指出，否则本文中对数值范围的引用仅旨在用作分别指代落入该范围内的每个单独值的简写方法，并且每个单独值都被并入说明书中，就如同其在本文中被单独叙述一样。除非另外指出或与上下文矛盾，否则术语“集”(例如“项目集”)或“子集”的使用应解释为包括一个或更多个成员的非空集合。此外，除非另外指出或与上下文矛盾，否则术语相应集的“子集”不一定表示对应集的适当子集，而是子集和对应集可以相等。

除非以其他方式明确指出或与上下文明显矛盾，否则诸如“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中的至少一个。另外，除非另有说明或与上下文矛盾，否则术语“多个”表示复数的状态(例如，“多个项目”表示多个项目)。多个项目中项目的数量至少为两个，但如果明确指示或通过上下文指示，则可以更多。此外，除非另有说明或从上下文中可以清楚得知，否则短语“基于”是指“至少部分基于”而不是“仅基于”。

除非本文另外指出或与上下文明显矛盾，否则本文描述的过程的操作可以任何合适的顺序执行。在至少一个实施例中，诸如本文所述的那些过程(或其变形和/或其组合)之类的过程在配置有可执行指令的一个或更多个计算机系统的控制下执行，并且被实现为代码(例如，可执行指令，一个或更多个计算机程序或一个或更多个应用程序)，该代码通过硬件或其组合在一个或更多个处理器上共同执行。在至少一个实施例中，代码以例如计算机程序的形式存储在计算机可读存储介质上，该计算机程序包括可由一个或更多个处理器执行的多个指令。在至少一个实施例中，计算机可读存储介质是非暂时性计算机可读存储介质，其排除了暂时性信号(例如，传播的瞬态电或电磁传输)，但包括非暂时性数据存储电路(例如，缓冲区、高速缓存和队列)。在至少一个实施例中，代码(例如，可执行代码或源代码)被存储在其上存储有可执行指令的一组一个或更多个非暂时性计算机可读存储介质(或用于存储可执行指令的其他存储器)上，该可执行指令在由计算机系统的一个或更多个处理器执行时(即，作为被执行的结果)，使得计算机系统执行本文所述的操作。在至少一个实施例中，一组非暂时性计算机可读存储介质包括多个非暂时性计算机可读存储介质，并且多个非暂时性计算机可读存储介质中的个体非暂时性存储介质中的一个或更多个缺少全部代码，而是多个非暂时性计算机可读存储介质共同存储全部代码。在至少一个实施例中，可执行指令被执行，以使得不同的指令由不同的处理器执行，例如，非暂时性计算机可读存储介质存储指令，并且主中央处理单元(“CPU”)执行一些指令，而图形处理单元(“GPU”) 执行其他指令。在至少一个实施例中，计算机系统的不同组件具有单独的处理器，并且不同的处理器执行指令的不同子集。

因此，在至少一个实施例中，计算机系统被配置为实现单独地或共同地执行本文所述的过程的操作的一个或更多个服务，并且这样的计算机系统被配置有使能实施操作的适用的硬件和/或软件。此外，实现本公开的至少一个实施例的计算机系统是单个设备，并且在另一实施例中是分布式计算机系统，其包括以不同方式操作的多个设备，使得分布式计算机系统执行本文所述的操作，并且使得单个设备不执行所有操作。

本文提供的任何和所有示例或示例性语言(例如，“诸如”)的使用仅旨在更好地阐明本公开的实施例，并且不对公开的范围构成限制，除非另有要求。说明书中的任何语言都不应被解释为表示任何未要求保护的要素对于实践公开内容是必不可少的。

本文引用的所有参考文献，包括出版物、专利申请和专利，均通过引用并入本文，其程度就如同每个参考文献被单独且具体地指示为以引用的方式并入本文并且其全部内容在本文中阐述一样。

在说明书和权利要求中，可以使用术语“耦合”和“连接”以及它们的派生词。应当理解，这些术语可能不旨在作为彼此的同义词。相反，在特定示例中，“连接”或“耦合”可用于指示两个或更多个元件彼此直接或间接物理或电接触。“耦合”也可能意味着两个或更多个元素彼此不直接接触，但仍彼此协作或交互。

除非另有明确说明，否则可以理解，在整个说明书中，诸如“处理”、“计算”、“计算”、“确定”等之类的术语，是指计算机或计算系统或类似的电子计算设备的动作和/或过程，其将计算系统的寄存器和/或存储器中表示为物理量(例如电子)的数据处理和/或转换为类似表示为计算系统的存储器、寄存器或其他此类信息存储、传输或显示设备中的物理量的其他数据。

以类似的方式，术语“处理器”可以指处理来自寄存器和/或存储器的电子数据并将该电子数据转换成可以存储在寄存器和/或存储器中的其他电子数据的任何设备或存储器的一部分。作为非限制性示例，“处理器”可以是CPU或GPU。“计算平台”可以包括一个或更多个处理器。如本文所使用的，“软件”进程可以包括例如随时间执行工作的软件和/或硬件实体，诸如任务、线程和智能代理。同样，每个过程可以指代多个过程，以连续地或间歇地顺序地或并行地执行指令。术语“系统”和“方法”在本文中可以互换使用，只要系统可以体现一种或更多种方法，并且方法可以被认为是系统。

在本文件中，可以参考获得、获取、接收或将模拟或数字数据输入子系统、计算机系统或计算机实现的机器中。可以通过多种方式来完成获得、获取、接收或输入模拟和数字数据的过程，例如通过接收作为函数调用或对应用程序编程接口的调用的参数的数据。在一些实现方式中，可以通过经由串行或并行接口传输数据来完成获得、获取、接收或输入模拟或数字数据的过程。在另一实现方式中，可以通过经由计算机网络将数据从提供实体传输到获取实体来完成获得、获取、接收或输入模拟或数字数据的过程。也可以参考提供、输出、传送、发送或呈现模拟或数字数据。在各种示例中，提供、输出、传送、发送或呈现模拟或数字数据的过程可以通过将数据作为函数调用的输入或输出参数、应用程序编程接口或进程间通信机制的参数进行传输来实现。

尽管上面的讨论阐述了所描述的技术的示例实现，但是其他架构可以用于实现所描述的功能，并且旨在落入本公开的范围内。此外，尽管出于讨论的目的在上面定义了具体的职责分配，但是根据情况，可以以不同的方式分配和划分各种功能和职责。

此外，尽管已经用特定于结构特征和/或方法动作的语言描述了主题，但是应当理解，所附权利要求书所要求保护的主题不必限于所描述的特定特征或动作。而是，公开了特定的特征和动作作为实现权利要求的示例性形式。

Claims (28)

1.一种处理器，包括：

一个或更多个电路，用于至少部分地基于一个或更多个对象从第一视频帧到第二视频帧的多个可能运动中的一个来生成第三视频帧。

2.根据权利要求1所述的处理器，其中所述一个或更多个电路进一步用于至少部分地基于与所述第二视频帧的像素相关联的后向运动向量来确定所述多个可能运动。

3.根据权利要求1所述的处理器，其中所述一个或更多个对象是像素。

4.根据权利要求1所述的处理器，其中所述一个或更多个电路进一步用于至少部分地基于相机视点的一个或更多个运动来生成所述第三视频帧。

5.根据权利要求1所述的处理器，其中所述一个或更多个电路用于至少部分地基于深度信息来选择所述多个可能运动中的所述一个。

6.根据权利要求1所述的处理器，其中所述一个或更多个电路用于至少部分地基于与所述第二视频帧的像素相关联的后向运动向量以及相机视点的一个或更多个运动来生成一个或更多个附加视频帧。

7.根据权利要求1所述的处理器，其中所述一个或更多个电路用于至少部分地基于从一个或更多个缓冲器接收到所述第一视频帧、所述第二视频帧、深度信息和后向运动向量来生成所述第三视频帧。

8.一种机器可读介质，其上存储有一组指令，如果由一个或更多个处理器执行所述一组指令，则使所述一个或更多个处理器至少：

至少部分地基于一个或更多个对象从第一视频帧到第二视频帧的多个可能运动中的一个来生成第三视频帧。

9.根据权利要求8所述的机器可读介质，其中所述多个可能运动中的每一个与和所述第二视频帧的像素深度值相关联的一组后向运动向量中的从所述第二视频帧到所述第一视频帧的后向运动向量对应，并且如果由所述一个或更多个处理器执行所述指令，则进一步使所述一个或更多个处理器：

至少部分地基于与所述一组后向运动向量中的一个相关联的深度值来识别所述多个可能运动中的所述一个；以及

至少部分地基于所识别的运动生成所述第三视频帧。

10.根据权利要求9所述的机器可读介质，其中所述一个或更多个对象是像素。

11.根据权利要求8所述的机器可读介质，其中如果由所述一个或更多个处理器执行所述指令，则进一步使所述一个或更多个处理器：

确定所述第一视频帧和所述第二视频帧之间的相机视点矩阵的变化；以及

至少部分地基于所确定的所述相机视点矩阵的变化来生成所述第三视频帧。

12.根据权利要求8所述的机器可读介质，其中如果由所述一个或更多个处理器执行所述指令，则进一步使所述一个或更多个处理器：

确定所述第三视频帧中的一组被遮挡像素位置；

确定所述第三视频帧中的一组被去遮挡像素位置；以及

至少部分地基于所述一组被遮挡像素位置和所述一组被去遮挡像素位置来生成所述第三视频帧。

13.根据权利要求8所述的机器可读介质，其中如果由所述一个或更多个处理器执行所述指令，则进一步使所述一个或更多个处理器：

至少部分地基于一组后向运动向量生成从所述第一视频帧到所述第二视频帧的一组估计的前向运动向量，其中所述一组后向运动向量是从所述第二视频帧到所述第一视频帧；以及

至少部分地基于所生成的所述一组估计的前向运动向量和所述一组后向运动向量生成所述第三视频帧。

14.根据权利要求13所述的机器可读介质，其中如果由所述一个或更多个处理器执行所述指令，则进一步使所述一个或更多个处理器：

至少部分地基于所生成的所述一组估计的前向运动向量生成从所述第三视频帧到所述第二视频帧的一组中间前向运动向量；

生成从所述第三视频帧到所述第一视频帧的一组中间后向运动向量；以及

至少部分地基于所述一组中间前向运动向量和所述一组中间后向运动向量生成所述第三视频帧。

15.根据权利要求8所述的机器可读介质，其中如果由所述一个或更多个处理器执行所述指令，则进一步使所述一个或更多个处理器：

至少部分地基于从一个或更多个缓冲器接收到所述第一视频帧、所述第二视频帧、深度信息和后向运动向量来生成所述第三视频帧。

16.一种方法，包括：

至少部分地基于一个或更多个对象从第一视频帧到第二视频帧的多个可能运动中的一个来生成第三视频帧。

17.根据权利要求16所述的方法，还包括：至少部分地基于与所述第二视频帧的像素相关联的后向运动向量确定所述多个可能运动，其中所述多个可能运动中的每一个对应于指向所述第一视频帧中相同像素位置的后向运动向量。

18.根据权利要求16所述的方法，其中所述一个或更多个对象是像素。

19.根据权利要求16所述的方法，还包括：

确定所述第一视频帧和所述第二视频帧之间的相机视点的一个或更多个运动；以及

至少部分地基于所确定的所述相机视点的一个或更多个运动来生成所述第三视频帧。

20.根据权利要求16所述的方法，其中选择所述多个可能运动中的一个是至少部分地基于所述第二视频帧中的像素的深度信息。

21.根据权利要求16所述的方法，还包括：

生成遮挡掩模；

生成去遮挡掩模；以及

至少部分地基于所述遮挡掩模和所述去遮挡掩模生成所述第三视频帧。

22.根据权利要求16所述的方法，其中生成所述第三视频帧是至少部分地基于从一个或更多个缓冲器接收到所述第一视频帧、所述第二视频帧、深度信息和后向运动向量。

23.一种系统，包括：

一个或更多个处理器，用于至少部分地基于一个或更多个对象从第一视频帧到第二视频帧的多个可能运动中的一个来生成第三视频帧；以及

一个或更多个存储器，用于存储所述第三视频帧。

24.根据权利要求23所述的系统，其中所述一个或更多个处理器进一步用于：

至少部分地基于与所述第二视频帧的像素相关联的后向运动向量来确定所述多个可能运动；以及

至少部分地基于深度信息选择所述多个可能运动之一。

25.根据权利要求23所述的系统，其中所述一个或更多个对象是像素。

26.根据权利要求23所述的系统，其中所述一个或更多个处理器进一步用于：

确定所述第一视频帧和所述第二视频帧之间的相机视点变化；以及

至少部分地基于所确定的相机视点变化来生成所述第三视频帧。

27.根据权利要求23所述的系统，其中所述一个或更多个处理器进一步用于：

识别所述第三视频帧的像素位置，该像素位置具有使用仅在所述第一视频帧和所述第二视频帧中的一个中的中间运动向量来识别的对应像素；并且仅对具有所识别的像素位置的对应像素的所述视频帧的像素数据进行采样。

28.根据权利要求23所述的系统，其中所述一个或更多个处理器用于至少部分地基于从一个或更多个缓冲器接收到所述第一视频帧、所述第二视频帧、深度信息和后向运动向量来生成所述第三视频帧。

Images