From a26e526945aba00220a99fb0d734b04257b3163b Mon Sep 17 00:00:00 2001 From: shamanDevel Date: Sun, 1 May 2016 18:38:48 +0200 Subject: [PATCH] Added an experimental Jocl binding. Note that Jogamp's Jocl only supports OpenCL1.1, some methods will throw an UnsupportedOperationException. --- .../java/jme3test/opencl/HelloOpenCL.java | 7 +- .../opencl/TestVertexBufferSharing.java | 1 + .../jme3test/opencl/TestWriteToTexture.java | 3 +- jme3-jogl/build.gradle | 1 + .../java/com/jme3/opencl/jocl/JoclBuffer.java | 225 ++++++++ .../jme3/opencl/jocl/JoclCommandQueue.java | 90 +++ .../com/jme3/opencl/jocl/JoclContext.java | 249 ++++++++ .../java/com/jme3/opencl/jocl/JoclDevice.java | 302 ++++++++++ .../java/com/jme3/opencl/jocl/JoclEvent.java | 103 ++++ .../java/com/jme3/opencl/jocl/JoclImage.java | 533 ++++++++++++++++++ .../java/com/jme3/opencl/jocl/JoclKernel.java | 261 +++++++++ .../com/jme3/opencl/jocl/JoclPlatform.java | 127 +++++ .../com/jme3/opencl/jocl/JoclProgram.java | 152 +++++ .../main/java/com/jme3/opencl/jocl/Utils.java | 160 ++++++ .../com/jme3/system/jogl/JoglContext.java | 98 ++++ .../com/jme3/opencl/lwjgl/LwjglImage.java | 9 +- 16 files changed, 2313 insertions(+), 8 deletions(-) create mode 100644 jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclBuffer.java create mode 100644 jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclCommandQueue.java create mode 100644 jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclContext.java create mode 100644 jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclDevice.java create mode 100644 jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclEvent.java create mode 100644 jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclImage.java create mode 100644 jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclKernel.java create mode 100644 jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclPlatform.java create mode 100644 jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclProgram.java create mode 100644 jme3-jogl/src/main/java/com/jme3/opencl/jocl/Utils.java diff --git a/jme3-examples/src/main/java/jme3test/opencl/HelloOpenCL.java b/jme3-examples/src/main/java/jme3test/opencl/HelloOpenCL.java index a4db6dd64..4100fd92b 100644 --- a/jme3-examples/src/main/java/jme3test/opencl/HelloOpenCL.java +++ b/jme3-examples/src/main/java/jme3test/opencl/HelloOpenCL.java @@ -58,6 +58,7 @@ public class HelloOpenCL extends SimpleApplication { AppSettings settings = new AppSettings(true); settings.setOpenCLSupport(true); settings.setVSync(true); +// settings.setRenderer(AppSettings.JOGL_OPENGL_FORWARD_COMPATIBLE); app.setSettings(settings); app.start(); // start the game } @@ -208,7 +209,11 @@ public class HelloOpenCL extends SimpleApplication { //query supported formats for (MemoryAccess ma : MemoryAccess.values()) { for (Image.ImageType type : Image.ImageType.values()) { - System.out.println("Formats for "+ma+" and "+type+": "+Arrays.toString(clContext.querySupportedFormats(ma, type))); + try { + System.out.println("Formats for " + ma + " and " + type + ": " + Arrays.toString(clContext.querySupportedFormats(ma, type))); + } catch (UnsupportedOperationException e) { + LOG.warning(e.getLocalizedMessage()); + } } } diff --git a/jme3-examples/src/main/java/jme3test/opencl/TestVertexBufferSharing.java b/jme3-examples/src/main/java/jme3test/opencl/TestVertexBufferSharing.java index 76c2e01f4..06e2684dc 100644 --- a/jme3-examples/src/main/java/jme3test/opencl/TestVertexBufferSharing.java +++ b/jme3-examples/src/main/java/jme3test/opencl/TestVertexBufferSharing.java @@ -64,6 +64,7 @@ public class TestVertexBufferSharing extends SimpleApplication { AppSettings settings = new AppSettings(true); settings.setOpenCLSupport(true); settings.setVSync(false); +// settings.setRenderer(AppSettings.JOGL_OPENGL_FORWARD_COMPATIBLE); app.setSettings(settings); app.start(); // start the game } diff --git a/jme3-examples/src/main/java/jme3test/opencl/TestWriteToTexture.java b/jme3-examples/src/main/java/jme3test/opencl/TestWriteToTexture.java index 43a6c64fe..5230fb69d 100644 --- a/jme3-examples/src/main/java/jme3test/opencl/TestWriteToTexture.java +++ b/jme3-examples/src/main/java/jme3test/opencl/TestWriteToTexture.java @@ -63,13 +63,13 @@ public class TestWriteToTexture extends SimpleApplication implements AnalogListe private Vector2f C; private Image texCL; private boolean dragging; - private int gcCounter; public static void main(String[] args){ TestWriteToTexture app = new TestWriteToTexture(); AppSettings settings = new AppSettings(true); settings.setOpenCLSupport(true); settings.setVSync(false); +// settings.setRenderer(AppSettings.JOGL_OPENGL_FORWARD_COMPATIBLE); app.setSettings(settings); app.start(); // start the game } @@ -87,7 +87,6 @@ public class TestWriteToTexture extends SimpleApplication implements AnalogListe guiNode.attachChild(pic); initCounter = 0; - gcCounter = 0; flyCam.setEnabled(false); inputManager.setCursorVisible(true); diff --git a/jme3-jogl/build.gradle b/jme3-jogl/build.gradle index df24c15ec..b71be027a 100644 --- a/jme3-jogl/build.gradle +++ b/jme3-jogl/build.gradle @@ -8,4 +8,5 @@ dependencies { compile 'org.jogamp.gluegen:gluegen-rt-main:2.3.2' compile 'org.jogamp.jogl:jogl-all-main:2.3.2' compile 'org.jogamp.joal:joal-main:2.3.2' + compile 'org.jogamp.jocl:jocl-main:2.3.2' } diff --git a/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclBuffer.java b/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclBuffer.java new file mode 100644 index 000000000..773b581b1 --- /dev/null +++ b/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclBuffer.java @@ -0,0 +1,225 @@ +/* + * 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.jocl; + +import com.jme3.opencl.*; +import java.nio.ByteBuffer; +import com.jogamp.opencl.*; +import com.jogamp.opencl.llb.CL; +import com.jogamp.opencl.llb.gl.CLGL; +import java.util.EnumSet; + +/** + * + * @author shaman + */ +public class JoclBuffer extends Buffer { + + final long id; + final CL cl; + + public JoclBuffer(long id) { + this.id = id; + this.cl = CLPlatform.getLowLevelCLInterface(); + OpenCLObjectManager.getInstance().registerObject(this); + } + + @Override + public long getSize() { + Utils.pointers[0].rewind(); + int ret = cl.clGetMemObjectInfo(id, CL.CL_MEM_SIZE, Utils.pointers[0].elementSize(), Utils.pointers[0].getBuffer(), null); + Utils.checkError(ret, "clGetMemObjectInfo"); + return Utils.pointers[0].get(); + } + + @Override + public MemoryAccess getMemoryAccessFlags() { + Utils.pointers[0].rewind(); + int ret = cl.clGetMemObjectInfo(id, CL.CL_MEM_TYPE, Utils.pointers[0].elementSize(), Utils.pointers[0].getBuffer(), null); + Utils.checkError(ret, "clGetMemObjectInfo"); + long flags = Utils.pointers[0].get(); + return Utils.getMemoryAccessFromFlag(flags); + } + + @Override + public void read(CommandQueue queue, ByteBuffer dest, long size, long offset) { + long q = ((JoclCommandQueue) queue).id; + int ret = cl.clEnqueueReadBuffer(q, id, CL.CL_TRUE, offset, size, dest, 0, null, null); + Utils.checkError(ret, "clEnqueueReadBuffer"); + } + + @Override + public Event readAsync(CommandQueue queue, ByteBuffer dest, long size, long offset) { + Utils.pointers[0].rewind(); + long q = ((JoclCommandQueue) queue).id; + int ret = cl.clEnqueueReadBuffer(q, id, CL.CL_FALSE, offset, size, dest, 0, null, Utils.pointers[0]); + Utils.checkError(ret, "clEnqueueReadBuffer"); + long event = Utils.pointers[0].get(0); + return new JoclEvent(event); + } + + @Override + public void write(CommandQueue queue, ByteBuffer src, long size, long offset) { + long q = ((JoclCommandQueue)queue).id; + int ret = cl.clEnqueueWriteBuffer(q, id, CL.CL_TRUE, offset, size, src, 0, null, null); + Utils.checkError(ret, "clEnqueueWriteBuffer"); + } + + @Override + public Event writeAsync(CommandQueue queue, ByteBuffer src, long size, long offset) { + Utils.pointers[0].rewind(); + long q = ((JoclCommandQueue)queue).id; + int ret = cl.clEnqueueWriteBuffer(q, id, CL.CL_FALSE, offset, size, src, 0, null, Utils.pointers[0]); + Utils.checkError(ret, "clEnqueueWriteBuffer"); + long event = Utils.pointers[0].get(0); + return new JoclEvent(event); + } + + @Override + public void copyTo(CommandQueue queue, Buffer dest, long size, long srcOffset, long destOffset) { + Utils.pointers[0].rewind(); + long q = ((JoclCommandQueue)queue).id; + long did = ((JoclBuffer) dest).id; + int ret = cl.clEnqueueCopyBuffer(q, id, did, srcOffset, destOffset, size, 0, null, Utils.pointers[0]); + Utils.checkError(ret, "clEnqueueCopyBuffer"); + ret = cl.clWaitForEvents(1, Utils.pointers[0]); + Utils.checkError(ret, "clWaitForEvents"); + } + + @Override + public Event copyToAsync(CommandQueue queue, Buffer dest, long size, long srcOffset, long destOffset) { + Utils.pointers[0].rewind(); + long q = ((JoclCommandQueue)queue).id; + long did = ((JoclBuffer) dest).id; + int ret = cl.clEnqueueCopyBuffer(q, id, did, srcOffset, destOffset, size, 0, null, Utils.pointers[0]); + Utils.checkError(ret, "clEnqueueCopyBuffer"); + long event = Utils.pointers[0].get(0); + return new JoclEvent(event); + } + + @Override + public ByteBuffer map(CommandQueue queue, long size, long offset, MappingAccess access) { + long q = ((JoclCommandQueue)queue).id; + Utils.errorBuffer.rewind(); + long flags = Utils.getMappingAccessFlags(access); + ByteBuffer b = cl.clEnqueueMapBuffer(q, id, CL.CL_TRUE, flags, offset, size, 0, null, null, Utils.errorBuffer); + Utils.checkError(Utils.errorBuffer, "clEnqueueMapBuffer"); + return b; + } + + @Override + public void unmap(CommandQueue queue, ByteBuffer ptr) { + long q = ((JoclCommandQueue)queue).id; + Utils.pointers[0].rewind(); + int ret = cl.clEnqueueUnmapMemObject(q, id, ptr, 0, null, Utils.pointers[0]); + Utils.checkError(ret, "clEnqueueUnmapMemObject"); + ret = cl.clWaitForEvents(1, Utils.pointers[0]); + Utils.checkError(ret, "clWaitForEvents"); + } + + @Override + public com.jme3.opencl.Buffer.AsyncMapping mapAsync(CommandQueue queue, long size, long offset, MappingAccess access) { + long q = ((JoclCommandQueue)queue).id; + Utils.pointers[0].rewind(); + Utils.errorBuffer.rewind(); + long flags = Utils.getMappingAccessFlags(access); + ByteBuffer b = cl.clEnqueueMapBuffer(q, id, CL.CL_FALSE, flags, offset, size, 0, null, Utils.pointers[0], Utils.errorBuffer); + Utils.checkError(Utils.errorBuffer, "clEnqueueMapBuffer"); + long event = Utils.pointers[0].get(0); + return new com.jme3.opencl.Buffer.AsyncMapping(new JoclEvent(event), b); + } + + @Override + public Event fillAsync(CommandQueue queue, ByteBuffer pattern, long size, long offset) { + throw new UnsupportedOperationException("Not supported by Jocl!"); + } + + @Override + public Event copyToImageAsync(CommandQueue queue, Image dest, long srcOffset, long[] destOrigin, long[] destRegion) { + if (destOrigin.length!=3 || destRegion.length!=3) { + throw new IllegalArgumentException("origin and region must both be arrays of length 3"); + } + Utils.pointers[0].rewind(); + Utils.pointers[1].rewind(); + Utils.pointers[2].rewind(); + Utils.pointers[1].put(destOrigin[0]).put(destOrigin[1]).put(destOrigin[2]).position(0); + Utils.pointers[2].put(destRegion[0]).put(destRegion[1]).put(destRegion[2]).position(0); + long q = ((JoclCommandQueue)queue).id; + long i = ((JoclImage) dest).id; + int ret = cl.clEnqueueCopyBufferToImage(q, id, i, srcOffset, Utils.pointers[1], Utils.pointers[2], 0, null, Utils.pointers[0]); + Utils.checkError(ret, "clEnqueueCopyBufferToImage"); + long event = Utils.pointers[0].get(0); + return new JoclEvent(event); + } + + @Override + public Event acquireBufferForSharingAsync(CommandQueue queue) { + Utils.pointers[0].rewind(); + Utils.pointers[1].rewind(); + Utils.pointers[1].put(0, id); + long q = ((JoclCommandQueue)queue).id; + ((CLGL) cl).clEnqueueAcquireGLObjects(q, 1, Utils.pointers[1], 0, null, Utils.pointers[0]); + long event = Utils.pointers[0].get(0); + return new JoclEvent(event); + } + + @Override + public Event releaseBufferForSharingAsync(CommandQueue queue) { + Utils.pointers[0].rewind(); + Utils.pointers[1].rewind(); + Utils.pointers[1].put(0, id); + long q = ((JoclCommandQueue)queue).id; + ((CLGL) cl).clEnqueueReleaseGLObjects(q, 1, Utils.pointers[1], 0, null, Utils.pointers[0]); + long event = Utils.pointers[0].get(0); + return new JoclEvent(event); + } + + @Override + public ObjectReleaser getReleaser() { + return new ReleaserImpl(id); + } + private static class ReleaserImpl implements ObjectReleaser { + private long mem; + private ReleaserImpl(long mem) { + this.mem = mem; + } + @Override + public void release() { + if (mem != 0) { + int ret = CLPlatform.getLowLevelCLInterface().clReleaseMemObject(mem); + mem = 0; + Utils.reportError(ret, "clReleaseMemObject"); + } + } + + } +} diff --git a/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclCommandQueue.java b/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclCommandQueue.java new file mode 100644 index 000000000..790bc152b --- /dev/null +++ b/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclCommandQueue.java @@ -0,0 +1,90 @@ +/* + * 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.jocl; + +import com.jme3.opencl.CommandQueue; +import com.jme3.opencl.OpenCLObjectManager; +import com.jogamp.opencl.CLCommandQueue; +import com.jogamp.opencl.CLPlatform; +import com.jogamp.opencl.llb.CL; +import com.jogamp.opencl.llb.CLCommandQueueBinding; + +/** + * + * @author shaman + */ +public class JoclCommandQueue implements CommandQueue { + + final CL cl; + final long id; + + public JoclCommandQueue(long id) { + this.id = id; + this.cl = CLPlatform.getLowLevelCLInterface(); + OpenCLObjectManager.getInstance().registerObject(this); + } + + @Override + public void flush() { + int ret = cl.clFlush(id); + Utils.checkError(ret, "clFlush"); + } + + @Override + public void finish() { + int ret = cl.clFinish(id); + Utils.checkError(ret, "clFinish"); + } + + @Override + public ObjectReleaser getReleaser() { + return new ReleaserImpl(id, cl); + } + private static class ReleaserImpl implements ObjectReleaser { + private long id; + private CLCommandQueueBinding cl; + + private ReleaserImpl(long id, CLCommandQueueBinding cl) { + this.id = id; + this.cl = cl; + } + + @Override + public void release() { + if (id != 0) { + int ret = cl.clReleaseCommandQueue(id); + id = 0; + Utils.reportError(ret, "clReleaseCommandQueue"); + } + } + } +} diff --git a/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclContext.java b/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclContext.java new file mode 100644 index 000000000..4e9b6e920 --- /dev/null +++ b/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclContext.java @@ -0,0 +1,249 @@ +/* + * 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.jocl; + +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.texture.FrameBuffer; +import com.jme3.texture.Texture; +import com.jogamp.opencl.CLContext; +import com.jogamp.opencl.CLImageFormat; +import com.jogamp.opencl.CLMemory; +import com.jogamp.opencl.CLMemory.Mem; +import com.jogamp.opencl.CLPlatform; +import com.jogamp.opencl.llb.CL; +import com.jogamp.opencl.llb.gl.CLGL; +import com.jogamp.opencl.llb.impl.CLImageFormatImpl; +import com.jogamp.opengl.GL; +import java.nio.ByteBuffer; +import java.nio.IntBuffer; +import java.util.List; +import java.util.logging.Level; +import java.util.logging.Logger; + +/** + * + * @author shaman + */ +public class JoclContext extends Context { + private static final Logger LOG = Logger.getLogger(JoclContext.class.getName()); + + final CLContext context; + final long id; + final CL cl; + private final List devices; + + public JoclContext(CLContext context, List devices) { + this.context = context; + this.id = context.ID; + this.cl = context.getCL(); + this.devices = devices; + OpenCLObjectManager.getInstance().registerObject(this); + } + + public CLContext getContext() { + return context; + } + + @Override + public List getDevices() { + return devices; + } + + @Override + @SuppressWarnings("element-type-mismatch") + public CommandQueue createQueue(Device device) { + assert (devices.contains(device)); //this also ensures that device is a JoclDevice + long d = ((JoclDevice) device).id; + long properties = 0; + long q = cl.clCreateCommandQueue(id, d, properties, Utils.errorBuffer); + Utils.checkError(Utils.errorBuffer, "clCreateCommandQueue"); + return new JoclCommandQueue(q); + } + + @Override + public Buffer createBuffer(long size, MemoryAccess access) { + long flags = Utils.getMemoryAccessFlags(access); + long mem = cl.clCreateBuffer(id, flags, size, null, Utils.errorBuffer); + Utils.checkError(Utils.errorBuffer, "clCreateBuffer"); + return new JoclBuffer(mem); + } + + @Override + public Buffer createBufferFromHost(ByteBuffer data, MemoryAccess access) { + long flags = Utils.getMemoryAccessFlags(access); + flags |= CL.CL_MEM_USE_HOST_PTR; + long mem = cl.clCreateBuffer(id, flags, data.capacity(), data, Utils.errorBuffer); + Utils.checkError(Utils.errorBuffer, "clCreateBuffer"); + return new JoclBuffer(mem); + } + + @Override + public Image createImage(MemoryAccess access, ImageFormat format, ImageDescriptor descr) { + if (descr.type != Image.ImageType.IMAGE_2D && descr.type != Image.ImageType.IMAGE_3D) { + throw new UnsupportedOperationException("Jocl only supports 2D and 3D images"); + } + long memFlags = Utils.getMemoryAccessFlags(access); + Utils.errorBuffer.rewind(); + //fill image format + CLImageFormatImpl f = CLImageFormatImpl.create(); + f.setImageChannelOrder(JoclImage.decodeImageChannelOrder(format.channelOrder)); + f.setImageChannelDataType(JoclImage.decodeImageChannelType(format.channelType)); + //create image + long mem; + if (descr.type == Image.ImageType.IMAGE_2D) { + mem = cl.clCreateImage2D(id, memFlags, f, descr.width, descr.height, + descr.hostPtr==null ? 0 : descr.rowPitch, descr.hostPtr, Utils.errorBuffer); + Utils.checkError(Utils.errorBuffer, "clCreateImage2D"); + } else { + mem = cl.clCreateImage3D(id, memFlags, f, descr.width, descr.height, descr.depth, + descr.hostPtr==null ? 0 : descr.rowPitch, descr.hostPtr==null ? 0 : descr.slicePitch, + descr.hostPtr, Utils.errorBuffer); + Utils.checkError(Utils.errorBuffer, "clCreateImage3D"); + } + return new JoclImage(mem); + } + + @Override + public ImageFormat[] querySupportedFormats(MemoryAccess access, Image.ImageType type) { + if (type != Image.ImageType.IMAGE_2D && type != Image.ImageType.IMAGE_3D) { + throw new UnsupportedOperationException("Jocl only supports 2D and 3D images"); + } + long memFlags = Utils.getMemoryAccessFlags(access); + CLImageFormat[] fx; + if (type == Image.ImageType.IMAGE_2D) { + fx = context.getSupportedImage2dFormats(Mem.valueOf((int) memFlags)); + } else { + fx = context.getSupportedImage3dFormats(Mem.valueOf((int) memFlags)); + } + //convert formats + ImageFormat[] formats = new ImageFormat[fx.length]; + for (int i=0; i devices; + private ReleaserImpl(long id, List devices) { + this.id = id; + this.devices = devices; + } + @Override + public void release() { + if (id != 0) { + int ret = CLPlatform.getLowLevelCLInterface().clReleaseContext(id); + id = 0; + devices.clear(); + Utils.reportError(ret, "clReleaseContext"); + } + } + + } +} diff --git a/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclDevice.java b/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclDevice.java new file mode 100644 index 000000000..d4d9b08e1 --- /dev/null +++ b/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclDevice.java @@ -0,0 +1,302 @@ +/* + * 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.jocl; + +import com.jme3.opencl.Device; +import com.jme3.opencl.Platform; +import com.jogamp.opencl.CLDevice; +import java.util.Arrays; +import java.util.Collection; + +/** + * + * @author shaman + */ +public final class JoclDevice implements Device { + + final long id; + final CLDevice device; + final JoclPlatform platform; + + public JoclDevice(CLDevice device, JoclPlatform platform) { + this.id = device.ID; + this.device = device; + this.platform = platform; + } + + public long getId() { + return id; + } + + public CLDevice getDevice() { + return device; + } + + @Override + public JoclPlatform getPlatform() { + return platform; + } + + @Override + public DeviceType getDeviceType() { + CLDevice.Type type = device.getType(); + switch (type) { + case ACCELERATOR: return DeviceType.ACCELEARTOR; + case CPU: return DeviceType.CPU; + case GPU: return DeviceType.GPU; + default: return DeviceType.DEFAULT; + } + } + + @Override + public int getVendorId() { + return (int) device.getVendorID(); + } + + @Override + public boolean isAvailable() { + return device.isAvailable(); + } + + @Override + public boolean hasCompiler() { + return device.isCompilerAvailable(); + } + + @Override + public boolean hasDouble() { + return hasExtension("cl_khr_fp64"); + } + + @Override + public boolean hasHalfFloat() { + return hasExtension("cl_khr_fp16"); + } + + @Override + public boolean hasErrorCorrectingMemory() { + return device.isErrorCorrectionSupported(); + } + + @Override + public boolean hasUnifiedMemory() { + return device.isMemoryUnified(); + } + + @Override + public boolean hasImageSupport() { + return device.isImageSupportAvailable(); + } + + @Override + public boolean hasWritableImage3D() { + return hasExtension("cl_khr_3d_image_writes"); + } + + @Override + public boolean hasOpenGLInterop() { + return hasExtension("cl_khr_gl_sharing"); + } + + @Override + public boolean hasExtension(String extension) { + return getExtensions().contains(extension); + } + + @Override + public Collection getExtensions() { + return device.getExtensions(); + } + + @Override + public int getComputeUnits() { + return device.getMaxComputeUnits(); + } + + @Override + public int getClockFrequency() { + return device.getMaxClockFrequency(); + } + + @Override + public int getAddressBits() { + return device.getAddressBits(); + } + + @Override + public boolean isLittleEndian() { + return device.isLittleEndian(); + } + + @Override + public long getMaximumWorkItemDimensions() { + return device.getMaxWorkItemDimensions(); + } + + @Override + public long[] getMaximumWorkItemSizes() { + int[] sizes = device.getMaxWorkItemSizes(); + long[] s = new long[sizes.length]; + for (int i=0; i 0) { + p2 = Utils.pointers[2].rewind(); + p2.put(workGroupSize.getSizes(), 0, workGroupSize.getSizes().length); + p2.position(0); + } + long q = ((JoclCommandQueue) queue).id; + int ret = cl.clEnqueueNDRangeKernel(q, kernel, + globalWorkSize.getDimension(), null, Utils.pointers[1], + p2, 0, null, Utils.pointers[0]); + Utils.checkError(ret, "clEnqueueNDRangeKernel"); + return new JoclEvent(Utils.pointers[0].get(0)); + } + + @Override + public ObjectReleaser getReleaser() { + return new ReleaserImpl(kernel); + } + private static class ReleaserImpl implements ObjectReleaser { + private long kernel; + private ReleaserImpl(long kernel) { + this.kernel = kernel; + } + @Override + public void release() { + if (kernel != 0) { + int ret = CLPlatform.getLowLevelCLInterface().clReleaseKernel(kernel); + kernel = 0; + Utils.reportError(ret, "clReleaseKernel"); + } + } + } +} diff --git a/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclPlatform.java b/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclPlatform.java new file mode 100644 index 000000000..56d551180 --- /dev/null +++ b/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclPlatform.java @@ -0,0 +1,127 @@ +/* + * 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.jocl; + +import com.jme3.opencl.Device; +import com.jme3.opencl.Platform; +import com.jogamp.opencl.CLDevice; +import com.jogamp.opencl.CLPlatform; +import com.jogamp.opencl.llb.CL; +import java.util.ArrayList; +import java.util.Arrays; +import java.util.Collection; +import java.util.List; + +/** + * + * @author shaman + */ +public final class JoclPlatform implements Platform { + + final CLPlatform platform; + List devices; + + public JoclPlatform(CLPlatform platform) { + this.platform = platform; + } + + public CLPlatform getPlatform() { + return platform; + } + + @Override + public List getDevices() { + if (devices == null) { + devices = new ArrayList<>(); + for (CLDevice d : platform.listCLDevices()) { + devices.add(new JoclDevice(d, this)); + } + } + return devices; + } + + @Override + public String getProfile() { + return platform.getProfile(); + } + + @Override + public boolean isFullProfile() { + return getProfile().contains("FULL_PROFILE"); + } + + @Override + public boolean isEmbeddedProfile() { + return getProfile().contains("EMBEDDED_PROFILE"); + } + + @Override + public String getVersion() { + return platform.getVendor(); + } + + @Override + public int getVersionMajor() { + return Utils.getMajorVersion(getVersion(), "OpenCL "); + } + + @Override + public int getVersionMinor() { + return Utils.getMinorVersion(getVersion(), "OpenCL "); + } + + @Override + public String getName() { + return platform.getName(); + } + + @Override + public String getVendor() { + return platform.getVendor(); + } + + @Override + public boolean hasExtension(String extension) { + return getExtensions().contains(extension); + } + + @Override + public boolean hasOpenGLInterop() { + return hasExtension("cl_khr_gl_sharing"); + } + + @Override + public Collection getExtensions() { + return platform.getExtensions(); + } + +} diff --git a/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclProgram.java b/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclProgram.java new file mode 100644 index 000000000..dcc32014a --- /dev/null +++ b/jme3-jogl/src/main/java/com/jme3/opencl/jocl/JoclProgram.java @@ -0,0 +1,152 @@ +/* + * 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.jocl; + +import com.jme3.opencl.Kernel; +import com.jme3.opencl.KernelCompilationException; +import com.jme3.opencl.OpenCLObjectManager; +import com.jme3.opencl.Program; +import com.jogamp.common.nio.PointerBuffer; +import com.jogamp.opencl.CLPlatform; +import com.jogamp.opencl.llb.CL; +import com.jogamp.opencl.util.CLUtil; +import java.nio.ByteBuffer; +import java.util.logging.Level; +import java.util.logging.Logger; + +import static com.jogamp.common.nio.Buffers.newDirectByteBuffer; +import static com.jogamp.opencl.CLException.newException; +import static com.jogamp.opencl.llb.CL.CL_SUCCESS; +/** + * + * @author shaman + */ +public class JoclProgram implements Program { + private static final Logger LOG = Logger.getLogger(JoclProgram.class.getName()); + + final long program; + final CL cl; + private final JoclContext context; + + public JoclProgram(long program, JoclContext context) { + this.program = program; + this.context = context; + this.cl = CLPlatform.getLowLevelCLInterface(); + OpenCLObjectManager.getInstance().registerObject(this); + } + + @Override + public void build(String args) throws KernelCompilationException { + int ret = cl.clBuildProgram(program, 0, null, args, null); + if (ret != CL.CL_SUCCESS) { + String log = Log(); + LOG.log(Level.WARNING, "Unable to compile program:\n{0}", log); + if (ret == CL.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(long device) { + Utils.pointers[0].rewind(); + int ret = cl.clGetProgramBuildInfo(program, device, CL.CL_PROGRAM_BUILD_LOG, 0, null, Utils.pointers[0]); + Utils.checkError(ret, "clGetProgramBuildInfo"); + int count = (int) Utils.pointers[0].get(0); + final ByteBuffer buffer = newDirectByteBuffer(count); + ret = cl.clGetProgramBuildInfo(program, device, CL.CL_PROGRAM_BUILD_LOG, buffer.capacity(), buffer, null); + Utils.checkError(ret, "clGetProgramBuildInfo"); + return CLUtil.clString2JavaString(buffer, count); + } + + private String Log() { + StringBuilder str = new StringBuilder(); + for (JoclDevice device : context.getDevices()) { + long d = device.id; + str.append(device.getName()).append(":\n"); + str.append(Log(d)); + str.append('\n'); + } + return str.toString(); + } + + @Override + public Kernel createKernel(String name) { + Utils.errorBuffer.rewind(); + long kernel = cl.clCreateKernel(program, name, Utils.errorBuffer); + Utils.checkError(Utils.errorBuffer, "clCreateKernel"); + return new JoclKernel(kernel); + } + + @Override + public Kernel[] createAllKernels() { + Utils.tempBuffers[0].b16i.rewind(); + int ret = cl.clCreateKernelsInProgram(program, 0, null, Utils.tempBuffers[0].b16i); + Utils.checkError(ret, "clCreateKernelsInProgram"); + int count = Utils.tempBuffers[0].b16i.get(0); + PointerBuffer buf = PointerBuffer.allocateDirect(count); + ret = cl.clCreateKernelsInProgram(program, count, buf, null); + Utils.checkError(ret, "clCreateKernelsInProgram"); + Kernel[] kx = new Kernel[count]; + for (int i=0; i 0) { + return MemoryAccess.READ_WRITE; + } + if ((flag & CL.CL_MEM_READ_ONLY) > 0) { + return MemoryAccess.READ_ONLY; + } + if ((flag & CL.CL_MEM_WRITE_ONLY) > 0) { + return MemoryAccess.WRITE_ONLY; + } + throw new OpenCLException("Unknown memory access flag: "+flag); + } + + public static long getMappingAccessFlags(MappingAccess ma) { + switch (ma) { + case MAP_READ_ONLY: return CL.CL_MAP_READ; + case MAP_READ_WRITE: return CL.CL_MAP_READ | CL.CL_MAP_WRITE; + case MAP_WRITE_ONLY: return CL.CL_MAP_WRITE; + case MAP_WRITE_INVALIDATE: return CL.CL_MAP_WRITE; //MAP_WRITE_INVALIDATE_REGION not supported + default: throw new IllegalArgumentException("Unknown mapping access: "+ma); + } + } + +} diff --git a/jme3-jogl/src/main/java/com/jme3/system/jogl/JoglContext.java b/jme3-jogl/src/main/java/com/jme3/system/jogl/JoglContext.java index 602298bb9..12c1d10cb 100644 --- a/jme3-jogl/src/main/java/com/jme3/system/jogl/JoglContext.java +++ b/jme3-jogl/src/main/java/com/jme3/system/jogl/JoglContext.java @@ -36,6 +36,11 @@ import com.jme3.input.JoyInput; import com.jme3.input.KeyInput; import com.jme3.input.MouseInput; import com.jme3.opencl.Context; +import com.jme3.opencl.DefaultPlatformChooser; +import com.jme3.opencl.Device; +import com.jme3.opencl.PlatformChooser; +import com.jme3.opencl.jocl.JoclDevice; +import com.jme3.opencl.jocl.JoclPlatform; import com.jme3.renderer.Renderer; import com.jme3.renderer.RendererException; import com.jme3.renderer.jogl.JoglGL; @@ -56,6 +61,10 @@ import com.jme3.system.JmeContext; import com.jme3.system.NanoTimer; import com.jme3.system.SystemListener; import com.jme3.system.Timer; +import com.jogamp.opencl.CLDevice; +import com.jogamp.opencl.CLPlatform; +import com.jogamp.opencl.gl.CLGLContext; +import com.jogamp.opencl.llb.CL; import java.nio.IntBuffer; import java.util.concurrent.atomic.AtomicBoolean; @@ -65,6 +74,8 @@ import java.util.logging.Logger; import com.jogamp.opengl.GL; import com.jogamp.opengl.GL2GL3; import com.jogamp.opengl.GLContext; +import java.util.ArrayList; +import java.util.List; public abstract class JoglContext implements JmeContext { @@ -218,8 +229,95 @@ public abstract class JoglContext implements JmeContext { } } + @SuppressWarnings("unchecked") protected void initOpenCL() { logger.info("Initialize OpenCL with JOGL"); + + //load platforms and devices + StringBuilder platformInfos = new StringBuilder(); + ArrayList platforms = new ArrayList(); + for (CLPlatform p : CLPlatform.listCLPlatforms()) { + platforms.add(new JoclPlatform(p)); + } + platformInfos.append("Available OpenCL platforms:"); + for (int i=0; i devices = platform.getDevices(); + platformInfos.append("\n * Available devices:"); + for (int j=0; j choosenDevices = chooser.chooseDevices(platforms); + List devices = new ArrayList<>(choosenDevices.size()); + JoclPlatform platform = null; + for (Device d : choosenDevices) { + if (!(d instanceof JoclDevice)) { + logger.log(Level.SEVERE, "attempt to return a custom Device implementation from PlatformChooser: {0}", d); + return; + } + JoclDevice ld = (JoclDevice) d; + if (platform == null) { + platform = ld.getPlatform(); + } else if (platform != ld.getPlatform()) { + logger.severe("attempt to use devices from different platforms"); + return; + } + devices.add(ld.getDevice()); + } + if (devices.isEmpty()) { + logger.warning("no devices specified, no OpenCL context created"); + return; + } + logger.log(Level.INFO, "chosen platform: {0}", platform.getName()); + logger.log(Level.INFO, "chosen devices: {0}", choosenDevices); + + //create context + try { + CLGLContext c = CLGLContext.create(GLContext.getCurrent(), devices.toArray(new CLDevice[devices.size()])); + clContext = new com.jme3.opencl.jocl.JoclContext(c, (List) choosenDevices); + } catch (Exception ex) { + logger.log(Level.SEVERE, "Unable to create OpenCL context", ex); + return; + } + + logger.info("OpenCL context created"); } public void internalCreate() { 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 904b3ca0b..1ed59641b 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 @@ -419,9 +419,7 @@ public class LwjglImage implements Image { Utils.pointerBuffers[1], Utils.pointerBuffers[2], Utils.pointerBuffers[3], Utils.pointerBuffers[4], null, null, Utils.errorBuffer); Utils.checkError(Utils.errorBuffer, "clEnqueueMapBuffer"); - long event = Utils.pointerBuffers[0].get(0); - return new ImageMapping(buf, Utils.pointerBuffers[3].get(0), Utils.pointerBuffers[4].get(0), - new LwjglEvent(q.getCLEvent(event))); + return new ImageMapping(buf, Utils.pointerBuffers[3].get(0), Utils.pointerBuffers[4].get(0)); } @Override @@ -441,9 +439,10 @@ public class LwjglImage implements Image { Utils.errorBuffer.rewind(); ByteBuffer buf = CL10.clEnqueueMapImage(q, image, CL10.CL_FALSE, flags, Utils.pointerBuffers[1], Utils.pointerBuffers[2], - Utils.pointerBuffers[3], Utils.pointerBuffers[4], null, null, Utils.errorBuffer); + Utils.pointerBuffers[3], Utils.pointerBuffers[4], null, Utils.pointerBuffers[0], Utils.errorBuffer); Utils.checkError(Utils.errorBuffer, "clEnqueueMapBuffer"); - return new ImageMapping(buf, Utils.pointerBuffers[3].get(0), Utils.pointerBuffers[4].get(0)); + long event = Utils.pointerBuffers[0].get(0); + return new ImageMapping(buf, Utils.pointerBuffers[3].get(0), Utils.pointerBuffers[4].get(0), new LwjglEvent(q.getCLEvent(event))); } @Override