TornadoVM：使用GPU和FPGA加速Java

TornadoVM是一个编程和执行框架，用于在异构硬件(多核CPU、GPU和FPGA)上卸载和运行JVM应用程序。

为TornadoVM编写的应用程序是单一源代码的-使用相同的代码来表示主机代码和加速代码。

去年3月，我在伦敦QCon做了一次关于TornadoVM的演讲，我在会上介绍了TornadoVM并解释了它的工作原理。在本文中，我将从QCon London演讲展开，并详细介绍开发人员如何通过在异构硬件上自动运行Java来从中受益。

首先，我将提供TornadoVM项目和架构的总体概述。其次，我会用一个实际的例子来解释TornadoVM的不同部分。

没有一种单一的计算机体系结构最适合高效地执行所有类型的工作负载。这导致了近年来异构硬件的激增，这意味着我们编写的每个系统都可能混合了计算元素。

这些元素中的每一个都有不同的硬件特性。硬件异构性使程序员能够在降低能耗的同时提高应用程序的性能。

这些用于计算的新型异构设备包括多核CPU、图形处理单元(GPU)和现场可编程门阵列(FPGA)。这种多样性是巨大的，但我们需要一种方法来有效地对这些新设备进行编程。

最好的例子是两种最流行的异构编程语言CUDA和OpenCL。但是，它们公开了API中的几个低级功能，使得非专家用户很难使用它们。作为示例，我突出显示了OpenCL 3.0标准中的以下引用：

OpenCL的目标是想要编写可移植但高效代码的专业程序员。[.]。因此，OpenCL提供了一个低级硬件抽象和一个框架来支持编程，并且暴露了底层硬件的许多细节。

前面的语句也适用于CUDA和类似的并行编程模型。在工业和学术界，开发人员倾向于使用高级的、面向对象的编程语言，而不是使用低级编程语言，这些语言通常在托管运行时环境(如Java、R、Python和JavaScript)上执行。尽管许多程序员可能认为这样的编程语言已经适用于在异构硬件上透明地执行，但现实情况是，它们的支持要么非常有限，要么不存在。

在本文中，我们将探讨TornadoVM，它是异构计算低级并行编程语言的替代方案。我们将展示开发人员如何在不需要了解并行计算体系结构或并行编程模型的情况下使用多核CPU和GPU。

简而言之，TornadoVM是一个针对JVM语言的并行编程框架，它可以透明地、动态地将Java字节码卸载到OpenCL中，并在异构硬件上执行生成的代码。此外，TornadoVM集成了一个优化运行时，可以重用设备缓冲区并保存跨设备的数据传输，以及一个新的动态应用程序重新配置组件，用于跨计算设备执行实时任务迁移。

下图显示了TornadoVM项目的高级概述。正如我们所看到的，TornadoVM由分层的微内核软件架构组成，其中核心组件是TornadoVM执行引擎。在顶层，TornadoVM向开发人员公开API。这是因为TornadoVM当前不检测并行性(自动并行化)。相反，它利用并行性。因此，TornadoVM需要一种方法来识别哪些方法或函数是在GPU和FPGA上运行的候选方法或函数。

此外，TornadoVM包含核心运行时，它分为几个组件：a)带有新字节码生成器的数据流优化器；b)运行新字节码的小字节码解释器；以及c)JIT编译器和内存管理。在本文中，我将重点介绍API、运行时以及对JIT编译器的总体概述。

最后，如上图所示，TornadoVM目前使用最新的JDK(U242)和JVMCI支持Java8，并通过GraalVM19.3.0支持OpenJDK 11。TornadoVM还与OpenCL 1.2兼容，因此可以在多种设备上运行，如GPU(AMD和NVIDIA)、FPGA(Xilinx和Intel)、集成GPU(如MariARM和Intel HD Graphics)以及多核CPU。

让我们用一个实际例子来详细说明一下。如下所示，我展示了如何在多核CPU、GPU和集成GPU上使用TornadoVM编程和运行矩阵乘法。矩阵乘法是用来说明TornadoVM中不同概念的简单代码，它构成了许多机器学习和深度学习应用程序的核心。

注意：虽然TornadoVM是用Java编程的，但是计算内核可以通过GraalVM(Truffle)的Polyglot编程框架向其他JVM语言公开。

