第三期有奖话题：“晒晒”你的CUDA开发经历！{活动结束获奖名单已公布}

bluesnowchen

一)纹理属性

(二)纹理拾取函数

(三)拾取纹理内存与读取全局或常量内存相比的优点





（一）纹理属性

纹理可以在线性内存或是CUDA数组（纹理内存）的任何区域。所以纹理拾取也就对存在与线性内存或CUDA数组中的纹理读取数据。

共用运行组件（既可以运行在host又可以运行在设备）中给出了纹理类型texture。纹理拾取的第一个参数就是纹理参考，纹理参考定义要拾取哪部分纹理内存，它必须通过宿主运行时（只运行在宿主上）函数绑定到一些内存区域（称为纹理(texture)），然后才能供内核使用。下面就来看看纹理参考的不可变属性和可变属性。

不可变属性

纹理参考的声明是texture<Type, Dim, ReadMode> texRef;Type、Dim和ReadMode都是不可变属性。其中：

Type指定拾取纹理时返回的数据类型；Type限制为基本的整数和浮点数类型，以及1-、

2-和4-分量的向量类型之一。

Dim指定纹理参考的维度，等于1、2或3；Dim是可选参数，缺省值为1；

ReadMode等于cudaReadModeNormalizedFloat或

cudaReadModeElementType。如果；ReadMode为cudaReadModeNormalizedFloat，且Type为16-位或8-位整数类型，则其值实际返回值为浮点数类型，即根据原整数类型的全范围进行归一化处理，结果被映射到 [0.0, 1.0]区间（对于无符号整数）或[-1.0, 1.0]为区间（对于有符号整数）；例如，值0xff的无符号8-位纹理元素返回值为1；如果ReadMode为cudaReadModeElementType，则不执行任何转换；ReadMode是可选参数，默认为cudaReadModeElementType.

可变属性

纹理参考的可变属性包括寻址模式、纹理过滤和纹理坐标是否归一化。这些属性都是执行宿主运行时改变的，如cudaBindTextureToArray（）执行时在硬件纹理单元（texture units）时。其中：

addressMode，即如何处理超出范围的纹理坐标。寻址模式是大小为2的数组，数组的第一个和第二个元素分别指定第一个第二个纹理坐标的寻址模式：当寻址模式等于cudaAddressModeClamp时，超出范围的纹理坐标将使用clamp寻址：非归一化纹理坐标的情况下，小于0的值设置为0，大于等于N的值设置为N-1；归一化纹理坐标的情况下，小于0.0或大于1.0的值设置到区间[0.0，1.0）中。当寻址模式是cudaAddressModeWarp时，超出范围的纹理坐标使用warp模式：warp寻址仅使用纹理坐标的小数部分：如1.25当作0.25处理，-1.25当作0.75处理。默认情况下是cudaAddressModeClamp。cudaAddressModeWarp只支持归一化的纹理坐标。

filterMode，过滤模式即当拾取纹理时，如何基于输入纹理坐标来计算返回的值。filterMode等于cudaFilterModePoint或cudaFilterModeLinear；如果它为cudaFilterModePoint，则返回值是纹理坐标最接近输入纹理坐标的纹理元素；如果它为cudaFilterModeLinear，则返回值是纹理坐标最接近输入纹理坐标的2个（对于1D纹理）、4个（对于1个2D纹理）或8给（对于3D纹理）的纹理元素的线性插值。cudaFilterModeLinear只对返回值是浮点类型有效。

normalize，指定纹理坐标是否是归一化的；如果其值非0，则纹理中的所有元素都使用区间[0,1]，而非区间[0,width-1]、[0,height-1]或[0,depth-1]中的纹理坐标来寻址，其中width、height和depth是纹理大小。

在线性内存中分配的纹理：

_维度只能等于1；

_不支持纹理过滤；

_只能使用非归一化的整数纹理坐标寻址；

_寻址模式单一：越界的纹理访问将返回0值。



（二）纹理拾取函数

纹理拾取函数是设备运行时函数。纹理存储的区间不同，拾取的方式也不同。

从线性内存中拾取，使用的函数是tex1Dfetch()的函数簇。如：

template<class Type>

Type tex1Dfetch(

texture<Type, 1, cudaReadModeElementType> texRef, int x);



float tex1Dfetch(

texture<unsigned char, 1, cudaReadModeNormalizedFloat> texRef,

int x);

这些函数用纹理坐标x拾取绑定到纹理参考texRef 的线性内存的区域。不支持纹理过滤和寻址模式。对于整数型，这些函数将会将整数型转化为单精度浮点型。除了这些函数，还支持2-和4-分量向量的拾取。如：

float4 tex1Dfetch(

texture<uchar4, 1, cudaReadModeNormalizedFloat> texRef,

int x);

用纹理坐标x拾取绑定到纹理参考的线性内存的区域。

从CUDA数组中拾取，是用函数tex1D()，tex2D()，tex3D()：

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);



