重点内容

• CUDA和GPU简介

• GPU工作原理与结构

• GPU应用领域

• GPU+CPU异构计算

• CUDA环境搭建

• CUDA向量加法

• CUDA矩阵乘法

• MPI与CUDA的区别

难点：

• grid、block、thread的关系

• grid、block、thread的一二三维的映射

• SM的运行逻辑

2 GPU工作原理与结构

GPU采用流式并行计算模式，可对每个数据行独立的并行计算。

GPU与CPU区别：

CPU基于低延时设计，由运算器（ALU，Arithmetic and Logic Unit 算术逻辑单元）和控制器

（CU,Control Unit），以及若干个寄存器和高速缓冲存储器组成，功能模块较多，擅长逻辑控

制，串行运算。

GPU基于大吞吐量设计，拥有更多的ALU用于数据处理，适合对密集数据进行并行处理，擅长

大规模并发计算，因此GPU也被应用于AI训练等需要大规模并发计算场景。

2.1 基础GPU架构

GPU为图形图像专门设计，在矩阵运算，数值计算方面具有独特优势，特别是浮点和并行计算上能

优于CPU的数十数百倍的性能。

GPU的优势在于快，而不是效果好。

比如用美团软件给一张图要加上模糊效果，CPU处理的时候从左到右从上到下进行处理。可以考虑

开多核，但是核数毕竟有限制，比如4核、8核 分块处理。

使用GPU进行处理，因为分块之前没有相互的关联关系，可以通过GPU并行处理，就不单只是4、8分块了，可以切换更多的块，比如16、64等

2.2 GPU编程模型

软件层面上不管什么计算设备，大部分异构计算都会分成主机代码和设备代码。整体思考过程就是

应用分析、内存资源分配、线程资源分配再到具体核函数的实现。

CUDA中线程也可以分成三个层次：线程、线程块和线程网络。

线程是CUDA中基本执行单元，由硬件支持、开销很小，每个线程执行相同代码；

线程块（Block）是若干线程的分组，Block内一个块至多512个线程、或1024个线程（根据不同的GPU规格），线程块可以是一维、二维或者三维的；

线程网络（Grid）是若干线程块的网格，Grid是一维和二维的。

线程用ID索引，线程块内用局部ID标记threadID，配合blockDim和blockID可以计算出全局ID，用于SIMT（Single Instruction Multiple Thread单指令多线程）分配任务。

首先需要关注的是具体线程数量的划分，在并行计算部分里也提到数据划分和指令划分的概念，GPU有很多线程，在CUDA里被称为thread，同时我们会把一组thread归为一个block，而多个block又会被组织成一个grid。

假如我们要对一个长度为1024的数组做reduce_sum（减少和求和），恰好我们正好有1024个thread，此时直接一一对应就行，但如果是一张很大的图片呢？

如果有很多核函数要处理不同的数据呢？GPU上有很多thread，但要完全和实际应用中需要处理的数据大小完全匹配是不可能的，事实上在满足规定的情况下我们可以给一个block内部分配很多thread，对于到硬件上也真的是相应数量的thread会自动归为一组直接在一个SM上实行吗？

答案当然不是，此时我们就要关注硬件，引入了wrap概念, GPU上有很多计算核心也就是Streaming Multiprocessor (SM)，在具体的硬件执行中，一个SM会同时执行一组线程，在CUDA里叫wrap，我们不用拘泥于称呼，直接可以理解这组硬件线程会在这个SM上同时执行一部分指令，这一组的数量一般为32或者64个线程。一个block会被绑定到一个SM上，即使这个block内部可能有1024个线程，但这些线程组会被相应的调度器来进行调度，在逻辑层面上我们可以认为1024个线程同时执行，但实际上在硬件上是一组线程同时执行，这一点其实就和操作系统的线程调度一样。 意思就是假如一个SM同时能执行64个线程，但一个block有1024个线程，那这1024个线程是分1024/64=16次执行。

解释完了执行层面，再来分析一下内存层面上的对应，一个block不光要绑定在一个SM上，同时一个block内的thread是共享一块share memory（一般就是SM的一级缓存，越靠近SM的内存就越快）。GPU和CPU也一样有着多级cache还有寄存器的架构，把全局内存的数据加载到共享内存上再去处理可以有效的加速。所以结合具体的硬件具体的参数（SM和寄存器数量、缓存大小等）做出合适的划分，确保最大化的利用各种资源（计算、内存、带宽）是做异构计算的核心。

2.2.1 软件和硬件的对应关系

GPU在管理线程(thread)的时候是以block(线程块)为单元调度到SM上执行。每个block中以warp(一般32个线程或64线程)作为一次执行的单位(真正的同时执行)。

一个 GPU 包含多个 Streaming Multiprocessor ，而每个 Streaming Multiprocessor 又包含多个 core 。

Streaming Multiprocessors 支持并发执行多达几百的 thread。

一个 thread block 只能调度到一个 Streaming Multiprocessor 上运行，直到 thread block 运行完毕。一个Streaming Multiprocessor 可以同时运行多个thread block （因为有多个core）。

通俗点讲：stream multiprocessor(SM)是一块硬件，包含了固定数量的运算单元，寄存器和缓存。

写cuda kernel的时候，跟SM对应的概念是block，每一个block会被调度到某个SM执行，一个SM可以执行多个block。

你的cuda程序就是很多的blocks(一般来说越多越好)均匀的喂给这80个SM来调度执行。具体每个block喂给哪个SM你没

法控制。

不同的GPU规格参数也不一样，比如 Fermi 架构（2010年的比较老）:

每一个SM上最多同时执行8个block。(不管block大小)

每一个SM上最多同时执行48个warp。

每一个SM上最多同时执行48*32=1,536个线程。

当warp访问内存的时候，processor(处理器)会做context switch(上下文切换)，让其他warp使用硬件资源。因为是硬件

来做，所以速度非常快。

2.2.2 网格（Grid）、线程块（Block）和线程（Thread）的组织关系以及线程索引的计算公式

2.2.2.1 格（Grid）、线程块（Block）和线程（Thread）的组织关系

CUDA的软件架构由网格（Grid）、线程块（Block）和线程（Thread）组成，相当于把GPU上的计算单元分为若干（2~3）个网格，每个网格内包含若干（65535）个线程块，每个线程块包含若干（512/1024）个线程，三者的关系如

下图：

Thread，block，grid是CUDA编程上的概念，为了方便程序员软件设计，组织线程。

thread：一个CUDA的并行程序会被以许多个threads来执行。

block：数个threads会被群组成一个block，同一个block中的threads可以同步，也可以通过shared memory通信。

grid：多个blocks则会再构成grid。

2.2.2.2 网格（Grid）、线程块（Block）和线程（Thread）的最大数量

CUDA中可以创建的网格数量跟GPU的计算能力有关，可创建的Grid、Block和Thread的最大数量参看以下表格：

