Problems in implementing adaptive thresholding using CUDA

huangapple
go评论107阅读模式
英文:

Problems in implementing adaptive thresholding using CUDA

问题

以下是代码的翻译部分：

  1. #include "cuda_runtime.h"
  2. #include "device_launch_parameters.h"
  3. #include <stdio.h>
  4. #include <iostream>
  5. #include <opencv2/core/core.hpp>
  6. #include <opencv2/highgui/highgui.hpp>
  7. #include <opencv2/imgproc/imgproc.hpp>
  9. using namespace cv;
  10. using namespace std;
  12. // CUDA核心函数
  13. __global__ void adaptiveThresholdCUDA(const uchar* image, uchar* binary, const int* iimage, int nl, int nc, int halfSize, int blockSize, int threshold) {
  14.     // 计算全局索引
  15.     int j = blockIdx.y * blockDim.y + threadIdx.y;
  16.     int i = blockIdx.x * blockDim.x + threadIdx.x;
  18.     // 检查边界条件
  19.     if (j >= nl - halfSize - 1 || i >= nc - halfSize - 1)
  20.         return;
  22.     // 获取行地址
  23.     const uchar* data = image + j * nc;
  24.     uchar* binaryRow = binary + j * nc;
  25.     const int* idata1 = iimage + (j - halfSize) * nc;
  26.     const int* idata2 = iimage + (j + halfSize + 1) * nc;
  28.     // 计算总和
  29.     int sum = (idata2[i + halfSize + 1] - idata2[i - halfSize] - idata1[i + halfSize + 1] + idata1[i - halfSize]) / (blockSize * blockSize);
  31.     // 应用自适应阈值
  32.     if (data[i] < (sum - threshold))
  33.         binaryRow[i] = 0;
  34.     else
  35.         binaryRow[i] = 255;
  36. }
  38. int main()
  39. {
  40.     Mat image = imread("image/test.jpg", 0);
  41.     if (!image.data) return 0;
  43.     resize(image, image, Size(), 1.0, 1.0);
  44.     namedWindow("Original Image");
  45.     imshow("Original Image", image);
  47.     /* 自适应阈值处理函数 */
  48.     int blockSize = 16; // 邻域大小
  49.     int threshold = 10; // 像素比较阈值
  51.     Mat binary = image.clone();
  53.     int nl = binary.rows; // 行数
  54.     int nc = binary.cols; // 每行的元素总数
  56.     Mat iimage;
  57.     integral(image, iimage, CV_32S);
  59.     // 在设备(GPU)上分配内存
  60.     uchar* imageDevice;
  61.     uchar* binaryDevice;
  62.     int* iimageDevice;
  63.     cudaMalloc((void**)&imageDevice, sizeof(uchar) * nl * nc);
  64.     cudaMalloc((void**)&binaryDevice, sizeof(uchar) * nl * nc);
  65.     cudaMalloc((void**)&iimageDevice, sizeof(int) * nl * nc);
  67.     // 从主机复制输入数据到设备
  68.     cudaMemcpy(imageDevice, image.data, sizeof(uchar) * nl * nc, cudaMemcpyHostToDevice);
  69.     cudaMemcpy(binaryDevice, binary.data, sizeof(uchar) * nl * nc, cudaMemcpyHostToDevice);
  70.     cudaMemcpy(iimageDevice, iimage.data, sizeof(int) * nl * nc, cudaMemcpyHostToDevice);
  72.     // 定义网格和块维度
  73.     dim3 blockDim(16, 16);  // 你可能需要调整块维度
  74.     dim3 gridDim((nc + blockDim.x - 1) / blockDim.x, (nl + blockDim.y - 1) / blockDim.y);
  76.     // 启动CUDA核心
  77.     adaptiveThresholdCUDA<<<gridDim, blockDim>>>(imageDevice, binaryDevice, iimageDevice, nl, nc, blockSize / 2, blockSize, threshold);
  78. }

希望这能帮助你。如果你有任何其他问题，请随时提出。

英文:

If you run the code below, the applied image will look like this: What do I need to fix to make it work correctly?

  1. #include &quot;cuda_runtime.h&quot;
  2. #include &quot;device_launch_parameters.h&quot;
  3. #include &lt;stdio.h&gt;
  4. #include &lt;iostream&gt;
  5. #include &lt;opencv2/core/core.hpp&gt;
  6. #include &lt;opencv2/highgui/highgui.hpp&gt;
  7. #include &lt;opencv2/imgproc/imgproc.hpp&gt;
  9. using namespace cv;
  10. using namespace std;
  12. // CUDA kernel function
  13. __global__ void adaptiveThresholdCUDA(const uchar* image, uchar* binary, const int* iimage, int nl, int nc, int halfSize, int blockSize, int threshold) {
  14.     // Calculate global indices
  15.     int j = blockIdx.y * blockDim.y + threadIdx.y;
  16.     int i = blockIdx.x * blockDim.x + threadIdx.x;
  18.     // Check boundary conditions
  19.     if (j &gt;= nl - halfSize - 1 || i &gt;= nc - halfSize - 1)
  20.         return;
  22.     // Get row addresses
  23.     const uchar* data = image + j * nc;
  24.     uchar* binaryRow = binary + j * nc;
  25.     const int* idata1 = iimage + (j - halfSize) * nc;
  26.     const int* idata2 = iimage + (j + halfSize + 1) * nc;
  28.     // Calculate sum
  29.     int sum = (idata2[i + halfSize + 1] - idata2[i - halfSize] - idata1[i + halfSize + 1] + idata1[i - halfSize]) / (blockSize * blockSize);
  31.     // Apply adaptive threshold
  32.     if (data[i] &lt; (sum - threshold))
  33.         binaryRow[i] = 0;
  34.     else
  35.         binaryRow[i] = 255;
  36. }
  38. int main()
  39. {
  40.     Mat image = imread(&quot;image/test.jpg&quot;, 0);
  41.     if (!image.data) return 0;
  43.     resize(image, image, Size(), 1.0, 1.0);
  44.     namedWindow(&quot;Original Image&quot;);
  45.     imshow(&quot;Original Image&quot;, image);
  47.     /*          Function for Adaptive Thresholding          */
  48.     int blockSize = 16; // Neighborhood size
  49.     int threshold = 10; // Pixel comparison threshold
  51.     Mat binary = image.clone();
  53.     int nl = binary.rows; // Number of lines
  54.     int nc = binary.cols; // Total number of elements per line
  56.     Mat iimage;
  57. &gt; integral(image, iimage, CV_32S);
  59.     // Allocate memory on device (GPU)
  60.     uchar* imageDevice;
  61.     uchar* binaryDevice;
  62.     int* iimageDevice;
  63.     cudaMalloc((void**)&amp;imageDevice, sizeof(uchar) * nl * nc);
  64.     cudaMalloc((void**)&amp;binaryDevice, sizeof(uchar) * nl * nc);
  65.     cudaMalloc((void**)&amp;iimageDevice, sizeof(int) * nl * nc);
  67.     // Copy input data from host to device
  68.     cudaMemcpy(imageDevice, image.data, sizeof(uchar) * nl * nc, cudaMemcpyHostToDevice);
  69.     cudaMemcpy(binaryDevice, binary.data, sizeof(uchar) * nl * nc, cudaMemcpyHostToDevice);
  70.     cudaMemcpy(iimageDevice, iimage.data, sizeof(int) * nl * nc, cudaMemcpyHostToDevice);
  72.     // Define grid and block dimensions
  73.     dim3 blockDim(16, 16);  // You may need to adjust the block dimensions
  74.     dim3 gridDim((nc + blockDim.x - 1) / blockDim.x, (nl + blockDim.y - 1) / blockDim.y);
  76.     // Launch the CUDA kernel
  77.     adaptiveThresholdCUDA &lt;&lt; &lt;gridDim, blockDim &gt;&gt; &gt; (imageDevice, binaryDevice, iimageDevice, nl, nc, blockSize / 2, blockSize, threshold);
  79. }

Problems in implementing adaptive thresholding using CUDA

Problems in implementing adaptive thresholding using CUDA

When changing the size of blockDim and applying it, if it is set too large or too small, it may not be applied.

答案1

得分: 3

以下是翻译好的部分：

问题在于iimage比image多了一行和一列。

根据cv::integral的文档：

