added dependency resolving and a random number library

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<String>(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 {
      *  <li>Some common OpenCL files used as libraries (Convention: file names end with {@code .clh}</li>
      *  <li>One main OpenCL file containing the actual kernels (Convention: file name ends with {@code .cl})</li>
      * </ul>
-     * Note: Files that can't be loaded are skipped.<br>
      * 
-     * 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 {
      *  <li>Some common OpenCL files used as libraries (Convention: file names end with {@code .clh}</li>
      *  <li>One main OpenCL file containing the actual kernels (Convention: file name ends with {@code .cl})</li>
      * </ul>
-     * Note: Files that can't be loaded are skipped.<br>
      * 
-     * 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<String> 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<String>(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<String> resources) {
         return createProgramFromSourceFilesWithInclude(assetManager, "", resources);

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: 
+ * <code>
+ *  __kernel void TestRandom(__global ulong* seeds) {
+ *    // ...
+ *    int i = randInt(seeds + get_global_id(0));
+ *    // ---
+ * }
+ * </code>
+ */
+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: 
+ * <code>
+ *  __kernel void TestRandom(__global ulong* seeds) {
+ *    // ...
+ *    int i = randIntN(n, seeds + get_global_id(0));
+ *    // ---
+ * }
+ * </code>
+ */
+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: 
+ * <code>
+ *  __kernel void TestRandom(__global ulong* seeds) {
+ *    // ...
+ *    long l = randLong(seeds + get_global_id(0));
+ *    // ---
+ * }
+ * </code>
+ */
+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: 
+ * <code>
+ *  __kernel void TestRandom(__global ulong* seeds) {
+ *    // ...
+ *    bool b = randBool(seeds + get_global_id(0));
+ *    // ---
+ * }
+ * </code>
+ */
+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: 
+ * <code>
+ *  __kernel void TestRandom(__global ulong* seeds) {
+ *    // ...
+ *    double d = randDouble(seeds + get_global_id(0));
+ *    // ---
+ * }
+ * </code>
+ */
+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: 
+ * <code>
+ *  __kernel void TestRandom(__global ulong* seeds) {
+ *    // ...
+ *    float f = randFloat(seeds + get_global_id(0));
+ *    // ---
+ * }
+ * </code>
+ */
+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: 
+ * <code>
+ *  __kernel void TestRandom(__global ulong* seeds) {
+ *    // ...
+ *    float2 f2 = randGausianf(seeds + get_global_id(0));
+ *    // ---
+ * }
+ * </code>
+ */
+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: 
+ * <code>
+ *  __kernel void TestRandom(__global ulong* seeds) {
+ *    // ...
+ *    double2 f2 = randGausian(seeds + get_global_id(0));
+ *    // ---
+ * }
+ * </code>
+ */
+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

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<count; ++i) {
+                seeds[i] = initRandom.nextLong();
+                randoms[i] = new Random(seeds[i]);
+                seeds[i] = (seeds[i] ^ 0x5DEECE66DL) & ((1L << 48) - 1); //needed because the Random constructor scrambles the initial seed
+            }
+            com.jme3.opencl.Buffer seedsBuffer = clContext.createBuffer(8 * count);
+            ByteBuffer tmpByteBuffer = BufferUtils.createByteBuffer(8 * count);
+            tmpByteBuffer.asLongBuffer().put(seeds);
+            seedsBuffer.write(clQueue, tmpByteBuffer);
+            
+            //test it
+            ByteBuffer resultByteBuffer = BufferUtils.createByteBuffer(8 * count);
+            IntBuffer resultIntBuffer = resultByteBuffer.asIntBuffer();
+            LongBuffer resultLongBuffer = resultByteBuffer.asLongBuffer();
+            FloatBuffer resultFloatBuffer = resultByteBuffer.asFloatBuffer();
+            DoubleBuffer resultDoubleBuffer = resultByteBuffer.asDoubleBuffer();
+            com.jme3.opencl.Buffer resultBuffer = clContext.createBuffer(8 * count);
+            //boolean
+            Kernel testBoolKernel = program.createKernel("TestBool");
+            testBoolKernel.Run1NoEvent(clQueue, ws, seedsBuffer, resultBuffer);
+            resultByteBuffer.rewind();
+            resultBuffer.read(clQueue, resultByteBuffer);
+            for (int i=0; i<count; ++i) {
+                assertEquals(randoms[i].nextBoolean() ? 1 : 0, resultByteBuffer.get(i), "randBool at i="+i);
+            }
+            testBoolKernel.release();
+            //int
+            Kernel testIntKernel = program.createKernel("TestInt");
+            testIntKernel.Run1NoEvent(clQueue, ws, seedsBuffer, resultBuffer);
+            resultByteBuffer.rewind();
+            resultBuffer.read(clQueue, resultByteBuffer);
+            for (int i=0; i<count; ++i) {
+                assertEquals(randoms[i].nextInt(), resultIntBuffer.get(i), "randInt at i="+i);
+            }
+            testIntKernel.release();
+            //int n
+            Kernel testIntNKernel = program.createKernel("TestIntN");
+            testIntNKernel.Run1NoEvent(clQueue, ws, seedsBuffer, 186, resultBuffer);
+            resultByteBuffer.rewind();
+            resultBuffer.read(clQueue, resultByteBuffer);
+            for (int i=0; i<count; ++i) {
+                assertEquals(randoms[i].nextInt(186), resultIntBuffer.get(i), "randInt at i="+i+" with n="+186);
+            }
+            testIntNKernel.Run1NoEvent(clQueue, ws, seedsBuffer, 97357, resultBuffer);
+            resultByteBuffer.rewind();
+            resultBuffer.read(clQueue, resultByteBuffer);
+            for (int i=0; i<count; ++i) {
+                assertEquals(randoms[i].nextInt(97357), resultIntBuffer.get(i), "randInt at i="+i+" with n="+97357);
+            }
+            testIntNKernel.release();
+            //long
+            Kernel testLongKernel = program.createKernel("TestLong");
+            testLongKernel.Run1NoEvent(clQueue, ws, seedsBuffer, resultBuffer);
+            resultByteBuffer.rewind();
+            resultBuffer.read(clQueue, resultByteBuffer);
+            for (int i=0; i<count; ++i) {
+                assertEquals(randoms[i].nextLong(), resultLongBuffer.get(i), "randLong at i="+i);
+            }
+            testLongKernel.release();
+            //float
+            Kernel testFloatKernel = program.createKernel("TestFloat");
+            testFloatKernel.Run1NoEvent(clQueue, ws, seedsBuffer, resultBuffer);
+            resultByteBuffer.rewind();
+            resultBuffer.read(clQueue, resultByteBuffer);
+            for (int i=0; i<count; ++i) {
+                assertEquals(randoms[i].nextFloat(), resultFloatBuffer.get(i), "randFloat at i="+i);
+            }
+            testFloatKernel.release();
+            //double
+            if (supportsDoubles) {
+                Kernel testDoubleKernel = program.createKernel("TestDouble");
+                testDoubleKernel.Run1NoEvent(clQueue, ws, seedsBuffer, resultBuffer);
+                resultByteBuffer.rewind();
+                resultBuffer.read(clQueue, resultByteBuffer);
+                for (int i=0; i<count; ++i) {
+                    assertEquals(randoms[i].nextDouble(), resultDoubleBuffer.get(i), "randLong at i="+i);
+                }
+                testDoubleKernel.release();
+            }
+            
+            seedsBuffer.release();
+            resultBuffer.release();
+            program.release();
+
+        } catch (AssertionError ex) {
+            LOG.log(Level.SEVERE, "kernel test failed with an assertion error");
+            return false;
+        } catch (Exception ex) {
+            LOG.log(Level.SEVERE, "kernel test failed with:", ex);
+            return false;
+        }
+        return true;
+    }
+    
+}