在单一维度上，程序的执行可以由多达365535512=100661760（一亿）个线程并行执行，这对在CPU上创建并行线程来说是不可想象的。

2.2.2.3 线程索引的计算公式

一个Grid可以包含多个Blocks，Blocks的组织方式可以是一维的，二维或者三维的。block包含多个Threads，这些

Threads的组织方式也可以是一维，二维或者三维的。

CUDA中每一个线程都有一个唯一的标识ID—ThreadIdx，这个ID随着Grid和Block的划分方式的不同而变化，这里给出

Grid和Block不同划分方式下线程索引ID的计算公式。

threadIdx是一个uint3类型，表示一个线程的索引。

blockIdx是一个uint3类型，表示一个线程块的索引，一个线程块中通常有多个线程。

blockDim是一个dim3类型，表示线程块的大小。

gridDim是一个dim3类型，表示网格的大小，一个网格中通常有多个线程块。

下面这张图比较清晰的表示的几个概念的关系：

1、 grid划分成1维，block划分为1维

int threadId = blockIdx.x *blockDim.x + threadIdx.x;

2、 grid划分成1维，block划分为2维

int threadId = blockIdx.x blockDim.x blockDim.y + threadIdx.y * blockDim.x + threadIdx.x;

3、 grid划分成1维，block划分为3维

int threadId = blockIdx.x blockDim.x blockDim.y * blockDim.z

• threadIdx.z blockDim.y blockDim.x

• threadIdx.y * blockDim.x + threadIdx.x;

4、 grid划分成2维，block划分为1维

int blockId = blockIdx.y * gridDim.x + blockIdx.x;

int threadId = blockId * blockDim.x + threadIdx.x;

5、 grid划分成2维，block划分为2维

int blockId = blockIdx.x + blockIdx.y * gridDim.x;

