TW201443783A - 在單一指令多執行緒處理器上以函數呼叫編譯或執行時間執行一分叉－會合資料平行程式的系統及方法 - Google Patents

在單一指令多執行緒處理器上以函數呼叫編譯或執行時間執行一分叉－會合資料平行程式的系統及方法 Download PDF

Info

Publication number
TW201443783A
TW201443783A TW102140064A TW102140064A TW201443783A TW 201443783 A TW201443783 A TW 201443783A TW 102140064 A TW102140064 A TW 102140064A TW 102140064 A TW102140064 A TW 102140064A TW 201443783 A TW201443783 A TW 201443783A
Authority
TW
Taiwan
Prior art keywords
thread
function
parallel
main
execution
Prior art date
Application number
TW102140064A
Other languages
English (en)
Inventor
Yuan Lin
Gautam Chakrabarti
Jaydeep Marathe
O-Kwan Kwon
Amit Sabne
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 TW201443783A publication Critical patent/TW201443783A/zh

Links

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/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/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/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

單一指令多執行緒處理器上以函數呼叫編譯或執行時間執行一分叉-會合資料平行程式的系統及方法 【相關申請案】
本申請案係主張美國專利臨時申請案案號61/722,661之優先權，其由Lin等人於2012年11月5日提出申請，標題為「使用一群組的執行緒執行循序碼(EXECUTING SEQUENTIAL CODE USING A GROUP OF THREADS)」，以及主張美國專利申請案案號13/724,359之優先權，其由Lin等人於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)」，兩案皆與本申請案共同讓與且併入本文作為參考。
本發明一般係有關平行處理器，特別是關於用以在單一指令多執行緒(SIMT)處理器上以函數呼叫編譯或執行時間執行一分叉-會合資 料平行程式的系統及方法。
如熟此技藝者所知，應用程式或程式可並行地執行以增加其效能。資料平行應用程式同時地對不同資料完成相同的程序。任務平行應用程式同時地對相同的資料完成不同的程序。靜態平行應用程式為具有在其執行前可被判定之平行度的應用程式。相反地，動態平行應用程式可達成的平行化僅可在其執行時被判定。不論應用程式為資料或任務平行、或是靜態或動態平行，其可在通常為用於圖形程式之情況的管線中被執行。
SIMT處理器特別擅長執行資料平行應用程式。SIMT處理器中的控制單元產生執行的執行緒群組並將其排程供執行，在其過程中，群組中的所有執行緒同時地執行相同的指令。在一特定的處理器中，每一群組或「執行包(warp)」具有32個執行緒，對應SIMT處理器中的32個執行管線或通道。
分叉-會合(fork-join)資料平行程式開始於單一執行緒主程式。程式在此階段係在循序狀態或區域中。在主程式之執行過程中的某一點，主(main或「master」)執行緒遇到一序列的平行狀態或區域。每一平行區域具有獨立的資料組且可由多個執行緒同時地執行。當平行區域開始且在平行區域期間沒有改變，判定在每一平行區域中之同時任務的數量。當遇到一平行區域，主執行緒將一組執行緒(稱作背景工作執行緒)分叉以平行地執行平行區域。接著程式進入平行區域。若一背景工作執行緒遇到新的平行區域，新的平行區域將被序列化，即平行區域將由遇到的背景工作執 行緒本身所執行。主執行緒等候直到平行區域結束。一旦離開平行區域，背景工作執行緒與主執行緒會合，其接著恢復主程式的執行，在此時程式進入循序區域。
以下的表格1提出分叉-會合資料平行程式的範例。
為了理解表格1及本發明揭露內容的其餘部份，術語「foo」 及「bar」為函數的任意名稱。因此，任何函數可代替「foo」或「bar」。
分叉-會合資料平行模型通常用於平行程式化中。舉例來說，OpenMP標準採用此模型作為其基本執行緒執行模型。OpenACC標準使用此模型用於在稱作「集體(gang)」之一群組中的背景工作執行緒。
一態樣提供一種以函數呼叫編譯或執行時間執行一分叉-會合資料平行程式的系統。在一具體實施例中，系統包含：(1)一分割器，可操作以將執行包(warps)分割為一主執行包及至少一背景工作執行包，以及(2)一執行緒指示器，關聯於分割器且可操作以從主執行包僅指定一個執行緒供執行並指定在至少一背景工作執行包中的所有執行緒供執行。
另一態樣提供一種以函數呼叫編譯或執行時間執行一分叉-會合資料平行程式的方法。在一具體實施例中，方法包含：(1)將執行包分割為一主執行包及至少一背景工作執行包，以及(2)從主執行包僅指定一個執行緒供執行並指定在至少一背景工作執行包中的所有執行緒供執行。
100‧‧‧SIMT處理器
102‧‧‧執行緒區塊
104-1~104-J‧‧‧執行緒群組
106‧‧‧核心
106-1~106-K‧‧‧核心
108‧‧‧管線控制單元
110‧‧‧共享記憶體
112-1~112-J‧‧‧本地記憶體
114‧‧‧資料匯流排
116‧‧‧記憶體匯流排
118-1~118-J‧‧‧本地匯流排
200‧‧‧系統
202‧‧‧分割器
204‧‧‧執行緒指示器
206‧‧‧執行緒排程器
208‧‧‧函數處理器
210‧‧‧程式
212‧‧‧入口函數
214‧‧‧非入口函數
216‧‧‧外部函數
218‧‧‧裝置執行時間程式庫
現在將參照以下描述並連同所附隨圖式，其中：圖1為可操作以包含或實現以函數呼叫編譯或執行時間執行一分叉-會合資料平行程式之系統或方法的SIMT處理器的方塊圖；圖2為以函數呼叫編譯或執行時間執行一分叉-會合資料平行程式之系統具體實施例的方塊圖；以及 圖3為以函數呼叫編譯或執行時間執行一分叉-會合資料平行程式之方法具體實施例的流程圖。
在一SIMT處理器中，執行的多個執行緒係劃分為群組。在群組中的所有執行緒在相同時間執行相同指令。在來自加州聖塔克拉拉之Nvidia公司的市售圖形處理單元(GPUs)中(其為SIMT處理器的一種類型)，群組係稱作「執行包(warps)」，且其以區塊為單位執行。
SIMT處理器的管線控制單元產生、管理、排程、執行及提供一機制以同步化群組。Nvidia GPUs提供bar.sync指令以同步化群組。Nvidia GPUs也支援「發散」條件分支由一群組的執行；群組的某些執行緒需採用分支(因為分支條件判定評估為「真」)，且其他執行緒需落到下一指令(因為分支條件判定評估為「假」)。管線控制單元追蹤群組中的有效執行緒。其首先執行其中一路徑(採用的分支或未採用的分支)並接著執行另一路徑；在每一路徑致能適當的執行緒。
本文中將理解到，在一GPU執行緒區塊內的所有執行緒開始於相同的程式位址，管線控制單元將從軟體機制中獲益以將執行緒劃分並排程為主執行緒及背景工作執行緒，使得其可在分叉-會合模型中執行。
更將理解到，軟體機制的某些具體實施例應按群組來管理並同步化執行緒，因為管線控制單元係按群組來管理執行緒。
更將理解到，由於主程式在分叉-會合模型中為單一線程執行，軟體機制的某些具體實施例應達成循序區域語意而不引入副作用。引 起副作用之指令的範例為使用共享資源者，例如共享記憶體讀取或寫入或可能引起共享異常處理程序(例如分割)的任何程式碼操作。
更將理解到，軟體機制的某些具體實施例應支援可在循序區域內及在平行區域內呼叫的函數。此外，這類函數本身可包含平行結構。更將理解到，軟體機制的某些具體實施例應支援可分叉-會合平行區域的函數呼叫。
更將理解到，軟體機制的某些具體實施例應支援外部函數，即未由相同編譯器編譯為程式的函數。舉例來說，現存GPU數學程式庫中的數學函數、以及像是malloc、free及print的系統函數。在某些具體實施例中，循序區域中的主執行緒及平行區域中的背景工作執行緒兩者皆應能夠呼叫一外部函數。
因此，本文描述在一SIMT處理器(如GPU)上以函數呼叫編譯或執行時間執行分叉-會合資料平行程式之系統及方法的各種具體實施例。
在描述系統及方法的特定具體實施例之前，將描述可操作以含有或實現以函數呼叫編譯或執行時間執行分叉-會合資料平行程式的系統或方法的SIMT處理器。
圖1為SIMT處理器100的方塊圖。SIMT處理器100包含多個執行緒處理器或核心106，其組織為執行緒群組104或「執行包(warps)」。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到104-J經由本地匯流排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的每一者係視為「背景工作(worker)」執行緒群組。主執行緒群組104-1包含許多核心，其中一者為主核心106-1，其基本上執行主執行緒。在SIMT處理器100上執行的程式係建構為一序列的核(kernels)。一般而言，每一核在下一核開始前完成執行。在某些具體實施例中，SIMT處理器 100可平行地執行多個核，其取決於核的尺寸。每一核係組織為在要核心106上執行之執行緒的一階層。
已經描述了可包含或實現本文所介紹之系統或方法於其中的SIMT處理器，將描述系統及方法的各種具體實施例。
本文所介紹之系統的一具體實施例包含一編譯器及一裝置執行時間程式庫。裝置執行時間程式庫實現了執行緒及群組管理功能。編譯器將分叉-會合資料平行程式轉譯為主執行緒程式及一組概括函數，其每一對應至一平行結構。轉譯的程式碼呼叫在裝置執行時間程式庫中的函數以進行執行緒及群組管理。
以下的表格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()為在裝置執行時間程式庫中的函數。
當一GPU執行緒區塊開始，區塊內的所有執行緒執行main()；然而，其採取不同路徑。執行緒0為主執行緒且執行init()、main_core()及signal_done()。群組0內的其他執行緒直接到main()函數的終點並在那裡等候。剩餘群組中的執行緒執行scheduler()。
針對一非入口函數，如foo()、bar()及main_core()，編譯器轉譯程式碼如同無平行結構存在。若非入口函數含有一平行結構，則針對每一平行結構，編譯器產生含有平行結構之本體的一函數(一概括函數)，並接著產生檢查執行的執行緒是否為主執行緒的條件分支。在假分支中，編譯器插入執行迴圈的程式碼。在真分支中，編譯器插入呼叫到裝置執行時間 程式庫以分配任務、喚醒背景工作執行緒、並執行一阻障。當非入口函數在平行區域之外被呼叫，條件為真。當非入口函數在平行區域內被呼叫，條件為假，在此情況中平行迴圈係由執行的執行緒依序地執行。
舉例來說，函數bar()的轉譯程式碼係顯示於以下的表格4中。
Signal_task()及barrier()為裝置執行時間程式庫中的函數。bar_par_frunc()為對應原始函數bar()中之平行結構的概括函數。
在此具體實施例中，裝置執行時間程式庫包含以下等函數：init()、scheduler()、signal_task()、signal_done()、及barrier()。程式庫亦 執行以下函數供內部使用：signal()、wait()、及fetch_task().
所有背景工作執行緒執行scheduler()函數。背景工作執行緒經歷休眠-喚醒-執行循環，直到被指示離開。
布林變數「exit_flag」係放入區塊共享記憶體中且可由執行緒區塊內的所有執行緒所存取。其係由主執行緒使用以對背景工作執行緒傳達其是否應全部退出執行。「exit_flag」在init()函數中係設定為假，且在signal_done()函數中係設定為真。兩個函數係由主執行緒呼叫。
區塊共享記憶體的另一部份係用以通訊當前任務。當前任務係由在signal_task()函數中的主執行緒所設定，並由在fetch_task()函數中的背景工作執行緒所提取。區塊共享記憶體包含到對應平行結構之概括函數的指標。
因為平行區域在執行緒區塊內依序執行，所以任何時候都只有一個任務為有效。若平行區域可非同步地執行，一般來說需要更複雜的資料結構(如堆疊、佇列或樹狀結構)以儲存有效任務。
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-1及共享記憶體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.離開。
在主群組中的其他執行緒基本上係在程式的終點等待，閒置。程式由主執行緒及背景工作執行緒交錯執行。這導致良好的指令快取覆蓋區，其比由主執行緒及背景工作執行緒兩者皆為有效且執行不同程式碼路徑的方法所產生之覆蓋區更佳。
圖3為以函數呼叫編譯或執行時間執行一分叉-會合資料平行程式之方法具體實施例的流程圖。方法開始於開始步驟310。在步驟320，在一執行緒區塊內的執行緒群組係分割為一主執行緒群組及至少一背景工作執行緒群組。在步驟330，主執行緒群組中的一執行緒係指定為主執行緒。主群組中的剩餘執行緒在整個執行過程中基本上是閒置的。同樣在步驟330，在至少一背景工作執行緒群組中的所有執行緒係指定為背景工作執行緒。方法結束於結束步驟340。
熟習本申請案相關技藝者將理解到，可對所述具體實施例做出其他及更多添加、刪減、替換或修改。
100‧‧‧SIMT處理器
104-1、104-2‧‧‧執行緒群組
106、106-1‧‧‧核心
108‧‧‧管線控制單元
110‧‧‧共享記憶體
114‧‧‧資料匯流排
118-1、118-2‧‧‧本地匯流排
200‧‧‧系統
202‧‧‧分割器
204‧‧‧執行緒指示器
206‧‧‧執行緒排程器
208‧‧‧函數處理器
210‧‧‧程式
212‧‧‧入口函數
214‧‧‧非入口函數
216‧‧‧外部函數
218‧‧‧裝置執行時間程式庫

