文摘

利用异构计算平台已成为普遍趋势,使重要的可移植性问题。OpenCL(开放计算语言)服务于这一目的通过使便携式执行在不同的架构。然而,不可预测的性能变化在不同的平台上已经成为负担的程序员编写OpenCL的应用程序。对传统的多核cpu尤其如此,因为一般OpenCL应用cpu的性能落后于同行的表现写在传统cpu并行编程模型。在本文中,我们评估OpenCL无序多核cpu上的应用程序的性能从架构的角度来看。我们评估OpenCL应用在各个方面,包括API开销,调度开销,指令级并行性,地址空间,位置数据,数据局部性,和向量化,比较OpenCL传统cpu的并行编程模型。我们的评估显示独特的性能特征OpenCL的应用程序,还提供了洞察cpu上获得更好的性能优化指标。

1。介绍

异构体系结构已经得到普及,可以看到来自AMD的融合和英特尔的桑迪大桥(1,2]。许多研究显示了高性能的异构体系结构的承诺和能源效率。然而,如何利用异构架构考虑性能和能源效率仍然是一个具有挑战性的问题。OpenCL并行编程是一个开放的标准在不同的架构,这使得它可以便携的方式表达并行性,这样在OpenCL编写的应用程序可以运行在不同的体系结构没有代码修改(3]。目前,许多厂商已经发布了自己的OpenCL框架(4,5]。

尽管OpenCL提供可移植性在多个架构,可移植性问题仍在性能方面。不可预知的性能变化在不同的平台上已经成为负担的程序员编写OpenCL的应用程序。有效优化技术是不同的根据执行内核的体系结构。特别是自OpenCL与CUDA共享许多相似之处,为NVIDIA gpu开发,许多OpenCL应用程序是现代多核cpu不优化。一般OpenCL应用cpu的性能落后于预期的性能程序员考虑传统并行编程模型。期待来自程序员的经验与传统编程模型。OpenCL应用程序显示cpu上非常贫穷的性能相比,在传统编写的应用程序编程模型。

OpenCL的原因我们认为cpu计算设备如下。(1)也可以用来增加CPU OpenCL应用程序的性能通过使用CPU和gpu(尤其是当CPU空闲)。(2)因为现代的cpu有更多的单位向量,cpu和gpu之间的性能差距已经降低了。例如,即使是大规模并行的内核,有时可以比cpu gpu,根据输入的大小。在某些工作负载高分支分歧或指令级并行度高(独立),CPU也可以比GPU。

使用OpenCL的一个主要好处是,相同的内核可以很容易地在不同的平台上执行。OpenCL,很容易在运行时动态地决定使用哪一个设备。OpenCL cpu和gpu之间选择一个计算装置的应用程序在运行时可以很容易地实现。然而,如果应用程序是写在OpenMP,例如,它不是简单将一个应用程序使用两个cpu和gpu。

在这里,我们评估OpenCL应用程序的性能在现代无序多核cpu从架构的角度来看,关于应用程序如何利用cpu的硬件资源。我们彻底评估OpenCL各个方面上的应用程序可能会改变他们的表现。我们重新审视通用性能指标已经轻轻评估在以前的作品,尤其是对OpenCL内核cpu上运行。使用这些指标,我们也验证OpenCL的电流限制和可能的改进的性能代价。总之,本文的贡献如下。(我)我们为程序员提供一个指导方针理解OpenCL cpu上的应用程序的性能。程序员可以验证OpenCL内核是否充分利用CPU的计算资源。(2)我们讨论的有效性OpenCL应用在多核cpu和可能的改进。

本文的主要目的是提供一种方法来理解OpenCL cpu性能。尽管OpenCL可以在cpu和GPU执行,大多数以前的工作都集中在只有GPU性能问题。我们相信我们的工作增加OpenCL cpu上的可理解性和有助于减少程序员的编程开销从头开始实现一个单独的CPU-optimized版本。一些关于OpenCL cpu上讨论一些先前的研究方面提出了,但他们都缺乏定量和定性评估,使他们难以使用当程序员想估算各方面性能的影响。

部分2描述了背景和架构方面了解OpenCL cpu性能。然后,我们评估OpenCL节中关于这些方面的应用3。我们在部分评审相关工作4并得出结论。

2。背景和标准

在本节中,我们描述的背景几个方面影响OpenCL cpu上应用程序的性能:API开销,线程调度开销,指令级并行性,数据传输,数据局部性和编译器自动向量化技术。这些方面一直强调在学术界和工业来提高应用程序的性能cpu上多个编程模型。尽管大多数在这一节中描述的架构方面是很好理解的基本概念,大多数OpenCL应用程序的写作目的不是考虑这些方面。

2.1。API的开销

启动内核OpenCL高开销,可以忽略其他传统cpu的并行编程模型。除了内核执行的计算设备,OpenCL OpenCL API函数调用在宿主代码需要协调执行的内核开销。OpenCL应用程序的一般步骤如下(3]:(1)打开一个OpenCL上下文。(2)创建一个命令队列接受执行和内存请求。(3)分配OpenCL的内存对象的输入和输出的内核。(4)在线编译和构建内核代码。(5)设置内核参数。(6)组工作项维度。(7)启动内核执行(排队内核执行命令)。(8)收集结果。

