TW201734770A - 計算設備和相應計算方法 - Google Patents

Abstract

本發明提供一種計算設備和相關計算方法。計算設備用於執行平行計算，包含一組綫程處理單元以及記憶體混移引擎，記憶體混移引擎包含儲存資料元素陣列的複數個暫存器以及輸入選擇器陣列，每一輸入選擇器耦接於複數個暫存器的對應子集，並經由輸出綫耦接於對應綫程處理單元，其中，依據第一控制訊號，每一輸入選擇器自複數個暫存器的對應子集發送至少第一資料元素至對應綫程處理單元，以及依據第二控制訊號，每一輸入選擇器自複數個暫存器的另一子集發送至少第二資料元素至對應綫程處理單元，其中另一子集經由其他複數個輸入綫耦接於另一輸入選擇器。本發明的計算設備和相關計算方法可以實現高效的資料訪問。

Description

計算設備和相應計算方法 【交叉引用】

本申請要求申請日為2016年3月24日，美國臨時申請號為62/312,567的美國臨時申請案的優先權，上述臨時申請案的內容一併併入本申請。

本發明的實施例有關於平行計算設備(parallel computing device)和該平行計算設備所執行的計算方法。

平行計算已經廣泛實現於現代計算系統中。為支持高效的平行計算，複數個平行編程模型被開發出來，用於程式設計員編寫跨異構平臺執行的代碼；例如，這樣的平臺可包含中央處理單元(CPU)、圖形處理單元(GPU)、數位訊號處理器(DSP)、現場可編程閘陣列(FPGA)、硬體加速器等的組合。常用平行編程模型包含開放計算語言(Open Computing Language，簡寫為OpenCLTM)、OpenCL的變體和擴展等。

通常，平行編程模型建立在平行計算平臺模型的基礎上。一個平行計算平臺模型，例如采用OpenCL的模型，包含耦接於一組計算資源(computational resource)的主機。計算資源進一步包含一組計算設備，且每一計算設備包含一組 計算單元。每一計算單元進一步包含一組處理元件。通常，主機執行串行代碼，並發布命令至計算設備。計算設備響應命令執行平行代碼，也稱為內核。內核是程式中聲明的功能，且可由複數個處理元件在複數個進程中執行。相同的內核可以作為同一工作組(workgroup)的複數個工作項執行。同一工作組的複數個工作項共享本地記憶體中的資料，並通過工作組障礙(barrier)彼此同步。

上述平行編程模型是程式設計員寫出高效的平行代碼的有力工具。然而，傳統的硬體平臺不適合具有特定資料訪問模式的某些內核，並且對於工作組的分配不靈活。因此，需要提高平行計算系統的效率。

有鑒於此，本發明特提供以下技術方案以解決上述問題。

依據本發明的一個實施例，提出一種計算設備，用於執行複數個平行計算，計算設備包含一組綫程處理單元；以及記憶體混移引擎，耦接於該組綫程處理單元，該記憶體混移引擎包含：複數個暫存器組成的暫存器陣列，儲存自記憶體緩衝器獲得的複數個資料元素的陣列；以及複數個輸入選擇器組成的輸入選擇器陣列，每一輸入選擇器經由複數個輸入綫耦接於該複數個暫存器的對應子集，並經由一個或複數個輸出綫耦接於一個或複數個對應綫程處理單元，其中，依據第一控制訊號，每一輸入選擇器自該複數個暫存器的該對應子集發送至少第一資料元素至該一個或複數個對應綫程處理單元，以及依 據第二控制訊號，每一輸入選擇器自該複數個暫存器的另一子集發送至少第二資料元素至該一個或複數個對應綫程處理單元，其中該另一子集經由其他複數個輸入綫耦接於另一輸入選擇器。

依據本發明的另一實施例，提出一種計算設備，用於執行複數個平行計算，計算設備包含控制單元，分配複數個工作組至一組批處理；以及該組批處理耦接於該控制單元，每一批處理包含：程式計數器，由分配給該批處理的M個工作組共享，其中M是依據可配置批處理設置决定的正整數；一組綫程處理單元，平行執行該M個工作組中的每一個工作組中的複數個工作項的子集；以及溢出記憶體，當該M個工作組中的一個或複數個工作組遇到同步障礙時，儲存該M個工作組的中間資料。

依據本發明的另一實施例，提出一種計算方法，用於利用複數個綫程處理單元執行複數個平行計算，計算方法包含：接收將複數個資料元素組成的資料元素陣列識別為輸入的指令；響應該指令，自記憶體緩衝器加載該資料元素陣列至複數個暫存器組成的暫存器陣列，其中該暫存器陣列耦接於複數個輸入選擇器組成的輸入選擇器陣列，且每一輸入選擇器經由複數個輸入綫耦接於該複數個暫存器的對應子集；依據第一控制訊號，每一輸入選擇器自該複數個暫存器的該對應子集發送至少第一資料元素至一個或複數個對應綫程處理單元；以及依據第二控制訊號，該每一輸入選擇器自該複數個暫存器的另一子集發送至少第二資料元素至該一個或複數個對應綫程處 理單元，其中該另一子集經由其他複數個輸入綫耦接於另一輸入選擇器。

依據本發明的另一實施例，提出一種計算方法，用於利用複數個綫程處理單元執行複數個平行計算，計算方法包含：分配複數個工作組至複數個批處理組成的一組批處理，其中該複數個批處理中的至少一個被分配給共享程式計數器的M個工作組，其中M是依據可配置批處理設置决定的正整數；通過一組綫程處理單元平行執行該M個工作組中的每一個工作組中的複數個工作項的子集；以及響應該M個工作組中的一個或複數個工作組遇到同步障礙的發現，將該M個工作組的中間資料儲存至溢出記憶體中，並將該M個工作組中的複數個工作項的下一子集加載至該複數個綫程處理單元中用於平行執行。

