如何在 CUDA 的核函数中将两个 openCV 矩阵相乘？

Question

我有以下最少量的代码，想知道如何在内核函数中将两个矩阵相乘？ 例如，我不能在内核函数中创建 Mat（如在 openCV 中）。

  __global__ void myMatKernel(int N, Mat *b)
  {
       Mat a;   // creates compilation error 1

      // b = a*b;      <---- what I would need


  }

  int main (void)
  {
        Mat a(10, 1, CV_64F);
        a.setTo(Scalar(2.2));
        Mat c(1, 10, CV_64F);
        c.setTo(Scalar(3.35));
        Mat d;

        d = a*c;    // works perfectly fine, but would like to do this operation on the GPU

        Mat *b;
        cudaMallocManaged(&b, sizeof(Mat));
        cudaDeviceSynchronize();
       //assign somehow values to matrix b before passing it to the function

        myMatKernel<<<1,256>>>(1, b) ;   
        cudaFree(b);
  }

编译错误 1：“错误：不允许从 __global__ 函数（“myKernel”）调用 __host__ 函数（“cv::Mat::Mat”）”

有人可以解释/展示我如何解决这些问题吗？

Answer 1

我遇到了这个问题，我无法将 Mat 对象发送到 CUDA 内核，我通过使用 Mat 类 (.data) 中图像的指针作为参数来克服这个问题，

所以代码将是：（我还没有编译它）

  __global__ void myMatKernel(double *d_a, double *d_b)
   {
     //index of this pixel
     int j = (blockIdx.x * blockDim.x) + threadIdx.x; //width
     int i = (blockIdx.y * blockDim.y) + threadIdx.y; //height

  // b = a*b;      <---- what I would need


   }

  int main (void)
  {
    Mat a(10, 1, CV_64F);
    a.setTo(Scalar(2.2));
    Mat c(1, 10, CV_64F);
    c.setTo(Scalar(3.35));
    Mat d;
    double* d_a;
    double* d_b;




    d = a*c;    // works perfectly fine, but would like to do this  operation on the GPU

    cudaMalloc((void**) &d_a, (a.rows)*(a.cols)*sizeof(double));
    cudaMalloc((void**) &d_b, (b.rows)*(b.cols)*sizeof(double));
    cudaMemcpy(d_a, a.data, (a.rows)*(a.cols)*sizeof(double), cudaMemcpyHostToDevice); 
    cudaMemcpy(d_b, b.data, (b.rows)*(b.cols)*sizeof(double), cudaMemcpyHostToDevice);

   //assign somehow values to matrix b before passing it to the function

    myMatKernel<<<1,256>>>(a,b) ;   
    cudaDeviceSynchronize();
    cudaFree(d_a);
    cudaFree(d_b);
  }

我一直在使用 CUDA 和 openCV，检查这个code，它用于从视差图中计算占用网格（它的功能无关紧要，检查它的结构，因为它可能会有所帮助）

Answer 2

由于 OpenCV 为 cv::Mat 分配主机内存，因此您不能像在主机代码中那样在内核中使用 Mat 和相关的 OpenCV API。因此，您必须为矩阵乘法编写自己的内核。

OpenCV 提供了一个名为cv::cuda::GpuMat 的类。 OpenCV 为它们分配设备内存。但是，与 GpuMat 相关的 API 旨在用于主机代码。对于矩阵乘法，无论如何您都必须编写自己的内核。

但是我有时会发现一些用于GpuMat 的 API 很方便，例如使用其构造函数分配设备内存以及使用download() 和upload() 在主机矩阵和设备矩阵之间复制数据。此外，Gpumat 类将矩阵的属性（例如 rows、cols、type()、step 等）保存在单个结构中。这在某些情况下可能会派上用场。

以下示例代码使用GpuMat。

int main (void)
{
    Mat a{ 10, 1, CV_64FC1 }; // 10x1 matrix
    Mat b{ 1, 10, CV_64FC1 }; // 1x10 matrix
    Mat c{ 10, 10, CV_64FC1 }; // multiplying a and b results in 10x10 matrix
    a.setTo(Scalar(2.2f));
    b.setTo(Scalar(3.35f));

    cv::cuda::GpuMat d_a{ a.rows, a.cols, CV_64FC1 };
    cv::cuda::GpuMat d_b{ b.rows, b.cols, CV_64FC1 };
    cv::cuda::GpuMat d_c{ c.rows, c.cols, CV_64FC1 };

    d_a.upload(a);
    d_b.upload(b);

    MatMul<<<1, dim3(c.cols, c.rows)>>>((double*)d_a.data, d_a.step,
                                        (double*)d_b.data, d_b.step,
                                        (double*)d_c.data, d_c.step,
                                        a.cols);

    d_c.download(c);
}

__global__ void MatMul(const double* const a, const int a_step,
                       const double* const b, const int b_step,
                       double* const c, const int c_step,
                       const int a_cols)
{
    int c_row = threadIdx.y;
    int c_col = threadIdx.x;

    double sum = 0;
    for (int i = 0; i < a_cols; i++)
        sum += ((double*)((unsigned char*)a + c_row * a_step))[i]
             * ((double*)((unsigned char*)b + i * b_step))[c_col];

    ((double*)((unsigned char*)c + c_row * c_step))[c_col] = sum;
    
}

请注意，如果c（结果矩阵）的元素数超过块中的最大线程数（cc >= 2.0 为 1024），则此代码将不起作用。内核的设计应该不同。

---

编辑

((double*)((unsigned char*)c + c_row * c_step))[c_col];

上述语句访问矩阵c的c_row-th行和c_col-th列元素。该矩阵为单通道矩阵，元素类型为双。它的步骤由c_step 给出。在 OpenCV 中，step 是指每行分配的字节数。大于或等于每行实际像素的总大小，以满足内存对齐，从而使内存访问更快。

上述语句首先将c（类型为double*）转换为unsigned char*，因为c_step 以字节为单位。将c_row * c_step 添加到(unsigned char*)c 会给出指向c_row-th 行的第0 列的指针。它现在将指向double* 的指针转换为使用标准数组访问运算符[] 访问c_col-th 列。