diff --git a/build/built-jar.properties b/build/built-jar.properties index e26f71e..01945bc 100644 --- a/build/built-jar.properties +++ b/build/built-jar.properties @@ -1,4 +1,4 @@ -#Tue, 11 Jul 2017 15:33:57 -0400 +#Tue, 25 Jul 2017 12:58:27 -0400 -C\:\\Users\\labramson\\Downloads\\Tutorial\\Tutorial= +C\:\\Users\\labramson\\Documents\\Tutorial= diff --git a/build/classes/tutorial/clTexture/DemoFractal$1.class b/build/classes/tutorial/clTexture/DemoFractal$1.class new file mode 100644 index 0000000..3b8cb0b Binary files /dev/null and b/build/classes/tutorial/clTexture/DemoFractal$1.class differ diff --git a/build/classes/tutorial/clTexture/DemoFractal$2.class b/build/classes/tutorial/clTexture/DemoFractal$2.class new file mode 100644 index 0000000..3042adf Binary files /dev/null and b/build/classes/tutorial/clTexture/DemoFractal$2.class differ diff --git a/build/classes/tutorial/clTexture/DemoFractal.class b/build/classes/tutorial/clTexture/DemoFractal.class new file mode 100644 index 0000000..2bda4ca Binary files /dev/null and b/build/classes/tutorial/clTexture/DemoFractal.class differ diff --git a/build/classes/tutorial/clTexture/Image.class b/build/classes/tutorial/clTexture/Image.class index 08fed33..ab1b44a 100644 Binary files a/build/classes/tutorial/clTexture/Image.class and b/build/classes/tutorial/clTexture/Image.class differ diff --git a/build/classes/tutorial/clTexture/Main.class b/build/classes/tutorial/clTexture/Main.class new file mode 100644 index 0000000..eea93f8 Binary files /dev/null and b/build/classes/tutorial/clTexture/Main.class differ diff --git a/build/classes/tutorial/clTexture/SharingMemCode.txt b/build/classes/tutorial/clTexture/SharingMemCode.txt new file mode 100644 index 0000000..277fdf1 --- /dev/null +++ b/build/classes/tutorial/clTexture/SharingMemCode.txt @@ -0,0 +1,42 @@ +// public CLMem createSharedMem(CLContext context, Texture texture) { +// // Create an OpenGL buffer +// int glBufId = GL15.glGenBuffers(); +// // Load the buffer with data using glBufferData(); +// // initialize buffer object +// +// GL15.glBindBuffer(GL_ARRAY_BUFFER, glBufId); +// +// //int size = (texture.getImageHeight() * texture.getImageWidth() * 4); +// GL15.glBufferData(GL_ARRAY_BUFFER, textureBuff, GL_DYNAMIC_DRAW); +// // Create the shared OpenCL memory object from the OpenGL buffer +// CLMem glMem = CL10GL.clCreateFromGLBuffer(context, CL10.CL_MEM_READ_WRITE, glBufId, null); +// +// return glMem; +// } +// +// public void useSharedMem(CLMem clmem, CLCommandQueue queue) { +// System.out.println("Acquiring mem lock"); +// // Acquire the lock for the 'glMem' memory object +// int error = CL10GL.clEnqueueAcquireGLObjects(queue, clmem, null, null); +// // Remember to check for errors +// if (error != CL10.CL_SUCCESS) { +// Util.checkCLError(error); +// } +// +// // Now execute an OpenCL command using the shared memory object, +// // such as uploading data to the memory object using 'CL10.clEnqueueWriteBuffer()' +// // or running a kernel using 'CL10.clEnqueueNDRangeKernel()' with the correct parameters +// // ... +// System.out.println("Doing Stuff"); +// // Release the lock on the 'glMem' memory object +// error = CL10GL.clEnqueueReleaseGLObjects(queue, clmem, null, null); +// if (error != CL10.CL_SUCCESS) { +// Util.checkCLError(error); +// } +// +// // Remember to flush the command queue when you are done. +// // Flushing the queue ensures that all of the OpenCL commands +// // sent to the queue have completed before the program continues. +// CL10.clFinish(queue); +// //System.out.println("Finished using CL & released shared mem"); +// } \ No newline at end of file diff --git a/build/classes/tutorial/clTexture/Texture.class b/build/classes/tutorial/clTexture/Texture.class index 6c5b592..6962b17 100644 Binary files a/build/classes/tutorial/clTexture/Texture.class and b/build/classes/tutorial/clTexture/Texture.class differ diff --git a/build/classes/tutorial/clTexture/TextureDriver.class b/build/classes/tutorial/clTexture/TextureDriver.class deleted file mode 100644 index a12f3d0..0000000 Binary files a/build/classes/tutorial/clTexture/TextureDriver.class and /dev/null differ diff --git a/build/classes/tutorial/clTexture/TextureMaker.class b/build/classes/tutorial/clTexture/TextureMaker.class deleted file mode 100644 index 909640e..0000000 Binary files a/build/classes/tutorial/clTexture/TextureMaker.class and /dev/null differ diff --git a/build/classes/tutorial/clTexture3d/Image.class b/build/classes/tutorial/clTexture3d/Image.class new file mode 100644 index 0000000..2d7f2b1 Binary files /dev/null and b/build/classes/tutorial/clTexture3d/Image.class differ diff --git a/build/classes/tutorial/clTexture3d/Main.class b/build/classes/tutorial/clTexture3d/Main.class new file mode 100644 index 0000000..e7b06f0 Binary files /dev/null and b/build/classes/tutorial/clTexture3d/Main.class differ diff --git a/build/classes/tutorial/clTexture3d/Texture.class b/build/classes/tutorial/clTexture3d/Texture.class new file mode 100644 index 0000000..46fe7ce Binary files /dev/null and b/build/classes/tutorial/clTexture3d/Texture.class differ diff --git a/nbproject/private/private.xml b/nbproject/private/private.xml index cdfc149..4b22223 100644 --- a/nbproject/private/private.xml +++ b/nbproject/private/private.xml @@ -3,10 +3,9 @@ - file:/C:/Users/labramson/Documents/Tutorial/src/tutorial/clTexture/Image.java - file:/C:/Users/labramson/Documents/Tutorial/src/tutorial/clTexture/Texture.java - file:/C:/Users/labramson/Documents/Tutorial/src/tutorial/clTexture/TextureMaker.java - file:/C:/Users/labramson/Documents/Tutorial/src/tutorial/clTexture/TextureDriver.java + file:/C:/Users/labramson/Documents/Tutorial/src/tutorial/clTexture3d/Texture.java + file:/C:/Users/labramson/Documents/Tutorial/src/tutorial/clTexture3d/Main.java + file:/C:/Users/labramson/Documents/Tutorial/src/tutorial/clTexture3d/Image.java diff --git a/nbproject/project.properties b/nbproject/project.properties index 22a892d..524bfb8 100644 --- a/nbproject/project.properties +++ b/nbproject/project.properties @@ -86,7 +86,7 @@ javadoc.splitindex=true javadoc.use=true javadoc.version=false javadoc.windowtitle= -main.class=tutorial.texture.TextureDriver +main.class=tutorial.clTexture.Main manifest.file=manifest.mf meta.inf.dir=${src.dir}/META-INF mkdist.disabled=false diff --git a/res/smileTexture2.jpg b/res/blank.jpg similarity index 55% rename from res/smileTexture2.jpg rename to res/blank.jpg index f4ce8ce..24e5c91 100644 Binary files a/res/smileTexture2.jpg and b/res/blank.jpg differ diff --git a/res/smileTexture.jpg b/res/smileTexture.jpg index 30e00d6..f4ce8ce 100644 Binary files a/res/smileTexture.jpg and b/res/smileTexture.jpg differ diff --git a/res/stripe.jpg b/res/stripe.jpg new file mode 100644 index 0000000..9813939 Binary files /dev/null and b/res/stripe.jpg differ diff --git a/src/tutorial/clTexture/DemoFractal.java b/src/tutorial/clTexture/DemoFractal.java new file mode 100644 index 0000000..c0590ca --- /dev/null +++ b/src/tutorial/clTexture/DemoFractal.java @@ -0,0 +1,917 @@ +/* + * Copyright (c) 2002-2010 LWJGL Project + * 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 'LWJGL' 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 tutorial.clTexture; + +import org.lwjgl.BufferUtils; +import org.lwjgl.LWJGLException; +import org.lwjgl.PointerBuffer; +import org.lwjgl.input.Keyboard; +import org.lwjgl.input.Mouse; +import org.lwjgl.opencl.*; +import org.lwjgl.opencl.api.Filter; +import org.lwjgl.opengl.*; +import org.lwjgl.util.Color; +import org.lwjgl.util.ReadableColor; + +import java.io.*; +import java.net.URL; +import java.nio.ByteBuffer; +import java.nio.IntBuffer; +import java.util.HashSet; +import java.util.List; +import java.util.Set; + +import static java.lang.Math.*; +import static org.lwjgl.opencl.CL10.*; +import static org.lwjgl.opencl.CL10GL.*; +import static org.lwjgl.opencl.KHRGLEvent.*; +import static org.lwjgl.opengl.AMDDebugOutput.*; +import static org.lwjgl.opengl.ARBCLEvent.*; +import static org.lwjgl.opengl.ARBDebugOutput.*; +import static org.lwjgl.opengl.ARBSync.*; +import static org.lwjgl.opengl.GL11.*; +import static org.lwjgl.opengl.GL15.*; +import static org.lwjgl.opengl.GL20.*; +import static org.lwjgl.opengl.GL21.*; + +/* + THIS DEMO USES CODE PORTED FROM JogAmp.org + Original code: http://github.com/mbien/jocl-demos + Original author: Michael Bien + ___ ___ ___ + / /\ / /\ ___ / /\ http://jocl.jogamp.org/ + / /:/ / /::\ /__/\ / /::\ a http://jogamp.org/ project. + /__/::\ / /:/\:\ \ \:\ / /:/\:\ + \__\/\:\ / /:/~/::\ \ \:\ / /:/~/::\ + \ \:\ /__/:/ /:/\:\ ___ \__\:\/__/:/ /:/\:\ + \__\:\\ \:\/:/__\//__/\ | |:|\ \:\/:/__\/ + / /:/ \ \::/ \ \:\| |:| \ \::/ + /__/:/ \ \:\ \ \:\__|:| \ \:\ + \__\/ \ \:\ \__\::::/ \ \:\ + \__\/ ~~~~ \__\/ + ___ ___ ___ ___ ___ + / /\ / /\ / /\ /__/\ / /\ + / /::\ / /::\ / /:/_ \ \:\ / /:/ + / /:/\:\ / /:/\:\ / /:/ /\ \ \:\ / /:/ ___ ___ + / /:/ \:\ / /:/~/:// /:/ /:/_ _____\__\:\ / /:/ ___ /__/\ / /\ + /__/:/ \__\:\/__/:/ /://__/:/ /:/ /\/__/::::::::\/__/:/ / /\\ \:\ / /:/ + \ \:\ / /:/\ \:\/:/ \ \:\/:/ /:/\ \:\~~\~~\/\ \:\ / /:/ \ \:\ /:/ + \ \:\ /:/ \ \::/ \ \::/ /:/ \ \:\ ~~~ \ \:\ /:/ \ \:\/:/ + \ \:\/:/ \ \:\ \ \:\/:/ \ \:\ \ \:\/:/ \ \::/ + \ \::/ \ \:\ \ \::/ \ \:\ \ \::/ \__\/ + \__\/ \__\/ \__\/ \__\/ \__\/ + _____ ___ ___ ___ ___ + / /::\ / /\ /__/\ / /\ / /\ + / /:/\:\ / /:/_ | |::\ / /::\ / /:/_ + / /:/ \:\ / /:/ /\ | |:|:\ / /:/\:\ / /:/ /\ + /__/:/ \__\:| / /:/ /:/_ __|__|:|\:\ / /:/ \:\ / /:/ /::\ + \ \:\ / /:/ /__/:/ /:/ /\ /__/::::| \:\ /__/:/ \__\:\ /__/:/ /:/\:\ + \ \:\ /:/ \ \:\/:/ /:/ \ \:\~~\__\/ \ \:\ / /:/ \ \:\/:/~/:/ + \ \:\/:/ \ \::/ /:/ \ \:\ \ \:\ /:/ \ \::/ /:/ + \ \::/ \ \:\/:/ \ \:\ \ \:\/:/ \__\/ /:/ + \__\/ \ \::/ \ \:\ \ \::/ /__/:/ + \__\/ \__\/ \__\/ \__\/ + */ +/** + * Computes the Mandelbrot set with OpenCL using multiple GPUs and renders the + * result with OpenGL. A shared PBO is used as storage for the fractal + * image.
+ * http://en.wikipedia.org/wiki/Mandelbrot_set + *

+ * controls:
+ * keys 1-9 control parallelism level
+ * space enables/disables slice seperator
+ * 'd' toggles between 32/64bit floatingpoint precision
+ * mouse/mousewheel to drag and zoom
+ * 'Home' to reset the viewport
+ *

