Project Panama外部函数调用OpenCL内核时local工作组大小配置无效？CLKernelWorkGroupSize与MemorySegment局部内存 - 智猿学院-前后端，数据库，人工智能，云计算等领域前沿技术讲座

Project Panama 与 OpenCL 内核调用：Local 工作组大小配置疑难解答

大家好，今天我们来深入探讨一个在使用 Project Panama 调用 OpenCL 内核时，经常遇到的一个问题：Local 工作组大小配置失效，以及它与 CLKernelWorkGroupSize 和 MemorySegment 分配的局部内存之间的关系。

在开始之前，我们先简单回顾一下 Project Panama 和 OpenCL 的相关概念。

Project Panama (Foreign Function & Memory API): Java 的一个孵化项目，旨在提高 Java 程序与本地代码（如 C/C++）的互操作性。它提供了更安全、更高效的方式来访问本地内存和调用本地函数，避免了传统 JNI 的一些缺陷。
OpenCL (Open Computing Language): 一个异构并行计算的开放标准，允许在各种平台上（CPU、GPU、FPGA 等）执行并行计算任务。OpenCL 程序通常由一个主机程序和一个或多个内核程序组成。主机程序负责管理 OpenCL 环境、加载和编译内核、分配内存、设置内核参数以及启动内核执行。内核程序则是在 OpenCL 设备上执行的并行计算代码。
Local 工作组 (Work-Group): OpenCL 执行模型中的一个重要概念。它将一个大的计算任务划分为更小的、可以并行执行的单元。一个 Global 范围的计算任务被划分为多个 Work-Group，每个 Work-Group 又包含多个 Work-Item。Work-Group 中的 Work-Item 可以通过 Local Memory 进行数据共享和协作。
CLKernelWorkGroupSize: 这是 OpenCL 内核中的一个属性，允许开发者指定内核倾向使用的 Work-Group 大小。但这仅仅是一个提示，OpenCL 运行时环境可以根据设备的限制和资源情况来调整实际的 Work-Group 大小。
MemorySegment: Project Panama 中用于表示本地内存区域的一个核心接口。它提供了安全且高效的方式来访问和操作本地内存，包括分配、释放、读取和写入数据。

问题描述：Local 工作组大小配置无效

在使用 Project Panama 调用 OpenCL 内核时，我们可能会遇到以下问题：

通过 clSetKernelArg 设置的 local 内存大小，与内核代码中声明的 __local 变量大小不匹配，导致内核执行错误。
尝试使用 clGetKernelWorkGroupInfo 查询内核的 CL_KERNEL_WORK_GROUP_SIZE 和 CL_KERNEL_LOCAL_MEM_SIZE，发现返回值与预期的不符，或者在内核执行过程中出现与局部内存访问相关的错误。
尽管设置了 CLKernelWorkGroupSize，但实际执行的 Work-Group 大小与设置的值不一致。

这些问题通常与以下几个方面有关：

OpenCL 运行时环境的限制: OpenCL 设备对 Work-Group 大小和 Local Memory 大小都有硬件限制。如果配置超过了这些限制，运行时环境可能会自动调整 Work-Group 大小或者拒绝执行内核。
内核代码中的内存访问模式: 如果内核代码中的内存访问模式不正确，例如发生越界访问或者并发访问冲突，即使 Work-Group 大小配置正确，也可能导致程序崩溃或者产生错误的结果。
Project Panama 的内存管理: 如果在使用 Project Panama 分配和管理本地内存时出现错误，例如内存泄漏或者访问无效内存，也可能导致内核执行错误。

案例分析：一个简单的向量加法内核

为了更好地理解这些问题，我们来看一个简单的向量加法内核的例子。

OpenCL 内核代码 (vector_add.cl):

__kernel void vector_add(__global const float *a,
                         __global const float *b,
                         __global float *c,
                         __local float *local_a,
                         __local float *local_b,
                         const int size) {
  int gid = get_global_id(0);
  int lid = get_local_id(0);
  int group_size = get_local_size(0);
  int group_id = get_group_id(0);

  if (gid < size) {
      local_a[lid] = a[gid];
      local_b[lid] = b[gid];

      barrier(CLK_LOCAL_MEM_FENCE); // Wait for all work-items to load data

      c[gid] = local_a[lid] + local_b[lid];
  }
}

