基于cuda流和共享內(nèi)存的dvh圖并行統(tǒng)計(jì)方法
【專利摘要】本發(fā)明公開了一種基于CUDA流和共享內(nèi)存的DVH圖并行統(tǒng)計(jì)方法，包括以下步驟：在主機(jī)端對(duì)器官進(jìn)行采樣，并將采樣點(diǎn)位置傳入設(shè)備端，每個(gè)器官的劑量統(tǒng)計(jì)分別用一個(gè)流進(jìn)行處理；步驟2：使用紋理存儲(chǔ)器載入劑量矩陣：步驟3：根據(jù)每個(gè)線程分配到的位置點(diǎn)，使用紋理拾取進(jìn)行拾取，紋理的濾波模式設(shè)置為線性插值，即對(duì)三維紋理的八個(gè)像元根據(jù)距離進(jìn)行線性插值，并返回線性插值得到的值；步驟4：使用共享內(nèi)存存儲(chǔ)統(tǒng)計(jì)結(jié)果，通過在共享內(nèi)存上開辟N個(gè)子劑量盒，解決共享內(nèi)存會(huì)出現(xiàn)的bank沖突問題，并加快了統(tǒng)計(jì)速度。
【專利說明】基于CUDA流和共享內(nèi)存的DVH圖并行統(tǒng)計(jì)方法
【技術(shù)領(lǐng)域】
[0001]本發(fā)明涉及醫(yī)學(xué)數(shù)據(jù)圖像處理領(lǐng)域，具體地，涉及一種基于CUDA流和共享內(nèi)存的DVH圖并行統(tǒng)計(jì)方法。
【背景技術(shù)】
[0002]劑量-體積直方圖(Dose Volume Histogram, DVH)是評(píng)價(jià)放射治療計(jì)劃方案優(yōu)劣的有力工具。在調(diào)強(qiáng)放療等逆向治療計(jì)劃系統(tǒng)中，對(duì)DVH圖統(tǒng)計(jì)的速度提出更高的要求。
[0003]DVH圖是在放療計(jì)劃中用二維圖表來(lái)統(tǒng)計(jì)三維劑量分布的方法，即表示某感興趣區(qū)域如腫瘤靶區(qū)或評(píng)估關(guān)鍵器官(organ at risk,OAR)的體積內(nèi)有多少體積受到多高劑量的照射。由于其直觀、有效地反應(yīng)了治療計(jì)劃的劑量分布與計(jì)劃的優(yōu)劣，成為評(píng)估放射治療計(jì)劃優(yōu)劣的主要依據(jù)。在DVH圖的統(tǒng)計(jì)中會(huì)根據(jù)靶區(qū)或關(guān)鍵器官的輪廓線和三維劑量數(shù)據(jù)的位置關(guān)系，對(duì)器官空間點(diǎn)對(duì)應(yīng)的劑量值進(jìn)行統(tǒng)計(jì)。通常會(huì)在器官范圍內(nèi)采樣，由劑量數(shù)據(jù)進(jìn)行三線性插值以求得采樣點(diǎn)的劑量值，并將其記錄在不連續(xù)的劑量盒(bin)中。隨著治療過程中病人CT切片的增加及調(diào)強(qiáng)放療的需要，統(tǒng)計(jì)過程中需要進(jìn)行大量的三線性插值運(yùn)算，CPU的計(jì)算速度不能滿足實(shí)時(shí)性要求，本發(fā)明考慮利用并行處理實(shí)現(xiàn)DVH圖的快速統(tǒng)計(jì)。
[0004]在并行處理領(lǐng)域，主要有多核中央處理單兀(Multicore Central ProcessingUnit, Multicore CPU)并行處理與圖像處理器單兀(Graphics Processing Unit, GPU)并行處理。由于CPU負(fù)責(zé)的是邏輯性較強(qiáng)的事務(wù)處理，并不擅長(zhǎng)于做高密集度計(jì)算，所以選擇專為計(jì)算密集型、高度并行化應(yīng)用而設(shè)計(jì)的高性能計(jì)算平臺(tái)NVIDIA GPU作為并行處理DVH圖的平臺(tái)。NVIDIA GPU的運(yùn)算能力和存儲(chǔ)器帶寬上相對(duì)于CPU有明顯的優(yōu)勢(shì)。通過計(jì)算統(tǒng)一設(shè)備構(gòu)架(Compute Unified Device Architecture, CUDA), GPU 可以在單指令多數(shù)據(jù)(Single Instruction Multiple Threads, SIMT)編程模型下發(fā)揮其強(qiáng)大的計(jì)算能力。
[0005]基于GPU并行實(shí)現(xiàn)DVH圖統(tǒng)計(jì)時(shí)有兩個(gè)難點(diǎn)。第一，對(duì)于每個(gè)器官對(duì)應(yīng)的輪廓線來(lái)說，判斷采樣點(diǎn)位置是否在其范圍內(nèi)時(shí)會(huì)用到大量的if-else語(yǔ)句，這在CPU中是很簡(jiǎn)單的任務(wù)，而在GPU中使用判斷語(yǔ)句則很可能會(huì)導(dǎo)致并行線程的串行執(zhí)行，極大地降低流多處理器內(nèi)部的執(zhí)行性能。第二，統(tǒng)計(jì)時(shí)會(huì)將采樣點(diǎn)結(jié)果分別保存到一百個(gè)bin中，尤其對(duì)于類似于重離子放療等具有特殊規(guī)律的劑量分布，會(huì)導(dǎo)致大量寫沖突，造成計(jì)算瓶頸。

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

