【代码札记】LWJGL's OpenCL HelloWorld

前几篇文章中简述了 OpenGL 的入门过程，本文将简述 LWJGL 提供的 OpenCL 绑定下，OpenCL 的入门过程。

OpenCL 展开目录

OpenGL 中的 GL 表示 Graphics Library，而 OpenCL 则是用于通用计算的：Open Computing Language。早期人们使用显卡进行图形渲染，但后来人们逐渐意识到，显卡上每一个运行着色器的单元就如同一个小 CPU—— 它不如正经 CPU 那么复杂，频率不快，可用内存也少，还缺乏高级的分支预测和乱序执行能力，但它数量多啊。以我新换的雷蛇笔记本为例：CPU 是英特尔的 i7-10875H，8 核 16 线程，如果在 Java 中处理 CPU 密集型的任务，一般会选择 17 个线程来执行，如果是 I/O 密集型任务，CPU 空闲时间比较多，则选择 33 个线程来执行。但是 GPU 呢？对于 RTX3070 来说，它有 5120 个 CUDA 核心，这意味着原则上我们可以同时在 GPU 中执行 5120 个代码，即 5120 个线程。虽然 CPU 能跑到 5.1GHz，而 GPU 最高 1.62GHz，但在需要高并发的任务中，GPU 确实脱颖而出。

前几天我在 Youtube 上看到一个视频：「Fully OpenCL accelerated space battle. (rts rpg game)」，视频作者利用 OpenCL 加速游戏。这个游戏是一个类似太空大战的游戏，是建立在二维平面上的，每艘船的行为使用 OpenCL 得出，而不再使用 CPU 遍历处理。这样得到的性能提升是什么呢？可以实时模拟接近 7 万个实体，游戏没有任何卡顿。

但使用 OpenCL 是有代价的。OpenCL 最初设计的目标并非只是为了在 GPU 上运算，实际上 OpenCL 的目标是异构计算，即同一段代码可以在 CPU、GPU、FPGA 中的任意一种上直接运行，而不必做出修改。这意味着 OpenCL 必须在设计上保证不依赖某一种平台的结构或特性，而提供一种通用的计算结构，因此我们平时习以为常的 CPU 内存结构和计算结构，在 OpenCL 中将不复存在。

另一方面说，GPU 虽然运算快，但是把数据从 CPU 内存送到 GPU 内存并不快。普通游戏显卡通常用于渲染游戏画面，即 CPU 把当前画面的数据送给显卡，显卡进行渲染，随后 CPU 送来新的一帧数据，数据是只进不出的；而为了实现通用计算，我们必须在运算结束后把数据从 GPU 内存读回 CPU 内存，游戏卡并非做不到，但没有经过特殊的设计与优化，会对速度造成一些影响，这一点在文章最后将代码跑起来后能够有所体现。

LWJGL's OpenCL 展开目录

由于 OpenCL 和 OpenGL 一样是用 C 写的，因此 Java 并不能直接调用他们。于是为了拉齐 Java 与 C 之间的鸿沟，和 OpenGL 一样，LWJGL 也提供了 OpenCL 的绑定，使得我们可以像在 C 语言中那样直接调用对应的函数就可以完成对应的功能。值得一提的是，LWJGL 的 OpenCL 绑定并非最好，也非唯一。实际上隔壁 JOCL 库做的更加面向对象一些，用起来也更加舒服，而且文档也很丰富。目前 LWJGL3 看起来正在重构，OpenCL 没有任何文档，只能去他们的 Github Repo 中读一些测试代码，或者参考 OpenCL 在 C 语言中的用法，结合 OpenGL 绑定中与 C 交互的部分去猜测代码应该如何编写。

所以我的建议是，如果不使用 LWJGL 提供的 OpenGL 绑定而只需要使用 OpenCL，取决于你是否对编写 C 语言内核（即着色器，实际跑在 GPU 上的代码）心存芥蒂，你可以使用 JOCL 来更加面向对象的加载你编写的 C 内核；或者直接使用 Aparapi 库将你的 Java 代码自动翻译成 OpenCL 需要的形式。

本文将使用 LWJGL 提供的 OpenCL 绑定，这个绑定更接近原始 C 接口，用起来更加复杂一些。如无意外，在未来我将使用 JOCL 进行 OpenCL 计算，届时会有对应的文章发表。