Claims (10)

  1. 一種以函數呼叫編譯或執行時間執行一分叉-會合資料平行程式的系統，包含：一分割器，可操作以將群組分割為一主群組及至少一背景工作群組；以及一執行緒指示器，關聯於該分割器且可操作以從該主群組僅指定一個執行緒供執行為主執行緒並指定在該至少一背景工作群組中的所有執行緒供執行為背景工作執行緒。
  2. 如申請專利範圍第1項所述之系統，更包含一執行緒排程器，該執行緒排程器關聯於該執行緒指示器且可操作以使一單一指令多執行緒處理器的一管線控制單元將該主執行緒的執行排程如下：當該主執行緒開始執行時，一程式離開旗標係設定為一第一狀態；該主執行緒執行直到達到一平行區域或該程式的一終點；在到達該平行區域，設定一平行任務、進入一第一阻障、進入一第二阻障且恢復該主執行緒的執行；以及在到達該終點，該程式離開旗標係設定為一第二狀態、進入一阻障且該主執行緒停止執行。
  3. 如申請專利範圍第1項所述之系統，更包含一執行緒排程器，該執行緒排程器關聯於該執行緒指示器且可操作以使一單一指令多執行緒處理器的一管線控制單元將該背景工作執行緒的執行排程如下：該背景工作執行緒進入一第一阻障；當該主執行緒到達一平行區域且該主執行緒進入一阻障時，該背景工作執行緒開始由該主執行緒所設定之一平行任務的執行；在該平行任務完成時該背景工作執行緒進入一第二阻障；以及若一程式離開旗標係設定為一第二狀態且該主執行緒進入一阻障，則該背景工作執行緒離開。
  4. 如申請專利範圍第1項所述之系統，更包含一執行緒排程器，該執行緒排程器關聯於該執行緒指示器且可操作以使用一單一指令多執行緒處理器之一管線控制單元的一阻障函數以控制該背景工作執行緒的執行及停止。
  5. 如申請專利範圍第1項所述之系統，更包含一函數處理器，該函數處理器關聯於該執行緒指示器且可操作以產生一入口函數的一複製複本並將該入口函數處理為一非入口函數。
  6. 如申請專利範圍第1項所述之系統，更包含一函數處理器，該函數處理器關聯於該執行緒指示器且可操作以藉由以下而轉譯具有一平行結構的一非入口函數： 產生含有該平行結構之一本體的一函數；以及若該函數在該主執行緒內執行，插入呼叫至一裝置執行時間程式庫。
  7. 如申請專利範圍第1項所述之系統，更包含一函數處理器，該函數處理器關聯於該執行緒指示器且可操作以依現況轉譯呼叫為外部函數。
  8. 如申請專利範圍第1項所述之系統，更包含一函數處理器，該函數處理器關聯於該執行緒指示器且可操作以使用一裝置執行時間程式庫，該裝置執行時間程式庫提供可從編譯使用者程式碼呼叫的函數以及內部函數。
  9. 如申請專利範圍第1項所述之系統，更包含一函數處理器，該函數處理器關聯於該執行緒指示器且可操作以使用一離開旗標，該離開旗標係儲存於共享記憶體中且用以允許該主執行緒傳達給該背景工作執行緒其何時應停止執行。
  10. 如申請專利範圍第1項所述之系統，其中該系統可操作以組態一單一指令多執行緒處理器的共享記憶體以識別一當前任務。
