系统gpu 调试_探索：如何在GPU上运行JAVA程序

图形处理单元(GPU)的编程似乎与Java编程相去甚远。这是可以理解的，因为Java的大多数用例并不适用于gpu。尽管如此，gpu提供了万亿次浮点运算的性能，所以让我们探索它们的可能性。

为了使这个主题更容易理解，我将花一些时间解释GPU架构以及一些历史，这将使深入研究硬件编程变得更容易。在展示了GPU与CPU计算的不同之后，我将展示如何在Java世界中使用GPU。最后，我将描述可用于编写Java代码并在gpu上运行它的主要框架和库，并提供一些代码示例。

关于GPU的一些背景：

GPU最早是在1999年由英伟达推广的。它是一种特殊的处理器，设计用来处理图形数据，然后将其传输到显示器上。在大多数情况下，它允许从CPU卸载一些计算，从而在加速这些卸载计算的同时释放CPU资源。其结果是，可以处理更多的输入数据，并以更高的输出分辨率呈现，使视觉表示更吸引人，帧速率更流畅。

2D/3D处理的本质主要是矩阵操作，因此可以使用大规模并行方法来处理。什么是有效的图像处理方法?为了回答这个问题，让我们比较一下标准cpu(如图1所示)和gpu的架构。 

 图1.CPU的架构

在CPU中，实际的处理元素——取数据器、算术逻辑单元(ALU)和执行上下文——只是整个系统的一小部分。为了加速以不可预测的顺序到达的不规则计算，存在一个大的、快速的、昂贵的缓存;不同种类的预取器;和分支预测。

你不需要所有这些在一个GPU上，因为数据是在一个可预测的方式接收和GPU执行一个非常有限的操作集的数据。因此，有可能使用类似于图2的块架构制作一个非常小且便宜的处理器。

 图2. GPU单个核心架构

因为这些类型的处理器很便宜，而且它们以并行块的方式处理数据，所以很容易让它们中的许多并行工作。这种设计被称为多指令、多数据或MIMD。

第二种方法关注这样一个事实:一条指令通常应用于多个数据项。这被称为单指令、多数据或SIMD。在本设计中，一个GPU包含多个ALUs和执行上下文，有一小块区域用于共享上下文数据，如图3所示。

 图3.对比MIMD风格GPU架构(左)和SIMD风格架构(右)

结合SIMD和MIMD处理可以提供最大的处理吞吐量，我稍后将对此进行讨论。在这样的设计中，您有多个SIMD处理器并行运行，如图4所示。

图4.并行运行多个SIMD处理器;这里有16个核，总共128个alu

因为你有一堆小而简单的处理器，你可以对它们进行编程以在输出中获得特殊效果。

在GPU上运行程序

游戏早期的大部分视觉效果实际上是在GPU上运行的硬编码的小程序，并应用于来自CPU的数据流。

很明显，即使在那时，硬编码算法仍然是不够的，特别是在游戏设计中，视觉表现实际上是主要卖点之一。作为回应，大供应商开放了对gpu的访问，然后第三方开发者可以为他们编码。

典型的方法是用一种特殊的语言(通常是C的一个子集)编写称为着色器的小程序，然后用对应架构的特殊编译器编译它们。选择着色器这个术语是因为着色器通常用于控制光照和阴影效果，但是没有理由它们不能处理其他特殊效果。

每个GPU供应商都有自己特定的编程语言和基础设施来为其硬件创建着色器。通过这些努力，已经创建了几个平台。主要包括:

• DirectCompute:微软专有的着色器语言/API，是Direct3D的一部分，从directx10开始

• AMD FireStream: ATI/Radeon专利技术，已被AMD终止

• OpenACC:多厂商联盟的并行计算解决方案

• C++ AMP:一个微软在c++中用于数据并行的专有库

• CUDA: Nvidia的专有平台，它使用C语言的一个子集

• OpenCL:一种最初由苹果公司设计但现在由Khronos集团管理的通用标准

大多数时候，使用gpu是低级的编程。为了使开发人员编写代码时更容易理解，提供了几个抽象。最著名的是来自微软的DirectX和来自Khronos Group的OpenGL。这些api是用来编写高级代码的，然后开发者可以将这些代码无缝地卸载给GPU。

