diff --git a/jme3-core/src/main/java/com/jme3/opencl/Context.java b/jme3-core/src/main/java/com/jme3/opencl/Context.java index 8877b5823..0be677e30 100644 --- a/jme3-core/src/main/java/com/jme3/opencl/Context.java +++ b/jme3-core/src/main/java/com/jme3/opencl/Context.java @@ -31,6 +31,7 @@ */ package com.jme3.opencl; +import com.jme3.asset.AssetManager; import com.jme3.scene.VertexBuffer; import com.jme3.scene.mesh.IndexBuffer; import java.nio.ByteBuffer; @@ -123,9 +124,9 @@ public interface Context { Image bindImage(com.jme3.texture.Image image); Program createProgramFromSourceCode(String sourceCode); - Program createProgramFromSourceFilesWithInclude(String include, String... resources); - Program createProgramFormSourcesWithInclude(String include, List resources); - Program createProgramFromSources(String... resources); - Program createProgramFromSources(List resources); + Program createProgramFromSourceFilesWithInclude(AssetManager assetManager, String include, String... resources); + Program createProgramFromSourceFilesWithInclude(AssetManager assetManager, String include, List resources); + Program createProgramFromSourceFiles(AssetManager assetManager, String... resources); + Program createProgramFromSourceFiles(AssetManager assetManager, List resources); } diff --git a/jme3-core/src/main/java/com/jme3/opencl/Program.java b/jme3-core/src/main/java/com/jme3/opencl/Program.java index 32c594a04..36203264a 100644 --- a/jme3-core/src/main/java/com/jme3/opencl/Program.java +++ b/jme3-core/src/main/java/com/jme3/opencl/Program.java @@ -41,5 +41,6 @@ public interface Program { void build() throws KernelCompilationException; Kernel createKernel(String name); - + Kernel[] createAllKernels(); + } diff --git a/jme3-examples/src/main/java/jme3test/gui/opencl/HelloOpenCL.java b/jme3-examples/src/main/java/jme3test/opencl/HelloOpenCL.java similarity index 87% rename from jme3-examples/src/main/java/jme3test/gui/opencl/HelloOpenCL.java rename to jme3-examples/src/main/java/jme3test/opencl/HelloOpenCL.java index d653c0934..3ac75b824 100644 --- a/jme3-examples/src/main/java/jme3test/gui/opencl/HelloOpenCL.java +++ b/jme3-examples/src/main/java/jme3test/opencl/HelloOpenCL.java @@ -30,15 +30,12 @@ * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. */ -package jme3test.gui.opencl; +package jme3test.opencl; import com.jme3.app.SimpleApplication; import com.jme3.font.BitmapFont; import com.jme3.font.BitmapText; -import com.jme3.opencl.Buffer; -import com.jme3.opencl.CommandQueue; -import com.jme3.opencl.Context; -import com.jme3.opencl.Event; +import com.jme3.opencl.*; import com.jme3.system.AppSettings; import com.jme3.util.BufferUtils; import java.nio.ByteBuffer; @@ -78,6 +75,7 @@ public class HelloOpenCL extends SimpleApplication { .append("\n Devices: ").append(clContext.getDevices()); str.append("\nTests:"); str.append("\n Buffers: ").append(testBuffer(clContext, clQueue)); + str.append("\n Kernel: ").append(testKernel(clContext, clQueue)); BitmapText txt1 = new BitmapText(fnt); txt1.setText(str.toString()); @@ -137,4 +135,17 @@ public class HelloOpenCL extends SimpleApplication { } return true; } + + private boolean testKernel(Context clContext, CommandQueue clQueue) { + String include = "#define TYPE float\n"; + Program program = clContext.createProgramFromSourceFilesWithInclude(assetManager, include, "jme3test/opencl/Blas.cl"); + program.build(); + Kernel[] kernels = program.createAllKernels(); + for (Kernel k : kernels) { + System.out.println("available kernel: "+k.getName()); + } + Kernel kernel = program.createKernel("Fill"); + System.out.println("number of args: "+kernel.getArgCount()); + return true; + } } \ No newline at end of file diff --git a/jme3-examples/src/main/resources/jme3test/opencl/Blas.cl b/jme3-examples/src/main/resources/jme3test/opencl/Blas.cl new file mode 100644 index 000000000..a39569258 --- /dev/null +++ b/jme3-examples/src/main/resources/jme3test/opencl/Blas.cl @@ -0,0 +1,4 @@ +__kernel void Fill (__global TYPE* data, TYPE a) +{ + data[get_global_id(0)] = a; +} \ No newline at end of file diff --git a/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglContext.java b/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglContext.java index 08321d502..4d7e3be57 100644 --- a/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglContext.java +++ b/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglContext.java @@ -31,18 +31,29 @@ */ package com.jme3.opencl.lwjgl; +import com.jme3.asset.AssetInfo; +import com.jme3.asset.AssetKey; +import com.jme3.asset.AssetManager; import com.jme3.opencl.*; import com.jme3.scene.VertexBuffer; import com.jme3.scene.mesh.IndexBuffer; +import java.io.BufferedReader; +import java.io.IOException; +import java.io.InputStreamReader; import java.nio.ByteBuffer; +import java.util.Arrays; import java.util.List; +import java.util.logging.Level; +import java.util.logging.Logger; import org.lwjgl.opencl.*; +import sun.misc.IOUtils; /** * * @author Sebastian Weiss */ public class LwjglContext implements Context { + private static final Logger LOG = Logger.getLogger(LwjglContext.class.getName()); private final CLContext context; private final List devices; @@ -56,7 +67,7 @@ public class LwjglContext implements Context { } @Override - public List getDevices() { + public List getDevices() { return devices; } @@ -125,27 +136,51 @@ public class LwjglContext implements Context { @Override public Program createProgramFromSourceCode(String sourceCode) { - throw new UnsupportedOperationException("Not supported yet."); + LOG.log(Level.INFO, "Create program from source:\n{0}", sourceCode); + Utils.errorBuffer.rewind(); + CLProgram p = CL10.clCreateProgramWithSource(context, sourceCode, Utils.errorBuffer); + Utils.checkError(Utils.errorBuffer, "clCreateProgramWithSource"); + return new LwjglProgram(p, this); } @Override - public Program createProgramFromSourceFilesWithInclude(String include, String... resources) { - throw new UnsupportedOperationException("Not supported yet."); + public Program createProgramFromSourceFilesWithInclude(AssetManager assetManager, String include, String... resources) { + return createProgramFromSourceFilesWithInclude(assetManager, include, Arrays.asList(resources)); } @Override - public Program createProgramFormSourcesWithInclude(String include, List resources) { - throw new UnsupportedOperationException("Not supported yet."); + public Program createProgramFromSourceFilesWithInclude(AssetManager assetManager, String include, List resources) { + StringBuilder str = new StringBuilder(); + str.append(include); + for (String res : resources) { + AssetInfo info = assetManager.locateAsset(new AssetKey(res)); + if (info == null) { + LOG.log(Level.WARNING, "unable to load source file ''{0}''", res); + continue; + } + try (BufferedReader reader = new BufferedReader(new InputStreamReader(info.openStream()))) { + while (true) { + String line = reader.readLine(); + if (line == null) { + break; + } + str.append(line).append('\n'); + } + } catch (IOException ex) { + LOG.log(Level.WARNING, "unable to load source file '"+res+"'", ex); + } + } + return createProgramFromSourceCode(str.toString()); } @Override - public Program createProgramFromSources(String... resources) { - throw new UnsupportedOperationException("Not supported yet."); + public Program createProgramFromSourceFiles(AssetManager assetManager, String... resources) { + return createProgramFromSourceFilesWithInclude(assetManager, "", resources); } @Override - public Program createProgramFromSources(List resources) { - throw new UnsupportedOperationException("Not supported yet."); + public Program createProgramFromSourceFiles(AssetManager assetManager, List resources) { + return createProgramFromSourceFilesWithInclude(assetManager, "", resources); } } diff --git a/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglKernel.java b/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglKernel.java new file mode 100644 index 000000000..4b962a3f7 --- /dev/null +++ b/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglKernel.java @@ -0,0 +1,138 @@ +/* + * Copyright (c) 2009-2016 jMonkeyEngine + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are + * met: + * + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * + * * Neither the name of 'jMonkeyEngine' nor the names of its contributors + * may be used to endorse or promote products derived from this software + * without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS + * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED + * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR + * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF + * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING + * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS + * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ +package com.jme3.opencl.lwjgl; + +import com.jme3.math.Quaternion; +import com.jme3.math.Vector2f; +import com.jme3.math.Vector4f; +import com.jme3.opencl.*; +import java.nio.ByteBuffer; +import org.lwjgl.opencl.CL10; +import org.lwjgl.opencl.CLKernel; + +/** + * + * @author Sebastian Weiss + */ +public class LwjglKernel extends Kernel { + + private final CLKernel kernel; + + public LwjglKernel(CLKernel kernel) { + this.kernel = kernel; + } + + public CLKernel getKernel() { + return kernel; + } + + @Override + public String getName() { + return kernel.getInfoString(CL10.CL_KERNEL_FUNCTION_NAME); + } + + @Override + public int getArgCount() { + return kernel.getInfoInt(CL10.CL_KERNEL_NUM_ARGS); + } + + @Override + public void setArg(int index, LocalMemPerElement t) { + throw new UnsupportedOperationException("Not supported yet."); + } + + @Override + public void setArg(int index, LocalMem t) { + throw new UnsupportedOperationException("Not supported yet."); + } + + @Override + public void setArg(int index, Buffer t) { + throw new UnsupportedOperationException("Not supported yet."); + } + + @Override + public void setArg(int index, byte b) { + throw new UnsupportedOperationException("Not supported yet."); + } + + @Override + public void setArg(int index, short s) { + throw new UnsupportedOperationException("Not supported yet."); + } + + @Override + public void setArg(int index, int i) { + throw new UnsupportedOperationException("Not supported yet."); + } + + @Override + public void setArg(int index, long l) { + throw new UnsupportedOperationException("Not supported yet."); + } + + @Override + public void setArg(int index, float f) { + throw new UnsupportedOperationException("Not supported yet."); + } + + @Override + public void setArg(int index, double d) { + throw new UnsupportedOperationException("Not supported yet."); + } + + @Override + public void setArg(int index, Vector2f v) { + throw new UnsupportedOperationException("Not supported yet."); + } + + @Override + public void setArg(int index, Vector4f v) { + throw new UnsupportedOperationException("Not supported yet."); + } + + @Override + public void setArg(int index, Quaternion q) { + throw new UnsupportedOperationException("Not supported yet."); + } + + @Override + public void setArg(int index, ByteBuffer buffer, long size) { + throw new UnsupportedOperationException("Not supported yet."); + } + + @Override + public Event Run(CommandQueue queue) { + throw new UnsupportedOperationException("Not supported yet."); + } + +} diff --git a/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglProgram.java b/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglProgram.java new file mode 100644 index 000000000..a3239e5fd --- /dev/null +++ b/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglProgram.java @@ -0,0 +1,111 @@ +/* + * Copyright (c) 2009-2016 jMonkeyEngine + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are + * met: + * + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * + * * Neither the name of 'jMonkeyEngine' nor the names of its contributors + * may be used to endorse or promote products derived from this software + * without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS + * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED + * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR + * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF + * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING + * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS + * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ +package com.jme3.opencl.lwjgl; + +import com.jme3.opencl.Kernel; +import com.jme3.opencl.KernelCompilationException; +import com.jme3.opencl.Program; +import java.util.ArrayList; +import java.util.logging.Level; +import java.util.logging.Logger; +import org.lwjgl.PointerBuffer; +import org.lwjgl.opencl.*; + +/** + * + * @author Sebastian Weiss + */ +public class LwjglProgram implements Program { + private static final Logger LOG = Logger.getLogger(LwjglProgram.class.getName()); + + private final CLProgram program; + private final LwjglContext context; + + public LwjglProgram(CLProgram program, LwjglContext context) { + this.program = program; + this.context = context; + } + + public CLProgram getProgram() { + return program; + } + + @Override + public void build(String args) throws KernelCompilationException { + int ret = CL10.clBuildProgram(program, (PointerBuffer) null, args, null); + if (ret != CL10.CL_SUCCESS) { + String log = Log(); + LOG.log(Level.WARNING, "Unable to compile program:\n{0}", log); + if (ret == CL10.CL_BUILD_PROGRAM_FAILURE) { + throw new KernelCompilationException("Failed to build program", ret, log); + } else { + Utils.checkError(ret, "clBuildProgram"); + } + } else { + LOG.log(Level.INFO, "Program compiled:\n{0}", Log()); + } + } + + @Override + public void build() throws KernelCompilationException { + build(""); + } + + private String Log() { + StringBuilder str = new StringBuilder(); + for (LwjglDevice device : context.getDevices()) { + CLDevice d = device.getDevice(); + str.append(device.getName()).append(":\n"); + str.append(program.getBuildInfoString(d, CL10.CL_PROGRAM_BUILD_LOG)); + str.append('\n'); + } + return str.toString(); + } + + @Override + public Kernel createKernel(String name) { + CLKernel kernel = CL10.clCreateKernel(program, name, Utils.errorBuffer); + Utils.checkError(Utils.errorBuffer, "clCreateKernel"); + return new LwjglKernel(kernel); + } + + @Override + public Kernel[] createAllKernels() { + CLKernel[] kernels = program.createKernelsInProgram(); + Kernel[] kx = new Kernel[kernels.length]; + for (int i=0; i