纹理内存（二）

纹理参考API请参考：https://blog.csdn.net/weixin_44444450/article/details/104217481

1.2.2纹理对象API

使用纹理对象包括
1.纹理对象的创建
2.访问纹理内存
3.纹理对象销毁

纹理对象的创建

1）创建纹理对象，cudaCreateTextureObject();

cudaCreateTextureObject(cudaTextureObject_t *pTexObject,
 	const struct cudaResourceDesc *pResDesc, 
 	const struct cudaTextureDesc *pTexDesc, 
 	const struct cudaResourceViewDesc *pResViewDesc);

这里我们需要定义cudaResourceDesc资源描述符和cudaTextureDesc纹理描述符。
分别如下所示：

//资源描述符
cudaResourceDesc resDesc;
memset(&resDesc, 0, sizeof(resDesc));//初始化
resDesc.resType = cudaResourceTypeArray;//指定对应设备内存的形式为 CUDA数组
resDesc.res.array.array = cuArray;//CUDA数组 对应的赋值形式

//纹理描述符
cudaTextureDesc texDesc;
memset(&texDesc, 0, sizeof(texDesc));
texDesc.addressMode[0] = cudaAddressModeClamp;
texDesc.addressMode[1] = cudaAddressModeClamp;
texDesc.filterMode = cudaFilterModePoint;//最近邻插值法
texDesc.readMode = cudaReadModeElementType;//若选用cudaFilterModeLinear,则readMode=cudaReadModeNormalizedFloat
texDesc.normalizedCoords = 1;//对坐标进行归一化

接着使用cudaCreateTextureObject()函数创建纹理对象

//创建纹理对象
cudaTextureObject_t tex = 0;
cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL);

访问纹理内存

从CUDA数组取纹理时，使用tex1D()或tex2D()。

template<class Type,enum cudaTextureReadMode readMode>
Type tex1D(texture<Type,1,readMode> texRef,float x);

template<class Type,enum cudaTextureReadMode readMode>
Type tex2D(texture<Type,2,readMode> texRef,float x,float y);

纹理对象销毁

使用cudaDestroyTextureObject();

cudaDestroyTextureObject(tex);

代码示例(CUDA+Opencv)

下面是用opencv读取一张图，然后对其压缩为原来的一半大小。

这里为resize.h文件

#pragma once
#ifndef _RESIZE_H
#define _RESIZE_H

typedef unsigned char ElementType;

typedef struct resize
{
	unsigned int resize_w;
	unsigned int resize_h;
	ElementType *resize_image;

}*P;

#endif // !_RESIZE_H

这里为.cu文件

#include "cuda_runtime.h"
#include "device_launch_parameters.h"

#include "opencv2/core.hpp"
//#include <opencv2/core/utility.hpp>
#include "opencv2/highgui.hpp"
#include "opencv2/imgproc.hpp"

#include<iostream>
#include<stdlib.h>
#include<stdio.h>

#include"resize.h"

using namespace std;
using namespace cv;

int iDivUp(int a, int b)
{
	return (a % b != 0) ? (a / b + 1) : (a / b);
}

__device__ size_t flatten_2d_index(size_t idx, size_t idy, size_t width)
{
	return idx + idy * width;
}

__global__ void resize_kernel(cudaTextureObject_t tex ,unsigned char *resize, unsigned int resize_h, unsigned int resize_w)
{
	size_t ix = threadIdx.x + blockIdx.x*blockDim.x;
	size_t iy = threadIdx.y + blockIdx.y*blockDim.y;
	if (ix < resize_w&&iy < resize_h)
	{
		float idx = ix * (1.0f / float(resize_w - 1));
		float idy = iy * (1.0f / float(resize_h - 1));

		size_t out_id = flatten_2d_index(ix, iy, resize_w);

		resize[out_id] = tex2D<unsigned char>(tex, idx, idy);

	}

}

P CreateResizeimage()
{
	P image = new struct resize;
	image->resize_h = 0;
	image->resize_w = 0;
	image->resize_image = NULL;
	return image;
}

void Init_resize(P image,unsigned int origin_w, unsigned int origin_h)
{
	unsigned int resize_h = origin_h / 2;
	unsigned int resize_w = origin_w / 2;
	image->resize_h = resize_h;//120
	image->resize_w = resize_w;//140
	image->resize_image = new unsigned char[resize_h*resize_w]();
	if (image->resize_image == NULL)
	{
		cout << "space error" << endl;
		exit(1);
	}
	
}

int main()
{
	//用opencv读取图片
	Mat image = imread("C:\\Users\\ASUS\\Desktop\\1.jpg", IMREAD_GRAYSCALE);
	namedWindow("origin_image");
	imshow("origin_image",image);
	//waitKey();

	//获取图片信息，并将数组赋值给Array数组
	unsigned int width = image.cols;
	unsigned int height = image.rows;
	size_t size = sizeof(unsigned char)*width*height;

	unsigned char *Array = new unsigned char[width*height]();

	Array = image.data;

	//cuda数组申请
	cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<unsigned char>();
	cudaArray *cuArray;
	cudaMallocArray(&cuArray, &channelDesc, width, height);
	cudaMemcpyToArray(cuArray, 0, 0, Array, size, cudaMemcpyHostToDevice);

	//资源描述符
	cudaResourceDesc resDesc;
	memset(&resDesc, 0, sizeof(resDesc));//初始化
	resDesc.resType = cudaResourceTypeArray;//指定对应设备内存的形式为 CUDA数组
	resDesc.res.array.array = cuArray;//CUDA数组 对应的赋值形式

	//纹理描述符
	cudaTextureDesc texDesc;
	memset(&texDesc, 0, sizeof(texDesc));
	texDesc.addressMode[0] = cudaAddressModeClamp;
	texDesc.addressMode[1] = cudaAddressModeClamp;
	texDesc.filterMode = cudaFilterModePoint;//最近邻插值法
	texDesc.readMode = cudaReadModeElementType;//若选用cudaFilterModeLinear,则readMode=cudaReadModeNormalizedFloat
	texDesc.normalizedCoords = 1;

	//创建纹理对象
	cudaTextureObject_t tex = 0;
	cudaCreateTextureObject(&tex, &resDesc, &texDesc, NULL);


	//为压缩后的图像申请内存空间
	P resize = CreateResizeimage();
	Init_resize(resize, width, height);
	unsigned int resize_h = resize->resize_h;
	unsigned int resize_w = resize->resize_w;
	unsigned char* d_resize;
	//unsigned char* h_resize = new unsigned char[resize_w*resize_h]();
	cudaMalloc((void**)&d_resize, sizeof(unsigned char)*resize_h*resize_w);

	//启动内核
	dim3 block(8, 8);
	dim3 grid(iDivUp(resize_w, block.x), iDivUp(resize_h, block.y));

	resize_kernel << <grid, block >> > (tex, d_resize, resize_h, resize_w);

	cudaMemcpy(resize->resize_image, d_resize, sizeof(unsigned char)*resize_h*resize_w, cudaMemcpyDeviceToHost);

	//销毁纹理对象和释放设备内存
	cudaDestroyTextureObject(tex);
	cudaFree(d_resize);
	cudaFreeArray(cuArray);

	//opencv显示压缩后的图片
	Mat resize_image = Mat(resize_h, resize_w, CV_8UC1, (resize->resize_image));
	namedWindow("resize_image");
	imshow("resize_image", resize_image);
	waitKey();


	system("pause");

	return 0;
}

参考：

https://blog.csdn.net/u012348774/article/details/80449121