采用GPU的提升紋理緩存命中光線投射方法

杜松江， 張思超

(1. 長江大學工程技術(shù)學院 信息工程學院， 湖北 荊州 434020；2. 中國礦業(yè)大學 機電工程學院， 江蘇 徐州 221116)

?

采用GPU的提升紋理緩存命中光線投射方法

杜松江1， 張思超2

(1. 長江大學工程技術(shù)學院 信息工程學院， 湖北 荊州 434020；2. 中國礦業(yè)大學 機電工程學院， 江蘇 徐州 221116)

提出一種改善紋理緩存命中率的方法.首先，分析圖形處理器(GPU)中三維紋理組織的布局特性；進而提出根據(jù)視點的變化動態(tài)選擇線程配置的策略，目的在于最小化warp級的投射光線紋理訪存跨距；最后，算法用CUDA(compute unified device architecture)實現(xiàn)并驗證.實驗結(jié)果表明：當視點分別圍繞x，y,z坐標軸旋轉(zhuǎn)時，改進后算法的幀速率分別為改進前的1.08，1.14,0.98倍.

三維紋理； 光線投射； 圖形處理單元； 紋理緩存

光線投射算法作為體繪制技術(shù)中的一種，在醫(yī)學、天文、地學等領(lǐng)域有著廣泛的應用.由于算法需要對屏幕上的每個像素做運算，因此,光線投射算法的計算量很大，很難滿足實時交互方面的應用.另一方面，投射光線之間是相互獨立的，該算法適合并行化實現(xiàn).Kruger 等[1]通過圖形應用程序編程接口(application programming interface,API)的功能調(diào)用，將體數(shù)據(jù)作為三維紋理保存在圖形處理器(graphics processing unit,GPU)中，利用圖形流水線的可編程著色器進行算法的實現(xiàn).之后，大量的研究工作都是針對圖形流水線模式下的算法進行改進和優(yōu)化[2-4].2007年以來，Nvidia公司推出支持計算統(tǒng)一設(shè)備架構(gòu) CUDA (compute unified device architecture)的GPU，該技術(shù)使GPU從圖形領(lǐng)域的應用進一步擴展到了更多的領(lǐng)域.CUDA編程模型可以使更多的通用算法在GPU上得到了實現(xiàn)，并取得可觀的加速效果[5].在支持CUDA的GPU上進行光線投射算法加速的研究工作中，Marsalekl等[6]首先實現(xiàn)了算法的移植，加速效果優(yōu)于基于Shader 的實現(xiàn).在不降低繪制速度的前提下，Zhang等[7]利用3次B樣條改善CUDA的光線投射算法的視覺效果.然而，通過觀察發(fā)現(xiàn)，當體數(shù)據(jù)規(guī)模較大時，繪制幀數(shù)率的變化受視點變換的影響嚴.本文從GPU的硬件體系結(jié)構(gòu)和訪存模型出發(fā)，最小化warp級投射光線訪問相鄰體數(shù)據(jù)時的跨距，從而提高繪制性能.

1　CUDA編程模型

CUDA編程模型將CPU作為主機，GPU作為協(xié)處理器.CPU負責進行邏輯性強的事務處理和串行計算，GPU則專注于執(zhí)行高度線程化的并行處理任務[8-10].CUDA計算流程通常包括CPU到GPU數(shù)據(jù)傳遞、Kernel函數(shù)執(zhí)行、GPU到CPU數(shù)據(jù)傳遞3個步驟.

CUDA采用單指令多線程(single instruction multiple thread,SIMT)執(zhí)行模式，即GPU上的所有線程并行執(zhí)行內(nèi)核函數(shù)Kernel[11].另外，CUDA將線程組織成塊網(wǎng)格、線程塊、線程3個不同的層次[12-13]，并采用多層次的存儲器結(jié)構(gòu).存儲器包括只對單個線程可見的寄存器和本地存儲器、對塊內(nèi)線程可見的共享存儲器、對所有線程可見的全局存儲器等.其中，全局內(nèi)存可以被綁定為紋理內(nèi)存，主要用在圖形圖像等應用中.

2　紋理訪存分析及線程配置策略

不同于全局內(nèi)存的2級緩存，紋理內(nèi)存提供的緩存主要是緩存空間上相鄰的數(shù)據(jù)[14].空間相鄰的投射光線對體數(shù)據(jù)采樣時，同樣會訪問空間上相鄰的體素.因此，使用紋理內(nèi)存保存體數(shù)據(jù)是合適的選擇.然而在交互過程中，視點并非靜止的.在不同位置訪問體數(shù)據(jù)所表現(xiàn)出的緩存效果也是不同的.

