【cuda学习笔记】2.纹理对象API的使用，实现sobel边缘检测

纹理对象API

使用cudaCreateTextureObject()函数从一个用来指定纹理并从纹理类型定义而来的资源描述类型结构（resource description of type struct）cudaResourceDesc建立纹理对象。

struct cudaTeextureDesc

{

enum cudaTextureAddressMode addess_mode[3];

enum cudaTextureFilterMode  filter_mode;

enum cudaTextureReadMode    read_mode;

enum cudaTexturefilterMode  mipmap_fileter_mod;

int                         sRGB;

int                         normalized_coords;

unsigned int                max_anisotropy;

float                       mipmap_level_bias;

float                       min_mipmap_level_clamp;

float                       max_mipmap_level_clamp;

};

下面程序是用纹理内存实现的图像sobel边缘检测的代码，作为纹理内存使用的例子：

int main()

{

int width, height;

float *image_buffer;

image_buffer=loadImage(width, height);

cudaChannelFormatDesc channel_format = cudaCreateChannelDesc(32,0,0,0,

cudaChannelFormatKindFloat);

//创建CUDA数组，作为纹理结构的Resource

cudaArray_t de_arr;

CheckErr(cudaMallocArray(&de_arr, &channel_format, width, height),
"malloc cuda arrary");

CheckErr(cudaMemcpyToArray(de_arr, 0, 0, image_buffer,sizeof(float)*width*height,
cudaMemcpyHostToDevice)

, "copy data from host to device arrary");

//创建纹理源Resource Descripte，类型为CUDA数组

struct cudaResourceDesc res_desc;

memset(&res_desc, 0,
sizeof(res_desc));

res_desc.resType =
cudaResourceTypeArray;

res_desc.res.array.array = de_arr;

//创建纹理对象的特征

struct cudaTextureDesc tex_desc;

memset(&tex_desc, 0,
sizeof(tex_desc));

tex_desc.addressMode[0] =
cudaAddressModeWrap;

tex_desc.addressMode[1] =
cudaAddressModeWrap;

tex_desc.filterMode =
cudaFilterModeLinear;

tex_desc.readMode =
cudaReadModeElementType;

tex_desc.normalizedCoords = 1;

cudaTextureObject_t tex_object = 1;

//创建纹理对象

CheckErr(cudaCreateTextureObject(&tex_object, &res_desc, &tex_desc,
NULL),
"create texture object");

float *out_buffer,*out_imag;

out_imag = new float[width*height];

CheckErr(cudaMalloc(&out_buffer,
sizeof(float)*width*height),
"malloc out buffer");

dim3 dimBlock = { 16, 16 };

dim3 dimGrid = { (width + dimBlock.x - 1) / dimBlock.x,

(height + dimBlock.y - 1) / dimBlock.y };

textureTest << <dimGrid,dimBlock >> >(out_buffer, tex_object, width, height);

CheckErr(cudaMemcpy(out_imag, out_buffer,
sizeof(float)*width*height,
cudaMemcpyDeviceToHost), "copy iamge out");

saveImgae(out_imag, width, height);

}

首先创建CUDA数组作为纹理的cudaResourceDesc类型，然后创建纹理对象特征，包括cudaTextureDesc 结构中的各个类型都定义好。然后使用cudaResourceDesc以及cudaTextureDesc 结构创建纹理对象。最后就剩下在kernel中使用纹理对象了。

Kernel函数如下：

__device__ float

ComputeSobel(float ul,
// upper left

float um,
// upper middle

float ur,
// upper right

float ml,
// middle left

float mm,
// middle (unused)

float mr,
// middle right

float ll,
// lower left

float lm,
// lower middle

float lr // lower right

)//这个函数用来计算sobel滤波的值

{

float Horz =
ur + 2.0 *
mr + lr -
ul - 2.0 * ml -
ll;

float Vert =
ul + 2.0 *
um + ur -
ll - 2.0 * lm -
lr;

float Sum = sqrtf(Horz*Horz+Vert*Vert);

return Sum;

}

__global__ void textureTest(float*
image,cudaTextureObject_t texture_obj,
int width,
int height)

{

const int x = threadIdx.x + blockIdx.x*blockDim.x;

const int y = threadIdx.y + blockIdx.y*blockDim.y;

//float u = x / (float)width;

//float v = y / (float)height;

float pix00 = tex2D<float>(texture_obj, (x - 1) / (float)width,
(y - 1) / (float)height);

float pix01 = tex2D<float>(texture_obj, (x + 0) / (float)width,
(y - 1) / (float)height);

float pix02 = tex2D<float>(texture_obj, (x + 1) / (float)width,
(y - 1) / (float)height);

float pix10 = tex2D<float>(texture_obj, (x - 1) / (float)width,
(y + 0) / (float)height);

float pix11 = tex2D<float>(texture_obj, (x + 0) / (float)width,
(y + 0) / (float)height);

float pix12 = tex2D<float>(texture_obj, (x + 1) / (float)width,
(y + 0) / (float)height);

float pix20 = tex2D<float>(texture_obj, (x - 1) / (float)width,
(y + 1) / (float)height);

float pix21 = tex2D<float>(texture_obj, (x + 0) / (float)width,
(y + 1) / (float)height);

float pix22 = tex2D<float>(texture_obj, (x + 1) / (float)width,
(y + 1) / (float)height);

// 从纹理对象读取对应位置的值，使用图像索引来索引纹理，因此处理的就是每个像素点的值

iamge[y *
width + x] = ComputeSobel(pix00, pix01, pix02,

pix10, pix11, pix12,

pix20, pix21, pix22);

}

使用图像的索引（根据线程和线程块的索引得到）作为纹理的坐标，因此纹理就可以获取每个像素点及周围滤波器窗口大小的像素点的值，然后调用__device__函数计算sobel滤波的结果，将结果写到global内存iamge中。tex2D<float>(tex_obj，coorx coory)函数用来获取纹理坐标（coorx，coory）处的值，由于使用的坐标是归一化以后的，因此每个坐标需要除以对应维度上的坐标最大值。