added querying of the program binaries and building the programs from these binaries.

TestVertexBufferSharing shows how this is used to build a simple program cache.

jme3-core/src/main/java/com/jme3/opencl/Context.java

@@ -367,4 +367,22 @@ public abstract class Context extends AbstractOpenCLObject {
     public Program createProgramFromSourceFiles(AssetManager assetManager, List<String> resources) {
         return createProgramFromSourceFilesWithInclude(assetManager, "", resources);
     }
+    
+    /**
+     * Creates a program from the specified binaries.
+     * The binaries are created by {@link Program#getBinary(com.jme3.opencl.Device) }.
+     * The returned program still needs to be build using
+     * {@link Program#build(java.lang.String, com.jme3.opencl.Device...) }.
+     * <b>Important:</b>The device passed to {@code Program.getBinary(..)},
+     * this method and {@code Program#build(..)} must be the same.
+     * 
+     * The binaries are used to build a program cache across multiple launches
+     * of the application. The programs build mach faster from binaries than
+     * from sources.
+     * 
+     * @param binaries the binaries
+     * @param device the device to use
+     * @return the new program
+     */
+    public abstract Program createProgramFromBinary(ByteBuffer binaries, Device device);
 }

jme3-core/src/main/java/com/jme3/opencl/Program.java

@@ -31,6 +31,8 @@
  */
 package com.jme3.opencl;
 
+import java.nio.ByteBuffer;
+
 /**
  * A wrapper for an OpenCL program. A program is created from kernel source code,
  * manages the build process and creates the kernels.
@@ -49,21 +51,27 @@ public abstract class Program extends AbstractOpenCLObject {
     }
     
     /**
-     * Builds this program with the specified argument string.
+     * Builds this program with the specified argument string on the specified
+     * devices.
      * Please see the official OpenCL specification for a definition of
-     * all supported arguments
+     * all supported arguments.
+     * The list of devices specify on which device the compiled program
+     * can then be executed. It must be a subset of {@link Context#getDevices() }.
+     * If {@code null} is passed, the program is built on all available devices.
+     * 
      * @param args the compilation arguments
+     * @param devices a list of devices on which the program is build.
      * @throws KernelCompilationException if the compilation fails
      * @see #build() 
      */
-	public abstract void build(String args) throws KernelCompilationException;
+	public abstract void build(String args, Device... devices) throws KernelCompilationException;
     /**
      * Builds this program without additional arguments
      * @throws KernelCompilationException if the compilation fails
      * @see #build(java.lang.String) 
      */
 	public void build() throws KernelCompilationException {
-        build("");
+        build("", null);
     }
 
     /**
@@ -82,4 +90,15 @@ public abstract class Program extends AbstractOpenCLObject {
      */
 	public abstract Kernel[] createAllKernels();
     
+    /**
+     * Queries a compiled binary representation of this program for a particular
+     * device. This binary can then be used e.g. in the next application launch
+     * to create the program from the binaries and not from the sources.
+     * This saves time.
+     * @param device the device from which the binaries are taken
+     * @return the binaries
+     * @see Context#createProgramFromBinary(java.nio.ByteBuffer, com.jme3.opencl.Device) 
+     */
+    public abstract ByteBuffer getBinary(Device device);
+    
 }

jme3-examples/src/main/java/jme3test/opencl/TestVertexBufferSharing.java

@@ -40,11 +40,19 @@ import com.jme3.scene.Geometry;
 import com.jme3.scene.VertexBuffer;
 import com.jme3.scene.shape.Box;
 import com.jme3.system.AppSettings;
