CN109964244A - 用于图形处理的本地图像块 - Google Patents

Abstract

本发明公开了涉及用于图形处理的硬件支持的柔性数据结构的技术。在一些实施方案中，所述数据结构的维度可在X方向、Y方向、每个像素的多个样本和每个样本的数据量中配置。在一些实施方案中，这些属性可使用硬件寄存器来配置。在一些实施方案中，所述数据结构在整个正在处理的图块上是持续的，由此使得渲染传递的渲染线程和中间渲染计算线程两者均可访问本地存储器上下文。

Description

用于图形处理的本地图像块

背景技术

技术领域

本公开整体涉及图形处理器，更具体地讲，涉及可编程着色器架构。

相关技术描述

图形处理通常涉及针对不同的图形元素(例如，像素或顶点)并行地执行相同的指令。此外，相同组的图形指令通常被多次执行(例如，用于在不同的时间针对不同的图形元素或针对相同的图形元素执行特定功能)。图形处理器(GPU)通常包括在移动设备诸如蜂窝电话、可穿戴设备等中，其中功率消耗和处理器区域是重要的设计问题。

图形单元通常利用具有专用于特定处理元件的高速缓存和存储器以及在多个处理元素之间共享的较高级别高速缓存和存储器的存储器分级结构。此外，一些存储器(诸如系统存储器)可与非图形处理元件诸如中央处理单元(CPU)共享。一些图形架构允许进行无序存储器访问，并且可以在共享高速缓存或本地高速缓存中高速缓存各种不同级别的数据。在这些架构中强制执行存储器一致性的硬件可能消耗相当大的功率，并且可能限制性能。

一般来讲，图形工作在概念上可分为三种类型：矢量任务、像素任务和计算任务。顶点处理涉及使用多边形来表示图像，其中矢量定义多边形。顶点着色的输出通常被光栅化以生成片段信息，该片段信息由像素/片段着色器操作以生成用于输出至显示器的像素数据。计算处理涉及诸如生成光列表、生成纹理图(或其他约简算法)等其他辅助任务。

通常，图形渲染传递(例如，对于产生像素数据的像素着色器线程)由着色器内核使用其本地存储器来执行。一旦完成渲染，结果就被写入设备存储器和可供其他核心或处理元件使用。因此，传统上，计算任务是在渲染器之间使用设备存储器中的数据来执行。换句话讲，设备存储器传统上用于在计算任务和像素渲染任务之间共享数据。然而，对设备存储器的访问可能是GPU设计的瓶颈，并且可能消耗大量电力。

通常期望GPU以特定顺序执行片段/像素操作(以图形程序向GPU提交操作的顺序)。一些GPU被配置为生成片段的“传递组”，其中操作顺序无关紧要(例如，给定传递组中的片段通常在屏幕空间中不重叠，因此组内的操作顺序不应当彼此影响)。传统上，对传递组排序由此使得较旧的传递组在执行较新的传递组前完成。强制传递组排序可能是性能瓶颈，可能消耗大量的电力。

发明内容

根据一些实施方案，公开了用于在具有宽松存储器访问排序的图形架构中保持存储器一致性的技术。在一些实施方案中，这些技术用于具有至少一个分体式高速缓存级别(例如，在图像和缓冲数据之间分割)的存储器分级结构中。在一些实施方案中，除非使用软件指令明确指定，否则存储器一致性和排序不宽松。这可将一致性的责任转移给编译器或开发器，这可増加GPU性能和/或降低功率消耗。各种操作、围栏和屏障操作可指定诸如获取、释放或顺序一致的存储器属性，并且还可指定实施属性的范围。可以根据指定的范围在存储器分级结构中的不同级别强制执行这些属性。

在一些实施方案中，在不访问共享存储器的情况下执行中间渲染的计算任务。这可减少存储器流量(例如，共享存储器传统上用于在计算内核和渲染传递之间共享数据)，并且计算任务可能能够使用本地存储器与片段着色器共享数据。

在一些实施方案中，硬件寄存器可配置为指定本地存储器中的柔性本地图像块数据结构的维度。可配置的示例性维度包括X维度和Y维度、每个像素的样本数以及每个样本的数据量。在一些实施方案中，本地图像块数据结构可促进中间渲染计算任务。例如，本地图像块可由计算内核声明并通过渲染传递来访问或反之亦然。

在一些实施方案中，图形单元被配置为用于基于明确的软件控制使对像素资源的访问同步。在这些实施方案中，图形单元可不被配置为隐式地实施绘制排序或API提交排序。相反，除非显式地同步，否则可能无序访问像素资源。图形单元可被配置为执行等待指令和释放指令来使对像素资源的访问同步，诸如颜色附件、深度缓存等等。可将像素资源映射到标识符(例如，由图形驱动程序)，以及用于线程的硬件寄存器可保持用于每个标识符的状态信息。这些同步技术可促进中间渲染计算(例如，它们可尤其可用于协调对本地图像块数据的访问，例如，因为计算任务可能是图块范围的，而片段着色器是像素范围的)。这些同步技术还可降低功率消耗并提高性能。

附图说明

图1A是示出示例性图形处理流的框图。

图1B是示出了图形单元的一个实施方案的框图。

图2是示出根据一些实施方案的示例性存储器分级结构的框图。

图3是示出了根据一些实施方案的用于保持存储器一致性的示例性方法的流程图。

图4A和图4B是示出根据一些实施方案的示例性命令缓冲器执行(包括中间渲染计算)的图示。

图5是示出根据一些实施方案的用于执行中间渲染辅助计算任务的示例性方法的流程图。

图6是示出根据一些实施方案的柔性本地图像块数据结构的图示。

图7是示出根据一些实施方案的用于使用本地图像块的示例性方法的流程图。

图8是示出根据一些实施方案的示例性像素资源同步电路系统和代码的图示。

图9是示出了根据一些实施方案的用于像素资源同步的示例性方法的流程图。

图10是示出了包括图形单元的设备的一个实施方案的框图。

图11是示出了根据一些实施方案的示例性计算机可读介质的框图。

本说明书包括对各种实施方案的参考，以指示本公开并非旨在提及一个特定具体实施，而是提及落入包括所附权利要求书的本公开的实质内的一系列实施方案。特定特征、结构或特性可以与本公开一致的任何合适的方式被组合。

在本公开内，不同实体(其可被不同地称为“单元”、“电路”、其他部件等)可被描述或声称成“被配置为”执行一个或多个任务或操作。此表达方式—被配置为[执行一个或多个任务]的[实体]—在本文中用于指代结构(即，物理的事物，诸如电子电路)。更具体地，此表达方式用于指示此结构被布置成在操作期间执行一个或多个任务。结构可被说成“被配置为”执行某个任务，即使该结构当前并非正被操作。“被配置为生成输出时钟信号的时钟电路”旨在覆盖例如在操作期间执行此功能的电路，即使讨论中的电路当前并非正被使用(例如，该电路并未连接到功率)。因此，被描述或表述成“被配置为”执行某个任务的实体是指用于实施该任务的物理的事物，诸如设备、电路、存储有可执行程序指令的存储器等等。该短语在本文中不被用于指代无形的事物。

术语“被配置为”并不旨在意指“可配置为”。例如，未经编程的FPGA不应当被认为是“被配置为”执行某个特定功能，虽然其可能“可配置为”执行该功能。在适当编程之后，FPGA然后可被配置为执行该功能。

所附权利要求书中的表述结构“被配置为”执行一个或多个任务明确地旨在对该权利要求要素不援引35U.S.C.§112(f)。于是，所提交的本申请中没有任何权利要求旨在要被解释为具有装置-加-功能要素。如果申请人在申请过程期间想要援引节段112(f)，则其将使用“用于”[执行功能]“的装置”结构来表述权利要求要素。

如本文所用，术语“基于”用于描述影响确定的一个或多个因素。此术语不排除可能有附加因素可影响确定。也就是说，确定可仅基于指定的因素或基于所指定的因素及其他未指定的因素。考虑短语“基于B确定A”。此短语指定B是用于确定A的因素或者B影响A的确定。此短语并不排除A的确定也可基于某个其他因素诸如C。此短语也旨在覆盖A仅基于B来确定的实施方案。如本文所用，短语“基于”与短语“至少部分地基于”是同义的。

具体实施方式

本公开初始参考图1A-图1B描述了图形处理流程和示例性图形单元的概览。参考图2-图3讨论了用于保持存储器一致性的示例性图形存储器分级结构和技术。参考图4A-图5讨论了中间渲染计算技术。参考图6-图7讨论了本地图像块数据结构和支持硬件。参考图8-图9讨论像素资源的同步。图10示出了示例性设备，图11示出了示例性计算机可读介质。相对于传统技术，各种所公开的实施方案可改善图形性能并降低功率消耗。

