CUDA纹理内存相关参数解释
来源:互联网 发布:凶宅数据库 编辑:程序博客网 时间:2024/06/06 05:47
纹理内存的读取模式:
cudaReadModeNormalizedFloat:如果纹理元素是一个16位或8位的整型,那么返回值将作为浮点数返回,
如果是无符号整型,将被映射到[0.0,1.0],如果是有符号整型,将被映射到[-1.0,1.0]
cudaReadModeElementType:那么将返回回原始类型,不做转换
纹理坐标归一化:
默认的,纹理坐标是使用浮点值进行访问:[0,N-1],N是纹理在相应维度上的值,例如,对于一个大小64X32的纹理,其x,y方向的坐标访问范围是:
[0,63],[0,31]
归一化的纹理坐标,将使相应的坐标访问在以下范围:[0,1-1/N],那么对于上面的64x32的纹理,其两个方向的坐标范围就分别是:[0,1-1/64],[0,1-1/32]
滤波模式:
基于输入的坐标值,给出其元素值的返回方式
cudaFilterModePoint:在该模式下,返回相应坐标下,最接近的元素值,类似与最近邻插值
cudaFilterModeLinear:只有当返回值的类型设定为浮点的时候才可以使用,其根据相应元素周围的元素值进行低精度的插值:一维的时候使用线性插值,二维 的时候使用双线性插值,三维的时候使用三线性插值
寻址模式:
寻址模式是一个有3个元素的数组,分别用于指定其相应维度上的寻址方式,用于定义当访问的坐标超出其坐标范围后的处理方式:
cudaAddressModeClamp:是默认的处理方式,对于未归一化的坐标,其截断到范围:[0,N),对于归一化后的坐标范围,截断到:[0.0,1.0)
cudaAddressModeBorder:对于超出范围的坐标,都返回0
cudaAddressModeWrap:只在归一化的坐标下有效,坐标变换公式如下:frac(x) = xfloor(x),其中floor(x)是不大于x的最大整数,
cudaAddressModeMirror:只在归一化坐标下有效,当floor(x)是奇数的时候,每一个坐标变换到frac(x),当floor(x)是偶数的时候,每一个坐标变换到1-frac(x)
举例:
对于一维信号c(k),k的范围是k=0,1,…,M-1
对于cudaAddressModeClamp,就是将超出的值截断为最大最小值
对于cudaAddressModeBorder,就是将其截断为0
对于cudaAddressModeWrap,对于超出范围的值,就是寻找一个整数p(可以是正负、0)满足:c[(k+p*M)/M] = c[k/M],也就是将其变成周期信号,周期为M
对于cudaAddressModeMirror,也是将其变成周期信号,周期为2*M-2,c[1/M]=c[k/M],对于任何1和k,满足:(1+k)mod(2*M-2)=0
下面给出示例代码:
#include <stdio.h>texture<float, 1, cudaReadModeElementType> texture_clamp;texture<float, 1, cudaReadModeElementType> texture_border;texture<float, 1, cudaReadModeElementType> texture_wrap;texture<float, 1, cudaReadModeElementType> texture_mirror;/********************//* CUDA ERROR CHECK *//********************/#define gpuErrchk(ans) { gpuAssert((ans), __FILE__, __LINE__); }inline void gpuAssert(cudaError_t code, char *file, int line, bool abort=true){ if (code != cudaSuccess) { fprintf(stderr,"GPUassert: %s %s %d\n", cudaGetErrorString(code), file, line); if (abort) exit(code); }}/******************************//* CUDA ADDRESS MODE CLAMPING *//******************************/__global__ void Test_texture_clamping(const int M) { printf("Texture clamping - i = %i; value = %f\n", -threadIdx.x, tex1D(texture_clamp, -(float)threadIdx.x)); printf("Texture clamping - i = %i; value = %f\n", M + threadIdx.x, tex1D(texture_clamp, (float)(M + threadIdx.x)));}/****************************//* CUDA ADDRESS MODE BORDER *//****************************/__global__ void Test_texture_border(const int M) { printf("Texture border - i = %i; value = %f\n", -threadIdx.x, tex1D(texture_border, -(float)threadIdx.x)); printf("Texture border - i = %i; value = %f\n", M + threadIdx.x, tex1D(texture_border, (float)(M + threadIdx.x)));}/**************************//* CUDA ADDRESS MODE WRAP *//**************************/__global__ void Test_texture_wrap(const int M) { printf("Texture wrap - i = %i; value = %f\n", -threadIdx.x, tex1D(texture_wrap, -(float)threadIdx.x/(float)M)); printf("Texture wrap - i = %i; value = %f\n", M + threadIdx.x, tex1D(texture_wrap, (float)(M + threadIdx.x)/(float)M));}/****************************//* CUDA ADDRESS MODE MIRROR *//****************************/__global__ void Test_texture_mirror(const int M) { printf("Texture mirror - i = %i; value = %f\n", -threadIdx.x, tex1D(texture_mirror, -(float)threadIdx.x/(float)M)); printf("Texture mirror - i = %i; value = %f\n", M + threadIdx.x, tex1D(texture_mirror, (float)(M + threadIdx.x)/(float)M));}/********//* MAIN *//********/void main(){ const int M = 4; // --- Host side memory allocation and initialization float *h_data = (float*)malloc(M * sizeof(float)); for (int i=0; i<M; i++) h_data[i] = (float)i; // --- Texture clamping cudaArray* d_data_clamping = NULL; gpuErrchk(cudaMallocArray(&d_data_clamping, &texture_clamp.channelDesc, M, 1)); gpuErrchk(cudaMemcpyToArray(d_data_clamping, 0, 0, h_data, M * sizeof(float), cudaMemcpyHostToDevice)); cudaBindTextureToArray(texture_clamp, d_data_clamping); texture_clamp.normalized = false; texture_clamp.addressMode[0] = cudaAddressModeClamp; dim3 dimBlock(2 * M, 1); dim3 dimGrid(1, 1); Test_texture_clamping<<<dimGrid,dimBlock>>>(M); printf("\n\n\n"); // --- Texture border cudaArray* d_data_border = NULL; gpuErrchk(cudaMallocArray(&d_data_border, &texture_border.channelDesc, M, 1)); gpuErrchk(cudaMemcpyToArray(d_data_border, 0, 0, h_data, M * sizeof(float), cudaMemcpyHostToDevice)); cudaBindTextureToArray(texture_border, d_data_border); texture_border.normalized = false; texture_border.addressMode[0] = cudaAddressModeBorder; Test_texture_border<<<dimGrid,dimBlock>>>(M); printf("\n\n\n"); // --- Texture wrap cudaArray* d_data_wrap = NULL; gpuErrchk(cudaMallocArray(&d_data_wrap, &texture_wrap.channelDesc, M, 1)); gpuErrchk(cudaMemcpyToArray(d_data_wrap, 0, 0, h_data, M * sizeof(float), cudaMemcpyHostToDevice)); cudaBindTextureToArray(texture_wrap, d_data_wrap); texture_wrap.normalized = true; texture_wrap.addressMode[0] = cudaAddressModeWrap; Test_texture_wrap<<<dimGrid,dimBlock>>>(M); printf("\n\n\n"); // --- Texture mirror cudaArray* d_data_mirror = NULL; gpuErrchk(cudaMallocArray(&d_data_mirror, &texture_mirror.channelDesc, M, 1)); gpuErrchk(cudaMemcpyToArray(d_data_mirror, 0, 0, h_data, M * sizeof(float), cudaMemcpyHostToDevice)); cudaBindTextureToArray(texture_mirror, d_data_mirror); texture_mirror.normalized = true ; texture_mirror.addressMode[0] = cudaAddressModeMirror; Test_texture_mirror<<<dimGrid,dimBlock>>>(M); printf("\n\n\n");}
最后的结果应该是:
index -7 -6 -5 -4 -3 -2 -1 0 1 2 3 4 5 6 7 8 9 10 11clamp 0 0 0 0 0 0 0 0 1 2 3 3 3 3 3 3 3 3 3border 0 0 0 0 0 0 0 0 1 2 3 0 0 0 0 0 0 0 0wrap 1 2 3 0 1 2 3 0 1 2 3 0 1 2 3 0 1 2 3mirror 1 2 3 3 2 1 0 0 1 2 3 3 2 1 0 0 1 2 3
- CUDA纹理内存相关参数解释
- cuda共享内存,全局内存,纹理等的解释
- CUDA使用纹理内存
- CUDA 纹理内存
- cuda纹理内存
- CUDA 纹理内存
- CUDA 纹理内存
- CUDA 纹理内存
- CUDA使用纹理内存
- cuda纹理内存使用
- CUDA 纹理内存
- CUDA总结:纹理内存
- CUDA纹理内存的使用
- CUDA一维纹理内存
- CUDA一维纹理内存
- CUDA学习之纹理内存
- CUDA:纹理内存及其使用
- CUDA编程——纹理内存
- Python 基础语法(二)
- Android中如何在“自定义布局”中“动态”添加控件
- Razor视图引擎
- nsis安装包_手把手教NIS Edit安装向导的使用
- webService
- CUDA纹理内存相关参数解释
- 2016即将结束,你的目标实现了吗?
- Python 基础语法(三)
- NVIDIA + docker + caffe/digits 配套使用
- HTTP请求和响应格式
- Java过滤html文档
- vitamio视频播放学习
- python 卡方检验原理及应用
- 光流法简单介绍