int threadId = blockId (blockDim.x blockDim.y） + (threadIdx.y * blockDim.x) + threadIdx.x;

6、 grid划分成2维，block划分为3维

int blockId = blockIdx.x + blockIdx.y * gridDim.x;

int threadId = blockId (blockDim.x blockDim.y blockDim.z) + (threadIdx.z (blockDim.x blockDim.y)) + (threadIdx.y blockDim.x) + threadIdx.x;

7、 grid划分成3维，block划分为1维

int blockId = blockIdx.x + blockIdx.y gridDim.x + gridDim.x gridDim.y * blockIdx.z;

int threadId = blockId * blockDim.x + threadIdx.x;

8、 grid划分成3维，block划分为2维

int blockId = blockIdx.x + blockIdx.y gridDim.x + gridDim.x gridDim.y * blockIdx.z;

int threadId = blockId (blockDim.x blockDim.y) + (threadIdx.y * blockDim.x) + threadIdx.x;

9、 grid划分成3维，block划分为3维

int blockId = blockIdx.x + blockIdx.y gridDim.x + gridDim.x gridDim.y * blockIdx.z;

int threadId = blockId (blockDim.x blockDim.y blockDim.z) + (threadIdx.z (blockDim.x blockDim.y)) + (threadIdx.y blockDim.x) + threadIdx.x;

cuda 通过<<< >>>符号来分配索引线程的方式，至少有15种索引方式。

下面程序展示了这15种索引方式：

#include "cuda_runtime.h"
#include "device_launch_parameters.h"
​
#include <stdio.h>
#include <stdlib.h>
#include <iostream>
​
using namespace std;
​
//thread 1D
__global__ void testThread1(int *c, const int *a, const int *b)
{
int i = threadIdx.x;
c[i] = b[i] - a[i];
}
​
//thread 2D
__global__ void testThread2(int *c, const int *a, const int *b)
{
int i = threadIdx.x + threadIdx.y*blockDim.x;
c[i] = b[i] - a[i];
}
​
//thread 3D
__global__ void testThread3(int *c, const int *a, const int *b)
{
int i = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;
c[i] = b[i] - a[i];
}
​
//block 1D
__global__ void testBlock1(int *c, const int *a, const int *b)
{
​
int i = blockIdx.x;
c[i] = b[i] - a[i];
}
​
//block 2D
__global__ void testBlock2(int *c, const int *a, const int *b)
{
int i = blockIdx.x + blockIdx.y*gridDim.x;
c[i] = b[i] - a[i];
}
​
//block 3D
__global__ void testBlock3(int *c, const int *a, const int *b)
{
int i = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;
c[i] = b[i] - a[i];
}
​
//block-thread 1D-1D
__global__ void testBlockThread1(int *c, const int *a, const int *b)
{
int i = threadIdx.x + blockDim.x*blockIdx.x;
c[i] = b[i] - a[i];
}
​
//block-thread 1D-2D
__global__ void testBlockThread2(int *c, const int *a, const int *b)
{
int threadId_2D = threadIdx.x + threadIdx.y*blockDim.x;
int i = threadId_2D+ (blockDim.x*blockDim.y)*blockIdx.x;
c[i] = b[i] - a[i];
}
​
//block-thread 1D-3D
__global__ void testBlockThread3(int *c, const int *a, const int *b)
{
int threadId_3D = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;
int i = threadId_3D + (blockDim.x*blockDim.y*blockDim.z)*blockIdx.x;
c[i] = b[i] - a[i];
}
​
//block-thread 2D-1D
__global__ void testBlockThread4(int *c, const int *a, const int *b)
{
int blockId_2D = blockIdx.x + blockIdx.y*gridDim.x;
int i = threadIdx.x + blockDim.x*blockId_2D;
c[i] = b[i] - a[i];
}
​
//block-thread 3D-1D
__global__ void testBlockThread5(int *c, const int *a, const int *b)
{
int blockId_3D = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;
int i = threadIdx.x + blockDim.x*blockId_3D;
c[i] = b[i] - a[i];
}
​
//block-thread 2D-2D
​
​
__global__ void testBlockThread6(int *c, const int *a, const int *b)
{
int threadId_2D = threadIdx.x + threadIdx.y*blockDim.x;
int blockId_2D = blockIdx.x + blockIdx.y*gridDim.x;
int i = threadId_2D + (blockDim.x*blockDim.y)*blockId_2D;
c[i] = b[i] - a[i];
}
​
//block-thread 2D-3D
__global__ void testBlockThread7(int *c, const int *a, const int *b)
{
int threadId_3D = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;
int blockId_2D = blockIdx.x + blockIdx.y*gridDim.x;
int i = threadId_3D + (blockDim.x*blockDim.y*blockDim.z)*blockId_2D;
c[i] = b[i] - a[i];
}
​
//block-thread 3D-2D
__global__ void testBlockThread8(int *c, const int *a, const int *b)
{
int threadId_2D = threadIdx.x + threadIdx.y*blockDim.x;
int blockId_3D = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;
int i = threadId_2D + (blockDim.x*blockDim.y)*blockId_3D;
c[i] = b[i] - a[i];
}
​
//block-thread 3D-3D
__global__ void testBlockThread9(int *c, const int *a, const int *b)
{
int threadId_3D = threadIdx.x + threadIdx.y*blockDim.x + threadIdx.z*blockDim.x*blockDim.y;
int blockId_3D = blockIdx.x + blockIdx.y*gridDim.x + blockIdx.z*gridDim.x*gridDim.y;
int i = threadId_3D + (blockDim.x*blockDim.y*blockDim.z)*blockId_3D;
c[i] = b[i] - a[i];
}
​
void addWithCuda(int *c, const int *a, const int *b, unsigned int size)
{
int *dev_a = 0;
int *dev_b = 0;
int *dev_c = 0;
​
cudaSetDevice(0);
​
cudaMalloc((void**)&dev_c, size * sizeof(int));
cudaMalloc((void**)&dev_a, size * sizeof(int));
cudaMalloc((void**)&dev_b, size * sizeof(int));
​
cudaMemcpy(dev_a, a, size * sizeof(int), cudaMemcpyHostToDevice);
cudaMemcpy(dev_b, b, size * sizeof(int), cudaMemcpyHostToDevice);
​
//testThread1<<<1, size>>>(dev_c, dev_a, dev_b);
​
//uint3 s;s.x = size/5;s.y = 5;s.z = 1;
//testThread2 <<<1,s>>>(dev_c, dev_a, dev_b);
​
//uint3 s; s.x = size / 10; s.y = 5; s.z = 2;
//testThread3<<<1, s >>>(dev_c, dev_a, dev_b);
​
​
//testBlock1<<<size,1 >>>(dev_c, dev_a, dev_b);
​
//uint3 s; s.x = size / 5; s.y = 5; s.z = 1;
//testBlock2<<<s, 1 >>>(dev_c, dev_a, dev_b);
​
//uint3 s; s.x = size / 10; s.y = 5; s.z = 2;
//testBlock3<<<s, 1 >>>(dev_c, dev_a, dev_b);
​
//testBlockThread1<<<size/10, 10>>>(dev_c, dev_a, dev_b);
​
//uint3 s1; s1.x = size / 100; s1.y = 1; s1.z = 1;
//uint3 s2; s2.x = 10; s2.y = 10; s2.z = 1;
//testBlockThread2 << <s1, s2 >> >(dev_c, dev_a, dev_b);
​
//uint3 s1; s1.x = size / 100; s1.y = 1; s1.z = 1;
//uint3 s2; s2.x = 10; s2.y = 5; s2.z = 2;
//testBlockThread3 << <s1, s2 >> >(dev_c, dev_a, dev_b);
​
//uint3 s1; s1.x = 10; s1.y = 10; s1.z = 1;
//uint3 s2; s2.x = size / 100; s2.y = 1; s2.z = 1;
//testBlockThread4 << <s1, s2 >> >(dev_c, dev_a, dev_b);
​
//uint3 s1; s1.x = 10; s1.y = 5; s1.z = 2;
//uint3 s2; s2.x = size / 100; s2.y = 1; s2.z = 1;
//testBlockThread5 << <s1, s2 >> >(dev_c, dev_a, dev_b);
​
//uint3 s1; s1.x = size / 100; s1.y = 10; s1.z = 1;
//uint3 s2; s2.x = 5; s2.y = 2; s2.z = 1;
//testBlockThread6 << <s1, s2 >> >(dev_c, dev_a, dev_b);
​
//uint3 s1; s1.x = size / 100; s1.y = 5; s1.z = 1;
//uint3 s2; s2.x = 5; s2.y = 2; s2.z = 2;
//testBlockThread7 << <s1, s2 >> >(dev_c, dev_a, dev_b);
​
//uint3 s1; s1.x = 5; s1.y = 2; s1.z = 2;
//uint3 s2; s2.x = size / 100; s2.y = 5; s2.z = 1;
//testBlockThread8 <<<s1, s2 >>>(dev_c, dev_a, dev_b);
​
uint3 s1; s1.x = 5; s1.y = 2; s1.z = 2;
uint3 s2; s2.x = size / 200; s2.y = 5; s2.z = 2;
testBlockThread9<<<s1, s2 >>>(dev_c, dev_a, dev_b);
​
cudaMemcpy(c, dev_c, size*sizeof(int), cudaMemcpyDeviceToHost);
​
cudaFree(dev_a);
cudaFree(dev_b);
cudaFree(dev_c);
​
cudaGetLastError();
}
​
int main()
{
const int n = 1000;
​
int *a = new int[n];
​
int *b = new int[n];
int *c = new int[n];
int *cc = new int[n];
​
for (int i = 0; i < n; i++)
{
a[i] = rand() % 100;
b[i] = rand() % 100;
c[i] = b[i] - a[i];
}
​
addWithCuda(cc, a, b, n);
​
FILE *fp = fopen("out.txt", "w");
for (int i = 0; i < n; i++)
fprintf(fp, "%d %d\n", c[i], cc[i]);
fclose(fp);
​
bool flag = true;
for (int i = 0; i < n; i++)
{
if (c[i] != cc[i])
{
flag = false;
break;
}
}
​
if (flag == false)
printf("no pass");
else
printf("pass");
​
cudaDeviceReset();
​
delete[] a;
delete[] b;
delete[] c;
delete[] cc;
​
getchar();
return 0;
}

这里只保留了3D-3D方式，注释了其余14种方式，所有索引方式均测试通过。对于后续要深入cuda的同学可以仔细研究。

2.2.3 CUDA程序结构

CUDA程序的结构大体是：{主机串行->GPU并行}+ -> 主机串行，这样的串并交叉结构。主机串行过渡到GPU并行时需要将数据从主机内存上拷贝到GPU设备内存上，GPU执行完毕时也需要把数据拷贝回来。

2.2.4 CUDA内核函数和配置

主机调用设备代码的唯一接口就是Kernel函数，使用限定符:global。

调用内核函数需要在内核函数名后添加<<<>>>指定内核函数配置，<<<>>>运算符完整的执行配置参数形式是<<<Dg, Db, Ns, S>>>

• 参数Dg用于定义整个grid的维度和尺寸，即一个grid有多少个block。为dim3类型。Dim3 Dg(Dg.x, Dg.y, 1)表示grid中每行有Dg.x个block，每列有Dg.y个block，第三维恒为1(目前一个核函数只有一个grid)。整个grid中共有Dg.x*Dg.y个block，其中Dg.x和Dg.y最大值为65535。

• 参数Db用于定义一个block的维度和尺寸，即一个block有多少个thread。为dim3类型。Dim3 Db(Db.x, Db.y, Db.z)表示整个block中每行有Db.x个thread，每列有Db.y个thread，高度为Db.z。Db.x和Db.y最大值为512，Db.z最大值为62。 一个block中共有Db.xDb.yDb.z个thread。计算能力为1.0,1.1的硬件该乘积的最大值为768，计算能力为1.2,1.3的硬件支持的最大值为1024。

• 参数Ns是一个可选参数，用于设置每个block除了静态分配的shared Memory以外，最多能动态分配的sharedmemory大小，单位为byte。不需要动态分配时该值为0或省略不写。

• 参数S是一个cudaStream_t类型的可选参数，初始值为零，表示该核函数处在哪个流之中。

如<<<DimGrid, DimBlock>>>指定线程网络和线程块维度。若当前硬件无法满足用户配置，则内

核函数不会被执行，直接返回错误。

2.2.5 CUDA限定符

函数限定符：（默认host、global异步、主机不能调device；设备上执行的函数参数数目固定、不能声明静态变量且不支持递归调用）

变量限定符：（shared共享一致性必须由显式线程同步保证）

2.2.6 同步

• CPU启动kernel函数是异步的，它并不会阻塞等到GPU执行完kernel函数才执行后面的CPU部分，因此如果后续程序立即需要用到上一个kernel函数的结果我们需要显式设置同步障来阻塞CPU程序。

• 一个线程块内需要同步共享存储器的共享变量（shared）时，需要在使用前显式调用__syncthreads()同步块内所有线程。

• 同一个Grid中不同Block之间无法设置同步。

2.2.7 CUDA运行时API

这里介绍最基础的内存管理函数，其他详见官网：http://docs.nvidia.com/cuda/cuda-runtime-api/index.html

cudaMemcpy

__host__ cudaError_t cudaMemcpy( void* dst, const void* src, size_t count, cudaMemcpyKind kind )

用于在主机和设备之间拷贝数据，其中cudaMemcpyKind枚举类型常用有cudaMemcpyHostToDevice表示把主机数据拷贝到内存以及逆向的cudaMemcpyDeviceToHost。

cudaMalloc

__host__ __device__ cudaError_t cudaMalloc( void** devPtr, size_t size )

在设备上分配动态内存，两个限定符表示可以在主机或设备上调用。

cudaFree

__host__ __device__ cudaError_t cudaFree( void* devPtr )

释放回收在设备上分配动态内存，两个限定符表示c可以在主机或设备上调用。

3 GPU应用领域

GPU适用于深度学习训练和推理，图像识别、语音识别等；计算金融学、地震分析、分子建模、基

因组学、计算流体动力学等；高清视频转码、安防视频监控、大型视频会议等；三维设计与渲染、

影音动画制作、工程建模与仿真（CAD/CAE）、医学成像、游戏测试等等。

GPU常见的应用领域如下所示：

1. 游戏：GeForce RTX/GTX系列GPU（PCs）、GeForce NOW（云游戏）、SHIELD（游戏主机）

2. 专业可视化：Quadro/RTX GPU（企业工作站）

3. 数据中心：基于GPU的计算平台和系统，包括DGX（AI服务器）、HGX（超算）、EGX（边缘计算）、AGX（自动设备）， 天气预报的预测 （精确到几百米）

4. 汽车：NVIDIA DRIVE计算平台，包括AGX Xavier（SoC芯片）、DRIVE AV（自动驾驶）、DRIVE IX（驾驶舱软件）、Constellation（仿真软件）

5. 消费电子：智能手机市场占据了全球GPU市场份额的主导地位，此外，智能音箱、智能手环/手表、VR/AR眼镜等移动消费电子都是GPU潜在的市场。比如拍照、导航地图的合成、UI图标、图像框、照片的后处理等都需要GPU来完成。

更详细的应用场景参考：一文看完GPU八大应用场景，抢食千亿美元市场。

GPU算力 TOPs： OPS是Tera Operations Per Second的缩写，1TOPS代表处理器每秒钟可进行一万亿次（10^12）操作。

这里简述部分应用场景：

边缘计算的应用场景

AI芯片主要承担推断任务，通过将终端设备上的传感器（麦克风阵列、摄像头等）收集的数据代入训练好的模型推理得出推断结果。由于边缘侧场景多种多样、各不相同，对于计算硬件的考量也不尽相同，对于算力和能耗等性能需求也有大有小。因此应用于边缘侧的计算芯片需要针对特殊场景进行针对性设计以实现最优的解决方案。

自动驾驶场景：

地平线公司CEO余凯曾在公开场合指出，“自动驾驶每提高一级，算力就增加一个数量级。L2级别大概需要2个TOPS的算力，L3需要24个TOPS，L4为320TOPS，L5为4000+TOPS。”

蔚来ET7集成了4颗英伟达 Orin 芯片，总算力1016TOPS，超过7个特斯拉FSD的算力总和。其系统中，两颗Orin芯片为主控芯片，负责自动驾驶系统的运算；1颗Orin芯片为实时冗余备份芯片，用于保证自动驾驶的安全性；另外还有1颗Orin芯片为群体智能与个性训练专用芯片，从而帮助蔚来的算法进行整体升级和单车个性化本地训练。据称，能实现高速、城区、泊车、换电场景全覆盖。

4 GPU+CPU异构计算

异构计算从常见的搭配有CPU+GPU、CPU+FPGA、CPU+DSP(多指令，矩阵乘法算子)，CPU +ASIC(专用集成电路， 阿里云硬件编解码，比GPU更专业 )等。

CPU的核心少但每一个核心的控制和计算能力都不弱，因此常作为主机。而GPU的计算核心很多，所以当遇到大数据量且逻辑简单的任务，CPU就会交给GPU来进行计算，同时CPU的核心虽少但也是有多个线程的，多线程可以调度并同时控制多张GPU同时完成多个任务，这本身也是一种并行思想，并且GPU也可以在接收到任务后让CPU的线程先去处理别的事情完成异步控制来进一步提高效率（这本质上也是一种时域上的并行）。

之所以出现GPU+CPU异构计算，因为CPU和GPU各自有优缺点：

CPU 适用于一系列广泛的工作负载，特别是那些对于延迟和单位内核性能要求较高的工作负载。作为强大的执行引擎，CPU 将它数量相对较少的内核集中用于处理单个任务，并快速将其完成。这使它尤其适合用于处理从串行计算到数据库运行等类型的工作。

GPU 最初是作为专门用于加速特定 3D 渲染任务的 ASIC 开发而成的。随着时间的推移，这些功能固定的引擎变得更加可编程化、更加灵活。尽管图形处理和当下视觉效果越来越真实的顶级游戏仍是 GPU 的主要功能，但同时，它也已经演化为用途更普遍的并行处理器，能够处理越来越多的应用程序。

5 CUDA环境搭建

参考：

Windows环境：https://www.yuque.com/linuxer/xngi03/ogz3ignhirrf8qgx?singleDoc# 《win10+cuda11.0+vs2019

安装教程》

Ubuntu环境：https://cloud.tencent.com/document/product/560/76423

特别需要注意：cuda不能在虚拟机下运行，需要在物理机运行。如果自己物理机是Windows的系统，建议先在Windows

环境进行测试（wsl环境也可以安装）。

6 CUDA向量加法

6.1 向量加法CPU实现

int i = 0;
for (i = 0; i < size; i += 1)
{
​
res[i] = a[i] + b[i];
}

6.2 向量加法GPU实现

global__ void addKernel(int *c, const int *a, const int *b)
​
{
​
int i = threadIdx.x;
​
c[i] = a[i] + b[i];
​
}

用cuda计算向量加法A+B=C

流程：

1.申请主机内存。向量A，向量B，计算结果C

2.初始化数据。用0-1之间的随机数初始化向量A，B，C

3.GPU内存申请。申请A，B，C需要的GPU内存空间

4.数据拷贝。把数据从主机内存拷贝至GPU内存

5.计算需要的线程数和线程块数。

6.调用GPU加法函数

7.数据拷贝。把结果从GPU内存拷贝至主机内存。

8.在CPU上重新运行一遍，与GPU结果进行对照。

9.释放GPU内存。

10.释放主机内存。

11.重置GPU状态。

要点：内存管理，数据拷贝。

sum_tutorial_timer.cu

#include <cuda_runtime.h>
#include <stdio.h>
#include "freshman.h"
​
#define LOOP_NUM 1
​
// CPU 加法
void sumArraysCPU(float* a, float* b, float* res, const int size, int loop)
{
for (int j = 0; j < loop; j++) {
int i = 0;
for (i = 0; i < size; i += 1)
{
res[i] = a[i] + b[i];
}
//printf("CPU res[%d] = %f\n", i - 1, res[i-1]);
}
}
​
// GPU 加法
__global__ void sumArraysGPU(float* a, float* b, float* res, int N, int loop)
{
for (int j = 0; j < loop; j++) {
int i = blockIdx.x * blockDim.x + threadIdx.x;
if (i < N)
res[i] = a[i] + b[i];
printf("GPU res[%d] = %f\n",i, res[i]);
}
}
// 测试多个循环，更容易分析gpu的性能影响
int main(int argc, char** argv)
{
// set up device 选择设备
initDevice(0);
​
//初始化向量维度
//int nElem = 512 * 512; // 262,144
// int nElem = 1024*1024;
int nElem = 2048*2048; // 4,194,304
​
printf("Vector addition of %d elements\n", nElem);
​
// 内存数据申请空间
int nByte = sizeof(float) * nElem;
​
printf("CPU nedd all memory size: %d\n", nByte *4);
​
// 对主机的a, b, res申请内存空间
float* a_h = (float*)malloc(nByte);
float* b_h = (float*)malloc(nByte);
float* res_h = (float*)malloc(nByte);
// 申请从gpu拷贝vector回主机的内存空间
float* res_from_gpu_h = (float*)malloc(nByte);
// 初始化为0
memset(res_h, 0, nByte);
memset(res_from_gpu_h, 0, nByte);
​
// 内存数据随机初始化 向量a和b的数据
initialData(a_h, nElem);
initialData(b_h, nElem);
​
// GPU显存申请空间
​
​
printf("GPU nedd all memory size: %d\n", nByte * 3);
float *a_d = NULL;
float *b_d = NULL;
float *res_d = NULL;
CHECK(cudaMalloc((float**)&a_d, nByte));
CHECK(cudaMalloc((float**)&b_d, nByte));
CHECK(cudaMalloc((float**)&res_d, nByte));
​
// 内存到显存数据拷贝
double iStart = 0;
double iElaps = 0;
iStart = cpuSecond();
CHECK(cudaMemcpy(a_d, a_h, nByte, cudaMemcpyHostToDevice));
CHECK(cudaMemcpy(b_d, b_h, nByte, cudaMemcpyHostToDevice));
iElaps = cpuSecond() - iStart;
printf("内存数据拷贝到GPU时间消耗\t%lf sec\n", iElaps);
​
//dim3 block(512);
//dim3 grid((nElem - 1) / block.x + 1);
​
//计算线程块与线程
#if 1
//每线程块线程数
int threadsPerBlock = 16;
//每网格线程块数
int blocksPerGrid = (nElem + threadsPerBlock - 1) / threadsPerBlock;
#else
int threadsPerBlock = 4; // 如果这么简单设置将只能测试部分结果
int blocksPerGrid = 4; // 比如4x4=16 其实只有16个向量进行加法
#endif
// printf("% Loop, CUDA kernel launch with %d blocks of %d threads\n", LOOP_NUM, blocksPerGrid,
threadsPerBlock);
// GPU 加法
iStart = cpuSecond();
// sumArraysGPU << <grid, block >> > (a_d, b_d, res_d, nElem, LOOP_NUM);
sumArraysGPU << <blocksPerGrid, threadsPerBlock >> > (a_d, b_d, res_d, nElem, LOOP_NUM);
iElaps = cpuSecond() - iStart;
printf("GPU计算时间 \t\t\t\t %lf sec\n", iElaps);
​
//显存到内存数据拷贝
iStart = cpuSecond();
CHECK(cudaMemcpy(res_from_gpu_h, res_d, nByte, cudaMemcpyDeviceToHost));
iElaps = cpuSecond() - iStart;
printf("结果从显存拷贝到内存时间消耗 %lf sec\n", iElaps);
​
// CPU 加法
iStart = cpuSecond();
sumArraysCPU(a_h, b_h, res_h, nElem, LOOP_NUM);
iElaps = cpuSecond() - iStart;
printf("CPU 计算时间\t\t\t\t %lf sec\n", iElaps);
​
// 检测GPU和CPU的计算结果
printf("检测GPU和CPU的计算结果是否相同\n");
checkResult(res_h, res_from_gpu_h, nElem);
cudaFree(a_d);
cudaFree(b_d);
cudaFree(res_d);
​
free(a_h);
free(b_h);
free(res_h);
free(res_from_gpu_h);
​
return 0;
}