图形处理概述

参见图1A，即示出用于处理图形数据的示例性处理流程100的流程图。在一个实施方案中，转换和照明步骤110可涉及基于限定的光源位置、反射率等处理从应用程序接收的顶点的照明信息，将顶点组装成多边形(例如三角形)，以及/或者基于三维空间中的位置将多边形转换为正确的尺寸和取向。剪辑步骤115可涉及丢弃在可视区域之外的多边形或顶点。光栅化步骤120可涉及在每个多边形内限定片段并且例如基于多边形顶点的纹理坐标来为每个片段分配初始色值。片段可指定它们重叠的像素的属性，但可基于组合多个片段(例如，在帧缓冲器中)和/或忽略一个或多个片段(例如，如果它们被其他对象覆盖)来确定实际像素属性。着色步骤130可涉及基于照明、阴影、隆起映射、半透明度等来改变像素分量。可将着色像素组装在帧缓冲器135中。现代GPU通常包括允许应用开发器定制着色和其他处理步骤的可编程着色器。因此，在各种实施方案中，图1A的示例性步骤可以各种顺序执行、并行执行或省略。还可实施另外的处理步骤。

现在参考图1B，示出了示出图形单元150的一个实施方案的简化框图。在图示实施方案中，图形单元150包括可编程着色器160、顶点管185、片段管175、纹理处理单元(TPU)165、图像写入单元170、存储器接口180和纹理状态高速缓存190。在一些实施方案中，图形单元150被配置为使用可编程着色器160来处理顶点数据和片段数据两者，该可编程着色器可被配置为使用多个执行流水线或实例来并行处理图形数据。

在图示实施方案中，顶点管185可包括被配置为处理顶点数据的各种固定功能硬件。顶点管185可被配置为与可编程着色器160通信，以便协调顶点处理。在图示实施方案中，顶点管185被配置为将经处理的数据发送至片段管175和/或可编程着色器160以用于进一步处理。

在图示实施方案中，片段管175可包括被配置为处理像素数据的各种固定功能硬件。片段管175可被配置为与可编程着色器160通信，以便协调片段处理。片段管175可被配置为在来自顶点管185和/或可编程着色器160的多边形上执行光栅化以生成片段数据。顶点管185和/或片断管175可耦接到存储器接口180(未示出耦接)以便访问图形数据。

在图示实施方案中，可编程着色器160被配置为接收来自顶点管185的顶点数据和来自片段管175和/或TPU 165的片段数据。可编程着色器160可被配置为对顶点数据执行顶点处理任务，该顶点处理任务可包括顶点数据的各种转换和/或调整。在例示的实施方案中，可编程着色器160还被配置为对像素数据执行片段处理任务，诸如像纹理和着色处理。可编程着色器160可包括用于并行处理数据的多个执行实例。

在例示的实施方案中，TPU 165被配置为调度来自可编程着色器160的片段处理任务。在一些实施方案中，TPU 165被配置为预先提取纹理数据并将初始颜色分配给片段，以供可编程着色器160进一步处理(例如，经由存储器接口180)。TPU 165可被配置为提供例如在规格化整数格式或浮点格式的片段分量。在一些实施方案中，TPU 165被配置为提供呈2x2格式的四个(“片段四元组”)组的片段，该片段由可编程着色器160中的一组四个执行流水线处理。

在一些实施方案中，图像写入单元(IWU)170被配置为存储处理图像图块并且可对渲染的图像执行操作然后传输用于显示或存储用于存储。在一些实施方案中，图形单元150被配置为执行基于图块的延迟渲染(TBDR)。在基于图块的渲染中，可单独处理屏幕空间的不同部分(例如，像素的正方形或矩形)。在各种实施方案中，存储器接口180可促进与各种存储器分级结构中的一个或多个的通信。

在各种实施方案中，可编程着色器诸如可编程着色器160可以各种适当构型中的任一种耦接到图形单元中的其他可编程和/或固定功能元件。图1B的示例性实施方案示出了图形单元150的一个可能的配置用于示例性目的。

示例性存储器分级结构

图2是示出根据一些实施方案的示例性存储器分级结构的框图。在例示的实施方案中，图形单元150包括本地存储器230、可编程着色处理内核260、L1指令高速缓存210、L1数据高速缓存220、IWU高速缓存PL1 275(在IWU 170中)、TPU L1高速缓存TL1 285(在TPU165中)和L2高速缓存250。如图所示，在一些实施方案中，L1高速缓存是分体式的。

在一些实施方案中，本地存储器230被实现为USC 160中的专用随机存取存储器(RAM)并且与全局存储器分离。在一些实施方案中，本地存储器用于实现本地图像块(LIB)存储装置245和线程组范围内的存储器，而全局存储器用于实现专用存储器，设备存储器和系统存储器。下面将更详细地讨论这些不同的存储器空间。在一些实施方案中，不同的存储器空间由应用程序编程接口(API)支持。在例示的实施方案中，L1高速缓存210、220、275和285以及L2高速缓存250包括在全局存储器具体实施中。在例示的实施方案中，全局存储器中存在多个路径，这取决于访问是缓冲，图像读取还是图像写入(例如，图像读取可由TPU165至TL1 285处理，并且图像写入可由IWU 170至PL1 275处理)。

在例示的实施方案中，本地图像块存储装置245是包括在本地存储器230中的存储器空间。然而，在各种实施方案中，本地图像块存储器空间可以或可以不被认为包括在本地存储器空间。在一些实施方案中，本地存储器230是片上存储器，其包括在与可编程着色器处理内核260相同的集成电路上。

TPU 165和IWU 170可被描述为协处理器，并且其相应的L1高速缓存可用于仅存储图像数据(例如，纹理、经处理的图块等)。一般来讲，TPU 165可用于读取图像数据和IWU以写入图像数据。因此，在一些实施方案中，TPU L1高速缓存285可以是只读高速缓存。在例示的实施方案中，用于数据的L1高速缓存水平在L1数据高速缓存220、IWU L1(PL1)高速缓存275和TPU L1(TL1)高速缓存285之间是分体式的，该高速缓存是不同的高速缓存。因此，为了实现L1高速缓存220、275和285之间的存储器一致性，在例示的实施方案中，L2高速缓存250是连贯点。在一些实施方案中，TPU 165和/或IWU 170被配置为执行图形程序而不是着色处理内核260中的某些指令。在一些实施方案中，由着色器处理内核260执行图形指令可使得命令被发送至这些单元，而不是这些单元实际执行图形指令。

在一些实施方案中，L1数据高速缓存220是回写高速缓存，由此使得结果不写入L2高速缓存250，直到将要修改或替换为新数据为止。在其他实施方案中，L1高速缓存可利用各种写入丢失策略中的任一种来实现其他高速缓存方案，诸如写入等。在一些实施方案中，L2高速缓存250可用于跨图形单元150的处理元件。在一些实施方案中，可实现比L1高速缓存较低水平的一个或多个高速缓存以供着色器处理元件使用。

在一些实施方案中，硬件架构支持三个硬件存储器空间(本地、全局和指令)，而API为私有、LIB、线程组、设备、系统存储器和调试指针提供指针类型，这些指针类型可被转换到硬件存储器空间。

在一些实施方案中，本地存储器230是API限定的线程组存储器空间的物理存储器实现方式，并且可访问以计算内核和片段处理任务(例如，分别来自计算数据主调度器和像素数据主调度器)，并且可能不可访问以用于其他任务，诸如顶点处理。线程组中的所有工作项(分配给同一着色器处理元件的一组线程并被调度为共享存储器上下文)在本地存储器中看到相同的分配。在一些实施方案中，对于片段处理，本地存储器是图块范围的，由此使得对应于图块的所有线程看到相同的本地存储器分配。如下文进一步详细讨论的，这些线程可包括片段线程和中间渲染计算线程。需注意，也可将线程分配到相同的单指令多数据(SIMD)组，该组可包括线程组中的线程的一部分，其中SIMD组中的线程执行相同的指令(在具有预测执行的实施方案中，被预测为关闭的指令除外)并且被调度为使用并行硬件并行执行。在不同实施方案中，分配给SIMD组的线程的数量可以变化。

在一些实施方案中，本地存储器230使用包括本地存储器分配标识符的虚拟地址来访问。该虚拟地址可由可编程着色器160转换为物理地址。本地存储器分配可能是持续的，直到与特定线程组或图块相关联的所有线程完成为止。在一些实施方案中，本地存储器230用于两个支持API的存储器空间：线程组存储器(通过负载和存储器访问)和image_block存储器(其经由本地图像块读取和写入访问，下文将参考图6更详细地讨论)。