OpenCL的复杂步骤应用程序由于OpenCL多个架构设计理念强调可移植性。因为OpenCL的目标是使一个应用程序运行在多个架构,使OpenCL尽可能灵活的编程模型。图1显示了OpenCL平台模型和OpenCL如何提供可移植性。OpenCL平台由一个主机和计算设备的列表。一个主机连接到一个或多个计算设备和负责管理资源计算设备。计算设备是一个抽象的处理器,它可以是任何类型的处理器,如传统的CPU、GPU, DSP。有一个单独的设备的内存和计算设备的列表计算单位。一个计算单元可以有多个处理元素(PEs)。通过这种抽象,OpenCL支持便携式执行。

相反,灵活性为各种平台支持不存在传统多核cpu的并行编程模型。许多OpenCL的api,这对OpenCL应用重要的执行时间不存在传统并行编程模型。OpenCL的计算设备和上下文隐含传统编程模型。用户不需要查询平台或计算设备和显式创建上下文。

OpenCL的独特特征的另一个例子比传统编程模型是“即时编译”[6在运行时)。在许多OpenCL应用程序,内核编译时间由JIT编译器开销的执行时间。相反,是静态编译完成,而不是在应用程序执行期间执行的其他编程模型编写的应用程序。

因此,要确定应用程序的实际表现,时间成本执行OpenCL API函数也应该被考虑。从评估,我们发现这个API的开销大于实际的计算在许多情况下。

2.2。线程调度

不像其他并行编程语言如TBB [7]和OpenMP [8),OpenCL的编程模型是一个单指令和多头螺纹(SIMT)模型就像CUDA (9]。一个OpenCL内核描述单个线程的行为,和主机应用程序显式地声明线程的数量表达应用程序的并行性。在OpenCL术语中,被称为一个线程工作项(一个线程在CUDA)。OpenCL程序员可以形成一套workitems作为工作组(一个threadblock在CUDA),程序员可以同步workitems之间障碍和mem_fence。一个工作组由workitems的多维数组。图2显示了OpenCL执行模型和如何OpenCL内核OpenCL计算设备上的映射。OpenCL,内核分配的计算设备上,工作组是一个计算单元上执行。处理一个工作项的处理元素(体育)。获得更好的性能,程序员可以调整workitems和工作组的大小变化。

OpenCL是很常见的应用程序启动大规模的内核期望加速并行执行的线程数量。然而,OpenCL应用程序的可移植性的性能不是维护在不同的体系结构。换句话说,一个最优的决定如何在gpu并行化(分区)内核cpu上通常不会保证良好的性能。内核的分区决策是通过改变workitems的数量和工作组的大小。

2.2.1。Workitems数量

首先,workitems的数量和所做的大量的工作对cpu和gpu workitem影响性能的不同。大量的短workitems伤害cpu性能但帮助gpu性能。性能差异来自不同的cpu和gpu之间的建筑特点。在gpu,处理一个工作项标量处理器(SP)或一个单一的SIMD车道。众所周知,gpu专门支持大量并发运行的线程,以及高线程级并行(张力腿平台)实现高绩效是至关重要的10- - - - - -13]。相反,cpu上,张力腿平台核心的数量有限,所以使用多个线程做同样多的工作不能帮助cpu性能但伤害由于模拟大量的并发执行的开销workitems少数的核心。

2.2.2。Workitems数量和指令级并行性(独立)

workitems影响指令级并行性的数量(独立)OpenCL内核的cpu。增加独立在GPU应用程序没有一个受欢迎的性能优化技术。的原因如下。首先,硬件可以独立探索张力腿平台所以不会影响显著的性能。第二,硬件不独立探索太多。GPU处理器是一个顺序调度处理器和还不支持分支预测增加独立。然而,在cpu,增加独立的硬件设计等多个特性超标量体系结构和分支预测执行。

现代超标量处理器并行执行多个指令调度多个独立的指令在一个时钟周期内利用cpu的多个功能单元。超标量体系结构的cpu使用的硬件检查数据指令在运行时之间的依赖关系和并行调度指令来运行14]。

OpenCL应用程序的性能问题之一cpu内核通常是主要利用TLP写的,而不是独立。OpenCL编程模型是一个SIMT模型,对OpenCL应用程序是很常见的一个巨大数量的线程。自独立计算不同元素分为不同的线程的指令,大多数指令在一个工作项在内核通常依赖于之前的指令,所以通常大多数OpenCL内核独立一个人;只能派出执行一条指令在一个工作项。相反,在传统编程模型如OpenMP,独立的指令之间存在不同的循环迭代。cpu上获得更好的性能,OpenCL内核应该写有更多独立的指令。

2.2.3。工作组的大小

第二个重要组成部分是工作组的大小。工作组的工作量大小决定在一个工作组,工作组的内核。在gpu,工作组或多个组流多处理器上执行(SM),这相当于一个物理核心在多核CPU上。同样,工作组是由一个逻辑CPU的核心处理15,16]。(尽管它取决于实现,许多实现这一特点共同点。)。每个工作组工作负载大小太小使调度开销更重要在cpu上总执行时间的线程上下文切换开销变得更大。

OpenCL程序员可以显式地设置工作组大小或让OpenCL的实现决定。如果零值传递工作组大小当主机应用程序调用clEnqueueNDRangeKernelOpenCL实现自动分区全球workitems适当数量的工作组。

2.2.4。提出解决方案和局限性

