TW200844853A - Virtual architecture and instruction set for parallel thread computing - Google Patents

Description

200844853 九、發明說明： 【發明所屬之技術領域】 >本發明大體上涉及平行處理’且更特定而言涉及用於平 仃緒計算的虛擬架構及指令集。 【先前技術】 、平订處理中，多個處理單元（例如，多個處理器晶片 或單個晶片内的多個處理核心)同時操作以處理資料。可 使用此種系統來解決便於分解成多個部分的問題。-個實 !是影像過遽’其中根據輸入影像的某-數目的像素來計 2輸出影像的每個像素。每個輸出像素㈣算―般獨立於 -它所有輸出像素，因此不同的處理單元可平行計算不同 ^輸出像素。其它許多類型的問題也適於平行分解。總的 來祝，N向（N_way)平行執行可將此種問題的解決方案 大約N倍。 另-類問題在平行的執行緒可彼此協調的情況下則適於 =理。一個實例是快速傳立葉變換㈣，其為一種 以便產峰斩^ 個，.及對别一級的輸出執行計算， 直到、，新的值用於對下-級的輸入，如此繼續 線二為止。單個執行緒可執行多個級，只要該 又付來自則一級的輸出資料即可。 =間劃分任務’那麼必須提供某種協調機制= ⑴如）線緒不會試圖讀取 12月# 询貝抖。（在2005年 月15日Μ的共同轉讓、制待決的美國 "一。號中描述了這個問題的一個解決方案广案弟 128625.doc 200844853 =里系統的程式設計可能較為困難。通常要 2 :::貝知道可用的處理單元的數目及其能力(指 …貝料暫存器數目、互連件等)，以便 供相當大的幫二：的編譯器可在此方面提 不同處理器時重新編譯代碼。 ）到 」匕夕卜二千仃處理架構的各種態樣正在迅速發展 的各種態樣（例如，程式設計模型或指 式:又成下一代’必須也相應地改變應用程 二二：/扁譯器及其它軟體及工具。這種不穩定性可 銷：〜里代碼的開發及維護增加相當大的額外開 程之間協調時，平行程式設計變得更加困難。 可十貝必須確定在特定處理器或電腦系統中有何種 可用於支援（或仿直）绩法pq、s π田… 的機制，且必須編寫利用 機制的代碼。由於不同的電腦系統上的可用及/或. 佳機制-般不同’所以這種類型的平行代碼—般不可: 植；必㈣對其執行於的絲硬體平Η編寫平行代碼。 此外，除了為處理器提供可執行代碼’程式設計人 :須為:主，，處理器提供控制代碼，控制代碼協調各種處理 早π的#作’例如向每個處理單元指示要執行哪個程式 及處理哪些輸人資料。此控制代碼通常專用於特 理器及處理器間的通信協定’且如果要替換不同的主處理 128625.doc 200844853 器則通常必須重編寫控制代碼。 編譯及重新編譯平 隨著吁曾姑1 & 代碼中的困難會使使用者不願 譯的;的發展而升級其系統。因此，將需要使已編 平行應用程式及工具的目伊^體平口，並提供一種作為 集。 π的知疋的平行處理架構及指令 【發明内容】 本發明的實施例提供一重 ^ ^ 虛擬指令隼。户糍卓μ 、千仃、、者计异的虛擬架構及 r且古 擬千仃架構Μ:—種虛擬處理器，其支 杈具有不同虛揆缺夕Μ αα ^ 如，同牛…， 同級別”料共用及協調（例 同乂）的夕個虛擬緒的併發執行；以及一種虛 驅動器，其控制虛擬處理器。虛擬，理哭 仃 構用來定義虛擬緒的行為：人：二、’㈣曰令集架 眘祖為 包含與平行緒行為（例如， 二/、及同步）有關的指令。藉由使用虛擬平行平台， 的=開發出其，虛擬緒併發執行以處理::料 八:;α私式可用高度可移植的中間形式儲存及 刀，遠形式例如為猫準虛擬平行平台的程式代碼 裝時間或執行時間，硬體特定的虛 女 動哭姑士叩 洲年為及執行驅 -吏：間形式的應用程式代碼適應於其執 用程式更力一更容易開發，因為= 心私獨立於特定處理硬體。 根據本發明的一個態樣，一種用 方法包含提供第-程式代碼，第-程式代碼:=: 作的虛擬绪的陣❹的若干虛擬緒中的每—者 128625.doc 200844853 f列。將第—程式代碼編譯成虛擬緒程式，虛擬緒程式定 要針對陣列中的代表性虛擬緒執行的每線緒指令的序 列且母線緒指令的序列包含至少一個指令，該指令定義 f虛擬、、者與陣财的—或多個其它虛擬緒之間的協作 仃為冑存虛擬緒程式(例如，儲存在記憶體中或碟片 )可P現後將虛冑緒程式轉言睪成符合目# +台架構的 指令序列。 匕外也可提供第二程式代碼以定義協作的虛擬緒的陣 列，陣列適合於處理輸入資料集以產生輸出資料集，其中 陣列中的每個虛擬緒併發執行虛擬緒程式。有利地將第二 Z式代碼轉換成虛擬函數庫中的函數則序列，其中庫包 3初始化及致使執行協作的虛擬緒的陣列的虛擬函數。也 …子k個函數凋用序列。接著，可將所儲存的虛擬緒程 式及函數調用序列轉譯成可在目標平台架構上執行的程式 代碼’其中可執行的程式代碼定義一或多個執行協作的虛 Ο

擬緒的陣列的平a綠# _ , A 口、、、、、。可在付&目標平台架構的電腦系 執行可執行的程式代碼，因而產生輸出資料集，可將 輸出資料集儲存在儲存媒體（例如，電腦記憶體、碟片 中。 7 々已提到的，虛擬緒程式代碼巾的每線緒指令的序列有 利地包含至少-個指令’該指令定義代表性虛擬緒與陣列 中的-或多個其它虛擬緒之間的協作行為。舉例而々，每 線緒指令的序料包含：在序財的特定點處暫停執行代 表性虛擬緒的操作直_如其它虛擬緒中的—或多者達到 128625.doc 200844853 特定點為止的時間的指令’使代表性虛擬緒將資料儲存在 其它虚擬緒中的-或多者可存取的共用記憶體中；;=在 性虛擬緒自動讀取及更新儲存在其它虛擬緒中的一 或夕者可存取的共用記憶體中的資料的指令，等等 =程=包含變數定義語句’語句在若干虛擬狀 二間之-中u變數，其中不同的虛擬狀態空間對應於 虛极緒之間的不同的資料共用模式。在_個實施例中 =支援每線緒不共用模式及全局共用模式。在其它實施例 “：也可支板額外模式，例如虛擬緒的—個陣列内的共用 杈式及/或複數個虛擬緒陣列之間 根據本發明的另'態樣，一種用於操:、二 法包含提供輸入程式代碼。輸入程式代碼包含第一部分， 該第-部分定義要對虛擬緒陣列中的若干虛擬緒中的:一 者執行的操作序列，該陣列適合於處理輪入資料隹 ==集，且輸入程式代碼還包含第二部分，：第二部 刀疋義虛擬緒的陣列的維度。將輸入程式代碼的第八 編譯成虛擬緒程式，虛擬緒程式定義要對陣列中的代= 虛擬緒執行的每線緒指令的序列。每線緒指令的序列包又人 至少-個指令’該指令定義代表性虛擬緒與陣列中的—I 多個其它虛擬緒之間的協作行為。將輸人程式代碼的第二 部分轉換成對虛擬函數庫的函數調用序列，其中該庫勺入 初始化及致使執行協作的虛擬緒的陣列的虛擬函數。= 擬緒程式及函數則序列轉譯成可在目標平台架構 亚 的程式代碼，其中可執行的程式代碼定義執行協作的虛擬丁 128625.doc 200844853 =::Γ:一或多個真實線緒。在符合目標平台架構的電 將幹出4 執行程式代碼，因而產生輸出資料集，可 將輸出貝料集儲存在儲存媒體中。 绪=實施例中’可在兩個或兩個以上維度上定義虛擬 =:::::^_二部分還可—義 口的或夕個維度的函數調用，其中柵袼 中的母個陣列均將被執行。 可使用任何目標平台架構。在—些實施例中，目標平台 木2包含主處判及共處理器。在轉譯期間，可將虛擬緒 程式轉譯成可由定義於共處理器上的若干線緒平行執行的 私式代碼’同時將函數調用序列轉譯成對在主處理器上執 行的用於共處理器的驅動程式的調用序列。在其它實施例 中’目標平台架構包含中央處理單元（CPU)。在轉譯期 2，將虛擬緒程式及函數調用序列的至少一部分轉譯成目 標程式代碼，所述目標程式代碼使用少於虛擬緒數目的數 目的CPU線緒來執行虛擬緒陣列。 根據本發明的又一實施例，一種用於操作目標處理器的 方法包含獲得虛擬緒程<，虛.疑緒程式定義要針對虛擬緒 陣列中的若干虛擬緒中的代表性虛擬緒執行的每線緒指令 序列，虛擬緒陣列適合於處理輸入資料集以產生輸出資料 集。每線緒指令序列包含至少一個指令，該指令定義代表 性虛擬緒與陣列中的一或多個其它虛擬緒之間的協作行 為。也獲得定義虛擬緒陣列的維度的額外程式代碼。將虛 擬緒程式及額外的程式代碼轉譯成可在目標平台架構上執 128625.doc 200844853 行的程式代碼，其中 列的一或多個平A錄 —的轾式代碼定義執行虛擬緒陣 上執行可執行的程式代石馬在^合目標平台架構的電腦系統 資料集儲存在記憶體中；而產生輸出資料集且將輸出 在-些實施例中’可藉由接收 的源程式代碼並編譯 纟-“吕、編寫 得虛擬緒程式。或者，”代碼以產生虛擬緒程式來獲 由網路從遠端電腦系统=存媒體頃取虛擬緒程式或經 或接收的虛擬緒代石馬可能之程式。應瞭解，所讀取 直接建立為符合虛擬指令集架構的代碼。 次者 以下詳細描述連同附圖—起將提 點的更好瞭解。 知π曰]注貝及優 【實施方式】 八：發:實繼供一種用於平行緒計算的趣 U。虛擬錢提供支援在不同緒之間具有多層資 及協調（例如’同步）的多個緒的併發執行的處理詩型， 以及控制模型處理器的虛擬執行驅動器。用於定義處理绪 的行為的虛擬指令集包含與平行緒行為有關的指令，例如 允許某些線緒間的資料共用的指令及要求不同緒在程式内 的由程式設計人員指定的某些點處變得同步的指令。使用 虛擬平台’程式設計人員可間恭—好 貝r開發在其中執行併發、協作的 緒以處理資料的應用程式。硬體特定的虛擬指令翻譯器及 執行驅動器使應用代碼適應其執行於的特定硬體。因此， 應用程式更具移植性並更易於開發，因為開發過程獨立於 128625.doc 200844853 特定處理硬體。 ι·系統概述 圖1為根據本發明實施例的電腦系統100的方塊圖。電腦 系統100包含中央處理單元（CPU)102，及包含記憶體橋1〇5 的經由匯流排路徑通信的系統記憶體104。記憶體橋