在一些实施方案中，全局存储器支持全局缓冲加载和存储(包括对专用和调试存储器的堆栈访问)、全局图像读取和写入(包括纹理样本操作)和指令提取。全局缓冲器访问可以将全局存储器视为字节可寻址的，并且可以是弱有序的(由此使得从单个线程来看，访问似乎按程序顺序发生，但不能保证不同线程之间的存储器排序)。下面讨论了用于明确指定排序以实现存储器一致性的附加技术。除非使用围栏指令(如下文更详细地讨论的)，否则即使在同一线程内也不能对全局图像访问进行排序。

在例示的实施方案中，有四条从着色器内核到全局存储器的不同路径：

1)通过用于缓冲加载和存储的L1数据高速缓存220

2)通过TPU 165进行图像读取和纹理样本操作

3)通过IWU 170进行图像写入

4)通过用于指令访问的L1指令高速缓存210

在图示实施方案中，这些不同路径全部会聚在L2高速缓存中。在一些实施方案中，L1级高速缓存被标记为对应的数据原模(像素、顶点或计算)的ID，其被映射到完整的上下文ID。在这些实施方案中，当改变上下文ID和数据原模之间的映射时，L1高速缓存可被刷新/失效。

在一些实施方案中，指令可包括字段以指定高速缓存行为。这些字段可被称为高速缓存控制字段，并且可包括在存储器访问指令中并与高速缓存行相关联。可由此类高速缓存控件指定的行为的实施例包括：(a)存储在L1高速缓存中并存储在L2高速缓存250中；(b)读入/写入刷新失效L1高速缓存，存储在L2高速缓存250中；以及(c)读入/写入刷新失效L1高速缓存和L2高速缓存250C两者。在一些实施方案中，每个物理高速缓存行可被标记为包含高速缓存数据(干净或脏)或读入/写入(RTI/WTI)数据。RTI/WTI模式可允许高速缓存用作存储数据的缓冲器，直到遇到同步指令，此时读入数据或写入数据分别从高速缓存中失效或从高速缓存中刷新。但是，传统的刷新操作可能不影响WTI/RTI数据。对于高速缓存控制字段，cacL1_cacL2值可指示存储在L1高速缓存和L2高速缓存中，bypL1_cacL2值可指示存储在L2中并且读入/写入刷新失效L1，并且bypL1_bypL2值可指示读入/写入刷新失效L1和L2两者。在一些实施方案中，高速缓存控制字段可使用两位来编码这三个值。在其他实施方案中，可使用更多的位来跟踪附加状态，例如，对于高速缓存分级结构中的更大数量的级别。

各种指令可基于其高速缓存控制字段以不同方式对高速缓存行进行操作。在一些实施方案中，刷新指令和失效指令影响标记为高速缓存的L1高速缓存中的行，而不影响标记为RTI或WTI的那些行(例如，如“byp”标签所指示)。在一些实施方案中，缓冲器数据高速缓存刷新(bdcf)指令将数据从L1数据高速缓存220中刷新。在一些实施方案中，缓冲数据高速缓存失效清除(bdcic)使L1数据高速缓存220失效和清除。在一些实施方案中，图像数据高速缓存刷新(idcf)指令从IWU L1高速缓存275中刷新数据。在一些实施方案中，图像数据高速缓存失效清除操作使TPU L1高速缓存285中的高速缓存数据失效(在一些实施方案中，这是不能容纳脏数据的只读高速缓存，因此该操作仅仅是失效的)。

例如，在一些实施方案中，遵循借助cacL1_cacL2执行的存储的bdcf指令应当导致存储数据从L1刷新到L2。另一方面，在这些实施方案中，在遵循借助bypL1_cacL2执行的存储的bdcf指令不应当影响存储数据，即使数据作为WTI数据存在于L1高速缓存中。相反，可使用syncbuf l2c.release指令(下文将进一步详细讨论)来访问存储数据，该指令应当迫使未完成的L2为目标的存储数据在L2高速缓存中排序(通过将存储数据从L1写到L2)。

这是RTI和WTI技术可允许L1高速缓存用作缓冲器的一种机制，其中数据存储在高速缓存中直到同步操作，但不被认为处于正常的高速缓存状态。在一些实施方案中，这可通过将写入分组到存储器分级结构中较高的层级而不是随时间推移将其扩展来减少存储器流量。

在一些实施方案中，指令可指定是否应当在L2高速缓存250(和/或其他高速缓存)中以低持续性分配数据。当已知(例如，如由编译器所指定的)其在不久的将来将不被使用时，这可允许标记数据进行替换。在一些实施方案中，使用指示高或低持续性的单个位来实现持续性字段。

在一些实施方案中，系统存储器空间在全局存储器的顶部上实现，并且地址空间被对齐，以使得可使用身份映射来执行系统存储器指针到全局存储器指针的转换。在一些实施方案中，系统存储器允许图形单元150以启动粒度与系统中的其他处理设备共享数据。在一些实施方案中，图形单元150的固件促进这种分享的定序。

在一些实施方案中，设备存储器空间也在全局存储器的顶部上实现。在一些实施方案中，设备存储器和全局存储器的地址空间完全对齐。在一些实施方案中，设备存储器可供在图形着色器上执行的所有线程访问。因此，在一些实施方案中，L2高速缓存250是设备存储器的连贯点(这是在图2的实施方案中，所有存储器访问路径一起到存储器分级结构中的公共点的最低级别)。在一些实施方案中，系统存储器还可用于其他非GPU处理元件，诸如CPU。因此，系统存储器可以是GPU和CPU的连贯点。

在一些实施方案中，线程地址空间包括专用实例堆栈存储器和调试存储器。这些空间可映射到GPU驱动器所提供的并且经由硬件寄存器(例如，使用基地址和大小)指定的全局存储器的连续块。

示例性存储器同步和一致性

在一些实施方案中，syncbuf和syncimg指令分别允许同步缓冲器操作和图像操作。在一些实施方案中，这些指令包括指定(a)目标(例如，L1高速缓存，L2高速缓存，或设备存储器)和(b)同步指令是否获取和/或释放存储器副作用的字段。这些操作可以指定针对存储器分级结构的特定级别的所有先前存储已经在分级结构的该级别被排序，然后是针对该级别的任何后续加载或存储。L1高速缓存刷新操作可被视为针对下一高速缓存级别的存储。在这种情况下，排序相对于单个线程的程序顺序。因此，同步指令可不在线程之间强制执行全局顺序，除非与屏障操作结合使用，例如用于识别全局顺序(屏障指令将在下文进一步详细讨论)。在一些实施方案中，为了完全同步所有加载或存储，可能需要针对分级结构的每个水平(例如，在所公开的实施方案中，L1、L2和设备水平)的同步指令。

在一些实施方案中，存在可应用于原子操作或用作围栏操作的多个不同的存储器属性。在一些实施方案中，这些属性可触发适当的同步指令。这些属性的实施例包括：松弛、获取、释放、acq_rel和seq_cst。可以使用API中的语法“memory_order”来引用这些标志，例如，如“memory_order_relaxed”中那样。在一些实施方案中，memory_order_relaxed意味着无顺序约束。在一些实施方案中，memory_order_acquire从具有与获取同步的释放属性(下文讨论)的操作获得副作用(例如，来自存储)。因此，如果获取与释放同步，则获取执行单元将看到释放之前的所有副作用(以及可能的后续副作用)。在一些实施方案中，memory_order_release释放对具有获取属性的操作的副作用，并且释放前的所有副作用均包括在该释放中。因此，通常可使用释放属性来同步先前的存储，并且获取属性通常可用于同步后续加载。在一些实施方案中，memory_order_acq_release同时具有操作的获取和释放属性。在一些实施方案中，memory_order_seq_cst表示每个执行单元的加载和存储表现出按程序顺序执行，而来自不同执行单元的加载和存储看起来只是交错的。换句话讲，顺序一致的排序对如此标记的所有原子操作实施单一的总修改顺序。这意味着，一旦不使用顺序一致的存储器排序的操作进入该画面，那么可能丢失顺序一致性。顺序排序可用于多个生产者/多个消费者情况，其中消费者应当观察以相同顺序发生的所有生产者的动作。在一些实施方案中，访问变量的其他线程(例如，除具有特定获取和释放对的线程之外)可看到不同的存储器访问顺序。换句话讲，对于未明确指定诸如上述那些属性的属性的线程，可能不保持存储器一致性和排序。

