基于CUDA的异构并行计算

一、并行计算

目的：提高运算速度

原则：将大的问题划分为很多可以同时解决的小问题

要求：硬件方面（计算机架构支持并行执行多进程或多线程）、软件方面（充分利用硬件并行执行）

1. 并行编程与串行编程

        编程时我们会很自然的把问题划分为很多运算块，块之间依次执行——串行。如果有并行执行的运算块则为并行。有些块可以并行，有些块则必循串行执行，视情况而定。比如数据相关发生时，就会限制程序的并行性。

2. 并行性

        并行存在两种：任务并行、数据并行。任务并行：任务或函数可以独立地并行执行，重点在于利用多核系统对任务进行分配。数据并行：同时处理许多数据，重点在于利用多核系统对数据进行分配。

        数据并行是指当大量的数据进行相同的运算操作时，可以通过同时在多个计算单元上运行相同的指令，分别处理不同的数据，最终共同完成这些数据的运算。数据并行第一步就要对数据进行划分，划分的方式有两种，块划分和周期划分。

块划分：每个线程作用于一部分数据，通常数据具有相同大小。

周期划分：每个线程作用于数据的多部份。

3.计算机架构

弗林分类：SISD、SIMD（单指令多数据，并行架构）、MISD、MIMD（多个核心使用多个指令流异步处理多个数据流）。

内存组织方式分类：分布式内存的多节点系统（每个处理器有自己的本地内存，处理器之间通过网络进行通信）、共享内存的多处理器系统（与同一个物理内存相关联或者共用一个低延迟的链路）。

GPU核心和CPU核心

CPU核心比较重，用于处理比较复杂的控制逻辑，控制密集型。

GPU核心比较轻，用于处理具有简单控制逻辑的任务，计算密集型。在许多数据元素上并行执行相同的计算控制流开销低，浮点运算强度高。每次内存访问都要进行许多计算。

 

 二、 异构计算

1. 异构架构

        一个异构计算节点 = 多个多核CPU + 多个众核GPU。GPU是CPU的协处理器，二者通过PCIe相联，可将CPU视为主机端，GPU视为设备端。

        一个异构计算应用 = 主机代码 + 设备代码。

2. CUDA

        CUDA提供了两层API来管理GPU设备，组织线程，如图所示。CUDA驱动API是一种低级API，较难编程，但其对于GPU设备使用上提供了更多的控制。运行时API是一个高级的API，在驱动API上层实现，每个运行时 API被分解为 多个传递给驱动API的基本运算。两种API相互排斥，只能使用二者之一。 

         一个CUDA程序，包含两部分，主机代码和设备代码。nvcc编译时将设备代码从主机代码中拆分出来，主机代码为标准的C代码，由C编译器进行编译，设备代码不是标准C，由nvcc进行编译。链接时在内核程序调用和显示GPU设备操作中添加CUDA运行时库。

3.第一个CUDA程序

(1) __global__关键字表明以下函数将在 GPU 上运行并可全局调用，而在此种情况下，则指由 CPU 或 GPU 调用。
  - 注意返回类型为 `void`。使用 `__global__` 关键字定义的函数需要返回 `void` 类型。

cudaDeviceSynchronize()
(2) 核函数启动方式为异步：CPU 代码将继续执行而无需等待核函数完成启动。
  - 调用 CUDA 运行时提供的函数 `cudaDeviceSynchronize` 将导致主机 (CPU) 代码暂作等待，直至设备 (GPU) 代码执行完成，才能在 CPU 上恢复执行。

(3) nvcc 进行编译：

nvcc -arch=sm_70 -o out test.cu -run

• nvcc 是使用 nvcc 编译器的命令行命令。

• 将 test.cu 作为文件传递以进行编译。

• o 标志用于指定编译程序的输出文件。

• arch 标志表示该文件必须编译为哪个架构类型。本示例中，sm_70 将用于专门针对本实验运行的 Volta GPU 进行编译，但有意深究的用户可以参阅有关 arch 标志、虚拟架构特性 和 GPU特性 的文档。

• 为方便起见，提供 run 标志将执行已成功编译的二进制文件。

#include <stdio.h>

void helloCPU()
{
  printf("Hello from the CPU.\n");
}

__global__ void helloGPU()
{
  printf("Hello also from the CPU.\n");
}

int main()
{

  helloCPU();
  helloGPU<<<1,10>>>(); 
  //异步执行，CPU代码等待GPU代码执行结束
  cudaDeviceSynchronize();
}

 (4) 运行结果：