一、CUDA简介

  CUDA是并行计算的平台和类C编程模型,可以实现并行算法。电脑要配备NVIDIA GPU,就可以在许多设备上运行你的并行程序。

 二、CUAD编程

  CUDA编程允许程序执行在异构系统上,即CPU和GPU,并由PCL-Express总线区分开。

        Host:CPU and itsmemory(host memory)

   Device:GPU and its memory(device memory)

  在编程时,一般用h_前缀代表host memory,d_代表device memory

  kernel是CUDA编程的关键,是在GPU上运行的代码,用标识符 __global__注明。

  host可以独立于host进行大部分操作。当一个kernel启动后,控制权会立刻返还给CPU来执行其他额外的任务。所以,CUDA编程是异步的。一个典型的CUD程序包含由并行代码补足的串行代码,串行代码由host执行,并行代码在device中执行,host 端代码是标准C,device是CUDA C代码。我们可以把所有代码放到一个单独的源文件,也可以使用多个文件或库。NVIDIA C编译器(nvcc)可以编译host和device生成可执行程序。

  CUDA程序的处理流程:

    (1)分配host内存,进行数据初始化

         (2)分配device内存,将数据从host上复制到device上

    (3)调用CUDA的核函数在device上完成运算

    (4)将device上的运算结果复制到host上

    (5)释放device和host上分配的内存

三、memory操作

  CUDA程序将系统分为host和device,二者有各自的memory。kernel可以操作device memory,为了可以很好的控制device端内存,CUDA提供了几个内存操作函数:

  

cuda 多个gpu同时工做_二维

 

   CUDA C的风格跟C很接近:

   cudaError_cudaMalloc(void**  devPtr,size_t size)

     

  我们主要看看cudaMencpy,其函数原型为:

        cudaError_t cudaMemcpy ( void* dst, const void* src, size_t count,cudaMemcpyKind kind )

   其中cudaMemcpykijnd的可选类型有:

    1. cudaMemcpyHostToHost

    2. cudaMemcpyHostToDevice

    3.cudaMemcpyDeviceToHost

    4.cudaMemcpyDeviceToDevice

  对于返回类型cudaError_t,如果正确调用,则返回cudaSuccess,否则返回cudaErrorMemoryAllocation,可以使用char* cudaGetErrorString(cudaError_t error)将其转化为易于理解的格式。

四、组织线程

    CUDA线程分为Grid和Block两个层次。

    

cuda 多个gpu同时工做_cuda 多个gpu同时工做_02

  一个单独的kernel启动所有线程组成一个grid, grid中所有线程共享global  memory,  一个grid由许多block组成,block由许多线程组成,grid和block都可以使一维二维或三维,上图是一个二维grid和二维block.

  CUDA内置变量:

      (1)blockIdx: block的索引,blockIdx.x表示block的x坐标。

   (2)threadIdx: 线程索引,同理blockIdx.

   (3)blockDim: block维度,上图blockDim.x=5.

   (4)gridDim: grid维度,同理blockDim.

一般会把grid组织成2D,block为3D。grid 和 block都使用dim3作为声明,例如;

dim3  block(3);

dim3 grid((nElem+block.x-1)/block.x);

 

需要注意的是:dim3仅为host端可见,其对应的device端类型为uint3.

CUDA的线程模型从小往大总结就是:

        (1)Thread : 线程,并行的基本单位。

        (2)Thread Block:线程块,互相合作的线程组,线程块有如下几个特点:

                  允许彼此同步

     可以通过共享内存快速交换数据。

       以1 维、2维、3维组织

       Grid: 一组线程块

       以1维、2维或3维组织

      共享全局内存

         Kernel: 在GPU上执行的核心程序,这个kernel函数是运行在Grid上的。

                      one kernel<-> one Grid

         每一个block和每个thread都有自己的ID,我们通过相应的索引找到相应的线程和线程块。

           threadIdx,  blockIdx

      Block ID:1D or  2D

                     Thread ID: 1D 2D 3D

         GPU上很多并行化的轻量级线程。Kernel在device上执行时实际上是启动很多线程,一个kernel所启动的所有线程称为一个网格(grid),同一个网格上的线程共享相同的全局内存空间。

         grid是线程结构的第一层次,而网格可分为很多线程块(block)

          第二层次:一个线程块里面包含很多线程

    线程两层组织结构:这是一个grid、block均为2-dim的线程组织。grid和block都是定义为dim3类型的变量,dim3可以看成是包含是三个无符号整数(x,y,z)成员的结构体变量,在定义时,缺省值初始化为1.

           grid和block可以灵活地定义为1-dim,2-dim以及3-dim结构,kernel调用时也必须通过执行配置<<<grid,block>>>来指定kernel所使用的网格维度和线程块维度。

    CUDA的这种<<<grid,block>>>其实就是一个多级索引的方法,

                             第一级索引是(grid.xIdx,grid.yIdy), 对应上图例子就是(1,1)

        通过它我们能找到了这个线程块的位置

                             第二级索引是(block.xIdx,block.yIdx,block.zIdx)来定位指定的线程

             这就是CUDA的线程组织结构。

