CUDA C编程（三十九）将C程序移植到CUDA C的案例研究

最新推荐文章于 2023-08-23 16:30:24 发布

伴君

最新推荐文章于 2023-08-23 16:30:24 发布

阅读量1.6k

点赞数

分类专栏： # CUDA 文章标签： cuda

本文链接：https://blog.csdn.net/AAAA202012/article/details/123560780

版权

CUDA 专栏收录该内容

39 篇文章 172 订阅

订阅专栏

这一部分内容主要是讲一个传统的应用程序通过APOD全过程最终转化到优化的CUDA应用程序的例子。

从Wrox.com上可以下载到这个传统应用程序的代码，名为crypt.c。crypt实现了IDEA加密和解密。crypt应用程序由3个主要部分组成：

应用程序设置在main中。设置包括读取输入、预分配输出空间和读取密钥，密钥是一个二进制串，消息的发送者和接收者都必须知道该密钥以成功加密或解密信息。
加密和解密输入信息的关键是使用generateEncryptKey和generateDecryptKey产生共享密钥。
实际上输入数据的加密和解密是encrypt_decrypt中的8字节块完成的。

crypt实际输入的是文件，该文件既不是加密文件也不是解密文件。此外，crypt需要一个密钥文件，用于存储加密或解密输入数据的64位密钥。Wrox.com提供了产生示例输入数据(generate_data.c)和密钥(generate_userkey.c)的文件。下图通过crypt应用程序概述了高级数据流。花一点事件来熟悉crypt.c的执行情况。
在这里插入图片描述
评估 crypt

用于评估主机应用程序性能的工具有很多。在下面的实例中会使用gprof,因为它使用范围非常广、免费并且提供低开销的性能分析。

在评估crypt之前，需要生成样本密钥和1GB的样本数据，代码如下所示：

$ ./generate_userkey key
$ ./generate_data data 1073741824

使用提供的Makefile文件进行编译crypt后，通过gprof运行它，生成以下性能信息：
在这里插入图片描述
这里只考虑最左边和最右边的列。右边是crypt应用程序中不同函数的名称。左边是该函数的执行时间占应用程序总执行时间的百分比。可以预料，总执行时间的87.03%花费encrypt_decrypt函数中，该函数实现主要的加密和解密逻辑。从这些信息可以总结出，如果crypt被并行化，那么并行化策略应该应用于encypt_decrypt中。

发现性能热点仅仅是评估步骤的一半。还必须要分析这些热点是否适合并行化，也就是说，是否存在一种方法可以将一些循环或热点内部或周围的代码部分进行并行化，这个并行化策略在GPU上是否存在潜在的提速作用。

对crypt来说，这个步骤很简单：encrypt_decrypt在一个循环中执行大量的计算。在每次迭代中，这个函数会处理块列表中的一个数据块。因为在这个循环中读和写都在不同地块中进行，所以这个循环可以跨输入列表实现并行。但是也会带来一些问题。下一次迭代要处理的数据块通过当前元素被指出，因此，下一次迭代(i+1)和当前迭代(i)之间是存在依赖性的。决定如何解除这种依赖性是crypt并行阶段中很重要的部分。

并行 crypt

通过两个步骤可以将crypt并行运算。第一，需要改变传统应用程序的控制流和数据结构，使它更适合并行化。第二，需要将计算内核转化到CUDA C中，并插入必要的CUDA API调用(比如，cudaMalloc或cudaMemcpy)来建立主机和设备之间联系。

crypt并行化之后的结果可以在Wrox.com上的crypt.parallelized.cu文件中找到。有几个转化需要进一步解释，以助于理解最终的产品是如何为并行化做准备的。

首先，用于存储输入和输出数据的数据结构由链表改为数组。这有几个好处。首先，它消除了评估阶段发现的第i+1次迭代和第i次迭代之间的依赖关系。属于块i的数据现在可以用偏移索引在数组中被检索到，而不用遍历整个链表元素。另外，数组转化到GPU中也更简单了。因为链表依赖于指针，所以，将链表由主机地址空间转化到设备地址空间，意味着也要将这些指针指向正确设备的正确元素上。数组可以直接用cudaMemcpy来拷贝。

除了改变crypt的主要数据结构，核心的计算内核也要提取到一个独立地函数doCrypt中，以使并行化更明显。doCrypt使用全局指针来进行输入、输出和处理数据块。使用这个函数作为抽象，调用的内核可以跨数据块实现并行。