本發明的計算設備和相應計算方法可以實現高效的資料訪問。

100、700‧‧‧計算設備

110‧‧‧計算單元

115‧‧‧綫程處理單元

116‧‧‧本地記憶體

120‧‧‧控制器

140‧‧‧記憶體緩衝器

150‧‧‧主機

170‧‧‧系統記憶體

180、380、480‧‧‧MSE

210‧‧‧指令解碼器

230‧‧‧暫存器檔案

282‧‧‧暫存器陣列

283‧‧‧輸入選擇器

310‧‧‧多工器

410‧‧‧輸入選擇器

600、900‧‧‧方法

610~640、910~930‧‧‧步驟

701‧‧‧編譯器

710‧‧‧計算單元

720‧‧‧批處理設置

730‧‧‧溢出記憶體

750‧‧‧控制器單元

760‧‧‧程式計數器

810‧‧‧工作項

820‧‧‧工作組

第1圖示出根據一個實施例之耦合到主機和系統記憶體的計算設備的示意圖。

第2圖示出了根據一個實施例之計算設備的進一步細節。

第3圖示出了根據一個實施例之MSE的示例。

第4圖示出了根據一個實施例之MSE的另一示例。

第5A圖和第5B圖根據一個實施例結合矩陣乘法示出了MSE的資料傳輸模式的示例。

第6圖是示出根據一個實施例之由計算設備執行的方法的流程圖。

第7圖示出了根據一個實施例之具有可配置數量的工作組分配的計算單元。

第8A圖和第8B圖示出根據一個實施例之工作組分配到批處理中的示例。

第9圖示出根據一個實施例之由具有用於平行計算的可配置工作組分配的計算設備執行的方法的流程圖。

在說明書及後續的申請專利範圍當中使用了某些詞彙來指稱特定的元件。所屬領域中具有通常知識者應可理解，製造商可能會用不同的名詞來稱呼同樣的元件。本說明書及後續的申請專利範圍並不以名稱的差異來作為區分元件的方式，而是以元件在功能上的差異來作為區分的準則。在通篇說明書及後續的請求項當中所提及的「包含」係為一開放式的用語，故應解釋成「包含但不限定於」。另外，「耦接」一詞在此係包含任何直接及間接的電氣連接手段。因此，若文中描述一第一裝置耦接于一第二裝置，則代表該第一裝置可直接電氣連接于該第二裝置，或透過其他裝置或連接手段間接地電氣連接至該第二裝置。

在訊號和圖像處理中以及在科學和工程的各個領域中經常出現的計算問題涉及濾波器的操作。例如，濾波器可以是移動平均(moving average)濾波器、加權移動平均濾波器、有限脈衝響應(finite impulse response，簡寫為FIR)濾 波器等。濾波器可以具有K個符號(tap)，其中K是大於1的整數，並且濾波器可以是應用於資料系列中的每個資料點(data point)。平行計算設備可以將濾波操作劃分為複數個任務或綫程，其中每個計算單元處理一個綫程的計算。因此，計算單元在下文中被稱為閘值處理單元。計算的一個特徵是大多數資料點在計算中重複使用。在常規系統中，具有N個綫程處理單元的平行計算設備執行用於訪問其源運算元(operand)的每通道(per-lane)記憶體訪問；例如，每個綫程處理單元可以從用於執行K個符號濾波操作的記憶體單獨訪問其K個源運算元。這種“每通道”訪問使得從記憶體重複訪問相同的資料點。對於相同的資料的重複記憶體訪問不僅在速度方面，而且在功耗方面同樣是低效的。

低效的資料訪問問題通常不僅發生在資料過濾中，而且在對象發現、一般的圖像處理、矩陣乘法以及大範圍的其他計算中都可以看到。在下文中，使用資料過濾作為示例來簡化描述和說明；然而，應當理解，本發明的實施例可應用於涉及通過平行操作的不同綫程處理單元重複使用相同資料元素的廣泛範圍的計算。

本發明的實施例提供了一種用於平行計算系統中的高效工作執行的系統和方法。以下描述包括兩個部分。第一部分描述了一種記憶體混移引擎(memory shuffle engine，簡寫為MSE)，其對於複數個綫程處理單元重複使用的資料元素實現高效的資料訪問。第二部分描述了用於設置由平行執行的可配置數量的工作組共享的程式計數器批處理的機制。

第1圖示出根據一個實施例之耦合到主機150和系統記憶體170的計算設備100的示意圖。計算設備100的圖示已經被簡化；但是應當理解，為了便於說明，計算設備100可以包括從第1圖中省略的更多的組件。計算設備100的示例包括但不限於圖形處理單元(GPU)、數位訊號處理器(DSP)、圖像處理器、專用應用指令集處理器(application-specific instruction set processor，簡寫為ASIP)等。DSP和ASIP執行訊號、圖像和/或多媒體處理操作。DSP和ASIP都可以編程。ASIP的示例是執行由系統支持的專門功能的專用硬體加速器；例如，編碼和解碼。GPU執行圖形處理任務；例如，創建3D場景的2D光罩表示。圖形處理可以被稱為3D圖形流水綫或渲染流水綫。3D圖形流水綫可以通過為了加速計算而定制的固定功能硬體和通用可編程硬體的組合來實現，以允許圖形渲染中的靈活性。通用可編程硬體也稱為著色器硬體(shader hardware)。除了渲染圖形之外，著色器硬體還可以執行通用計算任務。