Hello, World 展开目录

闲话少说，书归正传。要使用 LWJGL 的 OpenCL 绑定，只需要在其官网配置项目的时候选定 OpenCL 就可以了。

OpenCL 内核展开目录

就一般的开发流程而言，应当是先搭出来程序的架子，然后再去谈什么样的代码跑在 GPU 上，但是出于演示的目的，这里先给出 OpenCL 内核：

__kernel void demo(__global float *a, __global float *b, __global float *answer) {
    int xid = get_global_id(0);
    answer[xid] = cos(a[xid]) + sin(b[xid]);
}

内核是同时运行在每一个 CUDA 核心上的程序，因此我们并不能像 CPU 那样使用循环来完成两个数组的按项求和。正相反，我们要展开这个操作：每一个 Cuda 核心只负责将数组中的一个索引位置上的数据相加。于是对于每一个核心，我们先要获取 OpenGL 给这个核心的编号，然后再根据这个编号将对应位置上的数据相加。为了增加操作难度，我还在两个位置上套了个 Sin 和 Cos，以此来拉开 GPU 和 CPU 的表现。

关于核心编号的安排，OpenCL 提供了最多三维的组织方式。这里 get_global_id(0) 是取第一维的编号，类似于 x，取决于我们在 Java 代码中如何编写，根据处理数据的不同，我们可以给核心按 (x, y) 方式编号，此时我们利用 0 来取 x，1 取 y；或者按 (x, y, z) 编号。至于运算所需的数据，我们定义两个存储于全局内存（GPU 中最慢的内存）的两个浮点型数组 a 和 b 作为要相加的数据，将计算结果写回存储于全局内存中的 answer 数组。

准备之准备展开目录

在正式准备 OpenCL 的运行时环境之前，我们得先知道系统中有哪些 OpenCL 设备可供使用。在双显卡环境中，通常是 CPU 的核显提供一个基于 Intel 平台的 OpenCL 设备，英伟达的独显提供一个基于英伟达平台的设备。为了查询系统中有哪些 OpenCL 设备，并读取他们相关的信息，有如下辅助函数：

fun getPlatformInfoStringUTF8(cl_platform_id: Long, param_name: Int): String {
    stackPush().use { stack ->
        val returnValueSize = stack.mallocPointer(1)
        // query the return value size first
        checkCLError(clGetPlatformInfo(cl_platform_id, param_name, null as ByteBuffer?, returnValueSize))
        val bytes = returnValueSize[0].toInt()
        val buffer = stack.malloc(bytes)
        // query the value
        checkCLError(clGetPlatformInfo(cl_platform_id, param_name, buffer, returnValueSize))
        return memUTF8(buffer, bytes - 1)
    }
}

fun getDeviceInfoStringUTF8(cl_device_id: Long, param_name: Int): String {
    stackPush().use { stack ->
        val returnValueSize = stack.mallocPointer(1)
        checkCLError(clGetDeviceInfo(cl_device_id, param_name, null as ByteBuffer?, returnValueSize))
        val bytes = returnValueSize[0].toInt()
        val buffer = stack.malloc(bytes)
        checkCLError(clGetDeviceInfo(cl_device_id, param_name, buffer, returnValueSize))
        return memUTF8(buffer, bytes - 1)
    }
}

fun getDeviceInfoInt(cl_device_id: Long, param_name: Int): Int {
    stackPush().use { stack ->
        val pl = stack.mallocInt(1)
        checkCLError(clGetDeviceInfo(cl_device_id, param_name, pl, null))
        return pl[0]
    }
}

fun getDeviceInfoLong(cl_device_id: Long, param_name: Int): Long {
    stackPush().use { stack ->
        val pl = stack.mallocLong(1)
        checkCLError(clGetDeviceInfo(cl_device_id, param_name, pl, null))
        return pl[0]
    }
}

fun checkCLError(code: Int) {
    require(code == CL_SUCCESS) { "OpenCL error code: $code" }
}

fun checkCLError(code: IntBuffer) {
    checkCLError(code[code.position()])
}

其中 getxxxStringUTF8 的函数负责将 C 返回的字符串数据翻译成 Java 能够识别的 UTF8 字符串。这个过程需要两次查询：第一次查询返回数据的大小，然后创建对应的缓存区域，第二次查询才将真正的字符串数据写入缓存，然后 Java 代码将缓存中的数据转换成字符串。

