C++中GPU加速的图形渲染：揭秘高性能秘密

#include <opencv2/opencv.hpp>
#include <cuda.h>
#include <cuda_runtime.h>

__global__ void convolve(const float* in, float* out, const float* filter, int rows, int cols, int filterSize) {
  int x = blockIdx.x * blockDim.x + threadIdx.x;
  int y = blockIdx.y * blockDim.y + threadIdx.y;

  if (x < rows && y < cols) {
    float sum = 0.0f;
    for (int i = 0; i < filterSize; i++) {
      for (int j = 0; j < filterSize; j++) {
        int offsetX = x + i - filterSize / 2;
        int offsetY = y + j - filterSize / 2;
        if (offsetX >= 0 && offsetX < rows && offsetY >= 0 && offsetY < cols) {
          sum += in[offsetX * cols + offsetY] * filter[i * filterSize + j];
        }
      }
    }
    out[x * cols + y] = sum;
  }
}

int main() {
  cv::Mat image = cv::imread("image.jpg");

  cv::Size blockSize(16, 16);
  cv::Mat d_image, d_filter, d_result;

  cudaMalloc(&d_image, image.rows * image.cols * sizeof(float));
  cudaMalloc(&d_filter, 9 * sizeof(float));
  cudaMalloc(&d_result, image.rows * image.cols * sizeof(float));

  cudaMemcpy(d_image, image.data, image.rows * image.cols * sizeof(float), cudaMemcpyHostToDevice);
  cudaMemcpy(d_filter, ((float*)cv::getGaussianKernel(3, 1.5, CV_32F).data), 9 * sizeof(float), cudaMemcpyHostToDevice);

  dim3 dimGrid(image.cols / blockSize.width, image.rows / blockSize.height);
  dim3 dimBlock(blockSize.width, blockSize.height);
  convolve<<<dimGrid, dimBlock>>>(d_image, d_result, d_filter, image.rows, image.cols, 3);

  cudaMemcpy(image.data, d_result, image.rows * image.cols * sizeof(float), cudaMemcpyDeviceToHost);

  cv::imshow("Blurred Image", image);
  cv::waitKey(0);

  cudaFree(d_image);
  cudaFree(d_filter);
  cudaFree(d_result);

  return 0;
}