float* x, 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 i = 0; i < N; i++) { x[i] = 1.0f; y[i] = 2.0f; } add_cpu(N, x, y); delete[] x; delete[] y; } An Even Easier Introduction to CUDA 4 |TEST_CASE("cppcon-1" 1 << 20; float* x; float* y; cudaMallocManaged(&x, N*sizeof(float)); cudaMallocManaged(&y, N*sizeof(float)); // … cudaFree(x); cudaFree(y); } An Even Easier Introduction to

float ，大小仅为 12 字节。 • 所以当修改 MyClass 的 x 属性时，必须读取整个 MyClass ，仅仅 修改其中的 x 属性，然后写回。这就浪费了： y 的写、 z 的写、 z 的读。实际用上的只有： x 的读， y 的读， x 的写。浪费了 50% 带宽。 • 而 SOA 把三个属性分开存，每个属性作为独立的数组，稠密存储。 这样当用不到 z 的时候， z 数组就完全不会被读取，不会占用内 一次要读哪里，等发现跳跃时已经来不及了，从而 计算的延迟无法隐藏。 如果每个属性都要访问到，那还是 AOS 比较好（ AOSOA 也不赖哦） 这是因为使用 SOA 会让 CPU 不得不同时维护很多条预取赛道（ mc_x, mc_y, mc_z ），当赛 道多了以后每一条赛道的长度就变短了，从而能够周转的余地时间比较少，不利于延迟隐藏。 而如果把这三条赛道合并成一条（ mc ），这样同样的经费（缓存容量）能铺出的赛道（预 取）就更长，从而 float a[nz][ny][nx]; • a[z][y][x]; • 等价于： • float a[nz * ny * nx]; • a[(z * ny + y) * nx + x]; • //////////////////////////////////// • float a[nw][nz][ny][nx]; • a[w][z][y][x]; • 等价于： • float a[nw

“x” points to “v”. Disjointness Given “(x ↦ v) ∗ (y ↦ w)”, we know that “x ≠ y”, i.e., they do not alias. P ∗ Q Separation Logic (Iris) P Q v x y P −∗ Q Magic wand Consuming resource “P”, we = 1) ∗ (y = 2) ⟺ (x = 1) ∧ (y = 2) Resources, unique Resources, not duplicable Non-resources Logics {True} let x = Box::new(1) {x ↦ 1} {(x ↦ 1) ∗ (y ↦ 1)} *x += *y {(x ↦ 2) ∗ (y ↦ 1)} data Rc::::clone Examples {⟦rc τ ⟧.own([x])} let y = x.clone() {⟦rc τ ⟧.own([x]) ∗ ⟦rc τ ⟧.own([y])} x.cnt += 1; ? x.cnt += 1; let y = ptr::read(x) Goal Proof sketch {∃ℓ.?=ℓ∗∃?∈ℤ+¿,?∈ℚ∩¿

Requires a size() method returning a size_t template concept C = requires (T& x, T& y) { { x.swap(y) } noexcept; { x.size() } -> std::convertible_to; ... };23 Concepts  Combining Operator <=>  Common case:  auto X::operator<=>(const Y&) const = default;  Compiler generates all 6 comparison operators to compare X with Y (memberwise)  Advanced:  Non-defaulted (then you also Spaceship Operator <=> class Point { int x; int y; public: friend bool operator==(const Point& a, const Point& b){ return a.x==b.x && a.y==b.y; } friend bool operator< (const Point& a, const Point&

month_day  year_month  year_month_day  …30 Dates  Define an std::year: year y1 { 2021 }; auto y2 { 2021y };  Define an std::month: month m1 { 10 }; auto m2 { October };  Define an std::day: a date 2021-10-27: year_month_day fulldate1 { 2021y, October, 27d }; auto fulldate2 { 2021y / October / 27d }; auto fulldate3 { 27d / October / 2021y };  Create a date for the 4th Wednesday of October auto oct27 { October / 27d };  Create a year_month_day for October 27, 2021: auto oct27_2021 { 2021y / oct27 };32 Dates  Create a month_day_last for the last day of an October of an unspecified year:

只要在三重尖括号内指定的参数改成 dim3 类型即可。 dim3 的构造函数就是接受三 个无符号整数（ unsigned int ）非常简单 。 • dim3(x, y, z) • 这样在核函数里就可以通过 threadIdx.y 获取 y 方向的线程编号，以此类推。 那二维呢？ • 需要二维的话，只需要把 dim3 最后一位 （ z 方向）的值设为 1 即可。这样就只有 xy 方向有大小，就相当于二维了，不会有 • 类似的这样的低精度內建函数还有 __expf 、 __logf 、 __cosf 、 __powf 等。 • 还有 __fdividef(x, y) 提供更快的浮点除法 ，和一般的除法有相同的精确度，但是在 2^216 < y < 2^218 时会得到错误的结果。 编译器选项： --use_fast_math • 如果开启了 --use_fast_math 选项，那么所有对 a * b + c 优化成乘加 (FMA) 指令。 • 开启 --use_fast_math 后会自动开启上述所有。 SAXPY （ Scalar A times X Plus Y ） • 即标量 A 乘 X 加 Y 。 • 这是很多 CUDA 教科书上的 Hello, world 。 • 却是小彭老师课第 5 章的结尾。 第 6 章： thrust 库 替换成 CUDA 官方提供的

#include #include using namespace std::chrono; int main() { year y{2021}; std::cout << y << "\n"; month m{October}; auto result = m + months{3}; std::cout << result year_month_day ymd1{October/28d/2021y}; year_month_day ymd2{2021y/October/28d}; std::cout << ymd1 << "\t" << ymd2 << "\n"; year_month_day_last ymdlast{2021y/(October + months{1})/last}; year_month_day example using operator/: year_month_day ymd1{October/28d/ 2021y}; year_month_day ymd2{2021y/October/ 28d}; Without the day and year types, these would look like year_month_day

没关系，可以用结构化绑定的语法： • auto [x, y, ...] = tup; • 利用一个方括号，里面是变量名列表，即 可解包一个 tuple 。里面的数据会按顺序 赋值给每个变量，非常方便。 tuple ：结构化绑定为引用 • 结构化绑定也支持绑定为引用： • auto &[x, y, ...] = tup; • 这样相当于解包出来的 x, y, ... 都是 auto & 推 导出来的引用类型。对引用的修改可以影响到原 & 绑定为常引用： • auto const &[x, y, ...] = tup; • 常引用虽然不能修改，但是可以避免一次不必要拷贝。 tuple ：结构化绑定为万能推导 • 不过要注意一下万能推导的 decltype(auto) ， 由于历史原因，他对应的结构化绑定是 auto && ： • auto &&[x, y, ...] = tup; // 正确！ 正确！ • decltype(auto) [x, y, ...] = tup; // 错误！ • 对的，是两个与号 && 。 结构化绑定：还可以是任意自定义类！ • 其实，结构化绑定不仅可以解包 std::tuple ， 还可以解包任意用户自定义类： • 配合打包的 {} 初始化表达式，真是太便利了！ • 惊不惊喜？意不意外？ • 可惜 std::get 并不支持自定义类。

X Yint a[10]; int f(int x, int y) { if (x > y) return f(y, x); if (x < 0) return -1; if (y + x > 9) return 0; return a[y - x]; } Queries Given: is true

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 便起见我们的 C++ 封装类用了 uint3 表示 大小。 • GPU 的多维数组有特殊的数据排布来保障 访存的高效，和我们 CPU 那样简单地行主 序或列主序（如 a[x + nx * y] ）的多维数组 不一样。 • 随后可用 cudaMemcpy3D 在 GPU 的三 维数组和 CPU 的三维数组之间拷贝数据。 CUDA 表面对象：封装 • 要访问一个多维数组，必须先创建一个表面对象 • 考虑到多维数组始终是需要通过表面对象来访问的，这 里我们让表面对象继承自多维数组。 • 在核函数中可以用 surf3Dread 和 surf3Dwrite 来读写 表面对象中的元素， x,y,z 参数指定要访问元素的坐标 ，要注意 x 必须乘以 sizeof( 元素类型 ) ，否则出错。 • 这里用了访问者模式（ Accessor ， GPU 编程常用）。 原来的 CudaSurface