src - 输入图像，大小为W×H，8位或浮点型（32f或64f）。
sum - 积分图像，大小为(W+1)×(H+1)，32位整数或浮点型（32f或64f）。

发布示例的图像大小为332列×302行。
图像大小为iimage是333列×303行。

我们需要将iimage复制到具有一行和一列更多的设备上：

  1. int iimage_nl = iimage.rows; // = nl + 1
  2. int iimage_nc = iimage.cols; // = nc + 1
  4. cudaMalloc((void**)&amp;iimageDevice, sizeof(int) * iimage_nl * iimage_nc);
  6. cudaMemcpy(iimageDevice, iimage.data, sizeof(int) * iimage_nl * iimage_nc, cudaMemcpyHostToDevice);

我们需要根据iimage的实际大小来修复内核：

  1. __global__ void adaptiveThresholdCUDA(...) {
  2.     ...
  3.     int iimage_nc = nc + 1; // iimage中的列数等于（nc+1）
  4.     
  5.     const int* idata1 = iimage + (j - halfSize) * iimage_nc;
  6.     const int* idata2 = iimage + (j + halfSize + 1) * iimage_nc;
  7.     ...
  8. }

更新的代码示例：

  1. #include "cuda_runtime.h"
  2. #include "device_launch_parameters.h"
  4. #include <stdio.h>
  5. #include <iostream>
  6. #include "opencv2/opencv.hpp"
  8. using namespace cv;
  9. using namespace std;
  11. // CUDA核函数
  12. __global__ void adaptiveThresholdCUDA(const uchar* image, uchar* binary, const int* iimage, int nl, int nc, int halfSize, int blockSize, int threshold) {
  13.     // 计算全局索引
  14.     int j = blockIdx.y * blockDim.y + threadIdx.y;
  15.     int i = blockIdx.x * blockDim.x + threadIdx.x;
  17.     // 检查边界条件
  18.     if (j >= nl - halfSize - 1 || i >= nc - halfSize - 1)
  19.         return;
  21.     // 获取行地址
  22.     int iimage_nc = nc + 1; // iimage中的列数等于(nc+1)
  23.     const uchar* data = image + j * nc;
  24.     uchar* binaryRow = binary + j * nc;
  25.     const int* idata1 = iimage + (j - halfSize) * iimage_nc;
  26.     const int* idata2 = iimage + (j + halfSize + 1) * iimage_nc;
  28.     // 计算总和
  29.     int sum = (idata2[i + halfSize + 1] - idata2[i - halfSize] - idata1[i + halfSize + 1] + idata1[i - halfSize]) / (blockSize * blockSize);
  31.     // 应用自适应阈值
  32.     if (data[i] < (sum - threshold))
  33.         binaryRow[i] = 0;
  34.     else
  35.         binaryRow[i] = 255;
  36. }
  38. int main()
  39. {
  40.     // 读取图像
  41.     Mat image = imread("test.png", 0);
  42.     if (!image.data) return 0;
  44.     resize(image, image, Size(), 1.0, 1.0);
  45.     namedWindow("Original Image");
  46.     imshow("Original Image", image);
  48.     // 自适应阈值参数
  49.     int blockSize = 16; // 邻域大小
  50.     int threshold = 10; // 像素比较阈值
  52.     Mat binary = image.clone();
  53.     int nl = binary.rows; // 行数
  54.     int nc = binary.cols; // 每行的元素总数
  56.     Mat iimage;
  57.     integral(image, iimage, CV_32S);
  59.     int iimage_nl = iimage.rows; // = nl + 1
  60.     int iimage_nc = iimage.cols; // = nc + 1
  62.     // 在设备上分配内存（GPU）
  63.     uchar* imageDevice;
  64.     uchar* binaryDevice;
  65.     int* iimageDevice;
  66.     cudaMalloc((void**)&amp;imageDevice, sizeof(uchar) * nl * nc);
  67.     cudaMalloc((void**)&amp;binaryDevice, sizeof(uchar) * nl * nc);
  68.     cudaMalloc((void**)&amp;iimageDevice, sizeof(int) * iimage_nl * iimage_nc);
  70.     // 从主机复制输入数据到设备
  71.     cudaMemcpy(imageDevice, image.data, sizeof(uchar) * nl * nc, cudaMemcpyHostToDevice);
  72.     cudaMemcpy(binaryDevice, binary.data, sizeof(uchar) * nl * nc, cudaMemcpyHostToDevice);
  73.     cudaMemcpy(iimageDevice, iimage.data, sizeof(int) * iimage_nl * iimage_nc, cudaMemcpyHostToDevice);
  75.     // 定义网格和块维度
  76.     dim3 blockDim(16, 16);  // 您可能需要调整块维度
  77.     dim3 gridDim((nc + blockDim.x - 1) / blockDim.x, (nl + blockDim.y - 1) / blockDim.y);
  79.     // 启动CUDA核函数
  80.     adaptiveThresholdCUDA<<<gridDim, blockDim>>>(imageDevice, binaryDevice, iimageDevice, nl, nc, blockSize / 2, blockSize, threshold);
  82.     Mat output = image.clone();
  83.     cudaMemcpy(output.data, binaryDevice, sizeof(uchar) * nl * nc, cudaMemcpyDeviceToHost); // 从设备复制到主机
  84.     imshow("Output Image", output);
  85.     waitKey();
  86.     destroyAllWindows();
  88.     imwrite("output.png", output);  // 保存输出
  89. }