在一些实施方案中，还可针对上述存储器属性指定存储器范围。范围参数的示例值可以包括线程、线程组和设备。在一些实施方案中，memory_scope_thread在线程内应用存储器排序约束(这可仅用于纹理围栏，在一些实施方案中)。在一些实施方案中，memory_scope_threadgroup对在线程组中执行的线程应用存储器排序约束。在一些实施方案中，线程组范围可在L1数据高速缓存220处实施，因为L1数据高速缓存220在着色器处理内核260之间共享，其中一者可执行线程组的不同线程。在一些实施方案中，memory_scope_device将存储器排序约束应用于在设备上执行的所有线程，其可包括由TPU 165、IWU 170等执行的线程。因此，在一些实施方案中，设备范围可在L2高速缓存250处实施。可以使用应用于存储器操作的同步指令和高速缓存控制的组合来实现存储器范围属性。在一些实施方案中，这些属性也可用于围栏操作，其可在整个围栏上(在指定的范围处)对所有原子操作和非原子操作之间强加排序。

例如，下表1示出了使用高速缓存控制和同步指令实现不同范围和存储器排序属性的示例。具体地讲，表1包括对全局存储器(具有设备范围的变量)的原子返回操作的扩展：

表1-对全局存储器(设备范围的变量)的原子返回操作的扩展

在表1中，“gm”是指全局存储器(其中“lm”是指本地存储器，例如，对于线程组范围的变量)。如表1所示，在不同的范围使用同步指令来在线程组级别或设备级别强制执行一致性。在表1的实施例中，设备范围的同步在L2高速缓存250处进行，而线程组范围的同步在L1数据高速缓存220处进行。此外，设备范围的扩展使用L1数据高速缓存220的写入高速缓存控制，而线程组范围的扩展不使用。请注意，线程组范围的操作可能不需要同步操作，但可能例如仅作为atomic.lm.OP来实现。

在一些实施方案中，不返回的原子操作可仅具有排序语义的子集。具体地讲，由于这些操作不提供加载功能，因此它们可能不具有获取语义。因此，acq_rel和seq_cst可以退化以释放这些操作的排序。

在一些实施方案中，类似的属性可用于指定要在屏障上应用的排序。在一些实施方案中，“屏障”操作允许执行同步，例如，线程组中需要与代码流中的相同位置同步的多个线程(例如，已到达屏障的线程可被停止，直到其他线程赶上)。在一些实施方案中，使用“围栏”操作来强制在围栏之前和围栏之后的存储器操作之间的关系(例如，在围栏之前的所有存储对于一些围栏操作的围栏之后的负载必须是可见的)，但可能不同步指令流。

因此，在一些示例性实施方案中，以下属性可用于包括围栏的屏障：mem_none；mem_device；mem_threadgroup；mem_threadgroup_imageblock；和mem_texture。在一些实施方案中，mem_none指定操作不提供存储器围栏。在一些实施方案中，mem_device指定所有memory_scope_device缓冲器加载和存储横跨屏障排序。在一些实施方案中，mem_threadgroup指定跨屏障对所有memory_scope_threadgroup缓冲器加载和存储进行排序。在一些实施方案中，mem_threadgroup_imageblock指定在整个屏障上对所有memory_scope_threadgroup本地图像块加载和存储进行排序。在一些实施方案中，mem_texture指定所有memory_scope_device和memory_scope线程组图像加载和存储在整个屏障上被排序。

对于屏障操作，可指定代码同步的屏障范围和存储器排序属性两者。表2列出了针对不同屏障范围和存储器排序属性的示例性扩展。可以将多个标志与给定屏障指令一起使用并且所述属性可以是添加的。

在一些实施方案中，线程组范围的屏障意味着在允许任何线程在屏障之外继续执行之前，线程组中的所有线程必须执行线程组范围的屏障。类似地，在一些实施方案中，SIMD组范围的屏障是指在允许任何线程在屏障之外继续执行之前，SIMD组中的所有线程必须执行SIMD组范围的屏障。本文参考屏障所述的各种存储器定序功能可使用非屏障围栏操作来实现，例如，无需屏障的代码同步。

表2—屏障操作的扩展

例如，用于threadgroup_barrier(mem_device)的扩展应当为syncbufl1c.realease；屏障，如表2的第一行所示。在一些实施方案中，具有mem_none或无标志的屏障仅发出屏障操作。在一些实施方案中，如果对屏障指定多个存储器属性，则可实现来自表2的多个扩展。

为了举例说明的目的，包括存储器属性、范围参数、操作类型、特定高速缓存分级结构等，但并不旨在限制本公开的范围。在其他实施方案中，类似的技术可用于在弱有序图形处理器中实施特定的存储器一致性和排序特征。

用于存储器一致性的示例性方法

图3是示出根据一些实施方案的用于强制执行存储器一致性的方法300的一个示例性实施方案的流程图。除了其他设备之外，图3所示的方法可结合本文公开的计算机系统、装置、元件或部件中的任一者来使用。在各种实施方案中，所示的方法要素中的一些可按与所示次序不同的次序并发执行，或者可被省去。也可根据需要执行附加的方法要素。流程开始于310。

在310处，在例示的实施方案中，图形单元150处理指定存储器访问属性和范围信息的第一指令。在一些实施方案中，存储器访问属性指定要针对由第一指令指定的操作实施的顺序一致性的类型。如上所述，在一些实施方案中，顺序一致性的类型包括：无、获取、释放、获取+释放和顺序一致。在第一指令是屏障操作的情况下，可暗示存储器访问属性(指定范围处的加载和存储通过屏障排序)。如上所述，范围的实施方案可包括线程、线程组和设备。对于屏障操作，也可指定屏障自身的附加范围。

在320处，在图示的实施方案中，图形单元150根据由范围信息指定的存储器分级结构中的高速缓存级别处的存储器访问属性强制执行指定类型的顺序一致性。例如，在图2的实施方案中，可在L1级实施“线程组”范围，而“设备”范围可在L2级实施。在一些实施方案中，对具有指定的相同范围的操作实施具有释放属性和获取属性的操作之间的一致性。例如，具有释放属性和线程组范围的第一操作可与具有获取属性和线程组范围的第二操作同步，由此使得第一操作之前在L1级别的存储器副作用对第二操作可见。在一些实施方案中，为存储器分级结构中的分割级别的不同部分指定不同的范围(例如，在一些实施方案中，用于L1级的图像和缓冲器同步的不同范围)。

示例性命令缓冲器执行

图4A是示出根据一些实施方案的图形命令缓冲器425A-425N的执行的示例性图。在图示实施方案中，命令流420A-420C对应于由不同数据原模(计算数据原模，顶点数据原模和像素数据原模)分配的工作。每个数据原模可处理对可编程着色器160的部分的相应工作的调度。

在图示实施方案中，每个命令缓冲器425包括一组一个或多个启动。如图所示，给定命令缓冲器可包括来自多个数据原模420的启动。另外每个数据原模可包括相同命令缓冲器中的多个启动。例如，在例示的实施方案中以及相对于“cbuf 0,”启动“k1”和启动“k2”分别在命令流420A和420B上执行，另外两次启动分组在单个命令缓冲器“cbuf0”中。

在一些实施方案中，“启动”是来自可编程着色器160的其他部分的视角的原子工作单元。因此，在一些实施方案中，从在启动开始时在多个着色器处理元件(例如，L1数据高速缓存220)之间共享的存储器读取给定启动所需的任何数据，并且在启动结束时将结果写回共享存储器。因此，其他硬件在完成启动之前看不到启动的结果，在该点，结果在共享存储器中可用并且可被其他启动(包括来自其他数据原模的启动)访问。启动可包括一组一个或多个渲染命令，所述一组一个或多个渲染命令可包括用于绘制过程几何形状的命令、用于设置阴影采样方法的命令、用于绘制网格的命令、用于检索纹理的命令、用于执行生成计算的命令等。可在帧的渲染期间在各种阶段之一执行启动。渲染阶段的实施例包括但不限于：相机渲染、光渲染、投影、纹理、片段着色等。

在一些实施方案中，响应于完成每个命令缓冲器的执行，发出刷新请求以刷新和使L2高速缓存250的一部分或全部失效。因此，在命令缓冲器完成时，数据可用于其他处理元件，诸如CPU。在一些实施方案中，在给定命令缓冲器最后启动之后出现刷新；然而，在其他实施方案中，可在执行命令期间在其他时段出现刷新。从包括待刷新的特定命令缓冲器的程序请求刷新事件。除了将与已完成的给定命令缓冲器相关联的高速缓存行作为目标之外，刷新还可使得刷新控制器基于附加信息来刷新和使其他高速缓存行失效。

