implemented clBuildProgram(...) and updated JUnit test.

8 files changed, 261 insertions, 53 deletions

diff --git a/resources/cl-if.cfg b/resources/cl-if.cfg
index 08e21dcc..4d7d80f1 100644
--- a/resources/cl-if.cfg
+++ b/resources/cl-if.cfg
@@ -16,15 +16,19 @@ Ignore CL_GL_.*|cl.*GL.*
 
 #custom implementations
 Ignore clCreateContext
+CustomJavaCode CL
 CustomJavaCode CL  /** Interface to C language function: <br> <code> cl_context clCreateContext(intptr_t * , uint32_t, cl_device_id * , void (*pfn_notify)(const char *, const void *, size_t, void *), void *, int32_t * ); </code>    */
 CustomJavaCode CL  public long clCreateContext(IntBuffer properties, int properties_offset, long[] devices, CreateContextCallback pfn_notify, Object userData, IntBuffer errcode_ret, int errcode_offset);
 
 Ignore clCreateContextFromType
+CustomJavaCode CL
 CustomJavaCode CL  /** Interface to C language function: <br> <code> cl_context clCreateContextFromType(cl_context_properties *properties, cl_device_type device_type, void (*pfn_notify)(const char *errinfo,  const void *private_info, size_t cb, void *user_data), void *user_data, cl_int *errcode_ret) ; </code>    */
 CustomJavaCode CL  public long clCreateContextFromType(IntBuffer properties, int properties_offset, long device_type, CreateContextCallback pfn_notify, Object userData, IntBuffer errcode_ret, int errcode_offset);
 
 Ignore clBuildProgram
-#TODO..
+CustomJavaCode CL
+CustomJavaCode CL  /** Interface to C language function: <br> <code> int32_t clBuildProgram(cl_program, uint32_t, cl_device_id * , const char * , void (*pfn_notify)(cl_program, void *user_data), void * ); </code>    */
+CustomJavaCode CL  public int clBuildProgram(long program, long[] devices, String options, BuildProgramCallback cb, Object userData);
 
 Ignore clEnqueueNativeKernel
 #TODO..
diff --git a/resources/clImplCustomCode.c b/resources/clImplCustomCode.c
index 83d25938..9c87c24a 100644
--- a/resources/clImplCustomCode.c
+++ b/resources/clImplCustomCode.c
@@ -1,3 +1,12 @@
+/*
+void checkStatus(const char* msg, int status) {
+    if (status != CL_SUCCESS) {
+        printf("%s; error: %d \n", msg, status);
+        exit(EXIT_FAILURE);
+    }
+}
+*/
+
 
 /*
 void createContextCallback(const char * c, const void * v, size_t s, void * o) {
@@ -22,14 +31,7 @@ Java_com_mbien_opencl_impl_CLImpl_clCreateContextFromType0(JNIEnv *env, jobject
 
     intptr_t * _props_ptr  = NULL;
     int32_t * _errcode_ptr = NULL;
-
-/*
-    printf("jlong: %zu \n", sizeof(jlong)   );
-    printf("intptr_t: %zu \n", sizeof(intptr_t));
-    printf("size_t: %zu \n", sizeof(size_t));
-*/
-
-    cl_context _res;
+    cl_context _ctx;
 
     if (props != NULL) {
         _props_ptr = (intptr_t *) (((char*) (*env)->GetDirectBufferAddress(env, props)) + props_byte_offset);
@@ -39,9 +41,21 @@ Java_com_mbien_opencl_impl_CLImpl_clCreateContextFromType0(JNIEnv *env, jobject
     }
 
     //TODO callback; payload
-    _res = clCreateContextFromType((intptr_t *) _props_ptr, (uint64_t) device_type, NULL, NULL, (int32_t *) _errcode_ptr);
+    _ctx = clCreateContextFromType((intptr_t *) _props_ptr, (uint64_t) device_type, NULL, NULL, (int32_t *) _errcode_ptr);
 
-    return (jlong) (intptr_t) _res;
+/*
+    printf(" - - - - test - - - -  \n");
+    cl_uint err;
+    size_t deviceListSize;
+
+    // get the list of GPU devices associated with context
+    err = clGetContextInfo(_ctx, CL_CONTEXT_DEVICES, 0, NULL, &deviceListSize);  checkStatus("getContextInfo1", err);
+    cl_uint count = (cl_uint)deviceListSize / sizeof(cl_device_id);
+    printf("devices: %d \n",  count);
+    printf(" - - - - test end - - - -  \n");
+*/
+
+    return (jlong)_ctx;
 }
 
 
