版權(quán)通告

本文包含的是受到版權(quán)保護的材料。當(dāng)它出現(xiàn)在專利和商標(biāo)局專利文件或記錄中時，版權(quán)擁有者不反對任何人復(fù)制該專利公開，但無論如何在其它方面保留對版權(quán)的所有權(quán)利。

本文描述的實施例一般涉及異構(gòu)計算系統(tǒng)。具體地說，本文描述的實施例一般涉及異構(gòu)計算系統(tǒng)中的工作竊取。

背景技術(shù)：

計算機除了部署為中央處理單元(cpu)的通用處理器之外還經(jīng)常具有一個或多個硬件加速器裝置。此類硬件加速器裝置的一個廣泛使用的示例是圖形處理單元(gpu)。gpu傳統(tǒng)上已經(jīng)用作主要驅(qū)動一個或多個顯示器的圖形子系統(tǒng)的一部分。gpu幫助從cpu卸載工作，由此允許cpu資源用于其它任務(wù)。cpu通常實現(xiàn)為通用處理器，并且一般將適合于執(zhí)行各種不同類型的算法(例如一般性質(zhì)的處理)。相比之下，gpu經(jīng)常實現(xiàn)為專用處理器或至少更專業(yè)化的處理器，其一般更適合于執(zhí)行有限數(shù)量的專業(yè)化算法(例如圖形相關(guān)算法)。gpu經(jīng)常具有高度并行的計算硬件，其傾向于允許它們快速處理圖形數(shù)據(jù)。

近來，已經(jīng)認(rèn)識到gpu可用于其它非圖形應(yīng)用。例如，gpu上的通用計算(gpgpu)一般表示利用gpu執(zhí)行傳統(tǒng)上已經(jīng)預(yù)留用于cpu的類型的通用型計算。經(jīng)?？刹捎胓pu的高度并行計算硬件來顯著地加速此類計算。當(dāng)以這種方式使用時，至少當(dāng)正確實現(xiàn)時，gpu可幫助顯著加速計算密集應(yīng)用。

然而，使用gpu加速非圖形應(yīng)用的一個挑戰(zhàn)(例如在gpgpu的情況下)涉及工作調(diào)度。為了有效地利用cpu和gpu，重要的是在cpu與gpu之間以有效方式調(diào)度和分布工作。在一個可能方法中，cpu可以僅僅向gpu卸載或指配特定任務(wù)。cpu然后可等待或執(zhí)行其它任務(wù)，而gpu完成指配的任務(wù)并提供結(jié)果。然而，在這種方法中，在處理相同工作負(fù)載時在cpu與gpu之間沒有實際協(xié)作。也就是說，cpu和gpu可能不同時對相同工作負(fù)載的不同部分工作。它實際上不是混合執(zhí)行。

根據(jù)另一方法，cpu和gpu可參與混合執(zhí)行，其中它們對在cpu和gpu上靜態(tài)調(diào)度的并行工作負(fù)載的不同部分協(xié)作地工作。此類混合執(zhí)行可提供在相同工作負(fù)載上同時利用cpu和gpu的潛在優(yōu)點。然而，傾向于有挑戰(zhàn)的是有效地在cpu與gpu之間調(diào)度工作負(fù)載的不同部分。例如，cpu和gpu可以不同時鐘速度操作，和/或可具有不同存儲器層級，和/或可具有不同的基于工作負(fù)載的性能特性?？稍赾pu和gpu之一上有效地執(zhí)行的代碼可能沒有在另一個上執(zhí)行得那么有效。其它競爭的工作負(fù)載和/或功率約束的存在可進一步傾向于使先驗有效地調(diào)度工作負(fù)載的能力復(fù)雜化。缺點是，如果工作負(fù)載未有效調(diào)度，則cpu和gpu之一可變得負(fù)擔(dān)過度，而另一個可能利用不足。

附圖說明

通過參考用于說明本發(fā)明實施例的如下描述和附圖可最好地理解本發(fā)明。在附圖中：

圖1是異構(gòu)計算或計算機系統(tǒng)的實施例的框圖。

圖2是同步工作竊取器系統(tǒng)的實施例的框圖。

圖3是在異構(gòu)計算機系統(tǒng)中竊取工作的方法的實施例的流程框圖。

圖4是圖示在cpu的實施例與gpu的實施例之間的雙向同步工作竊取的實施例的框圖。

圖5是可使用同步工作竊取器系統(tǒng)的實施例的環(huán)境的框圖。

圖6是異構(gòu)計算或計算機系統(tǒng)的示例實施例的框圖。

圖7是異構(gòu)計算或計算機系統(tǒng)的示例實施例的框圖。

具體實施方式

在如下描述中，闡述了大量特定細(xì)節(jié)(例如特定處理單元、異構(gòu)計算機系統(tǒng)、操作序列、邏輯分割/集成細(xì)節(jié)、系統(tǒng)組件的類型和相互關(guān)系等)。然而，要理解到，本發(fā)明的實施例可以在沒有這些特定細(xì)節(jié)的情況下實行。在其它實例中，眾所周知的電路、結(jié)構(gòu)和技術(shù)未詳細(xì)示出，以免模糊對此描述的理解。

圖1是異構(gòu)計算或計算機系統(tǒng)100的實施例的框圖。在各種實施例中，異構(gòu)計算系統(tǒng)可表示臺式計算機、膝上型計算機、筆記本計算機、上網(wǎng)本計算機、工作站、個人數(shù)字助理(pda)、智能電話、蜂窩電話、移動計算裝置、服務(wù)器、因特網(wǎng)器具或本領(lǐng)域已知的各種其它類型計算機系統(tǒng)或其它計算系統(tǒng)。

異構(gòu)計算系統(tǒng)包含至少兩個異構(gòu)(即不同)類型的硬件計算單元。硬件計算單元在本文也可被稱為計算單元。在所圖示的實施例中，異構(gòu)計算機系統(tǒng)包含第一類型的第一硬件計算單元101-1和第二不同類型的第二硬件計算單元101-2。第一和第二計算單元例如可通過互連111耦合在一起。其它實施例可包含三個或更多異構(gòu)計算單元。適當(dāng)計算單元的示例包含但不限于處理器、核、硬件線程、線程槽、能夠維持獨立執(zhí)行狀態(tài)的硬件等。

