(一)CUDA编程

CPU：擅长流程控制和逻辑处理，不规则数据结构，不可预测存储结构，单线程程序，分支密集型算法

GPU：数据并行计算，规则数据结构，可预测存储模式

一般而言，并行部分在GPU上运行，串行部分在CPU运行，CPU负责总体的程序流程，而GPU负责具体的计算任务，当GPU各个线程完成计算任务后，就将GPU计算结果拷贝到CPU端，完成一次计算任务。

1、CUDA线程模型

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

Thread: 线程，并行的基本单位

Thread Block：线程块，互相合作的线程块，线程块有如下几个特点：允许彼此同步；通过共享内存快速交换数据；

Grid：一组线程块，以1维、2维组织；共享全局内存

Kernel：在GPU上执行的核心程序，kernel函数是运行在某个Grid上的

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

理解kernel，必须要对kernel的线程层次结构有一个清晰的认识。首先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.xldx,grid.yldy),对应上图例子就是(1,1),通过它我们就能找到这个线程块的位置，然后启动二级索引(block.xldx,block.yldx,block.zldx)来定位到指定的线程。

2、CUDA内存模型

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

每个线程都用自己的registers（寄存器）

每个线程都有自己的local memory（局部内存）

每个线程块内都有自己的shared memory（共享内存），所有线程块内的所有线程共享这段内存资源

每个grid都有自己的global memory（全局内存），不同线程块的线程都可使用

每个grid都有自己的constant memory（常量内存）和texture memory（纹理内存），不同线程块的线程都可使用

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

3、代码表示线程组织模型

由于GPU的线程太多，为GPU的线程划分国(grid)-省(block)-市(thread)的分级。

在一个grid中有很多block，声明一个有4*3个block的grid：dim3 grid(4,3);
gridDim.x=4;     gridDim.y=3
深绿色block(blockIdx.x=0,blockIdx.y=0)有自己的位置：
blockIdx.x=0;      blockIdx.y=0;  //第一行 第一列
让定义一个 3*2个thread的block:     dim3 block(3,2);
thread也有自己的位置。浅绿色(threadId.x=0,threadId.y=0)的Thread的位置：
blockIdx.x=3;    blockIdx.y=0;       //block 第一行 第四列
threadId.x=0;  threadId.y=0;       //thread 第一行 第一列

4、CUDA和C++结合编程中错误异常

在代码实现过程中cuda的plus<<<>>>()函数总是报错。错误信息:语法错误:”<”

头文件plus.h:

#include <stdlib.h>
#include <device_launch_parameters.h>
#include <cuda_runtime.h>
#include <iostream>
extern "C" {
    void pl();
}

源文件plus.cu

#include "plus.h"
__global__ void plus(float a[], float b[], float c[], int n) {
  int i = blockDim.x * blockIdx.x + threadIdx.x;
  c[i] = a[i] + b[i];
};
void pl() {
  float *A, *Ad, *B, *Bd, *C, *Cd;
  int n = 1024 * 1024;
  int size = n * sizeof(n);
  A = (float*)malloc(size);  // CPU端分配内存
  B = (float*)malloc(size);
  for (size_t i = 0; i < n; i++)  //初始化数组
  {
    A[i] = 90;
    B[i] = 10;
  }
  cudaMalloc((void**)&Ad, size);  // GPU端分配内粗
  cudaMalloc((void**)&Bd, size);
  cudaMalloc((void**)&Cd, size);
  cudaMemcpy(Ad, A, size, cudaMemcpyHostToDevice);  // CPU的数据拷贝到GPU端
  cudaMemcpy(Bd, B, size, cudaMemcpyHostToDevice);
  //定义kernel执行配置,(1024*1024/512)个block,每个block里面有512个线程
  dim3 dimBlock(512);
  dim3 dimGrid(n / 512);
  plus<<<dimGrid, dimBlock>>>(Ad, Bd, Cd, n);  // kernel
  free(A);
  free(B);
  cudaFree(Ad);
  cudaFree(Bd);
  cudaFree(Cd);
}

test.cpp文件中:

#include "plus.h"
int main(){ pl(); }

在CMakeLists.txt文件中

cmake_minimum_required(VERSION 2.8)
find_package(CUDA REQUIRED)
cuda_add_executable(squaresum test.cpp plus.cu)

5、设置VScode支持.cu文件语法高亮及跳转的方法

插件：vscode插件商店搜索cudacpp进行安装后，可支持语法高亮以及<<<>>>等cuda专用符号。

参考《VScode 为 *.cu文件 添加高亮及c++ intelligence相关操作的方法》，设置settings文件，添加文件cu后缀文件与cpp的关联："files.associations":{"*.cu":"cpp"}。设置完成后可支持cpp的语法高亮与跳转。