前文记述了我使用Kotlin在JVM上实现了Conway's Game of Life，本篇将记述如何使用GPU加速游戏，并将将游戏的运行过程保存为视频。
上文结尾说我将使用HumbleVideo这个库来生成视频，这个库最后一次发布是2018年，并且看起来并不支持H265，但它的好处是原生支持BufferedIamge，我可以通过代码将WorldMap的当前状态渲染成BufferedImage并直接利用CPU编码成视频。但是这时候我发现，CPU好像有点不太够用啊？
尽管更新游戏状态和编码视频是交替进行的，但CPU归根结底只有16个线程。我想来想去，OpenCL不可避免。在之前一篇介绍OpenCL的文章中，我用OpenCL同时计算了25万个双摆，能够做到每秒更新1000次。倘若我用OpenCL计算这个生命游戏，岂不是能够快很多？
OpenCL加速
但是和双摆不同，在双摆计算中，所有参数是一个1维数组，也就是说我可以不必考虑因为设计而引入的限制——唯一的限制就是数组的宽度。而康威生命游戏则是一个二维的棋盘——好巧不巧，OpenCL并不支持二维数组。当然，我们可以把二维数组伸展成一维数组，但这样就不得不考虑数组的长度上限了。在C的世界中，数组并没有长度限制，程序的能力取决于size_t的大小，这通常是64位。在Java的世界中，数组只能被int索引，而int还是有符号的，也就是说，一个数组最多只能有2147483648个元素（0到2147483647）。听起来不小，哪怕光存储byte，这样一个数组也要至少占用2GB内存。但如果我们把二维数组伸展成一维数组，这一步在Java中就会出问题：我们没办法所引超过的部分：在二维数组中，我们可以索引2G个子数组，而每一个子数组可以有2G个元素，把他们展开成一维数组，我们只能展开总长度小于2G的二维数组，这相当于人为设置了一个上限。我认为这样并不优雅，也不好。
如果我们转换思路，问题就迎刃而解了：更新一个细胞无外乎需要8个状态，而这八个状态来自于与之相邻的8个细胞。如果我们按行输送数据，每次输入三行，输出一行，那么无论这个二维数组有多大，只要我们分多次将数据送进去，按行更新，这样就没有问题。每次传送数据耗时较大，为了追求效率，我们可以一次输入N+2行，输出N行。
不过还有一个问题：假如我们想要动态的调整输入输出，我们不光需要修改OpenCL Kernel的代码，还需要对应地修改程序的逻辑。虽然这个不难做到，但我尝试了半天，并没有找到一种很优雅的解决方案。所以这一块应该是一个小遗憾。
总的来说，OpenCL Kernel就相当于是把Java中更新游戏的代码用OpenCL C重写了一下。我了解了一下OpenCL C，它有一点很好：int类型就是32位的。但也有一点不好：64位的long并不保证在所有OpenCL设备上都受到支持。好在前文中实现WorldMap时使用了Int作为内部存储。并且Java的Int也是32位的。
首先当然是一些基本代码：
#define WORD_SIZE 32 // int has 32 bits
int get(__global const int *data, int x, int gameWidth) {
if (x < 0 || x >= gameWidth) return false;
int wordIndex = x / WORD_SIZE;
int wordOffset = x % WORD_SIZE;
int word = data[wordIndex];
return word & (1 << wordOffset);
}
void set(__global int *data, int x, int gameWidth) {
if (x < 0 || x >= gameWidth) return;
int wordIndex = x / WORD_SIZE;
int wordOffset = x % WORD_SIZE;
int word = data[wordIndex] | (1 << wordOffset);
data[wordIndex] = word;
}
void clear(__global int *data, int x, int gameWidth) {
if (x < 0 || x >= gameWidth) return;
int wordIndex = x / WORD_SIZE;
int wordOffset = x % WORD_SIZE;
int word = data[wordIndex] & ~(1 << wordOffset);
data[wordIndex] = word;
}
int countCell(
__global const int *top,
__global const int *middle,
__global const int *bottom,
int x, int gameWidth
) {
int result = 0;
// count y-1
if (get(top, x-1, gameWidth)) result++;
if (get(top, x, gameWidth)) result++;
if (get(top, x+1, gameWidth)) result++;
// count y
if (get(middle, x-1, gameWidth)) result++;
if (get(middle, x+1, gameWidth)) result++;
// count y+1
if (get(bottom, x-1, gameWidth)) result++;
if (get(bottom, x, gameWidth)) result++;
if (get(bottom, x+1, gameWidth)) result++;
return result;
}这些代码基本上和Java一样，无非是对一个Word进行操作，以及数格子。然而和Java不同的地方在于，OpenCL里面没办法上锁。如果我们让一个Kernel更新一个细胞，那么一个Word就会有32个Kernel同时更新，并且还是非原子性的更新，这样就会引发同步问题。为了解决这个问题，我们让一个Kernel独占地更新整个Word：
void calculateRow(
__global const int *top,
__global const int *middle,
__global const int *bottom,
__global int *out,
const int wordIndex,
const int gameWidth
) {
// here we will update the whole word, because of synchronization issues.
// if we update 1 cell/bit per kernel, the update is not atomic,
// thus the same word will conflict, generating bad result.
// by updating the whole word in one kernel, we avoid conflict.
int start = wordIndex * WORD_SIZE;
for (int x = start; x < start + WORD_SIZE; x++) {
if (x < 0 || x >= gameWidth) continue;
int cell = get(middle, x, gameWidth);
int count = countCell(top, middle, bottom, x, gameWidth);
if (cell) {
if (count == 2 || count == 3) {
set(out, x, gameWidth);
} else {
clear(out, x, gameWidth);
}
} else {
if (count == 3) {
set(out, x, gameWidth);
} else {
clear(out, x, gameWidth);
}
}
}
}之余整个Kernel的入口，由于传入和传出都是可变的，因此我采用了几个占位符：
__kernel void calculateNextTick(
%GENERATED_INPUT%
%GENERATED_OUTPUT%
const int gameWidth
) {
int idx = get_global_id(0);
%GENERATED_CALCULATION%
}其中%GENERATED_INPUT%会被替换成多个__global const int *row#,的拼接，其中#为行号，从1开始；类似地，%GENERATED_OUTPUT%是__global int *out#,的拼接，#为行号，从1开始；最后这个%GENERATED_CALCULATION%则是被替换为函数调用calculateRow(row#, row##, row###, out#, idx, gameWidth);的拼接，其中#、##和###为顺序递增的行号，例如输入行1、2、3可以计算出输出行1。
在程序中我们可以这样初始化Kernel：
val rowPerKernel: Int = 16
/**
* Each time we update [rowPerKernel] rows per kernel
* */
private val inputBuffers = Array(rowPerKernel + 2) {
context.createIntBuffer(cellStatus.wordWidth, CLMemory.Mem.READ_ONLY)
}
private val outputBuffers = Array(rowPerKernel) {
context.createIntBuffer(cellStatus.wordWidth, CLMemory.Mem.WRITE_ONLY)
}
/**
* The compiled OpenCL program.
* */
private val program = context.createProgram(
this::class.java.getResourceAsStream("/cl_kernels/conways_game.cl")!!
.bufferedReader().use { it.readText() }
.replace(
"%GENERATED_INPUT%",
(1..inputBuffers.size).joinToString("\n ") {
"__global const int *row${it},"
}
)
.replace(
"%GENERATED_OUTPUT%",
(1..outputBuffers.size).joinToString("\n ") {
"__global int *out${it},"
}
)
.replace("%GENERATED_CALCULATION%",
(1..outputBuffers.size).joinToString("\n ") {
"calculateRow(row${it}, row${it + 1}, row${it + 2}, out${it}, idx, gameWidth);"
}
)
// .also { println("Generated OpenCL kernel code: \n$it") }
).build()
private val kernel = program.createCLKernel("calculateNextTick")
.also {
it.putArgs(*inputBuffers, *outputBuffers)
.putArg(gameWidth)
}当然，这里还对WorldMap这个类做了一些改造，使得它可以更加方便地将内容导出到OpenCL中，以及从OpenCL的结果中读入。至于在程序中调用Kernel进行计算，则无外乎如下几步：
- 从WorldMap中导出N+2行
- 计算
- 导出N行结果，使用WorldMap.Builder读入结果
- 回到Step 1，继续计算
关于性能，我使用了一个边长为10k的正方形世界作为测试，在使用OpenCL时，每一步消耗的时间不超过420ms，大约有20ms的时间花费在了数据传送方面。而使用CPU时，每一步耗时都在1500ms以上。可以说，使用OpenCL有效地提高了游戏的模拟速度。
生成视频
现在我们把游戏那部分的运算扔到GPU上了，就可以用CPU来生成视频了对吧？虽然一开始用HumbleVideo做到了，但我发现编码视频好像还是一个挺密集的活动，具体而言就是GPU已经生成完了所有步骤，而CPU还在苦苦编码视频。总之，无论我怎么调整参数，编码视频这个活儿对于CPU来说都太重了。
最终经过一番探索，和来自New Bing的帮助，我得知JavaCV能够调用C编写的ffmpeg库进行视频编码，同时ffmpeg可以调用硬件进行编码。也就是说，我可以在GPU上同时进行游戏的模拟和视频编码。
关于这个库的使用也比较简单明了，你需要一个FFmpegFrameRecorder：
recorder = FFmpegFrameRecorder(fileName, videoWidth, videoHeight)
.also {
it.frameRate = framerate
it.format = formatName
it.videoCodecName = codec.codecName
it.pixelFormat = codec.pixelFormat
it.videoBitrate = bitrate
it.start()
}这里使用文件名和分辨率实例化一个FFmpegFrameRecorder对象，然后设置一些基本信息，例如帧率、格式（基本上就是文件的后缀名，例如mp4）、视频编码器的名称和像素格式，还有比特率（决定了视频的清晰度，比特率越高，视频越清晰，细节越多，同时占用空间就越大），最后调用start()开始编码。
这里要说明的部分是编码器和像素格式。在Java的BuffererImage中有好多常量决定了一个BufferedImage的颜色类型，我比较常用的是BufferedImage.TYPE_4BYTE_ABGR，注意这里需要使用4BYTE_ARGB而不是INT_ARGB，后者会使ffmpeg引起JVM崩溃，我猜测可能是Java的Int是大端序，而ffmpeg期待小端序。RGB格式的像素通常是3到4个字节，取决于有没有Alpha通道，每个通道占用1字节，取值从0到255，表示这个通道对应分量的强度，每个像素一共32位（4字节），因此也成为32位色深。而视频领域中，比较常见的是YUV420P，这个YUV表示的是亮度（Y）和色度（U和V）。其中Y通道表示图像在灰度模式下的亮度信息，而U描述了蓝色通道和亮度之间的差异，取值-128到127之间，U越大，表示像素越蓝；类似地，V表示了红色通道与亮度之间的差异。而后面的数字则表示了采样格式。例如YUV420表示4:2:0采样，即每个Y分量有4个像素，2表示每个U和V分量有2个像素，0表示U和V分量二选一。此外还有YUV422，最后一个2表示保留了U和V分量，而非420的二选一。在YUV420格式中，U和V分量在水平和垂直方向上的取样频率都被减半了，而YUV422中U、V分量在垂直方向上没有减半，因此后者可以实现较高的图像质量，并且在水平方向上的色度分辨率较高。当然还有一个YUV444，表示每一个分量都有4个像素，没有任何减半，它可以提供最高质量的图像。也就是说，ARGB在显示效果上和YUV444无异。而YUV420P中的P则表示像素数据按照平面方式连续存储（没有P的话成为交错存储，也就是按照YUVYUV这样的顺序交替存储三个分量），即先存储全部的Y，然后存储U，最后存储V，也就是YYYYY...YUUU...UVVV...V这样子。
在实际应用中，电脑显示器使用RGB表示颜色，而YUV则利用人眼对图像的感知来压缩视频。因为人眼对亮度更加敏感，因此视频采用YUV格式能够使用更少的比特率来提供质量相当的图像。不过即便是YUV444，它也是经过采样的，而ARGB则是几个颜色分量的原始信息。
虽然ARGB虽好，但不是所有编码器都支持这种像素格式。我采用H265（也就是HEVC）编码，运行ffmpeg -h encoder=hevc，得到了如下输出：
......
Encoder libx265 [libx265 H.265 / HEVC]:
General capabilities: dr1 delay threads
Threading capabilities: other
Supported pixel formats: yuv420p yuvj420p yuv422p yuvj422p yuv444p yuvj444p gbrp yuv420p10le yuv422p10le yuv444p10le gbrp10le yuv420p12le yuv422p12le yuv444p12le gbrp12le gray gray10le gray12le
libx265 AVOptions:
......
Encoder hevc_amf [AMD AMF HEVC encoder]:
General capabilities: dr1 delay hardware
Threading capabilities: none
Supported hardware devices: d3d11va d3d11va dxva2 dxva2
Supported pixel formats: nv12 yuv420p d3d11 dxva2_vld
hevc_amf AVOptions:
......
Encoder hevc_mf [HEVC via MediaFoundation]:
General capabilities: dr1 delay hybrid
Threading capabilities: none
Supported pixel formats: nv12 yuv420p
hevc_mf AVOptions:
......
Encoder hevc_nvenc [NVIDIA NVENC hevc encoder]:
General capabilities: dr1 delay hardware
Threading capabilities: none
Supported hardware devices: cuda cuda d3d11va d3d11va
Supported pixel formats: yuv420p nv12 p010le yuv444p p016le yuv444p16le bgr0 bgra rgb0 rgba x2rgb10le x2bgr10le gbrp gbrp16le cuda d3d11
hevc_nvenc AVOptions:
......
Encoder hevc_qsv [HEVC (Intel Quick Sync Video acceleration)]:
General capabilities: delay hybrid
Threading capabilities: none
Supported hardware devices: qsv qsv qsv
Supported pixel formats: nv12 p010le p012le yuyv422 y210le qsv bgra x2rgb10le vuyx xv30le
hevc_qsv encoder AVOptions:
......这个命令打印出了FFMpeg支持的hevc编码器，你会发现他们支持的像素格式各不相同。除了libx265之外，还有一些硬件编码器，例如hevc_nvenc是英伟达的编码器。他们支持的像素格式各不相同，作为一个顶配控，我自然要选择最好的。我查询了Google、StackOverflow、FFMpeg的文档，还有ChatGPT和New Bing，得出的结论就是，如果源是ARGB格式，那么最好不要转换成YUV，即便是YUV444也是要经过采样，因此会有损失。考虑到编码器支持的像素格式，我选用了BGRA格式（和ARGB一样，只不过调整了通道的顺序）。
开始录制后，渲染也比较容易，只不过面对FFMpeg时不能使用Java的BufferedImage了，需要使用JavaCV提供的Java2DFrameConverter来将BufferedImage转换为FFMpeg的Frame：
val frame = converter.getFrame(image)
for (j in 0 until framePerStep) {
recorder!!.record(frame, avutil.AV_PIX_FMT_ARGB)
}由于帧率比较快，为了让人看清楚每一步的变化，这里针对每一步的图像重复渲染几帧。当然，你也可以直接降低帧率。要结束视频，只需要调用recorder.stop()即可。
关于生成每一步的图像：
fun generateBufferedImage(
videoWidth: Int, videoHeight: Int, cellSize: Int,
background: Color, grid: Color?,
alive: Color, aboutToDie: Color, aboutToLive: Color
): BufferedImage {
val (currentStatus, nextStatus) = getWorldMap()
val result = BufferedImage(videoWidth, videoHeight, BufferedImage.TYPE_4BYTE_ABGR)
for (y in 0 until videoHeight) {
val cellY = y / cellSize
for (x in 0 until videoWidth) {
if (grid != null) {
// draw grid if and only if the cell size is big enough
if (x % cellSize == 0 || y % cellSize == 0) {
result.setRGB(x, y, grid.rgb)
continue // grid
}
}
val cellX = x / cellSize
val next = nextStatus[cellX, cellY]
if (currentStatus[cellX, cellY]) {
// current alive
if (next) result.setRGB(x, y, alive.rgb)
else result.setRGB(x, y, aboutToDie.rgb)
} else {
// current dead
if (next) result.setRGB(x, y, aboutToLive.rgb)
else result.setRGB(x, y, background.rgb)
}
}
}
return result
}这里无非就是根据状态设置像素的颜色。同时为了体现出变化，当一个格子要在下一步死亡时，我们把它的颜色换成亮灰色（要比白色淡），而当一个格子在下一步复活时，我们把它的颜色调成深灰色（要比纯黑色的背景亮），这样有一种过渡的感觉。
总结
在做最后测试的时候我发现原来GPU上，OpenCL和视频编码可以同时运行（仔细想想好像也没什么可惊讶的，毕竟是分开的电路）。这使得模拟并记录游戏的过程变快了许多，同时还减轻了CPU的负担。很有意思。关于这两篇文章设计的代码，可以在GitHub上的hurui200320/pocog找到。
-全文完-
【代码札记】使用GPU加速并记录将康威生命游戏 由 天空 Blond 采用 知识共享 署名 - 非商业性使用 - 相同方式共享 4.0 国际 许可协议进行许可。
本许可协议授权之外的使用权限可以从 https://skyblond.info/about.html 处获得。