Ο 1 05(其可以是（例如）Northbridge晶片）經由匯流排或其它通 信路徑106(例如，HyperTransport鏈路）連接到1/〇(輸入/輸 出）橋1 07。I/O橋1 〇7(其可以是（例如）s〇uthbridge晶片）從一 或多個使用者輸入裝置1 〇8(例如，鍵盤、滑鼠）接收使用者 輸入，並將輸入經由路徑1〇6及記憶體橋1〇5轉遞到cpu 1〇2。平行處理子系統112經由匯流排或其它通信路徑 113(例如，PCI Express或加速圖形埠鏈路）耦接到記憶體 橋105 ;在一個實施例中，平行處理子系統112是圖形子系 統，其將像素遞送到顯示器裝置11〇(例如，習知的基於 CRT或LCD的監視器）。系統碟片114也連接到㊈橋⑼。 切換器11 6提供I / 0橋1 〇 7與例如網路配接器i i 8及各種内插 卡120及121等其它組件之間的連接。其它組件（未明確展 示）包含刪或其它埠連接、⑶驅動器、DVD驅動器等， 也可連接到I/O橋107。使圖1中的 一 κ 口 1T的各個組件互連的通信路 徑可使用任何適宜的協定爽恭 疋不汽施，例如PCI(周邊組件互 連）、PCI ExpreSS(PCI-E)、Arw Λ 土 ’ AGP(加速圖形埠）、 HyperTransport，或任何其它 匕匯流排或點對點通信協定， 且不同裝置之間的連接可使用 — 使用此項技術中已知的不同協 疋0 128625.doc 200844853 平行處理子系統112包含平行處理單元（ppu)122及平行 處理（PP)記憶體124，其可（例如）使用例如可程式設計處理 器、特殊應用積體電路（ASIC)及記憶體裝置等的一或多個 積體電路裝置來實施。PPU 122有利地實施包含一或多個 處理核心的高度平行處理器，所述處理核心的每一者能夠 併發執行大量（例如，數百個）緒。PPU 122可經程式設計 以執行一大批計算，包含線性及非線性資料變換、視訊及/ 或音訊資料的過濾、模型化（例如，應用物理定律來確定 物件的位置、速度及其它屬性）、影像渲染等。ppu 122可 將資料從系統記憶體104及/或PP記憶體124轉移到内部記 憶體中，處理所述資料，並將結果資料寫回到系統記憶體 104及/或pp記憶體124，在其中此類資料可由其它系統組 件（包含，例如，CPU丨02)存取。在一些實施例中，ppu 122是圖形處理器，其還可經組態以執行與以下操作有關 的各種任務：依據由CPU 1〇2及/或系統記憶體1〇4經由記 憶體橋105及匯流排U 3供應的圖形資料產生像素資料、與 PP記憶體124(其可用作圖形記憶體，包含（例如）習知圖框 緩衝^ )父互以儲存及更新像素資料、將像素資料遞送到 顯不态裝置110等。在一些實施例中，ρρ子系統112可包含 一個作為圖形處理器操作的PPU 122，及另一用於進行通 用计开的PPU 1 22。PPU可相同或不同，且每一 ppu可具有 其自身的專用PP記憶體裝置。 CPU 1 〇2作為系統i 〇〇的主處理器操作，其控制及協調其 匕系、、先、、且件的操作。具體而言，Cpu 1 〇2發布控制ppu 122 128625.doc -14- 200844853 的操作的命令。在一些實施例中，CPU 102將針對PPU ! 22 的命令流寫入到命令緩衝器，命令緩衝器可處於系統記憶 體104、PP記憶體124或CPU 102及PPU 122兩者均可存取 的另一儲存位置中。PPU i22從命令緩衝器讀取命令流， 並與CPU 102的操作不同步地執行命令。 將瞭解’本文展示的系統是說明性的，且可能作出變化 及修改。可視需要修改連接拓撲，包含橋的數目及組態。 舉例而言，在一些實施例中，系統記憶體1〇4直接而不藉 由橋連接到CPU 102，且其它裝置經由記憶體橋1〇5及cpu I 02與系統記憶體1 〇4通信。在其它替代拓撲中，pp子系統 II 2連接到I/O橋1〇7而不連接到記憶體橋1〇5。在另外其它 實施例中，I/O橋1 07及記憶體橋！ 05可能整合到單一晶片 中。本文展示的特定組件是任選的；舉例而言，可能支援 任何數目的内插卡或周邊裝置。在一些實施例中，排除切 換器116,且網路配接器118及内插卡12〇、i2i直接連接到 I/O橋 107。 還可改變mi 122到系統100的其餘部分的連接。在一些 實施例中，PP系統112被實施為可插人到系統⑽的擴充插 槽中的内插卡。在其它實施例中’ ppu可與匯流排橋(例 如，記憶體橋105或1/〇橋107)一起整合在單一晶片上。在 另外其它實施例中，刚122的—些或所有元日^與⑽ 102整合。 PPU可具備任何量的區域PP記憶體，包含無區域記憶 體，且可以任何組合使用區域記憶體與系統記憶體。舉例 128625.doc -15 - 200844853 而言，在統一記憶體結構（UMa)實施例中，ppu i22可以 是圖形處理器；在此類實施例中，提供極少或不提供專門 的圖形記憶體，且PPU 122將專門地或幾乎專門地使用系 統記憶體。在UMA實施例中，ppu可整合到橋晶片中或提 供成離散晶片，其中高速鏈路（例如，pci-E)將ppu連接到 橋晶片及糸統記憶體。 還將瞭解，可（例如）藉由在單一内插卡上包含多個 PPU，藉由將多個内插卡連接到路徑113，及/或藉由將一 或多個PPU直接連接到系統主機板，而將任何數目的 包含在系統中。多個PPU可平行操作，從而以高於利用單 一 PPU可此貫現的處理量處理資料。 熟習此項技術者還將瞭解，CPU及叩^^可整合到單一穿 置中，且CPU及PPU可共用例如指令邏輯、緩衝器、高2 缓衝記憶體、記憶體、處理引擎等各種資源；或者可針對 平行處理及其它操作提供單獨的資源。因此，本文描述為 與PPU相關聯的任何或所有電路及/或功能性也可在適當裝 備的CPU中實施並由所述適當裝備的CPU執行。 併入有PPU的系統可以多種組態及形狀因素進行實施， 包含桌上型電腦、膝上型電腦，A掌上型個人電腦、、伺服 器、工作站、遊戲機、嵌入式系統等。 熟習此項技術者還將瞭解，本發明的一個優點是相對於 特定計算硬體的獨立性增加。因此，將瞭解，本發明實施 例可使用任何電腦系統（包含不提供PPU的系統）來實踐。 2·虛擬程式設計模型概述 128625.doc 16 200844853 在本i明實施例中，需要使用ppu 122或計算系統的其 ，它處理器來使用緒陣列執行通料算。如本文所使用， 陣列”是由對輸入資料集併發執行相同程式以產生輸出 貧料集的若干（n0)緒組成的群組。緒陣列中的每一緒被分 ‘ 派有唯：緒識別符(”緒ID”)，其可由緒在其執行期間存 可定義為維或多維數值（例如，〇到i)的緒控制 、.的處理行為的各個恶樣。舉例而言，緒①可用於確定緒 ( 將處理輸入資料集的哪-部分’及/或確定緒將產生或寫 入輸出資料集的哪一部分。 在一些實施例中，緒陣列是，，協作的”緒陣列或CTA。與 其咖的緒陣列一樣，CTA是對輸入資料集併發執行相 同私式（本文稱為”CTA程式”）以產生輸出資料集的多個緒 、羊:在C丁A中，緒可藉由以依賴於緒ID的方式彼此共 用資料而進行協作。舉例而言，在CTA中，資料可由一個 : 並由另緒消耗。在一些實施例中，可在將共用資 ( 料的點處將同步指令插人到CTA程式代碼中，以確保在消 耗資料的、、者σ式圖訪問資料之前已由產生資料的緒實際產生 貧料。CTA的緒之間#資料共用（如果存纟的話）的程度由 CTA私式確定；因此，將瞭解，在使用c丁a的特定應用 中，CTA的緒可能或可能不實際上彼此共用資料，這取決 於CTA私式’且本文中同義地使用術語，’CTA，，及，，緒陣列”。 在一些實施例中，CTA中的緒與同一 CTA中的其它緒共 用輸入資料及/或中間結果。舉例而言，CTA程式;能包: 用於計算共用記憶體中的特定資料將被寫入到的位址的指 128625.doc 200844853 令，其中位址是緒ID的函數。每一緒使用其自身的緒⑴來 計算函數，並寫人到相應位置。有利地定義位址函數使得 不同緒寫入到不同位I ;只|函數是確定性的，就可預測 任何緒寫入到的位置。CTA程式還可包含用於計算共用記 憶體中的將從中讀取資料的位址的指令，其中位址是緒⑴ 的函數。藉由定義適宜的函數及提供同步技術，資料可以 可預測方式由C TA的-個緒寫人到共用記憶體中的給定位 置’並由同-CTA的不同緒從該位置進行讀取。因此，可 支援緒之間的任何所需的f料共㈣式，且⑽中的任何 緒可與同-CTA中的任何其它緒共用資料。 有利地使用CTA(或其它類型的緒陣列）來執行便於資料 平打分解（data-parallel dec〇mp〇siu〇n)的計算。如本文所 使用’”資料平行分解”包含1 一 卜 八T糟由對輸入貧料平行執行 同一演算法多次以產生給Φ次 座生輸出貝枓來對計算問題求解的任何 情形；舉例而言，資料平杆八 、十仃刀解的一個普通實例涉及將同 一處理演算法應用於輸入眘 貝枓集的不同部分以便產生輸出 Μ料木的不同部分。適於資料 > 、於貝科千仃分解的問題的實例包含 矩陣代數、任何數目的_ 、，、又的線性及/或非線性變換（例 如，快速傳立葉變換），4人y 、 匕3任何數目的維度的卷積濾 波、多個維度的可分離濟 、〜 uW 4各種濾波演算法。在CTA程 式中扎疋待應用於輸入資料隹 ,ΓΤΔ . 貝抖巿的母一部分的處理演算法， 且CTA中的母一緒對輸入資

^ ΓΤΔ^ . 、抖木的一個部分執行同一 CTA 粒式。CTA程式可使用廣範 卜卜 ’、国的數學及邏輯運算來f妳、、當 算法，且所述程式可包含 、科連-不貝她冷 3有條件或分支執行路徑以及直接 128625.doc '18· 200844853 及/或間接記憶體存取。 上文引用的申請案第1 1/303,780號中進一步詳細描述了 CTA及其執行。

在些情形中，定義相關CTA(或更一般來說，緒陣歹g ) 的”柵袼”也是有用的。如本文所使用，CTA的，，柵格"是若 干（nl)個CTA的集合，其中所有CTA為相同大小（即，緒數 目）並執行相同的CTA程式。栅格内的nl個CTA有利地彼此 獨立，意味著栅格中任何CTA的執行不受柵格中的任何其 它CTA的執行的影響。如將瞭解，此特徵提供在可用的處 理核心之間分配CTA的顯著靈活性。 為了區分栅格内的不同CTA，有利地向柵格的每一cta 分派"CTA識別符"（或CTA ID)。與緒出一樣，任何唯一識 別符（包含但不限於數值識別符）均可用作cta ι〇。在一個 實施例中，CTA ID只是從_ηι]的連續（一維）指數值。在 其它實施例中，可使用多維指數標^方案。cta ι〇對於 CTA的所有緒疋共同的，且柵格内給定ΜΑ的緒可使用其 CTA ID並結合其緒邮確_如）料讀取輸人資料的源 位置及/或用於寫入輸出資料的目的地位置。以此方式， 同一柵格的不同CTA中的祛π斜门 如 ^ Μ、、者可對同-資料集併發操作，但 在一些貫施例中， 立 支挺柵格中的不同CTA之間的資料共 用0 疋義C ΤΑ拇格（例如）在 題的不同部分求解的情況 能需要執行渡波演算法 需要使用多個CTA來對單一大問 下可能是有用的。舉例而言，可 以產生馬清晰度電視（HDTV)影 128625.doc -19- 200844853 像。如此項技術中已知，HDTV影像可能包含2百萬個以上 的像素。如果每一緒產生一個像素，那麼待執行的緒數目 將超過單一 CTA中可處理的緒數目（假定為使用習知技術構 造的具有合理尺寸及成本的處理平台）。 可藉由在多個CTA之間劃分影像使每一 CTA產生輸出像 素的不同部分（例如，16xl6小片）來管理這種大處理任務。 所= CTA執行相同程式，且緒使ffiCTA m與緒①的組合來

確定用於讀取輸入資料及寫入輸出資料的位置，使得每一 CTA對輸人資料集的正確部分操作並將其輸出資料集部分 寫入到正確位置。 間的緒(其…貝―问，柵格内的 利地不彼此共η料或以另外的方式彼此依賴。也 二是說，可循序(以任一次序)或併發執行同—柵格的兩個 /A’且仍產生相同結果。因此’處理平台(例如，圖!的 糸統1〇〇)可藉由首先執行一個CTA接著執行下一 CM等等 =已：：柵格的所有CTA為止，來執行cta柵格並獲得 ::果二或者，如果有足夠的資源可用，處理平台可 订執订多個CTA來執行相同栅格並獲得 曰 二-些實例中，可能需要定義CTA的多個二 :母：柵格執行資料處理程式或任務的不同部分 而、 :貝料處理任務可能劃分為若干個求解步驟”，1 由執行CTA栅格來執行每—求解步驟。 /、中， 料處理任務可能包含對—連串輸：實例’貧 視訊資料圖框）執行相同1玄§/ 、科木（例如，連續的 )執灯㈣或類似的操作，·可針對每-輸入 128625.doc -20- 200844853 貢料集執行CTA柵格。虛擬程式設計模型有利地支援至少 這三個等級的工作定義（即，緒、CTA及CTA拇格）；視需 要還可支援額外的等級。 將瞭解，用於對特定問題求解的CT_大小（緒的數目 n〇)、拇格的大小（CTA的數目ηι)及柵格的數目將取決於 問題的參數及定義問題分解的程式設計人員或自動代理的 ’ ⑥好。因此，在-些實施例中，CTA的大小、柵格的大小 及栅格的數目有利地由程式設計人員定義。 ( ^於CTA方法的問題通常特徵在於，存在可被平行處 理的大量資料元素。在—些實例中，資料元素是輸出元 素，其母一者是藉由對輸入資料集的不同（可能重疊）部分 執行相同演算法而產生的。在其它實例中，資料元素可以 是輸入元素，其每一者將使用相同演算法予以處理。 此類問題可始終分解為至少兩個等級並映 的緒、CTA及拇格上。舉例而言，每一拇格可能表= 〇 冑料處理任務中的—個求解步驟的結果。每-柵格有利地

劃分為若干"區塊” ’其每一者可作為單— cta被處理。每 -區塊有利地含有多個，，元素”，即待求解的問冑的基本部 - 分（例如，單一輸入資料點或單一輸出資料點）。在cTA 内，每一緒處理一或多個元素。 圖2A及2B說明在本發明實施例中使用的虛擬程式設計 模型中柵格、CTA與緒之間的關係、。圖2a展示若干柵格 2〇〇，每一栅格由CTA 202的二維（2_D)陣列組成。（本文 中，相似物件的多個實例用表示所述物件的參考標號及 128625.doc -21 - 200844853 (需要時）表示實例的括弧標號來指示。）如圖2B中針對 202(0,0)所示，每一 cTa 202包含緒（Θ) 204的2-D陣列。對

於每一栅格200的每一CTA 202中的每一緒2〇4，可定義j = [ig，L，it]的形式的唯一識別符，其中柵格識別符ig唯一地 識別栅格，CTA ID ic唯一地識別柵格内的CTA，且緒m ι 唯一地識別CTA内的緒。在此實施例中，可由一維栅格識 別付Ig、二維CTA識別符ie及二維緒識別符it來構成識別符 I。在其它實施例中，唯一識別符丨是整數的三元組，其中〇 < b < n2; 〇 s ic < ηι ;且〇 g it< n〇。在另外其它實施例 中，柵格、CTA及緒識別符中的任一者或全部可能表達為 、隹正數、2D座標對、3D三元組等。唯_緒識別符j可用 於（例如）確定包含針對整個柵格或多個柵袼的輸入資料集 的陣列内的輸入資料的源位置，及/或確定用於在包含針 對整個柵格或多個栅格的輸出資料集的陣列内儲存輸出資 料的目標位置。 、 舉例而言，在HDTV影像的情況下，每一緒2〇4可能對應 於輸出影像的像素。CTA 202的大小（緒2〇4的數目）是問題 分解中的選擇事項’其僅受對單— CTA2G2中的最大緒數 目的約束的限制（其反映處理器f源的有限性）。栅格⑽可 對應於HDTV資料的整個„，或者多個柵格可映射到單 圖框。 在-些實施例中，問題分解是統一的，意味著所 _具有相同數目及組態的CTA加，且所有CM加且^ 相同數目及組態的緒綱。在其它實施例巾，分解可能不 128625.doc -22- 200844853 統-。舉例而言’不同柵格可能包含不同數目的“A，且 不同CTA(相同栅格或不同柵格中)可能包含不同 緒。 如上定義的CTA可包含數打或甚至數百個併發緒。上面 將執行CTA的平行處理系統可能或可能不支援如此大量的 併發緒。在-個態樣，本發明藉由允許程式設計人員在不 2慮實際硬體能力的情況下使用CTA及CTA柵格的模型來 疋義處理任務，使程式設計人員脫離於此類硬體限制。舉 例而言，程式設計人員可編寫代碼（"cta程式.,），发定義 待由CTA的單一代表性線緒執行的處理任務；將CTA定義 為若干此類緒，每一緒具有唯一識別符；並將拇格定義為 右干CTA，每一 CTA具有唯一識別符。如下文所描述，此 j代碼被自動轉譯為可在特定平台上執行的代碼。舉例而 » ’如果將CTA定義為包含若干(n。)個併發緒但目標平台 僅支援-個緒，那麼翻譯器可定義執行分派到所有η。個緒 的任務的-個實際緒。如果目標平台支援_個以上但少於 η〇個併發緒，那麼可顽雪 J祝而要在所述數目的可用緒之間劃分 任務。 因此’ CTA及栅格的程式設計模型應理解為虛擬模型，

即與任何特定實體實現脫離的作為對程式設計人員的概念 上的輔助的模型。可在對平行處理具有不同程度的硬體支 援的多種目標平台中實現CTA及柵格的虛擬模型。具體而 言，本文所使用的術語"CTA緒"是指離散處理任務（可能與 一或多個其它處理任務協作)的虛擬模型，且應瞭解，CTA 128625.doc -23 - 200844853 、、者可此或可旎不一對一地映射到目標平台上的緒。 3·虛擬架構 祀據本U的-個態樣，定義用於執行CTA及CTA拇格 的虛擬平㈣構。虛擬平行㈣是支援執行能夠在所需時 間實現彼此的協作行為（例如，共用資料及同步）的大量併 發CTA緒时行處理器及相關聯的記憶體空間的表示。此 虛擬平行架構可映射到多種實際處理器及/或處理系統 上，包含（例如如的系統1〇〇的ppu 122。虛擬架構有利 地定義支援不同等級的資料共用及存取類型的若干虛擬記 憶體空間’以及識別可由虛擬處理器執行的所有功能的虛 擬指令集架構(ISA)。虛擬架構還有利地(例如)藉由定義及 啟動CTA或CTA柵格來^義可用於控制CTA執行的虛擬執 4丁驅動。 圖3是根據本發明實施例的虛擬架構3〇〇的方塊圖。虛擬 架構300包含帶有虛擬核心3〇8的虛擬處理器3〇2，虛擬核 心308經組態以平行執行大量CTA緒。虛擬架構3〇〇還包含 可由虛擬處理器302存取的全局記憶體3〇4，幻共應控制虛 擬處理器302的操作的命令的虛擬驅動器32〇。虛擬驅動器 320也可存取全局記憶體304。 虛擬處理器302包含接收及解譯來自虛擬驅動器32〇的命 令的前端306,以及能夠併發執行單一 CTA的所有η。個緒的 執行核心308。虛擬核心308包含大量（n❶個或nQ個以上）虛 擬處理引擎310·，在一個實施例中，每一虛擬處理引擎31〇 執行一個CTA緒。虛擬處理引擎31〇併發執行其各自cta 128625.doc -24- 200844853 4不疋平行執行。在一個實施例中，虛擬架構3〇〇 ，疋虛擬處理引擎31〇的數目T(例如，384、500、768 等）’此數目a定CTA中的緒數目的上限。應瞭解，虛擬 木構300的實現可包含少於指^數目了的實體處理引擎，且 單一處理引擎可執行若干CTA緒，作為單一，，真實"（即，平 台支援的）緒或作為多個併發的真實緒。 虛擬處理裔302還包含虛擬指令單元3 12，其保持向虛擬 處理引擎310供應針對其各自CTA緒的指令；所述指令由 作為虛擬架構3GG的-部分的虛擬ISA定義。下文描述用於 平行緒計算的虛擬ISA的實例。指令單元312在向虛擬處理 引擎310供應指令的過程中管理CTA緒同步及a緒行為的 其它協作性態樣。 執行核心3 0 8提供具有不同等級的可存取性的内部資料 儲存。特殊暫存器3 11可由虛擬處理引擎3 1()讀取但不可由 其寫入，且用於儲存定義圖2的問題分解模型内每一 緒的”位置”的參數。在一個實施例中，特殊暫存器3ιι對 於每CTA緒（或每虛擬處理引擎31())包含—個儲存_的暫 存器，·每-細暫存器僅可由虛擬處理引擎31〇中的個別 虛擬處理引擎310存取。特殊暫存器311還可包含可由所有 CTA緒（或由所有虛擬處理引擎31〇)讀取的額外暫存哭，其 儲存CTA識別符、CTA維度、CTA所屬的栅格的維；，及 CTA所屬的柵格的識別符。在初始化期間響應於經:前端 306從虛擬驅動器320接收的命令而寫入特殊暫存器η卜 且特殊暫存器311在CTA執行期間不改變。 128625.doc -25- 200844853 區域虛擬暫存器314由每一 CTA緒用作臨時空間（s_ch space);分配每—暫存器專用於—個⑽緒（或—個虛擬處 理引擎310) ’且區域暫存器314的任一者中的資料僅可由 其被分配到的CTA緒存取。共用記憶體316可由所有cta緒 (單CTA内）存取’共用記憶體3 1 6中的任何位置可由同一 CTA内的任何CTA緒存取(或可由虛擬核心3〇8内的任何虛 擬處理引擎31G存取）。參數記憶體318儲存可由任何CTA緒 f (或任何虛擬處理5丨擎31G)讀取料可由其寫人的執行時間 ' I數(常數)。在-個實施例中，虛擬驅動器320在引導虛擬 處理器302開始執行使用此等參數的cta之前，將參數提 供到參數記憶體川。㈣CTA内的任何CTA緒（或虛擬核 心308内的任何虛擬處理引擎31〇)均可藉由記憶體介面π] 存取全局記憶體304。 在虛擬架構300中，虛擬處理器3〇2在虛擬驅動器32〇的 控制下作為共處理器操作。虛擬架構規範有利地包含虛擬 t 應用耘式”面（API) ’其識別由虛擬驅動器320認可的函數 調用及每一函數調用預期產生的行為。了文描述用於平行 緒計算的虛擬API的實例函數調用。 . 虛擬架構300可在多種硬體平台上實現。在一個實施例 • 中虛擬架構300在圖1的系統1 〇〇中實現，其中ppu 122實 施虛擬處理器302，且在CPU 1〇2上執行的ppu驅動器程式 實施虛擬驅動器320。全局記憶體3〇4可在系統記憶體丨〇4 及/或PP記憶體124中實施。 在一個實施例中，PPU 122包含一或多個處理核心，其 128625.doc -26- 200844853 使用單指令多資料（SIMD)及多線緒技術來支援來自單一指 令單元（實施虛擬指令單元312)的大量（例如，384或768個） 緒的併發執行。每一核心包含p個（例如，8、16等）平行處 理引擎302的陣列，其經組態以從指令單元接收並執行 SIMD指令’從而允許平行處理具有多達p個緒的群組。核 =為多線緒的，其中每一處理引擎能夠（例如）藉由維持與 每一緒相關聯的當前狀態信息使得處理引擎可從一個緒快

速切換到另一緒，而併發執行多達某一數目G(例如，24 個）的緒群組。因此，核心併發執行各有p個緒的〇個§1以〇 群組，總計為P*G個併發緒。在此實現過程中，只要"Ο ^ （虛擬）CTA緒與在真實ppu 122上執行的併發緒之間 就可存在一對一對應。 特殊暫存為311可藉由以下操作而實施在ppu 中：向 每一處理核心提供1>*(3輸入項暫存器堆，其中每一輸入項 能夠儲存一緒ID;以及提供一組全局可讀取的暫存：來儲 存CTA ID、柵格ID&CTA及柵格維度。或者，可使用其它 儲存位置來實施特殊暫存器3 j i。 匕 。區域暫存器314可實施在PPU 122中，作為在實體上或邏 輯上劃分為P個巷道（lane)的區域暫存器 且古甘 一巷道 ”有某-數目的輸入項(其中每一輸入項可能儲存(例如识 位70字組）。一個巷道被分派到P個處理引擎中的每一者， 二1巷道中的相應輸入項可用執行相同程式：有助於 MD執行的不同緒的資料填充。巷道的不同部分 到G個併發緒群組中的不同緒 尺侍^域暫存器堆中 128625.doc -27- 200844853 的給定輸入項僅可由特定緒存取。在—個實施例中，區域 暫存器堆内的某些輸入項保留用於儲存緒識別符，從而實 施特殊暫存器3 11中的一者。 共^記憶體316可實施在PPU 122中作為共用暫存器堆或 ,、 片上同速緩衝5己’丨思體，其具有允許任何處理引擎對

共用記憶體中的任何位置進行讀取或寫人的互連。參數記 憶體318可實施在PPU 122中作為實施共用記憶體316的同 一共用暫存器堆或共用高速緩衝記憶體内的指定區域，或 作為處理引擎具有唯項（read-Qnly)存取權的單獨共用暫存 器堆或晶片上高速緩衝記憶體。在—個實施例中，實施參 數記憶體的區域還用於儲存CTA m及栅格m以及cta及柵 格維度，從而實施特殊暫存器31丨的若干部分。 在一個貫施例中，在圖1的CPU 1〇2上執行的ppu驅動器 程式藉由將命令寫入到記憶體（例如，系統記憶體1〇4)中的 推播緩衝器（pushbuffer)(未明確展示）來響應虛擬Αρι函數 调用，PPU 1 22從推播緩衝器處讀取所述命令。命令有利 地與狀態參數相關聯，狀態參數例如CTa中緒的數目、待 使用CTA處理的輸入資料集在全局記憶體中的位置、待執 行的CTA程式在全局記憶體中的位置，及輸出資料將被寫 入到的全局兄憶體中的位置。響應於命令及狀態參數， PPU 122將狀態參數載入到其核心中的一者中，接著開始 啟動緒直到已啟動CTA參數中指定的數目的緒為止。在_ 個實施例中，PPU 122包含控制邏輯，其在啟動緒時將緒 ID依次分派到緒；可將緒ID儲存（例如）在區域暫存器堆内 128625.doc -28 - 200844853 或專用於此目的的特殊暫存器中的指定位置。 在替代實施例中’在單線緒處理核心中(例如，在一些 CPU中）實現虛擬架構3⑽，單線緒處理核心使用少於個 實際緒來執行所有„场；虛擬程切計模型使得與不同 C T A緒相關聯的處理任藤 务了（例如）藉由針對一個CTA緒接 著針對下-CTA緒等等執行任務（或其一部分）而組合到單 n·用向量執行'SIMD執行及/或機器中可用的

任何其匕形式的平行性來平行執行與多個CTA緒相關聯的 處理任務，或平行執行與同一 CTA緒相關聯的多個處理任 :。因此’可使用單一緒、n。個緒或任何其它數目的緒來 貝見TA如下文所描述，虛擬指令翻譯器有利地將編寫 為目標虛擬架構300的代碼轉譯為特定針對目標平台的指 令。 將瞭解，本文描述的虛擬架構是說明性的且可能作出變 化及修改。舉例而言，在一個替代實施例中，每一虛擬處 理引擎可具有儲存分派到其緒的唯一緒ID的專門的緒〗〇暫 存器’而不使用區域虛擬暫存器中的空間用於此目的。 作為另一實例，虛擬架構可指定關於虛擬核心3〇8的内 邛結構的或多或少的細節。舉例而言，可能指定虛擬核心 308包含用於執行卩向simd群組中的CTA緒的p個多線緒虛 擬處理引擎，其中多達(3個81]^1〇群組共存於核心3〇8中使 得P*G確定T(CTA中緒的最大數目）。還可指定不同類型的 5己體及共用等級。 虛擬架構可實現於使用硬體及/或軟體元件的任何組合 128625.doc -29- 200844853 的夕種電月自糸統中以定義及控制每一組件。雖然已以舉例 的方式描述使用硬體組件的一個實現方案，但應瞭解，本 發明涉及使程式設計任務與特定硬體實現脫離。 4·對虛擬架構進行程式設計 圖4是根據本發明實施例使用虛擬架構3 〇〇來操作目標處 理裔或平台440的概念模型400。如模型4〇〇所展示，虛擬 架構300的存在使經編譯的應用程式及Αρι與目標處理器或 平台的硬體實施脫離。

應用程式402定義資料處理應用，其利用上述虛擬程式 設計模型（包含單一CTA及/或CTA柵格）。一般來說，應用 程式402包含多個態樣。第一，程式定義單一cta緒的行 為。第二，程式定義CTA的維度（以CTA緒的數目計），且 如果將使用柵格，那麼定義柵格的維度（以cta的數目 4)第二，程式定義待由CTA(或柵格）處理的輸入資料集 及輸出貝料集將被儲存在的位置。第目，程式定義總體處 理订為’包含（例如）何時啟動每—cta或栅格。程式可包 含動態確定CTA或柵格的維度、是否不斷啟動新的cta或 柵格等的額外代碼。 :U C/(： + +、FORTRAN等高階程式設計語言編寫應 用私式402。在一個實施例中，c/c++應用程式直接指定一 擬）CTA緒的行為。在另-實施例中，使用資料平行 Poman 90、c*或資料平行c)編寫應用程式， &式指定對陣列及聚合資料結構的資料平行操作； ，可、、扁澤成私疋一個（虛擬）cta緒的行為的虛擬MA程 128625.doc >30- 200844853 為了允許定義CTA緒的行為’可提供語言擴充或 =庫，程式料人貞可經由所料言擴充或函數庫來指 疋平行CTA緒行為。舉例而古，可宁差、姑& σ 疋義特殊符號或變數以 1應於緒ID、CTA ID及柵格id，且可楹彳丑τ # 口 η υ丑j徒供函數，程式設計 人員可經由所述函數來指示CTA块廡 A、、者應何日$與其它CTA緒同 〇 $編譯應用程式4G2時，編譯器彻針對應用程式術的 疋義CTA緒行&的那些部分產生虛擬ISA代碼41〇。在一個 實施例中，以圖3的虛擬架構_的虛擬似表達虛擬似代 碼410。虛擬ISA代碼410是程式代碼，但不一定是可在特 定目標平台上執行的形式的代碼。如此，虛擬似代碼彻 可作為任何其它程式代碼被儲存及/或分布。在其它實施 例中二可將應用程式完全或部分指定為虛擬似代碼41〇, 且可完全或部分避免使用編譯器4〇8。 虛擬指令翻譯器412將虛擬ISA代碼41〇轉換為目標isa代 碼414。在—些實施例中，目標ISA代碼414是可由目標平 台440直接執行的代碼。舉例而言，如由圖4中的虛線框所 不，在一個實施例中，目標ISA代碼414可由ppu 122中的 =令單元430接收及正確解碼。視目標平台44()的特殊性而 定，虛擬ISA代碼410可能被轉譯成待由目標平台44〇上的 n〇個、者的4 |執行的每線緒代碼。或者，虛擬似代碼 410可能被轉譯成將在少於…個的緒中執行的程式代碼’， 其中每-緒包含與CTA緒中的—者以上有關的處理任務。 在些只靶例中，CTA及/或栅格的維度的定義以及定義 128625.doc -31 - 200844853 輸入資料集及輸出資料集是由虛擬API處理的。應用程式 402可包含對於虛擬API函數庫4〇4的調用。在一個實施例 中，將虛擬API的規範（包含（例如）函數名稱、輸入、輸出 及作用，但不包含實施細節）提供給程式設計人員，且程 式設計人員將虛擬API調用直接併入到應用程式4〇2中，藉 此直接產生虛擬API代碼406。在另一實施例中，藉由編^ 使用某一其它語法來定義CTA及柵格的應用程式4〇2來產 生虛擬API代碼406。

部分地藉由提供虛擬執行驅動器416來實現虛擬Αρι代碼 4〇6,虛擬執行驅動器416將代碼4〇6的虛擬Αρι命令轉譯為 可由目標平台440處理的目標Αρι命令MS。舉例而言，如 由圖4中的虛線框所示’在一個實施例中，目標a二命令 418可由m；驅動器432接收及處理，ppu驅動器432將相應 命令料到PPU 122前端434。（在此實施例中，虛擬執^ 驅動器416可以是PPU驅動器432的—態樣或部分。）在另一 實施例中，虛擬執行驅動ϋ可能不對應於用於共處理器的 。動如八可犯/、疋在執行所述虛擬執行驅動器的同一處 理器上啟動其它程式或緒的控制程式。 处 應瞭解，可針對能夠支援CTA執行的任何平台或結構創 建虛擬4日令翻澤盗412及虛擬執行驅動器4 i 6。就針對不同 平台或結構的虛擬指令翻譯器412可從同—虛擬似進行轉 譯而言，同-虛擬！SA代碼41阿與任何平台或結構— 用。因此，不需要針對每一可能的平台或 用程式402。 4、勒#應 128625.doc -32- 200844853 此外，目標平台440不必包含如圖4所示的PPU及/或PPU 驅動器。舉例而言，在一個替代實施例中，目標平台是使 用軟體技術仿真大量緒的併發執行的CPU，且目標ISA代 碼及目標API命令對應於待由目標cpu執行的程式（或互相 通信的程式的群組）中的指令，所述目標CPU可以是（例如） 單核心或多核心CPU。 5· 虛擬ISA實例

現在將描述根據本發明實施例的虛擬ISA的實例。如上 所述，虛擬ISA有利地對應於上述虛擬程式設計模型（CTA 及栅格）。因此，在本實施例中，由編譯器4〇8產生的虛擬 ISA代碼4 10疋義要由圖3的虛擬核心3⑽中的虛擬處理引擎 310之一執行的單個CAT緒的行為；所述行為可包含與其 匕CTA緒的協作性交互，例如同步及/或資料共用。 應瞭解，本文中描述的虛擬ISA只用於說明的用途，且 本文中彳田述的特定要素或要素組合並不限制本發明的範 圍在1實施例中，程式設計人員可用虛擬IsA編寫代 碼；在其它實施例巾’程式設計人員用另一高階語言(例 如FORTRAN，C，C++)編寫代碼，且編譯器4〇8產生虛擬 ISA代碼。程式設計人員也可編寫"混合”代碼，其中代碼 中有些部分是用高階語言編寫而其它部分是用虛擬ISA編 寫0 5·1特殊變數 圖^列舉由實W虛擬ISA定義的，，特殊，，變數的表格 (本文中用前綴"來表示特殊變數）。此等變數與圖2 128625.doc -33- 200844853 的程式設計模型有關，其中藉由每個緒2〇4在（：丁八2〇2内的 位置來識別線緒，而CTA 202又位於某一數目的柵格2〇〇中 的特疋者θ在些貝知例中，表格5〇〇的特殊變數對 應於圖3的虛擬架構3〇〇中的特殊暫存器311。 在表格500中，假設CTA及栅格各自用三維空間來定 義，且不同的栅格在一維空間中循序編號。虛擬Ι§Α預期 圖5的特殊變數將在啟動CTA時被初始化，且虛擬isa代碼 可簡單地使用此等變數而無需初始化。以下參看虛擬Αρι 論述特殊變數的初始化。 如圖5所示，特殊變數的第一個3向量。/〇ntid = (%ntid x， /〇 ntid.y，％ntid，z)定義CTA的維度（以緒數目計）。cta中的 所有緒將共用同一 ％ntid向量。在虛擬架構30〇中，預期 /〇ntid向；g：的值將經由虛擬API函數調用提供給虛擬處理器 3 02，所述虛擬API函數調用如下所述地建立cta的維度。 如圖5所示，特殊變數的第二個3向量％tid = (%tid x， % tid.y，％tid.z)是指CTA内的給定緒的緒id。在圖3的虛擬 架構300中，預期虛擬處理器302將在啟動CTA的每個 緒時分派滿足制約條件〇幺0/〇tid.x < %ntid.x、0幺0/〇tid.y < %ntid.y 及 0 幺 ％tid，z < 0/〇ntid.z的唯一 ％tid向量。在一個 實施例中，％tid向量可定義成使得其可儲存在壓縮的32位 元的字組中（例如，對M%tid.x為16個位元，對M%tid.y為 10個位元，且對K%tid.z為6個位元）。

如圖5所示，特殊變數的第三個3向量％nctaid= (%nctaid.x，％nctaid.y，％nctaid.z)定義栅格的維度（以 CTA 128625.doc -34- 200844853 的數目計）。在圖3的虛擬架構300中，預期。/❶nctaid向量的 值將經由虛擬API函數調用提供給虛擬處理器3 〇2，虛擬 API函數調用建立cta的柵格的維度。 如圖5所示，特殊變數的第四個3向量％ctaid=(%ctaid x， 0/〇ctaid.y，o/octaid.z)是指柵格内的給定c：丁八的c丁a ID。在圖 3的虛擬架構300中，預期當啟動CTA時將把滿足cta的限 制條件 0 S %ctaid.x < %nctaicLx、〇 s %ctaid.y < %nctaid y& 〇 S %ctaid.z < %nctaid，z的唯一 ％ctaid向量提供給虛擬處理 器 302 〇 特殊變數還包含提供CTA所屬的栅格的柵格識別符的純 i /ogridid變數。在圖3的虛擬架構3⑼中，預期將把 /ogndid值提供給虛擬處理器3〇2，以便識別包括當前◦丁a 的柵格。例如當正在使用多個柵格來解決較大問題的不同 部分時，有利地在虛擬ISA代碼中使用％^^1(1值。 5·2程式定義的變數及虛擬狀態空間 虛擬ISA使得程式設計人員（或編譯器）可定義任意數目 的變數以代表正被處理的資料項目。藉由類型及指示如何 使用文數以及變數在何種程度上共用的，，虛擬狀態空間，，來 定義變數。使用目標平台中可用的暫存器或其它記憶體結 構來貝現文數’在许多目標平台巾，狀態空間可能會影響 對用來實現特定變數的記憶體結構的選擇。 圖6是列舉實例虛擬1SA實施例中支援的變數類型的表袼 6〇0。支援四個類型：未指定類型的位元、有正負號的整 數、無正負號整數及浮點。未指的變數就是單個位 128625.doc -35 - 200844853 凡或具有指定長度的位元的群組。可根據習知格式（例 如，IEEE 754標準）來定義有正負號整數及無正負號整數 格式以及浮點格式。 在本實施例中，針對每個類型支援多個寬度，其中使用 參數<n>來指定寬度；因此，舉例而言，·sl6指示16個位 元的有正負號的整數，.f32指示32位元的浮點數字等。如 表袼600所示，有些變數類型限於特定的寬度；舉例而 言’浮點變數必須至少為1 6個位元；且整數類型必須至少 為8個位元。預期虛擬ISA的實現支援所有指定寬度；如果 處理器的資料路徑及/或暫存器比最寬寬度窄，那麼如此 項技術中已知的可使用多個暫存器及處理器循環來處置較 寬類型。 應瞭解，本文中使用的資料類型及寬度是說明而不是限 制本發明。 圖7疋列舉實例虛擬is A中支援的虛擬狀態空間的表格。 疋義了九個狀悲空間’其對應於不同的共用級別及在圖3 的虛擬架構300中的可能的儲存位置。 在緒級別上共用别二個狀態空間，這意味著每個C τ a緒 將具有單獨的變數實例，且沒有任何CTA緒將可存取任何 其它CTA緒的實例。有利地使用虛擬暫存器（reg)狀態空間 來定義要由每個CTA緒執行的計算的運算數、臨時值及/或 結果。程式可宣稱任何數目的虛擬暫存器。只能藉由靜態 編澤時間名稱而不是計算出來的位址來定址虛擬暫存器。 這個狀態空間對應於圖3的虛擬架構3〇〇中的區域虛擬暫存 128625.doc -36- 200844853 器 3 1 4。 特殊暫存ϋ(.⑽g)狀態空間對應於圖5的預定義 變數，其儲存在虛擬架構中的特殊暫存器3ιι中。在一 些實施财，虛擬ISA代碼可能不會宣稱叫空間中的任 何其它變數，而是可使用特殊變數作為對計算的輸入。所 有CTA緒均可讀取，Sreg狀態空間中的任何變數。對於 。偏（或其分量），每個CTA緒將讀取其唯_的緒識別符； 對於.sreg狀態空間中的其它變數，同一cTA中的所有cTA 緒將讀取相同值。 每線緒區域記憶體（.1〇Cal)變數對應於全局記憶體3〇4中 在每CTA緒基礎上分配及定址的區域。換句話說，當c丁a 緒存取.local變數時，其存取其自身的變數實例，且在一 個CTA緒中對.l〇cai變數作出的改變不會影響其它緒。 與.reg及.sreg狀態空間不同，每線緒區域記憶體可使用計 算出來的位址來定址。 接下來的兩個狀態空間定義每CTA變數，這意味著每個 CTA將具有變數的一個實例，其任何（虛擬）緒均可存取所 述實例。任何CTA緒均可讀取或寫入共用（shared)變數。 在一些實施例中，這個狀態空間映射到虛擬架構3〇〇(圖3) 的虛擬共用記憶體316。在虛擬架構300的實現中，.shared 狀態空間可映射到晶片上共用記憶體實施方案（例如，共 用暫存器堆或共用高速緩衝記憶體）上，而在其它實現 中’ .shared狀態空間可映射到晶片外記憶體的分配及定址 為任何其它可全局存取記憶體的每CTA區域上。 128625.doc -37- 200844853 參數Cparam)變數是唯讀的，且可由CTA中的任何（虛擬） 緒讀取。這個狀態空間映射到虛擬架構3〇〇的參數記憶體 3 18，且可（例如）在晶片上共用參數記憶體或高速緩衝記憶 體中或者在可全局存取的晶片外記憶體的分配及定址為其 它任何可全局存取記憶體的區域中實現。預期此等變數將 響應於來自虛擬驅動器320的驅動器命令來初始化。 常量（.const)狀態空間用來定義可由柵格中的任何cta中 的任何（虛擬）緒讀取（但不可修改）的每栅格常量。在虛擬 架構300中，.const狀態空間可映射到全局記憶體中可由 CTA緒進行唯讀存取的區域。·c〇nst狀態空間可在晶片上 共用參數記憶體或高速緩衝記憶體中或在可全局存取晶片 外記憶體的分配及定址為任何其它可全局存取記憶體的每 栅格區域中實現。與.param狀態空間一樣，預期·c〇nstw 態空間中的變數將響應於來自虛擬驅動器3 2 〇的驅動器命 令而被初始化。 其餘三個狀態空間定義”内容”變數，其可供與應用程式 相關聯的任何CTA中的任何（虛擬）緒存取。此等狀態空間 映射到虛擬架構300中的全局記憶體304。全局（gi〇bal)變 數可用於一般用途。在一些實施例中，也可定義用於共用 紋理（.tex)及表面（.surf)的特定狀態空間。此等可用於（例 如）圖形相關應用程式的狀態空間可用來定義並提供對圖 形紋理及像素表面資料結構的存取，所述資料結構提供對 應於2D(或在一些實施例中為3D)陣列的每個像素的資料 值。 128625.doc -38- 200844853 在圖4的虛擬ISA代碼4丨〇中，藉由指定狀態空間、類型 及名稱來宣稱變數。名稱是占位符，且可由程式設計人員 或編譯者來選擇。因此，例如： •reg .b32 vrl; 旦稱虛擬暫存器狀態空間中名為vrl的32個位元的未指 疋類型k數。虛擬ISA代碼的後續行可引用vrl(例如）作為 操作的來源或目的地。 實例虛擬IS A還支援虛擬變數的陣列及向量。舉例而 •global .f32 resultArray[l〇〇〇][i〇〇〇j ； 旦稱32位元浮點數字的虛擬可全局存取1〇〇〇χ 1〇〇〇陣 列。虛擬指令翻譯器412可將陣列映射到對應於分派的狀 態空間的可定址記憶體區域中。 可使用向量前綴.v<m>來定義一個實施例中的向量，其 中m是向量的分量的數目。舉例而言： •reg ·ν3 .f32 vpos; 宣稱每線緒虛擬暫存器狀態空間中的3 2位元浮點數字的 3分量向量。一旦宣稱了向量，便可使用例如vp〇s.x， vpos’y’ VpOS Z等後綴來識別其分量。在一個實施例中，允 許m=2、3或4，且可使用例如（·χ， y，，z，⑼、 (•0，·1，.2, .3)或（.r，.g，.b，.a)等後綴來識別分量。 由於變數是虛擬的，所以虛擬ISA代碼41〇可定義或引用 任何狀態空間（除了 .sreg，其中變數是預定義的）中的任何 數目的雙數。可能針對虛擬ISA代碼41 0中的特定狀態空間 128625.doc -39- 200844853 定義的變數的數目可能會超過特定硬體實施 類型的儲存量。虛擬指 木、目〜 合適的儲存管理指令(例 間移動資料），以^在暫存盗與晶片外記憶體之 哭412也可^ ’變數在需要時可用。虛擬指令翻譯 :甘 夠檢測到不再需要臨時變數的情況並允許 為其分配的空問^ _ … U、另一變數重新使用；可使用習知的用於 分配暫存器的編譯器技術。 用於 此外，雖'然實例虛擬IS々義向量變數類型，但不要求 =標平台支援向量變數。虛擬指令翻譯器412可將任何向 量變數實施為適當數目(例如’ 2、3或4)個純量的集合。 5 · 3 ·虛擬指令 圖8A-8H是列舉實例虛擬似中定義的虛擬指令的表 格。藉由指令的效果來定義指令，所述效果例如是使用 或多個運算數計算出特定結果並將所述結果放置在目的地 暫存·™中。又置暫存态值等。將大多數虛擬指令分類以識 別輸入及/或輸出的格式，且指令執行的各態樣可取決於 類型。指令的一般格式是： name.<type> result, operands; 其中name是指令名稱；·<type>是圖6中列舉的任一類型 的占位符；result是儲存結果的變數；且叩以⑽心是提供為 對指令的輸入的一或多個變數。在一個實施例中，虛擬架 構3 00是暫存器到暫存器處理器，且用於除記憶體存取之 外的操作的result及operands(圖8F)被要求是在虛擬暫存器 狀態空間.reg(或在一些運算數的情況下是特殊暫存器狀態 128625.doc -40- 200844853 空間.sreg)中的變數。 預期目標平台實現虛擬ISA中的每個指令。指令可實現 為產生指定效果的相應的機器指令（本文中稱為，，硬體支援Ί， 或實現為在執行時產生指定效果的機器指令序列(本文中 ί冉為車人體支板）。特定目標平台的虛擬指令翻譯器化有 利地經組態以識別對應於每個虛擬指令的機器指令或機器

ί: :下子章節描述圖8Α_8Η中列舉的各種類別的指令。應 瞭解本文中提供的指令列表是說明性的，且虛擬Μ可 包含本文中未明確描述的額外指《，且可不包括本文中描 述的一些或所有指令。 5·3·1·虛擬指令-算術 圖8Α疋列舉實例虛擬ISA中定義的算術運算的表格 _。在此實施例巾，虛擬架構只支援暫存器到暫存哭曾 術’且所有算術運算均操縱—或多個虛擬暫存器運算: (圖8A中表示為a、b、c)，以產生寫入到虛擬暫存器的社 果.因此’算術運算的運算數及目的地始終在虛擬暫存 器狀態空間.reg中，除了圖5的特殊暫存器(在特殊暫存器 狀態空間.sreg中）可用作運算數以外。 时 一表格8咐的算術運算的列表包含四個基本的算術運 算··加法（add)、減法（sub)、乘法（丨）及除法⑻小此等 運算可對所有整數及浮點資料_執行，且產生相同類型 的結果作為輸人；在-些實施财，也可將捨人模式的限 定符（qualifier)添加到指令中，以允許程式設計人員指定（ 128625.doc -41 - 200844853 應當如何對結果進行捨入，以及在整數運算數的情況下是 否應當強加飽和極限。 也支援a、b及c運算數的三個複合算術運算··乘加 (mad)、熔合乘加（fma)以及絕對差和（sad)。乘加計算& Α的 乘積（具有括號指示的捨入），且將結果與c相加。熔合乘加 與mad的區別在於，在與c相加之前，不對a><b的乘積進行 捨入。絕對差和計算絕對值卜讣丨，然後與c相加。 餘數（rem)運算只對整數運算數執行，且當將運算數&除 以運算數b時計算餘數（a mod b)。絕對值（abs)及求反0%) 是可應用於浮點格式或有正負號的整數格式的運算數^的 一元運算。可應用於整數或浮點運算數的最小值（mb)及 最大值（max)運算將目的地暫存器設置成較小運算數或較 大運算數；也可指定對一個或兩個運算數是非正規數（例 如，根據IEEE 754標準）的特殊情況的處置。 表格800中的其餘運算只對浮點類型執行。分數汴匀運 算返回其輸入的分數部分。正弦(sin)、餘弦(c〇s)及比率的 反正切（atan2)提供對應於三角函數的方便的指令。也支援 以2為底的對數（lg2)及取冪㈣）。也支援倒數（邮、平方 根（sqrt)及倒數平方根（rsqrt)。 應注意，這個算術運算列表是說明而不是限制本發明。 可支援其它運算或運算組合’其中包含任何預期會以足夠 的頻率被調用的運算。 在一些實施例中，虛擬ISA還定義向旦 I心我句里運异。圖8B是列 舉實例虛擬ISA支援的向量運算的声坆 磲"r扪表格81〇。向量運算包含 128625.doc -42- 200844853 點積（dot)運异’其什异運算數向量&及b的純量點積d ;又 積（cross)運算，其計算運算數向量a&b的向量叉積d;以 及量值（mag)運算，其計算運算數向量a的純量長度d。向 量歸約（vred)運算藉由以迭代方式對向量運算數&的要素執 行指定操作<〇P>來計算純量結果d。在一個實施例中，對 於浮點向Ϊ只支挺~約運算add、mul、min及max ;對於整 數向里，也可支挺頭外的歸約運算（例如，and、or及x〇r， 如下文所述）。 除了此等運算之外，也可在虛擬ISA中定義例如向量加 法、向量定標等其它向量運算（圖8B中未列舉）。 如上所述，虛擬架構300的一些硬體實現可能不支援向 量處理。此種實現的虛擬指令翻譯器4丨2有利地適合於產 生純量機器指令的適當序列以執行此等運算；熟習此項技 術者將能夠確定適當的序列。 5·3·2·虛擬指令-選擇及設置暫存器 圖8C是列舉實例虛擬ISA中定義的選擇及設置暫存器運 异的表格820。此等運算可對任何數值資料類型執行，且 基於比較運算的結果設置目的地暫存器。如果c是非零那 麼基本選擇（sel)運算選擇運算數a，且如果c是零則選擇運 算數b。比較與設置（set)對運算數a&b執行比較運算 以產生比較結果t，接著基於比較結果t是真（〜〇)還是假（〇) 而將目的地暫存器d設置成布林（B〇〇lean)真（〜〇)或假（〇)。 個實施例中允許的比較運算<cmp>包含等於（如果$ 為真）、大於（如果a>b則t為真）、小於（如果a<M,jt為真）、 128625.doc -43 - 200844853 大於或等於（如果cb則t為真）、小於或等於（如果t為 真），以及其它比較m含例如a及心是數字值還是 未定義的值。 setb運算是對比較與設置的變化形式，其在比較運算 < c m p >的結果t與第三運算數c之間執行進一步的布林運算 <bop> ;布林運,t<b〇p>cw結果確定是將目的地暫存器d 設置成布林真還是布林假。一個實施例中允許的布林運算 <bop>包含and、or及xor(見下述圖8C)。setp運算與化讣相 似，區別只是設置了兩個i位元”斷言，，（predicate)目的地暫 存器：將目的地暫存器dl設置成t<b〇p>c的結果，而將目 的地暫存器d2設置成（!t)<bop>c的結果。 5·3·3·虛擬指令-邏輯及位元操縱 圖8D是列舉實例虛擬13八_定義的邏輯及位元操縱運算 的表格830。按位元的布林運算and、〇Γ&χ〇Γ藉由以下方式 來執行：對運算數a及b的每個位元執行指定運算，且將暫 存器d中的相應位元設置成結果。按位元求反運算反 轉運异數a的每個位元，而如果a是零（布林假）則邏輯求反 (c η 〇 t)運异將目的地暫存器設置成1 (布林真），否則設置成 〇(布林假）。 藉由左移（shl)及右移（shr)運算來支援移位，所述運算將 運算數a中的位元欄位左移或右移運算數b指定的位元數。 對於有正負號的格式，右移有利地基於符號位元來填補引 導位元；對於無正負號格式，右移用零來填補引導位元。 5·3·4·虛擬指令-格式轉換 128625.doc -44- 200844853 圖8E是列舉實例纽似令定義的格式轉換運算的表格 84〇。格式轉換㈣)指令將第一類型〈叫⑽的運算數&轉 換成目標類型令的均等值,ye>，且將結果儲存在目的 地暫°圖6中列舉—個實施财的有效類型；無法 :未#曰定類型值（ b<n>)轉換成整數或浮點類型或者將整數 或浮點類型轉換成未指定類型值。格式轉換指令的變化形 式2得程式設計人員可指定捨入模式<mode> ;也可指定 對當表達為目標類型時飽和的數字的處置。 5·3·5·虛擬指令·資料移動及資料共用 匕圖8F是列舉實例虛擬ISA中定義的資料移動及資料共用 指令的表格850。移動（mov)操作將目的地暫存器d設置成 立即運算數a的值，或者如果運算數3是暫存器，則設置成 暫存器a的内容。移動操作可限於虛擬暫存器類型的狀態 空間’例如圖7中的.reg及.sreg。 載入⑽指令將來自記憶體中的源位置的值载入到目的 地暫存器d中’所述目的地暫存器d在一個實施例中必須在 虛擬暫存器（.reg)狀態空間中。.<spaee>限定符指定源位置 的狀態空間，且可限於圖7中的可定址狀態空間，例如W 及.叫之外的空間(其中可改為使用移動操作)。由於這個 實施例中的虛擬架構300是暫存器到暫存器處理哭，所以 載入指令有利地用來將變數從可定址的狀態空間轉移到虛 擬暫存器.reg狀態空間中，使其可用作運算數。 业 使用可用各種方式定義的來源參數< s r c >來識別特定的 源位置’以便支援不同的定址模式。舉例而言，在—些實 128625.doc -45- 200844853 施例中，來源參數<src>可能是以下中的任何一者：其值 口σ在+的、赵命名#可定址變數、對保存來源位址的暫 子。的引用、對保存將與偏移值（提供為立即運算數）相加 的位址的暫存裔的引用，或立即絕對位址。 • 類似地，儲存（st)操作將來源暫存器a中的值儲存到由目 、/數dst>識別的圯憶體位置。一個實施例中的來源 暫存Ha必須在·reg狀態空間中；目的地必須在可寫入及可 f , ffl 7t local > .global ^.shared) 中目的地苓數<dst>可用各種方式定義以支援不同的定 址扠式，這與載入指令中的來源參數<src>相似。例如， 可使用儲存指令將操作結果從暫存器轉移到可定址的狀態 空間。 “ 在提供紋理及表面狀態空間的實施例中，可使用額外的 虛擬指令從紋理記憶體狀態空間進行讀取（tex)及從表面記 憶體狀態空間進行讀取（suld)並對其進行寫入（sust)。紋理 ( 讀取的運算數（t，X，y)指定紋理識別符⑴及座標（X，y);同 樣，表面讀取或寫入的運算數（s，x，y)指定表面識別符⑷ 及座標（X，y)。 • CTA緒可藉由與其它CTA緒共用資料而與其它cTA緒協 • 作。舉例而言，為了在CTA内共用資料，CTA緒可使用載 入與儲存虛擬指令（以及下述原子更新指令at〇m)以將資料 寫入到每CTA虛擬狀態空間及從中讀取資料。因此，一個 CTA緒可使用具有合適定義的目的地位址的st· shared指令 將資料寫入到.shared狀態空間；相同CTA内的另一 CTA緒 128625.doc -46- 200844853 I随後藉由在Id.shared指令中使用相同位址來讀取所述資 料。下述同步指令（例如，bar&memba〇可用來確保CTA線 、者間的貝料共用操作的正確序列，例如，產生資料的CTA 者在肩耗貝料的CTA緒讀取資料之前先寫入資料。類似 地，st.global及ld.glob_令可用於相同cta中的cta緒、 相同柵‘中的CTA及/或相同應用程式中的不同柵格之間的 協作及資料共用。 5·3·6·虛擬指令_程式控制 圖8G是列舉實例虛擬ISA中提供的程式控制操作的表格 此等控制刼作疋熟習此項技術者所熟悉的，且其使 得㈣設計人貝可將程式執行重定向。分支（㈣將程式流 重定向到目標位置<target>。在—些實施例中，藉由將字 =籤放置在虛擬ISA代碼中的目標指令之前並使用所述 標籤作為分支指令的目標識別符如阶來定義分支目 標。舉例而言，在一個實施例中： label: add.int32 d，vrl, vr2; 將add指令識別為具有標籤iabel的分支目標。在代碼中 的其它位置的指令： bra label; 將執行重定向到帶有標籤的指令。 ⑽與返回㈣指令支援函數及子例行程式調用； 識別函數或子例行程式。（在_個實施例中，”子例行程式" 就是其返回值被忽略的函數。)可使用^指示⑷二二 來宣稱函數f_e’且也提供定義所述函數的虛擬ΜΑ代 128625.doc -47- 200844853 碼。波形括號{}或其它分組符號可用來將定義函數或子 行程式的代碼與其它虛擬ISA代碼隔開 對於函數來說，可指定失*々而丨主/ 十, 扣疋苓數列表<rv>來識別應當將 值儲存在何處。對於函數 山致及子例仃程式來說，在引數列 ，中指定輸入引數。當執行斗料下一 ' 位址’當執行ret時，進入到所儲存位址的分支。 ㈣指令終止遭遇其的CTA緒。㈣指令調用處理哭定 義的或使用者定義的㈣例行程式。斷點（㈣州指令使執 行暫停，且可用於（例如）調試用途。無操作（卿）是在執行 時沒有影響的指令。例如，其可用來控射在多久 行下一操作。 5·3·7·虛擬指令-平行緒 圖8 Η是列舉在根據本發明實施例的實例虛擬〗s a中提供 的明確平行的虛擬指令的表格87G。此等指令支援cta執 订所需要的協作緒行為，例如在CTA緒之間交換資料。

