20180323CUDA入门（一）

CUDA自学

3/24

参考了cuda手册和http://blog.chinaunix.net/uid-20620288-id-4655719.html

---

Chapter 1 Introduction

• GPU需要处理大量并行工作，FLOPS(floating-point operations per second)比CPU大得多
• 出于大量并行计算的需要，GPU的更多面积用来放计算单元而不是cache单元或者flow control
• 3 key abstractions: a hierarchy of thread groups, shared memories, and barrier synchronization
• 目的: solve the problems into coarse sub-problems that can be solved independently in parallel by blocks of threads, and each sub-problem into finer pieces that can be solved cooperatively in parallel by all threads within the block
• 一个编译好的程序可以跑在任何核数的GPU上，运行时才确定具体哪几个block分配给哪几个核

Chapter 2 Programming Model

• kernels：定义C函数，称作kernels。调用时，在n个进程中执行n次

  用_global声明，跑它的进程个数N在调用时用<<<…>>>标记

  跑他的进程ID在进程内部可见，为一个三维向量threadIdx，可以用三位组成各种标记

  

  寻址方式类似多维数组

  一个thread block最多1024个thread

  int, dim3类型：threads per block / blocks per grid,threadblock的位数在kernel函数中通过blockDim访问

  Dim是尺寸，Idx是具体的索引

  多个block的调用示例如下

  // Kernel definition
  _global_ void MatAdd(float A[N][N], float B[N][N], float C[N][N])
  {
  int i = blockIdx.x * blockDim.x + threadIdx.x;
  int j = blockIdx.y * blockDim.y + threadInx.y;
  
  if (i < N && j < N)
      C[i][j] = A[i][j] + B[i][j];
  }
  
  int main()
  {
  ...
  // Kernal invocation
  dim3 threadPerBlock(16, 16);
  dim3 numBlocks(N / threadPerBlock.x, N / threadPerBlock.y);
  MatAdd<<<numBlocks, threadPerBlock>>>(A, B, C);
    ...
  }
  //就是个分块矩阵

  同一个block内的threads可以通过shared memory或synchronize their execution to coordinate memory access共享数据。通过_syncthreads()同步

• 内存层次

  在GPU上CUDA线程可以访问到的存储资源有很多，每个CUDA线程拥有独立的本地内存（local Memory）；每一个线程块（block）都有其独立的共享内存（shared memory），共享内存对于线程块中的每个线程都是可见的，它与线程块具有相同的生存时间；同时，还有一片称为全局内存（global memory）的区域对所有的CUDA线程都是可访问的。

  除了上述三种存储资源以外，CUDA还提供了两种只读内存空间：常量内存（constant memory）和纹理内存（texture memory），同全局内存类似，所有的CUDA线程都可以访问它们。对于一些特殊格式的数据，纹理内存提供多种寻址模式以及数据过滤方法来操作内存。这两类存储资源主要用于一些特殊的内存使用场合。

  一个程序启动内核函数以后，全局内存、常量内存以及纹理内存将会一直存在直到该程序结束。下面是CUDA的内存层次图：
  

• 异构多相编程

  按照revision number分类