MENU

【代码札记】使用GPU加速并记录将康威生命游戏

March 15, 2023 • 瞎折腾

前文记述了我使用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进行计算,则无外乎如下几步:

  1. 从WorldMap中导出N+2行
  2. 计算
  3. 导出N行结果,使用WorldMap.Builder读入结果
  4. 回到Step 1,继续计算

关于性能,我使用了一个边长为10k的正方形世界作为测试,在使用OpenCL时,每一步消耗的时间不超过420ms,大约有20ms的时间花费在了数据传送方面。而使用CPU时,每一步耗时都在1500ms以上。可以说,使用OpenCL有效地提高了游戏的模拟速度。

生成视频

现在我们把游戏那部分的运算扔到GPU上了,就可以用CPU来生成视频了对吧?虽然一开始用HumbleVideo做到了,但我发现编码视频好像还是一个挺密集的活动,具体而言就是GPU已经生成完了所有步骤,而CPU还在苦苦编码视频。总之,无论我怎么调整参数,编码视频这个活儿对于CPU来说都太重了。

最终经过一番探索,和来自New Bing的帮助,我得知JavaCV能够调用C编写的ffmpeg库进行视频编码,同时ffmpeg可以调用硬件进行编码。也就是说,我可以在GPU上同时进行游戏的模拟和视频编码。

屏幕截图 2023-03-15 142921.png

关于这个库的使用也比较简单明了,你需要一个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 处获得。

Last Modified: March 28, 2023
Archives QR Code
QR Code for this page
Tipping QR Code