CN101231585B - 定义用于目标平台结构的并行处理操作方法 - Google Patents

定义用于目标平台结构的并行处理操作方法 Download PDF

Info

Publication number
CN101231585B
CN101231585B CN2008100070263A CN200810007026A CN101231585B CN 101231585 B CN101231585 B CN 101231585B CN 2008100070263 A CN2008100070263 A CN 2008100070263A CN 200810007026 A CN200810007026 A CN 200810007026A CN 101231585 B CN101231585 B CN 101231585B
Authority
CN
China
Prior art keywords
virtual
thread
cta
program
instruction
Prior art date
Legal status (The legal status is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the status listed.)
Active
Application number
CN2008100070263A
Other languages
English (en)
Other versions
CN101231585A (zh
Inventor
约翰·R·尼科尔斯
亨利·P·莫尔顿
拉尔斯·S·尼兰
伊恩·A·巴克
理查德·C·约翰逊
罗伯特·S·格兰维尔
贾扬特·B·科尔希
Current Assignee (The listed assignees may be inaccurate. Google has not performed a legal analysis and makes no representation or warranty as to the accuracy of the list.)
Nvidia Corp
Original Assignee
Nvidia Corp
Priority date (The priority date is an assumption and is not a legal conclusion. Google has not performed a legal analysis and makes no representation as to the accuracy of the date listed.)
Filing date
Publication date
Application filed by Nvidia Corp filed Critical Nvidia Corp
Publication of CN101231585A publication Critical patent/CN101231585A/zh
Application granted granted Critical
Publication of CN101231585B publication Critical patent/CN101231585B/zh
Active legal-status Critical Current
Anticipated expiration legal-status Critical

Links

Images

Classifications

    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00—Arrangements for software engineering
    • G06F8/40—Transformation of program code
    • G06F8/41—Compilation
    • G06F8/45—Exploiting coarse grain parallelism in compilation, i.e. parallelism between groups of instructions
    • G06F8/456—Parallelism detection
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00—Arrangements for software engineering
    • G06F8/40—Transformation of program code
    • G06F8/41—Compilation
    • G06F8/44—Encoding
    • G06F8/445—Exploiting fine grain parallelism, i.e. parallelism at instruction level
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F9/00—Arrangements for program control, e.g. control units
    • G06F9/06—Arrangements for program control, e.g. control units using stored programs, i.e. using an internal store of processing equipment to receive or retain programs
    • G06F9/44—Arrangements for executing specific programs
    • G06F9/455—Emulation; Interpretation; Software simulation, e.g. virtualisation or emulation of application or operating system execution engines
    • G06F9/45533—Hypervisors; Virtual machine monitors

Abstract

一种虚拟结构和指令集支持明确的并行线程计算。所述虚拟结构定义：虚拟处理器，其支持具有不同虚拟线程之间的不同级别的数据共享和协调(例如，同步)的多个虚拟线程的并发执行；以及虚拟执行驱动器，其控制所述虚拟处理器。所述虚拟处理器的虚拟指令集结构用来定义虚拟线程的行为，且包含与并行线程行为(例如，数据共享和同步)有关的指令。通过使用所述虚拟平台，编程人员能够开发出其中虚拟线程并发地执行以处理数据的应用程序；虚拟翻译器和驱动器以对于所述编程人员透明的方式来调适所述应用程序代码以适应其所执行于的特定硬件。

Description

定义用于目标平台结构的并行处理操作方法
技术领域
本发明大体上涉及并行处理，且确切地说涉及用于并行线程计算的虚拟结构和指令集。 
背景技术
在并行处理中，多个处理单元(例如，多个处理器芯片或单个芯片内的多个处理核心)同时操作以处理数据。可使用此种系统来解决便于分解成多个部分的问题。一个实例是图像过滤，其中根据输入图像的某一数目的像素来计算输出图像的每个像素。每个输出像素的计算一般独立于其它所有输出像素，因此不同的处理单元可并行计算不同的输出像素。其它许多类型的问题也适于并行分解。总的来说，N向(N-way)并行执行可将此种问题的解决方案加速大约N倍。 
另一类问题在并行的执行线程可彼此协调的情况下则适于并行处理。一个实例是快速傅里叶变换(FFT)，其是一种递归算法：其中在每个级，对前一级的输出执行计算，以便产生新的值，所述新的值用于对下一级的输入，如此继续直到到达输出级为止。单个执行线程可执行多个级，只要所述线程能可靠地获得来自前一级的输出数据即可。如果要在多个线程之间划分任务，那么必须提供某种协调机制，以便(例如)线程不会试图读取尚未写入的输入数据。(在2005年12月15日申请的共同转让、共同待决的第11/303,780号美国专利申请案中描述了这个问题的一个解决方案)。 
然而，并行处理系统的编程可能较为困难。通常要求编程人员知道可用的处理单元的数目及其能力(指令集、数据寄存器数目、互连件等)，以便建立处理单元可实际执行的代码。虽然机器特定的编译器可在这个方面提供相当大的帮助，但仍有必要在每次将代码移植(port)到不同处理器时重新编译代码。 
此外，并行处理结构的各种方面正在迅速发展。举例来说，正在陆续开发新的平台结构、指令集和编程模型。随着并行结构的各种方面(例如，编程模型或指令集)从一代改变成下一代，必须也相应地改变应用程序、软件库、编译器和其它软件及工具。这种不稳定性可能会为并行处理代码的开发和维护增加相当大的额外开销。 
当需要线程之间协调时，并行编程变得更加困难。编程人员必须确定在特定处理器或计算机系统中有可用于支持(或仿真)线程间通信的机制，且必须编写利用可用机制的代码。由于不同的计算机系统上的可用和/或最优机制一般不同，所以这种类型的并行代码一般不可移植；必须针对其运行于的每种硬件平台重编写并行代码。 
此外，除了为处理器提供可执行代码，编程人员还必须为“主”处理器提供控制代码，所述控制代码协调各种处理单元的操作，例如向每个处理单元指示要执行哪个程序和处理哪些输入数据。此控制代码通常专用于特定的主处理器和处理器间的通信协议，且如果要替换不同的主处理器则通常必须重编写所述控制代码。 
编译和重新编译并行处理代码中的困难会使用户不愿随着计算技术的发展而升级其系统。因此，将需要使已编译的并行处理代码脱离特定的硬件平台，并提供一种作为并行应用程序和工具的目标的稳定的并行处理结构和指令集。 
发明内容
本发明的实施例提供一种用于并行线程计算的虚拟结构和虚拟指令集。所述虚拟并行结构定义：一种虚拟处理器，其支持具有不同虚拟线程之间的不同级别的数据共享和协调(例如，同步)的多个虚拟线程的并发执行；以及一种虚拟执行驱动器，其控制所述虚拟处理器。虚拟处理器的虚拟指令集结构用来定义虚拟线程的行为，且包含与并行线程行为(例如，数据共享和同步)有关的指令。通过使用虚拟并行平台，编程人员能够开发出其中虚拟线程并发执行以处理数据的应用程序。应用程序可用高度可移植的中间形式存储和分布，所述形式例如为瞄准虚拟并行平台的程序代码。在安装时间或运行时间，硬件特定的虚拟指令翻译器和执行驱动器使中间形式的应用程序代码适应于其执行于的特定硬件。因此，应用程序更加可移植且更容易开发，因为开发过程独立于特定处理硬件。 
根据本发明的一个方面，一种用于定义并行处理操作的方法包含提供第一程序代码，所述第一程序代码定义要针对协作的虚拟线程的阵列中的若干虚拟线程中的每一者执行的操作序列。将第一程序代码编译成虚拟线程程序，所述虚拟线程程序定义要针对阵列中的代表性虚拟线程执行的每线程指令的序列，且所述每线程指令的序列包含至少一个指令，所述指令定义代表性虚拟线程与所述阵列中的一个或一个以上其它虚拟线程之间的协作行为。存储虚拟线程程序(例如，存储在存储器中或盘上)，且可随后将所述虚拟线程程序翻译成符合目标平台结构的指令序列。 
此外，也可提供第二程序代码以定义协作的虚拟线程的阵列，所述阵列适合于处理输入数据集以产生输出数据集，其中阵列中的每个虚拟线程并发执行虚拟线程程序。有利地将第二程序代码转换成虚拟函数库中的函数调用序列，其中所述库包含初始化和致 使执行协作的虚拟线程的阵列的虚拟函数。也可存储这个函数调用序列。接着，可将所存储的虚拟线程程序和函数调用序列翻译成可在目标平台结构上执行的程序代码，其中所述可执行的程序代码定义一个或一个以上执行所述协作的虚拟线程的阵列的平台线程。可在符合目标平台结构的计算机系统上执行可执行的程序代码，因而产生输出数据集，可将所述输出数据集存储在存储媒体(例如，计算机存储器、盘等)中。 
如已提到的，虚拟线程程序代码中的每线程指令的序列有利地包含至少一个指令，所述指令定义代表性虚拟线程与所述阵列中的一个或一个以上其它虚拟线程之间的协作行为。举例来说，所述每线程指令的序列可包含：在序列中的特定点处暂停执行代表性虚拟线程的操作直到例如其它虚拟线程中的一者或一者以上达到所述特定点为止的时间的指令，使代表性虚拟线程将数据存储在其它虚拟线程中的一者或一者以上可存取的共享存储器中的指令，使代表性虚拟线程自动读取和更新存储在其它虚拟线程中的一者或一者以上可存取的共享存储器中的数据的指令，等等。 
所述虚拟线程程序也可包含变量定义语句，所述语句在若干虚拟状态空间之一中定义变量，其中不同的虚拟状态空间对应于虚拟线程之间的不同的数据共享模式。在一个实施例中，至少支持每线程不共享模式和全局共享模式。在其它实施例中，也可支持额外模式，例如虚拟线程的一个阵列内的共享模式和/或多个虚拟线程阵列之间的共享模式。 
根据本发明的另一方面，一种用于操作目标处理器的方法包含提供输入程序代码。所述输入程序代码包含第一部分，所述第一部分定义要对虚拟线程阵列中的若干虚拟线程中的每一者执行的操作序列，所述阵列适合于处理输入数据集以产生输出数据集，且所述输入程序代码还包含第二部分，所述第二部分定义虚拟线程的阵列的维度。将所述输入程序代码的第一部分编译成虚拟线程程序，所述虚拟线程程序定义要对阵列中的代表性虚拟线程执行的每线程指令的序列。所述每线程指令的序列包含至少一个指令，所述指令定义所述代表性虚拟线程与所述阵列中的一个或一个以上其它虚拟线程之间的协作行为。将所述输入程序代码的第二部分转换成对虚拟函数库的函数调用序列，其中所述库包含初始化和致使执行协作的虚拟线程的阵列的虚拟函数。将所述虚拟线程程序和函数调用序列翻译成可在目标平台结构上执行的程序代码，其中所述可执行的程序代码定义执行协作的虚拟线程的阵列的一个或一个以上真实线程。在符合目标平台结构的计算机系统上执行所述可执行程序代码，因而产生输出数据集，可将所述输出数据集存储在存储媒体中。 
在一些实施例中，可在两个或两个以上维度上定义虚拟线程的阵列。此外，所述输入程序代码的第二部分还可包含定义虚拟线程阵列的栅格的一个或一个以上维度的函数调用，其中栅格中的每个阵列均将被执行。 
可使用任何目标平台结构。在一些实施例中，目标平台结构包含主处理器和协处理器。在翻译期间，可将虚拟线程程序翻译成可由定义于协处理器上的若干线程并行执行的程序代码，同时将函数调用序列翻译成对在主处理器上执行的用于协处理器的驱动程序的调用序列。在其它实施例中，目标平台结构包含中央处理单元(CPU)。在翻译期间，将虚拟线程程序和函数调用序列的至少一部分翻译成目标程序代码，所述目标程序代码使用少于虚拟线程数目的数目的CPU线程来执行虚拟线程阵列。 
根据本发明的又一实施例，一种用于操作目标处理器的方法包含获得虚拟线程程序，所述虚拟线程程序定义要针对虚拟线程阵列中的若干虚拟线程中的代表性虚拟线程执行的每线程指令序列，所述虚拟线程阵列适合于处理输入数据集以产生输出数据集。所述每线程指令序列包含至少一个指令，所述指令定义所述代表性虚拟线程与所述阵列中的一个或一个以上其它虚拟线程之间的协作行为。也获得定义虚拟线程阵列的维度的额外程序代码。将所述虚拟线程程序和额外的程序代码翻译成可在目标平台结构上执行的程序代码，其中所述可执行的程序代码定义执行所述虚拟线程阵列的一个或一个以上平台线程。在符合目标平台结构的计算机系统上执行可执行的程序代码，因而产生输出数据集且将所述输出数据集存储在存储器中。 
在一些实施例中，可通过接收用高级编程语言编写的源程序代码并编译所述源程序代码以产生虚拟线程程序来获得虚拟线程程序。或者，可从存储媒体读取虚拟线程程序或经由网络从远程计算机系统接收虚拟线程程序。应了解，所读取或接收的虚拟线程代码可能之前已经从高级语言编译，或者直接建立为符合虚拟指令集结构的代码。 
以下详细描述连同附图一起将提供对本发明的性质和优点的更好了解。 
附图说明
图1是根据本发明实施例的计算机系统的方框图。 
图2A和2B说明本发明实施例中使用的编程模型中的栅格、线程阵列和线程之间的关系。 
图3是根据本发明实施例的虚拟结构的方框图。 
图4是根据本发明实施例使用虚拟结构来操作目标处理器的概念模型。 
图5是列举由根据本发明实施例的虚拟指令集结构(ISA)定义的特殊变量的表格。 
图6是列举在根据本发明实施例的虚拟ISA中支持的变量类型的表格。 
图7是列举在根据本发明实施例的虚拟ISA中支持的虚拟状态空间的表格。 
图8A-8H是列举在根据本发明实施例的虚拟ISA中定义的虚拟指令的表格。 
图9是根据本发明实施例的用于使用虚拟指令翻译器的过程的流程图。 
图10是列举根据本发明实施例的在虚拟库中可供虚拟执行驱动器使用的函数的表格。 
具体实施方式
本发明实施例提供一种用于并行线程计算的虚拟结构和指令集。所述虚拟结构提供支持在不同线程之间具有多层数据共享和协调(例如，同步)的多个线程的并发执行的处理器模型，以及控制所述模型处理器的虚拟执行驱动器。用于定义处理线程的行为的虚拟指令集包含与并行线程行为有关的指令，例如允许某些线程间的数据共享的指令和要求不同线程在程序内的由编程人员指定的某些点处变得同步的指令。使用虚拟平台，编程人员可开发在其中执行并发、协作的线程以处理数据的应用程序。硬件特定的虚拟指令翻译器和执行驱动器使应用代码适应其执行于的特定硬件。因此，应用程序更具移植性并更易于开发，因为开发过程独立于特定处理硬件。 
1.系统概述
图1是根据本发明实施例的计算机系统100的方框图。计算机系统100包含中央处理单元(CPU)102，和包含存储器桥105的经由总线路径通信的系统存储器104。存储器桥105(其可以是(例如)Northbridge芯片)经由总线或其它通信路径106(例如，HyperTransport链路)连接到I/O(输入/输出)桥107。I/O桥107(其可以是(例如)Southbridge芯片)从一个或一个以上用户输入装置108(例如，键盘、鼠标)接收用户输入，并将输入经由路径106和存储器桥105转发到CPU 102。并行处理子系统112经由总线或其它通信路径113(例如，PCI Express或加速图形端口链路)耦合到存储器桥105；在一个实施例中，并行处理子系统112是图形子系统，其将像素递送到显示器装置110(例如，常规的基于CRT或LCD的监视器)。系统盘114也连接到I/O桥107。切换器116提供I/O桥107与例如网络适配器118和各种内插卡120及121等其它组件之间的连接。其它组件(未明确展示)包含USB或其它端口连接、CD驱动器、DVD驱动器等，也可连接到I/O桥107。使图1中的各个组件互连的通信路径可使用任何适宜的协议来实施，例如PCI(外围组件互连)、PCI Express(PCI-E)、AGP(加速图形端口)、HyperTransport，或任何其它总线或点对点通信协议，且不同装置之间的连接可使用此项 技术中已知的不同协议。 
并行处理子系统112包含并行处理单元(PPU)122和并行处理(PP)存储器124，其可(例如)使用例如可编程处理器、专用集成电路(ASIC)和存储器装置等的一个或一个以上集成电路装置来实施。PPU 122有利地实施包含一个或一个以上处理核心的高度并行处理器，所述处理核心的每一者能够并发执行大量(例如，数百个)线程。PPU 122可经编程以执行一大批计算，包含线性和非线性数据变换、视频和/或音频数据的过滤、建模(例如，应用物理定律来确定对象的位置、速度和其它属性)、图像渲染等。PPU 122可将数据从系统存储器104和/或PP存储器124转移到内部存储器中，处理所述数据，并将结果数据写回到系统存储器104和/或PP存储器124，在其中此类数据可由其它系统组件(包含，例如，CPU 102)存取。在一些实施例中，PPU 122是图形处理器，其还可经配置以执行与以下操作有关的各种任务：依据由CPU 102和/或系统存储器104经由存储器桥105和总线113供应的图形数据产生像素数据、与PP存储器124(其可用作图形存储器，包含(例如)常规帧缓冲器)交互以存储和更新像素数据、将像素数据递送到显示器装置110等。在一些实施例中，PP子系统112可包含一个作为图形处理器操作的PPU 122，和另一用于进行通用计算的PPU 122。PPU可相同或不同，且每一PPU可具有其自身的专用PP存储器装置。 
CPU 102作为系统100的主处理器操作，其控制和协调其它系统组件的操作。明确地说，CPU 102发布控制PPU 122的操作的命令。在一些实施例中，CPU 102将针对PPU122的命令流写入到命令缓冲器，所述命令缓冲器可处于系统存储器104、PP存储器124或CPU 102和PPU 122两者均可存取的另一存储位置中。PPU 122从命令缓冲器读取命令流，并与CPU 102的操作异步地执行命令。 
将了解，本文展示的系统是说明性的，且可能作出变化和修改。可视需要修改连接拓扑，包含桥的数目和配置。举例来说，在一些实施例中，系统存储器104直接而不通过桥连接到CPU 102，且其它装置经由存储器桥105和CPU 102与系统存储器104通信。在其它替代拓扑中，PP子系统112连接到I/O桥107而不连接到存储器桥105。在另外其它实施例中，I/O桥107和存储器桥105可能集成到单一芯片中。本文展示的特定组件是任选的；举例来说，可能支持任何数目的内插卡或外围装置。在一些实施例中，排除切换器116，且网络适配器118和内插卡120、121直接连接到I/O桥107。 
还可改变PPU 122到系统100的其余部分的连接。在一些实施例中，PP系统112被实施为可插入到系统100的扩充插槽中的内插卡。在其它实施例中，PPU可与总线桥(例 如，存储器桥105或I/O桥107)一起集成在单一芯片上。在另外其它实施例中，PPU 122的一些或所有元件可与CPU 102集成。 
PPU可具备任何量的本地PP存储器，包含无本地存储器，且可以任何组合使用本地存储器与系统存储器。举例来说，在统一存储器结构(UMA)实施例中，PPU 122可以是图形处理器；在此类实施例中，提供极少或不提供专门的图形存储器，且PPU 122将专门地或几乎专门地使用系统存储器。在UMA实施例中，PPU可集成到桥芯片中或提供成离散芯片，其中高速链路(例如，PCI-E)将PPU连接到桥芯片和系统存储器。 
还将了解，可(例如)通过在单一内插卡上包含多个PPU，通过将多个内插卡连接到路径113，和/或通过将一个或一个以上PPU直接连接到系统母板，而将任何数目的PPU包含在系统中。多个PPU可并行操作，从而以高于利用单一PPU可能实现的处理量处理数据。 
所属领域的技术人员还将了解，CPU和PPU可集成到单一装置中，且CPU和PPU可共享例如指令逻辑、缓冲器、高速缓冲存储器、存储器、处理引擎等各种资源；或者可针对并行处理和其它操作提供单独的资源。因此，本文描述为与PPU相关联的任何或所有电路和/或功能性也可在适当装备的CPU中实施并由所述适当装备的CPU执行。 
并入有PPU的系统可以多种配置和形状因素进行实施，包含台式计算机、膝上型计算机，或手提个人计算机、服务器、工作站、游戏机、嵌入式系统等。 
所属领域的技术人员还将了解，本发明的一个优点是相对于特定计算硬件的独立性增加。因此，将了解，本发明实施例可使用任何计算机系统(包含不提供PPU的系统)来实践。 
2.虚拟编程模型概述
在本发明实施例中，需要使用PPU 122或计算系统的其它处理器来使用线程阵列执行通用计算。如本文所使用，“线程阵列”是由对输入数据集并发执行相同程序以产生输出数据集的若干(n0)线程组成的群组。线程阵列中的每一线程被分派有唯一线程识别符(“线程ID”)，其可由线程在其执行期间存取。可定义为一维或多维数值(例如，0到n0-1)的线程ID控制线程的处理行为的各个方面。举例来说，线程ID可用于确定线程将处理输入数据集的哪一部分，和/或确定线程将产生或写入输出数据集的哪一部分。 
在一些实施例中，线程阵列是“协作的”线程阵列或CTA。与其它类型的线程阵列一样，CTA是对输入数据集并发执行相同程序(本文称为“CTA程序”)以产生输出数据集的多个线程的群组。在CTA中，线程可通过以依赖于线程ID的方式彼此共享数据 而进行协作。举例来说，在CTA中，数据可由一个线程产生并由另一线程消耗。在一些实施例中，可在将共享数据的点处将同步指令插入到CTA程序代码中，以确保在消耗数据的线程试图访问数据之前已由产生数据的线程实际产生数据。CTA的线程之间的数据共享(如果存在的话)的程度由CTA程序确定；因此，将了解，在使用CTA的特定应用中，CTA的线程可能或可能不实际上彼此共享数据，这取决于CTA程序，且本文中同义地使用术语“CTA”和“线程阵列”。 
在一些实施例中，CTA中的线程与同一CTA中的其它线程共享输入数据和/或中间结果。举例来说，CTA程序可能包含用于计算共享存储器中的特定数据将被写入到的地址的指令，其中所述地址是线程ID的函数。每一线程使用其自身的线程ID来计算所述函数，并写入到相应位置。有利地定义地址函数使得不同线程写入到不同位置；只要函数是确定性的，就可预测任何线程写入到的位置。CTA程序还可包含用于计算共享存储器中的将从中读取数据的地址的指令，其中所述地址是线程ID的函数。通过定义适宜的函数和提供同步技术，数据可以可预测方式由CTA的一个线程写入到共享存储器中的给定位置，并由同一CTA的不同线程从所述位置进行读取。因此，可支持线程之间的任何所需的数据共享模式，且CTA中的任何线程可与同一CTA中的任何其它线程共享数据。 
有利地使用CTA(或其它类型的线程阵列)来执行便于数据并行分解(data-paralleldecomposition)的计算。如本文所使用，“数据并行分解”包含其中通过对输入数据并行执行同一算法多次以产生输出数据来对计算问题求解的任何情形；举例来说，数据并行分解的一个普通实例涉及将同一处理算法应用于输入数据集的不同部分以便产生输出数据集的不同部分。适于数据并行分解的问题的实例包含矩阵代数、任何数目的维度的线性和/或非线性变换(例如，快速傅里叶变换)，和包含任何数目的维度的卷积滤波、多个维度的可分离滤波等各种滤波算法。在CTA程序中指定待应用于输入数据集的每一部分的处理算法，且CTA中的每一线程对输入数据集的一个部分执行同一CTA程序。CTA程序可使用广范围的数学和逻辑运算来实施算法，且所述程序可包含有条件或分支执行路径以及直接和/或间接存储器存取。 
上文引用的第11/303,780号申请案中进一步详细描述了CTA及其执行。 
在一些情形中，定义相关CTA(或更一般来说，线程阵列)的“栅格”也是有用的。如本文所使用，CTA的“栅格”是若干(n1)个CTA的集合，其中所有CTA为相同大小(即，线程数目)并执行相同的CTA程序。栅格内的n1个CTA有利地彼此独立，意味着栅格中任何CTA的执行不受栅格中的任何其它CTA的执行的影响。如将了解，此 特征提供在可用的处理核心之间分配CTA的显著灵活性。 
为了区分栅格内的不同CTA，有利地向栅格的每一CTA分派“CTA识别符”(或CTAID)。与线程ID一样，任何唯一识别符(包含但不限于数值识别符)均可用作CTA ID。在一个实施例中，CTA ID只是从0到n1-1的连续(一维)指数值。在其它实施例中，可使用多维指数标定方案。CTA ID对于CTA的所有线程是共同的，且栅格内给定CTA的线程可使用其CTA ID并结合其线程ID来确定(例如)用于读取输入数据的源位置和/或用于写入输出数据的目的地位置。以此方式，同一栅格的不同CTA中的线程可对同一数据集并发操作，但在一些实施例中，不支持栅格中的不同CTA之间的数据共享。 
定义CTA栅格(例如)在需要使用多个CTA来对单一大问题的不同部分求解的情况下可能是有用的。举例来说，可能需要执行滤波算法以产生高分辨率电视(HDTV)图像。如此项技术中已知，HDTV图像可能包含2百万个以上的像素。如果每一线程产生一个像素，那么待执行的线程数目将超过单一CTA中可处理的线程数目(假定为使用常规技术构造的具有合理尺寸和成本的处理平台)。 
可通过在多个CTA之间划分图像使每一CTA产生输出像素的不同部分(例如，16x16小片)来管理这种大处理任务。所有CTA执行相同程序，且线程使用CTAID与线程ID的组合来确定用于读取输入数据和写入输出数据的位置，使得每一CTA对输入数据集的正确部分操作并将其输出数据集部分写入到正确位置。 
应注意，与CTA内的线程(其可共享数据)不同，栅格内的CTA有利地不彼此共享数据或以另外的方式彼此依赖。也就是说，可循序(以任一次序)或并发执行同一栅格的两个CTA，且仍产生相同结果。因此，处理平台(例如，图1的系统100)可通过首先执行一个CTA接着执行下一CTA等等直到已执行栅格的所有CTA为止，来执行CTA栅格并获得结果。或者，如果有足够的资源可用，处理平台可通过并行执行多个CTA来执行相同栅格并获得相同结果。 
在一些实例中，可能需要定义CTA的多个(n2)栅格，其中每一栅格执行数据处理程序或任务的不同部分。举例来说，数据处理任务可能划分为若干个“求解步骤”，其中通过执行CTA栅格来执行每一求解步骤。作为另一实例，数据处理任务可能包含对一连串输入数据集(例如，连续的视频数据帧)执行相同或类似的操作；可针对每一输入数据集执行CTA栅格。虚拟编程模型有利地支持至少这三个等级的工作定义(即，线程、CTA和CTA栅格)；视需要还可支持额外的等级。 
将了解，用于对特定问题求解的CTA的大小(线程的数目n0)、栅格的大小(CTA 的数目n1)和栅格的数目(n2)将取决于问题的参数和定义问题分解的编程人员或自动代理的偏好。因此，在一些实施例中，CTA的大小、栅格的大小和栅格的数目有利地由编程人员定义。 
受益于CTA方法的问题通常特征在于，存在可被并行处理的大量数据元素。在一些实例中，数据元素是输出元素，其每一者是通过对输入数据集的不同(可能重叠)部分执行相同算法而产生的。在其它实例中，数据元素可以是输入元素，其每一者将使用相同算法予以处理。 
此类问题可始终分解为至少两个等级并映射到上文描述的线程、CTA和栅格上。举例来说，每一栅格可能表示复杂数据处理任务中的一个求解步骤的结果。每一栅格有利地划分为若干“区块”，其每一者可作为单一CTA被处理。每一区块有利地含有多个“元素”，即待求解的问题的基本部分(例如，单一输入数据点或单一输出数据点)。在CTA内，每一线程处理一个或一个以上元素。 
图2A和2B说明在本发明实施例中使用的虚拟编程模型中栅格、CTA与线程之间的关系。图2A展示若干栅格200，每一栅格由CTA 202的二维(2-D)阵列组成。(本文中，相似对象的多个实例用表示所述对象的参考标号和(需要时)表示实例的括弧标号来指示。)如图2B中针对202(0，0)所示，每一CTA 202包含线程(Θ)204的2-D阵列。对于每一栅格200的每一CTA 202中的每一线程204，可定义I＝[ig，ic，it]的形式的唯一识别符，其中栅格识别符ig唯一地识别栅格，CTA ID ic唯一地识别栅格内的CTA，且线程ID it 唯一地识别CTA内的线程。在此实施例中，可由一维栅格识别符ig、二维CTA识别符ic 和二维线程识别符it来构成识别符I。在其它实施例中，唯一识别符I是整数的三元组，其中0≤ig＜n2；0≤ic＜n1；且0≤it＜n0。在另外其它实施例中，栅格、CTA和线程识别符中的任一者或全部可能表达为一维整数、2D坐标对、3D三元组等。唯一线程识别符I可用于(例如)确定包含针对整个栅格或多个栅格的输入数据集的阵列内的输入数据的源位置，和/或确定用于在包含针对整个栅格或多个栅格的输出数据集的阵列内存储输出数据的目标位置。 
举例来说，在HDTV图像的情况下，每一线程204可能对应于输出图像的像素。CTA202的大小(线程204的数目)是问题分解中的选择事项，其仅受对单一CTA 202中的最大线程数目的约束的限制(其反映处理器资源的有限性)。栅格200可对应于HDTV数据的整个帧，或者多个栅格可映射到单一帧。 
在一些实施例中，问题分解是统一的，意味着所有栅格200具有相同数目和配置的 CTA 202，且所有CTA 202具有相同数目和配置的线程204。在其它实施例中，分解可能不统一。举例来说，不同栅格可能包含不同数目的CTA，且不同CTA(相同栅格或不同栅格中)可能包含不同数目的线程。 
如上定义的CTA可包含数打或甚至数百个并发线程。上面将执行CTA的并行处理系统可能或可能不支持如此大量的并发线程。在一个方面，本发明通过允许编程人员在不考虑实际硬件能力的情况下使用CTA和CTA栅格的模型来定义处理任务，使编程人员脱离于此类硬件限制。举例来说，编程人员可编写代码(“CTA程序”)，其定义待由CTA的单一代表性线程执行的处理任务；将CTA定义为若干此类线程，每一线程具有唯一识别符；并将栅格定义为若干CTA，每一CTA具有唯一识别符。如下文所描述，此类代码被自动翻译为可在特定平台上执行的代码。举例来说，如果将CTA定义为包含若干(n0)个并发线程但目标平台仅支持一个线程，那么翻译器可定义执行分派到所有n0个线程的任务的一个实际线程。如果目标平台支持一个以上但少于n0个并发线程，那么可视需要在所述数目的可用线程之间划分任务。 
因此，CTA和栅格的编程模型应理解为虚拟模型，即与任何特定物理实现脱离的作为对编程人员的概念上的辅助的模型。可在对并行处理具有不同程度的硬件支持的多种目标平台中实现CTA和栅格的虚拟模型。明确地说，本文所使用的术语“CTA线程”是指离散处理任务(可能与一个或一个以上其它处理任务协作)的虚拟模型，且应了解，CTA线程可能或可能不一对一地映射到目标平台上的线程。 
III.  虚拟结构
根据本发明的一个方面，定义用于执行CTA和CTA栅格的虚拟并行结构。所述虚拟并行结构是支持执行能够在所需时间实现彼此的协作行为(例如，共享数据和同步)的大量并发CTA线程的并行处理器和相关联的存储器空间的表示。此虚拟并行结构可映射到多种实际处理器和/或处理系统上，包含(例如)图1的系统100的PPU 122。虚拟结构有利地定义支持不同等级的数据共享和存取类型的若干虚拟存储器空间，以及识别可由虚拟处理器执行的所有功能的虚拟指令集结构(ISA)。虚拟结构还有利地(例如)通过定义和启动CTA或CTA栅格来定义可用于控制CTA执行的虚拟执行驱动器。 
图3是根据本发明实施例的虚拟结构300的方框图。虚拟结构300包含带有虚拟核心308的虚拟处理器302，所述虚拟核心308经配置以并行执行大量CTA线程。虚拟结构300还包含可由虚拟处理器302存取的全局存储器304，和供应控制虚拟处理器302的操作的命令的虚拟驱动器320。虚拟驱动器320也可存取全局存储器304。 
虚拟处理器302包含接收和解译来自虚拟驱动器320的命令的前端306，以及能够并发执行单一CTA的所有n0个线程的执行核心308。虚拟核心308包含大量(n0个或n0 个以上)虚拟处理引擎310；在一个实施例中，每一虚拟处理引擎310执行一个CTA线程。虚拟处理引擎310并发执行其各自CTA线程，但不一定并行执行。在一个实施例中，虚拟结构300指定虚拟处理引擎310的数目T(例如，384、500、768等)；此数目设定CTA中的线程数目n0的上限。应了解，虚拟结构300的实现可包含少于指定数目T的物理处理引擎，且单一处理引擎可执行若干CTA线程，作为单一“真实”(即，平台支持的)线程或作为多个并发的真实线程。 
虚拟处理器302还包含虚拟指令单元312，其保持向虚拟处理引擎310供应针对其各自CTA线程的指令；所述指令由作为虚拟结构300的一部分的虚拟ISA定义。下文描述用于并行线程计算的虚拟ISA的实例。指令单元312在向虚拟处理引擎310供应指令的过程中管理CTA线程同步和CTA线程行为的其它协作性方面。 
执行核心308提供具有不同等级的可存取性的内部数据存储。特殊寄存器311可由虚拟处理引擎310读取但不可由其写入，且用于存储定义图2的问题分解模型内每一CTA线程的“位置”的参数。在一个实施例中，特殊寄存器311对于每CTA线程(或每虚拟处理引擎310)包含一个存储线程ID的寄存器；每一线程ID寄存器仅可由虚拟处理引擎310中的个别虚拟处理引擎310存取。特殊寄存器311还可包含可由所有CTA线程(或由所有虚拟处理引擎310)读取的额外寄存器，其存储CTA识别符、CTA维度、CTA所属的栅格的维度，和CTA所属的栅格的识别符。在初始化期间响应于经由前端306从虚拟驱动器320接收的命令而写入特殊寄存器311，且特殊寄存器311在CTA执行期间不改变。 
局部虚拟寄存器314由每一CTA线程用作临时空间(scratch space)；分配每一寄存器专用于一个CTA线程(或一个虚拟处理引擎310)，且局部寄存器314的任一者中的数据仅可由其被分配到的CTA线程存取。共享存储器316可由所有CTA线程(单一CTA内)存取；共享存储器316中的任何位置可由同一CTA内的任何CTA线程存取(或可由虚拟核心308内的任何虚拟处理引擎310存取)。参数存储器318存储可由任何CTA线程(或任何虚拟处理引擎310)读取但不可由其写入的运行时间参数(常数)。在一个实施例中，虚拟驱动器320在引导虚拟处理器302开始执行使用这些参数的CTA之前将参数提供到参数存储器318。任何CTA内的任何CTA线程(或虚拟核心308内的任何虚拟处理引擎310)均可通过存储器接口322存取全局存储器304。 
在虚拟结构300中，虚拟处理器302在虚拟驱动器320的控制下作为协处理器操作。虚拟结构规范有利地包含虚拟应用程序接口(API)，其识别由虚拟驱动器320认可的函数调用和每一函数调用预期产生的行为。下文描述用于并行线程计算的虚拟API的实例函数调用。 
虚拟结构300可在多种硬件平台上实现。在一个实施例中，虚拟结构300在图1的系统100中实现，其中PPU 122实施虚拟处理器302，且在CPU 102上执行的PPU驱动器程序实施虚拟驱动器320。全局存储器304可在系统存储器104和/或PP存储器124中实施。 
在一个实施例中，PPU 122包含一个或一个以上处理核心，其使用单指令多数据(SIMD)和多线程技术来支持来自单一指令单元(实施虚拟指令单元312)的大量(例如，384或768个)线程的并发执行。每一核心包含P个(例如，8、16等)并行处理引擎302的阵列，其经配置以从指令单元接收并执行SIMD指令，从而允许并行处理具有多达P个线程的群组。核心为多线程的，其中每一处理引擎能够(例如)通过维持与每一线程相关联的当前状态信息使得处理引擎可从一个线程快速切换到另一线程，而并发执行多达某一数目G(例如，24个)的线程群组。因此，核心并发执行各有P个线程的G个SIMD群组，总计为P*G个并发线程。在此实现过程中，只要P*G≥n0，(虚拟)CTA线程与在真实PPU 122上执行的并发线程之间就可存在一对一对应。 
特殊寄存器311可通过以下操作而实施在PPU 122中：向每一处理核心提供P*G条目寄存器堆，其中每一条目能够存储一线程ID；以及提供一组全局可读取的寄存器来存储CTA ID、栅格ID和CTA及栅格维度。或者，可使用其它存储位置来实施特殊寄存器311。 
局部寄存器314可实施在PPU 122中，作为在物理上或逻辑上划分为P个巷道(lane)的局部寄存器堆，其中每一巷道具有某一数目的条目(其中每一条目可能存储(例如)32位字)。一个巷道被分派到P个处理引擎中的每一者，且不同巷道中的相应条目可用执行相同程序以有助于SIMD执行的不同线程的数据填充。巷道的不同部分可分配到G个并发线程群组中的不同线程群组，使得局部寄存器堆中的给定条目仅可由特定线程存取。在一个实施例中，局部寄存器堆内的某些条目保留用于存储线程识别符，从而实施特殊寄存器311中的一者。 
共享存储器316可实施在PPU 122中作为共享寄存器堆或共享芯片上高速缓冲存储器，其具有允许任何处理引擎对共享存储器中的任何位置进行读取或写入的互连。参数 存储器318可实施在PPU 122中作为实施共享存储器316的同一共享寄存器堆或共享高速缓冲存储器内的指定区域，或作为处理引擎具有只读(read-only)存取权的单独共享寄存器堆或芯片上高速缓冲存储器。在一个实施例中，实施参数存储器的区域还用于存储CTA ID和栅格ID以及CTA和栅格维度，从而实施特殊寄存器311的若干部分。 
在一个实施例中，在图1的CPU 102上执行的PPU驱动器程序通过将命令写入到存储器(例如，系统存储器104)中的推播缓冲器(pushbuffer)(未明确展示)来响应虚拟API函数调用，PPU 122从所述推播缓冲器处读取所述命令。所述命令有利地与状态参数相关联，所述状态参数例如CTA中线程的数目、待使用CTA处理的输入数据集在全局存储器中的位置、待执行的CTA程序在全局存储器中的位置，和输出数据将被写入到的全局存储器中的位置。响应于命令和状态参数，PPU 122将状态参数载入到其核心中的一者中，接着开始启动线程直到已启动CTA参数中指定的数目的线程为止。在一个实施例中，PPU 122包含控制逻辑，其在启动线程时将线程ID依次分派到线程；可将线程ID存储(例如)在局部寄存器堆内或专用于此目的的特殊寄存器中的指定位置。 
在替代实施例中，在单线程处理核心中(例如，在一些CPU中)实现虚拟结构300，所述单线程处理核心使用少于n0个实际线程来执行所有CTA线程；虚拟编程模型使得与不同CTA线程相关联的处理任务可(例如)通过针对一个CTA线程接着针对下一CTA线程等等执行任务(或其一部分)而组合到单一线程中。可利用向量执行、SIMD执行和/或机器中可用的任何其它形式的并行性来并行执行与多个CTA线程相关联的处理任务，或并行执行与同一CTA线程相关联的多个处理任务。因此，可使用单一线程、n0个线程或任何其它数目的线程来实现CTA。如下文所描述，虚拟指令翻译器有利地将编写为目标虚拟结构300的代码翻译为特定针对目标平台的指令。 
将了解，本文描述的虚拟结构是说明性的且可能作出变化和修改。举例来说，在一个替代实施例中，每一虚拟处理引擎可具有存储分派到其线程的唯一线程ID的专门的线程ID寄存器，而不使用局部虚拟寄存器中的空间用于此目的。 
作为另一实例，虚拟结构可指定关于虚拟核心308的内部结构的或多或少的细节。举例来说，可能指定虚拟核心308包含用于执行P向SIMD群组中的CTA线程的P个多线程虚拟处理引擎，其中多达G个SIMD群组共存于核心308中使得P*G确定T(CTA中线程的最大数目)。还可指定不同类型的存储器和共享等级。 
虚拟结构可实现于使用硬件和/或软件元件的任何组合的多种计算机系统中以定义和控制每一组件。虽然已以举例的方式描述使用硬件组件的一个实现方案，但应了解， 本发明涉及使编程任务与特定硬件实现脱离。 
4.对虚拟结构进行编程
图4是根据本发明实施例使用虚拟结构300来操作目标处理器或平台440的概念模型400。如模型400所展示，虚拟结构300的存在使经编译的应用程序和API与目标处理器或平台的硬件实施脱离。 
应用程序402定义数据处理应用，其利用上述虚拟编程模型(包含单一CTA和/或CTA栅格)。一般来说，应用程序402包含多个方面。第一，程序定义单一CTA线程的行为。第二，程序定义CTA的维度(以CTA线程的数目计)，且如果将使用栅格，那么定义栅格的维度(以CTA的数目计)。第三，程序定义待由CTA(或栅格)处理的输入数据集和输出数据集将被存储在的位置。第四，程序定义总体处理行为，包含(例如)何时启动每一CTA或栅格。程序可包含动态确定CTA或栅格的维度、是否不断启动新的CTA或栅格等的额外代码。 
可用例如C/C++、FORTRAN等高级编程语言编写应用程序402。在一个实施例中，C/C++应用程序直接指定一个(虚拟)CTA线程的行为。在另一实施例中，使用数据并行语言(例如，Fortran 90、C*或数据并行C)编写应用程序，且应用程序指定对阵列和聚合数据结构的数据并行操作；此程序可编译成指定一个(虚拟)CTA线程的行为的虚拟ISA程序代码。为了允许定义CTA线程的行为，可提供语言扩充或函数库，编程人员可经由所述语言扩充或函数库来指定并行CTA线程行为。举例来说，可定义特殊符号或变量以对应于线程ID、CTA ID和栅格ID，且可提供函数，编程人员可经由所述函数来指示CTA线程应何时与其它CTA线程同步。 
当编译应用程序402时，编译器408针对应用程序402的定义CTA线程行为的那些部分产生虚拟ISA代码410。在一个实施例中，以图3的虚拟结构300的虚拟ISA表达虚拟ISA代码410。虚拟ISA代码410是程序代码，但不一定是可在特定目标平台上执行的形式的代码。如此，虚拟ISA代码410可作为任何其它程序代码被存储和/或分布。在其它实施例中，可将应用程序完全或部分指定为虚拟ISA代码410，且可完全或部分避免使用编译器408。 
虚拟指令翻译器412将虚拟ISA代码410转换为目标ISA代码414。在一些实施例中，目标ISA代码414是可由目标平台440直接执行的代码。举例来说，如由图4中的虚线框所示，在一个实施例中，目标ISA代码414可由PPU 122中的指令单元430接收和正确解码。视目标平台440的特殊性而定，虚拟ISA代码410可能被翻译成待由目标 平台440上的n0个线程的每一者执行的每线程代码。或者，虚拟ISA代码410可能被翻译成将在少于n0个的线程中执行的程序代码，其中每一线程包含与CTA线程中的一者以上有关的处理任务。 
在一些实施例中，CTA和/或栅格的维度的定义以及定义输入数据集和输出数据集是由虚拟API处理的。应用程序402可包含对于虚拟API函数库404的调用。在一个实施例中，将虚拟API的规范(包含(例如)函数名称、输入、输出和作用，但不包含实施细节)提供给编程人员，且编程人员将虚拟API调用直接并入到应用程序402中，藉此直接产生虚拟API代码406。在另一实施例中，通过编译使用某一其它语法来定义CTA和栅格的应用程序402来产生虚拟API代码406。 
部分地通过提供虚拟执行驱动器416来实现虚拟API代码406，虚拟执行驱动器416将代码406的虚拟API命令翻译为可由目标平台440处理的目标API命令418。举例来说，如由图4中的虚线框所示，在一个实施例中，目标API命令418可由PPU驱动器432接收和处理，PPU驱动器432将相应命令传送到PPU 122前端434。(在此实施例中，虚拟执行驱动器416可以是PPU驱动器432的一方面或部分。)在另一实施例中，虚拟执行驱动器可能不对应于用于协处理器的驱动器；其可能只是在运行所述虚拟执行驱动器的同一处理器上启动其它程序或线程的控制程序。 
应了解，可针对能够支持CTA执行的任何平台或结构创建虚拟指令翻译器412和虚拟执行驱动器416。就针对不同平台或结构的虚拟指令翻译器412可从同一虚拟ISA进行翻译来说，同一虚拟ISA代码410可与任何平台或结构一起使用。因此，不需要针对每一可能的平台或结构重新编译应用程序402。 
此外，目标平台440不必包含如图4所示的PPU和/或PPU驱动器。举例来说，在一个替代实施例中，目标平台是使用软件技术仿真大量线程的并发执行的CPU，且目标ISA代码和目标API命令对应于待由目标CPU执行的程序(或互相通信的程序的群组)中的指令，所述目标CPU可以是(例如)单核心或多核心CPU。 
5.虚拟ISA实例
现在将描述根据本发明实施例的虚拟ISA的实例。如上所述，虚拟ISA有利地对应于上述虚拟编程模型(CTA和栅格)。因此，在本实施例中，由编译器408产生的虚拟ISA代码410定义要由图3的虚拟核心308中的虚拟处理引擎310之一执行的单个CAT线程的行为；所述行为可包含与其它CTA线程的协作性交互，例如同步和/或数据共享。 
应了解，本文中描述的虚拟ISA只用于说明的用途，且本文中描述的特定要素或要 素组合并不限制本发明的范围。在一些实施例中，编程人员可用虚拟ISA编写代码；在其它实施例中，编程人员用另一高级语言(例如，FORTRAN，C，C++)编写代码，且编译器408产生虚拟ISA代码。编程人员也可编写“混合”代码，其中代码中有些部分是用高级语言编写而其它部分是用虚拟ISA编写。 
5.1特殊变量
图5是列举由实例虚拟ISA定义的“特殊”变量的表格500(本文中用前缀“％”来表示特殊变量)。这些变量与图2的编程模型有关，其中通过每个线程204在CTA 202内的位置来识别所述线程，而CTA 202又位于某一数目的栅格200中的特定一者内。在一些实施例中，表格500的特殊变量对应于图3的虚拟结构300中的特殊寄存器311。 
在表格500中，假设CTA和栅格各自用三维空间来定义，且不同的栅格在一维空间中循序编号。虚拟ISA预期图5的特殊变量将在启动CTA时被初始化，且虚拟ISA代码可简单地使用这些变量而无需初始化。以下参看虚拟API论述特殊变量的初始化。 
如图5所示，特殊变量的第一个3向量％ntid＝(％ntid.x，％ntid.y，％ntid.z)定义CTA的维度(以线程数目计)。CTA中的所有线程将共享同一％ntid向量。在虚拟结构300中，预期％ntid向量的值将经由虚拟API函数调用提供给虚拟处理器302，所述虚拟API函数调用如下所述地建立CTA的维度。 
如图5所示，特殊变量的第二个3向量％tid＝(％tid.x，％tid.y，％tid.z)是指CTA内的给定线程的线程ID。在图3的虚拟结构300中，预期虚拟处理器302将在启动CTA的每个线程时分派满足制约条件0≤％tid.x＜％ntid.x、0≤％tid.y＜％ntid.y和0≤％tid.z＜％ntid.z的唯一％tid向量。在一个实施例中，％tid向量可定义成使得其可存储在压缩的32位的字中(例如，对于％tid.x为16个位，对于％tid.y为10个位，且对于％tid.z为6个位)。 
如图5所示，特殊变量的第三个3向量％nctaid＝(％nctaid.x，％nctaid.y，％nctaid.z)定义栅格的维度(以CTA的数目计)。在图3的虚拟结构300中，预期％nctaid向量的值将经由虚拟API函数调用提供给虚拟处理器302，所述虚拟API函数调用建立CTA的栅格的维度。 
如图5所示，特殊变量的第四个3向量％ctaid＝(％ctaid.x，％ctaid.y，％ctaid.z)是指栅格内的给定CTA的CTA ID。在图3的虚拟结构300中，预期当启动CTA时将把满足CTA的限制条件0≤％ctaid.x＜％nctaid.x、0≤％ctaid.y＜％nctaid.y和0≤％ctaid.z＜％nctaid.z的唯一％ctaid向量提供给虚拟处理器302。 
特殊变量还包含提供CTA所属的栅格的栅格识别符的标量％gridid变量。在图3的虚拟结构300中，预期将把％gridid值提供给虚拟处理器302，以便识别包括当前CTA的栅格。例如当正在使用多个栅格来解决较大问题的不同部分时，有利地在虚拟ISA代码中使用％gridid值。 
5.2程序定义的变量和虚拟状态空间
虚拟ISA使得编程人员(或编译器)可定义任意数目的变量以代表正被处理的数据项目。通过类型和指示如何使用变量以及变量在何种程度上共享的“虚拟状态空间”来定义变量。使用目标平台中可用的寄存器或其它存储器结构来实现变量；在许多目标平台中，状态空间可能会影响对用来实现特定变量的存储器结构的选择。 
图6是列举实例虚拟ISA实施例中支持的变量类型的表格600。支持四个类型：未指定类型的位、带符号的整数、无符号整数和浮点。未指定类型的变量就是单个位或具有指定长度的位的群组。可根据常规格式(例如，IEEE 754标准)来定义带符号整数和无符号整数格式以及浮点格式。 
在本实施例中，针对每个类型支持多个宽度，其中使用参数<n>来指定宽度；因此，举例来说，.s16指示16个位的带符号的整数，.f32指示32位的浮点数字等。如表格600所示，有些变量类型限于特定的宽度；举例来说，浮点变量必须至少为16个位；且整数类型必须至少为8个位。预期虚拟ISA的实现支持所有指定宽度；如果处理器的数据路径和/或寄存器比最宽宽度窄，那么如此项技术中已知的可使用多个寄存器和处理器循环来处置较宽类型。 
应了解，本文中使用的数据类型和宽度是说明而不是限制本发明。 
图7是列举实例虚拟ISA中支持的虚拟状态空间的表格。定义了九个状态空间，其对应于不同的共享级别和在图3的虚拟结构300中的可能的存储位置。 
在线程级别上共享前三个状态空间，这意味着每个CTA线程将具有单独的变量实例，且没有任何CTA线程将可存取任何其它CTA线程的实例。有利地使用虚拟寄存器(.reg)状态空间来定义要由每个CTA线程执行的计算的运算数、临时值和/或结果。程序可宣称任何数目的虚拟寄存器。只能通过静态编译时间名称而不是计算出来的地址来寻址虚拟寄存器。这个状态空间对应于图3的虚拟结构300中的局部虚拟寄存器314。 
特殊寄存器(.sreg)状态空间对应于图5的预定义的特殊变量，其存储在虚拟结构300中的特殊寄存器311中。在一些实施例中，虚拟ISA代码可能不会宣称.sreg空间中的任何其它变量，而是可使用特殊变量作为对计算的输入。所有CTA线程均可读取.sreg 状态空间中的任何变量。对于％tid(或其分量)，每个CTA线程将读取其唯一的线程识别符；对于.sreg状态空间中的其它变量，同一CTA中的所有CTA线程将读取相同值。 
每线程局部存储器(.local)变量对应于全局存储器304中在每CTA线程基础上分配和寻址的区域。换句话说，当CTA线程存取.local变量时，其存取其自身的变量实例，且在一个CTA线程中对.local变量作出的改变不会影响其它CTA线程。与.reg和.sreg状态空间不同，每线程局部存储器可使用计算出来的地址来寻址。 
接下来的两个状态空间定义每CTA变量，这意味着每个CTA将具有变量的一个实例，其任何(虚拟)线程均可存取所述实例。任何CTA线程均可读取或写入共享(.shared)变量。在一些实施例中，这个状态空间映射到虚拟结构300(图3)的虚拟共享存储器316。在虚拟结构300的实现中，.shared状态空间可映射到芯片上共享存储器实施方案(例如，共享寄存器堆或共享高速缓冲存储器)上，而在其它实现中，.shared状态空间可映射到芯片外存储器的分配和寻址为任何其它可全局存取存储器的每CTA区域上。 
参数(.param)变量是只读的，且可由CTA中的任何(虚拟)线程读取。这个状态空间映射到虚拟结构300的参数存储器318，且可(例如)在芯片上共享参数存储器或高速缓冲存储器中或者在可全局存取的芯片外存储器的分配和寻址为其它任何可全局存取存储器的区域中实现。预期这些变量将响应于来自虚拟驱动器320的驱动器命令来初始化。 
常量(.const)状态空间用来定义可由栅格中的任何CTA中的任何(虚拟)线程读取(但不可修改)的每栅格常量。在虚拟结构300中，.const状态空间可映射到全局存储器中可由CTA线程进行只读存取的区域。.const状态空间可在芯片上共享参数存储器或高速缓冲存储器中或在可全局存取芯片外存储器的分配和寻址为任何其它可全局存取存储器的每栅格区域中实现。与.param状态空间一样，预期.const状态空间中的变量将响应于来自虚拟驱动器320的驱动器命令而被初始化。 
其余三个状态空间定义“上下文”变量，其可供与应用程序相关联的任何CTA中的任何(虚拟)线程存取。这些状态空间映射到虚拟结构300中的全局存储器304。全局(.global)变量可用于一般用途。在一些实施例中，也可定义用于共享纹理(.tex)和表面(.surf)的特定状态空间。这些可用于(例如)图形相关应用程序的状态空间可用来定义并提供对图形纹理和像素表面数据结构的存取，所述数据结构提供对应于2D(或在一些实施例中为3D)阵列的每个像素的数据值。 
在图4的虚拟ISA代码410中，通过指定状态空间、类型和名称来宣称变量。名称 是占位符，且可由编程人员或编译者来选择。因此，例如： 
.reg.b32 vr1； 
宣称虚拟寄存器状态空间中名为vr1的32个位的未指定类型变量。虚拟ISA代码的后续行可引用vr1(例如)作为操作的来源或目的地。 
实例虚拟ISA还支持虚拟变量的阵列和向量。举例来说： 
.global.f32 resultArray[1000][1000]； 
宣称32位浮点数字的虚拟可全局存取1000×1000阵列。虚拟指令翻译器412可将阵列映射到对应于分派的状态空间的可寻址存储器区域中。 
可使用向量前缀.v<m>来定义一个实施例中的向量，其中m是向量的分量的数目。举例来说： 
.reg.v3.f32 vpos； 
宣称每线程虚拟寄存器状态空间中的32位浮点数字的3分量向量。一旦宣称了向量，便可使用例如vpos.x，vpos.y，vpos.z等后缀来识别其分量。在一个实施例中，允许m＝2、3或4，且可使用例如(.x，.y，.z，.w)、(.0，.1，.2，.3)或(.r，.g，.b，.a)等后缀来识别分量。 
由于变量是虚拟的，所以虚拟ISA代码410可定义或引用任何状态空间(除了.sreg，其中变量是预定义的)中的任何数目的变量。可能针对虚拟ISA代码410中的特定状态空间定义的变量的数目可能会超过特定硬件实施方案中的相应类型的存储量。虚拟指令翻译器412有利地经配置以包含合适的存储管理指令(例如，在寄存器与芯片外存储器之间移动数据)，以便使变量在需要时可用。虚拟指令翻译器412也可能能够检测到不再需要临时变量的情况并允许为其分配的空间供另一变量重新使用；可使用常规的用于分配寄存器的编译器技术。 
此外，虽然实例虚拟ISA定义向量变量类型，但不要求目标平台支持向量变量。虚拟指令翻译器412可将任何向量变量实施为适当数目(例如，2、3或4)个标量的集合。 
5.3.虚拟指令
图8A-8H是列举实例虚拟ISA中定义的虚拟指令的表格。通过指令的效果来定义指令，所述效果例如是使用一个或一个以上运算数计算出特定结果并将所述结果放置在目的地寄存器中、设置寄存器值等。将大多数虚拟指令分类以识别输入和/或输出的格式，且指令执行的各方面可取决于类型。指令的一般格式是： 
name.<type>result，operands； 
其中name是指令名称；.<type>是图6中列举的任一类型的占位符；result是存储结果的变量；且operands是提供为对指令的输入的一个或一个以上变量。在一个实施例中， 虚拟结构300是寄存器到寄存器处理器，且用于除存储器存取之外的操作的result和operands(图8F)被要求是在虚拟寄存器状态空间.reg(或在一些运算数的情况下是特殊寄存器状态空间.sreg)中的变量。 
预期目标平台实现虚拟ISA中的每个指令。指令可实现为产生指定效果的相应的机器指令(本文中称为“硬件支持”)，或实现为在执行时产生指定效果的机器指令序列(本文中称为“软件支持”)。特定目标平台的虚拟指令翻译器412有利地经配置以识别对应于每个虚拟指令的机器指令或机器指令序列。 
以下子章节描述图8A-8H中列举的各种类别的指令。应了解，本文中提供的指令列表是说明性的，且虚拟ISA可包含本文中未明确描述的额外指令，且可不包括本文中描述的一些或所有指令。 
5.3.1.虚拟指令-算术
图8A是列举实例虚拟ISA中定义的算术运算的表格800。在此实施例中，虚拟结构只支持寄存器到寄存器算术，且所有算术运算均操纵一个或一个以上虚拟寄存器运算数(图8A中表示为a、b、c)，以产生写入到虚拟寄存器的结果(d)。因此，算术运算的运算数和目的地始终在虚拟寄存器状态空间.reg中，除了图5的特殊寄存器(在特殊寄存器状态空间.sreg中)可用作运算数以外。 
表格800中的算术运算的列表包含四个基本的算术运算：加法(add)、减法(sub)、乘法(mul)和除法(div)。这些运算可对所有整数和浮点数据类型执行，且产生相同类型的结果作为输入；在一些实施例中，也可将舍入模式的限定符(qualifier)添加到指令中，以允许编程人员指定应当如何对结果进行舍入，以及在整数运算数的情况下是否应当强加饱和极限。 
也支持a、b和c运算数的三个复合算术运算：乘加(mad)、熔合乘加(fma)以及绝对差和(sad)。乘加计算a×b的乘积(具有括号指示的舍入)，且将结果与c相加。熔合乘加与mad的区别在于，在与c相加之前，不对a×b的乘积进行舍入。绝对差和计算绝对值|a-b|，然后与c相加。 
余数(rem)运算只对整数运算数执行，且当将运算数a除以运算数b时计算余数(amod b)。绝对值(abs)和求反(neg)是可应用于浮点格式或带符号的整数格式的运算数a的一元运算。可应用于整数或浮点运算数的最小值(min)和最大值(max)运算将目的地寄存器设置成较小运算数或较大运算数；也可指定对一个或两个运算数是非正规数(例如，根据IEEE 754标准)的特殊情况的处置。 
表格800中的其余运算只对浮点类型执行。分数(frc)运算返回其输入的分数部分。正弦(sin)、余弦(cos)和比率的反正切(atan2)提供对应于三角函数的方便的指令。也支持以2为底的对数(lg2)和取幂(ex2)。也支持倒数(rcp)、平方根(sqrt)和倒数平方根(rsqrt)。 
应注意，这个算术运算列表是说明而不是限制本发明。可支持其它运算或运算组合，其中包含任何预期会以足够的频率被调用的运算。 
在一些实施例中，虚拟ISA还定义向量运算。图8B是列举实例虚拟ISA支持的向量运算的表格810。向量运算包含点积(dot)运算，其计算运算数向量a和b的标量点积d；叉积(cross)运算，其计算运算数向量a和b的向量叉积d；以及量值(mag)运算，其计算运算数向量a的标量长度d。向量归约(vred)运算通过以迭代方式对向量运算数a的要素执行指定操作<op>来计算标量结果d。在一个实施例中，对于浮点向量只支持归约运算add、mul、min和max；对于整数向量，也可支持额外的归约运算(例如，and、or和xor，如下文所述)。 
除了这些运算之外，也可在虚拟ISA中定义例如向量加法、向量定标等其它向量运算(图8B中未列举)。 
如上所述，虚拟结构300的一些硬件实现可能不支持向量处理。此种实现的虚拟指令翻译器412有利地适合于产生标量机器指令的适当序列以执行这些运算；所属领域的技术人员将能够确定适当的序列。 
5.3.2虚拟指令-选择和设置寄存器
图8C是列举实例虚拟ISA中定义的选择和设置寄存器运算的表格820。这些运算可对任何数值数据类型执行，且基于比较运算的结果设置目的地寄存器。如果c是非零那么基本选择(sel)运算选择运算数α，且如果c是零则选择运算数b。比较与设置(set)对运算数a和b执行比较运算<cmp>以产生比较结果t，接着基于比较结果t是真(～0)还是假(0)而将目的地寄存器d设置成布尔(Boolean)真(～0)或假(0)。一个实施例中允许的比较运算<cmp>包含等于(如果a＝b则t为真)、大于(如果a>b则t为真)、小于(如果a<b则t为真)、大于或等于(如果a≥b则t为真)、小于或等于(如果a≤b则t为真)，以及其它比较，其中包含例如a和/或b是数字值还是未定义的值。 
Setb运算是对比较与设置的变化形式，其在比较运算<cmp>的结果t与第三运算数c之间执行进一步的布尔运算<bop>；布尔运算t<bop>c的结果确定是将目的地寄存器d设置成布尔真还是布尔假。一个实施例中允许的布尔运算<bop>包含and、or和xor(见 下述图8C)。setp运算与setb相似，区别只是设置了两个1位“断言”(predicate)目的地寄存器：将目的地寄存器d1设置成t<bop>c的结果，而将目的地寄存器d2设置成(！t)<bop>c的结果。 
5.3.3.虚拟指令-逻辑和位操纵
图8D是列举实例虚拟ISA中定义的逻辑和位操纵运算的表格830。按位的布尔运算and、or和xor通过以下方式来执行：对运算数a和b的每个位执行指定运算，且将寄存器d中的相应位设置成结果。按位求反(not)运算反转运算数a的每个位，而如果a是零(布尔假)则逻辑求反(cnot)运算将目的地寄存器设置成1(布尔真)，否则设置成0(布尔假)。 
通过左移(shl)和右移(shr)运算来支持移位，所述运算将运算数a中的位字段左移或右移运算数b指定的位数。对于带符号的格式，右移有利地基于符号位来填补引导位；对于无符号格式，右移用零来填补引导位。 
5.3.4.虚拟指令-格式转换
图8E是列举实例虚拟ISA中定义的格式转换运算的表格840。格式转换(cvt)指令将第一类型<atpye>的运算数a转换成目标类型中的均等值<dtpye>，且将结果存储在目的地寄存器d中。图6中列举一个实施例中的有效类型；无法将未指定类型值(.b<n>)转换成整数或浮点类型或者将整数或浮点类型转换成未指定类型值。格式转换指令的变化形式使得编程人员可指定舍入模式<mode>；也可指定对当表达为目标类型时饱和的数字的处置。 
5.3.5.虚拟指令-数据移动和数据共享
图8F是列举实例虚拟ISA中定义的数据移动和数据共享指令的表格850。移动(mov)操作将目的地寄存器d设置成立即运算数a的值，或者如果运算数a是寄存器，则设置成寄存器a的内容。移动操作可限于虚拟寄存器类型的状态空间，例如图7中的.reg和.sreg。 
加载(ld)指令将来自存储器中的源位置的值加载到目的地寄存器d中，所述目的地寄存器d在一个实施例中必须在虚拟寄存器(.reg)状态空间中。.<space>限定符指定源位置的状态空间，且可限于图7中的可寻址状态空间，例如.reg和.sreg之外的空间(其中可改为使用移动操作)。由于这个实施例中的虚拟结构300是寄存器到寄存器处理器，所以加载指令有利地用来将变量从可寻址的状态空间转移到虚拟寄存器.reg状态空间中，使其可用作运算数。 
使用可用各种方式定义的来源参数<src>来识别特定的源位置，以便支持不同的寻址模式。举例来说，在一些实施例中，来源参数<src>可能是以下中的任何一者：其值存储在d中的经命名的可寻址变量、对保存来源地址的寄存器的引用、对保存将与偏移值(提供为立即运算数)相加的地址的寄存器的引用，或立即绝对地址。 
类似地，存储(st)操作将来源寄存器a中的值存储到由目的地参数<dst>识别的存储器位置。一个实施例中的来源寄存器a必须在.reg状态空间中；目的地必须在可写入和可寻址的状态空间(例如，图7中的.local、.global或.shared)中。目的地参数<dst>可用各种方式定义以支持不同的寻址模式，这与加载指令中的来源参数<src>相似。例如，可使用存储指令将操作结果从寄存器转移到可寻址的状态空间。 
在提供纹理和表面状态空间的实施例中，可使用额外的虚拟指令从纹理存储器状态空间进行读取(tex)和从表面存储器状态空间进行读取(suld)和对其进行写入(sust)。纹理读取的运算数(t，x，y)指定纹理识别符(t)和坐标(x，y)；同样，表面读取或写入的运算数(s，x，y)指定表面识别符(s)和坐标(x，y)。 
CTA线程可通过与其它CTA线程共享数据而与其它CTA线程协作。举例来说，为了在CTA内共享数据，CTA线程可使用加载与存储虚拟指令(以及下述原子更新指令atom)以将数据写入到每CTA虚拟状态空间和从中读取数据。因此，一个CTA线程可使用具有合适定义的目的地地址的st.shared指令将数据写入到.shared状态空间；相同CTA内的另一CTA线程可随后通过在ld.shared指令中使用相同地址来读取所述数据。下述同步指令(例如，bar和membar)可用来确保CTA线程间的数据共享操作的正确序列，例如，产生数据的CTA线程在消耗数据的CTA线程读取数据之前先写入数据。类似地，st.global和ld.global指令可用于相同CTA中的CTA线程、相同栅格中的CTA和/或相同应用程序中的不同栅格之间的协作和数据共享。 
5.3.6.虚拟指令-程序控制
图8G是列举实例虚拟ISA中提供的程序控制操作的表格860。这些控制操作是所属领域的技术人员所熟悉的，且其使得编程人员可将程序执行重定向。分支(bra)将程序流重定向到目标位置<target>。在一些实施例中，通过将字母标签放置在虚拟ISA代码中的目标指令之前并使用所述标签作为分支指令的目标识别符<target>未定义分支目标。举例来说，在一个实施例中： 
label：add.int32 d， vr1，vr2； 
将add指令识别为具有标签label的分支目标。在代码中的其它位置的指令： 
bra label； 
将执行重定向到带有标签的指令。 
Call与返回(ret)指令支持函数和子例行程序调用；fname识别函数或子例行程序。(在一个实施例中，“子例行程序”就是其返回值被忽略的函数。)可使用func指示(directive)来宣称函数fname，且也提供定义所述函数的虚拟ISA代码。波形括号{}或其它分组符号可用来将定义函数或子例行程序的代码与其它虚拟ISA代码隔开。 
对于函数来说，可指定参数列表<rv>来识别应当将返回值存储在何处。对于函数和子例行程序来说，在自变量列表<args>中指定输入自变量。当执行call时，存储下一指令的地址；当执行ret时，进入到所存储地址的分支。 
Exit指令终止遭遇其的CTA线程。陷阱指令调用处理器定义的或用户定义的陷阱例行程序。断点(brkpt)指令使执行暂停，且可用于(例如)调试用途。无操作(nop)是在执行时没有影响的指令。例如，其可用来控制可在多久以后执行下一操作。 
5.3.7.虚拟指令-并行线程
图8H是列举在根据本发明实施例的实例虚拟ISA中提供的明确并行的虚拟指令的表格870。这些指令支持CTA执行所需要的协作线程行为，例如在CTA线程之间交换数据。 
屏障(bar)指令指示：到达其的CTA线程应当在执行任何进一步的指令之前等待，直到其它所有CTA线程(相同CTA中的)也已经到达同一屏障指令时为止。可在CTA程序中使用任何数目的屏障指令。在一个实施例中，屏障指令不需要任何参数(不论使用多少个屏障)，因为必须在所有CTA线程均到达第n个屏障之后，任何线程才可前进到第(n+1)个屏障，依此类推。 
在其它实施例中，可例如通过指定应当在特定屏障处等待的CTA线程(或特定CTA线程的识别符)的数目来使屏障指令参数化。 
另外其它实施例提供“等待”和“不等待”两种屏障指令。在等待屏障指令下，CTA线程一直等到其它相关CTA线程也已经到达屏障为止；在不等待指令处，CTA线程指示其已到达，但可在其它CTA线程到达之前继续。在给定屏障处，一些CTA线程可能等待而其它CTA线程不等待。 
在一些实施例中，bar虚拟指令可用来同步CTA线程，所述CTA线程正在使用共享存储器状态空间进行协作或共享数据。举例来说，假设一组CTA线程(其可包含CTA的一些或全部线程)各自在每线程变量(例如，.fp32虚拟寄存器变量myData)中产生一些数据，然后读取集合中的另一CTA线程产生的数据。指令序列： 
st.shared.fp32 myWriteAddress，myData； 
bar； 
ld. shared.fp32 myData，myReadAddress； 
提供所需的行为，其中myWriteAddress和myReadAddress是对应于.shared状态空间中的地址的每线程变量。在每个CTA线程将其产生的数据写入到共享存储器之后，其一直等到所有CTA线程已经存储其数据为止，然后继续从共享存储器中读取数据(其可能已经由不同的CTA线程写入)。 
存储器屏障(membar)指令指示每个CTA线程应当等待其先前请求的存储器操作(或至少所有写入操作)完成。这个指令保证在membar指令之后发生的存储器存取将看到在其之前的任何写入操作的结果。Membar指令在一个实施例中使用可选的状态空间名称<space>以将其范围限制于瞄准指定状态空间的存储器操作，所述指定状态空间应当是存储器状态空间(例如，不是.reg或.sreg状态空间)。如果未指定任何状态空间名称，那么CTA线程等待瞄准所有存储器状态空间的所有待决操作完成。 
原子更新(atom)指令致使对由引用<ref>识别的共享变量a的原子更新(读取-修改-写入)。所述共享变量a可处于任何共享状态空间中，且与其它存储器引用一样，可使用各种寻址模式。举例来说，<ref>可为以下中的任一者：命名的可寻址变量a、对保存变量a的地址的寄存器的引用、对保存将添加到偏移值(提供为立即操作数)以定位变量a的地址的寄存器的引用，或变量a的立即绝对地址。CTA线程将变量a从共享状态空间位置加载到目的地寄存器d中，接着使用对运算数a和(取决于操作)第二和第三运算数b、c执行的指定操作<op>来更新变量a，并将结果存储回由<ref>识别的位置。目的地寄存器d保持a的原先加载的值。原子地执行加载、更新和存储操作，保证没有其它任何CTA线程在第一CTA线程执行原子更新时存取变量a。在一个实施例中，变量a限于.global或.shared状态空间，且可用与上述加载和存储操作一样的方式来指定。 
在一些实施例中，只有特定的操作可作为原子更新来执行。举例来说，在一个实施例中，如果a是浮点类型那么只可指定以下操作<op>：将a与b相加；用a和b的最小值或最大值来替换a；以及三元比较与交换操作，其在a等于b的情况下用c替换a，且否则使a保持不变。对于整数a，可支持额外操作，例如运算数a与b之间的按位的and、or和xor，以及使运算数a递增或递减。也可支持其它原子操作或操作组合。 
Vote指令在预定义组的CTA线程间对布尔(例如，.bl类型)运算数a执行归约运算<op>。在一个实施例中，虚拟结构指定CTA线程在SIMD群组中执行，且预定义的群组对应于SIMD群组；在其它实施例中，其它群组的CTA线程可由虚拟结构或编程人员来定义。归约运算<op>要求基于在群组中的CTA线程间对运算数a的归约以及.<op>限定 符指定的归约运算将结果值d设置成布尔真或布尔假。在一个实施例中，所允许的归约运算是：(1).all，其中如果a对于群组中的所有CTA线程为真则d为真，且否则为假；(2).any，其中如果a对于群组中的任何CTA线程为真则d为真；以及(3).uni，其中如果a对于群组中的所有活动的CTA线程具有相同值(真或假)则d为真。 
5.3.8.虚拟指令-断言执行
在一些实施例中，虚拟ISA支持任何指令的断言执行。在断言执行中，将布尔“保护断言(guard predicate)”值与指令相关联，且指令只有当在执行的时候保护断言评估为真时才执行。 
在实例虚拟ISA中，保护断言可为任何1位布尔虚拟寄存器变量(本文中表示为P)。通过将断言保护@P或不断言保护@！P放置在指令的操作码前面来指示断言执行。例如通过将P识别为产生了布尔结果的指令(例如，表820(图8C)中的setp指令)的目的地寄存器来在断言寄存器中建立一个值。在遇到@P或@！P保护断言时，虚拟处理器读取P寄存器。对于@P保护，如果P为真，则执行指令，否则将其跳过；对于@！P保护，如果P为假则执行指令，否则跳过。在遇到断言指令的每个CTA线程的执行时间评估断言P；因此，一些CTA线程可能执行断言指令而其它CTA线程不执行。 
在一些实施例中，可将断言设置为指令执行。举例来说，表800-870(图8A-8H)中的特定虚拟指令可能接受指定断言寄存器作为输出的参数；此种指令基于指令结果的某种性质来更新指定的断言寄存器。举例来说，断言寄存器可用来指示算术运算的结果是不是特殊数字(例如，零、无穷大或IEEE 754浮点运算中的非数字)等等。 
6.虚拟指令翻译器。
如参看图4所述，虚拟指令翻译器412瞄准特定的平台结构。虚拟指令翻译器412可实施为(例如)在例如图1的CPU 102等的处理器上执行的软件程序，所述虚拟指令翻译器412接收虚拟ISA代码410并将其翻译成目标ISA代码414，所述目标ISA代码414可在虚拟指令翻译器412瞄准的特定平台结构上执行(例如，通过图1的PPU 122)。虚拟指令翻译器412将在虚拟ISA代码410中宣称的虚拟变量映射到可用的存储位置上，其中包含处理器寄存器、芯片上存储器、芯片外存储器等。在一些实施例中，虚拟指令翻译器412将每个虚拟状态空间映射到特定类型的存储设备上。举例来说，可将.reg状态空间映射到线程特定的数据寄存器上，将.shared状态空间映射到处理器的可共享存储器上，将.global状态空间映射到分配给应用程序的虚拟存储器区域，等等。其它映射也是可能的。 
将虚拟ISA代码410中的虚拟指令翻译成机器指令。在一个实施例中，虚拟指令翻译器412经配置以将每个虚拟ISA指令映射成相应的机器指令或机器指令序列，这取决于在将执行CTA线程的处理器的指令集中是否存在相应的机器指令。 
虚拟指令翻译器412也将CTA线程映射到目标平台结构中的“物理”线程或处理上。举例来说，如果目标平台结构支持至少n0个并发线程，那么每个CTA线程可映射到一个物理线程上，且虚拟指令翻译器412可为单个CTA线程产生虚拟指令代码，并预期目标平台440将为具有n0个唯一识别符的n0个线程执行代码。如果目标平台结构支持少于n0个线程，那么虚拟指令翻译器412可产生虚拟ISA代码410，所述虚拟ISA代码410并入有对应于多个CTA线程的指令，并预期这个代码将对每个CTA执行一次，因而将多个CTA线程映射到单个物理线程或处理。 
确切地说，将与数据共享(例如，存取.shared或.global状态空间的加载、存储和原子更新指令)和/或协作线程行为(例如，图8H中的屏障、原子更新和其它指令)有关的虚拟指令翻译成机器指令或机器指令序列。针对CTA执行而优化的目标平台结构有利地包含硬件支持的屏障指令，例如其中用指令单元中的计数器和/或寄存器来计数已经到达屏障指令的线程的数目，并设置旗标以防止在线程在屏障处等待时发布线程的进一步的指令。其它目标结构可能不提供对线程同步的直接硬件支持，在此情况下可使用其它线程间通信技术(例如，信号量、存储器中的状态阵列等)来创建所需的行为。 
也将断言指令翻译成机器指令。在一些实例中，目标硬件直接支持断言执行。在其它实例中，可将断言与有条件的分支指令等一起存储在(例如)处理器寄存器中，所述有条件的分支指令等用来询问寄存器并通过有条件地围绕断言指令产生分支而创建所需的运行时间行为。 
图9是使用根据本发明实施例的虚拟指令翻译器的过程900的流程图。在步骤902处，编程人员用高级语言编写CTA程序代码。在一个实施例中，所述CTA程序代码定义单个CTA线程的所需行为，且可使用线程ID(包含CTA ID和/或栅格ID)作为参数来定义或控制CTA线程的行为的各个方面。举例来说，可将要读取或写入的共享存储器位置确定为线程ID的函数，使得相同CTA中的不同CTA线程将对共享存储器中的不同存储器位置进行读取和/或写入。在一个实施例中，包含CTA程序代码作为应用程序代码(例如，图4的程序代码402)的一部分。除了定义CTA线程行为之外，应用程序代码还可定义CTA和/或栅格、设置输入和输出数据集等。 
在步骤904处，编译器(例如，图4的编译器408)根据高级语言代码产生定义单 个(虚拟)CTA线程的行为的虚拟ISA代码。如果代码包含CTA程序代码和其它代码两者，那么编译器408可将CTA程序代码与其余代码分开，以便只使用CTA程序代码来产生虚拟ISA代码。可使用常规的用于将用一种语言编写的程序代码编译成另一(虚拟)语言的技术。应注意，由于所产生的代码是使用虚拟语言，所以编译器不需要束缚于特定硬件或针对特定硬件而优化。编译器可优化根据特定的输入代码序列而产生的虚拟ISA代码(例如，更偏好虚拟ISA指令的较短序列)。可将虚拟ISA中的程序代码存储在盘上的存储器中和/或分布到各种各样的平台结构，其中包含在物理上不同于图3的虚拟结构300的结构。虚拟ISA中的代码是独立于机器的，并且可在任何有虚拟指令翻译器可用的目标平台上执行。在替代实施例中，编程人员可用虚拟ISA直接编写CTA程序代码，或者虚拟ISA代码可由程序自动产生；如果程序代码起初创建为虚拟ISA代码，那么可省略编译步骤904。 
在步骤906处，虚拟指令翻译器(例如，图4的翻译器412)读取虚拟ISA代码，并以目标ISA产生可在目标平台上执行的代码。与编译器不同的是，虚拟指令翻译器瞄准特定的(真实)平台结构且有利地经配置以调适并优化目标ISA代码以便在所述结构上实现最好的性能。在目标结构支持至少n0个线程的一个实施例中，虚拟指令翻译器产生目标线程程序，所述目标线程程序可由n0个线程中的每一者并发执行以实现CTA。在另一实施例中，虚拟指令翻译器产生目标程序，所述目标程序使用软件技术(例如，指令序列)来仿真n0个并发线程，所述线程各自执行对应于虚拟ISA代码的指令。翻译器可在程序安装时、在程序初始化期间或在程序执行期间在及时的基础上操作。 
在步骤908处，目标平台中的处理器(例如，图1的PPU 122)执行目标ISA代码以处理数据。在一些实施例中，步骤908可包含将命令和状态参数提供给处理器，以便控制其行为，如下文进一步描述的。 
将明白，过程900是说明性的，且可进行更改和修改。描述为循序的步骤可并行执行，步骤次序可更改，且步骤可被修改或组合。举例来说，在一些实施例中，编程人员可直接使用虚拟ISA来编写CTA程序代码，从而无需产生虚拟ISA代码的编译器。在其它实施例中，将CTA程序代码编写为较大的应用程序的一部分，且所述CTA程序代码还包含(例如)定义要执行以解决特定问题的CTA和/或CTA栅格的维度的代码。在一个实施例中，只将代码中那些阐述CTA程序的部分编译到虚拟ISA代码中；可将其它部分编译到其它(真实或虚拟)指令集中。 
在其它实施例中，一个虚拟指令翻译器可经配置以产生适合于不同目标平台的目标 代码的多个版本。举例来说，翻译器可产生高级语言(例如，C)的程序代码、用于PPU的机器代码和/或用于使用软件技术来仿真PPU行为的单核或多核CPU的机器代码。 
7.虚拟执行驱动器
在一些实施例中，使用虚拟ISA代码410和虚拟指令翻译器412来产生要对CTA的每个线程执行的CTA程序代码。就图2A-2B的编程模型来说，指定CTA程序会为每个CTA线程204定义处理任务。为了完成所述模型，也有必要定义CTA 202的维度、栅格中的CTA的数目、要处理的输入数据集等。此种信息在本文中称为“CTA控制信息”。 
如图4所示，在一些实施例中，应用程序402通过使用对虚拟库404中的函数的调用来指定CTA控制信息。在一个实施例中，虚拟库404包含各种函数调用，编程人员可经由所述函数调用来定义CTA或CTA栅格并指示何时应当开始执行。 
图10是列举实例虚拟库404中可用的函数的表格1000。第一群组的函数涉及定义CTA。具体来说，initCTA函数是第一个要调用以创建新CTA的函数。这个函数使得编程人员可定义CTA的维度(ntid.x，ntid.y，ntid.z)，并向新CTA分派识别符cname。SetCTAProgram函数指定要由CTA cname的每个线程执行的CTA程序；参数pname是对应于所需的CTA程序(例如，用虚拟ISA代码编写的程序)的逻辑程序识别符。SetCTAInputArray函数使得编程人员可指定全局存储器中CTA cname将从中读取输入数据的源位置(开始地址和大小)，且setCTAOutputArray函数使得编程人员可指定全局存储器中CTA cname将向其写入输出数据的目标位置(开始地址和大小)。使用setCTAParams函数来设置CTA cname的运行时间常量参数。编程人员向函数提供参数列表，例如作为(名称，值)对。 
在一个实施例中，setCTAParams函数也可由编译器408在产生虚拟ISA代码410时使用。由于setCTAParams函数定义CTA的运行时间参数，所以编译器408可将这个函数解译为将每个参数定义为.param状态空间中的虚拟变量。 
表格1000还列举与定义CTA的栅格有关的函数。InitGrid函数是第一个调用来创建新栅格的函数。这个函数使得编程人员可定义栅格的维度(nctaid.x，nctaid.y，nctaid.z)，识别将在栅格上执行的CTA cname，并将识别符gname分派给新定义的栅格。setGridInputArray和setGridOutputArray函数与CTA级函数相似，允许针对栅格中的所有CTA的所有线程定义单个输入和/或输出阵列。setGridParams函数用来为栅格gname中的所有CTA设置运行时间常量参数。编译器408可将这个函数解译为将每个参数定义为.const状态空间中的虚拟变量。 
launchCTA和launchGrid函数指示应开始指定CTA cname或栅格gname的执行。 
虚拟API也可包含其它函数。举例来说，一些实施例提供可用来协调多个CTA的执行的同步函数。举例来说，如果第一CTA(或栅格)的输出将用作第二CTA(或栅格)的输入，那么API可包含一个函数(或启动函数的参数)，可经由所述函数来指示虚拟执行驱动器：在第一CTA(或栅格)的执行完成之前不应启动第二CTA(或栅格)。 
根据本发明的实施例，表格1000中的任何或所有函数调用均可包含在应用程序中，所述应用程序也定义要执行的CTA程序(或者如果在应用程序中存在多个CTA则要定义多个CTA程序)。在编译时，将函数调用视为对应用程序接口(API)库404的调用，因而产生虚拟API代码406。 
使用实施虚拟库中的每个函数的虚拟执行驱动器418来实现虚拟API代码。在一个实施例中，虚拟执行驱动器418是在图1的CPU 102上执行的驱动器程序，所述驱动器程序控制实现CTA线程的PPU 122。实施图10的表格1000中的各种函数调用，使得其导致驱动器经由推送缓冲器(pushbuffer)向PPU 122提供命令。在另一实施例中，CPU执行一个或一个以上程序以实现CTA，且虚拟执行驱动器418设置参数并控制CPU对此种程序的执行。 
将了解，本文中描述的虚拟API是说明性的，且变更和修改是可能的。可支持其它函数或函数组合。此项技术中已知的虚拟API技术可适合于本发明的用途。 
进一步的实施例
虽然已经参照特定实施例描述了本发明，但所属领域的技术人员将认识到许多修改是可能的。举例来说，不需要本文中描述的特定虚拟结构、虚拟指令和虚拟API函数；可用其它支持并发、协作线程的虚拟结构、指令和/或函数来代替。此外，上述实施例可引用所有区块具有相同数目的元件、所有CTA具有相同数目的线程且执行相同CTA程序等的情况。在一些应用程序(例如，其中使用多个依赖性栅格的情况)中，可能需要使不同栅格中的CTA执行不同CTA程序或具有不同数目和/或大小的栅格。 
虽然本文中引用“协作线程阵列”，但应了解，一些实施例可使用其中不支持并发线程之间的数据共享的线程阵列；在支持此种数据共享的其它实施例中，针对给定应用程序定义的线程可能或可能不在实际上共享数据。 
此外，虽然上述实施例可将线程阵列称为具有多个线程，但应了解，在“退化”的情况下，线程阵列可能只具有一个线程。因此，本发明可应用于在要在具有一个或一个以上单线程或多线程核心的CPU上执行的程序中提供可扩展性。通过使用本文中描述的 技术，可用可在任何数目的可用CPU核心上分布线程(例如，使用操作系统功能性)而无需对虚拟ISA代码的修改或重新编译的方式编写程序。 
本文中使用术语“虚拟”和“真实”来反映编程人员用来描述问题解决方案的概念性编程模型与可在上面最终执行程序的实际计算机系统的脱离。“虚拟”编程模型及其相关联的结构使得编程人员可对并行的处理任务采用高级视角，且应了解，可能存在或可能不存在其组件一一对应地映射到本文中描述的虚拟结构组件的实际计算系统或装置。有利地将包含虚拟ISA代码和虚拟API代码的虚拟代码实现为用一种语言编写的可能或可能不与任何实际处理装置的指令集一一对应的代码。与所有程序代码一样，本文中所指的虚拟代码可存储在有形的媒体(例如，存储器或盘)中、通过网络传输等。 
并入有本发明的各种特征的计算机程序——包含但不限于虚拟ISA和/或虚拟API代码、虚拟指令翻译器、虚拟驱动器、编译器、虚拟函数库等——可编码在各种计算机可读媒体上以供存储和/或传输；合适的媒体包含磁盘或磁带、光学存储媒体，例如光盘(CD)或DVD(多功能数字光盘)、闪速存储器等。此种程序也可使用适合经由符合各种协议的有线、光学和/或无线网络(包含因特网)传输的载波信号。用程序代码编码的计算机可读存储媒体可用兼容装置封装，或者程序代码可与其它装置独立地提供(例如，经由因特网下载)。 
此外，本文中可将特定动作描述为由“编程人员”作出。预期编程人员可以是人类、通过很少或没有人为介入来产生程序代码的自动过程、或人类交互与自动或半自动处理的用以产生程序代码的任何组合。 
此外，虽然本文中描述的实施例可能引用了特定目标平台的特征，但本发明不限于这些平台。实际上，虚拟结构可用硬件和/或软件组件的任何组合来实现。所属领域的技术人员将明白，可预期相同虚拟结构的不同实现在效率和/或输出量方面不同；然而，此种差异与本发明无关。 
因此，虽然已经参考特定实施例描述了本发明，但应明白，本发明并不意图涵盖属于随附权利要求书范围内的所有修改和等效物。 

Claims (21)

1.一种定义并行处理操作的方法，所述方法包括：
提供第一程序代码，所述第一程序代码定义要针对协作的虚拟线程的阵列中的多个虚拟线程中的每一个执行的操作序列；
将所述第一程序代码编译成虚拟线程程序，所述虚拟线程程序定义要针对所述多个虚拟线程中的代表性虚拟线程执行的每个线程指令序列，所述每个线程指令序列包含至少一个指令，所述指令定义所述代表性虚拟线程与所述多个虚拟线程中的一个或一个以上其它虚拟线程之间的协作行为，其中所述代表性虚拟线程与所述多个虚拟线程中的每一个一样执行相同的处理任务；以及
存储所述虚拟线程程序。
2.根据权利要求1所述的方法，其进一步包括：
将所述存储的虚拟线程程序翻译成符合目标平台结构的指令序列。
3.根据权利要求1所述的方法，其进一步包括：
提供第二程序代码，所述第二程序代码定义协作的虚拟线程的阵列，所述阵列适合处理输入数据集以产生输出数据集，其中所述阵列中的每个虚拟线程并发地执行所述虚拟线程程序；
将所述第二程序代码转换成虚拟函数库中的函数调用序列，所述库包含初始化并致使执行协作的虚拟线程的阵列的虚拟函数；以及
存储所述函数调用序列。
4.根据权利要求3所述的方法，其进一步包括：
将所述存储的虚拟线程程序和所述函数调用序列翻译成可在目标平台结构上执行的程序代码，所述可执行的程序代码定义执行所述协作的虚拟线程的阵列的一个或一个以上平台线程。
5.根据权利要求4所述的方法，其进一步包括：
在符合所述目标平台结构的计算机系统上执行所述可执行的程序代码，因而产生所述输出数据集；以及
将所述输出数据集存储在存储媒体中。
6.根据权利要求1所述的方法，其中所述每个线程指令序列包含用以在所述每个线程指令序列中的特定点处暂停对所述代表性虚拟线程的操作的执行，直到所述其它虚拟线程中的一个或一个以上到达所述特定点时为止的指令。
7.根据权利要求1所述的方法，其中所述每个线程指令序列包含用以使所述代表性虚拟线程将数据存储在所述其它虚拟线程中的一个或一个以上可存取的共享存储器中的指令。
8.根据权利要求1所述的方法，其中所述每个线程指令序列包含用以使所述代表性虚拟线程自动读取和更新存储在所述其它虚拟线程中的一个或一个以上可存取的共享存储器中的数据的指令。
9.根据权利要求1所述的方法，其中所述虚拟线程程序包含变量定义语句，所述变量定义语句定义多个虚拟状态空间之一中的变量，其中所述多个虚拟状态空间中的不同状态空间对应于所述虚拟线程之间的数据共享的不同模式。
10.根据权利要求9所述的方法，其中所述数据共享的模式包含每个线程不共享模式和全局共享模式。
11.根据权利要求9所述的方法，其中所述数据共享的模式包含每个线程不共享模式、虚拟线程的一个阵列内的共享模式以及全局共享模式。
12.根据权利要求9所述的方法，其中所述数据共享的模式包含每个线程不共享模式、虚拟线程的一个阵列内的共享模式、虚拟线程的多个阵列之间的共享模式以及全局共享模式。
13.一种产生执行于目标平台结构上的程序代码的方法，所述方法包括：
提供输入程序代码，其包含第一部分，所述第一部分定义要对虚拟线程的阵列中的多个虚拟线程中的每一者执行的适于处理输入数据集以产生输出数据集的操作序列，
所述输入程序代码进一步包含第二部分，所述第二部分定义所述虚拟线程的阵列的维度；
将所述输入程序代码的所述第一部分编译成虚拟线程程序，所述虚拟线程程序定义要针对所述多个虚拟线程中的代表性虚拟线程执行的每个线程指令序列，所述每个线程指令序列包含至少一个指令，所述指令定义所述代表性虚拟线程与所述多个虚拟线程中的一个或一个以上其它虚拟线程之间的协作行为，其中所述代表性虚拟线程与所述多个虚拟线程中的每一个一样执行相同的处理任务；
将所述输入程序代码的所述第二部分转换成对虚拟函数库的函数调用序列，所述库包含虚拟函数，所述虚拟函数初始化和致使执行所述协作的虚拟线程的阵列；
将所述虚拟线程程序和所述函数调用序列翻译成可在所述目标平台结构上执行的程序代码，所述可执行的程序代码定义执行所述协作的虚拟线程的阵列的一个或一个以上真实线程；
在符合所述目标平台结构的计算机系统上执行所述可执行的程序代码，因而产生所述输出数据集；以及
将所述输出数据集存储在存储媒体中。
14.根据权利要求13所述的方法，其中所述输入程序代码的所述第二部分包含定义所述虚拟线程的阵列的两个或两个以上维度的程序代码。
15.根据权利要求14所述的方法，其中所述输入程序代码的所述第二部分进一步包含：函数调用，其定义虚拟线程阵列的栅格的一个或一个以上维度，其中所述栅格中的每个阵列均将被执行。
16.根据权利要求13所述的方法，其中所述目标平台结构包含主处理器和协处理器，且其中所述翻译的动作包含：
将所述虚拟线程程序翻译成可由所述协处理器上定义的多个线程并行执行的程序代码；以及
将所述函数调用序列翻译成对所述协处理器的驱动器程序的调用序列，其中所述驱动器程序在所述主处理器上执行。
17.根据权利要求13所述的方法，其中所述目标平台结构包含中央处理单元(CPU)，且其中所述翻译的动作包含：
将所述虚拟线程程序和所述函数调用序列的至少一部分翻译成目标程序代码，所述目标程序代码使用少于所述虚拟线程数目的数目的CPU线程执行所述虚拟线程阵列。
18.一种产生执行于目标平台结构上的程序代码的方法，所述方法包括：
获得虚拟线程程序，所述虚拟线程程序定义要对虚拟线程阵列中的多个虚拟线程中的代表性虚拟线程执行的每个线程指令序列，所述阵列适合于处理输入数据集以产生输出数据集，其中所述代表性虚拟线程与所述多个虚拟线程中的每一个一样执行相同的处理任务；
所述每个线程指令序列包含至少一个指令，所述指令定义所述代表性虚拟线程与所述多个虚拟线程中的一个或一个以上其它虚拟线程之间的协作行为；
获得定义所述虚拟线程阵列的维度的额外程序代码；
将所述虚拟线程程序和所述额外程序代码翻译成可在所述目标平台结构上执行的程序代码，所述可执行的程序代码定义执行所述虚拟线程阵列的一个或一个以上平台线程；
在符合所述目标平台结构的计算机系统上执行所述可执行的程序代码，因而产生所述输出数据集并将所述输出数据集存储在存储器中。
19.根据权利要求18所述的方法，其中所述获得所述虚拟线程程序的动作包含：
接收用高级编程语言编写的CTA源程序代码；以及
编译所述CTA源程序代码以产生所述虚拟线程程序。
20.根据权利要求18所述的方法，其中所述获得所述虚拟线程程序的动作包含：
从存储媒体中读取所述虚拟线程程序。
21.根据权利要求18所述的方法，其中所述获得所述虚拟线程程序的动作包含：
经由网络从远程计算机系统接收所述虚拟线程程序。
CN2008100070263A 2007-01-26 2008-01-25 定义用于目标平台结构的并行处理操作方法 Active CN101231585B (zh)

Applications Claiming Priority (2)

Application Number Priority Date Filing Date Title
US11/627,892 2007-01-26
US11/627,892 US8321849B2 (en) 2007-01-26 2007-01-26 Virtual architecture and instruction set for parallel thread computing

Publications (2)

Publication Number Publication Date
CN101231585A CN101231585A (zh) 2008-07-30
CN101231585B true CN101231585B (zh) 2010-12-01

Family

ID=39587517

Family Applications (1)

Application Number Title Priority Date Filing Date
CN2008100070263A Active CN101231585B (zh) 2007-01-26 2008-01-25 定义用于目标平台结构的并行处理操作方法

Country Status (7)

Country Link
US (1) US8321849B2 (zh)
JP (1) JP4999183B2 (zh)
KR (1) KR101026689B1 (zh)
CN (1) CN101231585B (zh)
DE (2) DE202008017916U1 (zh)
SG (1) SG144869A1 (zh)
TW (1) TWI363294B (zh)

Families Citing this family (142)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
WO2005072307A2 (en) * 2004-01-22 2005-08-11 University Of Washington Wavescalar architecture having a wave order memory
US7490218B2 (en) * 2004-01-22 2009-02-10 University Of Washington Building a wavecache
RU2312388C2 (ru) * 2005-09-22 2007-12-10 Андрей Игоревич Ефимов Способ организации многопроцессорной эвм
US7861060B1 (en) * 2005-12-15 2010-12-28 Nvidia Corporation Parallel data processing systems and methods using cooperative thread arrays and thread identifier values to determine processing behavior
EP2011018B1 (en) 2006-04-12 2016-07-13 Soft Machines, Inc. Apparatus and method for processing an instruction matrix specifying parallel and dependent operations
US8082289B2 (en) 2006-06-13 2011-12-20 Advanced Cluster Systems, Inc. Cluster computing support for application programs
US7836116B1 (en) * 2006-06-15 2010-11-16 Nvidia Corporation Fast fourier transforms and related transforms using cooperative thread arrays
US7640284B1 (en) 2006-06-15 2009-12-29 Nvidia Corporation Bit reversal methods for a parallel processor
EP2122461A4 (en) 2006-11-14 2010-03-24 Soft Machines Inc DEVICE AND METHOD FOR PROCESSING COMMUNICATIONS IN A MULTITHREAD ARCHITECTURE WITH CONTEXT CHANGES
US8549500B2 (en) * 2007-02-14 2013-10-01 The Mathworks, Inc. Saving and loading graphical processing unit (GPU) arrays providing high computational capabilities in a computing environment
US8533697B2 (en) * 2007-02-14 2013-09-10 The Mathworks, Inc. Graphical processing unit (GPU) arrays providing high computational capabilities in a computing environment
US8286196B2 (en) 2007-05-03 2012-10-09 Apple Inc. Parallel runtime execution on multiple processors
US8276164B2 (en) 2007-05-03 2012-09-25 Apple Inc. Data parallel computing on multiple processors
EP3413198A1 (en) 2007-04-11 2018-12-12 Apple Inc. Data parallel computing on multiple processors
US11836506B2 (en) 2007-04-11 2023-12-05 Apple Inc. Parallel runtime execution on multiple processors
US8341611B2 (en) 2007-04-11 2012-12-25 Apple Inc. Application interface on multiple processors
US7725518B1 (en) 2007-08-08 2010-05-25 Nvidia Corporation Work-efficient parallel prefix sum algorithm for graphics processing units
US7877573B1 (en) * 2007-08-08 2011-01-25 Nvidia Corporation Work-efficient parallel prefix sum algorithm for graphics processing units
US8225295B2 (en) * 2007-09-21 2012-07-17 Jens Palsberg Register allocation by puzzle solving
US20090204173A1 (en) 2007-11-05 2009-08-13 Zi-Ping Fang Multi-Frequency Neural Treatments and Associated Systems and Methods
US8200947B1 (en) * 2008-03-24 2012-06-12 Nvidia Corporation Systems and methods for voting among parallel threads
US8286198B2 (en) * 2008-06-06 2012-10-09 Apple Inc. Application programming interfaces for data parallel computing on multiple processors
US8225325B2 (en) 2008-06-06 2012-07-17 Apple Inc. Multi-dimensional thread grouping for multiple processors
US8434093B2 (en) 2008-08-07 2013-04-30 Code Systems Corporation Method and system for virtualization of software applications
US8776038B2 (en) 2008-08-07 2014-07-08 Code Systems Corporation Method and system for configuration of virtualized software applications
US8436862B2 (en) * 2008-12-11 2013-05-07 Nvidia Corporation Method and system for enabling managed code-based application program to access graphics processing unit
US8570333B2 (en) * 2008-12-11 2013-10-29 Nvidia Corporation Method and system for enabling managed code-based application program to access graphics processing unit
US8307350B2 (en) * 2009-01-14 2012-11-06 Microsoft Corporation Multi level virtual function tables
JP5149840B2 (ja) * 2009-03-03 2013-02-20 株式会社日立製作所 ストリームデータ処理方法、ストリームデータ処理プログラム、および、ストリームデータ処理装置
KR101572879B1 (ko) 2009-04-29 2015-12-01 삼성전자주식회사 병렬 응용 프로그램을 동적으로 병렬처리 하는 시스템 및 방법
US8542247B1 (en) 2009-07-17 2013-09-24 Nvidia Corporation Cull before vertex attribute fetch and vertex lighting
US8564616B1 (en) 2009-07-17 2013-10-22 Nvidia Corporation Cull before vertex attribute fetch and vertex lighting
CN102023844B (zh) * 2009-09-18 2014-04-09 深圳中微电科技有限公司 并行处理器及其线程处理方法
US8266383B1 (en) 2009-09-28 2012-09-11 Nvidia Corporation Cache miss processing using a defer/replay mechanism
US10360039B2 (en) * 2009-09-28 2019-07-23 Nvidia Corporation Predicted instruction execution in parallel processors with reduced per-thread state information including choosing a minimum or maximum of two operands based on a predicate value
US9665920B1 (en) * 2009-10-05 2017-05-30 Nvidia Corporation Simultaneous execution of compute and graphics applications
US8976195B1 (en) 2009-10-14 2015-03-10 Nvidia Corporation Generating clip state for a batch of vertices
US8384736B1 (en) 2009-10-14 2013-02-26 Nvidia Corporation Generating clip state for a batch of vertices
EP2519876A1 (en) * 2009-12-28 2012-11-07 Hyperion Core, Inc. Optimisation of loops and data flow sections
KR101613971B1 (ko) * 2009-12-30 2016-04-21 삼성전자주식회사 프로그램 코드의 변환 방법
US8954958B2 (en) 2010-01-11 2015-02-10 Code Systems Corporation Method of configuring a virtual application
US9104517B2 (en) 2010-01-27 2015-08-11 Code Systems Corporation System for downloading and executing a virtual application
US8959183B2 (en) 2010-01-27 2015-02-17 Code Systems Corporation System for downloading and executing a virtual application
US9229748B2 (en) 2010-01-29 2016-01-05 Code Systems Corporation Method and system for improving startup performance and interoperability of a virtual application
US8763009B2 (en) 2010-04-17 2014-06-24 Code Systems Corporation Method of hosting a first application in a second application
US8375373B2 (en) * 2010-04-19 2013-02-12 Microsoft Corporation Intermediate language support for change resilience
US8527866B2 (en) 2010-04-30 2013-09-03 Microsoft Corporation Multi-threaded sort of data items in spreadsheet tables
US20110276868A1 (en) * 2010-05-05 2011-11-10 Microsoft Corporation Multi-Threaded Adjustment of Column Widths or Row Heights
US9851969B2 (en) 2010-06-24 2017-12-26 International Business Machines Corporation Function virtualization facility for function query of a processor
US10521231B2 (en) * 2010-06-24 2019-12-31 International Business Machines Corporation Function virtualization facility for blocking instruction function of a multi-function instruction of a virtual processor
US8782106B2 (en) 2010-07-02 2014-07-15 Code Systems Corporation Method and system for managing execution of virtual applications
CN103250131B (zh) 2010-09-17 2015-12-16 索夫特机械公司 包括用于早期远分支预测的影子缓存的单周期多分支预测
US9529574B2 (en) * 2010-09-23 2016-12-27 Apple Inc. Auto multi-threading in macroscalar compilers
KR20120031756A (ko) * 2010-09-27 2012-04-04 삼성전자주식회사 Ｃｐｕ와 ｇｐｕ를 사용하는 이종 시스템에서 가상화를 이용한 어플리케이션 컴파일 및 실행 방법 및 장치
KR101649925B1 (ko) * 2010-10-13 2016-08-31 삼성전자주식회사 멀티 트레드 프로그램에서 변수의 단독 메모리 접근여부를 분석하는 방법
US20120096292A1 (en) * 2010-10-15 2012-04-19 Mosaid Technologies Incorporated Method, system and apparatus for multi-level processing
US9021015B2 (en) 2010-10-18 2015-04-28 Code Systems Corporation Method and system for publishing virtual applications to a web server
US9209976B2 (en) 2010-10-29 2015-12-08 Code Systems Corporation Method and system for restricting execution of virtual applications to a managed process environment
JP5490253B2 (ja) * 2010-11-02 2014-05-14 インターナショナル・ビジネス・マシーンズ・コーポレーション 数値集約計算における文字列集約方法
US8819700B2 (en) * 2010-12-22 2014-08-26 Lsi Corporation System and method for synchronous inter-thread communication
KR101157596B1 (ko) * 2010-12-24 2012-06-19 서울대학교산학협력단 개방형 범용 병렬 컴퓨팅 프레임워크(ＯｐｅｎＣＬ)에서의 메모리 접근영역 분석장치 및 그 방법
CN108108188B (zh) 2011-03-25 2022-06-28 英特尔公司 用于通过使用由可分区引擎实例化的虚拟核来支持代码块执行的存储器片段
US9842005B2 (en) 2011-03-25 2017-12-12 Intel Corporation Register file segments for supporting code block execution by using virtual cores instantiated by partitionable engines
KR101638225B1 (ko) 2011-03-25 2016-07-08 소프트 머신즈, 인크. 분할가능한 엔진에 의해 인스턴스화된 가상 코어를 이용한 명령어 시퀀스 코드 블록의 실행
WO2012157786A1 (ja) 2011-05-19 2012-11-22 日本電気株式会社 並列処理装置、並列処理方法、最適化装置、最適化方法、および、コンピュータ・プログラム
KR101639853B1 (ko) * 2011-05-20 2016-07-14 소프트 머신즈, 인크. 복수의 엔진에 의해 명령어 시퀀스들의 실행을 지원하기 위한 자원들 및 상호접속 구조들의 비집중 할당
CN103649931B (zh) 2011-05-20 2016-10-12 索夫特机械公司 用于支持由多个引擎执行指令序列的互连结构
US8990830B2 (en) * 2011-07-19 2015-03-24 International Business Machines Corporation Thread management in parallel processes
KR101894752B1 (ko) 2011-10-27 2018-09-05 삼성전자주식회사 가상 아키텍쳐 생성 장치, 런타임 시스템, 멀티 코어 시스템 및 그 동작 방법
US9009686B2 (en) * 2011-11-07 2015-04-14 Nvidia Corporation Algorithm for 64-bit address mode optimization
US9507638B2 (en) * 2011-11-08 2016-11-29 Nvidia Corporation Compute work distribution reference counters
US10191746B2 (en) 2011-11-22 2019-01-29 Intel Corporation Accelerated code optimizer for a multiengine microprocessor
WO2013077876A1 (en) 2011-11-22 2013-05-30 Soft Machines, Inc. A microprocessor accelerated code optimizer
US10255228B2 (en) * 2011-12-06 2019-04-09 Nvidia Corporation System and method for performing shaped memory access operations
US9507593B2 (en) * 2011-12-23 2016-11-29 Intel Corporation Instruction for element offset calculation in a multi-dimensional array
CN104137055B (zh) * 2011-12-29 2018-06-05 英特尔公司 点积处理器、方法、系统和指令
KR101885211B1 (ko) * 2012-01-27 2018-08-29 삼성전자 주식회사 Ｇｐｕ의 자원 할당을 위한 방법 및 장치
US9104421B2 (en) * 2012-07-30 2015-08-11 Nvidia Corporation Training, power-gating, and dynamic frequency changing of a memory controller
WO2014022980A1 (en) * 2012-08-08 2014-02-13 Intel Corporation Isa bridging including support for call to overidding virtual functions
US9274772B2 (en) 2012-08-13 2016-03-01 Microsoft Technology Licensing, Llc. Compact type layouts
CN102902576B (zh) * 2012-09-26 2014-12-24 北京奇虎科技有限公司 一种渲染网页的方法、服务器和系统
US8982124B2 (en) 2012-09-29 2015-03-17 Intel Corporation Load balancing and merging of tessellation thread workloads
US9123167B2 (en) 2012-09-29 2015-09-01 Intel Corporation Shader serialization and instance unrolling
US9904625B2 (en) 2013-03-15 2018-02-27 Intel Corporation Methods, systems and apparatus for predicting the way of a set associative cache
WO2014150806A1 (en) 2013-03-15 2014-09-25 Soft Machines, Inc. A method for populating register view data structure by using register template snapshots
US10275255B2 (en) 2013-03-15 2019-04-30 Intel Corporation Method for dependency broadcasting through a source organized source view data structure
US10140138B2 (en) 2013-03-15 2018-11-27 Intel Corporation Methods, systems and apparatus for supporting wide and efficient front-end operation with guest-architecture emulation
WO2014150991A1 (en) 2013-03-15 2014-09-25 Soft Machines, Inc. A method for implementing a reduced size register view data structure in a microprocessor
EP2972845B1 (en) 2013-03-15 2021-07-07 Intel Corporation A method for executing multithreaded instructions grouped onto blocks
US9891924B2 (en) 2013-03-15 2018-02-13 Intel Corporation Method for implementing a reduced size register view data structure in a microprocessor
WO2014150971A1 (en) 2013-03-15 2014-09-25 Soft Machines, Inc. A method for dependency broadcasting through a block organized source view data structure
KR102083390B1 (ko) 2013-03-15 2020-03-02 인텔 코포레이션 네이티브 분산된 플래그 아키텍처를 이용하여 게스트 중앙 플래그 아키텍처를 에뮬레이션하는 방법
US9811342B2 (en) 2013-03-15 2017-11-07 Intel Corporation Method for performing dual dispatch of blocks and half blocks
US9886279B2 (en) 2013-03-15 2018-02-06 Intel Corporation Method for populating and instruction view data structure by using register template snapshots
US9569216B2 (en) 2013-03-15 2017-02-14 Soft Machines, Inc. Method for populating a source view data structure by using register template snapshots
US9535686B2 (en) 2013-03-15 2017-01-03 International Business Machines Corporation Dynamic library replacement
US9772864B2 (en) * 2013-04-16 2017-09-26 Arm Limited Methods of and apparatus for multidimensional indexing in microprocessor systems
US9600852B2 (en) * 2013-05-10 2017-03-21 Nvidia Corporation Hierarchical hash tables for SIMT processing and a method of establishing hierarchical hash tables
US9507594B2 (en) * 2013-07-02 2016-11-29 Intel Corporation Method and system of compiling program code into predicated instructions for execution on a processor without a program counter
US10628156B2 (en) * 2013-07-09 2020-04-21 Texas Instruments Incorporated Vector SIMD VLIW data path architecture
GB2524063B (en) 2014-03-13 2020-07-01 Advanced Risc Mach Ltd Data processing apparatus for executing an access instruction for N threads
US9471283B2 (en) * 2014-06-11 2016-10-18 Ca, Inc. Generating virtualized application programming interface (API) implementation from narrative API documentation
CN104077233B (zh) * 2014-06-18 2017-04-05 百度在线网络技术（北京）有限公司 多通道卷积层处理方法和装置
CN106406810B (zh) * 2014-07-02 2019-08-06 上海兆芯集成电路有限公司 微处理器及其方法
US10067768B2 (en) * 2014-07-18 2018-09-04 Nvidia Corporation Execution of divergent threads using a convergence barrier
US20160026486A1 (en) * 2014-07-25 2016-01-28 Soft Machines, Inc. An allocation and issue stage for reordering a microinstruction sequence into an optimized microinstruction sequence to implement an instruction set agnostic runtime architecture
US9398019B2 (en) 2014-08-07 2016-07-19 Vmware, Inc. Verifying caller authorization using secret data embedded in code
US9411979B2 (en) * 2014-08-07 2016-08-09 Vmware, Inc. Embedding secret data in code
US10922402B2 (en) 2014-09-29 2021-02-16 Vmware, Inc. Securing secret data embedded in code against compromised interrupt and exception handlers
US9749548B2 (en) 2015-01-22 2017-08-29 Google Inc. Virtual linebuffers for image signal processors
CN107408035B (zh) * 2015-03-27 2021-11-09 英特尔公司 用于缕程间通信的装置和方法
US9772852B2 (en) 2015-04-23 2017-09-26 Google Inc. Energy efficient processor core architecture for image processor
US9965824B2 (en) 2015-04-23 2018-05-08 Google Llc Architecture for high performance, power efficient, programmable image processing
US10095479B2 (en) * 2015-04-23 2018-10-09 Google Llc Virtual image processor instruction set architecture (ISA) and memory model and exemplary target hardware having a two-dimensional shift array structure
US9756268B2 (en) 2015-04-23 2017-09-05 Google Inc. Line buffer unit for image processor
US9785423B2 (en) 2015-04-23 2017-10-10 Google Inc. Compiler for translating between a virtual image processor instruction set architecture (ISA) and target hardware having a two-dimensional shift array structure
US9769356B2 (en) 2015-04-23 2017-09-19 Google Inc. Two dimensional shift array for image processor
US10291813B2 (en) 2015-04-23 2019-05-14 Google Llc Sheet generator for image processor
US10313641B2 (en) 2015-12-04 2019-06-04 Google Llc Shift register with reduced wiring complexity
US9830150B2 (en) 2015-12-04 2017-11-28 Google Llc Multi-functional execution lane for image processor
US10115175B2 (en) * 2016-02-19 2018-10-30 Qualcomm Incorporated Uniform predicates in shaders for graphics processing units
US10204396B2 (en) 2016-02-26 2019-02-12 Google Llc Compiler managed memory for image processor
US10387988B2 (en) 2016-02-26 2019-08-20 Google Llc Compiler techniques for mapping program code to a high performance, power efficient, programmable image processing hardware platform
US10380969B2 (en) 2016-02-28 2019-08-13 Google Llc Macro I/O unit for image processor
US20180007302A1 (en) 2016-07-01 2018-01-04 Google Inc. Block Operations For An Image Processor Having A Two-Dimensional Execution Lane Array and A Two-Dimensional Shift Register
US10546211B2 (en) 2016-07-01 2020-01-28 Google Llc Convolutional neural network on programmable two dimensional image processor
US20180005346A1 (en) 2016-07-01 2018-01-04 Google Inc. Core Processes For Block Operations On An Image Processor Having A Two-Dimensional Execution Lane Array and A Two-Dimensional Shift Register
US20180005059A1 (en) 2016-07-01 2018-01-04 Google Inc. Statistics Operations On Two Dimensional Image Processor
WO2018094087A1 (en) * 2016-11-17 2018-05-24 The Mathworks, Inc. Systems and methods for generating code for parallel processing units
US10204394B2 (en) 2017-04-10 2019-02-12 Intel Corporation Multi-frame renderer
US10437593B2 (en) * 2017-04-27 2019-10-08 Nvidia Corporation Techniques for comprehensively synchronizing execution threads
US10656964B2 (en) * 2017-05-16 2020-05-19 Oracle International Corporation Dynamic parallelization of a calculation process
US10754746B2 (en) 2017-11-15 2020-08-25 General Electric Company Virtual processor enabling unobtrusive observation of legacy systems for analytics in SoC
CN110825514B (zh) * 2018-08-10 2023-05-23 昆仑芯(北京)科技有限公司 人工智能芯片以及用于人工智能芯片的指令执行方法
GB2580327B (en) * 2018-12-31 2021-04-28 Graphcore Ltd Register files in a multi-threaded processor
US20200264921A1 (en) * 2019-02-20 2020-08-20 Nanjing Iluvatar CoreX Technology Co., Ltd. (DBA "Iluvatar CoreX Inc. Nanjing") Crypto engine and scheduling method for vector unit
CN110321193B (zh) * 2019-05-05 2022-03-18 四川盛趣时代网络科技有限公司 一种基于Direct3D共享纹理的交互方法及系统
US11216281B2 (en) 2019-05-14 2022-01-04 International Business Machines Corporation Facilitating data processing using SIMD reduction operations across SIMD lanes
CN110209509B (zh) * 2019-05-28 2021-08-17 北京星网锐捷网络技术有限公司 多核处理器间的数据同步方法及装置
CN110765082B (zh) * 2019-09-06 2023-11-24 深圳平安通信科技有限公司 Hadoop文件处理方法、装置、存储介质及服务器
CN114816529A (zh) * 2020-10-21 2022-07-29 上海壁仞智能科技有限公司 配置向量运算系统中的协作线程束的装置和方法
US20230090604A1 (en) * 2021-09-16 2023-03-23 T-Head (Shanghai) Semiconductor Co., Ltd. Parallel processing unit virtualization

Family Cites Families (19)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
DE69416152T2 (de) * 1993-10-06 1999-07-01 Honeywell Inc Virtueller graphikprozessor und verfahren für eingebettete echtzeitanzeigesysteme
US5692193A (en) * 1994-03-31 1997-11-25 Nec Research Institute, Inc. Software architecture for control of highly parallel computer systems
US5828880A (en) * 1995-07-06 1998-10-27 Sun Microsystems, Inc. Pipeline system and method for multiprocessor applications in which each of a plurality of threads execute all steps of a process characterized by normal and parallel steps on a respective datum
JPH10228382A (ja) * 1997-02-14 1998-08-25 Nec Corp コンパイル方式
US5991794A (en) * 1997-07-15 1999-11-23 Microsoft Corporation Component integration system for an application program
GB2336919A (en) 1998-04-30 1999-11-03 Ibm Pre-emptive threading in a virtual machine
KR20010072477A (ko) 1998-08-13 2001-07-31 썬 마이크로시스템즈, 인코포레이티드 가상 머신 환경에서 네이티브 코드를 변환하고 실행하는방법 및 장치
US7234139B1 (en) * 2000-11-24 2007-06-19 Catharon Productions, Inc. Computer multi-tasking via virtual threading using an interpreter
US6779049B2 (en) * 2000-12-14 2004-08-17 International Business Machines Corporation Symmetric multi-processing system with attached processing units being able to access a shared memory without being structurally configured with an address translation mechanism
US6839828B2 (en) * 2001-08-14 2005-01-04 International Business Machines Corporation SIMD datapath coupled to scalar/vector/address/conditional data register file with selective subpath scalar processing mode
US7275249B1 (en) * 2002-07-30 2007-09-25 Unisys Corporation Dynamically generating masks for thread scheduling in a multiprocessor system
JP4487479B2 (ja) * 2002-11-12 2010-06-23 日本電気株式会社 Ｓｉｍｄ命令シーケンス生成方法および装置ならびにｓｉｍｄ命令シーケンス生成用プログラム
US7000233B2 (en) * 2003-04-21 2006-02-14 International Business Machines Corporation Simultaneous multithread processor with result data delay path to adjust pipeline length for input to respective thread
DE602004017879D1 (de) * 2003-08-28 2009-01-02 Mips Tech Inc Integrierter mechanismus zum suspendieren und endznem prozessor
US6897871B1 (en) * 2003-11-20 2005-05-24 Ati Technologies Inc. Graphics processing architecture employing a unified shader
US20060150165A1 (en) * 2004-12-30 2006-07-06 Intel Corporation Virtual microengine systems and methods
US7861060B1 (en) * 2005-12-15 2010-12-28 Nvidia Corporation Parallel data processing systems and methods using cooperative thread arrays and thread identifier values to determine processing behavior
US8914618B2 (en) * 2005-12-29 2014-12-16 Intel Corporation Instruction set architecture-based inter-sequencer communications with a heterogeneous resource
US8010953B2 (en) * 2006-04-04 2011-08-30 International Business Machines Corporation Method for compiling scalar code for a single instruction multiple data (SIMD) execution engine

Also Published As

Publication number Publication date
JP4999183B2 (ja) 2012-08-15
US8321849B2 (en) 2012-11-27
JP2008276740A (ja) 2008-11-13
KR101026689B1 (ko) 2011-04-07
TW200844853A (en) 2008-11-16
DE202008017916U1 (de) 2010-11-04
DE102008005515A1 (de) 2008-08-07
US20080184211A1 (en) 2008-07-31
CN101231585A (zh) 2008-07-30
KR20080070599A (ko) 2008-07-30
SG144869A1 (en) 2008-08-28
TWI363294B (en) 2012-05-01

Similar Documents

Publication Publication Date Title
CN101231585B (zh) 定义用于目标平台结构的并行处理操作方法
US11182138B2 (en) Compiler for translating between a virtual image processor instruction set architecture (ISA) and target hardware having a two-dimensional shift array structure
Mark et al. Cg: A system for programming graphics hardware in a C-like language
CN102016926B (zh) 具有混合精度指令执行的可编程串流处理器
Tarditi et al. Accelerator: using data parallelism to program GPUs for general-purpose uses
Wang et al. EXOCHI: architecture and programming environment for a heterogeneous multi-core multithreaded system
US10216487B2 (en) Virtual image processor instruction set architecture (ISA) and memory model and exemplary target hardware having a two-dimensional shift array structure
JP2008276740A5 (zh)
Membarth et al. Generating device-specific GPU code for local operators in medical imaging
KR102232723B1 (ko) 2차원 실행 레인 어레이 및 2차원 시프트 레지스터를 갖는 이미지 프로세서상의 블록 연산을 위한 코어 프로세서
Seinstra et al. A software architecture for user transparent parallel image processing
CN103870242A (zh) 优化线程栈存储器的管理的系统、方法和计算机程序产品
CN102648449A (zh) 用于并行处理单元的陷阱处理器架构
CN108734637A (zh) 图形控制流机制
Klemm et al. JaMP: an implementation of OpenMP for a Java DSM
Gaudiot et al. The Sisal model of functional programming and its implementation
US20180082397A1 (en) Static Data Sharing Mechanism for a Heterogeneous Processing Environment
Balfour Efficient embedded computing
Köster et al. Massively parallel rule-based interpreter execution on GPUs using thread Compaction
Sousa et al. Data-flow analysis and optimization for data coherence in heterogeneous architectures
Kuo et al. The design of LLVM-based shader compiler for embedded architecture
Yang et al. Fei teng 64 stream processing system: architecture, compiler, and programming
Muller et al. Caches with compositional performance
Mammeri et al. Vcomputelib: Enabling cross-platform gpgpu on mobile and embedded gpus
Kritikakou et al. Near-optimal and scalable intrasignal in-place optimization for non-overlapping and irregular access schemes

Legal Events

Date Code Title Description
C06 Publication
PB01 Publication
C10 Entry into substantive examination
SE01 Entry into force of request for substantive examination
C14 Grant of patent or utility model
GR01 Patent grant