许多建议减少序列化的调度开销已经提出(15- - - - - -17]。调度开销不是一个基本问题OpenCL的编程模型。OpenCL实现可以减少开销比其他非最优实现。序列化技术,多个workitems序列化到一个工作项。例如,SnuCL [15)克服大量的开销workitems通过序列化它们在运行时workitems较少。然而,即使序列化,多个OpenCL实现cpu仍然有很高的调度开销,因为编译器分析的复杂性。因此,而不是使用许多workitems,通常是对OpenCL gpu的应用程序,我们最好分配更多的工作要用更少的cpu上workitems每个工作项。从我们的实验结果同意上述的推论。

2.3。内存分配和数据传输

一般来说,一个并行编程模型可以有两种类型的地址空间选项:统一内存空间和独立的内存空间18]。传统编程模型对于cpu提供统一内存地址空间的顺序代码和并行代码。统一内存空间的好处是容易编程,没有明确的数据传输为内核执行。

相反,尽管很难程序员编程,OpenCL为程序员提供独立的内存空间。这是因为大多数异构计算平台有不相交的记忆系统由于不同的内存需求不同的架构。OpenCL假定的目标系统,主机之间的通信和计算设备执行明确的系统网络,如pci - express。但是,离散记忆系统的假设是不正确的,当我们使用cpu作为内核执行的计算设备。主机和计算设备共享相同的内存系统资源,比如最后一级缓存,片上互连,内存控制器和后发展出。