+import com.jme3.system.JmeSystem;
+import com.jme3.util.BufferUtils;
+import java.io.File;
+import java.io.IOException;
+import java.nio.ByteBuffer;
+import java.nio.file.Files;
+import java.util.logging.Level;
 import java.util.logging.Logger;
 
 /**
  * This test class tests the capability to read and modify an OpenGL vertex buffer.
- * 
+ * It is also shown how to use the program binaries to implement a simple program
+ * cache.
  * @author shaman
  */
 public class TestVertexBufferSharing extends SimpleApplication {
@@ -106,19 +114,48 @@ public class TestVertexBufferSharing extends SimpleApplication {
     
     private void initOpenCL1() {
         clContext = context.getOpenCLContext();
-        clQueue = clContext.createQueue();
+        Device device = clContext.getDevices().get(0);
+        clQueue = clContext.createQueue(device);
         clQueue.register();
         //create kernel
-        String source = ""
+        Program program = null;
-                + "__kernel void ScaleKernel(__global float* vb, float scale)\n"
+        File tmpFolder = JmeSystem.getStorageFolder();
-                + "{\n"
+        File binaryFile = new File(tmpFolder, getClass().getSimpleName()+".clc");
-                + "  int idx = get_global_id(0);\n"
+        try {
-                + "  float3 pos = vload3(idx, vb);\n"
+            //attempt to load cached binary
-                + "  pos *= scale;\n"
+            byte[] bytes = Files.readAllBytes(binaryFile.toPath());
-                + "  vstore3(pos, idx, vb);\n"
+            ByteBuffer bb = BufferUtils.createByteBuffer(bytes);
-                + "}\n";
+            program = clContext.createProgramFromBinary(bb, device);
-        Program program = clContext.createProgramFromSourceCode(source);
+            program.build();
-        program.build();
+            LOG.info("reuse program from cached binaries");
+        } catch (java.nio.file.NoSuchFileException ex) {
+            //do nothing, cache was not created yet
+        } catch (Exception ex) {
+            LOG.log(Level.INFO, "Unable to use cached program binaries", ex);
+        }
+        if (program == null) {
+            //build from sources
+            String source = ""
+                    + "__kernel void ScaleKernel(__global float* vb, float scale)\n"
+                    + "{\n"
+                    + "  int idx = get_global_id(0);\n"
+                    + "  float3 pos = vload3(idx, vb);\n"
+                    + "  pos *= scale;\n"
+                    + "  vstore3(pos, idx, vb);\n"
+                    + "}\n";
+            program = clContext.createProgramFromSourceCode(source);
+            program.build();
+            //Save binary
+            ByteBuffer bb = program.getBinary(device);
+            byte[] bytes = new byte[bb.remaining()];
+            bb.get(bytes);
+            try {
+                Files.write(binaryFile.toPath(), bytes);
+            } catch (IOException ex) {
+               LOG.log(Level.SEVERE, "Unable to save program binaries", ex);
+            }
+            LOG.info("create new program from sources");
+        }
         program.register();
         kernel = program.createKernel("ScaleKernel");
         kernel.register();

jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglContext.java

@@ -208,6 +208,17 @@ public class LwjglContext extends Context {
         return new LwjglProgram(p, this);
     }
 
+    @Override
+    public Program createProgramFromBinary(ByteBuffer binaries, Device device) {
+        Utils.errorBuffer.rewind();
+        Utils.tempBuffers[0].b16i.rewind();
+        CLProgram p = CL10.clCreateProgramWithBinary(context, ((LwjglDevice) device).device, 
+                binaries, Utils.tempBuffers[0].b16i, Utils.errorBuffer);
+        Utils.checkError(Utils.errorBuffer, "clCreateProgramWithSource");
+        Utils.checkError(Utils.tempBuffers[0].b16i, "clCreateProgramWithSource");
+        return new LwjglProgram(p, this);
+    }
+
     private static class ReleaserImpl implements ObjectReleaser {
         private CLContext context;
         private final List<LwjglDevice> devices;

jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglProgram.java

@@ -31,10 +31,8 @@
  */
 package com.jme3.opencl.lwjgl;
 
-import com.jme3.opencl.Kernel;
+import com.jme3.opencl.*;
-import com.jme3.opencl.KernelCompilationException;
+import java.nio.ByteBuffer;
-import com.jme3.opencl.OpenCLObjectManager;
-import com.jme3.opencl.Program;
 import java.util.logging.Level;
 import java.util.logging.Logger;
 import org.lwjgl.PointerBuffer;
@@ -61,8 +59,17 @@ public class LwjglProgram extends Program {
     }
 
     @Override
-    public void build(String args) throws KernelCompilationException {
+    public void build(String args, Device... devices) throws KernelCompilationException {
-        int ret = CL10.clBuildProgram(program, (PointerBuffer) null, args, null);
+        PointerBuffer deviceList = null;
+        if (devices != null) {
+            deviceList = PointerBuffer.allocateDirect(devices.length);
+            deviceList.rewind();
+            for (Device d : devices) {
+                deviceList.put(((LwjglDevice) d).device.getPointer());
+            }
+            deviceList.flip();
+        }
+        int ret = CL10.clBuildProgram(program, deviceList, args, null);
         if (ret != CL10.CL_SUCCESS) {
             String log = Log();
             LOG.log(Level.WARNING, "Unable to compile program:\n{0}", log);
@@ -104,6 +111,20 @@ public class LwjglProgram extends Program {
         return kx;
     }
 
+    @Override
+    public ByteBuffer getBinary(Device device) {
+        ByteBuffer[] binaries = program.getInfoBinaries((ByteBuffer[]) null);
+        CLDevice[] devices = program.getInfoDevices();
+        //find the requested one
+        assert (binaries.length == devices.length);
+        for (int i=0; i<devices.length; ++i) {
+            if (((LwjglDevice) device).device == devices[i]) {
+                return binaries[i];
+            }
+        }
+        throw new com.jme3.opencl.OpenCLException("Program was not built against the specified device "+device);
+    }
+
     private static class ReleaserImpl implements ObjectReleaser {
         private CLProgram program;
         private ReleaserImpl(CLProgram program) {