@@ -73,5 +87,40 @@ Java_com_mbien_opencl_impl_CLImpl_clCreateContext0(JNIEnv *env, jobject _unused,
 }
 */
 
+/**
+ * Entry point to C language function: <code> int32_t clBuildProgram(cl_program, uint32_t, cl_device_id * , const char * , void (*pfn_notify)(cl_program, void *user_data), void * );
+ */
+JNIEXPORT jint JNICALL
+Java_com_mbien_opencl_impl_CLImpl_clBuildProgram0(JNIEnv *env, jobject _unused,
+        jlong program, jint deviceCount, jobject deviceList, jint offset, jstring options, jobject cb, jobject data) {
+
+    const char* _strchars_options = NULL;
+    cl_int _res;
+    size_t * _deviceListPtr = NULL;
+
+    if (options != NULL) {
+        _strchars_options = (*env)->GetStringUTFChars(env, options, (jboolean*)NULL);
+        if (_strchars_options == NULL) {
+            (*env)->ThrowNew(env, (*env)->FindClass(env, "java/lang/OutOfMemoryError"),
+                           "Failed to get UTF-8 chars for argument \"options\" in native dispatcher for \"clBuildProgram\"");
+            return CL_FALSE;
+        }
+    }
+
+    if (deviceList != NULL) {
+        _deviceListPtr = (void *) (((char*) (*env)->GetPrimitiveArrayCritical(env, deviceList, NULL)) + offset);
+    }
+    
+    // TODO payload, callback...
+    _res = clBuildProgram((cl_program)program, (cl_uint)deviceCount, _deviceListPtr, _strchars_options, NULL, NULL);
 
+    if (deviceList != NULL) {
+        (*env)->ReleasePrimitiveArrayCritical(env, deviceList, _deviceListPtr, 0);
+    }
 
+    if (options != NULL) {
+        (*env)->ReleaseStringUTFChars(env, options, _strchars_options);
+    }
+    
+    return _res;
+}
diff --git a/resources/clImplCustomCode.java b/resources/clImplCustomCode.java
index 5ccc3b02..618b03c2 100644
--- a/resources/clImplCustomCode.java
+++ b/resources/clImplCustomCode.java
@@ -1,14 +1,29 @@
 
-      public long clCreateContext(IntBuffer properties, int offset1, long[] devices, CreateContextCallback cb, Object userData, IntBuffer errcode_ret, int offset2) {
-          return this.clCreateContext0(properties, offset1, devices, cb, null, errcode_ret, offset2);
-      }
-      private native long clCreateContext0(IntBuffer properties, int size, long[] devices, CreateContextCallback pfn_notify, Object userData, IntBuffer errcode_ret, int size2);
+        public long clCreateContext(IntBuffer properties, int offset1, long[] devices, CreateContextCallback cb, Object userData, IntBuffer errcode_ret, int offset2) {
+            throw new RuntimeException("not yet implemented, use clCreateContextFromType instead");
+//            return this.clCreateContext0(properties, offset1, devices, cb, null, errcode_ret, offset2);
+        }
+        private native long clCreateContext0(IntBuffer cl_context_properties, int size, long[] devices, CreateContextCallback pfn_notify, Object userData, IntBuffer errcode_ret, int size2);
 
 
+        public long clCreateContextFromType(IntBuffer properties, int offset1, long device_type, CreateContextCallback pfn_notify, Object userData, IntBuffer errcode_ret, int offset2) {
+            if(pfn_notify != null)
+                throw new RuntimeException("asynchronous execution with callback is not yet implemented, pass null through this method to block until complete.");
+            if(userData != null)
+                System.out.println("WARNING: userData not yet implemented... ignoring");
+            return this.clCreateContextFromType0(properties, offset1, device_type, pfn_notify, null, errcode_ret, offset2);
+        }
+        private native long clCreateContextFromType0(IntBuffer properties, int size, long device_type, CreateContextCallback pfn_notify, Object userData, IntBuffer errcode_ret, int size2);
 
