CUDA和OpenCV混合編程,使用CUDA的紋理內存,實現圖像的二值化以及濾波功能。
#include <cuda_runtime.h> #include <highgui/highgui.hpp>#include <imgPRoc/imgproc.hpp>using namespace cv;int width = 512;int height = 512;// 2維紋理texture<float, 2, cudaReadModeElementType> texRef;// 核函數__global__ void transformKernel(uchar* output, int width, int height){ // 根據tid bid計算歸一化的拾取坐標 unsigned int x = blockIdx.x * blockDim.x + threadIdx.x; unsigned int y = blockIdx.y * blockDim.y + threadIdx.y; float u = x / (float)width; float v = y / (float)height; //從紋理存儲器中拾取數據,并寫入顯存 output[(y * width + x)] = tex2D(texRef, u / 4, v / 4);}int main(){ // 分配CUDA數組 cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc(32, 0, 0, 0, cudaChannelFormatKindFloat); cudaArray* cuArray; cudaMallocArray(&cuArray, &channelDesc, width, height); //使用OpenCV讀入圖像 Mat image = imread("D://lena.jpg", 0); resize(image, image, Size(width, height)); imshow("原始圖像", image); cudaMemcpyToArray(cuArray, 0, 0, image.data, width*height, cudaMemcpyHostToDevice); // 設置紋理屬性 texRef.addressMode[0] = cudaAddressModeWrap; //循環尋址方式 texRef.addressMode[1] = cudaAddressModeWrap; texRef.filterMode = cudaFilterModeLinear; //線性濾波 texRef.normalized = true; //歸一化坐標 //綁定紋理 cudaBindTextureToArray(texRef, cuArray, channelDesc); Mat imageOutput = Mat(Size(width, height), CV_8UC1); uchar * output = imageOutput.data; cudaMalloc((void**)&output, width * height * sizeof(float)); dim3 dimBlock(16, 16); dim3 dimGrid((width + dimBlock.x - 1) / dimBlock.x, (height + dimBlock.y - 1) / dimBlock.y); transformKernel << <dimGrid, dimBlock >> > (output, width, height); cudaMemcpy(imageOutput.data, output, height*width, cudaMemcpyDeviceToHost); imshow("CUDA+OpenCV濾波", imageOutput); waitKey(); cudaFreeArray(cuArray); cudaFree(output);}原始lena圖像:
運行效果:
新聞熱點
疑難解答