屏障（㈣指令指* :到達其的CTA緒應當在執行任何進 一步的指令之前等待，直到其它所有CTA緒（相同CTA中 的）也已經到達同-㈣指令時為止。可在CTA程式中使用 任何數目的屏障指令。在一個實施例中，屏障指令不需要 任何參數（不論使用多少個屏障），因為必須在所有CTA緒 均到達第η個屏障之後，任何緒才可前進到第（η+ι)個屏 障，依此類推。 在其它實施例中，可例如藉由指定應當在特定屏障處等 待的CTA緒（或特定CTA緒的識別符）的數目來使屏障指令 128625.doc -48- 200844853 參數化。 另外其匕實施例提供”等待”及”不等待，，兩種屏障指令。 在等待屏障指令下，CTA緒一直等到其它相關CTA緒也已 I到達屏障為止；在不等待指令處，cta緒指示其已刻 達仁可在其它CTA緒到達之前繼續。在給定屏障處，〆 些CTA緒可能等待而其它ctA緒不等待。 在一些實施例中，bar虛擬指令可用來同步cta緒，所述 CTA緒正在使用共用記憶體狀態空間進行協作或共用資 料。舉例而言，假設一組CTA緒（其可包含cTA的一些或全 口P緒）各自在每線緒變數（例如，fp32虛擬暫存器變數 myData)中產生一些資料，然後讀取集合中的另一 cTa緒 產生的資料。指令序列： st.shared.fp32 myWriteAddress, myData; bar;

