principles from introductory CUDA examples to an existing project that has a meaningful amount of non-trivial code. • Provide some guidance to people about to embark on using CUDA to speed up existing software float* y) { for (int i = 0; i < n; i++) y[i] = x[i] + y[i]; } TEST_CASE("cppcon-0", "[CUDA]") { int N = 1 << 20; float* x = new float[N]; float* y = new float[N]; for (int add_cpu(N, x, y); delete[] x; delete[] y; } An Even Easier Introduction to CUDA 4 |TEST_CASE("cppcon-1", "[CUDA]") { int N = 1 << 20; float* x; float* y; cudaMallocManaged(&x, N*sizeof(float));

CUDA 开启的 GPU 编程 by 彭于斌（ @archibate ） 往期录播： https://www.bilibili.com/video/BV1fa411r7zp 课程 PPT 和代码： https://github.com/parallel101/course 前置条件 • 学过 C/C++ 语言编程。 • 理解 malloc/free 之类的概念。 • 熟悉 STL 中的容器、函数模板等。 中的容器、函数模板等。 • 英伟达 GTX900 及以上显卡。 • CUDA 11 及以上。 • CMake 3.18 及以上。 我负责监督你学习 第 0 章： Hello, world! CMake 中启用 CUDA 支持 • 最新版的 CMake （ 3.18 以上），只需在 LANGUAGES 后面加上 CUDA 即可启用 。 • 然后在 add_executable 里直接加你 cn/docs/IO/51635/NVIDIA_CUDA_Programming_Guide_1.1_chs.pdf CUDA 编译器兼容 C++17 • CUDA 的语法，基本完全兼容 C++ 。包括 C+ +17 新特性，都可以用。甚至可以把任何一个 C++ 项目的文件后缀名全部改成 .cu ，都能编 译出来。 • 这是 CUDA 的一大好处， CUDA 和 C++ 的关 系就像 C++ 和

CUDA C++ 流体仿真实 战 by 彭于斌（ @archibate ） 往期录播： https://www.bilibili.com/video/BV16b4y1E74f 课程 PPT 和代码： https://github.com/parallel101/course CUDA 纹理对象 https://docs.nvidia.com/cuda/cuda-c-programming-guide/index g-guide/index.html#texture-and-surface-memory CUDA 多维数组：封装 • cudaMalloc3DArray 用于分配一个三维数组。 各维度上的大小通过 cudaExtent 指定，方 便起见我们的 C++ 封装类用了 uint3 表示 大小。 • GPU 的多维数组有特殊的数据排布来保障 访存的高效，和我们 CPU 那样简单地行主 序或列主序（如 序或列主序（如 a[x + nx * y] ）的多维数组 不一样。 • 随后可用 cudaMemcpy3D 在 GPU 的三 维数组和 CPU 的三维数组之间拷贝数据。 CUDA 表面对象：封装 • 要访问一个多维数组，必须先创建一个表面对象 （ cudaSurfaceObject_t ）。 • 考虑到多维数组始终是需要通过表面对象来访问的，这 里我们让表面对象继承自多维数组。 •

Makefile 启动时会把每个文件都检测一遍， 浪费很多时间。特别是有很多文件，但是实 际需要构建的只有一小部分，从而是 I/O Bound 的时候， Ninja 的速度提升就很明 显。 然而某些专利公司的 CUDA toolkit 在 Windows 上只允许用 MSBuild 构建，不能 用 Ninja （怕不是和 Bill Gates 有什么交 易） 第 1 章：添加源文件 一个 .cpp 源文件用于测试 指定了该项目使用了哪些编程语言。 • 目前支持的语言包括： • C ： C 语言 • CXX ： C++ 语言 • ASM ：汇编语言 • Fortran ：老年人的编程语言 • CUDA ：英伟达的 CUDA （ 3.8 版本新增） • OBJC ：苹果的 Objective-C （ 3.16 版本新增） • OBJCXX ：苹果的 Objective-C++ （ 3.16 版本新增） CXX_STANDARD 或是全局变量 CMAKE_CXX_STANDARD 来 设置 -std=c++17 这个 flag ， CMake 会在配置阶段检测编译器是否支持 C++17 。 CUDA 的 -arch=sm_75 也是同理，请使用 CUDA_ARCHITECTURES 属 性。 再说了 -std=c++17 只是 GCC 编译器的选项，无法跨平台用于 MSVC 编 译器。 假如你一定要用动态链接库（ Windows

Visual Studio 2019 （ Windows 用户） GCC 9 及以上（ Linux 用户） CMake 3.12 及以上（跨平台作业） Git 2.x （作业上传到 GitHub ） CUDA Toolkit 10.0 以上（ GPU 专题） 关于作者 • 我是 Taichi 编译器的贡献者之一（ https://github.com/taichi-dev/taichi ） 关于作者（续） 添加要编译的源文件 • 以及可以通过下列指令（不推荐使用），把选项加到所有接下来的目标去： • include_directories(/opt/cuda/include) # 添加头文件搜索目录 • link_directories(/opt/cuda) # 添加库文件的搜索路径 • add_definitions(MY_MACRO=1)

Visual Studio 2019 （ Windows 用户） GCC 9 及以上（ Linux 用户） CMake 3.12 及以上（跨平台作业） Git 2.x （作业上传到 GitHub ） CUDA Toolkit 10.0 以上（ GPU 专题） 第 0 章：从并发到并行 摩尔定律：停止增长了吗？ • 晶体管的密度的确仍在指数增长，但处理器主 频却开始停止增长了，甚至有所下降。 • 结论：改进后的并行扫描的时间复杂度为 O(logn) ，工作复杂度为 O(nlogn) 。 可见，并行后虽然降低了时间复杂度，但是以提升工作复杂度为代价！ 更多细节，敬请期待 GPU 专题，我们会以 CUDA 为例详细探讨两全方案。 封装好了： parallel_scan 第 3 章：性能测试 案例： map 与 reduce 的组合 测试所花费时间： tbb::tick_count::now()

• Currently up to 12 VMs, each with 16 cores • VMs install VS (with Clang, CMake, Ninja), Python, CUDA • We enforce clang-format for product/test code • Saves an incredible amount of time • Our parallelize

（大底数版） • 这就是 half 类型，他只有 5 位指数， 10 位底 数，总共占据 16 位，所以又称 float16 。精度很 低，但是节省内存空间！ • 然而只有 GPU （比如 CUDA ）支持 half 类型 ， CPU 需要支持 AVX512fp16 这个扩展才能用 。 • 据说深度学习（很多都是 membound ）很喜欢 用 half ，因为可以省一半内存，从而加快一倍。

Visual Studio 2019 （ Windows 用户） GCC 9 及以上（ Linux 用户） CMake 3.12 及以上（跨平台作业） Git 2.x （作业上传到 GitHub ） CUDA Toolkit 10.0 以上（ GPU 专题） 第 0 章：汇编语言 x64 架构下的寄存器模型 通用寄存器： 32 位时代 • 32 位 x86 架构中的通用寄存器有： • eax

Visual Studio 2019 （ Windows 用户） GCC 9 及以上（ Linux 用户） CMake 3.12 及以上（跨平台作业） Git 2.x （作业上传到 GitHub ） CUDA Toolkit 10.0 以上（ GPU 专题） 为什么需要模板函数（ template ） • 避免重复写代码。 • 比如，利用重载实现“将一个数乘以 2” 这个 功能，需要： 为什么面向对象在