GPU Resource Management On JDOS 梁永清 liangyongqing1@jd.com 提供的服务 1. 用于实验的 GPU 容器 2.基于 Kubeflow 的机器学习训练服务 3.模型管理和模型 Serving 服务 Experiment Training Serving 均基于容器，不对业务方直接提供 GPU 物理机 GPU 实验 JDOS 常规的容器服务 常规的容器服务 ，使用 gpu 的 zone ， 自行设定相应的镜像即 可，有完善的周边服务 训练服务 • 提供基于 kubeflow 的分布式训练方案 – 界面化操作，用户提供代码地址和执行命令即可 – 系统内建支持安装 pip 依赖 – 自制存储插件支持分布式文件系统存储用户数据 – 支持官方镜像，不需要 JDOS 提前协助制作镜像 – 提供 tensorboard 作为训练监控实时查看训练状态 作为训练监控实时查看训练状态 – 用户训练完成后释放 GPU 资源，提高 GPU 利用率 – Job 调度 （部门 quota 限制 + 优先级） • 创建训练 – 用户选择集群提供代码地址和执行命令即可 – 选择所用框架（镜像）：支持官方，亦可自制 （提供 dockerfile 生成镜像服务） – 选择存储来源：对接了内部的存储 – 填写代码地址，执行的命令等 – 可以选择是否监控训练，提供

CUDA 开启的 GPU 编程 by 彭于斌（ @archibate ） 往期录播： https://www.bilibili.com/video/BV1fa411r7zp 课程 PPT 和代码： https://github.com/parallel101/course 前置条件 • 学过 C/C++ 语言编程。 • 理解 malloc/free 之类的概念。 • 熟悉 STL 中的容器、函数模板等。 做不到的。 编写一段在 GPU 上运行的代码 • 定义函数 kernel ，前面加上 __global__ 修 饰符，即可让他在 GPU 上执行。 • 不过调用 kernel 时，不能直接 kernel() ，而 是要用 kernel<<<1, 1>>>() 这样的三重尖括 号语法。为什么？这里面的两个 1 有什么用 ？稍后会说明。 • 运行以后，就会在 GPU 上执行 printf 了。 kernel 函数在 GPU 上执行，称为核 函数，用 __global__ 修饰的就是核函数。 没有反应？同步一下！ • 然而如果直接编译运行刚刚那段代码，是不会打印出 Hello, world! 的。 • 这是因为 GPU 和 CPU 之间的通信，为了高效，是异 步的。也就是 CPU 调用 kernel<<<1, 1>>>() 后，并不 会立即在 GPU 上执行完毕，再返回。实际上只是把

