CN101556543A - 由通用处理器执行重定目标的图形处理器加速代码 - Google Patents

由通用处理器执行重定目标的图形处理器加速代码 Download PDF

Info

Publication number
CN101556543A
CN101556543A CNA2009101178975A CN200910117897A CN101556543A CN 101556543 A CN101556543 A CN 101556543A CN A2009101178975 A CNA2009101178975 A CN A2009101178975A CN 200910117897 A CN200910117897 A CN 200910117897A CN 101556543 A CN101556543 A CN 101556543A
Authority
CN
China
Prior art keywords
application program
code
execution
thread
computing system
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.)
Granted
Application number
CNA2009101178975A
Other languages
English (en)
Other versions
CN101556543B (zh
Inventor
维诺德·格罗夫
巴斯蒂安·约翰·马特乌斯·阿特斯
迈克尔·墨菲
贾扬特·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 CN101556543A publication Critical patent/CN101556543A/zh
Application granted granted Critical
Publication of CN101556543B publication Critical patent/CN101556543B/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
    • 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
    • G06F9/45537—Provision of facilities of other operating environments, e.g. WINE
    • 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
    • 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
    • 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/46—Multiprogramming arrangements
    • G06F9/48—Program initiating; Program switching, e.g. by interrupt
    • G06F9/4806—Task transfer initiation or dispatching
    • G06F9/4843—Task transfer initiation or dispatching by program, e.g. task dispatcher, supervisor, operating system
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F11/00—Error detection; Error correction; Monitoring
    • G06F11/22—Detection or location of defective computer hardware by testing during standby operation or during idle time, e.g. start-up testing
    • G06F11/26—Functional testing
    • G06F11/261—Functional testing by simulating additional hardware, e.g. fault simulation
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F12/00—Accessing, addressing or allocating within memory systems or architectures
    • G06F12/02—Addressing or allocation; Relocation
    • G06F12/0223—User address space allocation, e.g. contiguous or non contiguous base addressing
    • G06F12/023—Free address space management
    • G06F12/0253—Garbage collection, i.e. reclamation of unreferenced memory

Abstract

本发明公开了一种由通用处理器执行重定目标的图形处理器加速代码。本发明的一个实施例提出了一种技术，用于将使用并行编程模型编写的用于在多核图像处理单元(GPU)上执行的应用程序翻译成用于由通用中央处理单元(CPU)执行。所述应用程序中依赖于多核GUP的特定特征的部分由翻译器转换成由通用CPU执行。所述应用程序被划分为同步独立指令区域。所述指令被分类为收敛的或发散的，并且在区域之间共享的发散存储器基准被复制。插入线程循环，以确保在由通用CPU执行期间各种线程之间存储器的正确共享。

Description

由通用处理器执行重定目标的图形处理器加速代码
相关专利申请的交叉引用
本申请要求于2008年4月9日提交的美国临时申请序列号为61/043,708(律师案卷号No.NVDA/SC-08-0007-US0)，题为“用于在多核体系结构上执行GPU加速代码的系统”的优先权。该相关申请的主题在此通过援引的方式纳入本说明书。
技术领域
本发明的实施例大体上涉及编译器程序，更具体地涉及一种应用程序，所述应用程序被编写用于由多核图形处理器执行，并且被重定目标用于由具有共享存储器的通用处理器执行。
背景技术
现代图像处理系统通常包括配置为以多线程方式执行应用程序的多核图像处理单元(GPU)。这种图像处理系统还包括存储器，该存储器具有在执行线程之间共享的部分和每个线程专用的部分。
NVIDIA的CUDATM(计算统一设备架构)技术提供了一种C语言环境，使得编程人员和开发人员能够编写解决复杂计算问题的软件应用程序，这些复杂计算问题诸如视频和音频编码、油气勘探建模和医学成像。应用程序配置为由多核GPU并行执行，并且通常依靠多核GPU的特定特征。由于同样的特定特征在通用中央处理单元(CPU)中是不可用的，因此用CUDA编写的软件应用程序可能无法移植到通用CPU上运行。
如前所述，在本领域中需要这样一种技术，这种技术不需要编程人员对应用程序进行修改就能让用并行编程模型编写的用于在多核GPU上执行的应用程序能够在通用CPU上运行。
发明内容
本发明的一个实施例给出了一种用于配置通用处理器以执行翻译后的应用程序的方法。该方法包括接收翻译后的应用程序的步骤，所述翻译后的应用程序是由用并行编程模型编写的用于在多核图像处理单元上执行的应用程序转换而来的，和编译翻译后的应用程序的步骤，以生成用于由通用处理器执行的编译后的代码。可用来执行编译后的代码的通用处理器内执行内核的数量被确定，并且将通用处理器配置成启用该数量的执行内核。编译后的代码被启动用于由包含该数量的执行内核的通用处理器来执行。
在此披露的方法的一个优势在于：用并行编程模型编写的用于由多核GPU执行的应用程序可以无需修改而移植到通用CPU上。为了能够由通用CPU执行，将应用程序中依赖于多核GPU的特定特征的部分由翻译器进行转换。将应用程序划分为同步独立指令的区域。将这些指令分类为收敛的或发散的，并且将这些区域之间共享的发散存储器基准进行复制。在由通用CPU执行的过程中，插入线程循环以确保在各种线程之间能够正确地共享存储器。
附图说明
为了详细地理解本发明的上述特征，对于以上简要说明的发明，将参照实施例进行更为具体的描述，其中对部分实施例结合附图进行了说明。然而，需要注意的是，附图中示出的只是本发明代表性的实施例，因此不能认为附图限制了本发明的范围，本发明可以适用于其他同样有效的实施例。
图1是说明计算机系统的框图；
图2是根据本发明的一个实施例对计算机系统进行说明的框图；
图3A是根据本发明的一个实施例，将用于由多核图像处理单元执行而编写的代码翻译为由通用处理器执行的代码的方法步骤的流程图；
图3B是根据本发明的一个实施例，说明翻译为划分后的代码的输入代码的概念图；
图3C是根据本发明的一个实施例，说明翻译为优化后的代码的输入代码的概念图；以及
图4是根据本发明的一个实施例，说明由通用处理器执行翻译后的代码的方法步骤的流程图。
具体实施方式
在下面的描述中，提出了诸多具体细节以提供对本发明更加彻底的理解。然而，对于本领域技术人员来说显而易见的是，即使缺少其中一个或多个这些具体细节也可以实施本发明。在其他例子中，为了避免引起与本发明的混淆，对一些公知的特征没有进行描述。
图1示出了配置为用于执行用CUDA编写的代码的计算机系统100的框图。计算机系统100包括CPU 102和系统存储器104，两者通过包含存储器桥105的总线路径互相通信。存储器桥105例如可以是北桥芯片，通过总线或其他通信路径106(例如，超传输链接)与I/O(输入/输出)桥107相连接。I/O桥107例如可以是南桥芯片，从一个或多个用户输入装置108(例如，键盘、鼠标)接收用户输入，并将该输入通过路径106和存储器桥105转发给CPU 102。多线程处理子系统112通过总线或其他通信路径113(例如，PCI Express、加速图像端口(AGP)或者超传输链接)与存储器桥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(加速图像端口)、超传输或其他任何一种总线或点对点通信协议，并且不同装置之间的连接可以使用本领域已知的不同协议。
CPU 102作为计算机系统100的控制处理器来运行，管理和协调其他系统部件的工作。特别是，CPU 102发出控制多线程处理子系统112中的并行处理器134工作的命令。在一些实施例中，CPU 102将用于并行处理器134的命令流写入到命令缓冲器(图中未示出)，该命令缓冲器可以位于系统存储器104中、子系统存储器138中、或者是CPU 102和并行处理器134都可以访问的其他存储位置中。并行处理器134从命令缓冲器中读出命令流，并且相对于CPU 102的工作异步地执行这些命令。
系统存储器104包括操作系统的执行映像、装置驱动器103和配置为用于由多线程处理子系统112执行的CUDA代码101。CUDA代码101包含了意在多线程处理子系统112上执行的编程指令。在本文的描述中，代码指的是任何计算机代码、指令、和/或可以用处理器执行的函数。例如，在各种实施例中，所述代码可以包含C代码、C++代码等等。在一个实施例中，所述代码可以包括一种计算机语言的扩展语言(例如，C、C++的扩展等等)。
操作系统提供了用于管理和协调计算机系统100工作的详细指令。装置驱动器103提供了用于管理和协调多线程处理子系统112，特别是并行处理器134工作的详细指令。另外，装置驱动器103可以提供编译功能，用来生成特别针对并行处理器134进行优化的机器代码。装置驱动器103可以结合由NVIDIA公司提供的CUDATM框架来提供。
在一个实施例中，多线程处理子系统112包含了一个或多个并行处理器134，该并行处理器134可以例如用一个或多个集成电路装置来实现，所述集成电路装置例如是可编程处理器、专用集成电路(ASICs)。并行处理器134可以包括针对图像和视频处理进行优化的电路，例如包括视频输出电路和图像处理单元(GPU)。在另一个实施例中，多线程处理子系统112可以和一个或多个其他的系统元件集成，例如存储器桥105、CPU 102和I/O桥107，以形成片上系统(SoC)。一个或多个并行处理器134可以将数据输出到显示装置110上，或者每一个并行处理器134可以将数据输出到一个或多个显示装置110上。
并行处理器134有利地实现包括一个或多个处理内核的高度并行处理器，每一个处理内核能够同时执行大量的线程，其中每一个线程是一例程序，比如代码101。并行处理器134能够被编程以执行与各种广泛的应用相关的处理任务，所述应用包括但不限于线性和非线性数据变换、视频和/或音频数据的过滤、建模运算(例如，应用物理法则来确定物体的位置、速度和其他属性)、图像渲染操作(例如，镶嵌着色器、顶点着色器、几何着色器和/或像素着色器编程)等等。并行处理器134可以将数据从系统存储器104和/或本地子系统存储器138传输到本地(片上)存储器，对数据进行处理，并将结果数据写回到系统存储器104和/或子系统存储器138中，在那里这些数据能够由包括CPU 102或另一多线程处理子系统112的其他系统部件来访问。
并行处理器134可以设置有任意数量的子系统存储器138，也可以不包括子系统存储器138，并且可以使用子系统存储器138和系统存储器104的任意组合。例如，在统一存储器架构(UMA)的实施例中，并行处理器134可以是图像处理器。在这些实施例中，将会设置极少的甚至不设置任何专用子系统存储器138，并行处理器134将只使用或者几乎只使用系统存储器104。在UMA实施例中，并行处理器134可以被集成到桥芯片或处理器芯片中，或者是设置为具有高速链接(例如PCI-E)的分离的芯片，该芯片通过桥芯片或其他通信装置将并行处理器134连接到系统存储器104。
如上所述，多线程处理子系统112可以包括任意数量的并行处理器134。例如，多个并行处理器134可以设置在单个的外插卡上，或者多个外插卡可以与通信路径113相连，或者一个或多个并行处理器134可以被集成到桥芯片中。当存在有多个并行处理器134时，那些并行处理器134可以以高于单个并行处理器134可能达到的数据吞吐量来并行工作处理数据。包含有一个或多个并行处理器134的系统可以实现为各种配置和形式，包括桌上型电脑、笔记本电脑、或是手持个人计算机、服务器、工作站、游戏控制台、嵌入式系统等等。
在并行处理器134的一些实施例中，使用单指令多数据(SIMD)指令发送技术来支持大量线程的并行执行，而无需提供多个独立指令单元。在其他实施例中，使用单指令多线程(SIMT)技术来支持大量大体上同步化的线程的并行执行。不同于其中所有处理引擎通常执行相同指令的SIMD执行方案，SIMT执行允许不同的线程更易于通过给定的线程程序来跟随发散的执行路径。本领域技术人员能够理解的是，SIMD处理方案代表的是SIMT处理方案的功能性子集。并行处理器134中的功能性单元支持多种运算，包括整数和浮点计算(例如，加法和乘法)、比较运算、布尔运算(AND、OR、XOR)、移位和各种代数函数的计算(例如，平面插值、三角、指数和对数函数等等)。
传输到位于并行处理器134的处理内核(未示出)中的特定处理单元(未示出)的指令序列构成了线程，如在此前所定义的，在跨越处于一个处理内核中的多个处理单元同时执行的一定数量的线程集合在这里被称作“线程组”。如在此所使用的，“线程组”指的是一组对不同的输入数据执行同一程序的线程，该组里的每一个线程被分配给处理内核中不同的处理单元。线程组可以包含比处理单元的数量更少的线程，在此情况下，在该线程组正在进行处理的周期内，一些处理单元将处于空闲状态。线程组也可以包含比处理单元的数量更多的线程，在此情况下，处理将发生在多个时钟周期上。
由于每一个处理内核能够同时支持多达G个线程组，因此遵循在任意给定的时间上，多达G×M个线程组可以在处理内核中执行，这里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。三元索引唯一地识别线程块中的线程。线程块自身由隐含定义的二元变量来区分。这些索引的范围在运行时间时被定义，并且运行时间环境检查这些索引是否符合任何硬件限制。每一个CTA可以和其他CTA一起由并行处理器134并行执行。很多CTA可以和每一个执行一个或多个CTA的并行处理器134并行运行。运行时间环境负责根据要求来同步或异步地管理CUDA代码101的执行。CTA内的线程通过使用共享存储器和被称为synchthreads()的壁垒同步图元互相通信并同步。CUDA确保在线程块内的线程将同时存活，并为线程块内的线程提供了架构以执行快速壁垒同步和本地数据共享。在(由一维或多维定义的)CTA内不同的线程块对于其创建、执行或退隐没有顺序的要求。另外，不允许并行CTA访问系统调用，包括I/O。CUDA编程模型仅仅执行并行CTA之间的全局同步，并为CTA中的块之间的有限通信提供内禀原子操作(intrinsic atomic operations)。
每个线程的主体被称作内核，用CUDA来指定，其在标准C中可以用存储器模型注释和壁垒同步图元来表示。CUDA程序的语义是，每一个内核由CTA内的全部线程以按照由壁垒同步图元所暗示的存储器排序的顺序来执行。特别是，在壁垒同步图元之前发生的CTA内的全部共享存储器基准必须在壁垒同步图元之后发生的任何共享存储器基准之前完成。
内核代码中的每一个壁垒同步图元实例在概念上代表了单独的逻辑壁垒，并应当被处理为静态的。当CUDA线程可以采取不同的构建分支时，在如果-否则(if-else)构建的两条路径上都调用壁垒同步图元是非法的。虽然线程块内的所有线程都会到达其中一个同步图元，但它们代表的是单独的壁垒，每个都要求要么全部线程到达壁垒，要么没有线程到达壁垒。因此，这样的内核不会正确地执行。更普遍地，如果同步图元是包含在对于线程块中不同的线程表现不一样的任何控制流程构建内的话，则CUDA代码不能确保被正确地执行。
图2是根据本发明的一个实施例对计算机系统200进行说明的框图。计算机系统200包括CPU 202和系统存储器204，两者通过包含存储器桥205的总线路径互相通信。存储器桥205例如可以是北桥芯片，与I/O(输入/输出)桥107通过总线或其他通信路径106(例如，超传输链接)相连接。CPU 202产生输出以在显示装置210(例如，常规的CRT或基于LCD的显示器)上显示。
多线程处理子系统112不包括在计算机系统200中，并且CUDA代码101没有被改写用于由通用处理器来执行，所述通用处理器比如CPU 202。CUDA代码101被改写用于由多线程处理子系统112执行，并且使用翻译器220进行翻译以产生翻译后的代码201，该翻译后的代码201不包括壁垒同步图元。为了让CPU 202运行由代码101表示的程序，代码101首先必须要被翻译成代码201。然后该翻译后的代码可以由编译器225编译成用于由CPU 202来执行。编译器225可以执行专门针对于CPU 202的优化。翻译代码指的是将由第一计算机语言编写的代码转换成由第二计算机语言编写的代码。编译代码指的是将由一种计算机语言(例如，源代码)编写的代码转换成由另一种计算机语言(例如，目标代码)编写的代码。翻译器220结合图3A进行描述，编译器225结合图4进行描述。编译器225可以包含在装置驱动器203中，该装置驱动器203配置为在代码101、代码201和CPU 202之间接口。将运行时间环境227配置成实现编译后的代码的功能，例如，输入和输出、存储器管理等等。运行时间环境227还启动编译后的代码用于由CPU 202执行。翻译器220执行优化转换，以便将跨越CUDA线程组的细粒度线程的操作串连成单个的CPU线程，而运行时间环境227将线程组调度为由CPU 202并行处理的工作单元。
阻碍被设计为在GPU上运行的CUDA应用程序用于由通用CPU执行的这种可移植性的首要障碍是并行性的粒度。常规的CPU不支持单个CUDA CTA所要求的上百个硬件线程的环境。因此，在通用CPU上实现CUDA编程模型的系统的首要目标是要将任务级的并行性分配给可用的CPU内核。同时，该系统必须将一个任务中的微线程整合成单个的CPU线程，以防止过度调度的开销和频繁的内核间同步化。
图3A是根据本发明的一个实施例，将用于由多核图像处理单元(例如多线程处理子系统112)执行而编写的代码101翻译为由通用处理器(例如CPU 202)执行的代码201的方法步骤的流程图。为了保存用于代码101中的壁垒同步图元语义，将翻译器220配置为执行如图3A所示的一个或多个步骤。通过将围绕在壁垒同步图元周围的代码101进行划分，翻译器220将并行的线程“展开”，降低了共享状态的使用，改进了用于存储器访问的基准的位置，并插入了线程循环以转换CUDA专用代码以便用于由通用处理器执行。在不改变其目标定为由多线程处理子系统112执行的CUDA代码101的情况下，用CUP 202来执行代码201有可能获得很好的执行效果。编译器225可以利用由CPU 202提供的向量指令容量，并且在为了执行而编译代码201时进行优化。
在步骤300中，翻译器220接收到用于由多核GPU执行而编写的代码101，所述多核GPU例如多线程处理子系统112或者包括一个或多个并行处理器134的处理器，所述代码101例如是CUDA代码101。在步骤300中接收到的代码可以表示为由边连接的基本块节点组成的控制流程图。每个基本块指定了由目标环境(例如CPU 202)执行的操作。在步骤305中，翻译器220将围绕在壁垒同步图元周围的代码101进行划分，以生成划分后的代码。划分后的代码在图3B和图3C中示出，并且所述划分过程结合那些图来进行描述。同步划分是一个代码区域，在该区域内的操作顺序完全由该划分内的基本块的控制流和数据流属性来决定。划分具有这样的属性，即线程循环能够在一个划分的周围被插入以运行并行线程。通过将每个同步线程图元替换成边，将基本块节点分成不同的划分区域，所述控制流程图可以用来生成同步划分控制流程图。
在步骤310中，划分后的代码被分类，以便每一个语句被识别为收敛的或发散的。划分后的代码可以包括表达式和语句。表达式是一种计算，该计算可以包含常量、隐式线程身份(threadID)和由编程人员创建命名的变量，但是没有副作用或赋值。简单的语句定义为导致单一赋值的计算表达式。通用语句也能代表壁垒、控制流有条件的或循环构建、或者是一系列语句块。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_void function(){
   int leftIndex，rightIndex；
   SharedMem[threadIdX.x]＝...；//将值存储到共享存储器中
   leftIndex＝...threadId.x...；
   rightIndex＝...threadId.x；
   _synchthreads()；
   ＝...(SharedMem[leftIndex]+SharedMem[rightIndex])/2.0；
}
将表1所示的程序划分成在同步线程图元前的第一分区和在同步线程图元后的第二分区。第二分区包括基准(leftIndex和rightIndex)，该基准在第一分区中进行计算并依赖于CTA维度。如果不提升发散基准，则第二分区将无法正确地使用由第一分区的最后一次迭代所计算的值。第二分区应该使用为第一分区的threaded.x的每一次相应的迭代所计算的值。为确保计算是正确的，将发散基准按表2所示的进行提升。
表2
void function(){
   ...
   for(int tid_x＝0；tid_x＜dimblock.X；tid_x++){
   SharedMem[tid.x]＝...；//将值存入存储器中
   leftIndexArray[tid_x]＝...threadId.x...；
   rightIndexArray[tid_x]＝...threadId.x；
   }
   for(int tid_x＝0；tid_x＜dimblock.X；tid_x++){
＝...(SharedMem[leftIndexArray[tid_x]]+
         SharedMem[rightIndexArray[tid_x]])/2.0；
   }
}
在步骤325中，为在其变量向量中包含线程身份维度的那些语句生成线程循环。使用自适应循环嵌套来同时评估相当于循环互换、循环分裂和循环不变式消除的转换，以获得最佳的冗余消除效果。嵌套循环在线程身份元组的每一个维度的值上动态地生成，以最佳地适应应用，而不是假定特定的循环嵌套并在该嵌套的基础之上对应用进行评估。当语句在步骤315中被排序之后，可以仅在一些语句周围对线程身份维度生成循环，这些语句在其变量向量中包含该维度。为了消除循环开销，翻译器220可以融合相邻的语句组，在这些相邻的语句组中一个语句组所具有的变量向量是另一个语句组的变量向量的子集。
图3B是根据本发明的一个实施例，说明被翻译成为划分后的代码350的输入代码101的概念图。输入代码330被配置为用于由多线程处理子系统112执行，并且包括由同步壁垒指令336进行分隔的代码序列331和332。CTA中的所有线程将在其中任一线程开始执行代码序列332之前完成代码序列331的执行。翻译器220将输入代码330进行划分以产生划分后的代码350，其中分区351包括由代码序列331代表的指令，分区352包括由代码序列332代表的指令。当划分后的代码350由本身不支持同步壁垒指令的通用处理器执行时，线程循环353被插入到分区352周围，以确保同步语义得到维持。此例中，代码分区351包含收敛基准，分区352可以包含发散基准。因此，线程循环353在分区352周围被插入。
在图3A的步骤325中，翻译器220将线程循环(比如线程循环353)插入到经过优化的代码中，以生成被翻译成由CPU 202执行的代码201。每一个分区可以具有为每一个CTA维度插入的线程循环。同步分区和线程循环插入的例子在表3和4中示出。表3示出的程序被翻译为表4示出的程序。
表3
_global_void function(){
      SharedMem[threadIDX.x]＝...；//将值存入共享存储器中
      _synchthreads()；
      ＝...(SharedMem[threadIDX.x]+SharedMem[threadIdX.x-1])/2.0；
}
表3的程序使用了明确的同步来确保存储器在CTA中各个线程之间的正确共享。翻译器220将程序划分为两个分区，每一个分区依赖于x CTA维度。因此，线程循环在这两个分区的每一个周围被插入以确保翻译后的程序按正确的顺序来执行操作。
表4
Void function(){
      for(int tid_x＝0；tid_x＜dimblock.X；tid_x++){
          SharedMem[tid_x]＝...；//将值存入共享存储器中
      }
    for(int tid_x＝0；tid_x＜dimblock.X；tid_x++){
    ＝...(SharedMem[tid_x]+SharedMem[tid_x-1])/2.0；
     }
用于将程序翻译为由通用处理器来执行的更加简单的技术是给每个CTA维度插入明确的线程循环，这样就不必为在同一分区内的基准而决定维度的依赖性。例如，表5示出的程序被翻译为如表6所示的程序。注意，表5中插入的一个或多个线程循环可能不是必须的，这是因为不用决定维度的依赖性就可以生成程序。
表5
_global_void function(){
  Shared1＝...
  ＝Shared1
}
表6
void function(){
   for(int tid_x＝0；tid_x＜dimblock.X；tid_x++){
      for(int tid_y＝0；tid_y＜dimblock.Y；tid_y++){
            for(int tid_z＝0；tid_z＜dimblock.Z；tid_z++){
                  Shared1＝...
                  ＝Shared1
              }
         }
    }
图3C是根据本发明的一个实施例，说明翻译为优化后的代码360的输入代码333的概念图。将输入代码333配置为用于由多线程处理子系统112执行，并且包括由同步壁垒指令335进行分隔的代码序列334和338。CTA中的所有线程会在其中任一线程开始执行代码序列338之前完成代码序列334的执行。翻译器220将输入代码333划分以产生划分后的代码360，其中分区361包括由代码序列334代表的指令，分区362、364和365包括由代码序列338代表的指令。
分区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，，const int
    size)
{
(2)float sum；
(3)int matrix_start，col，row，out_index，i；
(4)martrix_start＝blockIdx.x*size*size；
(5)col＝matrix_start+threadIDx.x；
(6)row＝matrix_start+threadIdx.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+(threadIdx.y*size)+threadIdx.x；
(12)A_list[out_index]＝sum；
}
要注意的是表7中位于第(9)行的语句具有变量向量(x，y)，因为col(列)依赖于x维度且row(行)依赖于y维度。z维度从来不使用，所以没有循环被插入在z上进行迭代。通常的成本分析技术可以用来决定如在表7中示出的示例性内核中的语句5和6这样的情况。由于每一个仅依赖于一个线程身份维度，因此选择x和y索引循环中的任一嵌套顺序将迫使语句的冗余执行，或者是分区的主循环嵌套之外的冗余循环。
表8翻译后的CUDA内核
(1) _global_small_mm_list(float*A_list，float*B_list，，const int size)
{
(2)  float sum[]；
(3)  int matrix_start[]，col[]，row[]，out_index，i；
(4)      matrix_start[threadID]＝blockIDx.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[row[threadID]+i]*
               B_list[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是根据本发明的一个实施例，由通用处理器(比如CPU 202)执行翻译后的代码201的方法步骤的流程图。在步骤400中，编译器225将翻译后的代码201进行编译，选择性地执行针对CPU的优化，以生成编译后的代码。在步骤405中，在CPU 202中可用的执行内核400的数量由设备驱动器203决定。为了提高性能，翻译后的代码201可以自动地进行缩放以用于在可用执行内核上执行。在步骤410中，运行时间环境227或装置驱动器203将CPU 202配置为能够使多个执行内核执行翻译后的代码201。
运行时间环境227可以生成多个操作系统(OS)运行时间线程，该运行时间线程能够由环境变量控制。默认状态下，系统中内核的数量可以用作OS运行时间线程的数量。在步骤410中，可以对将要启动的CUDA线程的数量进行评估，并且在统计上划分为运行时间线程的数量。每一个运行时间线程顺序地执行一部分编译后的代码，并在壁垒上等待。当所有的运行时间线程都到达了壁垒，则CTA完成。在步骤415中，运行时间环境227或装置驱动器203启动编译后的代码用于由CPU 202执行。
翻译器220、编译器225以及运行时间环境227被用来将CUDA应用程序转换为由通用CPU执行的代码。CUDA编程模型支持批量的同步任务并行，其中每个任务由细粒度SPMD线程组成。CUDA编程模型的使用已经被限制于那些愿意为由GPU执行而编写特殊代码的程序员。这些特殊代码可以转换为由通用CPU来执行，并不需要程序员来重新编写CUDA应用程序。由CUDA支持的三个关键的抽象概念是SPMD线程块、壁垒同步和共享的存储器。翻译器220将遍及CUDA线程块的细粒度线程串行化为单一的CPU线程，并且执行优化变形以转换CUDA应用程序。
虽然前面所述的是本发明的实施例，在不背离本发明基本范围的前提下可以设计出本发明更多的实施例。例如，本发明的某些方面可以由硬件或软件来实现，或者是由硬件与软件结合在一起来实现。本发明的一个实施例可以实现为计算机系统所使用的程序产品。程序产品的程序对实施例的功能(包括在此描述的方法)进行定义，并且能够被包含在各种各样的计算机可读存储介质内。说明性的计算机可读存储介质包括但不限于：(i)信息在其上永久保存的非可写存储介质(例如，计算机内的只读存储装置，如可被CD-ROM驱动器读出的CD-ROM盘、闪存、ROM芯片或者任意类型的固态非易失性半导体存储器)；以及(ii)其上存储有可改变的信息的可写存储介质(例如，磁盘驱动器内的软盘或硬盘驱动器或任意类型的固态随机存取半导体存储器)。当携带有引导本发明的功能的计算机可读指令时，这样的计算机可读存储介质就是本发明的实施例。因此，本发明的范围由以下的权利要求来界定。

Claims (10)

1.一种配置用于执行翻译后的应用程序的计算系统，所述计算系统包括：
通用处理器，所述通用处理器被配置用于执行编译器；
系统存储器，所述系统存储器与所述处理器连接并且被配置用于存储翻译后的应用程序和编译后的代码；并且
所述编译器被配置用于：
接收翻译后的应用程序，所述翻译后的应用程序由采用并行编程模型编写的用于多核图像处理单元上执行的应用程序转换而来；以及
编译翻译后的应用程序，以生成用于由通用处理器执行的编译后的代码；
装置驱动器，所述装置驱动器被配置用于：
确定在通用处理器中可用来执行翻译后的应用程序的执行内核的数量；
将所述通用处理器配置用于启用该数量的执行内核；以及
运行时间环境，所述运行时间环境被配置为启动所述编译后的代码用于由包括所述数量的执行内核的所述通用处理器来执行。
2.根据权利要求1的计算系统，其特征在于，所述翻译后的应用程序包括在划分后的应用程序的第一区域周围的第一循环嵌套，以确保在合作线程阵列中的任一线程开始执行所述翻译后的应用程序的第二区域之前，所述合作线程阵列中的所有线程完成所述划分后的应用程序的第一区域的执行。
3.根据权利要求2的计算系统，其特征在于，所述第一循环在所述合作线程阵列的一个或多个维度上进行迭代。
4.根据权利要求1的计算系统，其特征在于，通过将所述应用程序划分成为同步独立指令的区域以生成划分后的应用程序，并且在划分后的应用程序的至少一个区域周围插入循环从而生成所述翻译后的应用程序，其中所述循环在合作线程阵列维度上进行迭代，所述合作线程阵列维度对应于由所述多核图像处理单元内的并行处理器同时执行的多个线程。
5.根据权利要求4的计算系统，其特征在于，所述划分后的应用程序的第一区域包括在同步壁垒指令之前的指令，所述划分后的应用程序的第二区域包括在同步壁垒指令之后的指令。
6.根据权利要求5的计算系统，其特征在于，在所述划分后的应用程序的至少一个区域周围插入额外的循环以生成所述翻译后的应用程序，其中所述额外的循环在不同的合作线程阵列维度上进行迭代。
7.根据权利要求4的计算系统，其特征在于，将所述划分后的应用程序进行分类，以便识别每个语句相对于所述合作线程阵列维度是收敛的还是发散的。
8.根据权利要求1的计算系统，其特征在于，所述通用处理器被进一步配置用于执行编译后的代码。
9.根据权利要求1的计算系统，其特征在于，所述通用处理器被进一步配置用于执行针对于所述通用处理器的优化。
10.根据权利要求1的计算系统，其特征在于，所述采用并行编程模型编写的用于在多核图像处理单元上执行的应用程序是CUDA(计算统一设备架构)应用程序。
CN200910117897.5A 2008-04-09 2009-04-09 用于执行翻译后的应用程序的计算系统 Active CN101556543B (zh)

Applications Claiming Priority (4)

Application Number Priority Date Filing Date Title
US4370808P 2008-04-09 2008-04-09
US61/043,708 2008-04-09
US12/408,559 2009-03-20
US12/408,559 US9448779B2 (en) 2008-04-09 2009-03-20 Execution of retargetted graphics processor accelerated code by a general purpose processor

Publications (2)

Publication Number Publication Date
CN101556543A true CN101556543A (zh) 2009-10-14
CN101556543B CN101556543B (zh) 2014-04-23

Family

ID=41164945

Family Applications (2)

Application Number Title Priority Date Filing Date
CN200910117897.5A Active CN101556543B (zh) 2008-04-09 2009-04-09 用于执行翻译后的应用程序的计算系统
CN200910117898XA Active CN101556544B (zh) 2008-04-09 2009-04-09 为了由通用处理器执行而对应用程序重定目标的计算系统

Family Applications After (1)

Application Number Title Priority Date Filing Date
CN200910117898XA Active CN101556544B (zh) 2008-04-09 2009-04-09 为了由通用处理器执行而对应用程序重定目标的计算系统

Country Status (4)

Country Link
US (5) US9678775B1 (zh)
JP (2) JP4984306B2 (zh)
CN (2) CN101556543B (zh)
TW (2) TWI437491B (zh)

Cited By (8)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN101963916A (zh) * 2010-09-28 2011-02-02 中国科学院地质与地球物理研究所 编译处理方法及装置
CN103959238A (zh) * 2011-11-30 2014-07-30 英特尔公司 使用gpu/cpu体系结构的rsa的高效实现
CN104106049A (zh) * 2012-02-16 2014-10-15 微软公司 计算着色器的栅格化
CN104641352A (zh) * 2012-08-02 2015-05-20 西门子公司 用于周期控制系统的流水线操作
CN108541321A (zh) * 2016-02-26 2018-09-14 谷歌有限责任公司 将程序代码映射到高性能、高功效的可编程图像处理硬件平台的编译技术
CN109844802A (zh) * 2016-09-02 2019-06-04 英特尔公司 用于在图形处理器中提高线程并行性的机制
CN110209509A (zh) * 2019-05-28 2019-09-06 北京星网锐捷网络技术有限公司 多核处理器间的数据同步方法及装置
CN110866861A (zh) * 2017-04-24 2020-03-06 英特尔公司 计算优化机制

Families Citing this family (68)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US8515052B2 (en) 2007-12-17 2013-08-20 Wai Wu Parallel signal processing system and method
US8572590B2 (en) * 2008-09-17 2013-10-29 Reservoir Labs, Inc. Methods and apparatus for joint parallelism and locality optimization in source code compilation
US8930926B2 (en) * 2008-02-08 2015-01-06 Reservoir Labs, Inc. System, methods and apparatus for program optimization for multi-threaded processor architectures
US8661422B2 (en) * 2008-02-08 2014-02-25 Reservoir Labs, Inc. Methods and apparatus for local memory compaction
US9858053B2 (en) 2008-02-08 2018-01-02 Reservoir Labs, Inc. Methods and apparatus for data transfer optimization
US8572595B1 (en) 2008-02-08 2013-10-29 Reservoir Labs, Inc. Methods and apparatus for aggressive scheduling in source code compilation
US9678775B1 (en) * 2008-04-09 2017-06-13 Nvidia Corporation Allocating memory for local variables of a multi-threaded program for execution in a single-threaded environment
US8866827B2 (en) * 2008-06-26 2014-10-21 Microsoft Corporation Bulk-synchronous graphics processing unit programming
US8755515B1 (en) 2008-09-29 2014-06-17 Wai Wu Parallel signal processing system and method
US8688619B1 (en) 2009-03-09 2014-04-01 Reservoir Labs Systems, methods and apparatus for distributed decision processing
US9185020B2 (en) * 2009-04-30 2015-11-10 Reservoir Labs, Inc. System, apparatus and methods to implement high-speed network analyzers
US8667474B2 (en) * 2009-06-19 2014-03-04 Microsoft Corporation Generation of parallel code representations
US8892483B1 (en) 2010-06-01 2014-11-18 Reservoir Labs, Inc. Systems and methods for planning a solution to a dynamically changing problem
US8584103B2 (en) 2010-06-17 2013-11-12 International Business Machines Corporation Reducing parallelism of computer source code
US8756590B2 (en) * 2010-06-22 2014-06-17 Microsoft Corporation Binding data parallel device source code
JP5059174B2 (ja) * 2010-08-10 2012-10-24 株式会社東芝 プログラム変換装置、およびそのプログラム
WO2012037735A1 (en) * 2010-09-26 2012-03-29 Mediatek Singapore Pte. Ltd. Method for performing video processing based upon a plurality of commands, and associated video processing circuit
US8914601B1 (en) 2010-10-18 2014-12-16 Reservoir Labs, Inc. Systems and methods for a fast interconnect table
US9134976B1 (en) 2010-12-13 2015-09-15 Reservoir Labs, Inc. Cross-format analysis of software systems
US20120159090A1 (en) * 2010-12-16 2012-06-21 Microsoft Corporation Scalable multimedia computer system architecture with qos guarantees
KR101738641B1 (ko) 2010-12-17 2017-05-23 삼성전자주식회사 멀티 코어 시스템의 프로그램 컴파일 장치 및 방법
US20120177119A1 (en) * 2011-01-07 2012-07-12 Sony Corporation Faster motion estimation in an avc software encoder using general purpose graphic process units (gpgpu)
WO2012157786A1 (ja) * 2011-05-19 2012-11-22 日本電気株式会社 並列処理装置、並列処理方法、最適化装置、最適化方法、および、コンピュータ・プログラム
US8966461B2 (en) * 2011-09-29 2015-02-24 Advanced Micro Devices, Inc. Vector width-aware synchronization-elision for vector processors
US9489180B1 (en) 2011-11-18 2016-11-08 Reservoir Labs, Inc. Methods and apparatus for joint scheduling and layout optimization to enable multi-level vectorization
WO2013092292A1 (en) 2011-12-21 2013-06-27 Dolby International Ab Audio encoder with parallel architecture
US9830133B1 (en) 2011-12-12 2017-11-28 Significs And Elements, Llc Methods and apparatus for automatic communication optimizations in a compiler based on a polyhedral representation
US9817668B2 (en) 2011-12-16 2017-11-14 Nvidia Corporation Batched replays of divergent operations
US9715413B2 (en) 2012-01-18 2017-07-25 Nvidia Corporation Execution state analysis for assigning tasks to streaming multiprocessors
US9069609B2 (en) 2012-01-18 2015-06-30 Nvidia Corporation Scheduling and execution of compute tasks
US9361079B2 (en) * 2012-01-30 2016-06-07 Nvidia Corporation Method for compiling a parallel thread execution program for general execution
US10152329B2 (en) * 2012-02-09 2018-12-11 Nvidia Corporation Pre-scheduled replays of divergent operations
US9329877B2 (en) 2012-03-18 2016-05-03 Microsoft Technology Licensing, Llc Static verification of parallel program code
US9798588B1 (en) 2012-04-25 2017-10-24 Significs And Elements, Llc Efficient packet forwarding using cyber-security aware policies
US9292265B2 (en) * 2012-05-09 2016-03-22 Nvidia Corporation Method for convergence analysis based on thread variance analysis
US10936569B1 (en) 2012-05-18 2021-03-02 Reservoir Labs, Inc. Efficient and scalable computations with sparse tensors
US9684865B1 (en) 2012-06-05 2017-06-20 Significs And Elements, Llc System and method for configuration of an ensemble solver
US9142005B2 (en) * 2012-08-20 2015-09-22 Nvidia Corporation Efficient placement of texture barrier instructions
KR101382945B1 (ko) 2012-10-17 2014-04-09 인하대학교 산학협력단 Ｇｐｕ 병렬 검색을 통한 ｒ－트리 계열의 계층형 색인 구조의 범위 질의 가속화 방법
US20140143755A1 (en) * 2012-11-20 2014-05-22 Nvidia Corporation System and method for inserting synchronization statements into a program file to mitigate race conditions
US9134979B2 (en) * 2013-02-06 2015-09-15 Nvidia Corporation Convergence analysis in multithreaded programs
US9015656B2 (en) * 2013-02-28 2015-04-21 Cray Inc. Mapping vector representations onto a predicated scalar multi-threaded system
US9442755B2 (en) * 2013-03-15 2016-09-13 Nvidia Corporation System and method for hardware scheduling of indexed barriers
KR102062208B1 (ko) * 2013-05-03 2020-02-11 삼성전자주식회사 멀티스레드 프로그램 코드의 변환 장치 및 방법
US9223550B1 (en) * 2013-10-17 2015-12-29 Google Inc. Portable handling of primitives for concurrent execution
US10326448B2 (en) 2013-11-15 2019-06-18 Scientific Concepts International Corporation Code partitioning for the array of devices
US9698791B2 (en) 2013-11-15 2017-07-04 Scientific Concepts International Corporation Programmable forwarding plane
US9294097B1 (en) 2013-11-15 2016-03-22 Scientific Concepts International Corporation Device array topology configuration and source code partitioning for device arrays
CN103716574A (zh) * 2013-12-13 2014-04-09 乐视致新电子科技（天津）有限公司 一种信号处理方法及数字电视
US9612811B2 (en) * 2014-01-21 2017-04-04 Nvidia Corporation Confluence analysis and loop fast-forwarding for improving SIMD execution efficiency
US9658823B2 (en) * 2014-02-25 2017-05-23 Nec Corporation Source-to-source compiler and run-time library to transparently accelerate stack or queue-based irregular applications on many-core architectures
US9400685B1 (en) * 2015-01-30 2016-07-26 Huawei Technologies Co., Ltd. Dividing, scheduling, and parallel processing compiled sub-tasks on an asynchronous multi-core processor
US11550632B2 (en) * 2015-12-24 2023-01-10 Intel Corporation Facilitating efficient communication and data processing across clusters of computing machines in heterogeneous computing environment
EP3336692B1 (en) * 2016-12-13 2020-04-29 Arm Ltd Replicate partition instruction
US10489204B2 (en) 2017-01-31 2019-11-26 Samsung Electronics Co., Ltd. Flexible in-order and out-of-order resource allocation
US10191745B2 (en) * 2017-03-31 2019-01-29 Intel Corporation Optimized call-return and binary translation
US11353868B2 (en) * 2017-04-24 2022-06-07 Intel Corporation Barriers and synchronization for machine learning at autonomous machines
US10474458B2 (en) * 2017-04-28 2019-11-12 Intel Corporation Instructions and logic to perform floating-point and integer operations for machine learning
US10528479B2 (en) * 2017-06-02 2020-01-07 Huawei Technologies Co., Ltd. Global variable migration via virtual memory overlay technique for multi-version asynchronous dynamic software update
KR101948135B1 (ko) 2017-07-31 2019-02-14 서울대학교산학협력단 이종 시스템의 애플리케이션 실행 방법 및 이를 수행하기 위한 이종 시스템
US10599482B2 (en) * 2017-08-24 2020-03-24 Google Llc Method for intra-subgraph optimization in tuple graph programs
US10831505B2 (en) 2018-09-29 2020-11-10 Intel Corporation Architecture and method for data parallel single program multiple data (SPMD) execution
US11093250B2 (en) * 2018-09-29 2021-08-17 Intel Corporation Apparatus and method for gang invariant operation optimizations using dynamic evaluation
US10776110B2 (en) 2018-09-29 2020-09-15 Intel Corporation Apparatus and method for adaptable and efficient lane-wise tensor processing
US10915328B2 (en) 2018-12-14 2021-02-09 Intel Corporation Apparatus and method for a high throughput parallel co-processor and interconnect with low offload latency
EP3812885A1 (de) * 2019-10-24 2021-04-28 dspace digital signal processing and control engineering GmbH Integrierte simulationscode- und produktionscodegenerierung
CN111078412B (zh) * 2019-12-12 2023-03-14 中山大学 一种通过api截获对gpu进行资源管理的方法
US11593157B2 (en) 2020-02-05 2023-02-28 Nec Corporation Full asynchronous execution queue for accelerator hardware

Citations (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US20070038987A1 (en) * 2005-08-10 2007-02-15 Moriyoshi Ohara Preprocessor to improve the performance of message-passing-based parallel programs on virtualized multi-core processors

Family Cites Families (50)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US5161216A (en) 1989-03-08 1992-11-03 Wisconsin Alumni Research Foundation Interprocedural slicing of computer programs using dependence graphs
US5048018A (en) 1989-06-29 1991-09-10 International Business Machines Corporation Debugging parallel programs by serialization
US5488713A (en) 1989-12-27 1996-01-30 Digital Equipment Corporation Computer simulation technique for predicting program performance
US5660176A (en) 1993-12-29 1997-08-26 First Opinion Corporation Computerized medical diagnostic and treatment advice system
US5860009A (en) 1994-04-28 1999-01-12 Kabushiki Kaisha Toshiba Programming method for concurrent programs and program supporting apparatus thereof
EP0729097A1 (en) * 1995-02-07 1996-08-28 Sun Microsystems, Inc. Method and apparatus for run-time memory access checking and memory leak detection of a multi-threaded program
US6029002A (en) * 1995-10-31 2000-02-22 Peritus Software Services, Inc. Method and apparatus for analyzing computer code using weakest precondition
US5867649A (en) * 1996-01-23 1999-02-02 Multitude Corporation Dance/multitude concurrent computation
US5903889A (en) * 1997-06-09 1999-05-11 Telaric, Inc. System and method for translating, collecting and archiving patient records
US5903899A (en) * 1997-04-23 1999-05-11 Sun Microsystems, Inc. System and method for assisting exact Garbage collection by segregating the contents of a stack into sub stacks
US6292822B1 (en) 1998-05-13 2001-09-18 Microsoft Corporation Dynamic load balancing among processors in a parallel computer
US6205414B1 (en) * 1998-10-02 2001-03-20 International Business Machines Corporation Methodology for emulation of multi-threaded processes in a single-threaded operating system
US6230313B1 (en) 1998-12-23 2001-05-08 Cray Inc. Parallelism performance analysis based on execution trace information
US7865883B1 (en) 1999-11-12 2011-01-04 Oracle America, Inc. Parallel and asynchronous debugger and debugging method for multi-threaded programs
US7590644B2 (en) 1999-12-21 2009-09-15 International Business Machine Corporation Method and apparatus of streaming data transformation using code generator and translator
US6550020B1 (en) 2000-01-10 2003-04-15 International Business Machines Corporation Method and system for dynamically configuring a central processing unit with multiple processing cores
AU2000276404A1 (en) 2000-09-30 2002-04-15 Intel Corporation (A Corporation Of Delaware) Method, apparatus, and system for building a compact model for large vocabulary continuous speech recognition (lvcsr) system
US7058561B1 (en) * 2000-11-02 2006-06-06 International Business Machines Corporation System, method and program product for optimising computer software by procedure cloning
JP3810631B2 (ja) 2000-11-28 2006-08-16 富士通株式会社 情報処理プログラムを記録した記録媒体
US6966055B2 (en) 2001-03-02 2005-11-15 International Business Machines Corporation Optimizing post-link code
US6904590B2 (en) 2001-05-25 2005-06-07 Microsoft Corporation Methods for enhancing program analysis
US7577942B2 (en) * 2001-07-26 2009-08-18 International Business Machines Corporation Efficient monitoring of program variables under debug
US20050022173A1 (en) 2003-05-30 2005-01-27 Codito Technologies Private Limited Method and system for allocation of special purpose computing resources in a multiprocessor system
KR100518584B1 (ko) 2003-07-12 2005-10-04 삼성전자주식회사 공유 라이브러리 시스템 및 상기 시스템 구축 방법
US7119808B2 (en) 2003-07-15 2006-10-10 Alienware Labs Corp. Multiple parallel processor computer graphics system
CN100498757C (zh) * 2003-07-25 2009-06-10 Rmi公司 高级处理器
US20050108695A1 (en) * 2003-11-14 2005-05-19 Long Li Apparatus and method for an automatic thread-partition compiler
FR2865047B1 (fr) * 2004-01-14 2006-04-07 Commissariat Energie Atomique Systeme de generation automatique de codes optimises
CN100498711C (zh) * 2004-07-20 2009-06-10 明星游戏株式会社 自动翻译程序和程序翻译服务器
US7493606B2 (en) 2004-08-03 2009-02-17 Université du Québec à Chicoutimi (UQAC) Method for compiling and executing a parallel program
US20060132492A1 (en) 2004-12-17 2006-06-22 Nvidia Corporation Graphics processor with integrated wireless circuits
WO2006087728A1 (en) 2005-02-15 2006-08-24 Codito Technologies System for creating parallel applications
US20070219771A1 (en) * 2005-12-01 2007-09-20 Verheyen Henry T Branching and Behavioral Partitioning for a VLIW Processor
US7788468B1 (en) * 2005-12-15 2010-08-31 Nvidia Corporation Synchronization of threads in a cooperative thread array
US7728841B1 (en) 2005-12-19 2010-06-01 Nvidia Corporation Coherent shader output for multiple targets
US7376808B2 (en) * 2006-01-31 2008-05-20 International Business Machines Corporation Method and system for predicting the performance benefits of mapping subsets of application data to multiple page sizes
TWI343020B (en) 2006-05-12 2011-06-01 Nvidia Corp Antialiasing using multiple display heads of a graphics processor
JP4784827B2 (ja) 2006-06-06 2011-10-05 学校法人早稲田大学 ヘテロジニアスマルチプロセッサ向けグローバルコンパイラ
US7814486B2 (en) * 2006-06-20 2010-10-12 Google Inc. Multi-thread runtime system
GB2443277B (en) 2006-10-24 2011-05-18 Advanced Risc Mach Ltd Performing diagnostics operations upon an asymmetric multiprocessor apparatus
US20080244592A1 (en) 2007-03-27 2008-10-02 Kabushiki Kaisha Toshiba Multitask processing device and method
US8341611B2 (en) 2007-04-11 2012-12-25 Apple Inc. Application interface on multiple processors
CN101657795B (zh) 2007-04-11 2013-10-23 苹果公司 多处理器上的数据并行计算
US8286196B2 (en) 2007-05-03 2012-10-09 Apple Inc. Parallel runtime execution on multiple processors
US8789031B2 (en) 2007-09-18 2014-07-22 Intel Corporation Software constructed strands for execution on a multi-core architecture
US8776030B2 (en) * 2008-04-09 2014-07-08 Nvidia Corporation Partitioning CUDA code for execution by a general purpose processor
US9678775B1 (en) * 2008-04-09 2017-06-13 Nvidia Corporation Allocating memory for local variables of a multi-threaded program for execution in a single-threaded environment
JP4381459B1 (ja) 2008-06-27 2009-12-09 株式会社東芝 情報処理装置、粒度調整方法およびプログラム
US8615770B1 (en) * 2008-08-29 2013-12-24 Nvidia Corporation System and method for dynamically spawning thread blocks within multi-threaded processing systems
US8959497B1 (en) * 2008-08-29 2015-02-17 Nvidia Corporation System and method for dynamically spawning thread blocks within multi-threaded processing systems

Patent Citations (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US20070038987A1 (en) * 2005-08-10 2007-02-15 Moriyoshi Ohara Preprocessor to improve the performance of message-passing-based parallel programs on virtualized multi-core processors

Non-Patent Citations (1)

* Cited by examiner, † Cited by third party
Title
JOHN A. STRATTON等: "MCUDA: An Efficient Implementation of CUDA Kernels for Multi-core CPUs", 《IMPACT TECHNICAL REPORT 08-01》 *

Cited By (13)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN101963916A (zh) * 2010-09-28 2011-02-02 中国科学院地质与地球物理研究所 编译处理方法及装置
CN103959238B (zh) * 2011-11-30 2017-06-09 英特尔公司 使用gpu/cpu体系结构的rsa的高效实现
CN103959238A (zh) * 2011-11-30 2014-07-30 英特尔公司 使用gpu/cpu体系结构的rsa的高效实现
CN104106049A (zh) * 2012-02-16 2014-10-15 微软公司 计算着色器的栅格化
US9529575B2 (en) 2012-02-16 2016-12-27 Microsoft Technology Licensing, Llc Rasterization of compute shaders
CN104641352B (zh) * 2012-08-02 2018-06-29 西门子公司 用于周期控制系统的流水线操作
CN104641352A (zh) * 2012-08-02 2015-05-20 西门子公司 用于周期控制系统的流水线操作
US10281892B2 (en) 2012-08-02 2019-05-07 Siemens Aktiengesellschaft Pipelining for cyclic control systems
CN108541321A (zh) * 2016-02-26 2018-09-14 谷歌有限责任公司 将程序代码映射到高性能、高功效的可编程图像处理硬件平台的编译技术
CN109844802A (zh) * 2016-09-02 2019-06-04 英特尔公司 用于在图形处理器中提高线程并行性的机制
CN109844802B (zh) * 2016-09-02 2024-03-26 英特尔公司 用于在图形处理器中提高线程并行性的机制
CN110866861A (zh) * 2017-04-24 2020-03-06 英特尔公司 计算优化机制
CN110209509A (zh) * 2019-05-28 2019-09-06 北京星网锐捷网络技术有限公司 多核处理器间的数据同步方法及装置

Also Published As

Publication number Publication date
CN101556544A (zh) 2009-10-14
JP4984306B2 (ja) 2012-07-25
CN101556543B (zh) 2014-04-23
US8984498B2 (en) 2015-03-17
CN101556544B (zh) 2013-09-18
US9448779B2 (en) 2016-09-20
US20090259832A1 (en) 2009-10-15
TWI437491B (zh) 2014-05-11
US20090259997A1 (en) 2009-10-15
US8612732B2 (en) 2013-12-17
JP2009259240A (ja) 2009-11-05
US8572588B2 (en) 2013-10-29
US20090259829A1 (en) 2009-10-15
TW200947302A (en) 2009-11-16
JP5152594B2 (ja) 2013-02-27
US9678775B1 (en) 2017-06-13
JP2009259241A (ja) 2009-11-05
TWI423133B (zh) 2014-01-11
TW201007572A (en) 2010-02-16
US20090259828A1 (en) 2009-10-15

Similar Documents

Publication Publication Date Title
CN101556544B (zh) 为了由通用处理器执行而对应用程序重定目标的计算系统
US8776030B2 (en) Partitioning CUDA code for execution by a general purpose processor
US9361079B2 (en) Method for compiling a parallel thread execution program for general execution
US8099584B2 (en) Methods for scalably exploiting parallelism in a parallel processing system
CN102696023B (zh) 用于访问并行存储器空间的统一寻址和指令
Nugteren et al. Introducing'Bones' a parallelizing source-to-source compiler based on algorithmic skeletons
CN103870246A (zh) 用于线程的simd执行的编译器控制区调度
CN103460188A (zh) 用于基于活跃分析的再具体化以减少寄存器不足并提高并行度的技术
CN103870242A (zh) 优化线程栈存储器的管理的系统、方法和计算机程序产品
US8615770B1 (en) System and method for dynamically spawning thread blocks within multi-threaded processing systems
US9645802B2 (en) Technique for grouping instructions into independent strands
CN103870309A (zh) 用于集群多级寄存器堆的寄存器分配
CN103793206A (zh) 基于工作队列的图形处理单元工作创建
US8959497B1 (en) System and method for dynamically spawning thread blocks within multi-threaded processing systems
KR101118321B1 (ko) 리타게팅된 그래픽 프로세서 가속 코드의 범용 프로세서에 의한 실행
KR101117430B1 (ko) 범용 프로세서에 의해 실행하기 위한 어플리케이션 프로그램의 리타게팅
Chang et al. A translation framework for automatic translation of annotated llvm ir into opencl kernel function
US11119743B1 (en) Build time optimization using thread object variables
US9542192B1 (en) Tokenized streams for concurrent execution between asymmetric multiprocessors
CN116802605A (zh) 用于根据触发条件执行指令的电路和方法
Hjort Blindell Synthesizing software from a ForSyDe model targeting GPGPUs
White Modeling and Optimization of Parallel Matrix-based Computations on GPU
Rocker Hardware-aware solvers for large, sparse linear systems
Beach Application acceleration: An investigation of automatic porting methods for application accelerators
Moazeni Parallel Algorithms for Medical Informatics on Data-Parallel Many-Core Processors

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