+ * + * @author Michael Bien, Spasi + */ +public class DemoFractal { + + // max number of used GPUs + private static final int MAX_PARALLELISM_LEVEL = 8; + + private static final int COLOR_MAP_SIZE = 32 * 2 * 4; + + private Set params; + + private CLContext clContext; + private CLCommandQueue[] queues; + private CLKernel[] kernels; + private CLProgram[] programs; + + private CLMem[] glBuffers; + private IntBuffer glIDs; + + private boolean useTextures; + + // Texture rendering + private int dlist; + private int vsh; + private int fsh; + private int program; + + private CLMem[] colorMap; + + private final PointerBuffer kernel2DGlobalWorkSize; + + // max per pixel iterations to compute the fractal + private int maxIterations = 500; + + private int width = 512; + private int height = 512; + + private double minX = -2f; + private double minY = -1.2f; + private double maxX = 0.6f; + private double maxY = 1.3f; + + private boolean dragging; + private double dragX; + private double dragY; + private double dragMinX; + private double dragMinY; + private double dragMaxX; + private double dragMaxY; + + private int mouseX; + private int mouseY; + + private int slices; + + private boolean drawSeparator; + private boolean doublePrecision = true; + private boolean buffersInitialized; + private boolean rebuild; + + private boolean run = true; + + // EVENT SYNCING + private final PointerBuffer syncBuffer = BufferUtils.createPointerBuffer(1); + + private boolean syncGLtoCL; // true if we can make GL wait on events generated from CL queues. + private CLEvent[] clEvents; + private GLSync[] clSyncs; + + private boolean syncCLtoGL; // true if we can make CL wait on sync objects generated from GL. + private GLSync glSync; + private CLEvent glEvent; + + public DemoFractal(final String[] args) { + params = new HashSet(); + + for (int i = 0; i < args.length; i++) { + final String arg = args[i]; + + if (arg.charAt(0) != '-' && arg.charAt(0) != '/') { + throw new IllegalArgumentException("Invalid command-line argument: " + args[i]); + } + + final String param = arg.substring(1); + + if ("forcePBO".equalsIgnoreCase(param)) { + params.add("forcePBO"); + } else if ("forceCPU".equalsIgnoreCase(param)) { + params.add("forceCPU"); + } else if ("debugGL".equalsIgnoreCase(param)) { + params.add("debugGL"); + } else if ("iterations".equalsIgnoreCase(param)) { + if (args.length < i + 1 + 1) { + throw new IllegalArgumentException("Invalid iterations argument specified."); + } + + try { + this.maxIterations = Integer.parseInt(args[++i]); + } catch (NumberFormatException e) { + throw new IllegalArgumentException("Invalid number of iterations specified."); + } + } else if ("res".equalsIgnoreCase(param)) { + if (args.length < i + 2 + 1) { + throw new IllegalArgumentException("Invalid res argument specified."); + } + + try { + this.width = Integer.parseInt(args[++i]); + this.height = Integer.parseInt(args[++i]); + + if (width < 1 || height < 1) { + throw new IllegalArgumentException("Invalid res dimensions specified."); + } + } catch (NumberFormatException e) { + throw new IllegalArgumentException("Invalid res dimensions specified."); + } + } + } + + kernel2DGlobalWorkSize = BufferUtils.createPointerBuffer(2); + } + + public static void main(String args[]) { + DemoFractal demo = new DemoFractal(args); + demo.init(); + demo.run(); + } + + public void init() { + try { + CL.create(); + Display.setDisplayMode(new DisplayMode(width, height)); + Display.setTitle("OpenCL Fractal Demo"); + Display.setSwapInterval(0); + Display.create(new PixelFormat(), new ContextAttribs().withDebug(params.contains("debugGL"))); + } catch (LWJGLException e) { + throw new RuntimeException(e); + } + + try { + initCL(Display.getDrawable()); + } catch (Exception e) { + if (clContext != null) { + clReleaseContext(clContext); + } + Display.destroy(); + throw new RuntimeException(e); + } + + glDisable(GL_DEPTH_TEST); + glClearColor(0.0f, 0.0f, 0.0f, 1.0f); + + initView(Display.getDisplayMode().getWidth(), Display.getDisplayMode().getHeight()); + + initGLObjects(); + glFinish(); + + setKernelConstants(); + } + + private void initCL(Drawable drawable) throws Exception { + // Find a platform + List platforms = CLPlatform.getPlatforms(); + if (platforms == null) { + throw new RuntimeException("No OpenCL platforms found."); + } + + final CLPlatform platform = platforms.get(0); // just grab the first one + + // Find devices with GL sharing support + final Filter glSharingFilter = new Filter() { + public boolean accept(final CLDevice device) { + final CLDeviceCapabilities caps = CLCapabilities.getDeviceCapabilities(device); + return caps.CL_KHR_gl_sharing; + } + }; + int device_type = params.contains("forceCPU") ? CL_DEVICE_TYPE_CPU : CL_DEVICE_TYPE_GPU; + List devices = platform.getDevices(device_type, glSharingFilter); + if (devices == null) { + device_type = CL_DEVICE_TYPE_CPU; + devices = platform.getDevices(device_type, glSharingFilter); + if (devices == null) { + throw new RuntimeException("No OpenCL devices found with KHR_gl_sharing support."); + } + } + + // Create the context + clContext = CLContext.create(platform, devices, new CLContextCallback() { + protected void handleMessage(final String errinfo, final ByteBuffer private_info) { + System.out.println("[CONTEXT MESSAGE] " + errinfo); + } + }, drawable, null); + + slices = min(devices.size(), MAX_PARALLELISM_LEVEL); + + // create command queues for every GPU, setup colormap and init kernels + queues = new CLCommandQueue[slices]; + kernels = new CLKernel[slices]; + colorMap = new CLMem[slices]; + + for (int i = 0; i < slices; i++) { + colorMap[i] = clCreateBuffer(clContext, CL_MEM_READ_ONLY, COLOR_MAP_SIZE, null); + colorMap[i].checkValid(); + + // create command queue and upload color map buffer on each used device + queues[i] = clCreateCommandQueue(clContext, devices.get(i), CL_QUEUE_PROFILING_ENABLE, null); + queues[i].checkValid(); + + final ByteBuffer colorMapBuffer = clEnqueueMapBuffer(queues[i], colorMap[i], CL_TRUE, CL_MAP_WRITE, 0, COLOR_MAP_SIZE, null, null, null); + initColorMap(colorMapBuffer.asIntBuffer(), 32, Color.BLUE, Color.GREEN, Color.RED); + clEnqueueUnmapMemObject(queues[i], colorMap[i], colorMapBuffer, null, null); + } + + // check if we have 64bit FP support on all devices + // if yes we can use only one program for all devices + one kernel per device. + // if not we will have to create (at least) one program for 32 and one for 64bit devices. + // since there are different vendor extensions for double FP we use one program per device. + // (OpenCL spec is not very clear about this usecases) + boolean all64bit = true; + for (CLDevice device : devices) { + if (!isDoubleFPAvailable(device)) { + all64bit = false; + break; + } + } + + // load program(s) + programs = new CLProgram[all64bit ? 1 : slices]; + + final ContextCapabilities caps = GLContext.getCapabilities(); + + if (!caps.OpenGL20) { + throw new RuntimeException("OpenGL 2.0 is required to run this demo."); + } else if (device_type == CL_DEVICE_TYPE_CPU && !caps.OpenGL21) { + throw new RuntimeException("OpenGL 2.1 is required to run this demo."); + } + + if (params.contains("debugGL")) { + if (caps.GL_ARB_debug_output) { + glDebugMessageCallbackARB(new ARBDebugOutputCallback()); + } else if (caps.GL_AMD_debug_output) { + glDebugMessageCallbackAMD(new AMDDebugOutputCallback()); + } + } + + if (device_type == CL_DEVICE_TYPE_GPU) { + System.out.println("OpenCL Device Type: GPU (Use -forceCPU to use CPU)"); + } else { + System.out.println("OpenCL Device Type: CPU"); + } + for (int i = 0; i < devices.size(); i++) { + System.out.println("OpenCL Device #" + (i + 1) + " supports KHR_gl_event = " + CLCapabilities.getDeviceCapabilities(devices.get(i)).CL_KHR_gl_event); + } + + System.out.println("\nMax Iterations: " + maxIterations + " (Use -iterations to change)"); + System.out.println("Display resolution: " + width + "x" + height + " (Use -res to change)"); + + System.out.println("\nOpenGL caps.GL_ARB_sync = " + caps.GL_ARB_sync); + System.out.println("OpenGL caps.GL_ARB_cl_event = " + caps.GL_ARB_cl_event); + + // Use PBO if we're on a CPU implementation + useTextures = device_type == CL_DEVICE_TYPE_GPU && (!caps.OpenGL21 || !params.contains("forcePBO")); + if (useTextures) { + System.out.println("\nCL/GL Sharing method: TEXTURES (use -forcePBO to use PBO + DrawPixels)"); + System.out.println("Rendering method: Shader on a fullscreen quad"); + } else { + System.out.println("\nCL/GL Sharing method: PIXEL BUFFER OBJECTS"); + System.out.println("Rendering method: DrawPixels"); + } + + buildPrograms(); + + // Detect GLtoCL synchronization method + syncGLtoCL = caps.GL_ARB_cl_event; // GL3.2 or ARB_sync implied + if (syncGLtoCL) { + clEvents = new CLEvent[slices]; + clSyncs = new GLSync[slices]; + System.out.println("\nGL to CL sync: Using OpenCL events"); + } else { + System.out.println("\nGL to CL sync: Using clFinish"); + } + + // Detect CLtoGL synchronization method + syncCLtoGL = caps.OpenGL32 || caps.GL_ARB_sync; + if (syncCLtoGL) { + for (CLDevice device : devices) { + if (!CLCapabilities.getDeviceCapabilities(device).CL_KHR_gl_event) { + syncCLtoGL = false; + break; + } + } + } + if (syncCLtoGL) { + System.out.println("CL to GL sync: Using OpenGL sync objects"); + } else { + System.out.println("CL to GL sync: Using glFinish"); + } + + if (useTextures) { + dlist = glGenLists(1); + + glNewList(dlist, GL_COMPILE); + glBegin(GL_QUADS); + { + glTexCoord2f(0.0f, 0.0f); + glVertex2f(0, 0); + + glTexCoord2f(0.0f, 1.0f); + glVertex2i(0, height); + + glTexCoord2f(1.0f, 1.0f); + glVertex2f(width, height); + + glTexCoord2f(1.0f, 0.0f); + glVertex2f(width, 0); + } + glEnd(); + glEndList(); + + vsh = glCreateShader(GL_VERTEX_SHADER); + glShaderSource(vsh, "varying vec2 texCoord;\n" + + "\n" + + "void main(void) {\n" + + "\tgl_Position = ftransform();\n" + + "\ttexCoord = gl_MultiTexCoord0.xy;\n" + + "}"); + glCompileShader(vsh); + + fsh = glCreateShader(GL_FRAGMENT_SHADER); + glShaderSource(fsh, "uniform sampler2D mandelbrot;\n" + + "\n" + + "varying vec2 texCoord;\n" + + "\n" + + "void main(void) {\n" + + "\tgl_FragColor = texture2D(mandelbrot, texCoord);" + + "}"); + glCompileShader(fsh); + + System.out.println("HERE"); + program = glCreateProgram(); + glAttachShader(program, vsh); + glAttachShader(program, fsh); + glLinkProgram(program); + + glUseProgram(program); + glUniform1i(glGetUniformLocation(program, "mandelbrot"), 0); + } + + System.out.println(""); + } + + private void buildPrograms() { + /* + * workaround: The driver keeps using the old binaries for some reason. + * to solve this we simple create a new program and release the old. + * however rebuilding programs should be possible -> remove when drivers are fixed. + * (again: the spec is not very clear about this kind of usages) + */ + if (programs[0] != null) { + for (CLProgram program : programs) { + clReleaseProgram(program); + } + } + + try { + createPrograms(); + } catch (IOException e) { + throw new RuntimeException(e); + } + + // disable 64bit floating point math if not available + for (int i = 0; i < programs.length; i++) { + final CLDevice device = queues[i].getCLDevice(); + + final StringBuilder options = new StringBuilder(useTextures ? "-D USE_TEXTURE" : ""); + final CLDeviceCapabilities caps = CLCapabilities.getDeviceCapabilities(device); + if (doublePrecision && isDoubleFPAvailable(device)) { + //cl_khr_fp64 + options.append(" -D DOUBLE_FP"); + + //amd's verson of double precision floating point math + if (!caps.CL_KHR_fp64 && caps.CL_AMD_fp64) { + options.append(" -D AMD_FP"); + } + } + + System.out.println("\nOpenCL COMPILER OPTIONS: " + options); + + try { + clBuildProgram(programs[i], device, options, null); + } finally { + System.out.println("BUILD LOG: " + programs[i].getBuildInfoString(device, CL_PROGRAM_BUILD_LOG)); + } + } + + rebuild = false; + + // init kernel with constants + for (int i = 0; i < kernels.length; i++) { + kernels[i] = clCreateKernel(programs[min(i, programs.length)], "mandelbrot", null); + } + } + + private void initGLObjects() { + if (glBuffers == null) { + glBuffers = new CLMem[slices]; + glIDs = BufferUtils.createIntBuffer(slices); + } else { + for (CLMem mem : glBuffers) { + clReleaseMemObject(mem); + } + + if (useTextures) { + glDeleteTextures(glIDs); + } else { + glDeleteBuffers(glIDs); + } + } + + if (useTextures) { + glGenTextures(glIDs); + + // Init textures + for (int i = 0; i < slices; i++) { + glBindTexture(GL_TEXTURE_2D, glIDs.get(i)); + glTexImage2D(GL_TEXTURE_2D, 0, GL_RGBA, width / slices, height, 0, GL_RGBA, GL_UNSIGNED_BYTE, (ByteBuffer) null); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST); + + glBuffers[i] = clCreateFromGLTexture2D(clContext, CL_MEM_WRITE_ONLY, GL_TEXTURE_2D, 0, glIDs.get(i), null); + } + glBindTexture(GL_TEXTURE_2D, 0); + } else { + glGenBuffers(glIDs); + + // setup one empty PBO per slice + for (int i = 0; i < slices; i++) { + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, glIDs.get(i)); + glBufferData(GL_PIXEL_UNPACK_BUFFER, width * height * 4 / slices, GL_STREAM_DRAW); + + glBuffers[i] = clCreateFromGLBuffer(clContext, CL_MEM_WRITE_ONLY, glIDs.get(i), null); + } + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); + } + + buffersInitialized = true; + } + + // init kernels with constants + private void setKernelConstants() { + for (int i = 0; i < slices; i++) { + kernels[i] + .setArg(6, glBuffers[i]) + .setArg(7, colorMap[i]) + .setArg(8, COLOR_MAP_SIZE) + .setArg(9, maxIterations); + } + } + + // rendering cycle + private void run() { + long startTime = System.currentTimeMillis() + 5000; + long fps = 0; + + while (run) { + if (!Display.isVisible()) { + Thread.yield(); + } + + handleIO(); + display(); + + Display.update(); + if (Display.isCloseRequested()) { + break; + } + + if (startTime > System.currentTimeMillis()) { + fps++; + } else { + long timeUsed = 5000 + (startTime - System.currentTimeMillis()); + startTime = System.currentTimeMillis() + 5000; + System.out.println(fps + " frames in 5 seconds = " + (fps / (timeUsed / 1000f))); + fps = 0; + } + } + + clReleaseContext(clContext); + + if (useTextures) { + glDeleteProgram(program); + glDeleteShader(fsh); + glDeleteShader(vsh); + + glDeleteLists(dlist, 1); + } + + CL.destroy(); + Display.destroy(); + } + + public void display() { + // TODO: Need to clean-up events, test when ARB_cl_events & KHR_gl_event are implemented. + + // make sure GL does not use our objects before we start computing + if (syncCLtoGL && glEvent != null) { + for (final CLCommandQueue queue : queues) { + clEnqueueWaitForEvents(queue, glEvent); + } + } else { + glFinish(); + } + + if (!buffersInitialized) { + initGLObjects(); + setKernelConstants(); + } + + if (rebuild) { + buildPrograms(); + setKernelConstants(); + } + compute(doublePrecision); + + render(); + } + + // OpenCL + private void compute(final boolean is64bit) { + int sliceWidth = (int) (width / (float) slices); + double rangeX = (maxX - minX) / slices; + double rangeY = (maxY - minY); + + kernel2DGlobalWorkSize.put(0, sliceWidth).put(1, height); + + // start computation + for (int i = 0; i < slices; i++) { + kernels[i].setArg(0, sliceWidth).setArg(1, height); + if (!is64bit || !isDoubleFPAvailable(queues[i].getCLDevice())) { + kernels[i] + .setArg(2, (float) (minX + rangeX * i)).setArg(3, (float) minY) + .setArg(4, (float) rangeX).setArg(5, (float) rangeY); + } else { + kernels[i] + .setArg(2, minX + rangeX * i).setArg(3, minY) + .setArg(4, rangeX).setArg(5, rangeY); + } + + // acquire GL objects, and enqueue a kernel with a probe from the list + clEnqueueAcquireGLObjects(queues[i], glBuffers[i], null, null); + + clEnqueueNDRangeKernel(queues[i], kernels[i], 2, + null, + kernel2DGlobalWorkSize, + null, + null, null); + + clEnqueueReleaseGLObjects(queues[i], glBuffers[i], null, syncGLtoCL ? syncBuffer : null); + if (syncGLtoCL) { + clEvents[i] = queues[i].getCLEvent(syncBuffer.get(0)); + clSyncs[i] = glCreateSyncFromCLeventARB(queues[i].getParent(), clEvents[i], 0); + } + } + + // block until done (important: finish before doing further gl work) + if (!syncGLtoCL) { + for (int i = 0; i < slices; i++) { + clFinish(queues[i]); + } + } + } + + // OpenGL + private void render() { + glClear(GL_COLOR_BUFFER_BIT); + + if (syncGLtoCL) { + for (int i = 0; i < slices; i++) { + glWaitSync(clSyncs[i], 0, 0); + } + } + + //draw slices + int sliceWidth = width / slices; + + if (useTextures) { + for (int i = 0; i < slices; i++) { + int seperatorOffset = drawSeparator ? i : 0; + + glBindTexture(GL_TEXTURE_2D, glIDs.get(i)); + glCallList(dlist); + } + } else { + for (int i = 0; i < slices; i++) { + int seperatorOffset = drawSeparator ? i : 0; + + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, glIDs.get(i)); + glRasterPos2i(sliceWidth * i + seperatorOffset, 0); + + glDrawPixels(sliceWidth, height, GL_RGBA, GL_UNSIGNED_BYTE, 0); + } + glBindBuffer(GL_PIXEL_UNPACK_BUFFER, 0); + } + + if (syncCLtoGL) { + glSync = glFenceSync(GL_SYNC_GPU_COMMANDS_COMPLETE, 0); + glEvent = clCreateEventFromGLsyncKHR(clContext, glSync, null); + } + + //draw info text + /* + textRenderer.beginRendering(width, height, false); + textRenderer.draw("device/time/precision", 10, height - 15); + for ( int i = 0; i < slices; i++ ) { + CLDevice device = queues[i].getDevice(); + boolean doubleFP = doublePrecision && isDoubleFPAvailable(device); + CLEvent event = probes.getEvent(i); + long start = event.getProfilingInfo(START); + long end = event.getProfilingInfo(END); + textRenderer.draw(device.getType().toString() + i + " " + + (int)((end - start) / 1000000.0f) + "ms @" + + (doubleFP ? "64bit" : "32bit"), 10, height - (20 + 16 * (slices - i))); + } + textRenderer.endRendering(); + */ + } + + private void handleIO() { + if (Keyboard.getNumKeyboardEvents() != 0) { + while (Keyboard.next()) { + if (Keyboard.getEventKeyState()) { + continue; + } + + final int key = Keyboard.getEventKey(); + + if (Keyboard.KEY_1 <= key && key <= Keyboard.KEY_8) { + int number = key - Keyboard.KEY_1 + 1; + slices = min(number, min(queues.length, MAX_PARALLELISM_LEVEL)); + System.out.println("NEW PARALLELISM LEVEL: " + slices); + buffersInitialized = false; + } else { + switch (Keyboard.getEventKey()) { + case Keyboard.KEY_SPACE: + drawSeparator = !drawSeparator; + System.out.println("SEPARATOR DRAWING IS NOW: " + (drawSeparator ? "ON" : "OFF")); + break; + case Keyboard.KEY_D: + doublePrecision = !doublePrecision; + System.out.println("DOUBLE PRECISION IS NOW: " + (doublePrecision ? "ON" : "OFF")); + rebuild = true; + break; + case Keyboard.KEY_HOME: + minX = -2f; + minY = -1.2f; + maxX = 0.6f; + maxY = 1.3f; + break; + case Keyboard.KEY_ESCAPE: + run = false; + break; + } + } + } + } + + while (Mouse.next()) { + final int eventBtn = Mouse.getEventButton(); + + final int x = Mouse.getX(); + final int y = Mouse.getY(); + + if (Mouse.isButtonDown(0) && (x != mouseX || y != mouseY)) { + if (!dragging) { + dragging = true; + + dragX = mouseX; + dragY = mouseY; + + dragMinX = minX; + dragMinY = minY; + dragMaxX = maxX; + dragMaxY = maxY; + } + + double offsetX = (x - dragX) * (maxX - minX) / width; + double offsetY = (y - dragY) * (maxY - minY) / height; + + minX = dragMinX - offsetX; + minY = dragMinY - offsetY; + + maxX = dragMaxX - offsetX; + maxY = dragMaxY - offsetY; + } else { + if (dragging) { + dragging = false; + } + + if (eventBtn == -1) { + final int dwheel = Mouse.getEventDWheel(); + if (dwheel != 0) { + double scaleFactor = Keyboard.isKeyDown(Keyboard.KEY_LCONTROL) || Keyboard.isKeyDown(Keyboard.KEY_RCONTROL) ? 0.25 : 0.05; + double scale = dwheel > 0 ? scaleFactor : -scaleFactor; + + double deltaX = scale * (maxX - minX); + double deltaY = scale * (maxY - minY); + + // offset for "zoom to cursor" + double offsetX = (x / (double) width - 0.5) * deltaX * 2.0; + double offsetY = (y / (double) height - 0.5) * deltaY * 2.0; + + minX += deltaX + offsetX; + minY += deltaY - offsetY; + + maxX += -deltaX + offsetX; + maxY += -deltaY - offsetY; + } + } + } + + mouseX = x; + mouseY = y; + } + } + + private static boolean isDoubleFPAvailable(CLDevice device) { + final CLDeviceCapabilities caps = CLCapabilities.getDeviceCapabilities(device); + return caps.CL_KHR_fp64 || caps.CL_AMD_fp64; + } + + private void createPrograms() throws IOException { + final String source = getProgramSource("org/lwjgl/test/opencl/gl/Mandelbrot.cl"); + for (int i = 0; i < programs.length; i++) { + programs[i] = clCreateProgramWithSource(clContext, source, null); + } + } + + private String getProgramSource(final String file) throws IOException { + InputStream source = null; + URL sourceURL = Thread.currentThread().getContextClassLoader().getResource(file); + if (sourceURL != null) { + source = sourceURL.openStream(); + } + if (source == null) // dev-mode + { + source = new FileInputStream("src/java/" + file); + } + final BufferedReader reader = new BufferedReader(new InputStreamReader(source)); + + final StringBuilder sb = new StringBuilder(); + String line; + try { + while ((line = reader.readLine()) != null) { + sb.append(line).append("\n"); + } + } finally { + source.close(); + } + + return sb.toString(); + } + + private static void initColorMap(IntBuffer colorMap, int stepSize, ReadableColor... colors) { + for (int n = 0; n < colors.length - 1; n++) { + ReadableColor color = colors[n]; + int r0 = color.getRed(); + int g0 = color.getGreen(); + int b0 = color.getBlue(); + + color = colors[n + 1]; + int r1 = color.getRed(); + int g1 = color.getGreen(); + int b1 = color.getBlue(); + + int deltaR = r1 - r0; + int deltaG = g1 - g0; + int deltaB = b1 - b0; + + for (int step = 0; step < stepSize; step++) { + float alpha = (float) step / (stepSize - 1); + int r = (int) (r0 + alpha * deltaR); + int g = (int) (g0 + alpha * deltaG); + int b = (int) (b0 + alpha * deltaB); + colorMap.put((r << 0) | (g << 8) | (b << 16)); + } + } + } + + private static void initView(int width, int height) { + glViewport(0, 0, width, height); + + glMatrixMode(GL_MODELVIEW); + glLoadIdentity(); + + glMatrixMode(GL_PROJECTION); + glLoadIdentity(); + glOrtho(0.0, width, 0.0, height, 0.0, 1.0); + } + +} diff --git a/src/tutorial/clTexture/Image.java b/src/tutorial/clTexture/Image.java index 043cc9d..9c40066 100644 --- a/src/tutorial/clTexture/Image.java +++ b/src/tutorial/clTexture/Image.java @@ -1,39 +1,55 @@ -/* - * To change this license header, choose License Headers in Project Properties. - * To change this template file, choose Tools | Templates - * and open the template in the editor. - */ package tutorial.clTexture; +import java.awt.Color; +import java.awt.Graphics; +import java.awt.color.ColorSpace; import java.awt.image.BufferedImage; +import java.awt.image.ColorModel; +import java.awt.image.ComponentColorModel; +import java.awt.image.DataBuffer; +import java.awt.image.DataBufferByte; +import java.awt.image.Raster; +import java.awt.image.WritableRaster; import java.io.FileInputStream; import java.nio.ByteBuffer; +import java.nio.ByteOrder; +import java.util.Hashtable; import javax.imageio.ImageIO; -import org.lwjgl.BufferUtils; - /* - * @author labramson + * @author LAA */ public class Image { - private final int width; - private final int height; - private BufferedImage image; - private ByteBuffer buffer; + private final int width, height; + private final BufferedImage img; + + /* The color model including alpha for the GL image */ + private final ColorModel glAlphaColorModel; + /* The color model for the GL image */ + private final ColorModel glColorModel; public Image(String imgFile) { - System.out.println("Made image"); - BufferedImage img = this.getImageFile(imgFile); - ByteBuffer buff = this.createByteBuff(img); + this.img = getBuffImage(imgFile); this.width = img.getWidth(); this.height = img.getHeight(); - this.image = img; - this.buffer = buff; + this.glAlphaColorModel = new ComponentColorModel(ColorSpace.getInstance(ColorSpace.CS_sRGB), + new int[]{8, 8, 8, 8}, + true, + false, + ComponentColorModel.TRANSLUCENT, + DataBuffer.TYPE_BYTE); + + this.glColorModel = new ComponentColorModel(ColorSpace.getInstance(ColorSpace.CS_sRGB), + new int[]{8, 8, 8, 0}, + false, + false, + ComponentColorModel.OPAQUE, + DataBuffer.TYPE_BYTE); } - public BufferedImage getImageFile(String imageFile) { - System.out.println("read image file"); + private BufferedImage getBuffImage(String imageFile) { + System.out.println("Fetching image file"); try { BufferedImage img = ImageIO.read(new FileInputStream(imageFile)); return img; @@ -43,29 +59,37 @@ public BufferedImage getImageFile(String imageFile) { } } - private ByteBuffer createByteBuff(BufferedImage img) { - System.out.println("Made image buffer"); - try { - int pixels[] = new int[img.getWidth() * img.getHeight()]; - img.getRGB(0, 0, img.getWidth(), img.getHeight(), pixels, 0, img.getWidth()); - - ByteBuffer buffer = BufferUtils.createByteBuffer(img.getWidth() * img.getHeight() * 4); - - for (int x = 0; x < img.getWidth(); x++) { - for (int y = 0; y < img.getHeight(); y++) { - int pixel = pixels[y * img.getWidth() + x]; - buffer.put((byte) ((pixel >> 16) & 0xFF)); - buffer.put((byte) ((pixel >> 8) & 0xFF)); - buffer.put((byte) (pixel & 0xFF)); - } - } - - buffer.flip(); - return buffer; - } catch (Exception e) { - e.printStackTrace(); - return null; + public ByteBuffer getByeBuff() { + ByteBuffer byteBuffer; + WritableRaster raster; + BufferedImage texImage; + + System.out.println("Making image byte buffer"); + + // create a raster that can be used by OpenGL as a source for a texture + if (this.img.getColorModel().hasAlpha()) { + raster = Raster.createInterleavedRaster(DataBuffer.TYPE_BYTE, this.img.getWidth() + 1, this.img.getHeight(), 4, null); + texImage = new BufferedImage(glAlphaColorModel, raster, false, new Hashtable()); + } else { + raster = Raster.createInterleavedRaster(DataBuffer.TYPE_BYTE, this.img.getWidth() + 1, this.img.getHeight(), 3, null); + texImage = new BufferedImage(glColorModel, raster, false, new Hashtable()); } + + // copy the source image into the produced image + Graphics g = texImage.getGraphics(); + g.setColor(new Color(0f, 0f, 0f, 0f)); + g.fillRect(0, 0, this.img.getWidth() + 1, this.img.getHeight()); + g.drawImage(this.img, 0, 0, null); + + // build a byte buffer from the temporary image to produce a texture + byte[] data = ((DataBufferByte) texImage.getRaster().getDataBuffer()).getData(); + + byteBuffer = ByteBuffer.allocateDirect(data.length); + byteBuffer.order(ByteOrder.nativeOrder()); + byteBuffer.put(data, 0, data.length); + byteBuffer.flip(); + + return byteBuffer; } public int getWidth() { @@ -76,11 +100,7 @@ public int getHeight() { return height; } - public ByteBuffer getBuffer() { - return buffer; - } - public BufferedImage getImg() { - return image; + return img; } } diff --git a/src/tutorial/clTexture/Main.java b/src/tutorial/clTexture/Main.java new file mode 100644 index 0000000..265e934 --- /dev/null +++ b/src/tutorial/clTexture/Main.java @@ -0,0 +1,492 @@ +/* + * To change this license header, choose License Headers in Project Properties. + * To change this template file, choose Tools | Templates + * and open the template in the editor. + */ +package tutorial.clTexture; + +import java.nio.IntBuffer; +import java.util.List; +import org.lwjgl.*; +import org.lwjgl.input.Keyboard; +import org.lwjgl.opencl.CL10; +import org.lwjgl.opencl.CL10GL; +import org.lwjgl.opencl.CL; +import org.lwjgl.opencl.CLCapabilities; +import org.lwjgl.opencl.CLCommandQueue; +import org.lwjgl.opencl.CLContext; +import org.lwjgl.opencl.CLDevice; +import org.lwjgl.opencl.CLDeviceCapabilities; +import org.lwjgl.opencl.CLEvent; +import org.lwjgl.opencl.CLKernel; +import org.lwjgl.opencl.CLMem; +import org.lwjgl.opencl.CLPlatform; +import org.lwjgl.opencl.CLProgram; +import org.lwjgl.opencl.api.Filter; +import org.lwjgl.opengl.*; +import static java.lang.Math.min; +import static org.lwjgl.opencl.CL10.CL_PROGRAM_BUILD_LOG; +import static org.lwjgl.opencl.CL10.CL_QUEUE_PROFILING_ENABLE; +import static org.lwjgl.opencl.CL10.clEnqueueNDRangeKernel; +import static org.lwjgl.opencl.CL10.clEnqueueWaitForEvents; +import static org.lwjgl.opencl.CL10.clFinish; +import static org.lwjgl.opencl.CL10.clReleaseMemObject; +import static org.lwjgl.opencl.CL10GL.clEnqueueAcquireGLObjects; +import static org.lwjgl.opencl.KHRGLEvent.clCreateEventFromGLsyncKHR; +import static org.lwjgl.opengl.ARBCLEvent.glCreateSyncFromCLeventARB; +import static org.lwjgl.opengl.ARBSync.GL_SYNC_GPU_COMMANDS_COMPLETE; +import static org.lwjgl.opengl.ARBSync.glFenceSync; +import static org.lwjgl.opengl.ARBSync.glWaitSync; +import static org.lwjgl.opengl.GL11.*; +import static org.lwjgl.opengl.GL15.glBindBuffer; +import static org.lwjgl.opengl.GL15.glDeleteBuffers; +import static org.lwjgl.opengl.GL21.GL_PIXEL_UNPACK_BUFFER; + +/* + * @author labramson + */ +public class Main { + + public static Image img; + public static Image img2; + private Texture inTexture, outTexture; + + static final String source + = "kernel void imgTest(__read_only image2d_t inputTexture){\n" + + " int2 imgCoords = (int2)(get_global_id(0), get_global_id(1));\n" + + " //printf(\"Coord x:%d Coord y:%d \", imgCoords.x, imgCoords.y);\n\n" + + " const sampler_t smp = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;\n\n" + + " float4 imgVal = read_imagef(inputTexture, smp, imgCoords); \n" + + " printf(\"(%d, %d) RGBA(%f, %f, %f, %f)\\n\", imgCoords.x, imgCoords.y, imgVal.x, imgVal.y, imgVal.z, imgVal.w);\n" + + "}"; + static final String code + = "kernel void imgTest2(__read_only image2d_t inTexture, __write_only image2d_t outTexture){\n" + + " const sampler_t smp = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;\n\n" + + " uint offset = get_global_id(1)*0x4000 + get_global_id(0)*0x1000;\n" + + " int2 coord = (int2)(get_global_id(0), get_global_id(1));\n" + + " float4 pixel = read_imagef(inTexture, smp, coord);\n" + + " if (coord.x > 5 & coord.y > 5) { \n" + + " pixel.x /= 1.5;\n" + + " pixel.y /= 1.5;\n" + + " pixel.z /= 1.5;\n" + + " }\n" + + " write_imagef(outTexture, coord, pixel);\n" + + "}"; + + //VARS FOR CL CONVERSION + private CLCommandQueue[] queues; //array of cl command queues + private CLContext context; //CL context + private CLEvent glEvent; //cl event + private CLEvent[] clEvents; //array of cl events for + private CLKernel[] kernels; //array of cl kernels for + private CLMem[] glBuffers; //array of clm for + private CLMem[] glBuffersOut; //array of clm for + private CLProgram[] programs; //array of cl programs for + private GLSync glSync; //glsync so that cl & gl dont cause race condition + private GLSync[] clSyncs; //array of gl sync objects for + private IntBuffer glIDs; //int buffer for + private IntBuffer glIDs2; //int buffer for + private boolean buffersInitialized; //buffers for something initialized + private boolean drawSeparator; //no idea what this is + private boolean rebuild; //boolean for rerendering + private boolean syncCLtoGL; // true if we can make CL wait on sync objects generated from GL. + private boolean syncGLtoCL; // true if we can make GL wait on events generated from CL queues. + private final boolean doublePrecision = true; //doubles used instead of floats + private final boolean useTextures = true; //for something... + private final PointerBuffer kernel2DGlobalWorkSize = BufferUtils.createPointerBuffer(2); //the global work size of the kernel + private final PointerBuffer syncBuffer = BufferUtils.createPointerBuffer(1); //buffer for dealing with gl cl sync + private int deviceType = CL10.CL_DEVICE_TYPE_GPU; + private int slices; //dividing up the image for faster processing + private static final int MAX_GPUS = 8; //Max GPUs used at once + + //CONSTRUCTOR + public Main() { + } + + public static void main(String... args) throws Exception { + //IMAGE FILE + String imgDir = "C:\\Users\\labramson\\Documents\\Tutorial\\res\\"; + String imgName = "texture.png"; + String imgName2 = "blank.jpg"; + //CREATE IMAGE OBJECT + img = new Image(imgDir + "" + imgName); + img2 = new Image(imgDir + "" + imgName2); + + //INIT DISPLAY & GL CONTEXT + initDisplay(); + initGL(); + + //INSTANCE OF MAIN TO RUN FUNCTIONS + Main run = new Main(); + + //SET UP CL + run.initCL(); + + //CREATE THE TEXTURE IN GL CONTEXT + run.initGLTexture(img); + run.initWritableTexture(img2); + + //COMPLETE ALL GL + glFinish(); + + //SET THE KERNEL PARAMS FOR CL COMPUTAIONS + run.setKernelParams(); + + //DISPLAY LOOP + while (!Display.isCloseRequested()) { + //ALLOWS DISPLAYING ON SCREEN + run.display(img); + + //IF ESC THEN CLOSE + if (Keyboard.isKeyDown(Keyboard.KEY_ESCAPE)) { + Display.destroy(); + System.exit(0); + } + + //UPDATES THE DISPLAY WITH 60 FPS + Display.update(); + Display.sync(60); + } + + //DESTROYS THE DISPLAY + Display.destroy(); + } + + public static void initDisplay() { + try { + Display.setDisplayMode(new DisplayMode(300, 300)); + Display.setTitle("Texture Demo"); + Display.create(); + } catch (LWJGLException e) { + e.printStackTrace(); + } + Display.update(); + } + + public static void initGL() { + glMatrixMode(GL_PROJECTION); + glLoadIdentity(); + glOrtho(0, 300, 300, 0, 1, -1); + glMatrixMode(GL_MODELVIEW); + glLoadIdentity(); + } + + //GL SETTINGS FOR THE TEXTURE + public static void glSettings() { + glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); //CLEARS SCREEN EACH LOOP + glDisable(GL_DEPTH_TEST); //DISABLES DEPTH TEST + glEnable(GL_TEXTURE_2D); //ENABLES GL_TEXTURE_2D + glPixelStorei(GL_PACK_ALIGNMENT, 4); + } + + //INITIALIZE THE GL TEXTURE + public void initGLTexture(Image img) { + if (glBuffers == null) { + glBuffers = new CLMem[slices]; + + glIDs = BufferUtils.createIntBuffer(slices); + } else { + for (CLMem mem : glBuffers) { + clReleaseMemObject(mem); + } + + if (useTextures) { + glDeleteTextures(glIDs); + } else { + glDeleteBuffers(glIDs); + } + } + + if (useTextures) { + GL11.glGenTextures(glIDs); + for (int i = 0; i < slices; i++) { + Texture texture = new Texture(img, GL_TEXTURE_2D, glIDs.get(i)); + glBuffers[i] = CL10GL.clCreateFromGLTexture2D(context, CL10.CL_MEM_READ_ONLY, texture.getTarget(), 0, texture.getId(), null); + inTexture = texture; + } + glBindTexture(GL_TEXTURE_2D, 0); + } + } + + public void initWritableTexture(Image img) { + if (glBuffersOut == null) { + glBuffersOut = new CLMem[slices]; + + glIDs2 = BufferUtils.createIntBuffer(slices); + } else { + for (CLMem mem : glBuffersOut) { + clReleaseMemObject(mem); + } + + if (useTextures) { + glDeleteTextures(glIDs2); + } else { + glDeleteBuffers(glIDs2); + } + } + + if (useTextures) { + GL11.glGenTextures(glIDs2); + for (int i = 0; i < slices; i++) { + Texture texture = new Texture(img, GL_TEXTURE_2D, glIDs2.get(i)); + glBuffersOut[i] = CL10GL.clCreateFromGLTexture2D(context, CL10.CL_MEM_WRITE_ONLY, texture.getTarget(), 0, texture.getId(), null); + outTexture = texture; + } + glBindTexture(GL_TEXTURE_2D, 0); + } + buffersInitialized = true; + } + + //INITIALIZE CL PLATFORM, DEVICES, CONTEXT, COMMAND QUEUE, KERNEL + public void initCL() { + try { + // Initialize OpenCL + CL.create(); + System.out.println("\nCL created"); + + // Drawable context that OpenCL needs + Drawable drawable = Display.getDrawable(); + System.out.println("Drawable created"); + + // LWJGL CLPlatform object + CLPlatform platform = CLPlatform.getPlatforms().get(0); + System.out.println("Platform created"); + + final Filter glSharingFilter; + glSharingFilter = (final CLDevice device) -> { + final CLDeviceCapabilities abilities = CLCapabilities.getDeviceCapabilities(device); + return abilities.CL_KHR_gl_sharing; + }; + + // List of LJWGL CLDevice objects representing hardware or software contexts that OpenCL can use + List devices = platform.getDevices(CL10.CL_DEVICE_TYPE_GPU); + if (devices == null) { + deviceType = CL10.CL_DEVICE_TYPE_CPU; + devices = platform.getDevices(CL10.CL_DEVICE_TYPE_CPU, glSharingFilter); + if (devices == null) { + throw new RuntimeException("No OpenCL devices found with KHR_gl_sharing support."); + } + } + System.out.println("Devices obtained"); + + slices = min(devices.size(), MAX_GPUS); + System.out.println("Slices: " + slices); + + // Create the OpenCL context using the patform, devices, and the OpenGL drawable + context = CLContext.create(platform, devices, null, drawable, null); + System.out.println("Context created"); + + // Create an command queue using our OpenCL context + queues = new CLCommandQueue[slices]; + fillQueue(context); + System.out.println("Command Queue created"); + + buildProgram(context); + System.out.println("Program created"); + + kernels = new CLKernel[slices]; + initKernel(); + System.out.println("Kernel initialized"); + + //CHECK SYNC STATUS BTWN GL & CL + syncStatus(context); + + } catch (Exception e) { + System.out.println("*** Problem creating CL context"); + e.printStackTrace(); + } + } + + //CREATES THE CL PROGRAM WITH CORRECT OPTIONS + private void buildProgram(CLContext context) { + programs = new CLProgram[slices]; + if (programs[0] != null) { + for (CLProgram program : programs) { + CL10.clReleaseProgram(program); + } + } + + for (int i = 0; i < programs.length; i++) { + programs[i] = CL10.clCreateProgramWithSource(context, code, null); + } + + for (int i = 0; i < programs.length; i++) { + final CLDevice device = queues[i].getCLDevice(); + final StringBuilder options = new StringBuilder(useTextures ? "-D USE_TEXTURE" : ""); + final CLDeviceCapabilities caps = CLCapabilities.getDeviceCapabilities(device); + + if (doublePrecision) { + //cl_khr_fp64 + options.append(" -D DOUBLE_FP"); + + //amd's verson of double precision floating point math + if (!caps.CL_KHR_fp64 && caps.CL_AMD_fp64) { + options.append(" -D AMD_FP"); + } + } + System.out.println("\nOpenCL COMPILER OPTIONS: " + options); + + try { + CL10.clBuildProgram(programs[i], device, options, null); + } finally { + System.out.println("BUILD LOG: " + programs[i].getBuildInfoString(device, CL_PROGRAM_BUILD_LOG)); + } + } + } + + //FILLS THE COMMAND QUEUE APPROPRIATELY + private void fillQueue(CLContext context) { + for (int i = 0; i < slices; i++) { + // create command queue on each used device + queues[i] = CL10.clCreateCommandQueue(context, context.getInfoDevices().get(i), CL_QUEUE_PROFILING_ENABLE, null); + queues[i].checkValid(); + } + } + + //INITIALIZES THE KERNEL + private void initKernel() { + for (int i = 0; i < kernels.length; i++) { + kernels[i] = CL10.clCreateKernel(programs[min(i, programs.length)], "imgTest2", null); + } + } + + //PASSES IN THE ARGUMENTS TO THE KERNEL + private void setKernelParams() { + for (int i = 0; i < slices; i++) { + kernels[i].setArg(0, glBuffers[i]) + .setArg(1, glBuffersOut[i]); + } + } + + //CHECKS THE SYNC STATUS BTWN GL & CL AND VISE VERSA + private void syncStatus(CLContext context) { + final ContextCapabilities abilities = GLContext.getCapabilities(); + + syncGLtoCL = abilities.GL_ARB_cl_event; // GL3.2 or ARB_sync implied + if (syncGLtoCL) { + clEvents = new CLEvent[slices]; + clSyncs = new GLSync[slices]; + System.out.println("\nGL to CL sync: Using OpenCL events"); + } else { + System.out.println("\nGL to CL sync: Using clFinish"); + } + + // Detect CLtoGL synchronization method + syncCLtoGL = abilities.OpenGL32 || abilities.GL_ARB_sync; + if (syncCLtoGL) { + for (CLDevice device : context.getInfoDevices()) { + if (!CLCapabilities.getDeviceCapabilities(device).CL_KHR_gl_event) { + syncCLtoGL = false; + break; + } + } + } + + if (syncCLtoGL) { + System.out.println("CL to GL sync: Using OpenGL sync objects"); + } else { + System.out.println("CL to GL sync: Using glFinish"); + } + } + + //DISPLAYS THE RESULTS + public void display(Image img) { + //CHECKS TO MAKE SURE ALL GL EVENTS HAVE COMPLETED + if (syncCLtoGL && glEvent != null) { + for (final CLCommandQueue queue : queues) { + clEnqueueWaitForEvents(queue, glEvent); + } + } else { + glFinish(); + } + + //IF THE GL TEXTURE BUFFERS HAVE NOT BEEN INITIALIZED + if (!buffersInitialized) { + initGLTexture(img); + setKernelParams(); + } + + //IF CHANGES OCCURED, AND NEEDS TO REBUILD PROGRAM & KERNEL + if (rebuild) { + buildProgram(context); + setKernelParams(); + } + + //SETS THE WORKSIZE OF THE KERNEL + kernel2DGlobalWorkSize.put(0, img.getWidth()).put(1, img.getHeight()); + + //GETS THE GL OBJECTS + for (int i = 0; i < slices; i++) { + // acquire GL objects, and enqueue a kernel with a probe from the list + clEnqueueAcquireGLObjects(queues[i], glBuffers[i], null, null); + clEnqueueAcquireGLObjects(queues[i], glBuffersOut[i], null, null); + + clEnqueueNDRangeKernel(queues[i], kernels[i], 2, + null, + kernel2DGlobalWorkSize, + null, + null, null); + + CL10GL.clEnqueueReleaseGLObjects(queues[i], glBuffers[i], null, syncGLtoCL ? syncBuffer : null); + CL10GL.clEnqueueReleaseGLObjects(queues[i], glBuffersOut[i], null, syncGLtoCL ? syncBuffer : null); + if (syncGLtoCL) { + clEvents[i] = queues[i].getCLEvent(syncBuffer.get(0)); + clSyncs[i] = glCreateSyncFromCLeventARB(queues[i].getParent(), clEvents[i], 0); + } + } + + // block until done (important: finish before doing further gl work) + if (!syncGLtoCL) { + for (int i = 0; i < slices; i++) { + clFinish(queues[i]); + } + } + + //RENDER THE TEXTURE + render(img); + } + + private void render(Image img) { + if (syncGLtoCL) { + for (int i = 0; i < slices; i++) { + glWaitSync(clSyncs[i], 0, 0); + } + } + + //draw slices + int sliceWidth = img.getWidth() / slices; + + for (int i = 0; i < slices; i++) { + int seperatorOffset = drawSeparator ? i : 0; + + //BIND THE TEXTURE + glBindTexture(GL_TEXTURE_2D, outTexture.getId()); + + //SETS GL SETTINGS + glSettings(); + + //DRAW A SQUARE WITH MAPPED TEXTURE + glBegin(GL_QUADS); + glTexCoord2f(0, 0); + glVertex2i(100, 100); //upper left + + glTexCoord2f(0, 1); + glVertex2i(100, 200); //upper right + + glTexCoord2f(1, 1); + glVertex2i(200, 200); //bottom right + + glTexCoord2f(1, 0); + glVertex2i(200, 100); //bottom left + + + glEnd(); + } + + //CHECKING SYNC + if (syncCLtoGL) { + glSync = glFenceSync(GL_SYNC_GPU_COMMANDS_COMPLETE, 0); + glEvent = clCreateEventFromGLsyncKHR(context, glSync, null); + } + } +} diff --git a/src/tutorial/clTexture/SharingMemCode.txt b/src/tutorial/clTexture/SharingMemCode.txt new file mode 100644 index 0000000..277fdf1 --- /dev/null +++ b/src/tutorial/clTexture/SharingMemCode.txt @@ -0,0 +1,42 @@ +// public CLMem createSharedMem(CLContext context, Texture texture) { +// // Create an OpenGL buffer +// int glBufId = GL15.glGenBuffers(); +// // Load the buffer with data using glBufferData(); +// // initialize buffer object +// +// GL15.glBindBuffer(GL_ARRAY_BUFFER, glBufId); +// +// //int size = (texture.getImageHeight() * texture.getImageWidth() * 4); +// GL15.glBufferData(GL_ARRAY_BUFFER, textureBuff, GL_DYNAMIC_DRAW); +// // Create the shared OpenCL memory object from the OpenGL buffer +// CLMem glMem = CL10GL.clCreateFromGLBuffer(context, CL10.CL_MEM_READ_WRITE, glBufId, null); +// +// return glMem; +// } +// +// public void useSharedMem(CLMem clmem, CLCommandQueue queue) { +// System.out.println("Acquiring mem lock"); +// // Acquire the lock for the 'glMem' memory object +// int error = CL10GL.clEnqueueAcquireGLObjects(queue, clmem, null, null); +// // Remember to check for errors +// if (error != CL10.CL_SUCCESS) { +// Util.checkCLError(error); +// } +// +// // Now execute an OpenCL command using the shared memory object, +// // such as uploading data to the memory object using 'CL10.clEnqueueWriteBuffer()' +// // or running a kernel using 'CL10.clEnqueueNDRangeKernel()' with the correct parameters +// // ... +// System.out.println("Doing Stuff"); +// // Release the lock on the 'glMem' memory object +// error = CL10GL.clEnqueueReleaseGLObjects(queue, clmem, null, null); +// if (error != CL10.CL_SUCCESS) { +// Util.checkCLError(error); +// } +// +// // Remember to flush the command queue when you are done. +// // Flushing the queue ensures that all of the OpenCL commands +// // sent to the queue have completed before the program continues. +// CL10.clFinish(queue); +// //System.out.println("Finished using CL & released shared mem"); +// } \ No newline at end of file diff --git a/src/tutorial/clTexture/Texture.java b/src/tutorial/clTexture/Texture.java index c7736c6..328b9de 100644 --- a/src/tutorial/clTexture/Texture.java +++ b/src/tutorial/clTexture/Texture.java @@ -6,66 +6,66 @@ package tutorial.clTexture; import org.lwjgl.opengl.GL11; +import static org.lwjgl.opengl.GL11.GL_NEAREST; +import static org.lwjgl.opengl.GL11.GL_RGB; +import static org.lwjgl.opengl.GL11.GL_RGB8; +import static org.lwjgl.opengl.GL11.GL_TEXTURE_2D; +import static org.lwjgl.opengl.GL11.GL_TEXTURE_MAG_FILTER; +import static org.lwjgl.opengl.GL11.GL_TEXTURE_MIN_FILTER; +import static org.lwjgl.opengl.GL11.GL_TEXTURE_WRAP_S; +import static org.lwjgl.opengl.GL11.GL_TEXTURE_WRAP_T; +import static org.lwjgl.opengl.GL11.GL_UNSIGNED_BYTE; +import static org.lwjgl.opengl.GL11.glTexImage2D; +import static org.lwjgl.opengl.GL11.glTexParameteri; +import static org.lwjgl.opengl.GL12.GL_CLAMP_TO_EDGE; +/* + * @author LAA + */ public class Texture { - public int target, textureID, height, width, texWidth, texHeight; - private float widthRatio, heightRatio; + private final int width, height, target, id; + private final Image image; - public Texture(int target, int textureID) { + public Texture(Image image, int target, int id) { + this.image = image; + this.width = image.getWidth(); + this.height = image.getHeight(); this.target = target; - this.textureID = textureID; - } + this.id = id; - public void bind() { - GL11.glBindTexture(target, textureID); - } + // bind this texture + GL11.glBindTexture(target, id); - public void setWidth(int width) { - this.width = width; - setWidth(); - } + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_NEAREST); + glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_NEAREST); - public void setHeight(int height) { - this.height = height; - setHeight(); + glTexImage2D(GL_TEXTURE_2D, 0, GL_RGB8, image.getWidth() + 1, image.getHeight(), 0, GL_RGB, GL_UNSIGNED_BYTE, image.getByeBuff()); } - public int getImageWidth() { - return width; - } - - public int getImageHeight() { - return height; - } - - public float getWidth() { - return widthRatio; + public void bind() { + GL11.glBindTexture(target, id); } - public float getHeight() { - return heightRatio; + public int getTarget() { + return target; } - public void setTextureWidth(int texWidth) { - this.texWidth = texWidth; - setWidth(); + public Image getImage() { + return image; } - public void setTextureHeight(int texHeight) { - this.texHeight = texHeight; - setHeight(); + public int getId() { + return id; } - private void setWidth() { - if (texWidth != 0) { - widthRatio = ((float) width) / texWidth; - } + public int getWidth() { + return width; } - private void setHeight() { - if (texHeight != 0) { - heightRatio = ((float) height) / texHeight; - } + public int getHeight() { + return height; } } diff --git a/src/tutorial/clTexture/TextureDriver.java b/src/tutorial/clTexture/TextureDriver.java deleted file mode 100644 index 6f0ecb0..0000000 --- a/src/tutorial/clTexture/TextureDriver.java +++ /dev/null @@ -1,95 +0,0 @@ -/* - * To change this license header, choose License Headers in Project Properties. - * To change this template file, choose Tools | Templates - * and open the template in the editor. - */ -package tutorial.clTexture; - -import static org.lwjgl.opengl.GL11.*; -import org.lwjgl.opengl.*; -import org.lwjgl.*; -import org.lwjgl.input.Keyboard; - -/* - * @author labramson - */ -public class TextureDriver { - - //IMAGE FOR THE TEXTURE - private static final String imgDir = "C:\\Users\\labramson\\Documents\\Tutorial\\res\\"; - public static String imgName = "smileTexture2.jpg"; - - //CONSTRUCTOR - public TextureDriver() {} - - public static void main(String... args) throws Exception { - initDisplay(); - initGL(); - - Image img = new Image(imgDir + "" + imgName); - Texture texture = new Texture(GL_TEXTURE_2D, GL11.glGenTextures()); - - //OBJECT USED TO MAKE THE TEXTURE - TextureMaker tex = new TextureMaker(img.getImg(), texture); - texture = tex.makeTexture(); - tex.makeCLTexture(); - - while (!Display.isCloseRequested()) { - //CLEARS SCREEN EACH LOOP - glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); - //ENABLES GL_TEXTURE_2D - glEnable(GL_TEXTURE_2D); - - //BIND THE TEXTURE - texture.bind(); - - glPixelStorei(GL_PACK_ALIGNMENT, 4); - - //DRAW A SQUARE WITH MAPPED TEXTURE - glBegin(GL_QUADS); - glTexCoord2f(0, 0); - glVertex2i(100, 100); //upper left - - glTexCoord2f(0, 1); - glVertex2i(100, 200); //upper right - - glTexCoord2f(1, 1); - glVertex2i(200, 200); //bottom right - - glTexCoord2f(1, 0); - glVertex2i(200, 100); //bottom left - glEnd(); - - //IF ESC THEN CLOSE - if (Keyboard.isKeyDown(Keyboard.KEY_ESCAPE)) { - Display.destroy(); - System.exit(0); - } - - //UPDATES THE DISPLAY WITH 60 FPS - Display.update(); - Display.sync(60); - } - - //DESTROYS THE DISPLAY - Display.destroy(); - } - - public static void initDisplay() { - try { - Display.setDisplayMode(new DisplayMode(300, 300)); - Display.setTitle("Texture Demo"); - Display.create(); - } catch (LWJGLException e) { - e.printStackTrace(); - } - Display.update(); - } - - public static void initGL() { - glMatrixMode(GL_PROJECTION); - glLoadIdentity(); - glOrtho(0, 300, 300, 0, 1, -1); - glMatrixMode(GL_MODELVIEW); - } -} \ No newline at end of file diff --git a/src/tutorial/clTexture/TextureMaker.java b/src/tutorial/clTexture/TextureMaker.java deleted file mode 100644 index e83c893..0000000 --- a/src/tutorial/clTexture/TextureMaker.java +++ /dev/null @@ -1,248 +0,0 @@ -/* - * To change this license header, choose License Headers in Project Properties. - * To change this template file, choose Tools | Templates - * and open the template in the editor. - */ -package tutorial.clTexture; - -//Java imports -import java.awt.Graphics; -import java.awt.Color; -import java.awt.color.ColorSpace; -import java.awt.image.BufferedImage; -import java.awt.image.ColorModel; -import java.awt.image.ComponentColorModel; -import java.awt.image.DataBuffer; -import java.awt.image.DataBufferByte; -import java.awt.image.Raster; -import java.awt.image.WritableRaster; -import java.nio.ByteBuffer; -import java.nio.ByteOrder; -import java.util.HashMap; -import java.util.Hashtable; -import java.util.List; - -//GL imports -import org.lwjgl.opengl.GL11; -import org.lwjgl.opengl.GL15; -import static org.lwjgl.opengl.GL11.GL_LINEAR; -import static org.lwjgl.opengl.GL11.GL_RGB; -import static org.lwjgl.opengl.GL11.GL_RGB8; -import static org.lwjgl.opengl.GL11.GL_TEXTURE_2D; -import static org.lwjgl.opengl.GL11.GL_TEXTURE_MAG_FILTER; -import static org.lwjgl.opengl.GL11.GL_TEXTURE_MIN_FILTER; -import static org.lwjgl.opengl.GL11.GL_UNSIGNED_BYTE; -import static org.lwjgl.opengl.GL11.glTexImage2D; -import static org.lwjgl.opengl.GL11.glTexParameteri; -import static org.lwjgl.opengl.GL15.GL_ARRAY_BUFFER; -import static org.lwjgl.opengl.GL15.GL_DYNAMIC_DRAW; -import org.lwjgl.opengl.Display; -import org.lwjgl.opengl.Drawable; - -//CL imports -import org.lwjgl.opencl.CL; -import org.lwjgl.opencl.CL10; -import org.lwjgl.opencl.CL10GL; -import org.lwjgl.opencl.CLContext; -import org.lwjgl.opencl.CLDevice; -import org.lwjgl.opencl.CLPlatform; -import org.lwjgl.opencl.CLCommandQueue; -import org.lwjgl.opencl.CLMem; -import org.lwjgl.opencl.Util; - -/* - * @author labramson - */ -public class TextureMaker { - - /* The table of textures that have been loaded in this loader */ - private HashMap table = new HashMap(); - /* The color model including alpha for the GL image */ - private ColorModel glAlphaColorModel; - /* The color model for the GL image */ - private ColorModel glColorModel; - - private Texture texture; - private BufferedImage image; - private ByteBuffer textureBuff; - - public TextureMaker(BufferedImage image, Texture texture) { - this.image = image; - this.texture = texture; - this.glAlphaColorModel = new ComponentColorModel(ColorSpace.getInstance(ColorSpace.CS_sRGB), - new int[]{8, 8, 8, 8}, - true, - false, - ComponentColorModel.TRANSLUCENT, - DataBuffer.TYPE_BYTE); - - this.glColorModel = new ComponentColorModel(ColorSpace.getInstance(ColorSpace.CS_sRGB), - new int[]{8, 8, 8, 0}, - false, - false, - ComponentColorModel.OPAQUE, - DataBuffer.TYPE_BYTE); - } - - public Texture makeTexture() { - int textureID = this.texture.textureID; - int textureTarget = this.texture.target; - Texture texture = new Texture(textureTarget, textureID); - int srcPixelFormat = 0; - // bind this texture - GL11.glBindTexture(textureTarget, textureID); - - BufferedImage bufferedImage = this.image; - texture.setWidth(bufferedImage.getWidth()); - texture.setHeight(bufferedImage.getHeight()); - - if (bufferedImage.getColorModel().hasAlpha()) { - srcPixelFormat = GL11.GL_RGBA; - } else { - srcPixelFormat = GL11.GL_RGB; - } - - ByteBuffer texBuff = imgToTexBuff(bufferedImage, texture); - textureBuff = texBuff; - glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MIN_FILTER, GL_LINEAR); - glTexParameteri(GL_TEXTURE_2D, GL_TEXTURE_MAG_FILTER, GL_LINEAR); - - glTexImage2D(GL_TEXTURE_2D, 0, GL_RGB8, bufferedImage.getWidth(), bufferedImage.getHeight(), 0, GL_RGB, GL_UNSIGNED_BYTE, texBuff); - - return texture; - } - - private ByteBuffer imgToTexBuff(BufferedImage bufferedImage, Texture texture) { - ByteBuffer imageBuffer = null; - WritableRaster raster; - BufferedImage texImage; - - int texWidth = 2; - int texHeight = 2; - - // find the closest power of 2 for the width and height of the produced texture - while (texWidth < bufferedImage.getWidth()) { - texWidth *= 2; - } - while (texHeight < bufferedImage.getHeight()) { - texHeight *= 2; - } - - texture.setTextureHeight(texHeight); - texture.setTextureWidth(texWidth); - - // create a raster that can be used by OpenGL as a source for a texture - if (bufferedImage.getColorModel().hasAlpha()) { - raster = Raster.createInterleavedRaster(DataBuffer.TYPE_BYTE, texWidth, texHeight, 4, null); - texImage = new BufferedImage(glAlphaColorModel, raster, false, new Hashtable()); - } else { - raster = Raster.createInterleavedRaster(DataBuffer.TYPE_BYTE, texWidth, texHeight, 3, null); - texImage = new BufferedImage(glColorModel, raster, false, new Hashtable()); - } - - // copy the source image into the produced image - Graphics g = texImage.getGraphics(); - g.setColor(new Color(0f, 0f, 0f, 0f)); - g.fillRect(0, 0, texWidth, texHeight); - g.drawImage(bufferedImage, 0, 0, null); - - // build a byte buffer from the temporary image used by OpenGL to produce a texture - byte[] data = ((DataBufferByte) texImage.getRaster().getDataBuffer()).getData(); - - imageBuffer = ByteBuffer.allocateDirect(data.length); - imageBuffer.order(ByteOrder.nativeOrder()); - imageBuffer.put(data, 0, data.length); - imageBuffer.flip(); - - return imageBuffer; - } - - public void makeCLTexture() { - try { - // Initialize OpenCL and create a context and command queue - CL.create(); - System.out.println("\n****************"); - System.out.println("CL created"); - - // Drawable context that OpenCL needs - Drawable drawable = Display.getDrawable(); - System.out.println("Drawable created"); - - // LWJGL CLPlatform object - CLPlatform platform = CLPlatform.getPlatforms().get(0); - System.out.println("Platform created"); - - // List of LJWGL CLDevice objects representing hardware or software contexts that OpenCL can use - List devices = platform.getDevices(CL10.CL_DEVICE_TYPE_GPU); - System.out.println("Devices obtained"); - - // Create the OpenCL context using the patform, devices, and the OpenGL drawable - CLContext context = CLContext.create(platform, devices, null, drawable, null); - System.out.println("Context created"); - - // Create an command queue using our OpenCL context and the first device in our list of devices - CLCommandQueue queue = CL10.clCreateCommandQueue(context, devices.get(0), CL10.CL_QUEUE_PROFILING_ENABLE, null); - System.out.println("Command Queue created"); - - //CLMem clglMem = createSharedMem(context, texture); - //CLMem gltoCLTex = CL10GL.clCreateFromGLTexture(context, CL_MEM_READ_WRITE, GL_TEXTURE_BUFFER, 0,texture, errcode_ret); - - // use the shared memory in CL - //useSharedMem(clglMem, queue); - //System.out.println("Shared Mem used"); - - - } catch (Exception e) { - e.printStackTrace(); - System.out.println("*** Problem initializing OpenCL"); - } - } - - public CLMem createSharedMem(CLContext context, Texture texture) { - // Create an OpenGL buffer - int glBufId = GL15.glGenBuffers(); - // Load the buffer with data using glBufferData(); - // initialize buffer object - - GL15.glBindBuffer(GL_ARRAY_BUFFER, glBufId); - - //int size = (texture.getImageHeight() * texture.getImageWidth() * 4); - GL15.glBufferData(GL_ARRAY_BUFFER, textureBuff, GL_DYNAMIC_DRAW); - // Create the shared OpenCL memory object from the OpenGL buffer - CLMem glMem = CL10GL.clCreateFromGLBuffer(context, CL10.CL_MEM_READ_WRITE, glBufId, null); - - return glMem; - } - - public void useSharedMem(CLMem clmem, CLCommandQueue queue){ - System.out.println("Acquiring mem lock"); - // Acquire the lock for the 'glMem' memory object - int error = CL10GL.clEnqueueAcquireGLObjects(queue, clmem, null, null); - // Remember to check for errors - if(error != CL10.CL_SUCCESS) { Util.checkCLError(error); } - - // Now execute an OpenCL command using the shared memory object, - // such as uploading data to the memory object using 'CL10.clEnqueueWriteBuffer()' - // or running a kernel using 'CL10.clEnqueueNDRangeKernel()' with the correct parameters - // ... - System.out.println("Doing Stuff"); - // Release the lock on the 'glMem' memory object - error = CL10GL.clEnqueueReleaseGLObjects(queue, clmem, null, null); - if(error != CL10.CL_SUCCESS) { Util.checkCLError(error); } - - // Remember to flush the command queue when you are done. - // Flushing the queue ensures that all of the OpenCL commands - // sent to the queue have completed before the program continues. - CL10.clFinish(queue); - deleteSharedMem(clmem); - //System.out.println("Finished using CL & released shared mem"); - } - - public void deleteSharedMem(CLMem clmem){ - // Delete/release an OpenCL shared memory object - CL10.clReleaseMemObject(clmem); - } - - public void textureViaGLTex(){ - } -} diff --git a/src/tutorial/clTexture3d/Image.java b/src/tutorial/clTexture3d/Image.java new file mode 100644 index 0000000..357f2e0 --- /dev/null +++ b/src/tutorial/clTexture3d/Image.java @@ -0,0 +1,110 @@ +package tutorial.clTexture3d; + +import java.awt.Color; +import java.awt.Graphics; +import java.awt.color.ColorSpace; +import java.awt.image.BufferedImage; +import java.awt.image.ColorModel; +import java.awt.image.ComponentColorModel; +import java.awt.image.DataBuffer; +import java.awt.image.DataBufferByte; +import java.awt.image.Raster; +import java.awt.image.WritableRaster; +import java.io.FileInputStream; +import java.nio.ByteBuffer; +import java.nio.ByteOrder; +import java.util.Hashtable; +import javax.imageio.ImageIO; + +/* + * @author LAA + */ +public class Image { + private final int width, height, depth; + private final BufferedImage img; + + /* The color model including alpha for the GL image */ + private final ColorModel glAlphaColorModel; + /* The color model for the GL image */ + private final ColorModel glColorModel; + + public Image(String imgFile) { + this.img = getBuffImage(imgFile); + this.width = img.getWidth(); + this.height = img.getHeight(); + this.depth = getDepth(); + + this.glAlphaColorModel = new ComponentColorModel(ColorSpace.getInstance(ColorSpace.CS_sRGB), + new int[]{8, 8, 8, 8}, + true, + false, + ComponentColorModel.TRANSLUCENT, + DataBuffer.TYPE_BYTE); + this.glColorModel = new ComponentColorModel(ColorSpace.getInstance(ColorSpace.CS_sRGB), + new int[]{8, 8, 8, 0}, + false, + false, + ComponentColorModel.OPAQUE, + DataBuffer.TYPE_BYTE); + } + + private BufferedImage getBuffImage(String imageFile) { + System.out.println("Fetching image file"); + try { + BufferedImage img = ImageIO.read(new FileInputStream(imageFile)); + return img; + } catch (Exception e) { + e.printStackTrace(); + return null; + } + } + + public ByteBuffer getByeBuff() { + ByteBuffer byteBuffer; + WritableRaster raster; + BufferedImage texImage; + + System.out.println("Making image byte buffer"); + + // create a raster that can be used by OpenGL as a source for a texture + if (this.img.getColorModel().hasAlpha()) { + raster = Raster.createInterleavedRaster(DataBuffer.TYPE_BYTE, this.img.getWidth() + 1, this.img.getHeight(), 4, null); + texImage = new BufferedImage(glAlphaColorModel, raster, false, new Hashtable()); + } else { + raster = Raster.createInterleavedRaster(DataBuffer.TYPE_BYTE, this.img.getWidth() + 1, this.img.getHeight(), 3, null); + texImage = new BufferedImage(glColorModel, raster, false, new Hashtable()); + } + + // copy the source image into the produced image + Graphics g = texImage.getGraphics(); + g.setColor(new Color(0f, 0f, 0f, 0f)); + g.fillRect(0, 0, this.img.getWidth() + 1, this.img.getHeight()); + g.drawImage(this.img, 0, 0, null); + + // build a byte buffer from the temporary image to produce a texture + byte[] data = ((DataBufferByte) texImage.getRaster().getDataBuffer()).getData(); + + byteBuffer = ByteBuffer.allocateDirect(data.length); + byteBuffer.order(ByteOrder.nativeOrder()); + byteBuffer.put(data, 0, data.length); + byteBuffer.flip(); + + return byteBuffer; + } + + public int getWidth() { + return width; + } + + public int getHeight() { + return height; + } + + public int getDepth(){ + return depth; + } + + public BufferedImage getImg() { + return img; + } +} diff --git a/src/tutorial/clTexture3d/Main.java b/src/tutorial/clTexture3d/Main.java new file mode 100644 index 0000000..80a6081 --- /dev/null +++ b/src/tutorial/clTexture3d/Main.java @@ -0,0 +1,594 @@ +/* + * To change this license header, choose License Headers in Project Properties. + * To change this template file, choose Tools | Templates + * and open the template in the editor. + */ +package tutorial.clTexture3d; + +import com.sun.prism.impl.BufferUtil; +import java.nio.IntBuffer; +import java.util.List; +import org.lwjgl.*; +import org.lwjgl.input.Keyboard; +import org.lwjgl.opencl.CL10; +import org.lwjgl.opencl.CL10GL; +import org.lwjgl.opencl.CL; +import org.lwjgl.opencl.CLCapabilities; +import org.lwjgl.opencl.CLCommandQueue; +import org.lwjgl.opencl.CLContext; +import org.lwjgl.opencl.CLDevice; +import org.lwjgl.opencl.CLDeviceCapabilities; +import org.lwjgl.opencl.CLEvent; +import org.lwjgl.opencl.CLKernel; +import org.lwjgl.opencl.CLMem; +import org.lwjgl.opencl.CLPlatform; +import org.lwjgl.opencl.CLProgram; +import org.lwjgl.opencl.api.Filter; +import org.lwjgl.opengl.*; +import static java.lang.Math.min; +import java.nio.ByteBuffer; +import java.nio.ByteOrder; +import java.nio.ShortBuffer; +import static org.lwjgl.opencl.CL10.CL_PROGRAM_BUILD_LOG; +import static org.lwjgl.opencl.CL10.CL_QUEUE_PROFILING_ENABLE; +import static org.lwjgl.opencl.CL10.clEnqueueNDRangeKernel; +import static org.lwjgl.opencl.CL10.clEnqueueWaitForEvents; +import static org.lwjgl.opencl.CL10.clFinish; +import static org.lwjgl.opencl.CL10.clReleaseMemObject; +import static org.lwjgl.opencl.CL10GL.clEnqueueAcquireGLObjects; +import static org.lwjgl.opencl.KHRGLEvent.clCreateEventFromGLsyncKHR; +import static org.lwjgl.opengl.ARBCLEvent.glCreateSyncFromCLeventARB; +import static org.lwjgl.opengl.ARBSync.GL_SYNC_GPU_COMMANDS_COMPLETE; +import static org.lwjgl.opengl.ARBSync.glFenceSync; +import static org.lwjgl.opengl.ARBSync.glWaitSync; +import static org.lwjgl.opengl.ARBTextureRg.GL_R16; +import static org.lwjgl.opengl.ARBTextureRg.GL_R16F; +import static org.lwjgl.opengl.GL11.*; +import static org.lwjgl.opengl.GL12.*; +//import static org.lwjgl.opengl.GL15.glBindBuffer; +import static org.lwjgl.opengl.GL15.glDeleteBuffers; +import static org.lwjgl.opengl.GL30.GL_RGBA16F; +//import static org.lwjgl.opengl.GL21.GL_PIXEL_UNPACK_BUFFER; + +/* + * @author labramson + */ +public class Main { + //public static Image img; + //public static Image img2; + private Texture inTexture, outTexture; + + static final String kernel_Sample + = "kernel void imgTest(__read_only image2d_t inputTexture){\n" + + " int2 imgCoords = (int2)(get_global_id(0), get_global_id(1));\n" + + " //printf(\"Coord x:%d Coord y:%d \", imgCoords.x, imgCoords.y);\n\n" + + " const sampler_t smp = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;\n\n" + + " float4 imgVal = read_imagef(inputTexture, smp, imgCoords); \n" + + " printf(\"(%d, %d) RGBA(%f, %f, %f, %f)\\n\", imgCoords.x, imgCoords.y, imgVal.x, imgVal.y, imgVal.z, imgVal.w);\n" + + "}"; + static final String kernel_3DImage + = "kernel void imgTest2(__read_only image3d_t inTexture, __write_only image3d_t outTexture){\n" + + " #pragma OPENCL EXTENSION cl_khr_3d_image_writes : enable\n" + + " const sampler_t smp = CLK_NORMALIZED_COORDS_TRUE | CLK_ADDRESS_CLAMP_TO_EDGE | CLK_FILTER_NEAREST;\n\n" + + " int4 coord = (int4)(get_global_id(0), get_global_id(1), get_global_id(2), 0);\n" + + " int4 coordL = (int4)(get_global_id(0)-1, get_global_id(1), get_global_id(2), 0);\n" + + " int4 coordR = (int4)(get_global_id(0)+1, get_global_id(1), get_global_id(2), 0);\n" + + " int4 coordT = (int4)(get_global_id(0), get_global_id(1)-1, get_global_id(2), 0);\n" + + " int4 coordBo = (int4)(get_global_id(0), get_global_id(1)+1, get_global_id(2), 0);\n" + + " int4 coordF = (int4)(get_global_id(0), get_global_id(1), get_global_id(2)-1, 0);\n" + + " int4 coordBa = (int4)(get_global_id(0), get_global_id(1), get_global_id(2)-1, 0);\n" + + " float4 pixel = read_imagef(inTexture, smp, coord);\n" + + " float4 pixelL = read_imagef(inTexture, smp, coordL);\n" + + " float4 pixelR = read_imagef(inTexture, smp, coordR);\n" + + " float4 pixelT = read_imagef(inTexture, smp, coordT);\n" + + " float4 pixelBo = read_imagef(inTexture, smp, coordBo);\n" + + " float4 pixelF = read_imagef(inTexture, smp, coordF);\n" + + " float4 pixelBa = read_imagef(inTexture, smp, coordBa);\n" + + " pixel.x = (pixelR.x - pixelL.x); \n" + + " if (pixel.x < 0.0) {pixel.x = -pixel.x;} \n" + + " pixel.y = pixelT.x - pixelBo.x; \n" + + " if (pixel.y < 0.0) {pixel.y = -pixel.y;} \n" + + " pixel.z = pixelF.x - pixelBa.x; \n" + + " if (pixel.z < 0.0) {pixel.z = -pixel.z;} \n" + + " write_imagef(outTexture, coord, pixel);\n" + + "}"; + + //VARS FOR CL CONVERSION + private CLCommandQueue[] queues; //array of cl command queues + private CLContext context; //CL context + private CLEvent glEvent; //cl event + private CLEvent[] clEvents; //array of cl events for + private CLKernel[] kernels; //array of cl kernels for + private CLMem[] glBuffers; //array of clm for + private CLMem[] glBuffersOut; //array of clm for + private CLProgram[] programs; //array of cl programs for + private GLSync glSync; //glsync so that cl & gl dont cause race condition + private GLSync[] clSyncs; //array of gl sync objects for + private IntBuffer glIDs; //int buffer for + private IntBuffer glIDs2; //int buffer for + private boolean buffersInitialized; //buffers for something initialized + private boolean drawSeparator; //no idea what this is + private boolean rebuild; //boolean for rerendering + private boolean syncCLtoGL; // true if we can make CL wait on sync objects generated from GL. + private boolean syncGLtoCL; // true if we can make GL wait on events generated from CL queues. + private final boolean doublePrecision = true; //doubles used instead of floats + private final boolean useTextures = true; //for something... + private final PointerBuffer kernel2DGlobalWorkSize = BufferUtils.createPointerBuffer(3); //the global work size of the kernel + private final PointerBuffer syncBuffer = BufferUtils.createPointerBuffer(1); //buffer for dealing with gl cl sync + private int deviceType = CL10.CL_DEVICE_TYPE_GPU; + private int slices; //dividing up the image for faster processing + private static final int MAX_GPUS = 8; //Max GPUs used at once + + //CONSTRUCTOR + public Main() { + } + + public static void main(String... args) throws Exception { + //INIT DISPLAY & GL CONTEXT + initDisplay(); + initGL(); + + //INSTANCE OF MAIN TO RUN FUNCTIONS + Main run = new Main(); + + //SET UP CL + run.initCL(); + + //CREATE THE TEXTURE IN GL CONTEXT + run.initGLTexture(); + run.initWritableTexture(); + + //COMPLETE ALL GL + glFinish(); + + //SET THE KERNEL PARAMS FOR CL COMPUTAIONS + run.setKernelParams(); + + //DISPLAY LOOP + while (!Display.isCloseRequested()) { + //ALLOWS DISPLAYING ON SCREEN + run.display(); + + //IF ESC THEN CLOSE + if (Keyboard.isKeyDown(Keyboard.KEY_ESCAPE)) { + Display.destroy(); + System.exit(0); + } + + //UPDATES THE DISPLAY WITH 60 FPS + Display.update(); + Display.sync(60); + } + + //DESTROYS THE DISPLAY + Display.destroy(); + } + + public static void initDisplay() { + try { + Display.setDisplayMode(new DisplayMode(300, 300)); + Display.setTitle("Texture Demo"); + Display.create(); + } catch (LWJGLException e) { + e.printStackTrace(); + } + Display.update(); + } + + public static void initGL() { + glMatrixMode(GL_PROJECTION); + glLoadIdentity(); + glOrtho(0, 300, 300, 0, 10, -300); + glMatrixMode(GL_MODELVIEW); + glLoadIdentity(); + } + + //GL SETTINGS FOR THE TEXTURE + public static void glSettings() { + glClear(GL_COLOR_BUFFER_BIT | GL_DEPTH_BUFFER_BIT); //CLEARS SCREEN EACH LOOP + glEnable(GL_DEPTH_TEST); //DISABLES DEPTH TEST + glEnable(GL_TEXTURE_3D); //ENABLES GL_TEXTURE_3D + glPixelStorei(GL_PACK_ALIGNMENT, 4); + } + + //INITIALIZE THE GL TEXTURE + public void initGLTexture() throws Exception { + if (glBuffers == null) { + glBuffers = new CLMem[slices]; + + glIDs = BufferUtils.createIntBuffer(slices); + } else { + for (CLMem mem : glBuffers) { + clReleaseMemObject(mem); + } + + if (useTextures) { + glDeleteTextures(glIDs); + } else { + glDeleteBuffers(glIDs); + } + } + + if (useTextures) { + GL11.glGenTextures(glIDs); + for (int i = 0; i < slices; i++) { + ByteBuffer bbuf = BufferUtil.newByteBuffer(64 * 64 * 64 * 2); + ShortBuffer sbuf = bbuf.asShortBuffer(); + for (int z=0; z<64; z++) { + for (int y=0; y<64; y++) { + for (int x=0; x<64; x++) { + short value = 0; + if ( x>16 && x<50 && y>16 && y<50 && z>16 && z<50) { + value = 32767; + } + else if ( x>18 && x<48 && y>18 && y<48 && z>18 && z<48) { + value = 128; + } + sbuf.put(value); + } + } + } + + + bbuf.flip(); + bbuf.order(ByteOrder.LITTLE_ENDIAN); + + Texture texture = new Texture(64, 64, 64, GL_R16F, bbuf); + glBuffers[i] = CL10GL.clCreateFromGLTexture3D(context, CL10.CL_MEM_READ_ONLY, texture.getTarget(), 0, texture.getId(), null); + inTexture = texture; + } + //UNBINDS THE TEXTURE + glBindTexture(GL_TEXTURE_3D, 0); + } + } + + public void initWritableTexture() throws Exception { + if (glBuffersOut == null) { + glBuffersOut = new CLMem[slices]; + + glIDs2 = BufferUtils.createIntBuffer(slices); + } else { + for (CLMem mem : glBuffersOut) { + clReleaseMemObject(mem); + } + + if (useTextures) { + glDeleteTextures(glIDs2); + } else { + glDeleteBuffers(glIDs2); + } + } + + if (useTextures) { + GL11.glGenTextures(glIDs2); + for (int i = 0; i < slices; i++) { + Texture texture = new Texture(64, 64, 64, GL_RGBA16F, null); + glBuffersOut[i] = CL10GL.clCreateFromGLTexture3D(context, CL10.CL_MEM_WRITE_ONLY, texture.getTarget(), 0, texture.getId(), null); + outTexture = texture; + } + glBindTexture(GL_TEXTURE_3D, 0); + } + buffersInitialized = true; + } + + //INITIALIZE CL PLATFORM, DEVICES, CONTEXT, COMMAND QUEUE, KERNEL + public void initCL() { + try { + // Initialize OpenCL + CL.create(); + System.out.println("\nCL created"); + + // Drawable context that OpenCL needs + Drawable drawable = Display.getDrawable(); + System.out.println("Drawable created"); + + // LWJGL CLPlatform object + CLPlatform platform = CLPlatform.getPlatforms().get(0); + System.out.println("Platform created"); + + final Filter glSharingFilter; + glSharingFilter = (final CLDevice device) -> { + final CLDeviceCapabilities abilities = CLCapabilities.getDeviceCapabilities(device); + System.out.println(abilities.CL_KHR_3d_image_writes); + return abilities.CL_KHR_gl_sharing; + }; + + // List of LJWGL CLDevice objects representing hardware or software contexts that OpenCL can use + List devices = platform.getDevices(CL10.CL_DEVICE_TYPE_GPU); + if (devices == null) { + deviceType = CL10.CL_DEVICE_TYPE_CPU; + devices = platform.getDevices(CL10.CL_DEVICE_TYPE_CPU, glSharingFilter); + if (devices == null) { + throw new RuntimeException("No OpenCL devices found with KHR_gl_sharing support."); + } + } + System.out.println("Devices obtained"); + + slices = min(devices.size(), MAX_GPUS); + System.out.println("Slices: " + slices); + + // Create the OpenCL context using the patform, devices, and the OpenGL drawable + context = CLContext.create(platform, devices, null, drawable, null); + System.out.println("Context created"); + + // Create an command queue using our OpenCL context + queues = new CLCommandQueue[slices]; + fillQueue(context); + System.out.println("Command Queue created"); + + buildProgram(context); + System.out.println("Program created"); + + kernels = new CLKernel[slices]; + initKernel(); + System.out.println("Kernel initialized"); + + //CHECK SYNC STATUS BTWN GL & CL + syncStatus(context); + + } catch (RuntimeException | LWJGLException e) { + System.out.println("*** Problem creating CL context"); + e.printStackTrace(); + } + } + + //CREATES THE CL PROGRAM WITH CORRECT OPTIONS + private void buildProgram(CLContext context) { + programs = new CLProgram[slices]; + if (programs[0] != null) { + for (CLProgram program : programs) { + CL10.clReleaseProgram(program); + } + } + + for (int i = 0; i < programs.length; i++) { + programs[i] = CL10.clCreateProgramWithSource(context, kernel_3DImage, null); + } + + for (int i = 0; i < programs.length; i++) { + final CLDevice device = queues[i].getCLDevice(); + final StringBuilder options = new StringBuilder(useTextures ? "-D USE_TEXTURE" : ""); + final CLDeviceCapabilities caps = CLCapabilities.getDeviceCapabilities(device); + + if (doublePrecision) { + //cl_khr_fp64 + options.append(" -D DOUBLE_FP"); + + //amd's verson of double precision floating point math + if (!caps.CL_KHR_fp64 && caps.CL_AMD_fp64) { + options.append(" -D AMD_FP"); + } + } + System.out.println("\nOpenCL COMPILER OPTIONS: " + options); + + try { + CL10.clBuildProgram(programs[i], device, options, null); + } finally { + System.out.println("BUILD LOG: " + programs[i].getBuildInfoString(device, CL_PROGRAM_BUILD_LOG)); + } + } + } + + //FILLS THE COMMAND QUEUE APPROPRIATELY + private void fillQueue(CLContext context) { + for (int i = 0; i < slices; i++) { + // create command queue on each used device + queues[i] = CL10.clCreateCommandQueue(context, context.getInfoDevices().get(i), CL_QUEUE_PROFILING_ENABLE, null); + queues[i].checkValid(); + } + } + + //INITIALIZES THE KERNEL + private void initKernel() { + for (int i = 0; i < kernels.length; i++) { + kernels[i] = CL10.clCreateKernel(programs[min(i, programs.length)], "imgTest2", null); + } + } + + //PASSES IN THE ARGUMENTS TO THE KERNEL + private void setKernelParams() { + for (int i = 0; i < slices; i++) { + kernels[i].setArg(0, glBuffers[i]) + .setArg(1, glBuffersOut[i]); + } + } + + //CHECKS THE SYNC STATUS BTWN GL & CL AND VISE VERSA + private void syncStatus(CLContext context) { + final ContextCapabilities abilities = GLContext.getCapabilities(); + + syncGLtoCL = abilities.GL_ARB_cl_event; // GL3.2 or ARB_sync implied + if (syncGLtoCL) { + clEvents = new CLEvent[slices]; + clSyncs = new GLSync[slices]; + System.out.println("\nGL to CL sync: Using OpenCL events"); + } else { + System.out.println("\nGL to CL sync: Using clFinish"); + } + + // Detect CLtoGL synchronization method + syncCLtoGL = abilities.OpenGL32 || abilities.GL_ARB_sync; + if (syncCLtoGL) { + for (CLDevice device : context.getInfoDevices()) { + if (!CLCapabilities.getDeviceCapabilities(device).CL_KHR_gl_event) { + syncCLtoGL = false; + break; + } + } + } + + if (syncCLtoGL) { + System.out.println("CL to GL sync: Using OpenGL sync objects"); + } else { + System.out.println("CL to GL sync: Using glFinish"); + } + } + + //DISPLAYS THE RESULTS + public void display() throws Exception { + //CHECKS TO MAKE SURE ALL GL EVENTS HAVE COMPLETED + if (syncCLtoGL && glEvent != null) { + for (final CLCommandQueue queue : queues) { + clEnqueueWaitForEvents(queue, glEvent); + } + } else { + glFinish(); + } + + //IF THE GL TEXTURE BUFFERS HAVE NOT BEEN INITIALIZED + if (!buffersInitialized) { + initGLTexture(); + setKernelParams(); + } + + //IF CHANGES OCCURED, AND NEEDS TO REBUILD PROGRAM & KERNEL + if (rebuild) { + buildProgram(context); + setKernelParams(); + } + + //SETS THE WORKSIZE OF THE KERNEL + kernel2DGlobalWorkSize.put(0, 64).put(1, 64).put(2, 64); + + //GETS THE GL OBJECTS + for (int i = 0; i < slices; i++) { + // acquire GL objects, and enqueue a kernel with a probe from the list + clEnqueueAcquireGLObjects(queues[i], glBuffers[i], null, null); + clEnqueueAcquireGLObjects(queues[i], glBuffersOut[i], null, null); + + clEnqueueNDRangeKernel(queues[i], kernels[i], 3, + null, + kernel2DGlobalWorkSize, + null, + null, null); + + CL10GL.clEnqueueReleaseGLObjects(queues[i], glBuffers[i], null, syncGLtoCL ? syncBuffer : null); + CL10GL.clEnqueueReleaseGLObjects(queues[i], glBuffersOut[i], null, syncGLtoCL ? syncBuffer : null); + if (syncGLtoCL) { + clEvents[i] = queues[i].getCLEvent(syncBuffer.get(0)); + clSyncs[i] = glCreateSyncFromCLeventARB(queues[i].getParent(), clEvents[i], 0); + } + } + + // block until done (important: finish before doing further gl work) + if (!syncGLtoCL) { + for (int i = 0; i < slices; i++) { + clFinish(queues[i]); + } + } + + //RENDER THE TEXTURE + render(); + } + + private void render() { + if (syncGLtoCL) { + for (int i = 0; i < slices; i++) { + glWaitSync(clSyncs[i], 0, 0); + } + } + + //draw slices + //int sliceWidth = texture.getWidth() / slices; + + for (int i = 0; i < slices; i++) { + //int seperatorOffset = drawSeparator ? i : 0; + + //SETS GL SETTINGS + glSettings(); + + //BIND THE TEXTURE + glEnable(GL_TEXTURE_3D); + glBindTexture(GL_TEXTURE_3D, outTexture.getId()); + + //DRAW A CUBE WITH MAPPED TEXTURE + glBegin(GL_QUADS); + //System.out.println("Drawing front"); + glTexCoord3f(0f, 0f, 0.5f); + glVertex3f(100, 100, 0); //upper left + glTexCoord3f(0f, 1.0f, 0.5f); + glVertex3f(100, 200, 0); //upper right + glTexCoord3f(1.0f, 1.0f, 0.5f); + glVertex3f(200, 200, 0); //bottom right + glTexCoord3f(1.0f, 0f, 0.5f); + glVertex3f(200, 100, 0); //bottom left + +// // Top-face +// System.out.println("Drawing top"); +// glTexCoord3f(0f, 0f, 0f); +// glVertex3f(1.0f, 1.0f, -1.0f); +// glTexCoord3f(0f, 1.0f, 0f); +// glVertex3f(-1.0f, 1.0f, -1.0f); +// glTexCoord3f(1.0f, 1.0f, 0f); +// glVertex3f(-1.0f, 1.0f, 1.0f); +// glTexCoord3f(1.0f, 0f, 0f); +// glVertex3f(1.0f, 1.0f, 1.0f); +// +// // Bottom-face +// System.out.println("Drawing bottom"); +// glTexCoord3f(0f, 0f, 0f); +// glVertex3f(1.0f, -1.0f, 1.0f); +// glTexCoord3f(0f, 1.0f, 0f); +// glVertex3f(-1.0f, -1.0f, 1.0f); +// glTexCoord3f(1.0f, 1.0f, 0f); +// glVertex3f(-1.0f, -1.0f, -1.0f); +// glTexCoord3f(1.0f, 0f, 0f); +// glVertex3f(1.0f, -1.0f, -1.0f); +// +// // Front-face +// System.out.println("Drawing front"); +// glTexCoord3f(0f, 0f, 0f); +// glVertex3f(1.0f, 1.0f, 1.0f); +// glTexCoord3f(0f, 1.0f, 0f); +// glVertex3f(-1.0f, 1.0f, 1.0f); +// glTexCoord3f(1.0f, 1.0f, 0f); +// glVertex3f(-1.0f, -1.0f, 1.0f); +// glTexCoord3f(1.0f, 0f, 0f); +// glVertex3f(1.0f, -1.0f, 1.0f); +// +// // Back-face +// System.out.println("Drawing back"); +// glTexCoord3f(0f, 0f, 0f); +// glVertex3f(1.0f, -1.0f, -1.0f); +// glTexCoord3f(0f, 1.0f, 0f); +// glVertex3f(-1.0f, -1.0f, -1.0f); +// glTexCoord3f(1.0f, 1.0f, 0f); +// glVertex3f(-1.0f, 1.0f, -1.0f); +// glTexCoord3f(1.0f, 0f, 0f); +// glVertex3f(1.0f, 1.0f, -1.0f); +// +// // Left-face +// System.out.println("Drawing left"); +// glTexCoord3f(0f, 0f, 0f); +// glVertex3f(-1.0f, 1.0f, 1.0f); +// glTexCoord3f(0f, 1.0f, 0f); +// glVertex3f(-1.0f, 1.0f, -1.0f); +// glTexCoord3f(1.0f, 1.0f, 0f); +// glVertex3f(-1.0f, -1.0f, -1.0f); +// glTexCoord3f(1.0f, 0f, 0f); +// glVertex3f(-1.0f, -1.0f, 1.0f); +// +// // Right-face +// System.out.println("Drawing right"); +// glTexCoord3f(0f, 0f, 0f); +// glVertex3f(1.0f, 1.0f, -1.0f); +// glTexCoord3f(0f, 1.0f, 0f); +// glVertex3f(1.0f, 1.0f, 1.0f); +// glTexCoord3f(1.0f, 1.0f, 0f); +// glVertex3f(1.0f, -1.0f, 1.0f); +// glTexCoord3f(1.0f, 0f, 0f); +// glVertex3f(1.0f, -1.0f, -1.0f); + + glEnd(); + + glBindTexture(GL_TEXTURE_3D, 0); + glDisable(GL_TEXTURE_3D); + } + + //CHECKING SYNC + if (syncCLtoGL) { + glSync = glFenceSync(GL_SYNC_GPU_COMMANDS_COMPLETE, 0); + glEvent = clCreateEventFromGLsyncKHR(context, glSync, null); + } + } +} \ No newline at end of file diff --git a/src/tutorial/clTexture3d/Texture.java b/src/tutorial/clTexture3d/Texture.java new file mode 100644 index 0000000..f2d562d --- /dev/null +++ b/src/tutorial/clTexture3d/Texture.java @@ -0,0 +1,96 @@ +/* + * To change this license header, choose License Headers in Project Properties. + * To change this template file, choose Tools | Templates + * and open the template in the editor. + */ +package tutorial.clTexture3d; + +import java.nio.ByteBuffer; +import java.nio.ShortBuffer; +import static org.lwjgl.opengl.ARBTextureRg.GL_R16; +import static org.lwjgl.opengl.ARBTextureRg.GL_R16F; +import org.lwjgl.opengl.GL11; +import static org.lwjgl.opengl.GL11.GL_NEAREST; +import static org.lwjgl.opengl.GL11.GL_RED; +import static org.lwjgl.opengl.GL11.GL_RGB; +import static org.lwjgl.opengl.GL11.GL_RGB8; +import static org.lwjgl.opengl.GL11.GL_RGBA; +import static org.lwjgl.opengl.GL11.GL_SHORT; +import static org.lwjgl.opengl.GL11.GL_TEXTURE_2D; +import static org.lwjgl.opengl.GL11.GL_TEXTURE_MAG_FILTER; +import static org.lwjgl.opengl.GL11.GL_TEXTURE_MIN_FILTER; +import static org.lwjgl.opengl.GL11.GL_TEXTURE_WRAP_S; +import static org.lwjgl.opengl.GL11.GL_TEXTURE_WRAP_T; +import static org.lwjgl.opengl.GL11.GL_UNSIGNED_BYTE; +import static org.lwjgl.opengl.GL11.glTexImage2D; +import static org.lwjgl.opengl.GL11.glTexParameteri; +import static org.lwjgl.opengl.GL12.GL_CLAMP_TO_EDGE; +import static org.lwjgl.opengl.GL12.GL_TEXTURE_3D; +import static org.lwjgl.opengl.GL12.glTexImage3D; +import static org.lwjgl.opengl.GL30.GL_RGBA16F; + +/* + * @author LAA + */ +public class Texture { + private int width, height, depth, target, id; + + public Texture(int width, int height, int depth, int colorFormat, ByteBuffer buff) throws Exception { + this.width = width; + this.height = height; + this.depth = depth; + this.target = GL_TEXTURE_3D; + this.id = GL11.glGenTextures(); + + // bind this texture + GL11.glBindTexture(target, id); + + glTexParameteri(target, GL_TEXTURE_WRAP_S, GL_CLAMP_TO_EDGE); + glTexParameteri(target, GL_TEXTURE_WRAP_T, GL_CLAMP_TO_EDGE); + glTexParameteri(target, GL_TEXTURE_MIN_FILTER, GL_NEAREST); + glTexParameteri(target, GL_TEXTURE_MAG_FILTER, GL_NEAREST); + + // glTexImage2D(GL_TEXTURE_2D, 0, GL_RGB8, image.getWidth() + 1, image.getHeight(), 0, GL_RGB, GL_UNSIGNED_BYTE, image.getByeBuff()); + switch (colorFormat){ + case GL_RGB8: + //glTexImage3D(); + glTexImage3D(target, 0, colorFormat, width, height, depth, 0, GL_RGB, GL_UNSIGNED_BYTE, buff); + break; + case GL_RGBA16F: + glTexImage3D(target, 0, colorFormat, width, height, depth, 0, GL_RGBA, GL_UNSIGNED_BYTE, buff); + break; + case GL_R16F: + glTexImage3D(target, 0, colorFormat, width, height, depth, 0, GL_RED, GL_SHORT, buff); + break; + case GL_R16: + glTexImage3D(target, 0, colorFormat, width, height, depth, 0, GL_RED, GL_SHORT, buff); + break; + default: + throw new Exception("Unrecognized texture color format: " + colorFormat); + } + } + + public void bind() { + GL11.glBindTexture(target, id); + } + + public int getTarget() { + return target; + } + + public int getId() { + return id; + } + + public int getWidth() { + return width; + } + + public int getHeight() { + return height; + } + + public int getDepth(){ + return depth; + } +}