在一些实施方案中，响应于给定启动的最后执行，结果数据从本地存储器写入到共享存储器(例如，到L1数据高速缓存220和/或L2高速缓存250)。这可允许其他着色硬件(或执行另一启动的相同着色硬件)访问启动的结果。对共享存储器的这些访问可能是性能瓶颈并且可能消耗大量的电力。例如，如果需要在两个相关像素启动之间执行计算任务，则在一些具体实施中，第一像素启动必须将数据写入共享存储器，以便由计算任务的计算启动访问，其继而必须将数据写入共享存储器以便被第二像素启动访问。因此，在一些实施方案中，中间渲染计算机技术允许在未写入共享存储器的情况下在像素启动的上下文中执行计算任务，如下文进一步详细讨论的。

示例性中间渲染计算技术

图4B是示出命令缓冲器的执行的示意图，其中在像素渲染启动的上下文内执行辅助计算任务。这些计算任务可被称为“图块着色器”，并且在一些基于图块的渲染实施方案中可为图块范围。在各种实施方案中，这可避免需要将数据从片段处理刷新到存储器以执行辅助计算任务，这可降低功率消耗和/或改善性能。换句话讲，在一些实施方案中，这可允许图形处理的数据在片上停留更长时间，而不是在计算任务之前和之后写入片外。

在一些实施方案中，辅助计算任务访问本地存储器空间和/或本地图像块空间中的数据结构，诸如像素渲染启动的图块渲染目标存储器(例如，像素输出寄存器)、像素渲染启动的本地图像块和/或像素渲染启动的其它本地存储器结构以访问或共享数据。如上所述，本地存储器对于像素渲染启动可为图块范围的。当在用于中间渲染内核的线程组或threadgroup_imageblock地址空间中分配变量时，可以为访问内核的每个线程组分配这些变量，并且对于给定的图块，可以在函数(中间渲染计算和片段函数)之间保持持续。在一些实施方案中，API(例如，金属API、OpenGL、DirectX、Glide、WebGL等)允许变量作为参数传递或在特定的存储器空间中声明，例如，代码“threadgroup float x；”可以在线程组地址空间中分配浮点数。

本文讨论的启动是“渲染传递”的示例性具体实施，其旨在包括其在本领域中的熟知含义，其中包括在不将所生成的数据刷新到共享存储器的情况下执行的一组操作(因此，渲染传递期间执行的操作对于不执行渲染传递的其他着色器处理元件是原子的，并且当在渲染传递完成时将数据写入共享存储器时，渲染传递的结果可供其他处理元件使用)。因此，本文对启动的各种引用均可被认为是对渲染传递的引用。渲染传递自身可包括各种着色器操作。渲染传递可包括例如可部分地重叠的多个混合操作。

出于解释的目的，中间渲染计算任务的实施例包括但不限于：生成光列表、生成纹理图、解析多采样图像数据、图像统计信息收集(例如，直方图)、减小图像区域的深度边界等。

纹理映射是存储图像序列的一种公知技术，其中所述序列中的每个图像为相同图像相对于序列中的先前图像较低分辨率的表示。较低分辨率的纹理图水平可用于在场景中更远出现的对象。纹理映射可用于详细程度的细节技术、改善图像质量、减少渲染时间(例如，通过减少采样以渲染每个像素的纹理的数量)和/或减少图形单元150上的加载。为了生成每个纹理图图像，图形单元150可在先前水平上过滤更高分辨率图像以生成较低分辨率版本的图像。

因此，在一些中间渲染计算实施方案中，生成纹理图水平的滤波可由在像素启动的上下文中执行的辅助计算代码来执行，然后该辅助计算代码可在生成纹理图之后利用一个或多个纹理图水平进行片段操作。这可允许纹理图(针对正在处理的图块)在本地存储器中生成和用于片段处理任务而不访问共享存储器直到完成像素启动。

在一些实施方案中，处理光列表是可在中间渲染执行的另一个计算任务。例如，给定场景可具有大的光列表，但光的仅较小部分可实际影响正在处理的当前图块。因此，生成影响该图块的光的列表是可在渲染期间执行的辅助计算任务，而无需访问共享存储器，直到渲染传递完成。然后可将光列表用于正在执行启动的其余部分进行片段处理。

包括中间渲染计算任务的这些实施例是为了解释的目的，而并非意图要限制本实施方案的范围。在其他实施方案中，可在包括片段着色器的渲染中间执行各种示例性计算任务中的任一个。一般来讲，计算任务可以比着色器传递的片段着色器更高的粒度在屏幕空间中执行(例如，在一些实施方案中，计算任务是图块范围的，而片段着色器是片段范围的，对其他像素的数据的访问受限)。

示例性中间渲染计算方法

图5是示出根据一些实施方案的用于执行中间渲染辅助计算任务的方法500的一个示例性实施方案的流程图。除了其他设备之外，图5所示的方法可结合本文公开的计算机系统、装置、元件或部件中的任一者来使用。在各种实施方案中，所示的方法要素中的一些可按与所示次序不同的次序并发执行，或者可被省去。也可根据需要执行附加的方法要素。流程开始于510。

在510处，在例示的实施方案中，图形单元150执行渲染传递，包括从共享存储器访问图形数据。在一些实施方案中，包括在渲染传递的初始阶段从共享存储器访问数据和在渲染传递的最后阶段将结果写入共享存储器，并且不以其他方式访问共享存储器。可基于渲染传递本身中的指令执行或不执行对共享存储器的访问(例如，可使用其他命令通过访问共享存储器来设置渲染传递)。

在520处，在例示的实施方案中，图形单元150执行辅助任务以生成数据并将数据存储到本地存储器空间。在一些实施方案中，辅助任务还从本地存储器空间读取数据(例如，由渲染传递的先前部分生成)。在一些实施方案中，本地存储器空间可由执行渲染传递的第一着色器处理元件访问，但不可访问一个或多个其他着色器处理元件。在一些实施方案中，辅助任务根本不访问共享存储器，但辅助任务的结果可在渲染传递结束时写入共享存储器。

在530处，在所示的实施方案中，图形单元150继续渲染传递，包括访问辅助任务所存储的数据。在例示的实施方案中，渲染传递在辅助任务之前不写入共享存储器(在一些实施方案中，渲染传递不写入共享存储器，直到渲染传递完成为止)。

示例性本地图像块数据结构

图6示出了根据一些实施方案的被配置为将数据存储在本地存储器230中的示例性柔性数据结构。这种类型的支持硬件的结构在本文中被称为本地图像块(LIB)。在一些实施方案中，LIB数据结构的可配置方面促进上述中间渲染计算技术。

在例示的实施方案中，LIB的尺寸可在四个维度上配置：X方向(例如，借助像素单位)、Y方向(例如，借助像素单位)、每个像素的样本数和深度(每个样本的数据量)。在一些实施方案中，着色器程序包括通过写入图形单元150中的硬件寄存器来声明LIB结构并配置其尺寸的指令。因此，在例示的实施方案中，LIB结构的大小(以字节为单位)为X乘以Y乘以2^{samples_per_pixel}乘以深度，其中深度以每个样本的字节为单位。在一些实施方案中，LIB可被配置为附加维度(例如，三维显示器诸如全息图的Z维度)。

在一些基于图块的实施方案中，对于片段处理，X维度和Y维度被设置为对应于正在处理的图块的尺寸。对于诸如计算任务的其他任务，可为特定任务定制X维度和Y维度(例如，对于行缓冲器任务，X维度中的条目数大于Y维度，反之亦然)。在一些实施方案中，例如，每个像素的样本数量基于多样本抗混叠(MSAA)模式，并且可在每像素的单个样本至每个像素的四个样本(或更多)的范围内。在一些实施方案中，每像素存储多个样本也可用于其中不使用MSAA的其他图形任务。可调节深度以存储每个样本的不同数量的颜色附件或其他像素数据。LIB可被配置为实现例如g缓冲器，和/或其他类型的缓冲器。

术语“切片”可用于指图像块中的区域，该区域描述图像块中针对所有像素位置的给定元素的值(例如，所有X、Y和样本值的LIB的一部分)。

在一些实施方案中，片段函数可以作为颜色附件或使用在着色器中声明的显式布局来访问LIB切片。在一些实施方案中，显式布局元素可使用标量、整数或浮点数据类型或某些像素数据类型。

在一些实施方案中，API为计算内核提供函数来声明现有LIB的本地图像块，查询属性(例如，宽度、高度、样本的数量、给定x/y坐标上的颜色数量，等等)，获取针对LIB中特定位置的图像数据块的参考，写入LIB位置，从LIB切片的区域读取(例如，将数据从LIB写到特定纹理区域)等。在一些实施方案中，API允许在内核和片段功能之间共享结构以促进像素和计算功能之间的共享。这可促进上述的中间渲染计算功能。在一些实施方案中，LIB可以例如通过改变LIB在一个或多个方向上的维度来至少部分地重新配置中间渲染。在一些实施方案中，GPU被配置为执行重新配置而不破坏LIB中的数据。