在一些實施例中，第一計算單元101-1可以是通用計算單元(或至少比第二計算單元101-2相對更通用)，而第二計算單元101-2可以不是通用計算單元和/或可以是專用計算單元(或至少比第一計算單元101-1相對更通用)。在一些實施例中，第一計算單元101-1可以是cpu、通用處理器和通用核中的一個，而第二計算單元可以不是。例如，第二計算單元101-2可以是圖形處理器(例如gpu、圖形協(xié)處理器、圖形核等)、硬件加速器裝置(例如專用加速器、固定功能加速器等)、密碼處理器、通信處理器、網(wǎng)絡(luò)處理器、專用處理器、專用核、高度并行專用硬件、數(shù)字信號處理器(dsp)、現(xiàn)場可編程門陣列(fpga)等中的一個。

在一些實施例中，第一計算單元101-1可以是cpu，而第二計算單元可以是圖形處理器(例如gpu、圖形協(xié)處理器、圖形核等)，盡管本發(fā)明的范圍不如此限制。在此類實施例中，異構(gòu)計算系統(tǒng)可表示gpgpu系統(tǒng)、cpu-gpu協(xié)作系統(tǒng)等。在這個具體實施方式中，經(jīng)常cpu和gpu被用作第一和第二異構(gòu)計算單元的示例，不過要認(rèn)識到，在備選實施例中，可替代地使用異構(gòu)計算單元的其它不同組合(例如相對更通用處理器連同從專用處理器、專用核、硬件加速器裝置、dsp、fpga等之間選擇的相對更專用處理器)。

再次參考圖1，當(dāng)部署在異構(gòu)計算系統(tǒng)中時，第一計算單元101-1和第二計算單元101-2各可操作以例如通過互連108、109與共享存儲器102耦合。共享存儲器對第一和第二計算單元可存取并由它們共享。在一些實施例中，共享存儲器可表示共享虛擬存儲器。共享存儲器或共享虛擬存儲器可表示實現(xiàn)在一個或多個存儲器類型的一個或多個存儲器裝置中的物理存儲器的一部分。在一些實施例中，共享存儲器可實現(xiàn)在動態(tài)隨機存取存儲器(dram)中，盡管本發(fā)明的范圍不如此限制。

在一些實施例中，對應(yīng)于第一計算單元101-1的第一工作隊列104-1和對應(yīng)于第二計算單元101-2的第二工作隊列104-2可存儲在共享存儲器102中。第一工作隊列104-1可以可操作以對用于第一計算單元101-1的工作接收并排隊。第二工作隊列104-2可以可操作以對用于第二計算單元101-2的工作接收并排隊。在圖示中為了簡化，僅示出兩個工作隊列，不過在一些實施例中，可能存在用于第一計算單元(例如用于多個核中的每個)的多個工作隊列和/或用于第二計算單元(例如可選地用于多個多處理器或其它組核中的每個)的多個工作隊列。

再次參考圖1，在一些實施例中，第一計算單元101-1(例如cpu、通用處理器、通用核等)可包含第一工作調(diào)度器模塊103-1。第一工作調(diào)度器模塊103-1可以可操作以調(diào)度第一計算單元101-1上的工作。在一些實施例中，第二計算單元101-2(例如gpu、圖形核、硬件加速器、專用處理器、專用核等)可包含第二工作調(diào)度器模塊103-2。第二工作調(diào)度器模塊103-2可以可操作以調(diào)度第二計算單元101-2上的工作。在一些實施例中，第一和第二工作調(diào)度器模塊可以可操作以調(diào)度第一和第二計算單元上的工作，以實現(xiàn)異構(gòu)計算單元上的公共工作負(fù)載的不同部分(例如數(shù)據(jù)并行工作負(fù)載的不同部分)上的混合執(zhí)行和/或協(xié)作/合作計算。例如，在一些實施例中，cpu和gpu可一起工作在gpgpu中。

作為示例，工作負(fù)載可被分成組塊、工作單元或其它部分。這些組塊、工作單元或部分可在第一和第二計算單元之間調(diào)度，并在對應(yīng)的第一和第二工作隊列中排隊。在一些實施例中，此類隊列例如可通過編碼一系列最小和最大索引范圍的數(shù)據(jù)并行操作來實現(xiàn)，其中工作的單個單元可表示為元組(例如min1,max1)。數(shù)據(jù)并行操作的總索引范圍可被分成工作的組塊。在一些情況下，組塊的大小可選地可選擇成使得它對應(yīng)于計算單元(例如gpu)的多個所述數(shù)量的硬件線程、simd道或核。更進一步說，它可選地可大于硬件線程、simd道或核的總數(shù)除以計算單元同時支持的屏障的數(shù)量。盡管不需要，但這可幫助確保工作的組塊有效地利用計算資源，并且初始調(diào)度通過硬件線程調(diào)度器是有效的，如果存在一個的話。組塊然后例如可在工作隊列之間均勻分布，或者基于各種算法/準(zhǔn)則分布。在運行時間期間，每一個計算單元可從其對應(yīng)工作隊列中檢索并處理組塊或工作單元。

這可繼續(xù)直到工作負(fù)載已經(jīng)完成。作為示例，完成可由所有隊列為空來指示。在一些情況下，可選地可使用終止令牌。例如，終止令牌可將最大整數(shù)值表示為終止令牌termempty。在所有計算工作者按確定次序竊取并且沒有工作者生成更多工作的實施例中，返回值empty一般將指示終止。在允許隨機竊取操作的實施例中，在它自身用完工作之后，足以讓至少一個工作者(例如cpu線程)按確定次序穿過所有隊列。如果它發(fā)現(xiàn)所有隊列空，則它可向所有隊列或至少向另一計算單元(例如gpu)的所有工作隊列寫empty終止令牌。這可幫助終止數(shù)據(jù)并行工作負(fù)載。

在一些實施例中，例如可執(zhí)行工作竊取，以便幫助改進負(fù)載平衡，提高性能，降低功耗等。術(shù)語“工作竊取”是在本領(lǐng)域使用的術(shù)語。本文所使用的術(shù)語“工作竊取”廣義上用于指將工作從一個計算單元重新指配或重新任務(wù)分配給另一計算單元，將工作從一個計算單元的隊列移動到另一計算單元的隊列，允許計算單元主張或負(fù)責(zé)之前由另一計算單元主張或由其負(fù)責(zé)的工作等。

再次參考圖1，在一些實施例中，第一計算單元101-1(例如cpu、通用處理器、通用核等)可選地可包含第一同步工作竊取器系統(tǒng)105-1的實施例，盡管這不是要求的。第一同步工作竊取器系統(tǒng)105-1可以可操作以從第二計算單元101-2執(zhí)行同步工作竊取，以便第一計算單元101-1做該工作，盡管本發(fā)明的范圍不如此限制。在一些實施例中，第一同步工作竊取器系統(tǒng)105-1可存取第一工作隊列104-1和第二工作隊列104-2。第一同步工作竊取器系統(tǒng)可用軟件、固件和硬件的任何組合實現(xiàn)。

