Merge remote-tracking branch 'wwalker/master'

6 files changed, 161 insertions, 38 deletions

diff --git a/.classpath b/.classpath
index 6f661528..deaae283 100644
--- a/.classpath
+++ b/.classpath
@@ -15,6 +15,6 @@
 	<classpathentry combineaccessrules="false" kind="src" path="/gluegen"/>
 	<classpathentry combineaccessrules="false" kind="src" path="/jogl"/>
 	<classpathentry kind="lib" path="/gluegen/make/lib/android-sdk/15/android.jar" sourcepath="/gluegen/make/lib/android-sdk/15/android-java-src.zip"/>
-	<classpathentry kind="con" path="org.eclipse.jdt.junit.JUNIT_CONTAINER/4"/>
+	<classpathentry kind="lib" path="/gluegen/make/lib/junit.jar" sourcepath="/gluegen/make/lib/junit-sources.jar"/>
 	<classpathentry kind="output" path="build/classes"/>
 </classpath>
diff --git a/doc/HowToBuild.html b/doc/HowToBuild.html
index 2ae6f4cc..927ef72e 100644
--- a/doc/HowToBuild.html
+++ b/doc/HowToBuild.html
@@ -39,14 +39,15 @@
                     </ul>
                     <h3>External References</h3>
                     <ul>
-                        <li><a href="http://www.khronos.org/registry/cl/sdk/1.1/docs/man/xhtml/">OpenCL Reference Pages</a></li>
-                        <li><a href="http://www.khronos.org/registry/cl/specs/opencl-1.1.pdf">OpenCL 1.1 Specification</a></li>
-                        <li><a href="http://www.khronos.org/files/opencl-quick-reference-card.pdf">OpenCL Quick Reference</a></li>
+                        <li><a href="http://www.khronos.org/opencl/">OpenCL Home</a></li>
+                        <li><a href="http://www.khronos.org/registry/cl/sdk/1.2/docs/man/xhtml/">OpenCL 1.2 Reference Pages</a></li>
+                        <li><a href="http://www.khronos.org/registry/cl/specs/opencl-1.2.pdf">OpenCL 1.2 Specification</a></li>
+                        <li><a href="https://www.khronos.org/files/opencl-1-2-quick-reference-card.pdf">OpenCL 1.2 Quick Reference Card</a></li>
                     </ul>
                     <h3>External Tutorials</h3>
                     <ul>
-                        <li><a href="http://developer.amd.com/zones/OpenCLZone/universities/pages/default.aspx">AMD OpenCL University Kit</a></li>
-                        <li><a href="http://developer.amd.com/documentation/videos/OpenCLTechnicalOverviewVideoSeries/Pages/default.aspx">AMD Stream OpenCL Technical Overview Video Series</a></li>
+                        <li><a href="http://developer.amd.com/partners/university-programs/">AMD OpenCL University Kit</a></li>
+                        <li><a href="http://developer.amd.com/resources/documentation-articles/videos/ati-stream-opencl-technical-overview-video-series/">OpenCL Technical Overview Video Series</a></li>
                         <li><a href="http://www.macresearch.org/opencl">Mac Research OpenCL Tutorials</a></li>
                     </ul>
                     <a href="http://www.khronos.org/opencl/"><img src="http://www.khronos.org/opencl/images/opencl_100px.png" alt="OpenCL"/></a>
@@ -56,31 +57,30 @@
                     <h2>Prerequisites</h2>
                     <hr/>
                     <p>
-                        <b>Follow all steps</b> described in <a href="../../jogl/doc/HowToBuild.html">How to build JOGL</a>,
-                        since GlueGen and JOGL are required as compiletime dependencies.
+                        Follow all steps described in <a href="../../jogl/doc/HowToBuild.html">How to build JOGL</a>,
+                        since GlueGen and JOGL are required as compile-time dependencies.
                     </p>
 
                     <h2>Platform and Component Requirements</h2>
                     <hr/>