2.1GPU紋理內(nèi)存的訪存分析

將體數(shù)據(jù)保存為三維紋理時，數(shù)據(jù)在紋理內(nèi)存中的布局，如圖1所示.

令體數(shù)據(jù)的長寬高都為N，且N=2l.圖1中：三維紋理可以看成是二維紋理切片沿z方向的集合.在全局內(nèi)存中，數(shù)據(jù)以一維線性的方式保存；而在紋理內(nèi)存中，每個二維紋理切片在紋理存儲器中以Morton編碼的方式組織(箭頭所指的)[15-16].Morton編碼具有遞歸的特性，因此，第l層的紋元編碼由第l-1層的編碼決定，層次之間的結(jié)構(gòu)關(guān)系，如圖2所示.

圖1　三維紋理布局 圖2　莫頓編碼的層次結(jié)構(gòu)　　　Fig.1　3D texture layout Fig.2　Hierarchical structure of Morton code

該編碼方式雖然優(yōu)化了二維訪問的空間局部性，但是造成相鄰數(shù)據(jù)的間隔距離不均勻，如圖1中的紋元1和4，紋元2和3.將相鄰紋元之間的跨度大小分為以下兩種情況討論.

1) 相鄰紋元有不同的z坐標.這種情況下，相鄰紋元出現(xiàn)在兩個相鄰的二維切片上.因為有相同的x和y坐標，相鄰紋元之間的跨距都是N2.

(1)

沿著y坐標軸的最大跨距為

(2)

因為N=2l，故沿著x，y坐標軸的紋元最大訪問跨距分別為(N2+2)/6和(N2+2)/3.由以上分析可知，沿著x，y，或z坐標軸訪問相鄰紋元時，訪問跨距近似為1∶2∶6.由于紋理緩存的大小僅為幾十KB，當體數(shù)據(jù)切片過大時，對紋理內(nèi)存的大跨距訪問會造成頻繁的緩存命中失效，使得算法性能的下降.

圖3　不同視點訪問同一三維紋理切片的示意圖Fig.3　Sketch of accessing same slice at different viewpoints

不同視點訪問同一三維紋理切片的示意圖，如圖3所示.由圖3可知：c視點位置的訪存效率最高，因為相鄰的投射光線訪問的是沿x軸平行的相鄰體素；反之，b視點訪存性能最差，因為相鄰的投射光線訪問的是不同切片上的體數(shù)據(jù).為提高相鄰投射光線在訪問紋理內(nèi)存時的紋理緩存命中率，訪問紋理內(nèi)存數(shù)據(jù)時應盡量沿著x方向進行.

2.2線程塊及warp的幾何形狀

CUDA將所有并行線程等分為多個線程塊.一個線程塊中的線程由1個或多個warp組成，而一個warp是連續(xù)的32個線程.線程塊之間和線程塊內(nèi)部可以組織成一維或者二維的形狀.當將線程塊設(shè)置為二維形狀時，也間接決定了warp的形狀.線程塊和warp的形狀,如表1所示.

表1　線程塊和warp的形狀

由表1可知：線程塊的幾何形狀表現(xiàn)為從垂直到水平的過程；對應的warp形狀也做相應的變化.

(a)情況1　　　(b)情況2　　　(c) 情況3

(d)情況4　　　(e)情況5　　　(f)情況6　　圖4　體數(shù)據(jù)和成像平面之間的6種典型情況Fig.4　6 kinds of typical situation between volume axes and screen

2.3基于視點的線程塊形狀動態(tài)分配

假定視點位于坐標原點O，體坐標軸中的兩個軸代表的平面和成像屏幕平行時的6種典型情況，如圖4所示.

灰色平面為當體坐標軸中的兩個軸和成像屏幕平行時，體數(shù)據(jù)的二維切片和成像屏幕之間的6種狀態(tài).前文描述中，沿x軸訪問相鄰體數(shù)據(jù)時跨距最小，因此,令體數(shù)據(jù)的x軸為主軸.基于視點的動態(tài)線程形狀配置主要通過以下3個步驟確定.

步驟1平行平面檢測.在視點變換的過程中，從xy-,xz-及yz-平面中選擇和屏幕最為平行的平面,如圖4(a)和圖4(e)中的xy-平面.

