没有合适的资源?快使用搜索试试~ 我知道了~
119→→通过高级并行结构实现WilliamS.Moseswmoses@mit.eduMIT CSAIL美国远藤敏雄endo@is.titech.ac.jp东京理工大学日本Ivan R.Ivanovivanov@m.titech.ac.jpTokyo Tech日本JohannesDoerfertjdoerfert@llnl.gov美国Jens Domkejens.riken.jpRIKEN日本zinenko@google.comGoogle法国摘要虽然并行性仍然是性能的主要来源,但架构实现和编程模型会随着每一代新硬件而变化,这通常会导致成本高昂的应用程序重新设计。大多数用于性能可移植性的工具需要手动和昂贵的应用程序移植到另一个编程模型。我们提出了一种替代方法,自动翻译程序写在一个编程模型(CUDA),到另一个(CPU线程)的基础上Polygeist/MLIR。我们的方法包括一个表示的并行结构,允许传统的编译器转换,以应用程序的透明,无需修改,并使并行特定的优化。 我们评估我们的框架,通过编译和优化的CUDA Rodinia基准套件的多核CPU,并实现了58%的几何加速比手写的OpenMP代码。此外,我们还展示了PyTorch的CUDA内核如何在仅CPU的超级计算机Fugaku上高效运行和扩展,而无需用户干预。我们的PyTorch兼容性层使用编译的CUDA PyTorch内核,性能比PyTorch CPU原生后端高2.7倍。CCS概念:·软件及其工程编译器;·计算理论并行计算模型。关键词:Polygeist,MLIR,CUDA,屏障同步ACM参考格式:William S.伊万?摩西伊凡诺夫,延斯·多姆克,远藤敏夫,乔-汉尼斯·多费特,和亚历山大·齐年科. 2023. 高性能本作品采用知识共享署名国际4.0许可协议进行许可PPoPP©2023版权归所有者/作者所有。ACM ISBN979-8-4007-0015-6/23/02。https://doi.org/10.1145/3572848.3577475通过高级并行结构实现GPU到CPU的转换和优化。 第28届ACMSIGPLAN并行编程原理与实践年度研讨会(PPoPP '23),2023年2月25日至3月1日,加拿大蒙特利尔。ACM,纽约,纽约,美国,16页。https://doi.org/10.1145/3572848.35774751引言尽管x86 CPU和NVidia GPU仍然是计算的主要平台,但定制和新兴架构在计算领域发挥着重要作用ARM CPU的定 制 版 本 A64FX 甚 至 用 于 顶 级 超 级 计 算 机 Fugaku[49],其高带宽内存有望与GPU竞争。然而,这些架构经常被面向效率的框架和库所忽视。例如,PyTorch [44]针对英特尔的oneDNN[28]后端,由于架构差异,预期在ARM上表现不佳,甚至富士通的定制oneDNN[20]在某些内核上也没有产生竞争力。 这种情况要求性能可移植性。许多非库的性能移植已经被提出并且包括语言扩展(例如,OpenCL [14]、OpenACC [26])、并行编程框架(例如,Kokkos [3])、领域特定语言(例如, Spi-ra l [17],Halide [47]或TensorComprehensions [64])。由于语言或底层编程模型的差异,所有这些方法仍然需要移植遗留应用程序,有时甚至需要完全重写。我们探索了一种基于完全自动化的编译器的替代方法,该编译器采用一种编程模型(CUDA)中的代码,并生成针对另一种编程模型(CPU线程)的二进制代码 虽然GPU到CPU的转换在过去已经被探索过[9,23,58],但它很少能够产生有效的代码。事实上,CPU的优化,甚至是通用编译器的转换,如常见的子表达式消除或循环不变的代码运动,都受到编译器内部缺乏并行结构的可分析表示的阻碍[39]。由于主流编译器中的并行表示PPoPPW.S. Moses,I.R.Ivanov,J.Domke,T.Endo,J.Doerfert和O.济年科120()()()[10,12,32,50,55],现有的转换是有限的,往往适用于简单的CPU代码。我们提出了一个编译器模型,最常见的GPU结构:多级并行,级别范围内的同步,和本地内存水平。 与源代码和AST级方法(在优化管道之前操作)以及现有的编译器方法(将同步建模为黑盒优化障碍)相比,我们从内存语义建模同步。 这允许基于同步的代码与现有的优化进行互操作,并启用新颖的特定于并行的优化。我 们 的 模 型 使 用 MLIR [34] 和 LLVM [33] 实 现 , 并 利 用MLIR我们扩展了Polygeist [40] C/C++前端以支持CUDA并产生MLIR,该MLIR保留了高级park结构。 我们的原型编译器能够将PyTorch CUDA内核以及其他计算密集型基准编译到LLVM支持的任何CPU架构。除了转换占执行模型中的差异,我们还通过OpenMP利用CPU上的并行性。最后,我们的MocCUDA PyTorch集成允许我们编译和执行CUDA内核2背景像Clang和GCC这样的主流编译器缺乏一个统一的高级并行表示。在CUDA、OpenMP或SYCL等框架中编译并行构造,会强制并行区域的主体存在于由并行运行时调用的单独(闭包)函数中。然后,线程索引或同步等概念通常通过不透明的内部调用单独表示。由于编译器在历史上缺乏有关并行性和所涉及的运行时的影响的信息,任何并行构造也无意中成为优化的障碍。虽然近年来已经尝试[10,12,32,39,50,55,61]来改进CPU并行构造的表示,但加速器编程带来了额外的挑战。独特的编程模型和复杂的体系结构使得主流编译器中GPU并行性的高级表示尚未得到充分开发。__device__ floatsum(float*data,intn){... {\fnSimHei\bord1\shad1\pos(200,288)}__全球__publicintfindDuplicate(int[] nums){inttid=blockIdx.x+blockDim.x*threadIdx.x;//优化:每个块计算一次求和。在没有GPU的情况下,同时替换不支持的调用。//public int val;// if(threadIdx.x == 0)val = sum(in,n);我们在Rodinia CUDA基准测试[5]和PyTorch CUDA内核上评估我们的编译器。当目标是商品CPU时,我们的OpenMP加速CUDA代码产生了比较,//同步线程;intn= sum(n);if(n)(d_out,d_in,n);}图1.一个示例CUDA程序规范化,它规范化一个向量和CPU函数启动启动内核。每个GPU线程调用sum,导致2个工作。使用共享内存(已注释)将工作量减少到(2/ ),但需要额外的资源成本。在内核之前计算sum可以将工作量减少到()。2.1GPU编译考虑图中的CUDA程序1,对向量进行归一化。 当使用Clang编译时,GPU程序是一个单独的编译单元。 这可以防止GPU内核和CPU调用代码之间的任何优化。与图1,在传统的编译器中程序的总工作量是2,因为每个线程都要调用sum。但是,如果在内核调用之前只执行一次对 sum 的 调 用 , 例 如 , 通 过 执 行 循 环 不 变 码 运 动(LICM),工作将减少至()。 这种优化的一个不太有效的变体可能把工作减少到最少2通过使用共享资源,1我们使用术语transpilation来指在一个编程模型中获取程序并为另一个编程模型发送代码,类似于源到源CUDA到C的translators,尽管现在在IR上。 此过程还交叉编译代码。其指的是发出非本机指令。或Y。MLIR为GPU提供了一个(电子)动态演示支持主机/设备代码运动的程序[21],但并行代码运动尚未实现。在GPU到CPU的代码运动中,并行循环之外的LICM始终是合法的,因为主机上的任何先前设备内存也是可用····通过高级并行结构实现GPU到CPU的转换和优化PPoPP121}//内核启动在调用//函数,实现跨// GPU/CPU边界。func@launch(%h_out:memref<?xf32>,%h_in:memref<?xf32>,%n:i64){//在网格中的所有块之间并行。__global__f(){ codeA();barrier();codeB();}__global__f(){// 0<=t.x//在一个块中的所有线程之间并行。并联for(%tx,%ty,%tz)=(0,0,0)to(blk. x,blk. y,blk。(z){//控制流直接保存。如果%tx==0{%sum =函数call@sum(%d_in,%n)memref。存储%sum,%shared_val[]:memreff32>}//通过显式操作同步。多灵论者屏障(%tx,%ty,%tz)%tid = %gx+网格。x *%txif%tid%n {%res =.存储%res,%d_out[%tid]:memref<?联系我们图3.左:在两条任意指令之间包含屏障的程序。右:屏障语义可以被细化为内存地址,除了当前线程之外,所有线程中它2.3多灵论Polygeist是一个基于Clang的MLIR的C/C++前端[40]。它能够将广泛的C++程序转换为MLIR方言的混合,这些方言保留了程序的高级结构元素。具体来说,Polygeist将结构化控制流(循环和条件)保留为MLIR SCF方言操作,并通过以下方式简化分析:}保持多维数组构造,无论何时}可能通过依赖MLIR}存储引用(memref)类型。最后,Poylgeist能够图2. Polygeist/MLIR等效于图1中的启动/标准化代码1. 内核调用可以直接在调用它的宿主代码中获得。并行性是显式的,可以跨块和线程执行循环 共享内存被并行地放置在块中,允许来自同一块中的任何线程的访问,但不允许来自不同块的访问。2.2MLIR基础设施MLIR是一种最新的编译器基础设施,旨在实现重用和可扩展性[34]。MLIR不是提供一组预定义的指令和类型,而是对包含可互操作的用户定义操作、属性和类型的方言集合进行操作。操作是IR指令的概括,其可以是任意复杂的,特别地,包含具有更多IR的区域,从而创建嵌套表示。操作定义和使用服从单一静态赋值(SSA)的值[7]。例如,MLIR方言可以建模整个指令集,如NVVM(NVidia GPU的虚拟IR),其他IR,如LLVM IR [33],控制流,如循环,并行编程模型,如OpenMP和OpenACC,机器学习图等。MLIR支持GPU,这要归功于epsilon方言,它定义了高级SIMT编程模型、主机/设备通信和一组平台特定的参数:NVVM(CUDA)、ROCDL(ROCm)和SPIR-V。MLIRGPU编程的方法受益于统一的代码表示。由于MLIR模块可以包含其他模块,因此“主机”转换单元可以将“设备”转换单元嵌入为IR而不是文件引用或二进制blob。 这种方法提供了其他编译器无法提供的主机/设备优化操作,特别是在主机和设备之间移动代码[21]。识别程序中适合多面体优化的部分[16],并使用仿射方言表示它们3方法我 们 扩 展 了 Polygeist 编 译 器 [40] , 以 直 接 从CUDA发 出partialMLIR。 这利用了统一的CPU/GPU表示,以允许优化器了解主机/设备执行,并支持跨内核边界的优化。使用现有的MLIR的第一类并行结构(SCF。并行,仿射。并行)使我们能够针对前CPU和GPU后端。最后,MLIR的可扩展操作集允许我们使用相关属性和自定义优化来定义自定义指令。我们定义GPU内核启动的表示如下(如图所示)。(2):在网格中的所有块上的3D并行for循环任何共享内存的堆栈分配每个块的作用域都是唯一的。一个块中所有线程的3D并行for循环自定义Polygeist屏障操作,提供与CUDA同步等效的语义这个过程使我们能够以保留所需语义的形式 它完全理解的编译器,因此是服从编译器优化。此外,通过表示GPU程序与一般的并行性,分配和同步结构,我们不仅能够优化原始程序,但也重新定位它为不同的架构。3.1障碍语义CUDAsyncthreads函数保证块中的所有线程在执行之前已经完成了所有指令的执行。····PPoPPW.S. Moses,I.R.Ivanov,J.Domke,T.Endo,J.Doerfert和O.济年科122R→DD\I→D R R→AA函数调用,在任何线程执行调用后的任何指令之前。传统上,编译器将这些函数表示为可能触及所有内存的不透明优化屏障,并禁止涉及它们的任何转换在我们的系统中,我们选择通过一个新的polygeist.barrier操作来表示线程级同步与其他方法不同,polygeist.barrier(因此简称为barrier)的目标是只阻止会改变外部可见行为的转换我们可以通过定义barrier具有特定的内存属性(表示为collec),来实现所需的语义,而不是禁止任何代码移动通过barrier并行%i=0到10{%x=加载数据[%i]%y=加载数据[2*%i]%a= fmul%x,%x%b= fmul%y,%y%c= fsub%x,y屏障call@use(%a,%b,%c)...}%x_cache=memref<10xf32>%y_cache=memref<10xf32>并行%i=0到10{%x=加载数据[%i]%y= 加 载 数 据 [2*%i] 存储 %x , %x_cache[%i] 存储%y,%y_cache[%i]}并行%i=0到10{%x=加载%x_缓存[%i]%y=加载%y_cache[%i]%a= fmul%x,%y%b=fsub%y,%zcall@use(%a,%b)...}内存位置(包括未知)和内存效应类型(读,写,分配,释放),如MLIR中的标准。考虑图中的简单程序3(左)。 只有当codeA和codeB访问相同的内存时,才能观察到同步的影响。而且,如果两者只读取相同的内存位置,同步也是不必要的. 我们可以列举其余的情况:(1)codeA写,codeB加载;(2)codeA加载,codeB写;(3)codeA写,codeB写。具有代码A的写入行为的屏障将确保(1)的正确性:代码B中的负载不能提升到屏障上方,因为它看起来读取不同的值。对称地,具有代码B的写入行为的屏障确保(2)的正确性因此,代码A和代码B的写入行为的联合足以防止负载越过屏障的非法移动。但是,这不会阻止写入被移动。例如,在(3)中,代码B可以被复制到屏障之上,并且它将看起来具有相同的最终存储器状态,因为屏障之前的无关写入将永远不会被读取。因此,我们也定义了阅读障碍代码A和代码B的行为。该模型可以扩展到包括并行循环中可能在给定屏障之前或之后执行的所有操作的记忆效应。 在具有显式分支的控制流图上,这需要分别探索前导或后继中的操作。然而,在MLIR的结构化控制流级别上进行操作,使用循环和条件的显式操作,简化了分析。此外,如果在同一块中存在多个障碍,则不必越过它。给定一个充分表达的副作用模型,障碍的语义可以进一步扩展。虽然屏障强制从不同线程对同一位置进行读/写排序,但自然执行顺序在一个线程内就足够了。因此,在地址是线程标识符的单射函数的情况下,屏障不需要捕获操作的内存效应。 我们利用MLIR/Polygiest中的多面体框架实现了仿射形式的访问表达式的细化。 对于每一次访问,我们在一组可能的线程id值和一组访问的数组下标之间定义一个整数关系,图4.围绕屏障的并行循环分割:屏障上方的代码与屏障后面的代码放在一个单独的并行“for”循环中。 这种转换消除了障碍,同时保留了语义。 min-cut算法存储%x和%y,然后在第二个循环中使用它们重新计算%a、%b和%c。**然后,我们为相关操作组合正向和反向关系,以获得访问相同下标=−1:������′的线程索引之间的关系。最后,我们减去恒等关系: ′。������如果非空,则表示不同的线程可以访问相同的地址,并且需要屏障。 给定非仿射访问或非静态控制流,我们保守地假设整个数组维度的访问。在实践中,这在GPU代码上很少需要,其循环通常具有参数/静态边界。 当涉及多个基址时,必须检查别名保证。考虑图中的代码3(右)。由于访问的地址集不重叠,所以允许代码移动越过屏障。相反,如果对A的加载或存储偏移1,则屏障将是必要的,因为在屏障之后加载的数据将由不同的线程存储。3.2势垒降低为了使GPU程序能够在CPU上运行,我们必须有效地模拟GPU程序的同步行为。鉴于第3.1节中的内存语义使我们能够在优化过程中保持屏障的正确性,本节将讨论如何在CPU上实现屏障。CPU体系结构没有线程块的概念,也没有等待线程概念分组的屏障指令。相反,我们使用常规CPU线程和工作共享来跨它们分发线程块循环迭代。从概念上讲,这与GPU执行模型不同,在GPU执行模型中,每个线程执行一次迭代工作共享要求每个线程顺序地执行多个迭代,使得不可能在迭代的中间同步,而只能在循环的末尾同步为了解决这个问题,我们开发了一种新的障碍消除技术,我们的MLIR表示。我们的做法是通过高级并行结构实现GPU到CPU的转换和优化PPoPP123对于%id=0到N {对于%j=5到0{if(%id2^%j)A[%id] +=\[%id+2^%j]屏障}}对于%j=5到0{并行%id=0到N { if(%id2^%j)A[%id]+=A[%id+2^%j]屏障}}__global__voidbpnn_layerforward(.){__shared__ floatnode[HEIGHT];__shared__ floatweights[HEIGHT][WIDTH];if(x== 0)node[ty]= input[index_in]// Unnecessary Barrier#1syncthreads();//不必要的存储#1weights[ty][tx]=hidden[index];return();图5.左:一个共享内存添加,它由一个内核调用组成,其中包含一个内部带有屏障的for循环。右:相同,但在并行/串行循环互换后,屏障直接位于并行循环中。//不必要的加载#1weights[ty][tx]= weights[ty][tx]*node[ty];syncthreads();public intfindDuplicate(inti= 1; i<=nums){if(ty% power(2,i)== 0)weights[ty][tx]+= weights[ty+ pow(2,i-1)][tx];syncthreads();parallel for%i =0 to N{ do {运行(%i)障碍} while()}%helper= alloca memrefscf. 做什么?并行%i=0到N { run(%i)屏障%c= condition()if%i ==0{存储%c,%helper[]}hidden[index]= weights[ty][tx];// Unnecessary Barrier#2syncthreads();if(tx== 0)out[by* hid+ ty]= weights[tx][ty];}图7. 一个来自Rodinia的CUDA内核示例-%c= load%helper []} while(%c)图6. 围绕while循环的并行交换。 由于condition()函数调用必须在每个线程上执行以保持正确性,因此使用了一个helper变量来保存第一个线程上的调用值。循环分裂的扩展(见第7节),结合了两种变换:并行循环分裂和交换。3.2.1并行环路分割。 假设一个barrier有核函数(或者,在我们的表示中,parallel for loop)作为它的直接父函数。 它可以通过将屏障周围的循环分成两个并行的for循环来消除,这两个循环分别在屏障之前和之后运行代码。如果屏障之前的代码创建了在屏障之后使用的SSA值,则必须在第二个并行循环中存储或重新计算这些值。我们使用类似于[41]中的技术来确定需要存储的最小数据量。具体来说,我们创建一个所有SSA值的图然后,我们标记每个不能重新计算的值定义(例如,从重写的内存中加载),将屏障之前的值作为源,将屏障之后使用的值作为接收器。我们通过对这个图执行最小分支切割来3.2.2并行环路交换。并不是所有的屏障操作都有一个并行的for作为它们的直接父操作,有些可能嵌套在其他控制流操作中。我们创建了一个模型,指定哪些指令可以并行运行除了屏障之外,我们的表示不需要任何特定的顺序或并发程序。因此,增加额外的障碍是合法的(尽管可能会减少并行性) 我们可以使用这个属性来实现控制流的屏障降低。包含不必要的同步和不必要的共享内存使用的适当测试。考虑一个包含barrier并嵌套在并行for中的控制流构造C。直接在C周围添加障碍将导致C正上方和正下方的平行回路分裂。因此,C上面和下面的操作将被分离成它们自己的并行操作,并且C将是中间循环中的唯一操作然后,我们可以应用以下技术之一来将C与并行for交换,从而使屏障的父对象成为并行for。考虑一个包含屏障的串行for循环的情况图。5. 这种模式在GPU代码中很常见,例如,以实现跨线程的缩减[24]。 由于屏障必须等待所有线程,因此每个线程必须执行相同数量的屏障。因此,内部循环的迭代次数对于所有线程都是相同的,从而允许循环互换。虽然if语句可以被认为是一个零次或一次迭代的循环在必要时直接将其与周围的并行语句for交换会更有效。而对于MLIR中的循环有一个固定的行程计数,而循环支持动态退出条件,如图。第六章由于正确性要求在每个线程中执行condition(),因此直接交换是不合法的。然而,GPU同步语义要求行程计数在所有线程中相同因此,仍然可以使用帮助变量来存储条件的结果来执行交换。这说明了建立MLIR/Polygeist的优势之一。 通过保留高级程序结构,我们可以使用更有效的模式来消除障碍。4并行优化Polygeist/MLIR提供的并行性和GPU程序的高级表示实现了各种优化。其中包括常规优化,}PPoPPW.S. Moses,I.R.Ivanov,J.Domke,T.Endo,J.Doerfert和O.济年科124∩∅������∩∅������适用于任何并行程序以及GPU到CPU转换上下文中4.1障碍消除运动由于GPU风格的屏障必须经过特殊改造以支持CPU架构,因此可以消除或简化任何噢。平行{噢。wsloop%i =1 to10{ codeA(%i)}}噢。平行{噢。wsloop%i =1 to10{ codeA(%i)}噢。平行{噢。wsloop%i =1 to10{ codeA(%i)}噢。屏障噢。wsloop%i =1 to10{ codeA(%i)}}障碍会产生巨大的影响。此外,即使在GPU上运行GPU代码时,屏障消除也非常有用,因为任何同步都会降低并行性。障碍消除/简化的大部分基础设施直接来自第3.1节中定义的内存行为。 设������↑(������↓)为从障碍物B到平行区域边缘的所有记忆效应的并集。让������������在第一个障碍之前,而不是区域边缘。给定障碍B,如果没有存储器效应跨越除读后读(RAR)之外的屏障到达相同位置,即,↑ <$ ↓=,B具有被前一屏障包含的其行为。对称������↑可移除屏障的一个具体情况是完全没有记忆效应例 如 , 考 虑 图 1 中 的 代 码 7 , 它 来 自 于backpropRodinia 基 准 [5] 。 第 一 个 和 最 后 一 个syncthreads指令是不必要的。这可以从我们上面的基于存储器的屏障消除算法证明如下。对于第一个障碍,↑(一直到开始)只包含一个写节点和一个读输入。���↓<$(转到第二个同步线程)包含对权重的写入和对隐藏的读取。如果给定调用上下文,已知指针没有别名,则这些都不因此,消除障碍是安全的同样的记忆分析也可以应用于执行障碍物运动。人们只需要在预期的位置放置一个虚构的障碍物,并检查先前的记忆分析是否会推断出当前的障碍物是不必要的,从而允许障碍物运动。4.2跨越障碍从其内存行为定义屏障语义的目标之一是使内存优化能够在包含屏障的代码中正确有效地操作。如第3.1节所述,屏障具有其上方和下方代码的内存,但从当前线程访问的显着例外 这个漏洞很重要,因为它允许内存到寄存器提升(mem2reg)对线程本地内存(如局部变量)进行操作。这种优化可以用快速寄存器代替慢速内存例如,再次考虑图1中的代码。第七章 考虑加载和存储到标记为“不必要的存储#1”和“不必要的加载#1”的权重[ty][tx],以及两者之间的同步。在该点可以加载的唯一值是先前存储的相同值,一个包含从隐藏加载的值的寄存器。作为}图8. OpenMP并行区域融合示例。通过插入一个屏障,允许 线 程 初 始 化 一 次 而 不 是 两 次 , 从 而 将 两 个 相 邻 的OpenMP并行区域连接起来在任何其他人可以从权重读取之前,该相同的位置被重写,一旦加载被移除,第一存储也可以被安全地消除。在mem2reg期间,Polygeist可以导出此转发属性,因为第3.1节中描述的内存属性中的漏洞允许它推断屏障操作不会覆盖当前线程的存储。 因此,传统的加载和存储转发正确地对屏障码进行操作。4.3并行循环不变码运动传统的循环不变代码运动优化旨在将指令I移动到串行“for”循环之外,从而减少I被执行的次数如果I可能访问内存,或者有其他副作用,除了检查I的操作数本身是循环不变的之外,编译器还必须检查“for”循环中没有其他代码与I执行的内存访问冲突。在目前的编译器上,虽然有可能将循环不变代码运动应用于GPU内核内的串行for循环,但不可能将循环不变代码运动应用于内核调用外部的提升指令。 这在一定程度上是由于GPU内核被保存在一个独立的模块中,与调用它们的CPU代码分开,以及缺乏对并行性的理解(见图1)。1)。与直觉相反,使用正确的语义,我们可以将循环不变代码运动应用于并行for循环,即使我们无法将其应用于等效的串行循环。我们将依赖于这样一个事实,即我们的程序的语义允许我们任意交错并行“for”循环的迭代,只要我们保持障碍所需的顺序。因此,它是合法的,虽然不一定快,在锁步运行程序换句话说,如果一个并行for循环有10条指令,每个线程可以在任何线程执行指令2之前执行指令1,依此类推。因此,现在提升一条指令是合法的,只要它的操作数是不变的,并且并行for循环中没有前面的指令与I冲突。4.4块优化OpenMP是我们在CPU上并行执行的主要目标它将并行的“for”循环实现为两个构造。通过高级并行结构实现GPU到CPU的转换和优化PPoPP125for(i=0; i N; i++){#pragma omp parallelforfor(j=0; j<10; j++){return(i);}}#pragma omp parallelfor(i=0; i N; i++){#pragma ompfor(j=0; j<10; j++){ body(i,j);}#pragma omp barrier}图9.OpenMP并行区域提升示例。这可以被看作是跨对应于外循环的每次迭代的“区域”的并行区域融合的扩展首先,循环被概括为一个函数,每个线程调用一次,表示OpenMP的“并行”结构。然后,在概述的函数中,迭代空间分布在线程之间,表示OpenMP的“工作共享循环”构造。OpenMP 也 有 一 个 “barrier” 结 构 , 但 其 语 义 与 GPUbarrier不同。当在一行中执行多个并行循环时,例如,在第3.2节降低了 屏 障 之 后 , 线 程 管 理 的 开 销 可 以 通 过 融 合 相 邻 的OpenMP“并行”结构[11]而不融合工作共享循环来减少(见图12)。8),因此不解除屏障降低。 这可以扩展到将OpenMP并行区域���移动到图9中周围的“for”之外,初始化线程一次而不是多次。将这些应用于控制流构造,使得通过在块上执行并行循环裂变而生成的所有由于GPU程序往往在编写时考虑到高并行性,因此由不同块提供的并行性可能已经使可用核的数量饱和。如果不使用共享内存,则块和线程并行性可以折叠为单个OpenMP并行,这将在单个并行区域中均匀划分总迭代空间。但是,如果有共享内存,我们的工具将生成嵌套的并行区域来表示共享内存分配。 在这种情况下,来自嵌套的OpenMP并行区域的额外开销可能超过潜在的增加的并行性。此外,并行化内部循环可能会导致不利的内存影响,例如错误共享,进一步降低性能[63,65]。因此,我们还支持对任何嵌套的OpenMP并行区域进行序列化的优化。执行这样的串行化可以利用存储器局部性来提高性能。5MocCUDA:集成到PyTorch我们的目标之一是支持在只有CPU的超级计算机(如Fugaku [49])上执行原始GPU代码。我们关注的是尚未移植到图 10.PolygeistInnerPar 的 性 能 与 MCUDA 相 似 ;PolygeistInnerSer的性能优于MCUDA。PolygeistInnerSer类似于MCUDA禁用内部循环并行化,而PolygeistInnerPar保持块和线程并行。左:作为线程数函数的平均运行时间(对矩阵大小求平均)。右:平均运行时间作为矩阵大小的函数(对线程计数求平均值)。实现或A64FX特定的数学库[20]。 我们收集库调用的统计数据,并可以选择用Polygeist编译的CPU版本替换它们。6评价我们在两个著名的GPU基准测试套件上展示了我们的方法的优势和适用性:GPU Rodinia基准测试套件的子集[5]和Resnet-50神经网络的PyTorch实现。选择这些基准测试是为了:1)在具有手工编码CPU版本的基准测试套件(Ro-dinia)上提供GPU到CPU编译的粗略性能比较; 2)在没有任何GPU的超级计算机Fugaku上展示我们的系统成功地端到端集成到有用的真实应用程序(PyTorch Resnet-50)中。此外,我们还比较 了我们 的方法 与现 有MCUDA [58]工具在CUDA矩阵乘法上的性能。对于Rodinia,我们将翻译后的CUDA与CPU进行代码对基准测试的OpenMP版本,如果他们存在,以及在GPU上运行对于PyTorch Resnet-50,我们将其与“原生”和oneDNN后端进行比较。Polygeist2使用LLVM 15(git00a1258)编译。对于PyTorch Resnet-50,我们使用NVidia的CUDA11.6SDKforArm3,LLVM13和Fujitsu的SSL2v1.2.34库编译Pytorch v1.4.0。对于基线PyTorch测量,我们使用富士通我们在运行Ubuntu 20.04的AWS c6i.metal实例(2.9 GHz的双插槽Intel Xeon Platinum8375C CPU,每个CPU有32个内核和256 GB RAM)上评估Rodinia和矩阵乘法测试。 在第一个插槽上进行测量,禁用超线程和涡轮增压。每个数字是至少5次重复的中位数A64FX架构,因此使用简单的回退CPU玉米粒 观察到具有高带宽架构的CPU可能会从GPU风格的优化中受益,我们实现了MocCUDA,这是PyTorch的模拟GPU后端,它将对CUDA运行时和库的调用重定向到我们的2MocCUDA和Polygeist可在https://gitlab.com/domke/MocCUDA和https://github.com/llvm/Polygeist。3即使我们将在无GPU的系统上运行PyTorch,我们也必须编译在支持CUDA的系统上运行PyTorch,以确保发出正确的代码我们还阻止了三个Pytorch函数的内联PPoPPW.S. Moses,I.R.Ivanov,J.Domke,T.Endo,J.Doerfert和O.济年科126×××图11.左:相对加速(越高越好)应用第4节中提出的并行优化,在我们的流程中没有优化。右图:与运行32个线程的本机OpenMP代码(可用时)相比,CUDA-to-OpenMP转译的速度提升。星号表示基准中的障碍6.1与MCUDA首 先 , 我 们 与 MCUDA 中 的 先 前 工 作 进 行 比 较 [58] 。MCUDA是一个AST级工具,它产生新的CPU C/C++作为输出,并使用循环分裂来处理同步。作为一个源代码到源代码的工具,MCUDA只处理一小部分输入语言,使其无法在Rodinia程序上运行。相反,我们在图中比较了一系列线程(1-24)和矩阵大小(128 128 - 2048 2048)的矩阵乘法内核的运行时间。10. Polygeist使用所有优化,不包括内部循环的序列化(PolygeistInnerPar),生成的代码在1以内。平均为MCUDA的 3% PolygeistInnerPar有一个1。1个线程速度减慢5%,3. 在32个线程上加速2% 此行为是由处理嵌套并行构造时的OpenMP开销引起的。实际上,MCUDA只并行化最外层的循环。当Polygeist也序列化内部循环( PolygeistInnerSer ) 时 , 它 达 到 了 14 。 9% 的加 速 比MCUDA,与4。5%的加速1线程和21. 在32个线程上加速7%6.2用例1:Rodinia基准我们对Polygeist目前支持的14个基准测试进行了基准测试,并且有一个不平凡的运行时。4.通过比较nvcc编译后在GPU上执行的程序输出和我们的流程编译后在CPU上执行的程序输出,验证了程序的正确性。 我们还采用了基于CPU的并行和未定义行为分析工具,通过我们的工具,使我们能够成功地诊断和修复原始CUDA代码中的一个种族错误和几个未定义的内存错误。 我们在内核和/或包含内核的代码的计算部分中插入了时序测量,在某些情况下,4hybridsort、kmeans、leukocyte、mummergpu huffman和heartwall使用Polygeist中不支持的C++或CUDA功能(虚函数和纹理内存)。lavaMD和dwt2d基准测试使用格式错误的C++,由于从未初始化的内存中读取,因此具有未定义的行为nn
下载后可阅读完整内容,剩余1页未读,立即下载
![pdf](https://img-home.csdnimg.cn/images/20210720083512.png)
![pdf](https://img-home.csdnimg.cn/images/20210720083512.png)
![pdf](https://img-home.csdnimg.cn/images/20210720083512.png)
![-](https://csdnimg.cn/download_wenku/file_type_lunwen.png)
![-](https://csdnimg.cn/download_wenku/file_type_lunwen.png)
![-](https://csdnimg.cn/download_wenku/file_type_column_c1.png)
![](https://csdnimg.cn/download_wenku/file_type_ask_c1.png)
![](https://csdnimg.cn/download_wenku/file_type_ask_c1.png)
![](https://csdnimg.cn/download_wenku/file_type_ask_c1.png)
![](https://csdnimg.cn/download_wenku/file_type_ask_c1.png)
![](https://csdnimg.cn/download_wenku/file_type_ask_c1.png)
![](https://csdnimg.cn/download_wenku/file_type_ask_c1.png)
![](https://csdnimg.cn/download_wenku/file_type_ask_c1.png)
![](https://csdnimg.cn/download_wenku/file_type_ask_c1.png)
![](https://csdnimg.cn/download_wenku/file_type_ask_c1.png)
![](https://csdnimg.cn/download_wenku/file_type_ask_c1.png)
![](https://csdnimg.cn/download_wenku/file_type_ask_c1.png)
![](https://csdnimg.cn/download_wenku/file_type_ask_c1.png)
![](https://profile-avatar.csdnimg.cn/default.jpg!1)
cpongm
- 粉丝: 4
- 资源: 2万+
上传资源 快速赚钱
我的内容管理 收起
我的资源 快来上传第一个资源
我的收益
登录查看自己的收益我的积分 登录查看自己的积分
我的C币 登录后查看C币余额
我的收藏
我的下载
下载帮助
![](https://csdnimg.cn/release/wenkucmsfe/public/img/voice.245cc511.png)
会员权益专享
最新资源
- 京瓷TASKalfa系列维修手册:安全与操作指南
- 小波变换在视频压缩中的应用
- Microsoft OfficeXP详解:WordXP、ExcelXP和PowerPointXP
- 雀巢在线媒介投放策划:门户网站与广告效果分析
- 用友NC-V56供应链功能升级详解(84页)
- 计算机病毒与防御策略探索
- 企业网NAT技术实践:2022年部署互联网出口策略
- 软件测试面试必备:概念、原则与常见问题解析
- 2022年Windows IIS服务器内外网配置详解与Serv-U FTP服务器安装
- 中国联通:企业级ICT转型与创新实践
- C#图形图像编程深入解析:GDI+与多媒体应用
- Xilinx AXI Interconnect v2.1用户指南
- DIY编程电缆全攻略:接口类型与自制指南
- 电脑维护与硬盘数据恢复指南
- 计算机网络技术专业剖析:人才培养与改革
- 量化多因子指数增强策略:微观视角的实证分析
资源上传下载、课程学习等过程中有任何疑问或建议,欢迎提出宝贵意见哦~我们会及时处理!
点击此处反馈
![](https://img-home.csdnimg.cn/images/20220527035711.png)
![](https://img-home.csdnimg.cn/images/20220527035711.png)
![](https://img-home.csdnimg.cn/images/20220527035111.png)
安全验证
文档复制为VIP权益,开通VIP直接复制
![](https://csdnimg.cn/release/wenkucmsfe/public/img/green-success.6a4acb44.png)