在一些實施例中，第二計算單元101-2(例如gpu、圖形核、硬件加速器、專用處理器、專用核等)可包含第二同步工作竊取器系統(tǒng)105-2的實施例。在一些實施例中，第二同步工作竊取器系統(tǒng)105-2可以可操作以從第一計算單元101-1執(zhí)行同步工作竊取，以便第二計算單元101-2做工作。竊取的工作可從第一工作隊列104-1取得，并添加到第二工作隊列104-2。在一些實施例中，第二同步工作竊取器系統(tǒng)105-2可存取第一工作隊列104-1和第二工作隊列104-2。第二同步工作竊取器系統(tǒng)可用軟件、固件和硬件的任何組合實現(xiàn)。

在一些實施例中，工作竊取可基于第一工作隊列104-1和第二工作隊列104-2的當(dāng)前充滿。例如，在一些實施例中，如果第二工作隊列104-2變空，填充在閾值等級以下，或被第二同步工作竊取器系統(tǒng)105-2感知為未充分充滿，則第二同步工作竊取器系統(tǒng)105-2可從第一工作隊列104-1竊取工作，并將竊取的工作放入第二工作隊列104-2中。作為另一示例，在一些實施例中，如果第二同步工作竊取器系統(tǒng)103-2感知到第一工作隊列104-1是滿的，填充在閾值等級以上，或者另外太滿，則第二同步工作竊取器系統(tǒng)103-2可從過度填充的第一工作隊列104-1中竊取工作。在其它實施例中，當(dāng)?shù)谝还ぷ麝犃?04-1填充不足時和/或當(dāng)?shù)诙ぷ麝犃?04-2填充過度時，第一同步工作竊取器系統(tǒng)103-1可執(zhí)行類似或相互類型的工作竊取。還考慮了工作竊取的其它原因。

在一些實施例中，工作竊取在第一與第二計算單元之間在任一方向上可以是雙向的。例如，第一同步工作竊取器系統(tǒng)105-1可從第二計算單元101-2竊取工作用于第一計算單元101-1去做(例如將工作從第二工作隊列104-2移動到第一工作隊列104-1)，并且第二同步工作竊取器系統(tǒng)105-2可從第一計算單元101-1竊取工作用于第二計算單元101-2去做(例如將工作從第一工作隊列104-1移動到第二工作隊列104-2)。在一些實施例中，可執(zhí)行雙向相互同步的工作竊取。在一些實施例中，可選地可使用隨機基本上同時的雙向相互同步的工作竊取。在其它實施例中，第二計算單元101-2可選地可從第一計算單元101-1竊取工作，無需第一計算單元從第二計算單元竊取工作。這可能不提供為巨大優(yōu)勢，但如果期望的話，可幫助允許更簡單實現(xiàn)。在此類情況下，第一計算單元可選地可省略同步工作竊取器系統(tǒng)105-1。

再次參考圖1，在一些實施例中，可選地可執(zhí)行同步工作竊取操作106、107。在一些實施例中，可通過同步存取工作竊取隊列104-1、104-2和/或同步存取共享存儲器103來執(zhí)行同步工作竊取操作。例如，第二計算單元101-2可發(fā)出幫助實現(xiàn)同步工作竊取(例如同步工作竊取106)的同步(例如原子存取/操作110)。在一些實施例中，同步工作竊取操作可用于幫助防止兩個實體竊取相同工作和/或執(zhí)行相同工作。這還可幫助防止用于實現(xiàn)工作隊列的共享數(shù)據(jù)結(jié)構(gòu)的破壞，并且可幫助允許它們從一個有效狀態(tài)轉(zhuǎn)變到另一有效狀態(tài)。在一些實施例中，同步工作竊取操作可充分完成，因為工作可能不從隊列竊取，但不由竊取計算單元執(zhí)行。

在一些實施例中，可通過一個或多個存儲器存取同步原語和/或指令和/或操作執(zhí)行同步工作竊取操作。在一些實施例中，退出隊列和竊取操作可利用由第一和第二計算單元支持的一組存儲器存取排序原語/指令/操作。在一些實施例中，存儲器存取同步原語/指令/操作可在支持第一與第二計算單元(例如cpu與gpu)之間任何期望的此類存儲器存取同步原語/指令/操作的硬件上實現(xiàn)。適當(dāng)?shù)拇祟愒Z/指令/操作的示例包含但不限于存儲器柵欄和/或屏障宏指令、原子存儲器存取宏指令、opencl原子操作、cuda存儲器存取同步操作或在本領(lǐng)域已知的其它方法。

作為示例，在原子存儲器接入操作中，處理器可基本上同時(例如在相同總線循環(huán)中)讀存儲器位置和向存儲器位置寫。此類原子操作可幫助防止其它計算單元或外部裝置寫或讀存儲器位置直到原子操作完成之后。一般而言，原子存儲器存取操作將完全執(zhí)行或者根本不執(zhí)行。此類原子存儲器存取原語/指令/操作的示例包含但不限于讀-修改-寫、比較并互換、比較并交換、測試并設(shè)置、比較并設(shè)置、加載鏈路/有條件存儲指令等以及它們的各種組合。這些可通過宏指令、opencl原子操作、cuda存儲器存取同步操作或通過在本領(lǐng)域已知的其它方法進行。

在一些實施例中，這些可包含一個或多個存儲器存取柵欄指令。存儲器存取柵欄指令在本領(lǐng)域中(例如在一些架構(gòu)中)有時也被稱為存儲器存取屏障指令。此類存儲器存取柵欄和/或屏障指令的示例包含但不限于加載柵欄/屏障(例如lfence指令)、存儲柵欄/屏障(例如sfence指令)以及加載并存儲柵欄/屏障(例如mfence指令)等以及它們的各種組合。此類柵欄或屏障操作可通過宏指令、opencl操作、cuda操作或通過在本領(lǐng)域已知的其它方法實現(xiàn)。

