纹理内存(二)
纹理参考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;
}