TW102140064A 2012-11-05 2013-11-05 在單一指令多執行緒處理器上以函數呼叫編譯或執行時間執行一分叉－會合資料平行程式的系統及方法 TW201443783A (zh)

Applications Claiming Priority (2)

Application Number Priority Date Filing Date Title
US201261722661P 2012-11-05 2012-11-05
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
TW201443783A true TW201443783A (zh) 2014-11-16

Family

ID=50623483

Family Applications (4)

Application Number Title Priority Date Filing Date
TW102140064A TW201443783A (zh) 2012-11-05 2013-11-05 在單一指令多執行緒處理器上以函數呼叫編譯或執行時間執行一分叉－會合資料平行程式的系統及方法
TW102140061A TWI510919B (zh) 2012-11-05 2013-11-05 分配不同特性的記憶體給共享資料物件的系統及方法
TW102140063A TWI488111B (zh) 2012-11-05 2013-11-05 轉譯程式函數以正確處理本地範圍變數的系統及方法和應用其之計算系統
TW102140062A TWI494853B (zh) 2012-11-05 2013-11-05 使用一群組的執行緒執行循序碼的系統及方法和應用其之單一指令多執行緒處理器

Family Applications After (3)

Application Number Title Priority Date Filing Date
TW102140061A TWI510919B (zh) 2012-11-05 2013-11-05 分配不同特性的記憶體給共享資料物件的系統及方法
TW102140063A TWI488111B (zh) 2012-11-05 2013-11-05 轉譯程式函數以正確處理本地範圍變數的系統及方法和應用其之計算系統
TW102140062A TWI494853B (zh) 2012-11-05 2013-11-05 使用一群組的執行緒執行循序碼的系統及方法和應用其之單一指令多執行緒處理器