有利地，此類工作竊取可幫助允許第一與第二計算單元之間的更有效的工作調(diào)度和分布，這可導(dǎo)致更好地利用資源和提高性能。因為允許工作竊取，并且特別是當(dāng)允許雙向工作竊取時，不要求能夠以高度有效方式在異構(gòu)計算單元之間靜態(tài)調(diào)度工作負(fù)載。如在背景技術(shù)部分中所討論的，此類靜態(tài)工作負(fù)載調(diào)度先驗傾向于由于若干原因(例如，由于潛在不同時鐘速度、不同存儲器層級、不同代碼執(zhí)行有效性、其它工作負(fù)載的存在、功率極限等)是困難的。而是，如果由于某種原因工作最初以低效方式調(diào)度，并且它導(dǎo)致一個計算單元變得負(fù)擔(dān)過度(或利用不足)，則可執(zhí)行工作竊取連同負(fù)載平衡，以便幫助重新平衡工作負(fù)載，并且由此幫助減輕各種計算單元的利用過度和/或利用不足。在其它實施例中，可選地可執(zhí)行工作竊取而不是負(fù)載平衡以幫助實現(xiàn)其它目的，諸如例如最大化處理速度，降低功耗等。

圖2是同步工作竊取器系統(tǒng)205的實施例的框圖。還示出的是共享存儲器202中的第一工作隊列204-1和第二工作隊列204-2。在一些實施例中，圖2的同步工作竊取器系統(tǒng)可包含在圖1的第二硬件計算單元101-2和/或異構(gòu)計算系統(tǒng)100中。備選地，圖2的同步工作竊取器系統(tǒng)可包含在類似的或不同的計算單元或異構(gòu)系統(tǒng)中。而且，圖1的第二硬件計算單元和異構(gòu)系統(tǒng)可包含與圖2的類似的或不同的同步工作竊取器系統(tǒng)。

同步工作竊取器系統(tǒng)205包含確定模塊220。確定模塊可操作以確定是否從第一類型的第一硬件計算單元(例如第一計算單元101-1)竊取工作222用于第二不同類型的第二硬件計算單元(例如第二計算單元101-2)。在一些實施例中，第二類型可比第一類型更專用。工作222可在對應(yīng)于第一硬件計算單元的第一工作隊列204-1中排隊。第一工作隊列在共享存儲器202中，其由第一和第二硬件計算單元共享。如所示，在一些實施例中，確定模塊可與第一工作隊列204-1和/或第二工作隊列204-2耦合，或者以其他方式通信。在一些實施例中，確定模塊可確定是否基于充滿或一個或多個工作隊列竊取工作。這可如在本文其它地方所描述的進行。

同步工作竊取器系統(tǒng)205還包含同步工作竊取器模塊221。同步工作竊取器模塊可操作以從第一硬件計算單元和/或第一隊列204-1竊取工作222，并將它作為竊取的工作223提供用于第二硬件計算單元，和/或?qū)⑺砑拥降诙ぷ麝犃?04-2。在一些實施例中，同步工作竊取器模塊可操作以通過對第一工作隊列204-1的同步存儲器存取210竊取工作。在一些實施例中，同步存儲器存取210可相對于對來自第一硬件計算單元的第一工作隊列的存儲器存取同步。

圖3是在異構(gòu)計算機系統(tǒng)中竊取工作的方法335的實施例的流程框圖。在一些實施例中，圖3的操作和/或方法可由圖1的同步工作竊取系統(tǒng)105-2和/或圖2的同步工作竊取系統(tǒng)205執(zhí)行，和/或在其內(nèi)執(zhí)行。這些系統(tǒng)的本文描述的組件、特征或特定可選細(xì)節(jié)可選地還適用于在實施例中可由這些系統(tǒng)和/或在這些系統(tǒng)內(nèi)執(zhí)行的操作和/或方法。備選地，圖3的操作和/或方法可由類似或不同的同步工作竊取系統(tǒng)執(zhí)行和/或在其內(nèi)執(zhí)行。而且，圖1和/或圖2的同步工作竊取系統(tǒng)可執(zhí)行與圖3的類似或不同的操作和/或方法。

該方法包含在塊336確定從第一類型的第一硬件計算單元竊取工作用于比第一類型更專用的第二類型的第二硬件計算單元。在一些實施例中，工作可在對應(yīng)于第一硬件計算單元的第一工作隊列中排隊。在一些實施例中，第一工作隊列可存儲在由第一和第二硬件計算單元共享的共享存儲器中。

該方法還包含在塊337竊取工作。在一些實施例中，竊取工作可包含執(zhí)行對第一工作隊列的同步存儲器存取。在一些實施例中，同步存儲器存取可相對于對來自第一硬件計算單元的第一工作隊列的存儲器存取同步。

圖4是圖示在cpu401-1的實施例與gpu401-2的實施例之間的雙向同步工作竊取406、407的實施例的框圖。cpu的所圖示的示例具有四個核，即cpu核0、cpu核1、cpu核2和cpu核3。其它cpu可具有更少或更多的核。在所圖示的實施例中，這些核中的每個都具有對應(yīng)于cpu的第一組工作隊列404-1中的不同工作隊列。具體地說，在所圖示的示例中，工作隊列ws_q[0]對應(yīng)于cpu核0，工作隊列ws_q[1]對應(yīng)于cpu核1，工作隊列ws_q[2]對應(yīng)于cpu核2，并且工作隊列ws_q[3]對應(yīng)于cpu核3。在其它實施例中，可使用核與工作隊列之間的多對一或者一對多的對應(yīng)關(guān)系。

cpu的所圖示的示例具有四個工作組，即工作組0、工作組1、工作組2和工作組3。這些工作組還可被稱為流播單元或多處理器。在圖示中，每個工作組具有16個simd道，盡管其它實施例可具有更少或更多的smid道。而且，其它gpu可具有更少或更多的工作組。例如，一些gpu具有大約數(shù)十個多處理器或工作組。這些工作組中的每個可包含一個或多個核，或者潛在地包含許多核。例如，一些gpu每個多處理器或工作組具有大約數(shù)十個核。在所圖示的實施例中，這些工作組中的每個共享對應(yīng)于gpu的第二工作隊列404-2。在其它實施例中，多個工作隊列可各用于一個或多個工作組。工作組404-1、404-2可在共享存儲器中。

在圖示中示出了雙向同步工作竊取操作406、407的示例。雖然cpu核0、cpu核1和cpu核3分別從它們自己的對應(yīng)工作隊列ws_q[0]、ws_q[l]和ws_q[3]工作，但cpu核2正在從gpu隊列404-2執(zhí)行同步工作竊取操作407。類似地，工作組0(亦稱流播單元0)正在從對應(yīng)于cpu核3的工作隊列ws_q[3]執(zhí)行同步工作竊取操作406。這只是一個說明性示例。工作竊取操作406、407可類似于本文描述的其它工作竊取操作，或與之相同。