template<class Type, enum cudaTextureReadMode readMode>

Type tex3D(texture<Type, 3, readMode> texRef,float x, float y, float z);

这些函数用用纹理坐标x，y，z拾取绑定到纹理参考 texRef 的CUDA数组。纹理参考的可变和不可变属性决定了如何理解坐标，在纹理拾取和返回值将做怎样的处理。

（三）拾取纹理内存与读取全局或常量内存相比的优点

1.              有高速缓存，如果CUDA数组中的纹理在片上的高速缓存中，则可以潜在的获得较高带宽

2.              不受访问模式的约束。全局或常量内存读取必须遵循相应访存模式才能获得好的性能。如全局内存在单个指令中将32-位、64-位或128-位从全局内存读取到寄存器，单个指令读取的位数要尽量多；另外每个半warp中同时访问全局内存地址的每个线程应该进行排列，以便内存访问可以合并到单个邻近的、对齐的内存访问中。

3.              寻址计算的延迟隐藏的更好，有时候会改善应用程序执行随机访问数据的性能。

4.              打包的数据可以在单个操作中广播到多个独立变量中

5.              8-位和16-位整数输入数据可以有选择地转化为[0.0,1.0]或[-1.0,1.0]区间内的32位浮点值

6.              如果访问的是CUDA数组还有其他的功能，过滤、归一化纹理坐标、寻址模式

renhuiying

Ubuntu 9.04下CUDA编程：
前一阵子把蚁群算法和改进的K-Means算法都搞定了，然后一直在看CUDA编程，前面看CUDA的介绍，一直认为会C之后CUDA就很容易上手，其实不然，还需要了解一些GPU的体系结构相关的知识才能写出好的程序来。《GPU高性能运算之CUDA》这本书看完一遍之后感觉它更像一个手稿整理，把之前的恒多文档整理了一下出了一本书，因为是集大家的智慧，讲的还不错，就是顺序上安排的不是太好。有总比没有好，看过一遍之后，对CUDA编程还是有一些底气的。推荐新手也先看看。

　　看书归看书，写程序是另外一件事情，上一篇文章里把环境搭建起来了，可是我还是不知道怎么创建CUDA工程，怎么动手开始写程序。还好CUDA提供了一个SDK，里面有很多的实例可以供我们参考，于是乎，我的第一个CUDA程序就从这里开始了。

　　CUDA SDK的实例都在src目录下，每一个实例都有一个自己的目录,例如deviceuery，在它的目录下还有一个编译时候使用的Makefile文件，这是编译单个项目的。现在我们将所有实例都编译一遍，在CUDA_SDK根目录下运行sudo make之后，可以在 <CUDA_SDK_HOME>/bin/linux/release下看到编译之后的可执行程序，运行即可看到结果。

那么到这里相信读者应该想到了我们完全可以利用这些实例来创建我们自己的工程。再实例中有一个template，将该目录下src中的.cu、.cpp文件删除，将obj目录下的内容清空，这就成为一个空的CUDA工程，可以再src下编写程序，然后在Makefie中将编译的文件名修改正确，编译即可。所生成的执行文件在CUDA_SDK_HOME/bin/linux/release下。这里是一个测试代码，执行矩阵加法运算的：

　　1 #include <stdio.h>

　　2 #include <stdlib.h>

　　3 #include <time.h>

　　4 #include <cuda_runtime.h>

　　5 #include <cutil.h>

　　6

　　7 #define VEC_SIZE 16

　　8

　　9 //kernel function

　　10 __global__ void vecAdd(float* d_A,float* d_B,float* d_C)

　　11 {

　　12  int index=threadIdx.x;

　　13 d_C[index]=d_A[index]+d_B[index];

　　14 }

　　15

　　16 int main()

