2012-10-08 63 views
1

我正在研究一個我們在cuda中運行sift算法的學校項目。我在一個點來計算基於它的鄰居的值的圖像的每個像素(X)的幅度值(A,B,C,d):用共享內存計算cuda中像素的大小

A 
B X C 
    D 

我設法通過使用以使其全局內存,因爲我可以很容易地從我的輸入數組中獲得我想要的值。

但是現在我想通過首先將輸入數組放入共享內存中,但是如何讓線程在共享內存上放置正確的像素,我遇到了非常困難的時間。我必須考慮圖像邊框上的填充。

我知道我需要更多的共享內存比我想放在那裏的圖像部分,以便填充將包括但我不知道我的線程塊應該包含比共享內存空間更多或更少的線程和如何指定要閱讀的內容。 如果有人可以給我一個關於如何思考的一般想法,我可以從那裏拿...

謝謝!

+0

對於圖像處理應用程序,共享內存使用情況有點複雜。我會建議你使用CUDA紋理。紋理讀取被緩存,處理邊界情況非常簡單。紋理最適合於社區操作。 – sgarizvi

+0

哪種計算能力具有您的圖形設備? – pQB

回答

1

我提供了一種通過灰度圖像和底肥索貝爾濾波器去的碼:(索貝爾類似於你的鄰居(濾波器A,B,C,d)函數)

#define QUANTUM_TYPE short 
__global__ void sobel_gpu(QUANTUM_TYPE *img_out, QUANTUM_TYPE *img_in, int WIDTH, int HEIGHT){ 
    int x,y; 
    x=blockDim.x*blockIdx.x+threadIdx.x; 
    y=blockDim.y*blockIdx.y+threadIdx.y; 
    QUANTUM_TYPE LUp,LCnt,LDw,RUp,RCnt,RDw; 
    int pixel; 

    if(x<WIDTH && y<HEIGHT){ 
     LUp = (x-1>=0 && y-1>=0)? img_in[(x-1)+(y-1)*WIDTH]:0; 
     LCnt= (x-1>=0)? img_in[(x-1)+y*WIDTH]:0; 
     LDw = (x-1>=0 && y+1<HEIGHT)? img_in[(x-1)+(y+1)*WIDTH]:0; 
     RUp = (x+1<WIDTH && y-1>=0)? img_in[(x+1)+(y-1)*WIDTH]:0; 
     RCnt= (x+1<WIDTH)? img_in[(x+1)+y*WIDTH]:0; 
     RDw = (x+1<WIDTH && y+1<HEIGHT)? img_in[(x+1)+(y+1)*WIDTH]:0; 
     pixel = -1*LUp + 1*RUp + 
       -2*LCnt + 2*RCnt + 
       -1*LDw + 1*RDw; 
     pixel=(pixel<0)?0:pixel; 
     pixel=(pixel>MAXRGB)?MAXRGB:pixel; 
     img_out[x+y*WIDTH]=pixel; 
    } 
} 

該代碼是工作用於全局內存並安全地處理邊界。我的完整代碼讀取一個BMP圖像並在其上應用過濾器,並將生成的BMP存儲回磁盤。它可在here(CPU和GPU實現集成,無論是Linux和Windows)。

你可以用一點工作把它變成共享內存風格。首先,你應該決定你給每個區塊多少任務。然後把這個任務分解成多個共享內存sink/drain。 CUDA SDK中的Matrix Multiply示例爲您提供了一個完美的想法。