步驟2確定主軸.從平行平面中選擇訪存跨距最小的軸為主軸.例如,圖4(a),圖4(e)中的xy-平面中，由于沿x軸訪問比沿y軸的訪問跨距小,因此，選擇x軸為主軸.

步驟3確定線程塊的形狀.根據(jù)主軸被繪制在屏幕上的方向選擇線程塊的形狀.如果主軸在屏幕上用垂直線繪制，選擇垂直warp的線程塊，如圖4(d),圖4(e)，線程塊形狀為1×256；如果主軸在屏幕上用水平線繪制，選擇水平warp的線程塊，如圖4(a),圖4(b)，線程塊形狀為256×1.對于其他處于中間過度狀態(tài)的情況，選擇垂直和水平混合的warp形狀.為此，將0°～90°的旋轉(zhuǎn)區(qū)域再次細分成6組過渡區(qū)域，即每15°為一個過渡區(qū)域.

在交互過程中，視點的任意旋轉(zhuǎn)變換可以看做是分別繞3個坐標軸的旋轉(zhuǎn)變換組合而成.為更進一步說明過渡區(qū)域的線程塊選擇策略分別討論繞x軸旋轉(zhuǎn)、繞y軸旋轉(zhuǎn)、繞z軸旋轉(zhuǎn)的選擇策略.旋轉(zhuǎn)角度分別為Θx，Θy及Θz，細分后的過渡區(qū)域的選擇，如表2所示.

表2　線程塊幾何形狀的動態(tài)選擇

由表2可知：繞x軸旋轉(zhuǎn)時，主軸x-始終平行于屏幕.warp中的相鄰投射光線沿著x軸訪問體素時跨距最小，緩存命中率也就越大.因此，選擇水平狀warp的線程塊形狀，即256×1；繞y軸旋轉(zhuǎn)時，平行平面由xy平面逐步過渡到y(tǒng)z平面，主軸也由x軸變?yōu)閥軸.為了盡量減少訪存跨距增大引起的命中下降，warp的形狀也由水平狀逐漸過渡到垂直狀；同樣地，繞z軸旋轉(zhuǎn)時，雖然平行平面始終為xy平面，但主軸由水平狀變?yōu)榇怪睜?，warp的形狀也跟著相應的變化，與繞y軸旋轉(zhuǎn)不同的是，在這個過程中主軸沒有發(fā)生改變.

旋轉(zhuǎn)角度為90°～360°時，warp形狀及線程塊的形狀配置利用幾何的對稱關(guān)系得到.

3　算法框架

假設(shè)成像屏幕的高和寬分為W和H，整個屏幕成像所需要的線程數(shù)量為W×H.一般情況下，GPU中一個線程block中的線程數(shù)量遠遠低于繪制整個屏幕需要的線程數(shù)量.通過將屏幕分塊，采用屏幕塊對應線程塊的做法可以解決這一問題.線程塊采用二維布局，維度大小為w×h.同樣地，線程grid也采用二維結(jié)構(gòu)，總共需要的線程塊個數(shù)為ceil(W/w) ×ceil(H/h)個.線程grid中線程坐標和屏幕上每個像素坐標的對應關(guān)系為

u=blockIdx.x×BLOCK_SIZE+threadIdx.x,

v=blockIdx.y×BLOCK_SIZE+threadIdx.y.

除了將體數(shù)據(jù)作為三維紋理保存外，充分利用GPU中各種存儲器的特性，即傳遞函數(shù)主要用于將投射光線擊中的體素值轉(zhuǎn)換為顏色值和不透明度，將其綁定為類型為float4的一維紋理，只需保存少量顏色值，其余的值可以利用硬件支持的插值算法生成.體數(shù)據(jù)顯示到屏幕上是一個三維對象變換為二維圖像的過程，并且每個投射光線所代表的線程都會用到該變換，利用該存儲器的廣播功能，將變換矩陣保存在常量內(nèi)存中.投射光線在對體數(shù)據(jù)進行采樣時，使用寄存器變量保存臨時累加值.最后，將每個投射光線的計算結(jié)果寫入全局內(nèi)存.

4　實驗和分析

