 C++高性能并行编程与优化 -  课件 - 08 CUDA 开启的 GPU 编程当前板块的编号: blockIdx • 总的板块数量: gridDim • 线程 (thread) :并行的最小单位 • 板块 (block) :包含若干个线程 • 网格 (grid) :指整个任务,包含若干个板块 • 从属关系:线程<板块<网格 • 调用语法: << C++高性能并行编程与优化 -  课件 - 08 CUDA 开启的 GPU 编程当前板块的编号: blockIdx • 总的板块数量: gridDim • 线程 (thread) :并行的最小单位 • 板块 (block) :包含若干个线程 • 网格 (grid) :指整个任务,包含若干个板块 • 从属关系:线程<板块<网格 • 调用语法: <<- >> 区分板块和线程有点麻烦?“扁平化”他们! • 你可能觉得纳闷,既然已经有线程可以并行了 多个线程,并行地给数组赋值 • 刚刚的 for 循环是串行的,我们可以把线 程数量调为 n ,然后用 threadIdx.x 作为 i 索引。这样就实现了,每个线程负责给数 组中一个元素的赋值。 小技巧:网格跨步循环( grid-stride loop ) • 无论调用者指定了多少个线程 ( blockDim ),都能自动根据给定的 n 区间循环,不会越界,也不会漏掉几个元 素。 • 这样一个 的习惯,又能自动匹配不同 的 blockDim ,看起来非常方便。 从线程到板块 • 核函数内部,用之前说到的 blockDim.x + blockIdx.x + threadIdx.x 来获取线程在整个 网格中编号。 • 外部调用者,则是根据不同的 n 决定板块的 数量( gridDim ),而每个板块具有的线程 数量( blockDim )则是固定的 128 。 • 因此,我们可以用 n / 128 0 码力 | 142 页 | 13.52 MB | 1 年前3
 C++高性能并行编程与优化 -  课件 - 10 从稀疏数据结构到量化数据类型com/video/BV1fa411r7zp 课程 PPT 和代码: https://github.com/parallel101/course 本课涵盖:稀疏矩阵、 unordered_map 、空间稀 疏网格、位运算、浮点的二进制格式、内存带宽优 化 面向人群:图形学、 CFD 仿真、深度学习编程人 员 第 0 章:稀疏矩阵 稠密数组存储矩阵 用 foreach 包装一下枚举的过程 改用 map 按行压缩( Compressed Row Storage ) http://www.netlib.org/linalg/html_templates/node91.html 第 1 章:稀疏网格 稠密网格计算粒子经过的格点数量 改用更小的 char 存储 只用一个 bit 存储,一个 char 可以存储 8 个 bit 用 map 来存储 读取:如果不存在,则读到 0 写入:如果不存在,则创建该表项 块优化。 实际上空间局域 性正是稀疏网格 能够实现的一大 前提,稍后详细 讨论。 在 16x16 分块的基础上,只用一个 bit 存储 图片解释稀疏的好处 传统稠密二维数组 无边界稀疏分块哈希表 有了无边界的稀疏网格,再也不用担心二维数组要分配多大了。 坐标可以无限延伸,甚至可以是负数!比如 (-1,2) 等…… 他会自动在写入时分配 16x16 的子网格,称之为叶节点 (leaf node)0 码力 | 102 页 | 9.50 MB | 1 年前3 C++高性能并行编程与优化 -  课件 - 10 从稀疏数据结构到量化数据类型com/video/BV1fa411r7zp 课程 PPT 和代码: https://github.com/parallel101/course 本课涵盖:稀疏矩阵、 unordered_map 、空间稀 疏网格、位运算、浮点的二进制格式、内存带宽优 化 面向人群:图形学、 CFD 仿真、深度学习编程人 员 第 0 章:稀疏矩阵 稠密数组存储矩阵 用 foreach 包装一下枚举的过程 改用 map 按行压缩( Compressed Row Storage ) http://www.netlib.org/linalg/html_templates/node91.html 第 1 章:稀疏网格 稠密网格计算粒子经过的格点数量 改用更小的 char 存储 只用一个 bit 存储,一个 char 可以存储 8 个 bit 用 map 来存储 读取:如果不存在,则读到 0 写入:如果不存在,则创建该表项 块优化。 实际上空间局域 性正是稀疏网格 能够实现的一大 前提,稍后详细 讨论。 在 16x16 分块的基础上,只用一个 bit 存储 图片解释稀疏的好处 传统稠密二维数组 无边界稀疏分块哈希表 有了无边界的稀疏网格,再也不用担心二维数组要分配多大了。 坐标可以无限延伸,甚至可以是负数!比如 (-1,2) 等…… 他会自动在写入时分配 16x16 的子网格,称之为叶节点 (leaf node)0 码力 | 102 页 | 9.50 MB | 1 年前3
 C++高性能并行编程与优化 -  课件 - 07 深入浅出访存优化矢量化的话可能还是要 SOA 或 AOSOA ,比如 hw04 那种的。而 “ pos 和 vel 应该用 SOA 分开存”是没问题的。 • 而且 SOA 在遇到存储不是 vector ,而是稀疏的哈希网格之类索引有一定 开销的数据结构,可能就不适合了。这就是为什么王鑫磊最喜欢 AOSOA :在高层保持 AOS 的统一索引,底层又享受 SOA 带来的矢量化 和缓存行预取等好处……就是随机索引比较麻烦。 其实操作系统惰性分配的特性,也是 SPGrid ( Sparsely-Paged-Grid )得以实现的基础 ,他利用 mmap 分配比机器大得多的内存(比如 2048*2028*1024 的三维网格),然后 在里面索引,这样就相当于利用硬件的分页机制实现了稀疏数据结构,既能高效利用内存 ,随机访问和插桩又特别高效。有兴趣可以研究一下他们的论文,也用了莫顿序增强 TLB 和缓存的局域性,非常精彩。 元素。图 像处理中用到的模糊操作,也需要向上下左右扩散 w 个单位。 • 这样假如当前元素正好在二维网格的边界上,那再往外索引 w 个单位就会超出数组的界限,导致出错。对此有多种解决办法 : 1. 使用 std::max(n, std::min(0, i)) 限制索引不要超出二维网格的 大小,但会导致难以 SIMD 矢量化。 2. 使用 (i + n) % n 让索引在边界产生回绕,但是更加低效也难0 码力 | 147 页 | 18.88 MB | 1 年前3 C++高性能并行编程与优化 -  课件 - 07 深入浅出访存优化矢量化的话可能还是要 SOA 或 AOSOA ,比如 hw04 那种的。而 “ pos 和 vel 应该用 SOA 分开存”是没问题的。 • 而且 SOA 在遇到存储不是 vector ,而是稀疏的哈希网格之类索引有一定 开销的数据结构,可能就不适合了。这就是为什么王鑫磊最喜欢 AOSOA :在高层保持 AOS 的统一索引,底层又享受 SOA 带来的矢量化 和缓存行预取等好处……就是随机索引比较麻烦。 其实操作系统惰性分配的特性,也是 SPGrid ( Sparsely-Paged-Grid )得以实现的基础 ,他利用 mmap 分配比机器大得多的内存(比如 2048*2028*1024 的三维网格),然后 在里面索引,这样就相当于利用硬件的分页机制实现了稀疏数据结构,既能高效利用内存 ,随机访问和插桩又特别高效。有兴趣可以研究一下他们的论文,也用了莫顿序增强 TLB 和缓存的局域性,非常精彩。 元素。图 像处理中用到的模糊操作,也需要向上下左右扩散 w 个单位。 • 这样假如当前元素正好在二维网格的边界上,那再往外索引 w 个单位就会超出数组的界限,导致出错。对此有多种解决办法 : 1. 使用 std::max(n, std::min(0, i)) 限制索引不要超出二维网格的 大小,但会导致难以 SIMD 矢量化。 2. 使用 (i + n) % n 让索引在边界产生回绕,但是更加低效也难0 码力 | 147 页 | 18.88 MB | 1 年前3
 C++高性能并行编程与优化 -  课件 - 09 CUDA C++ 流体仿真实战迭代因为需要写入 pre 的同时读取 pre ,所以也要用双缓冲。 投影部分:计算未消除的散度 为了评估效果的好坏,额外加一个计算散度方差的核函数,看看是不是无散度(不可压缩流)了。 多重网格法 投影部分:多重网格实现 投影部分:红黑高斯 投影部分:计算残差 投影部分:缩小一倍 投影部分:清零数组 投影部分:扩大一倍 创建与导出 主函数:创建场景 导出 VDB :调用接口 导出 VDB OpenVDB 在 Blender 中查看导出的结果 边界条件 边界条件:初始化 边界条件:添加判断边界的版本 边界条件:仅在第一层额外判断边界条件 进一步改进 VDB 导出:支持导出多个网格,并指定名称 进一步改进 VDB 导出: P-IMPL 模式 进一步改进 VDB 导出: F-IMPL 模式 Blender 渲染结果 改进 改进边界条件:外部边界流出而不是反弹,内部边界可以流出速度0 码力 | 58 页 | 14.90 MB | 1 年前3 C++高性能并行编程与优化 -  课件 - 09 CUDA C++ 流体仿真实战迭代因为需要写入 pre 的同时读取 pre ,所以也要用双缓冲。 投影部分:计算未消除的散度 为了评估效果的好坏,额外加一个计算散度方差的核函数,看看是不是无散度(不可压缩流)了。 多重网格法 投影部分:多重网格实现 投影部分:红黑高斯 投影部分:计算残差 投影部分:缩小一倍 投影部分:清零数组 投影部分:扩大一倍 创建与导出 主函数:创建场景 导出 VDB :调用接口 导出 VDB OpenVDB 在 Blender 中查看导出的结果 边界条件 边界条件:初始化 边界条件:添加判断边界的版本 边界条件:仅在第一层额外判断边界条件 进一步改进 VDB 导出:支持导出多个网格,并指定名称 进一步改进 VDB 导出: P-IMPL 模式 进一步改进 VDB 导出: F-IMPL 模式 Blender 渲染结果 改进 改进边界条件:外部边界流出而不是反弹,内部边界可以流出速度0 码力 | 58 页 | 14.90 MB | 1 年前3
 C++高性能并行编程与优化 - 课件 - Zeno 中的现代 C++ 最佳实践 态的虚函数,这就是被小彭老师称为自动 虚克隆 (auto-vitrual-clone) 的大法。 Zeno 中对 OpenVDB 对象的封装 • 开源的体积数据处理库 OpenVDB 中有许多“网格”的类(可以理解为多维数组),例如: • openvdb::Vec3fGrid , FloatGrid , Vec3IGrid , IntGrid , PointsDataGrid • 我们并不 之类的任意表达式,只要语义上通过,就可以实例化。 • (把 sort 封装成虚函数,留作回家作业) Zeno 中对 OpenVDB 的类型擦除 • 结合类型擦除技术,自动虚克隆技术。 • VDBGrid 作为所有网格类的基类提供各个 操作做为虚函数, VDBGridWrapper 则是 那个实现了擦除的包装类。 Zeno 节点系统 • 节点在 Zeno 中所扮演的角色,实际上相当于函数式编程中的函数。0 码力 | 54 页 | 3.94 MB | 1 年前3 C++高性能并行编程与优化 - 课件 - Zeno 中的现代 C++ 最佳实践 态的虚函数,这就是被小彭老师称为自动 虚克隆 (auto-vitrual-clone) 的大法。 Zeno 中对 OpenVDB 对象的封装 • 开源的体积数据处理库 OpenVDB 中有许多“网格”的类(可以理解为多维数组),例如: • openvdb::Vec3fGrid , FloatGrid , Vec3IGrid , IntGrid , PointsDataGrid • 我们并不 之类的任意表达式,只要语义上通过,就可以实例化。 • (把 sort 封装成虚函数,留作回家作业) Zeno 中对 OpenVDB 的类型擦除 • 结合类型擦除技术,自动虚克隆技术。 • VDBGrid 作为所有网格类的基类提供各个 操作做为虚函数, VDBGridWrapper 则是 那个实现了擦除的包装类。 Zeno 节点系统 • 节点在 Zeno 中所扮演的角色,实际上相当于函数式编程中的函数。0 码力 | 54 页 | 3.94 MB | 1 年前3
 使用硬件加速Tokio - 戴翔Cathy.Lu Loong Dai • Intel 云原生工程师 • 微软 MVP • Dapr 、 Thanos 、 Golangci-lint 的 Maintainer • 现在主要专注于服务网格领域,探索云原生软硬件结 合新范式 • Github ID: daixiang0 自我介绍 Cathy Lu • Intel 软件工程师 • 专注于 NFV, 电信网络云化等方案 Rust0 码力 | 17 页 | 1.66 MB | 1 年前3 使用硬件加速Tokio - 戴翔Cathy.Lu Loong Dai • Intel 云原生工程师 • 微软 MVP • Dapr 、 Thanos 、 Golangci-lint 的 Maintainer • 现在主要专注于服务网格领域,探索云原生软硬件结 合新范式 • Github ID: daixiang0 自我介绍 Cathy Lu • Intel 软件工程师 • 专注于 NFV, 电信网络云化等方案 Rust0 码力 | 17 页 | 1.66 MB | 1 年前3
 唐刚 - Use Rust to Develop the Decentralized Open Data Application - RustChinaConf2023Decentralized Open Data Application Mike Tang daogangtang@gmail.com @daogangtang 2023-06-08 ➔ 裁员 ➔ 互联网格局定型 ➔ 平台倒闭,数据丢失 这是一个什么时代? 互联网的终局 创业 -> 种子 -> 天使 -> A -> B -> C -> … -> IPO 创业的目的是? ➔0 码力 | 30 页 | 2.53 MB | 1 年前3 唐刚 - Use Rust to Develop the Decentralized Open Data Application - RustChinaConf2023Decentralized Open Data Application Mike Tang daogangtang@gmail.com @daogangtang 2023-06-08 ➔ 裁员 ➔ 互联网格局定型 ➔ 平台倒闭,数据丢失 这是一个什么时代? 互联网的终局 创业 -> 种子 -> 天使 -> A -> B -> C -> … -> IPO 创业的目的是? ➔0 码力 | 30 页 | 2.53 MB | 1 年前3
 C++高性能并行编程与优化 -  课件 - 06  TBB 开启的并行编程之旅那一份,分配给 (x + y * 3) % 4 号线程。 这样总体来看每个线程分到的块的位置是随机的, 从而由于正太分布数量越大方差越小的特点,每个 线程分到的总工作量大概率是均匀的。 • GPU 上称为网格跨步循环( grid-stride loop )。 1 1 1 1 2 2 2 3 3 4 3 3 4 4 2 4 tbb::static_partitioner 创建了0 码力 | 116 页 | 15.85 MB | 1 年前3 C++高性能并行编程与优化 -  课件 - 06  TBB 开启的并行编程之旅那一份,分配给 (x + y * 3) % 4 号线程。 这样总体来看每个线程分到的块的位置是随机的, 从而由于正太分布数量越大方差越小的特点,每个 线程分到的总工作量大概率是均匀的。 • GPU 上称为网格跨步循环( grid-stride loop )。 1 1 1 1 2 2 2 3 3 4 3 3 4 4 2 4 tbb::static_partitioner 创建了0 码力 | 116 页 | 15.85 MB | 1 年前3
共 8 条
- 1













