使用 JOCL/OpenCL 将图像复制到另一个

Question

所以我的目标是将 GPU 用于我全新的 Java 项目，即创建游戏和游戏引擎本身（我认为这是深入了解其工作原理的好方法）。

我在 CPU 上使用 java.awt.Graphics2D 的多线程来显示我的游戏，但我在其他 PC 上观察到游戏运行速度低于 40FPS，所以我决定学习如何使用 GPU （我仍将在 for 循环中渲染所有对象，然后在屏幕上绘制图像）。

出于这个原因，我开始按照 OpenCL 文档编写代码，并且 JOCL 采样了一个简单的小测试，即将纹理绘制到背景图像上（让我们假设每个实体都有纹理） .

这个方法在每次渲染调用中被调用，它被赋予背景、纹理和实体的位置作为参数。

以下两个代码都已更新，以适应 @ProjectPhysX 的建议。

public static void XXX(final BufferedImage output_image, final BufferedImage input_image, float x, float y) {
        cl_image_format format = new cl_image_format();
        format.image_channel_order = CL_RGBA;
        format.image_channel_data_type = CL_UNSIGNED_INT8;

        //allocate ouput pointer
        cl_image_desc output_description = new cl_image_desc();
        output_description.buffer = null; //must be null for 2D image
        output_description.image_depth = 0; //is only used if the image is a 3D image
        output_description.image_row_pitch = 0; //must be 0 if host_ptr is null
        output_description.image_slice_pitch = 0; //must be 0 if host_ptr is null
        output_description.num_mip_levels = 0; //must be 0
        output_description.num_samples = 0; //must be 0
        output_description.image_type = CL_MEM_OBJECT_IMAGE2D;
        output_description.image_width = output_image.getWidth();
        output_description.image_height = output_image.getHeight();
        output_description.image_array_size = output_description.image_width * output_description.image_height;

        cl_mem output_memory = clCreateImage(context, CL_MEM_WRITE_ONLY, format, output_description, null, null);
        
        //set up first kernel arg
        clSetKernelArg(kernel, 0, Sizeof.cl_mem, Pointer.to(output_memory));
        
        //allocates input pointer
        cl_image_desc input_description = new cl_image_desc();
        input_description.buffer = null; //must be null for 2D image
        input_description.image_depth = 0; //is only used if the image is a 3D image
        input_description.image_row_pitch = 0; //must be 0 if host_ptr is null
        input_description.image_slice_pitch = 0; //must be 0 if host_ptr is null
        input_description.num_mip_levels = 0; //must be 0
        input_description.num_samples = 0; //must be 0
        input_description.image_type = CL_MEM_OBJECT_IMAGE2D;
        input_description.image_width = input_image.getWidth();
        input_description.image_height = input_image.getHeight();
        input_description.image_array_size = input_description.image_width * input_description.image_height;

        DataBufferInt input_buffer = (DataBufferInt) input_image.getRaster().getDataBuffer();
        int input_data[] = input_buffer.getData();

        cl_mem input_memory = clCreateImage(context, CL_MEM_READ_ONLY | CL_MEM_USE_HOST_PTR, format, input_description, Pointer.to(input_data), null);

        //loads the input buffer to the gpu memory
        long[] input_origin = new long[] { 0, 0, 0 };
        long[] input_region = new long[] { input_image.getWidth(), input_image.getHeight(), 1 };
        int input_row_pitch = input_image.getWidth() * Sizeof.cl_uint; //the length of each row in bytes
        clEnqueueWriteImage(commandQueue, input_memory, CL_TRUE, input_origin, input_region, input_row_pitch, 0, Pointer.to(input_data), 0, null, null);
        
        //set up second kernel arg
        clSetKernelArg(kernel, 1, Sizeof.cl_mem, Pointer.to(input_memory));

        //set up third and fourth kernel args
        clSetKernelArg(kernel, 2, Sizeof.cl_float, Pointer.to(new float[] { x }));
        clSetKernelArg(kernel, 3, Sizeof.cl_float, Pointer.to(new float[] { y }));
        
        //blocks until all previously queued commands are issued
        clFinish(commandQueue);

        //enqueue the program execution
        long[] globalWorkSize = new long[] { input_description.image_width, input_description.image_height };
        clEnqueueNDRangeKernel(commandQueue, kernel, 2, null, globalWorkSize, null, 0, null, null);

        //transfer the output result back to host
        DataBufferInt output_buffer = (DataBufferInt) output_image.getRaster().getDataBuffer();
        int output_data[] = output_buffer.getData();
        long[] output_origin = new long[] { 0, 0, 0 };
        long[] output_region = new long[] { output_description.image_width, output_description.image_height, 1 };
        int output_row_pitch = output_image.getWidth() * Sizeof.cl_uint;
        clEnqueueReadImage(commandQueue, output_memory, CL_TRUE, output_origin, output_region, output_row_pitch, 0, Pointer.to(output_data), 0, null, null);

        //free pointers
        clReleaseMemObject(input_memory);
        clReleaseMemObject(output_memory);
    }