-      public long clCreateContextFromType(IntBuffer arg0, int offset1, long device_type, CreateContextCallback pfn_notify, Object userData, IntBuffer errcode_ret, int offset2) {
-          return this.clCreateContextFromType0(arg0, offset1, device_type, pfn_notify, null, errcode_ret, offset2);
-      }
-      private native long clCreateContextFromType0(IntBuffer arg0, int size, long device_type, CreateContextCallback pfn_notify, Object userData, IntBuffer errcode_ret, int size2);
 
-      
\ No newline at end of file
+        /** Interface to C language function: <br> <code> int32_t clBuildProgram(cl_program, uint32_t, cl_device_id * , const char * , void * ); </code>    */
+        public int clBuildProgram(long program, long[] deviceList, String options, BuildProgramCallback cb, Object userData)  {
+            if(cb != null)
+                throw new RuntimeException("asynchronous execution with callback is not yet implemented, pass null through this method to block until complete.");
+            if(userData != null)
+                System.out.println("WARNING: userData not yet implemented... ignoring");
+            return clBuildProgram0(program, deviceList.length, deviceList, 0, options, cb, userData);
+        }
+        /** Entry point to C language function: <code> int32_t clBuildProgram(cl_program, uint32_t, cl_device_id * , const char * , void * ); </code>    */
+        private native int clBuildProgram0(long program, int devices, Object deviceList, int arg2_byte_offset, String options, BuildProgramCallback cb, Object userData);
+
diff --git a/src/com/mbien/opencl/BuildProgramCallback.java b/src/com/mbien/opencl/BuildProgramCallback.java
new file mode 100644
index 00000000..0b2b1b04
--- /dev/null
+++ b/src/com/mbien/opencl/BuildProgramCallback.java
@@ -0,0 +1,10 @@
+package com.mbien.opencl;
+
+/**
+ * @author Michael Bien
+ */
+// TODO implement callbacks
+public interface BuildProgramCallback {
+    
+    public void buildProgramCallback(long cl_program, Object user_data);
+}
diff --git a/src/com/mbien/opencl/CLContext.java b/src/com/mbien/opencl/CLContext.java
index bcfb09a0..203172b0 100644
--- a/src/com/mbien/opencl/CLContext.java
+++ b/src/com/mbien/opencl/CLContext.java
@@ -1,6 +1,7 @@
 package com.mbien.opencl;
 
 import com.mbien.opencl.impl.CLImpl;
+import com.sun.gluegen.runtime.PointerBuffer;
 import java.nio.IntBuffer;
 
 /**
@@ -62,10 +63,9 @@ public final class CLContext {
     /**
      * Gets the device with maximal FLOPS from this context.
      */
+    /*
     public CLDevice getMaxFlopsDevice() {
 
-        //TODO not finished
-
         long[] longBuffer = new long[1];
 //        ByteBuffer bb = ByteBuffer.allocate(8);
 //        bb.order(ByteOrder.nativeOrder());
@@ -96,6 +96,12 @@ public final class CLContext {
         return null;
     }
 
+    public CLDevice[] getCLDevices() {
+
+    }
+*/
+
+
     /**
      * Lists all available OpenCL implementaitons.
      * @throws CLException if something went wrong initializing OpenCL
diff --git a/src/com/mbien/opencl/CLException.java b/src/com/mbien/opencl/CLException.java
index e8ed771f..4539046d 100644
--- a/src/com/mbien/opencl/CLException.java
+++ b/src/com/mbien/opencl/CLException.java
@@ -125,7 +125,7 @@ public class CLException extends RuntimeException {
                 return "CL_INVALID_MIP_LEVEL";
 
             default:
-                return "unknown cause";
+                return "unknown cause: error "+error;
         }
     }
 
diff --git a/src/com/mbien/opencl/CreateContextCallback.java b/src/com/mbien/opencl/CreateContextCallback.java
index b25c05fd..e965370c 100644
--- a/src/com/mbien/opencl/CreateContextCallback.java
+++ b/src/com/mbien/opencl/CreateContextCallback.java
@@ -6,6 +6,7 @@ import java.nio.ByteBuffer;
  *
  * @author Michael Bien
  */