freshman.h

#ifndef FRESHMAN_H
#define FRESHMAN_H
#define CHECK(call)\
{\
const cudaError_t error=call;\
if(error!=cudaSuccess)\
{\
printf("ERROR: %s:%d,",__FILE__,__LINE__);\
printf("code:%d,reason:%s\n",error,cudaGetErrorString(error));\
exit(1);\
}\
}
​
#include <stdio.h>
#include <time.h>
#ifdef _WIN32
​
include <windows.h>
​
#else
​
include <sys/time.h>
​
#endif
#ifdef _WIN32
int gettimeofday(struct timeval* tp, void* tzp)
{
time_t clock;
struct tm tm;
SYSTEMTIME wtm;
GetLocalTime(&wtm);
tm.tm_year = wtm.wYear - 1900;
tm.tm_mon = wtm.wMonth - 1;
tm.tm_mday = wtm.wDay;
tm.tm_hour = wtm.wHour;
tm.tm_min = wtm.wMinute;
tm.tm_sec = wtm.wSecond;
tm.tm_isdst = -1;
clock = mktime(&tm);
tp->tv_sec = clock;
tp->tv_usec = wtm.wMilliseconds * 1000;
return (0);
}
#endif
double cpuSecond()
{
struct timeval tp;
gettimeofday(&tp, NULL);
return((double)tp.tv_sec + (double)tp.tv_usec * 1e-6);
​
}
void initialData(float* ip, int size)
{
time_t t;
srand((unsigned)time(&t));
for (int i = 0; i < size; i++)
{
ip[i] = (float)(rand() & 0xffff) / 1000.0f;
}
}
void initialData_int(int* ip, int size)
{
time_t t;
srand((unsigned)time(&t));
for (int i = 0; i < size; i++)
{
ip[i] = int(rand() & 0xff);
}
}
void printMatrix(float* C, const int nx, const int ny)
{
float* ic = C;
printf("Matrix<%d,%d>:", ny, nx);
for (int i = 0; i < ny; i++)
{
for (int j = 0; j < nx; j++)
{
printf("%6f ", C[j]);
}
ic += nx;
printf("\n");
}
}
​
void initDevice(int devNum)
{
int dev = devNum;
cudaDeviceProp deviceProp;
CHECK(cudaGetDeviceProperties(&deviceProp, dev));
printf("Using device %d: %s\n", dev, deviceProp.name);
CHECK(cudaSetDevice(dev));
​
}
void checkResult(float* hostRef, float* gpuRef, const int N)
{
double epsilon = 1.0E-8;
for (int i = 0; i < N; i++)
{
if (abs(hostRef[i] - gpuRef[i]) > epsilon)
{
printf("Results don\'t match!\n");
printf("%f(hostRef[%d] )!= %f(gpuRef[%d])\n", hostRef[i], i, gpuRef[i], i);
return;
}
}
printf("Check result success!\n");
}
#endif//FRESHMAN_H