圖5是可使用同步工作竊取器系統(tǒng)505的實施例的環(huán)境的框圖。如所示，在一些實施例中，同步工作竊取器系統(tǒng)可實現(xiàn)在常規(guī)運行時間之上，諸如例如在常規(guī)opencl運行時間542之上。在一些實施例中，同步工作竊取器系統(tǒng)可將工作負(fù)載諸如例如常規(guī)opencl工作負(fù)載內(nèi)核540包裝成附加工作竊取調(diào)度器代碼。在一些實施例中，opencl運行時間或其它運行時間(例如cuda運行時間)可能需要基本上被修改。

圖6是異構(gòu)計算或計算機系統(tǒng)600的示例實施例的框圖。系統(tǒng)包含cpu601-1和具有g(shù)pu601-2的卡652。cpu包含任何期望數(shù)量的核650。gpu包含任何期望數(shù)量的核651?？ㄟ€具有與gpu耦合的存儲器653。在一些實施例中，存儲器可包含dram，盡管這不是要求的。cpu和gpu耦合在一起，并且通過互連608與存儲器653耦合。任何已知互連是適合的，諸如例如外圍組件互連以及其衍生或擴展。

存儲器654包含共享存儲器602。共享存儲器包含cpu工作隊列604-1和gpu工作隊列604-2。cpu具有第一同步工作竊取器系統(tǒng)605-1。gpu具有第二同步工作竊取器系統(tǒng)605-2。在一些實施例中，這些同步工作竊取器系統(tǒng)中任一個或二者可類似于在本文其它地方所描述的，或者與之相同。例如，第二同步工作竊取器系統(tǒng)605-2可類似于圖1的第二同步工作竊取器系統(tǒng)105-2和/或圖2的同步工作竊取器系統(tǒng)205，或與之相同。

圖7是異構(gòu)計算或計算機系統(tǒng)700的示例實施例的框圖。系統(tǒng)包含通過互連與存儲器754耦合的芯片和/或管芯761。在一些實施例中，存儲器754可包含dram，盡管這不是要求的。任何已知互連是適合的，諸如例如外圍組件互連以及其衍生或擴展。芯片/管芯包含cpu701-1和集成圖形701-2。cpu包含任何期望數(shù)量的核750。集成圖形包含任何期望數(shù)量的核751。cpu和集成圖形與片上互連760耦合。在本領(lǐng)域已知的任何片上互連機制是適合的。集成存儲器控制器762也與片上互連耦合。存儲器控制器通過互連708將芯片/管芯與存儲器754耦合。

存儲器754包含共享存儲器702。共享存儲器包含cpu工作隊列704-1和集成圖形工作隊列704-2。cpu具有第一同步工作竊取器系統(tǒng)705-1。集成圖形具有第二同步工作竊取器系統(tǒng)705-2。在一些實施例中，這些同步工作竊取器系統(tǒng)中任一個或二者可類似于本文在別的地方所描述的，或者與之相同。例如，第二同步工作竊取器系統(tǒng)705-2可類似于圖1的第二同步工作竊取器系統(tǒng)105-2和/或圖2的同步工作竊取器系統(tǒng)205，或與之相同。

如下代碼說明了適合的同步工作竊取器系統(tǒng)的示例實施例的使用：

/*注釋：

1.全局工作竊取隊列數(shù)據(jù)結(jié)構(gòu)ws_q保存cpu核和gpu核的所有工作竊取隊列。這個數(shù)據(jù)結(jié)構(gòu)被分配在cpu與gpu之間的共享存儲器中。

2.工作竊取隊列ws_q[0],...,ws_q[num_cpu_threads-1]表示cpu核的隊列。ws_q[num_cpu_threads]表示gpu工作竊取隊列。num_cpu_threads表示異構(gòu)系統(tǒng)中的cpu核的總數(shù)。所有這些工作竊取隊列都被分配在cpu與gpu之間的共享存儲器中。每個工作竊取隊列在下面的"wsqueue"數(shù)據(jù)結(jié)構(gòu)中描述。

3.每個gpu計算內(nèi)核由子例程"actual_kernel"定義，子例程"actual_kernel"將它在其上操作的迭代空間的索引作為變元。這個子例程的細(xì)節(jié)是用戶應(yīng)用特定的。

4.下面描述的子例程"kernel_wrapper"是基于軟件的包裝器，其由編譯gpu內(nèi)核的編譯器或由應(yīng)用本身或由opencl運行時間生成。

5.下面描述的子例程"steal_wrapper"跨cpu和gpu核的工作竊取隊列發(fā)現(xiàn)工作。如果發(fā)現(xiàn)任何工作，則它在對應(yīng)工作竊取隊列上執(zhí)行竊取操作。如果在所有隊列中什么也沒發(fā)現(xiàn)，則返回empty。

6.子例程“steal”從給定工作竊取隊列執(zhí)行實際竊取。當(dāng)它竊取時，它首先確保在隊列中至少存在一個工作，并且然后原子地更新隊列的頂部指針以確保正確執(zhí)行同時竊取。

7.子例程“cas”在cpu與gpu之間的共享存儲器上執(zhí)行比較并設(shè)置操作。

*/

/*每個計算內(nèi)核，例如"actual_kernel"都用這個子例程包裝，以執(zhí)行cpu與gpu工作竊取隊列之間的工作竊??；這個包裝器代碼由用戶/編譯器或者由opencl運行時間執(zhí)行。*/

__kernelvoidkernel_wrapper(

globalwsqueue*ws_q,/*用于在共享存儲器中分配的cpu和gpu的工作竊取隊列數(shù)據(jù)結(jié)構(gòu)*/

intnum_cpu_threads/*cpu核的數(shù)量*/

){

intlocal_id=get_local_id(0);//工作組內(nèi)這個工作項的local_id

__localintwork_idx;//總迭代空間中的索引；工作組中的工作項共享這個

barrier(clk_local_mem_fence);

//本地存儲器屏障；這確保work_idx對工作組中的所有工作項可見；

if(work_idx==empty)return;//如果我們在本地或通過竊取未發(fā)現(xiàn)工作則返回

/*調(diào)用實際opencl內(nèi)核*/

actual_kernel(work_idx+local_id);

}

}

/*這個子例程順序地遍歷所有工作竊取隊列并嘗試尋找工作。如果在所有隊列中什么也沒發(fā)現(xiàn)，則返回empty，指示終止。*/

