驅(qū)動執(zhí)行單元的機器碼產(chǎn)生方法以及裝置的制造方法
【技術(shù)領(lǐng)域】
[0001]本發(fā)明涉及一種圖形處理單元技術(shù)，特別是一種驅(qū)動執(zhí)行單元的機器碼產(chǎn)生方法以及裝置。
【背景技術(shù)】
[0002]在通用計算圖形處理單元(GPGPU,General-Purpose Computing on GraphicsProcessing Unit)模型，例如，開放計算語言(OpenCL, Open Computing Language)、清染腳本(rendering scripts)、DirectX (Direct extens1n)及開放圖形語言(OpenGL, OpenGraphics Language)等計算著色器(CS, Compute Shader)，大部分的輸出輸入都由存儲器的加載(load)、儲存(store)、縮減(reduct1n)及采樣(sample)操作數(shù)完成。如何提升加載、儲存、縮減及取樣作業(yè)的效能，是個重要議題。因此，本發(fā)明提出一種驅(qū)動執(zhí)行單元的機器碼產(chǎn)生方法以及裝置，不僅減少存儲器地址的運算成本，更可以提升暫存數(shù)據(jù)的命中率。

【發(fā)明內(nèi)容】

[0003]本發(fā)明的實施例提出一種驅(qū)動執(zhí)行單元的機器碼產(chǎn)生方法，使用編譯器執(zhí)行，包含下列步驟。搜集核心程序中關(guān)聯(lián)于存儲器表面的多個數(shù)據(jù)存取指令；分析數(shù)據(jù)存取指令關(guān)聯(lián)的地址樣式，用以產(chǎn)生全域識別碼地址；以及產(chǎn)生包含全域識別碼地址的機器碼。
【附圖說明】
[0004]圖1是依據(jù)本發(fā)明實施例的驅(qū)動執(zhí)行單元的機器碼產(chǎn)生裝置的硬件架構(gòu)圖。
[0005]圖2是依據(jù)本發(fā)明實施例的驅(qū)動執(zhí)行單元的機器碼產(chǎn)生方法流程圖。
[0006]圖3A至3C是依據(jù)本發(fā)明實施例的驅(qū)動執(zhí)行單元的機器碼產(chǎn)生方法流程圖。
【具體實施方式】
[0007]以下說明為完成發(fā)明的較佳實現(xiàn)方式，其目的在于描述本發(fā)明的基本精神，但并不用以限定本發(fā)明。實際的
【發(fā)明內(nèi)容】
必須參考之后的權(quán)利要求范圍。
[0008]必須了解的是，使用于本說明書中的“包含”、“包括”等詞，用以表示存在特定的技術(shù)特征、數(shù)值、方法步驟、作業(yè)處理、組件以及/或組件，但并不排除可加上更多的技術(shù)特征、數(shù)值、方法步驟、作業(yè)處理、組件、組件，或以上的任意組合。
[0009]在權(quán)利要求中使用如“第一”、“第二”、“第三”等詞是用來修飾權(quán)利要求中的組件，并非用來表示之間具有優(yōu)先權(quán)順序，先行關(guān)系，或者是一個組件先于另一個組件，或者是執(zhí)行方法步驟時的時間先后順序，僅用來區(qū)別具有相同名字的組件。
[0010]在通用計算圖形處理單元(GPGPU)模型，執(zhí)行于電子裝置的主程序(mainfunct1n)稱為核心程序(Kernel)。以下為范例核心程序的偽碼，用以將二個存儲器平面(surfaces)相加并輸出:
[0011]_kernel void VectorAdd(_global float4*A, _global float4*B, _globalfloat4*C, int WidthA, int WidthB, int WidthC){
[0012]int x = get_global_id (0)；
[0013]int y = get_global_id (1)；
[0014]int addressA = (y+2) *ffidthA+x+5 ；
[0015]int addressB = (y+3)*ffidthB+x+6 ；
[0016]int addressC = y*ffidthC+x ；
[0017]C[addressC] = A[addressA]+B[addressB] ；}
[0018]核心程序的加載、儲存及取樣作業(yè)的地址以全域識別碼(global-1ds)為基礎(chǔ)，為二維/三維(2D/3D)空間地址。在一些實施方式中，編譯器直接依據(jù)二維/三維地址為加載、儲存及取樣作業(yè)計算一維地址，并產(chǎn)生機器碼。范例的機器碼如下所示:
[0019]IMUL R0, C[4], R0
[0020]IADD R0，R4，R0
[0021]IADD R0，C[7]，R0 //x = get_global_id (0)
[0022]IMUL Rl, C[5], R1
[0023]IADD R1，R5，R1
[0024]IADD Rl, C[8], Rl //y = get_global_id(1)
[0025]IADD R2, Rl, 2
[0026]IMUL R2, CB[16], R2
[0027]IADD R2，R2，R0 //(y+2) *ffidthA+x ；
[0028]IADD R2, R2, 5 //(y+2)*ffidthA+x+5 ；
[0029]IADD R2, C[13], R2?4
[0030]LD R4.xyzw, R2 // 從 A [addressA]加載數(shù)據(jù)
[0031]IADD R3, Rl, 3
[0032]IMUL R3, C[17], R3
[0033]IADD R3，R3，R0 //(y+3) *ffidth+x ；
[0034]IADD R3，R3，6 //(y+3) *ffidth+x+6 ；
[0035]IADD R3, C[14], R3?4
[0036]LD R8.xyzw, R3 // 從 B [addressA]加載數(shù)據(jù)
[0037]FADD.rp3 R4, R4, R8 //A[addressA]+B[addressB]
[0038]IMUL Rl, C[18], Rl
[0039]IADD Rl, Rl, R0 //y*ffidth+x ；
[0040]IADD Rl, C[15], Rl?4
[0041 ] ST R4.xyzw, Rl //儲存數(shù)據(jù)至 C [addressC]
[0042]IMUL代表整數(shù)乘法操作數(shù)，IADD代表整數(shù)加法操作數(shù)，F(xiàn)ADD代表浮點數(shù)加法操作數(shù)，LD代表載入操作數(shù)，ST代表儲存操作數(shù)。這些計算出來的存儲器地址會直接送到執(zhí)行單元，用以存取數(shù)據(jù)。然而，此編譯后的機器碼需要較多的地址運算作業(yè)。
[0043]圖1是依據(jù)本發(fā)明實施例的驅(qū)動執(zhí)行單元的機器碼產(chǎn)生裝置的硬件架構(gòu)圖。編譯器(compiler) 130搜集每一個存儲器表面(memory surface)的屬性(attributes)，存儲器表面指一段存儲器地址的數(shù)據(jù)。詳細(xì)來說，編譯器130掃描核心程序110的內(nèi)容，用以獲得關(guān)聯(lián)于同一段存儲器地址(亦即是同一個存儲器表面)的數(shù)據(jù)存取指令。編譯器130分析搜集到的數(shù)據(jù)存取指令的變量的數(shù)據(jù)形態(tài)，用以判斷此存儲器表面是否只讀以及是否為專屬形態(tài)緩存區(qū)或影像(type buffer or image)。編譯器130另判斷數(shù)據(jù)存取指令的變量的數(shù)據(jù)形態(tài)是否為可支持?jǐn)?shù)據(jù)形態(tài)，若不是，則將原始指令改變?yōu)榭芍С謹(jǐn)?shù)據(jù)形態(tài)的指令。編譯器130更分析搜集到的數(shù)據(jù)存取指令的地址計算過程，用以產(chǎn)生以全域識別碼地址(global-1d addresses)。最后，編譯器130根據(jù)此存儲器表面是否只讀以及是否為專屬形態(tài)緩存區(qū)或影像的信息、原始的以及/或改變后的指令、以及原始的或新產(chǎn)生的存儲器地址產(chǎn)生機器碼(又稱為二進制代碼)150。驅(qū)動器(driver) 170依據(jù)加載的機器碼150設(shè)定適當(dāng)?shù)募拇嫫?register)，用以驅(qū)動執(zhí)行單元(EU, Execut1n Unit) 200完成特定的運算。執(zhí)行單元200可為頂點著色器(VS, Vertex Shader)、外殼著色器(HS，Hull Shader)、區(qū)域著色器(DS, Domain Shader)、幾何著色器(GS, Geometry Shader)、畫素著色器(PS, PixelShader)，或在三維圖形管道中的其它硬件。驅(qū)動器170可以專用硬件電路或通用硬件(例如，單一處理器、具并行處理能力的多處理器或其它具運算能力的處理器)實施。
[0044]圖2是依據(jù)本發(fā)明實施例的驅(qū)動執(zhí)行單元的機器碼產(chǎn)生方法流程圖。此方法由編譯器130完成。編譯器130反復(fù)執(zhí)行一個循環(huán)，用以產(chǎn)生關(guān)聯(lián)于每一個存儲器表面的加載、儲存、縮減及采樣指令的機器碼150，并且指令中的存儲器地址為全域識別碼地址。在每一回合中，編譯器130搜集核心程序110中關(guān)聯(lián)于一個存儲器表面的多個數(shù)據(jù)存取指令(步驟S210)。存儲器表面可指一段存儲器地址中儲存的數(shù)據(jù)。接著，分析數(shù)據(jù)存取指令關(guān)聯(lián)的地址樣式，用以產(chǎn)生全域識別碼地址(步驟S230)，以及產(chǎn)生包含全域識別碼地址的機器碼150(步驟S250)。此后，流程判斷是否分析完核心程序(步驟S270) ?如果不是，則返回步驟S210循環(huán)操作；如果是，則結(jié)束該方法。
[0045]圖3A至3C是依據(jù)本發(fā)明實施例的驅(qū)動執(zhí)行單元的機器碼產(chǎn)生方法流程圖。此方法由編譯器130完成。編譯器130反復(fù)執(zhí)行一個循環(huán)，用以產(chǎn)生關(guān)聯(lián)于每一個存儲器表面的加載、儲存、縮減及采樣指令的機器碼150，并且指令中的存儲器地址為全域識別碼地址。在每一回合中，編譯器130可包含三個階段的分析。第一階段分析關(guān)聯(lián)于一個存儲器表面的所有數(shù)據(jù)存取指令的類型，第二階段分析數(shù)據(jù)存取指令中的變量的數(shù)據(jù)形態(tài)，以及第三階段分析數(shù)據(jù)存取指令關(guān)聯(lián)的地址樣式(address pattern)。編譯器130根據(jù)分析結(jié)果在需要的情況下改變原始指令為新指令，以及產(chǎn)生全域識別碼地址。最后，編譯器130根據(jù)原始的以及/或更新后的指令以及全域識別碼地址產(chǎn)生機器碼150。
[0046]詳細(xì)來說，首先，編譯器130搜集核心程序110中關(guān)聯(lián)于一個存儲器表面的多個數(shù)據(jù)存取指令(步驟S311)。存儲器表面可指一段存儲器地址中儲存的數(shù)據(jù)。在第一階段的分析中，編譯器130判斷數(shù)據(jù)存取指令是否關(guān)聯(lián)于采樣指令(步驟S312)。若是，判斷此存儲器表面為專屬形態(tài)影像(type image)(步驟S314)。否則，編譯器130更判斷此存儲器表面是否為只讀(步驟S313)。在步驟S313，詳細(xì)來說，當(dāng)關(guān)聯(lián)于此存儲器表面的指令只有加載指令時，代表此存儲器表面為只讀。在第二階段的分析中，編譯器130判斷數(shù)據(jù)存取指令中的所有變量的數(shù)據(jù)形態(tài)是否相同(步驟S315以及S317)。通過第一及第二階段的分析結(jié)果，編譯器130可判斷此存儲器表面的類型。詳細(xì)來說，當(dāng)存儲器表面為只讀及數(shù)據(jù)存取指令的所有變量的數(shù)據(jù)形態(tài)都相同時(步驟S313及步驟S315中“是”的路徑)，判斷此存儲器表面為專屬形態(tài)緩存區(qū)(type buffer)(步驟S321)。當(dāng)存儲器表面為只讀及數(shù)據(jù)存取指令的變量的數(shù)據(jù)形態(tài)不同時(步驟S313中“是”的路徑接步驟S315中“否”的路徑)，判斷此存儲器表面為原始緩存區(qū)(raw buffer)(步驟S391)。當(dāng)存儲器表面非只讀及數(shù)據(jù)存取指令中的變量的數(shù)據(jù)形態(tài)不同時(步驟S313及步驟S315中“否”的路徑)，判斷此存儲器表面為非專屬形態(tài)隨機存取視圖(un-Type UAV, Un-order Accessed View)(步驟S341)。當(dāng)存儲器表面非只讀及數(shù)據(jù)存取指令的所有變量的數(shù)據(jù)形態(tài)相同時(步驟S313中“否”的路徑接步驟S315中“是”的路徑)，判斷此存儲器表面為專屬形態(tài)隨機存取視圖(Type UAV)(步驟S331)。此外，當(dāng)判斷此存儲器表面為專屬形態(tài)緩存區(qū)或?qū)傩螒B(tài)隨機存取視圖后(步驟S331或S321)，判斷數(shù)據(jù)存取指令的變量的數(shù)據(jù)形態(tài)是否為可支持?jǐn)?shù)據(jù)形態(tài)(步驟S351)。若否，編譯器130改變原始的數(shù)據(jù)存取指令成為包含可支持?jǐn)?shù)據(jù)形態(tài)的新指令(步驟S353)。例如，一個數(shù)據(jù)存取指令包含數(shù)據(jù)形態(tài)為CHAR16的變量，用以一次加載四