请注意，以上解决方案仅演示了可能的解决方案（意图是显示为什么输出图像是“倾斜的”）。

英文:

The issue is that iimage has one extra row and one extra column compared to image.

According to the documentation of cv::integral:

>src - input image as W×H, 8-bit or floating-point (32f or 64f).
sum - integral image as (W+1)×(H+1) , 32-bit integer or floating-point (32f or 64f).

The image size of the posted sample is 332 columns by 302 rows.
The image size iimage is 333 columns by 303 rows.

We have to copy iimage to device with one more row and column:

  1. int iimage_nl = iimage.rows; // = nl + 1
  2. int iimage_nc = iimage.cols; // = nc + 1
  3. cudaMalloc((void**)&amp;iimageDevice, sizeof(int) * iimage_nl * iimage_nc);
  4. cudaMemcpy(iimageDevice, iimage.data, sizeof(int) * iimage_nl * iimage_nc, cudaMemcpyHostToDevice);

We have to fix the kernel according to the actual size of iimage:

  1. __global__ void adaptiveThresholdCUDA(...) {
  2. ...
  3. int iimage_nc = nc + 1; //Number of columns in iimage equals (nc+1)
  4. const int* idata1 = iimage + (j - halfSize) * iimage_nc;
  5. const int* idata2 = iimage + (j + halfSize + 1) * iimage_nc;
  6. ...

Updated code sample:

  1. #include &quot;cuda_runtime.h&quot;
  2. #include &quot;device_launch_parameters.h&quot;
  3. #include &lt;stdio.h&gt;
  4. #include &lt;iostream&gt;
  5. //#include &lt;opencv2/core/core.hpp&gt;
  6. //#include &lt;opencv2/highgui/highgui.hpp&gt;
  7. //#include &lt;opencv2/imgproc/imgproc.hpp&gt;
  8. #include &quot;opencv2/opencv.hpp&quot;
  9. using namespace cv;
  10. using namespace std;
  11. // CUDA kernel function
  12. __global__ void adaptiveThresholdCUDA(const uchar* image, uchar* binary, const int* iimage, int nl, int nc, int halfSize, int blockSize, int threshold) {
  13. // Calculate global indices
  14. int j = blockIdx.y * blockDim.y + threadIdx.y;
  15. int i = blockIdx.x * blockDim.x + threadIdx.x;
  16. // Check boundary conditions
  17. if (j &gt;= nl - halfSize - 1 || i &gt;= nc - halfSize - 1)
  18. return;
  19. // Get row addresses
  20. int iimage_nc = nc + 1; //Number of columns in iimage equals (nc+1)
  21. const uchar* data = image + j * nc;
  22. uchar* binaryRow = binary + j * nc;
  23. const int* idata1 = iimage + (j - halfSize) * iimage_nc;
  24. const int* idata2 = iimage + (j + halfSize + 1) * iimage_nc;
  25. // Calculate sum
  26. int sum = (idata2[i + halfSize + 1] - idata2[i - halfSize] - idata1[i + halfSize + 1] + idata1[i - halfSize]) / (blockSize * blockSize);
  27. // Apply adaptive threshold
  28. if (data[i] &lt; (sum - threshold))
  29. binaryRow[i] = 0;
  30. else
  31. binaryRow[i] = 255;
  32. }
  33. int main()
  34. {
  35. //Mat image = imread(&quot;image/test.jpg&quot;, 0);
  36. Mat image = imread(&quot;test.png&quot;, 0);
  37. if (!image.data) return 0;
  38. resize(image, image, Size(), 1.0, 1.0);
  39. namedWindow(&quot;Original Image&quot;);
  40. imshow(&quot;Original Image&quot;, image);
  41. /*          Function for Adaptive Thresholding          */
  42. int blockSize = 16; // Neighborhood size
  43. int threshold = 10; // Pixel comparison threshold
  44. Mat binary = image.clone();
  45. //nl = 332, nc = 302
  46. int nl = binary.rows; // Number of lines
  47. int nc = binary.cols; // Total number of elements per line
  48. Mat iimage;
  49. integral(image, iimage, CV_32S);
  50. //https://docs.opencv.org/3.4/d7/d1b/group__imgproc__misc.html#gadeaf38d7701d7ad371278d663c50c77d
  51. //src	input image as W&#215;H, 8-bit or floating-point (32f or 64f).
  52. //sum	integral image as (W+1)&#215;(H+1) , 32-bit integer or floating-point (32f or 64f).
  53. //iimage_nl = 333, iimage_nc = 303
  54. int iimage_nl = iimage.rows; // = nl + 1
  55. int iimage_nc = iimage.cols; // = nc + 1
  56. // Allocate memory on device (GPU)
  57. uchar* imageDevice;
  58. uchar* binaryDevice;
  59. int* iimageDevice;
  60. cudaMalloc((void**)&amp;imageDevice, sizeof(uchar) * nl * nc);
  61. cudaMalloc((void**)&amp;binaryDevice, sizeof(uchar) * nl * nc);
  62. cudaMalloc((void**)&amp;iimageDevice, sizeof(int) * iimage_nl * iimage_nc);
  63. // Copy input data from host to device
  64. cudaMemcpy(imageDevice, image.data, sizeof(uchar) * nl * nc, cudaMemcpyHostToDevice);
  65. cudaMemcpy(binaryDevice, binary.data, sizeof(uchar) * nl * nc, cudaMemcpyHostToDevice);
  66. cudaMemcpy(iimageDevice, iimage.data, sizeof(int) * iimage_nl * iimage_nc, cudaMemcpyHostToDevice);
  67. // Define grid and block dimensions
  68. dim3 blockDim(16, 16);  // You may need to adjust the block dimensions
  69. dim3 gridDim((nc + blockDim.x - 1) / blockDim.x, (nl + blockDim.y - 1) / blockDim.y);
  70. // Launch the CUDA kernel
  71. adaptiveThresholdCUDA&lt;&lt;&lt;gridDim, blockDim&gt;&gt;&gt;(imageDevice, binaryDevice, iimageDevice, nl, nc, blockSize / 2, blockSize, threshold);
  72. Mat output = image.clone();
  73. cudaMemcpy(output.data, binaryDevice, sizeof(uchar) * nl * nc, cudaMemcpyDeviceToHost); //Copy from device to host
  74. imshow(&quot;Output Image&quot;, output);
  75. waitKey();
  76. destroyAllWindows();
  77. imwrite(&quot;output.png&quot;, output);  //Save the output
  78. }

test.png (input image):
Problems in implementing adaptive thresholding using CUDA

output.png (output image):
Problems in implementing adaptive thresholding using CUDA

Note that the above solution just demonstrates a possible solution (the intention is to show why the output image is "slanted").

huangapple
  • 本文由 发表于 2023年6月13日 01:43:39
  • 转载请务必保留本文链接：https://go.coder-hub.com/76459109.html
  • c++
  • cuda
  • opencv
匿名

发表评论

匿名网友

:?: :razz: :sad: :evil: :!: :smile: :oops: :grin: :eek: :shock: :???: :cool: :lol: :mad: :twisted: :roll: :wink: :idea: :arrow: :neutral: :cry: :mrgreen:

确定