Get best of the week

GPU计算-原因，时间和方式。加上一些测试

大家早就知道，在视频卡上，您不仅可以玩玩具，还可以执行与游戏无关的事情，例如，训练神经网络，记住加密货币或进行科学计算。它是怎么发生的，您可以在这里阅读，但是我想谈谈为什么GPU对普通程序员（与GameDev不相关）可能很有趣，如何在不花费大量时间的情况下在GPU上进行开发，决定是否朝这个方向看，然后“ 弄清楚”您可以获得什么利润。 



这篇文章是根据我在HighLoad ++中的演示而写的。它主要讨论了NVIDIA提供的技术。我无意宣传任何产品，我仅以它们为例，并且可以肯定的是，在竞争厂商中也可以找到类似的产品。

为什么要依靠GPU？


可以根据不同的标准比较两个处理器，可能最受欢迎的是内核的频率和数量，高速缓存的大小等，但是最后，我们对一个处理器每单位时间可以执行多少次操作，这是什么操作感兴趣，但这是一个单独的问题一个常见的度量标准是每秒触发器的浮点操作数。当我们想比较暖和软，以及我们的GPU和CPU时，此指标非常有用。

下图显示了处理器和视频卡的相同触发器随时间的增长。


（数据是从开源收集的，没有2019-20年的数据，因为不是那里的一切都那么漂亮，但GPU仍然是赢家）

好吧，这很诱人，不是吗？我们将所有计算从CPU转移到GPU，并获得八倍的最佳性能！

但是，当然，并非一切都那么简单。您不能只将所有内容都转移到GPU上，为什么，我们将进一步讨论。

GPU架构及其与CPU的比较


我带来了许多有关CPU的体系结构和基本元素的熟悉图片：


CPU核心有

什么特别之处？一个核心和一堆辅助块。

现在让我们看一下GPU架构：


GPU核心

显卡具有很多处理核心，通常为数千个，但它们被组合为多个块；对于NVIDIA显卡，通常每个为32个，并且具有共同的要素，包括和寄存器。 GPU核心和逻辑元素的体系结构比CPU上的体系结构简单得多，也就是说，没有预取器，早午餐预测器等。

好吧，这些是CPU和GPU架构差异的关键点，实际上，它们施加了限制，或者相反，为我们可以在GPU上有效读取内容开辟了可能性。

我没有提到更重要的一点，通常情况下，视频卡和处理器之间不会相互“翻腾”，也不会将数据写入视频卡并读回结果-这些操作是独立的，可能会成为系统中的“瓶颈”，即抽水时间与尺寸的关系图数据将在本文后面给出。

GPU限制和功能


此体系结构对可执行算法有什么限制：

  • 如果我们在GPU上进行计算，则不能只选择一个内核，而是将分配整个内核块（对于NVIDIA为32）。
  • 所有内核执行相同的指令，但使用不同的数据（我们将在后面讨论），这种计算称为单指令多数据或SIMD（尽管NVIDIA对其进行了改进）。 
  • 由于逻辑块和通用寄存器的相对简单的设置，GPU确实不喜欢分支，而且实际上不喜欢算法中的复杂逻辑。

它打开了什么机会：

  • 实际上，这些相同SIMD计算的加速。最简单的例子是矩阵的元素加法，让我们对其进行分析。

将经典算法简化为SIMD表示


转型


我们有两个数组A和B，我们想将数组B中的一个元素添加到数组A中的每个元素。下面是C中的示例，尽管我希望对那些不讲这种语言的人来说很清楚：

void func(float *A, float *B, size)
{ 
   for (int i = 0; i < size; i++) 
   { 
       A[i] += B[i]
   } 
}

循环和线性运行时中元素的经典环回。

现在，让我们看看这样的代码在GPU中的外观：

void func(float *A, float *B, size) 
{ 
   int i = threadIdx.x; 
   if (i < size) 
      A[i] += B[i] 
}

在这里已经很有趣了，出现了threadIdx变量，我们似乎没有在任何地方声明它。是的，它的系统为我们提供了服务。想象一下，在前面的示例中，数组由三个元素组成，您想在三个并行线程中运行它。为此，您需要添加另一个参数-索引或流号。这是视频卡为我们所做的事情，尽管它将索引作为静态变量传递并可以同时处理多个尺寸-x，y，z。