+// TODO implement callbacks
 public interface CreateContextCallback {
 
     public void createContextCallback(String errinfo, ByteBuffer private_info, long cb, Object user_data);
diff --git a/test/com/mbien/opencl/JOCLTest.java b/test/com/mbien/opencl/JOCLTest.java
index 0eccf430..3ed51cbf 100644
--- a/test/com/mbien/opencl/JOCLTest.java
+++ b/test/com/mbien/opencl/JOCLTest.java
@@ -24,32 +24,23 @@ public class JOCLTest {
     }
 
     @Test
-    public void lowLevelTest() {
+    public void lowLevelTest1() {
 
-        out.println(" - - - lowLevelTest - - - ");
+        out.println(" - - - lowLevelTest; contextless binding - - - ");
 
-        CreateContextCallback cb = new CreateContextCallback() {
-            @Override
-            public void createContextCallback(String errinfo, ByteBuffer private_info, long cb, Object user_data) {
-                throw new RuntimeException(errinfo);
-            }
-        };
-
-        out.println("creating OpenCL context");
-
-        int ret = 0;
+        int ret = CL.CL_SUCCESS;
 
         CL cl = CLContext.getLowLevelBinding();
 
         int[] intBuffer = new int[1];
         // find all available OpenCL platforms
         ret = cl.clGetPlatformIDs(0, null, 0, intBuffer, 0);
-        assertEquals(CL.CL_SUCCESS, ret);
+        checkForError(ret);
         out.println("#platforms: "+intBuffer[0]);
 
         long[] platformId = new long[intBuffer[0]];
         ret = cl.clGetPlatformIDs(platformId.length, platformId, 0, null, 0);
-        assertEquals(CL.CL_SUCCESS, ret);
+        checkForError(ret);
 
         // print platform info
         long[] longBuffer = new long[1];
@@ -62,24 +53,24 @@ public class JOCLTest {
             out.println("platform id: "+platform);
 
             ret = cl.clGetPlatformInfo(platform, CL.CL_PLATFORM_PROFILE, bb.capacity(), bb, longBuffer, 0);
-            assertEquals(CL.CL_SUCCESS, ret);
+            checkForError(ret);
             out.println("    profile: "+new String(bb.array(), 0, (int)longBuffer[0]));
 
             ret = cl.clGetPlatformInfo(platform, CL.CL_PLATFORM_VERSION, bb.capacity(), bb, longBuffer, 0);
-            assertEquals(CL.CL_SUCCESS, ret);
+            checkForError(ret);
             out.println("    version: "+new String(bb.array(), 0, (int)longBuffer[0]));
 
             ret = cl.clGetPlatformInfo(platform, CL.CL_PLATFORM_NAME, bb.capacity(), bb, longBuffer, 0);
-            assertEquals(CL.CL_SUCCESS, ret);
+            checkForError(ret);
             out.println("    name: "+new String(bb.array(), 0, (int)longBuffer[0]));
 
             ret = cl.clGetPlatformInfo(platform, CL.CL_PLATFORM_VENDOR, bb.capacity(), bb, longBuffer, 0);
-            assertEquals(CL.CL_SUCCESS, ret);
+            checkForError(ret);
             out.println("    vendor: "+new String(bb.array(), 0, (int)longBuffer[0]));
 
             //find all devices
             ret = cl.clGetDeviceIDs(platform, CL.CL_DEVICE_TYPE_ALL, 0, null, 0, intBuffer, 0);
-            assertEquals(CL.CL_SUCCESS, ret);
+            checkForError(ret);
             out.println("#devices: "+intBuffer[0]);
 
             long[] devices = new long[intBuffer[0]];
@@ -89,11 +80,11 @@ public class JOCLTest {
             for (int j = 0; j < devices.length; j++) {
                 long device = devices[j];
                 ret = cl.clGetDeviceInfo(device, CL.CL_DEVICE_NAME, bb.capacity(), bb, longBuffer, 0);
-                assertEquals(CL.CL_SUCCESS, ret);
+                checkForError(ret);
                 out.println("    device: "+new String(bb.array(), 0, (int)longBuffer[0]));
 
                 ret = cl.clGetDeviceInfo(device, CL.CL_DEVICE_TYPE, bb.capacity(), bb, longBuffer, 0);
-                assertEquals(CL.CL_SUCCESS, ret);
+                checkForError(ret);
                 out.println("    type: " + CLDevice.Type.valueOf(bb.get()));
                 bb.rewind();
 
@@ -101,22 +92,132 @@ public class JOCLTest {
 
         }
 
-        Arrays.fill(longBuffer, 0);
+    }
+
+    @Test
+    public void lowLevelTest2() {
+
+        out.println(" - - - lowLevelTest2; VectorAdd kernel - - - ");
+
+//        CreateContextCallback cb = new CreateContextCallback() {
+//            @Override
+//            public void createContextCallback(String errinfo, ByteBuffer private_info, long cb, Object user_data) {
+//                throw new RuntimeException("not yet implemented...");
+//            }
+//        };
+
+        long[] longBuffer = new long[1];
+        ByteBuffer bb = ByteBuffer.allocate(1024);
+        bb.order(ByteOrder.nativeOrder());
+
+        CL cl = CLContext.getLowLevelBinding();
+
+        int ret = CL.CL_SUCCESS;
+        int[] intArray = new int[1];
 
-        long context = cl.clCreateContextFromType(null, 0, CL.CL_DEVICE_TYPE_ALL, cb, null, null, 0);
+        long context = cl.clCreateContextFromType(null, 0, CL.CL_DEVICE_TYPE_ALL, null, null, null, 0);
         out.println("context handle: "+context);
 
-        ret = cl.clGetContextInfo(context, CL.CL_CONTEXT_DEVICES, 0, null, longBuffer, 0);
-        assertEquals(CL.CL_SUCCESS, ret);
+        // TODO fix gluegen bug: array-buffer mixing... bb is a noop
+        ret = cl.clGetContextInfo(context, CL.CL_CONTEXT_DEVICES, 0, bb, longBuffer, 0);
+        checkError("on clGetContextInfo", ret);
 
-        out.println("CL_CONTEXT_DEVICES result: "+longBuffer[0]);
+        int sizeofLong = 8; // TODO sizeof long...
+        out.println("context created with " + longBuffer[0]/sizeofLong + " devices");
 
-        ret = cl.clGetContextInfo(context, CL.CL_CONTEXT_NUM_DEVICES, 0, null, longBuffer, 0);
-        assertEquals(CL.CL_SUCCESS, ret);
+        ret = cl.clGetContextInfo(context, CL.CL_CONTEXT_DEVICES, bb.capacity(), bb, null, 0);
+        checkError("on clGetContextInfo", ret);
 
-        out.println("CL_CONTEXT_NUM_DEVICES result: "+longBuffer[0]);
+        for (int i = 0; i < longBuffer[0]/sizeofLong; i++) {
+            out.println("device id: "+bb.getLong());
+        }
+
+        long firstDeviceID = bb.getLong(0);
+
+        // Create a command-queue
+        long commandQueue = cl.clCreateCommandQueue(context, firstDeviceID, 0, intArray, 0);
+        checkError("on clCreateCommandQueue", intArray[0]);
+
+        int iNumElements = 11444777;	// Length of float arrays to process (odd # for illustration)
+        int szLocalWorkSize = 256;      // set and log Global and Local work size dimensions
+        int szGlobalWorkSize = roundUp(szLocalWorkSize, iNumElements);  // rounded up to the nearest multiple of the LocalWorkSize
+        int sizeofFloat = 4; // TODO sizeof float ...
+        
+        // Allocate the OpenCL buffer memory objects for source and result on the device GMEM
+        long cmDevSrcA = cl.clCreateBuffer(context, CL.CL_MEM_READ_ONLY, sizeofFloat * szGlobalWorkSize, null, intArray, 0);
+        checkError("on clCreateBuffer", intArray[0]);
+        long cmDevSrcB = cl.clCreateBuffer(context, CL.CL_MEM_READ_ONLY, sizeofFloat * szGlobalWorkSize, null, intArray, 0);
+        checkError("on clCreateBuffer", intArray[0]);
+        long cmDevDst  = cl.clCreateBuffer(context, CL.CL_MEM_WRITE_ONLY, sizeofFloat * szGlobalWorkSize, null, intArray, 0);
+        checkError("on clCreateBuffer", intArray[0]);
+        
+        String src = 
+              " // OpenCL Kernel Function for element by element vector addition                                                \n"
+            + "__kernel void VectorAdd(__global const float* a, __global const float* b, __global float* c, int iNumElements) { \n"
+            + "    // get index into global data array                                                                          \n"
+            + "    int iGID = get_global_id(0);                                                                                 \n"
+            + "    // bound check (equivalent to the limit on a 'for' loop for standard/serial C code                           \n"
+            + "    if (iGID >= iNumElements)  {                                                                                 \n"
+            + "        return;                                                                                                  \n"
+            + "    }                                                                                                            \n"
+            + "    // add the vector elements                                                                                   \n"
+            + "    c[iGID] = a[iGID] + b[iGID];                                                                                 \n"
+            + "}                                                                                                                \n";
+
+
+        // Create the program
+        long program = cl.clCreateProgramWithSource(context, 1, new String[] {src}, new long[]{src.length()}, 0, intArray, 0);
+        checkError("on clCreateProgramWithSource", intArray[0]);
+        
+        // Build the program
+        ret = cl.clBuildProgram(program, new long[] { firstDeviceID }, null, null, null);
+        checkError("on clBuildProgram", ret);
+
+        // Check program status
+        Arrays.fill(longBuffer, 42);
+        bb.rewind();
+        ret = cl.clGetProgramBuildInfo(program, firstDeviceID, CL.CL_PROGRAM_BUILD_STATUS, bb.capacity(), bb, null, 0);
+        checkError("on clGetProgramBuildInfo1", ret);
+
+        out.println("program build status: " + getBuildStatus(bb.getInt(0)));
+        assertEquals("build status", CL.CL_BUILD_SUCCESS, bb.getInt(0));
+
+        // Read build log
+        // TODO fix gluegen bug: array-buffer mixing... bb is a noop
+        ret = cl.clGetProgramBuildInfo(program, firstDeviceID, CL.CL_PROGRAM_BUILD_LOG, 0, bb, longBuffer, 0);
+        checkError("on clGetProgramBuildInfo2", ret);
+        out.println("program log length: " + longBuffer[0]);
+
+        bb.rewind();
+        ret = cl.clGetProgramBuildInfo(program, firstDeviceID, CL.CL_PROGRAM_BUILD_LOG, bb.capacity(), bb, null, 0);
+        checkError("on clGetProgramBuildInfo3", ret);
+        out.println("log:\n" + new String(bb.array(), 0, (int)longBuffer[0]));
+
+        // Create the kernel
+        long kernel = cl.clCreateKernel(program, "VectorAdd", intArray, 0);
+        checkError("on clCreateKernel", intArray[0]);
+        
+        ret = cl.clReleaseContext(context);
+        checkError("on clReleaseContext", ret);
 
-        cl.clReleaseContext(context);
+    }
+
+    private String getBuildStatus(int status) {
+        switch(status) {
+            case CL.CL_BUILD_SUCCESS:
+                return "CL_BUILD_SUCCESS";
+            case CL.CL_BUILD_NONE:
+                return "CL_BUILD_NONE";
+            case CL.CL_BUILD_IN_PROGRESS:
+                return "CL_BUILD_IN_PROGRESS";
+            case CL.CL_BUILD_ERROR:
+                return "CL_BUILD_ERROR";
+// can't find this flag in spec...
+//            case CL.CL_BUILD_PROGRAM_FAILURE:
+//                return "CL_BUILD_PROGRAM_FAILURE";
+            default:
+                return "unknown status: " + status;
+        }
     }
 
     @Test
@@ -158,4 +259,26 @@ public class JOCLTest {
     }
 
 
+    
+
+
+    private final int roundUp(int groupSize, int globalSize) {
+        int r = globalSize % groupSize;
+        if (r == 0) {
+            return globalSize;
+        } else {
+            return globalSize + groupSize - r;
+        }
+    }
+
+    private final void checkForError(int ret) {
+        this.checkError("", ret);
+    }
+
+    private final void checkError(String msg, int ret) {
+        if(ret != CL.CL_SUCCESS)
+            throw new CLException(ret, msg);
+    }
+
+
 }
\ No newline at end of file