一種基于cuda技術(shù)對柵格化數(shù)據(jù)進(jìn)行抽階的方法
【專利說明】一種基于CUDA技術(shù)對柵格化數(shù)據(jù)進(jìn)行抽階的方法
[0001]
技術(shù)領(lǐng)域
[0002]本發(fā)明涉及直寫式光刻機(jī)數(shù)據(jù)處理技術(shù)，具體來說是一種基于CUDA技術(shù)對柵格化數(shù)據(jù)進(jìn)行抽階的方法。
[0003]
【背景技術(shù)】
[0004]CUDA是NVIDIA公司2007年提出的支持GPU進(jìn)行通用計(jì)算的編程模型和開發(fā)環(huán)境，CUDA編程的思想是用海量的線程來開發(fā)程序中的并行性，海量線程以層次化的方式組織，單個的線程被映射到標(biāo)量核SP上執(zhí)行，一組線程被組織成一個線程塊(Block)被映射到一個流處理單位SM上執(zhí)行，最后由線程塊組成的線程柵格(Grid)映射到一個GPGPU(GPU)上執(zhí)行。由于GPU具有遠(yuǎn)超CPU的計(jì)算核心數(shù)以及海量的并行計(jì)算資源，適合進(jìn)行計(jì)算密集型、高度并行化的計(jì)算任務(wù)。同時，由于GPU的價格遠(yuǎn)遠(yuǎn)低于同等性能的并行計(jì)算系統(tǒng)，由CPU和GPGPU (GPU)組成的異構(gòu)系統(tǒng)已經(jīng)越來越廣的應(yīng)用到生物醫(yī)學(xué)、流體力學(xué)等諸多工程應(yīng)用領(lǐng)域。
[0005]直寫式光刻機(jī)的數(shù)據(jù)處理過程是將用戶提供的矢量數(shù)據(jù)，轉(zhuǎn)化為圖形發(fā)生器能接受的圖像數(shù)據(jù)，數(shù)據(jù)處理過程中涉及到數(shù)據(jù)的分析、計(jì)算和傳輸。目前實(shí)際應(yīng)用處理得到的柵格化數(shù)據(jù)中，一個像素用一個字節(jié)來表示(8階灰度)，而下位機(jī)只需要其中的1、2、4位即可滿足顯示的灰度要求，因此如果能夠針對柵格化數(shù)據(jù)進(jìn)行抽階處理，去掉其中冗余的數(shù)據(jù)，提取出有效的灰度值，即可降低了數(shù)據(jù)規(guī)模，降低傳輸鏈路帶寬。如針對同樣一幅圖，抽階后的數(shù)據(jù)量變少了，傳輸這些數(shù)據(jù)要求的時間不變，所以需要的帶寬(傳輸速率)降低了。實(shí)際應(yīng)用中選擇成本低，速度慢的傳輸鏈路也可滿足傳輸時間要求，則相當(dāng)于降低了生產(chǎn)成本。
[0006]對柵格化數(shù)據(jù)進(jìn)行抽階處理是根據(jù)實(shí)際需要來處理，如針對4位的灰度要求，在對數(shù)據(jù)柵格化時便可以將需要的4位排列在8位字節(jié)(一個像素)的前4位，在做抽階工作時，直接抽取0-3位即可；或針對2位的灰度要求，在數(shù)據(jù)柵格化時將需要的2位排在8位字節(jié)的第2位和第3位，在做抽階工作時，直接抽取1-2位即可。但是目前柵格化數(shù)據(jù)的數(shù)據(jù)量過于龐大，導(dǎo)致抽階工作較慢，抽階工作的分析、計(jì)算和傳輸均比較耗時，難以滿足產(chǎn)能要求，如何利用CUDA技術(shù)的特點(diǎn)，實(shí)現(xiàn)柵格化數(shù)據(jù)抽階的多線程并行處理已經(jīng)成為急需解決的技術(shù)問題。
[0007]

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