[0006]本發(fā)明的目的在于，針對(duì)上述問題，提出一種基于CUDA流和共享內(nèi)存的DVH圖并行統(tǒng)計(jì)方法，以提高計(jì)算速度，并具有避免內(nèi)存寫沖突的優(yōu)點(diǎn)。
[0007]為實(shí)現(xiàn)上述目的，本發(fā)明采用的技術(shù)方案是:
一種基于CUDA流和共享內(nèi)存的DVH圖并行統(tǒng)計(jì)方法，包括以下步驟:
步驟1:在主機(jī)端對(duì)器官進(jìn)行采樣，并將采樣點(diǎn)位置傳入設(shè)備端，
DVH圖統(tǒng)計(jì)時(shí)需要判斷采樣點(diǎn)是否在器官的輪廓線內(nèi)，CPU判斷得到每個(gè)器官的所有采樣點(diǎn)位置并存儲(chǔ)到數(shù)組中，采樣點(diǎn)位置用一個(gè)三維向量表示即POS = (x, y, z);
利用CUDA的流機(jī)制對(duì)得到的位置數(shù)組傳入GPU中進(jìn)行計(jì)算，不同的流在執(zhí)行時(shí)是相互獨(dú)立的，即每個(gè)流處理器處理每個(gè)器官的位置數(shù)組；
步驟2:使用紋理存儲(chǔ)器載入劑量矩陣:
將已有劑量矩陣數(shù)據(jù)從主機(jī)端拷貝到設(shè)備端，并將劑量矩陣數(shù)據(jù)存儲(chǔ)到紋理存儲(chǔ)器
中；
步驟3:根據(jù)每個(gè)線程分配到的位置點(diǎn)，使用紋理拾取進(jìn)行拾取，紋理的濾波模式設(shè)置為線性插值，即對(duì)三維紋理的八個(gè)像元根據(jù)距離進(jìn)行線性插值，并返回線性插值得到的值；
步驟4:使用共享內(nèi)存存儲(chǔ)統(tǒng)計(jì)結(jié)果，通過在共享內(nèi)存上開辟N個(gè)子劑量盒，解決共享內(nèi)存會(huì)出現(xiàn)的bank沖突問題,并加快了統(tǒng)計(jì)速度。
[0008]根據(jù)本發(fā)明的優(yōu)選實(shí)施例，上述步驟4中使用共享內(nèi)存存儲(chǔ)統(tǒng)計(jì)結(jié)果主要為:在共享內(nèi)存上開辟N個(gè)子劑量盒,使子劑量盒之間對(duì)應(yīng)的bin不在同一 bank,同一 warp中的相鄰的若干個(gè)線程分別更新對(duì)應(yīng)的子劑量盒。
[0009]根據(jù)本發(fā)明的優(yōu)選實(shí)施例，上述N的取值范圍為]_<N S32。
[0010]根據(jù)本發(fā)明的優(yōu)選實(shí)施例，所述N=S。
[0011]本發(fā)明的技術(shù)方案具有以下有益效果:
本發(fā)明的技術(shù)方案，采用CUDA的異構(gòu)(heterogeneous)編程模型和CUDA流機(jī)制分別將各個(gè)器官的采樣點(diǎn)傳入GPU中進(jìn)行計(jì)算。實(shí)現(xiàn)了兩個(gè)層次的并行化，即器官之間的并行和插值運(yùn)算之間的并行處理。從而提高計(jì)算速度。采用共享內(nèi)存對(duì)數(shù)據(jù)進(jìn)行存儲(chǔ)，避免了內(nèi)存之間的寫沖突。
[0012]下面通過附圖和實(shí)施例，對(duì)本發(fā)明的技術(shù)方案做進(jìn)一步的詳細(xì)描述。
【專利附圖】

