CN103809936A - 编译或运行时执行分叉-合并数据并行程序的系统和方法 - Google Patents

编译或运行时执行分叉-合并数据并行程序的系统和方法 Download PDF

Info

Publication number
CN103809936A
CN103809936A CN201310538671.9A CN201310538671A CN103809936A CN 103809936 A CN103809936 A CN 103809936A CN 201310538671 A CN201310538671 A CN 201310538671A CN 103809936 A CN103809936 A CN 103809936A
Authority
CN
China
Prior art keywords
thread
function
worker
parallel
operate
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.)
Pending
Application number
CN201310538671.9A
Other languages
English (en)
Inventor
林�源
高塔姆·查克拉瓦蒂
杰迪普·马拉蒂
权冠
阿米特·萨布尼
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 CN103809936A publication Critical patent/CN103809936A/zh
Pending legal-status Critical Current

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/46—Multiprogramming arrangements
    • G06F9/50—Allocation of resources, e.g. of the central processing unit [CPU]
    • G06F9/5005—Allocation of resources, e.g. of the central processing unit [CPU] to service a request
    • G06F9/5011—Allocation of resources, e.g. of the central processing unit [CPU] to service a request the resources being hardware resources other than CPUs, Servers and Terminals
    • G06F9/5016—Allocation of resources, e.g. of the central processing unit [CPU] to service a request the resources being hardware resources other than CPUs, Servers and Terminals the resource being the memory
    • 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/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/38—Concurrent instruction execution, e.g. pipeline or look ahead
    • 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
    • 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/08—Addressing or allocation; Relocation in hierarchically structured memory systems, e.g. virtual memory systems
    • G06F12/0802—Addressing of a memory level in which the access to the desired data or data block requires associative addressing means, e.g. caches
    • G06F12/0844—Multiple simultaneous or quasi-simultaneous cache accessing
    • G06F12/0853—Cache with multiport tag or data arrays
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F8/00—Arrangements for software engineering
    • G06F8/40—Transformation of program code
    • 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/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/30098—Register arrangements
    • G06F9/3012—Organisation of register space, e.g. banked or distributed register file
    • G06F9/30134—Register stacks; shift registers
    • 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/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/38—Concurrent instruction execution, e.g. pipeline or look ahead
    • G06F9/3836—Instruction issuing, e.g. dynamic instruction scheduling or out of order instruction execution
    • G06F9/3851—Instruction issuing, e.g. dynamic instruction scheduling or out of order instruction execution from multiple instruction streams, e.g. multistreaming
    • 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/30—Arrangements for executing machine instructions, e.g. instruction decode
    • G06F9/38—Concurrent instruction execution, e.g. pipeline or look ahead
    • G06F9/3885—Concurrent instruction execution, e.g. pipeline or look ahead using a plurality of independent parallel functional units
    • G06F9/3888—Concurrent instruction execution, e.g. pipeline or look ahead using a plurality of independent parallel functional units controlled by a single instruction for multiple threads [SIMT] in parallel
    • 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/52—Program synchronisation; Mutual exclusion, e.g. by means of semaphores
    • 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/52—Program synchronisation; Mutual exclusion, e.g. by means of semaphores
    • G06F9/522—Barrier synchronisation
    • G—PHYSICS
    • G06—COMPUTING; CALCULATING OR COUNTING
    • G06F—ELECTRIC DIGITAL DATA PROCESSING
    • G06F2212/00—Indexing scheme relating to accessing, addressing or allocation within memory systems or architectures
    • G06F2212/45—Caching of specific data in cache memory
    • G06F2212/451—Stack data

Landscapes

  • Engineering & Computer Science (AREA)
  • Theoretical Computer Science (AREA)
  • Software Systems (AREA)
  • General Engineering & Computer Science (AREA)
  • Physics & Mathematics (AREA)
  • General Physics & Mathematics (AREA)
  • Multimedia (AREA)
  • Devices For Executing Special Programs (AREA)
  • Advance Control (AREA)
  • Multi Processors (AREA)
  • Stored Programmes (AREA)
  • Information Retrieval, Db Structures And Fs Structures Therefor (AREA)

Abstract

用于采用函数调用来编译或运行时执行分叉-合并数据并行程序的系统和方法。在一个实施例中，系统包括：（1）分区器，其可操作以将组分区成主组和至少一个工作者组，以及（2）线程指定器，其与分区器相关联并可操作以指定来自主组的仅一个线程用于执行和该至少一个工作者组中的所有线程用于执行。

Description

