通过高级并行结构实现GPU到CPU的高性能转译和优化

1190高性能GPU到CPU的跨模型编译和优化通过高级并行结构0William S. Moseswmoses@mit.eduMIT CSAIL 美国0Ivan R. Ivanovivanov@m.titech.ac.jp 东京工业大学日本0Jens Domkejens.domke@riken.jp RIKEN 日本0Toshio Endoendo@is.titech.ac.jp 东京工业大学日本0Johannes Doerfertjdoerfert@llnl.govLLNL 美国0Oleksandr Zinenkozinenko@google.comGoogle法国0摘要尽管并行性仍然是性能的主要来源，但随着每一代硬件的推出，架构实现和编程模型也在不断变化，这往往导致昂贵的应用程序重构。大多数性能可移植性工具都需要手动和昂贵的应用程序移植到另一种编程模型。我们提出了一种替代方法，根据Polygeist/MLIR自动地将一个编程模型（CUDA）的程序自动转换为另一个编程模型（CPU线程）。我们的方法包括一种并行结构的表示，使得传统的编译器转换可以透明地应用并且不需要修改，并且支持并行性特定的优化。我们通过对CUDARodinia基准测试套件进行跨编译和优化来评估我们的框架，在多核CPU上实现了58%的几何平均加速比，超过手写的OpenMP代码。此外，我们展示了如何在仅CPU的超级计算机Fugaku上高效运行和扩展PyTorch的CUDA内核，而无需用户干预。我们的PyTorch兼容层利用了转换后的CUDAPyTorch内核，性能优于PyTorch的CPU本地后端2.7倍。0CCS概念: • 软件及其工程 → 编译器；• 计算理论 →并行计算模型。0关键词: Polygeist, MLIR, CUDA, 屏障同步0ACM参考格式: William S. Moses, Ivan R. Ivanov, Jens Domke,Toshio Endo, Johannes Doerfert和Oleksandr Zinenko. 2023.高性能0PPoPP ’23, 2023年2月25日至3月1日, 加拿大蒙特利尔, ©2023 版权由所有者/作者持有. ACM ISBN979-8-4007-0015-6/23/02.https://doi.org/10.1145/3572848.35774750通过高级并行结构的GPU到CPU的转换和优化。In  第28届ACMSIGPLAN年度并行编程原理与实践研讨会（PPoPP ’23）,2023年2月25日至3月1日, 加拿大蒙特利尔, ACM, 纽约, NY, 美国,16页.  https://doi.org/10.1145/3572848.357747501 引言0尽管x86 CPU和NVidiaGPU仍然是计算的主要平台，但定制化和新兴的架构在计算领域中扮演着重要角色。定制版的ARM CPUA64FX甚至被用在顶级超级计算机Fugaku [ 49]中，其高带宽内存有望与GPU的内存竞争。然而，这些架构往往被面向效率的框架和库所忽视。例如，针对英特尔的oneDNN [ 28]后端的PyTorch [ 44]在ARM上的性能预计会下降，甚至富士通定制的oneDNN [ 20]在某些内核上的性能也不具竞争力。这种情况需要实现性能可移植性。已经提出了许多非库的性能可移植性方法，包括语言扩展（如OpenCL [ 14 ]，OpenACC [ 26 ]），并行编程框架（如Kokkos [ 3]），领域专用语言（如Spiral [ 17 ]，Halide [ 47 ]或TensorComprehensions [ 64]）。所有这些方法仍然需要将旧应用程序进行移植，有时甚至需要完全重写，因为语言或底层编程模型存在差异。我们探索了一种基于完全自动化编译器的替代方法，该编译器接受一种编程模型（CUDA）的代码，并生成针对另一种编程模型（CPU线程）的二进制代码。虽然过去已经研究过GPU到CPU的转换 [ 9 ,  23 ,  58]，但很少能够生成高效的代码。事实上，为了CPU进行优化，甚至是通用编译器转换，如公共子表达式消除或循环不变代码移动，都受限于编译器内部对于并行结构的可分析表示的缺乏 [ 39]。由于主流编译器内部的并行性表示只是最近才开始出现，因此我们的工作探索了一种新的方法，该方法基于Polygeist/MLIR自动地将代码从一种编程模型（CUDA）转换为另一种（CPU线程）。我们的方法包括一种并行结构的表示，使得传统的编译器转换可以透明地应用并且不需要修改，并且支持并行性特定的优化。我们通过将CUDARodinia基准测试套件和PyTorch中的内部CUDA内核进行跨编译和优化来评估我们的框架，在多核CPU上实现了58%的几何平均加速比，超过手写的OpenMP代码。此外，我们展示了如何在仅CPU的超级计算机Fugaku上高效运行和扩展PyTorch的CUDA内核，而无需用户干预。我们的PyTorch兼容层利用了转换后的CUDAPyTorch内核，性能优于PyTorch的CPU本地后端2.7倍。0本作品采用知识共享署名4.0国际许可协议进行许可。1200PPoPP ’23, 2023年2月25日至3月1日, 加拿大蒙特利尔, W.S. Moses, I.R. Ivanov, J. Domke, T. Endo, J. Doerfert, 和 O. Zinenko0已经有一些尝试 [ 10 ,  12 ,  32 ,  50 ,  55 ]来改进CPU并行结构的表示，但现有的转换仅限于简单的CPU代码。我们提出了一种针对最常见的GPU构造的编译器模型：多级并行性、级别范围的同步和级别本地内存。与源代码和AST级别的方法不同，这些方法在优化流水线之前进行操作，并且现有的编译器方法将同步建模为黑盒优化屏障，我们从内存语义建模同步。这使得基于同步的代码可以与现有的优化相互操作，并且可以实现新颖的并行特定优化。我们使用MLIR [ 34]和LLVM [ 33]实现了我们的模型，并利用MLIR的嵌套模块方法来支持GPU [21 ]。我们扩展了Polygeist [ 40 ]C/C++前端，以支持CUDA并生成保留高级并行结构的MLIR。我们的原型编译器能够编译PyTorch的CUDA内核，以及其他计算密集型基准测试，支持LLVM支持的任何CPU架构。除了针对执行模型的差异进行转换之外，我们还通过OpenMP利用CPU上的并行性。最后，我们的MocCUDAPyTorch集成允许我们在没有GPU的情况下编译和执行CUDA内核，并替代不支持的调用。我们在Rodinia CUDA基准测试 [5 ]和PyTorchCUDA内核上评估了我们的编译器。当针对一款商用CPU时，我们加速的OpenMPCUDA代码的性能与Rodinia套件中的参考OpenMP实现相当，并且具有更好的可扩展性。当使用我们的框架在仅CPU的Fugaku超级计算机上运行PyTorch时，与现有的PyTorchCPU后端相比，我们每秒处理的图像数量增加了大约两倍。总体而言，我们的论文做出了以下贡献：0•一种普适的高级、与平台无关的SIMT风格并行性表示，由语义定义的屏障同步来支持正确性，从而能够透明地应用现有的优化。0•利用我们的高级并行语义来优化程序的新颖的并行特定优化。0• Polygeist C/C++MLIR前端的扩展，能够直接将GPU和CPU的并行结构映0• 对Rodinia [ 5]基准测试套件的CUDA的端到端跨模型编译1和PyTorch [ 44]中的内部CUDA内核进行跨模型编译，以在仅CPU的Fugaku超级计算机上运行Resnet-50所需。01我们使用术语“跨模型编译”来指代将一个编程模型的程序转换为另一个编程模型的代码，类似于源到源的CUDA到C的转换器，但现在是在中间表示上进行。这个过程还会对代码进行交叉编译，也就是生成非本地指令。02 背景主流编译器如Clang和GCC缺乏统一的高级并行性表示。在CUDA、OpenMP或SYCL等框架中编译并行结构会强制并行区域的主体存在于一个单独的（闭包）函数中，该函数由并行运行时调用。诸如线程索引或同步等概念则通过不透明的内部调用单独表示。由于编译器在历史上缺乏关于并行性和相关运行时效果的信息，任何并行结构都会不经意地成为优化的障碍。虽然近年来已经有一些尝试 [ 10 ,  12 ,  32 ,  39 ,  50 ,  55 ,61 ]改进CPU并行结构的表示，但加速器编程带来了额外的挑战。独特的编程模型和复杂的内存层次结构使得主流编译器内部的GPU并行性的高级表示尚未得到充分发掘。0__device__ float sum(float *data, int n) { ... } __global__void normalize(float *out, float *in, int n) {0int tid = blockIdx.x + blockDim.x * threadIdx.x; //优化：每个块只计算一次总和。 // __shared__ int val; // if(threadIdx.x == 0) val = sum(in, n); // __syncthreads; floatval = sum(in, n); if (tid < n) out[tid] = in[tid] / val; } voidlaunch(int *d_out, int *d_in, int n) { normalize<<<(n +31)/32, 32>>>(d_out, d_in, n); }0图1.一个样本CUDA程序normalize，它对一个向量进行归一化，以及CPU函数launch调用内核。每个GPU线程调用sum，结果为�(� 2 )。使用共享内存（已注释）将工作量减少到�(� 2 / �)，但带来了额外的资源成本。在内核之前计算sum将工作量减少到�(�)。02.1 GPU编译考虑图1中的CUDA程序，它对一个向量进行归一化。使用Clang编译时，GPU程序是一个独立的编译单元。这阻止了GPU内核和CPU调用代码之间的任何优化。在图1的情况下，传统编译器中程序的总工作量为�(� 2)，因为每个线程都执行了�(�)次sum调用。然而，如果只在内核调用之前执行一次sum调用，例如通过进行循环不变代码移动(LICM)，则工作量将减少到�(�)。这种优化的一个不太有效的变种可以将工作量减少到�(� 20� )通过使用共享内存进行优化。MLIR为支持主机/设备代码移动的GPU程序提供了一种嵌套模块表示[ 21]，但并行代码移动还没有被实现。在GPU到CPU代码移动中，LICM出一个并行循环始终是合法的，因为任何以前的设备存储器也可以在主机上使用。}1210通过高级并行构造进行GPU到CPU的转译和优化PPoPP ’23, 2023年2月25日至3月1日，加拿大蒙特利尔0//内核启动在调用函数内部可用，可以在GPU/CPU边界上进行优化。func @launch(%h_out : memref,0%h_in: memref, %n : i64) {0// 并行for循环遍历网格中的所有块。parallel. for (%gx,%gy, %gz) = (0, 0, 0) to (grid.x, grid.y, grid.z) {0// 共享内存 = 块中的堆栈分配。%shared_val =memref.alloca: memref //对块中的所有线程进行并行for循环。parallel. for (%tx, %ty,%tz) = (0, 0, 0) to (blk.x, blk.y, blk.z) {0// 控制流直接保留。if %tx == 0 {0%sum = func.call @sum(%d_in, %n) memref.store %sum,%shared_val[]: memref } //通过显式操作进行同步。polygeist.barrier(%tx, %ty, %tz)%tid = %gx + grid.x * %tx if %tid < %n {0%res = ... store %res, %d_out[%tid]: memref} } } }0图2.Polygeist/MLIR中从图1的launch/normalize代码转换而来的等效代码。内核调用直接在调用宿主代码中可用。并行性在块和线程之间进行了明确的并行for循环。共享内存位于块并行for循环内部，允许来自同一块中的任何线程访问，但不允许来自不同块的线程访问。02.2 MLIR基础设施0MLIR是一个最近设计的可重用和可扩展的编译器基础设施[34]。MLIR不是提供一组预定义的指令和类型，而是在包含可互操作的用户定义操作、属性和类型的dialects集合上操作。操作是IR指令的一般化，可以是任意复杂的，特别是可以包含具有更多IR的region，从而创建一个嵌套表示。操作定义和使用符合单一静态赋值(SSA)的值[ 7 ]。例如，MLIRdialects可以模拟整个指令集，如NVVM（NVidiaGPU的虚拟IR），其他IR，如LLVM IR[ 33]，控制流，如循环，平行编程模型，如OpenMP和OpenACC，机器学习图等。MLIR通过MLIRdialect定义支持GPU，该dialect定义了高级SIMT编程模型、主机/设备通信和一组特定于平台的dialects：NVVM（CUDA）、ROCDL（ROCm）和SPIR-V。MLIR对GPU编程的方法受益于统一的代码表示。由于MLIR模块可以包含其他模块，因此“主机”转换单元可以将“设备”转换单元嵌入为IR，而不是文件引用或二进制blob。这种方法为主机/设备优化提供了其他编译器无法实现的机会，特别是在主机和设备之间移动代码[21]。0__global__ f () {codeA(); barrier();codeB(); }0__0... = A[threadIdx.x]; // R A[i]: i==t.x0图3. 左：一个包含两个任意指令之间的屏障的程序。右：屏障的语义可以通过在除当前线程以外的所有线程中对其上方/下方操作访问的内存地址进行细化。02.3 Polygeist0Polygeist是基于Clang的MLIR的C/C++前端[ 40]。它能够将广泛范围的C++程序转换为保留程序高级结构的MLIRdialects的混合形式。具体来说，Polygeist将结构化控制流（循环和条件语句）保留为MLIR SCFdialect操作，并通过依赖于MLIR的多维度内存引用（memref）类型尽可能地保留多维数组构造，从而简化分析。最后，Poylgeist能够识别适合多面体优化[16]的程序部分，并使用Affine dialect表示它们。03方法0我们扩展了Polygeist编译器[ 40]，以直接从CUDA中发出并行MLIR。这利用了统一的CPU/GPU表示，使优化器能够理解主机/设备执行，并在内核边界上进行优化。使用现有的MLIR的一流并行构造（scf.parallel，affine.parallel），我们能够针对现有的CPU和GPU后端进行优化。最后，MLIR的可扩展操作集允许我们定义具有相关属性和自定义优化的自定义指令。我们如下定义了GPU内核启动的表示（图2中所示）：0• 一个对网格中的所有块进行的三维并行for循环。 •一个对每个块唯一的共享内存进行的堆栈分配。0• 一个对块中的所有线程进行的三维并行for循环。 •一个提供与CUDA同步等效语义的自定义Polygeistbarrier操作。0这个过程使我们能够以保留所需语义的形式来表示任何GPU程序。编译器完全理解这种形式，并且因此可以进行编译器优化。此外，通过使用通用并行性、分配和同步构造来表示GPU程序，我们不仅能够优化原始程序，还可以为不同的体系结构重新定位它。03.1屏障语义0CUDA的__syncthreads函数保证了一个块中的所有线程在执行之前的所有指令后都已经执行完毕}}1220PPoPP ’23，2023年2月25日至3月1日，加拿大蒙特利尔，魏尔斯∙莫西斯，伊万诺夫∙伊万诺夫，多姆克，遠藤，多尔夫特和齐年科。0parallel  %i  =  0  to  10  {0%x  = load data[ %i ] %y= load data[ 2  *  %i ] %a= fmul  %x ,  %x %b  =fmul  %y ,  %y %c  = fsub%x , y barrier call  @use (%a ,  %b ,  %c ) ...0%x_cache  =  memref< 10 x f32> %y_cache  =  memref< 10 xf32 > parallel  %i  =  0  to  10  {0%x  = load data[ %i ] %y  =load data[ 2  *  %i ] store%x ,  %x_cache [ %i ] store%y ,  %y_cache [ %i ] }parallel  %i  =  0  to  10  {0%x  = load  %x_cache [ %i] %y  = load  %y_cache [%i ] %a  = fmul  %x ,  %y%b  = fsub  %y ,  %z call@use ( %a ,  %b ) ...0图4.在屏障周围进行并行循环拆分：将屏障上方的代码放在一个单独的并行“for”循环中，将屏障后面的代码放在另一个循环中。该转换消除了屏障，同时保持了语义。最小剪切算法存储了%x和%y，然后在第二个循环中用于重新计算%a、%b和%c。0R  :  �  →  � .然后，我们组合相关操作的直接和逆关系，获得访问相同下标的线程索引之间的关系，  D  =  R − 1 ◦ R  :  �  →  � ′ .最后，我们减去恒等关系  D \ I  :  �  →  � ′ . 如果非空，  D  ≠  �，不同的线程可能访问相同的地址，需要barrier。对于非仿射访问或非静态控制流，我们保守地假设整个数组维度都被访问。在实践中，这在GPU代码中很少需要，其循环通常具有参数化/静态边界。当涉及多个基地址时，必须检查别名保证。考虑图3（右）中的代码。由于访问的地址集不重叠，即  A �  ∩A �  =  � ，可以允许跨barrier进行代码移动。相反，如果对  A的加载或存储进行了1的偏移，那么barrier是必要的，因为在barrier之后加载的数据将由不同的线程存储。03.2 屏障降低0为了使GPU程序在CPU上运行，我们必须有效地模拟GPU程序的同步行为。而第3.1节的内存语义使我们能够在优化过程中保持barrier的正确性，本节讨论如何在CPU上实现barrier。CPU体系结构没有线程块的概念，也没有等待该概念中的线程组的barrier指令。相反，我们使用常规的CPU线程和工作共享来将线程块循环迭代分配给它们。从概念上讲，这与GPU执行模型不同，其中线程每次执行一个迭代。工作共享要求每个线程按顺序执行多个迭代，这使得在迭代中间进行同步是不可能的，只能在循环结束时进行同步。为了解决这个问题，我们在我们的MLIR表示中开发了一种新的barrier消除技术。我们的方法是1230GPU到CPU的转译和优化通过高级并行构造 PPoPP '23, 2023年2月25日至3月1日, 加拿大蒙特利尔0parallel for  %id = 0  to N {for  %j  =  5  to  0  { if ( %id  < 2 ^ %j ) A[ %id ] += \ A[ %id+  2 ^ %j ] barrier } }0for  %j  =  5  to  0  { parallel for%id = 0  to N { if ( %id  <  2 ^%j ) A[ %id ]+=A[ %id  +  2 ^%j ] barrier } }0图5.左：一个共享内存的加法，它包含了一个内核调用，其中包含一个带有屏障的循环。右：相同的代码，但在一个并行/串行循环替换之后直接在并行循环中使用了屏障。0parallel for  %i = 0  to N {do { run( %i ) barrier }while(condition()) }0%helper  = alloca  memrefscf. do { parallel for  %i = 0  toN { run( %i ) barrier %c  =condition() if  %i  ==  0  { store%c ,  %helper [] } } %c  = load%helper [] } while( %c )0图6.在while循环中进行并行交换。为了保持正确性，每个线程都必须执行condition()函数调用，因此使用了一个辅助变量，它保存了第一个线程上调用的值。0是循环分裂（见第7节）的扩展，结合了两个转换：并行循环拆分和交换。03.2.1并行循环拆分。假设一个屏障具有内核函数（或者在我们的表示中，具有并行循环）作为其直接父级。它可以通过将循环在屏障周围拆分为两个并行循环来消除，分别运行屏障之前和之后的代码。如果屏障之前的代码创建了在屏障之后使用的SSA值，那么这些值必须在第二个并行循环中存储或重新计算。我们使用与[ 41]相似的技术来确定需要存储的最小数据量。具体来说，我们创建所有SSA值的图。然后，我们将无法在屏障之前重新计算的每个值定义（例如从被覆盖的内存加载）标记为源，将在屏障之后使用的值标记为汇。通过在该图上执行最小分支剪切来导出需要存储的最小数据量。03.2.2 并行循环互换 不是所有的屏障操作都有并行 for作为其直接父级，有些可能嵌套在其他控制流操作中。我们创建了一个模型，指定哪些指令可以并行运行。除了屏障之外，我们的表示不需要任何特定的排序或并发性。因此，增加额外的屏障是合法的（尽管可能会减少并行性）。我们可以利用这个特性来实现控制流的屏障降低。0__0__shared__ float  node[HEIGHT]; __shared__ floatweights[HEIGHT][WIDTH]; if (tx == 0) node[ty] =input[index_in]; // 不必要的屏障 #1 __syncthreads ();// 不必要的存储 #1 weights[ty][tx] = hidden[index];__syncthreads ();0// 不必要的加载 #1 weights[ty][tx] = weights[ty][tx] *node[ty]; __syncthreads ();0fo0if (ty % pow(2, i) == 0) weights[ty][tx] += weights[ty + pow(2, i0__syncthreads (); }0hidden[index] = weights[ty][tx]; //不必要的屏障 #2 __syncthreads ();0if (tx == 0) out[by * hid + ty] = weights[tx][ty]; }0图7. 一个包含不必要同步和不必要使用共享内存的 Rodinia反向传播测试中的示例 CUDA 内核。0考虑一个包含屏障并嵌套在并行 for 中的控制流构造 C。在 C周围立即添加屏障将导致并行循环在 C的上方和下方直接分离。因此，上方和下方的操作将被分别放入自己的并行 for 中，而 C将是中间循环中的唯一操作。然后，可以使用以下技术之一对C 进行互换，从而使屏障的父级成为并行for。考虑一个包含屏障的串行 for循环的情况，如图5。这种模式在 GPU代码中很常见，例如，为了实现跨线程的归约[24]。由于屏障必须等待所有线程，每个线程必须执行相同数量的屏障。因此，内部循环的迭代次数对于所有线程都是相同的，允许进行循环互换。虽然 if语句可以被认为是具有零个或一个迭代的循环，但在必要时直接与周围的并行 for 进行互换更高效。而 MLIR 中的 for循环具有固定的迭代次数，而 while循环支持像图6中的动态退出条件一样的动态退出条件。由于正确性要求在每个线程中执行condition()，直接互换将不合法。然而，GPU的同步语义要求所有线程的迭代次数相同。因此，仍然可以使用一个辅助变量来存储条件的结果进行互换。这展示了基于MLIR/Polygeist构建的优势之一。通过保留高级程序结构，我们可以使用更高效的模式来消除屏障。04 并行优化 Polygeist/MLIR 提供的并行性和 GPU程序的高级表示使得各种优化成为可能。这些优化包括一般优化，适用于任何并行程序，以及在 GPU 转 CPU转换的上下文中的特定优化。1240PPoPP ’23，2023年2月25日至3月1日，加拿大蒙特利尔 W.S. Moses，I.R. Ivanov，J. Domke，T. Endo，J. Doerfert和O. Zinenko0在 GPU 转 CPU的上下文中，这些优化适用于任何并行程序以及特定的优化。04.1 屏障消除和移动 由于 GPU样式的屏障必须经过特殊的转换才能支持 CPU架构，消除或简化任何屏障都可能产生显著影响。此外，即使在 GPU 上运行 GPU代码时，消除屏障也非常有用，因为任何同步都会减少并行性。屏障消除/简化的大部分基础设施直接来自于其在第3.1节中定义的内存行为。让 � ↑ �（� ↓ �）表示从屏障 B到并行区域边缘的内存影响的并集。让 � •† � 表示� • �中在第一个屏障而不是区域边缘之前的影响的子集。给定一个屏障B，如果除了读后读（RAR）之外，跨越屏障的同一位置没有其他内存影响，即 � ↑† � ∩ � ↓ � = �，则 B的行为被前一个屏障所包含。对称地，� ↑ � ∩ � ↓† � = �表示该屏障被后续屏障所包含。可移除屏障的一个具体案例是没有任何内存影响的屏障。例如，考虑图7中的代码，它来自于 backprop Rodinia 基准测试[5]。第一个和最后一个__syncthreads指令是不必要的。根据我们上述基于内存的屏障消除算法，可以证明这一点。对于第一个屏障，� ↑（一直到开头）仅包含对node 的写入和对 input 的读取。� ↓†（到第二个__syncthreads）包含对 weights 的写入和对 hidden的读取。如果在给定调用上下文中已知指针不别名，那么这些都不会冲突。因此，可以安全地消除该屏障。同样的内存分析也可以应用于执行屏障移动。只需在预期的位置放置一个虚拟屏障，并检查前面的内存分析是否会推断出当前屏障是不必要的，从而允许屏障移动。04.2 穿越屏障的内存到寄存器提升定义屏障的语义从其内存行为方面的目标之一是使内存优化在包含屏障的代码中能够正确有效地操作。如第3.1节所述，屏障的内存行为是它们上方和下方代码的并集，但除了当前线程的访问之外。这个漏洞很重要，因为它使得内存到寄存器提升（mem2reg）能够操作线程本地内存，如局部变量。这种优化可以将慢速的内存读取替换为快速的寄存器。例如，再次考虑图7中的代码。考虑标记为“Unnecessary Store #1”和“Unnecessary Load #1”的对  weights[ty][tx]的加载和存储，以及两者之间的同步。在那个点上唯一可以加载的值是之前存储的相同值，即包含从 hidden加载的值的寄存器。因此，如果同一个位置的存储在其他任何人读取weights 之前被覆盖，第一个存储也可以安全地消除。在 mem2reg过程中，Polygeist可以推导出这种转发属性，因为在第3.1节中描述的内存属性中的漏洞使其能够推断出屏障操作不会覆盖当前线程的存储。因此，传统的加载和存储转发可以正确地在屏障代码上操作。0omp. 并行 {0omp. wsloop  %i =  1  to  10  {codeA( %i ) } } omp. 并行 {0omp. wsloop  %i =  1  to  10  {codeA( %i ) } }0omp. 并行 {0omp. wsloop  %i = 1  to  10  {codeA( %i ) } omp. barrieromp. wsloop  %i = 1  to  10  {codeA( %i ) } }0图8.OpenMP并行区域融合的示例。通过插入一个屏障来融合相邻的OpenMP并行区域，以便线程只需初始化一次而不是两次。0当同一个位置在其他人读取 weights之前被覆盖时，第一次存储也可以安全地消除一旦加载被移除。在 mem2reg 过程中，Polygeist可以推导出这种转发属性，因为在第3.1节中描述的内存属性中的漏洞使其能够推断出屏障操作不会覆盖当前线程的存储。因此，传统的加载和存储转发可以正确地在屏障代码上操作。04.3 并行循环不变代码运动传统的循环不变代码运动优化旨在将指令 I移动到串行“for”循环之外，减少 I 的执行次数。如果 I可能访问内存或具有其他副作用，并且除了检查 I的操作数本身是循环不变的之外，编译器还必须检查“for”循环内的其他代码是否与 I所执行的内存访问冲突。在现有的编译器中，尽管可以将循环不变代码运动应用于 GPU 内核中的串行 for循环，但不可能将循环不变代码运动应用于将指令提升到内核调用之外。这部分是因为 GPU 内核与调用它们的 CPU代码保持在一个单独的模块中，并且对并行性的理解不足（见图1）。令人费解的是，通过正确的语义，即使不能将等效的串行循环应用循环不变代码运动，我们仍然可以将循环不变代码运动应用于并行循环。我们将依赖于我们程序的语义允许我们任意交错并行“for”循环的迭代，只要我们保持由屏障所需的排序。因此，合法的，虽然不一定快速的运行程序的方式是以锁步方式运行。换句话说，如果一个并行 for循环有10个指令，每个线程可以在任何线程执行指令2之前执行指令1，依此类推。因此，现在可以提升一个指令，只要它的操作数是不变的，并且并行 for 循环中的任何先前指令与 I不冲突。04.4 块并行优化0OpenMP是CPU上并行执行的主要目标。它将并行“for”循环实现为两个结构。}1250通过高级并行构造进行GPU到CPU的转换和优化 PPoPP ’23，2023年2月25日至3月1日，加拿大蒙特利尔0#pragma omp parallel for for(j =0 ; j <10 ; j ++ ) { body(i, j);} }0#pragma omp parallel for(i 0#pragma omp for for  (j=0 ; j <10 ; j ++ ) { body(i,j); } #pragma omp barrier0图9.OpenMP并行区域提升示例。这可以看作是将并行区域融合到每个外部循环迭代对应的“区域”之间的扩展。0首先，将循环提取到一个函数中，该函数被每个线程调用，表示OpenMP的“parallel”结构。然后，在提取的函数中，将迭代空间分布在线程之间，表示OpenMP的“worksharingloop”结构。OpenMP还有一个“barrier”结构，但其语义与GPUbarrier不同。当连续执行多个并行循环时，例如从第3.2节的降低屏障开始，可以通过融合相邻的OpenMP“parallel”结构[11]来减少线程管理的开销，而不融合worksharing循环（参见图8），从而不会撤消降低屏障。这可以扩展到将OpenMP并行区域移动到图9中周围的“for”之外，只初始化线程一次而不是�次。将这些应用于控制流结构可以使通过在块上执行并行循环分裂生成的所有“for”循环都有其OpenMP“parallel”（但没有worksharing循环）融合。由于GPU程序往往以高并行性为目标，不同块提供的并行性可能已经饱和了可用核心的数量。如果没有使用共享内存，块和线程的并行性可以折叠到一个单独的OpenMP并行for中，该for会均匀地划分总迭代空间。但是，如果存在共享内存，我们的工具将生成嵌套的并行区域以表示共享内存分配。在这种情况下，嵌套的OpenMP并行区域的额外开销可能超过潜在的额外并行性。此外，并行化内部循环可能会导致不良的内存效果，如伪共享，进一步影响性能[63，65]。因此，我们还支持将任何嵌套的OpenMP并行区域序列化的优化。执行这样的序列化可能利用内存局部性来提高性能。05 MocCUDA：集成到PyTorch中的一个目标是支持在仅支持CPU的超级计算机（如Fugaku）上执行最初的GPU代码[49]。我们专注于PyTorch[44]，尚未在A64FX架构上进行移植，因此使用了简单的回退CPU内核。我们观察到具有高带宽内存的CPU可能受益于GPU风格的优化，因此我们实现了MocCUDA，这是一个用于PyTorch的模拟GPU后端，将对CUDA运行时和库的调用重定向到我们的0图10.PolygeistInnerPar与MCUDA表现相似；PolygeistInnerSer优于MCUDA。PolygeistInnerSer与MCUDA类似地禁用了内部循环并行化，而PolygeistInnerPar保持了块和线程的并行。左图：平均运行时间作为线程数量的函数（在矩阵大小上进行平均）。右图：平均运行时间作为矩阵大小的函数（在线程数量上进行平均）。0实现或者A64FX特定的数学库[20]。我们收集库调用的统计信息，并可以选择使用Polygeist转换的CPU版本替换它们。06 评估0我们在两个著名的GPU基准套件上展示了我们的方法的优点和适用性：GPURodinia基准套件的一个子集[5]和Resnet-50神经网络的PyTorch实现。我们选择这些基准是为了1）对我们的GPU和CPU编译进行性能粗略比较（Rodinia），该基准套件具有手工编码的CPU版本；2）在没有任何GPU的SupercomputerFugaku上，展示我们的系统成功整合到一个有用和真实的应用程序（PyTorchResnet-50）中。此外，我们还将我们的方法的性能与现有的MCUDA[58]工具在CUDA矩阵乘法上进行了比较。对于Rodinia，我们将我们转换后的CUDA到CPU代码与基准的OpenMP版本进行了比较（如果存在），以及在GPU上运行的代码。对于PyTorchResnet-50，我们与“native”和oneDNN后端进行了比较。我们使用LLVM 15（git 00a1258）编译了Polygeist2。对于PyTorch Resnet-50，我们使用NVidia的CUDA 11.6SDK for Arm 3，LLVM 13和Fujitsu的SSL2v1.2.34库编译了Pytorchv1.4.0。对于基线PyTorch测量，我们使用了Fujitsu预安装的PyTorch（v1.5.0）。我们在运行Ubuntu 20.04的AWSc6i.metal实例上评估了Rodinia和矩阵乘法测试（双插槽Intel Xeon Platinum 8375C CPU，每个插槽2.9GHz，每个插槽32个核心和256 GBRAM）。测量是在第一个插槽上执行的，且关闭了超线程和Turbo Boost。每个数字都是至少5次重复的中位数。02 MocCUDA和Polygeist可以在 https://gitlab.com/domke/ MocCUDA 和https://github.com/llvm/Polygeist获得。3尽管我们将在没有GP