Country Status (3)

Country Link
US (4) US9436475B2 (zh)
CN (4) CN103809963A (zh)
TW (4) TW201443783A (zh)

Cited By (1)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
TWI724505B (zh) * 2019-03-08 2021-04-11 開曼群島商創新先進技術有限公司 提升ｃｐｕ並行性能的方法及裝置和電子設備

Families Citing this family (31)

* 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
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
US10061591B2 (en) 2014-06-27 2018-08-28 Samsung Electronics Company, Ltd. Redundancy elimination in single instruction multiple data/thread (SIMD/T) execution processing
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
CN110990151A (zh) * 2019-11-24 2020-04-10 浪潮电子信息产业股份有限公司 一种基于异构计算平台的业务处理方法
US20210294673A1 (en) * 2020-03-19 2021-09-23 Nvidia Corporation Techniques for orchestrating stages of thread synchronization
CN112114877B (zh) * 2020-09-28 2023-03-14 西安芯瞳半导体技术有限公司 一种动态补偿线程束warp的方法、处理器及计算机存储介质
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
CN117009054B (zh) * 2023-07-27 2024-06-28 北京登临科技有限公司 一种simt装置、线程组动态构建方法及处理器

Family Cites Families (83)

* 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
WO2005096136A1 (en) * 2004-03-31 2005-10-13 Intel Corporation Stack caching using code sharing
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
US20080109795A1 (en) * 2006-11-02 2008-05-08 Nvidia Corporation C/c++ language extensions for general-purpose graphics processing unit
US7584335B2 (en) 2006-11-02 2009-09-01 International Business Machines Corporation Methods and arrangements for hybrid data storage
US7593263B2 (en) * 2006-12-17 2009-09-22 Anobit Technologies Ltd. Memory device with reduced reading latency
US8291400B1 (en) 2007-02-07 2012-10-16 Tilera Corporation Communication scheduling for parallel processing architectures
CN101030152A (zh) * 2007-03-20 2007-09-05 华为技术有限公司 基于伪同步方式的操作控制方法及装置
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 国际商业机器公司 程序代码的并行性的分析方法和系统
US8966488B2 (en) 2007-07-06 2015-02-24 XMOS Ltd. Synchronising groups of threads with dedicated hardware logic
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
US8615646B2 (en) 2009-09-24 2013-12-24 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
US8499305B2 (en) * 2010-10-15 2013-07-30 Via Technologies, Inc. Systems and methods for performing multi-program general purpose shader kickoff
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 (주)인디링스 파티션 정보를 기초로 호스트의 요청에 대한 처리 기법을 적응적으로 결정하는 스토리지 장치 및 상기 스토리지 장치의 동작 방법
JP5576305B2 (ja) * 2011-01-20 2014-08-20 インターナショナル・ビジネス・マシーンズ・コーポレーション コンピュータの動作制御方法、プログラム及びシステム
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