Java 代码 (使用 Project Panama 调用 OpenCL):


import jdk.incubator.foreign.*;
import java.lang.invoke.MethodHandle;
import java.nio.ByteOrder;

public class VectorAdd {

    private static final String KERNEL_FILE = "vector_add.cl";
    private static final int VECTOR_SIZE = 1024;
    private static final int WORK_GROUP_SIZE = 256; //尝试设置Work-Group大小

    // OpenCL 函数的地址 (假设已经加载)
    private static MemorySegment clCreateContext;
    private static MemorySegment clCreateCommandQueue;
    private static MemorySegment clCreateProgramWithSource;
    private static MemorySegment clBuildProgram;
    private static MemorySegment clCreateKernel;
    private static MemorySegment clSetKernelArg;
    private static MemorySegment clEnqueueNDRangeKernel;
    private static MemorySegment clFinish;
    private static MemorySegment clReleaseMemObject;
    private static MemorySegment clReleaseKernel;
    private static MemorySegment clReleaseProgram;
    private static MemorySegment clReleaseCommandQueue;
    private static MemorySegment clReleaseContext;
    private static MemorySegment clGetKernelWorkGroupInfo;

    //OpenCL 函数的链接器
    private static SymbolLookup openCLLinker;

    //OpenCL 相关对象的句柄
    private static MemorySegment context;
    private static MemorySegment commandQueue;
    private static MemorySegment program;
    private static MemorySegment kernel;

    // 函数句柄定义 (简化起见，省略错误处理)
    private static MethodHandle createContextMH;
    private static MethodHandle createCommandQueueMH;
    private static MethodHandle createProgramWithSourceMH;
    private static MethodHandle buildProgramMH;
    private static MethodHandle createKernelMH;
    private static MethodHandle setKernelArgMH;
    private static MethodHandle enqueueNDRangeKernelMH;
    private static MethodHandle finishMH;
    private static MethodHandle releaseMemObjectMH;
    private static MethodHandle releaseKernelMH;
    private static MethodHandle releaseProgramMH;
    private static MethodHandle releaseCommandQueueMH;
    private static MethodHandle releaseContextMH;
    private static MethodHandle getKernelWorkGroupInfoMH;

    // 定义 OpenCL 常量 (简化起见)
    private static final int CL_DEVICE_TYPE_GPU = 4;
    private static final int CL_MEM_READ_ONLY = 1;
    private static final int CL_MEM_WRITE_ONLY = 2;
    private static final int CL_MEM_READ_WRITE = 4;
    private static final int CL_SUCCESS = 0;
    private static final int CL_KERNEL_WORK_GROUP_SIZE = 0x11B0;
    private static final int CL_KERNEL_LOCAL_MEM_SIZE = 0x11B1;

