多处理器计算平台中的处理器间通信技术

摘要

本发明描述可用于多处理器计算平台内的通信技术。在一些实例中，所述技术可提供可用于在使用命令队列起始任务的多处理器计算平台内支持消息传递的软件接口。在额外的实例中，所述技术可提供可用于多处理器计算平台内的共享存储器处理器间通信的软件接口。在进一步的实例中，所述技术可提供图形处理单元GPU，所述图形处理单元包含用于支持所述GPU与主机CPU之间的消息传递和/或共享存储器通信的硬件。

说明书

技术领域

本发明涉及计算平台，且更特定来说，涉及包含多个处理器的计算平台。

背景技术

包含多个处理器的计算平台用于提高具有高计算密集要求和/或高数据处理量要求 的应用的性能。多处理器计算平台可包含可充当主机装置的通用中央处理单元(CPU)以 及主机CPU可用来卸载计算密集型任务的性能的一个或一个以上计算装置，进而提高 整个系统的性能。在一些情况下，所述一个或一个以上计算装置可经特别设计以比主机 CPU更高效地处理某些类型的任务，其可提供对整个系统的进一步的性能改进。举例来 说，所述一个或一个以上计算装置可经特别设计以比主机CPU更高效地执行平行算法。

可用于多处理器计算系统中的一种类型的计算装置是图形处理单元(GPU)。传统上， GPU包含固定功能硬件，其经特别设计以用于向显示装置实时地再现三维(3D)图形，但 是通常不可编程，即，不可将经编译的程序下载到GPU并在GPU上执行。然而，近来， 随着可编程着色器单元的发展，大多数GPU架构已经转变为可编程架构，所述可编程 架构包含许多并行的处理元件。所述可编程架构允许GPU促进不仅仅是图形操作的执 行，而且促进以高度并行的方式执行通用计算任务。

使用GPU来执行通用非图形专有计算任务可在本文中被称作图形处理单元上的通 用计算(GPGPU)，或者被称作GPU计算。在一些情况下，GPU可使并非图形专有的应 用编程接口(API)可用，进而减轻对GPU的编程以用于执行通用计算任务。GPU计算任 务可包含计算密集的任务和/或包含高度并行的任务，例如矩阵计算、信号处理计算、统 计算法、分子模型化应用、财务应用、医疗成像、密码分析应用等。

GPU是可用于多处理器计算平台中的仅一种类型的计算装置，且可使用其它类型的 计算装置来补充或取代GPU。举例来说，可用于多处理器计算平台中的其它类型的计算 装置包含(例如)额外的CPU、数字信号处理器(DSP)、小区宽带引擎(Cell/BE)处理器，或 任何其它类型的处理单元。

具有多个计算装置的多处理器计算平台可为同类平台或异类平台。在同类平台中， 所有计算装置共享共同指令集架构(ISA)。相比而言，异类平台可包含具有不同ISA的 两个或两个以上计算装置。一般来说，不同类型的计算装置可具有不同的ISA，且相同 类型的不同品牌的计算装置也可具有不同的ISA。

可通过利用多核计算装置和/或众核计算装置来进一步改进多处理器计算平台的性 能。多核计算装置的一实例是上文所描述的GPU，其含有具有多个处理核心的可编程着 色器单元。然而，CPU还可经设计以包含多个处理核心。一般来说，可将包含多个处理 核心的任何芯片或裸片视为多核处理器。处理核心可指代能够对特定块数据执行指令的 处理单元。举例来说，可将GPU内的单一算法逻辑单元(ALU)单元或向量处理器视为处 理核心。众核处理器一般指代具有相对大量的核心的多核处理器，例如大于十个核心， 且通常使用与用于设计具有少量核心的多核处理器的技术不同的技术来设计。多核处理 器通过允许在单一芯片上在多个核心上并行地(例如，同时地)执行软件程序来提供性能 改进。

并行编程模型指代经设计以允许在多个处理核心上同时地执行程序的编程模型。所 述程序可为多线程程序，在这种情况下，单一线程可操作于每一处理核心上。在一些实 例中，单一计算装置可包含用于执行程序的所有处理核心。在其它实例中，用于执行程 序的一些处理核心可位于相同类型或不同类型的不同计算装置上。

可使用跨平台、跨供应商、异类计算平台、并行编程模型应用编程接口(API)开提供 共同语言规范，以用于对包含由实施不同ISA的不同供应商可能制成的不同类型的计算 装置的异类多核计算平台进行并行编程。开放计算语言(OpenCL^TM)是跨平台、跨供应商、 异类计算平台、并行编程API的一实例。此类API可经设计以允许GPU上的更一般化 的数据处理。举例来说，除了经由计算语言暴露经扩展的着色器子系统能力之外，这些 API可以非图形专有的方式将数据流和控制路径一般化到GPU中。然而，目前，由此类 API提供的指令集是基于GPU的硬件架构，且因此受限于与现有GPU架构相容的功能 性。

发明内容

本发明描述可用于多处理器计算平台内的通信技术。在一些实例中，所述技术可提 供软件接口，所述软件接口可用于支持在使用命令队列起始任务的多处理器计算平台内 的消息传递。在额外的实例中，所述技术可提供软件接口，所述软件接口可用于多处理 器计算平台内的共享存储器处理器间通信。在进一步的实例中，所述技术可提供图形处 理单元GPU，所述图形处理单元包含用于支持所述GPU与主机CPU之间的消息传递和 /或共享存储器通信的硬件。

在一个实例中，本发明描述一种包含一个或一个以上处理器的主机装置。所述装置 进一步包含命令队列接口，所述命令队列接口在一个或一个以上处理器上执行且经配置 以响应于从在主机装置上执行的进程接收到一个或一个以上排队指令而将多个命令置 于命令队列中。所述多个命令包含第一命令，所述第一命令指令所述主机装置在与所述 主机装置相关联的第一存储器空间与和图形处理单元(GPU)相关联的第二存储器空间之 间传送数据。所述多个命令进一步包含第二命令，所述第二命令指令所述主机装置起始 GPU上的任务的执行。所述装置进一步包含消息传递接口，所述消息传递接口在一个或 一个以上处理器上执行且经配置以在于GPU上执行的任务正在GPU上执行时且响应于 从在主机装置上执行的进程接收到一个或一个以上消息传递指令而在于所述主机装置 上执行的进程与所述任务之间传递一个或一个以上消息。

在另一实例中，本发明描述一种方法，所述方法包含响应于从在主机装置上执行的 进程接收到一个或一个以上排队指令而用在主机装置的一个或一个以上处理器上执行 的命令队列接口将多个命令放置到命令队列中。所述多个命令包含第一命令，所述第一 命令指令所述主机装置在与所述主机装置相关联的第一存储器空间与和图形处理单元 (GPU)相关联的第二存储器空间之间传送数据。所述多个命令进一步包含第二命令，所 述第二命令指令所述主机装置起始GPU上的任务的执行。所述方法进一步包含在于GPU 上执行的任务正在GPU上执行时且响应于从在主机装置上执行的进程接收到一个或一 个以上消息传递指令而用在主机装置的一个或一个以上处理器上执行的消息传递接口 在于所述主机装置上执行的进程与所述任务之间传递一个或一个以上消息。

在另一实例中，本发明描述一种设备，所述设备包含用于响应于从在主机装置上执 行的进程接收到一个或一个以上排队指令而将多个命令放置到命令队列中的装置。所述 多个命令包含第一命令，所述第一命令指令所述主机装置在与所述主机装置相关联的第 一存储器空间与和图形处理单元(GPU)相关联的第二存储器空间之间传送数据。所述多 个命令进一步包含第二命令，所述第二命令指令所述主机装置起始GPU上的任务的执 行。所述设备进一步包含在于GPU上执行的任务正在GPU上执行时且响应于从在主机 装置上执行的进程接收到一个或一个以上消息传递指令而在于所述主机装置上执行的 进程与所述任务之间传递一个或一个以上消息的装置。

在另一实例中，本发明描述一种包含指令的计算机可读存储媒体，所述指令致使一 个或一个以上处理器响应于从在主机装置上执行的进程接收到一个或一个以上排队指 令而将多个命令放置到命令队列中。所述多个命令包含第一命令，所述第一命令指令所 述主机装置在与所述主机装置相关联的第一存储器空间与和图形处理单元(GPU)相关联 的第二存储器空间之间传送数据。所述多个命令进一步包含第二命令，所述第二命令指 令所述主机装置起始GPU上的任务的执行。所述计算机可读存储媒体进一步包含致使 所述一个或一个以上处理器在于GPU上执行的任务正在GPU上执行时且响应于从在主 机装置上执行的进程接收到一个或一个以上消息传递指令而在于所述主机装置上执行 的进程与所述任务之间传递一个或一个以上消息的指令。

在另一实例中，本发明描述一种图形处理单元(GPU)，其包含经配置以执行任务的 一个或一个以上处理器。所述GPU进一步包含可由主机装置存取的一个或一个以上寄 存器。所述GPU进一步包含消息传递模块，所述消息传递模块经配置以在于所述一个 或一个以上处理器上执行的任务正在所述一个或一个以上处理器上执行时且响应于从 在所述一个或一个以上处理器上执行的任务接收到一个或一个以上消息传递指令而经 由所述一个或一个以上寄存器在所述任务与在主机装置上执行的进程之间传递一个或 一个以上消息。

在另一实例中，本发明描述一种方法，所述方法包含用图形处理单元(GPU)的消息 传递模块从在所述GPU上执行的任务接收一个或一个以上消息传递指令。所述方法进 一步包含经由可由主机装置存取的所述GPU内的一个或一个以上寄存器在于所述GPU 上执行的任务正在所述GPU上执行时且响应于从在所述GPU上执行的任务接收到一个 或一个以上消息传递指令而在所述任务与在所述主机装置上执行的进程之间传递一个 或一个以上消息。

在另一实例中，本发明描述一种设备，所述设备包含用于从在图形处理单元(GPU) 上执行的任务接收一个或一个以上消息传递指令的装置。所述设备进一步包含用于经由 可由主机装置存取的所述GPU内的一个或一个以上寄存器在于所述GPU上执行的任务 正在所述GPU上执行时且响应于从在所述GPU上执行的任务接收到一个或一个以上消 息传递指令而在所述任务与在所述主机装置上执行的进程之间传递一个或一个以上消 息的装置。

在另一实例中，本发明描述一种包括指令的计算机可读媒体，所述指令致使一个或 一个以上处理器从在图形处理单元(GPU)上执行的任务接收一个或一个以上消息传递指 令。所述计算机可读存储媒体进一步包含致使所述一个或一个以上处理器经由可由主机 装置存取的所述GPU内的一个或一个以上寄存器在于所述GPU上执行的任务正在所述 GPU上执行时且响应于从在所述GPU上执行的任务接收到一个或一个以上消息传递指 令而在所述任务与在所述主机装置上执行的进程之间传递一个或一个以上消息的指令。

在另一实例中，本发明描述一种方法，所述方法包含用在主机装置的一个或一个以 上处理器上执行的存储器缓冲器接口来接收包含指定是否应针对可由主机装置以及由 图形处理单元(GPU)存取的共享存储器空间来启用直接模式的信息的指令。所述方法进 一步包含基于指定是否应启用所述直接模式的所述信息而用所述存储器缓冲器接口针 对所述共享存储器空间选择性地启用所述直接模式。

在另一实例中，本发明描述一种包含一个或一个以上处理器的主机装置。所述装置 进一步包含存储器缓冲器接口，所述存储器缓冲器接口在所述一个或一个以上处理器上 执行且经配置以接收包含指定是否应针对共享存储器空间来启用直接模式的信息的指 令，且基于指定是否应启用所述直接模式的所述信息而针对所述共享存储器空间选择性 地启用所述直接模式，所述共享存储器空间可由主机装置以及由图形处理单元(GPU)存 取。

在另一实例中，本发明描述一种设备，所述设备包含用于接收包含指定是否应针对 可由主机装置以及由图形处理单元(GPU)存取的共享存储器空间来启用直接模式的信息 的指令的装置。所述设备进一步包含用于基于指定是否应启用所述直接模式的所述信息 而针对所述共享存储器空间选择性地启用所述直接模式的装置。

在另一实例中，本发明描述一种包括指令的计算机可读媒体，所述指令致使一个或 一个以上处理器接收包含指定是否应针对可由主机装置以及由图形处理单元(GPU)存取 的共享存储器空间来启用直接模式的信息的指令。所述计算机可读存储器媒体进一步包 含致使一个或一个以上处理器基于指定是否应启用所述直接模式的所述信息而针对所 述共享存储器空间选择性地启用所述直接模式的指令。

在另一实例中，本发明描述一种图形处理单元(GPU)，其包含与存储器相关联的GPU 高速缓冲存储器。所述装置进一步包含一个或一个以上处理模块，所述一个或一个以上 处理模块经配置以响应于接收到指定是否应使用高速缓存服务来用于相对于存储器的 存储器空间执行读取操作和写入操作中的至少一者的信息而选择性地使用GPU高速缓 冲存储器的高速缓存服务来相对于所述存储器空间执行读取操作和写入操作中的至少 一者。

在另一实例中，本发明描述一种方法，所述方法包含响应于接收到指定是否应使用 高速缓存服务来用于相对于存储器的存储器空间执行读取操作和写入操作中的至少一 者的信息而选择性地使用与存储器相关联的图形处理单元(GPU)高速缓冲存储器的高速 缓存服务来相对于所述存储器空间执行读取操作和写入操作中的至少一者。

在另一实例中，本发明描述一种设备，其包含与存储器相关联的GPU高速缓冲存 储器。所述设备进一步包含用于响应于接收到指定是否应使用高速缓存服务来用于相对 于存储器的存储器空间执行读取操作和写入操作中的至少一者的信息而选择性地使用 GPU高速缓冲存储器的高速缓存服务来相对于所述存储器空间执行读取操作和写入操 作中的至少一者的装置。

在另一实例中，本发明描述一种包括指令的计算机可读媒体，所述指令致使一个或 一个以上处理器响应于接收到指定是否应使用高速缓存服务来用于相对于存储器的存 储器空间执行读取操作和写入操作中的至少一者的信息而选择性地使用与存储器相关 联的图形处理单元(GPU)高速缓冲存储器的高速缓存服务来相对于所述存储器空间执行 读取操作和写入操作中的至少一者。

附图说明

图1是说明根据本发明的可用于执行消息传递技术的实例性计算系统的方框图。

图2是说明根据本发明的可用于图1的计算系统中的实例性GPU的方框图。

图3是说明根据本发明的用于多处理器平台环境中的消息传递的实例性技术的流程 图。

图4是说明根据本发明的用于执行由在主机装置上执行的进程发布的发送指令的实 例性技术的流程图。

图5和6是说明根据本发明的可用于实施图4中所说明的技术的若干部分的实例性 技术的流程图。

图7是说明根据本发明的用于处理例如GPU等计算装置中的所接收的消息的实例 性技术的流程图。

图8是说明根据本发明的用于执行由在例如GPU等计算装置上执行的任务发布的 接收指令的实例性技术的流程图。

图9和10是说明根据本发明的可用于实施图8中所说明的技术的若干部分的实例 性技术的流程图。

图11是说明根据本发明的用于执行由在例如GPU等计算装置上执行的进程发布的 发送指令的实例性技术的流程图。

图12和13是说明根据本发明的可用于实施图11中所说明的技术的若干部分的实 例性技术的流程图。

图14是说明根据本发明的用于执行由在主机装置上执行的进程发布的寄存回调例 程指令的实例性技术的流程图。

图15是说明根据本发明的用于处理从计算装置接收到的中断的实例性技术的流程 图。

图16和17是说明根据本发明的可用于实施图15中所说明的技术的若干部分的实 例性技术的流程图。

图18是说明根据本发明的用于执行由在主机装置上执行的进程发布的读取指令的 实例性技术的流程图。

图19是说明根据本发明的可用于实施图18中所说明的技术的若干部分的实例性技 术的流程图。

图20是说明根据本发明的可促进直接存储器对象的使用的实例性计算系统的方框 图。

图21是说明根据本发明的用于执行由在主机装置上执行的进程发布的存储器对象 创建指令的实例性技术的流程图。

图22是说明根据本发明的用于执行由在主机装置上执行的进程发布的存储器对象 创建指令的另一实例性技术的流程图。

图23到26是说明根据本发明的用于处理高速缓存模式指令和直接模式指令的实例 性技术的流程图。

图27是说明根据本发明的可用于图20的计算系统中的实例性GPU的方框图。

图28是说明根据本发明的用于处理高速缓存模式指令和直接模式指令的实例性技 术的流程图。

图29是说明根据本发明的用于执行由在主机装置上执行的进程发布的存储器对象 创建指令的另一实例性技术的流程图。

图30是说明根据本发明GPU可如何处理根据第一编译技术而编译的指令序列的流 程图。

图31是说明根据本发明的用于编译用于任务的源代码的实例性技术的流程图。

图32是说明根据本发明的可由GPU用来选择性地使用高速缓存服务的实例性技术 的流程图。

具体实施方式

本发明描述可用于多处理器计算平台内的通信技术。在一些实例中，所述技术可提 供软件接口，所述软件接口可用于支持在使用命令队列起始任务的多处理器计算平台内 的消息传递。在额外的实例中，所述技术可提供软件接口，所述软件接口可用于多处理 器计算平台内的共享存储器处理器间通信。在进一步的实例中，所述技术可提供图形处 理单元GPU，所述图形处理单元包含用于支持所述GPU与主机CPU之间的消息传递和 /或共享存储器通信的硬件。