計算設備100包括用於執行單指令-多資料(single-instruction-multiple-data，簡寫為SIMD)和/或單指令-多綫程(single-instruction-multiple-thread，簡寫為SIMT)操作的平行執行硬體。在一個實施例中，計算設備100包括一個或複數個計算單元110。每個計算單元110還包括一組N個綫程處理單元115，諸如算術邏輯單元(Arithmetic Logic Unit，簡寫為ALU)，用於對具有複數個資料集(例如，N個資料集)或綫程的相同指令執行平行計算。每個計算單元110還包括本 地記憶體116，用於存儲在相同工作組的工作項之間共享的資料。在一個實施例中，N個綫程處理單元115可以是來自同一工作組的分配任務；或者，N個綫程處理單元115可以是來自多於一個工作組的分配任務。

在一個實施例中，計算設備100還包括控制器120，也稱為記憶體混移控制器。控制器120的一個職責是控制從記憶體緩衝器140到綫程處理單元115的資料加載。記憶體緩衝器140可以位於本地記憶體116中；或者，記憶體緩衝器140可以位於系統記憶體170中，如第1圖中虛綫框所示。系統記憶體170可以是動態隨機存取記憶體(dynamic random access memory，簡寫為DRAM)或其他揮發性或非揮發性記憶體，並且通常是晶片外的；即，在與計算設備100不同的晶片上。相反，本地記憶體116在晶片上；也就是說，它與計算設備100在相同的晶片上。本地記憶體116的示例是靜態隨機存取記憶體(static random access memory，簡寫為SRAM)。其他揮發性或非揮發性記憶體也可以用作本地記憶體116。

在一個實施例中，計算設備100還包括記憶體混移引擎(MSE)180。MSE 180支持涉及使用資料陣列用於平行計算的大範圍操作。MSE 180用作記憶體緩衝器140和綫程處理單元115之間的中間設備。響應於來自控制器120的控制訊號，MSE 180可從記憶體緩衝器140加載資料陣列，並將資料陣列的適當資料元素發送到每個綫程處理單元115。關於MSE 180的進一步細節將提供如下。

在一個實施例中，主機150可以包括一個或複數個 中央處理單元(CPU)。主機150可以向計算設備100發出命令以指示計算設備100執行平行計算。在一些實施例中，計算設備100和主機150可以集成到晶片上系統(system-on-a-chip，簡寫為SoC)平臺中。在一個實施例中，SoC平臺可以是移動計算和/或通信設備(例如，智慧手機、平板電腦、膝上型計算機、遊戲設備等)、桌面計算系統、服務器計算系統或雲計算系統的一部分。

第2圖示出了根據一個實施例之計算設備100的進一步細節。在該實施例中，示出了計算設備中的一個計算單元110。應當理解，計算設備100可以包括任何數量的計算單元110。指令解碼器210解碼指令以供綫程處理單元115執行。控制器120從記憶體緩衝器140加載指令的源運算元到綫程處理單元115中，其中記憶體緩衝器140可以是晶片外(例如，在系統記憶體170中)或在晶片上(例如，如第2圖所示的本地記憶體116中)。一些指令可以在暫存器檔案230中具有源運算元(一些指令的源運算元可以放在暫存器檔案230中)。在一個實施例中，響應於導致資料陣列訪問(例如，濾波操作指令)的指令，控制器120使用兩個步驟加載源運算元：首先，將源運算元從記憶體緩衝器140加載到MSE 180中，然後將源運算元從MSE 180加載到綫程處理單元115中。因為使用MSE 180去除了對記憶體緩衝器140中的相同資料元素的重複訪問，經由MSE 180的記憶體訪問比每通道記憶體訪問更高效。

在一個實施例中，MSE 180在控制器120的控制下把將在濾波操作中使用的所有資料元素從記憶體緩衝器140加 載到MSE 180中。例如，MSE 180可被用於3個符號(3-tap)濾波操作，其被公式化為：filter_output(i)=(d(i-1)+d(i)+d(i+1))/3，其中i=1，並且N是被分配以平行執行濾波操作的綫程處理單元115的數量。MSE 180將所有(N+2)資料元素d(0)、d(1)、...、d(N+1)加載到其內部正反器(flip-flop)(即MSE暫存器陣列282(也稱為“暫存器”))。當綫程處理單元115準備好執行濾波操作時，MSE 180在三個時刻將三個連續的資料元素作為源運算元輸出到每個綫程處理單元115。

在一個實施例中，對於K個符號濾波操作，MSE 180可將將在濾波操作中使用的所有資料元素加載到其內部暫存器陣列282中。儘管K可以是任何正整數，但實際上，K的數量受到MSE 180中的暫存器282的數量以及暫存器陣列282和輸入選擇器283之間的連接數量的限制。將參考第3圖和第4圖提供輸入選擇器的更多細節。

第3圖示出了根據一個實施例之MSE 380的示例。MSE 380是第1圖和第2圖的MSE 180的示例。在該示例中的MSE 380支持K個符號濾波器的操作，其中K=3。應當理解，K=3被用於簡化描述和說明的非限制性示例。還應當理解，MSE 380可以包括未示出以簡化圖示的附加電路。

