1/66Bridging the Gap: Writing Portable Programs for CPU and GPU using CUDA Thomas Mejstrik Sebastian Woblistin 2/66Content 1 Motivation Audience etc.. Cuda crash course Quiz time 2 Patterns Oldschool Motivation Patterns The dark path Cuda proposal Thank you Why write programs for CPU and GPU Difference CPU/GPU Algorithms are designed differently Latency/Throughput Memory bandwidth Number of cores Motivation Patterns The dark path Cuda proposal Thank you Why write programs for CPU and GPU Difference CPU/GPU Why it makes sense? Library/Framework developers Embarrassingly parallel algorithms User

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 上执行完毕，再返回。实际上只是把

http://wongmichael.com/about ● C++11 book in Chinese: https://www.amazon.cn/dp/B00ETOV2OQ We build GPU compilers for some of the most powerful supercomputers in the world 34 Nevin “:-)” Liber nliber@anl Attribution 4.0 International License SYCL Single Source C++ Parallel Programming GPU FPGA DSP Custom Hardware GPU CPU CPU CPU Standard C++ Application Code C++ Libraries ML Frameworks give better performance on complex apps and libs than hand-coding AI/Tensor HW GPU FPGA DSP Custom Hardware GPU CPU CPU CPU AI/Tensor HW Other BackendsSYCL 2020 is here! Open Standard for

performance claims, visit www.intel.com/PerformanceIndex or scan the QR code: © Intel Corporation. Intel, the Intel logo, and other Intel marks are trademarks of Intel Corporation or its subsidiaries about future Intel products. - I work in Intel’s research labs. Work described here will involve experimental prototypes and early research.Problem: writing parallel programs is hard - Multi-GPU, multi-CPU / execution necessary. CPU NIC GPU GPU GPU GPU Xe LinkMulti-GPU Systems - NUMA regions: - 4+ GPUs - 2+ CPUs CPU NIC GPU GPU GPU GPU Xe LinkMulti-GPU Systems - NUMA regions: - 4+ GPUs

transistors incorporated in a chip will approximately double every 24 months." --Gordon Moore, Intel co-founderMoore’s Law (2/2) 29 ● Around 1965 Gordon Moore predicted the number of transistors months." --Gordon Moore, Intel co-founderDennard Scaling (1/3) "The number of transistors incorporated in a chip will approximately double every 24 months." --Gordon Moore, Intel co-founder http://www-cs-faculty transistors incorporated in a chip will approximately double every 24 months." --Gordon Moore, Intel co-founder http://www-cs-faculty.stanford.edu/~eroberts/cs181/projects/2010-11/TechnologicalSing

iOS, Android, WASM*) Hybrid profiling capabilities (sampling and/or instrumentation) (CPU and GPU instrumentation) Tracing capabilities (values, messages, plots, allocations, …) Hassle-free integration spall https://handmade.network/p/333/spall/ geiger https://github.com/david-grs/geiger Intel IACA https://www.intel.com/content/www/us/en/developer/ articles/tool/architecture-code-analyzer.html“There is experience! Tracy can do it all!Tracy Profiler GUI 13Tracy Profiler GUI 14 Frame Info Menu bar GPU Timeline (per “device”) CPU Timeline (per-thread) Custom Plots & Allocation Trackers15Tracy Client

macOS, iOS, Android, WASM*) Hybrid profiling capabilities (sampling and/or instrumentation) (CPU and GPU instrumentation) Tracing capabilities (values, messages, plots, allocations, …) Hassle-free integration https://github.com/david-grs/geiger Xpedite https://github.com/morganstanley/Xpedite Intel IACA https://www.intel.com/content/www/us/en/developer/ articles/tool/architecture-code-analyzer.html“There is transforms the profiling experience!Tracy Profiler GUI 13Tracy Profiler GUI 14 Frame Info Menu bar GPU Timeline (per “device”) CPU Timeline (per-thread) Custom Plots & Allocation Trackers15Tracy Client

platform at Woven Planet Ilya Burylov Principle Engineer at Intel An architect of C++ software solutions for autonomous driving market in Intel Contribution into functional safety MISRA standard Contribution http://wongmichael.com/about ● C++11 book in Chinese: https://www.amazon.cn/dp/B00ETOV2OQ We build GPU compilers for some of the most powerful supercomputers in the world 4 © The Khronos® Group Inc Generation Safety Critical APIs for Graphics, Compute and Display Industry Need for CPU/GPU Acceleration APIs designed to ease system safety certification Rendering Compute Display • Khronos

进阶：模板元编程与函数式编程 4.编译器如何自动优化：从汇编角度看 C++ 5.C++11 起的多线程编程：从 mutex 到无锁并行 6.并行编程常用框架： OpenMP 与 Intel TBB 7.被忽视的访存优化：内存带宽与 cpu 缓存机制 8.GPU 专题： wrap 调度，共享内存， barrier 9.并行算法实战： reduce ， scan ，矩阵乘法等 10.存储大规模三维数据的关键：稀疏数据结构 位时代过去了） 至少 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 上 的缩并。 结论：改进后的并行缩并的时间复杂度为

: Ryzen: 2-64, Epyc 4-64 • Intel: Core i9: 6-18, Xeon: 4-56 • Apple: M1: 4+4Helmut Hlavacs – Combining Co-Routines and Functions into a Job System - CppCon 2021 6 / 39 GPU Reducing the Frametime Physics Physics UI N Logic AI Record S GPU Physics UI N Logic AI Record S Frametime = 1 / FPS Threads GPU Physics UI N Logic AI Record S Frametime Threads P UI N L AI P P L Frametime R R R R R R R R R S GPU Physics UI N Logic AI Record S Physics UI N Logic AI Record S Physics UI N Logic AI GPU P UI N L AI P P L R R R R R R R R S P UI N L