并行化的下一步是在crypt的合适位置处插入CUDA API调用。这一过程的变化可以分为两部分：内核实现和内核管理。

改变crypt的内核实现很简单。第一，将关键字__device__加到doCrypt中，表明他应该在GPU中被执行。第二，将encrypt_decrypt声明为__global__函数，它包含的循环基于线程ID被转化为相邻设备线程上执行的每一个数据块。第三，添加了一个名为encrypt_decrypt_driver的新函数，用于启用encrypt_decrypt内核，该内核的执行配置由输入数据块的数量来决定。

encrypt_decrypt_driver内核也为移植内核执行内存管理，包括：

给任何输入输出数据分配需要的所有内存；
将所有应用程序的数据传输到设备中；
释放所有已分配的设备内存；
释放所有已分配的设备内存；

有了这些简单的改变后，crypt应用程序可以在CUDA中执行大部分计算了。下一部分将关注这个并行化过程的结果并使用配置文件驱动优化提升性能。

优化 crypt

在配置文件驱动优化的第一阶段，使用CUDA性能分析工具深入了解应用程序的性能特点，有了这些信息，就可以确定在哪个地方进行优化了。一旦做出了改变，就可以重新分析应用程序以帮助确定下一步需要做的工作，不断迭代改进性能。

在开始的时候，使用nvvp的无导向模式生成一个总和性能分析，包括总体的提升建议。在nvvp中设置crypt，就像在"创建新对话"弹出窗口中指定可执行文件的名字和输入输出文件的位置一样简单，如下图所示：

通过nvvp运行crypt，收集性能分析数据后，时间轴如下图所示：
在这里插入图片描述
值得注意的是，内核占用了执行时间的很大一部分，调用cudaMemcpy也占用了应用程序运行时间的重要的一部分。此外，由于使用同步副本，所以没有重叠通信和计算。我们能注意到，从主机到设备的传输(HtoD)发生在内核启动前，而且主要分为了两次cudaMemcpy调用：一次用于plain数据，另一次用于crypt数据。而crypt数据仅仅是一个输出数据，所以在启动内核之前将其状态转移到设备上没有任何意义。在这种情况下，通信实际上可以被删除。

下图所示的为分析视图的时间轴标签，其中的建议来自于nvvp性能统计。这些建议表明低复制带宽和低计算利用率是限制性能的明显因素。一些nvvp建议提示我们可以将重叠计算和通信作为提高计算和内存性能的一种方式。
在这里插入图片描述
有了这些了解后，下一步便是实施一个重叠计划。在crypt这个例子中，可以将输入分为更小的块，并在同一时间将一块传送到独立流中的设备上来完成这一步骤。然后，对于每个块进行异步cudaMemcpyAsync调用和内核启动。因为这些操作将被放置在不同地CUDA流中，所以CUDA运行时可以以任何顺序执行他们，实现计算和通信之间的重叠，以及更好的利用率。从Wrox.com下载的crypt.overlap.cu文件中，包括加入这些改变的crypt新版本。为了方便，这里列出了一段核心代码：

CHECK(cudaEventRecord(start, streams[0]));
    CHECK(cudaMemcpyAsync(dKey, key, KEY_LENGTH * sizeof(int),
                    cudaMemcpyHostToDevice, streams[0]));
    CHECK(cudaStreamSynchronize(streams[0]));

    for (b = 0; b < nBlocks; b++)
    {
        int blockOffset = b * BLOCK_SIZE_IN_CHUNKS * CHUNK_SIZE;
        int localChunks = BLOCK_SIZE_IN_CHUNKS;

        if (b * BLOCK_SIZE_IN_CHUNKS + localChunks > nChunks)
        {
            localChunks = nChunks - b * BLOCK_SIZE_IN_CHUNKS;
        }

        CHECK(cudaMemcpyAsync(dPlain + blockOffset, plain + blockOffset,
                        localChunks * CHUNK_SIZE * sizeof(signed char),
                        cudaMemcpyHostToDevice, streams[b]));

        encrypt_decrypt<<<nThreadBlocks, nThreadsPerBlock, 0, streams[b]>>>(
            dPlain + blockOffset, dCrypt + blockOffset, dKey, localChunks);
        CHECK(cudaMemcpyAsync(crypt + blockOffset, dCrypt + blockOffset,
                        localChunks * CHUNK_SIZE * sizeof(signed char),
                        cudaMemcpyDeviceToHost, streams[b]));
        CHECK(cudaEventRecord(finishes[b], streams[b]));
    }