近年来，最初经设计以用于处理实时3D图形的处理器(例如，图形处理单元(GPU)) 被一般化，从而执行通用计算任务(GPGPU)。已通过采用业界标准(例如，开放计算语言 (OpenCL^TM)标准)来部分地证明GPGPU的价值。OpenCL是可用于在多处理器计算平台 上执行具有任务级并行度和/或数据级并行度的跨平台、跨供应商、异类计算平台、并行 编程API的一实例。所述API经特别设计以通过以非图形专有的方式来使GPU的数据 流和控制路径一般化而允许GPU上的更一般化的数据处理。此方法的一个限制是主机 CPU与计算装置(例如，GPU)之间的数据通信的粗糙粒度。

举例来说，OpenCL API提供支持主机装置与一个或一个以上计算装置之间的任务 级粒度的通信的命令队列接口。每一命令队列一般保持将由特定计算装置执行的命令。 在主机装置上执行的主机进程可通过将指令主机装置执行存储器传送的命令放置在命 令队列中而在主机存储器空间与装置存储器空间之间传送数据。类似地，主机进程可通 过将指令主机装置在计算装置上执行任务的命令放置在命令队列中而致使任务开始在 计算装置上执行。

所述命令队列接口可经配置以提供对命令的按序执行或对命令的无序执行。当命令 队列接口经配置以提供对命令的按序执行时，命令队列接口保证将以将命令放置到命令 队列中的次序来执行命令，且直到前一命令已完成执行之后才将开始对后续命令的执 行。因此，当主机进程将命令放置在命令队列中来执行任务时，命令队列等待任务完成 执行，之后执行可能被随后放置到命令队列中的任何额外的命令。

在涉及主机CPU和GPU以及按序命令队列的简单环境中，主机CPU与GPU之间 的通信方案可涉及以下操作：(1)主机CPU准备好数据且将所述数据放置到GPU可存取 的存储器中；(2)主机CPU命令GPU执行任务；(3)主机CPU等待GPU完成对所述任务 的执行；以及(4)主机CPU将数据从GPU可存取的存储器复制到主机存储器。在此类配 置中，将在GPU上执行任务所需的所有数据传送到GPU可存取的存储器，之后开始对 所述任务的执行，且由在GPU上执行的任务产生的数据不可用于主机CPU，直到在GPU 上执行的任务完成执行之后方可。主机CPU与GPU之间的数据共享上的此粗糙度可阻 止对用于基于并行的应用的许多有用操作的实施，例如，在于主机装置上执行的进程与 在GPU上执行的任务之间传递进程间消息。此些消息(例如)对于允许在GPU上运行的 任务具有在主机CPU上执行远程过程调用(RPC)的能力可为有用的。

当命令队列接口经配置以提供对命令的无序执行时，在特定任务的执行期间，主机 进程不能够控制何时将发生对特定命令的执行。因此，用于命令队列的无序执行模式实 际上也不允许在于主机装置上执行的进程与在GPU上执行的任务之间实施进程间消息 传递。

关于用于OpenCL中的存储器模型，API界定所谓的全局CL缓冲器和全局CL图像， 其可用于在主机CPU与GPU之间共享数据或用于在多个OpenCL计算装置之间共享数 据。然而，CPU和GPU无法同时从缓冲器进行读取或写入到缓冲器。通常，CPU准备 好含有源数据的一个或一个以上缓冲器，且将所述缓冲器传递到GPU以供处理。GPU 修改这些缓冲器或将结果放置在还曾由在CPU上执行的软件先验分配的其它缓冲器中， 以用于接收GPU数据修改。

虽然OpenCL中的存储器对象当前允许将主机存储器空间的区用于存储由计算装置 使用的缓冲器数据，但所述规范允许计算装置对此数据进行高速缓存以用于对任务的更 高效的执行。主机装置一般不能直接使用于对缓冲器数据进行高速缓存的计算装置高速 缓冲存储器无效。因此，即使主机装置将盖写存储于主机存储器空间中的某些存储器缓 冲器数据，也不能保证计算装置中的高速缓冲存储器将得到更新以向计算装置提供对经 修改数据的直接存取。另外，因为由计算装置执行的计算的结果可被存储在计算装置高 速缓冲存储器中，所以在主机装置上执行的主机进程不能从缓冲器读取任何部分结果， 因为此类数据可能归因于计算装置高速缓冲存储器中所存储的较新的数据而无效。因 此，OpenCL中的存储器管理模型未容易地经由共享存储器实现运行中的数据共享。

在一些实例中，可使用本发明中所描述的技术来克服OpenCL API的上文提及的限 制中的一者或一者以上。举例来说，本发明的技术可提供软件接口，所述软件接口可用 于支持在使用任务级粒度命令队列起始任务的多处理器计算平台内的进程间消息传递。 作为另一实例，本发明的技术可提供软件接口，所述软件接口可用于支持经由多处理器 计算平台内的共享存储器的运行中的数据共享。

在一些实例中，本发明的技术可提供促进软件级消息传递的GPU硬件架构。举例 来说，本发明的技术可提供经配置以支持对软件级消息传递指令的执行的GPU硬件架 构。在进一步的实例中，本发明的技术可提供促进GPU与主机CPU之间的共享存储器 通信的GPU硬件架构。举例来说，本发明的技术可提供经配置以针对共享存储器空间 选择性地启用和停用高速缓存服务且/或针对共享存储器空间选择性地启用和停用高速 缓冲存储器-相关性机制。

根据本发明的第一方面，提供一种消息传递接口，所述消息传递接口促进在由计算 装置执行任务期间在主机装置与一个或一个以上计算装置之间执行消息传递指令。消息 传递可指代一种形式的进程间以及潜在地装置间的通信，其中正通信的进程各自执行互 补组的操作以成功地传递消息。举例来说，根据消息传递协议进行通信的进程中的每一 者可实施发送操作和接收操作。本发明中的消息传递技术可允许CPU和计算装置(例如， GPU)在于计算装置上执行任务期间在彼此之间传递消息。以此方式，实施任务级粒度命 令队列通信方案的多处理器计算平台可以能够促进进程间和/或装置间通信。

在一些实例中，本发明中所描述的消息传递技术可被称作“带外信令”技术，因为 所述技术使用不同于命令队列接口的接口，命令队列接口通常用于OpenCL中以用于在 主机装置与计算装置(例如，GPU)之间的通信。换句话说，本发明的技术可包含新的带 外通信接口，其与OpenCL内所包含的带内命令队列接口在逻辑上分离。带外通信接口 可不经受命令队列接口所经受的相同任务级粒度，进而提供对上文相对于命令队列的任 务级粒度所描述的一个或一个以上限制的解决方案。

根据本发明的技术在CPU与GPU之间传送的消息可为任何类型的消息。不同类型 的消息的实例包含信号、存储器分配请求、存储器取消分配请求、通知消息、同步消息、 远程过程调用消息(例如，作为远程过程调用(RPC)的部分的消息)、数据包、报告消息、 断言机制消息，以及记录消息。

在当前的OpenCL范例中，从主机CPU到GPU的所有请求均在OpenCL命令队列 中排队等候，且随后被发送到GPU。具体来说，应用可能将大量内核执行和缓冲器操作 排在命令队列中。同时，如果首先被排队的任务(例如，内核执行)需要(例如)向CPU请 求额外的存储器分配，那么就出现问题。首先，GPU如何在运行内核中向CPU通知其 需要进行存储器分配？其次，CPU如何向GPU通知存储器分配的完成以及新分配的存 储器块的地址？然而，本发明的消息传递接口技术可能够通过允许含有上述通知和信息 的一个或一个以上消息在CPU与GPU之间传递来解决这些问题。

在一些实例中，可使用本发明的带外信令技术在主机CPU与一个或一个以上计算 装置(例如，OpenCL计算装置)之间实施信令。带外信令可例如使用推拉机制来提供快速 的带外通知。在一些实例中，带外信令技术可携载相对少量的数据。

根据本发明的第二方面，提供能够将消息发送到在不同于GPU的处理器上执行的 进程以及从所述进程接收消息的GPU。举例来说，GPU可包含经配置以实施用于发送 和接收消息的一个或一个以上操作的硬件。在一些实例中，根据本发明而设计的GPU 可包含一个或一个以上主机可存取的寄存器，所述寄存器经配置以存储与消息传递协议 相关联的状态和数据信息。所述一个或一个以上寄存器可经配置以促进在GPU上执行 的任务与在不同于GPU的装置上执行的进程之间的消息传递。在进一步的实例中，GPU 的ALU处理块(例如，可编程着色器单元)可通信地耦合到主机可存取的寄存器以经由所 述主机可存取的寄存器来发送和接收消息。GPU还可经设计以包含各种轮询和/或中断 机制以实施同步和/或异步消息传递技术。

根据本发明的第三方面，提供存储器缓冲器接口，其允许创建直接存储器对象。直 接存储器对象可用于实施非可高速缓冲共享存储器空间和/或高速缓冲存储器相干共享 存储器空间，以便在于计算装置上执行的任务正在计算装置上执行时在于主机装置上执 行的进程与所述任务之间共享数据。所述共享存储器空间可为可由主机装置以及计算装 置(例如，GPU)两者在计算装置执行任务期间存取的存储器空间。如本文中所使用的非 可高速缓存的共享存储器空间可指代针对所述存储器空间而停用主机装置和计算装置 中的一者或两者中的一个或一个以上对应高速缓冲存储器的共享存储器空间。如本文中 所使用的高速缓冲存储器相干的共享存储器空间可指代其中使用共享存储器高速缓冲 存储器相干技术来维持主机装置和计算装置中的一者或两者中的一个或一个以上对应 高速缓冲存储器内的高速缓冲存储器相干的共享存储器空间。所述非可高速缓存共享存 储器空间以及高速缓冲存储器相干性共享存储器空间可在任何时间允许数据共享。在一 些实例中，可将直接存储器对象实施为非可高速缓存易失性共享存储器和/或实施为高速 缓冲存储器相干的易失性共享存储器来用于主机装置和计算装置。

在一些实例中，本发明的直接存储器对象可被集成在包含存储器对象存储器管理方 案的跨平台、跨供应商、异类计算平台、并行编程API内。举例来说，可将直接存储器 对象集成到OpenCL中以作为OpenCL存储器对象的额外属性，例如OpenCL缓冲器对 象或OpenCL图像对象。在此些实例中，可通过修改存储器对象创建功能以包含一参数 或旗标来创建直接存储器对象，所述参数或旗标指定由功能调用创建的所得的存储器对 象是否应为标准模式存储器对象或直接模式存储器对象。以此方式，本发明的技术可允 许实施包含若干存储器对象存储器管理方案(例如，OpenCL)以经由不经受高速缓冲存储 器相干性问题的共享存储器空间来实施运行中的数据共享的API的多处理器计算平台。

在进一步的实例中，本发明的直接存储器对象可用于主机CPU与OpenCL计算装置 之间或者不同的OpenCL计算装置之间的运行中的数据共享。在额外的实例中，直接存 储器对象可含有内部同步标记。在进一步的实例中，可与带外信号一起使用直接存储器 对象以用于同步。

根据本发明的第四方面，提供包含对应于共享存储器空间的高速缓冲存储器的 GPU，所述共享存储器空间可针对特定存储器地址空间而被选择性地停用以便提供非可 高速缓存共享存储器空间。举例来说，GPU可响应于接收到指定是否应使用高速缓存服 务来相对于共享存储器空间执行读取操作和/或写入操作的信息而启用和停用由与共享 存储器空间相关联的高速缓冲存储器提供的高速缓存服务。在一些实例中，指定是否应 使用高速缓存服务来相对于共享存储器空间执行读取操作和/或写入操作的信息可为高 速缓存模式指令或直接模式指令，其指定是否应使用高速缓存模式或直接模式来执行特 定指令。在进一步的实例中，指定是否应使用高速缓存服务来相对于共享存储器空间执 行读取操作和/或写入操作的信息可为直接模式存储器对象属性，其指定是否针对存储器 对象启用直接模式。

在进一步的实例中，本发明的技术可提供包含高速缓冲存储器相干模式的GPU，所 述高速缓冲存储器相干模式可被选择性地启用以提供高速缓冲存储器相干的共享存储 器空间。在一些实例中，GPU可基于从主机装置接收到的一个或一个以上指令来选择性 地启用高速缓冲存储器相干模式以用于对应于共享存储器空间的高速缓冲存储器的一 部分。在主机装置基于由主机进程指定的直接模式参数而分配共享存储器空间之后，主 机装置可即刻向GPU发布一个或一个以上指令以选择性地启用共享存储器空间高速缓 冲存储器相干模式以用于对应于共享存储器空间的高速缓冲存储器的一部分。

与可通过单独使用OpenCL命令队列接口而获得的主机CPU与GPU之间或两个 OpenCL计算装置之间的任务耦合相比，本发明的带外信令和直接缓冲技术可提供更精 细粒度的任务耦合。本发明的技术可允许多处理器计算平台执行多种操作以便辅助并行 和/或多线程程序的高校执行。举例来说，本发明的技术可允许在GPU上执行的任务启 动RPC。作为另一实例，本发明的技术可允许在GPU上执行的任务经由CPU来启动另 一GPU任务。作为进一步的实例，本发明的技术可允许在GPU上执行的任务向CPU 和/或在CPU上执行的驱动器发布资源管理请求，例如存储器分配和/或存储器取消分配 请求。作为又一实例，本发明的技术可允许在GPU上执行的任务执行状态检查和到CPU 的一般消息传递，例如断言机制的实施、进展报告，和/或诊断记录。

图1是说明根据本发明的实例性计算系统10的方框图。计算系统10经配置以在多 个处理装置上处理一个或一个以上软件应用。在一些实例中，所述一个或一个以上软件 应用可包含主机进程，且计算系统10可经配置以执行主机进程且分布由在计算系统10 内的其它计算装置上的主机进程起始的一个或一个以上任务的执行。在进一步的实例 中，可根据并行编程模型来编程由计算系统10执行的主机进程和/或任务。举例来说， 所述应用可包含经设计以充分利用基础硬件系统的任务级并行度和/或数据级并行度的 指令。

计算系统10可为个人计算机、桌上型计算机、膝上型计算机、计算机工作站、视 频游戏平台或控制台、移动电话(例如，蜂窝式或卫星电话)、移动电话、陆线电话、因 特网电话、手持式装置(例如，便携式视频游戏装置或个人数字助理(PDA))、数字媒体播 放器(例如，个人音乐播放器)、视频播放器、显示装置，或电视、电视机顶盒、服务器、 中间网络装置、大型计算机或处理信息的任何其它类型的装置。

计算系统10包含主机装置12、图形处理单元(GPU)14、存储器16和互连网络18。 主机装置12经配置以提供用于执行用于多处理器计算平台API的主机进程和运行时模 块的平台。通常，主机装置12是通用CPU，但主机装置12可为能够执行程序的任何类 型的装置。主机装置12经由互连网络18通信地耦合到GPU14和存储器16。主机装置 12包含主机进程20和运行时模块22，主机进程20和运行时模块22中的每一者可在一 个或一个以上可编程处理器的任何组合上执行。

主机进程20包含形成用于在计算系统10的计算系统平台上执行的软件程序的一组 指令。所述软件程序可经设计以执行用于终端用户的一个或一个以上特定任务。在一些 实例中，此些任务可涉及可利用由计算系统10提供的多个处理装置和并行架构的计算 密集算法。

运行时模块22可为在主机装置12上执行的软件模块，其实施经配置以服务于主机 进程20中所包含的指令中的一者或一者以上的一个或一个以上接口。由运行时模块22 实施的接口包含命令队列接口24和主机消息传递接口26。在一些实例中，运行时模块 22可实施除了本发明中所描述的接口之外的标准多处理器系统API内所包含的一个或 一个以上接口。在一些实例中，所述标准API可为异类计算平台API、跨平台API、跨 供应商API、并行编程API、任务级并行编程API和/或数据级并行编程API。在进一步 的实例中，所述标准API可为OpenCL API。在此些实例中，可将运行时模块22设计成 遵照OpenCL规范中的一者或一者以上。在额外的实例中，可将运行时模块22实施为 驱动器程序(例如，GPU驱动器)的一部分或实施为驱动器程序。

命令队列接口24经配置以从主机进程20接收一个或一个以上排队指令，且执行由 所接收的指令指定的功能。在一些实例中，可根据OpenCL规范来设计命令队列接口24。 举例来说，命令队列接口24可实施OpenCL规范中所指定的排队指令中的一者或一者 以上以用于与命令队列交互。

根据本发明，主机消息传递接口26经配置以从主机进程20接收一个或一个以上消 息传递指令，且执行由所接收的指令指定的功能。在一些实例中，可将主机消息传递接 口26实施为对现有标准API(例如，OpenCL API)的扩展。在额外的实例中，可将主机消 息传递接口26集成到现有标准API(例如，OpenCL API)中。

GPU14经配置以响应于从主机装置12接收到的指令来执行一个或一个以上任务。 GPU14可为包含一个或一个以上可编程处理元件的任何类型的GPU。举例来说，GPU14 可包含经配置以并行地执行任务的多个执行实例的一个或一个以上可编程着色器单元。 可编程着色器单元可包含顶点着色器单元、片段着色器单元、几何着色器单元和/或统一 着色器单元。GPU14经由互连网络18通信地耦合到主机装置12和存储器16。GPU14 包含任务28和装置消息传递接口30。任务28和装置消息传递接口30可在一个或一个 以上可编程处理元件的任何组合上执行。

