文摘

在标准OpenCL编程中,主持人应该控制他们的计算设备。因为计算设备专用内核计算,只有主机可以执行几种数据传输节点之间的通信和文件访问等。这些数据传输需要一个主机同时扮演两个或两个以上的角色由于主机和设备之间的协作的必要性。这种数据传输的编码可能系统,导致较低的可移植性。本文提出一个OpenCL扩展,包含这些数据传输OpenCL事件管理机制。与当前OpenCL标准,主线程的主机上运行不了序列化的操作。因此,应用程序可以很容易地使用主机的重叠并行的活动的机会和计算设备。此外,数据传输的实现细节隐藏在扩展和应用程序员可以使用优化数据传输没有任何复杂的编程技术。评价结果表明,该扩展可以使用优化数据传输的实现,从而提高持续数据传输性能18%在一个真实的应用程序中访问一个大数据文件。

1。介绍

今天,许多高性能计算(HPC)系统配有图形处理单元(gpu)担任方式来表述数据并行处理加速器除了传统的通用处理器(cpu)。对于这样一个异构高性能计算系统,应用程序编程人员需要管理系统异构性,同时利用并行性参与他们的应用程序。对于论文的其余部分,我们将遵循OpenCL的术语和参考的cpu主机和方式来表述数据并行处理加速器计算设备。

编程这样的异构系统的一个难题就是,一个程序员必须负责任命合适的处理器的任务。在当前OpenCL标准,只有主人可以执行的一些任务,因为计算设备致力于内核计算。例如,只有主机可以访问文件和与其他节点进行通信。将内核的计算结果写入一个文件,结果必须首先从设备记忆转移到主机内存内核执行后,然后主机将结果写入文件。

从程序员的角度,加速器编程模型如CUDA (1]和OpenCL [2)用于设备内存之间的数据传输和主机内存,MPI (3)用于节间数据通信和文件功能的编程语言,如流和fscanf在C编程,用于文件I / O。因此,这三个类别的数据传输与不同的编程模型描述。一些数据传输通过不同的编程模型可以依赖;一定的数据传输可以做只有在其之前的数据传输。为了执行这样的依赖,一个流行的方法是阻止主机线程,直到前面的数据传输完成。这种阻塞通常抑制重叠并行的活动主机和设备和公开的数据传输延迟总执行时间。你可以创建一个专用的主机线程同步相关的数据传输。然而,这样的多线程将进一步增加编程的复杂性。因此,应用程序的性能很大程度上取决于开发人员的编程技巧和工艺。

另一个困难是,没有标准的方式来编码数据传输甚至对常见的数据传输模式。自应用程序员应该适当地结合这些数据传输充分利用异构高性能计算系统的潜力,代码通常是专门为特定的系统。例如,一个计算装置可能能够直接访问一个文件,和另一个可能不会。在这种情况下,文件访问的代码前设备将完全不同于后者。因此,数据传输可能是系统特定的代码和一些抽象要求实现功能的可移植性以及性能的可移植性。虽然OpenCL设计编程各种计算设备,它提供了接口只有主机内存之间的数据传输和设备内存,而不是其他类型的数据传输。

为了克服上述困难,我们需要一个“桥接”的编程模型,提供了一种标准的方式编码的各种内存空间之间的数据传输和存储异构并行系统以统一的方式。在本文中,我们关注OpenCL的加速器编程模型代码可移植性高,提出一个OpenCL扩展抽象的数据传输,虽然这个想法可以非常外推到其他诸如CUDA GPU编程模型。拟议中的OpenCL扩展名为clDataTransfer提供了一种错觉,直接计算设备传输数据文件或其他节点。本文尤其在节点之间的通信和文件访问的典型数据传输需要协作的主机和设备。扩展提供了一些OpenCL命令和数据传输的功能。节点之间的通信和文件访问命令以同样的方式执行其他OpenCL命令,因此自然OpenCL编程模型扩展,无缝地访问文件数据并提高MPI互操作性。

clDataTransfer扩展提供了一个可移植的,标准化的方式编程节间通信和文件从/到设备内存访问。尽管MPI和文件函数用于在内部执行这些主机的数据传输的帮助下,这些内部应用程序行为是无形的;它可以隐藏背后的使用系统优化实现函数调用。因此,我们也可以认为clDataTransfer扩展提高了性能的可移植性OpenCL应用程序在不同的系统类型,天平,一代又一代。

本文的其余部分组织如下。部分2简要回顾了相关工作。部分3讨论了OpenCL的联合编程的困难,MPI,和标准C库的I / O包,所谓的头。然后,部分4提出clDataTransfer,这是一个OpenCL协作与MPI和它的扩展。部分5讨论了通过一些评估结果clDataTransfer对性能的影响。最后,部分6给出了结论和未来的工作。

2。相关工作