例如，对于中间渲染纹理图生成，考虑带有64×64图像、32×32图块和32×32LIB的情况。在该实施例中，第一纹理图水平可为32×32图像并且可存储在LIB的切片中。剩余的纹理图水平(16×16、8×8、4×4、2×2和1)可被存储在LIB的另一切片中(在典型纹理映射具体实施中，这些剩余水平的尺寸收敛至小于前一水平的尺寸)。

在一些实施方案中，对LIB的访问在线程内排序。在一些实施方案中，异常是直接存储器访问(DMA)操作，其可能不被排序。这种DMA操作的一个实施例是imgwblk指令。此操作可将数据从LIB复制到设备或系统存储器中的某个表面。在通过DMA操作从LIB复制的数据通过后续操作(例如，通过lstpix指令)访问的情况下，读取后写入操作应当用围栏操作来保护，如以下示例代码所示：

imgwblk

syncimg local.release

lstpix

但是，请注意，此代码序列保护单个线程中的imgwblk。为了保护线程组中可能写入LIB结构的其余线程，可使用屏障指令，如以下示例代码所示：

barrier(none)

mfsr r8 lin_id

icndst.eq r0,r8,#0,ADJ1//if(thread_index_in_threadgroup＝＝0)

imgwblk

syncimg local.release

cndend r0 ADJ1

barrier(mem_none)

lstpix

因此，在从LIB存储器空间写入共享存储器时，上述存储器一致性技术也可适用。

示例性LIB访问方法

图7是示出根据一些实施方案的用于使用柔性本地图像块的方法700的一个示例性实施方案的流程图。除了其他设备之外，图7所示的方法可结合本文公开的计算机系统、装置、元件或部件中的任一者来使用。在各种实施方案中，所示的方法要素中的一些可按与所示次序不同的次序并发执行，或者可被省去。也可根据需要执行附加的方法要素。流程开始于710。

在710处，在例示的实施方案中，图形单元150在一个或多个硬件存储元件中存储多个值，以指定可由一个或多个着色器处理元件访问的本地存储器中的数据结构的属性。本地存储器空间可专用于一组特定的一个或多个处理元素，并且不可供其他人访问。在图示的实施方案中，属性包括：数据结构在X维度的尺寸、数据结构在Y维度中的尺寸、每个像素的样本数量和每个样本的数据量。例如，可针对像素启动或中间渲染计算任务执行该步骤。

在720处，在所示实施方案中，图形单元150使用一个或多个请求访问数据结构，所述请求指定：一个或多个X坐标、一个或多个Y坐标、一个或多个样本和一个或多个样本之一的数据的索引。在一些实施方案中，图形单元150基于该请求确定本地存储器空间中的物理地址。如上所述，访问可由中间渲染计算任务到渲染传递所声明的LIB来进行。类似地，渲染传递中的片段着色器可访问中间渲染计算任务所声明的LIB中的数据。以下部分讨论了同步此类访问以及访问其他像素资源的技术。

示例性像素资源同步

一些GPU架构已被配置为以特定顺序执行片段/像素操作(以操作被提交到GPU的次序，其可被称为API提交次序)。对于不重叠的片段，排序无关紧要，因为处理将不影响其他片段。因此，一些架构被配置为生成片段的“传递组”，其中操作顺序无关紧要(例如，给定传递组中的片段通常在屏幕空间中不重叠，因此组内的操作顺序不应当彼此影响)。这些架构可包括硬件，以确保较旧的传递组在执行较新的传递组前完成，以保持API提交排序(也称为提取排序)。例如，美国专利申请公布No.2016/0055610讨论了基于传递组信息来调度图形任务的技术。

在一些实施方案中，不同于GPU隐式地确定像素操作的顺序，除非由编译器和/或应用程序开发器明确指定，否则不保证在程序顺序中发生操作。这可通过减少不需要的停滞而实现图形单元150的性能，同时仍管理由程序员指定的排序。它还可允许图形单元150消耗更少的功率，例如，因为可能不需要用于跟踪传递组的硬件。

在一些实施方案中，图形单元被配置为执行pixwait指令，该指令迫使着色器等待下叠片段(覆盖相同区域但较早提交的片段)以释放资源。类似地，“pixrelease”指令可将资源释放至重叠片段(覆盖相同区域但稍后提交的片段)。在一些实施方案中，图形单元150被配置为维护映射到不同像素资源的一组ID的状态信息(例如，驱动器可将ID映射到资源)。像素资源的实施例包括颜色附件(例如，在颜色缓冲器或g缓冲器中)\深度缓冲器\模板缓冲器等。深度缓冲器通常用于存储所生成像素的z坐标。这可通过不遮挡在其他像素后面的像素来提高性能。模板缓存器为可用于限制渲染区域的另一熟知的结构。在所公开的实施方案中，在一些实施方案中，除非使用pixwait和pixrelease指令明确同步，否则对这些资源的访问可能超出程序次序发生。在各种实施方案中，为了举例说明的目的，本文讨论的像素资源被提及，但并不旨在限制可使用明确指令进行同步的像素资源的类型。

图8包括框图，其示出了基于下叠片段信息来保持的线程的pixlock状态810。在一些实施方案中，针对每个线程维护pixlock状态。在例示的实施方案中，停滞电路系统820被配置为基于线程的pixlock状态、pixwait范围和一个或多个pixwait ID来确定是否停滞线程。在一些实施方案中，pixwait指令指定其中集合位指示等待对应ID的掩模。在这些实施方案中，如果设置了线程的对应pixlock位，则停滞电路系统820不允许线程继续。Pixrelease指令可在相应的资源被释放后清除pixlock位，此时停滞电路系统820可允许线程继续。

在一些实施方案中，pixwait指令还可指定范围，其可影响同步的粒度。例如，在一些实施方案中，该范围指定样本或像素粒度(注意，每像素可存在多个样本)。在一些实施方案中，也可在四边形粒度(例如，2×2像素)或甚至更大的粒度下指定范围。

图8还示出了ID与像素资源的示例性映射830。例如，该映射可由图形驱动程序生成。因此，对于不同的图形程序，状态位和像素资源之间的映射可以改变。在图示实施方案中，ID 0-3分别分配给颜色附件0-3和ID4被分配给深度和模板缓冲器(在该实施例中，单个ID被分配给这两个缓冲器，单独的ID可在其他具体实施和/或情况中使用)。示例性映射830用于图8的实施例中的代码样本840和850。

示例代码块840和850执行g缓冲器创建和光累积，两者之间具有同步指令。考虑示例性g缓冲器，其包括多种格式的四种颜色附件(例如，指定颜色累积、漫照率、遮光、法线X/Y、运动矢量、光谱功率、光谱强度等)。可以在单次传递技术中创建这种g缓冲器用于延迟照明。对于g缓冲器创建，在例示的实施方案中，代码840读取和写入模板和/或深度缓冲器，并且写入到g缓冲器中的颜色附件1-3。对于光累积，在图示实施方案中，代码850读取颜色附件1-3并写入颜色附件0。

在代码840中，ID0的pixrelease可导致pixlock状态中的对应位改变极性(如果先前被锁定)，指示颜色附件0现在可用。该版本初始可包括在代码中，因为代码840不写入颜色附件0。在代码840中，ID4上的pixwait确保在访问时缓冲器资源是最新的(基于来自任何下叠片段的输出)。然后，代码一旦写入缓冲器，就释放ID4。然后代码在存储到颜色附件1-3之前等待颜色附件同步(如果需要)。

在代码850中，可执行初始释放，因为该代码不写入颜色附件1-3。然后，代码将等待g缓冲器的其余部分(直到代码840中的“stpix[2][3]”指令才将被释放)然后读取g缓冲器，执行光累积，并且写入颜色附件0。如图所示，在一些实施方案中，等待或释放可指定多个ID，例如使用掩模编码。

在一些实施方案中，GPU被配置为针对给定资源保持至少两个单独标识符。在一些实施方案中，一个ID可被称为资源的READ ID，而另一个ID可被称为资源的WRITE ID。在这些实施方案中，修改资源时可使用以下方法：

pixwait READ ID|WRITE ID //等待prev.读取器和写入器

修改资源

pixel READ ID|WRITE ID //解锁未来读取器和写入器

在此实施例中，两个ID都用于等待和解锁特定资源。另外，在这些实施方案中，可使用以下方法从资源中读取而不修改资源：

pixrel WRITE ID //不阻止等待写入器的任何人

pixwait WRITE ID //等待先前写入器

读取资源

pixrel READ //解锁未来读取器

分割读和写标识符可以通过等待先前写入器完成允许阅读器同时工作(而未来写入器直到阅读器完成之后才开始)。此外，阅读器线程可在访问数据(例如，在对所访问的值执行计算之前)时立即释放READ ID。

