初识cuda一文通
cuda学习博客
本文为本人cuda学习过程中的记录和理解，多参考@谭升等大佬前辈的博客，以及NVIDIA官方文档。如有错误烦请指正，如有侵权请联系删除。
0. 并行计算与计算机架构
计算机架构是并行计算或者说HPC领域中十分重要的内容。
。。看书后期记得补充。。
无论是OpenMP还是CUDA，写并行程序主要是分解任务，在软件层面上大体可以分为‘指令并行’和‘数据并行’。
1. 异构计算与CUDA
cuda基础
进入正题，我们从hello world开始学习cuda
#include<stdio.h>
__global__ void hello_world(void)
{
printf("hello world!);
}
int main(int argc, char **argv)
{
hello_world<<<1, 10>>>();
cudaDeviceReset();
return 0;
}
下面我们具体审视一下这个简单的cuda程序。
（1）__global__前缀，cuda的精髓，即核函数。核函数的代码将在成千上万个线程上执行。
（2）cudaDeviceReset(),实现CPU和Device（即GPU）的同步，保证cpu和gpu一起退出程序。
大家这里先不用着急，后续我们还会具体谈cuda中执行单元的划分，并NVIDIA为我们提供的若干工具来分析编写并行程序。
在异构环境中，主机host和设备device通过pcie总线通信。通信内容就包括了最重要的内存信息传递。CUDA上的内存管理API大体和cpu中相同，例如：
cudaMalloc()诸如此类
跑在总线上的一个重要API：cudaMemcpy()内存拷贝函数，可以实现DeviceToHost/HostTODevice等等过程。
GPU架构
线程管理
下面我们进入到线程管理。首先，一个核函数只能有一个grid，每个grid可以分为多个block，每个block又可以分为多个thread。CUDA还为同一个block中的线程提供了共享内存机制，在每个block中单独划分出一块内存空间用于共享，通过关键词__shared__访问，拥有比全局内存更好的性能表现，这点后续我们也会详谈。
核函数
kernel_name<<<grid,block>>>
注意host端核函数执行是异步的，并且当主机发出执行核函数的指令后，会立即收回控制权。此刻就需要方法来使host和device同步，最常见的是使用cudaMemcpy来使host等待device数据，从而实现同步。
在CUDA程序编写时，最常见的是将串行程序中的for循环并行化（这点在OpenMP和MPI中已经很熟悉了）
细谈线程束Warp
CUDA中的同步
无论是pthread中的🔓还是openmp，都有相应的同步机制，CUDA为避免内存竞争，也有同步方法。__syncthread()可以实现同block内的线程同步，想要同步不同block的线程，方法是借助核函数执行。
CUDA并行性能分析
这里我们使用简单的矩阵加法为例来进行实验分析，硬件平台为RTX3060
分析工具介绍
- nvprof：
动态并行
动态并行或者叫他嵌套并行，类似于父进程/子进程的概念。
2.CUDA内存
- 寄存器
- 共享内存
- 常量内存
- 全局内存
- 纹理内存
- 本地内存
内存管理
cudaMelloc() //分配内存
cudaMemset() //初始化内存
cudaFree() //释放内存
cudaemcpy() //传输内存
特殊内存
固定内存：固定内存传输速更快，但是分配和释放成本更高
零拷贝内存：
统一内存寻址：cudaMallocManaged()
内存访问模式
这部分内容至关重要，对于程序调优来说，记得代码验证
1.对齐合并内存访问，减少带宽浪费
2.足够的并发内存操作，隐藏内存延迟