OpenCL的编程模型,CPU是一名主机管理一个或多个计算设备比如gpu。管理主机和设备之间的交互,OpenCL提供各种资源等OpenCL对象实例化情况下,命令队列,内存对象和事件对象。一个独特的处理给每一个对象,用于访问资源。各种资源的上下文是一个容器,类似于一个CPU的过程。使用命令队列与相应的计算设备;主机可以命令其计算设备执行一个任务。内存对象表示一个内存块从主机和设备访问。一个事件对象绑定一个命令的命令队列来表示状态的命令,用于阻止其他命令的执行。因此,它是用来描述命令之间的依赖关系。此外,可以结合多个事件到事件表达几个以前的命令列表。

例如,clEnqueueReadBuffer是一个典型的排队OpenCL函数命令,将数据从设备内存传输到主机内存。函数签名算法1。

OpenCL事件管理命令队列函数接受三个参数:事件的数量在等待列表中(numevts等待列表(的),最初的地址wlist),进入队列的事件对象的地址命令传递(evtret)。时队列的命令可以执行前面的命令与事件相关联的所有对象在等待列表中已经完成。

在MPI的联合编程和OpenCL,程序员不仅需要考虑host-device通信使用OpenCL还使用MPI节点之间的通信。到目前为止,一些研究人员已经提出了几个MPI扩展gpu来缓解MPI和CUDA / OpenCL的联合编程。我们将把这些方法GPU-aware MPI实现。软件提出了cudaMPI [4),为远程gpu之间的通信提供了一个类似的界面。MPI-ACC [5)使用MPI_数据类型参数表明,内存缓冲区传递给一个MPI函数位于设备内存。MVAPICH2-GPU [6]假设统一的虚拟寻址(UVA),它提供了一个内存空间为主机和设备记忆,并检查内存缓冲区传递给一个MPI函数是否在设备内存。然后,MVAPICH2-GPU在内部使用不同的实现取决于设备的内存缓冲区内存或主机内存。斯图尔特等人已经讨论了各种设计方案的MPI扩展支持加速器(7]。Gelado等人提出的通用汽车金融服务公司,它提供了一个内存空间共享CPU和GPU,因此允许MPI函数来访问设备内存数据(8]。这些扩展允许应用程序使用GPU内存缓冲区的终点MPI通讯;扩展的MPI实现支持使用MPI函数节点之间的通信从/到GPU内存缓冲区的内部使用CUDA / OpenCL的数据传输功能。

通过使用GPU-aware MPI扩展,应用程序开发人员不需要显式地描述了host-device数据传输等clEnqueueWriteBuffer和clEnqueueReadBuffer。与clDataTransfer一样,这些扩展不需要复杂的编程技术实现高效的数据传输,因为他们隐藏背后的优化实现MPI函数调用。

GPU-aware MPI扩展,所有节间通信仍由主机管理线程应用程序开发人员可见。例如,如果数据需要通过执行内核其他节点,主机内核线程需要等待执行完成为了序列化内核执行和MPI通讯;托管线程被阻塞直到完成内核的执行。

此外,MPI扩展OpenCL并非易事,因为这个地方等人讨论(5]。保持OpenCL数据传输对MPI应用程序透明,MPI实现必须以某种方式获得有效命令队列。这个等人认为一个MPI的过程大多只使用一个命令队列,因此其处理缓存的MPI实现用于随后的通信,虽然这个假设可能是错误的。即使缓存的命令队列用于随后的通信,可能存在一个更合适的通信命令队列。clDataTransfer允许应用程序程序员指定的命令队列通信。应该强调,GPU-aware MPI扩展和clDataTransfer互利而不是冲突。例如,尽管这个工作实现了管线式数据传输使用标准MPI函数,有可能clDataTransfer使用MPI实施扩展。

斯图尔特和欧文斯提出DCGN [9]。clDataTransfer, DCGN交流提供了一种错觉,gpu的主人没有任何帮助。不像clDataTransfer DCGN提供了节点之间的通信API函数,称为从GPU内核。调用API时由内核运行在GPU,内核设置的区域设备内存由CPU监视线程。然后,CPU线程从设备中读取必要的数据内存,因此从GPU处理的通信请求。因此,DCGN允许内核启动节点之间的通信。然而,要求主机监控设备内存带来不可忽视的运行时开销。另一方面,在clDataTransfer,表示为节点之间的通信请求OpenCL的命令。因此,主机启动命令和clDataTransfer实现可以依靠OpenCL事件管理机制同步命令。

OpenCL内存对象在同一上下文是由多个设备共享。OpenCL内存一致性模型隐式地确保设备的内存对象可见的内容都是一样的只是在同步点。一次设备更新由多个设备共享的内存对象,新的记忆内容含蓄地复制到每个设备在同一上下文的记忆。一些OpenCL实现(10]支持建立一个由多个设备共享的上下文在不同的节点,从而达到远程设备之间的数据共享,同时符合OpenCL规范。然而,在这种方法中,多个设备共享一个上下文可以只有一个内存空间;即使他们不能有不同的内存内容的一些内容不需要由所有节点。结果,不必要的内容可以复制到设备上的记忆,每个节点增加聚合内存使用和复制还节间通信。