　　17 {

　　18 //得到分配空间的大小

　　19 size_t size=VEC_SIZE*sizeof(float);

　　20

　　21 //为本地分配内存

　　22 float* h_A=(float*)malloc(size);

　　23 float* h_B=(float*)malloc(size);

　　24 float* h_C=(float*)malloc(size);

　　25

那么到这里相信读者应该想到了我们完全可以利用这些实例来创建我们自己的工程。再实例中有一个template，将该目录下src中的.cu、.cpp文件删除，将obj目录下的内容清空，这就成为一个空的CUDA工程，可以再src下编写程序，然后在Makefie中将编译的文件名修改正确，编译即可。所生成的执行文件在CUDA_SDK_HOME/bin/linux/release下。这里是一个测试代码，执行矩阵加法运算的：

　　1 #include <stdio.h>

　　2 #include <stdlib.h>

　　3 #include <time.h>

　　4 #include <cuda_runtime.h>

　　5 #include <cutil.h>

　　6

　　7 #define VEC_SIZE 16

　　8

　　9 //kernel function

　　10 __global__ void vecAdd(float* d_A,float* d_B,float* d_C)

　　11 {

　　12  int index=threadIdx.x;

　　13 d_C[index]=d_A[index]+d_B[index];

　　14 }

　　15

　　16 int main()

　　17 {

　　18 //得到分配空间的大小

　　19 size_t size=VEC_SIZE*sizeof(float);

　　20

　　21 //为本地分配内存

　　22 float* h_A=(float*)malloc(size);

　　23 float* h_B=(float*)malloc(size);

　　24 float* h_C=(float*)malloc(size);

　　25

26 //初始化

　　27 for (int i=0;i<VEC_SIZE;++i)

　　28 {

　　29 h_A=1.0;

　　30 h_B=2.0;

　　31 }

　　32

　　33 //将本地内存的中的数据复制到设备中

　　34 float* d_A;

　　35 cudaMalloc((void**)&d_A,size);

　　36 cudaMemcpy(d_A,h_A,size,cudaMemcpyHostToDevice);

　　37

　　38 float* d_B;

　　39 cudaMalloc((void**)&d_B,size);

　　40 cudaMemcpy(d_B,h_B,size,cudaMemcpyHostToDevice);

　　41

　　42 //分配存放结果的空间

　　43 float* d_C;

　　44         cudaMalloc((void**)&d_C,size);

　　45

　　46         //定义16个线程

　　47         dim3 dimblock(16);

　　48         vecAdd<<<1,dimblock>>>(d_A,d_B,d_C);

　　49

　　50         //讲计算结果复制回主存中

　　51         cudaMemcpy(h_C,d_C,size,cudaMemcpyDeviceToHost);

　　52

　　53         //输出计算结果

　　54         for (int j=0;j<VEC_SIZE;++j)

　　55         {

　　56                printf("%f\t",h_C[j]);

　　57         }

　　58

　　59         //释放主机和设备内存

　　60         cudaFree(d_A);

　　61         cudaFree(d_B);

　　62         cudaFree(d_C);

　　63

　　64         free(h_A);

　　65         free(h_B);

　　66         free(h_C);

　　67

　　68         return 0;

　　69 }

chenjintaoivt

这两天测试CUDA编写Jacobi的迭代程序，终于把CUDA编程体系整理清晰了。
    基于GPU编程平台CUDA，是现在热门的领域，但要应用于科学计算，确实是很头痛的事情。虽然NVIDIA开发的CUDA是基于C语言的，只是语法是C语言，但编程思想及通用性跟通用C语言编程，还是差得很远。有很多限制，还不能调试，让人头痛。
    通过两天的测试，发现了很多问题，认识了CUDA编程。没想到以前对CUDA的认识有很大出入，我想主要原因还是CUDA的学习资料中，总是例举矩阵相乘的实例，把CUDA编程的线程处理思维限制在一个很小的框架，而科学计算编程，特别又要考虑通用编程，矩阵相乘的实例是远不够的。
    CUDA程序运行效率很有意思，不是随着矩阵增大计算速度成线性增加，这个优势比CPU强多了。
     现在CUDA计时的函数Clock（）好像有问题，计算时间有时为负，严重违背因果关系啊，O(∩_∩)O哈哈~
     正在调试多Block的并行计算，这个调试通了，以后任何计算程序都可以在GPU上计算了。

ichiaboss

刚接触CUDA编程不久，我认为：
CUDA编程界面简单明了而且非常稳定。程序代码易于向CUDA移植，其灵活性和能力使得程序可以非常快速地运行。
CUDA计算架构可以节省时间和成本
与传统的CPU系统相比，GPU重建可以极大地加快进程的运行速度，而且还可以减少硬件成本。
CUDA计算架构支持C语言编程，C语言也是行业标准。大家都知道如何使用C语言编写程序。

lovebcb

刚刚了解，还没入门，向各位大虾学习。

继续听课！谢谢分析经验。

huangshuojch

怎么感觉楼上的都像是复制的

shuazi0313

顶一下，大家多多学习

xuefushan

CUDA是什么概念都不知道呢。希望有人能在知道一下。

yyfn风辰

CUDA 是计算统一设备架构的简称,是nVidia于2007年6月正式推出的GPU计算环境，它统一了编程模型，引入了共享存储器空间，效率很好。我目前在做分子动力学程序的GPU化。CUDA编程相对简单，而且效率很高。要了解相关的内容可以参考我的博客http://space.itpub.net/23057064/viewspace-623802。

[ 本帖最后由 yyfn风辰 于 2009-12-29 16:40 编辑 ]

superatao

还没有过CUDA开发经历，不过有经历时会考虑研究一下。