参考：CUDA编程入门：向量加法和矩阵乘法 https://blog.csdn.net/u014030117/article/details/45952971

7 CUDA矩阵乘法

7.1 矩阵乘法CPU实现

比如（两个矩阵要是想相乘需要满足，第一个矩阵的列数等于第二个矩阵的行数）：

先观察矩阵乘法串行的算法如下：

void matrixMulCpu(float* A, float* B, float* C, int width){
float sum = 0.0f;
for(int i = 0; i < width; i++){
for(int j = 0; j < width; j++){
for(int l = 0; l < width; l++){
sum += A[i * width + l] * B[l * width + j];
}
C[i * width + j] = sum;
sum = 0.0f;
}
}
}

• 通过上面CPU代码的实验观察可以看出，总共的计算次数为：m n k

• 时间复杂度为：O(N3 )

7.2 矩阵乘法GPU实现

获得 C 矩阵的计算方法都是相同的，只不过使用的是矩阵 A、B 不同的元素来进行计算，即不同数据的大量相同计算操作，这种计算是特别适合使用GPU来计算，因为GPU拥有大量简单重复的计算单元，通过并行就能极大的提高计算效率。

在 GPU 中执行矩阵乘法运算操作：

1. 在 Global Memory 中分别为矩阵 A、B、C 分配存储空间；