而 getxxxInt 或类似的函数则简单许多了，知道返回值的类型（Int 或 Long），则可以一次查询获取到返回值。最后的 checkCLError 则是检查返回值是否标识成功。OpenCL 的返回值一般有两种返回方式。一些函数将返回值作为函数调用的返回值，另一些函数将返回值作为函数参数以缓冲区的方式传入，函数执行完毕时返回值将写入缓冲区。

除了能够 “听懂” OpenCL 返回的数据之外，我们还得自己定义一下 OpenCL 设备的描述：

data class CLDeviceInfo(
    val platformId: Long,
    val deviceId: Long,
    val deviceType: Long,
    val deviceName: String,
    val computeUnitCount: Int,
    val clockFrequency: Int,
    val localMemorySize: Long,
    val globalMemorySize: Long
)

OpenCL 设备按照平台定义，每个平台具有唯一编号，而平台内的设备也具有唯一的编号。除此之外我们还将查询设备的类型（CPU、GPU 或者其他硬件）、设备的名称（便于我们识别）、最大的运算核心数量（将依据此选择最佳设备），除此之外是一些演示用的信息，比如时钟频率、局部内存大小和全局内存大小。

如下函数将查询系统中所有的 OpenCL 设备，并择其中运算核心数量最大的一个返回给调用者：

/**
 * Choose best device by compute unit count.
 * */
fun getBestOpenCLDevice(): CLDeviceInfo {
    val deviceList = mutableListOf<CLDeviceInfo>()

    stackPush().use { stack ->
        // get total count
        val clCounter = stack.mallocInt(1)
        checkCLError(clGetPlatformIDs(null, clCounter))
        val platforms: PointerBuffer = stackPush().mallocPointer(clCounter[0])
        checkCLError(clGetPlatformIDs(platforms, clCounter))

        for (platformIndex in 0 until platforms.capacity()) {
            val platformId = platforms[platformIndex]
            // query count
            checkCLError(clGetDeviceIDs(platformId, CL_DEVICE_TYPE_ALL.toLong(), null, clCounter))
            // query value
            val devices = stack.mallocPointer(clCounter[0])
            checkCLError(clGetDeviceIDs(platformId, CL_DEVICE_TYPE_ALL.toLong(), devices, clCounter))

            for (deviceIndex in 0 until devices.capacity()) {
                val deviceId = devices[deviceIndex]
                deviceList.add(
                    CLDeviceInfo(
                        platformId = platformId,
                        deviceId = deviceId,
                        deviceType = getDeviceInfoLong(deviceId, CL_DEVICE_TYPE),
                        deviceName = getDeviceInfoStringUTF8(deviceId, CL_DEVICE_NAME),
                        computeUnitCount = getDeviceInfoInt(deviceId, CL_DEVICE_MAX_COMPUTE_UNITS),
                        clockFrequency = getDeviceInfoInt(deviceId, CL_DEVICE_MAX_CLOCK_FREQUENCY),
                        localMemorySize = getDeviceInfoLong(deviceId, CL_DEVICE_LOCAL_MEM_SIZE),
                        globalMemorySize = getDeviceInfoLong(deviceId, CL_DEVICE_GLOBAL_MEM_SIZE)
                    )
                )
            }
        }
    }

    return deviceList.maxByOrNull { it.computeUnitCount }!!
}

如同之前所说的，我们先查询系统中一共有多少个平台，然后创建对应大小的缓冲区，再读取每一个平台的编号。遍历平台编号，查询平台下的每一个设备，这里也是同理：先查询多少个设备，再根据大小创建缓冲区读入每个设备的编号。最终对每个设备我们需要查询出我们需要的信息。全部查询完毕后，我们选出核心数量最多的设备作为本例使用的设备。至此，准备之准备可以告一段落，开始着手 OpenCL 运行时的准备了。

准备数据展开目录

在主函数一开头我们需要准备本次计算需要的数据。由于 GPU 计算浮点型只能达到单精度（高贵的 Quadro 和 Telsa 用户请无视这句话），所以给 GPU 的数据应该是浮点型的，而给 CPU 的数据要再转换成双精度，因为 Java 中的 cos 和 sin 函数都是在汇编指令（指 JVM Bytecode）级别优化的，因此只接受双精度浮点，又由于将单精度转换为双精度耗时相对更多，所以转换用时自然不能算在 CPU 时间里，需要我们提前准备好：