intsteal_wrapper(globalwsqueue*ws_q,int

num_cpu_threads){

/*ws_q[num_cpu_threads]表示gpu工作竊取隊列，因此我們以相反次序遍歷，以首先從gpu隊列并且然后cpu隊列發(fā)現(xiàn)工作。*/

for(inti=num_cpu_threads;i>=0;i--){

__globalwsqueueeach_ws_q=ws_q[i];

work_idx=steal(each_ws_q);//從工作竊取隊列執(zhí)行實際竊取操作

if(work_idx>=0)returnwork_idx;

//如果發(fā)現(xiàn)了工作，則將索引返回到工作。

}

returnempty;//如果在所有隊列中都沒發(fā)現(xiàn)工作，則返回empty。

}

/*此子例程執(zhí)行用于從給定隊列竊取工作的實際竊取操作*/

intsteal(globalwsqueue*a_ws_q){

inttop=a_ws_q->top;//指向工作竊取隊列中的頂條目

intbottom=a_ws_q->bottom;//指向工作竊取隊列中的底條目

__globalcirculararray*array=a_ws_q->activearray;

//發(fā)現(xiàn)存儲該對(min,max)的數(shù)組

intsize=bottom-top;//發(fā)現(xiàn)工作竊取隊列中的條目數(shù)量

if(size<=0){//如果沒有工作，返回empty

returnempty;

}

into=(array->segment[top%(l<<array->log_size)]).min;

//發(fā)現(xiàn)在頂上第一組塊的索引

/*既然我們發(fā)現(xiàn)了工作并且準(zhǔn)備好竊取，我們就在頂部指針上使用比較并設(shè)置(cas)執(zhí)行原子竊取。*/

if(!cas(a_ws_q,top,top+1)){//在頂部指針上比較并設(shè)置

returnabort；//如果我們未成功執(zhí)行原子操作則abort

}

returno;//從指示實際工作索引的元組(min,max)返回最小索引

}

/*此子例程示出了共享存儲器上cpu與gpu之間的比較并設(shè)置(cas)操作的實現(xiàn)。*/

boolcas(__globalwsqueue*a_ws_q,intoldval,intnewval){

intresult;

result=atomic_cmpxchg((volatileglobalint*)&(a_ws_q->top),oldval,newval);

//在共享存儲器上執(zhí)行原子比較并交換操作

return(result==oldval);

};

/*用于工作竊取隊列實現(xiàn)的數(shù)據(jù)結(jié)構(gòu)列表*/

/*每個單獨工作竊取隊列數(shù)據(jù)結(jié)構(gòu)*/typedefstructws_q_s{

volatileintbottom;//工作竊取隊列的底部

volatileinttop;//工作竊取隊列的頂部

circulararray*activearray;//用于保存所有工作組塊的基礎(chǔ)數(shù)組

}wsqueue;

typedefstructcirculararray_s{

intlog_size;//用于實現(xiàn)工作組塊的循環(huán)數(shù)組的log大小

pair_t*segment;//用于工作組塊的數(shù)組

}circulararray;

typedefstructpair_s{

intmin;//從那開始工作的最小索引，如果被竊取的話

intmax;//工作直到的最大索引，如果被竊取的話

}pair_t;

enumwsq_tag{

empty=-2,//指示隊列是empty，并且因此終止

abort=-1，//指示竊取操作由于競爭而未成功繼續(xù)

};

對于圖1和圖4-7中任一個所描述的組件、特征和細(xì)節(jié)可選地也可用在圖2-3中的任一個中。而且，本文對于任何設(shè)備描述的組件、特征和細(xì)節(jié)可選地也可用在本文描述的任何方法中，所述方法在實施例中可由此類設(shè)備和/或用此類設(shè)備執(zhí)行。

示例實施例

如下示例涉及另外實施例。在一個或多個實施例中的任何地方都可使用示例中的細(xì)節(jié)。

示例1是工作竊取器設(shè)備。工作竊取器包含：確定模塊，用于確定從第一類型的第一硬件計算單元竊取工作用于不同于第一類型的第二類型的第二硬件計算單元。該工作在第一工作隊列中排隊，第一工作隊列對應(yīng)于第一硬件計算單元并且被存儲在由第一和第二硬件計算單元共享的共享存儲器中。工作竊取器設(shè)備還包含：同步工作竊取器模塊，其通過對第一工作隊列的同步存儲器存取來竊取所述工作，所述同步存儲器存取相對于對來自第一硬件計算單元的第一工作隊列的存儲器存取同步。

示例2包含示例1的主題，并且可選地，其中同步工作竊取器模塊將工作添加到第二工作隊列。第二工作隊列對應(yīng)于第二硬件計算單元，并且被存儲在共享存儲器中。第二類型可選地比第一類型更專用。

示例3包含示例1的主題，并且可選地，其中所述同步工作竊取器模塊通過包含對第一工作隊列執(zhí)行的原子操作的所述同步存儲器存取來竊取所述工作。

示例4包含示例3的主題，并且可選地，其中所述原子操作包括讀-修改-寫操作、比較并互換操作、比較并交換操作、測試并設(shè)置操作、比較并設(shè)置操作以及加載鏈路/有條件存儲操作之一。

示例2包含示例1的主題，并且可選地，其中原子操作包括開放計算語言(opencl)原子操作。

示例6包含示例3的主題，并且可選地，其中同步工作竊取器模塊通過cuda同步存儲器存取來竊取工作。

示例7包含示例1的主題，并且可選地，進一步包括：第二確定模塊，用于確定從第二硬件計算單元竊取第二工作用于第一硬件計算單元。第二工作在第二工作隊列中排隊，第二工作隊列對應(yīng)于第二硬件計算單元并且被存儲在共享存儲器中。

示例8包含示例7的主題，并且可選地，進一步包括：第二同步工作竊取器模塊，其通過對來自第一硬件計算單元的第二工作隊列的第二同步存儲器存取來竊取第二工作。第二同步存儲器存取相對于對來自第二硬件計算單元的第二工作隊列的存儲器存取同步。

示例9包含示例1-8中任一個的主題，并且可選地，其中當(dāng)?shù)诙ぷ麝犃袨榭蘸吞畛湓陂撝档燃壱韵轮械囊豁棔r，所述確定模塊確定竊取所述工作，第二工作隊列對應(yīng)于第二硬件計算單元并且被存儲在所述共享存儲器中。