任务28包括形成用于在计算系统10中的计算装置上执行的任务的一组指令。在一 些实例中，用于任务28的所述组指令可在主机进程20中界定，且在一些情况下，由在 主机装置12上执行的主机进程20中所包含的指令编译。在进一步的实例中，任务28 可为具有在GPU14上并行地执行的多个执行实例的内核程序。在此些实例中，主机进 程20可界定用于内核的索引空间，其将内核执行实例映射到用于执行内核执行实例的 相应处理元件，且GPU14可根据为内核界定的索引空间来执行用于任务28的多个内核 执行实例。

根据本发明，装置消息传递接口30经配置以从主机进程20接收一个或一个以上消 息传递指令，且执行由所接收的指令指定的功能。在一些实例中，可将装置消息传递接 口30实施为对现有标准API的扩展。举例来说，所述标准API可为标准计算装置API， 例如OpenCL C API。在额外的实例中，可将装置消息传递指令30集成到现有标准API(例 如，OpenCL C API)中。

存储器16经配置以存储数据以供主机装置12和GPU14中的一者或两者使用。存 储器16可包含一个或一个以上易失性或非易失性存储器或存储装置的任何组合，所述 易失性或非易失性存储器或存储装置例如为随机存取存储器(RAM)、静态RAM(SRAM)、 动态RAM(DRAM)、只读存储器(ROM)、可擦除可编程ROM(EPROM)、电可擦除可编 程ROM(EEPROM)、快闪存储器、磁性数据存储媒体或光学存储媒体。存储器16经由 互连网络18通信地耦合到主机装置12和GPU14。存储器16包含命令队列32。

命令队列32可为实施于存储器16中的数据结构，存储器16存储并检索从命令队 列接口24接收到的命令。在一些实例中，命令队列32可为以特定次序存储命令以用于 执行的缓冲器。

互连网络18经配置以促进主机装置12、GPU14与存储器16之间的通信。互连网 络18可为此项技术中已知的任何类型的互连网络。在图1的实例性计算系统10中，互 连网络18是总线。所述总线可包含多种总线结构中的任一者中的一者或一者以上，例 如第三代总线(例如，超传输总线或不限带宽总线)、第二代总线(例如，高级图形端口总 线、外围组件互连快递(PCIe)总线，或高级可扩展接口(AXI)总线)，或任何其它类型的 总线。互连网络18耦合到主机装置12、GPU14和存储器16。

现在将进一步详细地描述计算系统10中的组件的结构和功能性。如上文所论述， 主机进程20包含一组指令。所述组指令可包含(例如)一个或一个以上排队指令，以及一 个或一个以上主机消息传递指令。在额外的实例中，所述组指令可包含指定将在GPU14 上执行的任务或内核的指令、创建命令队列且使命令队列与特定装置相关联的指令、编 译并捆绑程序的指令、设置内核自变量的指令、界定索引空间的指令、界定装置背景的 指令，以及支持由主机进程20提供的功能性的其它指令。

主机进程20可通过向命令队列接口24发布指令命令队列接口24将一个或一个以 上命令放置到命令队列32中的一个或一个以上排队指令而与命令队列接口24交互。所 述一个或一个以上排队指令可包含指令命令队列接口24将存储器传送命令排到命令队 列32中的存储器传送排队指令。举例来说，所述一个或一个以上排队指令可包含将一 命令排队的指令，所述命令指令主机装置12(例如，在主机装置12上执行的运行时模块 22)在与主机装置12相关联的存储器空间与和GPU14相关联的存储器空间之间传送数 据。

如果存储器空间在主机装置12执行主机进程20期间可由主机装置12存取，那么 存储器空间可与主机装置12相关联。类似地，如果存储器空间在GPU14执行任务28 期间可由GPU14存取，那么存储器空间可与GPU14相关联。与主机装置12相关联的 存储器空间可在本文中被称作主机存储器空间，且与GPU14相关联的存储器空间可在 本文中被称作装置存储器空间。在一些实例中，存储器16可包含主机存储器空间和装 置存储器空间两者的部分。在进一步的实例中，主机存储器空间和装置存储器空间中的 一者或两者的部分可位于图1的计算系统10中未展示的一个或一个以上其它存储器装 置上。

在一些实例中，指令主机装置12在与主机装置12相关联的存储器空间与和GPU14 相关联的存储器空间之间传送数据的命令可为指令运行时模块22将存储于主机存储器 空间的一部分中的数据传送到分配于装置存储器空间中的缓冲器对象的命令。由主机进 程20发布以将此命令排队的指令可在本文中被称作写入缓冲器排队指令。在一些情况 下，写入缓冲器排队指令可采取由OpenCL API规范指定的clEnqueueWriteBuffer()功能 的形式。

在额外的实例中，指令主机装置12在与主机装置12相关联的存储器空间与和GPU 14相关联的存储器空间之间传送数据的命令可为指令运行时模块22将存储于分配于装 置存储器空间中的缓冲器对象中的数据传送到主机存储器空间的一部分的命令。由主机 进程20发布以将此命令排队的指令可在本文中被称作读取缓冲器排队指令。在一些情 况下，读取缓冲器排队指令可采取由OpenCL API规范制定的clEnqueueReadBuffer()功 能的形式。

所述一个或一个以上排队指令还可包含指令命令队列接口24将任务执行命令排队 到命令队列32中的任务执行排队指令。举例来说，所述一个或一个以上排队指令可包 含用以将一命令排队的指令，所述命令指令主机装置12(例如，在主机装置12上执行的 运行时模块22)在GPU14上执行任务。在一些实例中，用以执行任务的命令可为在GPU 14的多个处理元件上并行地执行任务的多个执行实例的命令。举例来说，所述任务可为 内核，主机进程20可界定用于内核的索引空间，其将内核执行实例映射到GPU14中的 用于执行内核执行实例的相应处理元件。在此实例中，用以执行任务的命令可为用以根 据为GPU14界定的索引空间在GPU14上执行内核的命令。在一些情况下，任务执行 排队指令可采取由OpenCL API指定的clEnqueueNDRangeKernel()功能的形式。

根据本发明，主机进程20还可通过向主机消息传递接口26发布指令主机消息传递 接口26在于主机装置12上执行的主机进程20与在GPU14上执行的任务28之间传递 一个或一个以上消息的一个或一个以上主机消息传递指令而与主机消息传递接口26交 互。所述主机消息传递指令可由主机装置12执行。

在一些实例中，主机消息传递指令可包含指令主机装置12将指定数据发送到指定 装置的发送指令。举例来说，所述发送指令可指令主机消息传递接口26将消息从在主 机装置12上执行的主机进程20发送到在GPU14上执行的任务28。在一些实例中，所 述发送指令可包含指定应将消息发送到其的特定装置的第一输入参数，以及指定将发送 的消息的内容的第二输入参数。

所述发送指令可为封锁发送指令或非封锁发送指令。在一些实例中，所述发送指令 可包含指定所述发送指令是封锁发送指令还是非封锁发送指令的第三输入参数。封锁发 送指令可在完成发送操作之前一直进行等待，之后返回到调用进程，例如在主机装置12 上执行的主机进程20。非封锁发送指令可返回到调用进程，而不在完成发送操作之前一 直等待。举例来说，非封锁发送指令返回到特定发送操作的句柄，可由调用进程随后询 问所述句柄以确定发送操作是否成功。非封锁发送操作可能失败，且在失败的情况下， 调用进程可需要再次发布发送指令以重试发送操作。

在一些实例中，用于发送指令的接口可采取以下形式：

其中clSendOutOfBandData是指令识别符，cl_device*deviceId是指定应将消息发送 到其的特定OpenCL装置的输入参数，int OOB_data是指定将发送的消息的内容的输入 参数，且bool blocking是指定指令是封锁发送指令还是非封锁发送指令的输入参数。在 封锁指令的情况下，指令可返回指示发送操作是否成功完成的参数。在非封锁指令的情 况下，指令可返回用于由调用进程进行后续状态询问的句柄参数。

在一些实例中，主机消息传递指令可包含寄存回调例程指令，寄存回调例程指令指 令主机装置12寄存回调，从而以异步的方式从指定装置接收数据。举例来说，寄存回 调例程指令可响应于从GPU14接收到指示在GPU14上执行的任务已将一消息发送到 主机进程20的信号而指令主机消息传递接口26调用回调例程。寄存回调例程指令可包 含指定应寄存回调例程的特定装置的第一输入参数，以及指定回调例程的存储器位置的 第二输入参数。

在一些实例中，用于寄存回调例程的接口可采取以下形式：

clRegisterOutOfBandDataCallback(

cl_device*deviceId,

void(*)(int)callBackPtr)

其中clRegisterOutOf BandDataCallback是指令识别符，cl_device*deviceId是指定应 将消息发送到其的特定OpenCL装置的输入参数，且void(*)(int)callBackPtr是指定回 调例程的存储器位置的输入参数。寄存回调例程指令可返回指示回调例程回调例程寄存 操作是否成功完成的参数。

在一些实例中，主机消息传递指令可包含指令主机装置12尝试从指定装置读取数 据的轮询指令。举例来说，轮询指令可指令主机消息传递接口26针对指示在GPU14上 执行的任务28是否已发送消息的消息状态信息来轮询GPU14。轮询指令可包含指定将 轮询的特定装置的输入参数，以及指定由于轮询而获得的数据的输出参数。

在一些实例中，用于轮询指令的接口可采取以下形式：

clTryReadOutOfBandData(

cl_device*deviceId,

int*OOB_data)

其中ClTryReadOutOfBandData是指令识别符，cl_device*deviceId是指定将轮询的 特定OpenCL装置的输入参数，且int*OOB_data是指定由于轮询而获得的数据的输出 参数。轮询指令可返回指示是否从轮询操作成功地获得数据的参数。

类似于主机进程20，任务28可包含由计算装置执行的一个或一个以上装置消息传 递指令。所述装置消息传递指令可包含指令计算装置将指定数据发送到主机装置12的 发送指令。举例来说，发送指令可指令GPU14将消息从在GPU14上执行的任务28发 送到在主机装置12上执行的主机进程20。

所述发送指令可为封锁发送指令或非封锁发送指令。在一些实例中，所述发送指令 可包含指定所述发送指令是封锁发送指令还是非封锁发送指令的第一输入参数。封锁发 送指令可中止调用进程(例如，在GPU14上执行的任务28)，且等待第二操作完成，之 后返回到调用进程。非封锁发送指令可返回到调用进程，而不在完成发送操作之前一直 等待。举例来说，非封锁发送指令返回到特定发送操作的句柄，可由调用进程随后询问 所述句柄以确定发送操作是否成功。非封锁发送操作可能失败，且在失败的情况下，调 用进程可需要再次发布发送指令以重试发送操作。发送指令可包含指定将发送到主机装 置的消息的内容的第二输入参数。

在一些实例中，用于发送指令的接口可采取以下形式：

send_oobdata(

bool blocking,

int data)

其中send_oobdata是指令识别符，bool blocking是指定指令是封锁发送指令还是非 封锁发送指令的输入参数，且int data是指定将发送的消息的内容的输入参数。在封锁 指令的情况下，指令可返回指示发送操作是否成功完成的参数。在非封锁指令的情况下， 指令可返回用于由调用进程进行后续状态询问的句柄参数。

在一些实例中，装置消息传递指令可包含指令计算装置从主机装置12接收数据的 接收指令。举例来说，接收指令可指令GPU14(例如，装置消息传递接口30)向在GPU14 上执行的任务28提供从在主机装置12上执行的主机进程20发送到任务28的消息(如果 可用)。此指令可用于支持轮询机制。

所述接收指令可为封锁接收指令或非封锁接收指令。在一些实例中，所述接收指令 可包含指定所述接收指令是封锁接收指令还是非封锁接收指令的输入参数。封锁接收指 令可中止调用进程(例如，在GPU14上执行的任务28)，且一直等到消息可用，之后返 回到调用进程。非封锁接收指令可返回到调用进程，而不一直等到消息可用。举例来说， 如果消息可用，那么非封锁接收指令可返回所述消息。然而，如果消息不可用，那么非 封锁接收指令可失败。在失败的情况下，调用进程可需要再次发布接收指令以重试接收 操作。接收指令可包含指定由于接收操作而获得的数据(如果有)的输出参数。

在一些实例中，用于接收指令的接口可采取以下形式：

receive_oobdata(

bool blocking,

int data)

其中receive_oobdata是指令识别符，bool blocking是指定指令是封锁接收指令还是 非封锁接收指令的输入参数，且int data是指定由于接收操作而获得的数据(如果有)的输 出参数。所述指令可返回指示接收操作是否成功的参数。

命令队列接口24经配置以将命令排到命令队列32中。举例来说，命令队列接口24 可从主机进程20接收一个或一个以上排队指令，且将响应于从主机进程20接收一个或 一个以上排队指令而将一个或一个以上命令放置到命令队列32中。所述一个或一个以 上排队指令可包含任务执行排队指令以及数据传送排队指令，其分别指令命令队列接口 24对任务执行命令和数据传送命令进行排队。

命令队列接口24还经配置以执行存储于命令队列32中的命令。对于数据传送命令， 命令队列接口24可在主机存储器空间与装置存储器空间之间传送数据。举例来说，对 于写入缓冲器命令，命令队列接口24可将存储于主机存储器空间的一部分中的数据传 送到在装置存储器空间中分配的缓冲器对象。作为另一实例，对于读取缓冲器命令，命 令队列接口24可将存储于在装置存储器空间中分配的缓冲器对象中的数据传送到主机 存储器空间的一部分。装置存储器空间可对应于命令队列32与其相关联的装置。

对于任务执行命令，命令队列接口24可致使在与命令队列相关联的装置上开始任 务的执行。举例来说，在图1的实例中，命令队列32与运行时模块22的背景内的GPU 14相关联。因此，当执行任务执行命令时，命令队列接口24可致使任务开始在GPU14 上执行。在一些实例中，命令队列接口24可通过将一个或一个以上命令放置到GPU14 内所包含的本地命令队列中而致使任务开始在GPU14上执行。在其它实例中，命令队 列接口24可通过将指令GPU14开始执行任务的一个或一个以上指令发送到GPU14而 致使任务开始在GPU14上执行。命令队列接口24可使用互连网络18来与GPU14、存 储器16和主机存储器空间和装置存储器空间通信。

在一些实例中，命令队列接口24可按序执行命令。在此些实例中，如果将第一命 令排在第二命令之前，那么第二命令的执行在第一命令已完成执行后开始。在进一步的 实例中，命令队列接口24可无序执行命令。在此些实例中，如果将第一命令排在第二 命令之前，那么第二命令的执行不一定在第一命令已完成执行后才开始。

主机消息传递接口26经配置以执行从主机进程20接收的一个或一个以上消息传递 指令。举例来说，响应于从主机进程20接收到一个或一个以上消息传递指令，主机消 息传递接口26可在于GPU14上执行的任务28正在GPU14上执行时在于主机装置12 上执行的主机进程20与任务28之间传递一个或一个以上消息。在一些实例中，主机消 息传递接口26可在未将任何命令放置到命令队列32中的情况下执行一个或一个以上消 息传递指令。

根据第一实例，响应于从主机进程20接收到发送指令，主机消息传递接口26可在 任务28正在GPU14上执行时将消息从主机进程20发送到任务28。举例来说，主机消 息传递接口26可基于在发送指令内所包含的消息数据来撰写传出消息，且经由互连网 络18将传出消息传送到在发送指令中指定的装置(例如，GPU14)，以用于递送到在指 定装置上执行的任务(例如，任务28)。

根据第二实例，响应于从主机进程20接收到寄存回调例程指令，主机消息传递接 口26可使指令中所指定的回调例程与来自指令中所指定的装置(例如，GPU14)的信号 相关联，从而指示在指定装置上执行的任务(例如，任务28)已发送消息。在一些实例中， 来自装置的信号可为中断信号。在一些实例中，可经由专用中断信号线来递送中断信号。 响应于从指定装置接收到指示在装置上执行的任务已发送消息的信号，主机消息传递接 口26可起始在寄存回调例程指令中指定的回调例程的执行。回调例程可从指定装置(例 如，GPU14)获得由任务(例如，任务28)发送的消息，且将所述消息返回到主机进程20 以供进一步处理。

根据第三实例，响应于接收到轮询指令，主机消息传递接口26可针对消息状态信 息来轮询指令中所指定的装置(例如，GPU14)。主机消息传递接口26可使用互连网络 18或另一基于硬件的通信路径来轮询装置。如果消息状态信息指示在指定装置(例如， GPU14)上执行的任务(例如，任务28)已发送消息，那么主机消息传递接口26可从指定 装置获得消息，且将所述消息返回到主机进程20以供进一步处理。

装置消息传递接口30经配置以执行从任务28接收的一个或一个以上装置消息传递 指令。举例来说，响应于从任务28接收到一个或一个以上装置消息传递指令，装置消 息传递接口30可在于GPU14上执行的任务28正在GPU14上执行时在任务28与在主 机装置12上执行的主机进程20之间传递一个或一个以上消息。

根据第一实例，响应于接收到发送指令，装置消息传递接口30可将消息从在GPU14 上执行的任务28发送到在主机装置12上执行的主机进程20。举例来说，装置消息传递 接口30可基于在发送指令内所包含的消息数据来撰写传出消息，且经由互连网络18将 传出消息传送到在主机装置12以供递送到主机进程20。