在該示例中，MSE 380包括用於存儲(N+2)個資料元素的(N+2)個MSE暫存器282和用於將資料元素輸出到相應的N個綫程處理單元115的一組N個多工器310。多工器310是用於3個符號濾波器的第2圖的輸入選擇器283的示例。每個多工器310具有連接到暫存器282中的三個的三條輸入綫和 連接到一個綫程處理單元115的一條輸出綫。更具體地，每個多工器310可以經由其輸入綫從暫存器陣列282的相應子集將資料元素經由其輸出綫傳遞到其對應的綫程處理單元115。例如，多工器M1可將來自暫存器陣列282的對應子集(包括R0、R1和R2)的資料傳送到其對應的綫程處理單元115，且多工器M2可將來自暫存器陣列282的對應子集(包括R1、R2和R3)的資料傳送到其對應的綫程處理單元115。因此，兩個相鄰多工器310可以重複使用相同的輸入資料元素。例如，M1和M2可以重複使用存儲在R1和R2中的資料元素。

請注意，多工器310是輸入選擇器的示例。多工器310可根據第一控制訊號(例如，M1從R0接收第一資料元素)經由其輸入綫接收來自暫存器陣列282的對應子集的資料，也可以根據第二控制訊號(例如，M1從R1接收第二資料元素)經由其它輸入綫接收來自耦合到另一多工器310的暫存器陣列282的另一子集的資料。可以生成額外的控制訊號以控制進一步的資料傳輸。在一個實施例中，這些控制訊號由控制器120產生。例如，控制器120可以在三個時刻發送第一、第二和第三控制訊號，以分別選擇每個多工器310的第一、第二和第三輸入綫，以將三個不同的資料元素傳送到多工器的對應綫程處理單元115。以此方式，可選擇每一多工器310的所有三個輸入綫，且可將三個對應暫存器282中的資料元素發送到對應綫程處理單元115中。

當將多工器310的集合作為整體考慮時，第一控制訊號自暫存器282選擇N個連續資料元素(例如d(0)、d(1)、...、 d(N-1))的第一序列，並將它們發送到N個綫程處理單元115，其中每個綫程處理單元115接收第一序列中的資料元素之一。類似地，第二控制訊號自暫存器282選擇N個連續資料元素(例如，d(1)、d(2)、...、d(N))的第二序列，並且將它們發送到N個綫程處理單元115，其中每個綫程處理單元115接收第二序列中的資料元素之一。第三控制訊號自暫存器282選擇N個連續資料元素(例如，d(2)、d(3)、...、d(N+1))的第三序列，並且將它們發送到N個綫程處理單元115，其中每個綫程處理單元115接收第三序列中的資料元素之一。在該示例中，第一序列和第二序列被移位一個資料元素，並且類似地，第二序列和第三序列被移位一個資料元素。也就是說，在該示例中，每個序列是其先前序列(通過移位一個資料元素)的移位版本。在替代實施例中，每個資料序列可以從其緊鄰的前導(immediate predecessor)資料序列移位多於一個資料元素。

從該示例可以看出，MSE 380的使用顯著減少了記憶體訪問量。在該示例中，(N+2)個資料元素d(0)、d(1)、...、d(N+1)從記憶體(即，本地記憶體116或系統記憶體170)加載到MSE 380中一次，且(N-1)個資料元素d(2)、d(3)、...、d(N)的子集由綫程處理單元115重複使用。如果每個綫程處理單元310對其用於3個符號濾波操作的源運算元執行每通道記憶體訪問，則這些資料元素d(2)、d(3)、...、d(N)中的每一個將從記憶體讀取三次。相反，MSE 380將(N+2)個資料元素的整個資料序列加載到其暫存器282中一次，而不 重複讀取相同的資料元素。MSE 380隨後可以經由其多工器310高效地從其暫存器282向綫程處理單元發送適當的源運算元。

在第3圖的實例中，MSE 380包括比用於在資料陣列的邊界(即，開始和結束)保持兩個額外資料元素的多工器310的數目多兩個暫存器282。在替代實施例中，暫存器282的數量可以與多工器310的數量相同或比多工器310的數量多(不同於2的預定數量)。這些額外數目的暫存器282(如果有的話)影響可在濾波操作中使用的符號的數目、可從一個多工器310移位到下一個的資料元素的數目以及可由綫程處理單元115執行的其它操作。

第4圖示出了根據一個實施例之MSE 480的另一示例。MSE 480包括耦合到輸入選擇器410的陣列的MSE暫存器282的陣列。在一個實施例中，每個輸入選擇器410包括至少一個多工器，例如第3圖所示的多工器310。在替代實施例中，每個輸入選擇器410包括複數個多工器、開關或路由元件以選擇來自暫存器282的相應子集的輸入，並將這些輸入傳送到對應的綫程處理單元115。

在第4圖的實施例中，指令解碼器210解碼指令並將關於已解碼指令的信息傳遞到控制器120。根據該信息，控制器120產生一個或複數個控制訊號至輸入選擇器410，以控制它們對資料陣列的操作，包括但不限於：移位、混移，選擇和傳遞。為了簡化圖示，第4圖僅示出作為控制訊號的接收方的最左邊的輸入選擇器410；但是應當理解，所有輸入選擇器410直接或間接地從控制器120接收控制訊號。

