Почему мое ядро ​​свертки (размытия) CUDA оставляет полосы серого при фильтрации 2D-изображений в оттенках серого, считы ⇐ C++

[Anonymous]

Я разрабатываю программу 2D-свертки (размытия), которая будет получать изображения в оттенках серого в качестве входных данных и выполнять алгоритм свертки параллельно с CUDA для C++. Мой фильтр (ядро) представляет собой двумерную матрицу 5 x 5 со всеми значениями, инициализированными как значения с плавающей запятой 1.0f/25.0f. В настоящее время я пытаюсь найти крайний случай, который мне не хватает либо в моем ядре CUDA, либо когда я запускаю ядро ​​CUDA на основе рассчитанных размеров сетки и блока. Я намеренно не обрабатываю границу в 2 пикселя. Я не получаю никаких ошибок CUDA, просто неправильный результат при запуске с прикрепленным файлом Tree-grayscale.jpg
Вот минимальный воспроизводимый пример

Код: Выделить всё

#include
#include
#include
#include
#include
#include
#include

using namespace std;
using namespace cv;

#define FILTER_RADIUS 2
const float F_h[2 * FILTER_RADIUS + 1][2 * FILTER_RADIUS + 1] = {
{1.0f / 25, 1.0f / 25, 1.0f / 25, 1.0f / 25, 1.0f / 25},
{1.0f / 25, 1.0f / 25, 1.0f / 25, 1.0f / 25, 1.0f / 25},
{1.0f / 25, 1.0f / 25, 1.0f / 25, 1.0f / 25, 1.0f / 25},
{1.0f / 25, 1.0f / 25, 1.0f / 25, 1.0f / 25, 1.0f / 25},
{1.0f / 25, 1.0f / 25, 1.0f / 25, 1.0f / 25, 1.0f / 25}
};
__constant__ float F[2*FILTER_RADIUS+1][2*FILTER_RADIUS+1];

#define CHECK(call) { \
const cudaError_t cuda_ret = call; \
if(cuda_ret != cudaSuccess){\
printf("Error: %s:%d, ", __FILE__, __LINE__); \
printf("code: %d, reason:%s\n", cuda_ret, cudaGetErrorString(cuda_ret)); \
exit(-1); \
} \
}

double myCPUTimer(){
struct timeval tp;
gettimeofday(&tp, NULL);
return ((double)tp.tv_sec + (double)tp.tv_usec/1.0e6);
}

__global__ void blurImage_Kernel(unsigned char * Pout, unsigned char * Pin, unsigned int width, unsigned int height){
int outCol = blockIdx.x*blockDim.x + threadIdx.x;
int outRow = blockIdx.y*blockDim.y + threadIdx.y;

if (outCol < width && outRow < height) {
float Pvalue = 0.0f;
int inRow = 0, inCol = 0;
int condition = 2*FILTER_RADIUS + 1;

for(int fRow = 0; fRow < condition; fRow++){
for(int fCol = 0; fCol < condition; fCol++){
inRow = outRow - FILTER_RADIUS + fRow;
inCol = outCol - FILTER_RADIUS + fCol;
if(inRow >= 0 && inRow < height && inCol >= 0 && inCol <  width){
Pvalue += F[fRow][fCol] * (float) Pin[inRow* width + inCol];
}
}
}

Pout[outRow*width + outCol] = (unsigned char) min(max(Pvalue, 0.0f), 255.0f);
}
}

void blurImage_d(cv::Mat Pout_Mat_h, cv::Mat Pin_Mat_h, unsigned int nRows, unsigned int nCols){

printf("\n\nblurImage_Kernel: \n");

// (1) allocate device memory for arrays p_d
unsigned char *Pin_d, *Pout_d;
double start_time_malloc = myCPUTimer();
cudaMalloc((void**) &Pin_d, sizeof(unsigned char) * nRows * nCols);
cudaMalloc((void**) &Pout_d, sizeof(unsigned char) * nRows * nCols);
double end_time_malloc = myCPUTimer();
double elapsed_time_malloc = end_time_malloc - start_time_malloc;

printf("\tcudaMalloc: \t\t\t\t\t\t\t\t%f s\n", elapsed_time_malloc);

// (2) copy image matrix Pin_h to device memory Pin_d
unsigned char * Pin_h = Pin_Mat_h.data;
double start_time_memcpy = myCPUTimer();
cudaMemcpy(Pin_d, Pin_h, sizeof(unsigned char)*nCols*nRows, cudaMemcpyHostToDevice);
double end_time_memcpy = myCPUTimer();
double elapsed_time_memcpy = end_time_memcpy - start_time_memcpy;

printf("\tcudaMemcpy: \t\t\t\t\t\t\t\t%f s\n", elapsed_time_memcpy);

// (3) call kernel to launch a grid of threads to perform the image convolution on GPU
dim3 gridDim((nRows + 32 - 1) / 32, (nCols + 32 - 1) / 32);
dim3 blockDim(32, 32);

double start_time = myCPUTimer();
blurImage_Kernel(Pout_d, Pin_d, nRows, nCols);
CHECK(cudaGetLastError());
CHECK(cudaDeviceSynchronize());
double end_time = myCPUTimer();
double elapsed_time = end_time - start_time;

printf("\tblurImage_Kernel: \t\t\t%f s\n", (nCols + 32 - 1) / 32, (nRows + 32 - 1) / 32, 32, 32, elapsed_time);

// (4) Copy the result data from device memory of array Pout_d to host memory of array Pout_h
Pout_Mat_h = cv::Mat::zeros(nRows, nCols, CV_8U);
unsigned char * Pout_h = Pout_Mat_h.data;
double start_time_memcpy2 = myCPUTimer();
cudaMemcpy(Pout_h, Pout_d, sizeof(unsigned char)*nCols*nRows, cudaMemcpyDeviceToHost);
double end_time_memcpy2 = myCPUTimer();
double elapsed_time_memcpy2 = end_time_memcpy2 - start_time_memcpy2;

printf("\tcudaMemcpy: \t\t\t\t\t\t\t\t%f s\n\n", elapsed_time_memcpy2);

double total_elapsed_time = elapsed_time_malloc + elapsed_time_memcpy + elapsed_time + elapsed_time_memcpy2;

printf("Total elapsed time for convolution without tiling: %f s\n", total_elapsed_time);

// (5) free device memory of Pin_d and Pout_d
cudaFree(Pin_d);
cudaFree(Pout_d);

}

int main(int argc, char* argv[]){

char* file_name = argv[1];
printf("Blurring file: \"%s\"\n", file_name);

Mat Pin_Mat_h = cv::imread(file_name, IMREAD_GRAYSCALE);
unsigned int nRows = Pin_Mat_h.rows, nCols = Pin_Mat_h.cols, nChannels = Pin_Mat_h.channels();

printf("\n\n\n Dimension of image: %d %d \n\n\n", nRows, nCols);

// for comparison purpose, implement a CUDA kernel but without tiling
Mat Pout_Mat_h(nRows, nCols, CV_8UC1);
cudaMemcpyToSymbol(F, F_h, (2*FILTER_RADIUS+1)*(2*FILTER_RADIUS+1)*sizeof(float));
blurImage_d(Pout_Mat_h, Pin_Mat_h, nRows, nCols);

std::time_t t = std::time(nullptr);
std::tm* now = std::localtime(&t);
std::ostringstream oss_kernel;
oss_kernel 

Подробнее здесь: [url]https://stackoverflow.com/questions/79183250/why-is-my-cuda-convolutionblurring-kernel-leaving-streaks-of-gray-when-filteri[/url]