    public static void main(String[] args) throws Throwable {

        // 1. 初始化 OpenCL 环境
        initOpenCL();

        // 2. 创建 OpenCL 缓冲区
        MemorySegment aBuffer = allocateFloatBuffer(VECTOR_SIZE, CL_MEM_READ_ONLY);
        MemorySegment bBuffer = allocateFloatBuffer(VECTOR_SIZE, CL_MEM_READ_ONLY);
        MemorySegment cBuffer = allocateFloatBuffer(VECTOR_SIZE, CL_MEM_WRITE_ONLY);

        // 3. 填充输入缓冲区 (这里简单地填充一些数据)
        fillFloatBuffer(aBuffer, VECTOR_SIZE, 1.0f);
        fillFloatBuffer(bBuffer, VECTOR_SIZE, 2.0f);

        // 4. 创建内核对象
        kernel = (MemorySegment) createKernelMH.invokeExact((MemorySegment)program.address(), MemorySegment.ofArray("vector_add".getBytes()), (MemorySegment)null);

        // 5. 设置内核参数
        long sizeOfFloat = 4;
        setKernelArgMH.invokeExact((MemorySegment)kernel.address(), 0, sizeOfFloat, (MemorySegment)aBuffer.address());
        setKernelArgMH.invokeExact((MemorySegment)kernel.address(), 1, sizeOfFloat, (MemorySegment)bBuffer.address());
        setKernelArgMH.invokeExact((MemorySegment)kernel.address(), 2, sizeOfFloat, (MemorySegment)cBuffer.address());

        //分配局部内存
        MemorySegment localA = MemorySegment.allocateNative(WORK_GROUP_SIZE * sizeOfFloat, ResourceScope.newConfinedScope());
        MemorySegment localB = MemorySegment.allocateNative(WORK_GROUP_SIZE * sizeOfFloat, ResourceScope.newConfinedScope());

        setKernelArgMH.invokeExact((MemorySegment)kernel.address(), 3, (long)WORK_GROUP_SIZE * sizeOfFloat, (MemorySegment)null); // Local Memory for a
        setKernelArgMH.invokeExact((MemorySegment)kernel.address(), 4, (long)WORK_GROUP_SIZE * sizeOfFloat, (MemorySegment)null); // Local Memory for b
        setKernelArgMH.invokeExact((MemorySegment)kernel.address(), 5, sizeOfFloat, (MemorySegment)MemorySegment.ofArray(new int[]{VECTOR_SIZE}));

        // 6. 执行内核
        MemorySegment globalWorkSize = MemorySegment.allocateNative(8,ResourceScope.newConfinedScope());
        globalWorkSize.setAtIndex(ValueLayout.JAVA_LONG, 0, VECTOR_SIZE);

        MemorySegment localWorkSize = MemorySegment.allocateNative(8,ResourceScope.newConfinedScope());
        localWorkSize.setAtIndex(ValueLayout.JAVA_LONG, 0, WORK_GROUP_SIZE);

        enqueueNDRangeKernelMH.invokeExact((MemorySegment)commandQueue.address(), (MemorySegment)kernel.address(), 1, (MemorySegment)null, (MemorySegment)globalWorkSize.address(), (MemorySegment)localWorkSize.address(), 0, (MemorySegment)null, (MemorySegment)null);

        finishMH.invokeExact((MemorySegment)commandQueue.address());

        // 7. 读取结果并验证
        float[] results = readFloatBuffer(cBuffer, VECTOR_SIZE);
        for (int i = 0; i < VECTOR_SIZE; i++) {
            if (Math.abs(results[i] - 3.0f) > 0.001f) {
                System.err.println("Error at index " + i + ": expected 3.0, got " + results[i]);
            }
        }
        System.out.println("Vector addition completed successfully.");

        // 8. 释放 OpenCL 资源
        releaseOpenCLResources(aBuffer, bBuffer, cBuffer);
    }