cudaFree(x); cudaFree(y); } An Even Easier Introduction to CUDA 5 |__global__ void add_gpu(int n, float* x, float* y) { for (int i = 0; i < n; i++) y[i] = x[i] + y[i]; } TEST_CASE("cppcon-1" TEST_CASE("cppcon-1", "[CUDA]") { // … } An Even Easier Introduction to CUDA 6 |__global__ void add_gpu(int n, float* x, float* y) { for (int i = 0; i < n; i++) y[i] = x[i] + y[i]; } TEST_CASE("cppcon-1" 20; float* x; float* y; // … add_gpu<<<1, 1>>>(N, x, y); // … } An Even Easier Introduction to CUDA 7 |__global__ void add_gpu(int n, float* x, float* y) { for (int i = 0;

编译器如何自动优化：从汇编角度看 C++ 5.C++11 起的多线程编程：从 mutex 到无锁并行 6.并行编程常用框架： OpenMP 与 Intel TBB 7.被忽视的访存优化：内存带宽与 cpu 缓存机制 8.GPU 专题： wrap 调度，共享内存， barrier 9.并行算法实战： reduce ， scan ，矩阵乘法等 10.存储大规模三维数据的关键：稀疏数据结构 11.物理仿真实战：邻居搜索表实现 位时代过去了） 至少 2 核 4 线程（并行课…） 英伟达家显卡（ GPU 专题） 软件要求： Visual Studio 2019 （ Windows 用户） GCC 9 及以上（ Linux 用户） CMake 3.12 及以上（跨平台作业） Git 2.x （作业上传到 GitHub ） CUDA Toolkit 10.0 以上（ GPU 专题） 第 0 章：从并发到并行 摩尔定律：停止增长了吗？ ，其中 n 是元素个数 改进的并行缩并（ GPU ） • 刚才那种方式对 c 比较大的情况不友好， 最后一个串行的 for 还是会消耗很多时间 。 • 因此可以用递归的模式，每次只使数据缩 小一半，这样基本每次都可以看做并行的 for ，只需 log2(n) 次并行 for 即可完成 缩并。 • 这种常用于核心数量很多，比如 GPU 上 的缩并。 结论：改进后的并行缩并的时间复杂度为

cudaExtent 指定，方 便起见我们的 C++ 封装类用了 uint3 表示 大小。 • GPU 的多维数组有特殊的数据排布来保障 访存的高效，和我们 CPU 那样简单地行主 序或列主序（如 a[x + nx * y] ）的多维数组 不一样。 • 随后可用 cudaMemcpy3D 在 GPU 的三 维数组和 CPU 的三维数组之间拷贝数据。 CUDA 表面对象：封装 • 要访问一个多维数组，必须先创建一个表面对象 surf3Dwrite 来读写 表面对象中的元素， x,y,z 参数指定要访问元素的坐标 ，要注意 x 必须乘以 sizeof( 元素类型 ) ，否则出错。 • 这里用了访问者模式（ Accessor ， GPU 编程常用）。 原来的 CudaSurface 管理资源，禁止拷贝。然后单独 弄一个访问者类 CudaSurfaceAccessor ，不管理资源 ，仅仅是指向资源的一个弱引用，可以随意拷贝。并把 表面对象保障了高效的访存，并且自动判断越界，体 现了 GPU 作为图形学专业硬件的能力。 CUDA 纹理对象：封装 • 表面对象访问数组是可读可写的。纹理对象也可以访问 数组，不过是只读的。好处是他可以通过浮点坐标来访 问，且提供了线性滤波的能力。 • 在核函数中可以通过 tex3D 来读取纹理中的值。 • 之所以纹理是因为 GPU 一开始是渲染图形的专用硬件 ，会用到一些贴图等，这就是二维的纹理。

，成功算对结果了，代价是精度损失了 不少。 • 其实 GPU 存储贴图一般也是用的定点数 uint8_t （范围从 0 到 255 ），着色器在读取的时候才会把他 转换成 float （范围从 0.0 到 1.0 ）。这就是浮点数的 量化，存储时转换成低精度的定点数，读取时再转换 回高精度的浮点数，从而节省 4 倍内存带宽，提升 GPU 性能。 有没有更小的浮点类型？ • 浮点数在接近 float16 （大底数版） • 这就是 half 类型，他只有 5 位指数， 10 位底 数，总共占据 16 位，所以又称 float16 。精度很 低，但是节省内存空间！ • 然而只有 GPU （比如 CUDA ）支持 half 类型 ， CPU 需要支持 AVX512fp16 这个扩展才能用 。 • 据说深度学习（很多都是 membound ）很喜欢 用 half ，因为可以省一半内存，从而加快一倍。 还支持自适应的网格 SPGrid 的利弊 • 优点：平坦直观，适合插桩，顺序访问，自适应网格 。 • 缺点：尺寸受限，操作系统挂钩，依赖 x86 硬件机 制。 • 顺便一提， GPU 也可以搞 SPGrid ，不过 GPU 的 页大小是 2MB ，王鑫磊最近研究过这个，因为太繁 琐而放弃了。 http://pages.cs.wisc.edu/~sifakis/papers/SPGrid.pdf

编译器如何自动优化：从汇编角度看 C++ 5.C++11 起的多线程编程：从 mutex 到无锁并行 6.并行编程常用框架： OpenMP 与 Intel TBB 7.被忽视的访存优化：内存带宽与 cpu 缓存机制 8.GPU 专题： wrap 调度，共享内存， barrier 9.并行算法实战： reduce ， scan ，矩阵乘法等 10.存储大规模三维数据的关键：稀疏数据结构 11.物理仿真实战：邻居搜索表实现 位时代过去了） 至少 2 核 4 线程（并行课…） 英伟达家显卡（ GPU 专题） 软件要求： Visual Studio 2019 （ Windows 用户） GCC 9 及以上（ Linux 用户） CMake 3.12 及以上（跨平台作业） Git 2.x （作业上传到 GitHub ） CUDA Toolkit 10.0 以上（ GPU 专题） 温馨提示： 1. 会用到第二讲（ RAII 与智能指针）里的知识 defer_lock 做参数 ， owns_lock() 判断等，同学们自己研究 。 只需一次性上锁，且符合 RAII 思想：访问者模式 Accessor 或者说 Viewer 模式，王鑫磊常用于设计 GPU 容器 OpenVDB 数据结构的访问，也是采用了 Accessor 的设计…… 并且还有 ConstAccessor 和 Accessor 两种，分别对应于读和 写 同学们可以想想看，如果这里的

编译器如何自动优化：从汇编角度看 C++ 5.C++11 起的多线程编程：从 mutex 到无锁并行 6.并行编程常用框架： OpenMP 与 Intel TBB 7.被忽视的访存优化：内存带宽与 cpu 缓存机制 8.GPU 专题： wrap 调度，共享内存， barrier 9.并行算法实战： reduce ， scan ，矩阵乘法等 10.存储大规模三维数据的关键：稀疏数据结构 11.物理仿真实战：邻居搜索表实现 位时代过去了） 至少 2 核 4 线程（并行课…） 英伟达家显卡（ GPU 专题） 软件要求： Visual Studio 2019 （ Windows 用户） GCC 9 及以上（ Linux 用户） CMake 3.12 及以上（跨平台作业） Git 2.x （作业上传到 GitHub ） CUDA Toolkit 10.0 以上（ GPU 专题） 关于作者 • 我是 Taichi 编译器的贡献者之一（

编译器如何自动优化：从汇编角度看 C++ 5.C++11 起的多线程编程：从 mutex 到无锁并行 6.并行编程常用框架： OpenMP 与 Intel TBB 7.被忽视的访存优化：内存带宽与 cpu 缓存机制 8.GPU 专题： wrap 调度，共享内存， barrier 9.并行算法实战： reduce ， scan ，矩阵乘法等 10.存储大规模三维数据的关键：稀疏数据结构 11.物理仿真实战：邻居搜索表实现 位时代过去了） 至少 2 核 4 线程（并行课…） 英伟达家显卡（ GPU 专题） 软件要求： Visual Studio 2019 （ Windows 用户） GCC 9 及以上（ Linux 用户） CMake 3.12 及以上（跨平台作业） Git 2.x （作业上传到 GitHub ） CUDA Toolkit 10.0 以上（ GPU 专题） 第 0 章：汇编语言 x64 架构下的寄存器模型

编译器如何自动优化：从汇编角度看 C++ 5.C++11 起的多线程编程：从 mutex 到无锁并行 6.并行编程常用框架： OpenMP 与 Intel TBB 7.被忽视的访存优化：内存带宽与 cpu 缓存机制 8.GPU 专题： wrap 调度，共享内存， barrier 9.并行算法实战： reduce ， scan ，矩阵乘法等 10.存储大规模三维数据的关键：稀疏数据结构 11.物理仿真实战：邻居搜索表实现 位时代过去了） 至少 2 核 4 线程（并行课…） 英伟达家显卡（ GPU 专题） 软件要求： Visual Studio 2019 （ Windows 用户） GCC 9 及以上（ Linux 用户） CMake 3.12 及以上（跨平台作业） Git 2.x （作业上传到 GitHub ） CUDA Toolkit 10.0 以上（ GPU 专题） 为什么需要模板函数（ template ） • 避免重复写代码。