编译或运行时执行分叉-合并数据并行程序的系统和方法
相关申请的交叉引用
本申请要求于2012年11月5日由Lin等人所提交的、序列号为61/722,661的、标题为“EXECUTING SEQUENTIAL CODE USING AGROUP OF THREADS”的美国临时申请以及于2012年12月21日由Lin等人所提交的、序列号为13/724,359的、标题为“SYSTEM AND METHODFOR COMPILING OR RNTME EXECUTING A FORK-JOIN DATAPARALLEL PROGRAM WITH FUNCTION CALLS ON ASINGLE-INSTRUCTION-MULTIPLE-THREAD PROCESSOR”的美国申请的优先权，在先申请与本申请共同受让，并在本文通过援引的方式对二者加以合并。
技术领域
本申请总地指向并行处理器，并且，更具体地，指向用于在单指令多线程（SIMT）处理器上采用函数调用来编译或运行时（runtime）执行分叉-合并（fork-join）数据并行程序的系统和方法。
背景技术
如相关领域技术人员意识到的，可并行地执行应用或程序以增加其性能。数据并行程序在不同数据上并发实行相同进程。任务并行程序在相同数据上并发实行不同进程。静态并行程序是具有可在其执行之前被确定并行度级别的程序。相反，由动态并行程序可达到的并行度仅可随着其执行而被确定。无论程序是数据或任务并行、或静态或动态并行，其可在管线中执行，这通常是用于图形程序的情况。
SIMT处理器尤其擅长执行数据并行程序。SIMT处理器中的控制单元创建执行的线程组并调度其用于执行，在执行期间组中的所有线程并发执行相同指令。在一个特定处理器中，每个组或“线程束（warp）”具有32个线程，与SIMD处理器中的32个执行管线或通道（lane）相对应。
分叉-合并数据并行程序采用单线程的主体（main）程序来开始。程序在该级的顺序阶段或区中。在主体程序的执行期间的某一点处，主体或“主（master）”线程遭遇并行阶段或区的序列。每个并行区具有独立数据集并可由多个线程并发执行。当并行区开始时每个并行区中的并发任务的数目被确定，并且在并行区期间不改变。当遭遇并行区时，主体线程将一队线程（称为工作者（worker）线程）分叉以并行地执行并行区。程序随后进入并行区。如果工作者线程遭遇新并行区，那么新并行区将被序列化，即将由遭遇的工作者线程本身执行并行区。主线程等待直到并行区结束为止。一退出并行区，工作者线程与主线程合并，其随后在程序进入顺序区的点处恢复主体程序的执行。
下文的表1阐述分叉-合并数据并行程序的示例。
Figure BDA0000408016220000021
表1–分叉-合并数据并行程序的示例
出于理解表1和本公开的剩余内容的目的，术语“foo”和“bar”是函数的任意名称。因此“foo”和“bar”可替代任何函数。
分叉-合并数据并行模型通常在并行编程中使用。例如，OpenMP标准采用该模型作为其基本线程执行模型。OpenACC标准针对组中称为“gang”的工作者线程使用该模型。
发明内容
一个方面提供用于采用函数调用来编译或运行时执行分叉-合并数据并行程序的系统。在一个实施例中，系统包括：（1）分区器，其可操作以将线程束分区成主线程束和至少一个工作者线程束，以及（2）线程指定器，其与分区器相关联并可操作以指定来自主线程束的仅一个线程用于执行和该至少一个工作者线程束中的所有线程用于执行。
另一方面提供采用函数调用来编译或运行时执行分叉-合并数据并行程序的方法。在一个实施例中，方法包括：（1）将线程束分区成主线程束和至少一个工作者线程束，以及（2）指定来自主线程束的仅一个线程用于执行和该至少一个工作者线程束中的所有线程用于执行。
附图说明
现在结合附图对下面的描述加以参考，其中：
图1是SIMT处理器的框图，该SIMT处理器可操作以包含或实行用于采用函数调用来编译或运行时执行分叉-合并数据并行程序；
图2是用于采用函数调用来编译或运行时执行分叉-合并数据并行程序的系统的一个实施例的框图；以及
图3是采用函数调用来编译或运行时执行分叉-合并数据并行程序的方法的一个实施例的流程图。
具体实施方式
在SIMT处理器中，执行的多个线程被隔离成组。组中的所有线程同时执行相同指令。在可从市场上买到的来自加利福尼亚州圣塔克拉拉市的Nvidia公司的图形处理单元（GPU）中，其为一个类型的SIMT处理器，组被称为“线程束”，其按块执行。
SIMT处理器的管线控制单元创建、管理、调度、执行并提供机制以将组同步。Nvidia GPU提供bar.sync指令以将组同步。Nvidia GPU还支持由组进行的“发散”条件分支的执行；组中的一些线程需要采取分支（因为分支条件断言评估为“真”），并且其他线程需要落到下一指令（因为分支条件断言评估为“伪”）。管线控制单元保持跟踪组中的活动线程。其首先执行路径中的一个（所采取的分支或未被采取的分支）并且随后其他路径；在每个路径上使能适当的线程。
本文认识到的是，虽然GPU线程块内的所有线程在相同程序地址开始，但管线控制单元将受益于软件机制来将线程分区并调度成主体线程和工作者线程使得其可实施在分叉-合并模型中。
进一步认识到的是，软件机制的某些实施例应按组管理和同步线程，因为管线控制单元按组管理线程。
又进一步认识到的是，由于主体程序在分叉-合并模型中被单线程地执行，因此软件机制的某些实施例应达到顺序区语义而不引入副作用。导致副作用的指令的实施例是采用共享资源的那些指令，诸如共享存储器读或写或者可唤起共享异常处置器的任何代码操作（例如划分）。
还进一步认识到的是，软件机制的某些实施例应支持可在顺序区内和并行区内被调用的函数。这类函数可包含并行构造（parallel construct）本身。仍进一步认识到的是，软件机制的某些实施例应支持可分叉-合并并行区的函数调用。
仍进一步认识到的是，软件机制的某些实施例应支持外来函数，即并非由同一编译器编译为程序的函数。例如，现有GPU数学库中的数学函数，以及系统函数如malloc、free和print。在某些实施例中，顺序区中的主线程和并行区中的工作者线程二者应均能够调用外来函数。
因此，本文所描述的是用于在诸如GPU的SIMT处理器上采用函数调用来编译和执行分叉-合并数据并行程序的系统和方法的各种实施例。
在描述系统和方法的某些实施例之前，将描述SIMT处理器，其可操作以包含或实行用于采用函数调用来编译或运行时执行分叉-合并数据并行程序的系统或方法。
图1是SIMT处理器100的框图。SIMT处理器100包括被组织成线程组104或“线程束”的多个线程处理器或核心106。SIMT处理器100包含J个线程组104-1到104-J，每组具有K个核心106-1到106-K。在某些实施例中，线程组104-1到104-J可进一步被组织成一个或多个线程块102。某些实施例包括每线程组104三十二个核心106。其他实施例可包括少如每线程组中四个核心或多如数万核心。某些实施例将核心106组织成单线程组104，而其他实施例可具有数百或甚至数千个线程组104。SIMT处理器100的替代实施例可将核心106仅组织成线程组104，省略线程块组织级别。
SIMT处理器100进一步包括管线控制单元108、块共享存储器110和与线程组104-1到104-J相关联的本地存储器112-1到112-J的阵列。管线控制单元108通过数据总线114将任务分布到各个线程组104-1到104-J。线程组106-j内的核心106相互并行地执行。线程组104-1到104-J通过存储器总线116与块共享存储器110进行通信。线程组104-1到104J通过本地总线118-1到118-J分别与本地存储器112-1到112-J进行通信。例如线程组104-J以通过总线118-J进行通信来利用本地存储器112-J。SIMT处理器100的某些实施例将块共享存储器110的共享部分分配到每个线程块102，并允许由线程块102内的所有线程组104访问块共享存储器110的共享部分。某些实施例包括仅使用本地存储器112的线程组104。许多其他实施例包括平衡本地存储器112和块共享存储器110的使用的线程组104。
图1的实施例包括主线程组104-1。其余线程组104-2到104-J中的每一个被视为“工作者”线程组。主线程组104-1包括许多核心，其中的一个是主核心106-1，该主核心106-1最终执行主线程。在SIMT处理器110上所执行的程序被构造为内核的序列。典型地，每个内核在下一内核开始之前完成执行。在某些实施例中，SIMT处理器100可并行执行多个内核，这取决于内核的大小。每个内核被组织为要在核心106上所执行的线程的层级。
已描述其内可包含或实行本文所引入的系统或方法的SIMT处理器，将描述系统和方法的各种实施例。
本文所引入的系统的一个实施例包括编译器和设备运行时库。设备运行时库实现线程和组管理功能性。编译器将分叉-合并数据并行程序转译成主体线程程序和所括出的函数的集合，每个函数与并行构造相对应。经转译的代码对设备运行时库中的函数进行调用以实施线程和组管理。
下文的表2示出示范性程序以例示编译器转译和设备运行时实现方案。
Figure BDA0000408016220000061
表2–用于编译器转译和设备运行时实现方案的示范性程序
表2的main()程序的流程采用单主线程开始。主线程调用具有对编译器可见并由该编译器所编译的体的函数foo()。主线程随后调用函数ext()，其是具有对该编译器不可见的体的外部或外来函数。对外来函数的调用按原状被转译，不用由编译器进行任何特殊处置。主线程随后遭遇第一并行区。工作者线程将执行并行区而主线程等待其完成。在并行区内，每个工作者线程调用函数foo()和bar()。函数bar()包含另一并行区；然而，bar()已经在并行区之中。因为bar()已经在并行区之中，所以bar()之中的并行区将由每个工作者线程顺序地执行。
在第一并行区之后，主线程遭遇第二并行区。在第二并行区内，每个工作者线程调用外部外来函数ext()。在第二并行区之后，主线程调用函数bar()。在bar()内，主线程遭遇将再次由工作者线程执行的第三并行区。
函数Main()被称为入口函数，因为其是程序开始之处。诸如foo()和bar()的函数是非入口函数。
针对入口函数，编译器首先制作命名为main_core()的克隆拷贝。克隆拷贝随后如下文所描述被作为非入口函数处理。针对main()函数，编译器生成类似于下文表3中所示的代码，其中groupID()返回线程组的ID，所述线程组包含执行语句的线程。threadID()返回线程的ID。init()、
signal_done()和scheduler()是设备运行时库中的函数。
Figure BDA0000408016220000071
表3–示范性编译器生成的代码
当GPU线程块开始时，块内的所有线程执行main()；然而，其采取不同路径。线程0是主线程并执行init()、main_core()和singal_done()。组0内的其他线程直奔main()函数的结束并在那里等待。其余组中的线程执行scheduler()。
对于非入口函数，比如foo()、bar()和main_core()，编译器就如没有并行构造存在那样来转译代码。如果非入口函数包含并行构造，那么针对每个并行构造，编译器创建包含并行构造的体的函数（所括出的函数），并随后创建条件分支，所述条件分支检查执行线程是否是主线程。在伪分支中，编译器插入执行循环的代码。在真分支中，编译器插入对设备运行时库的调用以指派任务、唤醒工作者线程并执行屏障。当非入口函数在并行区之外被调用时条件为真。当非入口函数在并行区之内被调用时条件为假，在该情况下并行循环由执行线程顺序地执行。
例如，下文表4中示出用于函数bar()的经转译的代码。
Figure BDA0000408016220000081
表4-用于函数bar()的经转译的代码
signal_task()和barrier()是设备运行时库中的函数。bar_par_frunc()是与原始函数bar()中的并行构造相对应的所括出的函数。
此外，在该实施例中，设备运行时库包括下面的函数：init()、sheduler()、signal_task()、signal_done()和barrier()。库还实现以下函数用于内部使用：signal()、wait()和fetch_task()。
所有工作者线程执行scheduler()函数。工作者线程进行睡眠-唤醒-执行的循环直到被指示退出为止。
Figure BDA0000408016220000082
表5–使用程序退出标记的示范性代码
布尔变量‘exit_flag’被放入块共享存储器中并可由线程块内的所有线程访问。其由主线程用来以向工作者线程传达其是否应全部退出执行。‘exit_flag’在init()函数中设为伪，并且在signal_done()函数中设为真。两个函数均由主线程所调用。
表6–用于改变程序退出标记级的示范性代码
另一个块共享存储器用来传达当前任务。由主线程在signal_task()函数中设置当前任务，并且由工作者线程在fetch_task()函数中加以获取。块共享存储器包含指向与并行构造相对应的所括出的函数的指针。
Figure BDA0000408016220000093
表7–用于使用指针标识当前任务的示范性代码
因为并行区在线程块内被按顺序执行，所以在任何时刻仅一个任务是活动的。如果并行区可被异步地执行，那么典型地需要较复杂的数据结构诸如堆栈、队列或树来存储活动任务。
使用硬件屏障来实现barrier()、signal()和wait()函数。
Figure BDA0000408016220000101
表8–示范性barrier()、signal()和wait()函数
图2是用于采用函数调用来编译或运行时执行分叉-合并数据并行程序210的系统200的一个实施例的框图。程序210包括入口函数212、非入口函数214和外来函数216。系统200包括分区器202、线程指定器204、线程调度器206、函数处理器208、设备运行时库218和图1的SIMT处理器100。
SIMT处理器100包括图1的管线控制单元108、数据总线114、本地总线118-1和118-2以及共享存储器110。在图2的实施例中，SIMT处理器100示出为具有单线程块，其包含两个线程组：主线程组104-1和工作者线程组104-2。线程组104-1和104-2中的每一个包含线程106。
分区器202指定线程组104-1作为主线程组并且指定其余线程组作为工作者线程组。在图2的实施例中，示出单工作者线程组104-2。在替代实施例中，可采用许多工作者线程组。线程指定器204指定主线程组104-1的主线程106-1。主线程组104-1中的所有其他线程实际上是空闲的。线程指定器204还指定工作者线程组104-2中的线程106中的每一个作为工作者线程。
线程调度器206转译程序210使得管线控制单元108适当地控制主线程106-1和工作者线程组104-2中的各工作者线程的执行。线程调度器206转译程序210使得当主线程执行开始时，程序退出标记被禁止。线程调度器206调度主线程106-1以执行，直到到达程序210的并行区或结束为止。当到达程序210的并行区时，线程调度器206设置并行任务，并且工作者线程组104-2中的工作者线程开始执行。线程调度器206还针对工作者线程中的每一个设置屏障，使得当进入屏障时主线程106-1的执行恢复。当到达程序210的结束时，程序退出标记被使能，这使所有工作者线程停止执行。
函数处理器208对程序210的函数进行操作。处理入口函数212包括创建随后被作为非入口函数所处理的入口函数的克隆拷贝。处理原始入口函数使得主线程106-1将在其他调用之中执行克隆拷贝，并且工作者线程将执行循环，该循环是睡眠、唤醒、获取和执行由线程调度器206所设置的并行任务。
函数处理器208以两种方式转译非入口函数214。如果在非入口函数中不存在并行构造，那么简单地按原状处理函数。当存在并行构造时，创建包含并行构造的体的所括出的函数。函数处理器208随后创建分支条件，该分支条件将顺序地执行并行构造或采用设备运行时库218以指派任务、唤醒工作者线程和执行屏障，如上文所述。唤醒和睡眠功能性通过使用运行时设备库218的硬件屏障功能来实现。在屏障处的线程不被调度用于由硬件执行，所以其不浪费执行循环。在主线程组104-1内，仅主线程106-1参与屏障。之所以如此是因为硬件屏障是基于组的。如果组内的任何线程在屏障处，那么组视为在屏障处。
类似于处理不具有并行构造的非入口函数，外来函数216由函数处理器208按原状处理。
分叉-合并数据并行程序被分区成主程序和并行任务的集合。主程序是将由主体线程所执行的程序。并行任务与由工作者线程所执行的并行区相对应。主程序包含调度在该处主线程将指派并行任务的点、唤醒工作者线程和等待工作者线程完成。
专用主组中的专用主线程将正在执行程序的顺序区。
可替代地，可在组中的所有线程正在执行代码的同时在顺序区中仿真单线程行为。然而，仿真方案具有性能和工程复杂性这二者的限制，使其用处较小。必要的断言和同步引入执行开销。此外，可能从顺序区和并行区二者所调用的所有函数需要被不同地克隆和编译。
给定线程和组分区，工作者线程和主线程假定下面的寿命循环：
工作者线程的一个实施例在寿命循环中经历以下级：
1）线程块开始；
2）睡眠直到由主线程所唤醒为止；
3）如果程序退出标记设为真则退出；
4）获取并执行由主线程所指派的任务；
5）进入屏障；以及
6）回到级2。
主线程的一个实施例在寿命循环中经历以下级：
1）线程块开始；
2）将程序退出标记设为伪；
3）执行主程序直到到达并行区或到达主程序的结束为止；
4）在并行区的开始处：
a.设置并行任务，
b.唤醒工作者线程，
c.进入屏障，以及
d.恢复主程序（级3）；以及
5）在主程序的结束处：
a.将程序退出标记设为真，
b.唤醒工作者线程，以及
c.退出。
主组中的其他线程实质上在程序的结束处等待，其处于空闲。由主线程和工作者线程交错地执行程序。这产生良好的指令高速缓存足迹（foot-print），其比由主线程和工作者线程二者均活动并执行不同代码路径的方法所产生的足迹更好。
图3是采用函数调用来编译或运行时执行分叉-合并数据并行程序的方法的一个实施例的流程图。方法开始于开始步骤310。在步骤320，将线程块内的线程组分区成主线程组和至少一个工作者线程组。在步骤330，指定来自主线程组的一个线程作为主线程。主组的其余线程实质上在执行的从始至终是空闲的。并且在步骤330，指定该至少一个工作者线程组中的所有线程作为工作者线程。方法结束于结束步骤340。
本申请相关领域技术人员应理解的是，可对所描述的实施例做出其他和进一步的添加、删除、替换和修改。

