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 23e38cbd1..2fe085dfe 100644 --- a/jme3-core/src/main/java/com/jme3/opencl/Context.java +++ b/jme3-core/src/main/java/com/jme3/opencl/Context.java @@ -39,6 +39,8 @@ import com.jme3.opencl.Image.ImageFormat; import com.jme3.opencl.Image.ImageType; import com.jme3.scene.VertexBuffer; import com.jme3.scene.mesh.IndexBuffer; +import com.jme3.texture.FrameBuffer; +import com.jme3.texture.Texture; import java.io.BufferedReader; import java.io.IOException; import java.io.InputStreamReader; @@ -78,9 +80,23 @@ public abstract class Context { public abstract ImageFormat[] querySupportedFormats(MemoryAccess access, ImageType type); //Interop - public abstract Buffer bindVertexBuffer(VertexBuffer vb); - public abstract Buffer bindIndexBuffer(IndexBuffer ib); - public abstract Image bindImage(com.jme3.texture.Image image); + public abstract Buffer bindVertexBuffer(VertexBuffer vb, MemoryAccess access); + + public abstract Image bindImage(com.jme3.texture.Image image, Texture.Type textureType, int miplevel, MemoryAccess access); + public Image bindImage(Texture texture, int miplevel, MemoryAccess access) { + return bindImage(texture.getImage(), texture.getType(), miplevel, access); + } + public Image bindImage(Texture texture, MemoryAccess access) { + return bindImage(texture, 0, access); + } + public Image bindRenderBuffer(FrameBuffer.RenderBuffer buffer, MemoryAccess access) { + if (buffer.getTexture() == null) { + return bindPureRenderBuffer(buffer, access); + } else { + return bindImage(buffer.getTexture(), access); + } + } + protected abstract Image bindPureRenderBuffer(FrameBuffer.RenderBuffer buffer, MemoryAccess access); public abstract Program createProgramFromSourceCode(String sourceCode); diff --git a/jme3-core/src/main/java/com/jme3/opencl/Image.java b/jme3-core/src/main/java/com/jme3/opencl/Image.java index d69d6d7a4..564920f64 100644 --- a/jme3-core/src/main/java/com/jme3/opencl/Image.java +++ b/jme3-core/src/main/java/com/jme3/opencl/Image.java @@ -233,4 +233,8 @@ public interface Image { Event fillAsync(CommandQueue queue, long[] origin, long[] region, int[] color); Event copyToBufferAsync(CommandQueue queue, Buffer dest, long[] srcOrigin, long[] srcRegion, long destOffset); + + + Event acquireImageForSharingAsync(CommandQueue queue); + Event releaseImageForSharingAsync(CommandQueue queue); } diff --git a/jme3-core/src/main/java/com/jme3/opencl/Kernel.java b/jme3-core/src/main/java/com/jme3/opencl/Kernel.java index c726ae27a..8e9de6d70 100644 --- a/jme3-core/src/main/java/com/jme3/opencl/Kernel.java +++ b/jme3-core/src/main/java/com/jme3/opencl/Kernel.java @@ -150,6 +150,8 @@ public abstract class Kernel { setArg(index, (LocalMem) arg); } else if (arg instanceof Buffer) { setArg(index, (Buffer) arg); + } else if (arg instanceof Image) { + setArg(index, (Image) arg); } else { throw new IllegalArgumentException("unknown kernel argument type: " + arg); } diff --git a/jme3-examples/src/main/java/jme3test/opencl/HelloOpenCL.java b/jme3-examples/src/main/java/jme3test/opencl/HelloOpenCL.java index 79455af74..95d052aed 100644 --- a/jme3-examples/src/main/java/jme3test/opencl/HelloOpenCL.java +++ b/jme3-examples/src/main/java/jme3test/opencl/HelloOpenCL.java @@ -46,9 +46,10 @@ import java.util.Objects; import java.util.logging.Level; import java.util.logging.Logger; -/** Sample 1 - how to get started with the most simple JME 3 application. - * Display a blue 3D cube and view from all sides by - * moving the mouse and pressing the WASD keys. */ +/** + * Simple test checking if the basic functions of the OpenCL wrapper work + * @author Sebastian Weiss + */ public class HelloOpenCL extends SimpleApplication { private static final Logger LOG = Logger.getLogger(HelloOpenCL.class.getName()); diff --git a/jme3-examples/src/main/java/jme3test/opencl/TestWriteToTexture.java b/jme3-examples/src/main/java/jme3test/opencl/TestWriteToTexture.java new file mode 100644 index 000000000..eb898278c --- /dev/null +++ b/jme3-examples/src/main/java/jme3test/opencl/TestWriteToTexture.java @@ -0,0 +1,164 @@ +/* + * Copyright (c) 2009-2012 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 jme3test.opencl; + +import com.jme3.app.SimpleApplication; +import com.jme3.input.MouseInput; +import com.jme3.input.controls.ActionListener; +import com.jme3.input.controls.AnalogListener; +import com.jme3.input.controls.MouseAxisTrigger; +import com.jme3.input.controls.MouseButtonTrigger; +import com.jme3.math.Vector2f; +import com.jme3.opencl.*; +import com.jme3.system.AppSettings; +import com.jme3.texture.Texture2D; +import com.jme3.ui.Picture; +import java.util.logging.Logger; + +/** + * This test class tests the capability to write to a GL texture from OpenCL + * @author Sebastian Weiss + */ +public class TestWriteToTexture extends SimpleApplication implements AnalogListener, ActionListener { + private static final Logger LOG = Logger.getLogger(TestWriteToTexture.class.getName()); + private static final float MOUSE_SPEED = 0.5f; + + private Texture2D tex; + private int initCounter; + private Context clContext; + private CommandQueue clQueue; + private Kernel kernel; + private Vector2f C; + private Image texCL; + private boolean dragging; + + public static void main(String[] args){ + TestWriteToTexture app = new TestWriteToTexture(); + AppSettings settings = new AppSettings(true); + settings.setOpenCLSupport(true); + settings.setVSync(true); + app.setSettings(settings); + app.start(); // start the game + } + + @Override + public void simpleInitApp() { + initOpenCL1(); + + tex = new Texture2D(settings.getWidth(), settings.getHeight(), 1, com.jme3.texture.Image.Format.RGBA8); + Picture pic = new Picture("julia"); + pic.setTexture(assetManager, tex, true); + pic.setPosition(0, 0); + pic.setWidth(settings.getWidth()); + pic.setHeight(settings.getHeight()); + guiNode.attachChild(pic); + + initCounter = 0; + + flyCam.setEnabled(false); + inputManager.setCursorVisible(true); + inputManager.addMapping("right", new MouseAxisTrigger(MouseInput.AXIS_X, false)); + inputManager.addMapping("left", new MouseAxisTrigger(MouseInput.AXIS_X, true)); + inputManager.addMapping("up", new MouseAxisTrigger(MouseInput.AXIS_Y, false)); + inputManager.addMapping("down", new MouseAxisTrigger(MouseInput.AXIS_Y, true)); + inputManager.addMapping("drag", new MouseButtonTrigger(MouseInput.BUTTON_LEFT)); + inputManager.addListener(this, "right", "left", "up", "down", "drag"); + dragging = false; + } + + @Override + public void simpleUpdate(float tpf) { + super.simpleUpdate(tpf); + + if (initCounter < 2) { + initCounter++; + } else if (initCounter == 2) { + //when initCounter reaches 2, the scene was drawn once and the texture was uploaded to the GPU + //then we can bind the texture to OpenCL + initOpenCL2(); + updateOpenCL(tpf); + initCounter = 3; + } else { + updateOpenCL(tpf); + } + } + + private void initOpenCL1() { + clContext = context.getOpenCLContext(); + clQueue = clContext.createQueue(); + //create kernel + Program program = clContext.createProgramFromSourceFiles(assetManager, "jme3test/opencl/JuliaSet.cl"); + program.build(); + kernel = program.createKernel("JuliaSet"); + C = new Vector2f(0.12f, -0.2f); + } + private void initOpenCL2() { + //bind image to OpenCL + texCL = clContext.bindImage(tex, MemoryAccess.WRITE_ONLY); + } + private void updateOpenCL(float tpf) { + //aquire resource + texCL.acquireImageForSharingAsync(clQueue); + //no need to wait for the returned event, since the kernel implicitely waits for it (same command queue) + + //execute kernel + kernel.Run1(clQueue, new WorkSize(settings.getWidth(), settings.getHeight()), texCL, C, 16); + + //release resource + texCL.releaseImageForSharingAsync(clQueue); + } + + @Override + public void onAnalog(String name, float value, float tpf) { + if (!dragging) { + return; + } + if ("left".equals(name)) { + C.x -= tpf * MOUSE_SPEED; + } else if ("right".equals(name)) { + C.x += tpf * MOUSE_SPEED; + } else if ("up".equals(name)) { + C.y -= tpf * MOUSE_SPEED; + } else if ("down".equals(name)) { + C.y += tpf * MOUSE_SPEED; + } + } + + @Override + public void onAction(String name, boolean isPressed, float tpf) { + if ("drag".equals(name)) { + dragging = isPressed; + inputManager.setCursorVisible(!isPressed); + } + } +} \ No newline at end of file diff --git a/jme3-examples/src/main/resources/jme3test/opencl/JuliaSet.cl b/jme3-examples/src/main/resources/jme3test/opencl/JuliaSet.cl new file mode 100644 index 000000000..d2a42d338 --- /dev/null +++ b/jme3-examples/src/main/resources/jme3test/opencl/JuliaSet.cl @@ -0,0 +1,99 @@ + + +//2 component vector to hold the real and imaginary parts of a complex number: +typedef float2 cfloat; + +#define I ((cfloat)(0.0, 1.0)) + +inline float real(cfloat a){ + return a.x; +} +inline float imag(cfloat a){ + return a.y; +} + +inline float cmod(cfloat a){ + return (sqrt(a.x*a.x + a.y*a.y)); +} + +inline cfloat cadd(cfloat a, cfloat b){ + return (cfloat)( a.x + b.x, a.y + b.y); +} + +inline float carg(cfloat a){ + if(a.x > 0){ + return atan(a.y / a.x); + + }else if(a.x < 0 && a.y >= 0){ + return atan(a.y / a.x) + M_PI_F; + + }else if(a.x < 0 && a.y < 0){ + return atan(a.y / a.x) - M_PI_F; + + }else if(a.x == 0 && a.y > 0){ + return M_PI_F/2; + + }else if(a.x == 0 && a.y < 0){ + return -M_PI_F/2; + + }else{ + return 0; + } +} + +inline cfloat cmult(cfloat a, cfloat b){ + return (cfloat)( a.x*b.x - a.y*b.y, a.x*b.y + a.y*b.x); +} + +inline cfloat csqrt(cfloat a){ + return (cfloat)( sqrt(cmod(a)) * cos(carg(a)/2), sqrt(cmod(a)) * sin(carg(a)/2)); +} + +inline float4 getColor(int iteration, int numIterations) { + //color transition: black -> red -> blue -> white + int step = numIterations / 2; + if (iteration < step) { + return mix( (float4)(0,0,0,1), (float4)(1,0,0,1), iteration / (float) step); + } else { + return mix( (float4)(1,0,0,1), (float4)(0,0,1,1), (iteration-step) / (float) (numIterations - step)); + } +} + +__kernel void JuliaSet(write_only image2d_t outputImage, const cfloat C, int numIterations) +{ + // get id of element in array + int x = get_global_id(0); + int y = get_global_id(1); + int w = get_global_size(0); + int h = get_global_size(1); + + cfloat Z = { ( -w / 2 + x) / (w/4.0f) , ( -h / 2 + y) / (h/4.0f) }; + int iteration = 0; + + while (iteration < numIterations) + { + cfloat Zpow2 = cmult(Z, Z); + cfloat Zn = cadd(Zpow2, C); + Z.x = Zn.x; + Z.y = Zn.y; + iteration++; + if(cmod(Z) > 2) + { + break; + } + } + + float4 color; + + // threshold reached mark pixel as white + if (iteration == numIterations) + { + color = (float4)(1,1,1,1); + } + else + { + color = getColor(iteration, numIterations); + } + + write_imagef(outputImage, (int2)(x, y), color); +} \ 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 d9f4ab0d7..b9f89d676 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,26 +31,22 @@ */ 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.opencl.Context; import com.jme3.opencl.Image.ImageDescriptor; import com.jme3.opencl.Image.ImageFormat; import com.jme3.scene.VertexBuffer; import com.jme3.scene.mesh.IndexBuffer; -import java.io.BufferedReader; -import java.io.IOException; -import java.io.InputStreamReader; +import com.jme3.texture.FrameBuffer; +import com.jme3.texture.Texture; import java.nio.ByteBuffer; import java.nio.IntBuffer; -import java.util.Arrays; import java.util.List; import java.util.logging.Level; import java.util.logging.Logger; import org.lwjgl.BufferUtils; import org.lwjgl.opencl.*; -import sun.misc.IOUtils; +import org.lwjgl.opengl.*; /** * @@ -154,23 +150,58 @@ public class LwjglContext extends Context { } @Override - public Buffer bindVertexBuffer(VertexBuffer vb) { - throw new UnsupportedOperationException("Not supported yet."); + public Buffer bindVertexBuffer(VertexBuffer vb, MemoryAccess access) { + int id = vb.getId(); + if (id == -1) { + throw new IllegalArgumentException("vertex buffer was not yet uploaded to the GPU or is CPU only"); + } + long flags = Utils.getMemoryAccessFlags(access); + Utils.errorBuffer.rewind(); + CLMem mem = CL10GL.clCreateFromGLBuffer(context, flags, id, Utils.errorBuffer); + Utils.checkError(Utils.errorBuffer, "clCreateFromGLBuffer"); + return new LwjglBuffer(mem); } @Override - public Buffer bindIndexBuffer(IndexBuffer ib) { - throw new UnsupportedOperationException("Not supported yet."); + public Image bindImage(com.jme3.texture.Image image, Texture.Type textureType, int miplevel, MemoryAccess access) { + int imageID = image.getId(); + if (imageID == -1) { + throw new IllegalArgumentException("image was not yet uploaded to the GPU"); + } + long memFlags = Utils.getMemoryAccessFlags(access); + int textureTarget = convertTextureType(textureType); + Utils.errorBuffer.rewind(); + CLMem mem = CL12GL.clCreateFromGLTexture(context, memFlags, textureTarget, miplevel, imageID, Utils.errorBuffer); + Utils.checkError(Utils.errorBuffer, "clCreateFromGLTexture"); + return new LwjglImage(mem); } @Override - public Image bindImage(com.jme3.texture.Image image) { - throw new UnsupportedOperationException("Not supported yet."); + protected Image bindPureRenderBuffer(FrameBuffer.RenderBuffer buffer, MemoryAccess access) { + int renderbuffer = buffer.getId(); + if (renderbuffer == -1) { + throw new IllegalArgumentException("renderbuffer was not yet uploaded to the GPU"); + } + long memFlags = Utils.getMemoryAccessFlags(access); + Utils.errorBuffer.rewind(); + CLMem mem = CL10GL.clCreateFromGLRenderbuffer(context, memFlags, renderbuffer, Utils.errorBuffer); + Utils.checkError(Utils.errorBuffer, "clCreateFromGLRenderbuffer"); + return new LwjglImage(mem); + } + + private int convertTextureType(Texture.Type textureType) { + switch (textureType) { + case TwoDimensional: return GL11.GL_TEXTURE_2D; + case TwoDimensionalArray: return GL30.GL_TEXTURE_2D_ARRAY; + case ThreeDimensional: return GL12.GL_TEXTURE_3D; + case CubeMap: return GL13.GL_TEXTURE_CUBE_MAP; + default: throw new IllegalArgumentException("unknown texture type "+textureType); + } } @Override public Program createProgramFromSourceCode(String sourceCode) { - LOG.log(Level.INFO, "Create program from source:\n{0}", sourceCode); + LOG.log(Level.FINE, "Create program from source:\n{0}", sourceCode); Utils.errorBuffer.rewind(); CLProgram p = CL10.clCreateProgramWithSource(context, sourceCode, Utils.errorBuffer); Utils.checkError(Utils.errorBuffer, "clCreateProgramWithSource"); diff --git a/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglImage.java b/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglImage.java index c847228c9..e9b698129 100644 --- a/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglImage.java +++ b/jme3-lwjgl/src/main/java/com/jme3/opencl/lwjgl/LwjglImage.java @@ -477,6 +477,9 @@ public class LwjglImage implements Image { long event = Utils.pointerBuffers[0].get(0); return new LwjglEvent(q.getCLEvent(event)); //TODO: why does q.getCLEvent(event) return null? + //This is a bug in LWJGL: they forgot to include the line + // if ( __result == CL_SUCCESS ) command_queue.registerCLEvent(event); + // after the native call } @Override @@ -522,4 +525,21 @@ public class LwjglImage implements Image { return new LwjglEvent(q.getCLEvent(event)); } + @Override + public Event acquireImageForSharingAsync(CommandQueue queue) { + Utils.pointerBuffers[0].rewind(); + CLCommandQueue q = ((LwjglCommandQueue) queue).getQueue(); + int ret = CL10GL.clEnqueueAcquireGLObjects(q, image, null, Utils.pointerBuffers[0]); + Utils.checkError(ret, "clEnqueueAcquireGLObjects"); + long event = Utils.pointerBuffers[0].get(0); + return new LwjglEvent(q.getCLEvent(event)); + } + public Event releaseImageForSharingAsync(CommandQueue queue) { + Utils.pointerBuffers[0].rewind(); + CLCommandQueue q = ((LwjglCommandQueue) queue).getQueue(); + int ret = CL10GL.clEnqueueReleaseGLObjects(q, image, null, Utils.pointerBuffers[0]); + Utils.checkError(ret, "clEnqueueReleaseGLObjects"); + long event = Utils.pointerBuffers[0].get(0); + return new LwjglEvent(q.getCLEvent(event)); + } }