[0008]本發(fā)明的目的是為了解決現(xiàn)有技術(shù)中柵格化數(shù)據(jù)抽階效率較低的缺陷，提供一種基于CUDA技術(shù)對柵格化數(shù)據(jù)進(jìn)行抽階的方法來解決上述問題。
[0009]為了實(shí)現(xiàn)上述目的，本發(fā)明的技術(shù)方案如下:
一種基于CUDA技術(shù)對柵格化數(shù)據(jù)進(jìn)行抽階的方法，包括以下步驟:
CPU分配顯存和計(jì)算資源，CPU根據(jù)進(jìn)行抽階數(shù)據(jù)的規(guī)模，結(jié)合當(dāng)前GPU可供使用的硬件資源，計(jì)算出最優(yōu)的GPU線程分配方式；
根據(jù)GPU線程分配方式申請顯存空間，將輸入數(shù)據(jù)由內(nèi)存拷貝到顯存中；
GPU在每一個線程Thread上進(jìn)行核函數(shù)計(jì)算，對每個字節(jié)進(jìn)行抽階操作；
待所有線程Thread的核函數(shù)計(jì)算完后，將顯存中的結(jié)構(gòu)數(shù)據(jù)拷貝回內(nèi)存，此結(jié)構(gòu)數(shù)據(jù)為抽階后的柵格化數(shù)據(jù)，完成抽階過程。
[0010]所述的CPU分配顯存和計(jì)算資源包括以下步驟:
輸入柵格化處理后的二維位圖像素陣列，其寬度定義為width，高度定義為height ； 將每個二維線程塊Block的寬度定義為blockDim.X、高度定義為blockDim.y ；
計(jì)算出線程柵格Grid的寬度gridDim.X，其計(jì)算公式如下: gridDim.x = width/blockDim.x ；
計(jì)算出線程柵格Grid的高度gridDim.y，其計(jì)算公式如下: gridDim.y = height/b1ckDim.y ；
計(jì)算出顯存分配總大小length，其計(jì)算公式如下:
Iength=Width 氺height 氺(1+N/8)，
其中，N=l、2或4 ;
獲得線程分配方式，
線程分配方式中二維線程塊Block為Block (blockDim.x, blockDim.y)；出線程柵格Grid 為 Grid(gridDim.x, gridDim.y)。
[0011]所述的GPU在每一個線程Thread上進(jìn)行核函數(shù)計(jì)算包括以下步驟:
計(jì)算當(dāng)前線程偏移量，根據(jù)當(dāng)前線程的柵格坐標(biāo)計(jì)算全局線程編號，通過全局線程編號計(jì)算出當(dāng)前線程在緩沖區(qū)中的偏移量；
根據(jù)線程的偏移量從顯存中取出數(shù)據(jù)存入線程塊內(nèi)共享顯存；
針對滿足tid.x=0的所有線程來執(zhí)行其所在二維線程塊Block輸入數(shù)據(jù)的抽階操作；根據(jù)柵格化處理規(guī)則從每一個字節(jié)取出特定的位，將抽階的結(jié)果數(shù)據(jù)暫存到結(jié)果緩存mask中，其中設(shè)緩存mask為4個字節(jié)；
設(shè)int型為4字節(jié)，使用tid.x<4的線程Thread將結(jié)果數(shù)據(jù)從結(jié)果緩存mask中拷貝至對應(yīng)的顯存中。
[0012]所述的計(jì)算當(dāng)前線程偏移量包括以下步驟:
計(jì)算當(dāng)前線程所在塊的編號Md，其計(jì)算公式如下: bid=gridDim.x^blockldx.y+blockldx.x ；
其中，blockldx.y為當(dāng)前線程所在線程塊Block中的列號，blockldx.x為當(dāng)前線程所在線程塊Block中的行號；
計(jì)算當(dāng)前塊內(nèi)線程編號cur_tid，其計(jì)算公式如下: cur_tid=b1ckDim.x*threadldx.y+threadldx.x ；
其中，blockldx.y為當(dāng)前線程所在線程塊Block中的列號，blockldx.x為當(dāng)前線程所在線程塊Block中的行號； 計(jì)算全局線程編號total_tid，其計(jì)算公式如下: total_tid=bid*b1ckDim.x*blockDim.y+cur_tid ；
根據(jù)全局線程編號確定當(dāng)前線程輸入輸出數(shù)據(jù)在緩沖區(qū)中的偏移量offset，其計(jì)算公式如下:
offset=total_tid*(blockDim.x*blockDim.y)*(N/8)，
其中8是一個字節(jié)的位數(shù)，N為抽取的階數(shù)，N=l、2或4。
[0013]
有益效果
本發(fā)明的一種基于CUDA技術(shù)對柵格化數(shù)據(jù)進(jìn)行抽階的方法，與現(xiàn)有技術(shù)相比通過并行化提高了計(jì)算效率，增加了直寫式光刻機(jī)的產(chǎn)能，同時降低了數(shù)據(jù)規(guī)模，減少了對計(jì)算能力及傳輸帶寬的依賴，降低了成本。
[0014]本發(fā)明在直寫式光刻機(jī)的數(shù)據(jù)處理過程中，發(fā)掘出了其中的計(jì)算密集型、高度并行化過程，并將這個過程并行化，通過CUDA將其部署到GPU上并行執(zhí)行，極大的提高了處理速度。并且，在實(shí)現(xiàn)并行化的過程中，充分地利用了 GPU以及CUDA框架的特性，實(shí)現(xiàn)了最大化的加速比；在分配線程資源時，根據(jù)當(dāng)前硬件的最大線程數(shù)、線程塊中的最優(yōu)線程數(shù)來確定Block和Grid的尺寸；核函數(shù)在處理時先將輸入數(shù)據(jù)拷貝到共享顯存，充分利用了共享顯存的高帶寬特性，提高了處理速度；讀寫全局顯存時根據(jù)線程編號同步操作，有效的屏蔽了訪存延時，進(jìn)一步提高了處理效率。
[0015]
【附圖說明】
圖1為本發(fā)明的方法順序圖；
圖2為本發(fā)明中CUDA線程柵格示意圖。
[0016]
【具體實(shí)施方式】
[0017]為使對本發(fā)明的結(jié)構(gòu)特征及所達(dá)成的功效有更進(jìn)一步的了解與認(rèn)識，用以較佳的實(shí)施例及附圖配合詳細(xì)的說明，說明如下:
如圖1所示，本發(fā)明所述的一種基于CUDA技術(shù)對柵格化數(shù)據(jù)進(jìn)行抽階的方法，抽階的過程在GPU中并行進(jìn)行，利用CUDA框架實(shí)現(xiàn)。其包括以下步驟:
第一步，CPU分配顯存和計(jì)算資源，CPU根據(jù)進(jìn)行抽階數(shù)據(jù)的規(guī)模，結(jié)合當(dāng)前GPU可供使用的硬件資源，計(jì)算出最優(yōu)的GPU線程分配方式，為后面的步驟分配顯存和計(jì)算資源。其具體包括以下步驟:
(I)輸入柵格化處理后的二維位圖像素陣列，二維位圖像素陣列即為柵格化數(shù)據(jù)，其寬度定義為width，高度定義為height。
[0018](2)根據(jù)CUDA的技術(shù)要求，每個二維線程塊Block的寬度為blockDim.X、高度為blockDim.yD
[0019](3)計(jì)算出線程柵格Grid的寬度gridDim.X，其計(jì)算公式如下: gridDim.x = width/blockDim.χ0
[0020]計(jì)算出線程柵格Grid的高度gridDim.y，其計(jì)算公式如下: gridDim.y = height/b1ckDim.y。
[0021]計(jì)算出顯存分配總大小length，其計(jì)算公式如下:
Iength=Width 氺height 氺(1+N/8)，
其中，N=l、2或4。
[0022]如圖2所示，在此步驟后獲得最優(yōu)的GPU線程分配方式，在GPU (device端)共有g(shù)ridDim.X * gridDim.y個二維線程塊Block并行執(zhí)行，每個二維線程塊Block中有BlockDim.x * BlockDim.y個線程Thread并行執(zhí)行，通過這種高度并行化極大地提高了執(zhí)行效率。
[0023](4)獲得線程分配方式，
線程分配方式中二維線程塊Block為Block (blockDim.x, blockDim.y)；出線程柵格Grid 為 Grid(gridDim.x, gridDi