Свертка, массив с фильтром, в CUDA

Question

Я пытаюсь свернуть массив данных, 256x256, с фильтром, 3x3 на GPU, используя разделяемую память.Я понимаю, что я должен разбить массив на блоки, а затем применить фильтр в каждом блоке.В конечном итоге это означает, что блоки с перекрытием по краям и некоторые отступы должны будут выполняться по краям, где нет данных, чтобы фильтр работал правильно.

int grid = (256/(16+3-1))*(256/(16+3-1)), где 256 - длина или ширина моего массива, 16 - длина или ширина моего блока в разделяемой памяти, 3 - длина или ширина моего фильтра, и я минус один, чтобы сделать еготак что это даже.

int thread = (16+3-1)*(16+3-1)

Теперь я называю свое ядро ​​<< >> (output, input, 256) input и output - это массив размером 256 * 256

__global__ void kernel(float *input, float *output, int size)
{
    __shared__ float tile[16+3-1][16+3-1];
    blockIdx.x = bIdx;
    blockIdy.y = bIdy;
    threadIdx.x = tIdx;
    threadIdy.y = tIdy

    //i is for input
    unsigned int iX = bIdx * 3 + tIdx;
    unsigned int iY = bIdy * 3 + tIdy;

    if (tIdx == 0 || tIdx == width || tIdy == 0 || tIdy == height)
    {
        //this will pad the outside edges
        block[tIdy][tIdx] = 0;
    }
    else 
    {
        //This will fill in the block with real data
        unsigned int iin = iY * size + iX;
        block[tIdy][tIdx] = idata[iin];
    }

    __syncthreads();

    //I believe is above is correct; below, where I do the convolution, I feel is wrong
    float result = 0;
    for(int fX=-N/2; fX<=N/2; fX++){
        for(int fY=-N/2; fY<=N/2; fY++){
            if(iY+fX>=0 && iY+fX<size && iX+fY>=0 && iX+fY<size)
                result+=tile[tIdx+fX][tIdy+fY];
        }
    }
    output[iY*size+iX] = result/(3*3);
}

Когда я запускаю код, если я запускаю часть свертки, я получаю ошибку ядра.Есть идеи?Или предложения?

Answer 1

Посмотрите пример SDK sobelFilter.

Он использует текстуру для обработки краевых случаев, слегка перегружает блоки (но кэш текстуры делает это более эффективным) и использует общую память для обработки.*

Тонкая вещь в разделяемой памяти состоит в том, что вы получаете 4-сторонние конфликты банков, если вы читаете соседние байты.Один из способов обойти это, проиллюстрированный в примере sobelFilter, - развернуть цикл 4 раза и получить доступ к каждому четвертому байту.