CUDA课程-CS-139G-2:GPU History & CUDA Programming Basics

 CUDA基础框架
 1，kernels基础以及在GPU上的执行
 2，基本存储管理
 3，CPU管理以及GPU执行
 4，参看API编程指南
 
 CUDA编程模型
 1，kernel：启动和运行在一个device上，通过很多线程运行的并行代码
 2，启动是分层次的：线程组成线程块；线程块又组成grids
 3,常见的串行代码用来写一个线程：每个线程都很自由的执行一个独一无二的代码路径
   内建的线程和线程块标识变量。
 
 高层次的组织图
 CPU和GPU之间通过PCIE接口传输数据，
 
 运行在一个SM(streaming Mutiprocessor)
 每个线程块共享一部分内存
     
 运行在GPU上的整个grid:
     
 线程的等级制度
 1，引发并行部分的线程被分为各个线程块：grid就是一个已知任务的所有线程块
 2，线程块是由一组线程组成的，可以
    使执行同步；并且动过共享内存进行通信。
 
 内存模型
 1，不同的kernel读取global Memory (GPU 内部的存储空间)
 2，函数cudaMemcpy()用于GPU与CPU之间的数据拷贝。

 向量相加例子：
 1,Device代码(GPU): "__global__" "int i = threadIdx.x+blockDim.x*blockIdx.x"
 2,Host代码（CPU）:"Function_name<<<size1,size2,size3>>>(argu1,argu2,...argun)"
 3,一些常见的代码特征 cudaMalloc;--数据空间分配
                                 cudaMemcpy;--数据拷贝
                              cudaMemcpyHostToDevice,cudaMemcpyDeviceToHost --数据拷贝方向
                              cudaFree--释放数据空间
      
 4,Kernel变化以及输出
 
     __global__ void kernel( int*a )
     {
        intidx= blockIdx.x*blockDim.x+ threadIdx.x;
        a[idx] =7                                //输出7 7 7 7 7 7 7 7...7 7 7 7
        //a[idx] = blockIdx.x;    //输出0 0 0 0 0 0 0 0...0 0 0 0
        //a[idx] = threadIdx.x;    //输出0 1 2 3 0 1 2 3...0 1 2 3
    }
    
 blockIdx：block的ID
 blockDim:block的size
 threadIdx:thread的ID
 
 执行在GPU上的代码：
 限制：
 1，只能访问GPU存储空间
 2，不能用可变参数列表（）
 3，不允许静态变量
 4，不允许递归
 5，不允许动态多态性
 
 必须声明的前缀：
 1，__global__:由CPU调用，不能从GPU内调用，必须返回void
 2, __device__:从GPU调用，不能被CPU调用。
 3，__host__可以被CPU调用，和__device__可以联合使用。
 
 存储空间
 1，CPU和GPU有相互独立的存储空间。
   数据通过PCIE总线传输
   使用一些专门函数实现分配/设置/拷贝GPU上的存储空间（与C语言很像）
 2，指针就是地址
   通过指针不能得知地址在CPU还是GPU
   解引用的时候必须小心：在GPU上解引用CPU的指针将很可能导致崩溃
   反之亦然。
 
 GPU存储分配和释放
 
 1，CPU管理GPU存储空间
   cudaMalloc(void** pointer,size_t nbytes)                    开辟一定空间
   cudaMemset(void* pointer, int value, size_t cout) 将空间设置成1或者0等初始化工作
   cudaFree(void* pointer)    回收空间

intn = 1024;
intnbytes= 1024*sizeof(int);
int* d_a= 0;
cudaMalloc( (void**)&d_a, nbytes);
cudaMemset( d_a, 0, nbytes);
cudaFree(d_a);

2,数据拷贝

cudaMemcpy(void* dst, void* src, size_t nbytes, enum cudaMemcpyKind direction);
 当拷贝完全的时候才返回。
 所有拷贝结束才释放CPU线程。
 只有当前一个调用结束了才开始拷贝。
enum cudaMemcpyKind direction 包括4个拷贝方向
cudaMemcpyHostToDevice
cudaMemcpyDeviceToHost
cudaMemcpyDeviceToDevice
还可实现非阻塞型的拷贝。

需要注意是拷贝的和设置之类的动作都是基于指针下的操作，所以需要事先定义好指针。才接着对这些指针进行操作。
如果需要参数是指向指针的指针类型（void**）那么传递的参数将是（void**）&pointer
如果需要参数是指针，则直接传递指针本身。即可。

通常分配完空间之后可以对分配好点数据空间进行检查:
if(0==host_data || 0==device_data)

然后是正常的初始设置和拷贝，注意到是设置一般全设置成0；拷贝是特别注意拷贝方向的指定
最后就调用kernel函数，最后记得需要释放空间，操作对象是指针。


3，唯一的ID标号和尺寸
thread:3维，在一个block内具有唯一的标号；
Block: 2维，在一个grid内具有唯一的标号；
引导时尺寸设定：对每一个grid都是唯一的；
内建的变量：threadIdx,blockIdx,blockDim,gridDim

具有2维索引的Kernel

int idx=blockIdx.x*blockDim.x+threadIdx.x;
int idy=blockIdx.y*blockDim.y+threadIdx.y

Block必须相互独立
  任何有可能交叉的blocks应该是有效的。运行结束没有提示符；无序运行，可并发和相继运行。
Blocks可协调，但是并不是同步的
  共享队列指针--OK
  共用锁定--BAD--容易锁死
因此相对独立的需求提供了可伸缩性。

GPU历史
简要绘图：
1.生成图片：复杂的形状，复杂的视觉作用，无缝运动
2，加速：创造更聪明的技术，运用所有想象得到的办法。建立更加高级的硬件
绘图管线：
管道变换和照明--照明设置和光栅化--纹理和像素着色--明暗测试和混合-绘图缓冲存储
顶点--光栅化--点阵化--测试混合--缓冲存储

繪圖管線的系統，從應用程式的場景(Scene)管理，到幾何(Geometry)資料的算術運算，再進入到三角形資料的設置，最後透過渲染 (Rendering)將圖形依貼圖座標或相關的著色(shading)演算法，讓畫面呈現在螢幕上。
上述的步驟可以稱為階段(Stage)。

1，实时化绘图的关键抽象点。
2，硬件看上去像
3，每阶段一个芯片
4，固定的数据流流通管道
5，对于特定分辨率有特定功能
6，每个阶段的分辨率数量超高速增长
7，硬件难以优化
8，开发者需要更大的灵活
9，保持一个关键抽象，硬件习惯于看上去像；顶点和像素处理要可编程，增加了新的阶段
10，GPU架构围绕着着色执行程序增加
11，为一些阶段陈列一个（最初限制）指令集
12，限制的指令和指令类型没有历程控制，
3，扩展到所有的工业标准

为什么GPU曲线很好，
1，工作和编程模式提供了大量的并行计算
2，应用程序一次性提供大量的节点：节点可并行处理，所有节点应用同样的转换
3，三角形包含大量的像素：一个三角形中的像素可以并行处理;所有的像素应用一样的着色
4，有效的硬件隐藏了串行化带来的瓶颈

更大的效能：
1，给大量的像素和节点做一模一样的操作。是一个控制单元对多个运算逻辑单元（ALU）
2，一个warp由32条线程组成。通常一起同时执行。

早期的通用图形处理器 （GPGPU）
1,所有这些绩效可以影响到开发者，
2，为了使用到GPU,要用图形计算来重新表达他们的算法。
3，乏味，有限的可用性
4，仍然有好多效果

这就给CUDA做好了准备