根据第二实例，响应于从任务28接收到接收指令，装置消息传递接口30可确定来 自主机进程20的消息是否可用。在一些实例中，装置消息传递接口30可检查一个或一 个以上主机可存取的寄存器以确定消息是否可用。如果来自主机进程20的消息可用， 那么装置消息传递接口30可将所述消息提供到任务28。

虽然将命令队列接口24和主机消息传递接口26说明为与图1中的主机进程20分 离的组件，但在一些实例中，命令队列接口24和主机消息传递接口26中的一者或两者 的功能性可部分地且/或完全编译到主机进程20中。类似地，在一些实例中，装置消息 传递接口30的功能性可部分地且/或完全编译到任务28中。

为了易于说明，图1中所说明的实例性计算系统10描述将GPU14用作计算装置的 本发明的消息传递技术。然而，应认识到，可将本发明的技术应用于具有不同于除了 GPU14之外或作为GPU14的替代的GPU的计算装置的多处理器计算系统。在一些实 例中，计算装置可为OpenCL计算装置。OpenCL计算装置包含一个或一个以上计算单 元。计算单元中的每一者包含一个或一个以上处理元件。举例来说，计算单元可为具有 可由计算单元中的所有处理元件使用的芯片上共享存储器的处理元件(例如，ALU)的群 集。工作项目可为由放置到命令队列中的命令在OpenCL计算装置上调用的内核或任务 的多个并行执行中的一者。每一工作项目可在计算单元中的个别处理元件上与在其它处 理元件上执行的其它工作项目并行地执行。工作群组可为在计算装置内的单一计算单元 上作为单一内核执行命令的部分而处理的一个或一个以上工作项目的集合。OpenCL主 机可为平台的用于运行OpenCL运行时层的中央CPU。

OpenCL API可提供一组共同接口以用于主机装置与不同类型的计算装置之间的交 互。举例来说，OpenCL API可提供共同接口以用于主机与GPU计算装置以及主机与非 GPU计算装置之间的交互。OpenCL API允许主机使用共同接口来用于在各种计算装置 上执行任务(例如，OpenCL内核)。在一些实例中，任务可为通用计算任务，且OpenCL API可允许主机致使通用计算任务在GPU计算装置上执行。

图1中所示的实例性计算系统10说明用于促进主机装置与计算装置之间的消息传 递和/或带外信令的基础结构和技术。然而，在其它实例性计算系统中，可容易将所述技 术扩展以提供在具有一个以上计算装置的计算系统中的不同计算装置(例如，OpenCL计 算装置)之间的运行中的消息传递。在此些实例中，可在不同的计算装置之间布线一个或 一个以上中断线。

图2是说明根据本发明的可用于图1的计算系统10中的实例性GPU40的方框图。 在一些实例中，GPU40可用于实施图1中所说明的GPU14。GPU40包含GPU处理块 42、主机可存取的GPU寄存器44和总线控制器46。GPU40可经由互连网络18通信地 耦合到一个或一个以上其它主机装置或计算装置。

GPU处理块42经配置以执行任务，且促进在GPU处理块42上执行的任务与在其 它主机或计算装置上执行的进程之间的消息传递。GPU处理块42(例如)经由一个或一个 以上控制和/或数据线而通信地耦合到主机可存取的GPU寄存器44。在一些实例中，GPU 处理块42可被称作算术逻辑单元(ALU)块。GPU处理块42包含任务48、消息传递模块 50、传入数据寄存器52和传出数据寄存器54。

主机可存取的GPU寄存器44经配置以存储可被传送到主机装置或从主机装置传送 的数据。主机可存取的GPU寄存器44包含消息状态寄存器56、消息计数寄存器58、 传入消息寄存器60、传出消息寄存器62、中断状态寄存器64和中断确认寄存器66。主 机可存取的GPU寄存器44中的每一者可为可由主机装置(例如，图1中的主机装置12) 存取的。在一些实例中，主机可存取的GPU寄存器44可为存储器映射的寄存器，即被 映射到主机装置的存储器空间且可在所述存储器空间中寻址的寄存器。在进一步的实例 中，主机可存取的GPU寄存器44可为输入/输出映射(I/O映射)的寄存器，即被映射到 主机装置的I/O空间的寄存器。主机可存取的GPU寄存器44经由一个或一个以上控制 和/或数据线而通信地耦合到GPU处理块42。主机可存取的GPU寄存器44还经由互连 网络18而通信地耦合到总线控制器46。

任务48可在一个或一个以上可编程处理器上执行。在一些实例中，GPU处理块42 可包含经配置以执行任务48的多个执行实例的多个处理器或处理元件。任务48可实质 上类似于上文相对于图1所描述的任务28，且因此将不进一步详细地描述。

消息传递模块50经配置以控制由GPU40执行的消息传递操作。可以硬件、软件、 固件或其任何组合来实施消息传递模块50。在一些实例中，如果消息传递模块50的功 能性中的一些或全部是以软件实施，那么用于此实施方案的软件指令可包含在与含有用 于任务48的软件指令的可执行文件相同的可执行文件内。消息传递模块50通信地耦合 到任务48、消息传递模块50、传入数据寄存器52和传出数据寄存器54。

消息传递模块50可在于一个或一个以上处理器上执行的任务48正在一个或一个以 上处理器上执行且响应于从任务48接收到一个或一个以上消息传递指令而经由主机可 存取的GPU寄存器44在任务48与在主机装置上执行的进程之间传递一个或一个以上 消息。在一些实例中，所述一个或一个以上消息传递指令可包含指令消息传递模块50 将消息从任务48发送到在主机装置上执行的进程的发送指令。在此些实例中，消息传 递模块50可将与所述消息相关联的消息数据存储在主机可存取的GPU寄存器44中的 一者中。在进一步的实例中，所述一个或一个以上消息传递指令可包含指令消息传递模 块50向任务48提供从在主机装置上执行的进程发送到任务48的消息的接收指令。在 此些实例中，消息传递模块50可从主机可存取的GPU寄存器44中的一者或一者以上 获得与所述消息相关联的消息数据。

在图2的实例中，传入数据寄存器52是存储经由传入消息寄存器60从外部装置接 收的传入数据的硬件寄存器。传入数据寄存器52还可存储指示传入数据寄存器52中的 数据是否已被消耗且/或传入数据寄存器52中的数据是否以用于读取的状态位。传入数 据寄存器52经由一个或一个以上数据线而通信地耦合到传入消息寄存器60。在一些实 例中，数据线的数目可等于传入数据寄存器52中的位的数目，这两个数目可等于消息 中的位的数目。在进一步的实例中，位的数目可为32位。在一些实例中，GPU处理块 42可实施内部先入先出(FIFO)缓冲器来存储从传入数据寄存器52接收的多个传入消息。

在图2的实例中，传出数据寄存器54是存储从由任务48发布的一个或一个以上消 息传递指令接收的传出数据的硬件寄存器。传出数据寄存器54经由一个或一个以上数 据线而通信地耦合到传出消息寄存器62。在一些实例中，数据线的数目可等于传出数据 寄存器54中的位的数目，这两个数目可等于消息中的位的数目。在一些实例中，传出 数据寄存器54和传出消息寄存器62可经配置以使得当消息传递模块50将数据写入到 传出数据寄存器54时，自动地用写入到传出数据寄存器54的数据来更新传出消息寄存 器62。在一些实例中，GPU处理块42可实施内部先入先出(FIFO)缓冲器来存储将写入 到传出数据寄存器54的多个传出消息。

在图2的实例中，消息状态寄存器56经配置以存储指示传入消息是否被GPU40接 受的数据。消息状态寄存器56可由主机装置使用以确定消息是否被成功传输，且在一 些实例中，消息状态寄存器56可由主机装置使用以实施后退和/或溢出机制。在接受传 入消息之后，消息传递模块50可将消息状态寄存器56设定为指示传入消息被接受的特 定值。

在图2的实例中，消息计数寄存器58经配置以存储指示传入消息寄存器60是否含 有传入消息的数据。在一些实例中，当消息计数寄存器58已被主机装置递增时，消息 计数寄存器58可将指示消息到达的信号发送到消息传递模块50。在一些情况下，所述 信号可为1位脉冲线。在进一步的实例中，消息传递模块50可在从传入数据寄存器52 读取消息之后递减消息计数寄存器58。

在图2的实例中，传入消息寄存器60经配置以存储传入消息数据。举例来说，主 机装置可将传入消息数据放置到传入消息寄存器60中以便将消息发送到任务48。传入 消息寄存器60通信地耦合到传入数据寄存器52。

在图2的实例中，传出消息寄存器62经配置以存储从传出数据寄存器54接收的传 出消息数据。当将新数据写入到传出数据寄存器54时，传出消息寄存器62可自动地更 新传出消息寄存器62中的数据以对应于传出数据寄存器54。在一些实例中，消息传递 模块50可响应于传出消息被写入到传出消息寄存器62而产生中断信号。中断信号可被 发送到主机装置且指示消息传递模块50已发送消息。

在图2的实例中，中断状态寄存器64经配置以存储指示传出消息是否已被写入到 传出消息寄存器62的状态位。举例来说，中断状态寄存器64和传出消息寄存器62可 经配置以使得当传出消息被写入到传出消息寄存器62时设定中断状态寄存器64中的状 态位。状态位可允许在主机装置上执行的进程轮询GPU40以查看消息是否可用。

在图2的实例中，中断确认寄存器66经配置以存储指示主机装置是否已读取存储 于传出消息寄存器62中的传出消息的确认位。举例来说，传出消息寄存器62和中断确 认寄存器66可经配置以使得当传出消息被写入到传出消息寄存器62时，设定中断确认 寄存器66中的确认位。在此实例中，在主机装置读取传出消息寄存器62时，主机装置 可清除确认位，进而指示主机装置已读取传出消息，且可将新的传出消息写入到传出消 息寄存器62。可使用确认位来实施用于传出消息数据的流控制方案。

在图2的实例中，总线控制器46经配置以允许外部装置经由互连网络18对主机可 存取的GPU寄存器44的存取。举例来说，总线控制器46可对总线信号进行多路复用 和多路分用，且执行由总线信号指定的各种接收和传输操作。总线控制器46可根据一 个或一个以上公共的或专有的总线标准来操作。

现在将根据本发明的某些方面来描述用于多处理器计算系统中的消息传递的各种 技术。在一些实例中，图1的计算系统10可用于实施图3到19中所示的实例性技术。 为了易于阐释，可相对于图1中所示的实例性计算系统10的组件来描述所述技术，但 应理解，可以相同或不同的配置在具有相同或不同的组件的其它系统上执行所述技术。 在额外的实例中，可相对于图2的GPU40的特定组件来描述图3到19中所示的技术中 的一些。再者，应理解，图2是可以能够实施本发明的技术的GPU的一个实例，且可 以相同或不同的配置在具有相同或不同的组件的其它GPU上执行所述技术。

图3说明根据本发明的用于多处理器平台环境中的消息传递的实例性技术。在一些 实例中，图1的计算系统10可用于实施图3中所示的实例性技术。命令队列接口24将 存储器传送命令放置到命令队列32中(70)。命令队列接口24将任务执行命令放置到命 令队列32中(72)。命令队列接口24执行任务执行命令以起始GPU14上的任务的执行 (74)。在任务28正在GPU14上执行时，主机消息传递接口26在主机装置12与GPU14 之间传递一个或一个以上消息(76)。举例来说，主机消息传递接口26可将源自由主机进 程20发布的一个或一个以上发送指令的消息传递到GPU14。所述一个或一个以上发送 指令可指定GPU14或在GPU14上执行的任务是消息的目的地。

图4是根据本发明的用于执行由在主机装置上执行的进程发布的发送指令的实例性 技术。在一些实例中，图1的计算系统10可用于实施图4中所示的实例性技术。主机 消息传递接口26从主机进程20接收发送指令(78)。主机消息传递接口26基于与发送指 令包含在一起的消息数据而产生传出消息(80)。在一些实例中，传出消息可等同于发送 指令中所包含的消息数据。在额外的实例中，主机消息传递接口26可将一条或一条以 上标头信息和/或路由信息附加到发送指令中所包含的消息数据以产生传出消息。在进一 步的实例中，主机消息传递接口26可对发送指令中所包含的消息数据执行一个或一个 以上译码或变换操作以产生传出消息。主机消息传递接口26可将传出消息发送到GPU 14(82)。

主机消息传递接口26可确定发送指令是封锁指令还是非封锁指令(84)。在一些实例 中，主机消息传递接口26可基于发送指令中指定的输入参数来确定发送指令是封锁指 令还是非封锁指令。在其它实例中，可使用两种不同类型的发送指令，且主机消息传递 接口26可基于指令的类型(例如，指令的操作码(操作码))来确定发送指令是封锁指令还 是非封锁指令。如果主机消息传递接口26确定发送指令是非封锁指令，那么主机消息 传递接口26可将句柄返回到调用进程(86)。所述句柄可允许调用进程询问句柄来确定是 否已在稍后时间成功发送所述消息。如果后续询问指示发送失败，那么调用进程可需要 发布后续发送指令来重试发送操作。在一些实例中，调用进程可响应于失败的发送操作 而实施后退例程或溢出机制。

如果主机消息传递接口26确定发送指令是封锁指令，那么主机消息传递接口26可 确定传出消息是否被GPU14成功接收(88)。如果主机消息传递接口26确定传出消息被 成功接收，那么主机消息传递接口26可将指示发送指令中所包含的消息被成功发送的 值返回到调用进程(90)。否则，如果主机消息传递接口26确定传出消息未被成功接收， 那么主机消息传递接口26可前进到过程方框82且向GPU14重发传出消息。在一些实 例中，当主机消息传递接口26确定消息被成功接收或已达到不成功递送尝试的阈值数 目时，封锁指令可完成。

图5是说明根据本发明的可用于实施图4中的过程方框82的实例性技术的流程图。 在一些实例中，图1的计算系统10和/或图2的GPU40可用于实施图5中所示的实例 性技术。主机消息传递接口26可将传出消息放置或存储在GPU40的传入消息寄存器 60中(92)。主机消息传递接口26可递增GPU40的消息计数寄存器58以向GPU14中 的消息传递模块50指示新的消息已到达(94)。在一些实例中，主机消息传递接口26可 使用此项技术中已知的存储器映射寄存器硬件和/或I/O映射寄存器硬件来执行过程方框 92和94中的一者或一者以上。

图6是说明根据本发明的可用于实施图4中的决策方框88的实例性技术的流程图。 在一些实例中，图1的计算系统10和/或图2的GPU40可用于实施图6中所示的实例 性技术。主机消息传递接口26可检查GPU40的消息状态寄存器56中的状态位(96)。 主机消息传递接口26可基于消息状态寄存器56中的状态位来确定所发送的消息是否被 GPU14接受(98)。如果所述状态位指示所发送的消息被GPU14接受，那么主机消息传 递接口26可确定传出消息被成功接收(100)。另一方面，如果所述状态位指示所发送的 消息未被GPU14接受，那么主机消息传递接口26可确定传出消息未被成功接收(102)。

图7是说明用于处理例如GPU等计算装置中的所接收的消息的实例性技术的流程 图。在一些实例中，图2的GPU40可用于实施图7中所示的实例性技术。GPU40中的 消息传递模块50接收消息到达信号(104)。举例来说，消息计数寄存器58可经配置以使 得每当主机装置递增消息计数寄存器58时，便将消息到达脉冲发送到消息传递模块50。 消息传递模块50可致使将存储于传入消息寄存器60中的数据传送到传入数据寄存器 52(106)。举例来说，消息传递模块50可向传入数据寄存器52发布控制信号，从而致使 传入数据寄存器52用存储于传入消息寄存器60中的数据来盖写存储于传入数据寄存器 52中的当前数据。消息传递模块50可设定传入数据寄存器52中的状态位以指示数据在 传入数据寄存器52中可用，例如，未被消耗(108)。消息传递模块50可设定消息状态寄 存器56中的状态位以指示传入消息已被GPU40接受(110)。

图8是说明根据本发明的用于执行由在计算装置上执行的任务发布的接收指令的实 例性技术的流程图。在一些实例中，图1的计算系统10可用于实施图8中所示的实例 性技术。装置消息传递接口30从任务28接收接收指令(112)。装置消息传递接口30确 定可从主机装置得到消息(114)。

如果消息传递模块50确定消息不可用，那么消息传递模块50可确定所述接收指令 是封锁接收指令还是非封锁接收指令(116)。在一些实例中，消息传递模块50可基于接 收指令中指定的输入参数来确定接收指令是封锁指令还是非封锁指令。在其它实例中， 可使用两种不同类型的接收指令，且消息传递模块50可基于指令的类型(例如，指令的 操作码(操作码))来确定接收指令是封锁指令还是非封锁指令。如果消息传递模块50确 定接收指令是封锁指令，那么消息传递模块50可返回到决策方框114以确定传入消息 是否可用。否则，如果消息传递模块50确定接收指令是非封锁指令，那么消息传递模 块50可将指示接收指令失败的值返回到调用进程(118)。