為驗證基于warp級紋理訪存優(yōu)化的線程塊動態(tài)配置方法，實驗部分主要通過繪制幀速率的提高說明方法的有效性.算法所用的計算平臺為Nvidia開普勒GK110架構(gòu)的GeforceGT740M型GPU；CUDASDK為5.5.GPU的硬件規(guī)格的計算能力為3.5；CUDA核心數(shù)量為384；處理器頻率為1.03GHz；SM數(shù)量為2；全局內(nèi)存為2GB；共享內(nèi)存為48KB.

所用數(shù)據(jù)為1 024×1 024×1 024的HydrogenAtoms體數(shù)據(jù)，每個體素大小為8bit.成像屏幕的大小為1 024×1 024.

將視點分別繞x軸、y軸、z軸旋轉(zhuǎn)360°，分別測量不采用動態(tài)線程配置時的幀速率和采用動態(tài)線程配置后的幀速率.每個線程塊的大小設(shè)定為256，不采用動態(tài)配置的線程塊形狀為16×16，采用動態(tài)線程塊配置的形狀根據(jù)前述方法進行變化.

體數(shù)據(jù)分別繞x,y,z軸旋轉(zhuǎn)360°的幀速率結(jié)果進行對比，如圖5所示.圖5中：ω為旋轉(zhuǎn)角度；v為幀速率.

(a) x軸　　　　　　　　　　　　(b) y軸　　　　　　　　　　　　　　(c) z軸圖5　靜態(tài)形狀和動態(tài)形狀的幀速率對比圖Fig.5　Comparison chart of frame rate between static shape and dynamic shape

由圖5(a)可知：雖然在旋轉(zhuǎn)過程中平行平面一直在變化，但是由于主軸x軸一直是水平無變化的，因此，整個旋轉(zhuǎn)范圍內(nèi)的繪制幀速率平均高于繞其他兩個軸的旋轉(zhuǎn).在采用動態(tài)配置優(yōu)化后，更能適應warp中相鄰投射光線訪問紋理內(nèi)存的特點，繪制性能有了進一步的提升.

由圖5(b)可知：幀速率表現(xiàn)出了很大的差異.在旋轉(zhuǎn)范圍為0～90°時，因為主軸由x軸逐漸變化成為y軸，該變化過程導致warp級投射光線訪問紋理緩存時的命中率降低，繪制性能隨旋轉(zhuǎn)角度的增加而降低.當旋轉(zhuǎn)角度由90°～180°改變時，主軸又逐漸變回為x軸，繪制幀速率也得到回升.由于體數(shù)據(jù)本身具有對稱性，當旋轉(zhuǎn)角度為180°～360°時，繪制幀速率同樣也表現(xiàn)出了對稱性.采用動態(tài)配置方法也起到了改善繪制性能的作用.

由圖5(c)可知：當繞z軸旋轉(zhuǎn)時，雖然平行平面沒有發(fā)生改變，但是主軸x-的方向卻一直在變化，算法的繪制性造成一些影響，總體情況比繞y軸時要好.值得注意的是，當繞z-軸旋轉(zhuǎn)時，在理論上動態(tài)線程配置是能夠適應主軸的旋轉(zhuǎn)改變，并提升性能，但實際效果相反，繪制速率表現(xiàn)出了震蕩效應.造成該現(xiàn)象的原因是線程塊的幾何形狀在旋轉(zhuǎn)角度為0°～15°時，選擇的是256×1的水平狀線程塊，warp形狀高寬比為32∶1.該高寬比下，線程塊的形狀有4個備選方案.

線程塊形狀改為32×8，0°～90°旋轉(zhuǎn)區(qū)間的運行效果，如圖6所示.由圖6可知：修改線程塊形狀后，0°～15°這一區(qū)間的幀速率變化得到了改善.

最后，體數(shù)據(jù)分別繞x，y，z坐標軸旋轉(zhuǎn)360°時，提出方法的總體改進效果，如表3所示.

圖6　改進前和改進后的幀速率對比圖Fig.6　Comparison chart of framerate Before and after improvement

坐標軸v(靜態(tài))v(動態(tài))加速比x35.1838.011.08y14.0715.971.14z26.0125.560.98

5　結(jié)束語

根據(jù)視點動態(tài)選擇線程塊，改善基于GPU的光線投射算法在訪問紋理內(nèi)存時的性能.分析體數(shù)據(jù)保存為三維紋理后的數(shù)據(jù)布局及線程塊形狀和warp形狀的關(guān)系.在給定視點下，在旋轉(zhuǎn)變換中確定與成像屏幕平行的平行平面，進而確定主軸的方法指導線程塊的幾何形狀選擇，可以改善warp級投射光線在訪問紋理緩存時的命中率失效的問題.實驗結(jié)果表明：該方法能夠改善光線投射算法的性能.下一步的工作將繼續(xù)優(yōu)化體數(shù)據(jù)繞z-坐標軸旋轉(zhuǎn)時的紋理緩存，繼續(xù)考慮關(guān)于線程塊block級的優(yōu)化.

