vcscsvcscs/tv-image-denoising
1
0
Fork 0
Code Issues Pull Requests Packages Projects Releases Wiki Activity

check if threadIdx.x > 0 && threadIdx.y > 0

Browse Source
This commit is contained in:
Vargha Csongor
2023-06-27 20:38:59 +02:00
parent 20d8872930
commit 4d8dd20b06
1 changed files with 13 additions and 5 deletions
Download Patch File Download Diff File Expand all files Collapse all files

18
src/tv_denoising.cu
View File

@@ -21,7 +21,7 @@ __global__ void tvDenoisingKernel(float* image, int width, int height, float lam
// Perform TV denoising iteratively // Perform TV denoising iteratively
for (int iteration = 0; iteration < maxIterations; ++iteration) for (int iteration = 0; iteration < maxIterations; ++iteration)
{ {
// Calculate the gradients using central differences // Calculate the gradients using central differences
gradientX[threadIdx.x][threadIdx.y] = image[index + 1] - image[index - 1]; gradientX[threadIdx.x][threadIdx.y] = image[index + 1] - image[index - 1];
gradientY[threadIdx.x][threadIdx.y] = image[index + width] - image[index - width]; gradientY[threadIdx.x][threadIdx.y] = image[index + width] - image[index - width];
@@ -29,10 +29,12 @@ __global__ void tvDenoisingKernel(float* image, int width, int height, float lam
__syncthreads(); __syncthreads();
// Apply TV denoising update rule // Apply TV denoising update rule
updatedImage[threadIdx.x][threadIdx.y] = image[index] + lambda * ( if (threadIdx.x > 0 && threadIdx.y > 0) {
gradientX[threadIdx.x][threadIdx.y] - gradientX[threadIdx.x - 1][threadIdx.y] + updatedImage[threadIdx.x][threadIdx.y] = image[index] + lambda * (
gradientY[threadIdx.x][threadIdx.y] - gradientY[threadIdx.x][threadIdx.y - 1] gradientX[threadIdx.x][threadIdx.y] - gradientX[threadIdx.x - 1][threadIdx.y] +
); gradientY[threadIdx.x][threadIdx.y] - gradientY[threadIdx.x][threadIdx.y - 1]
);
}
// Update the global image array with the updated pixel values // Update the global image array with the updated pixel values
image[index] = updatedImage[threadIdx.x][threadIdx.y]; image[index] = updatedImage[threadIdx.x][threadIdx.y];
@@ -65,6 +67,12 @@ extern "C" void TVDenoising(cv::Mat& image, float lambda, int maxIterations)
// Invoke the TV denoising kernel // Invoke the TV denoising kernel
tvDenoisingKernel<<<gridSize, blockSize>>>(d_image, width, height, lambda, maxIterations); tvDenoisingKernel<<<gridSize, blockSize>>>(d_image, width, height, lambda, maxIterations);
// Check for errors during kernel launch
cudaError_t err = cudaGetLastError();
if (err != cudaSuccess)
{
printf("Error: %s\n", cudaGetErrorString(err));
}
// Copy the denoised image data back from device to host // Copy the denoised image data back from device to host
cudaMemcpy(floatImage.ptr<float>(0), d_image, width * height * sizeof(float), cudaMemcpyDeviceToHost); cudaMemcpy(floatImage.ptr<float>(0), d_image, width * height * sizeof(float), cudaMemcpyDeviceToHost);