在一個實施例中，每個輸入選擇器410經由輸入綫耦合到暫存器陣列282的相應子集(例如，輸入選擇器S1耦合到RA1，輸入選擇器S2耦合到RA2等)用於接收輸入。在第4圖中，劃分暫存器陣列282的虛綫表示一些子集的邊界。此外，每個輸入選擇器410還耦合到其相鄰輸入選擇器410中的一個或兩個。兩個相鄰輸入選擇器410之間的連接(稱為“選擇器間連接”)可以是單向的或雙向的。因此，每個輸入選擇器410可以不僅經由輸入綫從其對應的子集接收資料輸入，而且還經由選擇器間連接從暫存器陣列282的其他子集接收資料輸入。例如，輸入選擇器S1可以從RA1以及RA2、RA3、...等接收輸入。

在一個實施例中，響應於傳播控制訊號，每個輸入選擇器410可以將資料元素從第一相鄰輸入選擇器410傳遞到第二相鄰輸入選擇器410(例如，資料可以從S3經由S2到S1，或從S4經由S2和S3到S1)。傳播控制訊號可以由控制器120或輸入選擇器410中指定的一個產生。

每個輸入選擇器410還經由輸出綫耦合到綫程處理單元115的子集。在第4圖中，每個輸出綫可以由輸入選擇器410用於將輸入資料傳送到對應的綫程處理單元115。輸入資料可以從所選輸入綫或從所選擇的選擇器間連接接收。根據來自控制器120的控制訊號，輸入選擇器410可以通過輸入綫或選擇器間連接從暫存器中選擇資料元素，並將該資料元素傳送到綫程處理單元115中的一些或全部，其中綫程處理單元115連接到其輸出。或者，輸入選擇器410可通過輸入綫和/或選擇器間連接的組合從不同暫存器282中選擇不同的資料元素，並將不同 的資料元素傳送到連接其輸出的不同綫程處理單元115。儘管第4圖中示出了特定數量的輸入/輸出綫和選擇器間連接，但是應當理解，在各種實施例中，每個輸入選擇器410可以連接到任何數量的輸入/輸出綫和選擇器間連接。

在一個實施例中，控制器120可以生成到輸入選擇器410的一個或複數個控制訊號，以控制它們對資料陣列的操作，包括但不限於：移位、混移、選擇和傳遞。在一些實施例中，控制器120可以在由暫存器282的數量、輸入綫的數量和選擇器間連接的數量所强加的限制內控制移位的方向和量，使得輸出資料陣列是輸入資料陣列的移位版本。控制器120還可以控制施加在輸入資料陣列上的混移模式，以及選擇哪些輸入綫和/或選擇器間連接以使能到輸出的資料傳輸。控制器120可以進一步命令輸入資料陣列通過輸入選擇器410，使得輸出資料陣列與輸入資料陣列相同。

在第4圖中，以粗綫示出了輸入綫和選擇器間連接中的一部分，以指示這些綫/連接被選擇或啟用。這些粗綫示出了將輸入資料陣列向左移位一個資料元素以產生輸出資料陣列的示例。假定輸入資料陣列是(d(0)、d(1)、d(2)、d(3)、...)，並且每個輸入綫將輸入資料陣列的一個資料元素傳送到輸入選擇器410。作為資料移位的結果，由輸入選擇器410生成的輸出資料陣列是(d(1)、d(2)、d(3)、d(4)...)。如第3圖的示例中所示，K個符號濾波操作可以通過移位操作和通過(pass-through)操作的組合來實現。在資料移位的示例中，傳播控制不被使能。

第5A圖和第5B圖根據一個實施例結合矩陣乘法示出了MSE 480的資料傳輸模式的示例。第5A圖示出了被乘數從所選輸入綫(示出為連接到S1的粗輸入綫)到所有輸入選擇器410的輸出綫的傳播。例如，如果被乘數是(d(0)、d(1)、...、d(7))，輸入選擇器S1可以通過其所選擇的輸入綫接收這些被乘數，並且相同的被乘數可以通過S1到達其輸出綫，並傳播到其他輸入選擇器410以到達它們各自的輸出綫。在乘法的不同階段，可以選擇不同組的輸入綫。第5B圖示出了乘法器的傳播：每個乘法器從輸入選擇器410的一個選定輸入綫(示為粗輸入綫)傳送到該輸入選擇器410的所有輸出綫。在乘法的不同階段，可以選擇每個輸入選擇器410的不同的行。

除了上述濾波操作和矩陣乘法之外，MSE(例如，第1圖和第2圖的MSE 180，第3圖的MSE 380和第4圖的MSE 480)可以具有其他訊號、圖像、多媒體處理應用，包括但不限於計算機視覺中的對象發現和模式識別。一種用於對象發現的已知技術將滑動窗應用於圖像以檢測給定對象(例如，人)的存在。滑動窗口可以具有固定尺寸，並且可以在圖像的任何方向上滑動，諸如水平、垂直或對角綫方向。滑動窗口的每個實例(即，圖像的特定位置處的滑動窗口)可以被分配給綫程處理單元115進行處理，並且滑動窗口的不同實例可以被分配給平行地不同的綫程處理單元115。在一個實施例中，當滑動窗口在一個方向上滑動以創建滑動窗口實例的序列時，連續滑動窗口實例具有大量重叠的資料元素。因此，提供給第一綫程處理單元的資料元素(源運算元)可以由第二綫程處理單元重 複使用。MSE可以用於向綫程處理單元115提供資料元素，包括重叠的資料元素，而不具有由綫程處理單元115的每通道訪問引起的重複的相同資料訪問。

