nVIDIA标准计算体系结构（编程指南）

   一、
简介
1.1    在数据上并行的图形处理单元 擅长计算特别是并行计算，源于硬件设计上更多的偏重于数据处理，而不是像CPU一样既重视数据缓存还要兼顾流水线控制。    

高算术强度（算术强度=算术操作次数/存储单元操作次数）。
    数据级并行处理映射数据元素到并行处理线程，以解决大规模的数据运算

一些小技巧： 
1.2    CUDA-一个新的GPU计算体系结构 CUDA是一个用于数据并行计算的新的体系结构，无需再映射原始的图形接口了，它支持GeForce8800系列及更高配置的GPU。
二、 编程模式
2.1 用作高级多线程协处理器的GPU 在CUDA模式下，GPU被用作一个可以并行处理大量线程的计算设备。

2.2 线程分批 应用程序的数据不相关一个核心（kernel）被组织在一个由线程块（thread block）组成的线程格（grid）内。划分应用程序为不相关的核心由用户程序员完成。

2.2.1 线程块 一组线程共享数据并且同步它们的执行来协调访存，同步点。 线程块内线程ID的编号与计算：（x，y）=x+y*Dx，（x，y，z）=x+y*Dx+z* Dx*Dy 2.2.2 线程格 线程块内的线程数目是有限制的，但是在相同核心内执行的线程块可以放到一起组成一个线程格，这样在单个核心上运行的线程数就会大大增加。但是这是以减少线程间的合作为代价的（格内的不同块间的线程是不能相互同步和通信的）。 线程格内线程块ID的编号计算：（x，y）=x+y*Dx 2.3 访存模式（6种）  线程级别R/W寄存器、线程级别R/W本地内存、线程块级别R/W共享内存、线程格级别R/W全局内存、线程格级别只读固定内存、线程格级别只读纹理内存。 三、硬件实现 3.1 SIMD多处理器和片上显存 每一个时钟周期在每个多处理器上执行一条指令，而多处理器上的每一个处理器用这一指令处理不同的数据（即SIMD）。     每一个多处理器含有四种不同的显存结构： 寄存器，多处理器中的每个处理器对应有一组32位寄存器； 共享显存，多处理器中所有处理器共享一个的显存； 常量缓存，利用固定的一部分设备显存，为所有处理器共享； 纹理缓存，利用固定的一部分设备显存，为所有处理器共享； 本地和全局的显存用设备（显卡）内存的一部分实现，并且不带缓存。 3.2 执行模式     每组SIMD线程称作一个warps，每个warps含有相同数目（warp size）的线程，并被多处理器以SIMD的方式执行，线程调度在warps之间切换。

四、API
4.1        C语言扩展
4.2        语言扩展 4
.2.1    函数类型限定词 __device__ ，在显卡上执行，只被设备调用； __global__ ,在显卡上执行，只被主机调用； __host__，在主机上执行，只被主机调用，如不加这一限定词，默认为__host__，也可与__device__合并使用，表示既可在主机也可在设备上执行。 一些限制：     __device__和__global__的函数不支持递归，并且函数体内不可定义静态变量，不含可变参数；__device__的函数不可取地址，使用函数指针来指向__global__修饰的函数，但反之可以。__host__和__global__不可同时使用。 __global__修饰的函数必须返回空类型，任何对__global__修饰函数的调用，必须先进行运行配置；对__global__修饰函数的同步调用，意味着可以被阻塞，__global__修饰函数的参数传递通过共享显存到设备，占用显存大小限制为256个字节。

4.2.2    变量类型限定词 __device__ ,__constant__ ,__shared__，如无任何限定词，则变量使用全局内存空间，并且任何线程包括主存中的都能访问它； 4.2.3    运行配置 如调用的函数的限定词是__global__，则必须进行运行配置。 以<<>>的形式，置于函数名和参数列表之间，

4.2.4    内部变量 gridDim ,blockIdx ,blockDim ,threadIdx; 4.2.5    编译器NVCC NVCC的基本作用是将设备代码（device code）和主机代码（host code）分开，并将设备代码编译成二进制代码。主机代码则留给其他的工具来完成。

4.3        公共运行时组件  
4.3.1 内部向量类型：                      char1, uchar1, char2, uchar2, char3, uchar3, char4, uchar4, short1, short1, short2, ushort2, short3, ushort3, short4, ushort4, int1, uint1, int2, uint2, int3, uint3, int4, uint4, long1, ulong1, long2, ulong2, long3, ulong3, long4, ulong4, float1, float2, float3, float4             这些向量类型是从基本的整型和浮点型演化而来的，可以通过一个基本的构造函数产生：                 typen make_typen(x ,y ,z ,w)；其中，typen表示上面提到的类型，x,y,z,w对应于n值取其中的1-4个，比如：int2 make_int2(int x,int y);就得到了一个二维整型向量，取值为(x,y)。             dim3               一个基于uint3的特定的单位，其初始值为(1,1,1)。   4.3.2 数学函数：

  4.3.3 时间函数：clock_t clock()  