Id.shared.fp32 myData, myReadAddress; 提供所需的行為，其中myWriteAddress及myReadAddress是對 應於.shared狀態空間中的位址的每線緒變數。在每個CTA 緒將其產生的資料寫入到共用記憶體之後，其一直等到所 有CTA緒已經儲存其資料為止，然後繼續從共用記憶體中 讀取資料（其可能已經由不同的CTA緒寫入）。 記憶體屏障（membar)指令指示每個CTA緒應當等待其先 前請求的記憶體操作（或至少所有寫入操作）完成。這個指 令保證在membar指令之後發生的記憶體存取將看到在其之 前的任何寫入操作的結果。meinbar指令在一個實施例中 128625.doc -49- 200844853 使用可選的狀態空間名稱<space>以將其範圍限制於瞒準 指定狀態空間的記憶體操作，所述指定狀態空間應當是圮 fe、體狀怨空間（例如，不是·reg或·sreg狀態空間）。如果未 指定任何狀態空間名稱，那麼CTA緒等待瞄準所有記憶體 狀悲空間的所有待決操作完成。

原子更新（atom)指令致使對由引用<ref>.別的共用變數 a的原子更新（讀取-修改_寫入）。所述共用變數&可處於任 何共用狀態空間中，且與其它記憶體引用一樣，可使用各 種疋址模式。舉例而言，<ref>可為以下中的任一者：命 名的可定址變數a、對保存變數a的位址的暫存器的引用、 對保存將添加到偏移值（提供為立即操作數）以定位變數^的 位址的暫存器的引用，或變數a的立即絕對位址。CTA緒 將變數a從共用狀態空間位置載入到目的地暫存器4中，接 著使用對運算數a及（取決於操作）第二及第三運算數b、e執 行的指定操作<〇P>來更新變數a，並將結果儲存回由 識別的位置。目的地暫存保持a的原先載入的值。原子 地執行載入、更新及儲存操作，保證沒有其它任何CTA绪 在緒執行原子更新時存取變數a。在_個實施例 中’交數α限於.global或.shared狀態空間，且可用與上述 載入及儲存操作一樣的方式來指定。 ^ 在一些實施财，只有料的㈣可料原子更新來執 订。舉例而言，在一個實施例中，如果a是浮點類型那麼 Γ指定以下操作<°P>:將咖相力w的最小值或 取大值來替換a ;以及三元比較與交換操作，其在&等於b 128625.doc -50- 200844853 的情況下用C替換a，且否則 r . “更3保持不變。對於整數a，可 又援頭外知作，例如遁曾 運异數3與{3之間的按位的ad、〇r及 以及使運算數磷增或遞 或操作組合。 也了支杈其匕原子㈣ vote指令在預定義組 ^ ^ 琛、、者間對布林（例如，.bl類 i )運异數a執行歸約運算< 。 在 個實她例中，虛擬架 構指定CTA緒在SIMD群组中勃杆 八 ςτλ/ίη„, f、、中執仃，且預定義的群組對應於 SIMD群組；在其它實施例中，复 、 一匕群組的CTA緒可由虛擬 罙構或程式設計人員來定義。纟 我卸約運异<op>要求基於在群 組中的CTA線緒間對運算數3的歸約以及.外㈣符指定 的歸約運算將結果值攸置成布林真或布林假。在一個實 施例中，所允許的歸約運算是：⑴aU，其中如果a對於群 組中的所有CTA緒為真則d為真，且否則為假；⑺叫， 其中如果a對於群組中的任何CTA緒為真則d為直；以及 (3).Uni ’其中如^對於群組中的所有活動的CTA緒具有相 同值（真或假）則d為真。 5·3·8·虛擬指令-斷言執行 在一些實施例中，虛擬ISA支援任何指令的斷言執行。 在斷言執行中，將布林”保護斷言（g樹d Predicate)"值與指 令相關聯’且指令只有當在執行的時候保護斷言評估為直 時才執行。 ^ 〃 π在實例虛擬ISA中，保護斷言可為任何丨位布林虛擬暫存 為變數（本文中表示為P)。II由將斷言保護@p或不斷言保 護®!P放置在指令的操作碼前面來指示斷言執行。例如藉 128625.doc 200844853 由將p識別為產生了布林結果的指八 中的’指令)的目的地暫存器來在斷言暫存哭表圖一8〇 值。在遇到财或_保護斷言時，虛擬處理器讀取= 器。對於@P保護，如果p為真，則執行指令 : 過；對於㈣保護，如果P為假則執 、]將，、跳 在遇到斷言指令的每―執行 此’-些CTA緒可能執行斷言指令而其它cta緒不：行。 f 些實施例中，可將❹設置為指令執行。舉例而 :二二⑽:：:8A’W特定虛擬指令可能接受指定 σ暫存為作為輸出的參數；此種指令基於指令結果的某 種性質來更新指定的斷言暫存器。舉例而言，斷言暫存哭 可用來指示算術運算的結果是不是特殊數字(例如，零:。 無窮大或IEEE 754浮點運算中的非數字）等等。 ’ 6·虛擬指令翻譯器。 如參看圖4所述，虛擬指令翻譯器412猫準特定的平△架 構。虛擬指令翻譯器412可實施為⑽如)在例如圖二 1〇2等的處理器上執行的軟體程式，所述虛擬指令翻譯哭 412接收虛擬1SA代碼並將其轉譯成目標！SA代碼41/ 斤述目仏IS A代碼41 4可在虛擬指令翻譯器4丨2瞒準的特定 木構上執行(例如，藉由圖…卿j叫。虛擬指令翻 厚J12將在虛擬ISA代碼41()中宣稱的虛擬變數映射到可 用的儲存位置上，其中包含處理器暫存器、晶片上記憶 體、嶋記憶體等。在一些實施例中，虛擬指令翻譯器 412將每個虛擬狀態空間映射到特定類型的儲存設備上。 128625.doc • 52 - 200844853 舉例而言，可將.reg狀態空間映射到緒特定的資料暫存器 上，將.shared狀態空間映射到處理器的可共用記憶體上， 將.global狀態空間映射到分配給應用程式的虛擬記憶體區 域，等等。其它映射也是可能的。 將虛擬ISA代碼410中的虛擬指令轉譯成機器指令。在一 個實施例中，虛擬指令翻譯器412經組態以將每個虛擬isa 指令映射成相應的機器指令或機器指令序列，這取決於在