这是在内核上运行的程序源代码。

const sampler_t sampler = CLK_NORMALIZED_COORDS_FALSE | CLK_ADDRESS_CLAMP | CLK_FILTER_NEAREST;

__kernel void drawImage(__write_only image2d_t dst_image, __read_only image2d_t src_image, float xoff, float yoff)
{
    const int x = get_global_id(0);
    const int y = get_global_id(1);

    int2 in_coords = (int2) { x, y };

    uint4 pixel = read_imageui(src_image, sampler, in_coords);
    pixel = -16184301;
    printf("%d, %d, %u\n", x, y, pixel);

    const int sx = get_global_size(0);
    const int sy = get_global_size(1);

    int2 out_coords = (int2) { ((int) xoff + x) % sx, ((int) yoff + y) % sy};
    
    write_imageui(dst_image, out_coords, pixel);
}

如果不调用 write_imageui，背景将被涂成黑色，否则为白色。 目前，我有点难以理解为什么 C 函数中的 pixel = 0，但我认为熟悉 JOCL 的人会很快发现我在这段代码中的错误。今天，也许明天，我对这段代码感到非常困惑，但我觉得我永远不会发现自己的错误。出于这个原因，我请求您帮助查看我的代码。我觉得自己像个白痴，我当时想不通。

Answer 1

试试

    const int sx = get_global_size(0);
    const int sy = get_global_size(1);
    int2 out_coords = (int2) { (xoff + x)%sx, (yoff + y)%sy};

以避免错误或未定义的行为。现在，如果坐标+偏移量位于图像区域之外，您正在写入 Nirwana。另外在调用内核之前没有clEnqueueWriteImage，所以GPU上的src_image是未定义的，可能包含随机值。

OpenCL 要求在global 内存空间中声明内核参数：

__kernel void drawImage(global image2d_t dst_image, global image2d_t src_image, global float xoff, global float yoff)

作为一个用 Java、C++ 和 OpenCL 并行 GPU 编写图形引擎的人，让我给你一些指导：在 Java 代码中，你可能会使用画家算法：列出所有绘制的对象及其近似 z 坐标，按 z 坐标对对象进行排序，并在单个 for 循环中将它们从后到前绘制。在 GPU 上，painter 的算法将不起作用，因为您无法对其进行并行化。相反，您在 3D 空间中有一个对象列表（线/三角形），并在此列表上进行并行化：每个 GPU 线程同时光栅化单个三角形，所有线程同时在帧上绘制像素。要解决排水顺序问题，您可以使用 z 缓冲区：由每个像素的 z 坐标组成的图像。在线/三角形的光栅化过程中，您计算​​每个像素的 z 坐标，并且只有当它大于该像素的 z 缓冲区中先前的坐标时，您才绘制新颜色。

关于性能：java.awt.Graphics2D 在 CPU 使用率方面非常高效，您可以在 60fps 下每帧处理约 40k 个三角形。使用 OpenCL，预计每帧 3000 万个三角形，速度为 60fps。