如上所述，在不存在pixwait和pixrelease指令的情况下，在一些实施方案中，图形单元150未被配置为在对像素资源的访问之间强制执行排序。这可允许图形单元150通过执行程序次序(相对于强制执行传递组排序的GPU)来降低复杂性并提高性能，从而降低功率消耗。

在一些实施方案中，所公开的像素同步技术还促进中间渲染计算技术。例如，pixwait和pixrelease指令可用于同步对用于在中间渲染计算内核和片段着色器线程之间共享数据的资源的访问。例如，如果中间渲染计算任务正在本地图像块或光列表中生成纹理图，则片段着色器可能需要等待访问所生成的数据，直到中间渲染计算任务完成。注意到片段着色器通常只能访问它们正处理的片段的数据，而中间渲染计算任务可查看对应于图块中任何位置的数据。因此，像素同步也促进在该混合粒度上下文中进行排序(尽管不限于中间渲染计算)。在一些实施方案中，较高水平的粒度可用于例如通过使用pixwait的图块范围来进行中间渲染计算同步。

示例性像素资源同步方法

图9是示出根据一些实施方案的用于使对像素资源的访问同步的方法900的一个示例性实施方案的流程图。除了其他设备之外，图9所示的方法可结合本文公开的计算机系统、装置、元件或部件中的任一者来使用。在各种实施方案中，所示的方法要素中的一些可按与所示次序不同的次序并发执行，或者可被省去。也可根据需要执行附加的方法要素。流程开始于910。

在910处，在例示的实施方案中，图形单元150使用一个或多个像素资源进行第一图形操作。

在920处，在例示的实施方案中，图形单元150等待(如等待指令所指定)将针对第二图形操作使用一个或多个像素资源。等待可基于像素资源的一个或多个标识符和由等待指令指定的范围。对于每个资源，可以使用单独的标识符来不同地同步只读访问和修改访问。

在930处，在例示的实施方案中，图形单元150使用释放指令来指示一个或多个像素资源的释放(这发生在针对第一图形操作使用一个或多个像素资源之后)。在具有分割的READ和WRITE标识符的实施方案中，这可能仅为未来的读取器(例如，使用READ ID)或针对未来的读取器和写入器(例如，使用READ ID和WRITE ID两者)释放资源。该方法元素可包括基于释放来更新线程的pixlock状态寄存器。

在940处，在所示实施方案中，图形单元150响应于释放指令，针对第二图形操作使用一个或多个像素资源。这可包括读取资源和/或修改资源。

在950处，在例示的实施方案中，在不存在等待指令和释放指令的情况下，图形单元150以API提交顺序执行一个或多个其他图形操作。这可提高性能并降低对于其排序不重要的操作的功率消耗。

在一些实施方案中，编译器可被配置为生成与本文所讨论的任何指令，包括用于中间渲染计算、LIB访问、存储器一致性、像素同步等的指令。编译器本身可存储在非暂态计算机可读介质上。此外，本文讨论的各种计算机程序中的任一种可被存储非暂态计算机可读介质中用于一个或多个处理元件执行。

示例性设备

现在参考图10，其示出了例示设备10的示例性实施方案的框图。在一些实施方案中，可将设备1000的元件包括在片上系统内。在一些实施方案中，可将设备1000包括在可以是电池供电的移动设备中。因此，设备10 00的功率消耗可能是重要的设计考虑因素。在图示实施方案中，设备1000包括织物1010、计算复合1020输入/输出(I/O)网桥1050、高速缓存/存储器控制器1045、图形单元1080和显示单元1065。在一些实施方案中，除所示的部件之外和/或代替所示的部件，设备1000可包括其他部件(未示出)，诸如视频处理器编码器和解码器、图像处理或识别元件、计算机视觉元件等。

织物1010可包括各种互连件、总线、MUX、控制器等，并且可被配置为促进设备1000的各种元件之间的通信。在一些实施方案中，织物1010的部分可被配置为实现各种不同的通信协议。在其他实施方案中，织物1010可实现单个通信协议，并且耦接到织物1010的元件可在内部从单个通信协议转换到其他通信协议。

在图示实施方案中，计算复合体1020包括总线接口单元(BIU)1025、高速缓存1030和内核1035和1040。在各种实施方案中，计算复合体1020可包括各种数量的处理器、处理器内核和/或高速缓存。例如，计算复合体1020可包括1个、2个或4个处理器内核或任何其它合适数目。在一个实施方案中，高速缓存1030是一组关联L2高速缓存。在一些实施方案中，内核1035和/或1040可包括内部指令和/或数据高速缓存。在一些实施方案中，结构体1010、高速缓存1030或设备1000中的其他地方的一致性单元(未示出)可被配置为维持设备1000的各个高速缓存之间的一致性。BIU 1025可被配置为管理计算复合体1020和设备1000的其他元件之间的通信。处理器内核诸如内核1035和1040可被配置为执行特定指令集架构(ISA)的指令，该指令集架构可包括操作系统指令和用户应用指令。

高速缓存/存储器控制器1045可被配置为管理结构体1010与一个或多个高速缓存和/或存储器之间的数据传输。例如，可将高速缓存/存储器控制器1045联接到L3高速缓存，继而可将该L3高速缓存联接到系统存储器。在其他实施方案中，可将高速缓存/存储器控制器1045直接联接到存储器。在一些实施方案中，高速缓存/存储器控制器1045可包括一个或多个内部高速缓存。

如本文所用，术语“联接到”可指示元件之间的一个或多个连接，并且联接件可包括中间元件。例如，在图10中，图形单元150可被描述为通过结构体1010和高速缓存/存储器控制器1045“联接到”存储器。相比之下，在图10的例示的实施方案中，图形单元150“直接联接”到结构体1010，因为不存在中间元件。

图形单元150可包括一个或多个处理器和/或一个或多个图形处理单元(GPU)。例如，图形单元150可接收面向图形的指令，诸如Metal或指令。图形单元150可基于所接收的面向图形的指令执行专用GPU指令或执行其他操作。图形单元150通常可被配置为并行处理大块数据，并且可在帧缓冲器中构建图像以输出至显示器。图形单元150可包括在一个或多个图形处理流水线中的变换、照明、三角形和/或渲染引擎。图形单元150可输出用于显示图像的像素信息。在一些实施方案中，图形单元150被配置为执行上述存储器一致性、中间渲染计算、本地图像块和/或像素资源同步技术中的一者或多者。

显示单元1065可被配置为从帧缓冲器读取数据并且提供像素值流以用于显示。在一些实施方案中，显示单元1065可被配置为显示流水线。另外，显示单元1065可被配置为将多个帧混合以产生输出帧。此外，显示单元1065可包括用于耦接至用户显示器(例如，触摸屏或外部显示器)的一个或多个界面(例如，或嵌入式显示端口(eDP))。

I/O桥接部1050可包括各种元件被配置为实现例如：通用串行总线(USB)通信、安全性、音频和/或低功率永远在线功能。I/O桥接部1050还可包括接口诸如像脉冲宽度调制(PWM)、通用输入输出(GPIO)、串行外围接口(SPI)和/或内部集成电路。可将各种类型的外围设备和设备经由I/O桥接器1050耦接到设备1000。

在一些实施方案中，设备1000的各种元件可包括分级布置的时钟门卫，包括被耦接以将时钟信号递送到时钟树的不同部分的各系列DET时钟门卫。所公开的技术可减少设备10 00中的切换功率消耗、将时钟延迟平衡到设备10 00的不同部分、减少设备10 00中的误差、实现更高的频率、在较低的电源电压下实现所需的频率、减少每个周期(或例如每个任务、每个像素或每个字节)内消耗的能量等。

示例性计算机可读介质

本公开已经在上文中详细描述了各种示例性电路。意图在于本公开不仅涵盖包括此类电路系统的实施方案，而且还涵盖包括指定此类电路系统的设计信息的计算机可读存储介质。因此，本公开旨在支持不仅涵盖包括所公开电路系统的装置、而且还涵盖以被配置为生成包括所公开电路系统的硬件(例如集成电路)的制造系统识别的格式指定电路系统的存储介质的权利要求。对此类存储介质的权利要求旨在涵盖例如生成电路设计但本身不制造该设计的实体。

图11是根据一些实施方案示出一种存储电路设计信息的示例性非暂态计算机可读存储介质的框图。在例示的实施方案中，半导体制造系统1120被配置为处理存储在非暂态计算机可读介质1110上的设计信息1115并基于设计信息1115制造集成电路1130。

