How to use shared memory in cuda-c for sobelx

The problem is when copying data from global memory to shared memory and applying the Sobel filter because the boundary conditions are defined for each shared memory block. As a result, when copying the processed data back to global memory, the boundary information is not correctly maintained so how to fix it

here is the code:

​

#include <opencv2/opencv.hpp>
#include <cuda_runtime.h>
const int FILTER_WIDTH = 3;
const int FILTER_HEIGHT = 3;
// Define kernel weights (modify these as needed)
__constant__ float Kx[FILTER_HEIGHT][FILTER_WIDTH] = {
1, 0, -1,
2, 0, -2,
1, 0, -1};
// Define tile dimensions for shared memory
const int TILE_WIDTH = 32;
const int TILE_HEIGHT = 32;
__global__ void simple_sobelX(const unsigned char *src, unsigned char *dest, int width, int height)
{
int x = blockIdx.x * blockDim.x + threadIdx.x;
int y = blockIdx.y * blockDim.y + threadIdx.y;
// Define shared memory for a tile
__shared__ float tile[TILE_HEIGHT + FILTER_HEIGHT - 1][TILE_WIDTH + FILTER_WIDTH - 1];
// Compute the index in the shared memory
int shared_x = threadIdx.x;
int shared_y = threadIdx.y;
// Load the tile of the input image into shared memory
int src_x = x - FILTER_WIDTH / 2;
int src_y = y - FILTER_HEIGHT / 2;
if (src_x >= 0 && src_x < width && src_y >= 0 && src_y < height)
{
tile[shared_y][shared_x] = src[src_y * width + src_x];
}
else
{
tile[shared_y][shared_x] = 0.0f;
}
__syncthreads();
if (x >= FILTER_WIDTH / 2 && x < (width - FILTER_WIDTH / 2) && y >= FILTER_HEIGHT / 2 && y < (height - FILTER_HEIGHT / 2))
{
// Compute the convolution in parallel using shared memory
float Gx = 0;
for (int ky = 0; ky < FILTER_HEIGHT; ky++)
{
for (int kx = 0; kx < FILTER_WIDTH; kx++)
{
Gx += tile[shared_y + ky][shared_x + kx] * Kx[ky][kx];
}
}
float Gx_abs = (Gx < 0) ? 0 : Gx;
Gx_abs = (Gx > 255) ? 255 : Gx;
dest[y * width + x] = Gx_abs;
}
}
int main()
{
cv::Mat image = cv::imread("cameraman.png", cv::IMREAD_GRAYSCALE);
int width = image.cols;
int height = image.rows;
unsigned char *d_image, *d_result;
cudaMalloc((void **)&d_image, width * height * sizeof(unsigned char));
cudaMalloc((void **)&d_result, width * height * sizeof(unsigned char));
cudaMemcpy(d_image, image.data, width * height * sizeof(unsigned char), cudaMemcpyHostToDevice);
dim3 threads(32, 32);
dim3 blocks((width + threads.x - 1) / threads.x, (height + threads.y - 1) / threads.y);
simple_sobelX<<<blocks, threads>>>(d_image, d_result, width, height);
unsigned char *h_result = new unsigned char[width * height];
cudaMemcpy(h_result, d_result, width * height * sizeof(unsigned char), cudaMemcpyDeviceToHost);
cv::Mat result(height, width, CV_8U, h_result);
cudaFree(d_image);
cudaFree(d_result);
cv::imshow("Result", result);
cv::waitKey(0);
cv::imwrite("result.png", result);
delete[] h_result;
return 0;
}

input

output