將執行CTA緒的處理器的指令集中是否存在相應的機器指 〇 虛擬指令翻譯器4丨2也將CTA緒映射到目標平台架構中 的”實體”緒或處理上。舉例而言，如果目標平台架構支援 至少…個併發緒，那麼每個CTA緒可映射到一個實體緒 上，且虛擬指令翻譯器412可為單個cta緒產生虛擬指令 代碼’並預期目標平台440將為具個唯—識別符的n〇個 緒執行代碼。如果目標平台架構支援少於η❶個緒，那麼虛 擬指令翻譯器412可產生虛擬ISA代碼41〇，戶斤述虛擬似: 碼410併入有對應於多個CTA緒的指令，並預期這個代碼 將對母個CTA執行一次，因而將多個cta緒映射到 體緒或處理。 貝 更：定而言，將與資料共用（例如，存取.sh一 或，global狀態空問的哉人 _ , 間的载入、儲存及原子更新指令）及/或協 作緒行為（例如，圖8 Η中的展 " a 宁的屏卩早、原子更新及其它指 關的虛擬指令韓等出德抑4t "轉厚成機益指令或機器指令序列

執行而優化的目桿平a加拔女& A + 口木構有利地包含硬體支援的屏障指 128625.doc -53 - 200844853 令，例如其中用指八 已經到達屏障指令的：’中的計數器及/或暫存器來計數 屏障處等待時發布緒的進一步二m止在緒在 不提供對緒同步的直接 ："”匕目‘結構可能 線緒間通^技術（例如 匕 ^ ^號$、記憶體中的狀離陳列笙、 來創建所需的行為。 單列專） 也將斷言指令轉譯成機 战械杰才日令。在一些實例中， 體直接支援斷言執行。在 不 的分支指令等一起儲在十 韦條件 &儲存在（例如）處理器暫存器中， 條件的分支指令等用來确門斬六 所述有 八 灼問暫存器並藉由有條件地圍繞斷 “"產生/刀支而創建所需的執行時間行為。 圖9是使用根據本於明每 I月貝苑例的虛擬指令翻譯器的過程 900的流程圖。在步驟咖處，程式設計人員用 寫CTA程式代碼。在_個杏 D〇 ° 隹個貝轭例中，所述CTA程式代碼定 Ο 義单個CTA緒的所需行為，且可使用細（包含CTA ID及/ 或柵格ID)作為參數來定義或控制cta緒的行為的各個能 樣。舉例而言’可將要讀取或寫人的共用記憶體位置確二 為緒ID的函數，使得相同CTA中的不同cta緒將對丘用記 憶體中的不同記憶體位置進行讀取及/或寫入。在二個實 細例中’包含CTA程式代碼作為應用程式代碼（例如，圖* 的程式代碼402)的-部分。除了定義CTA緒行為之外，库 用程式代碼還可定義CTA及/或栅格、設置輸入及輪出資料 集等。 、 在步驟904處，編譯器（例如，圖怕編譯器4〇8)根據高 128625.doc -54- 200844853 階語言代碼產生定義單個（虛擬）CTA緒的行為的虛擬ISA代 碼。如果代碼包含CTA程式代碼及其它代碼兩者，那麼編 譯器408可將CTA程式代碼與其餘代碼分㈤，以便只使用 CTA程式代碼來產生虛擬似代碼。可使用習知的用於將 用種σ編寫的私式代碼編譯成另一（虛擬）語言的技 術:應注意，由於所產生的代碼是使用虛擬語言，所以編 澤器不需要束缚於特定硬體或針對特定硬體而優化。編譯 器可優化根據特定的輸入代碼序列而產生的虛擬似代碼 ⑴女t偏好虛擬ISA指令的較短序列）。可將虛擬ISA中 的程式代碼儲存在碟片上的記憶體中及/或分布料種各 樣的平。架構，其中包含在實體上不同於圖3的虛擬架構 00的、、Ό構。虛擬ISA中的代碼是獨立於機器的，並且可在 任何有虛擬指令翻譯器可用的目標平台上執行。在替代實 她例中’程式設計人員可用虛擬isa直接編寫CM程式代 ，、，或者虛擬ISA代碼可由程式自動產生；如果程式代碼 初創建為虛擬IS A代碼’那麼可省略編譯步驟9〇4。 在f驟9〇6處，虛擬指令翻譯器（例如，圖4的翻譯器 4⑺項取虛擬ISA代碼’並以目標ISA產生可在目標平台上 執^ 丁的代碼。★盘★罢up 〆、、、' α睪的不同的是，虛擬指令翻譯器瞄 疋的（真實）平Α赵士塞 、 口 ’、冓且有利地經組態以調適並優化目許 ISA代碼以便在結 不 至少η。個緒的一個，實二取好的“。在目標結構支援 個灵靶例中，虛擬指令翻譯器產 程式，目標绪葙々叮丄 王曰知緒 切由〜個緒中的每—者併發執行以實現 在另一只施例中，虛擬指令翻譯器產生目標程式， 128625.doc -55- 200844853 所述目標程式使用軟體技術（例如，指令序列）來仿真n〇個 併發緒’線緒各自執行對應於虛擬IS A代碼的指令。翻譯 可在程式安裝時、在程式初始化期間或在程式執行期間 在及時的基礎上操作。 在步驟908處，目標平台中的處理器（例如，圖1的ppu 122)執行目標ISA代碼以處理資料。在一些實施例中，步 驟908可包含將命令及狀態參數提供給處理器，以便控制 其行為，如下文進一步描述的。 將明白，過程900是說明性的，且可進行更改及修改。 描述為循序的步驟可平行執行，步驟次序可更改，且步驟 可被修改或組合。舉例而言，在—些實施例中，程式設計 人員可直接使用虛擬ISA來編寫cta程式代碼，從而無需 產生虛擬ISA代碼的編譯器。在其它實施例中，將程 式代碼編寫為較大的應用程式的一部分，且所述〔Μ程式 代碼還包含（例如）定義要執行以解決特定問題的CTA及/或 格的維度的代碼。在—個實施例中，只將代碼中那 些闡述CTA程式的部分蝙譯到虛擬BA代碼中 部分編譯到其它（真實或虛擬）指令集中。 、 :其它實施例中，一個虛擬指令翻 Γ於不同9標平台的目榡代碼的多個版本。舉:ΐ 翻譯器可產生高階語言（例★ 牛例而口 機器代碼及/或用於使用軟^ C)的程式代碼1於ppu的 多核CPU的機器代碼。s術來仿真PPU行為的單核或 7·虛擬執行驅動器 128625.doc -56- 200844853 在一些實施例中，使用虛擬IS A代碼410及虛擬指令翻譯 器412來產生要對CTA的每個緒執行的CTA程式代碼。就圖 2A-2B的程式設計模型來說，指定CTA程式會為每個CTA 緒204定義處理任務。為了完成所述模型，也有必要定義 CTA 202的維度、柵格中的CTA的數目、要處理的輸入資 料集等。此種信息在本文中稱為”CTA控制信息’’。 如圖4所示，在一些實施例中，應用程式402藉由使用對 虛擬庫404中的函數的調用來指定CTA控制信息。在一個 實施例中，虛擬庫404包含各種函數調用，程式設計人員 可經由函數調用來定義CTA或CTA柵格並指示何時應當開 始執行。 圖10是列舉實例虛擬庫404中可用的函數的表格1000。 第一群組的函數涉及定義CTA。具體來說，initCTA函數是 第一個要調用以創建新CTA的函數。這個函數使得程式設 計人員可定義CTA的維度（ntid.x，ntid.y，ntid.z)，並向新 CTA分派識別符cname。SetCTAProgram函數指定要由CTA cname的每個緒執行的CTA程式；參數pname是對應於所需 的CTA程式（例如，用虛擬IS A代碼編寫的程式）的邏輯程式 識別符。SetCTAInputAiray函數使得程式設計人員可指定 全局記憶體中CTA cname將從中讀取輸入資料的源位置（開 始位址及大小），且setCTAOutputArray函數使得程式設計 人員可指定全局記憶體中CTA cname將向其寫入輸出資料 的目標位置（開始位址及大小）。使用setCTAParams函數來 設置CTA cname的執行時間常量參數。程式設計人員向函 128625.doc -57- 200844853 數提供參數列表，例如作為（名稱，值）對。 在個貝鈀例中，setCTAParams函數也可由編譯器4〇8 在產生虛擬ISA代碼41〇時使用。由於setCTAParams函數定 義CTA的執行時間參數，所以編譯器4〇8可將這個函數解 口睪為將每個參數定義為·param狀態空間中的虛擬變數。 表才u 1000還列舉與定義CTA的栅格有關的函數。