// GPU
    val n = 30000000
    val inputA = FloatArray(n) { (Random.nextFloat() * 2 * Math.PI).toFloat() }
    val inputB = FloatArray(n) { (Random.nextFloat() * 2 * Math.PI).toFloat() }
    val answer = FloatArray(n)
    // also choose the device
    val clDeviceInfo = getBestOpenCLDevice()
// CPU
    val a = DoubleArray(n) { inputA[it].toDouble() }
    val b = DoubleArray(n) { inputB[it].toDouble() }
    val c = DoubleArray(n)

我们准备了三千万个数据。

准备施法展开目录

数据已经准备妥当，现在可以专注于 OpenCL 的吟唱准备工作了。

首先我们最好拿到一个 MemoryStack，它是 LWJGL 提供的便于内存管理的类。利用它我们可以方便的创建缓冲区之类的对象：

    val stack = stackPush()
    val errCodeRet = stack.mallocInt(1)

这里我们利用 MemoryStack 对象创建了一个大小为 1 个整型的缓冲区，用于接收前面说的以参数形式将返回值写入缓冲区的函数的返回值。

和创建 OpenGL 窗口一样，我们不能上来就创建窗口，得先设置一些 Hint。类似地，OpenCL 需要 Context Properties：

    // Specifies a list of context property names and their corresponding values.
    // Each property name is immediately followed by the corresponding desired value.
    // The list is terminated with 0
    val contextProperties = stack.mallocPointer(3)
    contextProperties.put(0, CL_CONTEXT_PLATFORM.toLong())
    contextProperties.put(1, clDeviceInfo.platformId)
    contextProperties.put(2, 0L)

这里具有误导性的地方在于 mallocPointer，看名字可能以为是声明一个用来存放指针的缓冲区，但实际上指针就是 Long 型，所以将里面的数据视为长整型数据即可。唯一的区别就是这个函数返回一个 PointerBuffer，本意是将内部的数据作为指针数据解释；而 mallocLong 返回 LongBuffer，它的本意是将内部的数据作为长整型解读。

根据 OpenCL 官方手册，Context Properties 是一个列表。每一个属性的名字后面立刻跟着它对应的值，并最终以一个 0 结尾。这里我们要确定所使用的平台编号，之后的 OpenCL 将面向这个平台执行操作（不同的平台编译后产生的二进制将不一样）。

随后我们就可以创建上下文对象。上下文对象的作用是管理程序运行所需的数据。假如你一边玩 Minecraft 一边运行程序，显卡同一时间只能干同一件事，我们的程序将和 Minecraft 争抢显卡的使用权，如果我们的程序刚刚从 MC 那里抢过显卡，显卡里还是 MC 的数据，我们直接拿来做运算，结果一定是错的。因此 OpenCL 将根据上下文设置 GPU 内存中的数据，让我们无需关心数据的调度问题。创建上下文的代码如下：

    var contextCallback: CLContextCallback?
    val context: Long = clCreateContext(
        contextProperties,
        clDeviceInfo.deviceId,
        CLContextCallback.create { errInfo: Long, private_info: Long, _: Long, _: Long ->
            System.err.println("[LWJGL] cl_context_callback")
            System.err.println("\tInfo: " + memUTF8(errInfo))
            System.err.println("\tPrivateInfo: " + memUTF8(private_info))
        }.also { contextCallback = it },
        NULL, // no user data
        errCodeRet
    )
    checkCLError(errCodeRet)

这里有一个回调函数，将在创建失败时被调用。这里函数的返回值将写入到缓冲区中，而函数调用的返回值则给出上下文对象的编号。我们可以通过观察回调函数是否被调用知道上下文是否创建失败，但程序需要通过检查返回值确保继续前进之前，我们成功创建了上下文对象。

随后我们创建命令队列。由于命令在 GPU 上的执行独立于程序在 CPU 上执行，因此我们需要将命令放入队列中供 GPU 慢慢执行。

    // create command queue
    val commandQueue = clCreateCommandQueue(context, clDeviceInfo.deviceId, 0, errCodeRet)
    checkCLError(errCodeRet)