-                    To verify the build you will have to install a OpenCL implementation on your platform.
+                    To run the tests included with JOCL, you will have to install an OpenCL implementation on your platform.
                     <p>
-                        Some vendors ship OpenCL already with the graphics driver or operating system.
-                        So please make sure your system is up2date.
+                        Some vendors include OpenCL with the graphics driver or operating system.
+                        So please make sure your system is up to date.
                     </p>
                     OpenCL SDKs for the desktop:
                     <ul>
-                        <li> <b>GPU NVidia</b> Geforce >= 8<br/>
+                        <li> <b>NVidia GPU</b> Geforce >= 8</b>
                             <ul>
                                 <li> NVidia <a href="http://developer.nvidia.com/cuda-downloads">GPU Computing SDK</a></li>
                             </ul>
                         </li>
-                        <li> <b>GPU AMD or</b></li>
-                        <li> <b>CPU x86 x86_64 SSE3</b>
+                        <li> <b>AMD GPU or CPU</b>
                             <ul>
-                                <li> AMD <a href="http://developer.amd.com/gpu/AMDAPPSDK/Pages/default.aspx">Accelerated Parallel Processing SDK</a></li>
+                                <li> AMD <a href="http://developer.amd.com/tools-and-sdks/heterogeneous-computing/amd-accelerated-parallel-processing-app-sdk/">Accelerated Parallel Processing SDK</a></li>
                             </ul>
                         </li>
-                        <li> <b>CPU Intel</b>
+                        <li> <b>Intel CPU</b>
                             <ul>
                                 <li>Intel <a href="http://software.intel.com/en-us/articles/opencl-sdk/">OpenCL SDK</a></li>
                             </ul>
@@ -90,16 +90,13 @@
                     <h2>Build Steps</h2>
                     <hr/>
                     <p>
-                        Here are the steps that are required in order to build JOCL.
+                        Here are the steps required to build JOCL.
                     </p>
 
                     <ol>
-                        <li><b>Optain the source code</b> using git:
-                            <ul>
-                                <li><a href="http://jogamp.org/git/?p=jocl.git">JOCL Dev GIT Repo</a></li>
-                            </ul>
-
-                            It is important that you checkout the source project side by side to its dependencies:
+                        <li><b>Check out and build Gluegen and JOGL</b> as described <a href="../../jogl/doc/HowToBuild.html">here.</a></li>
+                        <li><b>Get the JOCL source code</b> from the <a href="http://jogamp.org/git/?p=jocl.git">JOCL Git repository</a>.
+                            It is important that you check out the source code side by side with its dependencies:
                             <pre>
     /home/dude/projects/jogamp> git clone git://jogamp.org/srv/scm/jocl.git jocl
                             </pre>
@@ -110,23 +107,17 @@
     /home/dude/projects/jogamp/jogl
     /home/dude/projects/jogamp/jocl
                             </pre>
-                            you can now open the projects with NetBeans or proceed with:
                         </li>
 
                         <li> <b>Build the source tree:</b> <br/>
-                            Open a command shell in the "jocl" directory of the source tree and type <code>"ant jar"</code>.
+                            Open a command shell in the "jocl/make" directory of the source tree and type
+                            <code>"ant clean"</code>, then <code>"ant"</code>.
                         </li>
-                        <li> <b>Test your build:</b> <br/> Stay in your command shell in the "jocl" directory of the source tree and type <code>"ant test"</code>.</li>
-                        <li> <b>Build Javadoc:</b> <br/> Stay in your command shell in the "jocl" directory of the source tree and type "<code>ant javadoc"</code>.
+                        <li> <b>Test your build:</b> <br/> Stay in your command shell in the "jocl/make" directory of the source tree and type <code>"ant junit.run"</code>.</li>
+                        <li> <b>Build Javadoc:</b> <br/> Stay in your command shell in the "jocl/make" directory of the source tree and type "<code>ant javadoc"</code>.
                             This will produce the end-user documentation for JOCL.
                         </li>
                     </ol>