采用GPU计算不仅为传统HPC应用程序还为数据密集型应用程序,例如,(11,12),数据规模大,因此存储在文件中。因为只有主机可以访问存储在文件中的数据,GPU计算需要附加的主机之间的数据传输和GPU。尽管如此,gpu加速内核执行和有效降低总执行时间在实际数据密集型的应用程序。重叠的内核执行各种数据传输文件访问和host-device等数据传输是一个关键技术,减少数据传输延迟和显然通用代码模式。然而,据我们所知,并没有标准的方法来开发这个模式的方式在其他应用程序中重用。最近的和未来的HPC系统层次化存储子系统,高速的地方可以使用非易失性存储记忆。在这些情况下,重叠将变得更加重要,因为host-device数据传输开销相对增加到文件访问的开销。

3所示。联合编程的困难

本节讨论一些困难在OpenCL的联合编程和其他库,如MPI,由主机调用线程。清单1显示了一个简单的MPI和OpenCL的联合编程代码。在这段代码中,一个命令执行内核是首先通过调用进入队列clEnqueueNDRangeKernel。另一个命令然后阅读内核执行结果clEnqueueReadBuffer。使用事件对象的第一个命令,evt第二个命令的执行被阻塞,直到第一个命令完成。第二个命令队列clEnqueueReadBuffer可以阻止或非阻塞。如果第三个参数是函数调用阻塞CL_假;否则它挡住。如果非阻塞,我们必须使用一个同步等功能clFinish确保数据已经从设备记忆转移到主机内存提前打电话MPI_Sendrecv。在这个天真的实现,数据交换与其他节点数据传输后必须执行从内存设备到主机内存;这些数据传输必须序列化。同样的,MPI_Sendrecv和clEnqueueWriteBuffer必须序列化。因此,内核执行和数据传输都是序列化,这导致了长时间的沟通接触的总执行时间。此外,主机线程被阻塞时MPI和OpenCL操作序列化。虽然清单1显示了一个示例的联合编程的MPI和OpenCL,同样的困难时出现结合OpenCL和Stdio(或任何其他文件访问编程接口)。

更糟的是,没有标准的方式联合编程。即使是两个远程设备之间简单的点对点通信,我们至少可以考虑以下三个实现。一个是天真的实现如清单所示1。在实现中,主机内存缓冲区应该page-locked(固定)高效的数据传输(尽管OpenCL标准并不提供任何具体方法来分配固定主机内存缓冲区,大多数供应商依赖的使用clEnqueueMapBuffer为程序员提供固定主机内存缓冲区)。这也可以使不同厂商利用固定内存需要不同的实现。另一个实现是设备内存对象映射到主机内存地址使用clEnqueueMapBuffer传输数据,然后调用MPI函数从/到地址。MPI沟通后,clEnqueueUnmapMemObject调用设备映射的内存对象。另一个实现重叠host-device数据传输与节点之间的数据传输。在这个实现中,一个设备的数据内存对象分为一个固定大小的数据块,称为管道的块大小,host-device数据传输的每个块重叠与节间数据传输块以流水线的方式(6]。本文前面所提到的三种实现被称为固定,映射,管线式数据传输。那些实现,最好的变化取决于几个因素,如消息大小,设备类型、设备供应商和设备。也host-device重叠的情况下数据传输文件访问有很多实现选项和参数由于各种各样的文件访问的速度分级存储子系统。因此,应用程序开发人员可能需要实现多个版本优化数据传输性能跨不同系统的应用程序可移植性。

另一个常见的方法隐藏通信开销是重叠的数据传输和计算通过双缓冲技术11,13]。为此,计算通常分为两个阶段。在执行计算,第一阶段第一阶段数据传输准备执行第二阶段计算。如果计算和数据传输是一个循环内,第二阶段第一阶段计算的数据传输在第二阶段执行下一次迭代计算当前的迭代。

在OpenCL编程中,这个重叠优化可以实现使用两个顺序执行命令队列。清单2显示了一个简化版的Himeno基准代码中描述(13),这是最初写在CUDA和MPI。在代码中,jacobi_kernel_*函数(9),(18)、(28)、(35)调用内核使用命令队列中cmd1更新内存对象由第二个参数指定。的代码假定一维域分解,每个分解域进一步减半分为上下两部分,A和b图1说明了域分解假设的代码。顶部和底部平面的平面B是光环每次迭代更新地区与邻近的节点通过交换数据。因此,如果MPI排名的过程是一个偶数,在计算过程中,流程更新光环地区包括在B。然后,它计算B在交换数据更新的光环a .另一方面,如果MPI排名过程是一个奇数,流程将首先计算B在更新a的光环,它计算在交换数据更新的光环B .因此,时间不是暴露的通信总执行时间,如图2(一个)除非沟通时间超过了计算时间。

MPI进程数量的增加,计算时间变得更短,因为域由每个GPU处理变得越来越小。然而,第二阶段的沟通不能即使完成第一阶段计算开始之前,因此准备第二阶段的数据通信,如图2 (b)。这是因为主机线程通常是阻塞和绑在第一阶段沟通为了序列化MPI和OpenCL操作。