不相交的内存地址空间的缺点是,它需要程序员显式管理主机之间的数据传输和计算设备内核执行。OpenCL应用程序一样,来回传输数据应该为了处理主机或计算设备(3而成为不必要的,我们只使用主机计算。最小化数据传输开销在特定架构,OpenCL程序员通常需要重写主机代码(3]。通常,他们需要改变内存分配旗帜或api使用不同的数据传输性能。例如,程序员应该在主机内存分配内存对象或设备内存取决于目标平台。这些重写努力程序员的负担,甚至是浪费时间由于缺少建筑或运行时的一个特定的知识系统在大多数情况下。

2.3.1。内存分配的旗帜

重写的努力之一是改变内存分配的旗帜。OpenCL为程序员提供了多个选择内存对象分配国旗当程序员调用clCreateBuffer可能影响数据传输和内核执行的性能。内存分配标记用于指定对象的访问内核和分配。

访问类型。首先,程序员可以指定如果内存对象是只读内存对象(CL_MEM_READ_ONLY),或只写一个(CL_MEM_WRITE_ONLY当内核内部引用。程序员可以使用内存对象作为输入到内核设置为只读和内存对象作为输出从内核只写。如果程序员没有指定访问类型,默认选择是创建一个可以读取和写入的内存对象由内核(CL_MEM_READ_WRITE)。 CL_MEM_READ_WRITE也可以由程序员显式指定。

分配的地方。程序员可以指定另一个选择是分配一个内存对象的地方。当程序员没有指定分配的位置,设备上的内存对象分配内存OpenCL的计算设备。OpenCL支持固定内存。当主机代码创建内存对象使用CL_MEM_ALLOC_HOST_PTR国旗,host-accessible上的内存对象分配内存驻留的主机上。不同于分配的内存对象设备内存,不需要转移内核执行的结果返回到主机内存当结果所需的主机。

2.3.2。不同的数据传输api

OpenCL还提供了不同的api主机之间的数据传输和计算设备。主机可以排队命令读取数据从一个由OpenCL内存对象clCreateBuffer调用内存对象主要是由malloc主机内存(通过调用clEnqueueReadBufferAPI)。主机还可以排队命令写入数据到OpenCL主机内存中的内存对象从内存对象(通过clEnqueueWriteBufferAPI)。程序员也可以映射一个OpenCL内存对象的host-accessible指针映射对象(通过clEnqueueMapBufferAPI)。

2.4。向量化和线程关联

2.4.1。向量化

利用SIMD单位一直是一个关键的cpu性能优化技术(19]。因为SIMD指令可以对多个数据项进行计算同时,SIMD利用率可能会使应用程序更高效。许多厂商已经发布了各种SIMD指令扩展指令集架构,如MMX [20.]。

提出了各种方法利用SIMD指令:使用优化的函数库,如英特尔IPP (21)和英特尔MKL (22),使用c++向量类与英特尔ICC (23),或者使用DSL编译器如英特尔SPMD程序编译器(24]。程序员也可以在装配或使用程序的内在功能。然而,所有这些方法假设重写代码。由于这个限制,并帮助程序员轻松地编写应用程序利用SIMD指令,自动向量化技术已经在许多现代编译器实现(19,23]。

很自然的程序员认为编程模型的差异没有影响编译器自动向量化技术在同一架构。例如,如果一个内核OpenCL和OpenMP编写和实现都是用类似的方式编写的,程序员希望代码都是矢量化以类似的方式,从而使类似的性能数据。虽然这取决于实现,通常这不是真的。不幸的是,今天的编译器是非常脆弱的关于vectorizable模式,这取决于编程模型。应用程序应该满足一定的条件下,为了充分利用编译器自动向量化技术(19]。我们的评估中3.5.1显示了一个示例的脆弱性和验证可能的编程模型对向量化的影响。

2.4.2。线程关联

将线程可以影响现代多核cpu上的性能。线程可以被放置在每个核心以不同的方式,它可以创建一个性能差异。位置对性能的影响会增加更多的处理器系统上。

的性能差异可能发生多个原因。例如,由于不同的互连网络延迟,遥远的线程将需要更长的时间来相互通信,而线程接近邻核心可以更快的交流。同样,一个应用程序,该应用程序需要相邻的线程之间的数据共享可以受益,如果我们这些相邻的线程分配给附近的核心。适当的位置也可以消除通信开销利用共享缓存。出于性能的原因,大多数传统的并行编程模型支持亲和力,等CPU_AFFINITY在OpenMP [8]。

不幸的是,在OpenCL不支持线程关联。OpenCL workitem是一个逻辑线程,并不是紧密耦合的一个物理线程尽管大多数并行编程语言提供了这一功能。这个功能缺乏的原因是,OpenCL设计理念强调效率可移植性。

我们现在缺乏亲和力的支持性能的局限性之一OpenCL cpu上相比其他编程语言的cpu。我们想提出一个潜在的解决方案来提高OpenCL cpu性能。我们发现的利益更好地利用缓存OpenCL应用程序线程关联。节中给出了一个例子3.5.2。

3所示。评价

鉴于以上背景架构方面的预期影响理解OpenCL cpu性能,我们研究的目标是定量探索这些影响。

3.1。方法

实验环境为我们的评估表中描述1。我们的评估进行异构计算平台组成的多核CPU和GPU;OpenCL内核执行在英特尔OpenCL平台(4)或NVidia OpenCL平台(5]。我们实现了一个执行框架,这样我们可以改变和控制在许多方面的应用程序,而无需更改代码。执行框架构建作为OpenCL全权代表从供应商库,调用OpenCL库:一个从英特尔cpu内核执行和其他从NVidia gpu内核执行。

我们使用不同的应用程序为每个评估。验证API开销,我们使用NVIDIA OpenCL标准(5]。其他方面,包括调度开销,内存分配,和数据传输,我们首先使用简单的应用程序进行评估。我们也改变每个应用程序的数据大小。应用程序移植到我们实现了执行框架。与简单的应用,评价后我们也使用煮半熟基准(25,26]。表2,3,4描述评价应用程序和它们的默认参数。

我们使用时钟执行时间。测量稳定执行时间没有变动,我们遍历内核执行直到应用程序的总执行时间达到一个足够长的运行时间,在我们评估90秒。这是足够长的有多个执行的内核数量为所有应用程序在我们的评估。使用平均每个内核调用内核执行时间计算,我们使用归一化吞吐量明显存在多个部分的性能差异。

3.2。API的开销

正如我们讨论的部分2.1,OpenCL应用程序API开销。验证API开销,我们测量每个API函数的时间成本在NVIDIA OpenCL标准执行OpenCL应用程序(5]。每个基准工作负载大小的应用程序提供了一个默认的大小。图3显示内核执行的执行时间的比率和辅助API函数,每个OpenCL的总执行时间基准。(辅助API函数OpenCL API函数被称为主机代码协调内核执行。)最后一列总意味着从每个基准数据的算术平均值。从图中,我们可以看到,大部分执行时间花在辅助API函数而不是内核执行。

的详细分析,我们分类OpenCL api为16个类别。我们组对能见度在以下多个类别。图4提供了一个详细的例子API管理费用通过展示每个API函数的执行时间分布类别oclReduction。 队列的命令包括内核执行时间和主机之间的数据传输时间和计算设备和占12.1%的执行时间。我们发现这个API的开销大于实际的计算。

3.2.1之上。开销各种平台支持

图5显示了每个类别的执行时间比每个OpenCL的总执行时间基准。图中显示的性能下降由于各种平台的灵活性。我们看到的API函数平台,设备,上下文类别消耗超过80%的总执行时间平均每个OpenCL的基准。这些类别的需要调用API函数来自这样一个事实:每个OpenCL应用程序需要建立一个执行环境的详细机制将会改变,这取决于平台。从我们的评估,我们也看到这些类别中的每个调用API函数需要一个长的执行时间。特别是,上下文管理api产生大量执行时间开销。图6显示了执行时间的分布clCreateContext和clReleaseContext在每个基准总执行时间。这些函数被称为最多一次在每个OpenCL基准。但在传统的并行编程模型,环境和设备是隐式的,所以没有必要叫这样的管理功能。

3.2.2。由于JIT编译开销

OpenCL内核列表在应用程序的表示cl_program对象。cl_program创建对象时使用clCreateProgramWithSource或clCreateProgramWithBinary。通过调用执行JIT编译clBuildProgram函数或一个序列clCompileProgram和clLinkProgram函数cl_program对象构建项目可执行一个或多个与项目相关的设备(3]。

JIT编译开销是另一个源API的开销。图7显示了执行时间的分布程序类别的所有类别,除了执行时间的总和平台,设备,上下文类别。我们排除这三个类别,我们在前一节评估。图清晰地显示了性能退化由于JIT编译。我们看到的API函数程序类别消耗约33%的总执行时间为13个类别的API函数包括内核执行。执行时间的开销clBuildProgram不是在大多数基准测试可以忽略不计。

缓存。缓存JIT编译的代码可以帮助减少开销。缓存的一些想法在OpenCL是可用的。程序员可以通过使用提取编译后的二进制文件clGetProgramInfo使用文件I / O API函数和存储功能。当内核代码不修改自缓存,程序员可以加载缓存的磁盘上的二进制和使用执行JIT编译的二进制而不是在每一个应用程序的执行。

3.2.3。总结

在本节中,我们可以看到高开销的显式上下文管理(部分3.2.1之上)和JIT编译(部分3.2.2在OpenCL应用程序)。这些独特的特点OpenCL比传统编程模型为便携式执行多个架构。

应该注意的是,工作量大小的评价部分3所示。2是应用程序提供的大小为默认负载大小、相对较小。因此,这些费用可以减少很大的工作量大小,因此内核执行时间很长。但也是事实,这些开销不是微不足道的小工作量大小,所以程序员时应该考虑工作量的大小决定是否使用OpenCL。

3.3。线程调度

3.3.1。Workitems数量

与讨论相关的部分2.2。1评估的效果每workitem workitems数量和工作量的大小,我们对OpenCL的应用程序执行一个实验每workitem分配更多的计算。我们多个workitems合并到一个工作项,形成一个循环内的内核。保持的总量计算相同,我们减少workitems执行内核的数量。workitems合并数量的增加1来1000年通过乘以workitems10对每一个步骤。图8显示的性能广场和Vectoraddition应用程序与不同数量的计算/工作项。表5显示workitems本评价中使用的数量。

从图8,我们发现性能增益分配更多的工作/ workitem cpu上。一个明显的例子就是一个很好的例子Vectoraddition,我们添加一个数字数组。如果我们创建尽可能多的workitems数组的大小,我们最终创造显著的cpu开销。当我们减少workitems的数量时,我们看到一个cpu的主要性能改进。我们也可以发现性能是饱和有时当工作负载分配每个工作项超过某个阈值。这表明当每个工作项都有足够的工作量,调度开销减少。

相比cpu开销高的处理许多workitems, gpu低开销维持大量workitems,作为我们的评估显示。此外,减少workitems gpu性能显著退化。大在gpu性能下降是因为我们再也不能利用gpu的许多处理单元。

性能改善的原因之一每workitem分配更多的工作负载是减少数量的指示。图9显示动态指令的数量广场和Vectoraddition应用程序与不同数量的计算/工作项。左边的图的图9显示了动态指令数包括指示OpenCL OpenCL内核api之上的指令。和正确的图的图9只代表指令从内核。

对于这个评价,我们实现一个基于销工具(27],指令的数量。该工具还标识函数指令所属。从图9,我们可以看到,指令的数量减少了更多的工作量workitem虽然计算量无论workitems的数量是一样的。指令的数量从内核OpenCL api以及增加,以便调度开销存在OpenCL api和JIT编译OpenCL内核二进制。图10显示减少开销OpenCL api增加工作量/工作项。OpenCL api是为调度的指令,而不是由程序员表示为一个OpenCL内核计算目的。所以减少数量的指示OpenCL api意味着减少开销。

图11与类似的实验显示了煮半熟的性能基准(25,26]。workitems合并的数量取决于不同的基准,因为我们不能增加工作量/工作项以同样的方式对所有内核。我们找到一个类似的性能每workitem分配更多的工作。图12代表的数量动态增加工作量/ workitem指令。

3.3.2。Workitems数量和指令级并行性(独立)

正如我们讨论的部分2.2。2,workitems的数量,因此如何并行化计算,也影响OpenCL的指令级并行性(独立)内核cpu上。合并多个workitems不仅可以减少调度开销也提高性能通过利用独立。

评估独立工党影响CPU和GPU,我们实现了一组计算密集型的共同特征的微基准测试。每一个指标都有一个相同数量的动态指令和内存访问。每个基准也有相同的指令混合物,如分支指令的数量比例超过总数的指令。每个基准之间唯一的区别是由不同数量的独立独立指示。从基线的实现,我们增加操作变量的数量,这样独立的指令可以增加的数量。例如,在的情况下独立1,下一个指令的输出取决于先前的指令,这样独立的指令的数量;但在的情况下独立2依赖,一个独立的指令之间存在两个指令。

图13显示了性能越来越独立。我们提供足够的workitems充分利用张力腿平台。workitems的数量是相同的微基准测试。左边设在代表cpu的吞吐量,正确的代表了gpu的吞吐量。从图中,我们发现性能的提高取决于独立工党OpenCL内核cpu上的价值。相反,没有性能变化在gpu上不同的指令级并行度。

3.3.3。工作组的大小

与讨论相关的部分2.2。3,workitems在工作组的数量会影响OpenCL的应用程序的性能。我们评估工作组的作用大小、cpu和gpu。我们不同的workitems工作组工作组大小(通过一个不同的理由local_work_size内核调用)。我们维持workitems内核的总数是一样的。表6显示了不同的工作组为每个指标大小,和数字14,16,18显示不同的工作组大小的应用程序的性能。当零参数是通过内核调用,工作组的大小是由OpenCL隐式定义的实现。

基准可以分为三个类别,根据不同的行为。第一组包括广场,Vectoraddition,天真的实现Matrixmul; Matrixmul属于第二组;和Blackscholes属于过去。

广场,Vectoraddition,和天真的实施Matrixmul显示性能提升与增加工作组大小在CPU上,可以看到在图14。在广场和Vectoraddition应用程序性能的实现零工作组大小小于我们达到的峰值性能。这意味着程序员应该显式地设置工作组大小的最大性能。性能小规模工作组也不好在gpu工作组分配/ SM以来,这样的小工作组大小使gpu无法利用许多扭曲SM。即使没有可用硬件TLP在逻辑核心CPU上(评估CPU是SMT处理器,所以多个逻辑内核共享一个物理核心),性能增加,工作组的大小。这是因为大量的工作组管理的开销,许多线程在很多实现中,会降低。我们还发现性能是饱和在某个工作组的大小。

左边的图的图15显示动态指令的数量广场,Vectoraddition,天真的实现Matrixmul与不同的工作组大小cpu上。正确的图的图15显示指令从内核指令的比率clEnqueueNDRangeKernel对于这些应用程序使用不同的工作组的大小。从左边的图的图15,我们可以看到较大的指令数量减少了工作组的大小。这是因为OpenCL的指示api的数量减少,我们可以看到从右边图的图15。OpenCL内核上的指令的数量保持不变,不管工作组的大小。

从图我们可以看出16,我们还看到一个显著的性能提升Matrixmul应用程序增加工作组的大小。这个应用程序的最佳工作组大小是不同的,这取决于平台。为输入1和2cpu上,工作组的最优规模8×8,但gpu的最佳大小16×16。这里的性能不仅取决于调度开销,而且缓存使用。 Matrixmul利用本地内存OpenCL的阻塞。工作组的大小可以改变当地的内核的内存使用。由于缓存的大小在cpu和gpu的暂时存储器是不同的,最优工作组可以是不同的大小。图17显示了减少数量的动态指令Matrixmul随着工作组的大小。

与其他应用程序不同,Blackscholes显示不同的cpu和gpu性能行为。我们可以看到在图18,工作组大小不会改变cpu上的性能,但它影响了gpu的性能显著。由于工作负载分配在一个工作项与其他应用程序相比,相对较长管理大量的开销工作组就可以忽略不计。相反,扭曲的数量定义的SM是工作组gpu的大小,使小型工作组大小gpu性能低。图19显示指令的数量不会改变Blackscholes,不管工作组的大小。

图20.显示与不同的工作组大小煮半熟的性能基准。我们从一个增加工作组的大小16次乘以2对每一个步骤。自从工作组大小CP: cenergy内核是二维的,我们增加了工作组内核在两个方向上的大小。CP: cenergy ()代表了性能与工作组大小1×8,2×8,4×8,8×8,16×8。 CP: cenergy ()代表了性能与工作组大小16×1,16×2,16×4,16×8,16×16。一般来说,我们发现性能有一个很大的工作组的大小。性能变得饱和,当有足够的计算在工作组。图21显示性能是由于减少了调度开销,这是由数量减少的动态指令。

3.3.4。总结

在这里,我们总结研究结果对OpenCL线程调度程序。(1)分配更多的工作/ workitem通过手动合并多个workitems减少cpu调度开销(部分3.3.1)。(2)高独立cpu上提高性能而不是在gpu(部分3.3.2)。(3)工作组在cpu和gpu大小影响性能。一般来说,大型工作组大小增加性能通过减少cpu调度开销,使得利用高TLP在gpu。工作组的大小也会影响缓存使用情况(部分3.3.3)。

3.4。内存分配和数据传输

与讨论相关的部分2.3评估的性能影响不同的内存数据传输对象分配旗帜和不同的api,我们对OpenCL的应用程序执行一个实验不同的以下选项的组合。测量准确的执行性能,我们使用一个阻塞调用内核执行命令和内存对象命令,这样没有命令重叠与其他命令。结合我们使用三维如下。

3.4.1。评估选择内存分配和数据传输

(1)不同的api数据传输:(我)明确的转移:clEnqueueReadBuffer和clEnqueueWriteBuffer明确的读和写;(2)映射:clEnqueueMapBuffer与CL_MAP_READ,CL_MAP_WRITE隐式的读和写。(2)内核访问在一个引用类型时内核:(我)内核访问内存对象为只读/只写:(一)CL_MEM_READ_ONLY内核的输入;(b)CL_MEM_WRITE_ONLY计算结果从内核;(2)内核和读/写访问内存对象:CL_MEM_READ_WRITE所有的内存对象。(3)分配一个内存对象:(我)在设备上分配内存;(2)在host-accessible内存分配主机(固定内存)。

3.4.2。度量:应用程序吞吐量

这里提供的吞吐量性能,包括数据传输时间、主机和计算设备之间,不仅仅是内核执行计算设备上的吞吐量。例如,一个应用程序的吞吐量成为吞吐量的一半,当我们只考虑内核执行时间如果主机之间的数据传输时间和计算设备=内核执行时间。我们的计算方式说明了应用程序的吞吐量

3.4.3。不同的数据传输api

我们比较不同的数据传输api的性能在所有可能的配置标志。(组合如下:(1)只读/只写内存对象+配置在设备上;(2)只读/只写内存对象+主机分配;(3)读写内存对象+配置在设备上;(4)读写内存对象+分配在主机。)图22显示的性能基准测试使用不同的api用于数据传输。的设在代表规范化应用程序吞吐量(Throughput_app)当我们使用映射为数据传输在基线当我们使用显式数据传输api。我们发现映射api的性能优于显式数据传输api,不管决定在其他维度。首先,映射api的性能优越的地方分配内存对象:在主机设备内存或固定内存上。第二,映射api还表现得更好无论决定分配内存对象为只读只写或读/写/对象。

不同的api更改数据传输时间。图23显示了规范化数据吞吐量从主机转移到不同的数据传输api之间的计算设备。图24显示了一个从计算设备主机。数据传输时间短与映射的api。的差异数据传输吞吐量增加而增加的工作量大小,因此增加数据传输大小。

我们也煮半熟的性能基准测试报告与不同的api数据传输(25,26]。由于数据传输时间远远短于内核执行时间煮半熟的基准,而不是使用应用程序吞吐量所示(1),我们报告的数据传输时间从主机到设备,从设备主机和数据传输时间与不同的api。图25显示了不同的数据传输时间煮半熟的基准不同的api用于数据传输。的设在代表了数据传输时间,以毫秒为单位。左边的图在图25显示了数据传输时间从主机到计算设备使用不同的数据传输api。正确的图显示了一个来自主机的计算设备。对于简单的应用程序,我们发现映射的数据传输时间短对这些标准api。

数据传输时间的差异是由于不同的行为不同的api。当主机代码显式主机和计算设备之间传输数据,OpenCL运行时库应该为设备分配一个单独的内存对象和分配的内存对象之间的数据复制malloc调用和为设备分配的内存对象的分配clEnqueueReadBuffer调用。然而,一个单独的内存对象时不需要主机代码使用映射;只返回一个指针的内存对象是必要的。所以,复制内存对象之间变得不必要的。

3.4.4。在内核引用时内核访问类型

我们也验证指定的性能影响内存对象为只读或只写或读/写。图26显示这个标志的性能含义。的设在表示归一化吞吐量时分配内存对象为只读或只写从基线当我们分配对象为读/写。OpenCL实现可以使用的内存对象的详细信息访问OpenCL内核的优化,而不是天真的假设所有对象中读取和修改OpenCL内核。然而,我们没有看到明显的性能差异与我们的评估工作负载。内核执行时间和主机之间的数据传输时间和计算设备不不同,不管这个内存分配国旗。

3.4.5。分配一个内存对象在哪里

最后,我们还验证的性能影响分配内存对象的位置。程序员可以在主机内存分配内存对象或设备内存。图27与不同的分配位置显示的性能基准。的设在表示归一化吞吐量当我们分配的内存对象从基线当我们使用设备主机内存内存。我们发现一个分配的位置没有对性能产生巨大影响对内核执行时间和数据传输时间。这是由于设备内存和主机内存引用相同的内存,系统的内存时,计算设备的CPU。因此,不同的内存分配的位置并不意味着性能差异。相反,当计算设备不是CPU、内存分配位置可以显著地影响性能。

3.4.6。总结

在本节中,我们发现映射api执行上级明确数据传输api相比减少数据传输时间通过消除复制的cpu开销。分配的位置和内核访问类型不影响cpu上的性能。

3.5。向量化和线程关联

3.5.1。向量化

我们评估可能的编程模型对向量化的影响,即使向量化是更多关于编译器实现。评价,我们OpenCL内核移植到被OpenMP同行执行相同的计算。我们的地图上的多个workitems OpenCL循环OpenCL内核移植到他们的OpenMP同行。我们利用英特尔C / c++ 12.1.3 1.5编译器和英特尔OpenCL平台对我们的评价。程序员的期望,当我们运行相同的计算在OpenCL OpenMP应用程序,运行应该给类似的性能数字。然而,结果表明,这一假设并不成立。评估基准,OpenCL内核胜过OpenMP同行。图28显示了不同性能的OpenMP和OpenCL实现。这种不匹配的原因是不同的方式OpenMP OpenCL编译器vectorize代码。

OpenCL向量化。向量化的OpenCL内核编译器接合workitems。OpenCL向量化使几个workitems由一个向量的执行指令。向量化使多个工作项在一个线程并发处理。例如,如果目标指令集是上交所4.2,计算是基于一个单精度浮点数,然后四workitems可以同时取得进展,因此它们合并成一个工作项。通过这样做,矢量化OpenCL代码会减少动态指令数相比nonvectorized代码。

OpenMP向量化。另一方面,OpenMP编译器vectorizes循环通过展开一个循环结合了SIMD指令的生成。矢量化,循环应该是可数名词,有单独的入口和出口,有一个循环内直接控制流图(28]。许多因素可以防止一个循环的向量化OpenMP。两个关键因素是(1)不连续的内存访问和(2)数据依赖。(1)非连续内存访问:(我)连续四个浮动可能直接从内存加载在一个SSE指令;但如果这四个浮动加载不是连续的,我们将有一个负载使用多个指令;循环与nonunit跨步上述场景的一个示例。(2)数据依赖:(我)向量化需要操作的顺序在一个循环的变化因为每个SIMD指令作用于多个数据元素;但这种改变顺序可能不可能由于数据依赖关系。

例子。算法1显示了一个示例的向量化机制不同于OpenMP和OpenCL编译器如何影响是否相同的代码将矢量化。当有一个真正的数据依赖一个OpenCL内核或在一个循环迭代OpenMP平行的OpenCL内核部分,是矢量化,而不是OpenMP代码。因此,他们表现出不同的性能,即使向量化的OpenMP循环似乎是可能的。OpenCL的向量化的内核是相对简单的,因为不需要依赖项检查的传统编译器。尽管我们只显示当OpenCL编译器的例子显示了好处,相反的情况也是可能的:当vectorizes OpenMP编译器代码,但OpenCL编译器失败。

新OpenMP编译器。我们也评估OpenMP向量化与OpenMP 4.0 SIMD扩展和更新的编译器(英特尔C / c++编译器15.0.1)。OpenMP的评估显示类似的性能和OpenCL实现。编译器向量化是依赖于编译器的实现。

3.5.2。线程关联

我们评估使用CPU的性能优势亲和力OpenMP。我们使用OMP_PROC_BIND和GOMP_CPU_AFFINITY控制线程的调度处理器(8]。当OMP_PROC_BIND设置为真时,线程之间不会移动处理器。GOMP_CPU_AFFINITY使我们能够控制的分配一个线程在一个特定的核心。

我们用一个简单的应用程序来评估。应用程序的目的是验证绑定线程的影响核心缓存的利用率。性能可以提高当OpenCL运行时库映射逻辑线程内核的物理核心,这样就可以利用前面的内核执行的缓存数据。我们使用的应用程序包含两个内核:向量加法和向量乘法。计算每个内核都是分布在八个核心:第二个内核的计算依赖于第一个,使用数据产生的第一个。

表7显示了我们使用的方法。上面的表的表7代表了(一个)对齐情况下,代表(b)和下表偏差的情况。表中的数字表示逻辑线程id。线程具有相同id的内核访问相同的数据。(一)对齐情况下,我们结合的第二个内核线程的内核线程的内核。(b)偏差情况下,我们这个绑定洗牌。算法2显示了这个简单的应用程序的代码片段。

正如我们期待的,(一个)对齐情况下显示了更高的性能比(b)不一致的情况。(b)偏差一个运行时间15%。这是因为执行期间第二个内核,内核的CPU缓存错过遇到他们的私人缓存。相反,(a)对齐情况下会缓存命中率比(b)偏差因为第二个内核访问的数据已经在缓存中第一个内核的执行(a)对齐的情况。

结果显示,尽管OpenCL强调便携性,增加亲和力OpenCL支持可能在某些情况下提供显著的性能改进。因此,我们认为耦合逻辑结构与物理线程的线程(核心CPU)需要在OpenCL,尤其是对CPU。任务的粒度可以工作组;换句话说,程序员可以指定一个特定的核心工作组将被执行。此功能将有助于改善OpenCL应用程序的性能。例如,来自不同内核的数据可以不共享内存请求如果程序员为具体分配核心工作组在考虑不同内核的数据共享。的数据可以通过私有缓存共享内核。

4所示。相关工作

多研究研究如何优化OpenCL gpu性能。GPGPU社区提供了张力腿平台(29日)作为优化的一般指导方针GPGPU应用程序因为GPGPU通常配备了大量的处理元素。自从OpenCL CUDA(一样的背景9),大多数OpenCL应用程序写入更好地利用张力腿平台。广泛使用的入住率度量表明张力腿平台的程度。然而,这个计划不能应用于应用程序的CPU因为即使TLP很大,物理张力腿平台上可用的CPU被CPU核的数量有限,所以在CPU上下文切换的开销远远高于在gpu的开销可以忽略不计。

一些出版物指OpenCL内核cpu上的性能。一些关注算法和一些参考性能差异通过比较它与GPU实现和OpenMP实现cpu上(16,30.,31日]。然而,我们所知,我们的工作是第一个提供了一个广泛的概述,结合应用程序的架构知识提供一般指导方针理解OpenCL性能在多核cpu上。

阿里等人比较OpenCL OpenMP和英特尔的TBB在不同的平台上(30.]。他们主要讨论缩放效果和编译器优化。但它遗漏了为什么本文中列出的优化给提到的性能优势和缺乏定量评价。我们也评估OpenCL的性能和OpenMP对于一个给定的应用程序。然而,我们的工作考虑各个方面,可以改变应用程序的性能,并提供定量评估来帮助程序员估算各方面性能的影响。

Seo等人讨论OpenCL NAS并行基准的性能影响和给一个很好的概述如何优化基准首先获得一个想法的数据传输和调度开销,然后想出方法来避免他们(31日]。他们还展示了如何重写一个好的OpenCL代码,给定一个OpenMP代码。斯垂顿等人描述一个方法来实现细粒度SPMD-thread程序编译器执行多核平台上(16]。细粒度的编程模型,他们从CUDA开始,说它将适用于OpenCL。他们关注的性能改善基线。我们的工作是更通用的和广泛的相比,这些先前的研究,还包括一些重要的点,不解决这些论文。

引用之一是非常有用的理解从英特尔OpenCL的性能行为是文档(32]。它广泛地展示了一些一般性的指导方针,遵循以获得更好的性能OpenCL英特尔处理器上的应用程序。然而,它并不讨论性能改进和国家也不可以实现多少好处。

5。结论

我们评估OpenCL应用程序的性能在现代多核CPU架构。理解建筑资源利用率的性能而言有利于程序员。在本文中,我们评估各个方面,包括API开销,线程调度,独立,数据传输,数据局部性,compiler-supported向量化。我们验证OpenCL应用程序通过比较它们的独特的特点与传统的并行编程模型如OpenMP。评价的重要发现如下。(1)OpenCL API开销不是cpu上可以忽略不计(部分3所示。2)。(2)分配更多的工作/ workitem因此减少workitems帮助cpu性能(部分3.3.1)。(3)大型独立帮助cpu性能(部分3.3.2)。(4)大型工作组大小有助于更好的性能在cpu(部分3.3.3)。(5)相比cpu上,映射api执行上级明确数据传输api。内存分配旗帜不改变性能(部分3所示。4)。(6)编程模型可以有可能影响compiler-supported向量化。条件可以复杂的矢量化的代码(部分3.5.1)。(7)增加亲和力OpenCL支持有助于性能在某些情况下(部分3.5.2)。

我们的评估显示,考虑CPU架构的特点,OpenCL的应用可以进一步优化CPU,和程序员需要考虑这些见解的便携性能。

利益冲突

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

确认

作者要感谢金王、Sudhakar Yalamanchili, Inchoon唷,佐治亚理工学院HPArch成员,匿名评论者的建议和反馈。我们感激地承认NSF的支持事业奖1139083和三星。