异构编程和oneAPI工具包。 英特尔专家即兴演讲回答您的问题



作为“向英特尔专家询问问题”专栏的一部分,我们请英特尔领先的专家康斯坦丁·弗拉基米罗夫(Konstantin Vladimirov)回答了与异构编程, oneAPI工具包以及相关有趣事物有关的问题。 结果超出了我们的所有预期。 康斯坦丁(Konstantin)不会浪费时间,给出了详细而充分的答案,而不必担心会引起争论。 实际上,我们进行了有关各种形式的跨体系结构编程的小型讲座:卸载细微差别,优化,标准等等。
我们将麦克风转让给专家。 好吧,评论是给听众的。

图片 问题Soarex16
从OpenCL到oneAPI的过渡将是多么艰巨,并且可以从中获得什么好处?

答案。 切换到DPC ++可能很棘手,但我认为这是值得的。 有两个主要阶段。

首先,这是您最有可能基于API的异构编程语言(OpenCL,Vulkan计算)的过渡。 在这里,您已经了解了主题领域,这是一个开端,而且困难在于将思维从通过API的直接控制转换为稍微隐式的语言构造。
其次,这是从您的宿主语言开始的过渡。 如果您已经从纯C转移了整个生命,那么输入阈值等于从C切换到C ++的阈值,这是相当高的。

为什么要尝试?

首先,DPC ++对于程序员来说做得很好。 您将很快像噩梦一样忘记对clXXXYYY的所有这些显式调用,以及第六个参数的含义,以及是否忘记了返回码。 许多面向对象的包装器将例程隐藏得更糟,但是通常是以从标准OpenCL API切换到非标准包装器API的代价(我也看到了这些自行车)。 对于DPC ++,您只需编写带有Intel扩展名的标准SYCL(可能很快也会成为标准SYCL)。

其次,DPC ++提供了联合编译功能,也就是说,您可以确定类型,并且在API的尺寸,填充,对齐方式的边界上不会有问题。 您将内核和主机代码编写在一个文件中,这是相同的代码。 使用USM,您还可以更轻松地处理复杂的数据结构。

第三,DPC ++是真正的C ++,也就是说,它允许通用编程。 例如,用于添加两个向量的最简单内核:

auto kern = [A, B, C](cl::sycl::id<1> wiID) { C[wiID] = A[wiID] + B[wiID]; //   A, B  C?  ! }; 

在OpenCL上也是如此:

 _kernel void vector_add(__global int *A, __global int *B, __global int *C) { int i = get_global_id(0); C[i] = A[i] + B[i]; } 

您知道,我被迫指向一个OpenCL类型的int。 如果需要浮点运算,则要么必须编写另一个内核,要么使用预处理器,或者需要外部代码生成。 如果您没有使用C ++的经验,那么掌握C ++的几乎所有功能可能会有点吓人。 但这是重大技术转变的常态。

并且所有好处不仅仅限于此。 我将在以下答案中提及其他内容。

因此,我将在您的位置下载该编译器并进行尝试,因为使用OneAPI包这样做并不难。

图片 问题公正
OpenVINO和oneAPI会以某种方式关联吗?

答案。 现在,OpenVINO发行版是OneAPI发行版的一部分。 学习和使用神经网络是困难的计算任务,这从异构编程中受益匪浅。 我相信,所有OneAPI组件迟早都会使您能够使用所有可用的计算资源:图形加速器和特殊的加速器(如Nervana和FPGA)。 所有这一切都无需离开C ++程序的语言范式和类型系统。

图片 邮件中的问题
我试图了解AI硬件加速器在3年后的外观,请提供帮助。 有一家有趣的公司Graphcore及其IPU-这种设备的效率不比FPGA低,但是编程起来却容易得多-支持TensorFlow和其他框架的Python。 事实证明,如果Graphcore的承诺得以兑现,则机器学习市场将不再需要FPGA。 对于数据科学家而言,Python比C ++更方便。
您是否同意与Python可编程解决方案相比,FPGA不适合机器学习市场? 如果失去了机器学习市场,您还会看到其他哪些广泛的FPGA应用?
在哪些应用程序中,您会看到对异构编程的必然需求,而在这些应用程序中无法使用Python等更方便的工具来满足需求?

答案。 我简要浏览了哪种IPU。 每个人都将卸下另一块铁。 这些人与GPU和专用加速器而不是与FPGA竞争。

在需要改进专用硬件的任务中,它将始终胜过FPGA,例如,在视频卡上渲染视频效果更好,等等。 但是,在世界范围内(包括ML领域),有许多任务尚未发明或发布出来,在这里,FPGA将永远是必不可少的。 例如,由于存在价格问题,并且要便宜,专用硬件必须庞大。

现在假设指定的IPU非常酷。 这不会取消异构编程,相反,如此出色的加速器的存在会刺激它的发展。 它还将为OneAPI和DPC ++带来巨大的优势,因为早晚有人会说“我想从一个程序中同时使用IPU和GPU”。 之所以这么早,是因为异构编程与此有关。 其含义是将适当的任务转移到适当的设备上。 任务可以来自任何地方。 这个设备可以是任何东西,甚至可以是运行程序的同一设备。 例如,如果您卸载了用ISPC编写的内核并最大程度地利用了Xeon的矢量功能,则您可以自己卸载它,并且仍然是一个很大的收获。 这里的主要标准是性能。 好吧,在这个世界上永远不会有太多的生产力。 即使拥有世界上最好的加速器。

至于Python及其便利性……我必须马上承认我不喜欢动态类型的语言:它们速度慢,并且由于正常的错误而不是正常的编译错误,您必须等待两个小时才能进入运行时。 但是我看不到在Python下执行相同的卸载有多么糟糕。 顺便说一句,OneAPI已经包含了适用于Python的英特尔发行版,对于各种评论来说都非常方便。

也就是说,在Python爱好者的梦想世界中,您可以在上面编写一个程序,然后将其卸载到使用OneAPI可以找到的所有加速器中,而不是一堆特定于供应商的库中。 另一件事是,使用这种方法,您会错过端到端的输入,而回到基于API的编程极其不稳定的世界。 DPC ++的开发也许会鼓励社区更加积极地使用更合适的工具,例如C ++。

图片 邮件中的问题
性能与OpenCL。 必须对奢侈品征税-即 间接费用。 有测量值吗?

答案。 在互联网上,您可以找到许多具有各种结果的度量,具体取决于编译器,任务和实现的质量。 作为一项个人研究,我测量了笔记本电脑(集成的Skylake图形)上的简单任务(SGEMM,DGEMM),并且到目前为止,我发现有一些不足(百分比)。 但是在我看来,这是迄今为止所有这些都是beta版本的事实的结果。

从理论上讲,结果应该是加速而不是减速,也就是说,原则上,所有这些奢侈都应该具有负值。 全部与编译器有关。 当您的程序由一个源代码组成并且作为一个程序处理时,编译器将获得绝妙的,令人难以置信的优化机会:布置通用代码,反转循环,重新排列代码段以及编译器在基于API的方法中无法完成的所有其他工作,但是迟早她肯定会使用单一源模型进行学习。

另外,就开发时间而言,DPC ++的成本为负。 一个简单的例子是SYCL访问器,编译器已经在使用它来安排事件和管理异步队列。

 deviceQueue.submit([&](cl::sycl::handler &cgh) { auto A = bufferA.template get_access<sycl_read>(cgh); auto B = bufferB.template get_access<sycl_read>(cgh); auto C = bufferC.template get_access<sycl_write>(cgh); .... deviceQueue.submit([&](cl::sycl::handler &cgh) { auto A = bufferA.template get_access<sycl_read>(cgh); auto B = bufferB.template get_access<sycl_read>(cgh); auto D = bufferD.template get_access<sycl_write>(cgh); 

在这里,编译器看到这两个包都只读取A和B并写入独立的缓冲区C和D,因此,他看到了在有足够大的全局大小的情况下并行发送它们的能力。

当然,一个用书呆子编写的OpenCL程序也可以做到这一点,但是用一个普通的内核所花费的开发时间将是可比的。

图片 邮件中的问题
所有为DPC ++优化OpenCL应用程序的方法都相关吗? 有什么新内容要添加到其中?

答案。 我要说的是,内核编写者正在完成的大多数微妙的手动优化可以并且应该由编译器完成。 例如,以同样的方式,我认为在C ++程序中手动安装嵌入式汇编程序是有害的做法,因为即使提供了战术上的好处,它也会干扰优化并在产品的开发和转让中充当不利因素。 好吧,OpenCL现在也是汇编程序。

至于更详细的答案,我怕这里深渊。 例如,有一个著名的英特尔文档“用于英特尔处理器图形的OpenCL开发人员指南”。 还有一介绍如何尝试,以免将多余的同步放在哪里。

因此,从我的角度来看,这原则上是非人工的任务。 人们对多线程同步的推理能力很差,并且倾向于保守地或错误地或一次或两者地雕刻同步-我这样输入逗号( 但我们已将其修正-编者注 )。

另一方面,在DPC ++中,不要像这样编写带有显式障碍的代码:

  for (t = 0; t < numTiles; t++) { const int tiledRow = TS * t + row; const int tiledCol = TS * t + col; Asub[col][row] = A[globalRow * AY + tiledCol]; Bsub[col][row] = B[tiledRow * BY + globalCol]; // Synchronise to make sure the tile is loaded barrier(CLK_LOCAL_MEM_FENCE); // .... etc .... 

您很可能会在parallel_for_work_group中编写一个显式迭代,其中group.parallel_for_work_item

 cgh.parallel_for_work_group<class mxm_kernel>( cl::sycl::range<2>{BIG_AX / TS, BIG_BY / TS}, cl::sycl::range<2>{TS, TS}, [=](cl::sycl::group<2> group) { // .... etc .... for (int t = 0; t < numTiles; t++) { group.parallel_for_work_item([&](cl::sycl::h_item<2> it) { // .... etc .... Asub[col][row] = A[globalRow][tiledCol]; Bsub[col][row] = B[tiledRow][globalCol]; }); //      ,    

因此,您完全不必手动设置同步,整个部分都可以扔掉。

因此,您可以走遍所有部分。 某些东西将生存,某些东西将离开。 我预见到将出现一个新文档“ DPC ++的优化”,但是时间应该流逝,因为所有真正有效的技术都是后来才开发出来的。

图片 邮件中的问题
OpenCL有一个局限性-您不能在内核中使用“远程数据”,例如,实现“宽过滤器”,该宽过滤器在一次计算中使用来自比OpenCL工作组大的一大组像素的输入数据。 DPC ++在这方面提供了什么?

答案。 好吧,那是不可能的。 当然,我并不是特别编写内核...但是可以肯定的是,您可以按原样使用所有全局内存,只需确保您使用原子操作(或在外部同步分层内核)即可。 您还可以连接System SVM(或者DPC ++中的USM)。

las,所有这些都是非常低效的,我不喜欢所有这些技巧。 此外,它们很难通过编译器进行优化。

因此,如果我们谈论直接有效的解决方案,那么DPC ++当然就没有魔力。 最后,您的程序仍然分为几部分:主机代码和设备代码,所有设备限制都会影响设备代码。 工作组的最大规模是您的硬件具备的真正并行能力。 最重要的是,这些都是解决问题的方法,会对性能产生极大的负面影响。 这就是DPC ++提供执行此操作的机会的原因: device.get_info <sycl :: info :: device :: max_work_group_size>() ,然后决定如何使用生成的数字。

当然,当程序员使用任意长度的循环按自己的意愿工作时,用DPC ++建立模型是很诱人的,而编译器会看下一步该怎么做,但这将是致命的错误,因为它将隐藏常量,甚至增加复杂性的渐近性计算无处不在。 出于另一个原因,Alexandrescu写道“封装复杂性应视为犯罪”,这也适用。

有时修改算法本身会有所帮助。 DPC ++在这里使事情变得更容易,因为结构化的代码更易于重构。 但这真是安慰。

图片 邮件中的问题
DPC ++基于SYCL。 但是,如果您深入了解该怎么办,与OpenCL的后端实现(如果有)有何不同? 例如,异构设备之间的分配机制是否与OpenCL相同?

答案。 如果您不了解,那就是OpenCL。 SYCL的所有优点和强项都是该语言(即前端)的优点和强项。 良好的老式SPIRV从前端传到后端,在那里已经针对特定的视频卡进行了优化(通常是在运行时,即JIT),其优化方式与OpenCL相同。

另一件事是,用于在异构设备之间分配工作的机制只是前端而不是后端,这是因为主机代码决定了发送什么以及在哪里发送。 并且主机代码是从DPC ++获得的。 我已经展示了一个更高的示例,编译器如何基于访问器决定并行化程序包。 这只是冰山一角。

图片 邮件中的问题
图书馆 是的,我们不是在谈论CUDA。 但是我们知道对于CUDA开发人员来说,有非常有用的库可以在GPU上以高性能运行。 OneAPI还包含一些库,但是,例如IPP-在oneAPI / OpenCL中使用图像没有任何归档有用的东西。 会有什么,在这种情况下如何从CUDA切换到oneAPI?

答案。 从CUDA过渡到单一开放标准将很困难,但不可避免。 当然,CUDA现在拥有更成熟的基础架构。 但是其许可的功能是一个阻碍性的缺点,因为越来越多的参与者出现在异构系统的市场上,越来越多来自不同制造商的有趣的卡和加速器。

现有API的多样性使具有经典CPU经验的程序员难以利用这种可能性。 导致OneAPI或类似的东西。 在这里,魔术并不是图形技术上英特尔的突破,而是英特尔向所有人开放DPC ++的大门。 我们甚至都不拥有SYCL标准,它属于Khronos组,所有Intel扩展都是Khronos中的扩展,任何人都可以提交(那里的所有主要参与者都有代表)。 这意味着(图书馆)和社区将出现(已经出现),并且在此方向上将出现大量空缺。

当然,IPP将被重写以适应新的现实。 我与IPP无关,但是使用DPC ++是常识,理智的人正坐在那里。

但更重要的是,现在只是历史上您可以编写自己的库的那一刻,它将超越IPP,然后全世界都将使用它。 因为开放标准总能取胜。

图片 邮件中的问题
如果比较在Nervana和FPGA上训练和推理神经网络算法的发布情况,那么在编程和效率上有什么区别?

答案。 我对FPGA编程的细节一无所知,我编写了编译器。 但是我有一个反问。 我们将如何比较? 在标准的基准测试中,Nervana表现得不像运动员。 但是,如果您有一些有趣的东西,FPGA会解开您的手,而把这些东西放到Nervana上可能会很长,很昂贵,仅此而已。

事实证明,这个问题本身来自“谁比大象或鲸鱼强”系列。 但这不是一个真正的问题。 真正的问题是:如何在一辆推车中同时利用大象和鲸鱼? 好吧,或者至少是分发,例如让大象在陆地上拉它,在海上拉一条鲸鱼。

对于OneAPI,通常在标准C ++中将具有相同的程序。 您可以自己编写它,并来回卸载它。 这将是您非常感兴趣的任务,您可以自己衡量和优化性能。 单一标准和单一接口可用于异构设备,这是将苹果与苹果在此类问题上进行比较的一步。

例如:“从易于编程和提高效率的角度来看,什么对我的任务%更好?将这部分放在FPGA上,将这部分放在Nervana上或将这部分分成两部分,然后再为GPU重写这部分?”

而有关OneAPI的整个故事-只是您要说,“为什么要考虑很长时间,我现在就快点尝试一下,这很简单”。

还不,不容易。 但是会有。



专家后记

谢谢大家的提问。 我有可能甚至很可能是错的,不准确的并且犯了错误。 确实发生了,在Internet上经常有人错了。

我希望我能够对异构编程和DPC ++感兴趣。 我想向所有人推荐sycl.tech网站,其中包含大量报告,包括来自世界知名专家的报告(需要英语)

对所有人都好!

发行人的PS。 这次,根据编辑委员会的一致决定,决定将最佳问题奖授予答案的作者。 我认为您会同意这是公平的。

Source: https://habr.com/ru/post/zh-CN479486/


All Articles