由于代码清单2很简单,有一些解决方案技术来解决这个问题。然而,在流水线等情况更先进的优化技术应用于数据传输,主机线程停滞更频繁地及时同步MPI和OpenCL操作应用程序的多个并行的活动。一般来说,至少有三个并行的活动在一个应用程序:主机计算、设备计算和非阻塞MPI通讯。如果有相关操作的MPI和OpenCL,主机线程通常是阻塞序列化操作,抑制重叠的并行活动。同时,主机经常使用线程阻塞甚至在一个串行应用程序如果主人线程需要从文件加载数据,送他们到设备内存,和检索计算结果从设备内存。多线程编程或复杂的异步I / O api需要妥善管理这些并行的活动。通过这种方式,应用程序代码变得更加复杂和系统,导致低代码的可读性、可维护性和可移植性。这激励我们设计一个过渡性的编程模型,可以显式地描述了MPI之间的依赖关系,OpenCL,为了启动数据传输和文件访问操作没有任何帮助的线程。

4所示。一个OpenCL扩展与MPI和Stdio合作

本文提出了clDataTransfer,一个OpenCL扩展促进和规范MPI的联合编程,Stdio, OpenCL。这个扩展的关键思想是为节点之间的数据传输,使用OpenCL命令文件访问和主机之间的数据传输和本地设备。

clDataTransfer总结了如下的主要优势。(1)性能的可移植性:节间的实现细节隐藏在数据传输和文件访问扩展命令通过一个简单的编程接口,可以使用类似于OpenCL标准接口。(2)事件管理:大量线程不负责序列化节间通信、文件操作、host-device通信。相反,一个事件对象是用来阻止随后的命令,直到前面的命令完成时结束。(3)合作为延迟隐藏:clDataTransfer可以与MPI和Stdio为了隐藏数据传输延迟以流水线的方式。

通过封装文件访问到OpenCL命令,clDataTransfer扩展提供了两个文件访问命令:clEnqueueReadBufferToStdioFile和clEnqueueWriteBufferFromStdioFile。 clEnqueueReadBufferToStdioFile从设备内存缓冲区读取数据并将数据写入一个文件,和clEnqueueWriteBufferFromStdioFile从文件中读取数据并将数据写入一个设备内存缓冲区。函数签名算法2。

同样,clDataTransfer扩展提供了clEnqueueSendBuffer和clEnqueueRecvBuffer排队命令的传输数据和设备内存缓冲区,分别。这些clDataTransfer函数直接对应MPI_发送和MPI_Recv(3],因此采取相同的等级参数,标签,和沟通这两个MPI函数。例如,函数的签名clEnqueueRecvBuffer在算法3。

当一个MPI流程调用这些函数发送一个命令一个设备,设备成为一个沟通者的设备一个MPI通讯和工作好像不是主机通信线程。数据发送到MPI排名由传播者接收装置,接收的数据是存储在内存空间的沟通设备,也就是说,缓冲区。MPI的发送方是给函数,和发送方可以是主机线程或沟通者设备与MPI等级有关。

的情况下发送方和接收方提交节点之间的通信命令他们的设备,这些设备相互通信。清单3显示了一个简单的远程设备之间的通信的例子。在这段代码中,沟通者设备级别0发送数据的内存缓冲区对象的沟通设备1级没有显式地调用MPI函数。因此,设备出现与远程通信设备没有帮助他们所在国的线程。节点之间的通信通过结合MPI的实现细节和OpenCL背后隐藏着OpenCL命令执行。因此,应用程序可以使用优化的实现高效的数据传输不使用复杂的编程技术。如果一个MPI过程需要使用多个沟通设备,一个独特的标签是给每个MPI通讯指定哪个沟通者设备处理它。

4.1。事件管理

clDataTransfer扩展允许程序员使用事件对象为了表达之间的依赖节点之间的通信命令,存储文件访问命令和其他OpenCL的命令。如果提供的数据传输命令clDataTransfer需要前一个命令的结果,程序员可以前面的命令的事件对象并使用它来阻止数据传输命令的执行。这确保数据传输前命令完成后执行。通过这种方式,数据传输命令的clDataTransfer纳入OpenCL以自然的方式执行模型。因此,函数调用MPI和Stdio封装在OpenCL命令的依赖与其他OpenCL命令准确地执行由命令队列。与传统的联合编程MPI, Stdio, OpenCL,主机线程不需要等待前面的命令完成。排队的非阻塞函数调用的命令后,主机线程立即可用其他计算和数据传输;应用程序程序员可以认为如果一个设备是能够独立于主机工作线程。OpenCL运行时将在适当的时候释放clDataTransfer命令及时执行MPI函数如图2 (c),尽管这两个通信可能是也可能不是并发执行。