据我所知，没有支持DirectX的Java基础设施，但是有一个很好的OpenGL绑定。JSR 231在2002年开始用于GPU编程，但在2008年被抛弃，只支持OpenGL 2.0。一个名为JOCL的独立项目(它也支持OpenCL)继续支持OpenGL，并且它是公开可用的。顺便说一下，著名的《我的世界》游戏就是在下面用JOCL写的。

GPGPU的出现

尽管如此，Java和gpu并不是天各一方，尽管它们应该如此。Java在企业、数据科学和金融领域被大量使用，这些领域需要大量的计算和处理能力。这就是通用GPU (GPGPU)的想法是如何产生的。

以这种方式使用GPU的想法始于视频适配器供应商开始以编程方式打开帧缓冲区，使开发人员能够读取内容。一些黑客意识到他们可以利用GPU的全部功能进行通用计算。食谱很简单:

1. 将数据编码为位图数组。

2. 写一个着色器来处理它。

3. 把它们都提交到视频卡上。

4. 从帧缓冲区检索结果。

5. 解码数据从位图数组。

这是一个非常简单的解释。我不确定这个过程是否在生产中被大量使用，但它确实有效。

之后，来自斯坦福大学的几位研究人员开始寻找一种使GPGPU更容易使用的方法。2005年，他们发布了BrookGPU，这是一个包含语言、编译器和运行时的小型生态系统。

BrookGPU用ANSI c的变体Brook stream编程语言编译程序，它可以针对OpenGL v1.3+， DirectX v9+，或者AMD接近Metal的计算后端进行编译，它可以在微软Windows和Linux上运行。为了调试，BrookGPU还可以在CPU上模拟一个虚拟显卡。

然而，由于当时可用的硬件，它没有起飞。在GPGPU世界中，您需要将数据复制到设备，等待GPU处理数据，然后将数据复制回主运行时。这造成了很多延迟。在2000年代中期，当该项目处于积极开发阶段时，这种延迟几乎阻止了gpu在通用计算中的广泛使用。

尽管如此，许多公司还是看到了这项技术的未来。一些视频适配器供应商开始提供gpgpu与他们的专有技术，和其他组成联盟提供更通用的，多功能编程模型，以运行更多种硬件设备。

现在我已经分享了这个背景，让我们研究一下GPU计算的两种最成功的技术——opencl和cuda，看看Java是如何使用它们的。

OpenCL 和 Java

与许多其他基础设施包一样，OpenCL提供了用c编写的基本实现。从技术上讲，它可以通过Java本机接口(JNI)或Java本机访问(JNA)进行访问，但是这种访问对于大多数开发人员来说工作量太大了。幸运的是，这些工作已经由几个库完成了:JOCL、JogAmp和JavaCL。不幸的是，JavaCL是一个死项目。但是JOCL项目是活跃的，而且是最新的。我将在下面的例子中使用它。

但首先我要解释什么是OpenCL。正如我前面提到的，OpenCL提供了一个非常通用的模型，适合于对各种设备进行编程——不仅是gpu和cpu，而且是数字信号处理器(DSPs)和现场可编程门阵列(FPGAs)。

让我们探索最简单的例子:向量加法，可能是最具代表性和最简单的例子。您有两个正在相加的整数数组和一个结果数组。从第一个数组和第二个数组中取出一个元素，然后将它们的和放入结果数组中，如图5所示。

图5.添加两个数组的内容并将其和存储在结果数组中

如您所见，该操作是高度并发的，因此非常可并行。你可以把每个添加操作推到一个单独的GPU核心。这意味着，如果您有2048个核，就像在Nvidia 1080显卡上一样，您可以同时执行2048个添加操作。这意味着有潜在的万亿次浮点运算能力在等着你。下面是取自JOCL网站的1000万整数数组代码:

public class ArrayGPU {
/**
     * The source code of the OpenCL program 
     */
private static String programSource =
"__kernel void "+
"sampleKernel(__global const float *a,"+
"             __global const float *b,"+
"             __global float *c)"+
"{"+
"    int gid = get_global_id(0);"+
"    c[gid] = a[gid] + b[gid];"+
"}";

public static void main(String args[]){
int n = 10_000_000;
float srcArrayA[] = new float[n];
float srcArrayB[] = new float[n];
float dstArray[] = new float[n];
for (int i=0; i        {
            srcArrayA[i] = i;
            srcArrayB[i] = i;
        }
        Pointer srcA = Pointer.to(srcArrayA);
        Pointer srcB = Pointer.to(srcArrayB);
        Pointer dst = Pointer.to(dstArray);


// The platform, device type and device number
// that will be used
final int platformIndex = 0;
final long deviceType = CL.CL_DEVICE_TYPE_ALL;
final int deviceIndex = 0;

// Enable exceptions and subsequently omit error checks in this sample
        CL.setExceptionsEnabled(true);

// Obtain the number of platforms
int numPlatformsArray[] = new int[1];
        CL.clGetPlatformIDs(0, null, numPlatformsArray);
int numPlatforms = numPlatformsArray[0];

// Obtain a platform ID
        cl_platform_id platforms[] = new cl_platform_id[numPlatforms];
        CL.clGetPlatformIDs(platforms.length, platforms, null);
        cl_platform_id platform = platforms[platformIndex];

// Initialize the context properties
        cl_context_properties contextProperties = new cl_context_properties();
        contextProperties.addProperty(CL.CL_CONTEXT_PLATFORM, platform);

// Obtain the number of devices for the platform
int numDevicesArray[] = new int[1];
        CL.clGetDeviceIDs(platform, deviceType, 0, null, numDevicesArray);
int numDevices = numDevicesArray[0];

// Obtain a device ID 
        cl_device_id devices[] = new cl_device_id[numDevices];
        CL.clGetDeviceIDs(platform, deviceType, numDevices, devices, null);
        cl_device_id device = devices[deviceIndex];

// Create a context for the selected device
        cl_context context = CL.clCreateContext(
            contextProperties, 1, new cl_device_id[]{device}, 
null, null, null);

// Create a command-queue for the selected device
        cl_command_queue commandQueue = 
            CL.clCreateCommandQueue(context, device, 0, null);

// Allocate the memory objects for the input and output data
        cl_mem memObjects[] = new cl_mem[3];
        memObjects[0] = CL.clCreateBuffer(context,
            CL.CL_MEM_READ_ONLY | CL.CL_MEM_COPY_HOST_PTR,
            Sizeof.cl_float * n, srcA, null);
        memObjects[1] = CL.clCreateBuffer(context,
            CL.CL_MEM_READ_ONLY | CL.CL_MEM_COPY_HOST_PTR,
            Sizeof.cl_float * n, srcB, null);
        memObjects[2] = CL.clCreateBuffer(context,
            CL.CL_MEM_READ_WRITE,
            Sizeof.cl_float * n, null, null);

// Create the program from the source code
        cl_program program = CL.clCreateProgramWithSource(context,
1, new String[]{ programSource }, null, null);

// Build the program
        CL.clBuildProgram(program, 0, null, null, null, null);

// Create the kernel
        cl_kernel kernel = CL.clCreateKernel(program, "sampleKernel", null);

// Set the arguments for the kernel
        CL.clSetKernelArg(kernel, 0,
            Sizeof.cl_mem, Pointer.to(memObjects[0]));
        CL.clSetKernelArg(kernel, 1,
            Sizeof.cl_mem, Pointer.to(memObjects[1]));
        CL.clSetKernelArg(kernel, 2,
            Sizeof.cl_mem, Pointer.to(memObjects[2]));

// Set the work-item dimensions
long global_work_size[] = new long[]{n};
long local_work_size[] = new long[]{1};

// Execute the kernel
        CL.clEnqueueNDRangeKernel(commandQueue, kernel, 1, null,
            global_work_size, local_work_size, 0, null, null);

// Read the output data
        CL.clEnqueueReadBuffer(commandQueue, memObjects[2], CL.CL_TRUE, 0,
            n * Sizeof.cl_float, dst, 0, null, null);

// Release kernel, program, and memory objects
        CL.clReleaseMemObject(memObjects[0]);
        CL.clReleaseMemObject(memObjects[1]);
        CL.clReleaseMemObject(memObjects[2]);
        CL.clReleaseKernel(kernel);
        CL.clReleaseProgram(program);
        CL.clReleaseCommandQueue(commandQueue);
        CL.clReleaseContext(context);

    }

private static String getString(cl_device_id device, int paramName) {
// Obtain the length of the string that will be queried
long size[] = new long[1];
        CL.clGetDeviceInfo(device, paramName, 0, null, size);

// Create a buffer of the appropriate size and fill it with the info
byte buffer[] = new byte[(int)size[0]];
        CL.clGetDeviceInfo(device, paramName, buffer.length, Pointer.to(buffer), null);

// Create a string from the buffer (excluding the trailing \0 byte)
return new String(buffer, 0, buffer.length-1);
    }
}

这段代码看起来不像Java，但实际上是Java。下面我将解释代码;现在不要在上面花太多时间，因为我将很快讨论不那么复杂的解决方案。

代码有很好的文档记录，但是让我们做一个小的演练。如您所见，代码非常类似于c语言。这很正常，因为JOCL只是到OpenCL的绑定。一开始，在一个字符串中有一些代码，而这些代码实际上是最重要的部分:它被OpenCL编译，然后发送到显卡并在那里执行。这段代码称为内核。不要将这个术语与操作系统内核混淆;这是设备代码。这个内核代码是在C的一个子集中编写的。

内核之后是Java绑定代码，用于设置和编排设备、对数据进行块处理、在将要存储数据的设备上创建适当的内存缓冲区以及用于生成数据的内存缓冲区。

总而言之:有“主机代码”(通常是语言绑定(在本例中是Java))和“设备代码”。“你总是能区分什么运行在主机上，什么应该运行在设备上，因为主机控制着设备。

上述代码应该被视为GPU等效的“Hello World!”如你所见。如果硬件支持SIMD扩展，则可以使算术代码运行得更快。例如，让我们看看矩阵乘法的核代码。这是Java应用程序的原始字符串中的代码。

__kernel void MatrixMul_kernel_basic(int dim,
                  __global float *A,
                  __global float *B,
                  __global float *C){

int iCol = get_global_id(0);
int iRow = get_global_id(1);
float result = 0.0;
for(int i=0; i< dim; ++i)
    {
          result +=
          A[iRow*dim + i]*B[i*dim + iCol];
    }
    C[iRow*dim + iCol] = result;
}

从技术上讲，这段代码将处理OpenCL框架为您设置的数据块，以及您在准备仪式中提供的说明。如果您的视频卡支持SIMD指令，并能够处理四个浮点数的向量，一个小的优化可能会把以前的代码变成以下代码:

#define VECTOR_SIZE 4    
__kernel void MatrixMul_kernel_basic_vector4(size_t dim, // dimension is in single floatsconst float4 *A,const float4 *B,
    float4 *C){
size_t globalIdx = get_global_id(0);
size_t globalIdy = get_global_id(1);
    float4 resultVec = (float4){ 0, 0, 0, 0 };
size_t dimVec = dim / 4;
for(size_t i = 0; i < dimVec; ++i) {
        float4 Avector = A[dimVec * globalIdy + i];
        float4 Bvector[4];
        Bvector[0] = B[dimVec * (i * 4 + 0) + globalIdx];
        Bvector[1] = B[dimVec * (i * 4 + 1) + globalIdx];
        Bvector[2] = B[dimVec * (i * 4 + 2) + globalIdx];
        Bvector[3] = B[dimVec * (i * 4 + 3) + globalIdx];
        resultVec += Avector[0] * Bvector[0];
        resultVec += Avector[1] * Bvector[1];
        resultVec += Avector[2] * Bvector[2];
        resultVec += Avector[3] * Bvector[3];
    }

    C[dimVec * globalIdy + globalIdx] = resultVec;
}

使用这段代码，可以将性能提高一倍。

您已经为Java世界解锁了GPU !但作为一名Java开发人员，您真的想要完成所有这些绑定、编写C代码并处理这些低级细节吗?我当然不喜欢。但是现在你已经对如何使用GPU架构有了一些了解，让我们看看我刚才介绍的JOCL代码之外的其他解决方案。

CUDA 和 Java

CUDA是Nvidia解决这些编码问题的方案。CUDA为标准的GPU操作提供了更多现成的库，比如矩阵、直方图，甚至是深度神经网络。新兴的库列表已经包含了许多有用的绑定。这些来自JCuda项目:http://www.jcuda.org/

• JCublas:关于矩阵

• JCufft:快速傅里叶变换

• JCurand:都是随机数

• JCusparse:稀疏矩阵

• JCusolver:分解

• JNvgraph:都是关于图形的

• JCudpp:CUDA数据并行原语库和一些排序算法

• JNpp:GPU上的图像处理

• JCudnn:深度神经网络库

我将使用JCurand(http://www.jcuda.org/jcuda/jcurand/JCurand.html)进行描述，它会生成随机数。您可以直接从Java代码中使用它，而不需要使用其他特定的内核语言。例如:

int n = 100;
    curandGenerator generator = new curandGenerator();
float hostData[] = new float[n];
    Pointer deviceData = new Pointer();
    cudaMalloc(deviceData, n * Sizeof.FLOAT);
    curandCreateGenerator(generator, CURAND_RNG_PSEUDO_DEFAULT); 
    curandSetPseudoRandomGeneratorSeed(generator, 1234);
    curandGenerateUniform(generator, deviceData, n);
    cudaMemcpy(Pointer.to(hostData), deviceData, 
            n * Sizeof.FLOAT, cudaMemcpyDeviceToHost);
    System.out.println(Arrays.toString(hostData));
    curandDestroyGenerator(generator);
    cudaFree(deviceData);

这里GPU是用来创建更多的高质量随机数，基于一些非常强大的数学。在JCuda中，您还可以编写通用的CUDA代码，并通过在类路径中添加一些JAR文件从Java调用它。更多示例请参见JCuda文档(http://www.jcuda.org/jcuda/JCuda.html)。

保持在低级别代码之上

所有这些看起来都很好，但是有太多的繁文缛节、太多的设置和太多不同的语言来运行它。有没有办法至少部分地使用GPU ?

如果你不想考虑所有这些OpenCL, CUDA，和其他内部的东西呢?如果您只想用Java编写代码，而不考虑内部内容，那该怎么办呢?Aparapi项目可以提供帮助。Aparapi代表“并行API”。“我认为它是一种Hibernate的GPU编程使用OpenCL的外壳。让我们看一个向量加法的例子。

public static void main(String[] _args) {
final int size = 512;
final float[] a = new float[size];
final float[] b = new float[size];

/* fill the arrays with random values */
for (int i = 0; i < size; i++){
        a[i] = (float) (Math.random() * 100);
        b[i] = (float) (Math.random() * 100);
    }
final float[] sum = new float[size];

    Kernel kernel = new Kernel(){
@Override public void run() {
I           int gid = getGlobalId();
            sum[gid] = a[gid] + b[gid];
        }
    };

    kernel.execute(Range.create(size));
for(int i = 0; i < size; i++) {
        System.out.printf("%6.2f + %6.2f = %8.2f\n", a[i], b[i], sum[i])
    }
    kernel.dispose();
}

这是纯Java代码(取自Aparapi文档)，不过您可以在这里或那里发现一些特定于GPU领域的术语，如Kernel和getGlobalId。您仍然需要了解GPU是如何编程的，但是您可以以一种对java更友好的方式来处理GPGPU。此外，Aparapi提供了一种简单的方法来将OpenGL上下文绑定到底层的OpenCL层——从而使数据完全保留在显卡上——从而避免了内存延迟问题。

如果需要进行许多独立的计算，考虑Aparapi。这个丰富的示例集提供了一些非常适合大规模并行计算的用例。

此外，还有几个项目，如tornado ovm，可以自动将适当的计算从CPU转移到GPU，从而实现了开箱即用的大规模优化。

总结

虽然在很多应用中gpu可以带来一些改变游戏规则的好处，但你可能会说仍然存在一些障碍。然而，Java和gpu可以一起做很多事情。在本文中，我只触及了这个庞大主题的表面。我的目的是展示从Java访问GPU的各种高级和低级选项。探索这个领域将带来巨大的性能优势，特别是对于需要大量计算、可以并行执行的复杂问题。