专利名称：：为了用于由通用处理器执行而对应用程序重定目标的制作方法为了用于由通用处理器执行而对应用程序重定目标相关专利申请的交叉引用本申请要求于2008年4月9日提交的美国临时申请序列号为61/043，708(律师案巻号No.NVDA/SC-08-0007-US0)的优先权，在此通过4爰引的方式纳入本说明书。才支术领域本发明的实施例大体上涉及编译器程序，更具体地涉及一种翻译器，所述翻译器用于将为了由多核图像处理器执行而编写的应用程序重定目标，以便由具有共享存储器的通用处理器执行。
背景技术：
：现代图像处理系统通常包括配置为以多线程方式执行应用程序的多核图像处理单元(GPU)。这种图像处理系统还包括存储器，该存储器具有在执行线程之间共享的部分和每个线程专用的部分。NVIDIA的CUDA(计算统一设备架构)技术提供了一种C语言环境，使得编程人员和开发人员能够编写解决复杂计算问题的软件应用程序，这些复杂计算问题诸如视频和音频编码、油气勘探建模和医学成像。应用程序配置为由多核GPU并行执行，并且通常依靠多核GPU的特定特4正。由于同样的特定特征在通用中央处理单元(CPU)中是不可用的，因此用CUDA编写的软件应用程序可能无法移植到通用CPU上运行。如前所述，在本领域中需要这样一种技术，这种技术不需要编程人员对应用程序进行修改就能让用并行编程模型编写的用于在多核GPU上执行的应用程序能够在通用CPU上运行。
发明内容本发明的一个实施例给出了一种将用并行编程模型编写的程序翻:^,为用于由通用CPU执行的方法。该方法包括接收用并行编程模型编写的用于由多核图像处理单元执行的应用程序的步骤，和将该应用程序划分成同步独立指令的区域以创建划分的应用程序的步骤。将一个循环插入到划分的应用程序中的至少一个区域附近，以创建用于由通用处理器执行的翻译后的应用程序。该循环在合作线程阵列维度上进行迭代，该线程阵列维度与由位于多核图像处理单元内的并行处理器同时执行的多个线程相对应。在此披露的方法的一个优势在于用并行编程模型编写的用于由多核GPU执行的应用程序可以无需修改而移植到通用CPU上。为了能够由通用CPU执行，将应用程序中依赖于多核GPU的特定特征的部分由翻译器进行转换。将应用程序划分为同步独立指令的区域。将这些指令分类为收敛的或发散的，并且将这些区域之间共享的发散存储器基准进行复制。在由通用CPU执行的过程中，插入线程循环以确保在各种线程之间能够正确地共享存储器。为了详细地理解本发明的上述特征，对于以上筒要说明的发明，将参照实施例进行更为具体的描述，其中对部分实施例结合附图进行了说明。然而，需要注意的是，附图中示出的只是本发明代表性的实施例，因此不能认为附图限制了本发明的范围，本发明可以适用于其他同样有效的实施例。图l是说明计算机系统的框图2是根据本发明的一个实施例对计算机系统进行说明的框图3A是根据本发明的一个实施例，将用于由多核图像处理单元执行而编写的代码翻译为由通用处理器执行的代码的方法步骤的流程图3B是根据本发明的一个实施例，说明翻译为划分后的代码的输入代码的概念图3C是根据本发明的一个实施例，说明翻译为优化后的代码的输入代码的概念图；以及图4是根据本发明的一个实施例，说明由通用处理器^丸行翻译后的代码的方法步骤的流程图。具体实施例方式解。然而，对于本领域:术人员来说显而易见的i,'即使缺少其中二个或多个这些具体细节也可以实施本发明。在其他例子中，为了避免引起与本发明的混淆，对一些公知的特征没有进行描述。5图l示出了配置为用于执行用CUDA编写的代码的计算机系统100的框图。计算机系统100包括CPU102和系统存储器104，两者通过包含存储器桥105的总线路径互相通信。存储器桥105例如可以是北桥芯片，通过总线或其他通信路径106(例如，超传输链接)与I/O(输7v/输出)桥107相连接。I/O桥107例如可以是南桥芯片，从一个或多个用户输入装置108(例如，键盘、鼠标)接收用户输入，并将该输入通过路径106和存储器桥105转发给CPU102。多线程处理子系统112通过总线或其他通信路径113(例如，PCIExpress、加速图像端口(AGP)或者超传输链接)与存储器桥105相连接。在一个实施例中，多线程处理子系统112是将像素输送到显示装置110(例如，传统的CRT或基于LCD的显示器)的图像子系统。系统盘114还与I/O桥107相连接。开关116为I/O桥107和诸如网络适配器118以及各种外插卡120和121的其他部件之间提供了连接。其他部件(图中没有示出)包括USB或其他端口连接、CD驱动器、DVD马区动器、电影刻录装置及类似的部件，也可以与1/0桥107相连接。在图1PCI(外设部件互连)、PCIExpress(PCI-E)、AGP(加速图像端口)、超传输或其他任何一种总线或点对点通信协议，并且不同装置之间的连接可以4吏用本领域已知的不同协议。CPU102作为计算机系统IOO的控制处理器来运行，管理和协调其他系统部件的工作。特别是，CPU102发出控制多线程处理子系统112中的并行处理器134工作的命令。在一些实施例中，CPU102将用于并行处理器134的命令流写入到命令緩冲器(图中未示出)，该命令緩冲器可以位于系统存储器104中、子系统存储器138中、或者是CPU102和并行处理器134都可以访问的其他存储位置中。并行处理器134从命令緩冲器中读出命令流，并且相对于CPU102的工作异步地执行这些命令。系统存储器104包括操作系统的执行映像、装置驱动器103和配置为用于由多线程处理子系统112执行的CUDA代码101。CUDA代码101包含了意在多线程处理子系统112上执行的编程指令。在本文的描述中，代码指的是任何计算机代码、指令、和/或可以用处理器执行的函数。例如，在各种实施例中，所述代码可以包含C代码、C十+代码等等。在一个实施例中，所述代码可以包括一种计算机语言的扩展语言(例如，C、C++的扩展等等)。操作系统提供了用于管理和协调计算机系统IOO工作的详细指令。装置驱动器103提供了用于管理和协调多线程处理子系统112,特别是并行处理器134工作的详细指令。另外，装置驱动器103可以提供编译功能，用来生成特别针对并行处理器134进行优化的机器代码。装置驱动器103可以结合由NVIDIA公司提供的CUD八tm框架来提供。在一个实施例中，多线程处理子系统112包含了一个或多个并行处理器134,该并行处理器134可以例如用一个或多个集成电路装置来实现，所述集成电路装置例如是可编程处理器、专用集成电路(ASICs)。并行处理器134可以包括针对图像和视频处理进行优化的电路，例如包括^L频输出电路和图像处理单元(GPU)。在另一个实施例中，多线程处理子系统112可以和一个或多个其他的系统元件集成，例如存储器桥105、CPU102和I/O桥107,以形成片上系统(SoC)。一个或多个并行处理器134可以将数据输出到显示装置110上，或者每一个并行处理器34可以将数据输出到一个或多个显示装置110上。并行处理器134有利地实现包括一个或多个处理内核的高度并行处理器，每一个处理内核能够同时执行大量的线程，其中每一个线程是一例程序，比如代码101。并行处理器134能够被编程以执行与各种广泛的应用相关的处理任务，所述应用包括但不限于线性和非线性数据变换、视频和/或音频数据的过滤、建模运算(例如，应用物理法则来确定物体的位置、速度和其他属性)、图像渲染操作(例如，镶嵌着色器、顶点着色器、几何着色器和/或像素着色器编程)等等。并行处理器134可以将数据从系统存储器104和/或本地子系统存储器138传输到本地(片上)存储器，对数据进行处理，并将结果数据写回到系统存储器104和/或子系统存储器138中，在那里这些数据能够由包括CPU102或另一多线程处理子系统112的其他系统部件来访问。并行处理器134可以设置有任意数量的子系统存储器138，也可以不包括子系统存储器138，并且可以使用子系统存储器138和系统存储器104的任意组合。例如，在统一存储器架构(UMA)的实施例中，并行处理器134可以是图像处理器。在这些实施例中，将会设置极少的甚至不设置任何专用子系统存储器138，并行处理器134将只使用或者几乎只使用系统存储器104。在UMA实施例中，并行处理器134可以被集成到桥芯片或处理器芯片中，或者是设置为具有高速链接(例如PCI-E)的分离的芯片，该芯片通过桥芯片或其他通信装置将并行处理器134连接到系统存储器104。7如上所述，多线程处理子系统112可以包括任意数量的并行处理器134。例如，多个并行处理器134可以设置在单个的外插卡上，或者多个外插卡可以与通信路径113相连，或者一个或多个并行处理器134可以被集成到桥芯片中。当存在有多个并行处理器134时，那些并行处理器134可以以高于单个并行处理器134可能达到的数据吞吐量来并行工作处理数据。包含有一个或多个并行处理器134的系统可以实现为各种配置和形式，包括桌上型电脑、笔记本电脑、或是手持个人计算机、服务器、工作站、游戏控制台、嵌入式系统等等。在并行处理器134的一些实施例中，使用单指令多数据(SIMD)指令发送技术来支持大量线程的并行执行，而无需提供多个独立指令单元。在其他实施例中，使用单指令多线程(SIMT)技术来支持大量大体上同步化的线程的并行执行。不同于其中所有处理引擎通常执行相同指令的SIMD执行方案，SIMT"l丸行允许不同的线程更易于通过给定的线程程序来跟随发散的执行路径。本领域技术人员能够理解的是，SIMD处理方案代表的是SIM丁处理方案的功能性子集。并行处理器134中的功能性单元支持多种运算，包括整数和浮点计算(例如，加法和乘法)、比较运算、布尔运算(AND、OR、XOR)、移位和各种代数函数的计算(例如，平面插值、三角、指数和对数函数等等)。传输到位于并行处理器134的处理内核(未示出)中的特定处理单元(未示出)的指令序列构成了线程，如在此前所定义的，在跨越处于一个处理内核中的多个处理单元同时执行的一定数量的线程集合在这里被称作"线程组"。如在此所使用的，"线程组"指的是一组对不同的输入数据执行同一程序的线程，该组里的每一个线程被分配给处理内核中不同的处理单元。线程组可以包含比处理单元的数量更少的线程，在此情况下，在该线程组正在进行处理的周期内，一些处理单元将处于空闲状态。线程组也可以包含比处理单元的数量更多的线程，在此情况下，处理将发生在多个时钟周期上。由于每一个处理内核能够同时支持多达G个线程组，因此遵循在任意给定的时间上，多达GxM个线程组可以在处理内核中执4亍，这里M是并行处理器134内的处理内核的数量。另外，处理内核中的多个相关的线程称为"合作线程阵列，，("CTA，，)。CTA的尺寸通常由编程人员和CTA可用的硬件资源的数量来决定，所述硬件资源诸如存储器或寄存器。CUDA编程模型反映的是GPU加速器的系统架构。每一个线程都有专用的本地地址空间，并且每一个CTA共享的地址空间被用来在CTA中的线程之间传递数据。处理内核还访问片外"全局"存储器，该存储器可以例如包括子系统存储器138和/或系统存储器104。CUDA应用程序的主机部分用常规的方法和工具来编译，而内核函数指定CTA处理。在最高层，CUDA存储器模型将主机和设备存储器空间分开，这样主机代码和内核代码就仅能直接访问它们各自的存储器空间。API(应用编程接口)函数允许在主机和设备存储器空间之间进行数据复制。在CUDA编程模型的共享存储器CPU执行过程中，控制CPU线程能够在没有潜在数据竟争的情况下与并行CTA并行执行。主机存储器空间由C编程语言定义，并且设备存储器空间被指定为全局的、恒定的、本地的、共享的和紋理的。所有线程可以访问这些全局的、恒定的和紋理的存储器空间。正如前面已经解释过的，对本地空间的访问限于单个线程，并且对共享空间的访问限于在CTA内的线程。这种存储器模型鼓励在对于低等待时间的访问时使用小存储器空间，并且鼓励对通常有更长等待时间的大存储器空间进行明智地使用。CUDA寿呈序，比如代码101,通常会纟皮组织成以一维、二维或三维形式(例如x、y和z)的一组同步或异步执行的CTA。三元索引唯一地识别线程块中的线程。线程块自身由隐含定义的二元变量来区分。这些索引的范围在运行时间时被定义，并且运行时间环境检查这些索S1是否符合任何硬件限制。每一个CTA可以和其他CTA—起由并行处理器134并行执行。很多CTA可以和每一个执行一个或多个CTA的并行处理器134并行运行。运行时间环境负责根据要求来同步或异步地管理CUDA代码101的执行。CTA内的线程通过使用共享存储器和被称为synchthreads()的壁垒同步图元互相通信并同步。CUDA确保在线程块内的线程将同时存活，并为线程块内的线程提供了架构以执行快速壁垒同步和本地数据共享。在(由一维或多维定义的)CTA内不同的线程块对于其创建、执行或退隐没有顺序的要求。另外，不允许并行CTA访问系统调用，包括1/0。CUDA编程模型仅仅执行并行CTA之间的全局同步，并为CTA中的块之间的有限通信提供内禀原子操作(intrinsicatomicoperations)。每个线程的主体被称作内核，用CUDA来指定，其在标准C中可以用存储器模型注释和壁垒同步图元来表示。CUDA程序的语义是，每一个内核由CTA内的全部线程以按照由壁垒同步图元所暗示的存储器排序的顺序来执行。特别是，在壁垒同步图元之前发生的CTA内的全部共享存储器基准必须在壁垒同步图元之后发生的任何共享存储器基准之前完成。内核代码中的每一个壁垒同步图元实例在概念上代表了单独的逻辑壁垒，并应当被处理为静态的。当CUDA线程可以采取不同的构建分支时，在如果-否则(if-else)构建的两条路径上都调用壁垒同步图元是非法的。虽然线程块内的所有线程都会到达其中一个同步图元，但它们代表的是单独的壁垒，每个都要求要么全部线程到达壁垒，要么没有线程到达壁垒。因此，这样的内核不会正确地执行。更普遍地，如果同步图元是包含在对于线程块中不同的线程表现不一样的任何控制流程构建内的话，则CUDA代码不能确保被正确地执行。图2是根据本发明的一个实施例对计算机系统200进行说明的框图。计算机系统200包括CPU202和系统存储器204,两者通过包含存储器桥.205的总线路径互相通信。存储器桥205例如可以是北桥芯片，与I/O(输入/输出)桥107通过总线或其他通信路径106(例如，超传输链接)相连接。CPU202产生输出以在显示装置210(例如，常规的CRT或基于LCD的显示器)上显示。多线程处理子系统112不包括在计算机系统200中，并且CUDA代码101没有被改写用于由通用处理器来执行，所述通用处理器比如CPU202。CUDA代码101被改写用于由多线程处理子系统112执行，并且使用翻i斧器220进行翻译以产生翻译后的代码201,该翻译后的代码201不包括壁垒同步图元。为了让CPU202运行由代码IOI表示的程序，代码101首先必须要被翻译成代码201。然后该翻译后的代码可以由编译器225编译成用于由CPU202来执行。编译器225可以执行专门针对于CPU202的优化。翻译代码指的是将由第一计算机语言编写的代码转换成由第二计算机语言编写的代码。编译代码指的是将由一种计算机语言(例如，源代码)编写的代码转换成由另一种计算机语言(例如，目标代码)编写的代码。翻译器220结合图3A进行描述，编译器225结合图4进行描述。编译器225可以包含在装置驱动器203中，该装置驱动器203配置为在代码101、代码201和CPU202之间接口。将运行时间环境227配置成实现编译后的代码的功能，例如，输入和输出、存储器管理等等。运行时间环境227还启动编译后的代码用于由CPU202执行。翻译器220执行优化转换，以便将跨越CUDA线程组的细粒度线程的操作串连成单个的CPU线程，而运行时间环境227将线程组调度为由CPU202并行处理的工作单元。10的这种可移植性的首要障碍是并行性的粒度。常规的CPU不支持单个CUDACTA所要求的上百个硬件线程的环境。因此，在通用CPU上实现CUDA编程模型的系统的首要目标是要将任务级的并行性分配给可用的CPU内核。同时，该系统必须将一个任务中的微线程整合成单个的CPU线程，以防止过度调度的开销和频繁的内核间同步化。图3A是根据本发明的一个实施例，将用于由多核图像处理单元(例如多线程处理子系统112)执行而编写的代码101翻译为由通用处理器(例如CPU202)执行的代码201的方法步骤的流程图。为了保存用于代码101中的壁垒同步图元语义，将翻译器220配置为执行如图3A所示的一个或多个步骤。通过将围绕在壁垒同步图元周围的代码101进行划分，翻译器220将并行的线程"展开"，降低了共享状态的使用，改进了用于存储器访;处i器执行。在^^改变其目标定为由多线程处理子系统、、U2执行的CUDA代码101的情况下，用CUP202来执行代码201有可能获得很好的执行效果。编译器225可以利用由CPU202提供的向量指令容量，并且在为了执行而编译代码201时进行优化。在步骤300中，翻译器220接收到用于由多核GPU执行而编写的代码101,所述多核GPU例如多线程处理子系统112或者包括一个或多个并行处理器134的处理器，所述代码101例如是CUDA代码101。在步骤300中接收到的代码可以表示为由边连接的基本块节点组成的控制流程图。每个基本块指定了由目标环境(例如CPU202)执行的才栗作。在步骤305中，翻译器220将围绕在壁垒同步图元周围的代码101进行划分，以生成划分后的代码。划分后的代码在图3B和图3C中示出，并且所述划分过程结合那些图来进行描述。同步划分是一个代码区域，在该区域内的操作顺序完全由该划分内的基本块的控制流和数据流属性来决定。划分具有这样的属性，即线程循环能够在一个划分的周围被插入以运行并行线程。通过将每个同步线程图元替换成边，将基本块节点分成不同的划分区域，所述控制流程图可以用来生成同步划分控制流程图。在步骤310中，划分后的代码被分类，以便每一个语句被识别为收敛的或发散的。划分后的代码可以包括表达式和语句。表达式是一种计算，该计算可以包含常量、隐式线程身份(threadID)和由编程人员创建命名的变量，但是没有副作用或赋值。简单的语句定义为导致单一赋值的计算表ii达式。通用语句也能代表壁垒、控制流有条件的或循环构建、或者是一系列语句块。CTA的维度x、y和z通过代码进行传播以决定每个操作是否依赖于一个或多个CTA维度。引用了维度x、y和/或z中的线程身份(线程标识符)的搡作被认为是发散的，这是因为引用了CTA维度的线程可以与执行期间同一个CTA内的其他线程所不同。例如，依赖于线程身份x(threadID.x)的操作对于x維度是发散的。另一个不依赖于线程身份x(threadID.x)的操作在x维度中是收敛的。发散的语句对于每一个它们所引用的CTA维度都要求有线程循环。在步骤315,使用分类信息来对划分后的代码进行性能优化，从而产生优化后的代码。例如，一个划分内的指令可以被重新排序以融合才喿作，这样具有同一类別的操作被放在一组，并能落入在步骤325中插入的同一个线程循环内。操作按照这样的方式排序，即在其变量向量中具有较少线程身份维度的操作放在依赖于较多线程身份维度的操作之前。这种重新排序是有效的，因为一个语句必须具有一个变量向量，该变量向量是它所依赖的语句的变量向量的超集。因此其变量向量中只有一个维度的语句不能依赖于任何在其变量向量中有不同的维度或有多于一个维度的语句。在步骤320中，根据需要将优化后的代码中的线程-本地存储器基准提升至阵列基准，以保证每一个对象实例有唯一的位置来存放一个值。特别的，从一个分区输送到另一个分区的数据需要进行复制，这样使其对每一个分区都是可用的。符合下面条件之一的变量将被提升至阵列基准具有交叉分区依赖性的本地变量(在一个分区中被赋值并且在另一个分区中被引用)。在步骤320中，翻译器220将线程-本地存储器基准提升至阵列基准。表1所示的程序包含同步化壁垒图元和发散基准。表1—global—voidfUnction(){intleftlndex,rightlndex;SharedMem[threadldX.x]=...;〃将值存储到共享存储器中leftlndex=...threadld.x..;rightlndex=...threadld.x;—synchthreads();=…(SharedMem[leftlndex]+SharedMem[rightlndex])/2.0;将表1所示的程序划分成在同步线程图元前的第一分区和在同步线程图元后的第二分区。第二分区包括基准(leftlndex和rightlndex),该基准在第一分区中进行计算并依赖于CTA维度。如果不提升发散基准，则第二分区将无法正确地使用由第一分区的最后一次迭代所计算的值。第二分区应该使用为第一分区的threaded,x的每一次相应的迭代所计算的值。为确保计算是正确的，将发散基准按表2所示的进行提升。表2voidfunction(){for(inttid一x=0;tid—x<dimblock.X;tid—x++){SharedMem[tid.x]=...;〃将值存入存储器中leftlndexArray[tid—x]=...threadld.x.rightlndexArray[tid一x]=…threadld.x;for(inttid一x=0;tid一x<dimblock.X;tid—x++){=…(SharedMem[leftlndexArray[tid—x]]+SharedMem[rightlndexArray[tid一x]])/20;在步骤325中，为在其变量向量中包含线程身份维度的那些语句生成线程循环。使用自适应循环嵌套来同时评估相当于循环互换、循环分裂和循环不变式消除的转换，以获得最佳的冗余消除效果。嵌套循环在线程身份元组的每一个维度的值上动态地生成，以最佳地适应应用，而不是假定特定的循环嵌套并在该嵌套的基础之上对应用进行评估。当语句在步骤315中被排序之后，可以仅在一些语句周围对线程身份维度生成循环，这些语句在其变量向量中包含该维度。为了消除循环开销，翻译器220可以融合相邻的语句组，在这些相邻的语句组中一个语句组所具有的变量向量是另一个语句组的变量向量的子集。图3B是根据本发明的一个实施例，说明被翻译成为划分后的代码350的输入代码101的概念图。输入代码330^皮配置为用于由多线程处理子系统112执行，并且包括由同步壁垒指令336进行分隔的代码序列331和332。CTA中的所有线程将在其中任一线程开始执行代码序列332之前完成代码序列331的执行。翻译器220将输入代码330进行划分以产生划分后的代码350,其中分区351包括由代码序列331代表的指令，分区352包括由代码序列332代表的指令。当划分后的代码350由本身不支持同步壁垒指令的通用处理器执行时，线程循环353^C插入到分区352周围，以确保同步语义得到维持。此例中，代码分区351包含收敛基准，分区352可以包含发散基准。因此，线程循环353在分区352周围被插入。在图3A的步骤325中，翻译器220将线程循环(比如线程循环353)插入到经过优化的代码中，以生成被翻译成由CPU202执4亍的代码201。每一个分区可以具有为每一个CTA维度插入的线程循环。同步分区和线程循环插入的例子在表3和4中示出。表3示出的程序被翻译为表4示出的程序。<table>tableseeoriginaldocumentpage14</column></row><table>表3的程序使用了明确的同步来确保存储器在CTA中各个线程之间的正确共享。翻译器220将程序划分为两个分区，每一个分区依赖于xCTA维度。因此，线程循环在这两个分区的每一个周围被插入以确保翻译后的程序按正确的顺序来执行操作。<table>tableseeoriginaldocumentpage14</column></row><table>=…(SharedMem[tid—x]+SharedMem[tid—x—1])/2.0;用于将程序翻译为由通用处理器来执行的更加简单的才支术是给每个CTA维度插入明确的线程循环，这样就不必为在同一分区内的基准而决定维度的依赖性。例如，表5示出的程序被翻译为如表6所示的程序。注意，表5中插入的一个或多个线程循环可能不是必须的，这是因为不用决定维度的依赖性就可以生成程序。表5—global—voidfonction(){Sharedl=...=Sharedl表6voidfunction()(for(inttid一x=0;tid_x<dimblock.X;tidx+十){for(inttid_y=0;tid_y<dimblock.Y;tid_y++){for(inttid—z=0;tid_z<dimblock.Z;tid—z++){Sharedl=...=Sharedl图3C是才艮据本发明的一个实施例，说明翻"^为优化后的代码360的输入代码333的概念图。将输入代码333配置为用于由多线程处理子系统112执行，并且包括由同步壁垒指令335进行分隔的代码序列334和338。CTA中的所有线程会在其中任一线程开始执行代码序列338之前完成代码其中分区361包括由代码序列334代表的指令，分区362、364和365包括由代码序列338代表的指令。15分区362包括在第一CTA维度上发散的第一部分指令。分区364包括收敛的第二部分指令。分区365包括在第二CTA维度上发散的第三部分指令。当划分后的代码360由本身不支持同步壁垒指令的通用处理器执行时，线程循环363被插入到分区362周围以确保同步语义得到维持。线程循环363在第一CTA维度上进行迭代。线程循环366在分区365周围被插入以便在第二CTA维度上进行迭代。表7示出了CUDA内核的例子，表8示出了用于由通用处理器执行的CUDA内核的翻译。该示例性内核与一系列小矩阵相乘。每一个线程块计算系列中一个小矩阵的相乘，而每一个线程为其块计算结果矩阵的一个元素。表7示例的CUDA内核(1)一global—small_mm—list(float*A—list,float*B—list,，constintsize){(2)floatsum;(3)intmatrix—start,col，row,out—index,i;(4)martrix—start=blockldx.x*size*size;(5)col=matrix—start+threadIDx.x;(6)row—matrix—start+threadldx.y*size);(7)sum=0.0;(8)for(i=0;i<size;i++)(9)sum+=A—list[row+i]*B_list[col+(i*size)];//在重写输入数据之前进行同步(10)—syncthread();(11)out—index=matrix—start+(threadldx.y*size)+threadldx.x;(12)A—list[out一index]=sum;要注意的是表7中位于第(9)行的语句具有变量向量(x,y),因为col(列)依赖于x维度且row(行)依赖于y维度。z维度从来不^f吏用，所以没有循环被插入在z上进行迭代。通常的成本分析技术可以用来决定如在表7中示出的示例性内核中的语句5和6这样的情况。由于每一个仅依赖于一个线程身份维度，因此选择x和y索引循环中的任一嵌套顺序将迫使语句的冗余执行，或者是分区的主循环嵌套之外的冗余循环。表8翻译后的CUDA内核(1)一global—small—mm—list(float*A—list,float*B—list,,constintsize)(2)floatsum[];(3)intmatrix一start[],col[]，row[],out—index,i;(4)matrix一start[threadID]=blocklDx.x*size*size;for(threadID.x=0;threadID.x<blockDim.x;threadID.x++){(5)col[threadID]=matrix—start+threadIDx.x;for(threadID.y=0;threadID.y<blockDim.y;threadID.y++){(6)row[threadID]=matrix_start[threadID]+(threadIDx.y*size);(7)sum[threadID]=0.0;(8)for(i[threadID]=0;i<size;i++)(9)sum[threadID]+=A—list[讓[threadID〗+i]*BJist[col[threadID]+(i*size)];(10)for(threadID.x=0;threadID.x<blockDim.x;threadID.x++){for(threadID.y=0;threadID.y<blockDim.y;threadID.y++){(11)out—index=matrix一start[threadID]+(threadID.y*size)+threadID.x;(12)A—list[out一index]=sum[threadID];图4是根据本发明的一个实施例，由通用处理器(比如CPU202)执行翻译后的代码201的方法步骤的流程图。在步骤400中，编译器225将翻译后的代码201进行编译，选择性地执行针对CPU的优化，以生成编译后的代码。在步骤405中，在CPU202中可用的执行内核400的数量由设备驱动器203决定。为了提高性能，翻译后的代码201可以自动地进行縮17放以用于在可用4丸行内核上执行。在步骤410中，运行时间环境227或装置驱动器203将CPU202配置为启用执行翻译后的代码201的该数量的执4亍内核。运行时间环境227可以生成多个操作系统(OS)运行时间线程，该运行时间线程能够由环境变量控制。默认状态下，系统中内核的数量可以用作OS运行时间线程的数量。在步骤410中，可以对将要启动的CUDA线程的数量进行评估，并且在统计上划分为运行时间线程的数量。每一个运行时间线程顺序地执行一部分编译后的代码，并在壁垒上等待。当所有的运行时间线程都到达了壁昼，则CTA完成。在步骤415中，运行时间环境227或装置驱动器203启动编译后的代码用于由CPU202执行。翻译器220、编译器225以及运行时间环境227被用来将CUDA应用程序转换为由通用CPU执行的代码。CUDA编程模型支持批量的同步任务并行，其中每个任务由细粒度SPMD线程组成。CUDA编程模型的使用已经被限制于那些愿意为由GPU执行而编写特殊代码的程序员。这些特殊代码可以转换为由通用CPU来执行，并不需要程序员来重新编写CUDA应用程序。由CUDA支持的三个关键的抽象概念是SPMD线程块、壁垒同步和共享的存储器。翻译器220将遍及CUDA线程块的细粒度线程串行化为单一的CPU线程，并且执行优化变形以转换CUDA应用程序。虽然前面所述的是本发明的实施例，在不背离本发明基本范围的前提下可以设计出本发明更多的实施例。例如，本发明的某些方面可以由^^更件或软件来实现，或者是由硬件与软件结合在一起来实现。本发明的一个实施例可以实现为计算机系统所使用的程序产品。程序产品的程序对实施例的功能(包括在此描述的方法)进行定义，并且能够被包含在各种各样的计算机可读存储介质内。说明性的计算机可读存储介质包括但不限于(i)信息在其上永久保存的非可写存储介质(例如，计算机内的只读存储装置，如可被CD-ROM驱动器读出的CD-ROM盘、闪存、ROM芯片或者任意类型的固态非易失性半导体存储器)；以及(ii)其上存储有可改变的信息的可写存储介质(例如，磁盘驱动器内的软盘或硬盘驱动器或任意类型的固态随机存取半导体存储器)。当携带有引导本发明的功能的计算机可读指令时，这样的计算机可读存储介质就是本发明的实施例。因此，本发明的范围由以下的权利要求来界定。18权利要求1.一种计算系统，所述计算系统被配置用于将应用程序翻译为由通用处理器来执行，所述计算系统包括处理器，所述处理器被配置用于执行翻译器；和系统存储器，所述系统存储器与所述处理器连接并被配置用于存储所述翻译器、第一应用程序和第二应用程序，所述第一应用程序使用并行编程模型来编写以用于在多核图像处理单元上执行，所述第二应用程序被配置用于由通用处理器执行，以及所述翻译器被配置用于接收所述第一应用程序；将所述第一应用程序划分为同步独立指令的区域以生成划分后的应用程序；以及在至少其中一个划分后的应用程序周围插入循环，以生成所述第二应用程序，其中所述循环在合作线程阵列维度上进行迭代，所述合作线程阵列维度对应于由多核图像处理单元内的并行处理器同时执行的多个线程。2.根据权利要求1的计算系统，其特征在于，所述翻译器进一步被配置用于识别所述应用程序内的同步壁垒指令。3.根据权利要求2的计算系统，其特征在于，划分后的应用程序的第一区域包括位于所述同步壁垒指令之前的指令，划分后的应用程序的第二区域包括位于所述同步壁垒指令之后的指令。4.根据权利要求3的计算系统，其特征在于，所述插入循环的步骤包括在所述划分后的应用程序的第一区域周围插入第一循环，以确保在所述合作线程阵列中的任一线程开始执行所述划分后应用程序的第二区区域的执行。5.根据权利要求2的计算系统，其特征在于，所述第一应用程序表示为包括由边连接的基本块节点的控制流程图。6.根据权利要求5的计算系统，其特征在于，所述翻译器进一步被配置为将所述同步壁垒指令替换成边，以便将所述基本块节点分隔成与第一区域对应的第一基本块节点和与第二区域对应的第二基本块节点。7.根据权利要求1的计算系统，其特征在于，所述翻译器进一步被配置用于将所述划分后的应用程序进行分类，以便识别每个语句相对于所述合作线程阵列维度是收敛的还是发散的。8.根据权利要求1的计算系统，其特征在于，所述翻译器进一步被配置用于在所述划分后的应用程序的至少一个区域周围插入额外的循环，以生成所述第二应用程序，其中所述额外的循环在不同的合作线程阵列维度上进行迭代。9.根据权利要求1的计算系统，其特征在于，所述第一应用程序是CUDA(计算统一设备架构)应用程序。10.根据权利要求1的计算系统，其特征在于，进一步包括运行时间库，所述运行时间库将所述通用处理器配置成能使所述通用处理器内的多个执行内核执行所述第二应用程序。全文摘要本发明公开了一种为了用于由通用处理器执行而对应用程序重定目标。本发明的一个实施例提出了一种技术，用于将使用并行编程模型编写的用于在多核图像处理单元(GPU)上执行的应用程序翻译成用于由通用中央处理单元(CPU)执行。所述应用程序中依赖于多核GUP的特定特征的部分由翻译器转换成由通用CPU执行。所述应用程序被划分为同步独立指令区域。所述指令被分类为收敛的或发散的，并且在区域之间共享的发散存储器基准被复制。插入线程循环，以确保在由通用CPU执行期间各种线程之间存储器的正确共享。文档编号G06F9/45GK101556544SQ200910117898公开日2009年10月14日申请日期2009年4月9日优先权日2008年4月9日发明者巴斯蒂安·约翰·马特乌斯·阿特斯,维诺德·格罗夫,贾扬特·B·科尔希,迈克尔·墨菲,道格拉斯·塞勒,鲍里斯·柏林申请人:辉达公司