-
-                    <h2> Common build problems </h2>
-
-                    <ul>
-                        <li>TODO</li>
-                    </ul>
                 </div>
             </div>
             <div id="footer">
diff --git a/src/com/jogamp/opencl/util/CLDeviceFilters.java b/src/com/jogamp/opencl/util/CLDeviceFilters.java
index a5057fb6..6281b844 100644
--- a/src/com/jogamp/opencl/util/CLDeviceFilters.java
+++ b/src/com/jogamp/opencl/util/CLDeviceFilters.java
@@ -70,7 +70,7 @@ public class CLDeviceFilters {
     }
     
     /**
-     * Accepts all devices which support OpenGL-OpenCL interoparability.
+     * Accepts all devices which support OpenGL-OpenCL interoperability.
      */
     public static Filter<CLDevice> glSharing() {
         return new Filter<CLDevice>() {
diff --git a/src/com/jogamp/opencl/util/CLPlatformFilters.java b/src/com/jogamp/opencl/util/CLPlatformFilters.java
index 48d20916..451c6b1f 100644
--- a/src/com/jogamp/opencl/util/CLPlatformFilters.java
+++ b/src/com/jogamp/opencl/util/CLPlatformFilters.java
@@ -67,7 +67,7 @@ public class CLPlatformFilters {
     }
     
     /**
-     * Accepts all platforms containing at least one devices of which supports OpenGL-OpenCL interoparability.
+     * Accepts all platforms containing at least one devices of which supports OpenGL-OpenCL interoperability.
      */
     public static Filter<CLPlatform> glSharing() {
         return new Filter<CLPlatform>() {
@@ -86,7 +86,7 @@ public class CLPlatformFilters {
 
     /**
      * Accepts all with the given OpenGL context compatible platforms containing at least one
-     * devices of which supports OpenGL-OpenCL interoparability.
+     * devices of which supports OpenGL-OpenCL interoperability.
      */
     public static Filter<CLPlatform> glSharing(final GLContext context) {
         return new Filter<CLPlatform>() {
@@ -94,10 +94,24 @@ public class CLPlatformFilters {
             public boolean accept(CLPlatform item) {
                 String glVendor = context.getGL().glGetString(GL.GL_VENDOR);
                 String clVendor = item.getVendor();
-                return clVendor.equals(glVendor) && glFilter.accept(item);
+                return areVendorsCompatible(glVendor,clVendor) && glFilter.accept(item);
             }
         };
     }
+    
+    /**
+     * We need this test because:
+     *  - On at least some AMD cards, the GL vendor is ATI, but the CL vendor is AMD.
+     *  - On at least some Macs, the GL vendor is Nvidia, but the CL vendor is Apple.
+     * @param glVendor OpenGL vendor string.
+     * @param clVendor OpenCL vendor string.
+     * @return true if the strings are either the same, or indicate that they're part of the same card.
+     */
+    private static boolean areVendorsCompatible(final String glVendor, final String clVendor) {
+        return(   clVendor.equals(glVendor)
+               || (glVendor.contains("ATI Technologies") && clVendor.contains("Advanced Micro Devices"))
+               || (glVendor.contains("NVIDIA Corporation") && clVendor.contains("Apple")));
+    }
 
     /**
      * Accepts all platforms supporting the given extensions.
diff --git a/test/com/jogamp/opencl/LowLevelBindingTest.java b/test/com/jogamp/opencl/LowLevelBindingTest.java
index a4cbd03e..ef14534a 100644
--- a/test/com/jogamp/opencl/LowLevelBindingTest.java
+++ b/test/com/jogamp/opencl/LowLevelBindingTest.java
@@ -253,7 +253,7 @@ public class LowLevelBindingTest extends UITestCase {
         // Was originally 4096, but had to make this bigger or it would crash in UITestCase.oneTimeTearDown(){ System.gc() }
         // without even dumping a stack when using AMD drivers. Presumably the drivers would write past the end
         // of the block and mess up GC info somehow.
-        ByteBuffer bb = newDirectByteBuffer(8192);
+        ByteBuffer bb = newDirectByteBuffer(32768);
         ret = cl.clGetContextInfo(context, CL.CL_CONTEXT_DEVICES, bb.capacity(), bb, null);
         checkError("on clGetContextInfo", ret);
 
@@ -267,6 +267,7 @@ public class LowLevelBindingTest extends UITestCase {
         offset *= (is32Bit() ? 4 : 8);
         long device = is32Bit()?bb.getInt(offset):bb.getLong(offset);
 
+        bb.rewind();
         ret = cl.clGetDeviceInfo(device, CL.CL_DEVICE_MAX_WORK_GROUP_SIZE, bb.capacity(), bb, null);
         checkError("on clGetDeviceInfo", ret);
         int maxWGS = bb.getInt();
@@ -352,11 +353,13 @@ public class LowLevelBindingTest extends UITestCase {
         out.println("program source length (cl): "+longBuffer.get(0));
         out.println("program source length (java): "+programSource.length());
 
+        bb.rewind();
         ret = cl.clGetProgramInfo(program, CL.CL_PROGRAM_SOURCE, bb.capacity(), bb, null);
         checkError("on clGetProgramInfo CL_PROGRAM_SOURCE", ret);
         out.println("program source:\n" + clString2JavaString(bb, (int)longBuffer.get(0)));
 
         // Check program status
+        bb.rewind();
         ret = cl.clGetProgramBuildInfo(program, device, CL.CL_PROGRAM_BUILD_STATUS, bb.capacity(), bb, null);
         checkError("on clGetProgramBuildInfo1", ret);
 
diff --git a/test/com/jogamp/opencl/gl/CLGLTest.java b/test/com/jogamp/opencl/gl/CLGLTest.java
index 2ad31aa3..0d1a8c3f 100644
--- a/test/com/jogamp/opencl/gl/CLGLTest.java
+++ b/test/com/jogamp/opencl/gl/CLGLTest.java
@@ -33,7 +33,10 @@
 package com.jogamp.opencl.gl;
 
 import com.jogamp.common.nio.Buffers;
+import com.jogamp.opencl.CLBuffer;
 import com.jogamp.opencl.CLCommandQueue;
+import com.jogamp.opencl.CLKernel;
+import com.jogamp.opencl.CLProgram;
 
 import javax.media.opengl.GL2;
 import javax.media.opengl.GLException;
@@ -51,6 +54,7 @@ import com.jogamp.opencl.util.CLDeviceFilters;
 import com.jogamp.opencl.util.CLPlatformFilters;
 
 import java.io.IOException;
+import java.nio.ByteBuffer;
 import java.nio.IntBuffer;
 
 import javax.media.opengl.GLCapabilities;
@@ -222,6 +226,117 @@ public class CLGLTest extends UITestCase {
 
     }
 
+    @Test(timeout=15000)
+    public void textureSharing() {
+
+        out.println(" - - - glcl; textureSharing - - - ");
+        if(MiscUtils.isOpenCLUnavailable())
+            return;
+
+        initGL();
+        makeGLCurrent();
+        assertTrue(glcontext.isCurrent());
+
+        @SuppressWarnings("unchecked")
+        CLPlatform [] clplatforms = CLPlatform.listCLPlatforms(glSharing(glcontext));
+        if(clplatforms.length == 0) {
+            out.println("no platform that supports OpenGL-OpenCL interoperability");
+            return;
+        }
+
+        for(CLPlatform clplatform : clplatforms) {
+    
+            @SuppressWarnings("unchecked")
+            CLDevice [] cldevices = clplatform.listCLDevices(CLDeviceFilters.glSharing());
+    
+            for(CLDevice cldevice : cldevices) {
+                out.println(cldevice);
+                textureSharingInner(cldevice);
+            }
+        }
+        
+        deinitGL();
+    }
+
+    public void textureSharingInner(CLDevice cldevice) {
+
+        CLGLContext clglcontext = CLGLContext.create(glcontext, cldevice);
+
+        try {
+            out.println(clglcontext);
+
+            GL2 gl = glcontext.getGL().getGL2();
+
+            // create and write GL texture
+            int[] id = new int[1];
+            gl.glGenTextures(id.length, id, 0);
+            gl.glActiveTexture(GL2.GL_TEXTURE0);
+            gl.glBindTexture  (GL2.GL_TEXTURE_2D, id[0]);
+            int texWidth = 2;
+            int texHeight = 2;
+            gl.glTexImage2D(GL2.GL_TEXTURE_2D, 0, GL2.GL_RGBA, texWidth, texHeight, 0, GL2.GL_RGBA, GL2.GL_UNSIGNED_BYTE, null );
+            gl.glBindTexture(GL2.GL_TEXTURE_2D, 0);
+            gl.glFinish();
+
+            // create CLGL buffer
+            ByteBuffer bufferCL = Buffers.newDirectByteBuffer(texWidth*texHeight*4);
+            CLGLTexture2d<ByteBuffer> clTexture = clglcontext.createFromGLTexture2d(bufferCL, GL2.GL_TEXTURE_2D, id[0], 0, CLBuffer.Mem.WRITE_ONLY);
+
+            // set texel values to a formula that can be read back and verified
+            String sourceCL = "__kernel void writeTexture (__write_only image2d_t imageTex, unsigned w, unsigned h ) \n" +
+                    "{                                                                        \n" +
+                    "    for(unsigned y=1; y<=h; ++y) {                                       \n" +
+                    "        for(unsigned x=1; x<=w; ++x) {                                   \n" +
+                    "            write_imagef(imageTex, (int2)(x-1,y-1), (float4)(((float)x)/((float)(4*w)), ((float)y)/((float)(4*h)), 0.0f, 1.0f)); \n" +
+                    "        }                                                                \n" +
+                    "    }                                                                    \n" +
+                    "}";
+            CLProgram program = clglcontext.createProgram(sourceCL);
+            program.build();
+            System.out.println(program.getBuildStatus());
+            System.out.println(program.getBuildLog());
+            assertTrue(program.isExecutable());
+
+            CLKernel clkernel = program.createCLKernel("writeTexture")
+                    .putArg(clTexture)
+                    .putArg(texWidth)
+                    .putArg(texHeight)
+                    .rewind();
+
+            CLCommandQueue queue = cldevice.createCommandQueue();
+
+            // write gl texture with cl kernel, then read it to host buffer
+            queue.putAcquireGLObject(clTexture)
+                .put1DRangeKernel(clkernel, 0, 1, 1)
+                .putReadImage(clTexture, true)
+                .putReleaseGLObject(clTexture)
+                .finish();
+
+            for(int y = 1; y <= texHeight; y++) {
+                for(int x = 1; x <= texWidth; x++) {
+                    byte bX = bufferCL.get();
+                    byte bY = bufferCL.get();
+                    byte bZero = bufferCL.get();
+                    byte bMinusOne = bufferCL.get();
+                    byte bXCheck = (byte)(((float)x)/((float)(4*texWidth))*256);
+                    byte bYCheck = (byte)(((float)y)/((float)(4*texHeight))*256);
+                    assertEquals(bXCheck, bX);
+                    assertEquals(bYCheck, bY);
+                    assertEquals(0, bZero);
+                    assertEquals(-1, bMinusOne);
+                }
+            }
+
+            out.println(clTexture);
+
+            clTexture.release();
+            gl.glDeleteBuffers(1, id, 0);
+        }
+        finally {
+            clglcontext.release();
+        }
+    }
+
     private void makeGLCurrent() {
         // we are patient...
         while(true) {