cuda学习重点笔记专业资料.doc

文章目录1. CUDA是什么2. 64位Ubuntu12.04安装CUDA5.53. 对CUDA C个人懵懂感觉4. 重要概念与名称4.1. 主机4.2. 设备4.3. 线程(Thread)4.4. 线程块(Block)4.5. 线程格(Grid)4.6. 线程束4.7. 核函数（Kernel）4.8. dim3构造类型5. 函数修饰符6. 惯用GPU内存函数6.1. cudaMalloc()6.2. cudaMemcpy()6.3. cudaFree()7. GPU内存分类7.1. 全局内存7.2. 共享内存7.3. 常量内存7.4. 纹理内存7.5. 固定内存8. 原子性9. 惯用线程操作函数10. 使用事件来测量性能11. 流12. 技巧 CUDA是什么CUDA，Compute Unified Device Architecture简称，是由NVIDIA公司创立基于她们公司生产图形解决器GPUs（Graphics Processing Units,可以通俗理解为显卡）一种并行计算平台和编程模型。通过CUDA，GPUs可以很以便地被用来进行通用计算（有点像在CPU中进行数值计算等等）。在没有CUDA之前，GPUs普通只用来进行图形渲染（如通过OpenGL，DirectX）。开发人员可以通过调用CUDAAPI，来进行并行编程，达到高性能计算目。NVIDIA公司为了吸引更多开发人员，对CUDA进行了编程语言扩展，如CUDA C/C+,CUDA Fortran语言。注意CUDA C/C+可以看作一种新编程语言，由于NVIDIA配备了相应编译器nvcc，CUDA Fortran同样。更多信息可以参照文献。对CUDA C个人懵懂感觉如果粗暴以为C语言工作对象是CPU和内存条（接下来,称为主机内存），那么CUDA C工作对象就是GPU及GPU上内存（接下来,称为设备内存），且充分运用了GPU多核优势及减少了并行编程难度。普通通过C语言把数据从外界读入，再分派数据，给CUDA C，以便在GPU上计算，然后再把计算成果返回给C语言，以便进一步工作，如进一步解决及显示，或重复此过程。重要概念与名称主机将CPU及系统内存（内存条）称为主机。设备将GPU及GPU自身显示内存称为设备。线程(Thread)普通通过GPU一种核进行解决。（可以表达到一维，二维，三维，详细下面再细说）。线程块(Block)1. 由各种线程构成（可以表达到一维，二维，三维，详细下面再细说）。2. 各block是并行执行，block间无法通信，也没有执行顺序。3. 注意线程块数量限制为不超过65535（硬件限制）。线程格(Grid)由各种线程块构成（可以表达到一维，二维，三维，详细下面再细说）。线程束在CUDA架构中，线程束是指一种包括32个线程集合，这个线程集合被“编织在一起”并且“步调一致”形式执行。在程序中每一行，线程束中每个线程都将在不同数据上执行相似命令。核函数（Kernel）1. 在GPU上执行函数普通称为核函数。2. 普通通过标记符_global_修饰，调用通过，用于阐明内核函数中线程数量，以及线程是如何组织。3. 以线程格（Grid）形式组织，每个线程格由若干个线程块（block）构成，而每个线程块又由若干个线程（thread）构成。4. 是以block为单位执行。5. 叧能在主机端代码中调用。6. 调用时必要声明内核函数执行参数。7. 在编程时，必要先为kernel函数中用到数组或变量分派好足够空间，再调用kernel函数，否则在GPU计算时会发生错误，例如越界或报错，甚至导致蓝屏和死机。/* * file_name HelloWorld.cu 后缀名称.cu */#include #include /头文献/核函数声明，前面核心字_global_global_ void kernel( void ) int main( void ) /核函数调用，注意，第一种1，代表线程格里只有一种线程块；第二个1，代表一种线程块里只有一种线程。 kernel(); printf( Hello，World!n ); return 0;dim3构造类型1. dim3是基亍uint3定义矢量类型，相称亍由3个unsigned int型构成构造体。uint3类型有三个数据成员unsigned int x；unsigned int y；unsigned int z;2. 可使用亍一维、二维或三维索引来标记线程，构成一维、二维或三维线程块。3. dim3构造类型变量用在核函数调用中。4. 有关几种内置变量4.1. threadIdx，顾名思义获取线程threadID索引；如果线程是一维那么就取threadIdx.x，二维还可以多取到一种值threadIdx.y，以此类推到三维threadIdx.z。4.2. blockIdx，线程块ID索引；同样有blockIdx.x，blockIdx.y，blockIdx.z。4.3. blockDim，线程块维度，同样有blockDim.x，blockDim.y，blockDim.z。4.4. gridDim，线程格维度，同样有gridDim.x，gridDim.y，gridDim.z。5. 对于一维block，线程threadID=threadIdx.x。6. 对于大小为（blockDim.x，blockDim.y） 二维 block，线程threadID=threadIdx.x+threadIdx.y*blockDim.x。7. 对于大小为（blockDim.x，blockDim.y，blockDim.z） 三维 block，线程threadID=threadIdx.x+threadIdx.y*blockDim.x+threadIdx.z*blockDim.x*blockDim.y。8. 对于计算线程索引偏移增量为已启动线程总数。如stride = blockDim.x * gridDim.x；threadId += stride。函数修饰符1. _global_，表白被修饰函数在设备上执行，但在主机上调用。2. _device_，表白被修饰函数在设备上执行，但只能在其她_device_函数或者_global_函数中调用。惯用GPU内存函数cudaMalloc()1. 函数原型： cudaError_t cudaMalloc (void *devPtr，size_t size)。2. 函数用处：与C语言中malloc函数同样，只是此函数在GPU内存为分派内存。3. 注意事项：3.1. 可以将cudaMalloc()分派指针传递给在设备上执行函数；3.2. 可以在设备代码中使用cudaMalloc()分派指针进行设备内存读写操作；3.3. 可以将cudaMalloc()分派指针传递给在主机上执行函数；3.4. 不可以在主机代码中使用cudaMalloc()分派指针进行主机内存读写操作（即不能进行解引用）。cudaMemcpy()1. 函数原型：cudaError_t cudaMemcpy (void *dst，const void *src，size_t count，cudaMemcpyKind kind)。2. 函数作用：与c语言中memcpy函数同样，只是此函数可以在主机内存和GPU内存之间互相拷贝数据。3. 函数参数：cudaMemcpyKind kind表达数据拷贝方向，如果kind赋值为cudaMemcpyDeviceToHost表达数据从设备内存拷贝到主机内存。4. 与C中memcpy()同样，以同步方式执行，即当函数返回时，复制操作就已经完毕了，并且在输出缓冲区中包括了复制进去内容。5. 相应有个异步方式执行函数cudaMemcpyAsync()，这个函数详解请看下面流一节关于内容。cudaFree()1. 函数原型：cudaError_t cudaFree ( void* devPtr )。2. 函数作用：与c语言中free()函数同样，只是此函数释放是cudaMalloc()分派内存。下面实例用于解释上面三个函数#include #include _global_ void add( int a，int b，int *c ) *c = a + b;int main( void ) int c; int *dev_c; /cudaMalloc() cudaMalloc( (void*)&dev_c，sizeof(int) ); /核函数执行 add( 2，7，dev_c )； /cudaMemcpy() cudaMemcpy( &c，dev_c，sizeof(int),cudaMemcpyDeviceToHost ) ; printf( 2 + 7 = %dn，c ); /cudaFree() cudaFree( dev_c ); return 0;GPU内存分类全局内存通俗意义上设备内存。共享内存1. 位置：设备内存。2. 形式：核心字_shared_添加到变量声明中。如_shared_ float cache10。3. 目：对于GPU上启动每个线程块，CUDA C编译器都将创立该共享变量一种副本。线程块中每个线程都共享这块内存，但线程却无法看到也不能修改其她线程块变量副本。这样使得一种线程块中各种线程可以在计算上通信和协作。常量内存1. 位置：设备内存2. 形式：核心字_constant_添加到变量声明中。如_constant_ float s10;。3. 目：为了提高性能。常量内存采用了不同于原则全局内存解决方式。在某些状况下，用常量内存替代全局内存能有效地减少内存带宽。4. 特点：常量内存用于保存在核函数执行期间不会发生变化数据。变量访问限制为只读。NVIDIA硬件提供了64KB常量内存。不再需要cudaMalloc()或者cudaFree(),而是在编译时，静态地分派空间。5. 规定：当咱们需要拷贝数据到常量内存中应当使用cudaMemcpyToSymbol()，而cudaMemcpy()会复制到全局内存。6. 性能提高因素：6.1. 对常量内存单次读操作可以广播到其她“邻近”线程。这将节约15次读取操作。（为什么是15，由于“邻近”指半个线程束，一种线程束包括32个线程集合。）6.2. 常量内存数据将缓存起来，因而对相似地址持续读操作将不会产生额外内存通信量。纹理内存1. 位置：设备内存2. 目：可以减少对内存祈求并提供高效内存带宽。是专门为那些在内存访问模式中存在大量空间局部性图形应用程序设计，意味着一种线程读取位置也许与邻近线程读取位置“非常接近”。如下图：3. 纹理变量（引用）必要声明为文献作用域内全局变量。4. 形式：分为一维纹理内存 和 二维纹理内存。4.1. 一维纹理内存4.1.1. 用texture类型声明，如texture texIn。4.1.2. 通过cudaBindTexture()绑定到纹理内存中。4.1.3. 通过tex1Dfetch()来读取纹理内存中数据。4.1.4. 通过cudaUnbindTexture()取消绑定纹理内存。4.2. 二维纹理内存4.2.1. 用texture类型声明，如texture texIn。4.2.2. 通过cudaBindTexture2D()绑定到纹理内存中。4.2.3. 通过tex2D()来读取纹理内存中数据。4.2.4. 通过cudaUnbindTexture()取消绑定纹理内存。固定内存1. 位置：主机内存。2. 概念：也称为页锁定内存或者不可分页内存，操作系统将不会对这块内存分页并互换到磁盘上，从而保证了该内存始终驻留在物理内存中。因而操作系统可以安全地使某个应用程序访问该内存物理地址，由于这块内存将不会破坏或者重新定位。3. 目：提高访问速度。由于GPU懂得主机内存物理地址，因而可以通过“直接内存访问DMA（Direct Memory Access)技术来在GPU和主机之间复制数据。由于DMA在执行复制时无需CPU介入。因而DMA复制过程中使用固定内存是非常重要。4. 缺陷：使用固定内存，将失去虚拟内存所有功能；系统将更快耗尽内存。5. 建议：对cudaMemcpy()函数调用中源内存或者目的内存，才使用固定内存，并且在不再需要使用它们时及时释放。6. 形式：通过cudaHostAlloc()函数来分派；通过cudaFreeHost()释放。7. 只能以异步方式对固定内存进行复制操作。原子性1. 概念：如果操作执行过程不能分解为更小某些，咱们将满足这种条件限制操作称为原子操作。2. 形式：函数调用，如atomicAdd（addr,y)将生成一种原子操作序列，这个操作序列涉及读取地址addr处值，将y增长到这个值，以及将成果保存回地址addr。惯用线程操作函数同步办法_syncthreads()，这个函数调用，将保证线程块中每个线程都执行完_syscthreads()前面语句后，才会执行下一条语句。使用事件来测量性能1. 用途：为了测量GPU在某个任务上耗费时间。CUDA中事件本质上是一种GPU时间戳。由于事件是直接在GPU上实现。因而不合用于对同步包括设备代码和主机代码混合代码设计。2. 形式：一方面创立一种事件，然后记录事件，再计算两个事件之差，最后销毁事件。如：cudaEvent_t start，stop;cudaEventCreate( &start );cudaEventCreate( &stop );cudaEventRecord( start，0 );/do somethingcudaEventRecord( stop，0 );float elapsedTime;cudaEventElapsedTime( &elapsedTime,start，stop );cudaEventDestroy( start );cudaEventDestroy( stop )；流1. 扯一扯：并发重点在于一种极短时间段内运营各种不同任务；并行重点在于同步运营一种任务。2. 任务并行性：是指并行执行两个或各种不同任务，而不是在大量数据上执行同一种任务。3. 概念：CUDA流表达一种GPU操作队列，并且该队列中操作将以指定顺序执行。咱们可以在流中添加某些操作，如核函数启动，内存复制以及事件启动和结束等。这些操作添加到流顺序也是它们执行顺序。可以将每个流视为GPU上一种任务，并且这些任务可以并行执行。4. 硬件前提：必要是支持设备重叠功能GPU。支持设备重叠功能，即在执行一种核函数同步，还能在设备与主机之间执行复制操作。5. 声明与创立：声明cudaStream_t stream;，创立cudaSteamCreate(&stream);。6. cudaMemcpyAsync()：前面在cudaMemcpy()中提到过，这是一种以异步方式执行函数。在调用cudaMemcpyAsync()时，只是放置一种祈求，表达在流中执行一次内存复制操作，这个流是通过参数stream来指定。当函数返回时，咱们无法保证复制操作与否已经启动，更无法保证它与否已经结束。咱们可以得到保证是，复制操作必定会当下一种被放入流中操作之前执行。传递给此函数主机内存指针必要是通过cudaHostAlloc()分派好内存。（流中规定固定内存）7. 流同步：通过cudaStreamSynchronize()来协调。8. 流销毁：在退出应用程序之前，需要销毁对GPU操作进行排队流，调用cudaStreamDestroy()。9. 针对各种流：9.1. 记得对流进行同步操作。9.2. 将操作放入流队列时，应采用宽度优先方式，而非深度优先方式，换句话说，不是一方面添加第0个流所有操作，再依次添加背面第1，2,个流。而是交替进行添加，例如将a复制操作添加到第0个流中，接着把a复制操作添加到第1个流中，再继续其她类似交替添加行为。9.3. 要牢紧记住操作放入流中队列中顺序影响到CUDA驱动程序调度这些操作和流以及执行方式。技巧1. 当线程块数量为GPU中解决数量2倍时，将达到最优性能。2. 核函数执行第一种计算就是计算输入数据偏移。每个线程起始偏移都是0到线程数量减1之间某个值。然后，对偏移增量为已启动线程总数。