注意下面的循环：

for(b = 0; b < nBlocks; b++){
   //块大小和偏移量的值计算如下：
   int blockOffset = b * BLOCK_SIZE_IN_CHUNKS * CHUNK_SIZE;
   int localChunks = BLOCK_SIZE_IN_CHUNKS;
   if(b * BLOCK_SIZE_IN_CHUNKS + localChuncks > nChunks){
      localChunks = nChunks - b * BLOCK_SIZE_IN_CHUNKS;
   }
}

这种优化实现了在cudaMemcpyAsync和encrypt_decrypt之间基于流的重叠，用于由blockoffset和localChunks定义的块。这些改变所带来的性能提升，如下表所示：
在这里插入图片描述
现在可以在代码的其他部分进行重新分析和重新定位优化的工作了。这个过程与之前相同，但需使用新的基于流的可执行文件。时间轴视图和时间轴分析的最新结果表明，所有第一次运行时产生的问题已经被消除或减少。程序执行时间轴清楚地显示了通信和计算的重叠，而不是展示了大规模的阻塞cudaMemcpy调用。
在这里插入图片描述

下一步的决定没有之前的结果明显。有几个突出的问题可能是下一阶段关注的重点。首先，时间轴分析显示了低内存吞吐量的警告。这是由于为每个块进行很多小的内存复制而不是一个大的复制而产生的。然而，导致这种改变的重叠变换很显著地提升了性能，所以低内存吞吐量地代价是可以接受的。

多处理器分析视图表明了寄存器的压力也可能是一个问题。然而，这可以通过修改内核代码得到改变，所以在早期的优化阶段进行寄存器的使用优化，可能是无用功。
在这里插入图片描述
时间轴分析视图表明SM的利用率仍然很低。这意味着SM很可能会花费很多时间，要么没有符合条件的任何线程块来调度，要么等待I/O完成。内核内存面板还警告全局内存存储效率低。从这两个指标可以总结出，这个应用程序的全局内存操作可能会限制它的性能。下一步更加依赖如何使用全局内存中应用程序的特定知识。

那么，当前存储和全局内存中的对象是什么？目前，输入text,输出crypt、加密/解密key的存储和访问都在全局内存中进行。因为线程块是跨线程ID进行处理的，所以每个线程都读和写text和crypt中相邻的8字节块。虽然4字节是最优的，但是因为访问是合并和对齐的，这仍然可以使缓存和带宽得到合理有效的利用。

key的访问模式是一个非常不同的故事。每一个线程在同一时间读取key的相同位置。这将产生更低的全局带宽利用率，因为线程的完整线程束从全局内存中读取相同的4字节时将被阻塞。因为key，text和crypt都在GPU多处理器中共享全局一级和二级缓存，如果通过text或crypt中写入或读取来删除缓存中的读操作，那么可能要进行多次读操作。基于这样的分析和由nvvp报告的指标，看起来下一步较好的选择是优化key的使用。

优化key的使用，方法之一就是改变存储key的内存。哪种CUDA内存类型支持只读数据结构且对于广播单一元素到所有线程是最优的？看起来这应该是常量内存！在常量内存中放置key，有可能提高全局内存带宽和全局内存缓存效率。

在crypt的新版本中key存储在常量内存中，这个新版本可以从Wrox.com上的crypt.constant.cu文件中获得。所进行的更改包括：

加入了一个__constant__ dkey变量：constant int dkey[KEY_LENGTH];
修改doCrypt内核以引用新的dkey变量；
调用cudaMemcpyToSymbolAsync，将key中的内容传输到设备中：CALL_CUDA(cudaMemcpyToSymbolAsync(dkey,key,KEY_LENGTH * sizeof(int),0,cudaMemcpyHostToDevice,streams[0]));

这一变化获得的性能改进在下表中进行了总结。与原始版本的CUDA实现相比，性能几乎增加了一倍。迭代时间和重新定位增加了更多的优化机会！
在这里插入图片描述
使用nvvp再一次进行性能分析，从时间轴和多处理器分析视图的指标可以看出了改进。