Cited By (2)

* Cited by examiner, † Cited by third party
Publication number Priority date Publication date Assignee Title
TWI724505B (zh) * 2019-03-08 2021-04-11 開曼群島商創新先進技術有限公司 提升ｃｐｕ並行性能的方法及裝置和電子設備
US11080094B2 (en) 2019-03-08 2021-08-03 Advanced New Technologies Co., Ltd. Method, apparatus, and electronic device for improving parallel performance of CPU

Also Published As

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

Similar Documents

Publication Publication Date Title
TW201443783A (zh) 在單一指令多執行緒處理器上以函數呼叫編譯或執行時間執行一分叉－會合資料平行程式的系統及方法
US8615646B2 (en) Unanimous branch instructions in a parallel thread processor
US10430190B2 (en) Systems and methods for selectively controlling multithreaded execution of executable code segments
CN101366004A (zh) 用于带有专用线程管理的多核处理的方法和设备
Agullo et al. Multifrontal QR factorization for multicore architectures over runtime systems
CN103547993A (zh) 通过使用由可分割引擎实例化的虚拟核来执行指令序列代码块
US10013290B2 (en) System and method for synchronizing threads in a divergent region of code
US9665354B2 (en) Apparatus and method for translating multithread program code
El-Haj-Mahmoud et al. Virtual multiprocessor: an analyzable, high-performance architecture for real-time computing
Forsell et al. An extended PRAM-NUMA model of computation for TCF programming
Huynh et al. Mapping streaming applications onto GPU systems
Gharajeh et al. Heuristic-based task-to-thread mapping in multi-core processors
US20090133022A1 (en) Multiprocessing apparatus, system and method
Bhagyanath et al. Buffer allocation for exposed datapath architectures
Jesshope et al. The implementation of an svp many-core processor and the evaluation of its memory architecture
Nanjundappa et al. Accelerating systemc simulations using gpus
Garcia et al. A dynamic schema to increase performance in many-core architectures through percolation operations
Bernard et al. Resource-agnostic programming for many-core microgrids
Rolf et al. Load-balancing methods for parallel and distributed constraint solving
KR20140081206A (ko) 컴퓨터 시스템
Kessler et al. Flexible scheduling and thread allocation for synchronous parallel tasks
Brenner et al. Executing PRAM Programs on GPUs
Jesshope et al. Evaluating CMPs and their memory architecture
Rohr On development, feasibility, and limits of highly efficient CPU and GPU programs in several fields
Ohno et al. Dynamic Task Scheduling Scheme for a GPGPU Programming Framework