Ошибка сегментации в медианном фильтре ядра Cuda

Я программирую ядро ​​для медианного фильтра 3×3 и хочу применить его к изображениям. Мои изображения хранятся как float *myImage = new float[pixelCount * channelCount] в RGB.

Я запускаю поток для каждого пикселя и вычисляю все 3 цвета в каждом потоке.

Я пробовал это на разных размерах изображения с разными результатами:

• 512×512: нет ошибок cuda, ошибки при работе с cuda-memcheck
• 1024×1024 и выше: ошибки cuda и memcheck-erros

Ошибка CUDA я получаю:

an illegal memory access was encountered

(Первый) вывод cuda-memcheck:

========= CUDA-MEMCHECK
========= Invalid __global__ read of size 4
=========     at 0x00001410 in   BackwardMappingCUDAUtils::parallelMedianInImage(float*, float*, unsigned int, unsigned int, int)
=========     by thread (257,0,0) in block (127,0,0)
=========     Address 0x7f535e5c0000 is out of bounds
=========     Saved host backtrace up to driver entry point at kernel launch time
=========     Host Frame:/usr/lib/x86_64-linux-gnu/libcuda.so.1   (cuLaunchKernel + 0x2cd) [0x22b3fd]
=========     Host Frame:/path/to/libcudart.so.9.1 [0x15f70]
=========     Host Frame:/path/to/libcudart.so.9.1 (cudaLaunch + 0x14e) [0x347be]
=========     Host Frame:/path/to/build_debug  /lib/libBackwardMappingCudaUtilsD.so [0x23fc]
=========     Host Frame:/path/to/build_debug  /lib/libBackwardMappingCudaUtilsD.so  (_Z75__device_stub__ZN24BackwardMappingCUDAUtils21parallelMedianInImageEPfS0_jjiPfS_jji + 0xd6) [0x20f2]
=========     Host Frame:/path/to/build_debug/lib/libBackwardMappingCudaUtilsD.so (_ZN24BackwardMappingCUDAUtils21parallelMedianInImageEPfS0_jji + 0x36) [0x2139]
=========     Host Frame:./CUDAStream (main + 0x1476) [0xf211]
=========     Host Frame:/path/to/libc.so.6 (__libc_start_main + 0xe7) [0x21b97]
=========     Host Frame:./CUDAStream (_start + 0x2a) [0xd63a]

Ядро:

__global__
void parallelMedianInImage(float *source, float *sink, unsigned int width, unsigned int pixelCount, int channelCount)
{
unsigned int pixelID = blockIdx.x * blockDim.x + threadIdx.x;
unsigned int colorIndexRed = pixelID * channelCount;
unsigned int colorWidth = width * channelCount;
unsigned int valueCount = pixelCount * channelCount;

if(pixelID<pixelCount)
{
int validValues = 0;
bool valid[9];
int indizes[9];

indizes[0] = colorIndexRed - colorWidth - channelCount;
indizes[1] = colorIndexRed - colorWidth;
indizes[2] = colorIndexRed - colorWidth + channelCount;
indizes[3] = colorIndexRed - channelCount;
indizes[4] = colorIndexRed;
indizes[5] = colorIndexRed + channelCount;
indizes[6] = colorIndexRed + colorWidth - channelCount;
indizes[7] = colorIndexRed + colorWidth;
indizes[8] = colorIndexRed + colorWidth + channelCount;

for(int u=0;u<9;u++)
{
valid[u] = true;

if(u/3==0&&((indizes[u] / colorWidth) != (colorIndexRed / colorWidth) - 1)) valid[u] = false;
if(u/3==1&&((indizes[u] / colorWidth) != (colorIndexRed / colorWidth))) valid[u] = false;
if(u/3==2&&((indizes[u] / colorWidth) != (colorIndexRed / colorWidth) + 1)) valid[u] = false;
if(indizes[u]<0 || indizes[u]>valueCount) valid[u] = false;

if(valid[u]) validValues++;
}

for(int channel=0;channel<channelCount;channel++)
{
float values[9];
for(int u=0;u<9;u++)
{
if(valid[u])
values[u]=source[indizes[u] + channel];
else
values[u]=0.0;

}
insertionSortFloatArray(values, 9);
int middleIndex = 8 - (validValues/2);
sink[colorIndexRed + channel] = values[middleIndex];
}
}
}

Допустимые переменные предназначены для проверки того, что все значения находятся внутри границ изображения.
Функция сортировки:

__device__
void insertionSortFloatArray(float array[], int length)
{
float swapper;
for(int i=1;i<length;i++)
{
swapper = array[i];
for(int u=i-1;u>=0;u--)
{
if(array[u]>swapper)
{
array[u+1] = array[u];
array[u] = swapper;
}
}
}
}

Вызов ядра выполняется в цикле, потому что у меня есть несколько изображений:

cudaMalloc((void**)&smallUndistortedDeviceImages[reducedIndex], sizeSmall);
parallelMedianShrinking<<<(pixelCountSmall+TPB-1)/TPB,TPB>>>(undistortedDeviceImages[reducedIndex], smallUndistortedDeviceImages[reducedIndex], widthSmall, pixelCountSmall, channelCount);
error = cudaGetLastError();
if(error != cudaSuccess)
{
printf(" ### CUDA error: %s\n", cudaGetErrorString(error));
}
//removed code that copies the result to the devices and stores it as an image
cudaFree(smallLightMaskStep1Images[reducedIndex]);

Я пробовал это с разными переменными потока на блок.
Поскольку ядро ​​находится в библиотеке, которая динамически связана, cuda-memcheck не сообщает точную строку segfault. Как это всегда происходит в потоке (1,0,0), я написал printf с if(threadIdx.x==1) это дало мне все адреса переменных, которые я мог придумать, чтобы сравнить их с адресом из выходных данных cuda-memcheck, но я так и не смог найти, какая это переменная. Из-за того, что некоторые строки были закомментированы, я мог отследить их где-то рядом с последними 3 строками ядра. Функция сортировки работает и используется в другом ядре. Но когда я комментирую сортировочный вызов и просто использую values[4], тогда это работает (как копирование изображения без фильтра).

Не могу обернуть мою голову вокруг этого. Заранее спасибо. Это мой первый пост, надеюсь, я включил все, что имеет отношение, извините, если я что-то пропустил.

Спецификации: Ubuntu 18.04, Cuda V9.1.85, Geforce GTX1080 с 8 ГБ оперативной памяти