OpenCL 基础概念
- 平台 (Platform): OpenCL 实现的顶层容器,通常对应于一个 OpenCL 的实现厂商。
- 设备 (Device): 执行 OpenCL 代码的硬件单元,如 CPU、GPU 或加速器。
- 上下文 (Context): 管理设备和相关资源的环境。一个上下文可以包含多个设备。
- 命令队列 (Command Queue): 向设备发送命令的队列。每个命令队列与一个特定的设备相关联。
- 程序 (Program): OpenCL C 代码及其编译后的二进制。它包含一个或多个内核。
- 内核 (Kernel): 在设备上执行的函数。这是 OpenCL 程序的核心部分。
- 工作项 (Work-item): 内核执行的一个实例,类似于一个线程。
- 工作组 (Work-group): 工作项的集合。同一工作组中的工作项可以共享局部内存和同步。
OpenCL 程序的基本结构
一个典型的 OpenCL 程序包括以下步骤:
- 获取平台和设备信息
- 创建上下文
- 创建命令队列
- 创建和构建程序
- 创建内核
- 创建内存对象
- 设置内核参数
- 执行内核
- 读取结果
- 清理资源
使用vcpkg安装opencl
vcpkg install opencl:x64-windows
C++代码
#include <CL/opencl.hpp>
#include <opencv2/opencv.hpp>
#include <iostream>
#include <fstream>
#include <vector>
// 读取OpenCL内核源代码
std::string readKernelSource(const char* filename) {
std::ifstream file(filename);
return std::string(std::istreambuf_iterator<char>(file),
std::istreambuf_iterator<char>());
}
int main(int argc, char** argv) {
// 读取图像
cv::Mat image = cv::imread("C:\\Users\\Pictures\\7.jpg", cv::IMREAD_GRAYSCALE);//图片大小:6144*8192
if (image.empty()) {
std::cerr << "Error: Could not read image." << std::endl;
return -1;
}
// 获取OpenCL平台
std::vector<cl::Platform> platforms;
cl::Platform::get(&platforms);
if (platforms.empty()) {
std::cerr << "No OpenCL platforms found." << std::endl;
return -1;
}
// 选择第一个平台
cl::Platform platform = platforms[0];
// 获取GPU设备
std::vector<cl::Device> devices;
platform.getDevices(CL_DEVICE_TYPE_GPU, &devices);
if (devices.empty()) {
std::cerr << "No OpenCL devices found." << std::endl;
return -1;
}
// 选择第一个设备
cl::Device device = devices[0];
// 创建上下文和命令队列
cl::Context context(device);
cl::CommandQueue queue(context, device);
// 读取并编译OpenCL程序
std::string kernelSource = readKernelSource("F:\\WORK\\demo\\opgncl-test\\opgncl-test\\edge_detection.cl");
cl::Program program(context, kernelSource);
if (program.build({ device }) != CL_SUCCESS) {
std::cerr << "Error building: " << program.getBuildInfo<CL_PROGRAM_BUILD_LOG>(device) << std::endl;
return -1;
}
// 创建内核
cl::Kernel kernel(program, "sobel_edge_detection");
// 创建输入和输出缓冲区
cl::Buffer inputBuffer(context, CL_MEM_READ_ONLY | CL_MEM_COPY_HOST_PTR,
image.total() * sizeof(uchar), image.data);
cl::Buffer outputBuffer(context, CL_MEM_WRITE_ONLY,
image.total() * sizeof(uchar));
auto t = clock();
// 设置内核参数
kernel.setArg(0, inputBuffer);
kernel.setArg(1, outputBuffer);
kernel.setArg(2, image.cols);
kernel.setArg(3, image.rows);
// 执行内核(工作组 工作项)
cl::NDRange global(image.cols, image.rows);
queue.enqueueNDRangeKernel(kernel, cl::NullRange, global, cl::NullRange);
std::cout << "time1 : " << clock() - t << std::endl;//0ms
// 读取结果
t = clock();
cv::Mat result(image.size(), CV_8UC1);
queue.enqueueReadBuffer(outputBuffer, CL_TRUE, 0,
image.total() * sizeof(uchar), result.data);
std::cout << "time2 : " << clock() - t << std::endl;//30ms
t = clock();
cv::Mat dst;
cv::Canny(image, dst, 100, 200);
std::cout << "time3 : " << clock() - t << std::endl;//65ms
// 显示原图和结果
cv::imshow("Original Image", image);
cv::imshow("Edge Detection Result", result);
cv::imshow("cv::Canny", dst);
cv::waitKey(0);
return 0;
}
核函数
__kernel void sobel_edge_detection(__global const uchar* input,
__global uchar* output,
int width,
int height)
{
int x = get_global_id(0);
int y = get_global_id(1);
if (x < width && y < height) {
int idx = y * width + x;
// 如果是边界像素,直接设置为0
if (x == 0 || x == width - 1 || y == 0 || y == height - 1) {
output[idx] = 0;
return;
}
// 定义Sobel算子
int Gx[3][3] = {{-1, 0, 1},
{-2, 0, 2},
{-1, 0, 1}};
int Gy[3][3] = {{-1, -2, -1},
{ 0, 0, 0},
{ 1, 2, 1}};
int sum_x = 0, sum_y = 0;
// 应用Sobel算子
for (int i = -1; i <= 1; i++) {
for (int j = -1; j <= 1; j++) {
int pixel = input[(y + i) * width + (x + j)];
sum_x += pixel * Gx[i+1][j+1];
sum_y += pixel * Gy[i+1][j+1];
}
}
// 计算梯度幅值
int sum = abs(sum_x) + abs(sum_y);
//output[idx] = (sum > 255) ? 255 : sum;
output[idx] = (sum > 200) ? 255 : 0;
}
}
实验结果:
原图
opencl
opencv
调试信息:
time1 : 0
time2 : 30
time3 : 65
时间:
opencl | opencv |
0ms | 65ms |
可以看出,对6144*8192的图片执行查找边缘,GPU执行速度很快,不到1ms;CPU执行要65ms,当然从显卡读取数据要30ms,如果不算数据拷贝的时间,opencv大概是35ms