类计算{public static void matrix(Final Float[]A，Final Float[]B，Final Float[]C，Final int Size){For(int i=0；i<；size；i++){For(int j=0；j<；size；j++){Float sum=0.0f；for(int k=0；k<；size；k++)sum+=A[(i*size)+k]*B[(k*size)+。C[(i*大小)+j]=SUM；}。

代码片段显示了用于GPU计算的经典和规范的矩阵乘法示例。要使用TornadoVM加速此代码片段，我们首先必须注释可以并行化的循环。在这种情况下，我们可以完全并行化最外面的两个循环，其中迭代之间没有依赖关系。我们通过使用TornadoVM注释@Parallel来注释代码，如下所示：

类计算{public static void matrix(Final Float[]A，Final Float[]B，Final Float[]C，Final int Size){For(@Parallel int i=0；i<；size；i++){For(@Parallel int j=0；j<；size；j++){Float sum=0.0f；for(int k=0；k<；size；k++)sum+=A[(i*size)+k]*B[(k。C[(i*大小)+j]=SUM；}。

@Parallel注释被TornadoVM JIT编译器(将Java字节码转换为OpenCL)用作提示。

TornadoVM JIT编译器不强制并行化。相反，它检查带注释的循环是否可以并行化，并替换OpenCL中等价的并行索引(get_global_id(Dimension))的for-循环。如果for循环不能并行化，TornadoVM就退出并执行顺序代码。

此外，开发人员必须确定要加速哪些Java方法。为此，TornadoVM公开了一个基于任务的轻量级API，该API设置要加速的方法列表-其中每个方法对应一个任务。开发人员可以通过任务调度器创建一组任务。下面的代码片段显示如何为矩阵乘法示例创建任务计划：

我们创建一个Task-Schedule对象(T)。在其构造函数中，我们传递任务的名称。可以是任何名字。此名称对于更改要执行所有任务的设备非常有用。然后我们定义一组任务。在本例中，我们只有一个任务，但它可以是任意数量的任务。

任务的参数如下：我们还传递一个名称(在本例中为“t0”)和对我们想要加速的方法的引用(在本例中，它指向Java类Compute中的matrixMultiply方法。其余参数对应于该方法的实际参数集。

最后，我们指出要与主机(CPU)同步的变量或数组。这是必需的，因为通常情况下，GPU和FPGA不与CPU共享相同的内存。因此，TornadoVM运行时将为目标设备上的所有变量分配空间，并执行从主机(CPU)到设备(例如，GPU)的数据传输。因此，为了最终获得结果，我们通过TornadoVMAPI调用stream Out同步变量列表。

到目前为止，我们已经声明了我们的任务，并将它们放在可以执行并行化的代码中。要使用TornadoVM执行应用程序，需要在TaskSchedule对象上调用`ecute()`方法。

这是一个阻塞调用，它将创建所有OpenCL缓冲区，创建执行图，将所有任务从Java字节码编译为OpenCL，最后在目标设备上执行生成的OpenCL程序。此外，TornadoVM可以将许多方法组合在一个编译单元中一起编译，并在同一设备上(例如，在同一GPU上)执行。这为优化主机和异构设备之间的数据传输创造了机会，因为它们通常不与主要主机共享内存(除非设备是集成的GPU，如AMD APU、ARM马里或英特尔HD图形GPU)。

请注意，我们在源代码中没有设置任何特定于设备的信息，并且我们共享在多核CPU、GPU和FPGA上运行的相同代码。TornadoVM运行时和JIT编译器将根据架构自动优化代码。

那么，让我们运行我们的代码示例。我将首先向您展示如何设置TornadoVM环境。在Github上有一个存储库，里面有所有这些示例。

我们将使用Graal19.3.0作为JDK运行TornadoVM。请注意，我们经常更新Graal版本。Graal 20.x计划在今年年底集成到TornadoVM中。要执行代码，我们假设安装了OpenCL。请在此处查看所有前提条件。

$mkdir-p TornadoVM$cd TornadoVM$wget https://github.com/graalvm/graalvm-ce-builds/releases/download/vm-19.3.0/graalvm-ce-java11-linux-amd64-19.3.0.tar.gz$TAR-XF graalvm-ce-java11-linux-amd64-19.3.0.tar.gz$export JAVA_HOME=$pwd/graalvm-ce-java11-19.3.0$git克隆--深度1 https://github.com/beehive-lab/TornadoVM$cd TornadoVM。$export path=$pwd/bin/bin：$path$export tornado_sdk=$pwd/bin/sdk$export CMAKE_ROOT=<；将路径设置为CMAKE root>；$make graal-jdk-11$export tornado_root=$PWD

$git克隆https://github.com/jjfumero/qconlondon2020-tornadovm$CD qconlondon2020-tornadovm/$export JAVA_HOME=/PATH/to/graalvm-ce-java11-19.3.0$export PATH="；${PATH}：${TORNADO_ROOT}/bin/bin/"；##先前定义的$export tornado_sdk=${TORNADO_ROOT}/bin/sdk$export CLASSPATH=TARGET/TORNADO-1.0-SNAPSHOT.jar$MVN全新安装。

现在我们已经做好了执行示例的一切准备。我们可以从研究TornadoVM中可用和可见的设备开始。

$Tornado--设备Tornado驱动程序数量：1设备总数：3 Tornado设备=0：0 NVIDIA CUDA--GeForce GTX 1050全局内存大小：3.9 GB本地内存大小：48.0 KB工作组维度：3最大工作组配置：[1024,1024，64]Device OpenCL C版本：OpenCL C 1.2 Tornado Device=0：1英特尔(R)OpenCL--英特尔(R)酷睿(TM)i7-7700HQ [email protected] GHz全局内存大小：31.0 GB本地内存大小：32.0 KB工作组维度：3最大工作组配置：[8192,8192，8192]Device OpenCL C版本：OpenCL C 1.2 Tornado Device=0：2英特尔(R)OpenCL HD显卡--英特尔(R)Gen9 HD显卡NEO全局内存大小：24.8 GB本地内存大小：64.0 KB工作组维度：3最大工作组配置：[256,256,256]设备OpenCL C版本：OpenCL C 2.0。

就我的情况而言，我的笔记本电脑上有三个设备可用：NVIDIA GPU、英特尔多核CPU和英特尔高清显卡(集成GPU)。TornadoVM默认选择设备0。但是，我们可以通过将任务与设备相关联来更改设备。让我们从默认配置开始。

此程序执行矩阵乘法100次，并报告每次迭代的总时间。这个方法是一个简单的示例来演示发生了什么-稍后我们将使用JMH进行适当的性能比较。

$Tornado qconlondon.矩阵乘法512龙卷风计算MxM为512x512总时间：77568790(Ns)，0.0776(S)总时间：3133182(Ns)，0.0031(S)总时间：3126146(Ns)，0.0031(S)…。

请注意，第一次迭代比其余迭代花费的时间更长-这是由于JIT编译预热造成的，当我们使用JMH时，它将消失。

第一次执行任务调度时，TornadoVM调用OpenCL JIT编译器从Java字节码优化并生成OpenCL C代码。然后，一旦生成了代码，TornadoVM就会将生成的代码安装在代码缓存中，如果在运行时的任何时候再次执行相同的任务，则可以重用二进制文件。为确保TornadoVM在GPU(设备0)上运行，我们可以按如下方式启用调试信息：

$tornado--debug qconlondon.MatrixMultiplex 512 tornado Computing MxM of 512x512任务信息：s0.t0平台：NVIDIA CUDA设备：GeForce GTX 1050 CL_DEVICE_TYPE_GPU(可用)dims：2全局工作偏移：[0，0]全局工作大小：[512,512]本地工作大小：[32，32，1]。

太棒了，TornadoVM正在NVIDIA GTX1050上运行我们用于矩阵乘法的Java代码。作为参考，我们还运行顺序应用程序。这不需要调用TornadoVM JIT编译器来加速代码。我们向程序传递一个额外的参数来表示这一点：

$Tornado qconlondon.矩阵乘法512顺序计算512x512的MxM总时间：259398036(Ns)，0.2594(S)总时间：247857535(Ns)，0.2479(S).。

我们看到的是，即使使用TornadoVM JIT编译器，第一次迭代的速度也要快3.3倍。然后，从第二次迭代中，我们获得了比Java顺序代码快80倍的速度。在下一节中，我们将介绍使用Java JMH进行的性能比较。

我们可以通过命令更改运行应用程序的设备。例如，要在英特尔集成显卡上运行，我们可以使用以下选项执行：

$tornado-Ds0.t0.device=0：2--debug qconlondon。矩阵乘法512龙卷风计算MxM共512x512任务信息：s0.t0平台：英特尔(R)OpenCL HD显卡设备：英特尔(R)Gen9 HD Graphics neo CL_DEVICE_TYPE_GPU(可用)dims：2全局工作偏移：[0，0]全局工作大小：[512,512]本地工作大小：[16，16。

使用这些选项，我们可以轻松地开始获得一些性能结果。下图显示了通过Java顺序实现在不同OpenCL设备上运行TornadoVM时，TornadoVM的加速比(越高越好)。报告的加速比对应于使用Java JMH框架进行基准测试的平均值。请注意，由于加速比很高，y轴以对数刻度表示。所有使用JMH的基准测试都包含在与示例相同的存储库中。正如我们所看到的，与Java HotSpot相比，运行在带有TornadoVM的多核CPU上可以获得3.6倍的性能。在GPU上运行时，与用于英特尔高清显卡和NVIDIA 1050的Java相比，我们可以实现高达39倍和270倍的性能。

到目前为止，我们已经简要介绍了TornadoVM API以及如何在用户级别使用TornadoVM运行应用程序。现在让我们深入一点，看看TornadoVM如何在目标设备上执行代码。

任务调度的定义和从TornadoVMAPI对Execute方法的调用在单个Java线程(例如，主线程)上运行。Execute方法是一个阻塞调用，当该方法的执行返回时，它保证并行设备上的执行已经完成。当调用Execute方法时，TornadoVM首先构建一个数据流图，该图表示数据如何跨任务调度内的不同任务进行通信。此图用于优化数据传输。

然后，TornadoVM生成新的字节码(用于编排目标设备上的执行的简单指令，如COPY_IN、LAUNT、COPY_OUT、BALAR等)。当代码第一次启动时(通过启动字节码)，TornadoVM调用OpenCL JIT编译器并将输入的Java字节码从每个任务(每个要加速的Java方法)转换为优化的OpenCL C代码。

TornadoVM根据目标设备的不同专门处理OpenCL C代码，这意味着为GPU生成的代码对于CPU和FPGA是不同的。这是因为OpenCL代码可以跨设备移植，但性能并不统一。因此，TornadoVM通过对每个设备进行专门化并应用不同的优化来提高性能。

注意：TornadoVM JIT编译器在单个线程中运行，因此在繁重的负载下可能会耗尽编译器资源，正如我们在HotSpot中看到的那样。

编译的最后一步是通过调用OpenCL驱动程序从优化和专用的OpenCL C代码编译到目标平台。例如，如果应用程序在NVIDIA GPU上执行，此步骤将生成相应的PTX代码。

一旦生成并编译了OpenCL代码，TornadoVM就会在目标设备上启动应用程序。为此，TornadoVM部署了许多线程来运行内核。要部署的线程数量取决于应用程序的输入大小和硬件特性。

例如，我们前面展示的矩阵乘法示例使用512×512个线程的块部署在GPU上。这意味着TornadoVM从编写的单线程Java应用程序部署了512x512个线程块。如果目标设备是多核CPU，则TornadoVM部署的线程数与可用的最大CPU核数相同。

并行设备上的执行完成后，TornadoVM将结果复制到Java堆(通过字节码copy_out使其对主机端可见)，最后将控制权返回给JVM中的主线程。

我们可以查询TornadoVM为每个应用程序生成的字节码。例如，下面的代码片段显示了使用TornadoVM的字节码的调试信息运行矩阵乘法时的简化输出：

$Tornado--printBytecodes qconlondon.MatrixMultiplication 512 Tornado VM：COPY_IN[F@3e694b3f on NVIDIA--GeForce GTX 1050 VM：COPY_IN[F@397fbdb on NVIDIA--GeForce GTX 1050 VM：COPY_IN[F@33d512c1 on NVIDIA--GeForce GTX 1050 VM：启动任务s0.t0。

我们前面介绍的矩阵乘法接收三个参数(矩阵A、B和C)。对于每个变量，TornadoVM执行从主机到设备的数据传输(COPY_IN)。然后，它使用启动字节码运行应用程序。

回想一下，在第一次执行启动时，TornadoVM调用OpenCL JIT编译器，在该编译器中，代码针对每个计算设备进行专门化和优化。最后，TornadoVM执行从设备到主要主机的复制(STREAM_OUT_BLOCKING)以获得结果。

让我们深入研究TornadoVM生成的OpenCL内核。使用TornadoVM，我们可以使用--printKernel标志调试和检查生成的内核，如下所示：

TornadoVM在任务调度内为每个任务生成一个内核。此外，它还生成一个名为lookupBufferAddress的内核，该内核在VM引导期间执行。该内核背后的原因是，TornadoVM只分配一个大缓冲区，该缓冲区充当目标设备上的堆。为此，它需要一个有效的指针，该指针将用作目标设备的基地址，TornadoVM可以在该设备中执行数据传输。lookupBufferAddress内核返回这个基指针。

第二个内核对应于Open。

..