Claims (10)

1.一种用于采用函数调用来编译或运行时执行分叉-合并数据并行程序的系统，包括：
分区器，其可操作以将组分区成主组和至少一个工作者组；以及
线程指定器，其与所述分区器相关联并可操作以指定来自所述主组的仅一个线程用于作为主线程执行和所述至少一个工作者组中的所有线程用于作为工作者线程执行。
2.根据权利要求1所述的系统，进一步包括线程调度器，其与所述线程指定器相关联并可操作以使单指令多线程处理器的管线控制单元调度所述主线程的执行如下：
当所述主线程开始执行时将程序退出标记设为第一状态；
所述主线程执行直到到达所述程序的并行区或结束为止；
一到达所述并行区，则设置并行任务、进入第一屏障、进入第二屏障并且恢复所述主线程的执行；以及
一到达所述结束，则将所述程序退出标记设为第二状态、进入屏障并且所述主线程停止执行。
3.根据权利要求1所述的系统，进一步包括线程调度器，其与所述线程指定器相关联并可操作以使单指令多线程处理器的管线控制单元调度所述工作者线程的执行如下：
所述工作者线程进入第一屏障；
当所述主线程到达并行区并且所述主线程进入屏障时所述工作者线程开始由所述主线程所设置的并行任务的执行；
所述并行任务一完成，则所述工作者线程进入第二屏障；以及
如果将程序退出标记设为第二状态并且所述主线程进入屏障那么所述工作者线程退出。
4.根据权利要求1所述的系统，进一步包括线程调度器，其与所述线程指定器相关联并可操作以采用单指令多线程处理器的管线控制单元的屏障函数以控制所述工作者线程的执行和停止。
5.根据权利要求1所述的系统，进一步包括函数处理器，其与所述线程指定器相关联并可操作以创建入口函数的克隆拷贝，并且将所述入口函数作为非入口函数处理。
6.根据权利要求1所述的系统，进一步包括函数处理器，其与所述线程指定器相关联并可操作以通过以下各项转译具有并行构造的非入口函数：
创建包含所述并行构造的体的函数；以及
如果所述函数正在所述主线程内执行，那么插入对设备运行时库的调用。
7.根据权利要求1所述的系统，进一步包括函数处理器，其与所述线程指定器相关联并可操作以按原状转译对外来函数的调用。
8.根据权利要求1所述的系统，进一步包括函数处理器，其与所述线程指定器相关联并可操作以采用设备运行时库，所述设备运行时库提供能从经编译的用户代码和内部函数所调用的函数。
9.根据权利要求1所述的系统，进一步包括函数处理器，其与所述线程指定器相关联并可操作以采用退出标记，所述退出标记存储在共享存储器中并被采用以允许所述主线程向所述工作者线程传达何时所述工作者线程应停止执行。
10.根据权利要求1所述的系统，其中所述系统可操作以配置单指令多线程处理器的共享存储器以标识当前任务。
CN201310538671.9A 2012-11-05 2013-11-04 编译或运行时执行分叉-合并数据并行程序的系统和方法 Pending CN103809936A (zh)