2. 由于矩阵 C 中每个元素的计算均相互独立，NVIDIA GPU 采用的 SIMT (单指令多线程)的体系结构来实现并行计算的, 因此在并行度映射中，让每个 thread 对应矩阵 C 中1个元素的计算；

3. 执行配置 (execution configuration)中 gridSize 和 blockSize 均有 x(列向)、y(行向)两个维度，

   

4. 其中，CUDA的kernel函数实现如下：

   

每个 thread 需要执行的 workflow 为：

• 从矩阵 A 中读取一行向量 (长度为width) ==> A[row * width + i]

• 从矩阵 B 中读取一列向量 (长度为width（图中为height）) ==> B[i * width + col]

• 对这两个向量做点积运算 (单层 width 次循环的乘累加)==> A[row width + i] B[i * width + col]

• 最后将结果写回矩阵 C。==> C[row * width + col] = Pervalue

__global__ void multiply(int* A, int* B, int* C, int width) {
int bx = blockIdx.x;
int by = blockIdx.y;
int tx = threadIdx.x;
int ty = threadIdx.y;
​
int col = bx * blockDim.x + tx;
int row = by * blockDim.y + ty;
​
int perValue = 0;
if (row < width && col < width) {
for (int i = 0; i < width; i++) {
perValue += A[row * width + i] * B[i * width + col];
}
C[row * width + col] = perValue;
}
}

源代码如下：

multiply_tutorial_timer.cu

