水線相結合，可以進一步將數據傳輸部分的時間也重疊起來，即形成CPU計算、數據傳輸、以及GPU計算三個階段的流水線。形象地說，也就是當CPU在進行第i張圖片的初始化計算時，同時將第1-Ι張圖片的初始化數據傳到GPU內存中去，而GPU此時正在處理第1-2張圖片的特征檢測和描述。通過這種異步流水線的協(xié)同機制，CPU計算，數據傳輸和GPU計算能夠有效地并行進行，從而提高性能。
[0018]另外，本發(fā)明還充分利用了CPU的剩余資源，讓多余的核也獨立完成部分算法，以提高性能。因多核處理器的普及，我們還考慮到CHJ的剩余計算資源也應得以利用。假設是4核的CPU，一個核用于控制GPU及數據傳輸，另一個核做CPU計算，那么另外的兩個核就可以獨立地完成部分圖片的特征提取計算，如圖2所示。
[0019]測試結果表明，當硬件配置為Intel Q8300的CPU和GTX260的GPU時，本發(fā)明的算法速度為172.33幀/秒，是串行的67X。而當硬件配置為為Intel 17的CPU和GTX295的GPU時，速度高達340.47幀/秒，能夠滿足實時的處理要求。
【附圖說明】
[0020]圖1為基于GPU的圖像特征提取算法細粒度實現(xiàn)示意圖。其中，灰色為GPU上的計算部分，白色為CPU上的計算部分。每個階段中，圓括號中的數目表示該階段在GPU上以多少個線程計算。若有方括號，則其中的數目表示該階段的計算分為多少個內核完成。
[0021 ]圖2為GPU與CPU異步流水協(xié)作機制示意圖。
[0022]圖3為性能測試圖。
[0023]圖4為buildDet方法的具體代碼。
[0024]圖5為bui IdDetKerne I方法的部分代碼。
[0025]圖6為流水線的部分代碼。
【具體實施方式】
[0026]下面將結合本發(fā)明的【附圖說明】及程序中的源代碼，對本發(fā)明中的技術進行詳細地描述。本發(fā)明主要利用CUDA編程模型，將圖像特征提取算法中的局部特征提取算法，以細粒度的并行方式映射到GHJ上計算，并通過GPU特性以及GPU與CPU的協(xié)同工作方式來進一步優(yōu)化GPU上算法的實現(xiàn)。本發(fā)明選取的局部特征檢索算法為目前主流的檢索算法SURF(Speeded-Up Robust Features)。下面我們將詳細描述該技術具體的實現(xiàn)，并測試該技術的性能。
[0027](I)細粒度的GPU實現(xiàn)
本發(fā)明將圖像檢索算法中的特征檢測和特征描述兩個階段映射到GPU上計算。如上文所述及圖1所示，這兩個階段可以分成幾個小步驟。計算特征值時我們以每個采樣像素點為并行粒度。而特征點定位時以每8個采樣像素點為并行粒度。計算哈爾小波變換時，每個特征點周圍有109個像素需要計算，此時每個周圍像素點為一個GPU線程。接著，該109個周圍點將對特征點的42個特征方向區(qū)域進行投票，因此特征方向計算階段以每個特征方向區(qū)域為一個GPU線程。最后，每個特征點生成向量時，由該特征點周圍4*4的區(qū)域中計算出來，此過程中以每個區(qū)域為一個GPU線程。
[0028]在具體的代碼實現(xiàn)中，我們用CUDA編程模型，此處僅以特征檢測時計算特征值步驟作為例子，其他步驟的實現(xiàn)方法皆相似。涉及特征值計算的主要有兩個方法:
void buildDet(f1at *m_det, int width, int height);
__global__ void buiIdDetKernel(float *m_det, int o, int border , intstep);
其中，buiIdDet方法內部會調用buiIdDetKerneI方法。由于buiIdDetKernel方法為__global__，因此該方法是在GPU上運行的。m_det是一個指向數組的指針，用來存放圖像的特征值，該數組空間應該在GPU方法調用此方法前就在GPU上分配好。在GPU上分配空間的方法為cudaMalloc。同時，有些數據如圖像的寬度和高度，以及原積分圖像的像素值也需要傳到GPU的內存上以供后續(xù)計算時使用，對于一些簡單的參數可用cudaMemcpyToSymbol來實現(xiàn)，對于數據量較大的數組來說可先用cudaMalloc方法在GPU上分配一段空間，然后再用cudaMemcpy方法把數據從CPU內存?zhèn)鞯紾PU內存上去。
[0029]buildDet方法的具體代碼實現(xiàn)如圖4所示，其中需要特別注意的語句用下劃線表示出來了。對于圖像的每個octave，需要算每個采樣(以step為間隔選取)像素點的特征值。特征值的具體計算在buiIdDetKernel方法中。我們把原先的圖像分成寬度*高度為BL0CK_W*BL0CK_H的一個個小圖像塊，我們的默認寬度和高度參數為16*8。每個圖像塊中又有4個interval，所以每個圖像塊中即有16*8*4個采樣像素點。每一個圖像塊對應⑶DA GPU編程中Block的概念，同一個Block中的GPU線程是共享GPU shared memory的。threads O方法定義了每一個圖像塊中的16*8*4個像素點都是一個獨立的線程，即實際上計算特征值是以采樣像素點為并行粒度的。而這樣，總共就有((ws + BL0CK_ff -1) / BL0CK_ff) * ((hs +BL0CK_H -1) / BL0CK_H)個Block，WS和hs為去掉邊界以及沒有采樣到的點的圖像寬度與高度。當數據都劃分好了以后，即可調用buildDetKernel方法。
[0030]buildDetKernel方法中的計算是在GPU上運行的，且每個GHJ線程都跑相同的代碼，只是計算的數據輸入不同。buildDetKernel方法的部分代碼如圖5所示，方法一開始就要計算該線程計算的具體是圖像中的哪個像素點，由C-COIumn，r-row，i_interval識別。
[0031](2)利用GPU特性的內存優(yōu)化
在GPU細粒度并行實現(xiàn)版本的基礎上，本發(fā)明還利用GPU的內存特性對算法做了進一步優(yōu)化。對GPU的內存優(yōu)化主要包含兩個方面:一是GPU紋理內存提供對于二維局部性的支持，可以針對算法中存在的二維局部性特點對數組訪問做優(yōu)化。二是減少圖像處理過程中對內存分配和釋放的開銷。
[0032]GPU紋理存儲器在硬件上提供對于二維，三維局部性的支持。而SURF算法中存在大量快速近似積分的計算，有明顯的二維局部性。利用GHJ的紋理存儲器對二維數據的內存訪問有很好的性能提升。將具體的二維數組綁定到紋理存儲器上，可以用CUDA中cudaBindTexture2D方法。本發(fā)明將積分圖像數組綁定到了 2D紋理存儲器上，從而能利用2D紋理存儲器的局部性支持。
[0033 ]同時，算法中在存儲原圖像、積分圖像、特征值、以及其他變量時，都是在每張圖像處理開始和結尾進行內存分配和釋放。在批量處理圖片時，這樣的分配和釋放GPU內存是冗余并且影響性能的。本發(fā)明在程序初始階段就分配好固定的內存，從而減少多余的內存分配與釋放。
[0034](3)CPU與GPU的協(xié)同機制
雖然GPU的計算能力強大，但GPU在計算時仍需要CPU來進行輔助控制，因此，如何最有效率地協(xié)同CPU與GPU的工作，則成為影響性能的重要因素。本發(fā)明的CPU與GPU協(xié)同機制主要包括兩個方面，如圖2所示:(I)用異步流水線的方式使CPU與GPU并行工作。(2)充分挖掘CPU的剩余計算能力。
[0035]所謂流水線的方式是指將整個算法分成兩個部分，CPU處理第一部分的計算，GPU處理第二部分的計算，數據以流的方式在CPU與GPU間傳輸，達到兩個硬件并行工作的目的。也就是說，CPU在做后一張圖片的計算時，GPU能夠同時處理前一張圖片的計算，以此達到并行計算的效果。我們的流水線實現(xiàn)中，CPU主要做圖像的初始化加載工作，而GPU做特征檢測和描述的計算，流水線的部分代碼如圖6所示。其中:
左圖為CPU加載圖片數據的部分代碼，右圖為GPU控制線程的部分