    // 初始化 OpenCL 环境
    private static void initOpenCL() throws Throwable {
         // 1. 加载 OpenCL 库
        openCLLinker = SymbolLookup.libraryLookup("OpenCL", SegmentScope.global());

        // 2. 获取 OpenCL 函数地址
        clCreateContext = openCLLinker.find("clCreateContext").orElseThrow(() -> new RuntimeException("clCreateContext not found"));
        clCreateCommandQueue = openCLLinker.find("clCreateCommandQueue").orElseThrow(() -> new RuntimeException("clCreateCommandQueue not found"));
        clCreateProgramWithSource = openCLLinker.find("clCreateProgramWithSource").orElseThrow(() -> new RuntimeException("clCreateProgramWithSource not found"));
        clBuildProgram = openCLLinker.find("clBuildProgram").orElseThrow(() -> new RuntimeException("clBuildProgram not found"));
        clCreateKernel = openCLLinker.find("clCreateKernel").orElseThrow(() -> new RuntimeException("clCreateKernel not found"));
        clSetKernelArg = openCLLinker.find("clSetKernelArg").orElseThrow(() -> new RuntimeException("clSetKernelArg not found"));
        clEnqueueNDRangeKernel = openCLLinker.find("clEnqueueNDRangeKernel").orElseThrow(() -> new RuntimeException("clEnqueueNDRangeKernel not found"));
        clFinish = openCLLinker.find("clFinish").orElseThrow(() -> new RuntimeException("clFinish not found"));
        clReleaseMemObject = openCLLinker.find("clReleaseMemObject").orElseThrow(() -> new RuntimeException("clReleaseMemObject not found"));
        clReleaseKernel = openCLLinker.find("clReleaseKernel").orElseThrow(() -> new RuntimeException("clReleaseKernel not found"));
        clReleaseProgram = openCLLinker.find("clReleaseProgram").orElseThrow(() -> new RuntimeException("clReleaseProgram not found"));
        clReleaseCommandQueue = openCLLinker.find("clReleaseCommandQueue").orElseThrow(() -> new RuntimeException("clReleaseCommandQueue not found"));
        clReleaseContext = openCLLinker.find("clReleaseContext").orElseThrow(() -> new RuntimeException("clReleaseContext not found"));
        clGetKernelWorkGroupInfo = openCLLinker.find("clGetKernelWorkGroupInfo").orElseThrow(() -> new RuntimeException("clGetKernelWorkGroupInfo not found"));

        // 3. 创建 MethodHandle
        createContextMH = MethodHandles.lookup().findStatic(OCLUtils.class, "clCreateContext", MethodType.methodType(MemorySegment.class, SymbolLookup.class, int.class));
        createCommandQueueMH = MethodHandles.lookup().findStatic(OCLUtils.class, "clCreateCommandQueue", MethodType.methodType(MemorySegment.class, SymbolLookup.class, MemorySegment.class));
        createProgramWithSourceMH = MethodHandles.lookup().findStatic(OCLUtils.class, "clCreateProgramWithSource", MethodType.methodType(MemorySegment.class, SymbolLookup.class, String.class));
        buildProgramMH = MethodHandles.lookup().findStatic(OCLUtils.class, "clBuildProgram", MethodType.methodType(int.class, SymbolLookup.class, MemorySegment.class));
        createKernelMH = MethodHandles.lookup().findStatic(OCLUtils.class, "clCreateKernel", MethodType.methodType(MemorySegment.class, SymbolLookup.class, MemorySegment.class, String.class));
        setKernelArgMH = MethodHandles.lookup().findStatic(OCLUtils.class, "clSetKernelArg", MethodType.methodType(int.class, SymbolLookup.class, MemorySegment.class, int.class, long.class, MemorySegment.class));
        enqueueNDRangeKernelMH = MethodHandles.lookup().findStatic(OCLUtils.class, "clEnqueueNDRangeKernel", MethodType.methodType(int.class, SymbolLookup.class, MemorySegment.class, MemorySegment.class, int.class, MemorySegment.class, MemorySegment.class, MemorySegment.class, int.class, MemorySegment.class, MemorySegment.class));
        finishMH = MethodHandles.lookup().findStatic(OCLUtils.class, "clFinish", MethodType.methodType(int.class, SymbolLookup.class, MemorySegment.class));
        releaseMemObjectMH = MethodHandles.lookup().findStatic(OCLUtils.class, "clReleaseMemObject", MethodType.methodType(int.class, SymbolLookup.class, MemorySegment.class));
        releaseKernelMH = MethodHandles.lookup().findStatic(OCLUtils.class, "clReleaseKernel", MethodType.methodType(int.class, SymbolLookup.class, MemorySegment.class));
        releaseProgramMH = MethodHandles.lookup().findStatic(OCLUtils.class, "clReleaseProgram", MethodType.methodType(int.class, SymbolLookup.class, MemorySegment.class));
        releaseCommandQueueMH = MethodHandles.lookup().findStatic(OCLUtils.class, "clReleaseCommandQueue", MethodType.methodType(int.class, SymbolLookup.class, MemorySegment.class));
        releaseContextMH = MethodHandles.lookup().findStatic(OCLUtils.class, "clReleaseContext", MethodType.methodType(int.class, SymbolLookup.class, MemorySegment.class));
        getKernelWorkGroupInfoMH = MethodHandles.lookup().findStatic(OCLUtils.class, "clGetKernelWorkGroupInfo", MethodType.methodType(int.class, SymbolLookup.class, MemorySegment.class, int.class, long.class, MemorySegment.class, MemorySegment.class));

        // 4. 创建 OpenCL Context 和 Command Queue
        context = (MemorySegment) createContextMH.invokeExact(openCLLinker, CL_DEVICE_TYPE_GPU);
        commandQueue = (MemorySegment) createCommandQueueMH.invokeExact(openCLLinker, (MemorySegment)context.address());

        // 5. 创建 OpenCL Program
        program = (MemorySegment) createProgramWithSourceMH.invokeExact(openCLLinker, KERNEL_FILE);

        // 6. 构建 OpenCL Program
        int buildStatus = (int) buildProgramMH.invokeExact(openCLLinker, (MemorySegment)program.address());
        if (buildStatus != CL_SUCCESS) {
            System.err.println("Error building program. Status: " + buildStatus);
            throw new RuntimeException("OpenCL program build failed.");
        }

    }

