commit d06791d2ce3a64bde8e676ac59dd2acfe0af0c67 Author: unknown Date: Sun Nov 14 21:03:09 2021 +0800 add capter1-2 diff --git a/.gitignore b/.gitignore new file mode 100644 index 0000000..5998c9a --- /dev/null +++ b/.gitignore @@ -0,0 +1,23 @@ +# 忽略.class的所有文件 +# *.class + +# 忽略名称中末尾为ignore的文件夹 +*ignore/ + +# 忽略名称中间包含ignore的文件夹 +*ignore*/ + +# 忽略指定文件 +# HelloWrold.class +AStyle.exe +*.orig + +# 忽略指定文件夹 +build/ +*/bin/ +*/result.txt +*/__pycache__/ +.vs/ +.vscode/ +*/x64/ +*/Debug/ \ No newline at end of file diff --git a/ReadMe.md b/ReadMe.md new file mode 100644 index 0000000..c65d1f6 --- /dev/null +++ b/ReadMe.md @@ -0,0 +1,19 @@ +# CUDA Study Steps + +CUDA gpu 编程学习，基于 《CUDA 编程——基础与实践》（樊哲勇）。 + +包含章节： + +1. [GPU 硬件与 CUDA 程序开发工具](./capter1/ReadMe.md) +2. [CUDA 中的线程组织](./capter2/ReadMe.md) +3. [简单 CUDA 程序的基本框架](./capter3/ReadMe.md) +4. [CUDA 程序的错误检测](./capter4/ReadMe.md) +5. [GPU 加速的关键](./capter5/ReadMe.md) + +CUDA 官方文档： +[CUDA c++编程指南](https://docs.nvidia.com/cuda/cuda-c-programming-guide/index.html) +[CUDA c++最佳实践指南](https://docs.nvidia.com/cuda/cuda-c-best-practices-guide/index.html) +[CUDA 运行时API手册](https://docs.nvidia.com/cuda/cuda-runtime-api/index.html) +[CUDA 数学函数库API手册](https://docs.nvidia.com/cuda/cuda-math-api/index.html) + + diff --git a/capter1/ReadMe.md b/capter1/ReadMe.md new file mode 100644 index 0000000..7ad321a --- /dev/null +++ b/capter1/ReadMe.md @@ -0,0 +1,73 @@ +# GPU 硬件与 CUDA 程序开发工具 + +------ + +## GPU 硬件 + +在由 CPU 和 GPU 构成的异构计算平台中，通常将起控制作用的 CPU 称为 **主机（host）**， +将起加速作用的 GPU 称为 **设备（device）**。 + +主机和设备都有自己的 DRAM，之间一般由 PCIe 总线连接。 + +GPU 计算能力不等价于计算性能；表征计算性能的一个重要参数是 **浮点数运算峰值（FLOPS）**。 +浮点数运算峰值有单精度和双精度之分。对于 Tesla 系列的 GPU，双精度下 FLOPS 一般是单精度下的 1/2;对于 GeForce 系列的 GPU，双精度下 FLOPS 一般是单精度下的 1/32。 + +影响计算性能的另一个参数是 **GPU 内存带宽（显存）**。 + +------ + +## CUDA 程序开发工具 + +1. CUDA； +2. OpenCL，更为通用的各种异构平台编写并行程序的框架，AMD 的 GPU 程序开发工具； +3. OpenACC，由多公司共同开发的异构并行编程标准。 + +CUDA 提供两层 API，即 CUDA 驱动API 和 CUDA 运行时API。 +CUDA 开发环境中，程序应用程序是以主机（CPU）为出发点的；应用程序可以调用 CUDA 运行时 API、CUDA 驱动 API 和一些已有的 CUDA 库。 + +------ + +## CUDA 开发环境搭建 + +linux 操作系统：[linux下cuda环境搭建](https://docs.nvidia.com/cuda/cuda-installation-guide-linux/index.html) + +windows10 操作系统：[windows10下cuda环境搭建](https://docs.nvidia.com/cuda/cuda-installation-guide-microsoft-windows/index.html) + +------ + +## nvidia-smi 检查与设置设备 + + >> nvidia-smi + +-----------------------------------------------------------------------------+ + | NVIDIA-SMI 462.30 Driver Version: 462.30 CUDA Version: 11.2 | + |-------------------------------+----------------------+----------------------+ + | GPU Name TCC/WDDM | Bus-Id Disp.A | Volatile Uncorr. ECC | + | Fan Temp Perf Pwr:Usage/Cap| Memory-Usage | GPU-Util Compute M. | + | | | MIG M. | + |===============================+======================+======================| + | 0 GeForce MX450 WDDM | 00000000:2B:00.0 Off | N/A | + | N/A 39C P8 N/A / N/A | 119MiB / 2048MiB | 0% Default | + | | | N/A | + +-------------------------------+----------------------+----------------------+ + + +-----------------------------------------------------------------------------+ + | Processes: | + | GPU GI CI PID Type Process name GPU Memory | + | ID ID Usage | + |=============================================================================| + | No running processes found | + +-----------------------------------------------------------------------------+ + +1. **CUDA Version**， 11.2； +2. **GPU Name**，GeForce MX450，设备号为 0；如果系统中有多个 GPU 且只要使用其中某个特定的 GPU， +可以通过设置环境变量 **CUDA_VISIBLE_DEVICES** 的值，从而可以在运行 CUDA 程序前选定 GPU; +3. **TCC/WDDM**，WDDM（windows display driver model），其它包括 TCC（Tesla compute cluster）； +可以通过命令行 `nvidia-smi -g GPU_ID -dm 0`，设置为 WDDM 模式（1 为 TCC 模式）； +4. **Compute mode**, Default，此时同一个 GPU 中允许存在多个进程；其他模式包括 E.Process， +指的是独占进程模式，但不适用 WDDM 模式下的 GPU； +可以通过命令行 `nvidia-smi -i GPU_ID -c 0`，设置为 Default 模式（1 为 E.Process 模式）; +5. **Perf**，p8（GPU 性能状态，最大p0~最小p12）； + +更多关于 nvidia-smi 的资料：[nvidia-smi](https://developer.nvidia.com/nvidia-system-management-interface) + +------ \ No newline at end of file diff --git a/capter2/ReadMe.md b/capter2/ReadMe.md new file mode 100644 index 0000000..8d59d29 --- /dev/null +++ b/capter2/ReadMe.md @@ -0,0 +1,169 @@ +# CUDA 中的线程组织 + +CUDA 虽然支持 C++ 但支持得并不充分，导致 C++ 代码中有很多 C 代码的风格。 + +CUDA 采用 nvcc 作为编译器，支持 C++ 代码；nvcc 在编译 CUDA 程序时， +会将纯粹的 c++ 代码交给 c++ 编译器，自己负责编译剩下的 cu 代码。 + +------ + +## C++ 的 Hello World 程序 + + >> g++ hello.cpp -o ./bin/hello.exe + >> ./bin/hello + msvc: hello world! + +------ + +## CUDA 的 Hello World 程序 + +### 使用 nvcc 编译纯粹 c++ 代码 + + >> nvcc -o ./bin/hello_cu.exe hello.cu + >> ./bin/hello_cu.exe + nvcc: hello world! + +在该程序中其实并未使用 GPU。 + +### 使用 核函数 的 CUDA 程序 + +一个利用了 GPU 的 CUDA 程序既有主机代码，又有设备代码（在设备中执行的代码）。 +主机对设备的调用是通过 **核函数（kernel function）** 实现的。 + + int main() + { + 主机代码 + 核函数的调用 + 主机代码 + + return 0； + } + +核函数与 c++ 函数的区别： +1. 必须加 `__global__` 限定； +2. 返回类型必须是空类型 `void`。 + + __global__ void hell_from__gpu() + { + // 核函数不支持 c++ 的 iostream。 + printf("gpu: hello world!\n"); + } + +调用核函数的方式： + + hello_from_gpu<<<1, 1>>> + +主机在调用一个核函数时，必须指明在设备中指派多少线程。核函数中的线程常组织为若干线程块： +1. 三括号中第一个数字是线程块的个数（number of thread block）； +2. 三括号中第二个数字是每个线程块中的线程数（number of thread in per block）。 + +一个核函数的全部线程块构成一个网格（grid），线程块的个数称为网格大小（grid size）。 +每个线程块中含有相同数目的线程，该数目称为线程块大小（block size）。 + +所以，核函数的总的线程数即 网格大小*线程块大小: + + hello_from_gpu<<>> + +调用核函数后，调用 CUDA 运行时 API 函数，同步主机和设备： + + cudaDeviceSynchronize(); + +核函数中调用输出函数，输出流是先存放在缓冲区的，而缓冲区不会自动刷新。 + +------ + +## CUDA 的线程组织 + +核函数的总线程数必须至少等于计算核心数时才有可能充分利用 GPU 的全部计算资源。 + + hello_from_gpu<<<2, 4>>> + +网格大小是2，线程块大小是4，总线程数即8。核函数中代码的执行方式是 “单指令-多线程”， +即每个线程执行同一串代码。 + +从开普勒架构开始，最大允许的线程块大小是 2^10 (1024)，最大允许的网格大小是 2^31 - 1（一维网格）。 + +线程总数可以由两个参数确定： +1. gridDim.x, 即网格大小； +2. blockDim.x, 即线程块大小； + +每个线程的身份可以由两个参数确定： +1. blockIdx.x, 即一个线程在一个网格中的线程块索引，[0, gridDm.x); +2. threadIdx.x, 即一个线程在一个线程块中的线程索引，[0, blockDim.x); + +网格和线程块都可以拓展为三维结构（各轴默认为 1）： + +1. 三维网格 grid_size(gridDim.x, gridDim.y, gridDim.z); +2. 三维线程块 block_size(blockDim.x, blockDim.y, blockDim.z); + +相应的，每个线程的身份参数： + +1. 线程块ID (blockIdx.x, blockIdx.y, blockIdx.z); +2. 线程ID (threadIdx.x, threadIdx.y, threadIdx.z); + +多维网格线程在线程块上的 ID； + + tid = threadIdx.z * (blockDim.x * blockDim.y) // 当前线程块上前面的所有线程数 + + threadIdx.y * (blockDim.x) // 当前线程块上当前面上前面行的所有线程数 + + threadIdx.x // 当前线程块上当前面上当前行的线程数 + +多维网格线程块在网格上的 ID: + + bid = blockIdx.z * (gridDim.x * gridDim.y) + + blockIdx.y * (gridDim.x) + + blockIdx.x + +一个线程块中的线程还可以细分为不同的 **线程束（thread warp）**，即同一个线程块中 +相邻的 warp_size 个线程（一般为 32）。 + +对于从开普勒架构到图灵架构的 GPU，网格大小在 x, y, z 方向的最大允许值为 （2^31 - 1, 2^16 - 1, 2^16 -1）； +线程块大小在 x, y, z 方向的最大允许值为 （1024， 1024， 64），同时要求一个线程块最多有 1024 个线程。 + +------ + +## CUDA 的头文件 + +CUDA 头文件的后缀依然是 “.h”；同时，采用 nvcc 编译器会自动包含必要的 cuda 头文件， +如 , ，同时前者也包含了c++头文件 。 + +------ + +## 使用 nvcc 编译 CUDA 程序 + +nvcc 会先将全部源代码分离为 主机代码 和 设备代码；主机代码完整的支持 c++ 语法，而设备代码只部分支持。 + +nvcc 会先将设备代码编译为 PTX（parrallel thread execution）伪汇编代码，再将其编译为二进制 cubin目标代码。 +在编译为 PTX 代码时，需要选项 `-arch=compute_XY` 指定一个虚拟架构的计算能力；在编译为 cubin 代码时， +需要选项 `-code=sm_ZW` 指定一个真实架构的计算能力，以确定可执行文件能够使用的 GPU。 + +真实架构的计算能力必须大于等于虚拟架构的计算能力，例如： + + -arch=compute_35 -code=sm_60 (right) + -arch=compute_60 -code=sm_35 (wrong) + +如果希望编译出来的文件能在更多的GPU上运行，则可以同时指定多组计算能力，例如： + + -gencode arch=compute_35, code=sm_35 + -gencode arch=compute_50, code=sm_50 + -gencode arch=compute_60, code=sm_60 + +此时，编译出来的可执行文件将包含3个二进制版本，称为 **胖二进制文件（fatbinary）**。 + +同时，nvcc 有一种称为 **实时编译（just-in-time compilation）**机制，可以在运行可执行文件时从其中保留的PTX +代码中临时编译出一个 cubin 目标代码。因此， 需要通过选项 `-gencode arch=compute_XY, code=compute_XY`， +指定所保留 PTX 代码的虚拟架构， 例如： + + -gencode arch=compute_35, code=sm_35 + -gencode arch=compute_50, code=sm_50 + -gencode arch=compute_60, code=sm_60 + -gencode arch=compute_70, code=compute_70 + +于此同时，nvcc 编译有一个简化的编译选项 `-arch=sim_XY`，其等价于： + + -gencode arch=compute_XY, code=sm_XY + -gencode arch=compute_XY, code=compute_XY + +关于 nvcc 编译器的更多资料： [nvcc](https://docs.nvidia.com/cuda/cuda-compiler-driver-nvcc/index.html)。 + +------ + diff --git a/capter2/bin/hello.exe b/capter2/bin/hello.exe new file mode 100644 index 0000000..7b4c4d3 Binary files /dev/null and b/capter2/bin/hello.exe differ diff --git a/capter2/bin/hello_cu.exe b/capter2/bin/hello_cu.exe new file mode 100644 index 0000000..a2c88f3 Binary files /dev/null and b/capter2/bin/hello_cu.exe differ diff --git a/capter2/bin/hello_cu.exp b/capter2/bin/hello_cu.exp new file mode 100644 index 0000000..0178fce Binary files /dev/null and b/capter2/bin/hello_cu.exp differ diff --git a/capter2/bin/hello_cu.lib b/capter2/bin/hello_cu.lib new file mode 100644 index 0000000..70d393e Binary files /dev/null and b/capter2/bin/hello_cu.lib differ diff --git a/capter2/hello.cpp b/capter2/hello.cpp new file mode 100644 index 0000000..1c4cdc7 --- /dev/null +++ b/capter2/hello.cpp @@ -0,0 +1,9 @@ +#include +using namespace std; + +int main() +{ + cout << "msvc: hello world!" << endl; + + return 0; +} \ No newline at end of file diff --git a/capter2/hello.cu b/capter2/hello.cu new file mode 100644 index 0000000..56340fb --- /dev/null +++ b/capter2/hello.cu @@ -0,0 +1,33 @@ +#include +using namespace std; + + +__global__ void hell_from__gpu() +{ + // 核函数不支持 c++ 的 iostream。 + + // 输出流的缓存顺序。 + // printf("gpu: hello world! "); + + const int bx = blockIdx.x; + const int by = blockIdx.y; + const int bz = blockIdx.z; + + const int tx = threadIdx.x; + const int ty = threadIdx.y; + const int tz = threadIdx.z; + + printf("gpu: hello world! block(%d, %d, %d) -- thread(%d, %d, %d)\n", bx, by, bz, tx, ty, tz); +} + + +int main() +{ + printf("nvcc: hello world!\n"); + + const dim3 block_size(2, 4); + hell_from__gpu<<<1, block_size>>>(); + cudaDeviceSynchronize(); // 同步主机和设备，否则无法输出字符串。 + + return 0; +} \ No newline at end of file