#include <stdio.h>
#include <stdlib.h>
#include <time.h>
#include <cuda.h>
#include <cuda_runtime.h>
#include "freshman.h"
#define RANDOM(x) (rand() % x)
​
#define MAX 100000
​
#define BLOCKSIZE 16
​
__global__ void multiply(int* A, int* B, int* C, int width) {
int bx = blockIdx.x; //块的序号
int by = blockIdx.y;
int tx = threadIdx.x;
int ty = threadIdx.y;
​
int col = bx * blockDim.x + tx;
int row = by * blockDim.y + ty;
printf("c:%d, r:%d,blockDim:(%d,%d), blockIdx:(%d,%d), threadIdx:(%d,%d)\n",col, row, blockDim.x, blockDim.y,
blockIdx.x, blockIdx.y, threadIdx.x, threadIdx.y);
​
int perValue = 0;
if (row < width && col < width) {
for (int i = 0; i < width; i++) {
perValue += A[row * width + i] * B[i * width + col];
}
C[row * width + col] = perValue;
} else {
printf("> col:%d or row:%d\n", col, row);
}
}
​
int main(int argc, char** argv) {
int n = 24; // 16/ 32/24测试对比
​
int i, j, k;
double start, finish;
​
if (argc == 2) {
n = atoi(argv[1]);
}
​
int* host_a = (int*)malloc(sizeof(int) * n * n);
int* host_b = (int*)malloc(sizeof(int) * n * n);
int* host_c = (int*)malloc(sizeof(int) * n * n);
int* host_c2 = (int*)malloc(sizeof(int) * n * n);
​
srand(time(NULL));
​
for (i = 0; i < n * n; i++) {
host_a[i] = RANDOM(MAX);
host_b[i] = RANDOM(MAX);
}
​
cudaError_t error = cudaSuccess;
​
int* device_a, * device_b, * device_c;
error = cudaMalloc((void**)&device_a, sizeof(int) * n * n);
error = cudaMalloc((void**)&device_b, sizeof(int) * n * n);
error = cudaMalloc((void**)&device_c, sizeof(int) * n * n);
​
if (error != cudaSuccess) {
printf("Fail to cudaMalloc on GPU");
return 1;
}
​
//GPU parallel start
start = cpuSecond();
cudaMemcpy(device_a, host_a, sizeof(int) * n * n, cudaMemcpyHostToDevice);
cudaMemcpy(device_b, host_b, sizeof(int) * n * n, cudaMemcpyHostToDevice);
​
// ceil 向上取整；pow(float x, float y)即x的y次幂；sqrt平方根
double num = ceil(pow((double)n, 2) / pow((double)BLOCKSIZE, 2));
int gridsize = (int)ceil(sqrt(num)); // 主要等下我们用二维block
​
dim3 dimBlock(BLOCKSIZE, BLOCKSIZE, 1); // 16*16*1个线程
dim3 dimGrid(gridsize, gridsize, 1); // gridsize * gridsize *1个block
printf("n:%d, dimGrid(%d,%d,%d), dimBlock(%d,%d,%d), total_threads:%d\n",n, dimGrid.x, dimGrid.y, dimGrid.z,
dimBlock.x, dimBlock.y, dimBlock.z,
dimGrid.x*dimGrid.y*dimGrid.z*dimBlock.x*dimBlock.y*dimBlock.z);
multiply << <dimGrid, dimBlock >> > (device_a, device_b, device_c, n);
cudaThreadSynchronize();
​
cudaMemcpy(host_c, device_c, sizeof(int) * n * n, cudaMemcpyDeviceToHost);
​
finish = cpuSecond();
double t = finish - start;
printf("GPU %lf s\n", t);
//GPU parallel finish
​
//CPU serial start
start = cpuSecond();
​
for (i = 0; i < n; i++) {
for (j = 0; j < n; j++) {
host_c2[i * n + j] = 0;
for (k = 0; k < n; k++) {
host_c2[i * n + j] += host_a[i * n + k] * host_b[k * n + j];
}
}
}
​
finish = cpuSecond();
​
t = finish - start;
printf("CPU %lf s\n", t);
//CPU serial start
​
//check
int errorNum = 0;
for (int i = 0; i < n * n; i++) {
if (host_c[i] != host_c2[i]) {
errorNum++;
printf("Error occurs at index: %d: c = %d, c2 = %d\n", i, host_c[i], host_c2[i]);
}
}
if (errorNum == 0) {
printf("Successfully run on GPU and CPU!\n");
}
else {
printf("%d error(s) occurs!\n", errorNum);
}
​
free(host_a);
free(host_b);
free(host_c);
free(host_c2);
​
cudaFree(device_a);
cudaFree(device_b);
cudaFree(device_c);
​
return 0;
}

freshman.h

#ifndef FRESHMAN_H
#define FRESHMAN_H
#define CHECK(call)\
{\
const cudaError_t error=call;\
if(error!=cudaSuccess)\
{\
printf("ERROR: %s:%d,",__FILE__,__LINE__);\
printf("code:%d,reason:%s\n",error,cudaGetErrorString(error));\
exit(1);\
}\
}
​
#include <stdio.h>
#include <time.h>
#ifdef _WIN32
include <windows.h>
​
#else
​
include <sys/time.h>
​
#endif
#ifdef _WIN32
int gettimeofday(struct timeval* tp, void* tzp)
{
time_t clock;
struct tm tm;
SYSTEMTIME wtm;
GetLocalTime(&wtm);
tm.tm_year = wtm.wYear - 1900;
tm.tm_mon = wtm.wMonth - 1;
tm.tm_mday = wtm.wDay;
tm.tm_hour = wtm.wHour;
tm.tm_min = wtm.wMinute;
tm.tm_sec = wtm.wSecond;
tm.tm_isdst = -1;
clock = mktime(&tm);
tp->tv_sec = clock;
tp->tv_usec = wtm.wMilliseconds * 1000;
return (0);
}
#endif
double cpuSecond()
{
struct timeval tp;
gettimeofday(&tp, NULL);
return((double)tp.tv_sec + (double)tp.tv_usec * 1e-6);
​
}
void initialData(float* ip, int size)
{
time_t t;
srand((unsigned)time(&t));
for (int i = 0; i < size; i++)
{
ip[i] = (float)(rand() & 0xffff) / 1000.0f;
}
}
void initialData_int(int* ip, int size)
{
time_t t;
srand((unsigned)time(&t));
for (int i = 0; i < size; i++)
{
ip[i] = int(rand() & 0xff);
}
}
void printMatrix(float* C, const int nx, const int ny)
    {
float* ic = C;
printf("Matrix<%d,%d>:", ny, nx);
for (int i = 0; i < ny; i++)
{
for (int j = 0; j < nx; j++)
{
printf("%6f ", C[j]);
}
ic += nx;
printf("\n");
}
}
​
void initDevice(int devNum)
{
int dev = devNum;
cudaDeviceProp deviceProp;
CHECK(cudaGetDeviceProperties(&deviceProp, dev));
printf("Using device %d: %s\n", dev, deviceProp.name);
CHECK(cudaSetDevice(dev));
​
}
void checkResult(float* hostRef, float* gpuRef, const int N)
{
double epsilon = 1.0E-8;
for (int i = 0; i < N; i++)
{
if (abs(hostRef[i] - gpuRef[i]) > epsilon)
{
printf("Results don\'t match!\n");
printf("%f(hostRef[%d] )!= %f(gpuRef[%d])\n", hostRef[i], i, gpuRef[i], i);
return;
}
}
printf("Check result success!\n");
}
#endif//FRESHMAN_H

7.3 测试结果

使用不同的矩阵nxn大小进行测试。

n:16, dimGrid(1,1,1), dimBlock(16,16,1), total_threads:256

矩阵16x16, 一个grid有一个block，每个block的线程为(16x16)， 比如threadIdx:(1,8)，意思是该线程在当前block的位置为x=1, y=8。