    // 创建 OpenCL 缓冲区
    private static MemorySegment allocateFloatBuffer(int size, int flags) throws Throwable {
       return OCLUtils.allocateFloatBuffer(openCLLinker, (MemorySegment) context.address(), size, flags);
    }

    // 填充浮点缓冲区
    private static void fillFloatBuffer(MemorySegment buffer, int size, float value) {
        try (MemorySegment.Access<Float> accessor = buffer.asFloatSegment().asAccess()) {
            for (int i = 0; i < size; i++) {
                accessor.set(i, value);
            }
        }
    }

    // 读取浮点缓冲区
    private static float[] readFloatBuffer(MemorySegment buffer, int size) {
        float[] results = new float[size];
        try (MemorySegment.Access<Float> accessor = buffer.asFloatSegment().asAccess()) {
            for (int i = 0; i < size; i++) {
                results[i] = accessor.get(i);
            }
        }
        return results;
    }

    // 释放 OpenCL 资源
    private static void releaseOpenCLResources(MemorySegment aBuffer, MemorySegment bBuffer, MemorySegment cBuffer) throws Throwable {
        releaseMemObjectMH.invokeExact(openCLLinker, (MemorySegment) aBuffer.address());
        releaseMemObjectMH.invokeExact(openCLLinker, (MemorySegment) bBuffer.address());
        releaseMemObjectMH.invokeExact(openCLLinker, (MemorySegment) cBuffer.address());
        releaseKernelMH.invokeExact(openCLLinker, (MemorySegment) kernel.address());
        releaseProgramMH.invokeExact(openCLLinker, (MemorySegment) program.address());
        releaseCommandQueueMH.invokeExact(openCLLinker, (MemorySegment) commandQueue.address());
        releaseContextMH.invokeExact(openCLLinker, (MemorySegment) context.address());
    }

    public static class OCLUtils {
        public static MemorySegment clCreateContext(SymbolLookup openCLLinker, int device_type) throws Throwable {
            // 1. 获取函数描述符
            MethodHandle clGetPlatformIDsMH = MethodHandles.lookup().findStatic(OCLUtils.class, "clGetPlatformIDs", MethodType.methodType(int.class, SymbolLookup.class, int.class, MemorySegment.class, MemorySegment.class));
            MethodHandle clGetDeviceIDsMH = MethodHandles.lookup().findStatic(OCLUtils.class, "clGetDeviceIDs", MethodType.methodType(int.class, SymbolLookup.class, MemorySegment.class, int.class, int.class, MemorySegment.class, MemorySegment.class));
            MethodHandle createContextMH = MethodHandles.lookup().findOrNull(MethodHandles.Lookup.PRIVATE, OCLUtils.class, "createContext", MethodType.methodType(MemorySegment.class, SymbolLookup.class, MemorySegment.class));

            // 2. 调用 clGetPlatformIDs 获取平台 ID
            MemorySegment platformId = MemorySegment.allocateNative(8,ResourceScope.newConfinedScope());
            MemorySegment numPlatforms = MemorySegment.allocateNative(4,ResourceScope.newConfinedScope());
            clGetPlatformIDsMH.invokeExact(openCLLinker, 1, (MemorySegment)platformId.address(), (MemorySegment)numPlatforms.address());

            // 3. 调用 clGetDeviceIDs 获取设备 ID
            MemorySegment deviceId = MemorySegment.allocateNative(8,ResourceScope.newConfinedScope());
            MemorySegment numDevices = MemorySegment.allocateNative(4,ResourceScope.newConfinedScope());
            int result = (int) clGetDeviceIDsMH.invokeExact(openCLLinker, (MemorySegment)platformId.address(), device_type, 1, (MemorySegment)deviceId.address(), (MemorySegment)numDevices.address());

            // 4. 创建上下文 (使用私有方法)
            return (MemorySegment) createContextMH.invokeExact(openCLLinker, (MemorySegment)deviceId.address());
        }