使用clDataTransfer扩展,代码在清单2可以简单地写成的代码在清单吗4。这是一个示例,演示了简化OpenCL的常见模式在联合编程和其他编程模型。在这种特殊情况下,clDataTransfer扩展可以减少一半的代码行数来描述相同的计算OpenCL和MPI的联合编程。因为有队列命令之间的依赖关系,他们表示通过使用事件对象绑定的命令。在清单2第二阶段计算,雅可比_甚至_一个和雅可比_奇怪的_B被使用事件对象的通信,e (1]。第二阶段通信阻塞使用事件对象的第一阶段计算,e (0]。另一方面,在清单4,函数调用之间的依赖关系由OpenCL事件管理机制,管理,因此主机线程释放控制计算和通信。在代码中,clEnqueueSendrecvBuffer利用一个OpenCL命令两个MPI进程之间交换数据的内部调用MPI_SendrecvOpenCL事件管理的控制。因此,宿主线程只是等待最后的迭代通过调用clFinish。

4.2。与现有的MPI函数进行互操作

在clDataTransfer, MPI流程使用clDataTransfer命令传输数据从/到设备内存缓冲区。如果一个MPI的过程需要传输数据从/到主机内存缓冲区,clDataTransfer允许MPI程序使用标准MPI等功能MPI_Isend和MPI_Irecv与远程通信设备以及远程主机。清单5表明,0级的MPI过程从一个远程设备接收数据管理的MPI 1级的过程。一个特殊的MPI_数据类型值,MPI_CL_MEM的第三个参数MPI_Irecv为了表示发送方应该是一个沟通者设备和数据在设备内存。如果MPI_CL_MEM,发送方和接收方主机和设备之间的数据传输效率合作记忆。使用类似的方法MPI_数据类型可以看到在5),即使他们只延长MPI但不是OpenCL。

见清单5非阻塞MPI函数可用于节点之间的通信从/到主机内存缓冲区。因此,需要收到之前的数据clEnqueueWriteBuffer线(14)执行写数据到设备内存级别0。另外,内核在执行线(11)节点之间的通信。表达之间的依赖非阻塞MPI函数调用和OpenCL命令,clDataTransfer扩展提供了一个函数来创建一个OpenCL事件对象相对应MPI_请求非阻塞的MPI函数调用。使用事件对象,另一个OpenCL命令可以执行非阻塞MPI函数完成后;MPI操作之间的依赖和OpenCL操作正确执行不需要主机干预。在清单5事件对象是用来确保MPI_Irecv是之前完成设备内存缓冲区写入数据。

MPI互操作性非常重要,因为许多应用程序已经开发了以这样一种方式,通过MPI函数调用cpu管理所有节点之间的通信。考虑到重要性,clDataTransfer扩展不是被设计作为一个独立的通信库但OpenCL扩展与MPI互操作。互操作性,可以逐步移植遗留应用程序异构计算系统逐渐取代了MPI函数调用clDataTransfer扩展。这并不意味着所有节间通信应该替换为clDataTransfer扩展。我们认为MPI和OpenCL需要延长有效的互操作。

尽管clDataTransfer扩展提供的节间点对点远程主机和设备之间的通信,它目前并不提供任何集体交流。这是因为MPI集体通信阻塞的函数调用,不需要OpenCL扩展描述集体之间的可靠通信和OpenCL的命令。如果优化集体通信设备内存对象是必需的,我们可以在MPI集体隐藏实现细节沟通功能,而不是开发的一组特殊的集体交流功能设备内存对象。mpi - 3.0标准将支持集体非阻塞通信,一些集体通信和同步机制之间的非阻塞OpenCL命令可能需要在未来。在这种情况下,它将有效的进一步扩展OpenCL使用其同步事件管理机制。

5。评估和讨论

