CUDA C 编程权指南-professional CUDA C programming chapter-one 基于CUDA的异构并行计算

1. Introduction — CUDA C Programming Guide (nvidia.com)

CUDA Runtime API :: CUDA Toolkit Documentation (nvidia.com)

CUDA C编程权指南 professional CUDA C programming​product.dangdang.com/25089854.html

以下的内容主要来自这个页面：

1. Introduction — CUDA C Programming Guide (nvidia.com)，因书籍的内容也还全面，所以这的话主要是做补充。

CUDA C++ Programming Guide

1.1. The Benefits of Using GPUs

The Graphics Processing Unit (GPU) 图形处理器，CPU主要用来执行串行指令，像较少的较高速度的线程，GPU则主要是用来大量的并行执行，线程的大量吞吐量掩盖了较低速的线程。GPU被设计用来高并行计算，并且因此设计了更多的晶体管来做数据处理，而不是数据缓存或者流程控制的。

GPU是基于一系列流多处理器 Streaming Multiprocessors (SMs)，多线程程序被划分到线程块，每个线程块独立执行。所以一个GPU有着更多的 multiprocessor的执行时间更短。

以下内容来自书籍

《CUDA C编程权威指南》([美]程润伟（John Cheng）)【简介_书评_在线阅读】 - 当当图书 (dangdang.com)

序

CUDA的全称就是 compute Unified Device architecture 统一计算设备架构，是Nvidia提出的并行计算架构，用来处理密集型并行计算，CPU和GPU是两个独立的处理器，通过单个计算节点的PCI-Express来connect，GPU用来提高计算密集型应用程序中并行程序段的执行速度，CPU则负责管理设备端。

基于CUDA的异构并行计算

并行计算的目的主要是提高运算速度，让计算机同时做很多的运算，主要依靠了CUDA的多核，并行计算涉及两个部分：1、计算机硬件的架构，2、并行程序设计软件；硬件必须提供支持并行执行多进程或多线程的端。

单核CPU不支持并发执行，多核CPU支持并发执行的；

1.1.1 串行编程和并行编程

串行执行按照时间顺序依次执行，并行执行按照时间顺序并行执行可以包括串行。程序包含两个基本的部分：指令和数据，一个计算问题被划分到很多小的计算unit，若一个计算unit处理的是另一个计算unit的输出，就说这两个计算unit是相关的。否则是独立的。

1.1.2 并行性

两类并行类型：1、数据并行，同时处理很多的数据；2、task并行：task或者函数可以同时运算，数据没有交互，一般使用多核。CUDA编程适合解决数据并行的问题，数据并行处理是将数据映射给并行线程。

数据编程第一步：按照线程分配数据，每个线程处理一部分数据，数据划分：分 块划分 和 周期划分，块划分：一组continuous的数据分到一个块内，每个数据块以任意次序分配给线程，每个线程作用一部分数据；周期划分：每个线程作用数据的多个部分。

1.1.3 计算机架构

是弗林分类法（Flynn’s Taxonomy），根据指令和数据进入CPU的方式，可以划分到几个不同的类型，single: 单个，multi: 多个，instruction：指令，data：数据。也就是SISD：串行，SIMD：多核，MIMD：多核，MISD：少见。CUDA架构遵循的是SIMD单指令多数据。

延迟： 一个操作从开始到完成需要的时间， 带宽： unit时间内可以处理的数据量，通常表示是：MB/s或者GB/s，存储基本unit是bit，每个字节byte包含了8bit，1KB包含了512x2个byte字节，1MB包含了512x2个KB，1GB包含了512x2个MB，1TB包含了512x2个GB，1PB包含了512x2个TB。 吞吐量： 是unit时间内成功处理的运算数量，通常表示是：gflops(每秒十亿次的浮点运算数量)。

根据内存组织方式划分：分布式内存的多节点系统：由网络connect很多的处理器，每个处理器有自己的本地内存，处理器之间可以通过网络通信，也就是集群；共享内存的多处理器系统：多核多处理器，共享内存的共享内存地址空间；GPU有着成百上千的核心，适合解决大规模的并行计算问题。

1.2 异构计算

异构计算节点包括：多核CPU和很多核GPU，GPU是不能独立运行的，是CPU的协作处理器，GPU必须要通过PCIe总线和基于CPU的主机connect，CPU是主机端，GPU是设备端。

异构应用包括：主机host program运行在CPU，设备端device program运行在GPU。GPU可看作是硬件加速器。

Tegra系列产品是专$移动和嵌入式设备而设计的，如平板电脑和手机，GeForce面向 图形用户，Quadro用于专业绘图设计，Tesla用于大规模的并行计算。Fermi是Tesla系列产 品中的一种，用作GPU加速器。

描述GPU容量的重要特性：CUDA的核心数量、内存的大小；GPU的性能：峰值计算性能、内存带宽的； 峰值计算性能 和吞吐量类似，都是每秒能处理的单精度或双精度浮点数运算的数量，GFlops或者TFlops。内存带宽是从内存读取或者写入的比率，GB/s。

CUDA使用计算能力做版本号

1.2.2 异构计算范例

CPU适合控制密集型运算，GPU适合包含数据并行的计算密集型运算。