示例10包含示例1-8中任一個的主題，并且可選地，其中第一硬件計算單元包括從通用處理器和中央處理單元(cpu)中選擇的一個。第二硬件計算單元包括從圖形處理器、硬件加速器裝置、密碼處理器、通信處理器、網(wǎng)絡(luò)處理器、專用處理器、專用核、高度并行專用硬件、數(shù)字信號處理器(dsp)和現(xiàn)場可編程門陣列(fpga)中選擇的一個。

示例11包含示例10的主題，并且可選地，其中第一硬件計算單元包括所述cpu，并且其中第二硬件計算單元包括從圖形處理單元(gpu)和集成圖形核中選擇的所述圖形處理器。

示例12包含示例11的主題，并且可選地，其中圖形處理器包括集成圖形核，并且其中集成圖形核和cpu共享相同的最后一級高速緩存。

示例13是一種在異構(gòu)計算系統(tǒng)中的方法。該方法包含：確定從第一類型的第一硬件計算單元竊取工作用于比第一類型更專用的第二不同類型的第二硬件計算單元。該工作在第一工作隊列中排隊，第一工作隊列對應(yīng)于第一硬件計算單元并且被存儲在由第一和第二硬件計算單元共享的共享存儲器中。該方法還包含：竊取所述工作包含執(zhí)行對存儲在所述共享存儲器中的第一工作隊列的同步存儲器存取，所述同步存儲器存取相對于對來自第一硬件計算單元的第一工作隊列的存儲器存取同步。

示例14包含示例13的主題，并且可選地，進一步包括：將所述工作添加到第二工作隊列，其對應(yīng)于第二硬件計算單元，并且其也存儲在所述共享存儲器中。

示例15包含示例13的主題，并且可選地，其中執(zhí)行同步存儲器存取包括執(zhí)行原子操作。

示例16包含示例15的主題，并且可選地，其中執(zhí)行原子操作包括執(zhí)行從讀-修改-寫操作、比較并互換操作、比較并交換操作、測試并設(shè)置操作、比較并設(shè)置操作以及加載鏈路/有條件存儲操作中選擇的原子操作。

示例17包含示例15的主題，并且可選地，其中執(zhí)行原子操作包括執(zhí)行開放計算語言(opencl)原子操作。

示例18包含示例15的主題，并且可選地，其中竊取包括通過執(zhí)行cuda同步存儲器存取來竊取工作。

示例19包含示例13的主題，并且可選地，進一步包括：確定竊取第二工作，并從第二硬件計算單元竊取第二工作用于第一硬件計算單元。第二工作在第二工作隊列中排隊，第二工作隊列對應(yīng)于第二硬件計算單元并且被存儲在共享存儲器中。

示例20包含示例19的主題，并且可選地，其中竊取第二工作包括執(zhí)行對來自第一硬件計算單元的第二工作隊列的同步存儲器存取，所述同步存儲器存取相對于對來自第二硬件計算單元的第二工作隊列的存儲器存取同步。

示例21包含示例13的主題，并且可選地，其中確定包括響應(yīng)于第二工作隊列為空和填充在閾值等級以下中的一項而確定竊取所述工作，第二工作隊列對應(yīng)于第二硬件計算單元并且被存儲在所述共享存儲器中。

示例22包含示例13的主題，并且可選地，其中第一硬件計算單元是從通用處理器、中央處理單元(cpu)以及具有多個通用核的片上系統(tǒng)中選擇的一個。還有，其中第二硬件計算單元是從圖形處理器、硬件加速器裝置、密碼處理器、通信處理器、網(wǎng)絡(luò)處理器、專用處理器、專用核、片上系統(tǒng)上的專用核、高度并行專用硬件、數(shù)字信號處理器(dsp)和現(xiàn)場可編程門陣列(fpga)中選擇的一個。

示例23包含示例22的主題，并且可選地，其中第一硬件計算單元包括所述cpu，并且其中第二硬件計算單元包括從圖形處理單元(gpu)和集成圖形核中選擇的所述圖形處理器。

示例24包含示例22的主題，并且可選地，其中圖形處理器包括集成圖形核，并且其中集成圖形核和cpu的核共享相同的最后一級高速緩存。

示例25是異構(gòu)計算機系統(tǒng)。異構(gòu)計算機系統(tǒng)包含互連。異構(gòu)計算機系統(tǒng)還包含與互連耦合的第一類型的第一硬件計算單元。異構(gòu)計算機系統(tǒng)還包含與互連耦合的第二不同類型的第二硬件計算單元。第二類型比第一類型更專用。異構(gòu)計算機系統(tǒng)還包含與互連耦合的動態(tài)隨機存取存儲器(dram)，dram包含由第一和第二硬件計算單元共享的共享存儲器。共享存儲器包含將第一硬件計算單元的工作排隊的第一工作隊列和將第二硬件計算單元的工作排隊的第二工作隊列。異構(gòu)計算機系統(tǒng)還包含：工作竊取器設(shè)備，其確定竊取并通過對第一工作隊列的同步存儲器存取從第一隊列竊取工作并將它添加到第二隊列。同步存儲器存取相對于對來自第一硬件計算單元的第一工作隊列的存儲器存取同步。

示例26包含示例25的主題，并且可選地，其中工作竊取器設(shè)備通過包含第一工作隊列上的原子操作的同步存儲器存取來竊取工作。

示例27包含示例25-26中任一個的主題，并且可選地，進一步包括：第二工作竊取器設(shè)備，其確定竊取并通過對第二工作隊列的第二同步存儲器存取從第二工作隊列竊取第二工作并將它添加到第一工作隊列，第二同步存儲器存取相對于對來自第二硬件計算單元的第二工作隊列的存儲器存取同步。

示例28是存儲指令的一個或多個計算機可讀存儲介質(zhì)，所述指令如果由機器執(zhí)行則將使機器執(zhí)行操作。操作包含確定從第一類型的第一硬件計算單元竊取工作用于比第一類型更專用的第二不同類型的第二硬件計算單元的操作。該工作在第一工作隊列中排隊，第一工作隊列對應(yīng)于第一硬件計算單元并且被存儲在由第一和第二硬件計算單元共享的共享存儲器中。操作包含通過對第一工作隊列執(zhí)行的同步存儲器存取來竊取工作的操作。同步存儲器存取相對于對來自第一硬件計算單元的第一工作隊列的存儲器存取同步。

示例29包含示例28的主題，并且可選地，其中機器可讀存儲介質(zhì)進一步提供指令，所述指令如果由機器執(zhí)行則將使機器執(zhí)行包含將工作添加到第二工作隊列的操作，第二工作隊列對應(yīng)于第二硬件計算單元并且被存儲在共享存儲器中。

