CUDA二維紋理內存+OpenCV圖像濾波

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圖像：

運行效果：