在本节中,提出了扩展的性能影响是通过展示讨论的影响隐藏host-device数据传输延迟和性能改进。在这项工作中,GPU史密斯沃特曼算法的程序(11是首先用来评估性能通过重叠host-device数据传输和文件访问。然后,Himeno基准(13)和技术增长模拟(14)采用MPI互操作性的评价,提出改进的扩展。

三个系统称为Masamune,丽鱼科鱼,RICC用于以下评价。Masamune与Intel Xeon e5单个节点的PC - 2670 CPU 2.60 GHz和NVIDIA GeForce GTX泰坦GPU。丽鱼科鱼是一种小型PC集群系统的四个节点,每一个都包含一个英特尔酷睿i7 930 2.8 GHz CPU和一个NVIDIA Tesla C2070 GPU。通过千兆以太网网络的节点连接。另一方面,这里集成多功能PC集群的集群的集群(RICC), 100个计算节点通过一个InfiniBand DDR连接网络。每个计算节点有两个英特尔至强5570 cpu和一个NVIDIA Tesla C1060 GPU。总结了系统规范表1。

5.1。实现

在这项工作中,我们实现了clDataTransfer扩展的NVIDIA的OpenCL和开放的MPI (15如表所示1。目前的大多数OpenCL实现专有,clDataTransfer扩展设计,这样它可以实现专有OpenCL的实现。在实现中,我们必须考虑至少三分。有一点是如何实现clDataTransfer命令,模仿标准OpenCL命令。另一个是如何实现非阻塞函数调用。另一种是如何实现管线式数据传输。

实现clDataTransfer命令的执行是由OpenCL的事件管理系统,用户内部事件对象是用于创建事件对象提供的那些额外的命令clDataTransfer扩展。因为有几个不同的行为标准事件对象和用户事件对象之间,运行时clDataTransfer扩展了,这样用户的事件对象的额外的命令可以模拟事件对象的标准OpenCL的命令。简化clDataTransfer函数的伪代码如清单所示6。在执行函数时,从应用程序的角度,clDataTransfer运行时出现如下工作。一个用户事件对象的执行状态CL_提交首先是当clDataTransfer命令队列的创建。然后,clDataTransfer运行时自动执行状态的变化CL_完整的当命令完成时结束。这允许其他命令等待完成clDataTransfer命令通过使用它的用户事件对象。因此,应用程序编程人员可以使用事件对象的clDataTransfer命令一样的标准OpenCL命令。

在清单的clDataTransfer函数6可以调用阻塞或非阻塞模式。调用clDataTransfer函数没有阻断宿主线程,clDataTransfer运行时内部产生另一个线程专门用于数据传输。因为大多数现有OpenCL实现已经生成一个CPU线程支持回调,同一个线程技术可以用来处理clDataTransfer函数调用。因此,不需要额外的线程如果clDataTransfer OpenCL供应商实现。

随着clDataTransfer实现需要调用MPI和文件访问函数从主机的线和专用的线程,它们的底层实现假定为线程安全的。文件访问函数通常是线程安全的。另一方面,在MPI,MPI_初始化_线程应与MPI_线程_多个。让InfiniBand开放MPI正常工作在一个多线程环境中,知识产权在InfiniBand (IPoIB)是用于RICC绩效评估。

在当前实现中,管线式数据传输实现自己通过引用一些论文GPU-aware MPI实现(5,6)和封装在clDataTransfer命令如清单所示6。到目前为止,包装器函数的文件I / O函数和一些主要的MPI等功能MPI_发送和MPI_Recv已经开发出来,这些函数可以执行管线式重叠的数据传输与节间host-device沟通沟通什么时候MPI_CL_MEM给出的MPI_数据类型参数。

5.2。文件访问性能的评价

5.2.1。评估持续的数据传输带宽

持续的数据传输带宽文件评估表明设备内存缓冲区clEnqueueWriteBufferFromStdioFile可以减少数据传输时间相比,传统的串行数据传输。评估持续带宽不同的存储带宽的固态硬盘(SSD)和硬盘驱动器(HDD) Masamune作为本地存储,和NFS共享文件系统用作全球存储和访问从丽鱼科鱼。

首先,我们评估多少clDataTransfer扩展可以改善持续带宽。的情况下使用clEnqueueWriteBufferFromStdioFile从文件中读取数据,然后发送到设备内存缓冲区。存储器的带宽低于主机和设备之间的数据传输通过pci - express总线。因此,持续数据传输的带宽是有限的存储带宽。自clEnqueueWriteBufferFromSdtioFile使host-device数据传输可以重叠与文件阅读,它可以减少数据传输时间,从而实现持续的带宽高于这两个数据传输的顺序执行。

图3显示了持续获得的带宽与改变大小和管道的数据缓冲区大小。纵轴显示的是持续的带宽,和水平轴是数据的大小。在图中,串行意味着数据传输时间的情况不是host-device数据传输延迟和隐藏管方式的数据传输时间安排实现的缓冲区字节管道。通过隐藏的延迟,文件读取的数据传输时间的方法,这是FileRead在图中。这些结果表明clDataTransfer扩展可以隐藏host-device数据传输延迟,因此从文件的持续性能数据传输设备内存缓冲区几乎与持续带宽的阅读文件,也就是说,FileRead。程序员可以使用的优化数据传输实现排队clDataTransfer命令。

对于Masamune从硬盘读取数据,文件读取时间变化很大,如图3。这可能是由于磁盘的带宽和预读线程在操作系统内核的行为。因此,性能是看不见的。FileRead性能有时甚至低于clEnqueueWriteBufferFromStdioFile由于固有的测量精度。

5.2.2。评估与史密斯沃特曼算法

在这项工作中,CUDA史密斯沃特曼算法的程序(11移植到OpenCL。OpenCL版本的性能是评价表明clDataTransfer可以隐藏host-device数据传输延迟的一个真正的应用程序通过重叠的文件访问延迟。在史密斯沃特曼程序,数据传输时间和计算时间可以重叠。然而,数据传输时间仍然是部分暴露在总执行时间如果计算时间短于数据传输时间。公开的数据传输时间取决于大小的问题。因此,在这个评估、计算和数据传输的重叠是禁用的,和完全暴露评估数据传输时间清楚地显示重叠host-device数据传输延迟的影响与文件访问延迟。

OpenCL程序反复读取文件到主机内存缓冲区中的数据并将它们发送给设备内存缓冲区。假设d_db和h_db是处理设备内存缓冲区和主机内存缓冲区,分别。他们的缓冲区大小readsz,文件指针《外交政策》。然后,原始代码有以下代码模式:从文件中读(h_db,readsz,1,《外交政策》);clEnqueueWriteBuffer (cmd,d_db,CL_真正的,0,readsz,h_db,0,零,空);

上述模式被替换为一个额外的OpenCL命令队列clEnqueueWriteBufferFromStdioFile(cmd,d_db,CL_真正的,0,readsz,《外交政策》,0,零,空);

评估的结果数据传输时间和改变管道缓冲区大小如图4。在这里,数据传输时间是数据传输的总时间从数据库文件到设备内存缓冲区。这些结果表明,clDataTransfer扩展可以减少数据传输时间如果管道缓冲区大小适当的配置。clDataTransfer扩展的性能改进减少管道如果缓冲区大小太小由于管道实现的运行时开销。也减少,如果管道缓冲区大小太大数据大小相比,因为流水线太大缓冲区不受益于重叠的数据传输。因此,最优管道缓冲区大小不仅取决于数据大小的存储性能,但也从一个文件被转移到一个设备内存缓冲区。管道必须动态调整缓冲区大小,因为数据大小通常是在运行时确定。图4讨论了改变管道缓冲区大小对性能的影响。自从clDataTransfer扩展隐藏了数据传输的实现细节,这在技术上是可行的使用经验参数调优或自动寻找最优的自动调谐管道缓冲区大小,如MVAPICH2-GPU CUDA的支持。

史密斯沃特曼计划,从文件读取数据大小范围从511字节到4 mb,因此相对较小。持续带宽的文件读取和host-device低数据传输成为一个小数据块的转移。如果程序是用于大型输入数据,我们相信clDataTransfer会更加显著的性能提升,显示在图3。

5.3。节点之间的通信性能的评估

5.3.1。点对点的通信性能

clDataTransfer扩展的一个优点超过传统的联合编程的MPI和OpenCL clDataTransfer扩展可以隐藏实现细节使用系统优化的高效的数据传输。

图5显示持续的差异之间的带宽固定,映射,管线式实现部分中描述3。在图中,“廉线()的结果表明管线式数据传输管道的缓冲区大小mb。评价结果图5(一个)表明之间的性能差异三个小丽鱼科鱼系统中实现。这是因为他们持续对带宽有限的带宽GbE互连网络。host-device沟通的时间远远短于节间的通信,因此管线式实现几乎提高了持续的带宽。另一方面,在图5 (b),有一个持续的带宽相差甚大,这三个实现。此外,持续带宽的管线式实现更改管道缓冲区大小。管道与一个相对较小的管道时最有效的缓冲消息大小很小,因为管道缓冲区大小需要小于消息大小。另一方面,一个大的管道缓冲区会导致更高的持续带宽大消息,因为每个管道发送缓冲区通常持续带宽的增加管道缓冲区大小。因此,最优管道缓冲区大小变化至少在消息大小不同。

从上面的结果,很明显,使用系统优化往往需要多节点GPU应用程序实现高绩效,因此一些抽象的节点之间的数据传输对于高performance-portability是必要的。例如,RICC,固定数据传输总是快于映射的一个,而小消息的映射数据传输速度更快的丽鱼科鱼由于短延迟实现。clDataTransfer扩展提供了接口,抽象的节点之间的数据传输,从而允许应用程序程序员使用优化数据传输没有复杂的编程技术。自动选择机制背后的数据传输的实现可以采用接口。当前实现clDataTransfer运行时可以使用固定或映射的数据传输为小消息,和管线式大消息的数据传输可以执行。管线式数据传输也可以使用固定或映射实现数据传输。在接下来的评估、映射和固定数据传输用于丽鱼科鱼和RICC,分别。当然,其他优化数据传输可以被纳入可用的运行时和应用程序无需更改代码,导致高performance-portability跨系统类型、尺度,可能代。

5.3.2。评估与Himeno基准

使用clDataTransfer扩展的性能影响是首先通过比较评估持续表现三个Himeno基准的实现。一种实现叫做hand-optimized实现了(13]。hand-optimized实现使用固定数据传输交换光环大约750 kb的数据。另一个称为串行实现,几乎是一样的hand-optimized实现但所有的计算和通信被序列化。串行实现的性能应该是最低的。另一种是实现使用clDataTransfer扩展,称为clDataTransfer实现。

图6显示了持续演出的三个Himeno基准的实现造数据。自从hand-optimized实现设计良好的重叠计算和通信,它总是可以实现更高的性能比串行执行;丽鱼科鱼的平均加速率分别为51.2%和15.2%,RICC分别。clDataTransfer实现的性能几乎总是与hand-optimized实现的,因为沟通的时候hand-optimized和clDataTransfer实现不暴露他们的总执行时间。因此,clDataTransfer扩展允许应用程序程序员容易重叠的通信和计算只需发送节点之间的通信命令设备和利用OpenCL事件对象来执行OpenCL命令之间的依赖关系。

结果在图6(一)获得使用丽鱼科鱼的网络性能较低相比,计算性能。计算时间的比值串行实现的沟通时间也显示在图。只有在丽鱼科鱼中有四个节点,计算比小于1,交流,因此通信时间和计算时间不能完全重叠固定数据传输用于通信。在这种情况下,hand-optimized实现的性能明显低于clDataTransfer实现。性能差异的主要原因是,clDataTransfer背后的映射数据传输实现速度比固定的数据传输。这些结果清楚地表明的重要性与系统相关的优化高效数据传输。clDataTransfer扩展的编程模型封装了数据传输过程中,一个应用程序程序员不需要知道实现细节,可以自动使用优化的实现从一个简单的写的代码如清单所示4。

5.3.3。评价实际应用

clDataTransfer扩展的性能影响是进一步讨论以避免增长模拟(14作为真正的应用程序的一个例子。仿真代码开发的整个生长过程的数值分析二元合金在热等离子体合成技术。虽然各种现象被认为是模拟技术发展过程中,约有90%的原始代码的总执行时间是花在模拟凝固在纳米粒子的过程。

在下面的评价,clDataTransfer扩展应用于并行版本的仿真代码,只有凝固的例程使用MPI并行,及其内核循环使用OpenCL进一步加快。成核和冷凝等其他现象由一个主机线程计算,和系数约42 mb的数据所需的凝血常规分布从主机线程每个节点在每个模拟步骤。对于仿真代码,实现了两个版本澄清的效果使用提供的优化数据传输clDataTransfer扩展。一个是基线实现只使用MPI_Isend和MPI_Recv数据分布系数。另一个是clDataTransfer实现,它使用MPI_Isend与MPI_CL_MEM发送主机内存缓冲区和系数clEnqueueRecvBuffer接收他们。

图7显示了比较结果RICC表演的两个实现。与Himeno基准不同,通信开销显然是暴露在这个模拟程序的总执行时间。由于MPI并行分解方法,节点的数目必须是40除数。因为穷人的并行性,性能降低当节点数的增加,超过8。

如图7clDataTransfer优于基准的实现,因为它可以利用一个优化的实现,重叠host-device通信节点之间的通信的足够大的管线式时尚信息。因此,这些结果表明,更高的性能可以通过适当交互MPI和OpenCL,和clDataTransfer使我们能够表达互操作在一个简单而有效的方法。

在上面的评价,通过取代的结合MPI_Recv和clEnqueueWriteBuffer与clEnqueueRecvBuffer管道数据传输是用于沟通和导致持续的高带宽。因此,研究结果还表明,应用程序员可以逐步提高MPI程序使用clDataTransfer扩展。这很重要,因为大多数现有的应用程序已经开发使用MPI。

6。结论

clDataTransfer,本文提出了一个OpenCL扩展允许OpenCL执行主机之间的数据传输需要协作和计算设备。clDataTransfer扩展,额外OpenCL命令定义封装常见的编程模式的数据传输从/到设备内存,比如节间通信和文件访问。额外的命令是相同的方式执行其他OpenCL的命令。使用OpenCL事件对象,我们可以表达依赖传统的和额外的命令。因此,数据传输的额外的命令是纳入OpenCL以自然的方式执行模型。

数据传输是抽象OpenCL命令,数据传输的实现细节隐藏从应用程序代码。因此,clDataTransfer能够利用新特性的最新设备没有任何用户代码更改。因此,clDataTransfer将使今天的应用程序受益于硬件的改进没有进行任何代码更改甚至没有重新编译应用程序。即clDataTransfer不仅可以提高性能,但也跨系统类型的性能可移植性,天平,一代又一代。

绩效评估的结果清楚地表明,clDataTransfer可以实现高效的数据传输,同时隐藏复杂的实现细节,导致更高的性能和可伸缩性。此外,使用clDataTransfer扩展,主机应用程序的线程不被序列化的相关业务数据传输。因此,clDataTransfer扩展允许应用程序程序员方便地使用重叠通信与计算和存储访问的机会。

虽然这工作关注OpenCL,我们相信这个想法本身可能适用于其他诸如CUDA编程模型。在未来,我们将进一步提高扩展,这样就可以支持其他类型的任务,需要帮助的主机线程,如系统调用。

利益冲突

作者宣称没有利益冲突有关的出版。

确认

作者要感谢教授Mayasa Shigeta Fumihiko伊诺的大阪大学教授允许他们使用他们的绩效评估的仿真代码。作者还要感谢理研综合集群的集群(RICC),这里的用户支持和计算机资源用于绩效评估。这项研究部分由JST波峰“进化的方法来建设大规模并行软件开发环境的异构系统”和科学研究补助金(B)号,25280041和25280012。能源部温哥华项目支持的工作也是部分(DE-SC0005515)。