有了命令队列之后，我们就可以将内核需要的数据送给 GPU 了。

    val inputABuffer: Long =
        clCreateBuffer(context, CL_MEM_READ_ONLY.toLong(), n.toLong() * Float.SIZE_BYTES, errCodeRet)
    checkCLError(errCodeRet)
    checkCLError(
        clEnqueueWriteBuffer(
            commandQueue, inputABuffer,
            true, 0, inputA,
            null as PointerBuffer?, null as PointerBuffer?
        )
    )

    val inputBBuffer: Long =
        clCreateBuffer(context, CL_MEM_READ_ONLY.toLong(), n.toLong() * Float.SIZE_BYTES, errCodeRet)
    checkCLError(errCodeRet)
    checkCLError(
        clEnqueueWriteBuffer(
            commandQueue, inputBBuffer,
            true, 0, inputB,
            null as PointerBuffer?, null as PointerBuffer?
        )
    )
    val answerBuffer: Long =
        clCreateBuffer(context, CL_MEM_WRITE_ONLY.toLong(), n.toLong() * Float.SIZE_BYTES, errCodeRet)
    checkCLError(errCodeRet)

我们依次针对两个输入和一个输出创建对应的缓冲区。这里和 OpenGL 不同，OpenCL 的缓冲区大小的单位是字节，而非元素个数。我们在创建缓冲区的时候还将指定它的读写属性，对于两个输入，我们想要只读，而对于输出，我们要只写。

随后针对两个输入缓冲区，我们要将数组数据写入进去：

        clEnqueueWriteBuffer(
            commandQueue, inputABuffer,
            true, 0, inputA,
            null as PointerBuffer?, null as PointerBuffer?
        )

这个函数需要指定命令队列和要写入的缓冲区。第三个参数表示是否为同步操作。为 True 表示函数将一直阻塞直到操作完成，如果我们使用数组传入数据，我们必须使用同步操作，而如果从各种 Buffer 中写入数据，则可以选为异步操作，GPU 将在其他线程拷贝数据。如果使用异步操作，我们可以在进行运算之前调用 clFinish(commandQueue) 确保开始运算之前所有操作已经就绪，该函数将一直阻塞直到命令队列全部执行完毕。

数据准备妥当之后，我们可以着手编译内核了：

    // compile program
    val program = clCreateProgramWithSource(context, readResourceFileContent("/kernels/demo.cl"), errCodeRet)
    checkCLError(errCodeRet)
    checkCLError(clBuildProgram(program, clDeviceInfo.deviceId, "", null, NULL))

    // build kernel
    val kernel = clCreateKernel(program, "demo", errCodeRet)
    checkCLError(errCodeRet)

首先从 resources 中读取源代码，然后以此创建一个程序，并编译之。编译成功后创建内核，创建内核的第二个参数是源码中的函数名。

这里还需要一个额外的步骤就是将前面创建的 CLBuffer 包装成 PointerBuffer：

    val inputAPointer = stack.mallocPointer(1)
    inputAPointer.put(0, inputABuffer)
    val inputBPointer = stack.mallocPointer(1)
    inputBPointer.put(0, inputBBuffer)
    val answerPointer = stack.mallocPointer(1)
    answerPointer.put(0, answerBuffer)

然后利用这些参数设置内核工作需要的参数：

    // Set the arguments of the kernel
    checkCLError(clSetKernelArg(kernel, 0, inputAPointer))
    checkCLError(clSetKernelArg(kernel, 1, inputBPointer))
    checkCLError(clSetKernelArg(kernel, 2, answerPointer))

这里的 0、1、2 分别表示内核程序的第一个参数、第二个参数和第三个参数。如果你不将 CLBuffer 包装成 PointerBuffer 而直接作为 clSetKernelArg 的第三个参数传入，那么它将被解读成「参数大小」。所以为了将其作为参数所在的缓冲区传入内核，必须将他们包装成 PointerBuffer。别说文档了，LWJGL 关于 OpenCL 连个 JavaDoc 都没有，我真服了。这块坑了我好一阵子。

最后我们等待全部操作完成，准备计算：

    // wait all memory copy is done
    clFinish(commandQueue)

开始施法展开目录

与其说通过一个函数开始计算，不如说开始计算的过程更像是一个安排计算方式的过程。

还记得一开始我说 OpenCL 提供多种核心的安排方式吗，一维的、二维的、三维的？在这里一步，我们将通过定义 NDRange 来开始计算。整体代码如下：

    // execute kernel
    val globalWorkerOffset = stack.mallocPointer(1)
    globalWorkerOffset.put(0, NULL)
    val globalWorkerCount = stack.mallocPointer(1)
    globalWorkerCount.put(0, n.toLong())

    val localWorkerCount = stack.mallocPointer(1)
    localWorkerCount.put(0, 64L)
    checkCLError(
        clEnqueueNDRangeKernel(
            commandQueue, kernel, 1,
            globalWorkerOffset, globalWorkerCount, localWorkerCount,
            null as PointerBuffer?, null as PointerBuffer?
        )
    )

其中最为关键的还是 clEnqueueNDRangeKernel 函数。首先我们需要确定在哪一个命令队列上运行哪一个内核，因此我们将 commandQueue 和 kernel 传进去。其次我们还想告诉 OpenCL 以几维的方式安排核心，这里我们选 1 维。如果是二维的话就写 2，最大是 3 维。下一个参数取决于设备支持的 OpenCL 版本，在 OpenCL 版本中它必须为 NULL，而在 2.0 及以上版本中可以指定全局编号的开始位置，如果是 NULL 的话就从 0 开始（对应二维的 (0, 0) 和三维的 (0, 0, 0,)），否则按照你传入的这 1/2/3 个数作为起点开始分配全局编号。

然后是全局编号的总大小，即我们一共需要多少个核心来参与运算。这个数量是多维的，所以可以按照维数的不同精确控制唯一维度上核心的数量。

下一个参数是局部核心数量，局部核心数量的定义与全局核心数量一样是多维的。局部核心数量决定了一个 Worker Group（核心构成的组）的大小，OpenCL 中是以 Worker Group 为单位进行计算的：要么整个组一起在设备上运算，要不谁也不能单独运算。这个大小的选择一般根据 clinfo 给出的 Preferred work group size multiple (kernel) 和 Max work group size 及全局核心数量得出的。其中 Preferred work group size multiple (kernel) 给出这个设备更倾向于将 Work Group 的数量定位谁的倍数，我这里给出的是 32，因此将其定义为 32 的倍数会好一些；Max work group size 则对设备最大的 Work Group 大小做出的限制，太大了显卡装不下，我这里的限制是 1024。最后一个条件最为关键：Work Group 的数量必须能整除全局的核心数（在对应维度上），否则将会得到 -54 错误：CL_INVALID_WORK_GROUP_SIZE。

最后两个参数是定义哪些事件必须在计算之前完成，由于我们之前显示调用了 clFinish，所以不必理会这两个参数，直接传入 null。

这条命令一结束，GPU 便开始执行运算。由于 clEnqueueNDRangeKernel 函数是非阻塞的，所以函数返回时并不保证运算完毕，因此在读回结果之前我们还得调用 clFinish 来确保运算完成。

取回数据展开目录

前面说过，和渲染画面不同，我们做通用计算是要拿回运算结果的。运算完毕后结果存在 GPU 内存中，我们需要将其读回 CPU 内存才能继续处理。和写入类似，读取的操作是这样：

    // read out result to answer
    checkCLError(
        clEnqueueReadBuffer(
            commandQueue, answerBuffer,
            true, 0, answer,
            null as PointerBuffer?, null as PointerBuffer?
        )
    )

我们将缓冲区的数据读回数组 answer 里。

清理展开目录

运算结果有了，但是先前创建的那些缓冲区什么的并不能置之不理，否则随着程序的运行将会产生内存泄漏。好在清理并不难，注意顺序：最先创建的最后释放。

    // clean up
    checkCLError(clFlush(commandQueue))
    checkCLError(clFinish(commandQueue))
    checkCLError(clReleaseKernel(kernel))
    checkCLError(clReleaseProgram(program))
    checkCLError(clReleaseMemObject(inputABuffer))
    checkCLError(clReleaseMemObject(inputBBuffer))
    checkCLError(clReleaseMemObject(answerBuffer))
    checkCLError(clReleaseCommandQueue(commandQueue))
    checkCLError(clReleaseContext(context))
    contextCallback?.free()
    CL.destroy()

CPU 性能对比展开目录

本节编写对应的 CPU 代码并与 GPU 性能进行对比：

    for (i in inputA.indices) {
        c[i] = cos(a[i]) + sin(b[i])
    }

这是内核对应的等效的代码。我们在 val stack = stackPush() 之前获取当前系统时间，在 clEnqueueNDRangeKernel 之前再获取一次系统时间，然后在 clEnqueueReadBuffer 之后获取一次系统时间。这样就能够通过时间差算出来从创建上下文环境到最后读出运算结果的总耗时，以及其中从开始运算到读回结果的耗时，由此我们能更加直观的感受将数据从 CPU 送到 GPU 的耗时。

同样的我们在 for 循环之前获取系统时间，再在循环结束后获取一次系统时间，得到 CPU 的纯运算时间，因为数据就在内存里，所以不需要将他们复制到其他地方。

最终得到的数据是这样的：

设备	耗时（ms）
CPU	641
GPU 全程	322
GPU 从开始运算计时	79

可见 GPU 全称耗时 322 毫秒，而真正的计算加上取回数据只用了 79 毫秒，剩下的二百多毫秒全部用在数据传送以及 Java 和 C 的交互上。而 CPU 是单线程运算，耗时 641 毫秒。虽然数据按项相加并不是什么露脸的多线程任务，但毫无疑问作为一个入门程序，它体现出了 GPU 在并行计算上的潜力。

主函数代码展开目录

package info.skyblond.antSimu

import info.skyblond.antSimu.opencl.checkCLError
import info.skyblond.antSimu.opencl.getBestOpenCLDevice
import info.skyblond.antSimu.utils.ResourceUtils.readResourceFileContent
import org.lwjgl.PointerBuffer
import org.lwjgl.opencl.CL
import org.lwjgl.opencl.CL10.*
import org.lwjgl.opencl.CLContextCallback
import org.lwjgl.system.MemoryStack.stackPush
import org.lwjgl.system.MemoryUtil.NULL
import org.lwjgl.system.MemoryUtil.memUTF8
import kotlin.math.cos
import kotlin.math.sin
import kotlin.random.Random


fun main() {
    // data set
    val n = 30000000
    val inputA = FloatArray(n) { (Random.nextFloat() * 2 * Math.PI).toFloat() }
    val inputB = FloatArray(n) { (Random.nextFloat() * 2 * Math.PI).toFloat() }
    val answer = FloatArray(n)

    // choose best device by core count
    val clDeviceInfo = getBestOpenCLDevice()

    val gpuStartTime = System.currentTimeMillis()
    val stack = stackPush()

    // Specifies a list of context property names and their corresponding values.
    // Each property name is immediately followed by the corresponding desired value.
    // The list is terminated with 0
    val contextProperties = stack.mallocPointer(3)
    contextProperties.put(0, CL_CONTEXT_PLATFORM.toLong())
    contextProperties.put(1, clDeviceInfo.platformId)
    contextProperties.put(2, 0L)

    val errCodeRet = stack.mallocInt(1)

    var contextCallback: CLContextCallback?
    val context: Long = clCreateContext(
        contextProperties,
        clDeviceInfo.deviceId,
        CLContextCallback.create { errInfo: Long, private_info: Long, _: Long, _: Long ->
            System.err.println("[LWJGL] cl_context_callback")
            System.err.println("\tInfo: " + memUTF8(errInfo))
            System.err.println("\tPrivateInfo: " + memUTF8(private_info))
        }.also { contextCallback = it },
        NULL, // no user data
        errCodeRet
    )
    checkCLError(errCodeRet)

    // create command queue
    val commandQueue = clCreateCommandQueue(context, clDeviceInfo.deviceId, 0, errCodeRet)
    checkCLError(errCodeRet)

    // n * sizeof(float) = n * 8
    val inputABuffer: Long =
        clCreateBuffer(context, CL_MEM_READ_ONLY.toLong(), n.toLong() * Float.SIZE_BYTES, errCodeRet)
    checkCLError(errCodeRet)
    checkCLError(
        clEnqueueWriteBuffer(
            commandQueue,
            inputABuffer,
            true,
            0,
            inputA,
            null as PointerBuffer?,
            null as PointerBuffer?
        )
    )

    val inputBBuffer: Long =
        clCreateBuffer(context, CL_MEM_READ_ONLY.toLong(), n.toLong() * Float.SIZE_BYTES, errCodeRet)
    checkCLError(errCodeRet)
    checkCLError(
        clEnqueueWriteBuffer(
            commandQueue,
            inputBBuffer,
            true,
            0,
            inputB,
            null as PointerBuffer?,
            null as PointerBuffer?
        )
    )
    val answerBuffer: Long =
        clCreateBuffer(context, CL_MEM_WRITE_ONLY.toLong(), n.toLong() * Float.SIZE_BYTES, errCodeRet)
    checkCLError(errCodeRet)

    // compile program
    val program = clCreateProgramWithSource(context, readResourceFileContent("/kernels/demo.cl"), errCodeRet)
    checkCLError(errCodeRet)
    checkCLError(clBuildProgram(program, clDeviceInfo.deviceId, "", null, NULL))

    // build kernel
    val kernel = clCreateKernel(program, "demo", errCodeRet)
    checkCLError(errCodeRet)

    val inputAPointer = stack.mallocPointer(1)
    inputAPointer.put(0, inputABuffer)
    val inputBPointer = stack.mallocPointer(1)
    inputBPointer.put(0, inputBBuffer)
    val answerPointer = stack.mallocPointer(1)
    answerPointer.put(0, answerBuffer)

    // Set the arguments of the kernel
    checkCLError(clSetKernelArg(kernel, 0, inputAPointer))
    checkCLError(clSetKernelArg(kernel, 1, inputBPointer))
    checkCLError(clSetKernelArg(kernel, 2, answerPointer))

    // wait all memory copy is done
    clFinish(commandQueue)

    // execute kernel
    val globalWorkerOffset = stack.mallocPointer(1)
    globalWorkerOffset.put(0, NULL)
    val globalWorkerCount = stack.mallocPointer(1)
    globalWorkerCount.put(0, n.toLong())

    val localWorkerCount = stack.mallocPointer(1)
    localWorkerCount.put(0, 640L)

    val gpuComputingStartTime = System.currentTimeMillis()
    checkCLError(
        clEnqueueNDRangeKernel(
            commandQueue, kernel, 1,
            globalWorkerOffset, globalWorkerCount, localWorkerCount,
            null as PointerBuffer?, null as PointerBuffer?
        )
    )
    clFinish(commandQueue)

    // read out result to answer
    checkCLError(
        clEnqueueReadBuffer(
            commandQueue,
            answerBuffer,
            true,
            0,
            answer,
            null as PointerBuffer?,
            null as PointerBuffer?
        )
    )

    val gpuEndTime = System.currentTimeMillis()
    println("${gpuEndTime - gpuStartTime}ms consumed by GPU")
    println("${gpuEndTime - gpuComputingStartTime}ms for GPU computing")

    // clean up
    checkCLError(clFlush(commandQueue))
    checkCLError(clFinish(commandQueue))
    checkCLError(clReleaseKernel(kernel))
    checkCLError(clReleaseProgram(program))
    checkCLError(clReleaseMemObject(inputABuffer))
    checkCLError(clReleaseMemObject(inputBBuffer))
    checkCLError(clReleaseMemObject(answerBuffer))
    checkCLError(clReleaseCommandQueue(commandQueue))
    checkCLError(clReleaseContext(context))
    contextCallback?.free()
    CL.destroy()

    val a = DoubleArray(n) { inputA[it].toDouble() }
    val b = DoubleArray(n) { inputB[it].toDouble() }
    val c = DoubleArray(n)

    val cpuStartTime = System.currentTimeMillis()
    for (i in inputA.indices) {
        c[i] = cos(a[i]) + sin(b[i])
    }
    val cpuEndTime = System.currentTimeMillis()
    println("${cpuEndTime - cpuStartTime}ms for CPU computing")
    answer.forEachIndexed { index, fl ->
        require(c[index] - fl < 1e-6) { "WRONG! ${c[index]} $fl" }
    }
}

最后为了保证运算结果正确，还对 GPU 运算结果做了校验。

- 全文完 -

【代码札记】LWJGL's OpenCL HelloWorld 由天空 Blond 采用知识共享署名 - 非商业性使用 - 相同方式共享 4.0 国际许可协议进行许可。
本许可协议授权之外的使用权限可以从 https://skyblond.info/about.html 处获得。

Catalog