另一个细微差别，如果您要立即启动大量并行流，则必须将这些流划分为多个块（视频卡的体系结构功能）。最大块大小取决于视频卡，并且我们需要对其进行计算的元素的索引需要如下获得：

int i = blockIdx.x * blockDim.x + threadIdx.x; // blockIdx –  , blockDim –  , threadIdx –

结果就是我们所拥有的：许多并行运行的线程，它们执行相同的代码，但具有不同的索引，并因此具有不同的数据，即 相同的SIMD。

这是最简单的示例，但是如果要使用GPU，则需要将任务转换为相同的形式。不幸的是，这并非总是可能的，并且在某些情况下可能成为博士学位论文的主题，但是，仍然可以将经典算法引入这种形式。

聚合


现在，让我们看一下聚合转换为SIMD表示形式的样子：
 

我们有n个元素的数组。在第一阶段，我们启动n / 2个线程，每个线程添加两个元素，即 在一次迭代中，我们将数组中一半的元素加在一起。然后在循环中，对新创建的数组重复相同的操作，直到聚合最后两个元素。如您所见，数组的大小越小，我们可以启动的并行线程就越少，即 在GPU上，聚合足够大的数组是有意义的。这样的算法可用于计算元素总和（顺便说一句，不要忘记正在使用的数据类型可能溢出），搜索最大值，最小值或仅搜索。

排序


但是排序已经看起来更加复杂了。

GPU上两种最受欢迎​​的排序算法是：

  • 双音排序
  • 基数排序

但是基数排序仍然被更频繁地使用，并且可以在某些库中找到可用于生产的实现。我不会详细分析这些算法的工作原理；有兴趣的人可以在https://www.codeproject.com/Articles/543451/Parallel-Radix-Sort-on-the-GPU-using-Cplusplus-中找到基数排序的描述。AMP和https://stackoverflow.com/a/26229897

但想法是，即使是排序这样的非线性算法，也可以简化为SIMD视图。

现在，在研究可以从GPU获得的实数之前，让我们弄清楚如何为这种技术奇迹编程。

从哪儿开始


可在GPU下进行编程的最常见的两种技术：

  • Opencl的
  • 库达

OpenCL是大多数视频卡制造商支持的标准，包括 在移动设备上，用OpenCL编写的代码也可以在CPU上运行。

您可以从C / C ++使用OpenCL，也有其他语言的活页夹。

对于OpenCL，我最喜欢《OpenCL in Action》一书。它还描述了GPU上的不同算法，包括 Bitonic排序和Radix排序。

CUDA是NVIDIA的专有技术和SDK。您可以用C / C ++编写或使用对其他语言的绑定。

比较OpenCL和CUDA有点不正确，因为 一个是标准，另一个是整个SDK。尽管如此，尽管该技术是专有的，但免费且仅在NVIDIA卡上有效，但许多人还是选择CUDA来开发视频卡。有几个原因：

  • API
  • , GPU, (host)
  • , ..

特性包括CUDA带有其自己的编译器，该编译器还可以编译标准C / C ++代码。

我遇到的最全面的CUDA书籍是《专业CUDA C编程》，尽管它已经有些过时了，但是它讨论了NVIDIA卡编程的许多技术细节。

但是，如果我不想花几个月的时间阅读这些书，编写自己的视频卡程序，进行测试和调试，然后发现这不适合我，该怎么办？ 

正如我所说的，有大量的库隐藏了GPU下开发的复杂性：XGBoost，cuBLAS，TensorFlow，PyTorch等，我们将考虑推力库，因为它比上面的其他库没有那么专业，但是同时它实现了基本算法，例如排序，搜索，聚合，并且很有可能适用于您的任务。

Thrust是一个C ++库，旨在用基于GPU的算法“替换”标准STL算法。例如，使用视频卡上的该库对数字数组进行排序将如下所示：