[1]KRUGER J,WESTERMANN R.Acceleration techniques for GPU-based volume rendering[C]∥Proceedings of the 14th IEEE Visualization.Washington D C:IEEE Computer Society,2003:287-292.

[2]SALAMA C R,KELLER M,KOHLMANN P.High-level user interfaces for transfer function design with semantics[J].Visualization and Computer Graphics,2006,12(5):1021-1028.

[3]GOBBETTI E,MARTON F,GUITIN J A I.A single-pass GPU ray casting framework for interactive out-of-core rendering of massive volumetric datasets[J].Visual Computer,2008，24(7/8/9):797-806.

[4]李國和,段忠祥,吳衛(wèi)江,等.針對全空子數(shù)據(jù)體的 GPU 體繪制[J].中國圖象圖形學報,2014.19(4):577-582.

[5]OWENS J D,HOUSTON M,LUEBKE D,et al.GPU computing[J].Proceedings of the IEEE，2008，96(5):879-899.

[7]ZHANG Changgong，XI Ping,ZHANG Chaoxin.CUDA-based volume ray-casting using cubic B-spline[C]∥International Conference on Virtual Reality and Visualization.Beijing：IEEE Press,2011:84-88.

[8]甘新標,沈立,王志英.基于CUDA的并行全搜索運動估計算法[J].計算機輔助設(shè)計與圖形學學報,2010,22(3):457-460.

[9]趙麗麗,張盛兵,張萌,等.基于CUDA的高速FFT計算[J].計算機應用研究,2011,28(4):155-159.

[10]肖江,胡柯良,鄧元勇.基于CUDA的矩陣乘法和FFT性能測試[J].計算機工程,2009,35(10):7-10.

[11]王蓓蕾,朱志良,孟琭.基于CUDA加速的SIFT特征提取[J].東北大學學報(自然科學版),2013,34(2):200-204.

[12]JENKINS J,ARKAKAR I,OWENS J D,et al.Lessons learned from exploring the backtracking paradigm on the GPU[J].Lecture Notes in Computer Science,2011,6853(2):425-437.

[13]周洪,樊曉椏,趙麗麗.基于CUDA的稀疏矩陣與矢量乘法的優(yōu)化[J].計算機測量與控制,2010,18(8):1906-1908.

[14]SANDERS J,KANDROT E.CUDA by example:an introduction to general-purpose GPU programming[M].Boston:Addison-Wesley Professional,2010:116-117.

[15]MONTRYM J,MORETON H.The geforce 6800[J].IEEE Micro,2005(2):41-51.

[16]MORTON G M.A computer oriented geodetic data base and a new technique in file sequencing[M].New York：International Business Machines Company,1966:56-60.

(責任編輯： 陳志賢英文審校： 吳逢鐵)

Improving Texture Cache-Hit Rate of GPU-Based Ray Casting

DU Songjiang1， ZHANG Sichao2

(1. College of Information Engineering， Yangtze University College of Engineering Technology, Jingzhou 434020, China；2. School of Mechanical and Electrical Engineering， China University of Mining and Technology， Xuzhou 221116， China)

This paper presents a method of improving the texture cache hitrate for GPU-based volume rendering. Firstly, we analyze the data layout of 3D texture in GPU. Based on it, a dynamic strategy of selecting the thread block shape according to the viewpoint is proposed. The strategy can minimize the access stride for the warp-level threads. Finally, we realize the method in CUDA (compute unified device architecture) and testify the effectiveness. The experimental results show that when the viewpoint rotates around thex-,y-,z- axis, the frame rates are 1.08, 1.14 and 0.98 time faster than that of static thread block shape configuration, respectively.

3D texture； ray casting； graphics processing unit； texture cache

10.11830/ISSN.1000-5013.201605020

2016-03-15

張思超(1973-)，男，教授，博士，主要從事數(shù)據(jù)庫應用、軟件工程的研究.E-mail:dusongjiang2014@163.com.

國家自然科學基金資助項目(51204186)

TP 391

A

1000-5013(2016)05-0627-06