CUDA |
您所在的位置：网站首页 › cuda例程 › CUDA |
CUDA
|
struct textureReference{
int normalized; //是否归一化 true or false
enum cudaTextureFilterMode filterMode; //滤波模式：cudaFilterModepoint和cudaFilterModeLinear
enum cudaTextureAddressMode addressMode[3]; //寻址模式：cudaAddressModeClump和cudaAddressModeWarp
struct cudaChannelFormatDesc channelDesc;
}
例如：
2、纹理缓存的作用 1）纹理内存中的数据可以被重复利用，当需要的数据存在于纹理缓存中，就不用再去显存读取了 2）纹理拾取可以读取纹理坐标附近的几个象元，提高局部性的访问效率，实现滤波模式。换言之，对于图像滤波具有很好的性能。 3、一维纹理存储器的使用方法 1）在host端声明线性内存空间 分配线性内存空间的方法常见有cudaMalloc(), 由cudaMallocPitch()或者cudaMalloc3D()分配的线性空间是经过填充对齐的线性内存。 2）声明纹理参考系 texture texRef; //Type指定数据类型，特别注意：不支持3元组 //Dim指定纹理参考系的维度，默认为1 //ReadMode可以是cudaReadModelNormalizedFloat或cudaReadModelElementType(默认)注意：type仅限于基本整型、单精度浮点类型和CUDA运行时提供的1元组、2元组和四元祖。 texture texRef;这样就声明了一个一维整数的纹理。 3）绑定数据到纹理 通过cudaBindTexture()函数将纹理参考连接到内存。将1维线性内存绑定到1维纹理。 cudaError_t cudaBindTexture( size_t *offset, const struct textureReference *texref, const void *devPtr,const struct cudaChannelFormatDesc * desc,size_t size = UINT_MAX ) 例如：cudaBindTexture(0, texRef, d_data); 4）设备端的纹理拾取 在kernel中对纹理存储器进行访问，要通过tex1Dfetch()函数,通过纹理参考和给出的坐标位置就可以取得数据了。 type tex1Dfetch(texture texRef, int x); 通过以上四步，我们就可以在程序中利用纹理内存进行数据处理了。例子如下： #include "cuda_runtime.h" #include "device_launch_parameters.h" #include #define N 256 //声明纹理参考系 texture texRef; //CPU串行向量复制 void copyCPU(int* source, int* target, int size) { for (int i = 0; i < size; i++) { target[i] = source[i]; } } //GPU使用全局内存实现向量复制 __global__ void kernel(int* source, int* target, int size) { int index = threadIdx.x + blockDim.x * blockIdx.x; if (index (d_data, d_Result, size); //拷贝数据回主机端 cudaMemcpy(target, d_Result, Size, cudaMemcpyDeviceToHost); cudaFree(d_data); cudaFree(d_Result); } //GPU使用纹理内存实现向量复制 __global__ void kernelTexture(int* target, int size) { int index = threadIdx.x + blockDim.x * blockIdx.x; if (index (d_Result, N); cudaMemcpy(target, d_Result, Size, cudaMemcpyDeviceToHost); //解除纹理绑定 cudaUnbindTexture(texRef); cudaFree(d_data); cudaFree(d_Result); } int main() { int *source, *target; source = (int *)malloc(sizeof(int) * N); target = (int *)malloc(sizeof(int) * N); for (int i = 0; i < N; i++) { source[i] = i + 1; target[i] = 0; } cudaEvent_t t1, t2; cudaEventCreate(&t1); cudaEventCreate(&t2); cudaEventRecord(t1, 0); int number = 2; switch (number) { case 0: copyCPU(source, target, N); break; case 1: copyGPU(source, target, N); break; case 2: copyGPUTexture(source, target, N); break; } cudaEventRecord(t2, 0); cudaEventSynchronize(t2); printf("\nSource data:\n"); for (int i = 0; i < N; i++) { printf("%5d", source[i]); } printf("\nTarget data:\n"); for (int i = 0; i < N; i++) { printf("%5d", target[i]); } printf("\n"); float gpuTime = 0; cudaEventElapsedTime(&gpuTime, t1, t2); printf("GPU Time is：%f\n", gpuTime); return 0; } 4、二维纹理存储器的使用方法1）声明CUDA数组 使用CUDA数组主要通过三个函数使用：cudaMallocArray(), cudaMemcpyToArray(),cudaFreeArray(). 在声明CUDA数组之前，必须先描述结构体cudaChannelFormatDes()的组件数量和数据类型。 结构体定义： struct cudaChannelFormatDesc { int x, y, z, w; enumcudaChannelFormatKind f; };x, y, z和w分别是每个返回值成员的位数，而f是一个枚举变量，可以取一下几个值： cudaChannelFormatKindSigned，如果这些成员是有符号整型 cudaChannelFormatKindUnsigned，如果这些成员是无符号整型 cudaChannelFormatKindFloat，如果这些成员是浮点型CUDA数组的创建方法： cudaChannelFormatDesc channelDesc = cudaCreateChannelDesc();//声明数据类型 cudaArray *cuArray; //分配大小为W*H的CUDA数组 cudaMallocArray(&cuArray, &channelDesc, Width, Height);CUDA 数组的复制： cudaMemcpyToArray(struct cuArray* dstArray, size_t dstX, size_t dstY, const void *src, size_t count, enum cudaMemcpyKind kind); //函数功能是：把数据src复制到CUDA 数组dstArray中，复制的数据字节大小为count，从src的左上角(dstX,dstY)开始复制。 //cudaMemcpyKind 复制方向有：hostTohost, hostTodevice, deviceTohost,device todevice(简写方式).2）声明纹理参考系 //声明一个float 类型的2维的纹理，读取模式为cudaReadModeElementType，也可以声明其他的读取模式 texture texRef;3）绑定CUDA数组到纹理 调用cudaBindTextureToArray()函数把CUDA数组和纹理连接起来。 4）设备端的纹理拾取 和一维纹理内存的tex1Dfetch()不同，需要使用tex1D()、tex2D()、tex3D()这三个函数，分别是用在1D/2D/3D的纹理。
二维纹理内存使步骤和一维区别不大，主要是利用CUDA数组来进行绑定。具体实例如下： #include "cuda_runtime.h" #include "device_launch_parameters.h" #include #define W 16 #define H 16 //声明一个float 类型的2维的纹理，读取模式为cudaReadModeElementType texture texRef; __global__ void addKernel(float *array, const int w, const int h) { int idx = threadIdx.x + blockDim.x * blockIdx.x; int idy = threadIdx.y + blockDim.y * blockIdx.y; if (idx < w && idy (d_data, W ,H); cudaMemcpy(h_result, d_data, W*H * sizeof(float), cudaMemcpyDeviceToHost); //解除绑定 cudaUnbindTexture(texRef); //释放CDUA数组 cudaFreeArray(cuArray); //释放显存 cudaFree(d_data); printf("Origin data:\n"); for (int i = 0; i < W*H; i++) { printf("%8.1f", h_data[i]); } printf("\nResult:\n"); for (int i = 0; i < W*H; i++) { printf("%8.1f", h_result[i]); } return 0; } 参考：http://blog.csdn.net/augusdi/article/details/12187159 《基于CUDA的并行程序设计》刘金硕 |
【本文地址】
今日新闻 |
推荐新闻 |