【附圖說明】
[0013]圖1為本發(fā)明實(shí)施例所述的基于CUDA流和共享內(nèi)存的DVH圖并行統(tǒng)計(jì)方法的流程圖；
圖2a和圖2b為器官與劑量矩陣的關(guān)系圖；
圖3為單線程DVH圖統(tǒng)計(jì)流程圖；
圖4為CUDA架構(gòu)運(yùn)行示意圖；
圖5為DVH圖統(tǒng)計(jì)時(shí)多線程投票更新示意圖；
圖6為共享存儲(chǔ)器的bank沖突示意圖；
圖7a為1%到100%劑量分布示意圖；
圖7b為90%到100%劑量分布示意圖；
圖8a為CPU生成的DVH示意圖；
圖8b為GPU生成的DVH示意圖。
【具體實(shí)施方式】
[0014]以下結(jié)合附圖對(duì)本發(fā)明的優(yōu)選實(shí)施例進(jìn)行說明，應(yīng)當(dāng)理解，此處所描述的優(yōu)選實(shí)施例僅用于說明和解釋本發(fā)明，并不用于限定本發(fā)明。
[0015]如圖1所示，一種基于CUDA流和共享內(nèi)存的DVH圖并行統(tǒng)計(jì)方法，包括以下步驟: 步驟1:在主機(jī)端對(duì)器官進(jìn)行采樣，并將采樣點(diǎn)位置傳入設(shè)備端，
DVH圖統(tǒng)計(jì)時(shí)需要判斷采樣點(diǎn)是否在器官的輪廓線內(nèi)，CPU判斷得到每個(gè)器官的所有采樣點(diǎn)位置并存儲(chǔ)到數(shù)組中，采樣點(diǎn)位置用一個(gè)三維向量表示即pos = (x, y, z)；
利用CUDA的流機(jī)制對(duì)得到的位置數(shù)組傳入GPU中進(jìn)行計(jì)算，不同的流在執(zhí)行時(shí)是相互獨(dú)立的，即每個(gè)流處理器處理每個(gè)器官的位置數(shù)組；
步驟2:使用紋理存儲(chǔ)器載入劑量矩陣:
將已有劑量矩陣數(shù)據(jù)從主機(jī)端拷貝到設(shè)備端，并將劑量矩陣數(shù)據(jù)存儲(chǔ)到紋理存儲(chǔ)器
中；
步驟3:根據(jù)每個(gè)線程分配到的位置點(diǎn)，使用紋理拾取進(jìn)行拾取，紋理的濾波模式設(shè)置為線性插值，即對(duì)三維紋理的八個(gè)像元根據(jù)距離進(jìn)行線性插值，并返回線性插值得到的值；
步驟4:使用共享內(nèi)存存儲(chǔ)統(tǒng)計(jì)結(jié)果，通過在共享內(nèi)存上開辟N個(gè)子劑量盒，解決共享內(nèi)存會(huì)出現(xiàn)的bank沖突問題,并加快了統(tǒng)計(jì)速度
其中，步驟4中使用共享內(nèi)存存儲(chǔ)統(tǒng)計(jì)結(jié)果主要為:在共享內(nèi)存上開辟N個(gè)子劑量盒，使子劑量盒之間對(duì)應(yīng)的bin不在同一bank,同一 warp中的相鄰的若干個(gè)線程分別更新對(duì)應(yīng)
的子劑量盒。N的取值范圍為I € N < 32。因?yàn)橹筮€需要將子劑量盒進(jìn)行相加，且
共享內(nèi)存是有限的，所以并不是N越大越好，通過實(shí)驗(yàn)N=S時(shí)效果較好，N>8時(shí)投票時(shí)間反而變長(zhǎng)。
[0016]具體處理步驟如下:
1.1對(duì)于異構(gòu)編程模型
CUDA并不是單純的GPU語(yǔ)言，它協(xié)調(diào)了 CPU和GPU兩種處理單元的并發(fā)計(jì)算，即CUDA架構(gòu)面向的是一個(gè)由CPU和GPU組成的異構(gòu)計(jì)算網(wǎng)絡(luò)。CPU具有分支預(yù)測(cè)、程序堆棧和循環(huán)優(yōu)化等針對(duì)控制采取的復(fù)雜邏輯，而GPU比較簡(jiǎn)單的結(jié)構(gòu)使它較適合順序、單一、少循環(huán)、少跳轉(zhuǎn)的語(yǔ)句。CPU端運(yùn)行串行的程序并控制GPU的啟動(dòng)(新的kepler GKllO架構(gòu)支持DYNAMIC PARALLELISM機(jī)制，能通過應(yīng)用不返回主機(jī)CPU的數(shù)據(jù)而動(dòng)態(tài)創(chuàng)建新線程，即任何內(nèi)核可以啟動(dòng)另一個(gè)內(nèi)核)，GPU端運(yùn)行可并行的任務(wù)，即一般來(lái)說，GPU需要CPU的控制和協(xié)調(diào)，所以GPU通常被稱為設(shè)備(device)，CPU稱為主機(jī)(host)。
[0017]1.2 線程
在CUDA中，NVIDIA用一種分層的編程模型來(lái)組織線程。用戶根據(jù)需要?jiǎng)?chuàng)建一定數(shù)量的線程并確定線程和數(shù)據(jù)的映射關(guān)系。在計(jì)算中線程被分為若干個(gè)線程塊(block)，與一個(gè)線程塊對(duì)應(yīng)的硬件實(shí)現(xiàn)為一個(gè)流多處理器(streaming multiprocessor, SM),又稱多處理器。Iv SM包含若干個(gè)標(biāo)量處理器(scalar processor),或稱CUDA核(CUDA core)。標(biāo)量處理器對(duì)應(yīng)一個(gè)具體執(zhí)行的線程。所有的block組成一個(gè)grid,即一塊顯卡對(duì)應(yīng)一個(gè)grid。
[0018]1.3存儲(chǔ)器
CUDA存儲(chǔ)器被組織為一種層次結(jié)構(gòu)，分別為可被程序中所有線程訪問的全局存儲(chǔ)器、常量存儲(chǔ)器及紋理存儲(chǔ)器；可被block內(nèi)所有線程訪問的共享存儲(chǔ)器；可被單個(gè)線程訪問的寄存器。
[0019](I)寄存器:用于存儲(chǔ)自動(dòng)變量，即在全局函數(shù)和設(shè)備函數(shù)中不使用限定符說明的變量。寄存器獨(dú)立地分布于每個(gè)標(biāo)量處理器上，為每個(gè)線程提供私有的存儲(chǔ)空間。容量較小，訪問速度快。
[0020](2)共享存儲(chǔ)器:位于每一個(gè)多處理器內(nèi)，屬于片上存儲(chǔ)器(on-chip memory),訪問速度與寄存器相仿，讀寫4B的數(shù)據(jù)大約需要兩個(gè)時(shí)鐘周期(clock cycle)。同一個(gè)多處理器上的線程可以訪問一片共享存儲(chǔ)器。
[0021](3)全局存儲(chǔ)器:屬于片外存儲(chǔ)器(off-chip memory),—次直接訪問需要400到800個(gè)時(shí)鐘周期,在計(jì)算能力3.X的多處理器上有片上緩存(on-chip cache),較計(jì)算能力
1.X速度快很多。
[0022](4)紋理存儲(chǔ)器:屬于片外存儲(chǔ)器。紋理存儲(chǔ)器是來(lái)自圖形學(xué)的一個(gè)概念，在之前的圖形圖像顯卡上就是作為核心的存儲(chǔ)器來(lái)進(jìn)行圖像顯示，因此其有較高的硬件支持，具有片上緩存，表現(xiàn)為尋址的計(jì)算由特定的硬件單元實(shí)現(xiàn)，不需要耗費(fèi)內(nèi)核的計(jì)算時(shí)間；訪問不必遵循順序的原則；對(duì)鄰近的數(shù)據(jù)表現(xiàn)出較大的數(shù)據(jù)帶寬；提供硬件加速的插值功能，但是精度只有9位。
[0023]2 DVH 圖統(tǒng)計(jì)
2.1現(xiàn)有的DVH圖統(tǒng)計(jì)算法
統(tǒng)計(jì)DVH圖時(shí)輸入為劑量數(shù)據(jù)與器官的輪廓數(shù)據(jù)，它們的關(guān)系如圖2a和圖2b所示。
[0024]圖2a中矩形代表三維劑量矩陣，實(shí)體部分代表器官直觀的范圍；圖2b中曲線代表了器官的輪廓線，設(shè)定每條輪廓線有一定的厚度，統(tǒng)計(jì)時(shí)每個(gè)采樣點(diǎn)代表了一定的體積。當(dāng)使用單線程來(lái)做DVH圖統(tǒng)計(jì)時(shí)，每個(gè)器官按照次序串行進(jìn)行統(tǒng)計(jì)，一般的流程是:如圖3所示。單線程統(tǒng)計(jì)DVH圖時(shí)，在得到劑量數(shù)據(jù)和輪廓數(shù)據(jù)后，需要對(duì)所有相關(guān)器官和腫瘤靶區(qū)進(jìn)行采樣。根據(jù)器官的排序順序統(tǒng)計(jì)，每個(gè)器官內(nèi)由其輪廓線與劑量數(shù)據(jù)的關(guān)系進(jìn)行采樣和插值運(yùn)算，最后將每個(gè)器官統(tǒng)計(jì)結(jié)果輸出為一個(gè)大小為100的統(tǒng)計(jì)數(shù)組。
[0025]3.基于CUDA的DVH圖并行計(jì)算，
DVH圖并行實(shí)現(xiàn)時(shí)首先需要將器官內(nèi)所有體積的采樣點(diǎn)坐標(biāo)傳入GPU全局存儲(chǔ)器中，將劑量矩陣存入紋理內(nèi)存，之后根據(jù)采樣點(diǎn)坐標(biāo)進(jìn)行插值運(yùn)算，將統(tǒng)計(jì)結(jié)果存入共享存儲(chǔ)器中，最后將統(tǒng)計(jì)結(jié)果進(jìn)行累加傳回主機(jī)，具體為:
步驟1:在主機(jī)端對(duì)器官進(jìn)行采樣，并將采樣點(diǎn)位置傳入設(shè)備端。
[0026]DVH圖統(tǒng)計(jì)時(shí)需要判斷采樣點(diǎn)是否在器官的輪廓線內(nèi)，這個(gè)任務(wù)需要執(zhí)行大量的判斷語(yǔ)句，適合在CPU端進(jìn)行處理。CPU判斷得到每個(gè)器官的所有采樣點(diǎn)位置并存儲(chǔ)到數(shù)組中，位置用一個(gè)三維向量表示POS = (X，y, z)。
[0027]將得到的位置數(shù)組用cudaMemcpyO函數(shù)直接拷貝進(jìn)設(shè)備端時(shí)比較低效，因?yàn)樾枰獙⑺袛?shù)據(jù)都拷貝進(jìn)設(shè)備端后，內(nèi)核才會(huì)處理這些數(shù)據(jù)。CUDA的流機(jī)制在此時(shí)可以發(fā)揮重要作用，會(huì)將GPU計(jì)算與CPU/GPU內(nèi)存?zhèn)鬏斚嗷ジ采w。
[0028]如圖4所示，一般來(lái)說可以將CUDA程序分為三個(gè)階段，分別為:主機(jī)一設(shè)備數(shù)據(jù)傳輸(h部分)，內(nèi)核執(zhí)行(i部分)和設(shè)備一主機(jī)數(shù)據(jù)傳輸(j部分)。圖中假設(shè)每個(gè)流的內(nèi)核執(zhí)行時(shí)間以及數(shù)據(jù)傳輸消耗時(shí)間都是等分的。不同的流的不同的部分可能在時(shí)間軸上相互重疊，如圖中黃色部分與綠色部分的重疊區(qū)域和綠色部分與藍(lán)色部分的重疊區(qū)域，都代表了不同流之間的并行部分。
[0029]由于不同的流在執(zhí)行時(shí)是相互獨(dú)立的，所以本發(fā)明在DVH統(tǒng)計(jì)中，將每個(gè)器官的統(tǒng)計(jì)分別用一個(gè)流來(lái)進(jìn)行處理，CUDA上的核心代碼摘取如下:
//根據(jù)器官和靶區(qū)數(shù)量建立流 cudaStream_t stream[0RGAN_C0UNT]; for(int i=0; i<0RGAN_C0UNT; i++) cudaStreamCreate(&stream[i]);
Il使用cudaMemcpyAsync ()函數(shù)實(shí)現(xiàn)異步數(shù)據(jù)傳輸 for(int i=0; i<0RGAN_C0UNT; i++) cudaMemcpyAsync();
步驟2:使用紋理存儲(chǔ)器載入劑量矩陣。
[0030]需要將已有劑量矩陣數(shù)據(jù)從主機(jī)端拷貝到設(shè)備端，此時(shí)有兩種選擇:可以將其存入全局存儲(chǔ)器或紋理存儲(chǔ)器。因?yàn)樵谥髸?huì)根據(jù)采樣點(diǎn)位置頻繁地對(duì)三維劑量矩陣進(jìn)行三線性插值，且訪問時(shí)通常都為劑量矩陣中鄰近的八個(gè)點(diǎn)。由上述存儲(chǔ)器的介紹可知比較可知，紋理存儲(chǔ)器較全局存儲(chǔ)器有很大優(yōu)勢(shì)，所以選擇紋理存儲(chǔ)器作為劑量矩陣數(shù)據(jù)的存儲(chǔ)器。
[0031]步驟3:實(shí)現(xiàn)內(nèi)核
在內(nèi)核中，根據(jù)每個(gè)線程分配到的位置點(diǎn)，使用紋理拾取函數(shù)進(jìn)行拾取。因?yàn)榧y理的濾波模式已設(shè)置為線性插值，即對(duì)三維紋理的八個(gè)像元根據(jù)距離進(jìn)行線性插值，返回線性插值得到的值。
[0032]步驟4:使用共享內(nèi)存存儲(chǔ)統(tǒng)計(jì)結(jié)果。
[0033]使用全局存儲(chǔ)器存儲(chǔ)每個(gè)線程的投票更新，僅需要在全局存儲(chǔ)器上開辟一個(gè)大小為100的數(shù)組，每個(gè)線程的投票直接更新對(duì)應(yīng)的數(shù)組元素。如果有多個(gè)線程同時(shí)對(duì)一個(gè)bin進(jìn)行投票時(shí)，就會(huì)出現(xiàn)寫沖突。若使用并行的N個(gè)線程對(duì)一個(gè)器官進(jìn)行統(tǒng)計(jì)，之后將結(jié)果寫入100個(gè)bin中。假如N為7時(shí)，如圖5所示。
[0034]這種同時(shí)對(duì)一個(gè)地址進(jìn)行訪問所造成的寫沖突，CUDA并不能保證其正確性，只能依靠CUDA提供的原子操作來(lái)進(jìn)行可靠的訪問，而原子操作的原理是對(duì)設(shè)備總線上鎖，必然會(huì)嚴(yán)重影響統(tǒng)計(jì)的速度。
[0035]本技術(shù)方案結(jié)合放射治療計(jì)劃中劑量分布的特點(diǎn)提出了一種方法，可以極大降低bank沖突帶來(lái)的延遲，具體方法為在共享內(nèi)存上開辟N個(gè)子劑量盒，使子劑量盒之間對(duì)應(yīng)的bin不在同一 bank。同一 warp中的相鄰的若干個(gè)線程分別更新對(duì)應(yīng)的子劑量盒。當(dāng)N=4時(shí)，共享內(nèi)存中有4個(gè)子劑量盒，若某warp相鄰四個(gè)線程的投票數(shù)據(jù)分別為98，97，98,97時(shí)，根據(jù)此算法就會(huì)將這四個(gè)數(shù)分別放入四個(gè)不同bank中，避免訪問沖突，如圖6所示。
[0036]本發(fā)明以統(tǒng)計(jì)重離子放射治療的劑量數(shù)據(jù)為例說明。
[0037]圖7a和圖7b為重尚子分層照射劑量分布不意圖
由于重離子照射時(shí)具有的布拉格峰值，可以使高劑量區(qū)集中到癌癥靶區(qū)，如圖7a所示的區(qū)域。放療時(shí)是采用分層照射方法，將高劑量區(qū)可視化后得到如圖7b，可見數(shù)據(jù)分布呈波紋狀,接近于退化分布,會(huì)導(dǎo)致嚴(yán)重的bank沖突。
[0038]實(shí)驗(yàn)以C++語(yǔ)言為基礎(chǔ)，在VS2010環(huán)境下編譯。實(shí)驗(yàn)平臺(tái)所使用的CPU為IntelCeleron E3300 @ 2.5GHz，GPU 為英偉達(dá) GeForce GTX 650 Ti，計(jì)算機(jī)系統(tǒng)為 MicrosoftWindows XP Professional Service Pack 3,內(nèi)存 2G。
[0039]本實(shí)驗(yàn)主要分析分別使用傳統(tǒng)的基于CPU和本發(fā)明基于GPUCUDA流統(tǒng)計(jì)DVH圖時(shí)的加速效果。實(shí)驗(yàn)數(shù)據(jù)為4個(gè)器官，皮膚、左肺、右肺和腫瘤靶區(qū)，共262條輪廓線。劑量數(shù)據(jù)為x、y、z軸各間隔0.4cm, 116*60*84的三維矩陣。首先用CPU統(tǒng)計(jì)10次，得到平均的運(yùn)行結(jié)果時(shí)間為436.56ms,如表一?；贑UDA流統(tǒng)計(jì)時(shí)利用GPU、紋理存儲(chǔ)器及共享存儲(chǔ)器做并行化后統(tǒng)計(jì)，并在此基礎(chǔ)上使用避免bank沖突方法進(jìn)行改進(jìn)，分別統(tǒng)計(jì)10次，對(duì)比傳統(tǒng)基于CPU方法的加速比如表二。如圖8a和圖8b所示，分別用兩種方法統(tǒng)計(jì)出的DVH圖是相同的。
[0040]表一、CPU統(tǒng)計(jì)DVH圖時(shí)間表:
【權(quán)利要求】
1.一種基于CUDA流和共享內(nèi)存的DVH圖并行統(tǒng)計(jì)方法，其特征在于，包括以下步驟: 步驟1:在主機(jī)端對(duì)器官進(jìn)行采樣，并將采樣點(diǎn)位置傳入設(shè)備端； DVH圖統(tǒng)計(jì)時(shí)需要判斷采樣點(diǎn)是否在器官的輪廓線內(nèi)，CPU判斷得到每個(gè)器官的所有采樣點(diǎn)位置并存儲(chǔ)到數(shù)組中，采樣點(diǎn)位置用一個(gè)三維向量表示即pos = (x, y, z)； 利用CUDA的流機(jī)制對(duì)得到的位置數(shù)組傳入GPU中進(jìn)行計(jì)算，不同的流在執(zhí)行時(shí)是相互獨(dú)立的，即每個(gè)流處理器處理每個(gè)器官的位置數(shù)組； 步驟2:使用紋理存儲(chǔ)器載入劑量矩陣: 將已有劑量矩陣數(shù)據(jù)從主機(jī)端拷貝到設(shè)備端，并將劑量矩陣數(shù)據(jù)存儲(chǔ)到紋理存儲(chǔ)器中； 步驟3:根據(jù)每個(gè)線程分配到的位置點(diǎn)，使用紋理拾取進(jìn)行拾取，紋理的濾波模式設(shè)置為線性插值，即對(duì)三維紋理的八個(gè)像元根據(jù)距離進(jìn)行線性插值，并返回線性插值得到的值； 步驟4:使用共享內(nèi)存存儲(chǔ)統(tǒng)計(jì)結(jié)果，通過在共享內(nèi)存上開辟N個(gè)子劑量盒，解決共享內(nèi)存會(huì)出現(xiàn)的bank沖突問題,并加快了統(tǒng)計(jì)速度。
2.根據(jù)權(quán)利要求1所述的基于CUDA流和共享內(nèi)存的DVH圖并行統(tǒng)計(jì)方法，其特征在于，上述步驟4中使用共享內(nèi)存存儲(chǔ)統(tǒng)計(jì)結(jié)果主要為:在共享內(nèi)存上開辟N個(gè)子劑量盒，使子劑量盒之間對(duì)應(yīng)的bin不在同一bank,同一warp中的相鄰的若干個(gè)線程分別更新對(duì)應(yīng)的子劑量盒。
3.根據(jù)權(quán)利要求2所述的基于CUDA流和共享內(nèi)存的DVH圖并行統(tǒng)計(jì)方法，其特征在于，上述N的取值范圍為ISiV S32。
4.根據(jù)權(quán)利要求3所述的基于CUDA流和共享內(nèi)存的DVH圖并行統(tǒng)計(jì)方法，其特征在于，所述N=8。
【文檔編號(hào)】G06T1/60GK103810670SQ201410033988
【公開日】2014年5月21日 申請(qǐng)日期:2014年1月24日 優(yōu)先權(quán)日:2014年1月24日
【發(fā)明者】王陽(yáng)萍, 黨建武, 蔣偑釗, 杜曉剛, 王松, 楊景玉, 陳永, 郭治成, 鄧沖, 趙庶旭, 閔永智, 張?chǎng)? 羅維薇 申請(qǐng)人:蘭州交通大學(xué)