diff --git a/jme3-core/src/main/java/com/jme3/opencl/Context.java b/jme3-core/src/main/java/com/jme3/opencl/Context.java index 33d8a6843..fed3cd2a3 100644 --- a/jme3-core/src/main/java/com/jme3/opencl/Context.java +++ b/jme3-core/src/main/java/com/jme3/opencl/Context.java @@ -34,6 +34,7 @@ package com.jme3.opencl; import com.jme3.asset.AssetInfo; import com.jme3.asset.AssetKey; import com.jme3.asset.AssetManager; +import com.jme3.asset.AssetNotFoundException; import com.jme3.opencl.Image.ImageDescriptor; import com.jme3.opencl.Image.ImageFormat; import com.jme3.opencl.Image.ImageType; @@ -43,6 +44,7 @@ import com.jme3.texture.Texture; import java.io.BufferedReader; import java.io.IOException; import java.io.InputStreamReader; +import java.io.StringReader; import java.nio.ByteBuffer; import java.util.Arrays; import java.util.List; @@ -282,6 +284,52 @@ public abstract class Context extends AbstractOpenCLObject { */ public abstract Program createProgramFromSourceCode(String sourceCode); + /** + * Resolves dependencies (using {@code #include } in the source code) + * and delegates the combined source code to + * {@link #createProgramFromSourceCode(java.lang.String) }. + * Important: only absolute paths are allowed. + * @param sourceCode the original source code + * @param assetManager the asset manager to load the files + * @return the created program object + * @throws AssetNotFoundException if a dependency could not be loaded + */ + public Program createProgramFromSourceCodeWithDependencies(String sourceCode, AssetManager assetManager) { + StringBuilder builder = new StringBuilder(sourceCode.length()); + BufferedReader reader = new BufferedReader(new StringReader(sourceCode)); + try { + buildSourcesRec(reader, builder, assetManager); + } catch (IOException ex) { + throw new AssetNotFoundException("Unable to read a dependency file", ex); + } + return createProgramFromSourceCode(builder.toString()); + } + private void buildSourcesRec(BufferedReader reader, StringBuilder builder, AssetManager assetManager) throws IOException { + String ln; + while ((ln = reader.readLine()) != null) { + if (ln.trim().startsWith("#import ")) { + ln = ln.trim().substring(8).trim(); + if (ln.startsWith("\"")) { + ln = ln.substring(1); + } + if (ln.endsWith("\"")) { + ln = ln.substring(0, ln.length()-1); + } + AssetInfo info = assetManager.locateAsset(new AssetKey(ln)); + if (info == null) { + throw new AssetNotFoundException("Unable to load source file \""+ln+"\""); + } + try (BufferedReader r = new BufferedReader(new InputStreamReader(info.openStream()))) { + builder.append("//-- begin import ").append(ln).append(" --\n"); + buildSourcesRec(r, builder, assetManager); + builder.append("//-- end import ").append(ln).append(" --\n"); + } + } else { + builder.append(ln).append('\n'); + } + } + } + /** * Creates a program object from the provided source code and files. * The source code is made up from the specified include string first, @@ -294,14 +342,15 @@ public abstract class Context extends AbstractOpenCLObject { *
  • Some common OpenCL files used as libraries (Convention: file names end with {@code .clh}
    • *
  • One main OpenCL file containing the actual kernels (Convention: file name ends with {@code .cl})
    • * - * Note: Files that can't be loaded are skipped.
    * - * The actual creation is handled by {@link #createProgramFromSourceCode(java.lang.String) }. + * After the files were combined, additional include statements are resolved + * by {@link #createProgramFromSourceCodeWithDependencies(java.lang.String, com.jme3.asset.AssetManager) }. * * @param assetManager the asset manager used to load the files * @param include an additional include string * @param resources an array of asset paths pointing to OpenCL source files * @return the new program objects + * @throws AssetNotFoundException if a file could not be loaded */ public Program createProgramFromSourceFilesWithInclude(AssetManager assetManager, String include, String... resources) { return createProgramFromSourceFilesWithInclude(assetManager, include, Arrays.asList(resources)); @@ -319,14 +368,15 @@ public abstract class Context extends AbstractOpenCLObject { *
  • Some common OpenCL files used as libraries (Convention: file names end with {@code .clh}
    • *
  • One main OpenCL file containing the actual kernels (Convention: file name ends with {@code .cl})
    • * - * Note: Files that can't be loaded are skipped.
    * - * The actual creation is handled by {@link #createProgramFromSourceCode(java.lang.String) }. + * After the files were combined, additional include statements are resolved + * by {@link #createProgramFromSourceCodeWithDependencies(java.lang.String, com.jme3.asset.AssetManager) }. * * @param assetManager the asset manager used to load the files * @param include an additional include string * @param resources an array of asset paths pointing to OpenCL source files * @return the new program objects + * @throws AssetNotFoundException if a file could not be loaded */ public Program createProgramFromSourceFilesWithInclude(AssetManager assetManager, String include, List resources) { StringBuilder str = new StringBuilder(); @@ -334,8 +384,7 @@ public abstract class Context extends AbstractOpenCLObject { for (String res : resources) { AssetInfo info = assetManager.locateAsset(new AssetKey(res)); if (info == null) { - LOG.log(Level.WARNING, "unable to load source file ''{0}''", res); - continue; + throw new AssetNotFoundException("Unable to load source file \""+res+"\""); } try (BufferedReader reader = new BufferedReader(new InputStreamReader(info.openStream()))) { while (true) { @@ -349,12 +398,13 @@ public abstract class Context extends AbstractOpenCLObject { LOG.log(Level.WARNING, "unable to load source file '"+res+"'", ex); } } - return createProgramFromSourceCode(str.toString()); + return createProgramFromSourceCodeWithDependencies(str.toString(), assetManager); } /** * Alternative version of {@link #createProgramFromSourceFilesWithInclude(com.jme3.asset.AssetManager, java.lang.String, java.lang.String...) } * with an empty include string + * @throws AssetNotFoundException if a file could not be loaded */ public Program createProgramFromSourceFiles(AssetManager assetManager, String... resources) { return createProgramFromSourceFilesWithInclude(assetManager, "", resources); @@ -363,6 +413,7 @@ public abstract class Context extends AbstractOpenCLObject { /** * Alternative version of {@link #createProgramFromSourceFilesWithInclude(com.jme3.asset.AssetManager, java.lang.String, java.util.List) } * with an empty include string + * @throws AssetNotFoundException if a file could not be loaded */ public Program createProgramFromSourceFiles(AssetManager assetManager, List resources) { return createProgramFromSourceFilesWithInclude(assetManager, "", resources); diff --git a/jme3-core/src/main/resources/Common/OpenCL/Random.clh b/jme3-core/src/main/resources/Common/OpenCL/Random.clh new file mode 100644 index 000000000..18127e153 --- /dev/null +++ b/jme3-core/src/main/resources/Common/OpenCL/Random.clh @@ -0,0 +1,185 @@ +//This is a port of java.util.Random to OpenCL + +//Because not all devices support doubles, the double returning functions +//must be explicit activated with the following preprocessor macro: +//#define RANDOM_DOUBLES + +#ifdef RANDOM_DOUBLES +#ifdef cl_khr_fp64 +#pragma OPENCL EXTENSION cl_khr_fp64 : enable +#elif defined(cl_amd_fp64) +#pragma OPENCL EXTENSION cl_amd_fp64 : enable +#else +#error "Double precision floating point not supported by OpenCL implementation." +#endif +#endif + +inline int randNext(int bits, __global ulong* seed) +{ + *seed = (*seed * 0x5DEECE66DL + 0xBL) & ((1L << 48) - 1); + return (int)(*seed >> (48 - bits)); +} + +/** + * Retrieves the next random integer value. + * The buffer used as seed must be read-write. + * Usage: + * + * __kernel void TestRandom(__global ulong* seeds) { + * // ... + * int i = randInt(seeds + get_global_id(0)); + * // --- + * } + * + */ +inline int randInt(__global ulong* seed) { + return randNext(32, seed); +} + +/** + * Retrieves the next random integer value between 0 (inclusive) and n (exclusive). + * The buffer used as seed must be read-write. + * Usage: + * + * __kernel void TestRandom(__global ulong* seeds) { + * // ... + * int i = randIntN(n, seeds + get_global_id(0)); + * // --- + * } + * + */ +inline int randIntN(int n, __global ulong* seed) { + if (n <= 0) + return 0; + + if ((n & -n) == n) // i.e., n is a power of 2 + return (int)((n * (long)randNext(31, seed)) >> 31); + + int bits, val; + do { + bits = randNext(31, seed); + val = bits % n; + } while (bits - val + (n-1) < 0); + return val; +} + +/** + * Retrieves the next random long value. + * The buffer used as seed must be read-write. + * Usage: + * + * __kernel void TestRandom(__global ulong* seeds) { + * // ... + * long l = randLong(seeds + get_global_id(0)); + * // --- + * } + * + */ +inline long randLong(__global ulong* seed) { + // it's okay that the bottom word remains signed. + return ((long)(randNext(32, seed)) << 32) + randNext(32, seed); +} + +/** + * Retrieves the next random boolean value. + * The buffer used as seed must be read-write. + * Usage: + * + * __kernel void TestRandom(__global ulong* seeds) { + * // ... + * bool b = randBool(seeds + get_global_id(0)); + * // --- + * } + * + */ +inline bool randBool(__global ulong* seed) { + return randNext(1, seed) != 0; +} + +#ifdef RANDOM_DOUBLES +/** + * Retrieves the next random double value. + * The buffer used as seed must be read-write. + * To use this function, the preprocessor define RANDOM_DOUBLES must be set. + * Usage: + * + * __kernel void TestRandom(__global ulong* seeds) { + * // ... + * double d = randDouble(seeds + get_global_id(0)); + * // --- + * } + * + */ +inline double randDouble(__global ulong* seed) { + return (((long)(randNext(26, seed)) << 27) + randNext(27, seed)) + / (double)(1L << 53); +} +#endif + +/** + * Retrieves the next random float value. + * The buffer used as seed must be read-write. + * Usage: + * + * __kernel void TestRandom(__global ulong* seeds) { + * // ... + * float f = randFloat(seeds + get_global_id(0)); + * // --- + * } + * + */ +inline float randFloat(__global ulong* seed) +{ + return randNext(24, seed) / ((float)(1 << 24)); +} + +/** + * Retrieves the next random float values with a gaussian distribution of mean 0 + * and derivation 1. + * The buffer used as seed must be read-write. + * Usage: + * + * __kernel void TestRandom(__global ulong* seeds) { + * // ... + * float2 f2 = randGausianf(seeds + get_global_id(0)); + * // --- + * } + * + */ +inline float2 randGaussianf(__global ulong* seed) { + float v1, v2, s; + do { + v1 = 2 * randFloat(seed) - 1; // between -1 and 1 + v2 = 2 * randFloat(seed) - 1; // between -1 and 1 + s = v1 * v1 + v2 * v2; + } while (s >= 1 || s == 0); + float multiplier = sqrt(-2 * log(s)/s); + return (float2) (v1 * multiplier, v2 * multiplier); +} + +#ifdef RANDOM_DOUBLES +/** + * Retrieves the next random double values with a gaussian distribution of mean 0 + * and derivation 1. + * The buffer used as seed must be read-write. + * To use this function, the preprocessor define RANDOM_DOUBLES must be set. + * Usage: + * + * __kernel void TestRandom(__global ulong* seeds) { + * // ... + * double2 f2 = randGausian(seeds + get_global_id(0)); + * // --- + * } + * + */ +inline double2 randGaussian(__global ulong* seed) { + double v1, v2, s; + do { + v1 = 2 * randDouble(seed) - 1; // between -1 and 1 + v2 = 2 * randDouble(seed) - 1; // between -1 and 1 + s = v1 * v1 + v2 * v2; + } while (s >= 1 || s == 0); + double multiplier = sqrt(-2 * log(s)/s); + return (double2) (v1 * multiplier, v2 * multiplier); +} +#endif \ No newline at end of file diff --git a/jme3-examples/src/main/java/jme3test/opencl/TestOpenCLLibraries.java b/jme3-examples/src/main/java/jme3test/opencl/TestOpenCLLibraries.java new file mode 100644 index 000000000..9f27203cf --- /dev/null +++ b/jme3-examples/src/main/java/jme3test/opencl/TestOpenCLLibraries.java @@ -0,0 +1,257 @@ +/* + * Copyright (c) 2009-2012 jMonkeyEngine + * All rights reserved. + * + * Redistribution and use in source and binary forms, with or without + * modification, are permitted provided that the following conditions are + * met: + * + * * Redistributions of source code must retain the above copyright + * notice, this list of conditions and the following disclaimer. + * + * * Redistributions in binary form must reproduce the above copyright + * notice, this list of conditions and the following disclaimer in the + * documentation and/or other materials provided with the distribution. + * + * * Neither the name of 'jMonkeyEngine' nor the names of its contributors + * may be used to endorse or promote products derived from this software + * without specific prior written permission. + * + * THIS SOFTWARE IS PROVIDED BY THE COPYRIGHT HOLDERS AND CONTRIBUTORS + * "AS IS" AND ANY EXPRESS OR IMPLIED WARRANTIES, INCLUDING, BUT NOT LIMITED + * TO, THE IMPLIED WARRANTIES OF MERCHANTABILITY AND FITNESS FOR A PARTICULAR + * PURPOSE ARE DISCLAIMED. IN NO EVENT SHALL THE COPYRIGHT OWNER OR + * CONTRIBUTORS BE LIABLE FOR ANY DIRECT, INDIRECT, INCIDENTAL, SPECIAL, + * EXEMPLARY, OR CONSEQUENTIAL DAMAGES (INCLUDING, BUT NOT LIMITED TO, + * PROCUREMENT OF SUBSTITUTE GOODS OR SERVICES; LOSS OF USE, DATA, OR + * PROFITS; OR BUSINESS INTERRUPTION) HOWEVER CAUSED AND ON ANY THEORY OF + * LIABILITY, WHETHER IN CONTRACT, STRICT LIABILITY, OR TORT (INCLUDING + * NEGLIGENCE OR OTHERWISE) ARISING IN ANY WAY OUT OF THE USE OF THIS + * SOFTWARE, EVEN IF ADVISED OF THE POSSIBILITY OF SUCH DAMAGE. + */ + +package jme3test.opencl; + +import com.jme3.app.SimpleApplication; +import com.jme3.font.BitmapFont; +import com.jme3.font.BitmapText; +import com.jme3.math.ColorRGBA; +import com.jme3.opencl.*; +import com.jme3.system.AppSettings; +import com.jme3.util.BufferUtils; +import java.nio.*; +import java.util.Arrays; +import java.util.Objects; +import java.util.Random; +import java.util.logging.Level; +import java.util.logging.Logger; + +/** + * Test class for the build in libraries + * @author shaman + */ +public class TestOpenCLLibraries extends SimpleApplication { + private static final Logger LOG = Logger.getLogger(TestOpenCLLibraries.class.getName()); + + public static void main(String[] args){ + TestOpenCLLibraries app = new TestOpenCLLibraries(); + AppSettings settings = new AppSettings(true); + settings.setOpenCLSupport(true); + settings.setVSync(true); +// settings.setRenderer(AppSettings.JOGL_OPENGL_FORWARD_COMPATIBLE); + app.setSettings(settings); + app.start(); // start the game + } + + @Override + public void simpleInitApp() { + BitmapFont fnt = assetManager.loadFont("Interface/Fonts/Default.fnt"); + Context clContext = context.getOpenCLContext(); + if (clContext == null) { + BitmapText txt = new BitmapText(fnt); + txt.setText("No OpenCL Context created!\nSee output log for details."); + txt.setLocalTranslation(5, settings.getHeight() - 5, 0); + guiNode.attachChild(txt); + return; + } + CommandQueue clQueue = clContext.createQueue(clContext.getDevices().get(0)); + + StringBuilder str = new StringBuilder(); + str.append("OpenCL Context created:\n Platform: ") + .append(clContext.getDevices().get(0).getPlatform().getName()) + .append("\n Devices: ").append(clContext.getDevices()); + str.append("\nTests:"); + str.append("\n Random numbers: ").append(testRandom(clContext, clQueue)); + + clQueue.release(); + + BitmapText txt1 = new BitmapText(fnt); + txt1.setText(str.toString()); + txt1.setLocalTranslation(5, settings.getHeight() - 5, 0); + guiNode.attachChild(txt1); + + flyCam.setEnabled(false); + inputManager.setCursorVisible(true); + } + + private static void assertEquals(byte expected, byte actual, String message) { + if (expected != actual) { + System.err.println(message+": expected="+expected+", actual="+actual); + throw new AssertionError(); + } + } + private static void assertEquals(long expected, long actual, String message) { + if (expected != actual) { + System.err.println(message+": expected="+expected+", actual="+actual); + throw new AssertionError(); + } + } + private static void assertEquals(double expected, double actual, String message) { + if (Math.abs(expected - actual) >= 0.00001) { + System.err.println(message+": expected="+expected+", actual="+actual); + throw new AssertionError(); + } + } + private static void assertEquals(Object expected, Object actual, String message) { + if (!Objects.equals(expected, actual)) { + System.err.println(message+": expected="+expected+", actual="+actual); + throw new AssertionError(); + } + } + + private boolean testRandom(Context clContext, CommandQueue clQueue) { + try { + //test for doubles + boolean supportsDoubles = clContext.getDevices().get(0).hasDouble(); + + //create code + String code = "" + + "#import \"Common/OpenCL/Random.clh\"\n" + + "__kernel void TestBool(__global ulong* seeds, __global uchar* results) {\n" + + " results[get_global_id(0)] = randBool(seeds + get_global_id(0)) ? 1 : 0;\n" + + "}\n" + + "__kernel void TestInt(__global ulong* seeds, __global int* results) {\n" + + " results[get_global_id(0)] = randInt(seeds + get_global_id(0));\n" + + "}\n" + + "__kernel void TestIntN(__global ulong* seeds, int n, __global int* results) {\n" + + " results[get_global_id(0)] = randIntN(n, seeds + get_global_id(0));\n" + + "}\n" + + "__kernel void TestLong(__global ulong* seeds, __global long* results) {\n" + + " results[get_global_id(0)] = randLong(seeds + get_global_id(0));\n" + + "}\n" + + "__kernel void TestFloat(__global ulong* seeds, __global float* results) {\n" + + " results[get_global_id(0)] = randFloat(seeds + get_global_id(0));\n" + + "}\n" + + "#ifdef RANDOM_DOUBLES\n" + + "__kernel void TestDouble(__global ulong* seeds, __global double* results) {\n" + + " results[get_global_id(0)] = randDouble(seeds + get_global_id(0));\n" + + "}\n" + + "#endif\n"; + if (supportsDoubles) { + code = "#define RANDOM_DOUBLES\n" + code; + } + Program program = clContext.createProgramFromSourceCodeWithDependencies(code, assetManager); + program.build(); + + int count = 256; + Kernel.WorkSize ws = new Kernel.WorkSize(count); + + //create seeds + Random initRandom = new Random(); + long[] seeds = new long[count]; + Random[] randoms = new Random[count]; + for (int i=0; i