CUDA学习笔记

1. CUDA的英文全称是：Compute Unified Device Arcthitecture （统一计算设备架构）。

2. CUDA使用类C语言（对ANSI C进行了最小的必要扩展）开发。

3. CUDA拥有了两个层次的并行：(1)线程级并行实现的细粒度数据并行 (2)任务级并行实现的粗粒度并行。

4. 有时我们说CUDA是指软件，有时又是指硬件，更多时是指CUDA C语言。

    4.1. 从软件的角度上说，CUDA软件栈包含两个层次，一个是驱动层的API,这类函数以cu开头，一个是运行层的API，以cuda开头，运行层API是建立在驱动层API之上的，是对驱动层API的封装，这种封装方便了使用，但是同时也牺牲了灵活性。一般而言，我们的开发都是优先使用运行时API。要提到的是：NVIDIA自身也提供了一些建立在这两者之上的库，供大家使用。

    4.2. 从硬件的角度看，CUDA架构应当包含TPC-->SM-->SP结构及它们的关系。各种存储器及其组织，以及硬件对线程的创建、执行、调度等。

    4.3. 从CUDA C语言的角度看，CUDA的架构包括函数修饰符，变量修饰符，以及一些内置函数等

5. CUDA C 不是 C 语言，而是对 C 语言进行扩展。CUDA 对 C 的扩展主要包括以下四个方面：

    5.1. 函数类型限定符，用来确定函数是在 CPU 还是在 GPU 上执行，以及这个函数是从CPU 调用还是从 GPU 调用。

        5.1.1. __device__ ， 表示从 GPU 上调用，在 GPU 上执行； 也就是说其可以被__global__ 或者__device__修饰的函数调用。此限定符修饰的函数使用有限制，比如在 G80/GT200 架构上不能使用递归，不能使用函数指针等，具体可参见我翻译的编程指南。

        5.1.2. __global__，表示在 CPU 上调用，在 GPU 上执行，也就是所谓的内核(kernel)函数；内核只能够被主机调用，内核并不是一个完整的程序，它只是一个数据并行步骤，其指令流由多个线程执行。

        5.1.3. __host__，表示在 CPU 上调用，在 CPU 上执行，这是默认时的情况，也就是传统的 C 函数。CUDA 支持__host__和__device__的联用，表示同时为主机和设备编译。

    5.2. 变量类型限定符，用来规定变量存储什么位置上。在传统的 CPU 程序上，这个任务由编译器承担。在 CUDA 中，不仅要使用主机端的内存，还要使用设备端的显存和GPU 片上的寄存器、共享存储器和缓存。在 CUDA 存储器模型中，一共抽象出来了8种不同的存储器。复杂的存储器模型使得必须要使用限定符要说明变量的存储位置。

        5.2.1. __device__，__device__表明声明的数据存放在显存中，所有的线程都可以访问，而且主机也可以通过运行时库访问；

        5.2.2. __shared__，__shared__表示数据存放在共享存储器在，只有在所在的块内的线程可以访问，其它块内的线程不能访问；

        5.2.3. __constant__，__constant__表明数据存放在常量存储器中，可以被所有的线程访问，也可以被主机通过运行时库访问；

        5.2.4. Texture，texture 表明其绑定的数据可以被纹理缓存加速存取，其实数据本身的存放位置并没有改变，纹理是来源于图形学的一介概念，CUDA 使用它的原因一部分在于支持图形处理，另一方面也可以利用它的一些特殊功能。

        5.2.5. 如果在 GPU 上执行的函数内部的变量没有限定符，那表示它存放在寄存器或者本地存储器中，在寄存器中的数据只归线程所有，其它线程不可见。

        5.2.6. 如果 SM 的寄存器用完，那么编译器就会将本应放到寄存器中的变量放到本地存储器中。

    5.3. 执行配置运算符<<< >>>，用来传递内核函数的执行参数。执行配置有四个参数，第一个参数声明网格的大小，第二个参数声明块的大小，第三个参数声明动态分配的共享存储器大小，默认为 0，最后一个参数声明执行的流，默认为 0。

    5.4. 五个内建变量，用于在运行时获得网格和块的尺寸及线程索引等信息

        5.4.1. gridDim,  gridDim 是一个包含三个元素 x,y,z 的结构体，分别表示网格在x,y,z 三个方向上的尺寸，虽然其有三维，但是目前只能使用二维；

        5.4.2. blockDim, blockDim 也是一个包含三个元素 x,y,z 的结构体，分别表示块在x,y,z 三个方向上的尺寸，对应于执行配置中的第一个参数，对应于执行配置的第二个参数；

        5.4.3. blockIdx, blockIdx 也是一个包含三个元素 x,y,z 的结构体，分别表示当前线程所在块在网格中 x,y,z 三个方向上的索引；

        5.4.4. threadIdx, threadIdx 也是一个包含三个元素 x,y,z 的结构体，分别表示当前线程在其所在块中 x,y,z 三个方向上的索引；

        5.4.5. warpSize，warpSize 表明 warp 的尺寸，在计算能力为 1.0 的设备中，这个值是24，在 1.0 以上的设备中，这个值是 32。

    其它的还有数学函数，原子函数，纹理读取、绑定函数，内建栅栏，内存 fence 函数等。一般而言，知道这些就应该能够写出 CUDA 程序了，当然要写好的话，必须知道很多其它的细节。