如果消息传递模块50确定可从主机装置得到消息，那么消息传递模块50可将消息 数据返回到调用进程(120)。消息传递模块50确定是否应将所述消息数据标记为被消耗 (122)。消息传递模块50可基于一个或一个以上消耗模式来确定是否应将所述数据标记 为被消耗。在一些实例中，可将所述消耗模式硬连线到GPU14中。在额外的实例中， 所述消耗模式可由任务28和/或主机进程20编程。举例来说，任务28或主机进程20 中的发送和/或接收指令可含有指定特定消耗模式的参数。举例来说，一个消耗模式可指 定当任务的至少一个执行实例已读取数据时应将消息数据标记为被消耗。作为另一实 例，一个消耗模式可指定当任务的至少阈值数目的执行实例已读取数据时应将消息数据 标记为被消耗。

如果消息传递模块50确定应将消息数据标记为被消耗，那么消息传递模块50可清 除消息数据(124)。举例来说，消息传递模块50可清除传入数据寄存器52中的状态位。 另一方面，如果消息传递模块50确定不应将消息数据标记为被消耗，那么消息传递模 块50可保持消息数据(126)。举例来说，消息传递模块50可不清除传入数据寄存器52 中的状态位。

图9是说明根据本发明的可用于实施图8中的决策方框114的实例性技术的流程图。 在一些实例中，图1的计算系统10和/或图2的GPU40可用于实施图9中所示的实例 性技术。消息传递模块50可读取GPU40的传入数据寄存器52中的状态位(128)。消息 传递模块50可确定所述状态位是否被设定(130)。如果传入数据寄存器52中的状态位被 设定，那么消息传递模块50可确定所述传入消息可用(132)。另一方面，如果传入数据 寄存器52中的状态位未被设定，那么消息传递模块50可确定所述传入消息不可用(134)。

图10是说明根据本发明的可用于实施图8中的过程方框120的实例性技术的流程 图。在一些实例中，图1的计算系统10和/或图2的GPU40可用于实施图10中所示的 实例性技术。消息传递模块50可从GPU40中的传入数据寄存器52检索传入消息数据 (136)。消息传递模块50可基于从传入数据寄存器52检索到的消息数据而产生用于任务 48的返回消息数据(138)。在一些实例中，所返回的消息数据可等同于传入数据寄存器 52中所包含的消息数据。在额外的实例中，消息传递模块50可从传入数据寄存器52中 所包含的消息数据移除一条或一条以上标头信息和/或路由信息以产生返回消息数据。在 进一步的实例中，消息传递模块50可对传入数据寄存器52中所包含的消息数据执行一 个或一个以上解码或变换操作以产生返回消息数据。消息传递模块50将消息数据提供 给任务48(140)。

图11是根据本发明的用于执行由在计算装置(例如，GPU14)上执行的进程发布的发 送指令的实例性技术。在一些实例中，图1的计算系统10可用于实施图11中所示的实 例性技术。消息传递模块50从任务28接收发送指令(142)。消息传递模块50基于与发 送指令包含在一起的消息数据而产生传出消息(144)。在一些实例中，传出消息可等同于 发送指令中所包含的消息数据。在额外的实例中，消息传递模块50可将一条或一条以 上标头信息和/或路由信息附加到发送指令中所包含的消息数据以产生传出消息。在进一 步的实例中，消息传递模块50可对发送指令中所包含的消息数据执行一个或一个以上 译码或变换操作以产生传出消息。消息传递模块50可将传出消息发送到主机装置 12(146)。

消息传递模块50可确定发送指令是封锁指令还是非封锁指令(148)。在一些实例中， 消息传递模块50可基于发送指令中指定的输入参数来确定发送指令是封锁指令还是非 封锁指令。在其它实例中，可使用两种不同类型的发送指令，且消息传递模块50可基 于指令的类型(例如，指令的操作代码(操作码))来确定发送指令是封锁指令还是非封锁指 令。如果消息传递模块50确定发送指令是非封锁指令，那么消息传递模块50可将句柄 返回到调用进程(例如，任务28)(150)。所述句柄可允许调用进程询问句柄来确定是否已 在稍后时间成功发送所述消息。如果后续询问指示发送操作失败，那么调用进程可需要 发布后续发送指令来重试发送操作。

如果消息传递模块50确定发送指令是封锁指令，那么消息传递模块50可确定传出 消息是否被主机装置12成功接收(152)。举例来说，消息传递模块50可轮询主机装置 12内所包含的指示消息是否被接受的状态寄存器。如果消息传递模块50确定传出消息 被成功接收，那么消息传递模块50可将指示发送指令中所包含的消息被成功发送的值 返回到调用进程(154)。否则，如果消息传递模块50确定传出消息未被成功接收，那么 消息传递模块50可前进到过程方框146且向主机装置12重发传出消息。在一些实例中， 当消息传递模块50确定消息被成功接收或已达到不成功递送尝试的阈值数目时，封锁 指令可完成。

图12是说明根据本发明的可用于实施图11中的过程方框146的实例性技术的流程 图。在一些实例中，图1的计算系统10和/或图2的GPU40可用于实施图12中所示的 实例性技术。消息传递模块50可将传出消息放置或存储在传出数据寄存器54中(156)。 响应于新数据被放置到传出数据寄存器54中，传出消息寄存器62可更新传出消息寄存 器62中的数据以对应于传出数据寄存器54(158)。消息传递模块50可产生中断信号且 将中断信号发送到主机装置12，从而指示可从GPU40的任务28得到消息(160)。

图13是根据本发明的可用于实施图11中的过程方框146的另一实例性技术。在一 些实例中，图1的计算系统10和/或图2的GPU40可用于实施图13中所示的实例性技 术。消息传递模块50可将传出消息放置或存储在传出数据寄存器54中(162)。响应于新 数据被放置到传出数据寄存器54中，传出消息寄存器62可更新传出消息寄存器62中 的数据以对应于传出数据寄存器54(164)。消息传递模块50可设定中断状态寄存器64 中的状态位以指示可从GPU40的任务28得到消息。所述状态位可经设定以允许主机装 置12轮询GPU40来确定消息是否可用(166)。

图14是说明根据本发明的用于执行由在主机装置上执行的进程发布的寄存回调例 程指令的实例性技术的流程图。在一些实例中，图1的计算系统10可用于实施图14中 所示的实例性技术。主机消息传递接口26从主机进程20接收寄存回调例程指令(168)。 主机消息传递接口26使寄存回调例程指令中所指定的回调例程与来自所述指令中所指 定的装置(例如，GPU14)的中断信号相关联(170)。在一些实例中，所述中断信号可指示 在指定装置上执行的任务(例如，在GPU14上执行的任务28)已发送消息。在一些实例 中，可经由耦合于主机装置12与GPU14之间的专用中断信号线来递送中断信号。在进 一步的实例中，中断信号可指示除了任务28发送消息之外的其它事件。在此些实例中， 主机消息传递接口26可在接收到中断信号之后执行额外的处理以确定多个事件中的哪 一者是由所述信号表示。

主机消息传递接口26确定回调例程是否与中断信号成功关联(172)。如果回调例程 与中断信号成功关联，那么主机消息传递接口26可将指示寄存回调例程操作成功完成 的值返回到调用进程(174)。否则，如果回调例程未与中断信号成功关联(例如，发生错 误)，那么主机消息传递接口26可将指示寄存回调例程操作失败的值返回到调用进程 (176)。

图15是说明根据本发明的用于处理从计算装置接收到的中断的实例性技术的流程 图。在一些实例中，图1的计算系统10可用于实施图15中所示的实例性技术。主机消 息传递接口26从计算装置(例如，GPU14)接收中断信号(178)。主机消息传递接口26确 定是否响应于消息接收事件而发送中断信号(180)。换句话说，主机消息传递接口26可 确定中断信号是否指示在装置上执行的任务(例如，在GPU14上执行的任务28)已发送 消息。

在一些实例中，中断信号可为用信号通知消息接收事件和无其它事件的专用中断信 号。在此些实例中，主机消息传递接口26可通过接收到中断信号本身且不一定需要执 行其它操作而确定曾响应于消息接收事件而发送了中断信号。在其中中断信号用信号通 知多个可能事件的实例中，主机消息传递接口26可能需要询问计算装置以确定用信号 通知哪一事件。

如果主机消息传递接口26确定未曾响应于消息接收事件而发送中断信号，那么主 机消息传递接口26可检查其它类型的事件(182)。否则，如果主机消息传递接口26确定 曾响应于消息接收事件而发送中断信号，那么主机消息传递接口26可执行与从其接收 到消息的装置相关联的回调例程(184)。

图16是说明根据本发明的可用于实施图15中的决策方框180的实例性技术的流程 图。在一些实例中，图1的计算系统10和/或图2的GPU40可用于实施图16中所示的 实例性技术。主机消息传递接口26可读取GPU40中的中断状态寄存器64(186)。主机 消息传递接口26可确定中断状态寄存器64中的状态位是否指示新的消息可用于主机装 置(188)。举例来说，消息传递模块50可在消息可用时设定中断状态寄存器64中的状态 位，且主机消息传递接口26可轮询中断状态寄存器64来确定所述状态位是否被设定以 便确定新的消息是否可用于主机装置。如果状态位指示新的消息可用于主机装置，那么 主机消息传递接口26可确定曾响应于消息接收事件而发送了中断信号(190)。另一方面， 如果状态位指示新的消息不可用于主机装置，那么主机消息传递接口26可确定未曾响 应于消息接收事件而发送中断信号(192)。

图17是说明根据本发明的可用于实施图15中的过程方框184的实例性技术的流程 图。在一些实例中，图1的计算系统10和/或图2的GPU40可用于实施图17中所示的 实例性技术。主机消息传递接口26可从GPU40中的传出消息寄存器62检索消息(194)。 主机消息传递接口26可清除中断确认寄存器66中的确认位(196)。清除确认位可辅助对 GPU40的流控制。举例来说，GPU40可在传出消息被写入到传出消息寄存器62时设 定中断确认寄存器66中的确认位，且一直等到确认位已被清除，之后将额外数据写入 到传出消息寄存器62。

图18是说明根据本发明的用于执行由在主机装置上执行的进程发布的读取指令的 实例性技术的流程图。在一些实例中，图1的计算系统10可用于实施图18中所示的实 例性技术。主机消息传递接口26接收指定从其读取数据的特定装置的读取指令(198)。 主机消息传递接口26轮询读取指令中所指定的装置(200)。主机消息传递接口26基于从 轮询操作接收的轮询数据而确定是否可从接收指令中所指定的装置得到消息(202)。如果 主机消息传递接口26确定可从接收指令中所指定的装置得到消息，那么主机消息传递 接口26可从读取指令中所指定的装置检索消息(204)。在一些实例中，主机消息传递接 口26可从主机装置12可存取的装置中的寄存器(例如，GPU40中的传出消息寄存器62) 检索消息。主机消息传递接口26可将消息数据返回到调用进程(例如，主机进程20)(206)。 如果主机消息传递接口26确定不可从接收指令中所指定的装置得到消息，那么主机消 息传递接口26可返回指示读取指令失败的值(208)。调用进程可能需要再次发布读取指 令以重试读取操作。

图19是说明根据本发明的可用于实施图18中的决策方框202的实例性技术的流程 图。在一些实例中，图1的计算系统10和/或图2的GPU40可用于实施图19中所示的 实例性技术。主机消息传递接口26可读取GPU40中的中断状态寄存器64(210)。主机 消息传递接口26可确定中断状态寄存器64中的状态位是否指示新的消息可用于主机装 置(212)。举例来说，消息传递模块50可在消息可用时设定中断状态寄存器64中的状态 位，且主机消息传递接口26可轮询中断状态寄存器64来确定所述状态位是否被设定以 便确定新的消息是否可用于主机装置。如果状态位被设定，那么主机消息传递接口26 可确定消息可用(214)。另一方面，如果状态位未被设定，那么主机消息传递接口26可 确定消息不可用(216)。

虽然已在上文将由主机消息传递接口26和装置消息传递接口30实施的消息传递技 术描述为提供主机装置12与GPU14之间的带外信令，但在其它实例性系统中，可使用 其它技术来提供带外信令。举例来说，在一些实例中，可界定特殊的高优先级队列，其 可用于发送带外消息。

图20是说明根据本发明的可促进直接存储器对象的使用的实例性计算系统310的 方框图。计算系统310经配置以在多个处理装置上处理一个或一个以上软件应用。在一 些实例中，所述一个或一个以上软件应用可包含主机进程，且计算系统310可经配置以 执行主机进程且分布由在计算系统310内的其它计算装置上的主机进程起始的一个或一 个以上任务的执行。在进一步的实例中，可根据并行编程模型来编程由计算系统310执 行的主机进程和/或任务。举例来说，所述应用可包含经设计以充分利用基础硬件系统的 任务级并行度和/或数据级并行度的指令。

计算系统310可包括个人计算机、桌上型计算机、膝上型计算机、计算机工作站、 视频游戏平台或控制台、移动电话(例如，蜂窝式或卫星电话)、移动电话、陆线电话、 因特网电话、手持式装置(例如，便携式视频游戏装置或个人数字助理(PDA))、数字媒体 播放器(例如，个人音乐播放器)、视频播放器、显示装置，或电视、电视机顶盒、服务 器、中间网络装置、大型计算机或处理信息的任何其它类型的装置。

计算系统310包含主机装置312、GPU314、存储器316和互连网络318。主机装置 312经配置以提供用于执行用于多处理器计算平台API的主机进程和运行时模块的平 台。通常，主机装置312是通用CPU，但主机装置12可为能够执行程序的任何类型的 装置。主机装置12经由互连网络318通信地耦合到GPU314和存储器316。主机装置 312包含主机进程320、运行时模块322、主机高速缓冲存储器324和主机高速缓冲存储 器控制模块326。主机进程320和运行时模块322可在一个或一个以上可编程处理器的 任何组合上执行。

主机进程320包含形成用于在计算系统310平台上执行的软件程序的一组指令。所 述软件程序可经设计以执行用于终端用户的一个或一个以上特定任务。在一些实例中， 此些任务可涉及可利用由计算系统310提供的多个处理装置和并行架构的计算密集算 法。

运行时模块322可为实施经配置以服务于主机进程320中所包含的指令中的一者或 一者以上的一个或一个以上接口的软件模块。由运行时模块322实施的接口包含存储器 缓冲器接口328。在一些实例中，运行时模块322除了存储器缓冲器接口328之外还可 实施图1中所示的命令队列接口24以及图1中所示的主机消息传递接口26中的一者或 一者以上。在进一步的实例中，运行时模块322可实施除了本发明中所描述的接口之外 的标准多处理器系统API内所包含的一个或一个以上接口。在一些实例中，所述标准 API可为异类计算平台API、跨平台API、跨供应商API、并行编程API、任务级并行编 程API和/或数据级并行编程API。在进一步的实例中，所述标准API可为OpenCL API。 在此些实例中，可将运行时模块322设计成遵照OpenCL规范中的一者或一者以上。在 额外的实例中，可将运行时模块322实施为驱动器程序(例如，GPU驱动器)的一部分或 实施为驱动器程序。

存储器缓冲器接口328经配置以从主机进程20接收一个或一个以上存储器对象创 建指令，且执行由所接收的指令指定的功能。在一些实例中，可将存储器缓冲器接口328 实施为对现有标准API(例如，OpenCL API)的扩展。在额外的实例中，可将主命令队列 接口24集成到现有标准API(例如，OpenCL API)中。

主机高速缓冲存储器324经配置以存储数据以供在主机装置312内执行的进程使 用。在一些实例中，与存储于主机高速缓冲存储器324中的数据相关联的存储器空间可 与存储器316中的存储器空间的一部分重叠。主机高速缓冲存储器324可为此项技术中 已知的任何类型的高速缓冲存储器。举例来说，主机高速缓冲存储器324可包含若干高 速缓冲存储器级(例如，L1、L2等)和/或若干映射机制(例如，直接映射、全关联、组关 联等)的任何组合。主机高速缓冲存储器控制模块326经配置以控制主机高速缓冲存储器 324的操作。

GPU314经配置以响应于从主机装置312接收到的指令来执行一个或一个以上任 务。GPU314可为包含一个或一个以上可编程处理器或处理元件的任何类型的GPU。举 例来说，GPU314可包含经配置以并行地执行任务的多个执行实例的一个或一个以上可 编程着色器单元。可编程着色器单元可包含顶点着色器单元、片段着色器单元、几何着 色器单元和/或统一着色器单元。GPU314经由互连网络318通信地耦合到主机装置312 和存储器316。GPU314包含任务330、GPU高速缓冲存储器332和GPU高速缓冲存储 器控制模块334。任务330可在一个或一个以上可编程处理元件的任何组合上执行。

任务330包括形成用于在计算系统310中的计算装置上执行的任务的一组指令。在 一些实例中，用于任务330的所述组指令可在主机进程320中界定，且在一些情况下， 由在主机进程320中所包含的指令编译。在进一步的实例中，任务330可为具有在GPU 314上并行地执行的多个执行实例的内核程序。在此些实例中，主机进程320可界定用 于内核的索引空间，其将内核执行实例映射到用于执行内核执行实例的相应处理元件， 且GPU314可根据为内核界定的索引空间来执行用于任务330的多个内核执行实例。

GPU高速缓冲存储器332经配置以存储数据以供在GPU314内执行的任务使用。 在一些实例中，与存储于GPU高速缓冲存储器332中的数据相关联的存储器空间可与 存储器316中的存储器空间的一部分重叠。GPU高速缓冲存储器332可为此项技术中已 知的任何类型的高速缓冲存储器。举例来说，GPU高速缓冲存储器332可包含若干高速 缓冲存储器级(例如，L1、L2等)和/或若干映射机制(例如，直接映射、全关联、组关联 等)的任何组合。GPU高速缓冲存储器控制模块334经配置以控制GPU高速缓冲存储器 332的操作。

存储器316经配置以存储数据以供主机装置312和GPU314中的一者或两者使用。 存储器316可包含一个或一个以上易失性或非易失性存储器或存储装置的任何组合，所 述易失性或非易失性存储器或存储装置例如为随机存取存储器(RAM)、静态 RAM(SRAM)、动态RAM(DRAM)、只读存储器(ROM)、可擦除可编程ROM(EPROM)、 电可擦除可编程ROM(EEPROM)、快闪存储器、磁性数据存储媒体或光学存储媒体。存 储器316经由互连网络318通信地耦合到主机装置312和GPU314。存储器316包含共 享存储器空间336。共享存储器空间336可为可由主机装置312和GPU314存取的存储 器空间。

互连网络318经配置以促进主机装置312、GPU314与存储器316之间的通信。互 连网络318可为此项技术中已知的任何类型的互连网络。在图20的实例性计算系统310 中，互连网络318是总线。所述总线可包含多种总线结构中的任一者中的一者或一者以 上，例如第三代总线(例如，超传输总线或不限带宽总线)、第二代总线(例如，高级图形 端口总线、外围组件互连快递(PCIe)总线，或高级可扩展接口(AXI)总线)，或任何其它 类型的总线。互连网络318耦合到主机装置312、GPU314和存储器316。

现在将进一步详细地描述计算系统310中的组件的结构和功能性。如上文所论述， 主机进程320包含一组指令。所述组指令可包含(例如)一个或一个以上存储器对象创建 指令。在额外的实例中，所述组指令可包含指定将在GPU314上执行的任务或内核的指 令、创建命令队列且使命令队列与特定装置相关联的指令、编译并捆绑程序的指令、设 置内核自变量的指令、界定索引空间的指令、界定装置背景的指令，排队指令、消息传 递指令以及支持由主机进程320提供的功能性的其它指令。

根据本发明，主机进程320可通过将一个或一个以上存储器对象创建指令发布到存 储器缓冲器接口328而与存储器缓冲器接口328交互，所述一个或一个以上存储器对象 创建指令指令存储器缓冲器接口328基于所述指令中所包含的指定是否针对存储器对象 启用直接模式的信息而创建所述存储器对象。如本文中所使用，存储器对象可指代表示 可由GPU314存取的存储器空间的区的软件对象。在一些实例中，存储器空间的所述区 还可由主机装置312存取。存储器对象可包含存储器空间中的与存储器对象相关联的数 据。存储器对象可进一步包含与存储器空间相关联的一个或一个以上特性。在一些实例 中，存储器对象可包含全局存储器(例如，存储器316)的参考计数区的句柄。

所述存储器对象可包含缓冲器对象和图像对象。缓冲器对象可为存储一维字节集合 的存储器对象。所述一维字节集合可为与存储器对象相关联的数据。缓冲器对象还可包 含信息，例如以字节计的与缓冲器对象相关联的存储器空间的大小、缓冲器对象的使用 信息，以及为缓冲器对象分配的存储器空间的区。图像对象存储数据的二维或三维阵列， 例如纹理、帧缓冲器或图像。图像对象还可包含信息，例如图像的尺寸、图像中的每一 元素的描述、图像对象的使用信息，以及为图像对象分配的存储器空间的区。

根据本发明的一些方面，存储器对象创建指令可包含指定是否应针对将创建的存储 器对象而启用直接模式的输入参数。如本文中进一步详细地论述，当启用直接模式时， 可将存储器对象实施为非可高速缓存共享存储器且/或实施为高速缓冲存储器相干的共 享存储器。当停用直接模式时，可没有必要将存储器对象实施为非可高速缓存共享存储 器或实施为高速缓冲存储器相干的共享存储器。

在一些实例中，存储器对象可包含指示存储器对象是否为直接模式存储器对象的直 接模式属性。在此实施例中，存储器缓冲器接口328可经配置以基于指定是否应针对存 储器对象启用直接模式的信息而将将创建的存储器对象的直接模式属性设定为指示是 否针对存储器对象启用直接模式的值。存储器对象的直接模式属性可由计算系统310用 来确定是否将存储器对象实施为非可高速缓冲共享存储器且/或实施为高速缓冲存储器 相干的共享存储器。

在一些实例中，存储器对象创建指令可包含缓冲器对象创建指令，所述缓冲器对象 创建指令基于指令中所包含的指定是否针对缓冲器对象启用直接模式的信息而指令存 储器缓冲器接口328创建缓冲器对象。在进一步的实例中，存储器对象创建指令可包含 图像对象创建指令，所述图像对象创建指令基于指令中所包含的指定是否针对图像对象 启用直接模式的信息而指令存储器缓冲器接口328创建图像对象。

在一些实例中，用于缓冲器对象创建指令的接口可采取以下形式：

其中clCreateBuffer是指令识别符，cl_context context是用于创建缓冲器对象的有效 上下文(例如，OpenCL上下文)，cl_mem_flags flags是用于指定缓冲器对象的分配和使 用信息的位字段，size_t size是指定将分配的缓冲器存储器对象的以字节计的大小的参 数，void*host_ptr是到可能已由应用分配的缓冲器数据的指针，且cl_int*errcode_ret 返回一个或一个以上错误代码。所述指令可将所创建的缓冲器对象返回为cl_mem存储 器对象。在此实例中，指定是否应针对图像对象启用直接模式的输入参数可为(例如)在 cl_mem_flags flags字段中指定的CL_IMMEDIATE旗标。

在进一步的实例中，用于图像对象创建指令的接口可采取以下形式：

其中cl_CreateImage2D是指令识别符，cl_context context是用于创建缓冲器对象的 有效上下文(例如，OpenCL上下文)，cl_mem_flags flags是用于指定图像对象的分配和 使用信息的位字段，const cl_image_format*image_format是到描述将分配的图像的格式 性质的结构的指针，size_t image_width是以像素计的图像的宽度，size_t image_height 是以像素计的图像的高度，size_t image_row_pitch是以字节计的扫描线间距，void *host_ptr是到可能已由应用分配的图像数据的指针，且cl_int*errcode_ret返回一个或一 个以上错误代码。所述指令可将所创建的图像对象返回为cl_mem存储器对象。在此实 例中，指定是否应针对图像对象启用直接模式的输入参数可为(例如)在cl_mem_flags flags字段中指定的CL_IMMEDIATE旗标。

在一些实例中，存储器对象创建接口可经配置以在读取/写入属性方面仅允许 WRITE_ONLY属性或READ_ONLY属性。换句话说，在此些实例中，存储器缓冲器接口328 可不允许READ_WRITE属性。非直接CL图像可已经具有由OpenCL规范提供的此特征。 不允许READ_WRITE属性可减小在维持高速缓冲存储器相干中的复杂性。

根据本发明，存储器缓冲器接口328经配置以接收指定是否应针对可由主机装置312 和GPU314存取的共享存储器空间336启用直接模式的指令，且基于指定是否应针对共 享存储器空间336启用直接模式的所接收的指令而针对共享存储器空间336选择性地启 用直接模式。举例来说，存储器缓冲器接口328可在所述指令指定应针对共享存储器空 间336启用直接模式的情况下针对共享存储器空间336启用直接模式，且在所述指令指 定应针对共享存储器空间336停用直接模式的情况下针对共享存储器空间336停用直接 模式。所述指令可为(例如)存储器对象创建指令、缓冲器对象创建指令或图像对象创建 指令中的一者。共享存储器空间336可对应于(例如)存储器对象、缓冲器对象或图像对 象。

在一些实例中，当存储器缓冲器接口328针对共享存储器空间336启用直接模式时， 存储器缓冲器接口328可致使停用用于共享存储器空间336的高速缓存服务。类似地， 当存储器缓冲器接口328针对共享存储器空间336停用直接模式时，存储器缓冲器接口 328可致使针对共享存储器空间336启用用于共享存储器空间336的高速缓存服务。高 速缓存服务可由主机高速缓冲存储器324和GPU高速缓冲存储器332中的一者或两者 执行。如本文中所使用，高速缓存服务可指代通常由此项技术中已知的高速缓冲存储器 执行的服务。

在进一步的实例中，存储器缓冲器接口328可通过将与共享存储器空间336相关联 的直接模式属性设定为指示是否针对共享存储器空间启用直接模式的值而针对共享存 储器空间336启用和停用直接模式。举例来说，存储器缓冲器接口328可通过将与共享 存储器空间336相关联的直接模式属性设定为指示针对共享存储器空间336启用直接模 式的值(即，直接模式属性=true)而针对共享存储器空间336启用直接模式。类似地，存 储器缓冲器接口328可通过将与共享存储器空间336相关联的直接模式属性设定为指示 针对共享存储器空间336停用直接模式的值(即，直接模式属性=false)而针对共享存储 器空间336停用直接模式。在一些情况下，直接模式属性可为可由在GPU314上执行的 任务330存取的全局变量，例如，布尔变量。在一些实例中，直接模式属性可存储于共 享存储器空间336内。在其它实例中，直接模式属性可存储于可由在GPU314上执行的 任务存取的除了共享存储器空间336之外的位置中。在共享存储器空间336是存储器对 象的部分的情况下，直接模式属性可存储于其中存储有存储器对象的其它属性的存储器 空间的位置中。

在其中存储器缓冲器接口328通过设定与共享存储器空间336相关联的直接模式属 性而针对共享存储器空间336启用和停用直接模式的实例中，在一些情况下，用于任务 330的源代码可经编译以使得在相对于共享存储器空间336执行存储器读取或写入操作 之前，任务330存取与共享存储器空间336相关联的直接模式属性，且基于用于共享存 储器空间336的直接模式属性而确定是否针对共享存储器空间336启用直接模式。如果 针对共享存储器空间336启用直接模式，那么任务330可经编程以执行直接模式读取或 写入指令以从共享存储器空间336读取数据或将数据写入到共享存储器空间336。另一 方面，如果未针对共享存储器空间启用直接模式，那么任务330可经编程以执行高速缓 存模式读取或写入指令(例如，高速缓存读取或写入指令)以从共享存储器空间336读取 数据或将数据写入到共享存储器空间336。

直接模式读取和写入指令可(例如)在不使用高速缓存服务的情况下分别执行读取和 写入操作。举例来说，直接模式读取指令可致使高速缓冲存储器无效，之后执行读取操 作，且/或可在执行读取操作时绕过高速缓冲存储器。直接模式写入指令(例如)可致使高 速缓冲存储器在执行写入操作时执行直接回写，且/或可在执行写入操作时绕过高速缓冲 存储器。高速缓存读取和写入指令可(例如)使用GPU高速缓冲存储器332中的一者或两 者的高速缓存服务来分别执行读取和写入操作。

在额外情况下，当编译用于任务330的源代码时，用于任务330的编译器可具有对 指示是否针对共享存储器空间336启用直接模式的信息的存取权。举例来说，用于任务 330的源代码(例如，内核源代码)可包含指示是否针对由任务330使用且与共享存储器 空间336相关联的存储器对象而启用直接模式的旗标。在一些实例中，所述旗标可采取 OpenCL属性限定符的形式，例如a_cl_immediate属性限定符。如果针对与共享存储器 空间336相关联的存储器对象启用直接模式，那么编译器可编译任务330以使得用于任 务330的经编译代码包含用于相对于共享存储器空间336发生的读取或写入操作的直接 模式读取和/或写入指令。否则，如果未针对与共享存储器空间336相关联的存储器对象 启用直接模式，那么编译器可编译任务330以使得用于任务330的经编译代码不包含用 于相对于共享存储器空间336发生的读取或写入操作的直接模式读取和/或写入指令。举 例来说，编译器可编译任务330以使得用于任务330的经编译代码包含用于相对于共享 存储器空间336发生的读取或写入操作的高速缓存读取和/或写入指令。

在进一步的实例中，存储器缓冲器接口328可通过启用和停用由主机装置312中的 主机高速缓冲存储器324以及GPU314中的GPU高速缓冲存储器332中的至少一者对 用于共享存储器空间336的高速缓存服务的执行而针对共享存储器空间336启用和停用 直接模式。举例来说，存储器缓冲器接口328可通过停用由主机装置312中的主机高速 缓冲存储器324以及GPU314中的GPU高速缓冲存储器332中的至少一者对用于共享 存储器空间336的高速缓存服务的执行而针对共享存储器空间336启用直接模式。类似 地，存储器缓冲器接口328可通过启用由主机装置312中的主机高速缓冲存储器324以 及GPU314中的GPU高速缓冲存储器332中的至少一者对用于共享存储器空间336的 高速缓存服务的执行而针对共享存储器空间336停用直接模式。

在此些实例中，存储器缓冲器接口328可通过配置与执行用于共享存储器空间336 的高速缓存服务的高速缓冲存储器相关联的基于硬件的高速缓冲存储器控制模块和/或 基于硬件的存储器管理单元而启用和停用对用于共享存储器空间336的高速缓存服务的 执行。举例来说，为了由主机高速缓冲存储器324启用对用于共享存储器空间336的高 速缓存服务的执行，存储器缓冲器接口328可配置主机高速缓冲存储器控制模块326以 使得由主机高速缓冲存储器328针对共享存储器空间336提供高速缓存服务。为了由主 机高速缓冲存储器324停用对用于共享存储器空间336的高速缓存服务的执行，存储器 缓冲器接口328可(例如)配置主机高速缓冲存储器控制模块326以使得主机高速缓冲存 储器328不针对共享存储器空间336提供高速缓存服务。类似地，为了由GPU高速缓 冲存储器332启用对用于共享存储器空间336的高速缓存服务的执行，存储器缓冲器接 口328可(例如)配置GPU高速缓冲存储器控制模块334以使得由主机高速缓冲存储器 324针对共享存储器空间336提供高速缓存服务。为了由GPU高速缓冲存储器332停用 对用于共享存储器空间336的高速缓存服务的执行，存储器缓冲器接口328可(例如)配 置GPU高速缓冲存储器控制模块334以使得GPU高速缓冲存储器332不针对共享存储 器空间336提供高速缓存服务。

在一些实例中，存储器缓冲器接口328可通过将与共享存储器空间336相关联的一 个或一个以上基于硬件的直接旗标设定为指示是否应针对共享存储器空间336提供高速 缓存服务的值来配置主机高速缓冲存储器控制模块326和GPU高速缓冲存储器控制模 块334中的一者或两者。在一些实例中，所述一个或一个以上基于硬件的直接旗标可为 一个或一个以上寄存器。在进一步的实例中，所述基于硬件的直接旗标可为直接旗标的 表格的部分，其中直接旗标的表格中的每一直接旗标对应于存储器316内的特定地址空 间。在任何情况下，当与共享存储器空间336相关联的一个或一个以上直接旗标被设定 为指示应提供高速缓存服务的值时，主机高速缓冲存储器控制模块326和/或GPU高速 缓冲存储器控制模块334可使用主机高速缓冲存储器324和/或GPU高速缓冲存储器332 来提供用于共享存储器空间336的高速缓存服务。类似地，当与共享存储器空间336相 关联的一个或一个以上直接旗标被设定为指示不应提供高速缓存服务的值时，主机高速 缓冲存储器控制模块326和/或GPU高速缓冲存储器控制模块334可使不提供用于共享 存储器空间336的高速缓存服务。

在此些实例中，GPU高速缓冲存储器控制模块334可经配置以处理用于存储器316 的地址空间内的存储器地址的读取指令和/或写入指令。所述读取和写入指令可(例如)由 在GPU314上执行的任务330发布到GPU高速缓冲存储器控制模块334。响应于接收 到用以从存储器316的给定地址空间内的存储器位置读取数据或将数据写入到所述存储 器位置的读取或写入指令，GPU高速缓冲存储器控制模块334可识别与地址空间相关联 的基于硬件的旗标，且确定在基于基于硬件的旗标的值来处理读取或写入指令时是否使 用GPU高速缓冲存储器332的高速缓存服务。如果GPU高速缓冲存储器控制模块334 确定使用GPU高速缓冲存储器332的高速缓存服务，那么GPU高速缓冲存储器控制模 块334可(例如)尝试从GPU高速缓冲存储器332读取数据(如果所述数据有效)且/或将数 据写入到GPU高速缓冲存储器332。如果GPU高速缓冲存储器控制模块334确定不使 用GPU高速缓冲存储器332的高速缓存服务，那么在一些实例中，GPU高速缓冲存储 器控制模块334可绕过GPU高速缓冲存储器332且直接从存储器316读取数据或将数 据写入到存储器316。在额外实例中，如果GPU高速缓冲存储器控制模块334确定不使 用GPU高速缓冲存储器332的高速缓存服务，那么GPU高速缓冲存储器控制模块334 可使与地址空间相关联的高速缓冲存储器332的一部分无效，之后执行读取指令且/或在 执行写入指令时执行高速缓冲存储器回写或高速缓冲存储器通写技术。主机高速缓冲存 储器控制模块334可响应于从在主机装置312上执行的进程320接收到的读取和写入指 令而以与主机高速缓冲存储器324类似的方式进行操作。