然而，下图所示的内核内存面板仍然报告全局存储带宽的利用率低：

这似乎需要更多的调查。首先，重要的是要了解数值12.5%来自于哪里。例如，考虑doCrypt里每个线程执行的第一次读操作，对plain输入的单字节进行访问：

x1 = (((unsigned int)plain[chunk * CHUNK_SIZE]) & oxff);

因为线程是跨块的，每块是8字节，所以线程束里的线程在每次加载过程中通常访问plain中的每8字节。然而，缓存硬件把这些稀疏的单字节加载调整为双字节，这样有128字节的加载来自全局内存。因此，从性能分析工具的角度来看，来自全局内存的每8字节加载中，仅有1字节真正被使用，所以1/8 =12.5%就是利用率。

然而，这并不是整个故事。作为这128字节加载的结果，以后的每个plain引用可能会命中一级缓存也可能会命中二级缓存，这取决于GPU的架构，所以全局内存引用不是必需的。每个加载到缓存中的字节都会被使用，但可能不是特定加载指令的一部分。因此，这就是nvvp报告的次优资源利用的情况，通过更多的调查得到实际上不是性能的问题，因为数据被缓存了。

然而，时间轴分析仍然显示出一个低计算利用率的警告。回想一下可知，之前的运行中多处理器分析也显示过一个占用警告，这是由于寄存器消耗产生的。这些警告都表明，用于这个内核的线程配置可能不是最优的。因此，每个SM寄存器通过块中的线程被稀疏传播，当值从寄存器溢出后会导致I/O上更多的阻塞，因此计算利用降低。通过分析代码很难验证这一结构。相反，可以使用不同的线程配置进行实验，从而查看是否可以提升性能。带有这些改变的crypt应用程序的新副本可以从Wrox.com上下载的crypt.config.cu文件中得到。这个新版本允许我们使用命令行参数配置每个块的线程数。在一组线程配置上测试新的代码，所产生的结果如下表所示：
在这里插入图片描述
虽然在每块512个线程的原始配置下，执行性能是可接受的，但是每块128个线程可以获得19%的性能提升。因为减少了线程块的大小改进了性能，一个合乎逻辑的结论是，与最佳情况相比，每块512个线程导致给每个线程分配了更少的寄存器。

现在可以重新进行性能分析，看看是否可以确定新的性能问题。请注意，使用每块128个线程重新分析crypt.config，需要在nvvp交互式会话配置中添加命令行参数。下图所示为时间轴分析视图，SM利用率的提升要归功于更好的寄存器分配。
在这里插入图片描述
在这一点上，nvvp将会提供不能解释的性能问题。crypt通过优化阶段已成功被转换。如下表总结的那样，相较于并行阶段的非优化实现，性能提升增加了一倍多。使用配置文件驱动优化，高效利用开发时间。
在这里插入图片描述
目前，已经开发出了一种高性能实现，现在可以进入到下一个也是APOD最后的阶段：部署。

部署 crypt

在主机和设备函数的错误处理方面，crypt已经准备好部署了。然而，还可以提高其能力以适应新的硬件平台，这个平台可以有不同数量的GPU，也可以没有GPU。

多GPU的crypt

云提供商提供了更多的GPU部署，组织机构越来越多地将他们的产品系统移动到云上，在执行环境中已经增加了灵活地支持重大改变的重要性。对crypt来说，这意味着增加了跨所有可用GPU分担工作量的能力，以及在没有GPU被检测到的情况下回到主机执行。

crypt.flexible.cu中的源代码是crypt应用程序灵活实现的例子，程序可以在任何数量的GPU上运行，包括GPU数量为0的情况。crypt.flexible根据有没有GPU来选择主机或设备执行，这由cudaErrorNoDevice错误代码来显示。注意，doCrypt使用__host__ 、__device__函数在两种实现间共享代码，减少了复制代码和维护相同算法的两个副本的开销。为了方便，这里列出了通过多GPU划分工作的核心逻辑：

for(d = 0; d < nDevices; d++){
   CALL_CUDA(cudaSetDevice(d));
   int start = d * chunksPerDevice * CHUNK_SIZE;
   int len = chunksPerDevice * CHUNK_SIZE;
   if(start + len > textLen){
      len = textLen - start;
   }
   encrypt_decrypt_driver(text + start, crypt + start, key, len, nThreadsPerBlock, ctxs + d);
}