示例30包含示例28-29中任一個的主題，并且可選地，其中機器可讀存儲介質(zhì)進一步提供指令，所述指令如果由機器執(zhí)行則將使機器執(zhí)行包含通過包含執(zhí)行第一隊列的原子操作的同步存儲器存取來竊取工作的操作。

示例31是存儲指令的機器可讀存儲介質(zhì)，所述指令如果由機器執(zhí)行則使所述機器執(zhí)行示例13-24中任一個的方法。

示例32是包括用于執(zhí)行示例13-24中任一個的方法的部件的設(shè)備。

示例33是用于執(zhí)行示例13-24中任一個的方法的設(shè)備。

示例34是用于基本上執(zhí)行本文所描述的方法的設(shè)備。

示例35是包括用于基本上執(zhí)行本文所描述的方法的部件的設(shè)備。

在說明書和權(quán)利要求書中，可已經(jīng)使用了術(shù)語“耦合”和“連接”，連同它們的派生詞。應(yīng)該理解，這些術(shù)語不打算作為彼此的同義詞。而是，在具體實施例中，“連接”可用于指示兩個或更多元件彼此直接物理接觸或電接觸。“耦合”可意味著兩個或更多元件直接物理接觸或電接觸。然而，“耦合”也可意味著，兩個或更多元件彼此不直接接觸，但仍彼此協(xié)同操作或交互。例如，處理器可通過一個或多個中間組件(例如一個或多個互連和/或芯片集)與裝置耦合。在附圖中，使用箭頭示出連接和耦合。

在說明書和權(quán)利要求書中，可能已經(jīng)使用了術(shù)語“邏輯”。本文所使用的邏輯可包含諸如硬件、固件、軟件或它們的組合的模塊。邏輯示例包含集成電路、專用集成電路、模擬電路、數(shù)字電路、編程的邏輯裝置、包含指令的存儲器裝置等。

可已經(jīng)使用了術(shù)語“和/或”。本文所使用的術(shù)語“和/或”意味著一個或另一個或二者(例如a和/或b意味著a或b，或者a和b)。

在以上描述中，為了說明目的，已經(jīng)闡述了大量特定細(xì)節(jié)以便提供本發(fā)明實施例的全面理解。然而，對于本領(lǐng)域技術(shù)人員將顯而易見的是，在沒有這些特定細(xì)節(jié)中的一些細(xì)節(jié)的情況下，也可實踐一個或多個其它實施例。所描述的具體實施例不是提供用于限制本發(fā)明，而是用于通過示例實施例說明它。本發(fā)明的范圍不由特定示例確定而僅由權(quán)利要求書確定。在其它實例中，已經(jīng)以框圖形式或沒有細(xì)節(jié)示出了眾所周知的電路、結(jié)構(gòu)、裝置和操作，以便避免模糊描述的理解。

在視為合適的情況下，在附圖中已經(jīng)重復(fù)了附圖標(biāo)記或附圖標(biāo)記的末端部分以指示對應(yīng)或類似元件，它們可選地可具有類似或相同特性，除非另外規(guī)定或清楚地顯而易見。在一些情況下，已經(jīng)描述了多個組件，它們可被合并成單個組件。在已經(jīng)描述了單個組件的其它情況下，它可被分割成多個組件。

已經(jīng)描述了各種操作和方法。在流程圖中以比較基本的形式描述了其中一些方法，但可選地可向方法添加操作和/或從方法移除操作。此外，雖然流程圖示出了根據(jù)示例實施例的操作的具體次序，但具體次序是示范性的。備選實施例可選地可按不同次序執(zhí)行操作，組合某些操作，重疊某些操作等。

一些實施例包含制品(例如計算機程序產(chǎn)品)，其包含機器可讀介質(zhì)。介質(zhì)可包含提供例如存儲機器可讀形式的信息的機制。機器可讀介質(zhì)可提供或其上存儲有一個或多個指令，指令如果和/或當(dāng)由機器執(zhí)行時可操作以使機器執(zhí)行和/或?qū)е聶C器執(zhí)行本文公開的一個或多個操作、方法或技術(shù)。

在一些實施例中，機器可讀介質(zhì)可包含有形和/或非暫時性機器可讀存儲介質(zhì)。例如，有形和/或非暫時性機器可讀存儲介質(zhì)可包含軟盤、光存儲介質(zhì)、光盤、光數(shù)據(jù)存儲裝置、cd-rom、磁盤、磁光盤、只讀存儲器(rom)、可編程rom(prom)、可擦除和可編程rom(eprom)、電可擦除和可編程rom(eeprom)、隨機存取存儲器(ram)、靜態(tài)ram(sram)、動態(tài)ram(dram)、閃存、相變存儲器、相變數(shù)據(jù)存儲材料、非易失性存儲器、非易失性數(shù)據(jù)存儲裝置、非暫時性存儲器、非暫時性數(shù)據(jù)存儲裝置等。非暫時性機器可讀存儲介質(zhì)不包括暫時性傳播信號。在另一實施例中，機器可讀介質(zhì)可包含暫時性機器可讀通信介質(zhì)，例如電、光、聲或其它形式的傳播信號，諸如載波、紅外信號、數(shù)字信號等。

適當(dāng)機器的示例包含但不限于臺式計算機、膝上型計算機、筆記本計算機、平板計算機、筆記本、智能電話、蜂窩電話、服務(wù)器、網(wǎng)絡(luò)裝置(例如路由器和交換機)、移動因特網(wǎng)裝置(mid)、媒體播放器、智能電視、上網(wǎng)機、機頂盒和視頻游戲控制器以及其它計算機系統(tǒng)、計算裝置或具有一個或多個處理器的電子裝置。

還應(yīng)該認(rèn)識到，此說明書通篇提到“一個實施例”、“實施例”或“一個或多個實施例”例如意味著具體特征可包含在本發(fā)明的實踐中。類似地，應(yīng)該認(rèn)識到，在描述中各種特征有時一起組合在單個實施例、附圖或其描述中以便使本公開流暢并幫助理解各個發(fā)明方面。然而，公開的這個方法不被解釋為反映本發(fā)明需要比在每個權(quán)利要求中明確闡述的特征更多的特征的意圖。而是，如隨附權(quán)利要求書反映的，發(fā)明方面可在于少于單個公開的實施例的所有特征。從而，遵循具體實施方式的權(quán)利要求書由此被明確合并到此具體實施方式中，其中每個權(quán)利要求獨立作為本發(fā)明的獨立實施例。