在额外实例中，存储器缓冲器接口328可通过针对主机装置312中的主机高速缓冲 存储器324以及GPU314中的GPU高速缓冲存储器332中的至少一者启用和停用共享 存储器高速缓冲存储器相干模式而针对共享存储器空间336启用和停用直接模式。举例 来说，为了针对共享存储器空间336启用直接模式，存储器缓冲器接口328可针对主机 装置312中的主机高速缓冲存储器324以及GPU314中的GPU高速缓冲存储器332中 的至少一者启用共享存储器高速缓冲存储器相干模式。类似地，为了针对共享存储器空 间336停用直接模式，存储器缓冲器接口328可针对主机装置312中的主机高速缓冲存 储器324以及GPU314中的GPU高速缓冲存储器332中的至少一者停用共享存储器高 速缓冲存储器相干模式。在此些实例中，在一些情况下存储器缓冲器接口328可通过配 置主机高速缓冲存储器控制模块326和GPU高速缓冲存储器控制模块334中的一者或 两者来启用共享存储器高速缓冲存储器相干模式而针对主机高速缓冲存储器324启用共 享存储器高速缓冲存储器相干模式，且通过配置主机高速缓冲存储器控制模块326和 GPU高速缓冲存储器控制模块334中的一者或两者来停用共享存储器高速缓冲存储器相 干模式而针对主机高速缓冲存储器324停用共享存储器高速缓冲存储器相干模式。

当针对主机高速缓冲存储器324启用共享存储器高速缓冲存储器相干模式时，主机 高速缓冲存储器控制模块326可根据已知方法相对于共享存储器空间336执行共享存储 器高速缓冲存储器相干技术。当针对主机高速缓冲存储器324停用共享存储器高速缓冲 存储器相干模式时，主机高速缓冲存储器324可不相对于共享存储器空间336执行共享 存储器高速缓冲存储器相干技术。类似地，当针对GPU高速缓冲存储器332启用共享 存储器高速缓冲存储器相干模式时，GPU高速缓冲存储器控制模块334可根据已知方法 相对于共享存储器空间336执行共享存储器高速缓冲存储器相干技术。当针对GPU高 速缓冲存储器332停用共享存储器高速缓冲存储器相干模式时，GPU高速缓冲存储器控 制模块334可不相对于共享存储器空间336执行共享存储器高速缓冲存储器相干技术。

为了易于说明，图20中所说明的实例性计算系统310描述将GPU314用作计算装 置的本发明的直接缓冲技术。应认识到，可将本发明的技术应用于具有不同于除了GPU 314之外或作为GPU314的替代的GPU的计算装置的多处理器计算系统。在一些实例 中，计算系统310中的计算装置可为OpenCL计算装置。另外，图20中所示的实例性 计算系统310说明用于促进主机装置与计算装置之间的运行中的数据共享的基础结构和 技术。然而，在其它实例性计算系统中，可容易将所述技术扩展以提供在具有一个以上 计算装置的计算系统中的不同计算装置(例如，OpenCL计算装置)之间的运行中的数据共 享。在此些实例中，可在不同的计算装置之间布线一个或一个以上中断线。

图21是说明根据本发明的用于执行由在主机装置上执行的进程发布的存储器对象 创建指令的实例性技术的流程图。在一些实例中，图20的计算系统310可用于实施图 21中所示的实例性技术。存储器对象创建指令可为缓冲器对象创建指令或图像对象创建 指令。存储器缓冲器接口328接收存储器对象创建指令(340)。存储器缓冲器接口328确 定存储器对象创建指令是否指定应针对存储器对象启用直接模式(342)。举例来说，存储 器缓冲器接口328可确定直接旗标参数是否包含于用于存储器对象创建指令的参数列表 中。

如果存储器缓冲器接口328确定存储器对象创建指令未指定应针对存储器对象启用 直接模式，那么存储器缓冲器接口328可针对将创建的存储器对象分配共享存储器空间 336(344)，致使针对共享存储器空间336启用由主机高速缓冲存储器324和GPU高速缓 冲存储器332中的一者或两者对高速缓存服务的执行(346)，且返回对所创建的存储器对 象的参考(348)。存储器对象创建指令可(例如)通过不包含直接旗标参数或通过用不应启 用直接模式的另一参数值进行指定而指定不应启用直接模式。相反，如果存储器缓冲器 接口328确定存储器对象创建指令指定应针对存储器对象启用直接模式，那么存储器缓 冲器接口328可针对将创建的存储器对象分配共享存储器空间336(350)，致使针对共享 存储器空间336停用由主机高速缓冲存储器324和GPU高速缓冲存储器332中的一者 或两者对高速缓存服务的执行(352)，且返回对所创建的存储器对象的参考(354)。存储器 对象创建指令可(例如)通过包含直接旗标参数或通过用应启用直接模式的另一参数值进 行指定而指定应启用直接模式。

在一些实例中，存储器缓冲器接口328可通过将与共享存储器空间336相关联的存 储器对象的直接模式属性设定为指示应针对与共享存储器空间336相关联的存储器对象 提供高速缓存服务的值而针对共享存储器空间336启用高速缓存服务的执行。同样，存 储器缓冲器接口328可通过将与共享存储器空间336相关联的存储器对象的直接模式属 性设定为指示不应针对与共享存储器空间336相关联的存储器对象提供高速缓存服务的 值而针对共享存储器空间336停用高速缓存服务的执行。所返回的存储器对象可包含直 接模式属性。在此些实例中，用于存储器对象的直接模式属性可为可由在主机装置312 上执行的主机进程320和在GPU314上执行的任务330中的一者或两者存取的。主机进 程320和/或任务330可基于与共享存储器空间336相关联的存储器对象的直接模式属性 而确定在相对于共享存储器空间336执行特定读取和写入指令时是否使用高速缓存服 务。

在进一步的实例中，存储器缓冲器接口328可通过将与共享存储器空间336相关联 的基于硬件的直接旗标配置为指示应针对共享存储器空间336提供高速缓存服务的值而 针对共享存储器空间336启用高速缓存服务的执行。同样，存储器缓冲器接口328可通 过将与共享存储器空间336相关联的基于硬件的直接旗标配置为指示不应针对共享存储 器空间336提供高速缓存服务的值而针对共享存储器空间336停用高速缓存服务的执 行。所述一个或一个以上基于硬件的直接旗标可位于主机高速缓冲存储器控制模块326 和GPU高速缓冲存储器控制模块334或另一本地或全局存储器管理单元(未图示)中的一 者或一者以上中。

在额外实例中，存储器缓冲器接口328可使存储器对象返回到调用进程(例如，主机 进程320)，之后分配存储器316中的物理存储器空间以存储数据。在此些实例中，存储 器缓冲器接口328可包含所返回的存储器对象中的直接模式属性。随后，当在稍后时间 为存储器对象分配存储器316时，存储器缓冲器接口328可基于存储器对象的直接模式 属性来配置一个或一个以上基于硬件的直接旗标。

图22是说明根据本发明的用于执行由在主机装置上执行的进程发布的存储器对象 创建指令的另一实例性技术的流程图。在一些实例中，图20的计算系统310可用于实 施图22中所示的实例性技术。存储器对象创建指令可为缓冲器对象创建指令或图像对 象创建指令。存储器缓冲器接口328接收存储器对象创建指令(356)。存储器缓冲器接口 328确定存储器对象创建指令是否指定应针对存储器对象启用直接模式(358)。举例来说， 存储器缓冲器接口328可确定直接旗标参数是否包含于用于存储器对象创建指令的参数 列表中。

如果存储器缓冲器接口328确定存储器对象创建指令未指定应针对存储器对象启用 直接模式，那么存储器缓冲器接口328可针对将创建的存储器对象分配共享存储器空间 336(360)，针对共享存储器空间336停用共享存储器高速缓冲存储器相干模式(362)，且 返回对所创建的存储器对象的参考(364)。相反，如果存储器缓冲器接口328确定存储器 对象创建指令指定应针对存储器对象启用直接模式，那么存储器缓冲器接口328可针对 将创建的存储器对象分配共享存储器空间336(366)，针对共享存储器空间336启用共享 存储器高速缓冲存储器相干模式(368)，且返回对所创建的存储器对象的参考(370)。

在一些实例中，存储器缓冲器接口328可使存储器对象返回到调用进程(例如，主机 进程320)，之后分配存储器316中的物理存储器空间以存储数据。在此些实例中，存储 器缓冲器接口328可包含所返回的存储器对象中的直接模式属性。随后，当在稍后时间 为存储器对象分配存储器316时，存储器缓冲器接口328或另一模块可基于存储器对象 的直接模式属性而启用或停用共享存储器空间高速缓冲存储器相干模式。

图23到28说明根据本发明的GPU可用来处理直接模式和高速缓存模式加载和存 储指令的实例性技术。如上文所论述，在一些实例中，用于任务330的源代码可经编译 以使得经编译的代码可包含高速缓存模式指令和直接模式指令两者，以便支持直接存储 器对象和高速缓冲存储器对象两者。高速缓存模式指令可使用与基础存储器相关联的高 速缓冲存储器的高速缓存服务来相对于存储器执行读取和写入操作，且直接模式指令可 不使用与基础存储器相关联的高速缓冲存储器的高速缓存服务来相对于存储器执行读 取和写入操作。高速缓存模式指令可或者在本文中被称作非直接模式指令。加载和存储 指令可或者在本文中分别被称作读取指令和写入指令。

在一些实例中，加载或存储指令的高速缓存模式版本以及加载或存储指令的直接模 式版本可为不同的指令，例如各自具有不同的操作代码，即操作码。在进一步的实例中， 加载或存储指令的高速缓存模式版本以及加载或存储指令的直接模式版本可为相同的 指令，例如两者都具有相同的操作码。在此实例中，向指令提供的参数可指定所述指令 是高速缓存模式还是直接模式。

图23是说明根据本发明的用于处理高速缓存模式指令和直接模式指令的实例性技 术的流程图。在一些实例中，图20的计算系统310可用于实施图23中所示的实例性技 术。在图23的实例中，直接模式被称作绕过高速缓冲存储器模式，且直接模式指令对 应于绕过高速缓冲存储器模式指令。GPU高速缓冲存储器控制模块334接收指定存储器 位置以及是否启用绕过高速缓冲存储器模式的加载指令(372)。GPU高速缓冲存储器控 制模块334确定加载指令是否指定启用绕过高速缓冲存储器模式(374)。在一些情况下， GPU高速缓冲存储器控制模块334可基于指令的类型(例如，指令的操作码)来确定加载 指令是否指定启用绕过高速缓冲存储器模式。在额外情况下，GPU高速缓冲存储器控制 模块334可基于与加载指令一起包含的指示是否启用绕过高速缓冲存储器模式的参数来 确定加载指令是否指定启用绕过高速缓冲存储器模式。如果GPU高速缓冲存储器控制 模块334确定不启用绕过高速缓冲存储器模式，那么GPU高速缓冲存储器控制模块334 从高速缓冲存储器(例如，GPU高速缓冲存储器332)在与加载指令中所指定的存储器位 置相关联的高速缓冲存储器位置处检索数据(376)。另一方面，如果GPU高速缓冲存储 器控制模块334确定启用绕过高速缓冲存储器模式，那么GPU高速缓冲存储器控制模 块334从存储器(例如，共享存储器空间336)在加载指令中所指定的存储器位置处检索 数据(378)。

图24是说明根据本发明的用于处理高速缓存模式指令和直接模式指令的另一实例 性技术的流程图。在一些实例中，图20的计算系统310可用于实施图24中所示的实例 性技术。在图24的实例中，直接模式被称作绕过高速缓冲存储器模式，且直接模式指 令对应于绕过高速缓冲存储器模式指令。GPU高速缓冲存储器控制模块334接收指定存 储器位置、要存储的数据以及是否启用绕过高速缓冲存储器模式的存储指令(380)。GPU 高速缓冲存储器控制模块334确定存储指令是否指定启用绕过高速缓冲存储器模式 (382)。在一些情况下，GPU高速缓冲存储器控制模块334可基于指令的类型(例如，指 令的操作码)来确定存储指令是否指定启用绕过高速缓冲存储器模式。在额外情况下， GPU高速缓冲存储器控制模块334可基于与加载指令一起包含的指示是否启用绕过高速 缓冲存储器模式的参数来确定存储指令是否指定启用绕过高速缓冲存储器模式。如果 GPU高速缓冲存储器控制模块334确定不启用绕过高速缓冲存储器模式，那么GPU高 速缓冲存储器控制模块334将存储指令中所指定的数据存储在高速缓冲存储器(例如， GPU高速缓冲存储器332)中在与存储指令中所指定的存储器位置相关联的高速缓冲存 储器位置处(384)。另一方面，如果GPU高速缓冲存储器控制模块334确定启用绕过高 速缓冲存储器模式，那么GPU高速缓冲存储器控制模块334将存储指令中所指定的数 据存储到存储器(例如，共享存储器空间336)在加载指令中所指定的存储器位置处(386)。

图25是说明根据本发明的用于处理高速缓存模式指令和直接模式指令的另一实例 性技术的流程图。在一些实例中，图20的计算系统310可用于实施图25中所示的实例 性技术。GPU高速缓冲存储器控制模块334接收指定存储器位置、要存储的数据以及是 否启用直接模式的存储指令(388)。GPU高速缓冲存储器控制模块334将存储指令中所 指定的数据存储在高速缓冲存储器(例如，GPU高速缓冲存储器332)中在与存储指令中 所指定的存储器位置相关联的高速缓冲存储器位置处(390)。GPU高速缓冲存储器控制 模块334基于存储指令中的指定是否启用直接模式的信息而确定是否启用直接模式 (392)。在一些实例中，是否启用直接模式的信息可为指令的类型(例如，用于指令的操 作码)和/或与指令一起包含的指定是否针对所述指令启用直接模式的参数。如果未启用 直接模式，那么GPU高速缓冲存储器控制模块334不执行直接高速缓冲存储器回写操 作(394)。另一方面，如果启用直接模式，那么GPU高速缓冲存储器控制模块334执行 直接高速缓冲存储器回写操作(396)。

图26是说明根据本发明的用于处理高速缓存模式指令和直接模式指令的另一实例 性技术的流程图。在一些实例中，图20的计算系统310可用于实施图26中所示的实例 性技术。GPU高速缓冲存储器控制模块334接收指定存储器位置以及是否启用直接模式 的加载指令(398)。GPU高速缓冲存储器控制模块334基于加载指令中的指定是否启用 直接模式的信息而确定是否启用直接模式(400)。在一些实例中，是否启用直接模式的信 息可为指令的类型(例如，用于指令的操作码)和/或与指令一起包含的指定是否针对所述 指令启用直接模式的参数。如果未启用直接模式，那么GPU高速缓冲存储器控制模块 334不清洗高速缓冲存储器且使高速缓冲存储器无效(402)。GPU高速缓冲存储器控制模 块334在可在高速缓冲存储器中得到数据的情况下从高速缓冲存储器(例如，GPU高速 缓冲存储器332)检索加载指令中所指定的数据，或在高速缓冲存储器中得不到数据的情 况下从基础存储器检索加载指令中所指定的数据(404)。如果启用直接模式，那么GPU 高速缓冲存储器控制模块334清洗高速缓冲存储器且使高速缓冲存储器无效(406)。GPU 高速缓冲存储器控制模块334从基础存储器检索加载指令中所指定的数据(408)。高速缓 冲存储器不返回数据，因为高速缓冲存储器已被清洗且无效。

图27是说明根据本发明的可用于图20的计算系统310中的实例性GPU420的方框 图。在一些实例中，GPU420可用于实施图20中所说明的GPU314。GPU420包含GPU 处理模块422、GPU高速缓冲存储器控制模块424、GPU高速缓冲存储器426、高速缓 冲存储器总线428和绕过总线430。GPU处理模块422经由高速缓冲存储器总线428通 信地耦合到GPU高速缓冲存储器控制模块424。GPU处理模块422还经由绕过总线430 通信地耦合到存储器316。GPU高速缓冲存储器控制模块424和GPU高速缓冲存储器 426实质上类似于图20中的GPU高速缓冲存储器控制模块334和GPU高速缓冲存储器 332，且将不进一步详细地描述。GPU处理模块422包含处理元件432和总线控制器434。 处理元件432经配置以将加载和存储指令发布到总线控制器434。

总线控制器434可经配置以经由高速缓冲存储器总线428和绕过总线430将加载和 存储指令转发到适当位置。总线控制器434可经配置以基于加载或存储指令中的指示指 令是直接模式指令还是高速缓存模式指令的信息而在直接模式或非直接模式中操作。当 总线控制器434经配置以在非直接模式(即，高速缓存模式)中操作时，总线控制器434 可使用高速缓冲存储器总线428将加载和存储指令转发到GPU高速缓冲存储器控制模 块424以供执行。另一方面，当总线控制器434经配置以在直接模式中操作时，总线控 制器434可使用绕过总线430将加载和存储指令转发到存储器316以供执行。

图28是说明根据本发明的用于处理高速缓存模式指令和直接模式指令的实例性技 术的流程图。在一些实例中，图27的GPU420可用于实施图28中所示的实例性技术。 总线控制器434接收加载或存储指令(440)。总线控制器434基于加载或存储指令中的指 定是否启用直接模式的信息而确定是否启用直接模式(442)。在一些实例中，是否启用直 接模式的信息可为指令的类型(例如，用于指令的操作码)和/或与指令一起包含的指定是 否针对所述指令启用直接模式的参数。如果总线控制器434确定未启用直接模式，那么 总线控制器434使用高速缓冲存储器总线428将所接收的指令转发到GPU高速缓冲存 储器控制模块424(444)。否则，如果总线控制器434确定启用直接模式，那么总线控制 器434使用绕过总线430将所接收的指令转发到存储器316(446)。

