用线程组执行顺序代码的系统和方法和包含其的simt处理器的制造方法
【专利摘要】本发明提供用线程组执行顺序代码的系统和方法和包含其的SIMT处理器。用于在单指令、多线程（SIMT）处理器的上下文中执行顺序代码的系统和方法。在一个实施例中，系统包括：（1）管线控制单元，可操作以创建顺序代码的副本线程组，副本线程中的一个是主线程，副本线程中的其余线程是从线程，以及（2）通道，可操作以：（2a）仅在主线程中执行顺序代码的某些指令，从线程中的相应指令依据该某些指令而被断言，以及（2b）将主线程中的分支条件广播到从线程。
【专利说明】用线程组执行顺序代码的系统和方法和包含其的SIMT处理器
[0001]相关串请的交叉引用
[0002]本申请要求于2012年11月5日由Lin等人所提交的、序列号为61/722，661的、标题为 “EXECUTING SEQUENTIAL CODE USING A GROUP OF THREADS” 的美国临时申请以及于2012年12月21日由Lin等人所提交的、序列号为13/723，981的、标题为“SYSTEM ANDMETHOD FOR EXECUTING SEQUENTIAL CODE USING A GROUP OF THREADS AND SIGLE-1NSTRUCT10N, MULTIPLE-THREAD PROCESSOR INCORPORATING THE SAME” 的美国申请的优先权，在先申请与本申请共同受让，并在本文通过援引的方式对二者加以合并。
【技术领域】
[0003]本申请总地涉及并行处理器，并且，更具体地，涉及用于使用线程组执行顺序代码的系统和方法以及包含系统或方法的单指令多线程(SMT)处理器。
【背景技术】
[0004]如相关领域技术人员意识到的，可并行地执行应用以增加其性能。数据并行应用在不同数据上并发实行相同进程。任务并行应用在相同数据上并发实行不同进程。静态并行应用是具有可在其执行之前被确定的并行度级别的应用。相反，由动态并行应用可达到的并行度仅可随着其执行而被确定。无论应用是数据或任务并行、或静态或动态并行，其可在管线中执行，这通常是用于图形应用的情况。
[0005]SIMT处理器尤其擅长执行数据并行应用。SMT处理器中的管线控制单元创建执行的线程组并调度其用于执行，在执行期间组中的所有线程并发执行相同指令。在一个特定处理器中，每个组具有32个线程，与SMT处理器中的32个执行管线或通道(lane)相对应。
[0006]并行应用典型地包含顺序代码和并行代码区。顺序代码不能并行执行，所以执行在单线程中。当遭遇并行代码时，管线控制单元将执行分开，创建用于并行代码的并行执行的工作者线程组。当再次遭遇顺序代码时，管线控制单元对并行执行的结果加以合并、创建用于顺序代码的另一单线程，并且执行继续。
[0007]使组中的线程同步是重要的。同步部分地涉及遵循与每个通道相关联的本地存储器的状态。已经发现，如果在执行顺序代码的同时在通道中的每一个中执行顺序代码的副本(counterpart)线程，那么可使同步更快。因此如果执行稍后被分开，那么假定已经加以遵循本地存储器状态。
【发明内容】
[0008]一个方面提供用于执行顺序代码的系统。在一个实施例中，系统包括:(1)管线控制单元，可操作以创建顺序代码的副本线程组，副本线程中的一个是主线程，副本线程中的其余线程是从线程，以及(2)通道，可操作以:(2a)仅在主线程中执行顺序代码的某些指令，从线程中的相应指令依据该某些指令而被断言(predicate)，以及(2b)将主线程中的分支条件广播到从线程。
[0009]另一方面提供执行顺序代码的方法。在一个实施例中，方法包括:(I)创建顺序代码的副本线程组，副本线程中的一个是主线程，副本线程中的其余线程是从线程，(2)仅在主线程中执行顺序代码的某些指令，从线程中的相应指令依据该某些指令而被断言，以及
(3)将主线程中的分支条件广播到从线程。
[0010]又一方面提供SMT处理器。在一个实施例中，SMT处理器包括:(I)通道，(2)与通道中的相应通道相关联的本地存储器；(3)由通道所共享的存储器设备，以及(4)管线控制单元，可操作以创建顺序代码的副本线程组并使组在通道中执行，副本线程中的一个是主线程，副本线程中的其余线程是从线程。通道可操作以:(I)仅在主线程中执行顺序代码的某些指令，从线程中的相应指令依据该某些指令而被断言，以及(2)将主线程中的分支条件广播到从线程。
【专利附图】
【附图说明】
[0011]现在结合附图对下面的描述加以参考，其中:
[0012]图1是SMT处理器的框图，该SMT处理器可操作以包含或实行用于使用线程组执行顺序代码的系统或方法；
[0013]图2是用于使用线程组执行顺序代码的系统的一个实施例的框图；以及
[0014]图3是使用线程组执行顺序代码的方法的一个实施例的流程图。
【具体实施方式】
[0015]如上文所阐明的，已经发现，如果在通道中的每一个中执行顺序代码的副本线程，那么可使SMT处理器的通道或核心之间的同步过程更快。因为副本线程是相同代码的(即以相同次序的相同指令)，并且因为当代码的副本线程开始执行时遵循本地存储器状态，所以本地存储器状态将保持被遵循的假定似乎几成定局。然而，本文认识到的是，可存在这样的条件，在该条件下存储器状态发散(diverge)。
[0016]作为一个示例，假定顺序代码的副本线程要执行相同加载指令。加载的存储器位置由寄存器或地址二者之一所指定。如果由寄存器所指定，那么寄存器的值可按线程变化，因为每个线程具有其自己的寄存器的拷贝。如果由地址所指定，那么地址值可指向系统中的不同的线程本地存储器位置。在二者中的任一情况下，每个线程可从多种存储器位置加载不同的值，这致使线程本地存储器状态发散。如果副本线程随后基于所加载的数据进行分支，那么所采取的一些分支将是正确的，其他的将是错误的。
[0017]类似地，假定顺序代码的副本线程要执行相同存储指令。被存储到的存储器出于如上文针对加载指令所描述的相同原因而按线程变化。在顺序执行中未被修改的存储器位置在并行执行中将被错误地修改。
[0018]作为另一示例，假定顺序代码的副本线程要将数据并发地存储到共享存储器中的相同位置。结果，可能再次淹没(overwhelm)和损坏共享存储器。在这两个示例中所强调的问题均有时会在矢量操作中经历。
[0019]作为又一示例，假定异常处置器是各种通道间的共享资源。顺序代码区通常包括可潜在地使异常发生的许多指令。在并行执行这些指令的同时如果出现异常，那么并行处理可能抛出同时的异常并淹没共享的异常处置器，该共享的异常处置器将预期最多一个异常，并且可能并未预期任何异常。
[0020]因此本文认识到的是，在执行顺序代码的副本线程的同时本地存储器状态将必然保持被遵循的假定是难以维系的。本文进一步认识到的是，某些操作可能损坏共享存储器或使本地存储器状态发散作为“副作用”，所述某些操作不仅包括加载自和存储到共享存储器而且包括划分和潜在地导致异常的其他指令。本文还进一步认识到的是，需要机制以确保顺序代码的语义不会经由发散线程本地存储器状态而被曲解。
[0021]因此，本文所引入的是用于使用线程组执行顺序代码的系统和方法的各种实施例。从很高级别来看，各种实施例使顺序代码的副本线程执行仿真顺序代码的主线程执行。
[0022]根据各种实施例，副本线程中的一个被指派为主线程，其他线程被指派为从线程。随后以主线程中的相应指令来断言从线程中的某些指令(典型地，那些可采用共享资源或确实采用共享资源的指令)，并且仅执行主线程中的相应指令。如果在主线程中遭遇分支指令，那么主线程中的分支条件随后被广播到从线程。
[0023]图1是SMT处理器100的框图，该SMT处理器100可操作以包含或实行用于使用线程组执行顺序代码的系统或方法。SMT处理器100包括被组织成线程组104或“线程束(warp)”的多个线程处理器或核心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，省略线程块组织级别。
[0024]SIMT处理器100进一步包括管线控制单元108、共享存储器110和与线程组104_1到104-J相关联的本地存储器112-1到112-J的阵列。管线控制单元108通过数据总线114将任务分布到各个线程组104-1到104-J。管线控制单元108创建、管理、调度、执行并提供机制以将线程组104-1到104-J同步。在图形处理单元(GPU)内找到SMT处理器100的某些实施例。一些GPU提供组同步指令，诸如由加利福尼亚州圣塔克拉拉市的Nvidia公司所制造的GPU中的bar.sync。某些实施例支持线程组的发散条件分支的执行。给定分支，线程组104内的一些线程将因为分支条件断言评估为“真”而采取分支，并且其他线程因为分支条件断言评估为“伪”而落到下一指令。管线控制单元108通过以下各项跟踪活动的线程:首先执行路径之一，其中采取了分支或未采取分支，随后执行替代路径，这针对每个路径使能适当的线程。
[0025]继续图1的实施例，线程组内的核心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。[0026]图1的实施例包括主线程组104-1。其余线程组104-2到104-J中的每一个被视为“工作者”线程组。主线程组104-1包括许多核心，其中的一个是主核心106-1，该主核心106-1最终执行主线程。在SMT处理器110上所执行的程序被构建为内核的序列。典型地，每个内核在下一内核开始之前完成执行。在某些实施例中，SIMT100可并行执行多个内核，这取决于内核的大小。每个内核被组织为要在核心106上所执行的线程的层级。
[0027]图2是用于使用线程组执行顺序代码的系统200的一个实施例的框图。系统200包括具有顺序区204和并行区206的程序202、存储器208、断言模块210、线程标识器212、线程启动器214和线程组104。图1的线程组104包括K个核心106-1到106-K或通道。
[0028]线程组104耦连到存储器208，该存储器208被分配到与核心106_1到106-K中的每一个有关的段(section)。线程启动器214在核心106-1到106-K中创建处理线程。指派通常是第一核心106-1的一个核心执行主线程。其余线程是工作者线程。传统上，主线程执行程序202的顺序区204，并且并行区206传统上在工作者线程中执行。当到达并行区206时，线程启动器214创建必要的工作者线程以执行并行处理。
[0029]在图2的实施例中，程序202的顺序区204由断言模块210所处理。断言模块指派某些操作仅在主线程上实行。由线程标识器212实现断言，该线程标识器202标识主线程用于处理该某些操作。顺序区204的平衡在线程组104中的所有线程中执行。当工作者线程到达顺序区204的所断言的片段(segment)时，工作者线程跳过所断言的片段并继续直到到达分支语句为止。当工作者线程到达分支语句时，其等待来自主线程的引导，因为仅主线程可能可靠地评估分支条件。一旦主线程处理所断言的片段、到达分支语句并评估分支条件，主线程将分支条件广播到工作者线程中的每一个。工作者线程可随后恢复通过程序202的顺序区204的继续进行。
[0030]图3是使用线程组执行顺序代码的方法的一个实施例的流程图。顺序代码可以是矢量操作的一部分、根据OpenMP或OpenACC编程模型所开发的程序的一部分、或与无论任何类型的另一应用相关联。
[0031]方法开始于开始步骤310。在步骤320，创建顺序代码的副本线程组，副本线程中的一个是主线程，副本线程中的其余线程是从线程。在步骤330，仅在主线程中执行顺序代码的某些指令，从线程中的相应指令依据该某些指令而被断言。在各种实施例中，某些指令可以是加载指令、存储指令、划分指令、或可产生或可被认为产生副作用的任何其他指令。在一个实施例中，使用基于线程标识器的条件来断言相应指令。
[0032]在步骤340，将主线程中的分支条件广播到从线程。在一个实施例中，在主线程中的分支指令的执行之前广播分支条件，并且仅在广播之后在从线程中执行相应分支指令。方法结束于结束步骤350。
[0033]本申请相关领域技术人员将理解的是，可对所描述的实施例做出其他和进一步的添加、删除、替换和修改。
【权利要求】
1.一种用于执行顺序代码的系统，包括: 管线控制单元，可操作以创建所述顺序代码的副本线程组；所述副本线程中的一个是主线程，所述副本线程中的其余线程是从线程；以及通道，可操作以: 仅在所述主线程中执行所述顺序代码的某些指令，所述从线程中的相应指令依据所述某些指令而被断言，以及 将所述主线程中的分支条件广播到所述从线程。
2.根据权利要求1所述的系统，其中与执行所述从线程的通道相关联的本地存储器进一步配置为存储所述分支条件。
3.根据权利要求1所述的系统，其中所述某些指令选择自包括以下各项的组中: 加载指令， 存储指令，以及 异常引出指令。
4.根据权利要求1所述的系统，其中执行所述主线程的通道进一步可操作以在所述主线程中的分支指令的执行之前广播所述分支条件，并且执行所述从线程的通道进一步可操作以仅在所述通道广播所述分支条件之后在所述从线程中执行相应分支指令。
5.根据权利要求1所述的系统，其中所述管线控制单元进一步可操作以使用基于线程标识器的条件来断言所述相应指令。
6.根据权利要求1所述的系统，其中所述顺序代码是矢量操作的一部分。
7.根据权利要求1所述的系统，其中所述顺序代码是选择自包括以下各项的组中的应用的一部分: OpenMP程序，以及 OpenACC 程序。
8.一种单指令、多线程(SMT)处理器，包括: 通道； 与所述通道中的相应通道相关联的本地存储器； 由所述通道所共享的存储器设备；以及 管线控制单元，可操作以创建所述顺序代码的副本线程组并使所述组在所述通道中执行，所述副本线程中的一个是主线程，所述副本线程中的其余线程是从线程，所述通道可操作以: 仅在所述主线程中执行所述顺序代码的某些指令，所述从线程中的相应指令依据所述某些指令而被断言，以及 将所述主线程中的分支条件广播到所述从线程。
9.根据权利要求8所述的SMT处理器，其中与执行所述从线程的通道相关联的所述本地存储器进一步配置为存储所述分支条件。
10.根据权利要求8所述的SMT处理器，其中所述某些指令选择自包括以下各项的组中: 加载指令， 存储指令，以及异常引出指 令。
【文档编号】G06F9/44GK103809964SQ201310538631
【公开日】2014年5月21日 申请日期:2013年11月4日 优先权日:2012年11月5日
【发明者】高塔姆·查克拉瓦蒂, 林 源, 杰迪普·马拉蒂, 权冠, 阿米特·萨布尼 申请人:辉达公司