fpga深度学习gpu加速_TornadoVM：使用GPU和FPGA加速Java
fpga深度学习gpu加速
重要要点
- TornadoVM是一个编程和执行框架，用于在异构硬件（多核CPU，GPU和FPGA）上卸载和运行JVM应用程序
- TornadoVM通过OpenCL的新后端扩展了Graal JIT编译器
- 为TornadoVM编写的应用程序是单源代码-相同的代码用于表示主机代码和加速的代码
- TornadoVM可以在计算设备之间执行实时任务迁移
去年三月，我在QCon-London上发表了有关TornadoVM的演讲，在其中我对TornadoVM进行了介绍并解释了其工作原理。 在本文中，我从QCon London演讲中进行了扩展，并展示了有关开发人员如何通过在异构硬件上自动运行Java来从中受益的更多详细信息。
首先，我将提供TornadoVM项目和体系结构的一般概述。 其次，我将通过一个实际的例子来说明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支持Java 8，并通过GraalVM 19.3.0支持OpenJDK 11。 TornadoVM还与OpenCL 1.2兼容，因此可以在多种设备上运行，例如GPU（AMD和NVIDIA），FPGA（Xilinx和Intel），集成GPU（例如Mali ARM和Intel HD Graphics）以及多核CPU。
实践中的TornadoVM
让我们通过一个实际的例子来详细介绍。 如下所示，我演示了如何在多核CPU，GPU和集成GPU上使用TornadoVM编程和运行矩阵乘法。 矩阵乘法是一个简单的代码，用于说明TornadoVM中的不同概念，它构成了许多机器学习和深度学习应用程序的核心。
注意：尽管TornadoVM是用Java编程的，但可以通过GraalVM（ Truffle ）的Polyglot编程框架将计算内核公开给其他JVM语言。
以下代码段显示了用Java编程的矩阵乘法：
class Compute {
public static void matrixMultiplication(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) + j];
C[(i * size) + j] = sum;
}
}
}
} 该代码段显示了用于GPU计算的经典且规范的矩阵乘法示例。 为了使用TornadoVM加速此代码段，我们首先必须注释可以并行化的循环。 在这种情况下，我们可以完全并行化两个最外面的循环，其中迭代之间没有依赖关系。 我们使用TornadoVM注释@Parallel注释@Parallel ，如下所示：
class Compute {
public static void matrixMultiplication(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 * size) + j];
C[(i * size) + j] = sum;
}
}
}
} @Parallel批注由TornadoVM JIT编译器（将Java字节码转换为OpenCL）用作提示。
TornadoVM JIT编译器不会强制并行化。 相反，它检查带注释的循环是否可以并行化，并在OpenCL（ get_global_id(dimension) ） get_global_id(dimension) for循环替换等效的并行索引。 如果不能并行处理for循环，则TornadoVM退出并执行顺序代码。
此外，开发人员必须确定要加速的Java方法。 为此，TornadoVM公开了一个轻量级的基于任务的API，该API设置了要加速的方法列表-每个方法都对应一个任务。 开发人员可以通过任务计划程序创建一组任务。 以下代码段显示了如何为矩阵乘法示例创建任务计划：
TaskSchedule t = new TaskSchedule("s0")
.task("t0", Compute::matrixMultiplication, matrixA, matrixB, result, size)
.streamOut(result);我们创建一个任务计划对象（t）。 在其构造函数中，我们传递任务的名称。 可以是任何名称。 此名称对于更改将要在其中执行所有任务的设备很有用。 然后，我们定义了一组任务。 在此示例中，我们只有一个，但是它可以是任意数量的任务。
任务的参数如下：我们还传递一个名称（在本例中为“ t0 ”）和对要加速的方法的引用（在本例中，它指向Java类Compute的方法matrixMultiplication 。其余参数对应于该方法的实际参数集。
最后，我们指出要与主机（CPU）同步的变量或数组。 这是必需的，因为通常，GPU和FPGA与CPU共享的内存不同。 因此，TornadoVM运行时将为目标设备上的所有变量分配空间，并且它将执行从主机（CPU）到设备（例如GPU）的数据传输。 因此，为了最终获得结果，我们通过TornadoVM API调用streamOut同步变量列表。
到目前为止，我们已经声明了任务，并将其放置在可以执行并行化的代码中。 要使用TornadoVM执行应用程序，我们需要在TaskSchedule对象上调用`execute（）`方法。
这是一个阻塞调用，它将创建所有OpenCL缓冲区，创建执行图，将所有任务从Java字节码编译到OpenCL，并最终在目标设备上执行生成的OpenCL程序。 另外，TornadoVM可以将许多方法组合在一起，以在一个编译单元中一起编译，并在同一设备（例如，同一GPU）上执行。 这为优化主机和异构设备之间的数据传输创造了机会，因为它们通常不与主要主机共享内存（除非设备是集成GPU，例如AMD APU，ARM Mali或Intel HD Graphics GPU）。
请注意，我们没有在源代码中设置任何特定于设备的信息，并且我们共享相同的代码以在多核CPU，GPU和FPGA上运行。 TornadoVM运行时和JIT编译器将根据体系结构自动优化代码。
因此，让我们运行我们的代码示例。 我将首先向您展示如何设置TornadoVM环境。 Github上有所有这些示例的存储库。
运行矩阵乘法：设置TorandoVM
我们将使用Graal 19.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 clone --depth 1 https://github.com/beehive-lab/TornadoVM
$ cd TornadoVM
$ export PATH=$PWD/bin/bin:$PATH
$ export TORNADO_SDK=$PWD/bin/sdk
$ export CMAKE_ROOT=<SET YOUR PATH TO CMAKE ROOT>
$ make graal-jdk-11
$ export TORNADO_ROOT=$PWD现在，我们将下载带有示例的存储库。
$ git clone 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/" ## Defined previously
$ export TORNADO_SDK=${TORNADO_ROOT}/bin/sdk
$ export CLASSPATH=target/tornado-1.0-SNAPSHOT.jar
$ mvn clean install现在我们已经准备好执行示例。 我们可以从TornadoVM探索哪些设备可用和可见开始。
$ tornado --devices
Number of Tornado drivers: 1
Total number of devices : 3
Tornado device=0:0
NVIDIA CUDA -- GeForce GTX 1050
Global Memory Size: 3.9 GB
Local Memory Size: 48.0 KB
Workgroup Dimensions: 3
Max WorkGroup Configuration: [1024, 1024, 64]
Device OpenCL C version: OpenCL C 1.2
Tornado device=0:1
Intel(R) OpenCL -- Intel(R) Core(TM) i7-7700HQ CPU @ 2.80GHz
Global Memory Size: 31.0 GB
Local Memory Size: 32.0 KB
Workgroup Dimensions: 3
Max WorkGroup Configuration: [8192, 8192, 8192]
Device OpenCL C version: OpenCL C 1.2
Tornado device=0:2
Intel(R) OpenCL HD Graphics -- Intel(R) Gen9 HD Graphics NEO
Global Memory Size: 24.8 GB
Local Memory Size: 64.0 KB
Workgroup Dimensions: 3
Max WorkGroup Configuration: [256, 256, 256]
Device OpenCL C version: OpenCL C 2.0就我而言，我的笔记本电脑上有三台设备可用：NVIDIA GPU，Intel多核CPU和Intel HD Graphics（集成GPU）。 TornadoVM默认选择设备0。 但是，我们可以通过将任务与设备关联来更改设备。 让我们从默认配置开始。
$ tornado qconlondon.MatrixMultiplication 512 tornado该程序执行100次“矩阵乘法”方法，并报告每次迭代的总时间。 该方法是一个简单的示例，用于演示正在发生的事情-稍后，我们将使用JMH进行正确的性能比较。
$ tornado qconlondon.MatrixMultiplication 512 tornado
Computing MxM of 512x512
Total time: 77568790 (ns), 0.0776 (s)
Total time: 3133182 (ns), 0.0031 (s)
Total time: 3126146 (ns), 0.0031 (s)
…请注意，第一次迭代要比其余的迭代花费更长的时间-这是由于JIT编译预热所致，并且在我们使用JMH时会消失。
我们第一次执行任务计划时，TornadoVM会调用OpenCL JIT编译器从Java字节码优化并生成OpenCL C代码。 然后，一旦生成代码，TornadoVM将生成的代码安装在代码缓存中，并且如果在运行时的任何时候再次执行相同的任务，则二进制文件可以重用。 为确保TornadoVM在GPU（设备0）上运行，我们可以启用调试信息，如下所示：
$ tornado --debug qconlondon.MatrixMultiplication 512 tornado
Computing MxM of 512x512
task info: s0.t0
platform : NVIDIA CUDA
device : GeForce GTX 1050 CL_DEVICE_TYPE_GPU (available)
dims : 2
global work offset: [0, 0]
global work size : [512, 512]
local work size : [32, 32, 1]太好了，TornadoVM在NVIDIA GTX 1050上运行我们的Java代码以实现矩阵乘法。作为参考，我们还运行顺序应用程序。 无需调用TornadoVM JIT编译器即可加速代码。 我们向程序传递一个额外的参数来表明这一点：
$ tornado qconlondon.MatrixMultiplication 512 sequential
Computing MxM of 512x512
Total time: 259398036 (ns), 0.2594 (s)
Total time: 247857535 (ns), 0.2479 (s)
...我们看到的是，即使使用TornadoVM JIT编译器，第一次迭代的速度也要快3.3倍。 然后，从第二次迭代开始， 我们将Java顺序代码的速度提高了80倍 。 但是，请谨慎使用此数字。 在下一节中，我们将介绍使用Java JMH的性能比较。
如何更换设备？
我们可以通过命令更改在其中运行应用程序的设备。 例如，要在Intel Integrated Graphics上运行，我们可以使用以下选项执行：
$ tornado -Ds0.t0.device=0:2 --debug qconlondon.MatrixMultiplication 512 tornado
Computing MxM of 512x512
task info: s0.t0
platform : Intel(R) OpenCL HD Graphics
device : Intel(R) Gen9 HD Graphics NEO CL_DEVICE_TYPE_GPU (available)
dims : 2
global work offset: [0, 0]
global work size : [512, 512]
local work size : [16, 16, 1] 语法如下-D<taskScheduleName>:<taskName>.device=0:<deviceIndex>
在Dell XPS 15笔记本电脑上运行的TornadoVM for MxM的性能
使用这些选项，我们可以轻松地开始获得一些性能结果。 下图显示了通过Java顺序实现在不同的OpenCL设备上运行TornadoVM时，TornadoVM的速度提高。 报告的提速对应于使用Java JMH框架进行基准测试的平均值。 注意，由于高速化，y轴以对数标度表示。 所有使用JMH的基准与示例都包含在同一存储库中。 如我们所见，与Java Hotspot相比，在具有TornadoVM的多核CPU上运行可以达到3.6倍。 在GPU上运行时，与Java相比，我们可以将Intel HD显卡和NVIDIA 1050分别提高39倍和270倍。
执行模型与编译
到目前为止，我们已经简要说明了TornadoVM API以及如何在用户级别使用TornadoVM运行应用程序。 现在，让我们更深入地了解TornadoVM如何在目标设备上执行代码。
下图显示了JVM和TornadoVM之间的执行流。
任务计划的定义以及对TornadoVM API的execute方法的调用在单个Java线程（例如主线程）上运行。 execute方法是一个阻塞调用，当该方法的执行返回时，它保证并行设备上的执行已完成。 调用execute方法时，TornadoVM首先构建一个数据流图，该图表示如何在任务计划中的不同任务之间传递数据。 该图用于优化数据传输。
然后，TornadoVM生成新的字节码（用于在目标设备上协调执行的简单指令，例如COPY_IN，LAUNCH，COPY_OUT，BARIER等）。 首次启动代码（通过LAUNCH字节码）时，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 x 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: LAUNCH task s0.t0-matrixMultiplication on NVIDIA -- GeForce GTX 1050
vm: STREAM_OUT_BLOCKING [F@33d512c1 on NVIDIA -- GeForce GTX 1050我们前面介绍的矩阵乘法方法接收三个参数（矩阵A，B和C）。 对于每个变量，TornadoVM执行从主机到设备的数据传输（COPY_IN）。 然后，它使用LAUNCH字节码运行该应用程序。
回想一下，第一次执行LAUNCH时，TornadoVM会调用OpenCL JIT编译器，其中的代码针对每个计算设备进行了专门化和优化。 最后，TornadoVM执行从设备到主要主机的复制（STREAM_OUT_BLOCKING）以获取结果。
分析OpenCL生成的代码
让我们深入研究TornadoVM生成的OpenCL内核。 使用TornadoVM，我们可以使用--printKernel标志来调试和检查生成的内核，如下所示：
$ tornado --printKernel qconlondon.MatrixMultiplication 512 tornado TornadoVM在任务计划中为每个任务生成一个内核。 此外，它lookupBufferAddress生成一个名为lookupBufferAddress的内核，该内核将在VM引导期间执行。 该内核背后的原因是，TornadoVM仅分配一个大缓冲区来充当目标设备上的堆。 为此，它需要一个有效的指针，该指针将用作目标设备的基本地址，TornadoVM可以在该目标设备中执行数据传输。 lookupBufferAddress内核返回此基本指针。
第二个内核对应于我们加速的Java方法中的OpenCL代码。 以下代码段显示了生成的内核的简化，并带有Java和OpenCL代码要点的注释。 请注意，生成的内核可能会因目标体系结构而异。 还要注意，TornadoVM从静态单一分配（ SSA ）表示生成OpenCL C代码，其中每个变量只分配一次。 这是因为TornadoVM是Graal-IR的扩展，它以SSA表示形式工作（HotSpot的主流JIT编译器C2也是如此）。
__kernel void lookupBufferAddress(...parameters) {
__global ulong *_frame = (__global ulong *) &_heap_base[_frame_base];
_frame[0] = (ulong) _heap_base;
}
__kernel void matrixMultiplication(...parameters) {
// Variables declaration …
// Access to the stack-frame
__global ulong *_frame = (__global ulong *) &_heap_base[_frame_base];
// Access elements within the stack-frame
ul_0 = (ulong) _frame[6]; // base address of input matrix A
ul_1 = (ulong) _frame[7]; // base address of input matrix B
ul_2 = (ulong) _frame[8]; // base address of input matrix C
i_3 = get_global_id(1); // Parallel OpenCL indexing (2nd dimension)
i_4 = i_3;
for(;i_4 < 512;) {
i_5 = get_global_id(0); // Parallel OpenCL indexing (1st dimension)
i_6 = i_5;
for(;i_6 < 512;) {
i_7 = i_4 << 9;
f_8 = 0.0F;
i_9 = 0;
for(;i_9 < 512;) {
i_10 = i_9 + 1;
i_11 = i_7 + i_9;
l_12 = (long) i_11;
l_13 = l_12 << 2;
l_14 = l_13 + 24L; // Skip Java object header
ul_15 = ul_0 + l_14;
f_16 = *((__global float *) ul_15); // Load element from matrix A
i_17 = i_9 << 9;
i_18 = i_17 + i_6;
l_19 = (long) i_18;
l_20 = l_19 << 2;
l_21 = l_20 + 24L;
ul_22 = ul_1 + l_21;
f_23 = *((__global float *) ul_22);// Load element from matrix B
f_24 = fma(f_16, f_23, f_8); // Computation (fuse-multiple-add)
f_8 = f_24;
i_9 = i_10;
}
i_25 = i_6 + i_7;
l_26 = (long) i_25;
l_27 = l_26 << 2;
l_28 = l_27 + 24L;
ul_29 = ul_2 + l_28;
*((__global float *) ul_29) = f_8; // Store the result in Matrix C
i_30 = get_global_size(0);
i_31 = i_30 + i_6;
i_6 = i_31;
}
i_32 = get_global_size(1);
i_33 = i_32 + i_4;
i_4 = i_33;
}
}如何使用TornadoVM？
在本文中，我们集中于一个简单的示例，即矩阵乘法，以轻松展示TornadoVM运行时和JIT编译器的不同部分。 但是，使用TornadoVM，您不仅可以使用简单的数据类型对多个任务进行编程。 TornadoVM已被用于通过Microsoft Kinect Fusion 加速SLAM （同步本地化和映射）应用程序，与NVIDIA GPU上的Java相比，每秒可加速90帧。 该应用程序包含大约7k行Java代码，这些代码可以使用TornadoVM加速，并且强调了TornadoVM能够生成的Java结构的复杂性。
通常，TornadoVM适用于加速遵循SIMD（单指令多数据）模式的工作负载以及管道应用程序。 令人惊讶的是，该分类包括了广泛的应用，例如深度学习，机器学习，数学和物理模拟，计算摄影，计算机视觉，金融应用，信号处理和化学。
另外，开发人员可以在GraalVM之上从Python，R，Ruby，Javascript或任何其他语言调用TornadoVM（正如我在QCon-London中展示的，用于加速Node.js应用程序 ）。
TornadoVM诞生于学术界（目前在曼彻斯特大学正在开发中），但是已经有一些公司使用TornadoVM来加速深度学习应用程序。
一个例子是总部位于伦敦的科技公司Exus Ltd.，该公司目前正在改进英国NHS（医疗保健）系统，以预测患者的住院再住院次数，这成功地改善了2个数据集的训练阶段的绩效。使用TornadoVM将14倍的百万患者
卢森堡的NEUROCOM是行业中早期采用TornadoVM的另一个例子，该公司正在GPU上使用TornadoVM将自然语言处理中使用的一些关键计算加速了10倍和28倍（具体而言，分别使用余弦相似性度量算法的Levenshtein距离和层次分类） 。
摘要
TornadoVM是OpenJDK和GraalVM的插件，允许开发人员将JVM应用程序脱机到异构硬件中，包括多核CPU，GPU和FPGA。 此外，TornadoVM在设备之间执行实时任务迁移，以最大化整个应用程序的性能。 本文通过一个示例探讨了TornadoVM的功能。 我们探索了TornadoVM的执行方式，并发现了生成的代码。
本文仅介绍TornadoVM是什么以及它可以做什么。 在这篇介绍性文章中，我们无法涵盖许多重要主题。 例如，对每种架构的编译器专业化的描述，如何有效运行以减少计算量，FPGA编译管道和实时任务迁移。 您可以通过以下链接找到有关这些主题的更多信息：
更多参考
致谢
TornadoVM开发得到了欧盟Horizon 2020 E2Data 780245的部分支持。
fpga深度学习gpu加速