Applications Claiming Priority (4)

Application Number Priority Date Filing Date Title
US201261722661P 2012-11-05 2012-11-05
US61/722,661 2012-11-05
US13/724,359 2012-12-21
US13/724,359 US9747107B2 (en) 2012-11-05 2012-12-21 System and method for compiling or runtime executing a fork-join data parallel program with function calls on a single-instruction-multiple-thread processor

Publications (1)

Publication Number Publication Date
CN103809936A true CN103809936A (zh) 2014-05-21

Family

ID=50623483

Family Applications (4)

Application Number Title Priority Date Filing Date
CN201310538671.9A Pending CN103809936A (zh) 2012-11-05 2013-11-04 编译或运行时执行分叉-合并数据并行程序的系统和方法
CN201310538631.4A Active CN103809964B (zh) 2012-11-05 2013-11-04 用线程组执行顺序代码的系统和方法和包含其的simt处理器
CN201310538507.8A Pending CN103809963A (zh) 2012-11-05 2013-11-04 转译程序函数以正确处置局部作用域变量的系统和方法
CN201310538409.4A Pending CN103885751A (zh) 2012-11-05 2013-11-04 将区别属性的存储器分配给共享数据对象的系统和方法

Family Applications After (3)

Application Number Title Priority Date Filing Date
CN201310538631.4A Active CN103809964B (zh) 2012-11-05 2013-11-04 用线程组执行顺序代码的系统和方法和包含其的simt处理器
CN201310538507.8A Pending CN103809963A (zh) 2012-11-05 2013-11-04 转译程序函数以正确处置局部作用域变量的系统和方法
CN201310538409.4A Pending CN103885751A (zh) 2012-11-05 2013-11-04 将区别属性的存储器分配给共享数据对象的系统和方法

