cuda实践2

内存分配：

cudaMemcpy 进行主机与设备端的数据内存交换。

CUDA程序的处理流程：

1. 从CPU拷贝数据到GPU。
2. 调用kernel来操作存储在GPU的数据。
3. 将操作结果从GPU拷贝至CPU。

CUDA线程分成Grid和Block两个层次：

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

__global__ 、__device__、__host__  修饰：

__global__     函数运行在device上，可被host和device函数调用，必须返回void；

__device__    函数运行在device上，只可被device函数调用；

 __host__       函数运行在host上，只可被host函数调用；

 

线程索引：

在kernel里，线程的唯一索引非常有用，为了确定一个线程的索引，我们以2D为例：

• 线程和block索引
• 矩阵中元素坐标
• 线性global memory 的偏移

首先可以将thread和block索引映射到矩阵坐标：

ix = threadIdx.x + blockIdx.x * blockDim.x

iy = threadIdx.y + blockIdx.y * blockDim.y

 

 询和管理GPU device：

查询所有关于GPU device 的信息  

cudaError_t cudaGetDeviceProperties(cudaDeviceProp *prop, int device);

使用nvidia-smi来查询GPU信息

nvidia-smi是一个命令行工具，可以帮助你管理操作GPU device，并且允许你查询和更改device状态。

nvidia-smi用处很多，比如，下面的指令：

$ nvidia-smi -L

GPU 0: Tesla M2070 (UUID: GPU-68df8aec-e85c-9934-2b81-0c9e689a43a7)

GPU 1: Tesla M2070 (UUID: GPU-382f23c1-5160-01e2-3291-ff9628930b70)

然后可以使用下面的命令来查询GPU 0 的详细信息：

$nvidia-smi –q –i 0

 

GPU架构：

　　GPU中每个SM都设计成支持数以百计的线程并行执行，并且每个GPU都包含了很多的SM，所以GPU支持成百上千的线程并行执行，当一个kernel启动后，thread会被分配到这些SM中执行。大量的thread可能会被分配到不同的SM，但是同一个block中的thread必然在同一个SM中并行执行。

　　CUDA采用Single Instruction Multiple Thread（SIMT）的架构来管理和执行thread，这些thread以32个为单位组成一个单元，称作warps。warp中所有线程并行的执行相同的指令。每个thread拥有它自己的instruction address counter和状态寄存器，并且用该线程自己的数据执行指令。

　　SIMT和SIMD（Single Instruction, Multiple Data）类似，SIMT应该算是SIMD的升级版，更灵活，但效率略低，SIMT是NVIDIA提出的GPU新概念。二者都通过将同样的指令广播给多个执行官单元来实现并行。一个主要的不同就是，SIMD要求所有的vector element在一个统一的同步组里同步的执行，而SIMT允许线程们在一个warp中独立的执行。SIMT有三个SIMD没有的主要特征：

• 每个thread拥有自己的instruction address counter
• 每个thread拥有自己的状态寄存器
• 每个thread可以有自己独立的执行路径

一个block只会由一个SM调度，block一旦被分配好SM，该block就会一直驻留在该SM中，直到执行结束。一个SM可以同时拥有多个block。

 

需要注意的是，大部分thread只是逻辑上并行，并不是所有的thread可以在物理上同时执行。这就导致，同一个block中的线程可能会有不同步调。

并行thread之间的共享数据回导致竞态：多个线程请求同一个数据会导致未定义行为。CUDA提供了API来同步同一个block的thread以保证在进行下一步处理之前，所有thread都到达某个时间点。不过，我们是没有什么原子操作来保证block之间的同步的。

同一个warp中的thread可以以任意顺序执行，active warps被SM资源限制。当一个warp空闲时，SM就可以调度驻留在该SM中另一个可用warp。在并发的warp之间切换是没什么消耗的，因为硬件资源早就被分配到所有thread和block，所以该新调度的warp的状态已经存储在SM中了。

SM可以看做GPU的心脏，寄存器和共享内存是SM的稀缺资源。CUDA将这些资源分配给所有驻留在SM中的thread。因此，这些有限的资源就使每个SM中active warps有非常严格的限制，也就限制了并行能力。所以，掌握部分硬件知识，有助于CUDA性能提升。

Fermi架构与Kepler 架构

warp：

逻辑上，所有thread是并行的，但是，从硬件的角度来说，实际上并不是所有的thread能够在同一时刻执行，接下来我们将解释有关warp的一些本质。