我的显卡配置��
显卡类型：独立显��
显卡芯片：NVIDIA GeForce GT 520M
显存容量��1GB
显存类型：DDR3 纠错
显存位宽��64bit
流处理器数量��48
DirectX��11
如果显卡槽可以插多个显卡，那么为了提高效率，可以多插几个显卡。在�€个新的环境中编程�€发，必须先了解这个环境的配置，包括处理器、存储等信息，这样编写出来的程序才能具备�€优秀的效率�€�所以本文开篇给出了我的显卡信息��
1 cuda编程�€述：
�€ CUDA是以C语言扩展的方式提供给程序员的，这样程序员也可以定义C语言的函数，称之为kernel。当调用kernel的时候，会由N个CUDA线程并行地执行N次，这与通常C函数只执行一次是很不�€样的。Kernel的定义需要使用_global�€声明标志符来声明，同时还要使��<<<��>>>语法来说明需要的CUDA线程数�€�
�€个简单的kernel定义的语法：
�€�€_global_voidvecAdd(fIoat*A,float*B,float*C){}
�€�€int main(){vecAdd<<<1，N>>>(A,B,C);}
�€�€每个执行kernel的线程都给予了一个独特的线程lD，在kernel中可以�€�过使用内置的threadldx变量来访问这个线程ID，作为说明，下面的范例代码把两个矢量A��8相加，并把结果保存到矢量C中：
�€�€_global_voidvecAdd(fIoat*A,fIoat*B,fIoat*C)
�€�€{
int i=threadldx.x��
C[i]=A[i]+B[i];
}
�€�€int mainO{vecAdd<<>>(A,B,C);}
�€�€每个执行vecAdd()的线程做�€个元素的加法��
CUDA表示ComputeUnifiedDeviceArchitecture(统一计算设备架构)。是NVIDIA为自家的GPu编写的一套编译器及相关的库文件�€�使�–UDA可以发挥显卡的大规模并行计算能力。CUDA采用C语言作为编程语言提供大量的高性能计算指令�€发能力，使开发�€�能够在GPU的强大计算能力的基础上建立起�€种效率更高的密集教据计算解决方案��
线程的索引号与其的线程ID以一种直观的方式联系在一起：对于�€维的块，他们是一样的；对于二维的��(Dx，Dy)，下标为(X，Y)的线程其线程ID��(x+yDx)：对于三维的��(Dx，Dy，Dz)，下标为(x，y，z)的线程其线程ID��(x+yDx+zDxDy)��
2 cuda线程体系结构
�€�€threadldx是一个三维的向量，这样，线程就可以方便地使用�€维�€�二维或者三维的下标来表示�€�这对于�€维向量，二维矩阵，或者三维数组的计算来说，是很方便的。在下面的例子中，把两个NXN的矩阵A和B相加，并把结果保存到矩阵C中：
�€�€_global_voidmatAdd(fIoatA[N][N],floatB[N][N],floatC[N][N])
�€�€{
�€�€int i=threadldx.x;
�€�€int j=threadldx.y;
�€�€C[i][j]=A[i][j]+B[i][j];
�€�€}
�€�€int main()
�€�€{
�€�€dim3dimBIock(N,N);
�€�€matAdd<<<1，dimBIock>>>(A,B,C);
�€�€}
�€�€线程的索引号与其的线程ID以一种直观的方式联系在一起：对于�€维的块，他们是一样的；对于二维的��(Dx，Dy)，下标为(X，Y)的线程其线程ID��(x+yDx)：对于三维的��(Dx，Dy，Dz)，下标为(x，y，z)的线程其线程ID��(x+yDx+zDxDy)��
�€�€在同�€个block中的thread通过“共享存储器”共享数据来相互协作，�€�过合理安排对内存的访问来同步�€�程序员可以明�‘的指定kernel中的同步点，则可以�€�过调用内置函数_syncthreadsO来实现；_syncthreadso就像�€个屏障，在block中的�€有thread必须在这里等待�€�
�€�€为了高效的相互协作，“共享存储器”是低延迟的，并且接近处理器核心，很类似与CPU中的�€级缓存�€�一syncthreadsO是轻量级的，block中的�€有thread都在同一个处理核心上运行。因此每个block的thread数目受到处理核心�€拥有的存储器资源的限制，在NVIDIATesla架构中，�€个线程block�€多可以包��512个线程�€�
�€�€然�€�，�€个kernel可以被多个线程block执行，因为thread的�€�数等于每block的thread数乘以block数�€�这许多个block组织成一维或者二维线程块，称之为grid�� grid的维度是��<<<��>>>中的第�€�个参数来指定的。每个grid中的block可以通过�€维或者二维的下标来标志，这可以�€�过访问内建的blockldx变量。线程block的维度在kernel中是可以通过内建的blockDim变量来访问，这样，前面的例子可以改写为：
�€�€__global__ void matAddl(floatA[N][N],floatB[N][N],floatC[N][N])
�€�€{
�€�€inti=blockldx.x*blockDimx+threadldx.x;
�€�€intj=blockldx*y*blockDim.y+threadldx.y;
�€�€if(i
�€�€C[i][j]=A[i][j]+B[i][j];
�€�€}
�€�€int main()
�€�€{
�€�€dim3dimBIock(16,16);
�€�€dim3dimGrid(N+dimBlock.x 1)/dimBlock.x,
�€�€(N+dimBlock.y 1)/dimBlock.y);
�€�€matAdd<<>>>(A,B,C);
�€�€}
�€�€线程block的大��1 6×1 6=256个线程是任意选取的�€�与前面�€样，我们创建了一个grid，grid中包含了足够多的block，使得每个线程处理一个矩阵元素�€�
�€�€线程block必须能够独立执行，它们必须可以以任意的持续执行，既可以并行也可以串行执行。只有这样，线程block才可以在任意数量的核心上以任意次序调度运行，从�€�支持程序员写出具有申缩性的代码��
3 cuda文件系统
首先创建.cu文件，这是cuda的相当于c++的cpp文件，�€�.cuh相信你可以领悟�€� .cuh文件类似于C/C++中的.h头文件，你可以将核函数写在里面，然后��.cu 文件中包含该.cuh文件后就可以调用核函数了��
4 API
cuda api分两个级别：running 和device。device是对running的一层封装�€�
4.1 通用函数
时间函数 clock()
同步函数 __syncthreads()
4.2 设备管理函数
runing api : cudaGetDeviceCount() �� cudaGetDeviceProperties()
device api : cuDeviceGetCount()和cuDeviceGet()
4.3 内存管理
cudaMalloc() 或�€� cudaMallocPitch() 来分配线性内存，通过cudaFree()释放内存.
4.5 runing api
cuda包的常用头文件如下：
#include "cuda_runtime.h"
4.6 device api
#include "device_launch_parameters.h"
5 函数限定��
__device__ 在设备上执行，仅可�€�过设备调用��
__global__ 在设备上执行，仅可�€�过主机调用��
__host__ 在主机上执行，仅可�€�过主机调用��
_device_ �� _global_ 函数不支持�€�归��
_device_ �� _global_ 函数的函数体内无法声明静态变量�€�
_device_ �� _global_ 函数不得有数量可变的参数��
_device_ 函数的地�€无法获取，但支持 _global_ 函数的函数指针�€�
_global_ �� _host_ 限定符无法一起使用�€�
_global_ 函数的返回类型必须为空�€�
�� _global_ 函数的任何调用都必须按第 4.2.3 节介绍的方法指定其执行配置�€�
_global_ 函数的调用是异步的，也就是说它会在设备执行完成之前返回�€�
_global_ 函数参数将同时�€�过共享存储器传递给设备，且限制�� 256 字节��
6 变量限定��
_device_ 限定符声明位于设备上的变量�€�位于全�€存储器空间中；与应用程序具有相同的生命周期；可�€�过网格内的�€有线程访问，也可通过运行时库从主机访��
_constant_ 限定符可选择�� _device_ 限定符一起使用，位于固定存储器空间中；与应用程序具有相同的生命周期；可�€�过网格内的�€有线程访问，也可通过运行时库从主机访问�€�
_shared_ 限定符可选择�� _device_ 限定符一起使用，�€声明的变量具有以下特征：位于线程块的共享存储器空间中；与块具有相同的生命周期；尽可�€�过块内的所有线程访问�€�
知道这些基本的知识，那么，我们可以写出cuda的hello world 程序了�€�
