【代码札记】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 处获得。