        private static MemorySegment createContext(SymbolLookup openCLLinker, MemorySegment device_id) throws Throwable{
            // 函数描述符
            MethodHandle clCreateContextMH = MethodHandles.lookup().findOrNull(MethodHandles.Lookup.PRIVATE, OCLUtils.class, "clCreateContextInner", MethodType.methodType(MemorySegment.class, SymbolLookup.class, MemorySegment.class));

            return (MemorySegment) clCreateContextMH.invokeExact(openCLLinker, device_id);
        }

        private static MemorySegment clCreateContextInner(SymbolLookup openCLLinker, MemorySegment device_id) throws Throwable {
            // 获取 clCreateContext 函数地址
            MemorySegment clCreateContextAddress = openCLLinker.find("clCreateContext").orElseThrow(() -> new RuntimeException("clCreateContext not found"));

            // 定义函数类型
            MethodType clCreateContextType = MethodType.methodType(MemorySegment.class, MemorySegment.class, int.class, MemorySegment.class, MemorySegment.class, MemorySegment.class);

            // 创建函数句柄
            MethodHandle clCreateContextMH = MethodHandles.lookup().findVirtual(ForeignLinker.class, "downcallHandle", MethodType.methodType(MethodHandle.class, MemorySegment.class, MethodType.class, FunctionDescriptor.class));

            // 创建 ForeignLinker 实例
            ForeignLinker foreignLinker = ForeignLinker.getInstance();

            // 定义函数描述符
            FunctionDescriptor functionDescriptor = FunctionDescriptor.of(ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.JAVA_INT,
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)));

            // 创建 downcallHandle
            MethodHandle handle = (MethodHandle) clCreateContextMH.invoke(foreignLinker, clCreateContextAddress, clCreateContextType, functionDescriptor);