Country Status (3)

Country Link
US (4) US9436475B2 (zh)
CN (4) CN103809936A (zh)
TW (4) TWI510919B (zh)

Cited By (3)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN112114877A (zh) * 2020-09-28 2020-12-22 西安芯瞳半导体技术有限公司 一种动态补偿线程束warp的方法、处理器及计算机存储介质
CN117009054A (zh) * 2023-07-27 2023-11-07 北京登临科技有限公司 一种simt装置、线程组动态构建方法及处理器
CN117009054B (zh) * 2023-07-27 2024-06-28 北京登临科技有限公司 一种simt装置、线程组动态构建方法及处理器

Families Citing this family (29)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US9436475B2 (en) 2012-11-05 2016-09-06 Nvidia Corporation System and method for executing sequential code using a group of threads and single-instruction, multiple-thread processor incorporating the same
US9519668B2 (en) * 2013-05-06 2016-12-13 International Business Machines Corporation Lock-free creation of hash tables in parallel
US9250877B2 (en) * 2013-09-20 2016-02-02 Cray Inc. Assisting parallelization of a computer program
US9207979B1 (en) * 2014-05-28 2015-12-08 Freescale Semiconductor, Inc. Explicit barrier scheduling mechanism for pipelining of stream processing algorithms
US10061591B2 (en) 2014-06-27 2018-08-28 Samsung Electronics Company, Ltd. Redundancy elimination in single instruction multiple data/thread (SIMD/T) execution processing
US10061592B2 (en) 2014-06-27 2018-08-28 Samsung Electronics Co., Ltd. Architecture and execution for efficient mixed precision computations in single instruction multiple data/thread (SIMD/T) devices
US9804883B2 (en) * 2014-11-14 2017-10-31 Advanced Micro Devices, Inc. Remote scoped synchronization for work stealing and sharing
US9886317B2 (en) * 2015-02-02 2018-02-06 Oracle International Corporation Fine-grained scheduling of work in runtime systems
US10318307B2 (en) 2015-06-17 2019-06-11 Mediatek, Inc. Scalarization of vector processing
US9594512B1 (en) 2015-06-19 2017-03-14 Pure Storage, Inc. Attributing consumed storage capacity among entities storing data in a storage array
GB2539958B (en) * 2015-07-03 2019-09-25 Advanced Risc Mach Ltd Data processing systems
GB2540937B (en) * 2015-07-30 2019-04-03 Advanced Risc Mach Ltd Graphics processing systems
US10996989B2 (en) * 2016-06-13 2021-05-04 International Business Machines Corporation Flexible optimized data handling in systems with multiple memories
US10255132B2 (en) 2016-06-22 2019-04-09 Advanced Micro Devices, Inc. System and method for protecting GPU memory instructions against faults
US10310861B2 (en) 2017-04-01 2019-06-04 Intel Corporation Mechanism for scheduling threads on a multiprocessor
GB2560059B (en) 2017-06-16 2019-03-06 Imagination Tech Ltd Scheduling tasks
TWI647619B (zh) * 2017-08-29 2019-01-11 智微科技股份有限公司 用來於一電子裝置中進行硬體資源管理之方法以及對應的電子裝置
CN109471673B (zh) * 2017-09-07 2022-02-01 智微科技股份有限公司 用来于电子装置中进行硬件资源管理的方法及电子装置
US10990394B2 (en) * 2017-09-28 2021-04-27 Intel Corporation Systems and methods for mixed instruction multiple data (xIMD) computing
US11068247B2 (en) 2018-02-06 2021-07-20 Microsoft Technology Licensing, Llc Vectorizing conditional min-max sequence reduction loops
US10606595B2 (en) * 2018-03-23 2020-03-31 Arm Limited Data processing systems
CN108804311B (zh) * 2018-05-07 2022-06-03 微梦创科网络科技（中国）有限公司 一种执行测试文件的方法及装置
US11061742B2 (en) * 2018-06-27 2021-07-13 Intel Corporation System, apparatus and method for barrier synchronization in a multi-threaded processor
CN110032407B (zh) * 2019-03-08 2020-12-22 创新先进技术有限公司 提升cpu并行性能的方法及装置和电子设备
CN110990151A (zh) * 2019-11-24 2020-04-10 浪潮电子信息产业股份有限公司 一种基于异构计算平台的业务处理方法
US11861403B2 (en) 2020-10-15 2024-01-02 Nxp Usa, Inc. Method and system for accelerator thread management
CN112214243B (zh) * 2020-10-21 2022-05-27 上海壁仞智能科技有限公司 配置向量运算系统中的协作线程束的装置和方法
CN112579164B (zh) * 2020-12-05 2022-10-25 西安翔腾微电子科技有限公司 一种simt条件分支处理装置及方法
US11361400B1 (en) 2021-05-06 2022-06-14 Arm Limited Full tile primitives in tile-based graphics processing