cuda 多个gpu同时工做_多线程_03

 

    每个线程由每个线程处理器(SP)、线程块由多核处理器(SM)执行、一个kernel其实由一个grid来执行,一个kernel一次只能在一个GPU上执行

 

       block是软件概念,一个 block  只会由一个sm调整,程序员在开发时,通过设定block的属性,告诉GPU硬件,我有多少线程,线程怎么组织。

       而具体怎么调整由sm的warp scheduler负责,block一旦被分配好SM,该block就会一直驻留该SM中,直到执行结束。一个SM可以拥有多个Blocks,但需要序列执行。

       下图显示了GPU内部的硬件架构:

       

cuda 多个gpu同时工做_CUDA_04

 

 3. CUDA内存模型

     CUDA中的内存模型分为以下几个层次:

                每个线程都用自己的registers(寄存器)

      每个线程都有自己的local memory(局部内存)

      每个线程块都有自己的shared memory(共享内存),所有线程块内的所有线程共享这                   段内存资源。

      每个grid都有自己的global memory(全局内存),不同线程块的线程都可使用

      每个grid都有自己的constant memory(常量内存)和 texture memory(纹理内存),不                     同线程块的线程都可使用。

 

      线程访问这几类存储器的速度是register > local  memory  >  shared memory >global memory

      计算机架构中的所在层次:

        

cuda 多个gpu同时工做_cuda 多个gpu同时工做_05

 

      4.CUDA编程模型

         CUDA是怎么写程序的了:

         CUDA术语:

         

cuda 多个gpu同时工做_多线程_06

 

        重要:

         1. 通过关键字就可以表示某个程序在CPU上跑还是GPU上跑!

    我们用__global__定义一个kernel函数,就是CPU上调用,GPU上执行,

    注:__global__函数的返回值必须设置为void.

         

cuda 多个gpu同时工做_cuda 多个gpu同时工做_07

 

         2.  CPU和GPU间的数据传输:

            GPU内存分配回收内存的函数接口:

                   cudaMemcpyHostToDevice(CPU到GPU)

      cudaMemcpyDeviceToHost(GPU到CPU)

                   cudaMemcpyDeviceToDevice (GPU到GPU)

         3. 用代码表示线程组织模型?

              可以用dim3类来表示网格和线程块的组织方式,

              网格grid可以 表示为一维和二维格式,

              线程块block可以表示为一维、二维、三维的数据格式。

              dim3   DimGrid(100,50)    //5000个线程块,维度是100*50

              dim3   DimBlock(4,8,8);   //每个线层内包含256个线程,线程块内的维度4*8*8        

              

   4. 计算线程号

               (1)使用N个线程块,每一个线程块只有一个线程,即:

                        dim3  dimGrid (N);

                        dim3 dimBlock (1);

                        此时的线程号的计算方式是:

                         threadId = blockIdx.x;

        其中threadId的范围为0到N-1。对于这种情况,我们可以将其看做是一个列向量,                         列向量中的每一行对应一个线程块。列向量中每一行只有一个元素,对应一个线                              程。

                  (2)使用M*N个线程块,每个线程块1个线程

                           由于线程块是2维的,故可以看做是一个M*N的2维矩阵,其线程号有两个维                                  度, 即:

                           dim3  dimGrid(M,N);

                           dim3   dimBlock(1);

        其中

         blockIdx.x 取值0到M-1

        blockIdx.y取值0到N-1

                          这种情况一般用于处理2维数据结构,比如2维图像,,每一个像素用一个线程来处理处理,此时需要线程号来映射图像像素的对应位置,如:

                          pos= blockIdx.y *  blockDim.x+blockIdx.x;

                     (3)使用一个线程块,该线程具有N个线程

                           dim3 dimGrid(1);

                          dim3 dimblock(N)

                          此时线程号的计算方式是:

                            threadId = threadIdx.x;

                         其中threadId的范围是0到N-1,对于这种情况,可以看做是一个行向量,行向量中的每一个元素对应着一个线程。

                        (4)使用M个线程块,每个线程块内含有N个线程,即:

         dim3 dimGrid(M);

                                dim3 dimBlock(N);

                               这可以想象成二维矩阵,矩阵的行与线程块相对应,矩阵的列与线程编号对应,那线程号的计算方式是:

                              threadId = threadIdx.x + blockIdx*blockDim.x;

              就是把二维的索引空间转换为一维索引空间的过程。

      (5)使用M*N的二维线程块,每一个线程块具有P*Q个线程,即

                              dim3 dimGrid(M, N);

         dim3   dimBlock(P,Q);

        这种适合于处理具有二维数据结构的算法,比如图像处理领域。

                            索引有两个维度:

                            threadId.x = blockIdx.x* blockDim.x+ threadIx.x;

                            threadId.y =  blockIdx.y*blockDim.y+threadIdx.y;

                            上述是把线程和线程块的索引映射为图像像素坐标的计算方法.