4.3.4 纹理类型：      纹理内存只能通过纹理引用访问，纹理引用是内存中的一段特殊的区域，称为“纹理”，并且需要定义特殊的访问模式。纹理是有维度的，一维或者二维，其中的元素称为纹素，通过纹理引用从输入的纹理坐标读取数据的过程称为纹理获取(texture fetching)。 全局范围的纹理引用定义如下： Texture texRef; 这里面，Type指的是纹理获取时得到的数据类型，可取为4.3.1里提到的类型；Dim指的是纹理引用的维度，取值为1或2，默认值是1；ReadMode可取值有： cudaReadModeNormalizeFloat、cudaReadModeElementType。      纹理类型是一个结构体，它有下面几个域：      channelDesc描述了获取纹理时返回值的格式，本身的类型定义如下：          struct  cudaChannelFormatDesc{              int x, y, z, w;              enum cudaChannelFormatKind f; };f的取值有：cudaChannelFormatKindSigned、cudaChannelFormatKindUnsigned、cudaChannelFormatKindFloat.      normalized      addressMode      filterMode指定了获取的纹理值是怎样被返回的，可取的值有：   cudaFilterModePoint和cudaFilterModeLinear，如果是前者，返回值是最接近输入纹理的纹素坐标；后者则返回的是最接近输入纹理的两次或四次现行插值，而且后者只对浮点类型的数有效。      所有这些域，除了channelDesc以外，都可以在主机代码中指定。 4.4        设备运行时组件 4.5          主机运行时组件 五、G80系列的技术指标 5.1 一般性指标 每个线程块内允许最大512个线程 线城格的每个维上允许最大65535个线程块 多处理器是： GTX有16个时钟频率为675MHZ的多处理器 GTS有12个时钟频率为600MHZ的多处理器 设备显存：     GTX为768MB，GTS为640MB 每一个多处理器上的共享显存大小为16KB，并且被细分成16段； 每一个多处理器拥有64KB的常量显存，其中的8KB用来作cache工作集； 一个多处理器上的一维纹理显存的cache工作集大小为8KB； Warp（经线）大小为32个线程； 纹理过滤权值保存为一个9位的定点数格式，这个数有8位是小数值； 在公测0.8版（Beta version 0.8）中，系统内存和设备显存之间的最大观测带宽为2GB/s； 每一个多处理器上载有8个以两倍于时钟速度运行的处理器，因而能在两个时钟周期内处理一个含有32个线程的warp（经线）。

5.2 浮点运算标准 G80系列遵循IEEE-754单精度二进制浮点数运算标准，但是也有下面一些例外：     加法和乘法常被合并为一条加乘指令；     除法的实现使用的是非标准的倒数方法；
    平方根的实现也是使用非标准的反平方根方法；     对于加法和乘法，只提供通过静态舍入方法实现的四舍五入到最近的偶数和四舍五入到零方法，不支持直接对正/负无穷大的四舍五入；
    不提供动态舍入处理；
    异常的源操作数被当成零参与运算；     向下溢出的结果变为零；
    不提供检测浮点运算异常的机制，但这种异常可以被屏蔽，并且对异常的屏蔽遵循IEEE的标准；     不支持无效的信号；
    一个或几个NaN的数参与运算的结果不再是普通的NaN，而是一个标准的代表NaN的数（0x7fffffff）
。另外根据IEEE-754R标准，如果求极大极小值函数min（），max（）的输入参数中有一个为NaN，那计算结果就是那个非NaN的参数。     在浮点数到整数的转换时，如出现浮点数值超出表示范围，IEEE-754没有定义，而在G80中，将依照表示范围作截取处理，这种处理方式也不同于x86架构。 六、性能指南 6.1 指令性能     为一个warp执行一条指令，多处理器需做以下工作：        读每一个线程的操作数，执行指令，为每个线程写回结果。     因而指令性能依赖于指令吞吐量和访存等待时间与带宽。

6．1.1 指令吞吐量     算术指令     流程控制指令     存储指令 多处理器需要2个时钟周期来处理一个warp的存储指令，当访问到主机内存时，额外需要200~300个时钟周期的访存等待时间。     同步指令        __syncthreads指令在没有线程等待时，一个warp需要2个时钟周期； 6.1.2 存储器带宽     设备显存的带宽显然没有主存那么大，因而因该尽量减少显存的访问，典型的编程模式是将数据从显存中提取出来，放到共享显存里，以供程序运行。对每一个线程，经历如下几个步骤：     将数据从设备显存取到共享显存中；     与线程块中的其他线程同步，以便每一个线程都能安全的读到正确的数据；     对共享显存里的数据进行处理；     对共享显存里的数据进行同步更新；     将结果写回设备显存。

来自 “ ITPUB博客 ” ，链接：http://blog.itpub.net/14741601/viewspace-371996/，如需转载，请注明出处，否则将追究法律责任。

转载于:http://blog.itpub.net/14741601/viewspace-371996/