Citations (4)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN101030152A (zh) * 2007-03-20 2007-09-05 华为技术有限公司 基于伪同步方式的操作控制方法及装置
US20090013323A1 (en) * 2007-07-06 2009-01-08 Xmos Limited Synchronisation
CN102135916A (zh) * 2010-10-15 2011-07-27 威盛电子股份有限公司 同步方法以及图形处理系统
US20120191947A1 (en) * 2011-01-20 2012-07-26 International Business Machine Corporation Computer operation control method, program and system

Family Cites Families (79)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
US5875464A (en) * 1991-12-10 1999-02-23 International Business Machines Corporation Computer system with private and shared partitions in cache
US6058465A (en) * 1996-08-19 2000-05-02 Nguyen; Le Trong Single-instruction-multiple-data processing in a multimedia signal processor
JP3849951B2 (ja) * 1997-02-27 2006-11-22 株式会社日立製作所 主記憶共有型マルチプロセッサ
US6009517A (en) * 1997-10-06 1999-12-28 Sun Microsystems, Inc. Mixed execution stack and exception handling
GB9825102D0 (en) 1998-11-16 1999-01-13 Insignia Solutions Plc Computer system
US6389449B1 (en) * 1998-12-16 2002-05-14 Clearwater Networks, Inc. Interstream control and communications for multi-streaming digital processors
US7020879B1 (en) * 1998-12-16 2006-03-28 Mips Technologies, Inc. Interrupt and exception handling for multi-streaming digital processors
US6574725B1 (en) * 1999-11-01 2003-06-03 Advanced Micro Devices, Inc. Method and mechanism for speculatively executing threads of instructions
US6609193B1 (en) 1999-12-30 2003-08-19 Intel Corporation Method and apparatus for multi-thread pipelined instruction decoder
JP3632635B2 (ja) 2001-07-18 2005-03-23 日本電気株式会社 マルチスレッド実行方法及び並列プロセッサシステム
TWI245221B (en) * 2002-08-22 2005-12-11 Ip First Llc Apparatus and method for selective memory attribute control
US7159211B2 (en) * 2002-08-29 2007-01-02 Indian Institute Of Information Technology Method for executing a sequential program in parallel with automatic fault tolerance
US7086063B1 (en) * 2003-03-25 2006-08-01 Electric Cloud, Inc. System and method for file caching in a distributed program build environment
US8683132B1 (en) 2003-09-29 2014-03-25 Nvidia Corporation Memory controller for sequentially prefetching data for a processor of a computer system
US7464370B2 (en) 2003-12-04 2008-12-09 International Business Machines Corporation Creating a method from a block of code
US7380086B2 (en) * 2003-12-12 2008-05-27 International Business Machines Corporation Scalable runtime system for global address space languages on shared and distributed memory machines
CN100461090C (zh) * 2004-03-31 2009-02-11 英特尔公司 利用代码共享进行堆栈高速缓存的系统、方法和设备
US7822779B2 (en) 2004-04-23 2010-10-26 Wal-Mart Stores, Inc. Method and apparatus for scalable transport processing fulfillment system
US20060095675A1 (en) 2004-08-23 2006-05-04 Rongzhen Yang Three stage hybrid stack model
US7469318B2 (en) 2005-02-10 2008-12-23 International Business Machines Corporation System bus structure for large L2 cache array topology with different latency domains
US7574702B2 (en) * 2005-03-18 2009-08-11 Microsoft Corporation Method and apparatus for hybrid stack walking
US8516483B2 (en) 2005-05-13 2013-08-20 Intel Corporation Transparent support for operating system services for a sequestered sequencer
US8397013B1 (en) * 2006-10-05 2013-03-12 Google Inc. Hybrid memory module
KR101257848B1 (ko) * 2005-07-13 2013-04-24 삼성전자주식회사 복합 메모리를 구비하는 데이터 저장 시스템 및 그 동작방법
US20070136523A1 (en) * 2005-12-08 2007-06-14 Bonella Randy M Advanced dynamic disk memory module special operations
US20070143582A1 (en) 2005-12-16 2007-06-21 Nvidia Corporation System and method for grouping execution threads
US7478190B2 (en) 2006-02-10 2009-01-13 University Of Utah Technology Commercialization Office Microarchitectural wire management for performance and power in partitioned architectures
JP4900784B2 (ja) 2006-04-13 2012-03-21 株式会社日立製作所 ストレージシステム及びストレージシステムのデータ移行方法
US8108844B2 (en) 2006-06-20 2012-01-31 Google Inc. Systems and methods for dynamically choosing a processing element for a compute kernel
US8606998B2 (en) 2006-08-24 2013-12-10 Advanced Micro Devices, Inc. System and method for instruction-based cache allocation policies
US7584335B2 (en) 2006-11-02 2009-09-01 International Business Machines Corporation Methods and arrangements for hybrid data storage
US20080109795A1 (en) * 2006-11-02 2008-05-08 Nvidia Corporation C/c++ language extensions for general-purpose graphics processing unit
US7593263B2 (en) * 2006-12-17 2009-09-22 Anobit Technologies Ltd. Memory device with reduced reading latency
US8250556B1 (en) 2007-02-07 2012-08-21 Tilera Corporation Distributing parallelism for parallel processing architectures
US8095782B1 (en) * 2007-04-05 2012-01-10 Nvidia Corporation Multiple simultaneous context architecture for rebalancing contexts on multithreaded processing cores upon a context change
US8286196B2 (en) * 2007-05-03 2012-10-09 Apple Inc. Parallel runtime execution on multiple processors
US7856541B2 (en) 2007-04-18 2010-12-21 Hitachi, Ltd. Latency aligned volume provisioning methods for interconnected multiple storage controller configuration
KR101458028B1 (ko) * 2007-05-30 2014-11-04 삼성전자 주식회사 병렬 처리 장치 및 방법
DE102007025397B4 (de) * 2007-05-31 2010-07-15 Advanced Micro Devices, Inc., Sunnyvale System mit mehreren Prozessoren und Verfahren zu seinem Betrieb
CN101329638B (zh) * 2007-06-18 2011-11-09 国际商业机器公司 程序代码的并行性的分析方法和系统
TW200917277A (en) * 2007-10-15 2009-04-16 A Data Technology Co Ltd Adaptive hybrid density memory storage device and control method thereof
US20090240930A1 (en) * 2008-03-24 2009-09-24 International Business Machines Corporation Executing An Application On A Parallel Computer
US9477587B2 (en) * 2008-04-11 2016-10-25 Micron Technology, Inc. Method and apparatus for a volume management system in a non-volatile memory device
US8161483B2 (en) * 2008-04-24 2012-04-17 International Business Machines Corporation Configuring a parallel computer based on an interleave rate of an application containing serial and parallel segments
US8291427B2 (en) * 2008-06-09 2012-10-16 International Business Machines Corporation Scheduling applications for execution on a plurality of compute nodes of a parallel computer to manage temperature of the nodes during execution
US20100079454A1 (en) 2008-09-29 2010-04-01 Legakis Justin S Single Pass Tessellation
US9672019B2 (en) * 2008-11-24 2017-06-06 Intel Corporation Systems, apparatuses, and methods for a hardware and software system to automatically decompose a program to multiple parallel threads
US8528001B2 (en) * 2008-12-15 2013-09-03 Oracle America, Inc. Controlling and dynamically varying automatic parallelization
US20100169540A1 (en) * 2008-12-30 2010-07-01 Sinclair Alan W Method and apparatus for relocating selected data between flash partitions in a memory device
US8321645B2 (en) 2009-04-29 2012-11-27 Netapp, Inc. Mechanisms for moving data in a hybrid aggregate
US8914799B2 (en) * 2009-06-30 2014-12-16 Oracle America Inc. High performance implementation of the OpenMP tasking feature
US8561046B2 (en) * 2009-09-14 2013-10-15 Oracle America, Inc. Pipelined parallelization with localized self-helper threading
US8271729B2 (en) * 2009-09-18 2012-09-18 International Business Machines Corporation Read and write aware cache storing cache lines in a read-often portion and a write-often portion
US9798543B2 (en) 2009-09-24 2017-10-24 Nvidia Corporation Fast mapping table register file allocation algorithm for SIMT processors
US8677106B2 (en) 2009-09-24 2014-03-18 Nvidia Corporation Unanimous branch instructions in a parallel thread processor
US8335892B1 (en) 2009-09-28 2012-12-18 Nvidia Corporation Cache arbitration between multiple clients
EP2499576A2 (en) 2009-11-13 2012-09-19 Richard S. Anderson Distributed symmetric multiprocessing computing architecture
US8612978B2 (en) * 2009-12-10 2013-12-17 Oracle America, Inc. Code execution utilizing single or multiple threads
US9696995B2 (en) 2009-12-30 2017-07-04 International Business Machines Corporation Parallel execution unit that extracts data parallelism at runtime
US20110191522A1 (en) 2010-02-02 2011-08-04 Condict Michael N Managing Metadata and Page Replacement in a Persistent Cache in Flash Memory
US9235531B2 (en) 2010-03-04 2016-01-12 Microsoft Technology Licensing, Llc Multi-level buffer pool extensions
CN101819675B (zh) 2010-04-19 2011-08-10 浙江大学 一种基于gpu的层次包围盒的快速构造方法
US8650554B2 (en) 2010-04-27 2014-02-11 International Business Machines Corporation Single thread performance in an in-order multi-threaded processor
US8898324B2 (en) 2010-06-24 2014-11-25 International Business Machines Corporation Data access management in a hybrid memory server
US8751771B2 (en) 2010-09-29 2014-06-10 Nvidia Corporation Efficient implementation of arrays of structures on SIMT and SIMD architectures
US8547385B2 (en) 2010-10-15 2013-10-01 Via Technologies, Inc. Systems and methods for performing shared memory accesses
US9552206B2 (en) 2010-11-18 2017-01-24 Texas Instruments Incorporated Integrated circuit with control node circuitry and processing circuitry
KR101713051B1 (ko) 2010-11-29 2017-03-07 삼성전자주식회사 하이브리드 메모리 시스템, 및 그 관리 방법
US8996842B2 (en) 2010-12-09 2015-03-31 Seagate Technology Llc Memory stacks management
KR20120082218A (ko) * 2011-01-13 2012-07-23 (주)인디링스 파티션 정보를 기초로 호스트의 요청에 대한 처리 기법을 적응적으로 결정하는 스토리지 장치 및 상기 스토리지 장치의 동작 방법
US9195550B2 (en) 2011-02-03 2015-11-24 International Business Machines Corporation Method for guaranteeing program correctness using fine-grained hardware speculative execution
JP5668573B2 (ja) 2011-03-30 2015-02-12 日本電気株式会社 マイクロプロセッサ、メモリアクセス方法
US8769537B1 (en) * 2011-06-08 2014-07-01 Workday, Inc. System for partitioning batch processes
KR101895605B1 (ko) * 2011-11-21 2018-10-25 삼성전자주식회사 플래시 메모리 장치 및 그것의 프로그램 방법
US9921873B2 (en) * 2012-01-31 2018-03-20 Nvidia Corporation Controlling work distribution for processing tasks
US9063759B2 (en) 2012-03-28 2015-06-23 International Business Machines Corporation Optimizing subroutine calls based on architecture level of called subroutine
US9575806B2 (en) 2012-06-29 2017-02-21 Intel Corporation Monitoring accesses of a thread to multiple memory controllers and selecting a thread processor for the thread based on the monitoring
TWI479314B (zh) * 2012-08-30 2015-04-01 Phison Electronics Corp 系統資料儲存方法、記憶體控制器與記憶體儲存裝置
US9436475B2 (en) 2012-11-05 2016-09-06 Nvidia Corporation System and method for executing sequential code using a group of threads and single-instruction, multiple-thread processor incorporating the same