非暂态计算机可读介质1110可包括各种适当类型的存储器设备或存储设备中的任何设备。介质1110可以是安装介质，例如CD-ROM、软盘或磁带设备；计算机系统存储器或随机存取存储器诸如DRAM、DDR RAM、SRAM、EDO RAM、Rambus RAM等；非易失性存储器诸如闪存、磁介质，例如，硬盘驱动器或光学存储装置；寄存器，或其他类似类型的存储器元件等。介质1110也可包括其他类型的非暂态存储器或它们的组合。介质1110可包括可驻留在不同位置例如通过网络连接的不同计算机系统中的两个或更多个存储器介质。

设计信息1115可利用各种适当的计算机语言中的任何语言来指定，包括硬件描述语言诸如但不限于：VHDL、Verilog、SystemC、SystemVerilog、RHDL、M、MyHDL等。设计信息1115可以能被半导体制造系统1120用来制造集成电路1130的至少一部分。设计信息1115的格式可被至少一个半导体制造系统1120识别。在一些实施方案中，设计信息1115还可包括指定集成电路1130的综合和/或布局的一个或多个单元库。在一些实施方案中，设计信息整体或部分地以指定单元库元素及其连接性的网表的形式来指定。单独获取的设计信息1115可包括或可不包括用于制造对应集成电路的足够信息。例如，设计信息1115可指定要制造的电路元件，但不指定它们的物理布局。在这种情况下，设计信息1115可需要与布局信息结合以实际制造指定电路系统。

半导体制造系统1120可包括被配置为制造集成电路的各种适当元件中的任何元件。这可包括例如用于(例如在可包括掩膜的晶片上)沉积半导体材料、移除材料、改变所沉积材料的形状、(例如通过掺杂材料或利用紫外处理修改介电常数)对材料改性等等的元件。半导体制造系统1120还可被配置为对于正确操作执行所制造电路的各种测试。

在各种实施方案中，集成电路1130被配置为根据设计信息1115指定的电路设计来操作，这可包括执行本文所述的功能性中的任何功能性。例如，集成电路1130可包括图1B、图2、图8和/或图10中所示的各种元件中的任何元件。另外，集成电路1130可被配置为执行本文结合其他部件所述的各种功能。另外，本文所述的功能性可由多个连接的集成电路来执行。

如本文所用，形式为“指定被配置为…的电路的设计的设计信息”的短语并不暗示为了满足该要素就必须制造所涉及的电路。相反，该短语表明设计信息描述了一种电路，该电路在被制造时将被配置为执行所指示的动作或者将包括所指定的部件。

***

尽管上文已经描述了具体实施方案，但这些实施方案并非要限制本公开的范围，即使仅相对于特定特征描述单个实施方案的情况下也是如此。本公开中提供的特征示例意在进行例示，而非限制，除非做出不同表述。上述说明书意在涵盖此类替代形式、修改形式和等价形式，这对知晓本公开有效效果的本领域技术人员应当是显而易见的。

本公开的范围包括本文(明确或暗示)公开的任意特征或特征组合或其任意推广，而无论其是否减轻本文解决的任何或所有问题。因此，在本专利申请(或要求享有其优先权的专利申请)进行期间可针对特征的任何此类组合作出新的权利要求。具体地，参考所附权利要求书，可将从属权利要求的特征与独立权利要求的特征进行组合，并可通过任何适当的方式而不是仅通过所附权利要求书中所列举的特定组合来组合来自相应独立权利要求的特征。

Claims (20)

1.一种设备，包括：

多个着色器处理元件；

本地存储器，所述本地存储器可被一组着色器处理元件中的一个或多个访问并且不可被所述着色器处理元件的其他着色器处理元件访问；

一个或多个存储元件，所述一个或多个存储元件被配置为存储由程序声明并存储在所述本地存储器中的数据结构的属性，其中所述属性包括所述数据结构在X维度中的尺寸、所述数据结构在Y维度中的尺寸、每个像素的样本数量和每个样本的数据量；和

转换电路系统，所述转换电路系统被配置为：

接收访问所述数据结构的请求，其中所述请求指定：一个或多个X坐标、一个或多个Y坐标、一个或多个样本和一个或多个样本之一的数据的索引；以及

将所述请求转换为所述本地存储器中对应于所述存储的数据结构的地址。

2.根据权利要求1所述的设备，其中用于所述设备的应用程序编程接口包括一个或多个功能，所述一个或多个功能允许图形程序查询所述一个或多个存储元件之一以确定所述数据结构的维度。

3.根据权利要求1所述的设备，其中所述数据结构可被一个或多个片段着色器和一个或多个计算线程组两者访问。

4.根据权利要求1所述的设备，其中所述本地存储器包括在与所述多个着色器处理元件相同的集成电路上。

5.根据权利要求1所述的设备，其中每个像素的所述样本数量可被配置用于多样本反混叠(MSAA)处理。

6.根据权利要求1所述的设备，其中所述数据结构在整个对应于正在渲染的图像的一部分的线程上是持续的。

7.根据权利要求1所述的设备，其中所述设备还包括多组所述一个或多个存储元件，所述多组所述一个或多个存储元件可配置用于所述本地存储器中的多个不同数据结构。

8.一种非暂态计算机可读介质，所述非暂态计算机可读介质具有在其上存储的指令，所述指令可被计算设备执行以执行包括以下项的操作：

将多个值存储在一个或多个存储元件中，以指定可由一个或多个着色器处理元件访问的本地存储器中的数据结构的属性，其中所述属性包括：

所述数据结构在X维度中的尺寸；

所述数据结构在Y维度中的尺寸；

每个像素的样本数量；和

每个样本的数据量；以及

使用一个或多个请求访问所述数据结构，所述请求指定：一个或多个X坐标、一个或多个Y坐标、一个或多个样本和一个或多个样本之一的数据的索引。

9.根据权利要求8所述的非暂态计算机可读介质，其中所述操作还包括：

存储对应于正在处理的图块宽度的X维度属性；

存储对应于小于所述X维度属性的尺寸的Y维度属性；以及

使用所述数据结构作为滚动行缓冲器来处理所述图块。

10.根据权利要求8所述的非暂态计算机可读介质，其中所述操作还包括：

确定正在处理的图块的尺寸；以及

存储对应于所述图块的所述尺寸的X维度属性和Y维度属性。

11.根据权利要求8所述的非暂态计算机可读介质，其中所述X维度属性和所述Y维度属性在屏幕空间中以像素为单位来指定。

12.根据权利要求8所述的非暂态计算机可读介质，其中所述操作还包括：

查询所述一个或多个存储元件之一以确定所述数据结构的维度。

13.根据权利要求8所述的非暂态计算机可读介质，其中所述操作还包括：

使用一个或多个片段着色器和一个或多个中间渲染计算图块着色器两者来访问所述数据结构。

14.根据权利要求8所述的非暂态计算机可读介质，还包括：

每个像素分配大于一个的样本数并使用所述数据结构用于多样本反混叠(MSAA)处理。

15.根据权利要求8所述的非暂态计算机可读介质，其中所述数据结构在整个对应于正在渲染的图像的特定图块的线程上是持续的。

16.一种非暂态计算机可读存储介质，所述非暂态计算机可读存储介质在其上存储有设计信息，所述设计信息以半导体制造系统识别的格式指定硬件集成电路的至少一部分的设计，所述半导体制造系统被配置为使用所述设计信息生成根据所述设计的所述电路，包括：

多个着色器处理元件；

本地存储器，所述本地存储器可被一组着色器处理元件中的一个或多个访问并且不可被所述着色器处理元件的其他着色器处理元件访问；

一个或多个存储元件，所述一个或多个存储元件被配置为存储由程序声明并存储在所述本地存储器中的数据结构的属性，其中所述属性包括所述数据结构在X维度中的尺寸、所述数据结构在Y维度中的尺寸、每个像素的样本数量和每个样本的数据量；和

转换电路系统，所述转换电路系统被配置为：

接收访问所述数据结构的请求，其中所述请求指定：一个或多个X坐标、一个或多个Y坐标、一个或多个样本和一个或多个样本之一的数据的索引；以及

将所述请求转换为所述本地存储器中对应于所述存储的数据结构的地址。

17.根据权利要求16所述的非暂态计算机可读存储介质，其中所述数据结构可被一个或多个片段着色器和一个或多个中间渲染计算任务两者访问。

18.根据权利要求16所述的非暂态计算机可读存储介质，其中所述设计信息指定所述本地存储器包括在与所述多个着色器处理元件相同的集成电路上。

19.根据权利要求16所述的非暂态计算机可读存储介质，其中所述数据结构在整个对应于正在渲染的图像的图块的线程上是持续的。

20.根据权利要求16所述的非暂态计算机可读存储介质，其中所述设计信息还指定所述电路包括可配置用于所述本地存储器中的多个不同数据结构的多组存储元件。