例如，梯度方向直方圖(histogram of oriented gradients，簡寫為HOG)是在計算機視覺和圖像處理中使用的用於對象發現的目的的特徵描述符。該技術對圖像的局部部分中的梯度定向的發生進行計數。圖像被分成稱為單元的小連接區域。對於每個單元內的像素，編譯梯度方向的直方圖。描述符是這些直方圖的級聯。HOG計算的第一步是梯度值的計算。一種用於計算梯度值的常用方法是在水平和垂直方向中的一個或兩個方向上應用離散導數掩模(discrete derivative mask)。該掩模是濾波掩模，其被應用於圖像單元以過濾圖像的顏色或强度資料。因此，類似於對於K個符號濾波器的上述濾波操作，MSE也可以用於HOG計算的梯度值的計算。

如前所述，計算設備100響應主機150或計算設備100的指令集架構(instruction set architecture，簡寫為ISA)中定義的指令來執行操作。在一個實施例中，濾波器指令指定指示資料陣列的基本地址的第一運算元，指示資料陣列的大小的第二運算元，以及指示訪問順序(例如，綫性、塊、3D等)的第三運算元。以下偽代碼提供了具有3個符號濾波器的濾波操作指令的示例：//定義mse緩衝區大小_local float mse[workgroup_size+2]；*mse=(image[f(workgroup_ID)]，workgroup_size+2， linear)；//定義如下：//1. 由於訪問是基於工作組的，因此以workgroup_ID作為在記憶體中的起始地址//2. 訪問大小//3. 訪問順序(綫性/塊/3D/...)//使用mse計算結果float result=(mse[woki_id]+mse[woki_id+1]+mse[woki_id+2])/3；在一個實施例中，ISA還為可以利用上述MSE進行對象發現、矩陣乘法等的其他操作定義指令。

第6圖是示出根據一個實施例之由計算設備100執行的方法600的流程圖。在一個實施例中，方法600可以由平行計算系統或設備(例如第1圖的計算設備100)執行。在第3圖和第4圖中提供了輸入選擇器的示例；也可以使用其他類型的輸入選擇器。

方法600開始於計算設備100接收識別資料元素的陣列作為輸入的指令(步驟610)。計算設備100響應於指令將資料元素的陣列從記憶體緩衝器加載到暫存器的陣列中(步驟620)。暫存器陣列耦合到輸入選擇器陣列，並且每個輸入選擇器經由輸入綫耦合到暫存器的相應子集。

根據第一控制訊號，每個輸入選擇器將至少第一資料元素從暫存器的相應子集傳送到一個或複數個對應的綫程處理單元(步驟630)。根據第二控制訊號，每個輸入選擇 器將來自暫存器的另一子集的至少第二資料元素傳送到一個或複數個對應的綫程處理單元(步驟640)，該另一子集經由其他輸入綫耦合到另一輸入選擇器。如上該，不僅從輸入選擇器的相應子暫存器集合，而且從另一輸入選擇器的相應子暫存器集合選擇輸入的能力促進了資料重用，並減少了從記憶體重複加載相同的資料的不必要的記憶體流量。

在上面的描述中，假設N個綫程處理單元被分配以平行執行濾波操作。在OpenCL中，分配給每個綫程處理單元的計算任務稱為工作項，並且共享資料和同步障礙的相關工作項形成工作組。通常，同一工作組的工作項使用相同的程式計數器(PC)並形成單批處理；因此，這些工作項將在鎖步中逐步(step through)執行程式的指令。在某些情況下，工作組可能包含比計算單元中的綫程處理單元數少的工作項。或者，工作組可以包含比計算單元中的綫程處理單元數量更多的工作項，並且這些工作項的數量不能被綫程處理單元的數量均分(evenly divisible)。在傳統系統中的任何這些情況下，不能與工作項匹配的綫程處理單元將被保留未使用。當平行計算中未使用某些綫程處理單元時，將浪費系統資源。

根據本發明的一個實施例，計算單元中的綫程處理單元被劃分為複數個批處理。分配到批處理的工作組數量(M)是可配置的。假設計算單元中的綫程處理單元數為N，批處理(batch)中的綫程處理單元數為P。M的取值範圍為1MP。對於M個工作組中的每個工作組，工作組中可平行執行的工作項數為(P/M)，小於工作組的大小(即工作組中 工作項的總數)。因此，它需要多次迭代來處理工作組中的所有工作項。

第7圖示出了根據一個實施例之具有可配置數量的工作組分配的計算單元710。作為計算設備700的一部分的計算單元710是第1圖的計算單元110的一個示例。根據第7圖，計算單元710包括附加的和/或替代的元件，其在第1圖的實施例中未示出。在該實施例中，計算單元710支持複數個PC批處理和每個批處理中的可配置數量的工作組。計算單元710包括控制器單元750，在一個實施例中，控制器單元750可以執行第1圖的控制器120的操作。附加地或替代地，控制器單元750可以根據編譯器701所决定的批處理設置720將M個工作組分配給批處理。在下面的描述中，跨不同批處理使用相同的可配置值M。在替代實施例中，不同批處理可以用不同數量的工作組分配。

在一個實施例中，可以在編譯時决定可配置數量M(如第7圖所示)，或者在任務調度時决定可配置數量M。例如，編譯器701可以分析指示要執行的任務的維度的輸入NDRange，並且决定哪些工作組可以被打包到同一批處理中。可以使用的一個標準是具有相同同步障礙的工作組可以被打包到同一批處理中(即可配置批處理設置是從複數個工作組的編譯時的分析决定的)。如果M個工作組被打包到一個批處理中，則這些工作組可以根據相同的程式計數器760(第7圖中分別標示為CP0和PC1)執行指令。

