C++高性能并行编程与优化 - 课件 - 08 CUDA 开启的 GPU 编程如需总的线程编号: blockDim * blockIdx + threadIdx 分离 __device__ 函数的声明和定义:出错 • 默认情况下 GPU 函数必须定义在同一个文件里。 如果你试图分离声明和定义,调用另一个文件里 的 __device__ 或 __global__ 函数,就会出错 。 分离 __device__ 函数的声明和定义:解决 • 开启 CMAKE_CUDA_ CMAKE_CUDA_SEPARABLE_COMPILATION 选 项(设为 ON ),即可启用分离声明和定义的支持。 • 不过我还是建议把要相互调用的 __device__ 函数放在 同一个文件,这样方便编译器自动内联优化(第四课讲 过)。 两种开启方式:全局有效 or 仅针对单个程序 只对 main 这个程序启用: 对下方所有的程序启用(推荐): 顺便一提, CXX_STANDARD 和 CUDA_ARCHITECTURES cudaMallocManaged 、 cudaFree • 如果我没记错的话,统一内存是从 Pascal 架构开始支持的,也就是 GTX9 开头及以上 。 • 虽然方便,但并非完全没有开销,有条件的话还是尽量用分离的设备内存和主机内存吧。 第 3 章:数组 分配数组 • 如 malloc 一样,可以用 cudaMalloc 配 合 n * sizeof(int) ,分配一个大小为 n 的 整型数组。这样就会有0 码力 | 142 页 | 13.52 MB | 1 年前3
谈谈MYSQL那点事lock , Rows level lock , 读写性能都非常优秀 读写性能都非常优秀 • 能够承载大数据量的存储和访问 能够承载大数据量的存储和访问 • 拥有自己独立的缓冲池,能够缓存数据和索引 拥有自己独立的缓冲池,能够缓存数据和索引 MySQL 架构设计—应用架构 强一致性 对读一致性的权衡,如果是对读写实时性要求非常高的话, 就将读写都放在 M1 上面, M2 只是作为 standby SAN ) ) 设计合理架构,如果 设计合理架构,如果 MySQL MySQL 访问频繁,考虑 访问频繁,考虑 Master/Slave Master/Slave 读写分离;数据库分表、数据库切片(分 读写分离;数据库分表、数据库切片(分 布式),也考虑使用相应缓存服务帮助 布式),也考虑使用相应缓存服务帮助 MySQL MySQL 缓解访问 缓解访问 压力 压力 系统优化 系统优化0 码力 | 38 页 | 2.04 MB | 1 年前3
C++高性能并行编程与优化 - 课件 - 04 从汇编角度看编译器优化而我们可以用 const 禁止写入访问。 结论:所有非 const 的指针都声明 __restrict 。 禁止优化: volatile 结论:加了 volatile 的对象,编 译器会放弃优化对他的读写操作 。 做性能实验的时候非常有用。 注意一下区别 1. volatile int *a 或 int volatile *a 2. int *__restrict a • 语法上区别: volatile a 总是对齐到 16 字节,在 GCC 编译器中这样 写: 但这样不通用,因此 C++20 引入了标准化的 std::assume_aligned : movups 变成了 movaps 对齐的读写可能 带来微乎其微的 性能提升…… 数组求和: reduction 的优化 你看懂了吗?没关系!小彭老师也没看 懂!总之非常高效就对了! 第 5 章:循环 循环中的矢量化:还在伺候指针别名 结构体的内存布局: AOS 与 SOA • AOS ( Array of Struct )单个对象的属性紧挨着存 • xyzxyzxyzxyz • SOA ( Struct of Array )属性分离存储在多个数组 • xxxxyyyyzzzz • AOS 必须对齐到 2 的幂才高效, SOA 就不需要。 • AOS 符合直觉,不一定要存储在数组这种线性结构, 而 SOA 可能无法保证多个数组大小一致。0 码力 | 108 页 | 9.47 MB | 1 年前3
C++高性能并行编程与优化 - 课件 - 07 深入浅出访存优化com/parallel101/course 为什么往 int 数组里赋值 1 比赋值 0 慢一倍? 第 1 章:内存带宽 cpu-bound 与 memory-bound • 通常来说,并行只能加速计算的部分,不能加速内存读写的部分 。 • 因此,对 fill 这种没有任何计算量,纯粹只有访存的循环体,并 行没有加速效果。称为内存瓶颈( memory-bound )。 • 而 sine 这种内部需要泰勒展开来计算,每次迭代计算量很大的 在太少了。 • 计算太简单,数据量又大,并行只带来了多线程调度的额外开销 。 • 小彭老师经验公式: 1 次浮点读写 ≈ 8 次浮点加法 • 如果矢量化成功( SSE ): 1 次浮点读写 ≈ 32 次浮点加法 • 如果 CPU 有 4 核且矢量化成功: 1 次浮点读写 ≈ 128 次浮点加 法 常见操作所花费的时间 • 图中加法 (add) 和乘法 (mul) 都指的整数。 • 用了 6 核才饱和。 • 结论:要想利用全部 CPU 核心,避免 mem-bound ,需要 func 里有足够的计算 量。 • 当核心数量越多, CPU 计算能力越强,相 对之下来不及从内存读写数据,从而越容 易 mem-bound 。 1 2 4 6 8 10 0 50 100 150 200 250 300 350 funcA funcB funcC 内存信息查看工具:0 码力 | 147 页 | 18.88 MB | 1 年前3
C++高性能并行编程与优化 - 课件 - 05 C++11 开始的多线程编程。 • 所以, download 函数才会出师未捷身先死 ——还没开始执行他的线程就被销毁了。 解构函数不再销毁线程: t1.detach() • 解决方案:调用成员函数 detach() 分离该 线程——意味着线程的生命周期不再由当 前 std::thread 对象管理,而是在线程退 出以后自动销毁自己。 • 不过这样还是会在进程退出时候自动退出 。 解构函数不再销毁线程:移动到全局线程池 上仍是 const 的。因此,为了让 this 为 const 时仅仅给 m_mtx 开后门,可以用 mutable 关键字修饰他,从而所有成员里 只有他不是 const 的。 为什么需要读写锁? • 刚才说过 mutex 就像厕所,同一时刻只有一个人能上。但是如果“上”有两种方式呢? • 假设在平行世界,厕所不一定是用来拉的,还可能是用来喝的(只是打个比方,请勿尝试) • 喝厕所里的水时,可以多个人插着吸管一起喝。 结论:读可以共享,写必须独占,且写和读不能共存。 • 针对这种更具体的情况,又发明了读写锁,他允许的状态有: 1. n 个人读取,没有人写入。 2. 1 个人写入,没有人读取。 3. 没有人读取,也没有人写入。 读写锁: shared_mutex • 为此,标准库提供了 std::shared_mutex 。 • 上锁时,要指定你的需求是拉还是喝,负 责调度的读写锁会帮你判断要不要等待。 • 这里 push_back()0 码力 | 79 页 | 14.11 MB | 1 年前3
C++高性能并行编程与优化 - 课件 - 09 CUDA C++ 流体仿真实战cudaSurfaceObject_t )。 • 考虑到多维数组始终是需要通过表面对象来访问的,这 里我们让表面对象继承自多维数组。 • 在核函数中可以用 surf3Dread 和 surf3Dwrite 来读写 表面对象中的元素, x,y,z 参数指定要访问元素的坐标 ,要注意 x 必须乘以 sizeof( 元素类型 ) ,否则出错。 • 这里用了访问者模式( Accessor , GPU 编程常用)。 编程常用)。 原来的 CudaSurface 管理资源,禁止拷贝。然后单独 弄一个访问者类 CudaSurfaceAccessor ,不管理资源 ,仅仅是指向资源的一个弱引用,可以随意拷贝。并把 读写访问的方法( surf3Dread )定义在访问者类。 CUDA 表面对象:封装 • 此外,表面对象还支持自动判断 x,y,z 坐标是否越界 , surf3Dread/write 的最后一个参数,用于指定出现 新采样。核函数的 gridDim 通过上整除技巧保证每个元素都能访问到, blockDim 为 8x8x8=512 。 • 如果在 resample_kernel 需要读取 clr ,然后再写入 clr ,并且读写是不同的坐标位置。 • 因此对 clr 和 vel 使用了双缓冲,写入 clrNext 的同时读取 clr 没有冲突,写入完毕后对调 clrNext 和 clr 。 投影部分 投影部分 •0 码力 | 58 页 | 14.90 MB | 1 年前3
Rust分布式账务系统 - 胡宇TPS 的流量 可演化性:业务逻辑与底层 API 解耦,当业务发生改变 时,底层 API 不用改变 分布式账务系统 设计理念 - Rust 是我们可靠的基石 分布式账务系统 存算分离 API 解耦 读写分离 层级账号 Rust ● 事务层与账户层分 离 ● 独立水平扩展 ● CQRS ● Event Sourcing ● 针对读场景,写场 景分别优化 ● 稳定的底层 API0 码力 | 27 页 | 12.60 MB | 1 年前3
新一代分布式高性能图数据库的构建 - 沈游人,可扩展的分析引擎支持更复 杂的数据挖掘和机器学习场景 MPP Massively Parallel Processing 架构,大规模集群 分布式存储及并行计 算, Shared Nothing 模式支 持存储计算分离 高性能 基于 Rust 开发的分布式存储引 擎及图计算引擎,精细的内存 管理设计,内置索引系统,支 持毫秒级的并发查询响应速度 易用 AQL(Atlas Graph Query Language) u32 u32 string string 定长 变长 高可用技术方案 基于 Chain Replication ( CRAQ ) 算法实现,进行数据副本处理,头 结点写,多结点读,支持读写分离 ,提供更好的并发查询能力 数据高可用实现 Chain Replication 数据高可用方案 服务高可用实现 系统中 Meta , TS 服务采用主备架 构,基于 Raft 算法实现租约,进行0 码力 | 38 页 | 24.68 MB | 1 年前3
C++高性能并行编程与优化 - 课件 - 10 从稀疏数据结构到量化数据类型疏网格、位运算、浮点的二进制格式、内存带宽优 化 面向人群:图形学、 CFD 仿真、深度学习编程人 员 第 0 章:稀疏矩阵 稠密数组存储矩阵 用 foreach 包装一下枚举的过程 改用 map 来存储 分离 read/write/create 三种访问模式 foreach 直接给出当前坐标指向的值 改用 unordered_map 来存储 unordered_map 手动 read(i, j) 也一样速度 int64_t :每个占据 8 字节 • 如果用更大的数据类型,用时会直接提升两倍! • 这是因为 i % 2 的计算时间,完全隐藏在内存 的超高延迟里了。 • 可见,当数据量足够大,计算量却不多时,读写 数据量的大小唯一决定着你的性能。 • 特别是并行以后,计算量可以被并行加速,而访 存却不行。 使用 int8_t :每个占据 1 字节 • 因此我们可以把数据类型变小,这样所需的内存 量就变小,从而内存带宽也可以减小! 类操作系统的 mmap 函数,他位于 sys/mman.h 头文件里。 • Windows 可以用 VirtualAllocateEx 之类。 • mmap 出来的起始地址保证是对齐到 4KB 的,读写访问其 中偏移地址时,会按页的粒度自动分配和释放内存,从而满 足稀疏数据结构“按需分配”的需求。且由于分页是硬件自动 来做的,比我们软件哈希和指针数组的稀疏更高效,写起来 就和普通的二维数组没什么两样,就好像顺序访问。也用不0 码力 | 102 页 | 9.50 MB | 1 年前3
C++高性能并行编程与优化 - 课件 - 02 现代 C++ 入门:RAII 内存管理这种情况出现时,就意味着你需要把成员变量的读写封装为成员函数 不变性:请勿滥用封装 • 仅当出现“修改一个成员时,其他也成员要 被修改,否则出错”的现象时,才需要 getter/setter 封装。 • 各个成员之间相互正交,比如数学矢量类 Vec3 ,就没必要去搞封装,只会让程序员 变得痛苦,同时还有一定性能损失:特别 是如果 getter/setter 函数分离了声明和定 义,实现在另一个文件时!0 码力 | 96 页 | 16.28 MB | 1 年前3
共 19 条
- 1
- 2