thrust::host_vector<DataType> h_vec(size); //    
std::generate(h_vec.begin(), h_vec.end(), rand); //   
thrust::device_vector<DataType> d_vec = h_vec; //         
thrust::sort(d_vec.begin(), d_vec.end()); //    
thrust::copy(d_vec.begin(), d_vec.end(), h_vec.begin()); //   ,

（不要忘记该示例必须由NVIDIA的编译器编译）

如您所见，推力::排序与STL中的类似算法非常相似。该库隐藏了许多困难，尤其是子程序（更确切地说是内核）的开发，该子程序将在视频卡上执行，但同时又缺乏灵活性。例如，如果我们要对几GB的数据进行排序，则将一条数据发送到卡上以开始排序是合乎逻辑的，而在进行排序时，将更多的数据发送到卡上。这种方法称为延迟隐藏，可以更有效地利用服务器映射资源，但是，不幸的是，当我们使用高级库时，这种机会仍然被隐藏。但是对于原型设计和性能评估来说，它们是相同的，尤其是在推力方面，您可以测量数据传输所带来的开销。

我写了一个小基准 使用此库，该库在GPU上运行几种具有不同数据量的流行算法，让我们看看结果如何。

GPU算法结果


为了测试GPU，我在AWS上使用Tesla k80显卡作为实例，它不是迄今为止功能最强大的服务器卡（功能最强大的Tesla v100），但价格却最便宜，并且具有以下功能：

  • 4992 CUDA内核
  • 24 GB内存
  • 480 Gb / s-内存带宽 

为了测试CPU，我以2.30GHz @ Intel Xeon处理器CPU E5-2686 v4作为实例

转型





您可以看到，GPU和CPU上的转换执行时间以毫秒为单位，您可以看到，在GPU和CPU上，数组元素的常规转换时间大致相同。又为什么呢 因为将数据发送到卡并返回的开销占用了整个性能提升（我们将分别讨论开销），并且卡上的计算相对较少。同样，不要忘记处理器也支持SIMD指令，并且在简单情况下的编译器可以有效地使用它们。 

现在，让我们看看如何在GPU上高效地进行聚合。

聚合



GPU和CPU上的聚合执行时间（以ms为单位）

在聚合示例中，我们已经看到随着数据量的增加，性能有了显着提高。还值得关注的事实是我们将大量数据泵入卡的内存，并且仅取回一个汇总值，即 将数据从卡传输到RAM的开销最小。

让我们继续进行最有趣的示例-排序。

排序



以毫秒为单位对GPU和CPU进行排序的时间

尽管我们将整个数据阵列都发送到了视频卡，反之亦然，但在GPU上进行800 MB数据的排序大约比在处理器上快25倍。

数据传输开销


从转换示例可以看出，即使在并行性很好的任务中，GPU是否有效仍然并不总是很明显。这样做的原因是将数据从计算机的RAM传输到视频卡的内存的开销很大（顺便说一句，在游戏机中，内存是在CPU和GPU之间共享的，因此无需传输数据）。视频卡的特征之一是内存带宽或内存带宽，它确定了卡的理论带宽。对于Tesla k80，它是480 GB /秒，对于Tesla v100，它已经是900 GB /秒。此外，PCI Express版本以及将数据传输到卡的方式的实现也会影响吞吐量，例如，可以通过多个并行流来完成。

让我们看看在亚马逊云中为Tesla k80显卡获得的实际结果：


将数据传输到GPU，将数据分类并传输回RAM的时间（以毫秒为

单位）-将数据​​传输到

GPU 视频卡执行-在视频卡上分类

DtoH-将数据从视频卡复制到RAM

首先要注意的是，从视频卡读取数据的速度比把它们写下来。

第二个-使用视频卡时，您可以获得350微秒的延迟，这对于某些低延迟的应用程序可能已经足够了。

下图显示了更多数据的开销：


将数据传输到GPU，将数据分类并传输回RAM的时间（以毫秒为单位）

服务器使用


最常见的问题是游戏视频卡与服务器视频卡有何不同？根据特性，它们非常相似，但价格差异很大。


服务器（NVIDIA）和游戏卡之间的主要区别是：

  • 制造商的保修（游戏卡不适用于服务器）
  • 消费类显卡可能存在的虚拟化问题
  • 服务器卡上的纠错机制的可用性
  • 并行线程（不是CUDA内核）的数量或对Hyper-Q的支持，这使您可以从CPU上的多个线程处理卡，例如，从一个线程将数据上传到一个卡并从另一个线程开始计算