為了說明的簡單起見，第7圖在計算單元710中 僅示出了一個計算單元710和兩個批處理。應當理解，可以使用多於一個計算單元710和多於兩個批處理來在一個輸入中處理所有的工作組。每個批處理包括由批處理中的所有工作組、P個綫程處理單元115、溢出記憶體730和本地記憶體116共享的程式計數器760。屬同一工作組的工作項可以使用本地記憶體116用於共享資料。在一個實施例中，控制器單元750管理用於計算單元110中的不同批處理和批處理中的不同工作組的本地記憶體116的分配和訪問。

屬同一工作組的工作項也可以使用溢出記憶體730來存儲中間資料或上下文，例如當遇到同步障礙並且暫時中止(例如，通過進入等待狀態)時。當該工作組中的所有工作項達到同步障礙時，等待狀態結束。在一個實施例中，當批處理中的給定工作組遇到同步障礙時，給定工作組中的工作項可以被保存在溢出記憶體730中，並且給定工作組中的其餘工作項將被處理。假設分配給給定工作組的綫程處理單元115的數量等於K(其中K=P/M)，在一個實施例中，在給定的工作組可以移動通過同步障礙之前，K個綫程處理單元115可循環遍歷給定的工作組中所有剩餘的工作項。在替代實施例中，批處理中的所有P個綫程處理單元115可以專用於處理給定工作組中的所有其餘工作項，以便整個給定工作組快速遇到同步障礙。由於同步障礙而暫停的批處理中的工作項可以被暫時保存在溢出記憶體730中。

第8A圖和第8B圖示出根據一個實施例之工作組分配到批處理中的示例。在這些示例中，四個工作組820 (WG0、WG1、WG2、WG3)被分配給兩個批處理(批處理0和批處理1)，其中每個工作組820包含複數個工作項810。每個批處理包括兩個工作組820，並且具有平行處理四個工作項的能力。第8A圖示出了平行處理的前四個工作項目，第8B圖示出了平行處理的下四個工作項目。在一些實施例中，不同工作組中的工作項可以是不同步的；因此，在一些實施例中，不同的工作組可以以不同的步速(pace)前進。然而，由於共享程式計數器，同一批處理中的工作組以相同的步速前進。

第9圖示出了根據一個實施例之由具有用於平行計算的可配置工作組分配的計算設備執行的方法900的流程圖。在一個實施例中，方法900可以由平行計算系統或設備(諸如第7圖的計算設備700)執行。

方法900開始於控制器單元750將工作組分配給批處理的集合(步驟910)。至少一個批處理(“給定批處理”)被分配了共享程式計數器的M個工作組，M是根據可配置批處理設置决定的正整數。給定批處理包括一組綫程處理單元。該組綫程處理單元平行地執行M個工作組中的每個工作組中的工作項的子集(步驟920)。響應於M個工作組中的一個或複數個遇到同步障礙的發現，M個工作組的中間資料被存儲在溢出記憶體中，並且M個工作組中的下一個工作項的子集被加載到綫程處理單元中以平行執行(步驟930)。

前文已經描述了平行計算系統。平行計算系統可以克服常規硬體平臺的限制，以實現高效的資料訪問和重用，以及工作組的靈活分配。因此，可以提高系統的總體能量和計 算效率。

已經參考第1、3、4圖和第7圖的示例性實施例描述了第6圖和第9圖的流程圖的操作。然而，應當理解，除了參考第1、3、4圖和第7圖討論的那些實施例實施第6圖和第9圖的流程圖的操作之外，也可以參考第1、3、4圖和第7圖討論的實施例執行不同於這些流程圖的操作。雖然第6圖和第9圖的流程圖示出了本發明的某些實施例執行的操作的特定順序，但是應當理解，這種順序是示例性的(例如，替代實施例可以以不同順序執行操作，組合某些操作，重叠某些操作等)。

以上所述僅為本發明之較佳實施例，凡依本發明申請專利範圍所做之均等變化與修飾，皆應屬本發明之涵蓋範圍。

100‧‧‧計算設備

110‧‧‧計算單元

115‧‧‧綫程處理單元

116‧‧‧本地記憶體

120‧‧‧控制器

140‧‧‧記憶體緩衝器

150‧‧‧主機

170‧‧‧系統記憶體

180‧‧‧MSE

Claims (22)

