有时候需要使用cuda处理多帧的图像,需要把多帧时间序列的图像传入到显存中,这个时候就可以把保存图像的全局内存绑定到二维纹理内存,核函数通过纹理拾取来访问输入的图像数据。不过二维纹理内存的宽是有限制的:
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<float>();
cudaMallocArray((cudaArray**)&arr_mat_x, &channelDesc, img_size, IIR_N_X);
在以上定义二维cuda数组的代码中,img_size是有限制的,如果超过64K就会出错。所以如果每帧图像的大小超过了64K,则不能使用二维纹理内存,这时候可以使用三维纹理内存。下面将详细说明三维纹理内存的使用。
- 定义三维cuda数组:
cudaArray *arr_mat_x; //定义device端指针
cudaExtent extent; //定义图像的尺寸和帧数结构体
cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc<uchar>(); //定义数据类型
extent.width = col; //创建extent时,在旧版本中宽度w以字节为单位,即必须乘上sizeof(DTYPE),新版本已经不用乘以sizeof(DTYPE),否则反而会出错!文档和函数说明都没有改过来,这里是坑!
extent.height = row; //每帧图像的行数
extent.depth = picnum; //图像的总帧数
cudaMalloc3DArray((cudaArray**)&arr_mat_x, &channelDesc, extent); //创建picnum帧row*col的存储空间
- 将数据从CPU拷贝到cuda数组中:
cudaMemcpy3DParms HostToDev = {0}; //定义数据传输的结构体
HostToDev.dstArray = arr_mat_x; //指定数据传输的目标地址为cuda数组
HostToDev.extent = make_cudaExtent(col, row, 1); //创建extent时,在旧版本中宽度w以字节为单位,即必须乘上sizeof(DTYPE),新版本已经不用乘以sizeof(DTYPE),否则反而会出错!文档和函数说明都没有改过来,这里是坑!
HostToDev.kind = cudaMemcpyHostToDevice; //定义传输方向为CPU到GPU显存
HostToDev.srcPos = make_cudaPos(0, 0, 0); //定义数据传输的源地址的偏移量
for(int i = begin; i < begin+picnum; i++) //拷贝多帧图像到cuda数组
{
HostToDev.srcPtr = make_cudaPitchedPtr((void *)pic[i].data, col*sizeof(uchar), col, row); //指定数据传输的源地址,注意这里的第二个参数需要乘以数据类型所占的字节数
HostToDev.dstPos = make_cudaPos(0, 0, i-begin); //指定目标地址的偏移量,分别为x,y,z地址
cudaMemcpy3D(&HostToDev); //根据以上设置的参数实行拷贝
}
- 绑定纹理内存
texture<uchar, cudaTextureType3D, cudaReadModeElementType> tex_mat; //定义为3D类型的纹理内存
tex_mat.normalized = 0; //不归一化
tex_mat.filterMode = cudaFilterModePoint; //filterMode:滤波模式。仅对绑定 CUDA 数组的纹理有效。当使用浮点型的坐标寻址纹理时,将根据设定返回不同类型的值。设定可以有:cudaFilterModePoint 和 cudaFilterModeLinear。分别表示最近邻插值和线性插值
tex_mat.addressMode[0] = cudaAddressModeClamp; //寻址模式,即如何处理越界的纹理坐标。可设置:cudaAddressModeClamp 和 cudaAddressModeWrap。Clamp 即钳位模式,Wrap 为循环模式。循环模式只支持归一化的纹理坐标
tex_mat.addressMode[1] = cudaAddressModeClamp;
tex_mat.addressMode[2] = cudaAddressModeClamp;
tex_mat.channelDesc = channelDesc; //描述纹理返回值类型,同cuda数组部分的内容
cudaBindTextureToArray(tex_mat, (cudaArray *)arr_mat_x, channelDesc); //绑定纹理内存
- 纹理拾取:
tex3D(tex_mat, (float)(j-1), (float)(i-1), (float)k); //后面三个参数分别是x,y,z坐标