n:16, dimGrid(1,1,1), dimBlock(16,16,1), total_threads:256
c:0, r:8,blockDim:(16,16), blockIdx:(0,0), threadIdx:(0,8)
c:1, r:8,blockDim:(16,16), blockIdx:(0,0), threadIdx:(1,8)
c:2, r:8,blockDim:(16,16), blockIdx:(0,0), threadIdx:(2,8)
c:3, r:8,blockDim:(16,16), blockIdx:(0,0), threadIdx:(3,8)
c:4, r:8,blockDim:(16,16), blockIdx:(0,0), threadIdx:(4,8)
c:5, r:8,blockDim:(16,16), blockIdx:(0,0), threadIdx:(5,8)
c:6, r:8,blockDim:(16,16), blockIdx:(0,0), threadIdx:(6,8)
c:7, r:8,blockDim:(16,16), blockIdx:(0,0), threadIdx:(7,8)
c:8, r:8,blockDim:(16,16), blockIdx:(0,0), threadIdx:(8,8)
c:9, r:8,blockDim:(16,16), blockIdx:(0,0), threadIdx:(9,8)
c:10, r:8,blockDim:(16,16), blockIdx:(0,0), threadIdx:(10,8)
.........
c:8, r:9,blockDim:(16,16), blockIdx:(0,0), threadIdx:(8,9)
c:9, r:9,blockDim:(16,16), blockIdx:(0,0), threadIdx:(9,9)
c:10, r:9,blockDim:(16,16), blockIdx:(0,0), threadIdx:(10,9)
c:11, r:9,blockDim:(16,16), blockIdx:(0,0), threadIdx:(11,9)
c:12, r:9,blockDim:(16,16), blockIdx:(0,0), threadIdx:(12,9)
c:13, r:9,blockDim:(16,16), blockIdx:(0,0), threadIdx:(13,9)
c:14, r:9,blockDim:(16,16), blockIdx:(0,0), threadIdx:(14,9)
c:15, r:9,blockDim:(16,16), blockIdx:(0,0), threadIdx:(15,9)
c:0, r:10,blockDim:(16,16), blockIdx:(0,0), threadIdx:(0,10)
c:1, r:10,blockDim:(16,16), blockIdx:(0,0), threadIdx:(1,10)
c:2, r:10,blockDim:(16,16), blockIdx:(0,0), threadIdx:(2,10)
.........
c:15, r:10,blockDim:(16,16), blockIdx:(0,0), threadIdx:(15,10)
c:0, r:11,blockDim:(16,16), blockIdx:(0,0), threadIdx:(0,11)
c:1, r:11,blockDim:(16,16), blockIdx:(0,0), threadIdx:(1,11)
c:2, r:11,blockDim:(16,16), blockIdx:(0,0), threadIdx:(2,11)
c:3, r:11,blockDim:(16,16), blockIdx:(0,0), threadIdx:(3,11)
c:4, r:11,blockDim:(16,16), blockIdx:(0,0), threadIdx:(4,11)
c:5, r:11,blockDim:(16,16), blockIdx:(0,0), threadIdx:(5,11)
c:6, r:11,blockDim:(16,16), blockIdx:(0,0), threadIdx:(6,11)
c:7, r:11,blockDim:(16,16), blockIdx:(0,0), threadIdx:(7,11)
.........
c:7, r:1,blockDim:(16,16), blockIdx:(0,0), threadIdx:(7,1)
c:8, r:1,blockDim:(16,16), blockIdx:(0,0), threadIdx:(8,1)
c:9, r:1,blockDim:(16,16), blockIdx:(0,0), threadIdx:(9,1)
c:10, r:1,blockDim:(16,16), blockIdx:(0,0), threadIdx:(10,1)
c:11, r:1,blockDim:(16,16), blockIdx:(0,0), threadIdx:(11,1)
c:12, r:1,blockDim:(16,16), blockIdx:(0,0), threadIdx:(12,1)

n:32, dimGrid(2,2,1), dimBlock(16,16,1), total_threads:1024

矩阵32x32, 1个grid有4个block（二维2x2），每个block的线程为(16x16)， 比如threadIdx:(0,8)，意思是该线程在当前block的位置为x=1, y=8。

n:32, dimGrid(2,2,1), dimBlock(16,16,1), total_threads:1024
​
c:16, r:24,blockDim:(16,16), blockIdx:(1,1), threadIdx:(0,8)
​
c:17, r:24,blockDim:(16,16), blockIdx:(1,1), threadIdx:(1,8)
​
c:18, r:24,blockDim:(16,16), blockIdx:(1,1), threadIdx:(2,8)
​
c:19, r:24,blockDim:(16,16), blockIdx:(1,1), threadIdx:(3,8)
​
c:20, r:24,blockDim:(16,16), blockIdx:(1,1), threadIdx:(4,8)
​
c:21, r:240,blockDim:(16,16), blockIdx:(1,1), threadIdx:(5,8)
​
c:22, r:24,blockDim:(16,16), blockIdx:(1,1), threadIdx:(6,8)
​
c:23, r:24,blockDim:(16,16), blockIdx:(1,1), threadIdx:(7,8)
​
###### .....
​
c:30, r:8,blockDim:(16,16), blockIdx:(1,0), threadIdx:(14,8)
​
c:31, r:8,blockDim:(16,16), blockIdx:(1,0), threadIdx:(15,8)
​
c:16, r:9,blockDim:(16,16), blockIdx:(1,0), threadIdx:(0,9)
​
c:17, r:9,blockDim:(16,16), blockIdx:(1,0), threadIdx:(1,9)
​
c:18, r:9,blockDim:(16,16), blockIdx:(1,0), threadIdx:(2,9)
​
## ......
​
c:10, r:23,blockDim:(16,16), blockIdx:(0,1), threadIdx:(10,7)
​
c:11, r:23,blockDim:(16,16), blockIdx:(0,1), threadIdx:(11,7)
​
c:12, r:23,blockDim:(16,16), blockIdx:(0,1), threadIdx:(12,7)
​
c:13, r:23,blockDim:(16,16), blockIdx:(0,1), threadIdx:(13,7)
​
c:14, r:23,blockDim:(16,16), blockIdx:(0,1), threadIdx:(14,7)
​
c:15, r:23,blockDim:(16,16), blockIdx:(0,1), threadIdx:(15,7)

8 MPI与CUDA的区别

全称Massage Passing Interface

是支持c、c++等语言的并行编程的拓展库，主要是负责多进程之间的通信。用于编写并行计算程序。我们通过MPI并行库来编写并行化的程序。

由于“天河二号”等高性能计算机在运行的时候是同一个程序会运行在很多节点上，每个节点上都是一个进程。这些进程也就是这些节点之间需要相互通信来达到程序的并行。因此想要利用“天河二号”的计算能力来帮助自己运行程序，就需要将自己的程序改为MPI的并行程序，至于超级计算机分配哪些任务给哪些节点是我们不需要知道的，以及节点之间如何通信，利用中心架构进行通信还是非中心架构进行通信也是我们不需要知道的，我们要了解的就是如何将自己在个人计算机上运行的普通程序改成可以在超级计算机上运行的MPI程序即可。

MPI框架下，同一个程序在多个节点中以进程形式存在，这些进程组成一个group，每个进程都有唯一的进程号，MPI的点对点通信有两种，一种是消息发送，一种是消息的接收，最简单的为MPI_Send()和MPI_Recv()。相对应的，还有另外三种通信方式缓存通信、同步通信、就绪通信。

（1）缓存通信：用户提供通信缓冲区，避免了系统内存拷贝，提高了通信效率，但是缓冲区需用

户自己管理。

（2）同步通信：发送进程只有当接受进程开始接收（不需要全部接收）的时候才返回。

（3）就绪通信：发送进程的发送操作只有当接受进程已经开启了接收操作的时候才能够成功调

用，否则发送操作将会出错。

参考

https://www.zhihu.com/question/35063258/answer/108012477

https://www.zhihu.com/question/35063258

https://www.zhihu.com/question/21231074/answer/20701124

https://www.zhihu.com/question/303525717/answer/1029812564

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