Patent Citations (4)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN101030152A (zh) * 2007-03-20 2007-09-05 华为技术有限公司 基于伪同步方式的操作控制方法及装置
US20090013323A1 (en) * 2007-07-06 2009-01-08 Xmos Limited Synchronisation
CN102135916A (zh) * 2010-10-15 2011-07-27 威盛电子股份有限公司 同步方法以及图形处理系统
US20120191947A1 (en) * 2011-01-20 2012-07-26 International Business Machine Corporation Computer operation control method, program and system

Non-Patent Citations (2)

* Cited by examiner, † Cited by third party
Title
孙洪迪、高柱: "基于OpenMP技术的多核处理器程序的开发实现", 《北京工业职业技术学院学报》, vol. 9, no. 1, 31 January 2010 (2010-01-31), pages 19 - 22 *
滕人达、刘青昆: "CUDA、MPI和OpenMP三级混合并行模型研究", 《微计算机应用》, vol. 31, no. 9, 30 September 2010 (2010-09-30), pages 63 - 69 *

Cited By (3)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
CN112114877A (zh) * 2020-09-28 2020-12-22 西安芯瞳半导体技术有限公司 一种动态补偿线程束warp的方法、处理器及计算机存储介质
CN117009054A (zh) * 2023-07-27 2023-11-07 北京登临科技有限公司 一种simt装置、线程组动态构建方法及处理器
CN117009054B (zh) * 2023-07-27 2024-06-28 北京登临科技有限公司 一种simt装置、线程组动态构建方法及处理器