函數疋苐個调用來創建新柵格的函數。這個函數使得程 式叹计人員可疋義柵格的維度（⑽aid』，⑽心^ nctaid.z)，識別將在栅格上執行的cTA cname 並將識別 符gname分派給新定義的柵格。setGridlnputArray及 setGridOutputArray函數與CTA級函數相似，允許針對栅格 中的所有CTA的所有緒定義單個輸入及/或輸出陣列。 setGridParams函數用來為柵格gname中的所有cta設置執 行時間常量參數。料器4〇8可將這個函數解譯為將每個 參數定義為.const狀態空間中的虛擬變數。 laimchCTA及launchGrid函數指示應開始指定c丁a㈣㈣ 或柵格gname的執行。 虛擬API也可包含其它函數。舉例而言，一些實施例提 供可用來協調多個CTA的執行的同步函數。舉例而言，如 果第一CTA(或柵格）的輸出將用作第二CTA(或柵格）的輸 入，那麼API可包含一個函數（或啟動函數的參數），可經 由所述函數來指示虛擬執行驅動器：在第一CTA(或栅格） 的執行完成之前不應啟動第二CTA(或柵格）。 根據本發明的實施例，表格1000中的任何或所有函數調 128625.doc -58- 200844853 用均可包含在應用程式中，所述應用程式也定義要執行的 CTA程式（或者如果在應用程式中存在多個CTA則要定義多 個CTA程式）。在編譯時，將函數調用視為對應用程式介面 (API)庫404的調用，因而產生虛擬Αρι代碼4〇6。 使用實施虛擬庫中的每個函數的虛擬執行驅動器418來 實現虛擬API代碼。在-個實施例中，虛擬執行驅動器川 是在圖1的CPU 102上執行的驅動器程式，驅動器程式控制

V 實現CTA緒的PPU 122。實施圖1〇的表格1〇〇〇中的各種函 數調用，使得其導致驅動器經由推送緩衝器（州心邮向 PPU 122提供命令。在另-實施例中，cpu執行_或多個 程式以實現CTA，且虛擬執行驅動器418設置參數並控制 CPU對此種程式的執行。 工 ’且變更及 此項技術中 將瞭解’本文中描述的虛擬API是說明性的 修改是可能的。可支援其它函數或函數組合。 已知的虛擬API技術可適合於本發明的用途。 進一步的實施例 雖然已經參照特定實施例描述了本發明，但熟習此項技 術者將認識到許多修改是可能的。舉例而言，不需要本文 中描述的特定虛擬架構、虛擬指令及虛擬Αρι函數；可用 其它支援併發、協作緒的虛擬架構、指令及數來代 替。此外，上述實施例可引用所有區塊具有相同數目的元 件、所有CTA具有相同數目的緒且執行相㈣八程式等的 情況；在一些應用程式(例如1中使用多個依賴性栅格 的f月況）中’可能需要使不同柵格中的CTA執行不同⑽程 128625.doc -59- 200844853 式或具有不同數目及/或大小的柵格。 雖然本文中引用，，協作緒陣列”，但應瞭解…實施例 可使用其中不支援併發緒之間的資料共用的緒陣列；在支 援此種：料共用的其它實施例中’針對給定應用程式定義 的緒可能或可能不在實際上共用資料。 此外’雖然上述實施例可將緒陣列稱為具有多個緒，但 應瞭解，在”退化，，的情況下’緒陣列可能只具有_個緒。 2此，本發明可應用於在要在具有—或多個單線緒或多線 1心的CPU上執行的程式中提供可擴展性。藉由使用本 文中描述的技術，可用可在任何數目的可用CPU核心上分 布緒（例如，使用操作系統功能性）而無需對虛擬ISA代碼 的修改或重新編譯的方式編寫程式。 本文中使用術語”虛擬”及”真實，，來反映程式設計人員用 來描述問題解決方案的概念性程式設計模型與可在上面最 終執行程式的實際電腦系統的脫離。”虛擬,，程式設計模型 ^其_聯的結構使得程式設計人員可對平行的處理任務 =π p白視角’且應瞭解，可能存在或可能不存在其組件 :一對應地映射到本文中描述的虛擬架構組件的實際計算 :統或裝置。有利地將包含虛擬ISA代碼及虛擬納代碼的 业擬代馬貝現為用—種語言編寫的可能或可能不與任何· =處理裝置的指令集―對應的代碼。與所有程式代碼一 :本文中所指的虛擬代碼可儲存在有形的媒體(例如， _體或W )中、藉由網路傳輸等。 併入有本發明的各種特徵的電腦程式—包含但不限於虛 128625.doc -60- 200844853 擬ISA及/或虛擬ΑΡΙ伐民 ^ ^ ,. 哭 才日令翻譯器、虛擬驅動 亚擬函數庫等__可編碼在各種電腦可讀媒髀 上以供儲存及7或傳輸；合適的媒體包含磁碟或磁帶^ 學儲存媒體’例如光碟（CD)或DVD(多功能數位光碟 =憶：等。此種程式也可使用適合經由符合各種協定的 有，'[先學及/或無線網路(包含網際網路)傳輸的载波信 T。用程式代碼編碼的電腦可讀儲存媒體可用相容裝 裝，或者程式代碼可與其它裝置獨立地提供（例如， 網際網路下載）。 '、、工田 此外’本文中可將特定動作描述為由”程式設計人員 出。預期程式設計人員可以是人類、藉由很少或沒有人為 介入來產生程式代碼的自動過程、或人類交互與自動 自動處理的用以產生程式代碼的任何組合。 Ο 此外’雖然本文中描述的實施例可能引用了特定目標平 台的特徵，但本發明不限於此等平台。實際上，虛擬:構 可用硬體及/或軟體組件的任何組合來實現 術者將明自，可預期㈣虛擬㈣的不同實現在;^支/ 或輸出量態樣不同；’然而，此種差異與本發明無關。 因此’雖，然已經參考特定實施例描述了本發明，但應明 白，本發明並不意圖涵蓋屬於隨附申請專利範圍内的所 修改及等效物。 【圖式簡單說明】 圖1是根據本發明實施例的電腦系統的方塊圖。 圖2A及職明本發明實施例中使用的程式料模型中 128625.doc 61 200844853 的栅格、緒陣列及緒之間的關係。 圖3是根據本發明實施例的虛擬架構的方塊圖。 圖4是根據本發明實施例使用虛擬架構來操作目標處理 器的概念模型。 圖5疋列舉由根據本發明實施例的虛擬指令集架構（ι§Α) 定義的特殊變數的表格。 圖6是列舉在根據本發明實施例的虛擬ISA中支援的變數 類型的表格。 f :丨 ^ 圖7是列舉在根據本發明實施例的虛擬ISA中支援的虛擬 狀態空間的表格。 圖8A-8H是列舉在根據本發明實施例的虛擬ISA中定義 的虛擬指令的表格。 圖9是根據本發明實施例的用於使用虛擬指令翻譯器的 過程的流程圖。 圖10是列舉根據本發明實施例的在虛擬庫中可供虛擬執 行驅動器使用的函數的表格。 【主要元件符號說明】 100 電腦系統 102 CPU 104 系統記憶體 105 記憶體橋 106 通信路徑 107 I/O橋 108 使用者輸入裝置 128625.doc -62- 200844853

C 110 顯示器裝置 112 平行處理子系統 113 路徑 114 系統碟片 116 切換器 118 網路配接器 120 内插卡 121 内插卡 122 平行處理單元 124 平行處理（PP)記憶體 202 CTA 204 緒 300 虛擬架構 302 虛擬處理器 304 全局記憶體 306 如端 308 虛擬核心 310 處理引擎 311 特殊暫存器 312 指令單元 314 區域暫存器 316 共用記憶體 318 參數記憶體 320 虛擬驅動器 -63 - 128625.doc 200844853

C 322 記憶體介面 400 概念模型 402 應用程式 404 虛擬API函數庫 406 虛擬A PI代碼 408 編譯器 410 虛擬IS A代碼 412 虛擬指令翻譯器 414 目標ISA代碼 416 虛擬執行驅動器 418 目標API命令 430 指令單元 432 PPU驅動器 434 前端 440 平台 500 表格 600 表格 700 表格 800 表格 810 表格 820 表格 830 表格 840 表格 850 表格 -64- 128625.doc 200844853 860 870 900 1000 表格 表格 過程 表格 128625.doc -65

Claims (1)

1. 200844853 十、申請專利範圍： 一種用於定義一平件泠 提供笛操作的方法，該方法包括： k供弟一程式代碼， 一 μ 口 耘式代碼定義待針對協作 者的一陣列中的複數個虛擬緒中的每一者執行的 操作序列； 7母考執仃的 式定羞料 譯成—虛擬緒程式，該虛擬緒程 針對該複數個虛擬財的—代妹虛擬緒執行 的赠緒指令的-序列，該每線緒指令之序列包含至少 =指令’該指令定義該代表性虛擬緒與該複數個虛擬 中的—或多個其它虛擬緒之間的協作行為；以及 儲存該虛擬緒程式。 2. 如請求項1的方法，其進一步包括： 目標平台架構的 將忒儲存的虛擬緒程式轉譯成符合一 一指令序列。 3·如請求項1的方法，其進一步包括： 名提供弟二程式代碼，該第二程式代碼定義協作的虛擬 、的陣列’该陣列經調適以處理一輸入資料集以產生 輸出貪料集，其中該陣列中的每個虛擬緒併發地執行 該虛擬緒程式； 、〜罘二程式代碼轉換成一虛擬函數庫中的一函數調 序歹〗邊庫包含初始化並致使執行協作的虛擬緒的該 陣列的虛擬函數；以及 Λ 儲存該函數調用序列。 4·如明求項3的方法，其進一步包括： 128625.doc 200844853 5. =儲存的虛擬緒程式及該函數則序列轉譯成可在 平台架構上執行的程式代碼，該可執行的程式代 義執行該協作的虛擬緒陣列的一或多個：。' 如味未項4的方法，其進一步包括： 爷 在符合該目標平台架構的—電m 的程式代碼，因而產生該輸出資料集；以及 丁 6. 將該輪出資料集儲存在一儲存媒體中。 如明求項1的方法，其中琴. 母線緒指令序列包含用以在 =列中的-特定點處暫停對該代表性虛擬緒的操作的 =止Γ該等其它虛擬緒中的—或多者到達該特定點 日守馮止的指令。 如明求項1的方法’其中該每線緒指令序列包含一用以 8. ㈣代表性虛擬緒將資料儲存在該等其它虛擬緒中的一 或夕者可存取的一共用記憶體中的指令。 9. 士明求項1的方法’其中該每線緒指令序列包含一用以 ^該^表性虛擬緒自動讀取及更新儲存在該等其它虛擬緒 a、、或夕者可存取的-共用記憶體中的資料的指令。 :明求項1的方法，其中該虛擬緒程式包含一變數定義 σ句。亥變數定義語句定義複數個虛擬狀態空間之一者 2的父數，其中該複數個虛擬狀態空間中的不同狀態 工:對應於該等虛擬緒之間的資料共用的不同模式。 d項9的方法中該等資料共用模式包含一每線 、、者不共用模式及一全局共用模式。 η·如請求項9的方法，其中該等#料共用模式包含一每線 128625.doc 200844853 、者不,、用杈式、虛擬緒的一個 及-全局共用模式。 ，、用抵式，以 12.:請：項9的方法，其中該等資料共用模式包含每線 緒不共用模式、卢 、匕s —母線 挪心、- 緒的一個陣列内的一共用楔式、Θ 擬緒的m數個陣列、式虛 模式。 1的、用拉式，以及-全局共用 13 · —種用於操作一 乍目標處理器的方法，該方法包括. 提供輪入程式代碼，其包含一第一· 定義待對虛擬緒的 μ弟一部分 執行的_&4/^的複數個虛擬緒中的每一者 的一操作序；；:處理-輸入刚以產生-輪出資料集 = 呈式代碼進一步包含一第二部分 疋義该虛擬緒陣列的一維度； 1刀 將:輸广…的該V部分編譯成—虛擬緒程 虛擬緒程式定義待針對該複數個虛擬緒中的一代 表性虛擬緒執行的一每 列包含至小4母線以曰々序列，該每線緒指令序 夕J已3至乂 一個指令，該 複數個虛擬緒令的一ρΜ義錢表性虛擬緒與該 為； 的或夕個其它虛擬緒之間的-協作行 座Γ輸人以代碼的㈣二部㈣換成對-虛擬函數 庫的一函數調料列，該庫包含虛擬函數 數初始化及致使執行該協作的虛擬緒陣列. 函 =虛擬緒程式及該函數調料列轉 平台架構上執行的程式代碼，該可執行的程式代碼= 128625.doc 200844853 執行該協作的虛擬緒陣列的一或多個真實緒； f符合該目標平台架構的_電腦系統上執行該可執行 的程式代碼，因而產生該輸出資料集；以及 將。亥輸出資料集儲存在一儲存媒體中。 14. 如請求項1 3的方法， 包含定義該虛擬緒陣 碼。 其中該輸入程式代碼的該第二部分 列的兩個或兩個以上維度的程式代

   15·如請求項14的方法’其中該輸入程式代碼的該第二部分 進一步包含： ^函數呑周用，其定義虛擬緒陣列的一柵格的一或多個 、准度，其中該柵格中的每一陣列均將被執行。 16’ ：明求項13的方法，其中該目標平台架構包含一主處理 卯及共處理器，且其中該翻譯的動作包含： 將該虛擬緒程式轉譯成可由該共處理器上定義的複數 個緒平行執行的程式代碼；以及 、將該函數則序列轉譯成對該共處理H的-驅動器程 、周用序列’其中該驅動器程式在該主處理器上執 行。 17·如 的方法’其中該目標平台架構包含一中央處 理單元（CPU)，且其中該轉譯的動作包含： 成^虛擬緒程式及該函數調用序列的至少—部分轉譯 目' 矛式代碼，该目標程式代碼使用少於該虛擬緒數 一的-數目之CPU線緒執行該虛擬緒陣列。 18. -種用於操作一目標處理器的方法，該方法包括： 128625.doc 200844853 獲得虛擬緒程式，該虛擬緒程式定義待對 列中的複數個虛擬緒中的一代表性虛二 :: 指令的-序列，該陣列 丁的母線緒 生一輸出資料集； 周適以處理-輸入資料集以產 該每線緒指令序列包含至少_個指令，該至少 令定義該代表性虛擬緒與該複數個虛擬料的—或多： 其它虛擬緒之間的一協作行為； U 獲得定義該虛擬緒陣列的維度的額外程式代碼； 將該虛擬緒程式及該額外程式代碼轉譯成可在該目標 平台架構上執行的程式代碼，該可執行的程式代碼定義 執行該虛擬緒陣列的一或多個平台線緒；及 在符合該目標平台架構的一電腦系統上執行該可執行 的程式代碼，因而產生該輸出資料集並將該輪出資料集 儲存在一記憶體中。 19.如請求項18的方法，其中該獲得該虛擬緒程式的動作包 含： 接收用一高階程式設計語言編寫的源程式代碼；以及 編譯該源程式代碼以產生該虛擬緒程式。 20·如請求項18的方法，其中該獲得該虛擬緒程式的動作包 含： 從一儲存媒體中讀取該虛擬緒程式。 21·如請求項18的方法，其中該獲得該虛擬緒程式的動作包 含： 級由一網路從一遠端電腦系統接收該虛擬緒程式。 128625.doc