diff options
-rw-r--r-- | resources/cl-if.cfg | 6 | ||||
-rw-r--r-- | resources/clImplCustomCode.c | 69 | ||||
-rw-r--r-- | resources/clImplCustomCode.java | 33 | ||||
-rw-r--r-- | src/com/mbien/opencl/BuildProgramCallback.java | 10 | ||||
-rw-r--r-- | src/com/mbien/opencl/CLContext.java | 10 | ||||
-rw-r--r-- | src/com/mbien/opencl/CLException.java | 2 | ||||
-rw-r--r-- | src/com/mbien/opencl/CreateContextCallback.java | 1 | ||||
-rw-r--r-- | test/com/mbien/opencl/JOCLTest.java | 183 |
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 |