Ошибка индексации 2D-изображений в ядре CUDA

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;

И я хочу вернуть четыре границы (я обслужу их позже).Я делаю это следующим образом:

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

Где N - количество пикселей на каждой границе, которое я не хочу вычислять.

Проблема:

Код отлично работает на всех стандартных размерах изображений.Но для некоторых случайных размеров изображения он показывает диагональную линию (и).Например, в моем случае изображение 500x333 (даже если ни одно измерение не кратно 16) показывает правильный вывод, тогда как 450x365 показывает диагональные линии на выходе.Проблема остается, даже если я просто возвращаю дополнительные потоки сетки и ничего больше:

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

Код остается прежним, некоторые входные данные работают нормально, а другие нет.Кто-нибудь может обнаружить ошибку?Я прикрепил здесь образцы ввода и вывода: ИЗОБРАЖЕНИЯ Спасибо!

Обновление:

Код ядра (упрощено для возврата входного изображения, но даетта же проблема)

__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:

Я также удалил оператор return , изменив , если состояние.Но проблема сохраняется.

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

  /*Kernel Code*/

}

Answer 1

Есть довольно много вещей, которые вы до сих пор не очень хорошо описали, но основываясь на информации, которую вы опубликовали, я построил то, что, я предполагаю, является разумным случаем воспроизведения с параметрами, которые соответствуют случаю, который вы считаете неудачным (450x 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();
}

При запуске он делает именно то, что я ожидал, перезаписывая память вывода везде, кроме первой и последней двух строк и первой и последней двух записейво всех линиях между ними.Это работает с CUDA 3.2 на OS X 10.6.5 с графическим процессором 1.2.Поэтому, что бы ни происходило в вашем коде, это не происходит в моем случае с репро, что означает либо то, что я неверно истолковал то, что вы написали, либо что-то еще, что вы не описали, что вызывает проблему.