图29是说明根据本发明的用于执行由在主机装置上执行的进程发布的存储器对象 创建指令的另一实例性技术的流程图。在一些实例中，图20的计算系统310可用于实 施图29中所示的实例性技术。存储器对象创建指令可为缓冲器对象创建指令或图像对 象创建指令。存储器缓冲器接口328接收存储器对象创建指令(448)。存储器缓冲器接口 328确定存储器对象创建指令是否指定应针对存储器对象启用直接模式(450)。举例来说， 存储器缓冲器接口328可确定直接旗标参数是否包含于用于存储器对象创建指令的参数 列表中。

如果存储器缓冲器接口328确定存储器对象创建指令未指定应针对存储器对象启用 直接模式，那么存储器缓冲器接口328将用于所创建的存储器对象的直接模式属性设定 为指示不启用直接模式的值(例如，“false”)(452)。另一方面，如果存储器缓冲器接口 328确定存储器对象创建指令指定应针对存储器对象启用直接模式，那么存储器缓冲器 接口328将用于所创建的存储器对象的直接模式属性设定为指示启用直接模式的值(例 如，“true”)(454)。在一些实例中，存储器对象的直接模式属性可由主机装置312和/或 GPU314使用以在存取存储于特定存储器对象中的数据时确定是执行高速缓存模式还是 直接模式读取和写入操作。

在一些实例中，主机进程320和/或任务330可能希望将一些存储器对象编程为直接 存储器对象，且将其它对象编程为高速缓冲存储器对象，即，非直接存储器对象。在一 些实例中，本发明的技术可包含专用编译技术，所述专用编译技术允许经编译任务330 相对于高速缓冲存储器对象以及直接存储器对象两者执行读取和写入操作。第一实例性 编译技术可将给定读取操作或写入操作编译为指令序列。所述指令序列可检查从其进行 读取或对其进行写入的存储器对象的直接模式属性的值，且基于所述直接模式属性的值 来确定是执行高速缓存模式指令还是直接模式指令。第二实例性编译技术可使用源代码 中的指示存储器对象是否为直接模式对象的信息，以选择高速缓存模式指令或直接模式 指令来用于经编译代码中以用于存取存储器对象。

根据第一实例性编译技术，编译器可编译用于任务330的源代码，以使得用于任务 330的经编译代码包含根据以下实例性伪代码的读取序列：

其中"isImmediate"表示将从其读取数据的存储器对象的布尔直接模式属性， "immediate_read(...)"表示直接模式读取指令，且"cached_read(...)"表示高速缓存模式读 取指令。

GPU高速缓冲存储器控制模块334可通过(例如)在从GPU高速缓冲存储器332读取 数据之前使GPU高速缓冲存储器332(在使用的情况下)无效而处理immediate_read(...) 指令。GPU高速缓冲存储器控制模块334可通过(例如)以正常方式从GPU高速缓冲存储 器332读取数据(例如，而不在执行读取之前使高速缓冲存储器无效)来处理cached_read (...)指令。

根据第一实例性编译技术，编译器可编译用于任务330的源代码，以使得用于任务 330的经编译代码包含根据以下实例性伪代码的写入序列：

其中"isImmediate"表示将向其写入数据的存储器对象的布尔直接模式属性， "immediate_write(...)"表示直接模式写入指令，且"cached_write(...)"表示高速缓存模式写 入指令。

在一些实例中，GPU高速缓冲存储器控制模块334可通过在使用高速缓冲存储器的 情况下针对GPU高速缓冲存储器332使用通写模式来处理immediate_write(...)指令。在 进一步的实例中，GPU高速缓冲存储器控制模块334可通过在使用高速缓冲存储器的情 况下将数据写入到GPU高速缓冲存储器332来处理immediate_write(...)指令，且响应于 将数据写入到GPU高速缓冲存储器332而针对GPU高速缓冲存储器332执行高速缓冲 存储器清洗。GPU高速缓冲存储器控制模块334可通过以正常方式将数据写入到GPU 高速缓冲存储器332(例如，而不响应于写入操作而清洗高速缓冲存储器)来处理 cached_write(...)指令。

图30是说明根据GPU可如何处理根据上述第一编译技术而编译的指令序列的流程 图。在一些实例中，图30中所说明的技术可用于实施上文针对读取和写入序列而提供 的实例性伪代码。任务330开始读取序列或写入序列(456)。举例来说，任务330可在任 务330到达任务330的执行中的一点时(其中应发生对特定存储器对象的读取或写入指令) 开始读取序列或写入序列。任务330存取与将从其读取数据或将向其写入数据的存储器 对象相关联的直接模式属性(458)。任务330确定用于存储器对象的属性是否被设定为指 示启用直接模式的值(例如，"true")(460)。如果任务330确定用于存储器对象的属性被设 定为指示不启用直接模式的值，那么任务330使用高速缓存读取或写入指令针对存储器 对象执行高速缓存读取或写入操作(462)。否则，如果任务330确定用于存储器对象的属 性被设定为指示启用直接模式的值，那么任务330使用直接读取或写入指令针对存储器 对象执行直接读取或写入操作(464)。

根据第二实例性编译技术，当编译源代码时，编译器可具有对指示是否针对任务330 从其进行读取或对其进行写入的特定存储器对象而启用直接模式的信息的存取权。当任 务330从所述特定存储器对象进行读取或向其进行写入时，编译器使用此信息来编译用 于任务330的源代码，以在高速缓存模式读取和写入指令或直接模式读取和写入指令之 间进行选择。

在一些实例中，指示是否针对特定存储器对象启用直接模式的信息可为指示是否针 对由任务330的源代码存取的一个或一个以上存储器对象启用直接模式的编译时间属 性。举例来说，用于任务330的源代码(例如，内核源代码)可包含指示是否针对由任务 330使用的一个或一个以上存储器对象启用直接模式的编译时间属性。在一些情况下， 编译时间属性可采取OpenCL属性限定符(例如，_cl_immediate)的形式。所述属性限定 符可与一个或一个以上特定存储器对象和/或存储于所述一个或一个以上存储器对象内 的一个或一个以上变量相关联。当属性限定符与特定存储器对象相关联时，那么编译器 可确定针对存储器对象启用直接模式。类似地，当属性限定符不与特定存储器对象相关 联时，那么编译器可确定不针对存储器对象启用直接模式。使用此属性可减少编译器的 工作且潜在地减小内核大小。在一些实例中，软件应用可将直接缓冲器的使用限制到需 要此些缓冲器的情况。在此些实例中，是否使用直接缓冲器的决策可为编译时间决策。

如果编译时间属性指示针对与共享存储器空间336相关联的存储器对象启用直接模 式，那么编译器可编译任务330以使得用于任务330的经编译代码包含用于相对于共享 存储器空间336发生的读取或写入操作的直接模式读取和/或写入指令。否则，如果未针 对与共享存储器空间336相关联的存储器对象启用直接模式，那么编译器可编译任务330 以使得用于任务330的经编译代码不包含用于相对于共享存储器空间336发生的读取或 写入操作的直接模式读取和/或写入指令。举例来说，编译器可编译任务330以使得用于 任务330的经编译代码包含用于相对于共享存储器空间336发生的读取或写入操作的高 速缓存读取和/或写入指令。

图31是说明根据本发明的用于编译用于任务的源代码的实例性技术的流程图。在 一些实例中，使用图31中的技术编译的所得代码对应于图20中的任务330。在图31的 实例性技术中，任务330被称作内核。编译器处理由存储器对象实施的内核自变量(466)。 编译器确定存储器对象是否为直接模式存储器对象(468)。在一些实例中，编译器可基于 内核的源代码中所包含的信息(例如，与内核自变量相关联的编译时间属性)来确定存储 器对象是否为直接模式存储器对象。如果编译器确定存储器对象不是直接模式存储器对 象，那么编译器使用高速缓存读取和写入指令来编译与特定内核自变量相关联的读取操 作和写入操作(470)。另一方面，如果编译器确定存储器对象是直接模式存储器对象，那 么编译器使用直接模式读取和写入指令来编译与特定内核自变量相关联的读取操作和 写入操作(472)。

图32是说明根据本发明的可由GPU用来选择性地使用高速缓存服务的实例性技术 的流程图。举例来说，所述技术可允许GPU响应于接收到指定是否应使用高速缓存服 务来用于相对于存储器的存储器空间执行读取操作和写入操作中的至少一者的信息而 选择性地使用与存储器相关联的GPU高速缓冲存储器来相对于所述存储器空间执行读 取操作和写入操作中的至少一者。在一些实例中，图20中所说明的GPU314和/或图27 中所说明的GPU420可用于实施图32中所说明的技术。

GPU314接收读取指令或写入指令中的至少一者进行处理(474)。所接收的指令可指 令GPU314相对于存储器的存储器空间执行读取操作和写入操作中的至少一者。GPU 314接收指定是否应使用高速缓存服务来用于相对于存储器空间执行读取操作和写入操 作中的至少一者的信息(476)。在一些实例中，高速缓冲存储器模式信息可包含在所接收 的指令内。在进一步的实例中，高速缓冲存储器模式信息可为与存储器空间相关联的存 储器对象的直接模式属性。GPU314确定是否基于高速缓冲存储器模式信息(478)来使用 高速缓存服务。响应于接收到指定应使用高速缓存服务来用于执行所接收的指令的信 息，GPU314可使用高速缓存服务来执行所接收的指令(480)。响应于接收到指定不应使 用高速缓存服务来用于执行所接收的指令的信息，GPU314可不使用高速缓存服务来执 行所接收的指令(482)。在一些实例中，GPU314可使用图23到28以及30中所说明的 技术中的一者或一者以上来实施决策框478以及过程框480和482中的一者或一者以上。 在一些情况下，可使用GPU高速缓冲存储器控制模块或存储器管理单元(例如，图20 中所说明的GPU高速缓冲存储器控制模块334)来实施图32中所示的技术。在额外情况 下，可使用总线控制器(例如，图27中所说明的总线控制器434)来实施图32中所示的 技术。

在一些实例中，为了实施直接存储器对象，GPU ALU可经设计以执行ALU指令， 所述ALU指令使指令中所指定的全局存储器高速缓冲存储器和/或全局存储器高速缓冲 存储器的特定部分无效。一般来说，主机装置312可使用现有的GPU能力来实施直接 存储器对象。

现在将进一步详细地论述本发明中所描述的带外信令技术(例如，本文中所描述的消 息传递技术)和本发明中所描述的直接存储器对象的各种使用情况。根据第一使用情况， 可将带外信令用作独立特征，而没有必要除了带外信令技术之外还使用直接存储器对 象。带外信令可用于同步以及快速地传递相对少量的数据。在一些实例中，带外信令可 具有比直接存储器对象低的等待时间，但具有比直接存储器对象低的带宽。

还可根据存储器分配操作的第一使用情况使用带外信令。举例来说，GPU可使用带 外信令来请求主机CPU分配新的缓冲器。GPU还可使用还使用带外信令来向主机CPU 指定所请求的缓冲器长度。作为另一实例，CPU可在分配缓冲器之后使用带外信令以向 GPU发送指定用于缓冲器的存储器位置的指针。

还可根据远程过程调用的第一使用情况(其中将交换少量的数据)使用带外信令。举 例来说，在其中在计算装置内的计算单元上执行的内核使用RPC以在同一计算装置或在 另一计算装置中的另一计算单元上启动另一内核的情况下，用于所述RPC的数据可能被 存储在启动的计算单元的本地存储器中。本发明的带外信令技术可用于将数据从启动的 计算单元的本地存储器传送到执行新启动的内核的计算单元的本地存储器。

还可根据进展报告的第一使用情况使用带外信令。举例来说，GPU可使用带外信令 来向主机CPU报告当前任务的完成百分比。

还可根据错误报告的第一使用情况使用带外信令。举例来说，GPU可使用带外信令 来向主机CPU报告错误代码。

还可根据辅助上下文切换的第一使用情况使用带外信令。举例来说，主机CPU可 使用带外信令来请求GPU保存状态以准备上下文切换。

根据第二使用情况，可将直接存储器对象用作独立特征，而没有必要除了直接存储 器对象之外还使用带外信令。举例来说，可使用直接缓冲器来完成相对大量数据的交换。 直接缓冲器可不仅含有数据，而且还有同步标记。在此情况下，数据产生器可首先将数 据写入到缓冲器，且随后写入同步标记，所述同步标记向消费者指示数据的就绪和/或位 置。消费者在先验决定的位置中寻找同步数据，例如通过轮询此存储器位置而在缓冲器 的标头区段中寻找同步数据。一旦获得同步标记，消费者便读取数据。可将类似技术应 用于直接图像对象。

可针对这些技术使用多种同步协议。举例来说，可将同步标记嵌入在数据缓冲器中， 或可定位在单独的缓冲器中。可将此些技术应用于使用可变长度编码或行程长度编码方 案压缩的经压缩数据的传输。

根据第三使用情况，可与带外信令联合地使用直接存储器对象(例如)以完成对相对 大量数据的交换。在此情况下，可在直接存储器对象存储数据的同时将带外信令用于同 步。举例来说，数据产生器可将数据放置到直接缓冲器中，且使用带外信令向消费者通 知数据的就绪以及位置和/或大小。在流控制情形中，消费者读取数据且通知产生器可再 次使用缓冲器。还可使用带外信令来完成所述通知。

此些技术可用于需要流控制数据流水线的算法中。对于主机CPU和GPU，可使用 此些技术(例如)用于诊断记录。对于多个OpenCL计算装置，可使用这些技术将多个装 置连接到异步流控制数据管线中。这可允许应用被分解为更适合于每一CPU或GPU的 若干块，启动多个装置上的各种管线处理级，且/或从主机CPU卸载大多数或甚至全部 的数据同步。

在一些实例中，本发明的技术可为使用命令队列起始任务的多处理器计算平台提供 消息传递接口，所述消息传递接口促进在主机装置上执行的进程与在计算装置上执行的 任务之间的消息的发送和接收。在一些情况下，所述计算装置可为GPU。在额外情况下， 所述计算装置可为由跨平台、跨供应商、异类计算平台API界定的任何类型的计算装置。

在进一步的实例中，本发明的技术可提供包含可由主机装置存取的一个或一个以上 寄存器的GPU。所述一个或一个以上寄存器可经配置以促进在GPU上执行的任务与在 不同于GPU的装置上执行的进程之间的消息传递。

在额外实例中，本发明的技术可提供允许创建直接存储器对象的存储器缓冲器接 口。所述直接存储器对象可用于实施非可高速缓存共享存储器空间和/或高速缓冲存储器 相干共享存储器空间，以便在于计算装置上执行的任务正在计算装置上执行时在于主机 装置上执行的进程与所述任务之间共享数据。在一些情况下，所述计算装置可为图形处 理单元(GPU)。在额外情况下，所述计算装置可为由跨平台、跨供应商、异类计算平台 API界定的任何类型的计算装置。

在又进一步的实例中，本发明的技术可提供包含用于共享存储器空间的高速缓冲存 储器的GPU，所述共享存储器空间可被选择性地停用以便提供非可高速缓存共享存储器 空间。在额外实例中，本发明的技术可提供包含高速缓冲存储器相干模式的GPU，所述 高速缓冲存储器相干模式可被选择性地启用以提供高速缓冲存储器相干的共享存储器 空间。

本发明中所描述的技术可至少部分以硬件、软件、固件，或其组合来实施。举例来 说，所描述技术的各种方面可实施于以下各项内：一个或一个以上处理器(包含一个或一 个以上微处理器)、数字信号处理器(DSP)、专用集成电路(ASIC)、现场可编程门阵列 (FPGA)或任何其它等效集成或离散逻辑电路，以及此些组件的任何组合。术语“处理器” 或“处理电路”通常可指代单独或与其它逻辑电路组合的前述逻辑电路中的任一者，或 任何其它等效电路，例如执行处理的离散硬件。

此类硬件、软件和固件可实施于相同装置内或单独装置内，以支持本发明中所描述 的各种操作和功能。另外，所描述的单元、模块或组件中的任一者可一起实施，或单独 地实施为离散但可互操作的逻辑装置。将不同特征描绘为模块或单元意在突出不同功能 方面，且未必暗示必须通过单独的硬件或软件组件来实现此些模块或单元。而是，与一 个或一个以上模块或单元相关联的功能性可由单独的硬件、固件和/或软件组件执行，或 集成于共用或单独的硬件或软件组件内。

本发明中所描述的技术还可存储、体现或编码于计算机可读媒体(例如，存储指令的 计算机可读存储媒体)中。嵌入或编码于计算机可读媒体中的指令(例如)在所述指令被一 个或一个以上处理器执行时可致使一个或一个以上处理器执行本文中所描述的技术。计 算机可读存储媒体可包含：随机存取存储器(RAM)、只读存储器(ROM)、可编程只读存 储器(PROM)、可擦除可编程只读存储器(EPROM)、电可擦除可编程只读存储器 (EEPROM)、快闪存储器、硬盘、CD-ROM、软盘、盒式磁带、磁性媒体、光学媒体或 有形的其它计算机可读存储媒体。

计算机可读媒体可包含对应于有形存储媒体(例如上文所列举的有形存储媒体)的计 算机可读存储媒体。计算机可读媒体还可包括通信媒体，通信媒体包括(例如)根据通信 协议促进将计算机程序从一处传送到另一处的任何媒体。以此方式，短语“计算机可读 媒体”一般可对应于(1)非暂时性的有形计算机可读存储媒体和(2)例如暂时性信号或载波 等非有形的计算机可读通信媒体。