1. 一種計算設備，用於執行複數個平行計算，其中該計算設備包含：一組綫程處理單元；以及一記憶體混移引擎，耦接於該組綫程處理單元，該記憶體混移引擎包含：複數個暫存器組成的暫存器陣列，儲存自一記憶體緩衝器獲得的複數個資料元素的一陣列；以及複數個輸入選擇器組成的輸入選擇器陣列，每一輸入選擇器經由複數個輸入綫耦接於該複數個暫存器的對應子集，並經由一個或複數個輸出綫耦接於一個或複數個對應綫程處理單元，其中，依據一第一控制訊號，每一輸入選擇器自該複數個暫存器的該對應子集發送至少第一資料元素至該一個或複數個對應綫程處理單元，以及依據一第二控制訊號，每一輸入選擇器自該複數個暫存器的另一子集發送至少第二資料元素至該一個或複數個對應綫程處理單元，其中該另一子集經由其他複數個輸入綫耦接於另一輸入選擇器。
2. 如申請專利範圍第1項所述之計算設備，其中每一輸入選擇器包含一多工器，經由複數個輸入綫耦接於該複數個暫存器的該對應子集，並經由一個輸出綫耦接於一個綫程處理單元。
3. 如申請專利範圍第1項所述之計算設備，其中每一輸入選擇器更經由一個或複數個選擇器間連接耦接於至少一個相鄰 輸入選擇器，並自該至少一個相鄰輸入選擇器接收一個或複數個該資料元素。
4. 如申請專利範圍第1項所述之計算設備，其中響應一傳播控制訊號，每一輸入選擇器更自一第一相鄰輸入選擇器傳送資料元素至一第二相鄰輸入選擇器。
5. 如申請專利範圍第1項所述之計算設備，其中每一輸入選擇器自複數個不同暫存器選擇複數個不同資料元素，並將該複數個不同資料元素傳送至複數個不同綫程處理單元。
6. 如申請專利範圍第1項所述之計算設備，其中每一輸入選擇器自該複數個暫存器選擇一個資料元素，並將該資料元素傳送至複數個不同綫程處理單元。
7. 如申請專利範圍第1項所述之計算設備，其中更包含：一記憶體混移控制器，利用一個或複數個控制訊號控制該輸入選擇器陣列，以執行包含移位、混移、選擇以及傳遞該複數個資料元素的陣列的至少一個操作。
8. 一種計算設備，用於執行複數個平行計算，其中該計算設備包含：一控制單元，分配複數個工作組至一組批處理；以及該組批處理，耦接於該控制單元，每一批處理包含：一程式計數器，由分配給該批處理的M個工作組共享，其中M是依據可配置批處理設置决定的正整數；一組綫程處理單元，平行執行該M個工作組中的每一個工作組中的複數個工作項的子集；以及一溢出記憶體，當該M個工作組中的一個或複數個工作組遇 到同步障礙時，儲存該M個工作組的中間資料。
9. 如申請專利範圍第8項所述之計算設備，其中該可配置批處理設置是從該複數個工作組的編譯時的分析决定的。
10. 如申請專利範圍第8項所述之計算設備，其中該M個工作組中的每一個工作組包含比平行執行的該複數個工作項的該子集多的工作項。
11. 如申請專利範圍第10項所述之計算設備，其中當該複數個工作項的該子集完成或者暫停時，該組綫程處理單元更平行執行該M個工作組中的該複數個工作項的下一子集。
12. 一種計算方法，用於利用複數個綫程處理單元執行複數個平行計算，其中該計算方法包含：接收將複數個資料元素組成的一資料元素陣列識別為輸入的一指令；響應該指令，自記憶體緩衝器加載該資料元素陣列至複數個暫存器組成的一暫存器陣列，其中該暫存器陣列耦接於複數個輸入選擇器組成的一輸入選擇器陣列，且每一輸入選擇器經由複數個輸入綫耦接於該複數個暫存器的一對應子集；依據一第一控制訊號，每一輸入選擇器自該複數個暫存器的該對應子集發送至少一第一資料元素至一個或複數個對應綫程處理單元；以及依據一第二控制訊號，該每一輸入選擇器自該複數個暫存器的一另一子集發送至少一第二資料元素至該一個或複數個對應綫程處理單元，其中該另一子集經由其他複數個輸 入綫耦接於另一輸入選擇器。
13. 如申請專利範圍第12項所述之計算方法，其中該每一輸入選擇器包含一多工器，經由複數個輸入綫耦接於該複數個暫存器的該對應子集，並經由一個輸出綫耦接於一個綫程處理單元。
14. 如申請專利範圍第12項所述之計算方法，其中更包含：經由一個或複數個選擇器間連接，該每一輸入選擇器自至少一個相鄰輸入選擇器接收一個或複數個該資料元素。
15. 如申請專利範圍第12項所述之計算方法，其中更包含：響應一傳播控制訊號，該每一輸入選擇器自一第一相鄰輸入選擇器傳送資料元素至一第二相鄰輸入選擇器。
16. 如申請專利範圍第12項所述之計算方法，其中更包含：該每一輸入選擇器自複數個不同暫存器選擇複數個不同資料元素；以及將該複數個不同資料元素傳送至複數個不同綫程處理單元。
17. 如申請專利範圍第12項所述之計算方法，其中更包含：該每一輸入選擇器自該複數個暫存器選擇一個資料元素；以及將該資料元素傳送至複數個不同綫程處理單元。
18. 如申請專利範圍第12項所述之計算方法，其中更包含：利用一個或複數個控制訊號控制該輸入選擇器陣列，以使該複數個輸入選擇器對該資料元素陣列執行包含移位、混移、選擇以及傳遞操作的至少一個操作。
19. 一種計算方法，用於利用複數個綫程處理單元執行複數個 平行計算，其中該計算方法包含：分配複數個工作組至複數個批處理組成的一組批處理，其中該複數個批處理中的至少一個被分配給一共享程式計數器的M個工作組，其中M是依據一可配置批處理設置决定的一正整數；通過一組綫程處理單元平行執行該M個工作組中的每一個工作組中的複數個工作項的一子集；以及響應該M個工作組中的一個或複數個工作組遇到同步障礙的發現，將該M個工作組的中間資料儲存至一溢出記憶體中，並將該M個工作組中的複數個工作項的下一子集加載至該複數個綫程處理單元中用於平行執行。
20. 如申請專利範圍第19項所述之計算方法，其中更包含：從該複數個工作組的編譯時的分析决定該可配置批處理設置。
21. 如申請專利範圍第19項所述之計算方法，其中該M個工作組中的每一個包含比平行執行的該複數個工作項的該子集多的工作項。
22. 如申請專利範圍第21項所述之計算方法，其中更包含：當該複數個工作項的該子集完成或者暫停時，該組綫程處理單元平行執行該M個工作組中的該複數個工作項的該下一子集。