CUDA系列一：基本概念

       太久没有写博客了，最近一直在整cuda c这块的东西，学了一阵子了但是水平还是一般般。作为一个半入门的人，写点记录一来是可以帮助自己记忆，更多的是希望可以帮助他人，做技术这块写点东西能够帮助到他人，我想这大概是自己最开心的事情，以后尽量多写些有用的东西吧。

1. 基本概念

        常用的操作主要包括函数声明，变量声明，内存类型声明，纹理内存，原子操作等。

• 主机：CPU以及系统内存称为主机
• 设备：GPU以及GPU本身的显存称为设备
• 线程：一般通过GPU的一个核进行处理
• 线程块：

     (1)一个线程块包含多个线程；

     (2)每个block之间是并行执行的，block之间无法通信，因而无执行顺序可言；

     (3)线程块的数量限制为不超过65536；

• 线程格

     由多个线程块组成

                                                 

• 线程束

    在CUDA架构中，线程束是指一个包含32个线程的集合，这个线程集合被编织在一起且执行步调一致。

• 核函数

    (1)在GPU上具体执行操作过程往往是通过核函数来实现的。

    (2)一般通过标识符__global__修饰，通过<<<param1,param2>>>调用，用于说明核函数中的线程数和线程块的组织分布情况。

    (3)以线程格的形式组织，每个线程格由若干个线程块组成，而且每个线程块又由若干个线程组成。

     (4)以block为单元执行。

     (5)能在主机端进行调用。

     (6)调用时必须声明内核函数的执行参数。

     (7)在编程时，必须先为kernel函数中用到的数组或者变量分配好足够的设备端空间，再调用kernel函数，否则GPU计算时会发生错误。

     (8)函数修饰符：

         __global__，表明被修饰的函数在设备上执行，但在主机上调用。

        __device__，表明被修饰的函数在设备上执行，但智能在其他__device__函数或者__global__函数中调用，即设备函数上调用。

2.常用的GPU内存函数

• cudaMalloc()

    (1)函数原型：cudaError_t cudaMalloc(void **devPtr, size_t size)

    (2)函数用处：用于设备端变量分配显存

    (3)注意：

        cudaMalloc()分配的指针传递在设备上执行的函数

        可以在设备代码中使用cudaMalloc()分配的指针进行设备内存读写操作

        可以将cudaMalloc()分配的指针传递给主机上执行的函数

        不可以在主机代码中使用cudaMalloc()分配的指针进行主机内存的读写操作

• cudaMemcpy()

    (1)函数原型：cudaError_t cudaMemcpy(void *dst,const void *src,size_t count,cudaMemcpyKind kind)

    (2)函数作用：与c语言中的memcpy函数一样，该函数可以在主机内存和GPU内存之间互相拷贝数据

•     cudaFree()

    (1)函数原型：cudaError_t cudaFree(void *devPtr)

    (2)函数作用：该函数释放的时cudaMalloc分配的内存

3.GPU内存分类

• 全局内存

     通俗意义上的设备内存

•  共享内存

     (1)位置：设备内存

     (2)形式：关键字__shared__添加到变量声明中，如__shared__ int arr[10]。

     (3)目的：对于GPU上启动的每个线程块，CUDA c编译器都将创建该共享变量的一个副本。线程块中的每个线程都共享这块内存，但线程却无法看到也不能修改其他线程块的变量副本。这样使得一个线程块中的多个线程中的多个线程能够在计算上通信和协作。

• 常量内存

    (1)位置：设备内存

    (2)形式：关键字__constant__添加到变量声明中。如__constant__ int arr[10];

    (3)目的：为了提升性能。常量内存采取了不同于标准全局内存的处理方式。在某些情况下，用常量内存替换全局内存能有效地减少内存带宽。

    (4)常量内存用于保存在该核函数执行期间不会发生变化的数据。变量的访问限制为只读。NVIDIA硬件提供了64KB的常量内存。不再需要cudaMalloc()或者cudaFree(),而是在编译时，静态地分配空间。

    (5)要求：当需要拷贝数据到常量内存中应该使用cudaMemcpyToSymbol(),而cudaMemcpy()会幅值到全局内存。

• 纹理内存

    (1)位置：设备内存

    (2)目的：能够减少对内存的请求并提供有效的内存带宽。专门为哪些内存访问模式中存在大量空间局部性的图形应用程序设计，意味着一个线程读取的位置可能与邻近线程读取的位置非常接近。

                    

     (3)纹理内存必须声明为文件作用域内的全局变量

     (4)形式：

        一维纹理内存：texture<float>texin

        通过cudaBindTexture()绑定到纹理内存中的数据

       通过tex1Dfetch()来读取纹理内存中的数据

       通过cudaUnbindTexture()取消绑定纹理内存

       二维纹理内存：texture<类型，数字>

       通过cudaBindTexture2D()绑定到纹理内存中

      通过tex2D()来读取纹理内存中的数据

      通过cudaUnbingTexture()取消绑定纹理内存

• 原子

    (1)如果操作的执行过程不能分解为更小的部分，则将满足这种条件限制的操作称为原子操作

    (2)形式：函数调用，如atomicAdd(addr,y)将生成一个原子的操作序列，这个操作训练包括读取地址addr的值，将y增加到这个值，以及将结果保存回地址addr。

• 常用线程操作

    同步方法__syncthreads()，这个函数的调用，将确保线程块中的每个线程都执行完。

 

其实这些基本概念也是参考网上其他人写的，结合个人理解觉得写的还不错的，所以重新写了一遍，水平有限，如有不当之处，请指明，谢谢！