从两个方面来区分CPU和GPU的应用区间：并行数量级；数据的规模。问题的数据规模较小、复杂的控制逻辑和或很少的并行性，此时应该选择CPU处理；若是问题包含大规模的数据，并表现出大量的数据并行性，使用GPU是最好的选择，因GPU内有大量的可编程core，可以支持大规模的多线程运算，峰值带宽也比CPU大。

CPU和GPU功能互补导致了CPU+GPU的异构计算架构的发展，最佳性能的：CPU执行串行或者任务并行部分，GPU执行数据密集型运算。

CPU线程与GPU线程

CPU线程：重量级实体，上下文切换缓慢且开销大，被设计用来减少线程延迟，可支持的线程数量少；

GPU线程：大量成千上万，处理大量并发轻量级线程

1.2.3 CUDA：一种异构计算平台

是一种通用的并行计算平台和编程模型，利用GPU并行计算引擎有效解决复杂计算问题，CUDA平台可以通过CUDA加速库、编译器指令、应用编程接口、以及行业标准的扩展语言来使用（像c, c++, python等）。

CUDA C是标准ANSI C语言的一个扩展，能通过API来管理设备、内存或者其他的任务

CUDA提供了两层API来管理GPU设备和组织线程，1、CUDA运行时API：高级API，基于驱动API，分解到驱动API，2、CUDA驱动API：低级API，较难实现控制的更多。运行时API和驱动API没有明显的性能差异，内核如何使用内存以及你是如何组织线程的，对性能有更显著的影响。

两类API互斥，只能使用一种，不能混用的。通常CUDA程序混合了：1、在CPU上运行的host program，2、在GPU上运行的device program。

CUDA NVCC编译器在编译时将device program从host program分离出来的，host program由C编译器编译，像g++。device program或者核函数，通过nvcc编译的，link阶段会使用到运行时库。

1.3 用GPU输出Hello World

检查cuda编译器是否安装好： which nvcc ，检查机器是否安装了GPU显卡： ls -l /dev/nv，source* 扩展名.cu。编译的方式是：

// hello.cu
/*
nvcc hello.cu -o hello
./hello
*/
#include <stdio.h>

int main(void)
{
    printf("Hello World!\n");
    return 0;
}

下面的核函数分配了1个网格block=grid size，10个线程thread=block size，也就是运行了10次核函数的。__global__修饰符告诉编译器这个函数从CPU调用，在GPU执行，启动内核函数用的是helloFromGPU<<<1,10>>>();三重尖括号也是标识符，用来表示从主机端到设备端，三重尖括号内的是执行配置，也就是<<<grid_size, block_size, 共享内存大小, *>>> 。编译的方式是：nvcc -arch sm_60

hello.cu -o hello，-o是重命名a.out可执行档案到hello，-arch则是指定了架构。完整的写作方式应该是：

核函数没有返回值，所以返回类型是 void

#include "../common/common.h"
#include <stdio.h>

/*
 * A simple introduction to programming in CUDA. This program prints "Hello
 * World from GPU! from 10 CUDA threads running on the GPU.
 */

__global__ void helloFromGPU()
{
    printf("Hello World from GPU!\n");
}

int main(int argc, char **argv)
{
    printf("Hello World from CPU!\n");

    helloFromGPU<<<1, 10>>>();
    CHECK(cudaDeviceReset());  // 显式的释放和清空当前进程和设备有关的source
    return 0;
}
/*
Hello World from GPU!
Hello World from GPU!
Hello World from GPU!
......
*/

完整的GPU编程的过程：1、分配GPU内存的，2、从CPU内存拷贝数据到GPU内存，3、调用CUDA内核函数完成运算，3+1、将数据从GPU拷贝回CPU内存，3+2、释放GPU的内存空间

使用 CUDA C 编程难的么

CPU编程和GPU编程的主要区别是programmer对GPU架构的了解程度。数据局部性是在并行编程中非常重要的概念，指的是数据重用，以降低内存访问的延迟，包括了两种类型：时间局部性是指在相对较短的时间段内数据和或资源的重用，空间局部性是指在相比较接近的存储空间内数据element的重用。

CUDA内有内存层次和线程层次的概念，使用这些层次结构可以进行更高层次的控制和调度。1、内存层次结构，2、线程层次结构。在CUDA编程模型中使用的共享内存能节省带宽。共享内存是在线程块以内的内存，由线程块以内的线程共享。CUDA C是C语言程序的扩展，通常可以直接移植C到CUDA C内，剥离program内的循环后产生CUDA C的内核函数。CUDA抽象了硬件细节，不需要将应用程序映射到图形API上。

CUDA开发环境

nsight 集成开发环境，CUDA-GDB命令行调试器，用于性能分析的可视化和命令行分析器，CUDA-MEMCHECK内存分析器，GPU设备管理工具。

1.5 总结

CPU+GPU的异构系统在高性能计算领域已经是main streaming，GPU执行数据并行工作，CPU执行串行和任务并行工作。

要掌握CUDA编程，需要对GPU架构相当的了解，并且掌握内存层次分层，因不同的内存的带宽不相同，数据传输的延迟也不相同，内存层次主要是：

全局内存：数量最多基本就是显存的大小，

https://zhuanlan.zhihu.com/p/659389698