Also Published As

Publication number Publication date
TW201439907A (zh) 2014-10-16
US20140129812A1 (en) 2014-05-08
US9727338B2 (en) 2017-08-08
CN103809964A (zh) 2014-05-21
US20140129783A1 (en) 2014-05-08
US9747107B2 (en) 2017-08-29
TW201439905A (zh) 2014-10-16
TW201443783A (zh) 2014-11-16
TWI510919B (zh) 2015-12-01
CN103809964B (zh) 2017-06-16
US20140130021A1 (en) 2014-05-08
TW201443639A (zh) 2014-11-16
US9710275B2 (en) 2017-07-18
US20140130052A1 (en) 2014-05-08
CN103809963A (zh) 2014-05-21
TWI488111B (zh) 2015-06-11
TWI494853B (zh) 2015-08-01
US9436475B2 (en) 2016-09-06
CN103885751A (zh) 2014-06-25

Similar Documents

Publication Publication Date Title
CN103809936A (zh) 编译或运行时执行分叉-合并数据并行程序的系统和方法
Pérache et al. MPC: A unified parallel runtime for clusters of NUMA machines
KR101759266B1 (ko) 프로세서들에 걸쳐 데이터-병렬 쓰레드들을 지닌 프로세싱 로직을 매핑하는 방법
Grasso et al. LibWater: heterogeneous distributed computing made easy
CN101366004A (zh) 用于带有专用线程管理的多核处理的方法和设备
EP2601577B1 (en) A method and apparatus for a compiler and related components for stream-based computations for a general-purpose, multiple-core system
Meng et al. The Uintah framework: A unified heterogeneous task scheduling and runtime system
Kamal et al. FG-MPI: Fine-grain MPI for multicore and clusters
WO2007128168A1 (en) Thread scheduling on multiprocessor systems
Imam et al. Cooperative scheduling of parallel tasks with general synchronization patterns
Chen et al. Mgmr: Multi-gpu based mapreduce
Yazdanpanah et al. Analysis of the task superscalar architecture hardware design
Buono et al. Optimizing message-passing on multicore architectures using hardware multi-threading
JP7450728B2 (ja) 協調ワークスティーリングスケジューラ
Gharajeh et al. Heuristic-based task-to-thread mapping in multi-core processors
Garg et al. Share-a-GPU: Providing simple and effective time-sharing on GPUs
Jesshope et al. The implementation of an svp many-core processor and the evaluation of its memory architecture
CN116775265A (zh) 协作组阵列
CN116774914A (zh) 分布式共享存储器
Yang et al. Managing asynchronous operations in Coarray Fortran 2.0
Prell et al. Go's Concurrency Constructs on the SCC
Singh Communication Coroutines For Parallel Program Using DW26010 Many Core Processor
Bernard et al. Resource-agnostic programming for many-core microgrids
Zhang et al. DMR: A deterministic MapReduce for multicore systems
Uddin Microgrid-The microthreaded many-core architecture

Legal Events

Date Code Title Description
C06 Publication
PB01 Publication
C10 Entry into substantive examination
SE01 Entry into force of request for substantive examination
WD01 Invention patent application deemed withdrawn after publication

Application publication date: 20140521

WD01 Invention patent application deemed withdrawn after publication