gpu host to device的拷贝速度最多有多块

CUDA程序的编写

数据传输过程：把输入数据从CPU内存复制到GPU显存，在执行芯片上缓存数据，加载GPU程序并执行，将计算结果从GPU显存中复制到CPU内存中。

__global__ 和 __device__ 都在GPU内执行。
__device__只能由device调用，即在__global__和__device__修饰的函数内被调用
__global__都可以。
__host___执行位置和调用位置都是host （不标注的情况下默认为__host__）

NVPROF

分析程序时间的强大工具

nvprof a.exe (summary的模式，测得一些函数的时间)
nvprof --print-gpu-trace a.exe （查看的是占用GPU的资源，比如Grid Size，Block Size等等）
nvprof --print-api-trace a.exe（调用api（cuda）的顺序，以及时间）

如何使用核心

一个核心就是一个兵，管理核心就是练兵。

CUDA线程层次以及执行单位：

• Thread——CUDA Core级别的调度单位
  所有线程执行相同的核函数（但处理的数据是不同的）、并行执行的基本单位。
  SM基本的执行单元是线程束，允许同一个warp中的thread读取其他thread的值
• Thread Block——SM级别的调度单位
  一个Block必须在一个Streaming Multiprocessor（SM）执行，一个SM中有多个Block
  同一个Block中的线程可以协作，同一个SM同一个SM（shared memory）
• Thread Grid——Device（GPU级别的调度单位）
  一个Grid中的Block可以在多个SM中执行
  grid共享资源kernel和执行上下文context

每个线程层次可以展开成三维坐标

• idx类
  -threadIdx是指当前kernel函数的线程在block中的索引值
  -blockIdx是指当前kernel函数的线程所在block，在grid中的索引值
• dim类
  -blockDim表示一个block包含多少个线程
  -gridDim表示当前grid包含多少个block

<<<gidDim,blockDim>>>，<<<1, 4>>>。1表示一个block，4表示1个block四个线程。blockDim设置为32的倍数。

确定index线程索引

确定索引值是从x，y，z三维分别单独确定的，考虑问题也要从某一特定维度考虑。首先是确定要求的是在哪一维上的索引值，以x维为例：
即给定的是blockDim.x，blockIdx.x，threadIdx.x

int index = blockDim.x * blockIdx.x + threadIdx.x

CUDA的线程分配

以warp为基本单位 Warp is successive 32 threads in a block.

所以blockDim尽量设置为32的整数倍

如果数据过大，线程不够用怎么办？那就一轮一轮的算。

n为数据总数
__global__ add() {
	int index = blockDim.x * blcokIdx.x + threadIdx.x;
	int stride = blockDim.x * gridDim.x;
	for (;index < n;index += stride)
		z[index] = x[index] + y[index];
}