这些也许是我发现的主要重要区别。

多线程


在我们弄清楚如何在视频卡上运行最简单的算法以及可以预期的结果之后，下一个逻辑问题是视频卡在处理多个并行请求时的行为。作为回答，我有两个GPU图形计算以及一个具有4和32核的处理器：


在GPU和CPU上以1000 x 60 in ms

进行矩阵计算的时间，该图以1000 x 60个元素的矩阵进行计算。计算从几个程序流开始，为每个CPU流为GPU创建一个单独的流（使用了Hyper-Q）。 

如您所见，处理器可以很好地应对这种负载，而每个GPU的一个请求的延迟会随着并行请求数量的增加而显着增加。


在10,000 x 60单位为毫秒的GPU和CPU上执行数学计算的时间。

在第二张图上，相同的计算，但是矩阵长10倍，并且在这种负载下GPU的性能要好得多。这些图非常具有指示性，我们可以得出结论：负载下的行为取决于负载本身的性质。处理器还可以在一定程度上高效地处理矩阵计算。对于视频卡，其特征是，对于较小的计算负载，性能大约呈线性下降。随着负载的增加和并行线程数的增加，视频卡的处理效果会更好。 

很难假设GPU在各种情况下的行为，但是如您所见，在某些情况下，服务器卡可以非常有效地处理来自多个并行流的请求。

如果您仍然决定在项目中使用GPU，我们将讨论您可能会有的其他问题。

资源限制


正如我们已经说过的，视频卡的两个主要资源是计算核心和内存。

例如，我们有多个使用视频卡的进程或容器，我们希望能够在它们之间共享视频卡。不幸的是，对此没有简单的API。NVIDIA提供了vGPU技术，但是我在支持的列表中没有找到Tesla k80卡，据我从描述中可以了解到，该技术更侧重于虚拟显示而非计算。也许AMD提供了更合适的产品。

因此，如果计划在项目中使用GPU，则应依赖于该应用程序将专门使用视频卡的事实，或者将以编程方式控制分配的内存量和用于计算的内核数。

容器和GPU


如果确定了资源限制，那么将出现以下逻辑问题：如果服务器中有多个视频卡，该怎么办？

同样，您可以在应用程序级别决定它将使用哪个GPU。

另一个更方便的方法是Docker容器。您可以使用常规容器，但是NVIDIA提供了NGC容器，其中包含各种软件，库和驱动程序的优化版本。对于一个容器，您可以限制使用的GPU数量及其对容器的可见性。容器使用的间接费用约为3％。

在集群中工作


另一个问题，如果要在同一服务器或群集中的多个GPU上执行一项任务，该怎么办？

如果选择类似于推力或较低级别解决方案的库，则必须手动解决任务。例如，用于机器学习或神经网络的高级框架通常支持开箱即用使用多个卡的功能。

另外，我想指出，例如，NVIDIA提供了一个用于卡之间直接数据交换的接口-NVLINK，它比PCI Express快得多。还有其他技术可以直接从其他PCI Express设备访问卡的内存-GPUDirect RDMA，包括。和网络。

推荐建议


如果您正在考虑在项目中使用GPU，则在以下情况下GPU最适合您：

  • 您的任务可以简化为SIMD视图
  • 可以在计算（缓存）之前在地图上加载大多数数据
  • 挑战涉及密集计算

您还应该提前提出问题：

  • 多少个并行查询 
  • 您期望什么延迟
  • 您需要一张卡来装载吗？您需要一台带有多张卡的服务器还是一台GPU服务器集群 

就是这样，我希望这些材料对您有用，并帮助您做出正确的决定！

参考文献


github上的基准测试和结果-https: //github.com/tishden/gpu_benchmark/tree/master/cuda

除主题外，还记录了报告“ GPU数据库-使用的体系结构，性能和前景”

NVIDIA NGC Containers网络研讨会-http ：//bit.ly/2UmVIVt或http://bit.ly/2x4vJKF

More articles:


All Articles