MIC优化
9
第 9章 MIC性能优化
本章将介绍 MIC 平台应用程序的性能优化方法，包括并行度、内存管理、数据传输、存 储器访问、向量化、负载均衡和扩展性等优化手段，并且通过大量的示例代码展示每种优化手 段的使用方法。
本章目标
通过本章的学习，你可以：
l了解 MIC 性能优化策略。
l了解如何通过并行度、内存管理、向量化、cache优化等手段优化 MIC 程序的性能。
l了解如何使用 nocopy、异步、scif技术优化 CPU与 MIC 的通信。
l了解如何对 MIC 的负载均衡和线程扩展性进行优化。
9.1 MIC 性能优化策略
MIC 程序性能优化包含 CPU 与 MIC 端的通信优化以及 MIC 内核计算、访存优化。MIC 程序的优化是一个循环反复的过程，可以用性能优化循环图表示，MIC 性能优化的基本循环 过程如图 9-1 所示，包括以下几个步骤：
（1）获取程序的性能数据作为基准。
（2）分析性能数据，通过 VTune等工具找出性能瓶颈。
（3）根据瓶颈进行分析，并找到相应的优化手段。
（4）实施优化手段，对代码做相应的修改。
（5）测试结果是否正确，如果正确并且性能获得提升就完成一次优化的循环。
（6）进入下次优化循环。
MIC 高性能计算编程指南
186
9
C h a p t e r
图9-1 性能优化循环
MIC 性能优化循环中常见问题： （1）选择合适的测试用例，即 workload 。合适的 workload 需要满足几个要求：可测试性、 可重现性、稳定性和代表性。合适的测试用例是可以“用例测试”的，即：不能测试时间过短 或过长； 测试是可以重复进行的， 可以复现的； 每次复现的性能结果是稳定的， 不能偏移太大， 而且测试用例具有广泛的代表性，不是和大部分用例大相径庭的。例如，如果所选的测试用例 执行的是算法的第 1条路径， 而其他大部分用例执行的是算法的第 2 条路径， 则该测试用例的 选择是不合适的。
（2）如何得到性能指标。性能指标有很多种，最简单最常用的是计算程序执行时间。在 MIC 程序中，我们可以通过时间函数获取内核计算执行时间，以及数据传递时间，也可以借 助 VTune 等工具获取 MIC 内核中每个线程时间等。
（3）分析问题的时候主要考虑热点和关键路径。非主要矛盾可以忽略，把精力放在关键 问题的优化上。考虑基准性能指标是什么，最优的性能指标能到多少，可能的潜力有多少。回 答这些问题需要考虑制约 MIC 性能的关键因素，如 GFLOPS 值、CPI 、程序的并行度、数据 的局部性、带宽的压力、向量化程度、IO 是否为主要的瓶颈等。这些关键数据可以借助 VTune 测试得到。
（4）在实施优化过程中，对代码的修改还需兼顾代码的质量。要保证代码的可移植性、 可读性、可维护性、可靠性。采用的方法可以是修改编译选项、各种数学库，手动修改热点代 码等。
（5）测试用例要完备，要尽可能覆盖所有的情况，并且只有在程序保证正确的前提下测 得的 MIC 性能指标才有效。
（6）最后，决定是否进行新一轮的循环优化。考虑的因素有：现有性能是否已经逼近极
MIC 性能优化 第 9 章
187
9 Chapter
限？是否实现了 MIC 硬件的利用率最大化？
MIC 程序的性能优化主要包括系统级和内核级：系统级优化包括节点之间、CPU与 MIC 之间的负载均衡优化；MIC 内存空间优化；计算与 IO 并行优化；IO 与 IO 并行优化；数据传 递优化；网络性能优化；硬盘性能优化等。内核级优化包括并行度优化；负载均衡优化；进程 /线程的同步优化；线程扩展性优化；向量化优化；cache优化；数据对齐优化；库函数的选择 等。下面小节我们将详细介绍这些优化手段。
9.2 MIC 优化方法
MIC 作为众核协处理器，具有众多的核，只有程序高度并行化，才能充分发挥出 MIC 的 性能。并行化使得程序可以在 MIC 的各个核上同时执行，能够显著提高程序的性能。并行程 序常用的两种并行方法为数据并行和任务并行。
数据并行以数据分割为依据，扩展性良好，随着数据规模的增大，线程数也可以增多，并 且每个线程的任务量不减少。可以说，数据并行是天然的并行方式。需要注意的是分割不均匀 会导致严重的负载不均衡。MIC 是共享内存并行系统众核处理器，因此，数据并行是MIC 上 并行优化的很好的选择。
任务并行是指以多任务为基础，多个任务并行操作，不同任务处理的方式可能不一样，处 理的数据本身没有太多的依赖关系。因此，在 CPU与 MIC 协同计算时，可以采用任务并行的 方式，让 CPU和 MIC 执行不同的任务，以充分发挥各自的优势。
并行程序执行时，由于存在多个任务/线程之间的相互通信和影响使得应用程序行为变得 复杂。因此，需要优化MIC 上的各个方面才能达到性能的最优，MIC 上的并行程序的优化主 要包括并行度、内存空间、数据通信/传递、Cache 访问、向量化、负载均衡、线程扩展性等 方面的优化。下面对这些优化方法进行展开介绍。
9.2.1 并行度优化
在计算机体系结构中，并行度是指指令并行执行的最大条数。在设计并行程序时，我 们可以简单地把并行度认为是在多核/众核处理器上能同时执行的线程数/进程数。对于同一个 程序，并行度设计方法的不同将会严重影响到程序的性能。MIC 上的并行度优化主要涉及并 行线程/进程的数目、并行层级、并行粒度等方面。
9.2.1.1 并行度
MIC 卡包含众多的物理核，同时每个核上可以开启 4 个线程，因此，程序员只有设计足 够多的线程/进程才可以把所有的核利用起来。例如一块 60 个核的 MIC 卡上，我们最多可以 开启 240 个线程，最佳线程数一般是每个核设置 3 个或 4 个线程，图 9-2 展示的是某一实际高
MIC 高性能计算编程指南
188
9
C h a p t e r
性能应用程序在 60 个核的 MIC 卡上设置不同线程数的性能扩展性结果图，从该图可以看出， 只有让 MIC 卡上的所有核都充分利用起来才能发挥 MIC 的最大性能。当然，也不是在 MIC 卡上设置的线程数越多越好，线程数太多的话，线程开销比较大，我们只需要让设置的线程数 可以保证程序并发度和 MIC 核的高利用率即可。
图9-2 某一高性能应用程序在MIC 上的性能扩展性
9.2.1.2 并行粒度
并行程序是否选择了合适的层级实现并行， 是性能优化中需要关心的重要问题。 根据并行 程序尽可能使用粗粒度的并行原则， 尽可能在最上层并行化代码。 在外层上并行除了带来易编 程的好处之外，还可以带来好的性能：增加粒度，减少线程调度和销毁的次数，也就是减少线 程本身的开销所占的比例，尤其对于 MIC 平台要开启上百个线程，减少线程的开启对性能影 响更为重要；同时，隐藏了底层的线程交互，减少了不必要的同步带来的损耗。
下面通过简单的例子说明并行层级， 例如程序中有两层循环， 并且每层循环都没有数据依
赖性，即两层循环都可以并行，根据并行程序尽可能使用粗粒度的并行原则我们可以采用在 i 层循环并行的方式。
1 #pragma omp parallel for num_threads(THREAD_NUM)
2 for (i=0? i 3 { 4 for (j=0? j 5 { 6 … 7 } 8 } MIC 性能优化 第 9 章 189 9 Chapter 当然， 并不是所有的应用程序都是在外层循环并行效果最佳， 外层循环的并行可能会导致 线程之间访问的数据跨度比较大，可能会引起 Cache miss，这种情况下采取内层循环的并行效 果更佳，同时为了减少线程的开销，我们可以在外层 for 之前开启多线程，在内层 for 进行任 务分发，如上面的代码采用下面的并行方式。 1 #pragma omp parallel num_threads(THREAD_NUM) 2 for (i=0? i 3 { 4 #pragma omp for 5 for (j=0? j 6 { 7 … 8 } 9 } 在实际的应用程序中也可能出现某一层循环无法达到 MIC 的并行度要求，针对这种情 况，我们可以采取多层循环合并的方式。例如上面的代码中 M=20，N=30，无论我们并行哪 层 for 都无法达到 MIC 的并行度要求，我们可以合并两层 for，合并之后的循环次数为 600 次，显然可以满足 MIC 平台上的要求。当然，我们也可以采用嵌套并行的方式满足 MIC 的 并行度要求。 合并循环： 1 #pragma omp parallel for num_threads(THREAD_NUM) 2 for (k=0? k 3 { 4 i = k/M? 5 j = k%M? 6 … 7 } 嵌套并行： 1 omp_set_nested(true)? //允许嵌套并行 2 #pragma omp parallel for num_threads(THREAD_NUM1) 3 for (i=0? i 4 { 5 #pragma omp parallel for num_threads(THREAD_NUM2) 6 for (j=0? j 7 { 8 … 9 } 10 } MIC 高性能计算编程指南 190 9 C h a p t e r 9.2.2 内存管理优化 MIC 作为协处理器，卡上的内存空间相比主机端十分有限，且不可扩展。因此，如何充 分利用有限的内存空间，用尽量少的内存，完成尽量多的计算，成为 MIC 优化的重点。另外， 由于卡上内存空间有限，导致一些程序在 MIC 移植可行性上，也会面临很大的挑战。 高度并行，也会对内存空间带来新的问题。以前的串行程序，每次迭代所用到的临时内存 空间都是可以重复利用的。发展为并行程序以后，需要为每个线程分配独立的临时空间，但由 于计算单元核数不多，因此并发线程不多，内存问题并不突出。可是在 MIC 上，由于并发线 程达到了 200 个甚至更多，导致所需内存空间急剧增加。假设一个线程私有变量占用空间为 1MB ，在 CPU 端，总内存占用可能仅有 4MB 或 8MB 而已，但在 MIC 上，则会根据线程数 瞬间扩展到 200MB 以上！因此对内存空间占用的优化成了一个不得不面对的挑战。 另外，MIC 每个核心的时钟频率比 CPU 要低，因此在创建内存空间的效率上，设备端也 与主机端有一定的差距。所以对 MIC 卡上的内存空间优化，不仅仅是对空间上的优化，同时 还要考虑时间上的优化。 正如上文所述，MIC 内存空间的优化分为两个方面：内存使用容量和申请次数。内存使 用容量是指程序对卡上内存空间的使用量， 即占用内存的大小， 通常关注的是最大情况下的大 小。 申请次数是指程序运行过程中， 静态或动态开辟卡上内存空间的次数。 优化关注使用量时， 有时会与其他注重性能的优化方法产生冲突。但是，优化空间占用通常事关移植的可行性，因 此通常会重视对空间占用的优化并因此牺牲一定的性能。 9.2.2.1 内存占用 卡上内存占用量出现瓶颈，通常有两种情况。其一是，任务本身需要占用大量内存，由于 MIC 端内存容量较小，所以移植中出现困难。其二是，每线程占用临时空间较大，当移植到 MIC 上时，由于线程数较多，导致内存不足。减少内存占用的方法包括： 1．任务分块 解决内存占用最根本的办法，就是任务分块。如果能把大任务划分为小任务，把一次需要 占用的内存空间降下来，则可以解决几乎全部的内存空间不足的问题。但是面对不同的原因， 通常会采用不同的任务划分方式，下面分情况进行讨论。 第一种情况，任务本身需要占用大量内存，则通常需要将任务本身分块，一次只处理任务 的一个子集。 第二种情况，每个线程需要较多的临时空间，通常可以选择降低并发度，减少线程数，自 然可以降低总体内存占用。由于 MIC 卡上可以并发的线程较多，降低一定的并发度，仍然有 可能保持性能在可以接受的范围内。 如果任务整体完全无法分割，或 MIC 卡上内存无法满足即使是最小分块的需求，那么 MIC 性能优化 第 9 章 191 9 Chapter 只能认为该程序（通常认为是该算例）无法移植到 MIC 上运行了。即使是主机端，内存也 是有限的，如果任务实在太大，在 CPU 端也会无法运行。严格说起来，这种情况也并非完 全不能移植，只是需要将任务分成极小的部分，虽然每个部分都仅仅是一个片段，但可以不 断地传输－计算－传输－计算……重复这一过程，将整个任务完成。但是，在绝大多数情况 下，这种方式所产生的性能都不足以抵消急剧膨胀的开发成本。列此方法，仅供一些极端情 况下的参考。 我们显然可以发现，任务分块无论如何也会增加传输次数，并且很有可能减少并发度，这 会降低程序的性能。但是正如前文所述，如果不进行分块，程序很有可能无法在 MIC 上运行， 因此牺牲一定性能，使程序具有可行性，也是可以接受的。至于其中的“度” ，则需要程序员 根据具体情况进行把握。 2．临时空间复用 有些程序中用到的一些临时空间，是可以合并或者节省的。例如：程序前半部分用到数组 a，大小为 100MB，后半部分用到数组b，大小为 150MB，使用数组 b 时，数组 a已不再使用。 那么可以只开辟 150MB 大小的空间，前面用作数组 a，后面用作数组 b。虽然减少这些临时空 间可能会对代码的可读性造成一些不利影响， 但是如果这部分对性能影响比较大， 则需要权衡 可维护性和性能之间的比重，尽量寻找其中的平衡点。 还有一种情况，如代码所示： 1 for(i=0?i 2 { 3 c[i]=a[i]+N? 4 d[i]=b[i]+N? 5 } a 和 b 为中间变量， c 和 d 为所求数组。由于 a、c 和 b、d 之间没有依赖关系，因此完全 可以将循环分为两个，一个计算 c，一个计算 d。在计算 c 结束，且计算 d 尚未开始时，即可 将数组 a的空间释放，以节省内存占用。 3．改变算法 从程序算法来看， 通常粗粒度并行， 即并行层次较高的循环， 会占用相对较多的内存资源， 而细粒度并行占用的内存资源较少。例如：一个计算N 个 A*B 的矩阵乘应用（互相无依赖）。 如果并行计算 N 个矩阵乘法，则内存占用会比较大。但是如果每次仅并行 A*B，N 个乘法间 是串行执行的，这样内存占用就会小很多。与任务分块法不同的是，这里需要改变并行的算法 （以前一个线程的功能是串行计算一个矩阵乘法，现在一个线程的功能是计算矩阵的一个元 素）和粒度，而任务分块法通常只需要改变任务大小和传输的方式。 从细节代码的算法来看， 曾经有一道经典的面试题—— “如何在不使用临时变量的情况下， 交换两个整型变量的值” 。虽然这种方法在MIC 上未必适用，但也给我们提供了一个思路，即 可以试着找寻别的算法，以节省内存空间，完成相同的功能。 MIC 高性能计算编程指南 192 9 C h a p t e r 除了上述几种方法以外，在编程史的“上古时期” ，那时硬件资源比较匮乏，程序员们会 竭尽所能，减少哪怕 1 个字节的占用。之后随着硬件的发展，这些技巧被逐渐遗忘，或是不再 使用。但是，现在协处理器的出现，面临着与当年同样的问题。因此，在一些比较古老的书籍 和文献中，也许会有仍然适用于现在的技巧和方法，读者也不妨一试。 9.2.2.2 申请次数 对内存申请次数的优化不一定会减少内存占用， 但是如果注意运用一些技巧， 则可以使性 能有所提高。 对于内存空间方面的性能优化来说，最关键的一点是：把开辟空间的操作放到循环外面。 无论是简单的循环，还是循环中调用的子函数，都有可能根据需要开辟自己私有的空间，尤其 是使用 malloc 等函数开辟的大块内存空间。由于 MIC 的时钟频率等原因，开辟空间的操作比 主机端要慢。因此，如果在循环内开辟较大的内存空间的话，每次开辟时都会耽误一些时间， 而循环次数一多，时间累积起来就会很可观了。虽然这是我们平时不注意的小时间，但是在一 些情况下也会造成很大的性能损失。 例如我们在某软件移植优化的过程中， 就遇到了这种情况。 该代码片段非常简单：一个循环，循环体内容为调用计算函数，计算函数内部首先是开辟内存 空间，然后计算。我们的任务就是把这个循环移植到 MIC 上。移植本身很简单，但是测试时 发现计算函数仅仅占总体运行时间的一半！最开始怀疑是 offload 传输的问题，但自己写测试 用例测试时并没有重现问题。 后来测试才偶然发现是计算函数内部开辟空间的问题。 由于每次 开辟的时间比较长，又根据循环开辟了多次，导致运行时间急剧增加。将内存空间全部移到外 面一次开辟，程序运行时间即大幅缩短。在转移内存开辟的过程中，我们使用在主机端声明指 针，在 offload 时使用 nocopy 开辟空间的方式，进一步节省了不必要的传输时间。在开辟空间 的大小上， 我们采用线程数乘以单次循环需要的内存大小的容量， 运行时各线程根据自己的线 程号查找自己“私有”的内存地址。 当需要多次调用 offload 函数， 进行一系列操作时， 如果不同 offload 函数中有公用的数组， 也可以使用 nocopy 等方式一次申请，多次使用。一方面减少了数据传输时间，另一方面也避 免了多次申请空间的开销。这种优化方式详见下一节“数据传输优化” 。 9.2.3 数据传输优化 数据传输是一个通信的过程。 数据传输对并行计算的性能有很大的影响。 对于单机节点来 说，频繁进行的发送/接收操作将大大降低并行程序的执行性能。对于集群而言，巨大的通信 开销对并行效率的影响是致命的。处理器之间的通信是并行程序执行时重要的时间开销来源。 作为并行计算额外开销的主要组成部分，降低通信成本可以有效地缩短并行程序的执行时间。 所以我们在做并行编程的时候要尽可能地降低 I/O 传输操作的开销。 一般对数据传输会用到的优化方法有：nocopy ，offload 异步，SCIF 模型，4K 倍数等。下 MIC 性能优化 第 9 章 193 9 Chapter 面我们将逐一介绍数据传输优化方法。 9.2.3.1 nocopy CPU与 MIC 之间通过 PCI-E 通信，PCI-E 的速度较慢，因此，我们需要尽量减少 CPU与 MIC 之间的数据通信，通过 nocopy技术可以有效地减少 CPU与 MIC 之间的通信次数。 offload 中的 in、out 语句默认为每次 offload 开始时申请空间，结束时释放空间，然而在 很多应用程序中，数据或空间是可以重复利用的，并不需要每次 offload 时都申请空间、释放 空间。nocopy 主要应用在多次调用 offload 语句时，可以减少 CPU 与 MIC 的通信次数，即主 要应用在有迭代调用 MIC 的程序中。 没有采用 nocopy的 MIC 程序可以表示成下面的伪代码： 1 … 2 p_c = …? //p_c在每次迭代中值不改变 3 for(i=0? i 4 { 5 p_in = …? //每次迭代计算时，p_in的值变化 6 #pragma offload target(mic)\ 7 in(p_in:length(…))\ 8 in(p_c: length(…))\ 9 out(p_out: length(…)) 10 { 11 kernel(p_in, p_c, p_out)? 12 } 13 } 14 … = p_out? //CPU端在所有迭代完成之后才用到p_out的值 上面的代码中，每次迭代 p_in、p_c、p_out 都会在 MIC 端申请空间、进行 CPU 与 MIC 的数据传递、最后释放空间，计算过程如图 9-3所示。而实际运行中，p_c的值只需要传递给 MIC 端一次即可，p_out 的值只在迭代完成之后回传给 CPU 端即可，显然上面的代码有很多 不必要的 CPU 与 MIC 通信。针对这种情况，我们可以采用 nocopy 的模式减少 CPU 与 MIC 的通信次数。 nocopy 技术往往与 alloc_if()、free_if()联用，nocopy 的使用方法在 MIC 编程章节中进行 了详细的介绍。针对上面的 MIC 程序采用 nocopy之后的伪代码如下： 1 … 2 p_c = …? //p_c在每次迭代中值不改变 3 #pragma offload target(mic)\ 4 in(p_c: length(…) alloc_if(1) free_if(0))\ 5 nocopy(p_in:length(…)alloc_if(1) free_if(0))\ 6 nocopy(p_out: length(…) alloc_if(1) free_if(0)) 7 { MIC 高性能计算编程指南 194 9 C h a p t e r p_c = ... p_in = ... MIC 端给p_c 分配空间，并由CPU 给 MIC 端传递p_c 的值 MIC 端给p_out 分配空间，并由CPU 给MIC 端传递p_out 的值 kernel(p_in,p_c,p_out)? MIC 把结果p_out 回传给CPU 端 MIC 端给p_in 分配空间，并由CPU 给MIC 端传递p_in 的值 迭代完成 否 是 MIC …= p_out 释放MIC 端已申请的空间 p_in 、p_c 、p_out 图9-3 MIC 程序执行过程 8 } //申请空间，并且不释放；传递p_c 的值 9 for(i=0? i 11 p_in = …? //每次迭代计算时，p_in 的值变化 12 #pragma offload target(mic)\ 13 in(p_in:length(…) alloc_if(1) free_if(0) )\ 14 nocopy(p_c)\ 15 nocopy(p_out) //每次迭代传递p_in 的值，p_c 和p_out 采用nocopy 16 { 17 kernel(p_in, p_c, p_out)? MIC 性能优化 第 9 章 195 9 Chapter 18 } 19 } 20 #pragma offload target(mic)\ 21 nocopy (p_c: length(…) alloc_if(0) free_if(1))\ 22 nocopy(p_in:length(…)alloc_if(0) free_if(1))\ 23 out(p_out: length(…) alloc_if(0) free_if(1)) 24 { 25 } //回传p_out值到CPU端，并释放MIC端申请的空间 26 … = p_out? //CPU端在所有迭代完成之后才用到p_out的值 采用 nocopy 的 MIC 程序的执行过程如图 9-4 所示，从该图可以看出，采用 nocopy 技术 减少了 CPU 与 MIC 的通信次数。nocopy 技术对程序的性能提升起着至关重要的作用，主要 应用在需要迭代计算的应用程序中。 图9-4 采用nocopy的MIC程序执行过程 MIC 高性能计算编程指南 196 9 C h a p t e r 9.2.3.2 offload 异步传输 对于 MIC 的 offload 模式来说，异步有两种含义，其一是数据传输与计算的异步，即MIC 卡与主机端的数据传输与 MIC 卡上计算的异步，其二是计算与计算的异步，即 MIC 卡与 CPU 计算的异步。 1．数据传输与计算异步 数据传输层面的异步，通常用于需要多次调用 MIC 函数，且相邻调用之间没有依赖关系 的情况。出现这种情况有可能是原始算法就是这么写的，也有可能是因为某种原因（例如卡上 内存空间不足）对数据进行了分块处理，使得程序需要多次调用没有数据依赖的 MIC 函数。 数据传输层面的异步，能够带来的明显好处是：以流水线方式执行传输和计算，可以将绝大部 分数据传输时间隐藏，如图 9-5 所示。 图9-5 异步执行示意图 因此，在需要对算法进行分块处理的时候，或者数据传输时间比较明显，已经严重影响到 程序执行效率的时候，可以使用数据传输异步的方式，对程序进行优化。 数据传输优化主要使用两个 offload 语句的变种：offload_transfer 和 offload_wait 。 offload_transfer 的作用是传输数据，并在数据传输完成时发送信号。其参数是与传统的 offload 语句完全一致的，区别在于 offload_transfer 语句后面没有代码段，仅有 offload 一条语句。 offload_wait 的作用是暂停程序的执行，直到接收到 offload_transfer 发送的信号。举例如下： //C/C++ 1 #pragma offload_attribute(push, target(mic)) 2 int count = 25000000? 3 int iter = 10? 4 float *in1, *out1? 5 float *in2, *out2? 6 #pragma offload_attribute(pop) 7 8 void do_async_in() 9 { 10 int i? 11 #pragma offload_transfer target(mic:0) \ 12 in(in1 : length(count) alloc_if(0) free_if(0) ) signal(in1) MIC 性能优化 第 9 章 197 9 Chapter 13 14 for (i=0? i 15 { 16 if (i%2 == 0) { 17 18 #pragma offload_transfer target(mic:0) if(i!=iter-1)\ 19 in(in2 : length(count) alloc_if(0) free_if(0) ) signal(in2) 20 21 #pragma offload target(mic:0) nocopy(in1) \ 22 wait(in1) out(out1 : length(count) alloc_if(0) free_if(0) ) 23 24 compute(in1, out1)? 25 26 } else { 27 28 #pragma offload_transfer target(mic:0) if(i!=iter-1)\ 29 in(in1 : length(count) alloc_if(0) free_if(0) ) signal(in1) 30 31 #pragma offload target(mic:0) nocopy(in2) \ 32 wait(in2) out(out2 : length(count) alloc_if(0) free_if(0) ) 33 34 compute(in2, out2)? 35 } 36 } 37 } !Fortran 1 integer, parameter :: iter = 10 2 integer,parameter :: count = 25000 3 !dir$ options /offload_attribute_target=mic 4 real(4), allocatable :: in1(:), in2(:), out1(:), out2(:) 5 integer :: sin1, sin2 6 !dec$ end options 7 8 contains 9 10 subroutine do_async_in() 11 integer i 12 13 !dir$ offload_transfer target(mic:0) & 14 in(in1 : alloc_if(.false.) free_if(.false.) ) signal(sin1) 15 MIC 高性能计算编程指南 198 9 C h a p t e r 16 do i = 0, (iter -1) 17 if (mod(i,2) == 0) then 1819 !dir$ offload_transfer target(mic:0) if(i/=iter-1) & 20 in(in2 : alloc_if(.false.) free_if(.false.) ) signal(sin2) 2122 !dir$ offload target(mic:0) nocopy(in1) wait(sin1) & 23 out(out1 : length(count) alloc_if(.false.) free_if(.false.) ) 2425 call compute(in1, out1)? 2627 else 2829 !dir$ offload_transfer target(mic:0) if(i/=iter-1) & 30 in(in1 : alloc_if(.false.) free_if(.false.) ) signal(sin1) 3132 !dir$ offload target(mic:0) nocopy(in2) wait(sin2) & 33 out(out2 : length(count) alloc_if(.false.) free_if(.false.) ) 3435 call compute(in2, out2)? 3637 endif 38 enddo 3940 end subroutine do_async_in 本例仅展示了 do_async_in 一个函数， 但这个函数已足以展示异步传输的应用。 在本例中， 我们开辟了两个数组空间，并等待使用第1 个空间 in1、out1 计算时，同时在传输第 2 个空间 in2、out2 的数据，当第 2 个空间数据传输完成，且第 1 个空间的计算完成时，开始第 2 个空 间的计算，并且传输第 1个空间的数据，循环往复，直到全部计算完成。即在循环变量为奇数 时，传输第 2 个空间数据，并等待第 1个空间的数据计算完成，循环变量为偶数时反之。由于 二者交替进行， 因此最终的执行时间大致等于传输和计算二者中较长的时间， 而非同步时的二 者时间之和。 在进入循环以前，首先使用 offload_transfer 将第 1 个空间填充（C ：11～12 行，Fortran ： 13～14 行），这样在进入循环体时（C ：14 行，Fortran ：16 行），可以首先启动填充第 2 个空 间的工作（C ：18～19 行，Fortran ：19～20 行），然后等待第 1 个空间传输完毕（C ：21～22 行，Fortran ：22～23 行），再对第一个空间的内容进行计算（C ：24 行，Fortran ：25 行）。由 于 MIC 卡只有一个，因此对计算来说（也对循环来说）是串行执行的，而只有传输和计算之 间才是异步的。代码中需要注意的是，由于本段代码并非完整程序，因而在第一次传输时，并 没有开辟空间（alloc_if(0)），这是因为在该函数外，已经在 MIC 卡上开辟空间了（代码中未 MIC 性能优化 第 9 章 199 9 Chapter 体现）。这种方式可以避免在循环中重复开辟空间，节省了时间。 另外一个值得注意的地方是 signal 的参数。在 C/C++中，参数是需要传递的一个数组的指 针，在 Fortran中，则是一个 integer 类型的变量。这里需要注意两点，其中一点是在不同语言 中参数的区别，另一点是 signal 中能且只能使用一个参数，即使在 C/C++中，无论同时传输多 少数组，只需要其中任意 1 个作为 signal 的参数即可。 在使用数据传输异步优化时， 最主要的是需要注意对卡上内存空间的占用。 由于卡上要同 时拥有传输和计算两部分的空间，即内存占用成倍增长，因此如果分块过大，很有可能造成卡 上内存空间不够用，导致程序运行失败。因此，即使是传输时间较长的应用，也不一定适宜使 用过多的数组空间进行异步。 2．计算与计算异步 计算异步多用于 CPU与 MIC 协同计算的情况，让 CPU与 MIC 各分担一部分任务，以充 分利用节点内的计算资源，达到加速程序运行的目的。通常意义上的协同计算，是指 MIC 计 算函数与 CPU 计算函数分属不同线程，甚至不同进程当中，二者之间是相对较严格的并行方 式。这种情况也可以算是计算异步的一种方式，但在本节所述的计算异步，则更多地是指另一 种方式，即二者处于同一线程，启动函数是串行的，但函数的执行是并行的。 传统 offload方式的 MIC 程序中，当调用 offload语句，即 MIC 函数后，MIC 函数即接管 控制权，CPU 线程处于等待状态，直到 MIC 函数执行完成返回时，才将控制权交回给 CPU 线程，此时 CPU 线程才能继续向下执行。在计算异步的 MIC 程序中，调用 offload 语句后， 代码段在 MIC 卡上启动以后，驱动即刻将控制权交还给 CPU 线程，CPU 线程继续下面的工 作，当 MIC 函数执行完毕返回时，会给 CPU线程发送一个信号。可以很显然地发现，异步模 式下，CPU线程和 MIC 函数在一部分时间内是并行执行的，这样自然节省了运行时间。 这种计算异步方式相对前文所述的协同计算， 在代码上比较简单， 也不需要启动多个线程， 虽然并发性可能略有不如，但灵活性却有很大提高。MIC 端的计算函数不需要很大改变，而 CPU 端的函数也未必与 MIC 计算函数完成同样的任务（也许规模不同），CPU 端的函数也可 以做一些数据准备，或是其他的工作，以便为下次 MIC 计算做好准备。这是狭义的协同计算 方式所无法做到的。 计算异步方式用到的除了传统的 offload 语句以外，也会使用数据传输异步中介绍的 offload_wait 语句等待计算完成的信号。以下是一个简单示例： 1 int counter? 2 float *in1? 3 counter = 10000? 4 __attributes__((target(mic))) mic_compute? 5 while(counter>0) 6 { 7 #pragma offload target(mic:0) signal(in1) 8 { MIC 高性能计算编程指南 200 9 C h a p t e r 9 mic_compute()? 10 } 11 cpu_compute() //此时本函数与上面的MIC 函数并行执行 12 #pragma offload_wait target(mic:0) wait(in) 13 counter--? 14 } 1 integer signal_var 2 integer counter 3 counter = 10000 4 !DIR$ A TTRIBUTES OFFLOAD:MIC :: mic_compute 5 do while (counter .gt. 0) 6 !DIR$ OFFLOAD TARGET(MIC:0) SIGNAL(signal_var) 7 call mic_compute() 8 call cpu_compute()！此时本函数是与上面的MIC 函数并行执行的 9 !DIR$ OFFLOAD_WAIT TARGET(MIC:0) W AIT (signal_var) 10 counter = counter -1 11 end do 12 end 如上例所示，在调用 mic_compute 后，程序马上返回，将控制权交还给 CPU ，CPU 可以 继续执行 cpu_compute ，并在 offload_wait 处等待，直到 MIC 函数计算完成后，offload_wait 才返回，继续执行下面的代码。 9.2.3.3 SCIF 传输优化 我们在前面章节已经知道CPU 与MIC 之间是通过PCI-E 总线进行数据通信， 而通过PCI-E 的数据传输速度较慢，因此，我们需要尽量减少 CPU 与 MIC 之间的数据通信，但有时候 CPU 与 MIC 之间进行频繁的小数据通信或者在内核计算完后， MIC 与 MIC 或 MIC 与 CPU 之间需 要小数据通信等情况是避免不了的，遇到这些情况该如何解决呢？可以考虑 SCIF 这种数据传 输方式，这种数据传输方式可以有效地提高在频繁地进行小数据通信的情况下数据传输的性 能。下面我们通过一个简单的例子来看一下这种通信是如何进行的。 在 host 端运行的程序（作为请求端）： 1 #include 2 #include 3 #include 4 #include 5 #include 6 #include 7 #include 8 #include MIC 性能优化 第 9 章 201 9 Chapter 9 #include 10 int main(int argc, char **argv) 11 { 12 scif_epd_t epd? 13 int err = 0? 14 int req_pn = 11? 15 int con_pn? 16 struct scif_portID portID? 17 int i, num_loops = 0, total_loop = 10? 18 char *senddata? 19 char *recvdata? 20 int msg_size? 21 int block? 22 int node? 23 struct timeval tv? 24 struct timeval tv1? 25 if (argc != 4) { 26 printf("Usage ./scif_connect_send_recv* 27 exit(1)? 28 } 29 msg_size = atoi(argv[1])? 30 block = atoi(argv[2])? 31 node = atoi(argv[3])? 32 portID.node = node? 33 portID.port = 10? 34 printf("Open the scif driver\n")? 35 if ((epd = scif_open()) < 0) {//在请求端建立一个新的endpoint。 36 printf("scif_open failed with error %d\n", (int)epd)? 37 exit(1)? 38 } 39 printf("scif_bind to port 11\n")? 40 if ((con_pn = scif_bind(epd, req_pn))< 0) {//请求端：绑定该endpoint到该endpoint所在的端口上。 41 printf("scif_bind failed with error %d\n", con_pn)? 42 exit(2)? 43 } 44 printf("req_pn=%d",req_pn)? 45 retry: 46 if ((scif_connect(epd, &portID)) != 0) {//此处使用retry去试着和监听端建立通信， 如果监听端 已经做好连接准备，此处的连接则会直接连通，否则会进行几次retry连接操作。如果已经连 通则可以进行下面的数据交换等操作。 47 if (ECONNREFUSED == errno) { 48 printf("scif_connect failed with error %d retrying\n", errno)? MIC 高性能计算编程指南 202 9 C h a p t e r 49 goto retry? 50 } 51 printf("scif_connect failed with error %d\n", errno)? 52 exit(3)? 53 } 54 printf("scif_connect success\n")? 55 printf("node=%d,port=%d\n",portID.node,portID.port)? 56 while (num_loops < total_loop) { 57 senddata = (char *)malloc(msg_size)? 58 if (!senddata) { 59 perror("malloc failed")? 60 err = ENOMEM? 61 } 62 memset(senddata, 0x25, msg_size)? 63 err = 0? 64 gettimeofday(&tv,0)? 65 while ((err = scif_send(epd, senddata, msg_size, block))<= 0) {//此处两个endpoints 之间 已经建立连通，请求端向监听端发送数据。 66 if (err < 0) { 67 printf("scif_send failed with err %d\n", errno)? 68 fflush(stdout)? 69 free(senddata)? 70 goto close? 71 } 72 } 73 gettimeofday(&tv1,0)? 74 printf("err=%d",err)? 75 printf(" total = %f\n", (https://www.sodocs.net/doc/8f6323467.html,_https://www.sodocs.net/doc/8f6323467.html,_sec)*1e6+(https://www.sodocs.net/doc/8f6323467.html,_https://www.sodocs.net/doc/8f6323467.html,_usec))? 76 err = 0? 77 free(senddata)? 78 num_loops++? 79 } 80 close: 81 printf("Close the scif driver\n")? 82 close(epd)? 83 if (!err) { 84 printf("Test success\n")? 85 } 86 else 87 printf("Test failed\n")? 88 return (err)? 89 } MIC 性能优化 第 9 章 203 9 Chapter 在 MIC 端运行的程序（作为监听端） ： 1 #include 2 #include 3 #include 4 #include 5 #include 6 #include 7 #include 8 #include 9 int main(int argc, char **argv) 10 { 11 int epd? 12 int newepd? 13 int err = 0? 14 int req_pn = 10? 15 int con_pn? 16 int backlog = 16? 17 struct scif_portID portID? 18 int i, num_loops = 0, total_loop = 10? 19 char *senddata? 20 char *recvdata? 21 int msg_size? 22 int block? 23 if (argc != 3) { 24 printf("Usage ./scif_accept_send_recv* 25 exit(1)? 26 } 27 msg_size = atoi(argv[1])? 28 block = atoi(argv[2])? 29 portID.node = 2? 30 portID.port = 11? 31 if ((epd = scif_open())< 0) {//监听端建立新的endpoint。 32 printf("scif_open failed with error %d\n", errno)? 33 exit(1)? 34 } 35 printf("scif_bind to port 10\n")? 36 if ((con_pn = scif_bind(epd, req_pn))< 0) {//监听端：绑定已建立的endpoint到该endpoint所 在的端口上。 37 printf("scif_bind failed with error %d\n", errno)? 38 exit(2)? 39 } 40 printf("scif_listen with backlog of 16\n")? MIC 高性能计算编程指南 204 9 C h a p t e r 41 if ((scif_listen(epd, backlog))< 0) {//监听端已经建立好了监听的准备，并且设置了可以在监 听端最大的等待连接请求。 42 printf("scif_listen failed with error %d\n", errno)? 43 exit(3)? 44 } 45 printf("scif_accept in syncronous mode\n")? 46 if (((scif_accept(epd, &portID, &newepd, SCIF_ACCEPT_SYNC))< 0) && (errno != EAGAIN)){ //监听端接受请求端的连接请求，并且此处设置了接受方式，是同步还是异步的方式。 47 printf("scif_accept failed with errno %d\n", errno)? 48 exit(4)? 49 } 50 printf("scif_accept complete\n")? 51 printf("node=%d,port=%d",portID.node,portID.port)? 52 while (num_loops < total_loop) { 53 recvdata = (char *)malloc(msg_size)? 54 if (!recvdata) { 55 free(senddata)? 56 perror("malloc failed")? 57 err = ENOMEM? 58 } 59 memset(recvdata, 0x00, msg_size)? 60 err = 0? 61 while ((err = scif_recv(newepd, recvdata, msg_size, block))<= 0) {//接受请求端发送的数据。 62 if (err < 0) { 63 printf("scif_recv failed with err %d\n", errno)? 64 fflush(stdout)? 65 free(senddata)? 66 free(recvdata)? 67 goto close? 68 } 69 } 70 printf("err=%d",err)? 71 err = 0? 72 free(recvdata)? 73 num_loops++? 74 } 75 close: 76 printf("Connection is complete\n")? 77 fflush(stdout)? 78 scif_close(newepd)? 79 fflush(stdout)? 80 if (!err) {