            // 调用 clCreateContext
            MemorySegment properties = MemorySegment.allocateNative(0,ResourceScope.newConfinedScope());//null
            MemorySegment pfn_notify = MemorySegment.allocateNative(0,ResourceScope.newConfinedScope());//null
            MemorySegment errcode_ret = MemorySegment.allocateNative(4,ResourceScope.newConfinedScope());//int
            return (MemorySegment) handle.invokeExact((MemorySegment)properties.address(), 1, (MemorySegment)device_id.address(), (MemorySegment)pfn_notify.address(), (MemorySegment)errcode_ret.address());
        }

        private static int clGetPlatformIDs(SymbolLookup openCLLinker, int num_entries, MemorySegment platforms, MemorySegment num_platforms) throws Throwable {
            // 获取 clGetPlatformIDs 函数地址
            MemorySegment clGetPlatformIDsAddress = openCLLinker.find("clGetPlatformIDs").orElseThrow(() -> new RuntimeException("clGetPlatformIDs not found"));

            // 定义函数类型
            MethodType clGetPlatformIDsType = MethodType.methodType(int.class, int.class, MemorySegment.class, MemorySegment.class);

            // 创建函数句柄
            MethodHandle clGetPlatformIDsMH = MethodHandles.lookup().findVirtual(ForeignLinker.class, "downcallHandle", MethodType.methodType(MethodHandle.class, MemorySegment.class, MethodType.class, FunctionDescriptor.class));

            // 创建 ForeignLinker 实例
            ForeignLinker foreignLinker = ForeignLinker.getInstance();

            // 定义函数描述符
            FunctionDescriptor functionDescriptor = FunctionDescriptor.of(ValueLayout.JAVA_INT,
                    ValueLayout.JAVA_INT,
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)));

            // 创建 downcallHandle
            MethodHandle handle = (MethodHandle) clGetPlatformIDsMH.invoke(foreignLinker, clGetPlatformIDsAddress, clGetPlatformIDsType, functionDescriptor);

            // 调用 clGetPlatformIDs
            return (int) handle.invokeExact(num_entries, (MemorySegment)platforms.address(), (MemorySegment)num_platforms.address());
        }

        private static int clGetDeviceIDs(SymbolLookup openCLLinker, MemorySegment platform, int device_type, int num_entries, MemorySegment devices, MemorySegment num_devices) throws Throwable {
            // 获取 clGetDeviceIDs 函数地址
            MemorySegment clGetDeviceIDsAddress = openCLLinker.find("clGetDeviceIDs").orElseThrow(() -> new RuntimeException("clGetDeviceIDs not found"));

            // 定义函数类型
            MethodType clGetDeviceIDsType = MethodType.methodType(int.class, MemorySegment.class, int.class, int.class, MemorySegment.class, MemorySegment.class);

            // 创建函数句柄
            MethodHandle clGetDeviceIDsMH = MethodHandles.lookup().findVirtual(ForeignLinker.class, "downcallHandle", MethodType.methodType(MethodHandle.class, MemorySegment.class, MethodType.class, FunctionDescriptor.class));

            // 创建 ForeignLinker 实例
            ForeignLinker foreignLinker = ForeignLinker.getInstance();

            // 定义函数描述符
            FunctionDescriptor functionDescriptor = FunctionDescriptor.of(ValueLayout.JAVA_INT,
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.JAVA_INT,
                    ValueLayout.JAVA_INT,
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)));

            // 创建 downcallHandle
            MethodHandle handle = (MethodHandle) clGetDeviceIDsMH.invoke(foreignLinker, clGetDeviceIDsAddress, clGetDeviceIDsType, functionDescriptor);

            // 调用 clGetDeviceIDs
            return (int) handle.invokeExact((MemorySegment)platform.address(), device_type, num_entries, (MemorySegment)devices.address(), (MemorySegment)num_devices.address());
        }

        public static MemorySegment clCreateCommandQueue(SymbolLookup openCLLinker, MemorySegment context) throws Throwable {
            // 获取 clCreateCommandQueue 函数地址
            MemorySegment clCreateCommandQueueAddress = openCLLinker.find("clCreateCommandQueue").orElseThrow(() -> new RuntimeException("clCreateCommandQueue not found"));

            // 定义函数类型
            MethodType clCreateCommandQueueType = MethodType.methodType(MemorySegment.class, MemorySegment.class, MemorySegment.class, int.class, MemorySegment.class);

            // 创建函数句柄
            MethodHandle clCreateCommandQueueMH = MethodHandles.lookup().findVirtual(ForeignLinker.class, "downcallHandle", MethodType.methodType(MethodHandle.class, MemorySegment.class, MethodType.class, FunctionDescriptor.class));

            // 创建 ForeignLinker 实例
            ForeignLinker foreignLinker = ForeignLinker.getInstance();

            // 定义函数描述符
            FunctionDescriptor functionDescriptor = FunctionDescriptor.of(ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.JAVA_INT,
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)));

            // 创建 downcallHandle
            MethodHandle handle = (MethodHandle) clCreateCommandQueueMH.invoke(foreignLinker, clCreateCommandQueueAddress, clCreateCommandQueueType, functionDescriptor);

            // 调用 clCreateCommandQueue
            MemorySegment properties = MemorySegment.allocateNative(0,ResourceScope.newConfinedScope());//null
            MemorySegment errcode_ret = MemorySegment.allocateNative(4,ResourceScope.newConfinedScope());//int

            return (MemorySegment) handle.invokeExact((MemorySegment)context.address(), (MemorySegment)properties.address(), 0, (MemorySegment)errcode_ret.address());
        }

        public static MemorySegment clCreateProgramWithSource(SymbolLookup openCLLinker, String kernelFile) throws Throwable {
            // 1. 读取内核代码
            String kernelCode = readKernelFile(kernelFile);

            // 2. 获取 clCreateProgramWithSource 函数地址
            MemorySegment clCreateProgramWithSourceAddress = openCLLinker.find("clCreateProgramWithSource").orElseThrow(() -> new RuntimeException("clCreateProgramWithSource not found"));

            // 3. 定义函数类型
            MethodType clCreateProgramWithSourceType = MethodType.methodType(MemorySegment.class, MemorySegment.class, MemorySegment.class, int.class, MemorySegment.class, MemorySegment.class);

            // 4. 创建函数句柄
            MethodHandle clCreateProgramWithSourceMH = MethodHandles.lookup().findVirtual(ForeignLinker.class, "downcallHandle", MethodType.methodType(MethodHandle.class, MemorySegment.class, MethodType.class, FunctionDescriptor.class));

            // 5. 创建 ForeignLinker 实例
            ForeignLinker foreignLinker = ForeignLinker.getInstance();

            // 6. 定义函数描述符
            FunctionDescriptor functionDescriptor = FunctionDescriptor.of(ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.JAVA_INT,
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)));

            // 7. 创建 downcallHandle
            MethodHandle handle = (MethodHandle) clCreateProgramWithSourceMH.invoke(foreignLinker, clCreateProgramWithSourceAddress, clCreateProgramWithSourceType, functionDescriptor);

            // 8. 调用 clCreateProgramWithSource
            MemorySegment context = VectorAdd.context;
            int count = 1;
            MemorySegment strings = MemorySegment.allocateNative(8,ResourceScope.newConfinedScope());//String[]
            strings.setAtIndex(ValueLayout.ADDRESS, 0, MemorySegment.ofArray(kernelCode.getBytes()).address());
            MemorySegment lengths = MemorySegment.allocateNative(4,ResourceScope.newConfinedScope());//size_t[]
            lengths.setAtIndex(ValueLayout.JAVA_INT, 0, kernelCode.length());
            MemorySegment errcode_ret = MemorySegment.allocateNative(4,ResourceScope.newConfinedScope());//int*

            return (MemorySegment) handle.invokeExact((MemorySegment)context.address(), count, (MemorySegment)strings.address(), (MemorySegment)lengths.address(), (MemorySegment)errcode_ret.address());
        }

        private static String readKernelFile(String kernelFile) {
            try {
                return new String(VectorAdd.class.getClassLoader().getResourceAsStream(kernelFile).readAllBytes());
            } catch (Exception e) {
                throw new RuntimeException("Failed to read kernel file: " + kernelFile, e);
            }
        }

        public static int clBuildProgram(SymbolLookup openCLLinker, MemorySegment program) throws Throwable {
            // 获取 clBuildProgram 函数地址
            MemorySegment clBuildProgramAddress = openCLLinker.find("clBuildProgram").orElseThrow(() -> new RuntimeException("clBuildProgram not found"));

            // 定义函数类型
            MethodType clBuildProgramType = MethodType.methodType(int.class, MemorySegment.class, int.class, MemorySegment.class, MemorySegment.class, MemorySegment.class, MemorySegment.class);

            // 创建函数句柄
            MethodHandle clBuildProgramMH = MethodHandles.lookup().findVirtual(ForeignLinker.class, "downcallHandle", MethodType.methodType(MethodHandle.class, MemorySegment.class, MethodType.class, FunctionDescriptor.class));

            // 创建 ForeignLinker 实例
            ForeignLinker foreignLinker = ForeignLinker.getInstance();

            // 定义函数描述符
            FunctionDescriptor functionDescriptor = FunctionDescriptor.of(ValueLayout.JAVA_INT,
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.JAVA_INT,
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)),
                    ValueLayout.ADDRESS.withTargetLayout(MemoryLayout.ofValue(ValueLayout.ADDRESS)));

            // 创建 downcallHandle
            MethodHandle handle = (MethodHandle) clBuildProgramMH.invoke(foreignLinker, clBuildProgramAddress, clBuildProgramType, functionDescriptor);

            // 调用 clBuildProgram
            MemorySegment device_list = MemorySegment.allocateNative(0,ResourceScope.newConfinedScope());//null
            MemorySegment pfn_notify = MemorySegment.allocateNative(0,ResourceScope.newConfinedScope());//null
            MemorySegment user_data = MemorySegment.allocateNative(0,ResourceScope.newConfinedScope());//null

            return (int) handle.invokeExact((MemorySegment)program.address(), 0, (MemorySegment)device_list.address(), (MemorySegment)null, (MemorySegment)pfn_notify.address(), (MemorySegment)user_data.address());
        }

        public static MemorySegment clCreateKernel(SymbolLookup openCLLinker, MemorySegment program, String kernelName) throws Throwable {

            // 获取 clCreateKernel 函数地址
            MemorySegment clCreateKernelAddress = openCLLinker.find("clCreateKernel").orElseThrow(() -> new RuntimeException("clCreateKernel not found"));

            // 定义函数类型
            MethodType clCreateKernelType = MethodType.methodType(MemorySegment.class, MemorySegment.class, MemorySegment.class, MemorySegment.class);

            // 创建函数句柄
            MethodHandle clCreateKernelMH = MethodHandles.lookup().findVirtual(ForeignLinker.class, "downcallHandle", MethodType.methodType(MethodHandle.class, MemorySegment.class, MethodType.

Project Panama 与 OpenCL 内核调用：Local 工作组大小配置疑难解答

问题描述：Local 工作组大小配置无效

案例分析：一个简单的向量加法内核

发表回复 取消回复

发表回复取消回复