CALL_CUDA(cudaEventRecord(finishEvent));

for(d = 0; d < nDevices; d++){
   CALL_CUDA(cudaSetDevice(d));
   CALL_CUDA(cudaDeviceSynchronize());
}

在这段代码中，修改后的encrypy_decrypt是在每个GPU中被调用的，用来一部初始化数据传输和内核启动。然后，一旦每个设备开始工作，主机暂停并等待每个设备完成工作。
混合OpenMP-CUDA Crypt

因为没有可用于执行的GPU，所以前面的crypt.flexibile.cu例子中使用了CPU，如果发现了GPU，则将CPU闲置。虽然在一个系统的GPU上执行所有的计算可能会提高性能，但是这同时也会导致可用的硬件未被充分利用。一些应用程序支持混合并行：CPU和GPU在同一个问题上协同并行。

在一般情况下，有两种类型的混合并行：

数据并行的混合并行：CPU与GPU执行相同的数据并行计算，但跨越的是CPU核心而不是GPU的SM。本质上CPU成为了系统中的另一个设备。在这种情况下，可以使用__host__、__device__函数在两个处理器上执行相同的逻辑。
任务并行的混合并行：CPU与GPU执行不同的计算，该计算更适合于基于主机的体系结构。例如，CPU可以执行具有更复杂的控制流或不规则访问模式的任务。

crypt.openmp.cu包含了一个例子，这个示例中，在单一的应用程序中使用了CPU上的OpenMP并行和GPU上的CUDA并行。OpenMP是一种针对主机的并行编程模型，它使用编译器指令标记并行区域，类似于OpenACC。只有加入了OpenMP的特定代码才是对omp_set_num_threads的调用，用来配置CPU核心使用的数量，并且在主机端的计算函数h_encrypy_decrypt中加入OpenMP编译指示。编译器提示#pragma omp parallel for标记了下述可并行化的循环，并指示OpenMP在多个CPU线程上运行它。

#pragma omp parallel for
for(c = 0; c < nChunks; c++){
   doCrypt(c, plain, crypt, key);
}

crypt.openmp.cu还增加了跨CPU和GPU划分工作量的逻辑，在这里使用了一个新的命令行参数cpu-percent,它指定了在CPU上被加密或是解密的字节百分比，以及转到GPU上的剩余工作量。

CPU和GPU的计算是并行执行的，通过在不同的流中为每个设备排队数据传输和内核执行，然后一旦异步CUDA调用返回控制到主机就启动CPU线程：

CALL_CUDA(cudaEventRecord(startEvent));

for(d = 0; d < nDevices; d++){
   CALL_CUDA(cudaSetDevice(d));
   int start = d * chunksPerDevice * CHUNK_SIZE;
   int len = chunksPerDevice * CHUNK_SIZE;
   if(start + len > textLen){
      len = textLen - start;
   }
   encrypt_decrypt_driver(text + start, crypt + start, key, len, nThreadsPerBlock, ctxs + d);
}
int cpuStart = gpuLen;
h_encrypt_decrypt(text + cpuStart, crypt + cpuStart, key, textLen - cpuStart);
CALL_CUDA(cudaEventRecord(finishEvent));

crypt.openmp.cu必须在OpenMP的支持下进行编译和链接。如果NVIDIA编译器使用的是gcc主机编译器，那么可以使用如下语句：

$ nvcc -Xcompiler -fopenmp -arch=sm_20 crypt.openmp.cu -o crypt.openmp -lgomp

注意，crypt.openmp添加了两个新的命令行参数：使用的CPU内核的数量(ncpus)和在CPU上处理数据的百分比(cpu-percent)。应该使用0.0~1.0之间的数值来指定cpu-percent命令行的选项。

通过使用cpu-percent命令行参数，可以研究随CPU上工作量的增加，性能是如何变化的。在下表中，一系列工作量的分配结果显示，将工作放置在CPU上是不利于crypt的。对于这个特定的应用程序，将产生新的CPU线程的开销和较慢的CPU计算性能，这意味着在CPU上运行任何数量的工作都将导致性能降低。
在这里插入图片描述
其他的应用程序在某种程度上可以同时使用CPU和GPU并产生互补的效果，这比处理器单独运行时实现了更好的性能。例如，用于分类世界排名前500的超级计算机的高性能LINPACK（HPL）标准检查程序在混合执行体系下表现得最好。