CUDA內核中的2D圖像索引問題

Question

我正在使用CUDA對圖像進行線性過濾。我使用2D線程塊和2D網格來使問題變得自然。下面是我的索引：（高度和寬度是圖像尺寸）

dim3 BlockDim(16,16); 

dim3 GridDim; 
GridDim.x = (width + 15)/16; 
GridDim.y = (height + 15)/16; 

在內核我訪問的地點如下：

unsigned int xIndex = blockIdx.x*16+ threadIdx.x; 
unsigned int yIndex = blockIdx.y*16+ threadIdx.y; 
unsigned int tid = yIndex * width + xIndex; 

而且我想返回四至界線（I」稍後會迎合他們）。我這樣做：

if(yIndex>=height-N || xIndex>=width-N || yIndex<N || xIndex<N) 
    return; 

其中N是每個邊界的像素數目，我不想計算。

問題：

的代碼運行在所有標準圖像尺寸細。但對於一些隨機圖像尺寸，它顯示對角線。例如，在我的情況下，500x333圖像（即使沒有尺寸是16的倍數）顯示正確的輸出，而450x365顯示輸出中的對角線。該問題仍然存在，即使我只返回電網，並沒有其他的額外的線程是這樣的：

if(yIndex>=height || xIndex>=width) 
return; 

代碼保持不變，一些投入運行良好，而有的則沒有。任何人都可以發現錯誤嗎？我在這裏附上了輸入和輸出示例：IMAGES謝謝！

更新：

內核代碼（簡體返回輸入圖像，但給出了同樣的問題）

__global__ void filter_8u_c1_kernel(unsigned char* in, unsigned char* out, int width, int height, float* filter, int fSize) 
{ 
    unsigned int xIndex = blockIdx.x*BLOCK_SIZE + threadIdx.x; 
    unsigned int yIndex = blockIdx.y*BLOCK_SIZE + threadIdx.y; 
    unsigned int tid = yIndex * width + xIndex; 

    unsigned int N = filterSize/2; 

    if(yIndex>=height-N || xIndex>=width-N || yIndex<N || xIndex<N) 
     return; 

     /*Filter code removed, still gives the same problem*/ 

    out[tid] = in[tid]; 
} 

更新2：

我也刪除回報聲明通過逆轉如果條件。但問題依然存在。

if(yIndex<=height-N && xIndex<=width-N && yIndex>N && xIndex>N){ 

    /*Kernel Code*/ 

} 

Answer 1

有你還沒有描述得很好，但基於您發佈的信息相當多的東西，我建我猜是合理的攝製情況相匹配你說的情況下，它的參數它失敗（450×364 filterSize=5）：

#include <stdio.h> 
#include <assert.h> 

template<int filterSize> 
__global__ void filter_8u_c1_kernel(unsigned char* in, unsigned char* out, int width, int height, float* filter, int fSize) 
{ 
    unsigned int xIndex = blockIdx.x*blockDim.x + threadIdx.x; 
    unsigned int yIndex = blockIdx.y*blockDim.y + threadIdx.y; 
    unsigned int tid = yIndex * width + xIndex; 

    unsigned int N = filterSize/2; 

    if(yIndex>=height-N || xIndex>=width-N || yIndex<N || xIndex<N) 
     return; 

    out[tid] = in[tid]; 
} 

int main(void) 
{ 
    const int width = 450, height = 365, filterSize=5; 
    const size_t isize = sizeof(unsigned char) * size_t(width * height); 
    unsigned char * _in, * _out, * out; 

    assert(cudaMalloc((void **)&_in, isize) == cudaSuccess); 
    assert(cudaMalloc((void **)&_out, isize) == cudaSuccess); 
    assert(cudaMemset(_in, 'Z', isize) == cudaSuccess); 
    assert(cudaMemset(_out, 'A', isize) == cudaSuccess); 

    const dim3 BlockDim(16,16); 
    dim3 GridDim; 
    GridDim.x = (width + BlockDim.x - 1)/BlockDim.x; 
    GridDim.y = (height + BlockDim.y - 1)/BlockDim.y; 

    filter_8u_c1_kernel<filterSize><<<GridDim,BlockDim>>>(_in,_out,width,height,0,0); 
    assert(cudaPeekAtLastError() == cudaSuccess); 

    out = (unsigned char *)malloc(isize); 
    assert(cudaMemcpy(out, _out, isize, cudaMemcpyDeviceToHost) == cudaSuccess); 

    for(int i=0; i<width; i++) { 
     fprintf(stdout, "%d: ", i); 
     for(int j=0; j<height; j++) { 
      unsigned int idx = i + j*width; 
      fprintf(stdout, "%c", out[idx]); 
     } 
     fprintf(stdout, "\n"); 
    } 

    return cudaThreadExit(); 
} 

在運行時它不正是我所期望的，處處覆蓋與輸入輸出內存除了第一個和最後兩行，並在第一和最後兩個條目在所有的行之間。這是在OS X 10.6.5上運行的CUDA 3.2與計算1.2 GPU。因此，無論您在代碼中發生什麼，它都不會出現在我的重複案例中，這意味着我錯誤地解釋了您寫的內容，或者還有其他您沒有描述過的問題。